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 on arXiv at arXiv:2608.12700 and archived 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 open implementations still run 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. Priority on the MIMO backward is TileLang's, which ships the official one. Full GPU suite: 1,191 passed / 3 skipped / 0 failed 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. The tensor-memory half of the failure is addressed by a merged Triton change (PR #9093) that has not shipped in any Triton 3.7.x 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. Three filters then reduce that set to the headline denominator, in order: 2,970 carry a positive recorded speedup (final_speedup > 0), the corpus's own record that the kernel cleared its pipeline and was timed; the release ships no separate correctness column, so this is the strongest acceptance predicate the artifact supports. 2,919 of those expose a floating-point tensor input for the gates to perturb, the rest marked not-applicable rather than counted as failures; and 2,638 remain after excluding toolchain and compile artifacts. This accepted-only set is the headline denominator. The operator classifier and the artifact predicate are both stated exactly in the paper's appendix, so each of the three counts is reconstructible rather than asserted.
Fairness is enforced in code by five rules, pinned by nineteen unit tests. 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 that aliases a shared buffer, crashes, or does not run at all, is wrong under any tolerance. The floor is the conservative claim, and the 62.1% is the wider, tolerance-dependent superset above it.
The floor is the union of five channels that admit no tolerance argument: EXC-01, ORD-02, output aliasing, a CUDA error during gating, and a non-artifact failure to run. EXC-02, subnormal flush-to-zero, is tolerance-free by the same standard but is deliberately held out, since it is the one member of the group a reader can reasonably call a convention rather than a defect; including it would raise the floor to 43.7%. The reported 39.5% is the conservative reading of our own gate set, not the maximal one. It is also not carried by any one operator class: every one of the seven audited classes independently exceeds 25%, and removing any single class leaves the corpus floor between 31.6% and 41.9%.
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 channel-mismatch guard added to both public ops and the reference (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). That is 39% of the accepted set: the sharded run was stopped once agreement had stabilized, at 98.9% over 458 pairs and 98.5% over 1,030. Of the 15 disagreements, 7 are our replica being stricter, which cannot inflate the finding, and 8 are rows our replica accepts and their code rejects: 2 where their harness raised before any comparison happened, and 6 input-draw-dependent value mismatches. The two patches we needed were scoped to loading, not checking, so the correctness logic under test is byte-for-byte theirs. Their code ran on 1,030 of the 2,638; the differential below uses our replica over all of them, so what this check establishes is agreement wherever both ran.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 is the sharpest single element of the contribution: one tensor-memory reservation held across the entire unrolled reverse scan, both matrix-multiply accumulators addressed at fixed column offsets inside it, and a single release at the end. The fp32 dS carry itself is held outside the tensor memory, updated by the kernel's own threads between chunks.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 carried in fp32 across the whole loop, never demoted to the fp16 operand dtype, 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. Priority on the MIMO backward is TileLang's, which ships the official one. All six pass every applicable gate on B200, and the committed record of the full GPU test suite is 1,191 passed, 3 skipped, 0 failed.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 open implementations still run 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 on arXiv at arXiv:2608.12700 and archived 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 of Cooper Veit (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 by Liu et al., ICML 2026 (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.