Part I · Why the classic model breaks
What makes a kernel fast now
Lesson 00 said the tensor core is starving. This lesson makes "fed vs starved" a number you can compute, and replaces the old lever — "add more warps to hide latency" — with the two that actually feed a Blackwell tensor core: the roofline and explicit async overlap.
__syncthreads() for the next tile to arrive. The obvious fix from the gpu_kernels era was occupancy: launch more warps so that while one warp stalls on memory, another has work ready. On a B200 that lever is broken — and we need to know why before we can pick a real one.
First: turn "fed vs starved" into a number
You can't optimize what you can't measure. The roofline model is the measurement. It answers one question for a given kernel: what is the highest FLOP/s this kernel could possibly attain on this GPU, given how much data it has to move? The bound is a single inequality:
attainable FLOP/s ≤ min( peak FLOP/s, bandwidth × arithmetic intensity )
Two terms compete inside the min, and each is a "roof" you cannot punch through:
- The compute roof — a flat ceiling at the hardware's peak FLOP/s. No kernel goes faster than the silicon can multiply-add.
- The memory roof — a sloped ceiling, bandwidth × AI. If you only move bytes at BW and each byte buys you AI FLOPs, that product caps your FLOP/s. Double the bytes per useful FLOP and this roof drops by half.
The lever is arithmetic intensity (AI): useful FLOPs done per byte moved.
AI = useful FLOPs / bytes moved [FLOP/byte]
Plot attainable FLOP/s (y) against AI (x) on log–log axes and you get the classic shape: the sloped memory roof rises with AI until it hits the flat compute roof. They cross at the ridge point:
AIridge = peak FLOP/s / bandwidth
Left of the ridge, the memory roof is lower — you are memory-bound (starved): bytes arrive too slowly to keep the math busy. Right of the ridge, the compute roof is lower — you are compute-bound (fed): the math is the bottleneck, which is exactly where a tensor-core kernel wants to live.
The B200 ridge: ≈250 FLOP/byte
Put the numbers in. A B200 sustains roughly 2 PFLOP/s of fp16 tensor-core throughput, fed by roughly 8 TB/s of HBM3e. The HBM ridge is their ratio:
AIridge = 2×10¹⁵ FLOP/s ÷ 8×10¹² B/s ≈ 250 FLOP/byte
So the bar is brutally high. To be compute-bound on a B200, a kernel must do ~250 useful FLOPs for every byte it reads from HBM. Anything less and the sloped roof — not the tensor core — sets the speed. (All three numbers are vendor-class approximations; the exact peak depends on clocks, the exact ridge on your HBM SKU. The point is the order of magnitude: hundreds of FLOPs per byte, not tens.)
Arithmetic intensity by workload — who is starved?
Now classify the operations in a transformer block by where they fall relative to that ridge.
Elementwise and reductions: memory-bound by nature
A GELU reads a big activation tensor, does a handful of FLOPs per element, and writes it back. Read one element, write one element, do a few FLOPs on it — an arithmetic intensity of order one FLOP per byte. RMSNorm is a reduction: read the row, sum of squares, scale, write — again a few FLOPs per element over two tensor traversals. Both sit at AI of order one, hundreds of times below the ridge. They are pinned to the far left of the roofline; the tensor core is irrelevant to them. These ops are memory-bound no matter how you write them — their job is just to move bytes, so the only win is to move fewer of them.
GEMM: compute-bound at the sizes that matter
A square fp16 matmul with M = N = K = N does N³ multiply-adds = 2N³ FLOPs, and (in the naive accounting) touches three N × N matrices at 2 bytes each. So:
AI = 2N³ / (3 · N² · 2) = N / 3 FLOP/byte
The numerator 2N³ is N³ multiply-adds × 2 FLOPs each; the denominator is 3 matrices × N² elements × 2 bytes. AI grows linearly with N. At N = 4096:
AI ≈ 4096 / 3 ≈ 1365 FLOP/byte ≫ 250
That is more than 5× past the ridge — GEMM at this size is deeply compute-bound. The algorithm permits living against the compute roof. Which sets up the single most important sentence in this lesson:
Attention: the score-matrix trap
Attention is the cautionary middle case. The matmuls inside it (QKᵀ and P·V) are high-AI like any GEMM. The problem is the score matrix in between: a naive implementation writes the full N × N scores out to HBM, reads them back to softmax, writes again, reads again for the value matmul. That round-trip moves an enormous intermediate across the HBM boundary for almost no extra FLOPs — it drags attention's effective AI down to memory-bound territory even though its matmuls are not. (FlashAttention's whole reason to exist is to keep that score matrix on-chip and never round-trip it — a later lesson.)
Two responses, depending on which roof you hit
The roofline doesn't just diagnose; it tells you which knob to turn. There are exactly two situations, and they call for opposite work.
If you're memory-bound (left of ridge): two ways out
(A) Raise AI — slide right, toward the compute roof. More FLOPs per byte:
- Fusion — the highest-leverage move. If an intermediate (the GELU output, the attention scores) is produced and immediately consumed, don't round-trip it through HBM — keep it in registers / SMEM / TMEM and fuse the producer and consumer into one kernel. You kill the bytes entirely, so AI can jump by an order of magnitude. This is why
matmul + bias + GELUfused beats three separate kernels. - Blocking for reuse — load a tile of an operand into SMEM once and use it for many output elements. Each byte read from HBM now serves many FLOPs instead of one, so AI rises. This is the entire point of tiled GEMM.
- Smaller dtype — fp32 → fp16 → fp8 → fp4 halves the bytes per element (and doubles the tensor-core FLOP rate) at each step, lifting both AI and the compute roof. Honest caveat: the real speedup is always less than the raw 2×, because low precision drags along scale factors, block metadata, and conversion work that the clean ratio ignores.
(B) Saturate the memory roof — if AI genuinely can't rise (a pure copy, a normalization), then accept memory-bound and get as close to the sloped roof as possible. Move each byte exactly once; eliminate redundant reads; coalesce and vectorize so every transaction is full-width; use a hardware copy engine (TMA) for regular bulk tiles; and keep enough requests in flight that the HBM channels never sit idle.
If you're compute-bound (right of ridge): keep the math units busy
Bytes are no longer the constraint — issue is. The job becomes scheduling: make sure a tensor-core instruction is always ready to fire, with its operands already on-chip, so the math pipe never bubbles. That is the rest of this track. And it starts with a scheduling idea that the occupancy mindset gets backwards.
Latency hiding: overlap, not occupancy
Every kernel that streams tiles has the same dependency chain per step: load a tile from HBM, compute on it, store the result. Written the obvious way, those happen in series:
// naive synchronous schedule — one engine works at a time
for k in 0..K:
load(k) // HBM busy, tensor core IDLE
compute(k) // tensor core busy, copy engine IDLE
store(k) // store path busy, the rest IDLE
At every moment exactly one hardware path is doing useful work and the others stall. The tensor core spends most of the loop waiting for load(k). That is the idle we measured on the roofline.
The fix is to software-pipeline: while the tensor core chews on tile k, the copy engine is already fetching tile k+1, and the store path is draining tile k−1. Same dependencies — you still can't compute a tile before it lands — but they're scheduled so different tiles occupy different engines at the same instant:
// pipelined schedule — engines run concurrently on different tiles
load(0) // prologue: prime the pipe
for k in 0..K:
load(k+1) // copy engine fetches NEXT tile ...
compute(k) // ... while tensor core computes THIS tile
store(k-1) // ... while the store path drains the LAST tile
k+1 runs while the MMA for tile k is in flight. The dependency (compute needs its tile loaded first) is untouched; what changed is that the long-latency load is overlapped with independent compute instead of blocking it.
Why occupancy stopped being the answer
Here is the crux of the lesson. The classic way to hide that load latency was occupancy — pack many warps onto each SM so the warp scheduler always has some warp whose operands are ready while others stall. That worked when a "compute step" was a few CUDA-core FMAs and you needed dozens of warps in flight to paper over memory latency.
But occupancy is capped by physical resources — registers, SMEM, warp slots, and CTA slots per SM — and a modern tensor-core kernel deliberately spends all of them on a small number of CTAs:
- Multi-stage SMEM pipelines eat SMEM — you stage several tiles ahead, so each CTA's SMEM footprint balloons and fewer CTAs fit.
- Big register fragments (and large accumulators) eat registers — fewer warps fit per SM.
- TMEM (Blackwell's dedicated accumulator memory) eats tensor-memory, another scarce per-SM resource.
- Warp specialization reserves whole warps for a single role (a producer warp that only issues copies, a consumer warp that only issues MMAs) — capacity spent on coordination, not on more parallel copies of the same work.
So these kernels run at low occupancy on purpose. They don't hide latency by having a spare warp ready — they hide it by explicit overlap inside a few resident CTAs: a copy engine, the tensor core, and the store path all working at once on different tiles. The hardware paths stay busy not because there are many warps, but because the one resident CTA keeps every engine fed.
Interactive · the B200 roofline calculator
Pick a square matmul size N and a dtype. The widget computes the time to do 2N³ FLOPs at that dtype's tensor-core peak (t_compute) versus the time to move the three N × N matrices across HBM at 8 TB/s (t_HBM), reports the arithmetic intensity against the ≈250 FLOP/byte ridge, and draws both as bars — the kernel's wall-clock is whichever bar is longer. Slide N down to watch a compute-bound GEMM tip over into memory-bound; switch fp16 → fp8 to watch the compute roof drop and the crossover move.
Notice what the calculator makes concrete. At fp16, the crossover N (where the bars are equal) is small — only a couple hundred — so any real matmul is compute-bound, and the only way to go faster is to keep the tensor core busy. Drop to fp8 and the compute roof doubles: the same N is now more compute-bound, the crossover pushes higher, and the AI-vs-ridge gap widens because the ridge itself moved (peak doubled, bandwidth didn't). And at fp32 — CUDA cores only, no tensor core — the "peak" is ~30× lower, so even huge matmuls struggle to be impressively fast: that is the throughput the whole async paradigm exists to leave behind.
The six async mechanisms this track will build
Explicit overlap needs hardware that can run independent work on independent paths and coordinate them without thread-level barriers. The Hopper/Blackwell stack provides exactly six pieces, and the rest of Part II derives each one as a forced move. A preview, so you can see where we're going:
| Mechanism | What it does | Path it keeps busy |
|---|---|---|
| TMA | hardware copy engine: GMEM → SMEM bulk tiles, no math threads involved | the copy path |
| mbarriers | async barriers that signal completion and hand work between roles | coordination |
| tcgen05 | async tensor-core MMA issued by one elected thread | the math path |
| TMEM | dedicated memory for long-lived accumulators (off the register file) | accumulator storage |
| warp specialization | split the SM into producer (memory) and consumer (math) warps | both paths, concurrently |
| clusters | pool several SMs into bigger cooperative tiles + operand multicast | cross-SM reuse |
Six different mechanisms, one scheduling goal: keep useful work running on more than one hardware path at once. Every one of them is a way to make the pipelined schedule above real instead of aspirational — and every one of them costs occupancy, which is exactly why we had to retire occupancy as the thing we maximize.
gpu_kernels/10 · GPU as a bandwidth machine; this lesson is that model upgraded for a tensor core fast enough that occupancy no longer feeds it. The roofline reasoning here is the same lens used to budget an entire training run end-to-end in cs336 · Language Modeling from Scratch, where "is this op memory- or compute-bound?" decides what to fuse, what to shard, and what precision to train in.
So the workflow is fixed: estimate AI → locate the roof → decide memory- vs compute-bound → optimize the resource that actually sets the ceiling. But notice what both responses quietly assumed. Fusion keeps an intermediate "in registers / SMEM / TMEM"; overlap fires an MMA "with its operands already on-chip." Both depend on getting the right bytes into the right place in a tensor-core-readable form — and a tensor core does not accept bytes in any old arrangement. Before we can fuse or overlap anything, we need the rules for where bytes live and how they're shaped: a layout algebra.