gpu register pressure
**GPU Register Pressure** is the **conflict between a kernel's per-thread register demand and the GPU's fixed register file capacity** — where each additional register per thread reduces the number of concurrent threads (occupancy), potentially hiding less memory latency, while reducing registers may cause spills to slow local memory, creating a critical optimization tradeoff for GPU kernel performance.
**GPU Register File Architecture**
- Each NVIDIA SM (Streaming Multiprocessor) has a **fixed register file**: 65,536 32-bit registers (typical).
- Registers are **partitioned** among all active threads on the SM.
- More registers per thread = fewer threads per SM = lower occupancy.
**Occupancy Example (NVIDIA A100)**
| Registers/Thread | Max Threads/SM | Occupancy (of 2048 max) |
|-----------------|---------------|------------------------|
| 32 | 2048 | 100% |
| 64 | 1024 | 50% |
| 128 | 512 | 25% |
| 255 (max) | 256 | 12.5% |
- At 255 registers/thread: Only 256 threads active (8 warps) — very little latency hiding.
- At 32 registers/thread: Full 2048 threads (64 warps) — maximum latency hiding potential.
**Register Spilling**
- When kernel needs more registers than allocated → compiler **spills** excess to local memory.
- Local memory is actually device DRAM (L1/L2 cached) — 100x slower than register access.
- Spilling causes significant performance degradation: 2-10x slowdown for spill-heavy kernels.
**Optimization Strategies**
- **Limit register count**: `__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor)` or `--maxrregcount=N` compiler flag.
- **Reduce live variables**: Recompute values instead of storing them. Reorder operations to reduce simultaneous live registers.
- **Use shared memory**: Move some per-thread data to shared memory (explicitly managed cache).
- **Loop unrolling control**: Aggressive unrolling increases register usage — `#pragma unroll` factor tuning.
**Profiling Register Usage**
- `nvcc --ptxas-options=-v` reports register count per kernel.
- NVIDIA Nsight Compute shows register usage, spills, and occupancy.
- CUDA occupancy calculator: Interactive tool to find optimal register/thread configuration.
**Register Pressure vs. ILP**
- Some kernels benefit from low occupancy + high ILP (instruction-level parallelism per thread).
- Heavy compute kernels (matrix math): Fewer threads with more registers can outperform many threads with spilling.
- **Principle**: Occupancy is not the only metric — achieved throughput is what matters.
GPU register pressure is **one of the most impactful performance-limiting factors in GPU programming** — understanding and managing the register-occupancy-spill tradeoff is essential for extracting peak performance from GPU hardware.