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**.