Occupancy: why 'higher is better' is a misunderstanding
The three occupancy limiters, how to compute them, and when you should deliberately lower occupancy
- Work out how registers, shared memory, and block count jointly cap occupancy
- Use the CUDA Occupancy API to pick a block size automatically
- Understand that ILP can replace part of what occupancy does
- Know when and how to use `__launch_bounds__`
Occupancy = warps actually resident on each SM / maximum warps the hardware supports. The previous stage said a GPU hides memory latency with oversubscription; occupancy measures how many extra warps you gave the scheduler to switch among. It is a means, not a goal. A kernel at 50% occupancy that saturates bandwidth beats a kernel at 100% occupancy that sits idle.
▸Open the occupancy calculatorPick an architecture, tune block size, register count, and shared-memory use, and see which resource is capping your occupancy. This is the fastest way to see how the three interact.The three limiters
Each SM has a fixed pool of resources. How many blocks can reside at once is the minimum of the three. On Ampere (sm_80), each SM has at most 2048 threads (64 warps), 65536 32-bit registers, 164 KB of shared memory, and 32 blocks.
| Limiter | How to compute | Example (block=256) |
|---|---|---|
| Thread cap | 2048 / blockDim | 2048 / 256 = 8 blocks |
| Registers | 65536 / (regs × blockDim) | 32 registers per thread: 65536 / 8192 = 8 blocks |
| Shared memory | 164 KB / use per block | 16 KB per block: 10 blocks |
| Block cap | Hardware constant | 32 blocks |
| Actual resident | Take the minimum | min(8, 8, 10, 32) = 8 blocks = 64 warps = 100% |
Let CUDA compute it for you
1// 1) let the runtime recommend a block size that can reach peak occupancy2int minGridSize, blockSize;3CUDA_CHECK(cudaOccupancyMaxPotentialBlockSize(4 &minGridSize, &blockSize, myKernel,5 0, // dynamic shared memory in bytes (if it depends on blockSize, use the callback overload)6 0)); // block size cap; 0 means unlimited7printf("recommended blockSize = %d\n", blockSize);8 9// 2) query how many blocks a given config can actually keep resident10int numBlocks;11CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(12 &numBlocks, myKernel, blockSize, 0));13 14cudaDeviceProp prop;15CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));16float occupancy = (numBlocks * blockSize / (float)prop.warpSize)17 / (prop.maxThreadsPerMultiProcessor / (float)prop.warpSize);18printf("theoretical occupancy = %.1f%%\n", occupancy * 100.0f);Constrain the compiler with `__launch_bounds__`
The compiler by default uses as many registers as it can to cut instruction count, but it does not know what block size you plan to use. __launch_bounds__ is how you tell it: the compiler then derives a per-thread register budget and, if needed, uses fewer registers (the cost may be a little spilling, in exchange for higher occupancy).
1__global__ void __launch_bounds__(256, 4) myKernel(float* data) {2 // promise: blockDim is at most 256, and you want at least 4 blocks resident per SM.3 // the compiler then caps the per-thread register budget at 65536 / (256 * 4) = 64.4 ...5}6 7// you can also hard-cap the register count without __launch_bounds__8// nvcc -maxrregcount=32 ... (global, too coarse; usually not recommended)Key: high occupancy is not high performance
Vasily Volkov's 2010 talk *Better Performance at Lower Occupancy* made the counterintuitive case: give each thread more independent work (ILP, instruction-level parallelism) and you can get higher performance at lower occupancy. Latency hiding can come from more warps (TLP) or from more independent instructions inside one warp (ILP).
1// low ILP: one element per thread. the add chain is fully serial, so occupancy must be high to fill the pipeline2__global__ void lowILP(const float* a, const float* b, float* c, int n) {3 int i = blockIdx.x * blockDim.x + threadIdx.x;4 if (i < n) c[i] = a[i] * b[i];5}6 7// high ILP: 4 elements per thread; the 4 multiplies are independent and can fly in the pipeline together8__global__ void highILP(const float4* a, const float4* b, float4* c, int n4) {9 int i = blockIdx.x * blockDim.x + threadIdx.x;10 if (i < n4) {11 float4 va = a[i], vb = b[i]; // one 128-bit load covers four scalars12 float4 vc;13 vc.x = va.x * vb.x; vc.y = va.y * vb.y;14 vc.z = va.z * vb.z; vc.w = va.w * vb.w;15 c[i] = vc;16 }17}- ▸Memory-bound kernels: occupancy usually needs to be above 50% to hide latency. Protect that first.
- ▸Compute-bound kernels: 30% to 50% is often enough. Giving resources to register tiling (a larger tile) is usually the better trade.
- ▸Kernels that use many registers (register-tiled GEMM): occupancy may sit at 25%, and that is the intentional, correct design.
- ▸There is only one test: look at
Achieved Occupancyand the warp stall reasons in Nsight Compute. If stalls are mostlyLong Scoreboard(waiting on memory), raising occupancy helps; if not, raising it does nothing.
Self-check
Answer in your head first, then open the solution. If you cannot, reread that section.