{"id":"65fe4654-ab91-400c-a3fb-bafc2bc595fe","arxiv_id":"2608.12700","paper_version":1,"verdict":"CONDITIONAL","confidence":"MODERATE","novelty_score":7.0,"correctness_risk":"medium","formal_verification":"none","parameter_count":2,"one_line_summary":"A twelve-gate contract verifier finds that 39.5% of 2,638 LLM-generated kernels accepted by a standard harness fail tolerance-free checks, then clears the authors' own native tcgen05 GDN backward.","lead":"This paper builds a 12-gate verifier for LLM-generated GPU kernels and finds that 39.5% of kernels previously accepted as correct fail a tolerance-free check. It also reports a new native GPU backward pass for gated linear recurrence models, verified against a double-precision oracle.","discovery_kind":"new_method","skeptic_critique":{"model":"deepseek-v4-flash","headline":"The 39.5% tolerance-free floor assumes a single fp64/non-FTZ reference semantics; under common GPU FTZ/fast-math contracts many EXC-01/EXC-02 failures are not bugs, so the floor's 'no tolerance can excuse' claim is not yet independently supported.","rationale":"The reader's weakest assumption identifies exactly the load-bearing point: the audit's reference semantics are defined, not argued for, and the tolerance-free gates inherit that normativity. I agree with that identification. The paper is unusually transparent about its choices and even lists reference correctness as a threat to validity, but it does not provide an independent benchmark for the semantics. The practical question is whether a generated kernel that flushes subnormals or fails to propagate a NaN under fast-math is 'broken.' For training linear recurrences, NaN propagation matters, so EXC-01 is valuable; but the 'tolerance-free' framing overclaims because the exactness is relative to a chosen environment. The 62.1% number is less affected since it includes tolerance-dependent gates, and the differential's one-directionality (1,487 vs 14) is strong evidence of a systematic blind spot. Still, the headline floor specifically rests on the normative semantics, and a concrete re-audit with an environment-matched reference would settle whether the floor survives. Because this concern is the same one the reader flagged and does not overturn the paper's overall conditional assessment, I keep the verdict unchanged.","tokens_in":16574,"tokens_out":7984,"duration_ms":89527,"concrete_test":"From the 1,043 tolerance-free failures, sample 200 stratified by gate and operator. Re-run each with an environment-matched reference: same declared precision and dtype, FTZ enabled to match the candidate's actual Triton/CUDA compilation, fast-math flags as used, and inputs restricted to the candidate's representable range; keep the other gates unchanged, and pre-commit that if fewer than 90% of the sampled EXC-01/EXC-02 failures persist, the floor is not robust to legitimate semantics. If persistence is instead near 100%, the concern is refuted and the floor stands.","verdict_should_be":"UNCHANGED","load_bearing_attack":"Section 3.1 defines the reference as 'a plain high-precision loop that is defined to be ground truth and refuses to execute in reduced precision.' The tolerance-free floor (1,043/2,638 = 39.5%) is carried mainly by EXC-01 (34.2%) and EXC-02 (9.3%): exact NaN/Inf position/sign and subnormal flush-to-zero behavior must match the fp64 non-FTZ reference exactly, with no tolerance. That is a normative contract, not a mathematical fact. On Blackwell/CUDA, flush-to-zero is a default for single-precision operations, and fast-math compilation legitimately changes NaN propagation; a kernel written for that environment is correct under its platform contract yet fails these gates. The paper's eighteen fairness rules move ambiguous cases toward the candidate, and the four defenses (7/7 positive control on the authors' own kernels, calibration sweep, 98.5% agreement with KernelBench's own correctness code, hand-audit of 31 disputes) show the verifier is internally consistent and not merely a stricter allclose. None of them independently validates the chosen reference semantics as the right contract for the audited systems. The claim that 39.5% are 'broken in a way no tolerance can excuse' therefore holds only if exact fp64/non-FTZ/full-NaN-propagation semantics is the one legitimate correctness contract; the paper asserts this rather than benchmarks it. This is the load-bearing assumption, and it is not settled by the paper's own evidence.","agreement_with_reader":"agree"},"referee_report":{"model":"deepseek-v4-flash","summary":"The paper presents a twelve-gate 'contract-grade' verifier for GPU kernels, operationalizing a prior Kernel Contracts taxonomy, and uses it both to audit 2,638 kernels that the Dr. Kernel/KernelGYM corpus's own harness had accepted and to certify the authors' native Blackwell tcgen05 training backward for the GDN family. The audit reports 62.1% of accepted kernels with at least one contract violation and a tolerance-free floor of 39.5% (1,043/2,638), with a differential showing the standard KernelBench-style check certifying 1,487 kernels the verifier rejects and only 14 in the reverse direction. Four defenses are offered: a 7/7 positive control, a threshold-calibration sweep, 98.5% agreement with KernelBench's own correctness code, and a hand-audit of 31 disputed cases. The native backward is verified against an fp64 oracle with worst relative error about 3.3e-3 (with one disclosed 5.21e-3 arm), is deterministic, and trains five GDN-family variants, though it is slower than the fla Triton baseline. The paper also reports supporting results: six contract-verified Mamba-3 Triton kernels, a reproduction of the #904 TMEM bug, a 1.1B-parameter training demonstration, and a GRPO autotuning study with a negative edit-RL result.","tokens_in":16868,"tokens_out":10109,"duration_ms":94979,"significance":"If the audit's reference semantics is accepted, this is a substantial and carefully defended result. The paper does what the field's loose allclose checks cannot: it makes silent wrongness measurable, and the near-unidirectional disagreement (1,487 vs 14) is a strong signature of a systematic blind spot in the standard acceptance signal. The four defenses are genuinely falsifiable, the red team of nineteen broken kernels is a useful two-sided check, and the paper is unusually transparent about thin margins (ORD-03), idle gates (CMP-02/RES-02), the dv=128 accuracy overrun, and the native kernel's speed deficit. The native tcgen05 GDN backward is a credible systems contribution even without a speed win, and the fp64-oracle verification ladder with machine-exact tile spec-pins gives the correctness claim independent support. My reservation is that the headline 'broken beyond tolerance' floor rests on a normative choice of reference semantics (exact NaN/Inf propagation, non-FTZ subnormal handling, bitwise determinism) that the paper asserts rather than benchmarks; this should be fixed before the headline rates are published as facts.","major_comments":[{"comment":"The tolerance-free floor of 39.5% (1,043/2,638) is carried by EXC-01 (34.2%), EXC-02 (9.3%), and ORD-02 (4.9%), whose pass criteria encode a normative reference semantics and not a mathematical fact. Section 3.1 defines the reference as 'a plain high-precision loop that is defined to be ground truth and refuses to execute in reduced precision,' and EXC-02 requires flush-to-zero handling to match that loop, while EXC-01 requires exact NaN/Inf positions and signs. On CUDA, FTZ is the default for single-precision operations and fast-math compilation legitimately changes NaN propagation, so a kernel written for that platform contract can be correct under its own semantics and still fail both gates. The four defenses in §3.3 show the verifier is internally consistent and not merely a stricter allclose, but none of them independently validates the chosen reference as the only legitimate correctness contract. The 'no tolerance can excuse' claim therefore requires either benchmarking the reference semantics against common CUDA compilation modes (e.g., an FTZ/fast-math variant of the reference) or explicitly scoping the headline floor to the paper's contract rather than to correctness simpliciter.","section":"§3.1, Table 1, Figure 1"},{"comment":"ORD-02's byte-for-byte determinism requirement is likewise a normative contract, not a tolerance-free mathematical property. A reduction using atomics may produce run-to-run results that differ by small rounding amounts yet lie well within every meaningful acceptance tolerance; labeling such a kernel 'broken beyond any tolerance argument' presupposes determinism as a universal requirement. The 126 kernels failing ORD-02 (4.9%) are included in the 1,043 floor, so the floor's interpretation depends on this presupposition as well. The paper should either justify determinism as a correctness contract for all audited systems or move ORD-02 out of the tolerance-free floor.","section":"§3.1, ORD-02 row, Table 4"},{"comment":"The hand-audit of 31 disputed cases classifies 7 as out of scope and says they are 'dropped from the strict floor,' yet the headline floor of 1,043 is reported in Section 3.2 as the automated gate count. The paper should clarify whether those 7 were excluded from the 1,043 and, if so, how the floor was recomputed for the entire corpus; if they were not excluded, the sentence is misleading. More generally, the hand-audit's 16/31 genuine-broken rate is not shown to extrapolate to the load-bearing cell, so the defense does not by itself confirm the 39.5% floor.","section":"§3.3(4), Section 3.2"},{"comment":"The claim that 'the benchmark's own harness agrees (98.5%)' is reported over 1,030 pairs, but the accepted corpus has 2,638 kernels. Please specify how those 1,030 pairs were selected and, if they are a sample, demonstrate representativeness with respect to operator class, gate outcomes, and the load-bearing cell. Without this, the otherwise valuable check that the replica matches KernelBench's own code cannot be extended to the full differential of §3.4.","section":"§3.3(3)"}],"minor_comments":[{"comment":"The abstract and Section 1 state the 39.5% floor as 'broken beyond any tolerance argument' before Section 3.1 defines the normative reference semantics; consider adding a one-sentence caveat at first use so readers do not over-read the floor.","section":"Abstract, Section 1"},{"comment":"The phrase 'thesamebattery in thesameModel/ModelNewconvention' is missing spaces; the manuscript has several similar spacing typos that should be corrected.","section":"§3.3(1)"},{"comment":"The 'pre-registered kill criterion' is mentioned without a pre-registration timestamp or document identifier; if it was pre-registered, cite it, otherwise use 'prespecified'.","section":"§3.5"},{"comment":"The note that printed margins are not exactly reproducible from printed values because of unrounded measurements is likely to confuse; round the threshold and onset columns consistently or include the unrounded values in the supplement.","section":"Appendix B, Table 5"}],"recommendation":"major_revision","confidential_remarks":"The paper is within scope and potentially important. The authors are transparent and the artifact work is impressive, but the normative reference-semantics concern directly affects the headline floor and should be resolved before publication. I recommend major revision rather than rejection; the central claim is defensible with additional benchmarking or careful scoping."},"author_rebuttal":null,"desk_editor":{"model":"deepseek-v4-flash","letter":"The paper does two real things. First, it ships a twelve-gate contract verifier and audits 2,638 kernels that a public system had already accepted as correct, reporting a 39.5% tolerance-free failure floor and 62.1% with at least one violation. Second, it presents a native Blackwell tcgen05 backward for the GDN family, closing the reverse-state fallback gap. Both are genuinely new and neither is a toy: the verifier has a nineteen-kernel red team, derived tolerances, and eighteen fairness rules; the backward is checked against an fp64 oracle and trains five model variants.\n\nCredit where it's due: the four defenses are serious. The 7/7 positive control caught a real bug in the authors' own kernel, the calibration sweep is disclosed honestly (ORD-03 thin at 1.2×), the 98.5% agreement with KernelBench's own correctness code is a strong check, and the hand-audit discards rather than rubber-stamps disputed cases. The differential is one-directional — 1,487 to 14 — which is hard to explain away. The native kernel is scoped carefully: slower than fla by a lot, verified on a narrow envelope, with one accuracy overrun reported rather than hidden.\n\nThe soft spot is the one the stress-test flags, and it's real. The tolerance-free floor is carried mostly by EXC-01 (34.2%) and EXC-02 (9.3%), which require exact NaN/Inf propagation and subnormal flush-to-zero behavior to match an fp64, non-FTZ reference. That is a normative contract, not a mathematical fact. Under common CUDA fast-math and FTZ settings, subnormal behavior and NaN propagation legitimately differ; a kernel written for that environment would fail EXC-02 without being broken on its own terms. The paper asserts the fp64/non-FTZ/full-propagation reference is ground truth but doesn't benchmark that choice against what kernel-generation systems actually target. This makes the \"no tolerance can excuse\" floor contract-dependent, and the exact 39.5% number should be read with that caveat until independent replication pins down the contract.\n\nThat said, this doesn't sink the paper. The stricter contract is defensible for correctness-critical use, the fairness rules move ambiguous cases toward the candidate, and the qualitative conclusion — loose allclose acceptance certifies many genuinely wrong kernels — is robust even if the floor would shrink under a fast-math contract. The paper also discloses its own limitations thoroughly, which is more than most.\n\nRecommendation: yes, send to peer review. It's an honest, well-evidenced paper with a result that should change how benchmark correctness is reported. I'd bring it to reading group and cite it. The revision should address the reference-semantics question directly.","headline":"A serious, well-defended audit with a real artifact, but the tolerance-free floor is contract-dependent because the reference semantics are asserted rather than benchmarked.","tokens_in":17441,"tokens_out":3355,"would_cite":true,"duration_ms":33497,"reading_group":"yes","serious_thinker":"yes","would_accept_peer_review":true},"rs_alignment":null,"lean_confirmation":null,"pith_extraction":{"msc":[],"pacs":[],"model":"deepseek-v4-flash","headline":"This paper claims that the standard one-shot closeness check certifies thousands of LLM-generated GPU kernels that are still broken, and that a twelve-gate tolerance-free verifier exposes the gap.","keywords":["GPU kernel generation","LLM-generated code","correctness verification","tolerance-free contracts","floating-point semantics","gated linear recurrence","Blackwell tcgen05","kernel benchmarks"],"falsifier":"Re-run the tolerance-free gates on the 1,043 failing kernels under a reference that follows the kernel's declared fast-math mode, allowing flush-to-zero subnormal handling and non-propagating NaNs, and then count how many of the 868 EXC-01 and 238 EXC-02 failures still reproduce. If a substantial share pass under that reference, the claimed 39.5% floor is an artifact of the fixed reference semantics rather than a property of the kernels.","tokens_in":16329,"feed_emoji":"⚙️","tokens_out":8347,"duration_ms":71034,"temperature":0.7,"pith_summary":"Correctness reports for LLM-generated GPU kernels rest on a single weak test: a few random inputs at one fixed shape, accepted if the output is close to a reference. This paper builds a twelve-gate verifier that checks properties a correct kernel must have, several of them with no numerical tolerance at all, and runs it over 2,638 kernels that a public system's own harness had already accepted as correct. It finds 39.5% fail a tolerance-free gate and 62.1% carry at least one contract violation, while the standard check certifies 1,487 kernels the verifier rejects against only 14 the other way. The same verifier also certifies the paper's own artifact, a native Blackwell tensor-core training backward for the gated-linear-recurrence family, verified independently against a double-precision oracle, which serves as a positive control that the battery is not tuned to fail foreign code.","feed_headline":"39.5% of accepted AI-written GPU kernels fail zero-tolerance checks","feed_subtitle":"The standard closeness test certifies 1,487 kernels this 12-gate verifier rejects, with only 14 going the other way.","key_machinery":"The load-bearing instrument is a twelve-gate contract verifier that operationalizes a correctness taxonomy for GPU kernels. Several gates are tolerance-free—exact-mask or byte-equality comparisons for non-finite propagation (EXC-01), subnormal flush-to-zero behavior (EXC-02), and run-to-run determinism and buffer aliasing (ORD-02)—so no threshold can excuse a failure; the remaining gates use error-model-derived bounds, such as $\\text{atol} \\approx 4\\varepsilon\\sqrt{N}\\cdot\\text{scale}$ for reordered reductions. On the systems side, the central object is the gated linear recurrence $S_t = (I-k_t(b_t\\odot k_t)^\\top)\\operatorname{Diag}(e^{g_t})S_{t-1} + k_t(w_t\\odot v_t)^\\top$, whose backward splits into two hard stages: a reverse-time inter-chunk state scan and a triangular-inverse vector-Jacobian product. The native kernel keeps the state gradient $dS$ resident in tensor memory across the whole reverse scan, handling the 512-column tensor-memory lifecycle that a known compiler failure trips when a kernel requests 544 columns.","core_discovery":"On the paper's own terms, the central discovery is a measured rigor gap: among kernels a public generation system had accepted as correct, 62.1% violate at least one of twelve correctness contracts and 39.5% fail a tolerance-free gate, meaning no choice of acceptance threshold can explain the failure away. The modal defect is a kernel silently returning an ordinary finite number where the reference returns a NaN or infinity, which turns a detectable training error into silent data corruption. The paper defends the rates with a positive control that passes all seven trusted kernels, a threshold-calibration sweep, 98.5% agreement with the benchmark's own correctness code, and a stratified hand-audit, and it reports a second, native-CUDA corpus showing a weaker but related pattern. Its second claim is that a single native Blackwell tensor-core backward for the gated-linear-recurrence family, including the reverse-state scan stage other implementations still run in a fallback, is correct: it matches a double-precision oracle to about $3.3\\times10^{-3}$ end to end, is bit-for-bit deterministic, and trains five family members through the real kernels.","pith_inferences":["Beyond the paper: if other public kernel-generation corpora share this acceptance signal, their reported speedups should be re-audited with tolerance-free contracts before being relied on in production or downstream training.","Beyond the paper: a benchmark standard that pays a speed bonus only after every contract passes could shift what language models optimize for, since the paper's edit-based reinforcement-learning attempt shows source-level exploration is the weak action space while launch-knob autotuning succeeds.","Beyond the paper: the verifier's tolerance-free gates could serve as a cheap runtime health monitor in training loops—surfacing swallowed NaNs early—rather than only as a post-hoc audit instrument.","Beyond the paper: the observed one-directional disagreement (1,487 to 14) is itself a diagnostic signature; a similarly asymmetric disagreement between a new checker and an established one elsewhere would indicate a systematic blind spot rather than stricter calibration."],"forward_implications":["If the audit is correct, published correctness rates for LLM-generated GPU kernels are substantially overstated: the standard check accepts 1,487 kernels the verifier rejects, versus only 14 in the reverse direction.","A small set of tolerance-free contracts—non-finite propagation, determinism, and shape polymorphism—would close most of the gap, since 958 of the disputed acceptances fail on a tolerance-free gate.","The gated-linear-recurrence family (linear attention, gated linear attention, SSD/Mamba-2, KDA, and gated DeltaNet) can be trained through one native tensor-core backward, verified against a double-precision oracle, rather than through separate implementations.","The same tensor-memory lifecycle error behind a known 38.7x slowdown is detectable by the verifier's resource gate, so future kernel-generation benchmarks can screen for it without waiting for an upstream compiler fix.","Because the verifier passed all seven independently verified kernels and caught a real defect in the authors' own kernel, its acceptances and rejections carry weight for judging both foreign and self-authored kernels."],"supporting_citations":[{"why":"Supplies the 8,920-trajectory Triton kernel corpus from which the 2,638 accepted kernels are drawn and audited.","marker":"[2]"},{"why":"Supplies the standard allclose acceptance test and the reference correctness code used for the differential and the 98.5% agreement check.","marker":"[8]"},{"why":"Supplies the correctness-class taxonomy that the twelve verifier gates operationalize.","marker":"[12]"},{"why":"Documents the tensor-memory budget failure (544 columns requested versus 512) that the resource gate detects and the native backward avoids.","marker":"[11]"},{"why":"Supplies the second, native-CUDA kernel corpus whose weaker but related failure pattern supports cross-corpus robustness.","marker":"[4]"},{"why":"Supplies the Triton linear-attention library that is the speed baseline for the native backward's structural slowdown.","marker":"[17]"},{"why":"Supplies the Mamba-3 recurrence derivations that ground the mathematical appendix and the GDN-family unification.","marker":"[7]"},{"why":"Supplies the earlier claim that loose checks accept broken kernels, which the audit turns into measured rates.","marker":"[9]"}],"fun_headline_variants":["39.5% of AI-written GPU kernels fail zero-tolerance checks","62% of AI-generated GPU kernels flunk 12-gate contract check","Silent NaN failures: 39.5% of 'correct' GPU kernels are broken","Native Blackwell GDN backward: correct, bit-for-bit deterministic","AI kernel 'correctness' hides 39.5% silent failures"],"cache_read_input_tokens":3200,"weakest_assumption_plain":"The audit's headline rates assume that the high-precision reference loops—including their exact NaN/Inf propagation and subnormal flush-to-zero behavior—are the only legitimate contract for every kernel, so a kernel that intentionally targets different but sanctioned IEEE/GPU fast-math semantics would be counted as broken even when it is not.","fun_headline_variants_meta":{"raw":{"variants":["39.5% of AI-written GPU kernels fail zero-tolerance checks","62% of AI-generated GPU kernels flunk 12-gate contract check","Silent NaN failures: 39.5% of 'correct' GPU kernels are broken","Native Blackwell GDN backward: correct, bit-for-bit deterministic","AI kernel 'correctness' hides 39.5% silent failures"]},"model":"deepseek-v4-flash","effort":"low","cost_usd":0.00138,"raw_usage":{"total_tokens":5689,"prompt_tokens":1146,"completion_tokens":4543,"prompt_tokens_details":{"cached_tokens":384},"prompt_cache_hit_tokens":384,"prompt_cache_miss_tokens":762,"completion_tokens_details":{"reasoning_tokens":4444}},"tokens_in":762,"tokens_out":4543,"duration_ms":29095,"temperature":1.0,"reasoning_tokens":4444,"cache_read_input_tokens":384,"cache_creation_input_tokens":0},"cache_creation_input_tokens":0},"created_at":"2026-08-16T04:16:46.648609+00:00","model_set":{"reader":"deepseek-v4-flash"},"falsifier":"Re-run the tolerance-free gates on the 1,043 failing kernels under a reference that follows the kernel's declared fast-math mode, allowing flush-to-zero subnormal handling and non-propagating NaNs, and then count how many of the 868 EXC-01 and 238 EXC-02 failures still reproduce. If a substantial share pass under that reference, the claimed 39.5% floor is an artifact of the fixed reference semantics rather than a property of the kernels.","supporting_citations":[{"cited_title":"Kernel Contracts: A Specification Language for ML Kernel Correctness Across Heterogeneous Silicon","cited_arxiv_id":"2604.22032","evidence_quote":"Supplies the correctness-class taxonomy that the twelve verifier gates operationalize."},{"cited_title":"Issue #904: Mamba-3 SISO backward is 38.7× slower on GB200/B200","cited_arxiv_id":null,"evidence_quote":"Documents the tensor-memory budget failure (544 columns requested versus 512) that the resource gate detects and the native backward avoids."},{"cited_title":"FLA:Atriton-basedlibraryforhardware-efficientimplementations of linear attention mechanisms.https://github.com/fla-org/flash-linear-attenti on, 2024","cited_arxiv_id":null,"evidence_quote":"Supplies the Triton linear-attention library that is the speed baseline for the native backward's structural slowdown."}],"review_version":1}