What is occupancy in CUDA?
The fraction of an SM's warp slots your kernel actually fills, which is a means of hiding latency and not a goal in itself.
An SM has a fixed number of warp slots, 32 on a T4, 64 on an H100, 48 on an RTX 50, and your kernel fills as many as its resource appetite allows. Registers, shared memory and the block-count cap each impose a ceiling, and the lowest one wins. Theoretical occupancy is that arithmetic; achieved occupancy is what the profiler saw. The warp scheduler needs some warp ready to issue every cycle, and occupancy is the size of the pool it chooses from, so more occupancy means more chances that somebody is ready. That is all it means.
Which is why the question "does higher occupancy make my kernel faster" has a measured answer: only while the pool is what limits latency hiding. What the hardware needs is memory requests in flight, and warps are one way to get them; independent loads inside one thread are another. A kernel at low occupancy issuing four independent loads per thread can keep as many requests in flight as one at full occupancy issuing one. Past the point where the memory system is saturated, extra warps buy nothing at all.
The API half is cudaOccupancyMaxPotentialBlockSize, which asks the driver rather than guessing from the tables, and asking the driver matters: shared-memory allocation granularity is not published in the compute-capability tables, so a hand calculation can be off by a block. On day 45's kernel the API suggested 1024 threads per block; the fastest measured configuration used 256. The suggestion maximizes residency, not speed, and treating it as a tuning answer is the standard misreading.
Measured
Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -std=c++17 -O3 -arch=sm_75, captured 2026-09-01. Day 45 ran one axpy-style kernel body at ten configurations, dialing occupancy with a never-read dynamic shared-memory reservation so the work stayed identical. Three rows tell the story ("in flight" is resident warps times the loads one thread issues before its first multiply):
| label | occupancy | loads in flight | ms | GB/s |
|---|---|---|---|---|
| tlp1 | 100% | 64 | 0.785 | 256.5 |
| ilp4-b1 | 25% | 64 | 0.798 | 252.2 |
| tlp1-b1 | 25% | 16 | 1.139 | 176.7 |
The first two rows hold loads in flight at 64 and differ by 4x in occupancy: the times land within a ratio of 1.02. The last row keeps the 25% occupancy and cuts the loads in flight to 16: 0.69 of the best row's bandwidth. Occupancy mattered exactly as far as it changed the number of requests in flight, and not one step further.
Diagram
occupancy-stepper, preset residency-dials: one SM's 32 warp slots, with the shared-memory reservation shrinking the resident block count step by step while the work per thread stays fixed.
Alt text: "Filling fewer of an SM's 32 warp slots leaves the kernel's speed unchanged until the total loads in flight fall, which is the number that was doing the work."
Related terms
- resident warps and blocks per SM
- latency hiding
- register file
- shared memory
- warp scheduler
- Little's law
- register pressure
Where you meet this
- Day 45, occupancy, the lesson that owns this term and produced the table above.
- Day 10, picking a block size, where 9 of 13 block sizes sat within a few percent of the best.
- Day 17, registers and spills, where
__launch_bounds__restored 100% occupancy and made the kernel slower. - Day 42, Nsight Compute, where achieved occupancy is read rather than computed.
Sources
- CUDA Programming Guide, compute capabilities appendix, Table 30, for the per-architecture resident warp and block limits: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html (checked 2026-08-30)
- CUDA Programming Guide, "Hardware Multithreading", for the resident execution context that makes a warp slot cheap to fill: https://docs.nvidia.com/cuda/cuda-programming-guide/03-advanced/advanced-kernel-programming.html (checked 2026-08-30)
Byline
Written by: pending. Reviewed by: pending. Written on: pending. Last checked on: pending. The numbers came off the verification node on 2026-09-01, and this entry stays a draft until a named author and a different named reviewer sign it.