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

Abstract: Systems that generate GPU kernels with LLMs 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.

• The paper introduces a twelve-gate, contract-grade verifier that tests value accuracy, shapes, precision, exceptional values, determinism, aliasing, and device behavior against FP64 and adversarial references.
• The audit finds that 62.1% of 2,638 previously accepted machine-generated kernels violate at least one contract, while 39.5% fail tolerance-free checks that ordinary allclose tests cannot detect.
• The paper presents a native Blackwell tcgen05 backward for five gated-linear-recurrence families, achieving roughly 3.3×10⁻³ worst relative error and deterministic training behavior, but remaining slower than mature Triton implementations.

Contract-Grade Verification and Native Blackwell Backpropagation for Gated Linear Recurrences

"A Contract-Grade Verifier for LLM-Generated GPU Kernels, and a Native Blackwell Backward for the Gated-Linear-Recurrence Family" (2608.12700) addresses two closely related systems problems: the inadequacy of prevailing correctness tests for LLM-generated GPU kernels and the absence of a fully native Blackwell training backward for the gated-linear-recurrence family. The paper’s central methodological claim is that kernel-generation benchmarks substantially overstate correctness when they rely on a single approximate-equality test at one fixed shape. Its systems contribution is a hand-written tcgen05 backward for GDN and several related recurrence models, including the reverse-state stage that existing implementations leave on a Triton fallback.

The two contributions are linked through a common verification standard. The verifier is used first as an external auditing instrument against 2,638 kernels already accepted by a public generation system, and then as an internal acceptance test for a kernel whose correctness is independently established against an FP64 oracle. This dual use is intended to address the principal threat to validity in an audit of this kind: that a stricter checker merely rejects foreign implementations by construction.

The Correctness Problem in Generated GPU Kernels

Existing kernel-generation benchmarks typically evaluate a candidate by executing it on several random inputs at one fixed shape and applying an allclose criterion. In the KernelBench setting considered by the paper, the common configuration uses absolute and relative tolerances of $10^{-2}$. Such a test is inexpensive, but it probes only a narrow portion of a kernel’s behavioral contract.

The paper identifies several failure modes that can survive this procedure:

• A kernel can return a finite value where the reference produces a NaN or infinity.
• It can be nondeterministic across repeated executions.
• It can work at the benchmarked shape while failing under a different sequence length, batch size, or feature dimension.
• It can accumulate in FP16 even when the reference maintains an FP32 accumulator.
• It can mishandle subnormals, device placement, reduction ordering, or aliasing.
• It can produce numerically plausible outputs while implementing the wrong gradient.

These are not merely disagreements over an appropriate numerical tolerance. Several are categorical violations of the operator’s semantics. In particular, replacing a non-finite reference value with an ordinary finite number can convert an observable training failure into silent state corruption.

The paper therefore distinguishes a conservative, tolerance-free correctness floor from a broader rate that includes tolerance-dependent violations. This distinction is important: the former cannot be dismissed by arguing that the verifier’s numerical thresholds are overly strict.

The Twelve-Gate Verifier

The verifier operationalizes the Kernel Contracts taxonomy (Veit, 23 Apr 2026) as twelve adversarial gates. They cover value correctness, gradients, shape polymorphism, reduction behavior, precision regimes, exceptional values, device residency, determinism, aliasing, and hardware resource constraints.

Seven gates are treated as load-bearing for the forward-only audit:

• CMP-01: value correctness over random and adversarial inputs;
• CMP-03: correctness across shapes;
• ORD-01: reduction-order error within a derived bound;
• ORD-02: deterministic, non-aliased output;
• EXC-01: exact NaN and infinity propagation;
• EXC-02: subnormal and flush-to-zero behavior;
• PRC-01: correctness across FP32, FP16, and BF16.

The remaining gates are either unavailable or less informative for portions of the audited corpus. Gradient correctness cannot be tested when the candidate is forward-only, and resource metadata is unavailable for some generated kernels. The verifier nonetheless exercises these dimensions on the authors’ own kernels and on deliberately constructed controls.

Its numerical policy is designed to avoid arbitrary threshold selection. Tolerances are derived from floating-point error models and scale with the relevant operation. For example, reduction-order error is bounded approximately according to the expected $\sqrt{N}$ accumulation behavior rather than by an empirically convenient constant. Exact-mask comparisons are used for non-finite values, and repeated executions are compared byte-for-byte for determinism. Inputs are generated from fixed random seeds, ensuring that verdicts are reproducible and not dependent on stochastic sampling luck.

