Elfsong/KernelBench-M

KernelBench-M

The measurement artifact for Measuring the Checker: Mutation Analysis for GPU-Kernel Benchmark Oracles: the mutation operators, the verified CUDA substrates they mutate, the kill witnesses, and the pipeline that produced every number in the paper.

Layout

rules/          124 mutation rules, six families (mutator.py loads all of them)
substrates/     208 gate-verified CUDA implementations, one per KernelBench
                problem: the mutation targets. Never used as oracle.
witnesses/      per-problem kill records: for every mutant, its rule, family,
                mutated site, whether the official protocol kills it, and — when
                it survives — a suite that does
pipeline/       the measurement code (screening, kill matrix, audit, set cover)
summary.json    per-problem mutant / witnessed / official-kill counts

Supporting records: mechanism_curves.json (Fig. 2b–c), holdout.json (dev/test split), ladder_summary.json (knowledge ladder), gate_report.json (substrate admission), invalid_suites.json (suites rejected for crossing the validity ceiling).

What a substrate is, and is not

KernelBench ships PyTorch references whose GPU execution bottoms out in closed cuDNN/cuBLAS binaries, so there is no source to mutate. A substrate is a correct CUDA implementation of the same computation, used only as a mutation target. The oracle stays the benchmark's own PyTorch reference. Each substrate is admitted by pipeline/gate_candidates.py, which compares it against that reference on every suite before it may enter the pool (gate_report.json records admissions and rejections).

Witness format

"softmax:barrier-drop:3": {
  "rule": "barrier-drop", "family": "sync",
  "site": "__syncthreads();",
  "cls": "high_value",
  "killed_by_official": false,
  "witness": "T2_misaligned_batch"
}

cls is killed_by_T0 (the official inputs detect it), high_value (survives the official inputs, a targeted input detects it), or no_witness (nothing we have detects it — an equivalence or oracle-blindness candidate, excluded from every denominator in the paper).

Reproducing

Requires an NVIDIA GPU with CUDA 12.x and PyTorch. From pipeline/:

python screen_mutants.py     # generate, filter, and screen mutants
python full_matrix.py        # kill matrix over all suites
python audit_matrix.py       # score competing protocols (§5)
python analyze_cover.py      # set-cover suite synthesis (§6)
python analyze_holdout.py    # dev/test holdout (§6)

Each stage writes append-only JSONL journals and resumes from them, so a crashed or preempted run can simply be restarted. SHARD_ID/SHARD_N split problems across GPUs; STEAL=1 lets an idle worker pick up unclaimed problems.

Caveats

Mutation score is adequacy relative to this fault model, not absolute correctness: these operators cannot express faults living in structures the substrates do not contain (tensor-core paths, double-buffered pipelines), nor multi-site interactions. The artifact supports comparative claims between protocols and existence claims about specific faults. It cannot certify that a kernel passing these suites is correct.