Matmul is not born compute-bound. It is made compute-bound.

Lesson 6 · tiling and reuse · where arithmetic intensity actually comes from

Everyone repeats that matmul is the compute-bound op. Written naively it is nothing of the sort — it runs at 0.3% of an A100's peak. The gap between 0.3% and 84% is one idea, and it is the same idea that makes BLOCK_SIZE worth autotuning.

The arithmetic that should worry you

Take C[M,N] = A[M,K] @ B[K,N]. The work is fixed: 2·M·N·K FLOPs. Now count the traffic if every output element fetches its own inputs from HBM:

Naive matmul: no reuse, so intensity collapses to a constant.

Listen first: each of the M times N output elements reads a row of A of length K and a column of B of length K, so the bytes moved are two times M times N times K times bytes per element. Divide the FLOP count by that and almost everything cancels — you are left with one over bytes per element. In fp16 that is one half of a FLOP per byte, against a machine that wants a hundred and fifty-three. Naive matmul is not slightly memory-bound; it is three hundred times short.

FLOPs  = 2·M·N·K
bytes  = 2·M·N·K · bytes_per_element        # every output fetches its own row and column

intensity = FLOPs / bytes = 1 / bytes_per_element
          = 0.5 FLOP/byte in fp16           vs machine balance 153   ← 306x short

The M, N and K all cancel. Problem size does not save you. This is the same kind of cancellation you saw in Lesson 2 when the parameter count vanished out of decode intensity, and it means the same thing: the ratio is structural, and only a change of structure fixes it.

The structural change: compute a tile, not an element

Instead of one output element at a time, one program computes a BM × BN tile of C. It loads a strip of A and a strip of B into fast on-chip memory once, then every loaded element gets used across the whole tile:

Tiled matmul: each loaded element serves an entire row or column of the output tile.

Listen first: when you compute a tile of BM by BN outputs, each element of the A strip you loaded is used BN times — once for every column of the tile — and each element of the B strip is used BM times. The bytes drop by that reuse factor, so intensity becomes two divided by the quantity one over BM plus one over BN, all over bytes per element. For a square tile of side T in fp16 that simplifies to T over two. A 128-wide tile gives 64 FLOPs per byte; a 256-wide tile gives 128. The tile size is not a tuning detail — it is the term that sets your arithmetic intensity.

intensity = 2 / ((1/BM + 1/BN) · bytes_per_element)

square tile, side T, fp16:   intensity = T / 2

T =  16  →    8 FLOP/byte      (still hopeless)
T = 128  →   64 FLOP/byte      (now you are in the game)
T = 256  →  128 FLOP/byte      (near the 153 ridge)

That is the entire trick. Reuse is the only thing that creates arithmetic intensity, and a tile is how you buy reuse. Elementwise ops like vector add can never do this — there is nothing to reuse, which is why they sit permanently on the memory-bound side.

Two honest caveats

1. This model counts HBM traffic only. It assumes every re-read of an A-strip goes all the way to DRAM. In reality the L2 cache catches many of them, so a real kernel lands between the tiled line and the "perfect reuse" line below. Treat the tiled figure as a pessimistic bound and the ideal figure as an unreachable one.

2. Perfect reuse would read each matrix exactly once — (M·K + K·N + M·N) elements total. For a square 4096 matmul in fp16 that is an intensity of about 1365 FLOP/byte: massively compute-bound. That number is why the folklore says "matmul is compute-bound." The folklore is describing a well-tiled matmul and forgetting to say so.

What Triton does, and what it leaves to you

Lesson 1 listed shared-memory allocation and synchronization among the things the compiler handles. This is where that matters: you write a tile, and Triton stages it through on-chip memory, picks tensor-core instructions, and schedules the async copies. What it does not do is pick BLOCK_M, BLOCK_N and BLOCK_K for you — those are tl.constexpr parameters, and choosing them is choosing your intensity. That is what triton.autotune is searching over, and now you know what it is searching for.

Micro-world: buy intensity with tile size

Square matmul on one A100 80GB

Three strategies, same FLOPs, same answer. Only the traffic differs — and traffic is the whole story.

Matrix size n (n×n @ n×n)

Tile size T (T×T of C)

Element type

Arithmetic intensity, log scale — where each strategy lands against the 153 ridge

naive (no reuse) tiled T×T perfect reuse (each matrix read once)

Table view

Look at the naive row first. 0.5 FLOP/byte, about 1 TFLOP/s — 0.3% of the card. Change n to anything you like; it does not move. Problem size cannot rescue a bad ratio.
Walk the tile size up: 16 → 32 → 64 → 128 → 256. Intensity doubles each step and so does the achieved throughput. This is the single highest-leverage knob in a matmul kernel.
Find where tiled crosses the ridge. In fp16 it takes T = 256 to reach 128 F/B and still not quite clear 153. Now switch to float32 — the same tile buys you half the intensity, because intensity is per byte, not per element.
Compare the tiled and perfect-reuse rows. That gap is what the L2 cache is for. A real kernel lives in between, which is why measured matmul beats this pessimistic model — and why you should never quote a roofline number as a prediction.
Set T = 16 and n = 8192. A huge, expensive matmul running at a few percent of peak. This is exactly what an untuned BLOCK_SIZE looks like, and why triton.autotune exists.

Check yourself

Cold-recall defense (one breath):
Matmul FLOPs are fixed at 2·M·N·K; only the traffic is negotiable. Naive, every output fetches its own row and column, so M, N, K cancel and intensity is just 1/bytes_per_element — 0.5 F/B in fp16, hopeless. A T×T output tile makes each loaded element serve T outputs, giving intensity = 2/((1/BM + 1/BN)·bytes) = T/2 in fp16. Reuse is the only source of arithmetic intensity, which is why elementwise ops can never escape the memory-bound side. Triton stages the tile through shared memory for you, but BLOCK_M/N/K are yours — picking them is picking your intensity.

Primary source

Read the Triton tutorial — Matrix Multiplication (~20 min). Read it with one question in mind: where does BLOCK_SIZE_M/N/K appear, and what does each one do to the reuse ratio? The autotune configs at the top of that file are a list of intensity choices, and you can now compute what each one is worth before running it.

💬 I'm your teacher for this — ask me followups any time. The question this lesson sets up: "if bigger tiles are always better, why not T = 1024?" The answer is shared-memory capacity and occupancy, and it is a good next lesson whenever you want it.