occupancy optimization
**Occupancy Optimization** — maximizing the number of active warps on a GPU Streaming Multiprocessor (SM) to hide memory latency through warp-level parallelism.
**What Is Occupancy?**
$$Occupancy = \frac{\text{Active warps per SM}}{\text{Max warps per SM}}$$
- Each SM can hold a maximum number of concurrent warps (e.g., 64 on A100)
- Higher occupancy → more warps to schedule → better latency hiding
**What Limits Occupancy?**
1. **Registers per thread**: More registers per thread → fewer threads fit on SM
- SM has 65536 registers. Thread using 64 regs → 65536/64 = 1024 threads max
2. **Shared memory per block**: More SMEM per block → fewer blocks fit on SM
3. **Block size**: Must be multiple of 32 (warp size). Max 1024 threads per block
4. **Blocks per SM**: Hardware limit (e.g., 32 blocks per SM on Ampere)
**CUDA Occupancy Calculator**
```bash
# Launch configuration for 75%+ occupancy:
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize, kernel);
```
**Best Practices**
- Start with 256 threads per block (good default)
- Reduce register usage: `__launch_bounds__(maxThreads, minBlocks)`
- Profile with Nsight Compute → check achieved occupancy
- Higher occupancy doesn't always mean higher performance (compute-bound kernels may not need it)
**Typical targets**: 50-75% occupancy is usually sufficient. 100% is often impossible and unnecessary.
**Occupancy** is a key metric in GPU optimization — but always measure actual performance, not just theoretical occupancy.