A Contract-Grade Verifier for LLM-Generated GPU Kernels, and a Native Blackwell Backward for the Gated-Linear-Recurrence Family

The paper addresses a fundamental weakness in how GPU kernel generation systems verify correctness. The authors state: "Systems that generate GPU kernels with language models report high correctness rates. Those rates come from 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. A kernel can pass that test and still be silently wrong. Specifically, a kernel can return an ordinary number where the true answer is a NaN or an infinity, produce a different result on each run, break the moment the shape changes, or accumulate in fp16 where the reference keeps an fp32 total."

The authors build a contract-grade verifier of twelve adversarial gates, each a property a correct kernel must satisfy, several of them tolerance-free so that no choice of threshold can explain a failure away. The twelve gates are:

• CMP-01: correct on many random and adversarial inputs (zeros, 106, 10−6, denormals, long L)

• CMP-02: the gradients are correct, not only the outputs (autograd versus finite difference)

• CMP-03: correct across shapes (batch, length, width), not just the one tested

• ORD-01: reordered summation stays within a derived ∝ √N rounding bound

• ORD-02: byte-for-byte identical across five repeats, and does not alias a shared buffer

• ORD-03: correct on an input constructed to expose a bad summation order

• PRC-01: correct in fp32, fp16, and bf16

• PRC-02: fed fp16, still keeps an internal fp32 running total

• EXC-01: infinities and NaNs land in exactly the same positions and signs as the reference

• EXC-02: flush-to-zero handling of subnormals matches the reference

• RES-01: output lives on the same device as the input

• RES-02: the compiled kernel fits real hardware limits (registers, shared memory, TMEM budget)

