| name | compute-sanitizer |
| description | Run NVIDIA compute-sanitizer (memcheck, racecheck, initcheck, synccheck) against a Triton/TLX kernel to find runtime memory and synchronization bugs. Use when a kernel produces wrong results, crashes with an illegal/misaligned access, or is suspected of a shared-memory data race or invalid barrier usage — especially warp-specialized (WS) kernels using mbarriers, named barriers, TMA copies, or MMA accumulators. This is a runtime check: it runs the real kernel via its reproduce command, so it needs a working GPU and is 10-100x slower than a normal run.
|
Compute Sanitizer
compute-sanitizer is NVIDIA's runtime correctness checker. It instruments the
real kernel launch and reports memory and synchronization errors that static IR
analysis cannot see. Use it to confirm (or rule out) out-of-bounds accesses,
uninitialized reads, shared-memory data races, and illegal barrier usage in
Triton/TLX kernels.
It complements the static investigations: barrier-visualization reasons about
the intended barrier protocol from the IR, while compute-sanitizer observes
what actually happens at runtime. When a race or sync error fires, cross-check
the two.
The four tools
Select with --tool; memcheck is the default.
- memcheck — out-of-bounds and misaligned global/local/shared accesses, plus
device-side allocation/leak errors. First tool to run; catches the classic
TMA-past-end and bad-index bugs.
- racecheck — hazards on shared memory (RAW / WAR / WAW). Highly relevant
to WS producer/consumer kernels that stage data through SMEM buffers. Does NOT
check global-memory races.
- synccheck — illegal use of barriers and synchronization primitives
(divergent
bar.sync, mismatched arrive/wait counts, invalid cluster/async
barrier usage). Directly relevant to mbarrier / named-barrier WS code.
- initcheck — reads of uninitialized global memory. Useful when output
has nondeterministic garbage rather than a hard crash.
Setup
- Locate the binary. It is often not on PATH. Prefer
$COMPUTE_SANITIZER_BIN if set, else command -v compute-sanitizer, else
/usr/local/cuda/bin/compute-sanitizer, else the newest
/usr/local/cuda-*/bin/compute-sanitizer.
- Get the reproduce command (a
pytest/python invocation). Use the
smallest failing shape/config — sanitizer overhead makes large grids
impractical. Prefer a single test id over a whole file.
- Pin a healthy GPU. Run
bash third_party/tlx/find_working_gpu.sh, take a
WORKING_GPUS index, and prepend CUDA_VISIBLE_DEVICES=<idx>. See the
debug-failing-gpu skill. If a run hangs past your timeout, run
third_party/tlx/killgpu.sh.
- Enable source mapping. Triton emits line info by default; keep it on
(do not set
TRITON_DISABLE_LINE_INFO=1) so errors map to file:line. Pass
--show-backtrace yes for host+device backtraces.
Invocation
General form (the sanitizer wraps the whole command, env vars go before it):
CUDA_VISIBLE_DEVICES=<idx> \
<compute-sanitizer> --tool memcheck \
--error-exitcode 1 \
--show-backtrace yes \
--log-file <out_dir>/memcheck.log \
--target-processes all \
python -m pytest -s -x "<test_id>"
Useful flags:
--error-exitcode 1 — return non-zero when any error is found (scripting).--log-file <path> — capture the report (%p expands to PID if needed).--target-processes all — follow child processes (pytest workers, subprocs).--kernel-name-exclude / --kernel-name <regex> — filter to the Triton
kernel (names look like _attn_fwd_..., matmul_kernel, etc.) to cut noise.--launch-timeout <s> — bound a single launch.- memcheck:
--leak-check full, --padding <bytes> (catch off-by-a-few OOB).
- racecheck:
--racecheck-report all (hazards + analysis).
--print-limit 0 — do not truncate the error list while triaging.
Run the tools in order — memcheck → racecheck → synccheck → initcheck — each
into its own log. Stop early only with a stated reason (e.g. memcheck already
found the crashing OOB, or an earlier tool hung).
Triton/TLX specifics
- First run JIT-compiles the kernel (slow, unrelated to sanitizer). Let the
compile complete; do not mistake compile time for a hang.
- racecheck is shared-memory only. WS kernels stage operands through SMEM, so
a missing producer/consumer barrier shows up here as a WAR/RAW hazard on the
SMEM buffer. A global-memory race will NOT be reported.
- synccheck + mbarriers. Invalid named-barrier / mbarrier / cluster-barrier
usage (e.g. wrong arrive count, divergent participation) surfaces here. Pair
with
barrier-visualization to map the offending barrier to a partition.
- TMA / async copies. Out-of-bounds TMA tile accesses crash with an illegal
instruction at runtime rather than masking — memcheck localizes the launch;
also see the
tma-illegal-instruction skill for the structural launcher-bug
pattern.
- Possible noise. Some async/TMA/
cp.async paths may emit benign warnings or
unsupported-feature notes on certain driver/toolkit versions. Note them as
caveats; do not treat a warning as a confirmed bug without corroboration.
Interpreting output
- Each tool ends with
========= ERROR SUMMARY: N errors. 0 errors = clean for
that tool. racecheck prints RACECHECK SUMMARY.
- A real finding includes the access kind (e.g. Invalid global read of size
16), the address, the kernel name, and — with line info — the
file:line.
Capture the first error; later ones are often cascades.
- "Clean" is a result, not a failure. Report it plainly: a clean memcheck +
racecheck + synccheck materially narrows a wrong-result bug toward
compiler/logic rather than memory/sync.
Reporting
Return a compact matrix plus triage:
tool | exit | errors | first finding (kernel @ file:line) | mapped?
memcheck | 0 | 0 | - | n/a
racecheck | 1 | 3 | WAR hazard on smem buf @ k.py:142 | yes
synccheck | 0 | 0 | - | n/a
initcheck | 0 | 0 | - | n/a
Then state the most likely root cause (OOB vs. shared-memory race vs.
uninitialized read vs. illegal sync), the implicated TLX/Triton construct, and
the next action (e.g. narrow with a kernel filter, or hand a race/sync hit to
barrier-visualization).
Do not run performance benchmarks.
