Occupancy (CUDA)

Occupancy is the ratio of warps resident on a multiprocessor to the maximum it can hold. It matters because resident warps are how the GPU hides memory latency: when one warp stalls on a load, another issues in its place. Too few warps and the multiprocessor sits idle.

Three resources cap it, and the tightest one wins:

• Registers. Each multiprocessor has a fixed register file, typically 65536 32-bit registers. A kernel using 64 registers per thread allows 1024 threads, so 32 warps.
• Shared memory. A block requesting 32 KB out of a 64 KB budget limits the multiprocessor to two resident blocks.
• Block count. There are hard limits on resident blocks and warps per multiprocessor, so very small blocks waste capacity.

Register usage is not something you write, it is something the compiler decides. Ask it what it did:

$ nvcc --ptxas-options=-v -o prog prog.cu
ptxas info: Used 40 registers, 8192 bytes smem, 380 bytes cmem[0]

__launch_bounds__ caps the register count by telling the compiler the block size you intend to use, trading spills for residency:

__global__ __launch_bounds__(256, 4) void kernel(...) { /* ... */ }

Higher occupancy is not automatically better, and this is the part that catches people out. A kernel with high arithmetic intensity and heavy register use can beat a low-register version at half the occupancy, because the extra registers avoided spilling to local memory. Occupancy is a means to hiding latency, not a goal, and once latency is hidden more warps buy nothing.

The practical order is: get coalescing right first, then check occupancy, then tune. ncu reports achieved occupancy alongside the theoretical figure, and a large gap between the two usually points at load imbalance or divergence rather than at a resource limit.