Part I · Why the classic model breaks

Feeding the tensor core

On a Blackwell B200 the math unit got so fast that supplying it with data is now the whole kernel. Everything Hopper and Blackwell added — TMA, TMEM, async barriers, warp specialization, clusters — is one forced move after another to stop starving a tensor core that can multiply faster than memory can feed it.

You already know the classic CUDA story from the gpu_kernels track: a thread block loads a tile of A and a tile of B into shared memory, calls __syncthreads(), every thread does its slice of the multiply with a tensor-core MMA (the WMMA / mma.sync path), syncs again, and loops over K. That kernel is correct. On an A100 it is even respectable. On a B200 it leaves roughly 99% of the tensor core idle, and this track is the story of why — and of the entirely different kernel shape that fixes it.

The two numbers that set everything

A Blackwell B200 has two headline rates, and the entire paradigm falls out of their ratio:

• Compute: the 5th-generation tensor core (tcgen05) sustains ≈ 2 PFLOP/s of dense fp16 / bf16 matrix-multiply — that is 2 × 10¹⁵ multiply-adds every second.
• Memory: HBM3e delivers ≈ 8 TB/s — 8 × 10¹² bytes per second from off-chip DRAM into the GPU.

Divide them. The roofline ridge — the arithmetic intensity at which a kernel stops being limited by memory and starts being limited by math — is:

ridge ≈ 2 × 10¹⁵ FLOP/s ÷ 8 × 10¹² B/s ≈ 250 FLOP / byte

Read that as a threshold. To break even with the tensor core, every byte you pull from HBM must be reused in about 250 floating-point operations. Touch a byte fewer times than that and you are memory-bound — the tensor core sits waiting. (These are round, order-of-magnitude ceilings; the exact figures wander with clock, power cap, and how you measure. The ratio, and the wall it implies, is what matters.)

For scale: a square matmul of size N has 2N³ FLOPs and moves ~3N² elements, so its arithmetic intensity is ≈ N/3 FLOP/byte (fp16). A 4096³ GEMM sits at ≈ 1365 — comfortably past the ridge, so it can be compute-bound. An RMSNorm or a GELU reuses each byte a handful of times — intensity ~1 — and lives permanently on the memory roof. The whole craft is getting the matmuls that deserve to be compute-bound to actually run compute-bound. The widget below lets you slide N and watch a point cross the ridge.

Interactive · the B200 roofline and the GEMM ladder

Log–log roofline. The sloped line is the memory roof (slope = 8 TB/s); the flat line is the compute roof (2 PFLOP/s); they meet at the ridge ≈ 250 FLOP/byte. Drag matmul size N to move the large-GEMM point (intensity ≈ N/3) and watch it climb the memory roof, hit the ridge, and saturate. The faint stepping markers are the optimization ladder this track builds — naive kernels live under the compute roof even when their math says they shouldn't, because they are starved, not slow.

Why the synchronous kernel starves

The classic tiled GEMM runs its stages in series. One iteration of the K-loop looks like this, and each engine is idle while it waits on the one before:

loop over K tiles:
    load   A_tile, B_tile  from GMEM into SMEM   // memory path busy, tensor core IDLE
    __syncthreads()                              // everyone waits
    compute  acc += A_tile · B_tile  (mma.sync)  // tensor core busy, memory path IDLE
    __syncthreads()                              // everyone waits
store acc to GMEM                                // epilogue busy, everything else IDLE

On an A100 the math was slow enough that the load overlapped well in practice and this was fine. On a B200 the tcgen05 MMA finishes its slice of work so fast that the load it is waiting on dominates the clock — the tensor core spends most of its cycles parked, and the threads that should be issuing math are instead computing addresses and copying bytes, stealing the very issue slots the math needs. Worse, the accumulator for a big modern tile outgrew the register file entirely. The synchronous loop has three separate problems, and Blackwell added one feature to kill each:

• The math warps touch the load. Generating addresses and copying tiles burns issue slots and registers that belong to the MMA. → fixed by TMA, a hardware copy engine (lesson 04).
• The accumulator no longer fits in registers. Bigger tiles → bigger accumulator fragments → register pressure collapses occupancy. → fixed by TMEM, a dedicated accumulator memory (lesson 05).
• __syncthreads() can't wait on an engine. TMA and the async MMA finish on their own schedule; a thread-barrier can't track "did 49 152 bytes land yet?" → fixed by mbarriers, async barriers with a byte budget and a phase bit (lesson 06).

The fix: turn the kernel into a supply chain

Stop running the stages in series. Run all three hardware paths at once, each working on a different tile, like a 3-stage pipeline in a factory. While the tensor core computes tile k, the TMA engine is already fetching tile k+1, and the epilogue is draining the result of tile k−1:

That is the mental model for the entire track: the kernel is a supply chain, not a loop. The math unit is the bottleneck machine you must never let go hungry, and every other piece of hardware exists to keep its input queue full and its output queue drained. Once you commit to that, the rest of the modern feature set stops looking like a grab-bag of acronyms and starts looking inevitable: you need a way to run producer and consumer warps simultaneously (warp specialization, lesson 10), a way to pool several SMs so one operand load feeds more math (clusters, lesson 07), and a way to keep the chain running across output tiles without paying launch overhead (persistent kernels, lessons 09–10).

The payoff this track builds: a 744× ladder to cuBLAS

The capstone (lessons 08–10) optimizes one kernel — a 4096 × 4096 × 4096 fp16 GEMM on a B200 — across a ladder of steps. Each step removes exactly one thing starving the tensor core. The numbers are the destination of this whole track:

Read the table as a story, not a leaderboard. The single biggest jump lands at ~142× over naive — 53.6 ms drops to 0.49 ms — and it is just adding the TMA copy engine. That one number is the thesis of the track in miniature: the math was never the problem; feeding it was. By step 8 the kernel is genuinely compute-bound (the tensor core is the bottleneck, exactly as it should be), and step 9's final 0.094 ms is cuBLAS parity — the same kernel NVIDIA's own library runs.

The hardware vocabulary you'll need

Two hierarchies recur in every later lesson. First, the execution hierarchy — who runs an instruction, from one lane up to the whole GPU. Hopper inserted a new level (the warpgroup) and a new scope above the SM (the cluster):

Second, the memory spaces — where data can live, from slowest/biggest to fastest/smallest. Blackwell added a brand-new one (TMEM) purely to house the accumulator:

The three knobs every decision turns

Underneath all the acronyms, every design choice in this track is some combination of three recurring elements. Naming them now means later lessons can say "this is a dispatch decision" and you'll know exactly what kind of move it is:

scope layout dispatch

Two generations, one invariant

Everything here targets two NVIDIA architectures, and they share one invariant the tensor core has held since Volta: D = A · B + C. What changes each generation is how the multiply is issued, where the operands live, and where the accumulator lives — which is exactly the scope/layout/dispatch trio above.

The 15 lessons, as one argument

The track is one chain of forced moves in four parts. Each part hands the next a bottleneck it can't solve, and the next part's mechanism is the only way out:

So how do we actually measure "fed" versus "starved", and tell which kernels even deserve to be compute-bound? That's the roofline as a working tool — and the async overlap that lets a real kernel reach it. Lesson 01.