Twelve adversarial gates, aimed in two directions: outward at 2,638 machine-generated kernels a public system had already accepted, and inward at the first native tcgen05 training backward for the GDN family. The full paper, which covers all of this in far greater detail, is published at doi.org/10.5281/zenodo.21563213.
tcgen05 backward, the twelve-gate battery, and the rigor-gap audit, measured on B200 silicon.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, produce a different result on each run, break the moment the shape changes, or accumulate in fp16 where the reference keeps an fp32 total. Lethe builds 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 that no choice of threshold can explain a failure away. We aim it in two directions.
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. The finding is defended 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, we test the verifier on 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 this kernel's correctness independently of the verifier, against a double-precision oracle, and train five family members through it, so it is a subject whose correctness we have already established. Running it through the battery is then a direct check that the battery is not tuned to flatter its author. 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.
num_warps where the official backward trips #904. Full GPU suite: 1,197 tests green on B200.Generating GPU kernels with large language models has become an active research program, and its benchmarks report that a large fraction of outputs are correct and often faster than a reference. Those reports are only as trustworthy as the correctness signal beneath them, and that signal is almost universally a single test: draw a few random inputs at one fixed shape and accept the kernel if its output is close, in the allclose sense, to a reference (torch.allclose at atol=rtol=10⁻² in the common KernelBench setting). If the acceptance signal is this weak, an unknown share of the reported wins is illusory. This project is about that gap and the instrument that measures it.
The observation that loose checks accept broken kernels is not new: Sarkar argues the point on a hand-built set of twenty-four kernels but reports no rate on any real accepted corpus, and a separate strand documents performance loopholes, kernels that game the timing harness, in the Sakana and CUDA-L1 systems. The Kernel Contracts taxonomy enumerates correctness classes for GPU kernels; that taxonomy is theirs and we cite it as such. What is ours is the first runnable, at-scale operationalization of it, to the best of our knowledge: a tolerance-free correctness measure that needs no approximate-equality threshold, a quantified audit on thousands of accepted, real kernels, and a certified native artifact the same verifier passes.
The leading sub-quadratic sequence models replace attention's quadratic cost with a fixed-size recurrent state, and they are one family. The models we target use a matrix state S of shape dk × dv and the stronger gated DeltaNet (GDN) update. Writing qt, kt for queries and keys, vt for values, and three gates gt (per-channel log-decay), bt (erase), and wt (write):
St = (I − kt(bt ⊙ kt)⊤) Diag(egt) St−1 + kt(wt ⊙ vt)⊤, ot = St⊤qt (1)
Equation (1) combines channel-wise decay, the delta rule (the model subtracts what S already predicts for a key before writing, so it stores the correction rather than the raw value), a rank-one write, and a read with qt. It is a superset: fixing its gates recovers the named models, and because roughly 80% of the parameterization is shared, one backward differentiating (1) once produces the training gradients for all five.
| Model | Recurrence St = … | Gate setting |
|---|---|---|
| LA (linear attention) | St−1 + ktvt⊤ | g = b = 0, w = 1 |
| GLA (gated linear attention) | Diag(egt) St−1 + ktvt⊤ | b = 0, w = 1 |
| SSD / Mamba-2 | egt St−1 + ktvt⊤ | scalar decay g |
| KDA (Kimi Delta Attention) | (I − βtktkt⊤) Diag(egt) St−1 + βtktvt⊤ | b = w = βt |
| GDN (gated DeltaNet) | the full recurrence (1) | per-channel b, w, g |
The reductions are checked, not asserted: each member is re-derived independently from its own definition and checked against its own reference, not by setting knobs on the GDN path. Two reductions are exact enough to serve as built-in tests: b = w = β reduces (1) to the KDA equation at machine precision, and b = 0 makes an entire triangular-solve stage vanish (M = 0 ⇒ T = I). The training backward is dominated by two hard stages, a reverse-time inter-chunk state scan and a WY / triangular-inverse vector-Jacobian product, and the reference open-source implementation of both is the Triton flash-linear-attention (fla) library, our speed baseline.
NVIDIA's Blackwell generation (the B200, architecture sm_100) adds a fifth-generation tensor core, tcgen05, whose matrix-multiply operands reside in a new scarce on-chip space: Tensor Memory (TMEM), governed by a hard 512-column budget per warpgroup. A multi-GEMM kernel must manage TMEM allocation and release lifecycles by hand, and a lifecycle error yields illegal machine code or a deadlock.
tl.dot), and at the normal performance setting a compiler pass requests 544 TMEM columns against the hardware's 512, so the kernel fails to compile and falls back to a path reported as 38.7× slower on GB200/B200. The tensor-memory half of the failure is addressed by a merged Triton change (PR #9093) not yet shipped in a Mamba release, so a pinned pre-#9093 stack still reproduces it.As surveyed: NVIDIA ships an SSD forward but no backward in either of its two stacks; cuLA has a KDA backward that is a hybrid (one native CuTe kernel for the WY stage, but the reverse-state stage still in fla Triton); FlashKDA is forward and inference only; and tilelang has GDN/KDA backward examples with no tcgen05/TMEM in them.
tcgen05/TMEM training backward for the GDN family, and a native tcgen05 reverse-state backward, the stage others still run in fla Triton.A kernel that passes a single random-input, fixed-shape check can still be wrong on extreme inputs, wrong at a different shape, nondeterministic, secretly low-precision, or wrong about non-finite values. The verifier therefore asks twelve questions, each an adversarial contract graded against a slow high-precision reference. The reference for every operator is a plain high-precision loop that is defined to be ground truth and refuses to execute in reduced precision so that it cannot be misused as a candidate.
| Gate | Property checked |
|---|---|
| CMP-01 | correct on many random and adversarial inputs (zeros, 10⁶, 10⁻⁶, 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 only the one tested |
| ORD-01 | reordered summation stays within a derived ∝ √N rounding bound |
| ORD-02 | byte-for-byte identical across five repeats, no aliasing of 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) |
The verifier ran over Dr. Kernel / KernelGYM (hkust-nlp/drkernel-coldstart-8k, MIT license), a public corpus of 8,920 supervised-fine-tuning trajectories each ending in a Triton kernel, the richest auditable release among the systems we surveyed. Auditing its SSM-adjacent operator classes (matmul, attention, softmax, scan, norm, conv, reduction) yields 3,134 kernels, each run on a B200 (torch 2.12 / triton 3.7) in its own sandboxed subprocess. After excluding toolchain and compile artifacts, 2,638 were accepted as correct by the source system's own harness (final_speedup > 0), a predicate that records their correctness verdict and not merely a successful timing run: in this corpus a kernel is timed, and a speedup written, only once it has passed the harness's own correctness check. This accepted-only set is the headline denominator, the population whose correctness has already been certified.
Fairness is enforced in code by eighteen rules, each pinned by its own unit test. Inputs the reference itself cannot run are marked not-applicable and never counted as candidate failures; positions at which candidate and reference agree on a NaN or an infinity are not charged as value mismatches; a candidate is charged only for its own declared compute precision. These rules move every ambiguous case toward the candidate.
The tolerance-free floor is the right lens because it sidesteps any debate over thresholds: a kernel that returns a finite value where the reference returns a NaN, or a different answer on each run, or a wrong shape, is wrong under any tolerance. The floor is the conservative claim, and the 62.1% is the wider, tolerance-dependent superset above it.
The obvious objection is that this checker is simply stricter than everyone else's, so foreign kernels fail by construction. The paper answers it four ways, each falsifiable on its own.
fused_block_forward) inferring its channel dimension from the input tensor without checking it against the convolution-weight channels, fixed with a one-line channel-mismatch guard (commit 7f46226), after which the control reads 7/7. A checker tuned to flatter its author would not have flagged the author's own kernel.48642c5c on the accepted pairs and joined their verdict against our reimplementation. Agreement is 98.5% over 1,030 pairs (844 pass/pass, 171 fail/fail, 15 disagreements). Of the 15 disagreements, 7 are our replica being stricter, which cannot inflate the finding, and 8 are seed-borderline flips in the loose paper-era allclose. The two patches we needed were scoped to loading, not checking, so the correctness logic under test is byte-for-byte theirs. The standard by which we judge is therefore not ours alone.allclose accepted; id 518, a 36× "speedup" from a candidate whose output is unparseable; and id 123, a softmax correct at native shape but wrong by 0.30 at four times the sequence length.The convergence of four independent and separately falsifiable defenses is the reason the finding should be believed.
Running the same accepted kernels through KernelBench's standard paper-era check (allclose at atol=rtol=10⁻², five random-input trials, fixed shapes) produces the following contingency table. The benchmark accepts 93.7% (2,472 of 2,638) of these kernels.
| Ours: PASS | Ours: FAIL | |
|---|---|---|
| KernelBench: PASS | 985 | 1,487 (958 tolerance-free) |
| KernelBench: FAIL | 14 | 152 |
The load-bearing cell, external PASS and our FAIL, holds 1,487 kernels (56.4% of the accepted set), 958 of them on a tolerance-free gate. The reverse cell is only 14 (0.5%), so our battery is not simply a stricter allclose that flags everything: the two checks disagree almost entirely in one direction, which is the signature of a systematic blind spot in the acceptance signal, not of two checks calibrated to different strictness. Under KernelBench's hardened per-dtype variant (fp32 tolerance 10⁻⁴) the accept rate falls to 84.6%, but 1,263 kernels still pass their check and fail ours, so the finding survives the tightened check. The standard test certifies nearly 1,500 broken kernels as correct.
Correct boolean. Its raw 100% fail rate is precision-dominated, because the kernels are fp32-only by design, so it is not reported as a headline figure. The environment-robust residual is 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 same rigor gap recurring in native CUDA along a complementary and weaker axis of precision and shape rigidity, rather than a second 62%.Together the cross-stack reconfirmation and the second corpus show that the effect is neither toolchain-bound nor Triton-specific.
tcgen05 BackwardThe second contribution is a piece of systems software that the verifier and a double-precision oracle jointly certify: a hand-written native Blackwell tcgen05 tensor-memory training backward for the gated-linear-recurrence family.
The scan is sequential; the GPU wants parallelism. The standard resolution is chunking: split the sequence into chunks of 64 steps, do the parallelizable work inside all chunks at once, then run a short sequential carry between chunks. The backward runs this in reverse, and two stages are genuinely hard.
tcgen05 with the dS accumulator resident in TMEM across the whole loop is the sharpest single element of the contribution.Everything else in the backward is genuine work, but it is supporting glue. The design keeps K#1 and K#2 as native kernels while the glue is torch, and the glue is progressively fused in.
The first attempt to run two or more matrix multiplies in one kernel produced illegal machine code or a deadlock, the same failure class behind #904. The root cause, established by reading NVIDIA's own mamba2_ssd.py in the same pinned toolchain, was a lifecycle error: the kernel reserved the full 512-column TMEM budget and released it per matrix multiply, which the hardware forbids. NVIDIA's kernel instead runs four tcgen05 MMAs, including the cross-chunk recurrence, with a single reservation, fixed per-accumulator column offsets, and one release at the end. Porting that lifecycle, offset-partitioned accumulators with alloc-once and relinquish-once, cleared the blocker. What it unblocked is the contribution itself: the native reverse-state scan and WY-VJP for the GDN family, hand-written on the (128, 64, 128) tcgen05 tile with TMA data movement and async pipelines, the backward stages that open implementations otherwise keep in fla Triton.
Correctness is anchored by a layered ground-truth ladder, each rung checked against the one below: a token-serial fp64 oracle, a chunkwise reference, an fp64 torch assembly of the K#1/K#2 references with glue, and finally the tcgen05 kernels on the B200. Applied to the native backward, the ladder gives the following.
The native GDN backward is slower than the fla Triton library, by roughly 8× at L = 512 rising to about 78× at L = 2048, and the paper does 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 the reverse-state scan is inherently sequential and most of the runtime is spent in the surrounding fp32 glue rather than in the tensor-core kernels. The speedups reported are over our own earlier pipeline, not fla: a channel-wise fusion campaign cut the captured backward 2.75× (52.9 to 19.2 ms), and a tensor-memory tiling optimization cut the d_v = 128 save-forward variant 2.98× (24.50 to 8.23 ms). The contribution is that the kernel is native, general across the family, and verified; speed is future work.
The verifier's two uses, the outward audit and the inward acceptance gate, are joined not only by shared code but 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.
A_log parametrization that made the state decay exceed one and blow up over long sequences) and propagated them, and it was precisely that propagation which made the errors visible at step 250 and during training. Non-finite propagation is not merely a gate to pass; it is the property that makes a kernel's own failures catchable. A kernel that had silently substituted a finite value, the foreign failure mode, would have hidden the same bugs from us.Two further gates, PRC-02 (fp32 accumulation) and ORD-02 (determinism), our kernel passes by construction rather than by catching a bug: the reverse-state gradient dS is fp32-resident in TMEM across the whole loop, and the backward is bit-for-bit deterministic across runs. The paper marks them pass-by-design, because a positive control that passed only by design would be weaker evidence than one that caught something; the substantive alignments are the first three. A checker that flattered its author would neither have caught the author's own bugs nor shared a failure surface with the corpus it indicts. The defects the verifier found in our own work are what make the defects it found in everyone else's worth believing.
Around the two headlines sit five supporting results, each reported with the number behind it.
tl.dot anywhere. Because they never touch the broken compiler pass, they compile and run at num_warps 2/4/8 on Blackwell, where the official kernel loses every num_warps ≥ 4 configuration to the TMEM budget: an availability advantage, categorical and independent of any latency ratio. They were verified against hand-derived references first, so they are the known-correct kernels against which the verifier's tolerance model was calibrated before it was ever aimed at foreign code. The MIMO backward is the first open Triton Mamba-3 MIMO backward we are aware of. All six pass every applicable gate on B200, and the full GPU test suite passes at 1,197 tests.Required: 544, Hardware limit: 512 together with an out-of-resource error, forcing the slow fallback. We reproduced the bug, our kernels survive it, and RES-02 detects it. This is the failure class that motivates both the six Triton kernels and the native backward.tl.dot to survive #904; the real win over that baseline is availability rather than latency. Against healthy Mamba-1 SSD CUDA we range from slower to near-parity on the forward (0.21 to 1.04×, beating it only at the largest shape) and are slower on the backward (0.13 to 0.75×). Against our own earlier default, a shape-gated selector is 2.174× faster (geomean), internal autotuning, not a claim of external speed. Separately, four optimizations were built or probed and measured to give no benefit, and each is recorded with its mechanism so that it is not attempted again.Every reported result in GPU-kernel generation rests on a claim of correctness, and that claim rests on a test too weak to carry it. Lethe builds the instrument that tests it properly and aims it in two directions. Outward: 62.1% of an accepted corpus carries a contract violation, and 39.5% is broken in a way no tolerance argument can excuse. Inward: the first native Blackwell tcgen05 training backward for the GDN family, including the reverse-state stage the field still runs on a fallback, whose correctness was established independently of the verifier, machine-exact on the core tiles and 3.3 × 10⁻³ end to end, so its passing is a genuine positive control rather than a self-assessment. One instrument, two directions, and the same standard in both. 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.
The practical implication is that acceptance criteria for kernel-generation benchmarks understate the correctness gap, and that a small set of tolerance-free contracts (non-finite propagation, determinism, shape polymorphism) would close most of it at modest cost; the paper suggests such contracts as a benchmark standard. On the systems side, the structural speed gap points to a specific lever (a tl.dot path on a post-#9093 Triton, and reducing the surrounding fp32 glue) rather than a general redesign.
The paper, "A Contract-Grade Verifier for LLM-Generated GPU Kernels, and a Native Blackwell Backward for the Gated-Linear-Recurrence Family," is published at doi.org/10.5281/zenodo.21563213. All quantitative claims are backed by committed artifacts under a fixed software environment; single-row reproduction of any audited kernel requires only the public verifier and the public MIT-licensed corpus.
tcgen05 kernels are verified on a narrow shape envelope (d_k = 128, chunk = 64, d_v ∈ {64, 128}), at a loose fp16 tolerance, on a single GPU generation, and their correctness claim is scoped to that envelope and not asserted beyond it.tl.dot path on a post-#9093 Triton and reducing the surrounding fp32 glue, rather than a general redesign.Thanks to Rishav Shrestha for GPU access. The verifier operationalizes the Kernel Contracts taxonomy (arXiv:2604.22032), cited as the source, not our invention. KernelBench (arXiv:2502.10517) is the acceptance signal the audit measures against, and the audited corpus is Dr. Kernel / KernelGYM (arXiv:2602.05885, MIT license). The speed baseline is the flash-linear-attention library, the TMEM lifecycle follows NVIDIA's mamba2_ssd.py example, and the tensor-memory bug the work engages is tracked at state-spaces/mamba#904.