gpu occupancy optimization
**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:
| Resource | H100 SM Limit | How It Limits Occupancy |
|----------|---------------|------------------------|
| **Registers** | 65,536 per SM | Kernel using 64 regs/thread × 256 threads/block = 16,384 regs/block → max 4 blocks/SM |
| **Shared Memory** | 228 KB per SM | Kernel using 48KB/block → max 4 blocks (192KB used) |
| **Thread Blocks** | 32 per SM | Hard limit regardless of resource usage |
| **Warps** | 64 per SM | Maximum occupancy = 64 warps × 32 threads = 2048 threads/SM |
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.