ttsem

Run the tests of your Triton kernels on a machine with no GPU and get the faults a GPU run hides: a write past the end of a tensor, two program instances that touch the same element with at least one of them writing it, an operation with an undefined result. Each one fails its test at the line of your kernel. Memory nobody wrote is filled with NaN instead, so a read of it shows up in the values. ttsem compiles every launch with Triton's own frontend and executes the IR it produces on the CPU.

$ pip install --index-url https://download.pytorch.org/whl/cpu --extra-index-url https://pypi.org/simple \
      'ttsem[triton] @ git+https://github.com/alepot55/ttsem'
$ python -m ttsem.sanitize pytest tests/ -q        # your suite, written for device="cuda"

The first command installs ttsem with Triton 3.8.0 and the CPU build of torch 2.14.0 (Linux, Python 3.12 or 3.13); once ttsem is on PyPI, 'ttsem[triton]' replaces the git requirement. The second runs a pytest suite written for device="cuda", unchanged. On suite_no_mask.py, a test file whose kernel forgets the mask on its store, the size that is a multiple of the block passes and the other one fails at line 15 (from a clone; the last lines of the output):

$ python -m ttsem.sanitize pytest tests/fixtures/agents/suite_no_mask.py -q
...
E   KernelFault: kernel `add_kernel`, launch 1: `suite_no_mask.py:15: tl.store(out_ptr + offs, x + y)` writes 24 element(s) past the end of the tensor passed as `out_ptr`. A GPU run usually hides this (allocations are rounded up). Check the mask of this access against the tensor's real size, for every program id and for sizes that are not a multiple of the block.
...
FAILED tests/fixtures/agents/suite_no_mask.py::test_add[1000] - KernelFault: ...
1 failed, 1 passed in 5.58s

python -m ttsem.sanitize script.py does the same for a script. In CI, the action runs your suite this way on a CPU runner (pin it to @v0.1.0 once that release is tagged):

- uses: alepot55/ttsem@main
  with:
    install: pip install -e .        # your project
    args: tests/kernels -q           # pytest arguments

When ttsem runs the module every compiler pass receives, it names the first pass that changes what a kernel computes. The demo replays the per-pass verdicts recorded on a real GPU launch of triton#11519 (Triton main at e1f944a7a); the recording keeps the verdicts, so nothing is executed. The end of its output:

$ python -m ttsem demo
  launch 0 (kernel), 88 stages
  stages    1-23  mismatch     n_diff=320   inline .. tritongpu-fuse-nested-loops
  stages      24  unsupported  n_diff=-     canonicalize  (arith.ceildivsi)
  stages   25-57  match        n_diff=0     triton-licm .. triton-nvidia-gpu-remove-tmem-tokens
  stages      58  unsupported  n_diff=-     canonicalize  (type: !ttg.async.token)
  stages   59-87  match        n_diff=0     triton-loop-aware-cse .. convert-triton-gpu-to-llvm
  stages      88  unsupported  n_diff=-     initialize-ws-cluster-barriers  (<llvm dialect>)

