Explain occupancy and register pressure: launch bounds, spills, the calculator, and why 50% occupancy can beat 100%.

TL;DR: An H100 SM has 65,536 registers, 228 KB of shared memory, and room for 64 warps (2,048 threads); occupancy is the fraction of those 64 warps a kernel can keep resident, and it is limited by whichever resource runs out first. A kernel using 128 registers per thread allows 65,536 / 128 = 512 threads = 16 warps = 25%; at 32 registers it allows 64 warps = 100%. Occupancy matters because the SM hides memory latency by switching among resident warps: with 16 warps and a 600-cycle HBM latency, each warp must issue enough independent work to cover the gap or the SM idles. __launch_bounds__(threads, minBlocks) tells the compiler to cap registers so a target occupancy is reachable; if the kernel needs more, the compiler spills to local memory (L1-cached, but still loads and stores), and past a point the spills cost more than the occupancy gains. A kernel with high instruction-level parallelism per thread (a register-blocked GEMM with 64 independent accumulators, or a loop unrolled to keep 8 loads in flight) hides latency within a warp and runs faster at 25 to 50% occupancy with more registers than at 100% with fewer.

How to approach it

Give the SM resource numbers and the occupancy arithmetic with two examples. Explain why occupancy hides latency. Then launch bounds, spills and how to read them in the compiler output. Then the counter-example (ILP) with numbers, and the profiler view. Close with the tuning procedure.

A strong answer

A typical situation: a candidate's reduction kernel reaches 100% occupancy and 60% of bandwidth; a colleague's version at 50% occupancy with each thread handling 8 elements in flight hits 90%; the candidate is asked to explain why theirs lost.

The arithmetic:

H100 SM: 65,536 32-bit registers; 228 KB shared memory (up to 227 KB per block); 64 warps max; 32 blocks max;
  2,048 threads max; register allocation granularity is per warp in units of 256 registers (8 per thread)
occupancy = resident warps / 64, limited by:
  registers: warps ≤ 65,536 / (regs_per_thread × 32)   → 32 regs: 64 warps (100%); 64: 32 (50%); 128: 16 (25%); 255: 8 (12.5%)
  shared memory: blocks ≤ 228 KB / smem_per_block        → 48 KB per block: 4 blocks; at 256 threads per block that is
    32 warps (50%)
  block size: blocks ≤ 32 and warps ≤ 64                 → 64-thread blocks: 32 blocks × 2 warps = 64 warps (fine);
    32-thread blocks: 32 blocks × 1 warp = 32 warps (50%), a block-count limit
example 1: 256 threads per block, 40 registers, 16 KB shared → registers allow 51 warps (rounded to whole blocks: 6
  blocks × 8 warps = 48), shared allows 14 blocks, block limit 32 → 48 warps = 75%
example 2: 128 threads per block, 168 registers, 96 KB shared → registers allow 12 warps (3 blocks), shared allows 2
  blocks = 8 warps → 8 warps = 12.5%; the shared memory is the binding limit
