Occupancy is one number with an unusually bad reputation: how many warps are resident on a streaming multiprocessor, divided by the most that SM can hold. It is not a performance metric and not a score to maximize — it is a resource accounting result. A kernel gets the warps its scarcest per-SM resource permits, and touching the other resources changes nothing. Knowing which of the three ceilings — registers per thread, shared memory per block, or block and warp slots — binds tells you exactly what to change, and often tells you the honest answer is nothing.
What occupancy actually measures
Occupancy is defined per SM: resident warps divided by the maximum resident warps the SM supports. A warp is 32 threads that issue together; ‘resident’ means the warp has been assigned a slot, has its registers allocated, and is eligible for the scheduler to issue from. It does not mean the warp is executing right now.
That distinction is where most confusion starts. Occupancy is not GPU utilization, which merely says a kernel was running, nor issue-slot utilization, which says how often the SM found an instruction to issue. Occupancy is a capacity figure: how many warps the scheduler could choose from. A kernel can sit at 100% occupancy and stall constantly if every resident warp is blocked on the same uncoalesced load; it can sit at 25% and saturate the machine if each thread carries enough independent work. Occupancy bounds the scheduler’s options; it does not describe what the scheduler did with them.
Why resident warps hide latency
The reason capacity matters is latency. A global-memory load costs on the order of several hundred cycles — time in which the SM’s arithmetic pipelines could have retired an enormous amount of work. GPUs do not attack this with the out-of-order machinery a CPU uses. They attack it with oversubscription: keep many warps resident, and when one stalls on a dependency, issue from another.
What makes this cheap is that a warp switch costs nothing. Every resident warp’s registers already live in the SM’s register file and its block’s data already sits in shared memory, so there is no context to save or restore — the scheduler just picks a different eligible warp next cycle. That is the whole trade: you pay for latency hiding in resources permanently reserved rather than in switching time. The consequence is Little’s-law shaped — you need enough memory requests in flight to cover the latency-bandwidth product, and resident warps buy them.
Ceiling one — registers per thread
Each SM has a fixed register file, and every thread’s live values are allocated out of it for the thread’s entire lifetime. Multiply registers per thread by 32 to get registers per warp, divide the register file by that, and you have the warps the register file can hold. A kernel allocated at 32 registers per thread fits twice as many warps as one at 64 — a straight reciprocal.
Two details make that arithmetic slightly optimistic. Registers are handed out in granules per warp, so the real allowance rounds the clean division down to a granule boundary. And warps are admitted a block at a time: if a block is eight warps and the register file allows 51, you get six blocks — 48 warps, not 51. You do not set the register count directly; the compiler picks it from the kernel’s live ranges. You influence it with __launch_bounds__, with -maxrregcount, and by simplifying the inner loop — and you read the result from --ptxas-options=-v.
Ceiling two — shared memory per block
Shared memory is reserved per block, not per thread, and it comes from a fixed on-chip pool per SM. If a block asks for a third of the pool, three blocks fit — and the warps in those three blocks are all the warps you get, however frugal the kernel is with registers. Tiling algorithms that stage large working sets on chip hit this ceiling first, which is why tile size is an occupancy decision and not only a reuse decision.
One lever is worth knowing: where shared memory and L1 share a physical block of SRAM, the split between them is a tunable carveout. Giving shared memory a larger share raises the block ceiling but shrinks L1; giving L1 more helps irregular access at the cost of resident blocks. Large per-block allocations may also require explicitly opting in to the extended limit. The on-chip layout and bank behaviour is a separate subject; here it matters only as a divisor.
Ceiling three — block slots and warp slots
The last two ceilings are pure hardware bookkeeping. Each SM has a maximum number of resident blocks and a maximum number of resident warps, both fixed per generation and not tunable. The warp cap is the denominator of occupancy, so it can never push you below 100% — but the block cap absolutely can, and it is the ceiling people forget.
Consider 32-thread blocks: one warp each. Against a 32-block-per-SM limit that is 32 resident warps, full stop. On a part whose warp cap is 64, you are pinned at 50% occupancy however few registers the kernel uses, because you ran out of block slots. This is the strongest argument for block sizes of 128 to 512 threads: they amortize the block slot across four to sixteen warps. The effect recurs subtly when a block size does not divide the warp cap evenly — a 768-thread block against a 64-warp cap admits two blocks and strands sixteen warp slots.
Working the min() — a concrete launch configuration
Take an SM with illustrative capacities — 65,536 registers, 128 KB of shared memory, 64 resident warps, 32 resident blocks; every one of these varies by generation, so the arithmetic is the point, not the constants. Launch 256-thread blocks (8 warps) using 8 KB of shared memory each. Shared memory allows 16 blocks; the block cap allows 32; the warp cap allows 8 blocks. At 40 registers per thread the register file allows 65,536 / (40 × 32) = 51 warps, which rounds to 6 whole blocks, or 48 warps — 75% occupancy. Registers bind; everything else has slack, so shrinking the shared-memory tile would change nothing.
Now let the inner loop grow. At 64 registers per thread the register file holds 32 warps: 50%. At 96 registers it holds 21 warps, rounded to two blocks: 25%. Occupancy fell threefold while the launch configuration never changed a character. That is the shape of a register cliff — a modest source edit takes a whole block’s worth of warps with it, which is why the register report belongs in your build output.
Theoretical versus achieved occupancy
Everything computed so far is theoretical occupancy: what the resource math permits. Achieved occupancy is what the hardware averaged over the kernel’s actual execution, and it is usually lower. Nsight Compute reports both side by side, and the gap between them is diagnostic on its own.
The gap has three usual causes. Wave quantization: if the grid launches 40 blocks onto a machine that runs 32 at a time, the second wave is 8 blocks wide and most SMs idle through it. Tail effects: blocks finish at different times, and a block’s resources are not released until its last warp retires, so the end of a kernel drains toward empty. Load imbalance: divergent work per block leaves some SMs long since finished. None of these are fixed by touching registers — they are fixed by sizing the grid to the SM count, by grid-stride or persistent-kernel patterns, or by balancing work.
Why higher occupancy is not automatically faster
Occupancy buys latency tolerance, and latency tolerance has a saturation point. Once you have enough warps to cover the memory latency at the achieved bandwidth, the next warp adds nothing — you are bandwidth- or issue-bound, not latency-bound, and the roofline says so. Past that point occupancy is only free if it costs nothing, and it usually costs something.
Three costs recur. First, instruction-level parallelism substitutes for occupancy: a thread holding four independent accumulators covers its own latency without help, and that thread needs registers. Second, cache pressure: doubling resident warps halves the effective L1 and L2 footprint per warp, and a kernel that was hitting cache starts missing it. Third and worst, the spill cliff: force the register count down to buy warps and the compiler spills live values to DRAM-backed local memory. You traded a scheduling win for memory traffic on the hot path — higher occupancy, slower kernel.
When low occupancy is right, and how to tune
Some of the fastest kernels ever written run at low occupancy deliberately. A tensor-core GEMM or attention kernel keeps a large accumulator tile resident in registers precisely so the math units never wait on memory; that tile costs registers, occupancy drops, and throughput goes up, because the kernel is compute-bound and hides latency through deep pipelining rather than warp count. The same holds for any high-ILP inner loop.
So tune in this order. Establish from the profiler whether the kernel is latency-bound at all — long memory-dependency stalls with few eligible warps is the one signature that says more occupancy would help. Only then find which ceiling binds and move that one: block size for slot limits, tile size for shared memory, __launch_bounds__ for registers. The occupancy API can suggest a block size, but treat it as a starting point and re-measure. Occupancy is a dial, not a target.