Part IV · the second capstone and the craft · the finale

Synthesis — the supply chain and the real toolchain

Fourteen lessons, one sentence: the tensor core got so fast that feeding it became the entire kernel. Every primitive you learned was a forced move to stop one more thing from starving it. This lesson collapses the whole paradigm into one picture, names the invariant that ran through nearly every page, and places it all in the real-world toolchain — so you know what to actually reach for, and what knowing all this is for.

The bottleneck lesson 13 left us — and why it closes the loop

Lesson 13's lesson was sobering: the async paradigm puts the entire correctness burden on you — no compiler pass verifies your phases, your byte counts, or your fences — and then it showed how the compiler does lower tile primitives into PTX. That is the last mechanism. We now have every piece: a copy engine, an off-register accumulator, async barriers, warp roles, clusters, persistent schedulers, two assembled capstones, and a debugging discipline. The only thing left is to see them as one machine and to be honest about where that machine actually lives. So the forced move of the finale is not a new mechanism — it is compression: fourteen forced moves into one supply chain, and one supply chain into the libraries you will really call.

Recall the engine: every primitive was a forced move against starvation

Walk the whole chain compactly. Read it top to bottom and notice that no link is "another option" — each is the only way out of the wall the previous one hit. (FLOP/byte ridge ≈ 250 on a B200 is the number that started it: ~2 PFLOP/s of fp16 against ~8 TB/s of HBM. All the ridge / SM-count / TMEM-capacity figures are computed, as flagged in their lessons.)

00 · 01The roofline says no. A single tcgen05 MMA is so large and fast that the classic synchronous loop — load tile → __syncthreads() → compute — cannot keep the tensor core busy. Threads spent generating addresses and copying tiles are pure overhead stealing issue slots from the math. Modern kernels deliberately lower occupancy and hide latency with explicit async overlap instead.

02You need a layout algebra to even place the bytes. Operands have rigid layout requirements and SMEM bank conflicts serialize. Layout = shape:stride; transpose is a layout change, not a copy; swizzle (XOR-with-row) is what makes a column access conflict-free. Layout is part of the instruction interface.

03The operand/accumulator contract changed across generations. D = A·B + C is invariant since Volta; what changed is how it's issued and where the operands and accumulator live — Ampere mma.sync (register fragments) → Hopper wgmma (SMEM descriptors, accumulator still in registers) → Blackwell tcgen05.mma (one elected thread, accumulator in TMEM).

04TMA frees the math warps from the load. A hardware copy engine moves a rectangular tile from GMEM to SMEM and applies swizzle on write, so the math warps never touch the load. Load completion is an mbarrier + expect_tx byte count; store completion is commit_group — the load/store asymmetry.

05TMEM frees the accumulator from registers. Larger MMA tiles grew the accumulator past the register file → occupancy collapse. Blackwell breaks the link: a dedicated 128-lane × ≤512-column scratchpad (~256 KB/SM, computed), read back with tcgen05.ld (lane l → row l/4) so the epilogue reuses Ampere-era cast-and-store code.

06mbarriers coordinate the now-asynchronous engines. TMA and tcgen05 finish on their own time, and __syncthreads() is thread-only and stateless — it can't wait on an engine or track a byte count. An mbarrier has an arrival counter + phase bit; completion flips the phase; you track the expected phase in a register. One barrier per stage, not per iteration — forget to flip the phase and the consumer reads last round's data.

07Clusters pool SMs and CLC kills the tail. One CTA's SMEM can't stage a big-enough tile or reuse what a neighbor already loaded. Clusters co-schedule CTAs with distributed SMEM + multicast; 2-CTA cooperative MMA doubles arithmetic intensity; clusterlaunchcontrol turns the static persistent schedule into hardware work-stealing.

08 · 10GEMM assembles it all to cuBLAS parity. The ladder climbed from a correct-but-slow synchronous tile to warp-specialized, clustered, multi-consumer SOTA — 70 ms → 0.094 ms, ≈ 744×, compute-bound by the cluster step. Same hops every time (GMEM→SMEM→TMEM→reg→GMEM); each step only changed how a hop happens.

