GPU Occupancy is the ratio of active warps on a Streaming Multiprocessor (SM) to the maximum number of warps the SM can support — a key performance metric that determines the GPU's ability to hide memory latency through warp switching, where insufficient occupancy (too few active warps) leaves the SM idle during memory stalls while excessive resource usage per thread (registers, shared memory) is the primary factor that limits occupancy.

Why Occupancy Matters

GPU performance relies on latency hiding through massive multithreading. When one warp stalls on a memory access (~400 cycles), the SM instantly switches to another ready warp at zero cost (hardware warp scheduling). But this only works if there are enough warps ready to execute. If occupancy is too low (e.g., 25%), the SM exhausts ready warps and stalls.

Occupancy Limiters

Each SM has fixed resources. The occupancy is the MINIMUM imposed by any resource:

1. Registers: Each SM has a register file (e.g., 65,536 registers on Ampere). If a kernel uses 64 registers/thread and threads come in warps of 32: 64 × 32 = 2,048 registers per warp. Max warps = 65,536 / 2,048 = 32 (but SM max may be 48). Reducing register usage to 48/thread: 48 × 32 = 1,536/warp → 42 warps. Higher occupancy.

2. Shared Memory: If a block uses 48 KB of shared memory, and the SM has 164 KB configured as shared, max 3 blocks per SM. If block size is 256 threads (8 warps): 3 × 8 = 24 active warps out of 48 max = 50% occupancy.

3. Thread Block Size: If block size is 64 (2 warps) and max blocks per SM is 16: 16 × 2 = 32 warps. Larger blocks (256 threads) may allow higher occupancy if other resources permit.

The Occupancy Trap

Higher occupancy does NOT always mean higher performance:

• A kernel at 50% occupancy using more registers per thread may outperform 100% occupancy with register spilling (register values stored to/loaded from slow local memory).
• A kernel with extensive shared memory reuse at 25% occupancy may be compute-bound and fully utilizing the ALUs.
• The goal is ENOUGH occupancy to hide latency — typically 40-60% is sufficient for many kernels.

Tuning Tools

• CUDA Occupancy Calculator: Given kernel resource usage, computes theoretical occupancy. Available as spreadsheet and cudaOccupancyMaxActiveBlocksPerMultiprocessor() API.
• Nsight Compute: Reports achieved occupancy, active warps, and identifies the limiting resource (registers, shared memory, or block count).
• Launch Configuration: __launch_bounds__(maxThreadsPerBlock, minBlocksPerSM) hints to the compiler to limit register usage for target occupancy.

GPU Occupancy is the resource-constrained balancing act of GPU programming — trading per-thread resource richness (registers, shared memory) against parallelism (active warps), where the optimal balance depends on whether the kernel is memory-latency-bound, compute-bound, or bandwidth-bound.