occupancy
**GPU occupancy** measures **the ratio of active warps to maximum possible warps on a streaming multiprocessor (SM)** — indicating how well a kernel utilizes GPU parallel resources, with higher occupancy generally (but not always) correlating with better performance for memory-bound workloads.
**What Is Occupancy?**
- **Definition**: Active warps ÷ Maximum warps per SM.
- **Range**: 0% to 100%.
- **Unit**: Warps (groups of 32 threads).
- **Goal**: Keep GPU execution units busy.
**Why Occupancy Matters**
- **Latency Hiding**: More warps = better memory latency hiding.
- **Utilization**: Higher occupancy often means better GPU use.
- **Memory-Bound**: Particularly important for memory-bound kernels.
- **Not Always Key**: Compute-bound kernels may not need high occupancy.
**Occupancy Calculation**
**Factors Limiting Occupancy**:
```
Resource | Limit Per SM | Impact
------------------|-------------------|------------------
Registers | 65,536 (typical) | More regs → fewer threads
Shared memory | 48-164 KB | More shmem → fewer blocks
Block size | 1024 threads max | Limits parallelism
Warp slots | 64 warps (2048 threads)| Hardware maximum
```
**Example Calculation**:
```
GPU: A100 (64 warps max per SM)
Kernel uses:
- 64 registers per thread
- 256 threads per block
- 8 KB shared memory per block
Registers: 65,536 / (64 × 256) = 4 blocks
Shared memory: 164 KB / 8 KB = 20 blocks
Thread limit: 2048 / 256 = 8 blocks
Bottleneck: Registers (4 blocks)
Active warps: 4 × (256/32) = 32 warps
Occupancy: 32/64 = 50%
```
**Checking Occupancy**
**NVIDIA Tools**:
```bash
# Nsight Compute profiling
ncu --metrics sm__warps_active.avg.pct_of_peak_sustained_active
./my_cuda_program
# CUDA Occupancy Calculator (spreadsheet tool)
# Also available as API
```
**CUDA API**:
```cuda
int blockSize = 256;
int minGridSize;
int maxBlockSize;
cudaOccupancyMaxPotentialBlockSize(
&minGridSize, &maxBlockSize,
myKernel, 0, 0
);
// Use maxBlockSize for kernel launch
```
**PyTorch Kernel Info**:
```python
import torch
from torch.utils.benchmark import Timer
# Profile to see occupancy
with torch.profiler.profile(
activities=[torch.profiler.ProfilerActivity.CUDA],
) as prof:
result = model(input)
print(prof.key_averages().table())
```
**Improving Occupancy**
**Strategies**:
```
Issue | Solution
-------------------|----------------------------------
Too many registers | Use -maxrregcount compiler flag
| Spill to local memory (slower)
| Reduce kernel complexity
|
Too much shared mem| Reduce shared memory usage
| Use global memory (slower)
| Split kernel
|
Block size too small| Increase threads per block
| Aim for multiple of 32
|
Block size too large| Reduce to allow more blocks
```
**Register Limiting**:
```cuda
// Limit registers per thread
__launch_bounds__(256, 4) // 256 threads, 4 blocks/SM
__global__ void myKernel() {
// Compiler will limit registers to achieve this
}
```
**Occupancy vs. Performance**
**Not Always Correlated**:
```
Scenario | High Occupancy | Performance
----------------------|----------------|------------
Memory-bound kernel | Helps | Improves
Compute-bound kernel | May not help | Depends
High ILP | Less important | Good anyway
Low latency needed | Very important | Critical
```
**When Low Occupancy Is OK**:
```
- Kernel is compute-bound
- High instruction-level parallelism (ILP)
- Data fits in cache
- Register usage enables optimizations
```
**Occupancy Guidelines**
```
Occupancy | Interpretation
----------|----------------------------
>75% | Good for memory-bound
50-75% | Usually acceptable
25-50% | May leave performance on table
<25% | Likely suboptimal
```
**Balance With**:
```
Higher occupancy trades:
- Registers (more spills)
- Shared memory (less per block)
- Block flexibility
Lower occupancy allows:
- More registers (faster compute)
- More shared memory
- Compiler optimization
```
GPU occupancy is **one metric among many for kernel optimization** — while important for memory-bound workloads, blindly maximizing occupancy without understanding the kernel's characteristics can actually hurt performance.