Mutant Replay

Does your test suite really protect your Triton kernel? Plant small bugs, let Nemotron write the missing tests, and accept a test only if running it proves it works.

Inspiration

A Triton kernel can pass its tests and still hide a dropped mask, a shifted loop boundary or a half-precision accumulator, because the tests never exercise that line. Mutation testing finds such gaps, but a surviving mutant leaves you with a question: is my test weak, or is the mutant harmless? We wanted a loop that answers it without trusting a language model's word.

What it does

mutant-gate run kernel.py takes an ordinary Triton file and:

1. Plants mutants: small rewrites such as dropping mask or other=, < to <=, or an fp32 accumulator to fp16.
2. Nemotron writes extra tests (stage 1). It never sees the mutants. A generated test is kept only if it passes your original kernel in a sandbox; if it fails there, the error goes back to Nemotron and it tries again (up to a few attempts).
3. Runs every mutant against the kernel's own check plus that test inside a disposable Nebius sandbox (Triton CPU interpreter, out-of-bounds checks on). Mutants that get through are the survivors.
4. Nemotron targets each survivor (stage 2). It sees one survivor's diff and either writes a killer test or answers EQUIVALENT. A killer test counts only if it passes the original and fails the mutant, checked by running the code.

This is an agentic loop with execution as the judge: Nemotron proposes a test, the sandbox runs it, a failure on the original kernel is fed back for another attempt, and nothing is accepted on the model's say-so.

Two extra checks keep the report honest:

• --compile-check compiles every survivor for CUDA (sm_75, no GPU needed) in the same sandbox and sets aside survivors that are not valid GPU programs.
• --diff-oracle involves no model. It runs each survivor next to the original with the same seeds and compares every tensor after every launch, so the report says "outputs identical on all N launches" or "outputs differ in Rstd (max abs 316); the test passes anyway".

The replay demo (https://alexxony.github.io/triton-mutant-gate/) shows recorded runs: forty planted bugs across four kernels as a grid, where a click on a red square shows the change and the test Nemotron wrote for it. It is a static page, so you can explore without a key.

How we used Nebius and NVIDIA

• NVIDIA Nemotron on Nebius Token Factory writes the tests (stage 1) and the per-survivor verdicts (stage 2). Default model nvidia/nemotron-3-super-120b-a12b; Nano-30B-A3B and Ultra-550b-a55b were also run. Every call checks that the response names the requested model.
• Token Factory made the study repeatable. The original five repeats of all three models used about 1.5M tokens; the Super-only part (343,276 tokens) cost about $0.17 at list price (our calculation, as stated in the repo README).
• Nebius Sandboxes (through contree-sdk) run every test, original and mutant. Generated code never runs on the user's machine.

Results (and what they do not show)

• Forty planted bugs, four tutorial-style kernels, one run per kernel: the kernel's own check plus the stage-1 test Nemotron wrote (accepted on the first try for two kernels, after retries for the other two) caught 31 of 40 (77.5%); the kernel's own check alone was not measured separately, and all verdicts come from the CPU interpreter. Adding Nemotron's verified killer tests for the survivors: 33 of 40 (82.5%). One of those two killer tests targets an fp16-accumulator mutant that does not compile in the T4 environment (Triton 3.6.0; it compiles in the 3.8.0 sandbox image). Of the remaining 7 survivors, Nemotron called 5 equivalent and found no accepted test for 2. This illustrates the loop; it is not a benchmark.
• The Triton CPU interpreter is not a compiler. In a pre-registered check on a Colab T4 (Triton 3.6.0), 3 of the 9 survivors were fp16-accumulator programs that do not compile. They were not test gaps. That finding became --compile-check, which agreed with the T4 run on 43 of 44 programs; the one disagreement was a Triton version difference (3.8.0 in the sandbox image).
• The output oracle was pre-registered before its first run (4 of 5 predictions right). It flagged a gap the model's stage 2 did not close: in layernorm, a < to <= mutant moves one Rstd element by 316 while the test still passes. For the 5 survivors Nemotron called equivalent, no claim was contradicted. "Identical" only means no contradiction on the launches the test makes; it is not a proof.
• Files we did not write. Unmodified tutorials from triton-lang/triton, run with small sandbox shims: 04 ran (7 mutants, 7 caught), 03 ran (13 mutants, 10 caught, 3 survived; compile check skipped for this file), 02 is out of scope (it builds a GPU-compiled object). @triton.autotune files run, but only the first config is exercised in the sandbox.
• ** A compiled run on a small plain file. For autotune_axpy (9 mutants, 2 survivors in the sandbox) we ran the original and every mutant on a Colab T4 (Triton 3.6.0), compiled and in the interpreter, with predictions committed before the run. Both survivors (other= changed on masked-out lanes) also passed compiled, as predicted: consistent with equivalence on this hardware, not a proof. Of the 6 dropped-mask mutants, 4 passed compiled although the sandbox's bounds checks had caught them; 2 were caught. One file, one GPU, one Triton build. Tutorial 03's original needs about 700 s to compile on that T4 (mostly compiling autotune configs), so its 13 mutants take about 2.5 hours; the first attempt lost its Colab VM twice (cause unknown) and finished 8 of 14 programs, so we re-ran the other 6 in a third session with the change pre-registered (Amendment F); that session finished. All 14 programs now have a compiled result, from two sessions on the same Triton build. The 3 sandbox survivors behaved as pre-registered: two other= changes compiled and passed, and the fp16-accumulator one failed to compile. Of the 10 mutants the sandbox had caught, 9 passed compiled (all 7 dropped-mask mutants and both drop_modulo_wrap mutants) and 1 was caught compiled (k_loop_floor_div). Each of those 9 is a mutant a compiled run on this T4 did not flag with one generated test; we have not shown any of them to be equivalent. Raw results and the pre-registration are in experiments/compiled_external/.

What we are careful about

Mutants come from fixed rewrite rules, not real bug reports. "Equivalent" is a model's opinion plus sampled inputs. A model that only read the test and the diff predicted catch/survive correctly 86.9% of the time against a 79.4% always-"caught" baseline, so execution is not shown to be strictly necessary; it matters most for numeric behaviour. Every part of the idea (LLM test generation against mutants, mutation analysis of GPU kernels, equivalent-mutant detection) exists in earlier work (docs/PRIOR_ART.md). This is a combination, not a new method. We did not rank the three Nemotron models, because the data does not support it.

Challenges

• The unmodified file must run in a CPU sandbox first. Real files hit GPU-driver queries, triton.testing benchmarks and missing packages, so the sandbox loads a small shim; the cost is that a file branching on device type takes the CUDA branch with CPU tensors.
• Reproducibility: a generated test that seeds itself with hash(str) made the original unstable across processes (fixed with PYTHONHASHSEED=0), and the sandbox API returned transient HTTP 5xx (retried).
• Keeping claims below the evidence. The pre-registrations record what we guessed before each run, including the wrong guesses.

What's next

Tests against real bug reports instead of fixed rewrite rules, mutants that compile and run on a GPU, and more files from real repositories (so far only tutorials and two small plain files).

Built With

• c++
• performance-profiling
• reverse-engineering
• triton