11 · 12Flash Attention 4 does it again, harder. GEMM is one MMA repeated; FA4 is two MMAs with softmax wedged between, on the critical path — so it needs six roles across 512 threads, ~12 barriers in three flavors, a scalar mailbox, and S/P/O sharing one TMEM allocation by fp32/fp16 aliasing.

13And you debug it by hand, because the compiler won't. Four failure classes (deadlock / poisoned context / wrong results / correct-but-slow), each with a symptom→cause→fix; compute-sanitizer + Nsight to see stalls and overlap; then the lowering itself — tile primitives → dispatch + layout application → plain IR → PTX.

The supply-chain model: one picture for the whole paradigm

Here is the synthesis in a single sentence. A fast Blackwell kernel is an asynchronous supply chain:

TMA feeds SMEM — the copy engine streams operand tiles in, applying swizzle on the way.
mbarriers coordinate the independent engines — each handoff (load→MMA, MMA→softmax, MMA→epilogue) is a phase-tracked barrier, because the engines finish on their own clocks.
the tensor core (tcgen05) computes into TMEM — the accumulator lives off-register, so the math tile can be huge without crushing occupancy.
warps specialize into producer (memory) and consumer (math) roles — and in FA4, into softmax and correction roles too — so the engines run at once instead of one team multiplexing all jobs.
clusters pool several SMs — distributed SMEM + multicast + 2-CTA cooperative MMA let one operand load feed more math.
persistent kernels + CLC amortize launch and keep operands hot in L2, with hardware work-stealing to kill the tail.

The payoff of arranging it this way is the whole point of the track: the copy engine, the tensor core, and the epilogue all stay busy at the same time instead of taking turns. The widget below animates exactly that — three engines overlapping across pipeline stages, with the three contracts that must agree drawn as the gates between them.

Widget · the asynchronous supply chain

Three engine lanes — TMA load, tcgen05 compute → TMEM, epilogue store — running across pipeline stages of a persistent GEMM. In the synchronous schedule the engines take turns and idle. Flip on async overlap and the lanes slide under each other: while the tensor core chews stage k, TMA is already loading stage k+1 and the epilogue is draining stage k−1. The three contract gates (▸ layout · ◆ accumulator location · ● completion signal) light at each handoff — they are the agreements that let one engine hand work to the next. Watch the engine-idle KPIs collapse as overlap turns on.

The supply chain · three engines, the contract gates, sync vs async

Lanes: TMA (load GMEM→SMEM) · tcgen05 (MMA→TMEM) · epilogue (TMEM→reg→GMEM). Gates between handoffs: ▸ layout (operand swizzle agrees) · ◆ accumulator (TMEM region agrees) · ● signal (mbarrier phase + byte count agree). Toggle overlap to see the engines stop taking turns.

overlap: async stages: 5

schedule

async overlap

tensor-core idle

—

wall-clock

—

speedup vs sync

—

That is the entire track in one diagram. The synchronous schedule is the gpu_kernels model — correct, and leaving most of the GPU idle. The async schedule is everything lessons 04–10 added. The gates are the next idea.

The contract triad: the invariant that recurred in nearly every lesson

If you keep one thing from this track, keep this. Every tensor-core kernel must hold three contracts in agreement at once. They appeared, separately, in almost every lesson — and the reason the async paradigm is hard is that the hardware does not check them. Get one wrong and the kernel still runs; it just reads the wrong bytes, or reads them slowly.

1 · SMEM operand layout Where the operand tiles sit in shared memory and in what arrangement — including the swizzle mode. The TMA descriptor, the SMEM layout, and the MMA descriptor must all name the same swizzle, or the tensor core reads scrambled bytes. (Lessons 02, 03, 04.)

