GPU fundamentals
Occupancy and launch overhead
Resident warps hide HBM latency, and a 5–10 µs launch tax wrecks decode when every op is tiny.
Two different "the GPU is idle" stories show up in serving. One is a kernel that cannot hide HBM latency because too few warps fit on the SM. The other is the CPU launching hundreds of tiny kernels per token and paying ~5–10 µs each time. Occupancy is the first. Launch overhead is the second.
Concept
Occupancy is the fraction of maximum warps that can run on an SM given a kernel's register and shared-memory footprint. Higher occupancy means more warps to switch to when one stalls on HBM. It is not the same as performance — some kernels peak at moderate occupancy — but very low occupancy often means you cannot hide latency.
The big idea
A shop with one cashier cannot hide a slow card reader. More cashiers (warps) keep the line moving while one waits. The analogy breaks when the shop is empty because the manager (the CPU) takes 8 µs to unlock the door for every customer. That is launch overhead, and CUDA graphs and fusion are how serving fixes it.
Walking through occupancy
Suppose an SM allows at most 64 warps resident (architecture-dependent; use the number for your compute capability). Your kernel uses so many registers that only 16 warps fit.
Occupancy ≈ 16 / 64 = 25%If each warp regularly stalls ~200 cycles on HBM and the math between stalls is short, 25% occupancy may not hide latency. Reducing register pressure (smaller tile, -maxrregcount, rewrite) to fit 32 warps raises occupancy to 50% and often improves bandwidth-bound kernels more than micro-optimizing arithmetic.
| Use more… | Benefit | Cost |
|---|---|---|
| Registers per thread | Fewer spills, more ILP | Fewer resident warps |
| Shared memory per block | Better tiling | Fewer concurrent blocks |
| Threads per block | Amortize sync | Must fill warps; diminishing returns |
128–256 threads (4–8 warps) per block is a common starting band. Very small blocks underuse warp slots. Very large blocks may limit how many blocks coexist when shared memory is tight. A custom kernel that blows the register file can still surprise you.
Launch overhead, streams and sync
Training of huge matmuls hides the launch tax. Decode may run hundreds of small ops per token, so graphs, fusion and compiled runtimes matter. CUDA streams overlap a copy in one stream with a kernel in another. Forget streams and you leave bandwidth on the table.
__syncthreads() is required for shared-memory handoff, but each barrier ends warp overlap inside the block. Host-side cudaDeviceSynchronize() or .item() in PyTorch forces the CPU to wait. Avoid device-host sync per token.
Trade-off
Many small kernels are easier to write (PyTorch eager) and lose on launch tax and HBM round trips. Fused kernels win for decode. Default path: libraries + torch.compile / engine graphs. Custom CUDA or Triton when a profile shows a hot, unfused chain.
Common mistakes
- Chasing 100% occupancy. Some kernels are faster with more registers and fewer warps.
- Synchronizing every token.
.item()orcudaDeviceSynchronize()kills the async pipeline. - Ignoring launch tax. A beautiful 2 µs kernel that you launch 200 times per token is a 1 ms overhead bill.
Check yourself
- 16 resident warps on an SM that can hold 64. What is occupancy? Answer: 16 / 64 = 25%.
- Why does launch overhead matter more for decode than for training? Answer: decode runs many tiny ops per token; training hides the tax inside huge matmuls.
Key takeaways
- Occupancy trades registers and shared memory against latency hiding. More is not always better.
- Launch overhead of ~5–10 µs wrecks decode. Fuse, compile and graph.
- Avoid host-device sync on the token path.
Further learning
The compile and graph levers live in kernels and compilation. For writing fused kernels in Python, see Triton.
Related
The CUDA programming model
Kernels, grids, blocks and threads, memory coalescing and occupancy, plus Triton as the higher-level path.
Kernels, fusion and compilation
Kernel fusion, torch.compile, CUDA graphs and the libraries that make models fast.
Coalescing and shared-memory tiling
Adjacent warp loads become few HBM transactions, and tiles in shared memory let you reuse each loaded byte.