On measuring CUDA kernels without lying to yourself

I had a CUDA kernel that failed numerical validation about six times in thirty runs. In my own notes, I had already explained it: accumulated FP32 rounding error over an unoptimised summation order. It was a plausible sentence. It sat there for a while.

It was also wrong in a way I could have caught immediately, because rounding error is deterministic. The same seeded input, through the same binary, must produce the same output. Mine did not. Every run gave something different. Whatever was happening, accumulation order could not explain it, and one line of reasoning about the shape of the failure would have ruled the explanation out before any tooling was involved.

compute-sanitizer --tool racecheck found it in one command: a genuine shared-memory data race, where a timestep’s write of the hidden state and the next timestep’s read of it were racing across an intervening __syncthreads(). Double-buffering the hidden state so a timestep’s read and write never target the same shared array fixed it. Zero hazards under racecheck, where it had been reporting thousands. 100/100 runs passing at the tolerance it had previously been scraping through at 6/30 — and 20/20 passing at a tolerance three orders of magnitude tighter, which is the part that mattered. It was not squeaking past a loose threshold; it was genuinely close to the reference.

The lesson I actually took is not “run the sanitizer”, though you should. It is that a plausible explanation for a failure is a hypothesis with a testable shape, and I had not asked what the shape was. Nondeterministic output falsifies a deterministic cause. That check is free.

The second problem: your benchmark is less stable than its standard deviation says

The other finding came from re-running the same statistical harness later the same day. Not a different machine, not a different build — the same n=100-trial suite, minutes to hours apart.

Configuration	Session 1	Session 2	Delta
Transposed recurrent weights, no graphs	804 µs	1,023 µs	+27%
Transposed weights + CUDA Graphs	789 µs	905 µs	+15%
FP16 half2 packing	602 µs	548 µs	−9%

Each individual session looked fine from the inside. Within-session coefficient of variation ran 6.8%–24.4%, which for a laptop GPU is not alarming. And yet the means moved by up to 27% between sessions. The full framework benchmark suite showed the same thing at the model level: torch.compile and TensorRT each swung 14–17% between back-to-back runs.

So within-session variance understates the real uncertainty. Thermal state, background load, the WSL2 scheduler — something is drifting on a timescale longer than one benchmark session, and no single run, however tight its own standard deviation, will show it to you.

The consequence is that a single-run latency table is not a measurement. It is one sample from a distribution you have not characterised, presented in a format that implies you have.

What I do now

• Report ranges across independent sessions, not a mean with a standard deviation from one run. If two sessions disagree, that disagreement is the finding — publishing the more flattering one is a choice, and it is not a neutral one.
• Make both sides of a ratio real distributions. I had a period where a kernel’s speedup was computed against a single-run PyTorch reference, and two competing single-run estimates for that baseline (740.7 µs and 943.6 µs) bracketed the true value. A ratio is only as trustworthy as its denominator; the reference now comes from the same style of harness as the numerator.
• Keep the negative results in the table. One hyperparameter configuration in a distillation sweep collapsed on the two minority classes while looking excellent on the majority ones. It stays in the table. A sweep with the failures pruned out is a marketing table.
• Tie every published number to an artifact. Each headline figure in the project traces to a JSON file, and a verifier re-derives it from that file on demand. A second script fails the build if a withdrawn number reappears anywhere in the docs. That tooling exists because I withdrew a number once, and I would rather the repository catch the next one than rely on me remembering.

None of this makes the work faster. It makes the claims survivable, which for anything that ends up in a paper — or in front of someone deciding whether to trust the rest of it — is the part that actually has to hold.