Occupancy as a Mental Model

Occupancy as a Mental Model

Thesis

Occupancy is the fraction of a multiprocessor’s theoretical warp capacity that your kernel actually resides. It is a ceiling on latency-hiding — not a score to maximise blindly. Registers, shared memory, and block size are the knobs.

Mental model

  SM resources:
    - register file (per-thread regs × threads)
    - shared memory (per-block)
    - max warps / blocks / threads per SM

  occupancy ≈ resident_warps / max_warps
  limited by whichever resource runs out first

Worked example

Docs/calculator pattern (CUDA Occupancy Calculator / cudaOccupancyMaxActiveBlocksPerMultiprocessor — no GPU here):

Kernel A: 64 regs/thread, 0 shared, blockDim=256
  → often high occupancy on modern SMs

Kernel B: 64 regs/thread, 48 KB shared / block, blockDim=256
  → shared memory may cap blocks/SM → lower occupancy

API sketch:

int blocks;
cudaOccupancyMaxActiveBlocksPerMultiprocessor(
    &blocks, myKernel, blockSize, dynamicSharedBytes);
float occ = blocks * blockSize / (float)(maxThreadsPerSM);

Trap

Chasing 100% occupancy. A kernel with 50% occupancy and better instruction-level parallelism / cache behaviour often beats a register-starved 100% occupancy variant. Occupancy is necessary for latency hiding when the kernel is latency-bound; it is not sufficient for performance.

Boring rule

Report occupancy beside runtime. Change one resource (regs via launch bounds, shared, block size) at a time and re-measure. Prefer Nsight Compute’s achieved occupancy vs theoretical.

Try this

1. Sweep block sizes 128/256/512 for one kernel and plot occupancy vs time.
2. Add __launch_bounds__ and watch register allocation change.
3. Read the occupancy section of the CUDA Best Practices Guide.

Sources

• CUDA Best Practices Guide — occupancy
• CUDA Runtime API: cudaOccupancy*
• Nsight Compute occupancy charts