2 · Accumulator / scale location Where the running sum lives and in what layout: registers on Ampere/Hopper, TMEM on Blackwell — plus where the block-scaling factors sit for MXFP8/NVFP4. The epilogue's readback (tcgen05.ld, lane l → row l/4) has to match this exactly. (Lessons 03, 05.)

3 · Async completion signal How "done" is announced: an mbarrier + byte count for loads and MMA results (expect_tx, tcgen05.commit), a commit-group for stores. The waiter must track the matching phase. (Lessons 04, 05, 06.)

Why this is the spine, not a footnote

Trace it: lesson 02 said “layout is part of the instruction interface.” Lesson 03 said descriptors “don't remove layout work — they make the contract explicit.” Lesson 06's reuse bug is contract 3 violated (phase desync). FA4's three-gate value MMA (lesson 12) is all three at once — V in SMEM (layout), an O slot that's been rescaled (accumulator), and the barriers that gate them (signal). Lesson 13's failure taxonomy is literally “which of the three contracts did you break?” — row-stripe corruption is a signal (phase) bug, NaN is a layout/descriptor bug, finite-but-wrong is a stale tile. Every bug in the async paradigm is one of these three contracts silently disagreeing.

The lens behind every decision: scope · layout · dispatch

One more compression. Behind every primitive and every kernel choice in this track were the same three design elements — the lens introduced in lesson 00. When you read a new kernel or a profile, these are the three questions to ask.

scope

Which threads run an op. One elected thread for tcgen05; a warp for ldmatrix/tcgen05.ld; a warpgroup for wgmma; a cluster for cooperative MMA. Picking the wrong scope (elect_sync where you meant tid==0) is a whole class of lesson-13 bugs.

layout

Where data lives and in what arrangement. The shape:stride algebra, the swizzle mode, registers vs SMEM vs TMEM. This is contracts 1 and 2 of the triad — and the thing CuTe exists to make composable.

dispatch

Which hardware path executes it. Copy via TMA or cp.async; math via tcgen05 or CUDA cores; signal via mbarrier or commit-group. A "correct but slow" kernel is usually a dispatch that silently fell back (no tcgen05 in the generated CUDA).

The roofline payoff, revisited

The GEMM ladder is the proof that the supply chain works. It went 70 ms → 0.094 ms — cuBLAS parity, ≈744× — purely by climbing from memory-bound to compute-bound: TMA removed the copy overhead, warp specialization made the engines overlap, and the cluster step doubled arithmetic intensity until the kernel sat on the compute roof. Nothing about the math changed. The same supply chain is what closes the gap on any tensor-core-dense kernel; FA4 is the proof it generalizes past plain GEMM to a kernel with real CUDA-core work on the critical path.

Where it lives in the real world — what to actually use

Now the honest part. You will rarely hand-write tcgen05 PTX. Everything in this track exists in production libraries, written and tuned by people who do this full-time. The value of having walked the supply chain by hand is not so you can re-derive it — it is so you can read a profile, pick the right library, set tile and stage sizes, and judge when a custom fused kernel is worth writing. Here is the map.

