cs.LGSep 2, 2026

Measuring the Checker: Mutation Analysis for GPU-Kernel Benchmark Oracles

Authors: Mingzhe Du, Anh Tuan Luu, Dong Huang, See-Kiong Ng

Organizations: National University of Singapore · Nanyang Technological University · CAIR, VinUniversity

Read PDF
Ask AI
Cite

Copy citation

DOINot available for this paper

Abstract

Benchmarks for LLM-generated GPU kernels decide correctness with a few random inputs and a loose floating-point tolerance, and their verdicts now feed leaderboards and reinforcement-learning rewards. Recent work agrees these checkers are weak and patches them by hand---extra input distributions, fuzzing recipes, tighter tolerances---with no way to measure whether any patch suffices. We introduce mutation analysis as an adequacy metric for kernel-benchmark oracles: deterministic rules inject 10{,}303 compilable faults into verified CUDA implementations of 188 KernelBench problems, 7{,}384 of them with an independent kill witness; any test protocol is scored by the fraction it detects. The official check misses one in six witnessed faults (16.9%), deterministically, and the misses are skewed by family: 8.7% of arithmetic faults escape, but 78.6% of precision faults do. The metric explains why (a tolerance blind band growing with reduction size; a measured ceiling on input aggressiveness set by legitimate floating-point variance), audits the strongest existing patch (KernelBench-Verified's gain splits into +4.0+4.0 points from hidden inputs and +4.5+4.5 from tighter tolerance, a split its authors could not compute), and exposes a published fuzzing recipe that rejects correct kernels 107 times. Optimizing suites over the kill matrix reaches 98.0% detection with two inputs per problem (94.8% held-out), and the measurement's fault taxonomy teaches a test generator more than the raw faults themselves. Across 48 whole architectures, the blindness grows with scale, concentrating in deep homogeneous pipelines, and two problems prove unrefereeable: their official references violate the benchmark's own tolerance against fp64. We release everything as KernelBench-M.

Figures & tables

Figure 1: The measurement pipeline. Block height is proportional to mutant count. Deterministic rules mutate 188 gate-verified CUDA substrates into 10,303 mutants; compilation, compiled-image hashing, and crash quarantine remove 2,050 (the red strip cut from beneath each block, split by cause), and a kill-witness search over validity-gated inputs admits 7,384 into the scoring denominator, quarantining the 869 that no valid input kills. The witnessed kill matrix (mutants ×\times inputs, filled = detected) records only inputs that pass the validity gate; an input that rejects the correct kernel (red column) is discarded. The matrix then scores any protocol, synthesizes minimal suites, and yields the fault taxonomy.
Figure 2: What the official check misses, and why. (a) Share of witnessed faults missed by the official protocol, by family ( nn at right): textbook arithmetic mutations are mostly caught; precision faults escape at 9×9\times the average rate. (b) The check accepts any error below atol+rtol ∣f∣\texttt{atol}+\texttt{rtol}\,|f| (black line; grey region invisible to it; filled markers: the check fires, open: it accepts). Vacuity: the fault “return all zeros” has error ∣f∣|f| (blue); as the softmax reduction dd grows it slides into the grey region and passes. Validity ceiling: the fp32 accumulation noise between two correct kernels grows with input scale ss ; on sign-mixed inputs the output cancels and the noise climbs out of the region, rejecting a correct kernel ( s≥300s\geq 300 ); on same-sign inputs it stays deep inside (green).
 Operators (levels 1–2)Architectures (level 3)
Problems (substrates)18948
Mutation rules124124
Mutants generated10,30312,589
−- non-compiling1,59438
−- duplicate compiled image279155
−- equivalent to original17735
Table 1: The mutation pipeline of Fig. 1 in numbers, at operator scale (levels 1–2) and architecture scale (level 3). “Distinct” counts mutants surviving compilation, compiled-image deduplication, and equivalence filtering; “witnessed” additionally requires a validity-gated input on which the mutant verifiably fails. The 124 rules split boundary 39, precision 30, semantic 22, synchronization 14, indexing 14, arithmetic 5; 96 fire with a witness (Table 2 ). Substrates total 44,270 lines of CUDA across 360 device kernels; at operator scale a problem carries 37 mutants at the median (199 at most) and 5 survivors (81 at most). The kill matrix comprises over 120,000 mutant–input evaluations, each mutant compiled in 84 ms by NVRTC versus 200 s through the torch extension path.
Figure 3: Population structure of the official check’s blindness at operator scale. (a) Per-problem miss rate, sorted: half the problems lose under 10% of witnessed faults, while a tail of transposed-convolution and reduction problems—exactly the shapes tolerance vacuity predicts—lose 40–73%. (b) Pooled miss rate under bootstrap subsampling of problems: the 90% band contracts onto 16.9%, so the headline is a property of the benchmark, not of our sample.
FamilyRulesWitnessedMissedExample operators
Arithmetic53,3498.7%+ ⁣→ ⁣−+\!\to\!- , ∗ ⁣→ ⁣/*\!\to\!/ , off-by-one constant
Indexing1374914.4%axis swap, stride confusion, transposed access
Semantic1885814.9%tie-breaking, fused-op reordering, flag flips
Boundary271,73522.9%guard deletion, ceil →\to floor grid, tail drop
Synchronization640027.8%barrier removal, __syncwarp weakening
Precision2727178.6%fp16 accumulators, no-stabilizer, fast-math
Table 2: Operator taxonomy and where the official check fails. Witnessed mutants per family and the share the official five-input protocol misses. The gradient is the paper’s central empirical fact: textbook arithmetic mutations are mostly caught, while the implementation-level families real kernel bugs occupy—boundary, synchronization, precision—escape at 3×3\times to 9×9\times the rate.
Figure 4: Protocol audit on the unified denominator (8,215 witnessed mutants across 235 operator- and architecture-level problems). “Ours” is the official dense-random input plus one per-problem input chosen by set cover (§ 6 ). Hardened protocols improve on the official check (dashed line) but plateau below optimized selection at half the input budget; the fuzz baseline’s score comes with 107 false rejections of correct kernels. Fig. 5 dissects the blind band into its miss-patterns.
Figure 5: Anatomy of the blind band: the kill matrix collapsed to its distinct miss-patterns (UpSet form). Each column is one pattern of “which protocols miss this mutant” (filled = missed), its bar the number of operator-scale mutants with that pattern; left margins give each protocol’s total misses (cf. Fig. 4 ). The band’s core is the leftmost column: 643 mutants that every baseline misses and only our targeted suite detects. Each KBV scaling removes a different, partially overlapping slice (columns 2–5); the fuzzer removes more; our suite misses 10 of 7,320 in the residual mixed tail.
 Optimized, budget bbOfficial
Scored onb=1b{=}1b=2b{=}2b=3b{=}3b=5b{=}55 random
Full pool87.1%98.0%99.5%100.0%83.1%
Held-out mutants84.5%94.8%96.4%96.6%82.9%
Table 3: Detection versus input budget on the witnessed level-1/2 pool. Held-out: suites selected on a 50% dev split of each problem’s mutants and scored only on the other half.
Figure 6: Synthesis results. (a) Detection versus input budget: greedy set cover on the kill matrix, scored on the full pool (solid) and on held-out mutants under dev-half selection (dashed). (b) Minimal full-coverage suite size per problem: median two inputs, maximum six. (c) Killing power of generated suites on 1,154 official-input survivors as generator knowledge grows (Table 4 ): the measurement-derived taxonomy beats both ignorance and white-box fault listings.
RungGenerator knowsSurvivors killedValidity
R0 random fuzznothing21.6%107 false kills
R1 naive prompt“tests may miss subtle bugs”42.9%60/60
R2 ++ taxonomy & ceilingfault families, safe magnitudes61.0%94% (accum.)
R3 ++ white-box diffsthe mutated source sites57.3%55/56
Table 4: The knowledge ladder: identical task (kill 1,154 official-input survivors), increasing generator knowledge. Validity = share of generated suites that do not reject the correct kernel. The taxonomy rung accumulated more suites per problem than the single-shot rungs, so its edge over the naive prompt is an upper bound; the ordering is unaffected.
ScaleProblemsDistinctSurvival (distinct)Missed (witnessed)
Operators (levels 1–2)1898,25325.7%16.9%
Architectures (level 3)4812,36143.0%≥\geq 17.3%
Table 5: Operator scale versus architecture scale. Distinct = behaviourally distinct compilable mutants; survival = share not killed by the official inputs; witnessed = mutants with a validity-gated kill witness. The level-3 witness search is far shallower (2–4 targeted suites per problem versus the full matrix at operator scale), so its witnessed miss rate is a conservative lower bound and its distinct-survival an upper bound.
Figure 7: Architecture-scale results (distinct-mutant accounting). (a) Survival by injection-depth quartile across 48 architectures: faults with more network downstream survive more. (b) Per-architecture survival versus pipeline depth: deep homogeneous pipelines (VGG, LSTM, MLP chains) approach opacity, while normalization- and branch-dense networks (ResNet-18, SwinMLP) remain checkable at any depth.
Appendix figures & tables3 assets

Supplementary material from the paper’s appendix.

Appendix

OperatorFamilyMissed / witnessedRateExample site
store-fp16precision169 / 18392%out[i] = (half)acc
mul2plusarithmetic99 / 99010%delta * delta
loop-start-1boundary82 / 41720%for (int i = 0;
mul2divarithmetic78 / 9688%value * norm_weight[c]
loop-bound-minus1boundary77 / 38420%i < hidden_size; i++
plus2minusarithmetic75 / 1,1736%variance + 1.0e-5
Appendix
Table 6: Top mutation operators by officially-missed faults (operator scale, witnessed accounting). “Example site” shows a representative mutated source location.
RuleDescription
Arithmetic (4 rules)
minus2plusAOR: - →\to +
mul2divAOR: * →\to /
negate-storeUOI: negate stored expression
plus2minusAOR: + →\to -
Indexing (14 rules)
Appendix
Table 7: The mutation-rule inventory loaded by the released mutator, by family. [TODO: the released rule files load 120 unique rules (7 definitions duplicate an existing pattern and are skipped); the main text counts 124. Reconcile.]
RungKilledEvaluatedRate
R0 random fuzz2491,15421.6%
R1 naive prompt4951,15442.9%
R2 ++ taxonomy & ceiling7041,15461.0%
R3 ++ white-box diffs6261,09257.3%
Appendix
Table 8: Knowledge ladder, raw counts: survivors killed over survivors evaluated at each rung.
More in this topic

GPU Kernel Correctness Verification

See all 14 papers

Explore similar work

CardsList
Jun 18, 2026

The Correctness Illusion in LLM-Generated GPU Kernels

Benchmarks for LLM-generated GPU kernels (KernelBench, TritonBench, GEAK) score correctness through fixed-shape, small-sample allclose-style checks. The number of inputs varies between benchmarks. The shape, dtype, and tolerance are fixed for each kernel. We test that oracle empirically. We construct a controlled corpus of 24 Triton and CPU stand-in kernels (15 correct controls and 9 LLM-style buggy variants seeded with documented transcription errors) and re-evaluate it under op-schema-aware seeded fuzzing with a high-precision (fp64) CPU reference and per-(op, dtype) absolute tolerances. The seeded oracle flags 9 of 9 buggy kernels and passes 15 of 15 correct controls, at zero precision cost on controls. We extend the corpus to 26 ops (adding a flash-attention pair) and re-run the same protocol on five GPU classes (RTX 3060, A10, L40S, A100 SXM4, H100 NVL). The verdicts are identical across all five GPUs: 10 of 10 illusions caught and 16 of 16 controls clean. The corpus result is about LLM-style transcription bugs that the allclose-on-one-shape oracle certifies as correct, not about the bug rate of any specific deployed LLM. Every flagged failure replays byte-for-byte from a stored seed.

Dipankar Sarkar

Arizona State University, USA

Automated Software TestingDifferential TestingGPU Kernel GenerationGPU Kernel Correctness VerificationCode Generation Evaluation
AskPDF
Oct 5, 2026

RESOLVE: Language-Agnostic Validation of GPU Kernels Through Testing, Reduction, and Proof

AI systems can now write and optimize production GPU kernels, but validating them remains an important challenge. Evaluating the kernel on a few random inputs and checking that its outputs match a trusted reference kernel within numeric tolerances is not sufficient: races can cause nondeterministic behavior that fails to manifest in tests, and numeric tolerances can hide bugs and cause false positives even after extensive calibration. To address this challenge, we present RESOLVE, which combines testing and formal verification to build a comprehensive kernel validation pipeline. It operates in three steps: First, it tests for nondeterminism using binary instrumentation that perturbs execution timing to expose races. Second, an agent rewrites the candidate and reference kernels to obtain "reduced-concurrency" versions that are simpler to analyze but still produce bitwise-identical outputs in all tests. Third, the reduced kernels are formally analyzed in the F*/Pulse framework and prove that they perform the same computation on real numbers. This sidesteps the need for numeric tolerances. We show that RESOLVE can validate a broad selection of kernels using KernelBench, and prove equivalence across fused GEMMs in three state-of-the-art frameworks and languages: CUTLASS, Triton, and Gluon. It also analyzes mega-kernels, notoriously difficult to validate, and finds four previously unreported issues, including two clear bugs. We show that agents can use RESOLVE to repair the issues, with minimal performance impact, highlighting that agents can optimize aggressively when they can rigorously check their results.

Ashkan Vedadi Gargary, Guido Martínez, Sebastian Burckhardt +4

University of California, Riverside Riverside, California, USA · Math, Inc. New York City, New York, USA · Microsoft Research Redmond, Washington, USA

Formal VerificationGPU Kernel Correctness Verification
AskPDF
Aug 13, 2026

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

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, 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.

Rishi Shah, Rishav Shrestha

E3A Healthcare

Automated Software TestingFormal VerificationGPU Kernel GenerationGPU Kernel Correctness Verification
AskPDF