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

1. Problem Statement

The field of generating GPU kernels using Large Language Models (LLMs) reports high correctness rates, yet these claims rely on a fragile validation signal. Current benchmarks (e.g., KernelBench, TritonBench) typically accept a kernel if its output is "close" to a reference implementation on a few random inputs at a single fixed shape, using a loose tolerance (e.g., torch.allclose with atol=rtol=10⁻²).

The paper argues that this "single loose test" creates a correctness illusion. A kernel can pass this test while being fundamentally broken in ways the test never probes:

• Silent Data Corruption: Returning a finite number where the reference returns NaN or Infinity.
• Non-Determinism: Producing different results on each run.
• Shape Fragility: Breaking immediately when sequence length or batch size changes.
• Precision Drift: Accumulating reductions in fp16 while the reference maintains an fp32 total.

The authors posit that a significant portion of reported "wins" in kernel generation are illusory because the acceptance criteria are too weak to detect these silent failures.

2. Methodology

The Contract-Grade Verifier

To address the weak validation signal, the authors built a contract-grade verifier that operationalizes the "Kernel Contracts" taxonomy. This system subjects kernels to twelve adversarial gates, graded against a slow, high-precision reference (ground truth).

The gates are divided into two categories:

1. Tolerance-Free Gates: These require exact matches (byte equality or exact masks) and admit no tolerance arguments. They check for:
   • EXC-01/02: Correct propagation of NaN/Infinity and subnormal handling.
   • ORD-02: Bit-for-bit determinism across runs and absence of buffer aliasing.
   • RES-02: Adherence to hardware resource limits (e.g., Tensor Memory budget).
2. Thresholded Gates: These use derived tolerances based on floating-point error models (e.g., atol ≈ 4ε√N) rather than arbitrary hand-tuned values. They check for:
   • CMP-01/03: Correctness across diverse inputs and varying shapes.
   • ORD-01: Summation order bounds.
   • PRC-01/02: Precision regimes and internal accumulator precision.

Fairness Mechanisms:

• Red Team: The verifier is validated against 19 deliberately broken kernels (e.g., caching answers, returning inputs, silent NaN swallowing). It must reject all 19 while passing honest references.
• Positive Control: The verifier is tested on the authors' own hand-written kernels to ensure it is not tuned to fail foreign code.
• Calibration: Thresholds are set at the loosest defensible point, separating the "honest noise floor" of correct kernels from the "real error" of broken ones.

The Audit

The verifier was applied to 2,638 machine-generated kernels from the Dr. Kernel / KernelGYM corpus. Crucially, these kernels had already been accepted as "correct" by the source system's own harness. The audit measured how many of these "certified" kernels actually violated the contract.

The Native Backward Kernel

To provide an independent positive control, the authors developed a native Blackwell tcgen05 training backward for the Gated-Linear-Recurrence (GDN) family. This kernel:

• Implements the reverse-state scan and WY/triangular-inverse Vector-Jacobian Product (VJP) stages natively on Blackwell hardware.
• Solves the Tensor Memory (TMEM) constraint (Issue #904), where official kernels fail to compile due to exceeding the 512-column TMEM budget.
• Is verified against a double-precision (fp64) oracle and bit-for-bit determinism.

3. Key Contributions

1. A Twelve-Gate Contract-Grade Verifier: A scalable operationalization of correctness contracts featuring tolerance-free gates and derived thresholds, validated by a red team of broken kernels and a positive control.
2. The Rigor-Gap Audit: A large-scale audit of 2,638 accepted kernels revealing a massive correctness gap:
   • 39.5% (1,043 kernels) fail a tolerance-free gate, meaning they are broken in ways no tolerance argument can excuse.
   • 62.1% carry at least one contract violation.
   • The standard benchmark accepts 1,487 of these broken kernels, while the verifier rejects them; conversely, the verifier rejects only 14 kernels that the benchmark accepts.
3. First Native Blackwell tcgen05 Backward for GDN: A hand-written training backward for the GDN family (covering Mamba-3, GLA, KDA, etc.) that correctly manages the TMEM lifecycle, clearing the #904 compilation blocker.
4. Supporting Results:
   • Six contract-verified Mamba-3 Triton kernels that bypass #904.
   • Training of a 1.1B-parameter model end-to-end on these kernels.
   • A reinforcement learning (GRPO) autotuner that successfully tunes launch knobs but fails to discover source edits that improve speed without breaking correctness.

4. Results

The Rigor Gap

The audit demonstrates that the standard "allclose" check is a systematic blind spot.

• Modal Defect: The most common failure (34.2%) is EXC-01: a kernel silently replacing a NaN or Infinity with a finite number. This converts a detectable training error into silent data corruption.
• Shape Rigidity: 18.1% of kernels fail when the shape changes (CMP-03), despite passing on the fixed shape used in benchmarks.
• Differential: The disagreement between the standard check and the verifier is almost entirely unidirectional (1,487 vs. 14), indicating the standard check misses specific classes of errors rather than just being "stricter."

Verification of the Native Kernel

The authors' native backward kernel serves as a positive control:

• It passes all 12 gates of the verifier.
• It achieves a worst-case relative error of 3.3 × 10⁻³ against an fp64 oracle.
• It is bit-for-bit deterministic.
• It successfully trains five members of the GDN family.
• Speed Note: The native kernel is currently slower (structural latency) than the existing Triton fallback library (fla), primarily due to the sequential nature of the reverse-state scan and fp32 glue code. The authors explicitly state this speed deficit rather than obscuring it.

Alignment of Failure Modes

The paper highlights a critical alignment: the same gates that detected real bugs in the authors' own kernel during development (e.g., NaN propagation issues, shape inference errors, TMEM lifecycle errors) are the exact gates that detected the failures in the foreign corpus. This confirms the verifier is not biased against foreign code but is a general diagnostic tool.

5. Significance and Claims

The paper claims that the correctness signal behind reported progress in kernel generation is far weaker than the numbers suggest.

• Qualitative Gap: The gap is not merely a matter of "how close is close enough" (a threshold debate) but a qualitative failure to detect silent errors like NaN swallowing and non-determinism.
• Actionable Standard: A small set of tolerance-free contracts (non-finite propagation, determinism, shape polymorphism) would close most of the gap at modest cost. The authors propose these as a new benchmark standard.
• Methodological Integrity: By using the same instrument to audit foreign code and validate their own work (including admitting their own speed deficits), the authors demonstrate a methodology where the verifier's rejections are diagnostic rather than arbitrary.
• Systemic Impact: The work exposes that a significant share of "correct" LLM-generated kernels are actually broken, and that the field's reliance on loose, single-shape tests is insufficient for production-grade reliability.

The paper concludes that the twelve contracts identified are sufficient to begin closing this gap, shifting the focus from "passing a loose test" to "satisfying a contract."