GPU Occupancy Optimization is the performance tuning discipline that maximizes the number of active warps per Streaming Multiprocessor (SM) — measured as the ratio of active warps to the SM's maximum supported warps — to ensure that the GPU's warp scheduler always has warps ready to execute, hiding memory latency through context switching between warps rather than stalling on any single memory request.

Why Occupancy Matters

GPU SMs operate by rapidly switching between active warps. When one warp stalls on a global memory access (~500 cycles), the scheduler immediately switches to another ready warp. With enough active warps (high occupancy), the SM stays busy while stalled warps wait for data. With too few warps (low occupancy), all warps may be stalled simultaneously → the SM sits idle.

Resources That Limit Occupancy

Each SM has fixed quantities of three resources shared among all active thread blocks:

Occupancy is limited by whichever resource is exhausted first.

Register Pressure

Registers are the most common occupancy limiter. A complex kernel with many variables may use 128 registers per thread, limiting occupancy to 2 blocks of 256 threads (25% occupancy). Reducing register usage (via __launch_bounds__, algorithmic simplification, or register spilling to local memory) increases occupancy — but spilling registers to memory adds latency. The optimum is usually 50-75% occupancy, not maximum occupancy.

Diminishing Returns

Occupancy beyond 50% often provides minimal additional performance because: 1. Enough warps already exist to hide memory latency. 2. Cache thrashing increases as more blocks compete for the same L1/shared memory. 3. Register spilling to achieve higher occupancy adds local memory traffic that offsets the latency-hiding benefit.

The right approach: start at maximum occupancy, benchmark, then systematically trade occupancy for more registers/shared memory per thread if it improves IPC.

Tooling

• CUDA Occupancy Calculator: Excel spreadsheet or cudaOccupancyMaxActiveBlocksPerMultiprocessor() API. Takes kernel register count, shared memory, and block size → reports achievable occupancy.
• Nsight Compute: Profiles actual vs. theoretical occupancy and identifies the limiting resource. Shows achieved occupancy (affected by workload) vs. theoretical (resource-limited).

GPU Occupancy Optimization is the art of balancing resource allocation per thread against total active parallelism — giving each thread enough registers and shared memory to work efficiently while ensuring enough warps exist to keep the SM continuously busy.