« All posts

New Verifier Exposes Hidden Bugs in LLM-Generated GPU Kernels

A 12-gate contract verifier finds that 39.5% of LLM-generated GPU kernels are broken despite passing standard correctness benchmarks.

Correctness claims for LLM-generated GPU kernels typically rest on a single loose test: run the kernel on a few random inputs at one fixed shape and accept it if the output is close to a reference. That test can miss kernels that return an ordinary number where a NaN or infinity is expected, vary from run to run, break under shape changes, or drift in fp16 while the reference accumulates in fp32.

Researchers built a contract-grade verifier composed of 12 adversarial gates, each checking a property a correct kernel must satisfy, with several gates being tolerance-free so no threshold choice can excuse a failure. Applied to 2,638 machine-generated kernels already accepted by a public benchmark's own test harness, the verifier found 39.5% broken beyond any tolerance argument and 62.1% carrying at least one violation. The standard test accepted 1,487 kernels the verifier rejected, versus only 14 cases going the other way.

The result was validated four independent ways: a 7/7 positive control, a threshold-calibration sweep, 98.5% agreement with the reference benchmark's own correctness code, and a stratified manual audit. The team also turned the verifier on their own work, checking the first native Blackwell tcgen05 training backward pass for the gated-linear-recurrence (GDN) family, including its reverse-state stage. They confirmed its correctness independently against a double-precision oracle and trained five family members using it.