The repository also contains nineteen deliberately defective kernels implementing behaviors such as input return, cached outputs, FP16 accumulation, nondeterminism, shape specialization, non-finite suppression, and device migration. The verifier rejects all nineteen while accepting the honest reference operator. This two-sided test establishes that the battery is sensitive to targeted violations without rejecting a known-correct implementation.

The Audit of Accepted Machine-Generated Kernels

The external audit targets the Dr. Kernel/KernelGYM corpus (Liu et al., 5 Feb 2026), specifically an accepted subset of Triton kernels evaluated on a B200 using PyTorch 2.12 and Triton 3.7. Of 3,134 kernels in the selected operator classes, 2,638 had already passed the source system’s correctness predicate and had recorded positive speedups.

The verifier reports two headline rates:

• 62.1% of the accepted kernels violate at least one contract.
• 39.5%, or 1,043 of 2,638 kernels, fail a tolerance-free gate.

The second figure is the paper’s principal conservative result. These 1,043 kernels exhibit failures that cannot be explained by a choice of approximate-equality threshold.

The modal defect is EXC-01: non-finite non-propagation. It affects 868 kernels, or 34.2% of applicable cases. Other substantial failure rates include value errors on varied inputs, shape rigidity, missing FP32 accumulation, reduction-order errors, precision-regime failures, subnormal mishandling, and nondeterminism.

Figure 1: The audit of 2,638 previously accepted kernels, showing the contract-violation rate, tolerance-free failure floor, and per-gate failure distribution.

The findings remain stable under several re-slicings. Including shape rejection in the tolerance-free subset raises the floor to 41.1%. Removing the matrix-multiplication class, where TF32-related tolerance arguments might be most relevant, leaves a 61.9% violation rate. Removing all contested channels still yields 60.6%.

The result is reproduced on a second software stack: a 300-kernel re-audit under PyTorch 2.11 and CUDA 12.8 produces a 68.6% violation rate. The paper does not treat this as a new headline estimate because the smaller sample overweights the reduction class, but it demonstrates that the result is not tied to a single software environment. A separate native-CUDA corpus shows a weaker residual defect rate of 20.2%, driven primarily by shape rigidity, subnormal handling, and nondeterminism rather than widespread value corruption.

Validation Against the Benchmark Harness

The authors provide four defenses against the claim that the verifier is simply excessively strict.

First, seven independently known-correct kernels pass all applicable gates, including six Mamba-3 Triton kernels and the native GDN backward. The native backward was not used to calibrate the thresholds, making it a particularly important positive control. The verifier did identify a real input-validation defect in one of the authors’ own kernels: a fused block inferred its channel dimension from the input without checking compatibility with the convolution weights. The defect was fixed rather than reclassified.

Second, the threshold-calibration sweep places the selected thresholds between the noise floor of correct kernels and the error magnitudes induced by deliberate perturbations. The margins are broad for most gates. The exception is the deliberately relaxed ORD-03 gate, whose upper margin is only $1.2\times$; its coverage is therefore supported by the stronger CMP-01 and ORD-01 gates.

Figure 2: Calibration of thresholded gates, showing the separation between correct-kernel noise and deliberately injected errors.

Third, the authors compare their reimplementation with KernelBench’s own correctness code (Ouyang et al., 14 Feb 2025). The two implementations agree on 98.5% of 1,030 pairs. Of the fifteen disagreements, seven involve the verifier being stricter and eight are seed-sensitive borderline cases. This comparison supports the claim that the large discrepancy arises from the narrow behavioral scope of the benchmark test rather than from a faulty replica.

Fourth, a stratified hand audit of 31 disputed cases rejects the possibility that every disagreement represents a genuine correctness failure. Sixteen cases are genuinely broken under the tolerance-free core, eight are real but tolerance-dependent, and seven are excluded as out of scope. The willingness to remove seven disputed cases strengthens the methodological credibility of the audit.

The direct differential is strongly asymmetric. KernelBench’s paper-era check accepts 2,472 of 2,638 kernels, or 93.7%. Of the accepted set, 1,487 are rejected by the contract verifier, including 958 tolerance-free failures. In the reverse direction, only fourteen kernels are rejected by KernelBench but accepted by the verifier.

Figure 3: The directional disagreement between the standard benchmark test and the contract-grade verifier, dominated by kernels accepted externally but rejected under adversarial contracts.

