Most kernel speedups are measurement bugs

Lesson 9 · warmup, synchronization, L2, and quantiles · how a benchmark becomes evidence

This is the lesson that makes the other eight count. Your mission says lessons teach but proof comes from measurement — and a measurement you cannot defend is not proof, it is a story. Every trap below has produced a confident, wrong, publicly-posted speedup number.

Trap 1: the GPU has not finished

CUDA launches are asynchronous. Your Python call queues the kernel and returns immediately. If you wrap it in time.time() you are timing the launch, not the work.

The wrong way and the right way to put a clock around a GPU kernel.

Listen first: the broken version records a start time, calls the kernel, and records an end time. Because the launch returns before the GPU has done anything, the elapsed time is microseconds of queueing, and you conclude your kernel is a hundred times faster than it is. The correct version calls torch.cuda.synchronize before the first timestamp — to drain anything already queued — and again before the second, which blocks until the GPU has actually finished. Only then does the elapsed time describe the kernel.

# WRONG — measures the launch, not the work
t0 = time.time();  my_kernel[grid](x, y, out);  t1 = time.time()
print(t1 - t0)                          # ~8 microseconds. It is not 8 microseconds.

# RIGHT
torch.cuda.synchronize();  t0 = time.time()
my_kernel[grid](x, y, out)
torch.cuda.synchronize();  t1 = time.time()

Trap 2: the first call is not the kernel

Triton is a JIT. The first invocation of a @triton.jit function compiles it, and if you used @triton.autotune it also benchmarks every config in your search space. That first call can be hundreds of milliseconds against a kernel that runs in one. Discard it — and discard a few more, because the caching allocator, clocks and cuDNN heuristics all settle over the first several calls.

Trap 3: you are measuring the L2 cache

This is the subtle one, and it is the one that inflates bandwidth claims specifically. Run the same kernel on the same buffers in a tight loop and after the first iteration the input is already sitting in L2 — 40 MB of it on an A100. Your kernel is now reading from cache at several times HBM bandwidth, and you report a number the memory system cannot actually sustain.

This matters enormously for everything in this lab, because per Lesson 2 your kernels are memory-bound — so a benchmark that accidentally removes the memory bottleneck is measuring nothing you care about. triton.testing.do_bench writes over a cache-sized buffer between repetitions for exactly this reason.

Trap 4: the mean of a skewed distribution

GPU timings are noisy and right-skewed — an OS hiccup or a clock dip adds time, nothing ever subtracts it. One 3 ms outlier in a hundred 1 ms runs drags the mean but not the median. Report the median with quantiles, which is what do_bench's quantiles argument gives you. Minimum is not the answer either: it is the best case you ever saw under ideal cache and clock conditions, and production never runs there.

The tool that does all four for you

The standard harness, with the four traps handled.

Listen first: triton.testing.do_bench takes a zero-argument callable. Warmup and rep are durations in milliseconds — it estimates how many iterations fit in each window rather than making you pick a count. Passing a list of quantiles returns the median and the bounds you asked for instead of a single number. It synchronises, discards the warmup window, and clears the L2 cache between repetitions by default. The line above it is the one people skip: assert that your kernel agrees with the reference before you time it, because a wrong kernel can be arbitrarily fast.

import triton, torch

torch.testing.assert_close(mine(x), reference(x), atol=1e-2, rtol=0)   # CORRECTNESS FIRST

ms, lo, hi = triton.testing.do_bench(
    lambda: mine(x),
    warmup=25, rep=100,              # DURATIONS in ms, not iteration counts
    quantiles=[0.5, 0.2, 0.8],       # median, and the spread around it
)
gbps = total_bytes_moved / (ms * 1e-3) / 1e9     # memory-bound kernel → report GB/s

And always against a real baseline. "2.1× faster than PyTorch eager on an A100 at this shape, median of 100 reps" is evidence. "Very fast" is not, and neither is a speedup over a naive implementation you wrote yourself to lose.

Micro-world: measure a kernel wrong, then right

A kernel whose true runtime is exactly 1.000 ms

You know the truth here — that is the point. Turn the safeguards off one at a time and watch what you would have reported. The PyTorch baseline for this shape is 2.000 ms, so the honest speedup is 2.0×.

cuda.synchronize()

Flush L2 between reps

Warmup calls discarded: 25

Repetitions: 30

Statistic reported

Per-repetition timings (after warmup)

Table view

Everything on: 25 warmup, 30 reps, L2 flushed, median. Reported ≈ 1.00 ms against a truth of 1.00 ms, and an honest 2.0× over the baseline. This is the only configuration you may publish.
Untick synchronize. You now report microseconds and a speedup in the hundreds. This is the single most common way a GPU benchmark lies, and it lies in your favour, which is why it survives review.
Sync back on; set warmup to 0 and reps to 1. Your single measurement is the JIT compile: 250× too slow. Measurement error is not always flattering.
Now raise reps back to 30, warmup still 0. The reported number returns to ~1 ms — the median quietly absorbed the 250 ms first call. Then switch to mean and watch it explode. Two safeguards covering for each other, which is exactly why you want both rather than either.
Untick "flush L2". A quiet ~2.4× overstatement: the input never leaves cache, so you are benchmarking L2 rather than HBM. On a memory-bound kernel that invalidates the entire result.
Switch median → mean with 30 reps. One outlier rep drags the mean noticeably while the median ignores it. Then try minimum, and notice it reports a time the kernel only achieves on its luckiest run.

Check yourself

Cold-recall defense (one breath):
Four traps. Async — CUDA launches return immediately, so synchronize() on both sides or you time the queue. Warmup — Triton JIT-compiles and autotunes on the first call; discard the early ones. L2 — repeat the same buffers and you benchmark a 40 MB cache instead of HBM, which destroys any memory-bound result. Statistics — timings are right-skewed, so report the median with quantiles, never the mean and never the min. triton.testing.do_bench does all four (warmup/rep are durations in ms). And check correctness before you time anything, against a real baseline — a wrong kernel is arbitrarily fast.

The reporting checklist

A benchmark in this lab is only evidence if it states all of these. Copy this into every experiment README:

Field	Why it is required
GPU model and clocks	A T4 result does not transfer to an A100. Boost vs sustained clocks differ.
Exact shapes and dtypes	Kernel performance is shape-dependent; one shape is an anecdote.
Baseline, named and versioned	"Faster than PyTorch 2.x eager" — not faster than nothing.
Correctness check with tolerances	Stated `atol`/`rtol`, run before timing.
Warmup and rep settings	So the reader knows the traps above were handled.
Median + spread, not a single number	A point estimate hides the distribution.
GB/s for memory-bound, TFLOP/s for compute-bound	Per Lesson 2 — quoting the wrong one hides the real limit.
What you did not test	Honesty about scope is what makes the rest believable.

Primary source

Read the source of triton.testing.do_bench (~50 lines, 10 min). It is short, and reading it is worth more than any blog post: you can see the cache flush, the synchronize calls, the warmup estimation and the quantile handling in code. Every one of this lesson's four traps is a line in that file.

💬 I'm your teacher for this — ask me followups any time. When Experiment 0001 produces numbers, bring them here and we will audit them against the checklist together — that is the highest-value thing this lesson can do for you.