Modern GPU Programming

The Hopper & Blackwell async tensor-core paradigm, derived as one chain of forced moves. The engine of the whole track: the tensor core got so fast that feeding it became the entire kernel. A B200 sustains ≈2 PFLOP/s of fp16 but HBM delivers only ≈8 TB/s — so the classic synchronous "load, sync, compute" loop leaves the tensor core idle most of the clock, and every modern feature (TMA, TMEM, mbarriers, warp specialization, clusters) exists to remove one more thing that starves it. We build two kernels that assemble all of it: a GEMM that climbs from 70 ms to cuBLAS parity, and Flash Attention 4.

How to read this track

Strictly linear. Part I (00–02) establishes the forcing function: the roofline, why occupancy stopped feeding the tensor core, and the layout algebra everything else depends on. Part II (03–07) derives the five hardware primitives one at a time — each is the forced answer to a bottleneck the previous one exposed: tensor-core generations → TMA → TMEM → mbarriers → clusters/CLC. Part III (08–10) assembles them into a GEMM, growing it from the smallest correct tile to a warp-specialized, clustered, cuBLAS-parity kernel. Part IV (11–14) does it again for Flash Attention 4 (two MMAs with a softmax wedged between), then teaches how to debug this class of kernel and where the whole paradigm lives in the real toolchain.

Mental stance

This is the advanced sequel to gpu_kernels. That track teaches the classic synchronous SIMT model (load tile → __syncthreads() → compute) and the WMMA-era tensor core; this track is what happens when that model breaks down on Hopper/Blackwell and the kernel must be rebuilt as an asynchronous supply chain. Assume you know warps, shared memory, coalescing, and the roofline. Code is framed vendor-real — CUTLASS/CuTe abstractions and PTX (cp.async.bulk.tensor, wgmma, tcgen05.mma, mbarrier.*) — because that is the toolchain you will actually meet.

The dependency chain

Not a wheel — a directed line. Each primitive is forced by a bottleneck the previous step could not remove, and the two capstones simply assemble the line. If you skip a step, the next one reads like a folk trick instead of the only way out of the wall.

Part I · the forcing function

Why a Blackwell tensor core can't be fed by the classic loop, how to measure it, and the layout algebra every later lesson depends on.

Feeding the tensor core

The engine of the whole track. The B200 roofline (≈2 PFLOP/s vs ≈8 TB/s → ridge ≈250 FLOP/byte), why the synchronous tiled GEMM idles the tensor core, and the supply-chain model that fixes it.

modern_gpu/lessons/00_feeding_the_tensor_core.html

What makes a kernel fast now

Roofline, arithmetic intensity by workload, and the shift from "raise occupancy to hide latency" to explicit async overlap. Modern tensor kernels run at lower occupancy on purpose.

modern_gpu/lessons/01_what_makes_a_kernel_fast.html

Data layout and its algebra

Shape·stride layouts (CuTe), nested tiling, named hardware axes, and the XOR swizzle that kills bank conflicts. "Layout is part of the instruction interface."

modern_gpu/lessons/02_data_layout.html

Part II · the five primitives

Each one is the forced answer to a bottleneck the previous left behind — the supply chain, built component by component.

Tensor cores across generations

mma → wgmma → tcgen05. The operand-layout contract, and why operands and the accumulator kept migrating out of the register file.

modern_gpu/lessons/03_tensor_cores_generations.html

TMA — the async bulk copy engine

A hardware copy engine: one thread issues a tensor-map descriptor, the engine moves a swizzled tile async. Loads complete via mbarrier + byte count; stores via commit groups.

modern_gpu/lessons/04_tma.html

TMEM — the accumulator's new home

Blackwell's per-CTA 128×512 tensor-memory scratchpad. Why the accumulator left the register file, how it's allocated in 32-column units, and tcgen05.ld/.st/.cp.

modern_gpu/lessons/05_tmem.html

mbarriers and the phase model

The correctness centerpiece. Arrival counter + phase bit, expect_tx byte budgets, the reuse bug, and one barrier per pipeline stage — what __syncthreads() can't do.

modern_gpu/lessons/06_mbarriers.html

Clusters, DSMEM & cluster launch control

CTAs cooperating across SMs: distributed shared memory, 2-CTA cooperative MMA, TMA multicast, and CLC work-stealing to kill the static-schedule tail.

modern_gpu/lessons/07_clusters_clc.html

Part III · GEMM, tiled to SOTA

Assemble the five primitives into one kernel, grown one forced step at a time from 70 ms to cuBLAS parity.

GEMM I — the tiled baseline

The memory model (GMEM→SMEM→TMEM→reg→GMEM), the smallest correct tile, K-loop accumulation with the phase-flip table, and the multi-CTA grid that plants the operand-reuse debt.

modern_gpu/lessons/08_gemm_tiled.html

GEMM II — pipelining with TMA

Hand the load to TMA (tid==0, not elect_sync), double-buffer into a software pipeline (prefetch ≠ overlap), then go persistent with a tile scheduler that pays off the L2 reuse debt.

modern_gpu/lessons/09_gemm_pipelined.html

GEMM III — warp specialization & clusters

The climax. Split the SM into producer/consumer warps (four barriers, opposite initial phases), add a 2-CTA cluster, then a second consumer — 70 ms → 0.094 ms = cuBLAS parity.

modern_gpu/lessons/10_gemm_warp_specialized.html

Part IV · the second capstone & the craft

Flash Attention 4 reuses every primitive, then raises the difficulty; then how to debug this class of kernel and where the paradigm lives.

Flash Attention 4 — the algorithm

Online softmax (can't materialize the 64 MB score matrix), the exp2 trick, and why FA4 is harder than GEMM: two MMAs with a softmax wedged between, sharing one TMEM allocation by fp32/fp16 aliasing.

modern_gpu/lessons/11_flash_attention_algorithm.html

Flash Attention 4 — warp-role choreography

Five warp roles (one warp drives both MMAs), a 12-barrier graph plus a scalar mailbox, the online-softmax correction warp, causal masking, GQA, and LPT tile scheduling.

modern_gpu/lessons/12_flash_attention_choreography.html

Debugging & how the compiler lowers it

No pass verifies your phases. The roles/storage/handoff/lifetime worksheet, four failure classes (deadlock, poisoned context, wrong results, slow), compute-sanitizer & Nsight, and the lowering pipeline.

modern_gpu/lessons/13_debugging_lowering.html

Synthesis — the supply chain

The whole paradigm on one page: the contract triad that must always agree, the scope/layout/dispatch lens, and where it lives in the real world — CUTLASS/CuTe, cuBLASLt, cuDNN, Triton.

modern_gpu/lessons/14_synthesis.html

Where this sits in the library

Advanced sequel to gpu_kernels (the classic synchronous model) and triton_kernels (the higher-level DSL). The "how SOTA kernels are built" counterpart to ai_compilers ("how kernels are generated"). The kernels here are what cs336 plugs into the LLM stack, and what vllm / sglang call at serving time. Distilled from MLC.ai's Modern GPU Programming for MLSys, reframed in CUTLASS/CuTe + PTX terms.