This $1,487$-to-$14$ asymmetry is theoretically more informative than a simple difference in acceptance rates. If the verifier were merely a uniformly stricter allclose, disagreements would be expected in both directions. Instead, the result indicates a systematic blind spot in the standard test: it omits behavioral dimensions on which many kernels fail.

The Native Blackwell GDN Backward

The paper’s second contribution is a native Blackwell training backward for the gated-linear-recurrence family. The target recurrence maintains a matrix state $S_t$ and combines channel-wise decay, delta-rule correction, rank-one writes, and query-based reads. The general GDN formulation subsumes five model families:

• linear attention;
• gated linear attention;
• SSD/Mamba-2 with scalar decay;
• Kimi Delta Attention;
• gated DeltaNet.

The authors argue that approximately 80% of the parameterization is shared across these models. Consequently, a single differentiated implementation of the general recurrence can serve all five, provided each reduction is independently checked against its own reference rather than inferred solely from gate settings.

The backward is organized around two difficult stages. The first is a reverse inter-chunk state scan. Chunking makes the forward recurrence partially parallel, but the backward must propagate a state-gradient accumulator backward through the chunks, including the rank-one delta correction. This is the stage that existing open implementations leave in the Flash Linear Attention Triton fallback.

The second is the WY/triangular-inverse vector-Jacobian product. Within each chunk, the delta rule requires a triangular solve. The backward reuses the forward triangular factor and expresses the VJP through triangular matrix multiplications rather than recomputing an inverse.

The implementation uses Blackwell’s fifth-generation tensor cores through tcgen05, Tensor Memory, TMA movement, asynchronous pipelines, and a $(128,64,128)$ tile. The main hardware difficulty is the 512-column TMEM budget per warpgroup. The paper connects this constraint to issue #904 in the Mamba repository, where a compiler-generated kernel requests 544 columns and falls back to a path reported as 38.7 times slower on B200/GB200 hardware.

The authors reproduce the failure and attribute their own initial illegal-code and deadlock failures to incorrect TMEM lifecycle management. The successful implementation reserves TMEM once, uses fixed accumulator offsets, and releases the allocation once after all matrix multiplications. This allocation-once/relinquish-once discipline permits multiple tcgen05 MMAs within the kernel.

Numerical Verification and Training Behavior

The native backward is verified through a layered oracle chain:

1. a token-serial FP64 specification;
2. a chunkwise reference;
3. an FP64 assembly of the two difficult backward stages;
4. the native tcgen05 implementation on B200 hardware.

The core tiles agree with their closed-form specifications to approximately $7\times10^{-16}$. The assembled backward reaches worst relative errors of $3.29\times10^{-3}$ for the scalar path and $3.31\times10^{-3}$ for the channel-wise path, both below the stated $\sqrt{N}$0 acceptance bound. The implementation is bit-for-bit deterministic across repeated executions.

The paper discloses an important exception rather than suppressing it: the widest $\sqrt{N}$1 save-forward arm reaches $\sqrt{N}$2, exceeding the acceptance bound by approximately 4%. The individual kernel itself measures $\sqrt{N}$3 against the same FP64 oracle, so the overrun is attributed to the assembled pipeline rather than to the native kernel. The formal correctness claim is correspondingly scoped.

The backward trains all five supported recurrence variants without enabling the reference fallback. Real 300-step training loops exhibit no numerical blow-ups after fixing an exponent-overflow bug through a masked exp2 implementation. The native path also supports graph replay with bit-exact agreement against eager execution.

The positive-control result is significant: the native backward passes the complete twelve-gate verifier and is clean on every tolerance-free channel. This does not prove universal correctness outside the tested domain, but it demonstrates that the audit battery can accept a complex, independently established Blackwell kernel.

Performance and Engineering Trade-offs

The native backward is not competitive with the mature FLA Triton implementation in latency. It is approximately eight times slower at sequence length $\sqrt{N}$4 and approximately 78 times slower at $\sqrt{N}$5. The authors attribute this disparity to the inherently sequential reverse-state scan and to substantial unfused FP32 glue surrounding the native tensor-core stages. The FLA implementation also benefits from reusing intermediates computed during the forward pass.