Library / DSL	What it is, in this track's terms	Exposes vs abstracts
CUTLASS / CuTe	The production version of everything here. CuTe's `Layout`/`Swizzle` are the layout algebra of lesson 02; its TMA copy atoms are lesson 04; the warp-specialized collective mainloop is lesson 10; persistent tile schedulers are lessons 07 & 09; the `SM100` / `tcgen05` MMA atoms are lessons 03 & 05.	TMAwarp-specTMEMclusters — exposes all five primitives, as composable C++ templates.
cuBLASLt	The GEMM you'd actually call. It is the lesson 08–10 ladder, shipped and auto-tuned — the 0.094 ms endpoint, behind one API.	everything — you pass dtypes, shapes, and an epilogue; it picks the kernel.
cuDNN	The Flash Attention you'd actually call. The lessons 11–12 choreography, productionized — fused attention with masking, GQA, and scheduling handled.	everything — a fused-attention op; you don't see the 12 barriers.
Triton	The higher-level DSL where most engineers actually write fused kernels. It abstracts most of this paradigm — you write a tile program, it generates the schedule.	most of it, but historically couldn't reach explicit TMAwarp-specTMEM mgmt — newer Triton is starting to expose some. This track's value is partly seeing exactly what it abstracts and where the abstraction historically ended.
ThunderKittens & tile DSLs (and the MLC course's TIRx)	The same niche as Triton: tile-level DSLs that name the supply chain's pieces (tiles, loads, MMAs) and lower them to PTX. TIRx is the pseudocode this track used for the GEMM/FA4 step kernels.	tilessome async — more explicit about the schedule than Triton, less than raw CuTe.

The honest framing

Reach for the library that matches the job: cuBLASLt for a GEMM, cuDNN for attention, Triton for a custom fusion, CUTLASS/CuTe when you need to control the schedule the others won't expose. Drop to raw tcgen05 PTX essentially never — only inside a kernel library, or to understand a Nsight counter. Knowing the supply chain is what lets you tell why Triton left throughput on the table, which tile/stage sizes to hand cuBLASLt's heuristics, and when a hand-written fused kernel beats stitching library calls together. That judgement is the deliverable of this whole track.

What to read next

Each of these is a sibling track that this one points into — one line on why you'd go there.

gpu_kernels — the classic synchronous model this track supersedes: the WMMA-era tensor core, the roofline from first principles, and the serving & training kernels. Start here if the async paradigm felt like it skipped steps.
triton_kernels — the higher-level DSL where you'll actually write fused kernels; now you know what it abstracts and where the abstraction historically reached its edge.
ai_compilers — how kernels get generated: the graph→IR→fusion→codegen pipeline that lesson 13's lowering previewed, from the compiler's side.
cs336 — where these kernels plug into the whole LLM training stack: data, scaling laws, FSDP/TP/PP, alignment — the level above the kernel.
vllm + sglang — serving: PagedAttention, continuous batching, and how the fused attention kernel you just studied gets fed at inference time.

Primary sources, when you want the ground truth: the CUTLASS / CuTe documentation (the real layout algebra and atoms), the FlashAttention papers (FA2, FA3, FA4) for the algorithm lineage, the PTX ISA for the exact semantics of cp.async.bulk.tensor / tcgen05.* / mbarrier.*, and NVIDIA's Hopper and Blackwell tuning guides for the architecture numbers this track quoted.

Takeaway — the whole track, in five sentences

Feeding the tensor core is the whole kernel. The roofline made the classic synchronous loop impossible, so the kernel was rebuilt as an asynchronous supply chain — TMA feeds SMEM, mbarriers coordinate the engines, the tensor core computes into TMEM, warps specialize, clusters pool SMs, and persistent kernels amortize launch — and every primitive you learned was a forced move to remove one more thing that starved the math. Three contracts must always agree — operand layout, accumulator location, completion signal — and every bug in this paradigm is one of them silently disagreeing, because the hardware checks none of them. GEMM (70 ms → 0.094 ms, cuBLAS parity) and Flash Attention 4 proved the supply chain assembles to peak throughput, on a plain matmul and on a kernel with softmax on the critical path alike. And that same supply chain is the foundation every production kernel library — CUTLASS, cuBLASLt, cuDNN, Triton — is built on: you'll call them, not tcgen05 by hand, but knowing the chain is exactly how you read the profile and choose.

Where this connects

This track is the advanced sequel to gpu_kernels (the synchronous model it supersedes) and the sibling of triton_kernels (the DSL that abstracts this paradigm) and ai_compilers (how kernels get generated). Above it sits cs336 (the LLM training stack these kernels power); beside it, the serving tracks vllm and sglang (where the fused attention kernel is fed at inference). The ground truth lives in the CUTLASS/CuTe docs, the FlashAttention papers, the PTX ISA, and NVIDIA's Hopper/Blackwell tuning guides. To start over with fresh eyes, the track index maps all fifteen lessons.