sanity: the calculator (Nsight Compute's occupancy section, or the spreadsheet) does this arithmetic; the interviewer
        wants the candidate to do it by hand once.

Occupancy and Register Pressure has the resource tables per architecture; GPU Execution Model explains the warp scheduler that makes occupancy matter.

Why occupancy hides latency, with the number:

an HBM load has ~600 to 800 cycles of latency on H100; the SM's 4 schedulers each issue one instruction per cycle
  from a ready warp
if each warp issues one load and then waits, the SM needs ~600 / (instructions per warp between loads) warps to keep
  issuing; with one instruction between loads, hundreds of warps (impossible); with 10, ~60 warps (100% occupancy);
  with 40 independent instructions per load, ~15 warps (25%)
so: the product (resident warps × independent instructions per warp) must cover the latency; occupancy is one factor,
  ILP within the warp is the other, and they trade

Launch bounds and spills:

__global__ void __launch_bounds__(256, 4) kern(...)   // 256 threads per block; want ≥ 4 blocks per SM resident
// → the compiler caps registers at 65,536 / (4 × 256) = 64 per thread; if the kernel needs 80, it spills 16

what a spill is: the compiler stores a live value to local memory (thread-private, in global address space, cached
  in L1) and reloads it later; each spill is a store and a load through L1, which is faster than HBM but far slower
  than a register, and spills often land inside the inner loop
how to see it: nvcc -Xptxas -v prints "Used 64 registers, 128 bytes spill stores, 128 bytes spill loads"; Nsight
  Compute shows local memory traffic and "registers per thread" with the occupancy it allows
the trade: a few bytes of spill outside the hot loop cost nothing; spills inside a loop that runs a million times
  cost more than the occupancy they bought; the rule is to look at where the spills are (the SASS view) before
  deciding
the alternative to a cap: reduce live values (smaller per-thread tiles, recompute cheap values instead of holding
  them, split a long kernel), or accept lower occupancy and check that ILP covers the latency
maxrregcount: a compile-wide cap; launch bounds are per kernel and preferred

The counter-example, with the scenario's numbers:

reduction over 1 GB of floats, H100
version A: 256-thread blocks, each thread loads 1 element, adds to a running sum, strided grid loop with 1 element in
  flight per thread; 24 registers → 100% occupancy; per warp, 1 load then dependent adds → the SM needs all 64 warps
  to hide latency and still idles → 2.0 TB/s (60%)
version B: each thread loads 8 elements (unrolled, 8 independent loads issued back to back, or 2 × float4), 48
  registers → 50% occupancy; per warp, 8 loads in flight before any use → 32 warps × 8 loads = 256 loads in flight
  per SM vs A's 64 → latency covered → 3.0 TB/s (90%)
the general form: bytes in flight per SM needed ≈ latency × bandwidth per SM = 700 cycles × (3.35 TB/s / 132 SMs /
  1.7 GHz) ≈ 700 × 15 B/cycle ≈ 10 KB per SM; A has 64 warps × 32 lanes × 4 B = 8 KB in flight; B has 32 × 32 × 32
  B = 32 KB; B is comfortably over the line, A is under it
sanity: occupancy is a way to get bytes in flight; ILP is another; the target is bytes in flight, and either path
        that reaches it is fine.

Reduce a 100M-Element Array Fast (the companion question) is the full kernel; Memory-Bound vs Compute-Bound Kernels is why bytes in flight is the measure for this kernel.

The profiler view: Nsight Compute's Occupancy section shows theoretical occupancy (from the resource arithmetic) and achieved occupancy (measured; lower when blocks finish unevenly or the grid has a tail); the Launch Statistics section shows registers per thread and shared memory per block; the Warp State Statistics section shows what warps were waiting on (long scoreboard stalls mean memory latency not hidden, which is the signal that either occupancy or ILP is short); and the Scheduler Statistics section shows issue slots used, which is the number that says whether more warps would help (no eligible warp most cycles means yes; eligible warps unused means no).

Profiling with Nsight walks those sections in order.

The tuning procedure: read achieved occupancy and the stall reasons; if stalls are memory latency and the scheduler has no eligible warps, raise occupancy (cap registers with launch bounds, cut shared memory, resize blocks) or raise ILP (unroll, more elements per thread); if the scheduler has eligible warps it is not issuing, occupancy is not the problem and the kernel is bound elsewhere (bandwidth, a bank conflict, the tensor pipe); measure each change, because the register cap that helps one kernel spills another into slowness.

The reversal condition: a compute-bound tensor-core kernel wants few, fat warps with many registers (wgmma accumulators are hundreds of registers per thread) at 25% or lower occupancy; pushing it to 50% by capping registers destroys it. Occupancy is a means; the end is issue slots used.

What interviewers probe next

• "What is achieved versus theoretical occupancy?" Theoretical is the resource limit; achieved is the average resident warps measured over the kernel, lower with tail effects, uneven block durations, or a grid too small to fill the SMs.
• "Does higher occupancy ever hurt directly?" It can: more resident blocks share L1 and shared memory, and cache thrashing rises; and forcing it through register caps causes spills.
• "How many blocks should a grid have?" At least SMs × blocks-per-SM × a few waves so the tail is small; 132 SMs × 8 blocks = 1,056 resident, so a few thousand blocks minimum for a short kernel.
• "What does __launch_bounds__ do if the kernel already fits?" Nothing; it is an upper bound on registers, not a request for them.

Common mistakes

• Treating 100% occupancy as the goal.
• Capping registers and ignoring the spill report.
• Reading theoretical occupancy and not the stall reasons.
• Forgetting that shared memory per block limits occupancy as much as registers do.

Key takeaways

• Occupancy = resident warps ÷ 64, limited by registers (65,536 per SM), shared memory (228 KB) and block counts; do the arithmetic.
• It exists to hide latency; ILP within a warp does the same job; the target is bytes (or instructions) in flight.
• Launch bounds cap registers; spills go to local memory; check where they land before accepting them.
• Read achieved occupancy, stall reasons and issue-slot utilization together; tune to issue slots, not to occupancy.