7. Occupancy, registers, and spills

Occupancy is resident active warps divided by the architectural maximum. It is a latency-hiding resource, not a performance score.

GPUs cannot make a DRAM read finish faster; a round trip to device memory still takes hundreds of cycles no matter how the kernel is written. Instead, an SM hides that latency by keeping many warps resident at once, so that when one warp's next instruction has an unmet dependency — an operand not yet computed, a value not yet returned from memory — the scheduler can switch to issuing an instruction from a different, already-ready warp on the very next cycle instead of leaving the pipeline idle. Occupancy measures how much of that switching inventory is present. This is why the book calls it "latency-hiding" rather than "more parallelism is always better": once enough resident warps exist to keep the scheduler supplied with an eligible one on every cycle, additional occupancy has nothing left to hide and stops helping, while forcing it higher can even cost performance if it comes at the expense of registers or shared memory a kernel needed for something else (the warp-specialized GEMM later in this chapter is a case where low occupancy is by design). Occupancy has a ceiling because only a finite number of thread slots, warp slots, registers, and shared-memory bytes exist per SM; whichever of those four resources is exhausted first caps how many warps can be simultaneously resident. The formula below computes exactly that bound, one term per resource.

Resource limits combine

For threads/block $T_b$, warps/block $W_b=\lceil T_b/32\rceil$, registers per thread $R_t$, and shared memory per block $S_b$, an approximate block limit is

$$B_{resident}=\min\left( B_{arch}, \left\lfloor\frac{T_{SM}}{T_b}\right\rfloor, \left\lfloor\frac{W_{SM}}{W_b}\right\rfloor, \left\lfloor\frac{R_{SM}}{R_tT_b}\right\rfloor, \left\lfloor\frac{S_{SM}}{S_b}\right\rfloor \right).$$

Real register and shared-memory allocation is rounded in architecture-specific granularities, so use Nsight Compute's occupancy section/calculator for exact limits.

Figure 7.1 — The first exhausted resource determines theoretical residency; achieved issue still depends on eligible warps and pipeline demand.

Read theoretical and achieved occupancy together

Theoretical occupancy is a ceiling computed from resources; it says nothing about whether that ceiling is reached or whether reaching it would even help. Compare theoretical against achieved before drawing a conclusion:

On a warp-specialized GEMM, one 384-thread CTA may intentionally consume about 80–90 KiB shared memory and 168 registers/thread. One CTA/SM can still keep the Tensor pipe near saturation. Reducing registers simply to raise occupancy can slow it down by breaking the software pipeline.

Registers versus local memory

“Local” memory is thread-private in the programming model but physically lives in device memory and is cached. Compiler-generated stack slots, arrays with dynamic indexing, and spilled live values use it.

Compile CUDA code with resource reporting and line information:

nvcc -O3 -lineinfo -Xptxas=-v kernel.cu -o kernel

Then inspect in Nsight Compute:

• Registers Per Thread;
• Local Memory Total;
• local load/store instructions and sectors;
• L1/L2 sectors attributable to local traffic;
• source lines with local-memory instructions.

A register-spill diagnosis

Put registers, local memory, and occupancy together on one concrete kernel. Consider a 128-thread backward kernel using 255 registers/thread. Its stall profile is dominated by long-scoreboard stalls (waiting on an L1TEX-tracked memory result; the full stall taxonomy follows in ch. 8):

Registers limit residency and the kernel still spills. This is not an argument to apply -maxrregcount blindly—that usually creates more spills. It is evidence to reduce live ranges, recompute cheap intermediates, shrink per-thread tiles, reorganize state, or split the kernel at a boundary with a favorable intermediate size.

Safe experiments

Turn that evidence into experiments, one variable at a time:

1. Inspect the exact SASS/source lines producing local traffic.
2. Change one tile/live-state parameter.
3. Verify numerical behavior.
4. Re-check registers, spills, occupancy, and duration together.
5. Measure the complete workload.

Step 2 covers more than shrinking a tile size. One concrete version of it is the kernel split already named above: splitting the kernel at a boundary with a favorable intermediate size trades one extra read/write pass — the intermediate now has to be written out and read back in — for lower live range per thread and, potentially, higher occupancy in each half. That tradeoff is not free, and it is not automatically favorable. Ch. 13's Case G tried exactly this hypothesis on this exact kernel: materializing an intermediate between phases of the backward pass was numerically valid and did lower register pressure, but it ran slower — the added memory traffic from writing and re-reading the intermediate outweighed the spill savings it was meant to buy. Treat a kernel split the same way as any other step-2 change: a hypothesis to run through the full ledger discipline from ch. 12, not a technique that pays off just because it addresses the diagnosis on paper.

__launch_bounds__ and maximum-register flags are experiments, not fixes. A higher occupancy number is a win only if relevant elapsed time falls.