first pass that changes the meaning: tritongpu-fuse-nested-loops (triton#11519, fixed upstream in #11521)

A stage is the module its pass receives: stage 23 is what tritongpu-fuse-nested-loops receives and stage 24 is what it returns. The device ran the miscompiled kernel, so it disagrees with every stage before the faulty pass. In this recording stage 24, the module canonicalize receives, is unsupported, so the recording alone cannot tell tritongpu-fuse-nested-loops from the canonicalize after it (the demo names the first of the two); only the live run separates them. python -m ttsem demo --live compiles the same program with the Triton 3.8.0 wheel, which still has the bug, and runs every stage here against Triton's interpreter on the source: stages 1-20 match, 21-59 mismatch, and it names the same pass (12.9 s and 13.3 s on the machine of "Limits" below, two runs on 9 Oct 2026). For torch.compile, python -m ttsem demo --torch runs Inductor's Triton code for two PyTorch bugs and, on torch 2.14, stops at the race in each (docs/INDUCTOR.md). For coding agents: docs/AGENTS.md and the skill skills/triton-verify.

How it compares

Each cell is either measured here (the last column says so) or taken from what the tool's authors published (documentation, source code, a talk), linked there. A blank cell is one we could not check.

What ttsem executes is the IR Triton's compiler produces: TTIR for a test run, and under the validator the module every pass receives, down to the LLVM conversion, so the run that checks a kernel can also name the pass that broke it. Triton's interpreter simulates the kernel's Python with numpy instead of compiling it, and Triton-San lowers it through triton-shared, a different pipeline that the program has to select in its code.

Evidence

Against the device. On Triton's own test_core.py (14,235 launches, RTX PRO 6000, 5 Sep 2026), 12,548 launches are bit-exact with the device, 1,484 within the float policy written in docs/SEMANTICS.md, 173 mismatch, 29 are unsupported (28 of them inline PTX) and 1 is undefined (MIN % -1). Most of the 173 are the order in which an atomic returns old values, which depends on the order the lanes arrive in; each class of mismatch is described in docs/SEMANTICS.md. The recordings behind these numbers are not in this repository; docs/RESULTS.md has the method and the other corpora.

TritonBench, judged by value. TritonBench's check compares printed output, and 342 of its 350 tests print nothing. Judged by value with python -m ttsem judge --tritonbench, 245 of the 763 published answers (models of 2024-25) that run and can be judged return wrong values (147) or are unsafe (98: 62 out-of-bounds reads, 23 out-of-bounds writes, 7 races, 6 undefined values), 32%. With no GPU, the public judge redid the whole audit from scratch on 7 Oct in about 2 hours on a 4-vCPU cloud VM (2 jobs, then 4), and found 245 of 767; on 9 Oct the 22 answers with a launch too large for the race check's access log were judged again by this release, which no longer calls such a run verified unless the launch's address ranges rule a race out, and 4 answers verified before are now not_judged (races_unchecked). The private harness of 19 Sep had taken about 75 minutes on a 20-core desktop. The audit of 19-20 Sep 2026, with the judge of that date, ran the same files on an RTX 4070 as a control: of the answers both sides ran and compared, 636 of 637 right or wrong decisions agree; of the 95 answers ttsem stops for a fault and the device decides, 86 crash on the device or return wrong values and 9 return the reference's values. The commands that redo the audit from a clone are in audits/tritonbench; those that judge two public sets of KernelBench answers the same way are in audits/kernelbench.

torch.compile. On torch 2.14, over five families of generated in-place programs (3,292), an RTX 4070 gets 856 wrong, 855 of them only at a large shape. Inductor emitted each program's Triton for the GPU at small or medium shapes (on the same RTX 4070), and ttsem, judging that code with no GPU, flags all 856, plus 90 programs the device does not confirm at those shapes. The judge caps Inductor's block at 16, then 4, so that a small shape spans several program instances; Inductor masks its loads and stores for any power-of-two block, so the cap changes the schedule, not the values, and 16 need not be a size the autotuner would pick. These five families are also where the judge's choices were tuned (the block cap, a second block size, a second input, the medium shapes, not counting a store of the value already there), so they are not a held-out test. On programs drawn at random from a separate generator, with no GPU at all (Inductor's Triton for CPU tensors, judged by ttsem, set against the RTX 4070), ttsem flags 74 of the 88 programs the device gets wrong in seed 1 and 81 of the 92 in seed 2 (82 when the CPU code is scheduled as CUDA code is). Seed 1 also flags 50 programs the device gets right: races in the CPU code that the CUDA code need not have. In seed 2 a rerun of 7 Oct, with the block capped at 16 and then 4, flags 135 programs, so at least 43 of them are programs the device gets right. Seed 1 is not fully held out either: it was drawn before the generator had its two rules (keep a program only if eager agrees with its functionalized form, and drop an update that reads a view of its own destination), and it is where those rules came from. With two fixes loaded into torch, our #198010 (merged) and #185450 (another contributor's fix for index_put_, not merged when this was measured), the three families they address, views, alias and scatter (2,092 programs), give 0 wrong and 0 flagged. The generators of these programs are not in this repository.

Repair from the verdict, a negative result. On 145 TritonBench-T tasks, 67 answers of AutoTriton-8B pass a check that only runs the file and can be judged; 31 of them (46%) are wrong, unsafe or launch no kernel. Given ttsem's message for three more rounds, the model repaired none of the 31. The harness of that experiment is not in this repository.

Upstream bugs from this project

State on 8 October 2026; LEDGER.md has every row with its dates and, for the reported ones, a link to the report that holds its reproducer. The copy of this page on PyPI keeps the states of its release date. PyTorch, since 20 Sep: 18 defects reported (17 issues, one reported with its fix); 7 fixed on main, 6 of them by 5 PRs of ours and one by another contributor; 3 more fixes of ours are open PRs, one of them approved. Triton: 4 bug-fix PRs of ours merged, 7 merged in all (the other 3 are a documentation change and two optimizations). "ours" marks a PR by the author.