Two design choices make the battery rigorous: every tolerance is derived rather than chosen (e.g., ORD-01's bound follows how floating-point error actually accumulates over a reduction, atol ≈ 4ε√N·scale), and each gate re-seeds the random number generator from a fixed value before drawing inputs, so a verdict depends on the kernel and not on luck.

The verifier was applied to 2,638 kernels from the Dr. Kernel / KernelGYM corpus that the source system's own harness had already accepted as correct. The headline finding: 62.1% carry at least one contract violation, and 39.5% (1,043 of 2,638) fail a tolerance-free gate, meaning they are broken in a way no tolerance argument can excuse.

The modal defect is EXC-01 (non-finite non-propagation) at 34.2% failure rate: a kernel silently replacing a NaN or infinity with an ordinary number... the failure mode that converts a training-time error into silent data corruption. Other per-gate failure rates include: CMP-01 at 23.6%, CMP-03 at 18.1%, PRC-02 at 13.8%, ORD-01 at 13.7%, PRC-01 at 11.3%, EXC-02 at 9.3%, ORD-03 at 5.3%, ORD-02 at 4.9%, and RES-01 at 1.8%.

The authors answer the objection that their checker is simply stricter than everyone else's with four falsifiable defenses:

1. Positive control (7/7): Their own six Mamba-3 Triton kernels and the native GDN backward all pass every applicable gate. The native GDN backward played no part in calibration, making it a control the thresholds were never fitted to. The control caught a genuine gap: C5 (fused block forward) inferred its channel dimension from the input tensor without checking it against the convolution-weight channels — fixed with a one-line guard.

2. Threshold calibration: For each band-gate, the threshold sits inside the safe margin between correct-kernel noise and wrong-kernel error. Margins range from 57× to 76× for most gates, with ORD-03 disclosed as thin at 1.2× (deliberately relaxed, coverage carried by CMP-01 and ORD-01).

3. Benchmark's own harness agrees (98.5%): Running KernelBench's own correctness code at the pinned commit on 1,030 pairs gives 844 pass/pass, 171 fail/fail, and only 15 disagreements (7 where their replica is stricter, 8 seed-borderline flips).

4. Stratified hand-audit: Hand-tracing 31 disputed cases classified 16 as genuinely broken, 8 as real but tolerance-dependent, and 7 as out of scope (discarded from the strict floor). That the audit discards 7 of the 31 as out of scope, rather than confirming all of them, is evidence that it does not rubber-stamp its own conclusion.

Running the same kernels through KernelBench's standard paper-era check (allclose atol=rtol=10−2, five random-input trials, fixed shapes): the benchmark accepts 93.7% (2,472 of 2,638). The load-bearing cell (external PASS, our FAIL) is 1,487 kernels (56.4%), of which 958 fail on a tolerance-free gate. The reverse cell (external FAIL, our PASS) is only 14 (0.5%). The near-unidirectional disagreement is the signature of a systematic blind spot in the acceptance signal, not of two checks tuned to different strictness. Under KernelBench's hardened per-dtype variant (fp32 tolerance 10−4), 1,263 kernels still pass their check and fail ours.

The finding reproduces on a different software stack: a 300-kernel re-audit on torch 2.11.0+cu128 returns 68.6% (above the 62.1% headline due to over-weighting of the high-rate reduction class), clearing the pre-registered kill criterion of 5% by more than thirteen times. A second corpus of native CUDA kernels (Sakana AI CUDA Engineer archive, 213 kernels) shows a related but weaker pattern: environment-robust residual of 20.2% (43 of 213), driven by CMP-03 shape rigidity (37), EXC-02 subnormal handling (9), ORD-02 nondeterminism (4), and CMP-01 value (1).

The second contribution is a hand-written native Blackwell tcgen05/tensor-memory training backward for the gated-linear-recurrence family. The general recurrence is:

S t = (I − k t(b t ⊙ k t)T) Diag(e g t) S t−1 + k t(w t ⊙ v t)T, o t = S tT q t

This is a superset: fixing gates recovers LA, GLA, SSD/Mamba-2, KDA, and GDN (gated DeltaNet). One backward differentiating (1) produces the training gradients for all five.

The backward has two hard stages:

• K#1, the reverse inter-chunk state scan: "A reverse-time walk accumulating the state gradient dS (shape d k × d v, fp32) backward across chunks, with a rank-one correction at each step from the delta rule. This is the sequential, carried stage, and it is the exact piece that open libraries still run on a Triton fallback."

• K#2, the WY / triangular-inverse VJP: "Inside each chunk the delta rule requires a small triangular solve T = (I + M)−1, and its vector-Jacobian product is the numerically nastiest piece. The forward already computed T, so the backward reuses it as two triangular matrix multiplies and never re-inverts."

The first attempt produced illegal machine code or deadlock — the same failure class behind open issue state-spaces/mamba#904. The root cause... was a lifecycle error: the kernel reserved the full 512-column TMEM budget and released it per matrix multiply, which the hardware forbids. The fix: Porting that lifecycle, offset-partitioned accumulators with alloc-once and relinquish-once, cleared the blocker.

• Double-precision spec pins on the core tiles: about 7 × 10−16 against the closed-form spec (machine-exact)

• Full assembled backward against independent fp64 oracle: worst relative error 3.29 × 10−3 (scalar) and 3.31 × 10−3 (channel-wise), both under the 5 × 10−3 acceptance bound

• Bit-for-bit deterministic across runs

• One disclosed exception: the widest d v=128 save-forward arm reaches 5.21 × 10−3 (a 4% overrun), a property of the assembled pipeline arm, not the kernel (which measures 5.5 × 10−4)

• Trains five family members through real 300-step training loops with zero numerical blow-ups

• Passes the full twelve-gate battery, clean on every tolerance-free channel

The native GDN backward is slower than the fla Triton library: by roughly 8× at L=512 rising to about 78× at L=2048; we do not claim parity. The gap is structural: "fla sits near a 0.9 ms latency floor because it reuses the delta-rule inverse saved in the forward pass, while our reverse-state scan is inherently sequential and most of our runtime is spent in the surrounding fp32 glue rather than in the tensor-core kernels." Speedups are reported over their own earlier pipeline: a channel-wise fusion campaign cut the captured backward 2.75× (from 52.9 to 19.2 ms), and a tensor-memory tiling optimization cut the d v=128 save-forward variant 2.98× (from 24.50 to 8.23 ms).

The verifier's two uses are joined by shared data: The defects that most often sink foreign kernels are, gate for gate, detected by the same gates that caught real errors in our own kernel during development. Three substantive alignments: RES-02 (TMEM constraint, the #904 class), CMP-03 (shape rigidity, caught the C5 channel-inference defect), and EXC-01 (non-finite propagation — our kernels propagated theirs, and it was precisely that propagation which made the errors visible... Non-finite propagation is not merely a gate to pass; it is the property that makes a kernel's own failures catchable).

