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?

• 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

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