GPU roofline

Two ceilings, measured

A roofline has two inputs, and both are measurements on your device:

ms_copy = triton.testing.do_bench(lambda: y.copy_(x))
gbps = 2 * n * 4 * 1e-9 / (ms_copy * 1e-3)
ms_mm = triton.testing.do_bench(lambda: a @ b)
tflops = 2 * M * N * K * 1e-12 / (ms_mm * 1e-3)

do_bench returns milliseconds, so the 1e-3 is not optional. Record the dtype, the clock state and triton.__version__ beside both figures. A copy kernel is the right bandwidth probe because it has no arithmetic to hide behind; a matmul is the right compute probe because it is the densest thing you will run.

The ridge point

Arithmetic intensity is flops per byte moved, and the model says the reached rate is min(peak_flops, bandwidth * AI). The two lines meet at ridge = peak_flops / peak_bandwidth, in FLOP/byte. Below the ridge a kernel is bandwidth-limited; above it, compute-limited. For a square f32 matmul the ideal intensity is N/6 FLOP/byte, because 2N^3 flops move 12N^2 bytes. That is algebra, not a measurement, and it says a 512 matmul is far hungrier per flop than a 4096 one, which is why tile size changes the classification.

Classify before you optimise

Compute intensity from the traffic the kernel actually causes. A vector add reads two streams and writes one for a single flop, so it sits far left of any GPU ridge and no unrolling makes it compute-bound. Softmax and attention live in the same region until fusion and tiling raise reuse. The classification decides the next move: bandwidth-limited means delete bytes, compute-limited means feed the tensor cores.

Things that break

The plot is drawn with spec-sheet peaks, so every kernel lands at a fraction of a ceiling the part cannot sustain at your clock and the picture says nothing. The other failure is ideal traffic: a kernel with uncoalesced loads moves several times the bytes its shape suggests, so its true intensity is lower and it plots right of where it belongs. Measure dram__bytes.sum in ncu, divide by elapsed time, and check that the dot you drew is the dot you ran.