Chapter 7: Matrix Multiplication
Write the kernel that powers deep learning. Move from a memory-drowning naive matmul to a tiled kernel that stages blocks in SRAM, accumulates across the inner-K loop, reorders programs for L2 cache reuse (super-grouping), and approaches cuBLAS-level performance.
- Topics
- 5
- Videos
- 3
Chapter Overview
Matrix multiplication is the beating heart of deep learning — every linear layer, attention projection, and convolution-as-matmul runs through it. It's also the canonical compute-bound kernel, the opposite of softmax: with the right tiling it reuses each loaded value many times, pushing arithmetic intensity high enough to saturate the GPU's ALUs and Tensor Cores.
The journey of this chapter is the journey of GPU optimization in miniature. A naive matmul re-reads A and B from DRAM for every output element, drowning in memory traffic. A tiled matmul loads a block of A and a block of B into SRAM once and reuses them for a whole tile of outputs, accumulating the result across an inner-K loop. That single change moves the kernel from memory-bound to compute-bound. Then super-grouping — reordering which output tiles neighboring programs compute — keeps the reused A/B tiles hot in L2 cache, squeezing out the last big factor and bringing a hand-written Triton matmul within striking distance of cuBLAS.
This is the chapter where everything connects: the memory hierarchy (Ch 2), the tile/program model (Ch 3–4), and the roofline (Ch 5) all converge in one kernel.
This chapter covers:
- Naive vs Tiled Matmul: why the obvious version is slow
- Blocking & SRAM Accumulation: load once, reuse many times
- The Inner-K Loop: accumulating partial products
- L2 Cache Super-Grouping: program order for cache reuse
- Approaching cuBLAS: how close can you get?
Chapter Roadmap
Click any topic to jump in
Naive vs Tiled Matmul
Naive re-reads A and B from DRAM per output; tiling reuses on-chip blocks and raises intensity.
Blocking & SRAM Accumulation
Each program owns a BMxBN tile; a float32 accumulator stays on-chip across the K loop.
The Inner-K Loop
Stream BK-blocks, acc += tl.dot(a, b); num_stages pipelines loads behind compute.
L2 Cache Super-Grouping
Reorder programs so shared A/B blocks stay hot in L2 — fewer DRAM re-reads, no extra FLOPs.
Approaching cuBLAS
Tiling + mixed precision + pipelining + grouping + autotune → within a few % of cuBLAS.
Naive vs Tiled Matmul
Blocking & SRAM Accumulation
The Inner-K Loop
Sign up to unlock this chapter
This chapter is part of PixelBank Premium. Create a free account, then upgrade to read the full lesson — concepts, walkthroughs, and exercises.