This performance result qualifies the systems contribution. The native implementation closes an availability and coverage gap, rather than delivering a speed record. Its primary advantages are that it is genuinely native to Blackwell, handles the reverse-state stage, generalizes across the five recurrence families, and correctly manages the TMEM lifecycle. The paper reports internal optimizations, including a 2.75-times reduction in captured backward time through channel-wise fusion and a 2.98-times reduction for a $\sqrt{N}$6 save-forward variant through TMEM tiling, but these are improvements over the authors’ own earlier pipeline, not over FLA.

The supporting Mamba-3 kernels provide a related availability result. Six Triton kernels avoid tl.dot and therefore compile under configurations where the official implementation encounters the 544-column TMEM request. This is a categorical compatibility advantage, although the paper explicitly avoids turning it into an unconditional latency claim.

Broader Experimental Implications

The paper’s applied demonstrations reinforce, but do not independently establish, the central audit result. The authors train a 1.1-billion-parameter Mamba-3 SISO model on PTB-XL using the verified kernels and report a peak macro-AUC of 0.880. The experiment demonstrates end-to-end trainability and zero NaNs after fixing an unstable $\sqrt{N}$7 parameterization and a device-placement issue; it is not presented as a state-of-the-art physiological-signal result.

The authors also construct a GRPO-based autotuning system whose reward grants speed bonuses only after the full contract battery passes. Configuration-only RL improves performance by 1.167 times over the shipped default, although a deterministic shape-gated selector achieves a larger 2.174-times geometric-mean improvement over the earlier serial default. Source-edit RL performs poorly: 208 of 320 edits fail to apply, 37 applied edits fail correctness, and no successful edit exceeds the baseline. This negative result illustrates a practical benefit of contract-grade verification: incorrect optimizations are rejected before they can receive performance credit.

More generally, the work suggests that kernel-generation benchmarks should treat correctness as a multi-dimensional contract rather than a scalar approximate-equality score. At minimum, benchmark protocols should test non-finite propagation, determinism, shape variation, mixed-precision accumulation, and device/resource constraints. These tests are relatively inexpensive compared with the cost of training models or deploying kernels whose silent errors may be difficult to diagnose.

Limitations and Theoretical Significance

The audit is corpus-dependent. Although the result survives a second stack and a second native-CUDA corpus, the headline rate is still derived primarily from one Triton corpus. Gate coverage is also incomplete: gradient correctness and resource metadata are unavailable for the forward-only audit set. The verifier’s reference implementations are another potential dependency, although the FP64 design, adversarial red team, positive controls, and independent hand audit mitigate this concern.

The native backward is evaluated on a narrow envelope: B200 hardware, $\sqrt{N}$8, chunk length 64, and $\sqrt{N}$9. It remains partially native because normalization gradients, packing, masks, and other glue execute through PyTorch. Consequently, neither the numerical verification nor the performance characterization should be generalized to arbitrary dimensions, future Blackwell variants, or fully fused training pipelines.

The paper’s stronger theoretical implication concerns the semantics of evaluation. A conventional tolerance test treats correctness as proximity on sampled finite-valued outputs. The contract formulation instead treats correctness as preservation of a structured execution contract across input domains, shapes, precision modes, exceptional values, executions, devices, and resource constraints. This reframing is especially appropriate for GPU kernels, where undefined behavior, silent precision changes, aliasing, and hardware-specific compilation failures can invalidate an implementation without producing a large error on ordinary random inputs.

Conclusion

The paper presents a unified argument for stronger evaluation of generated GPU kernels and for disciplined implementation of specialized GPU backward passes. Its audit of 2,638 previously accepted kernels finds a 39.5% tolerance-free failure floor and a 62.1% overall contract-violation rate, with the standard benchmark check accepting 1,487 kernels that the verifier rejects and only fourteen disagreements in the opposite direction. The result is supported by positive controls, calibration experiments, agreement with the benchmark’s own code, and a stratified manual audit.

Its native Blackwell contribution supplies a tcgen05/TMEM backward for the GDN family, including the reverse-state scan absent from existing native implementations. The backward is independently checked against FP64 references, achieves approximately $1.2\times$0 worst relative error in the principal configurations, is deterministic, and trains five recurrence variants. It is nevertheless substantially slower than FLA, and its verified operating envelope remains narrow.

The principal consequence is methodological: reported progress in LLM-based GPU-kernel generation should not be interpreted without examining the contract used to certify correctness. A small set of tolerance-free behavioral checks would eliminate a large fraction of the current acceptance gap, while the Blackwell implementation demonstrates how the same verification discipline can guide the development of complex, hardware-specific kernels.