TL;DR: Occupancy is resident warps divided by the architectural maximum, not the fraction of compute units doing useful work. Registers, shared memory, threads and block limits jointly constrain residency. A register cap helps only when the extra ready work outweighs spills, recomputation or lost instruction-level parallelism. Measure runtime as well as occupancy.
How to approach it
Compute an illustrative resource bound, then name what the real occupancy query adds. Separate resident, eligible and issued warps. Make a register-cap decision from compiler output and profiler evidence, with a clear condition for reverting it.
A strong answer
A compiler change increases occupancy but makes the kernel slower. That is possible: more warps now fit, but each may need extra memory loads, recompute values it previously kept in registers, or have fewer independent operations. The first question is what limited useful progress before the change, not whether the percentage rose.
Calculate the tightest limit
For an ordinary SM, NVIDIA documents 64 resident warps, 65,536 32-bit registers and at most 32 blocks. The guide lists 228 KB shared memory per SM and 227 KB per block. Actual usable shared memory depends on configuration and reservations; query the target rather than treating those headline capacities as freely allocatable bytes.
The diagram deliberately uses a simplified 228 KiB shared-memory budget to show the calculation. With 256 threads per block, there are eight warps per block. Forty registers per thread consumes 10,240 registers per block before allocation rules; 65,536 / 10,240 permits six whole blocks. The illustrated shared and thread limits are looser, so six blocks × eight warps / 64 = 75% theoretical occupancy.
Illustrative register-only bound: 65,536 registers; 256 threads/block
48 registers/thread → floor(65,536 / (48 × 256)) = 5 blocks
5 blocks × 8 warps/block / 64 warps = 62.5% occupancy
64 registers/thread → 4 blocks → 50%
128 registers/thread → 2 blocks → 25%
Check: halving registers from 128 to 64 doubles this block limit.
Forty-eight registers does not imply 50% occupancy. Real register allocation uses hardware allocation units and compiler-generated resource counts, so small source changes can cross a step rather than cause a smooth percentage change.
Use cudaOccupancyMaxActiveBlocksPerMultiprocessor for the compiled kernel, block size and dynamic shared-memory request. Include static shared memory and the actual launch configuration. Cluster launches need cluster-aware reasoning, including cudaOccupancyMaxActiveClusters; the simple independent-block picture is incomplete there. GPU Execution Model provides the underlying thread/block/SM vocabulary.
Resident, eligible and issued are different states
A resident warp occupies resources. An eligible warp has an instruction ready for a scheduler. An issued warp was selected to execute. Many resident warps can all wait on the same memory or barrier dependency. Conversely, fewer warps can sustain a busy pipeline if each exposes enough independent work.
A simple latency-hiding estimate illustrates the distinction. Suppose an SM's workload needs an assumed 25 GB/s with a 600 ns memory round trip. Bandwidth × latency requires roughly 15,000 bytes in flight. If a warp issues one independent 512-byte load footprint, that is about 30 warps' worth. Eight warps with four such independent loads each would expose 16,384 bytes instead. These are illustrative outstanding-byte counts, not a promise that either schedule reaches the bandwidth: coalescing, cache hits, queue limits and dependencies also matter.
Occupancy and Register Pressure connects this to instruction-level parallelism. A larger per-thread register working set can support more independent operations, reducing the number of warps needed to cover a delay. Register count alone does not tell you whether those operations are independent.
Decide whether to limit registers
Record the unmodified kernel's runtime, compiler register count, shared-memory usage and local-memory accesses. Use Nsight Compute's occupancy and scheduler statistics to examine active and eligible warps, issue behavior and stalls. Local memory is backed by device memory and cached; it is not an on-chip register extension. Local traffic can also come from explicit thread-local arrays, so its presence is not by itself proof of spilling.
Test a register limit when resource-bound residency leaves insufficient ready work. Compare the generated code and duration. Keep it if the extra useful concurrency improves representative workloads with acceptable numerical behavior. Revert it if extra instructions, local traffic or reduced per-thread parallelism outweigh that benefit. If shared memory already limits residency, lowering registers may not add a single block.
__launch_bounds__ tells the compiler launch constraints and a residency target used in its register heuristic. It does not guarantee that the runtime places that many blocks on each SM: shared memory and other limits still apply. A compiler can reorganize or recompute work rather than spilling a fixed “needed minus allowed” number of registers. When already below a bound, it can also use more registers to reduce instructions. Inspect the result instead of assuming that the annotation always lowers register count.
What interviewers probe next
- Why can lower occupancy win? It may preserve independent work and avoid expensive local accesses.
- What would justify a cap? Evidence that register-limited residency prevents enough eligible work, followed by an actual runtime improvement.
- Why can achieved occupancy be lower than theoretical? Short launches, tail waves, workload imbalance and runtime behavior can leave theoretical slots unused.
Common mistakes
- Reporting occupancy as Tensor Core utilization or treating every resident warp as ready.
- Rounding the 48-register example down to four blocks without a specific allocation constraint.
- Claiming a universal “best occupancy” range or applying independent-block arithmetic to clusters unchanged.
Key takeaways
- The illustrated 40/48/64-register cases produce 75%/62.5%/50% register-limited occupancy.
- The bottleneck is the minimum across resources; another resource can make a register cap pointless.
- Keep the cap only when the measured workload improves, not merely the displayed percentage.
Primary references: Hopper occupancy limits, Nsight Compute profiling guide, and CUDA launch bounds.