TL;DR: With 256-thread blocks on an H100 (compute capability 9.0), registers allow 2 blocks per SM (65,536 / 32,768), shared memory allows 4 (228 KB / 49 KB, counting the 1 KB CUDA reserves per block) and the warp limit allows 8, so registers bind: 2 blocks, 16 of 64 warps, 25% occupancy. Cutting to 64 registers doubles it to 50%, where shared memory becomes the limit. Whether that is worth doing depends on whether the kernel is latency-bound, not on the percentage.
How to approach it
Occupancy needs a block size, so ask for it; the question as stated does not fix it, and the answer below assumes 256 threads (8 warps). Then compute each per-SM limit separately (registers, shared memory, warps, blocks) and take the minimum. Finish by asking what the kernel is bound by, because 25% can be the right number.
A strong answer
The H100 per-SM budgets from NVIDIA's Hopper tuning guide and the CUDA compute-capability tables: 65,536 32-bit registers, at most 255 per thread; 228 KB of shared memory per SM, at most 227 KB per block, with 1 KB reserved by CUDA for each resident block; at most 64 warps (2,048 threads) and 32 blocks. Registers are allocated per warp in units of 256 registers, the value the occupancy calculator uses.
registers per warp = 128 regs x 32 threads = 4,096 (multiple of 256)
per block = 4,096 x 8 warps = 32,768
blocks/SM = floor(65,536 / 32,768) = 2
shared mem per block = 48 KB + 1 KB reserved = 49 KB
blocks/SM = floor(228 / 49) = 4
warps blocks/SM = floor(64 / 8) = 8
blocks limit = 32
resident = min(2, 4, 8, 32) = 2 blocks = 16 warps
occupancy = 16 / 64 = 25%
Registers bind. The driver also chooses the shared-memory carve-out from the supported sizes (0, 8, 16, 32, 64, 100, 132, 164, 196 or 228 KB): two resident blocks need 98 KB, so it can pick the 100 KB configuration and leave the rest of the unified 256 KB array to L1.
What moves the answer:
| Change | Reg limit | Smem limit | Blocks | Warps | Occupancy |
|---|---|---|---|---|---|
| As given (128 regs, 48 KB) | 2 | 4 | 2 | 16 | 25% |
__launch_bounds__(256, 3): 80 regs | 3 | 4 | 3 | 24 | 37.5% |
| 64 regs | 4 | 4 | 4 | 32 | 50% |
| 64 regs, 64 KB smem | 4 | 3 | 3 | 24 | 37.5% |
| 168 regs | 1 | 4 | 1 | 8 | 12.5% |
3 blocks: regs <= 65,536 / (3 x 256) = 85.3, rounded down to a multiple of 8 = 80
4 blocks: regs <= 65,536 / (4 x 256) = 64; smem 4 x 49 KB = 196 KB <= 228 KB
5 blocks: smem 5 x 49 KB = 245 KB > 228 KB, so 48 KB caps the kernel at 50%
The 168-register row is the cliff to know: one extra register per thread past a threshold can drop a whole block, which is why a small code change can halve occupancy. The same calculation as code, for checking any configuration:
def occupancy(regs, smem_kb, threads=256):
warps = threads // 32
regs_per_warp = -(-regs * 32 // 256) * 256 # 256-register granularity
by_regs = 65536 // (regs_per_warp * warps)
by_smem = int(228 // (smem_kb + 1)) if smem_kb else 32 # +1 KB reserved
blocks = min(by_regs, by_smem, 64 // warps, 32)
return blocks, blocks * warps, blocks * warps / 64
print(occupancy(128, 48)) # (2, 16, 0.25)
print(occupancy(64, 48)) # (4, 32, 0.5)
Does 25% matter? Occupancy exists to hide latency: while one warp waits on memory, others issue. A kernel whose warps each have many independent loads in flight (unrolled, 16-byte vector loads, a multi-stage shared-memory pipeline) can saturate HBM at 16 warps, and Hopper GEMMs run one large block per SM by design. Forcing 64 registers onto a kernel that needs 128 makes the compiler spill to local memory, and spills are DRAM traffic through L1, often worse than the occupancy lost.
The reversal condition: if the profiler shows warps stalled on long-scoreboard (waiting for global loads) with idle issue slots and low DRAM throughput, the kernel is latency-bound with too little parallelism, and raising occupancy is the fix. If it shows high DRAM throughput or tensor-pipe utilization already, leave occupancy where it is.
What interviewers probe next
"The shared memory is 48 KB of static __shared__ arrays. Can you raise it to 100 KB?" Not statically: 48 KB is the static limit, and more requires dynamic shared memory plus cudaFuncSetAttribute(..., cudaFuncAttributeMaxDynamicSharedMemorySize, ...).
"Switch to 128-thread blocks with the same per-thread resources." Registers allow 4 blocks of 4 warps (still 16 warps), shared memory still allows 4, so occupancy stays 25%, but smaller blocks mean less tail waste when blocks finish unevenly.
"How do you see the real number?" Nsight Compute reports theoretical and achieved occupancy side by side; a large gap points to load imbalance or a short grid, not to resource limits.
How it's scored
Weak: divides 2,048 threads by 256 and answers 100%, or computes only the register limit and never mentions shared memory or the reserved 1 KB.
Hire: computes all four limits with the 256-thread assumption stated, identifies registers as binding, and shows what change reaches 50%.
Strong Hire: knows the allocation granularity and the 48 KB static ceiling, predicts the shared-memory limit taking over at 64 registers, and argues from stall reasons whether higher occupancy would help or would cost spills.
Common mistakes
- Using 64 KB or 256 KB as the H100 shared-memory budget per SM; the configurable maximum is 228 KB.
- Forgetting that CUDA reserves 1 KB per block, which turns 228 / 48 = 4.75 into 228 / 49 = 4.65 here and changes the answer for other sizes.
- Treating higher occupancy as the goal and forcing register caps that introduce spills.
- Computing occupancy for a kernel whose grid has fewer blocks than SMs, where the per-SM limit is irrelevant.