A new audit of machine-generated GPU code suggests the field has been grading its own homework with a broken answer key. Researchers built a verification instrument that re-examined 2,638 GPU kernels produced by language models, all of which had already been accepted as correct by the generating system's own test harness. The verifier judged 39.5% of them broken in ways no tolerance setting could excuse, and found at least one contract violation in 62.1%.
The paper, posted to arXiv as preprint 2608.12700 by Rishi Shah and a co-author, argues the problem is not that models write bad code but that the standard test is too permissive to notice. Current practice runs a candidate kernel on a handful of random inputs at one fixed tensor shape, compares the output against a reference implementation, and accepts anything numerically close.
How a Kernel Passes and Still Fails
That single check leaves several failure modes invisible. A kernel can return an ordinary floating-point number where the correct answer is NaN or infinity. It can produce different results on repeated runs of identical input. It can work at the tested shape and break at every other one. It can quietly accumulate in fp16 while the reference maintains an fp32 running total.
The most common defect the audit found was non-finite non-propagation, the silent substitution of a normal number for a NaN or an infinity. The authors stress this is a correctness failure rather than a rounding artifact, since no threshold choice makes a suppressed NaN acceptable.
Their answer is a set of twelve adversarial gates, each encoding a property any correct kernel must satisfy. Several gates are tolerance-free by design, which removes the usual escape hatch of blaming a failure on a badly chosen threshold.
Defending the Number
A 39.5% failure rate is a large claim, and the paper anticipates the pushback with four independent defenses. The verifier was run against a kernel the authors could vouch for and cleared it 7 out of 7. A threshold-calibration sweep tested whether the result was an artifact of tuning. The tool agreed with the reference benchmark's own correctness code 98.5% of the time. A stratified hand-audit then examined disputed cases individually.
The disagreement between the two testing regimes is lopsided. The field's standard check accepts 1,487 kernels the verifier rejects, while only 14 go the other way. The loose test is not merely noisier, it is systematically permissive in one direction.
Turning the Instrument Inward
The team also pointed the verifier at their own work: what they describe as the first native Blackwell tcgen05 training backward pass for the gated-linear-recurrence family, covering the reverse-state stage that the field still runs on a fallback path. They validated it against a double-precision oracle and trained five family members through it.
The work lands alongside related efforts including KernelBench-Verified and a paper titled The Correctness Illusion in LLM-Generated GPU Kernels, pointing to a growing recognition that kernel generation benchmarks need stricter instrumentation.
Why It Matters
AI-assisted kernel authoring has become a real efficiency lever for teams training large models, where hand-tuned CUDA expertise remains scarce and expensive. If nearly two in five accepted kernels carry defects, reported progress in the area reflects the weakness of the measuring stick as much as the strength of the models.
The authors' conclusion is pointed but constructive. The correctness signal behind headline results in kernel generation is far weaker than the numbers imply, and adopting a set of tolerance-free contracts would close most of that gap without waiting for better models.