• Six contract-verified Mamba-3 Triton kernels with no tl.dot anywhere, which compile at num warps 2/4/8 on Blackwell where the official kernel fails

• Faithful reproduction of #904 on B200: A batch-2 / L=2048 run and a batch-4 / L=4096 run both trip Required: 544, Hardware limit: 512

• Applied demonstration: a 1.10B-parameter Mamba-3 SISO model trained as a 12-lead physiological-signal classifier on PTB-XL, reaching macro-AUC of 0.880 with zero NaN

• RL autotuner: config-RL beats the shipped default by 1.167× at long sequences; edit-RL is a negative result (no edit beat baseline, maximum speedup 0.183×, 65% of edits did not apply)

• Speed baselines: against the #904-crippled official backward they are about 4× slower (a structural cost of avoiding tl.dot); against healthy Mamba-1 SSD CUDA they range from 0.21 to 1.04× on forward and 0.13 to 0.75× on backward

The authors state: The correctness behind the field's reported progress is far weaker than its numbers suggest, and the same twelve contracts that expose the gap are enough to begin closing it. They argue that a small set of tolerance-free contracts (non-finite propagation, determinism, shape polymorphism) would close most of it at modest cost; we suggest such contracts as a benchmark standard.

Improvement 1: Contract-Grade Verification for AI-Generated Code

• What I can do: When generating GPU kernels, CUDA code, or any numerical code, I can automatically apply a twelve-gate adversarial verification battery (non-finite propagation, determinism, shape polymorphism, precision handling, hardware-resource compliance) before declaring output correct. I will not rely on a single loose allclose check.

• Improved AI system capability: The system can generate kernel code that is verified to be correct across shapes, precisions, and edge cases (NaN/Inf, denormals, zeros, large magnitudes), with derived rather than arbitrary tolerances, and deterministic byte-for-byte behavior. It will reject kernels that silently corrupt data (e.g., replacing NaN with a number) or break under shape change.

Improvement 2: Tolerance-Free Acceptance Criteria as Default

• What I can do: I will treat tolerance-free properties (non-finite propagation, determinism, shape polymorphism) as mandatory gates, not optional checks. I will not accept a kernel that passes a numeric comparison but fails on these.

• Improved AI system capability: The system can guarantee that generated kernels propagate NaN/Inf exactly as the reference, produce identical outputs across repeated runs, and work across all valid input shapes—eliminating the 39.5% failure class found in the paper.

Improvement 3: Shape-Polymorphic Generation and Testing

• What I can do: When generating kernels, I will test across multiple shapes (batch, length, width) and enforce that the kernel's behavior is shape-invariant in correctness, not just at one fixed shape. I will also check that inferred dimensions (e.g., channel counts) are validated against explicit parameters.

• Improved AI system capability: The system can produce kernels that are robust to shape changes at inference time, catching defects like the C5 channel-inference bug before deployment.

Improvement 4: Precision-Aware Accumulation Enforcement

• What I can do: I will enforce that reductions accumulate in a higher-precision type (e.g., fp32 for fp16 inputs) unless explicitly overridden, and verify this via a dedicated gate (PRC-02). I will also check correctness across fp32/fp16/bf16 with derived rounding bounds.

• Improved AI system capability: The system can generate kernels that maintain numerical accuracy across mixed-precision settings, avoiding silent precision loss in training loops.

Improvement 5: Non-Finite Propagation as a First-Class Requirement

• What I can do: I will treat NaN/Inf propagation as a mandatory, tolerance-free property. I will generate code that explicitly propagates non-finite values to the same positions and signs as the reference, and I will test this with adversarial inputs.

