GEMM Steps 1 to 9, Cycle by Cycle
A companion to Part III of Modern GPU Programming for MLSys (MLC Community), which builds a Blackwell GEMM from a single sequential tile up to a kernel that matches cuBLAS in nine steps. The book has the code and the reasoning; this page has the timeline. Pick a step below, read the book's section (each step links to it), then press play.
How to read a trace: one column is one logical cycle, not a clock cycle; lanes are warp roles or hardware units; hatched blocks are a role parked on a barrier; the chips underneath show which barriers fire in that cycle, and the panels show what each SMEM stage and TMEM accumulator holds. The Left and Right keys step, Space plays, and #step-N in the URL opens a given step.
One CTA, one 128 x 128 tile, K=64. The whole GMEM to SMEM to TMEM to registers to GMEM path runs once, in order. Watch the Tensor Core lane: the MMA is asynchronous, and the mma_bar wait at cycle 4 is all that separates it from the threads reading TMEM.
One 128 x 64 fp16 tile each for A and B, 128-byte swizzled, placed from byte offset 1024; the TMEM address and mbarrier sit below.
128 x 128 fp32 inside the 512 columns reserved by tcgen05.alloc. One MMA, so one segment.
One CTA, one warpgroup. The SMEM pool hands out a 4-byte slot for the TMEM address and an 8-byte mbarrier, then jumps to offset 1024 for the A and B tiles. Warp 0 initialises mma_bar with an arrival count of 1 and allocates 512 TMEM columns; the fences and a cta_sync publish both to the whole CTA.