GPU Occupancy Optimization
GPU Occupancy Optimization is the practice of maximizing the number of active warps per Streaming Multiprocessor (SM) relative to the hardware maximum — balancing register usage, shared memory allocation, and thread block configuration to keep the GPU's warp scheduler fully utilized and hide memory access latency.
Occupancy Definition:
- Theoretical Occupancy: ratio of active warps per SM to maximum warps supported by the hardware; e.g., A100 supports 64 warps (2048 threads) per SM; if a kernel achieves 32 active warps, occupancy is 50%
- Achieved Occupancy: actual runtime average of active warps per cycle, accounting for block launch timing and completion; typically 5-15% lower than theoretical due to partial waves and resource fragmentation
- Sufficient Occupancy: diminishing returns above ~50-60% occupancy for compute-bound kernels; memory-bound kernels benefit from higher occupancy (more warps to hide memory latency); the exact threshold is workload-dependent
Resource Limiters:
- Register Usage: each SM has a fixed register file (e.g., 65536 registers on A100); kernel using 64 registers per thread limits occupancy to 1024 threads (32 warps, 50% of max); reducing to 32 registers enables full occupancy but may increase register spilling to local memory
- Shared Memory Per Block: each SM has limited shared memory (e.g., 164 KB on A100 configurable vs L1); a kernel using 48 KB shared memory per block can fit 3 blocks per SM; increasing to 96 KB limits to 1 block per SM
- Thread Block Size: block size must be a multiple of warp size (32); small blocks (32 threads) may not fill SM due to max-blocks-per-SM limits; large blocks (1024 threads) may underutilize SM if resource usage per block is high
- Block Count Limitation: each SM supports maximum blocks (e.g., 32 on A100); very small blocks (32 threads each) with low resource usage may still be limited by block count
Optimization Strategies:
- CUDA Occupancy Calculator: NVIDIA provides API (cudaOccupancyMaxActiveBlocksPerMultiprocessor) and spreadsheet tool; input register count, shared memory, block size → output occupancy percentage and limiting factor
- Launch Bounds: __launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor) directive hints the compiler to limit register usage to achieve target occupancy; may increase instruction count due to spilling but improves parallelism
- Register Pressure Reduction: restructuring code to reduce live variable count; using shared memory for intermediate results; compiler flag --maxrregcount limits register allocation globally
- Dynamic Shared Memory: using extern __shared__ with dynamic allocation rather than fixed arrays allows block size flexibility; combined with occupancy API to select optimal configuration at runtime
Beyond Occupancy:
- Latency vs Throughput: some kernels achieve peak performance at low occupancy by maximizing per-thread register usage and instruction-level parallelism; ILP can hide latency as effectively as thread-level parallelism
- Memory Bandwidth Saturation: memory-bound kernels may saturate bandwidth at 50-75% occupancy; higher occupancy adds warps that compete for the same bandwidth without improving throughput
- Instruction Mix: compute-bound kernels with high arithmetic intensity need fewer warps to saturate compute pipelines; memory-bound kernels need maximum warps to generate enough outstanding memory requests
GPU occupancy optimization is a crucial but nuanced aspect of CUDA performance tuning — high occupancy is necessary for memory-bound kernels to hide latency, but blindly maximizing occupancy at the expense of per-thread efficiency can hurt compute-bound kernels — the optimal balance requires understanding the kernel's arithmetic intensity and profiling with Nsight Compute.