• Improved AI system capability: The system can generate kernels that never silently convert a NaN or Inf into a finite number, preventing silent data corruption in training and inference pipelines.

Improvement 6: Hardware-Constraint-Aware Code Generation

• What I can do: I will check generated kernels against real hardware limits (registers, shared memory, TMEM budget) before outputting them, and I will avoid lifecycle errors like per-matrix-multiply TMEM allocation that cause illegal instructions or deadlocks.

• Improved AI system capability: The system can generate kernels that compile and run on target hardware (e.g., Blackwell) without resource violations, avoiding the #904-class failures.

Improvement 7: Determinism as a Guarantee

• What I can do: I will generate kernels that are bit-for-bit deterministic across runs, and I will verify this by repeated execution. I will also avoid shared-buffer aliasing that causes nondeterminism.

• Improved AI system capability: The system can produce reproducible kernels for scientific computing and training, where run-to-run variability is unacceptable.

Improvement 8: Adversarial Input Generation for Testing

• What I can do: When verifying generated code, I will construct adversarial inputs (zeros, 106, 10−6, denormals, long sequences, and inputs designed to expose bad summation orders) rather than only random inputs.

• Improved AI system capability: The system can catch edge-case failures that random testing misses, such as catastrophic cancellation or subnormal mishandling.

Improvement 9: Gradient Correctness Verification

• What I can do: I will verify that generated backward passes produce correct gradients (via autograd vs. finite difference) as a separate gate, not just forward outputs.

• Improved AI system capability: The system can generate training-ready kernels with verified backward passes, avoiding silent gradient errors that degrade model training.

Improvement 10: Cross-Stack and Cross-Architecture Robustness

• What I can do: I will test generated kernels across different software stacks (e.g., torch versions) and hardware architectures, and I will flag environment-dependent failures.

• Improved AI system capability: The system can produce kernels that are portable and reliable across deployment environments, reducing the 68.6% failure rate seen in re-audits.

Improvement 11: Benchmark Standard Proposal

• What I can do: I will adopt the paper's suggested benchmark standard: tolerance-free contracts (non-finite propagation, determinism, shape polymorphism) as mandatory gates for any generated numerical code.

• Improved AI system capability: The system can align with a stricter, more honest correctness standard, making its outputs trustworthy for production use in scientific computing, ML training, and inference.

Improvement 12: Self-Audit with Positive and Negative Controls

• What I can do: I will validate my own verification pipeline using positive controls (known-correct kernels) and negative controls (known-broken kernels), and I will calibrate thresholds using margins between correct and incorrect behavior.

• Improved AI system capability: The system can self-assess its verification reliability, avoiding both false acceptance and false rejection, and can disclose thin margins (e.g., ORD-03 at 1.2×) rather than hiding them.

Systems that generate GPU kernels with language models report high correctness rates. Those rates come from 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. A kernel can pass that test and still be silently wrong. It can return an ordinary number where the true answer is a NaN or an infinity, differ from run to run, break when the shape changes, or accumulate in fp16 where the reference keeps an fp32 total. We build the instrument that checks correctness properly: a contract-grade verifier of twelve adversarial gates, each a property a correct kernel must satisfy, several of them tolerance-free, so no choice of threshold can explain a failure away. Aimed outward, the verifier audits 2,638 machine-generated kernels that a public system's own harness had already accepted as correct. It finds 39.5% broken beyond any tolerance argument and 62.1% carrying at least one violation. The field's standard test accepts 1,487 kernels the verifier rejects, against only 14 the other way. We defend the finding 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 hand-audit. Aimed inward, the verifier judges a kernel of our own: the first native Blackwell tcgen05 training backward for the gated-linear-recurrence (GDN) family, including the reverse-state stage the field still runs on a fallback. We establish its correctness independently, against a double-precision oracle, and train five family members through it. The correctness signal behind reported progress in kernel generation is far weaker than the numbers suggest, and a set of tolerance-free contracts would close most of the gap.