The race detectors, the validator, the minimiser and the memory model are in this package. The PyTorch program generator (whose fourth tier adds aliased inputs, integers, freezing and autograd), the OpInfo layout sweep, the program fuzzer, the static witness and the LLVM continuation are companion tools that are not; LEDGER.md describes each.

Questions a skeptic asks

• How many of the listed bugs did ttsem find itself? Going by the "found by" column of the 53 upstream rows of LEDGER.md, 6: pytorch#197829, #198031 and #198033 (the race detector), triton#11752 (the validator), #11738 (the memory model) and #11751 (found while modelling the semantics). 7 more surfaced while Triton's own suites ran under the validator, but fail without it too: triton#11730, the two defects of #11733, #11746, #11747, #11831 and #11859. The other 40 were found by hand or by the companion tools. In 5 of those a tool of this package also took part: the validator named the pass of #11407 and #11519, the level-2 gather check reproduced the boundary of #11600, and the minimiser produced the witnesses of #11601 (with the validator's read set) and #11612.
• Does it raise false alarms on correct code? On KernelBook's Triton code generated by Inductor, 16,858 rows ran and raised 4 reports, all true: three races, all the same Inductor miscompile, and one undefined value, an index_put with duplicate indices (20-21 Sep, at the code's own block sizes; the harness is not in this repository). Liger-Kernel's own kernel tests, unmodified (c5b9f31, 19 Sep, the semantics of that day): 173 passed, 0 faults, 3 failures at tolerances chosen for two GPU runs, here against torch on the CPU; four files with 4096 x 32000 tensors did not finish.
• What is the TritonBench denominator? 7,752 published answers, of which 1,002 ran under ttsem and 763 could be judged by value; the 32% is 245 of those 763. Comparing values on the RTX 4070 alone finds 220 of the 816 answers it ran and compared wrong (27%, audit of 19-20 Sep); ttsem, with no GPU, finds 147 wrong and 98 unsafe of 763 and names the fault and its line (audits/tritonbench).
• How long does a verdict take? About 13 s for a small KernelBench answer on the machine of "Limits" (12.6 s to 14.1 s over two runs each of tests/fixtures/judge/task_relu.py with answer_unsafe.py and with answer_ok.py): two sandboxed processes, the task's reference and the answer. The 7,752 TritonBench answers took about 2 hours under python -m ttsem judge --tritonbench, with the judge of e9b2025, on a 4-vCPU cloud VM (7 Oct, 2 jobs, then 4), and about 75 minutes under the private harness of 19 Sep on a 20-core desktop.
• Which targets and versions? Kernels are compiled for CUDA capability 9.0 (sm_90a in Triton's naming); a script run can pick another with --cc, the judge always uses 9.0. There is no AMD back end; only Triton 3.8.0 and torch 2.14.0 are tested, the versions the triton extra pins.
• Is the calibration a held-out test? No. Triton's test_core.py was also the development set: its first run had 1,008 mismatches and 3,659 errors, and the semantics and the harness were fixed until what was left had a documented cause (docs/RESULTS.md). Two other corpora: the 2,387 runnable programs of three fuzzer families, bit-exact at TTIR and at the final TTGIR, and the 3.8.0 release's own test_core.py on sm_120 (14,979 launches on 12 Sep, no unexplained mismatch; triton#11730 in LEDGER.md). That one is the release branch's copy of the same suite, so it is not independent of the development set either.

Limits

• Linux only. Triton 3.8.0 publishes wheels for Linux (x86_64 and aarch64) only, and the judge's sandbox is bwrap.
• An interpreter, not a GPU. No timing. On a 4-core cloud VM (Intel Xeon at 2.8 GHz), under python -m ttsem.sanitize, a vector add over 16.8 million float32 elements took 20.1 s and 20.4 s and a 1024 x 1024 x 1024 fp16 matmul in 64 x 64 x 32 tiles 12.5 s in both runs (9 Oct 2026; the add took 17.8 s to 20.8 s in four earlier runs on the same VM).
• Past 4,000,000 element accesses in one launch, races are checked by address range only. Only the range of bytes each access touched is kept, up to 2,000,000 ranges per launch: where no program instance writes into a range another instance touches, no race is possible and the launch counts as checked; where the ranges meet, or a launch leaves more ranges than that, the launch still runs and is checked for out-of-bounds accesses, but not for races. A script run ends with the number of launches not checked: of the two launches above, the vector add is ruled race-free by its ranges, while the matmul's 64 x 64 output tiles are not contiguous, so the ranges of neighbouring tiles meet, and the run ends with 1 launch(es) too large to be checked for races. Under python -m ttsem.sanitize pytest the terminal summary prints the same count. Under the judge, a run with such a launch is never verified: it is not_judged (races_unchecked).
• A race needs two program instances on the same data. A block that covers a whole small tensor leaves one instance and nothing to race, so test at sizes several blocks long. For Inductor's code, sanitize.inductor_names(block=16) caps the block (docs/INDUCTOR.md).
• TTIR and TTGIR, not the GPU's code. What happens below the LLVM conversion (LLVM, ptxas, the driver) is not executed. The LLVM-level numbers in docs/RESULTS.md come from a companion interpreter that is not shipped here.
• A verdict holds for the inputs that ran. A kernel right at the tested sizes can be wrong at others.
• The calibration is reported, not reproducible from a clone: the device recordings behind it are not published.
• Out of scope: inline PTX (tt.elementwise_inline_asm), clusters of more than one CTA, and at level 1 warp-specialized partitions that wait on barriers. These and the few ops listed in docs/OVERVIEW.md stop as unsupported, by name, rather than with a guess.

Judging model-written kernels

python -m ttsem judge task.py answer.py judges a KernelBench answer written with Triton on a CPU, by KernelBench's own rule (five seeded trials, allclose at 1e-2), with every launch executed by ttsem. With --tritonbench it judges a TritonBench answer against its task's reference, by value. The verdict is verified (exit 0), wrong (where memory nobody wrote shows up, as NaN in the values), unsafe (an out-of-bounds access, a race, or an operation with an undefined result, at the answer's line), no_kernel (no Triton kernel launched), error, too_slow or not_judged. Each answer is model-written code and runs in its own bwrap sandbox, and the verdict is computed outside it, from the run's raw outputs. wrong versus verified tells whether the values are right and holds against an adversarial answer, since the comparison rests on a reference the answer cannot touch (verified means right values; a kernel launched with the work then done in PyTorch still verifies). unsafe and no_kernel rest on two stop-gaps, a source scan and a run-time integrity check of the whole checker, that block every attack in the test suite but are not a proof: until the answer and the checker run in separate processes (planned for 0.2), use the judge as an auditor and an evaluation tool, and as an RL reward against adversarial optimisation with that caveat. The rules, the batch mode, the classes, the sandbox and what each verdict guarantees are in docs/JUDGE.md.

How it was built

Built by Alessandro Potenza with Claude Code (Anthropic's coding agent) as the engineering agent throughout; the author set the direction and the constraints and reviewed the results. Every PyTorch patch was A/B tested on PyTorch's own suites before it was sent (a first draft of the #198033 fix failed 200 optimizer tests there). On 23 Sep a lerp fix passed its A/B, then the fuzzer under dynamic shapes corrupted the heap through it: the cause was the copy_strided bug above, in stock PyTorch.

Related work

Alive2 (Lopes et al., PLDI 2021) validates LLVM's passes by bounded translation validation with an SMT solver; ttsem executes a tile IR on concrete inputs and compares the results, a different method on a different IR. GPUVerify (Betts et al., OOPSLA 2012) verifies statically that CUDA and OpenCL kernels are free of data races and barrier divergence, and GKLEE (Li et al., PPoPP 2012) explores CUDA programs by concolic execution to find races and to generate tests; both work on the threads of a CUDA kernel (GPUVerify also OpenCL), where ttsem runs Triton's tile IR on the inputs a test gives. The judge applies KernelBench's correctness rule (Ouyang et al., 2025) and runs TritonBench's tasks (Li et al., 2025) as that benchmark runs them, comparing values where it compares printed output. KernelBench-Verified (Zhang et al., 2026) and robust-kbench (Lange et al., 2025) harden KernelBench's check with more inputs. In February 2025 Sakana AI walked back the speedups claimed for its AI CUDA Engineer after the system was found to exploit its evaluation harness (TechCrunch).

Reading further

A report on the semantics, the validator and what they found is in preparation; until then, cite this repository with the metadata in CITATION.cff and the commit you used. MIT licence, see LICENSE.