Skip to content

[FIX] Keep traces from mutating Autotuner/Heuristics and isolate client hooks - #494

Open
mark14wu wants to merge 1 commit into
mainfrom
split/trace-fixes
Open

mark14wu wants to merge 1 commit into
mainfrom
split/trace-fixes

Conversation

@mark14wu

@mark14wu mark14wu commented Oct 4, 2026 •

Copy link
Copy Markdown
Collaborator

Summary

Fixes in the existing trace path; no IR-mode code. This is the bottom of the stack #494 → #479 → #495 → #480 → #482 and can be reviewed and merged on its own.

Bugs fixed (reproduced on main, Triton 3.8)

  • A @triton.heuristics kernel: "Heuristics.run() got multiple values for keyword argument 'warmup'".
  • @triton.autotune over @triton.heuristics: "missing 1 required positional argument" (the constexpr the Heuristics layer adds never reached the kernel).
  • A client that votes for a real compile (the profiler) on a kernel that calls a traced helper: "Unsupported function referenced: <TritonTrace ...>" when the helper is called by name; when it is passed as a constexpr argument, the compile ran the interpreter and left triton.language patched.
  • Every launch of a trace appended the same Launch object to launches.
  • patch_warmup stopped asking clients at the first True vote, and left a bound method in the JITFunction's instance dict.
  • Registering the sanitizer and the race detector on one kernel (two loop overriders) failed after patching the interpreter's ops, and left them patched.
  • One client raising in finalize kept the others from finalizing.

How

TritonTrace no longer assigns to the user's Autotuner / Heuristics. Each layer is shallow-copied over the interpreted function or the JITFunction: an Autotuner copy gets its own config cache, never uses Triton's on-disk autotune cache, has reset_to_zero / restore_value rebound to it and (interpreted) uses the dummy benchmarker. During a voted real compile, the traced helpers the kernel reaches are bound back to their JITFunctions. A warmup scope takes its own gate out on exit, so overlapping scopes can close in any order.

Behaviour change

A TritonTrace called outside any interpreter (neither a traced launch's nor Triton's own under TRITON_INTERPRET=1) raises TypeError instead of silently running the interpreter. Under TRITON_INTERPRET=1 it runs as before.

Testing

  • New tests/unit/test_trace_lifecycle.py: 19 tests; 18 of them fail on main.
  • Triton 3.8, CPU only: the eager test files give the same results as on main (the only failures need installed tile-* scripts); 198 passed at this branch's tip.

…nt hooks

TritonTrace built its interpreted runner by assigning to the user's
Autotuner/Heuristics (runner.fn, _do_bench) and its warmup runner by
deep-copying the result. On Triton 3.8 this crashes three ways on main:

- a @triton.heuristics kernel: "Heuristics.run() got multiple values for
  keyword argument 'warmup'" (the launch's warmup=False reached
  KernelInterface.warmup, which adds warmup=True);
- @triton.autotune over @triton.heuristics: "k() missing 1 required
  positional argument" (the Heuristics layer was replaced by the
  interpreted function, so its constexpr never reached the kernel);
- a client that votes for a real compile (the profiler) on a kernel that
  calls a traced helper, as under the CLI wrappers: the compile got the
  TritonTrace. Called by name, it was rejected by the dependency walk
  that keys the compile ("Unsupported function referenced:
  <TritonTrace ...>"); passed as a constexpr argument, it was called by
  the code generator, which ran the interpreter and left
  triton.language patched.

Every Autotuner/Heuristics layer is now shallow-copied over the
interpreted function or the JITFunction. An Autotuner copy gets its own
config cache, never uses Triton's on-disk autotune cache, has Triton's
reset_to_zero/restore_value hooks rebound to it, and (interpreted) uses
the dummy benchmarker even if the user's autotuner cached a real one. A
Heuristics copy warms up through fn.warmup, so the clients' warmup vote
applies to it. During a voted real compile, the traced helpers the
kernel's code reaches by name are bound back to their JITFunctions, and
traced arguments and parameter defaults are passed as JITFunctions. A
TritonTrace called outside any interpreter (neither a traced launch's
nor Triton's own under TRITON_INTERPRET=1) raises a TypeError instead of
running the interpreter; an untraced kernel run by Triton's interpreter
still calls it as before.

Client hooks:
- every traced launch gets its own Launch; two launches used to append
  the same object to `launches`;
- patch_warmup asks every client before compiling (main stopped at the
  first True vote, yet every client got post_warmup_callback). On exit a
  scope takes its own gate out wherever it sits, so overlapping scopes
  (e.g. on two host threads) may close in any order, and jit_fn ends
  with exactly the warmup it had (main left a bound method in its
  instance dict);
- patch_run undoes partial patching when registering client hooks fails:
  the sanitizer plus the race detector on one kernel refuse with two
  loop overriders and used to leave the interpreter's ops patched;
- ClientManager.finalize finalizes every client even if one raises, then
  re-raises the first failure.
@mark14wu
mark14wu added this pull request to stack #496 October 4, 2026 23:48
@mark14wu mark14wu changed the title split/trace fixes [FIX] Keep traces from mutating Autotuner/Heuristics and isolate client hooks Oct 4, 2026
@github-actions

github-actions Bot commented Oct 4, 2026

Copy link
Copy Markdown

Performance Benchmark

Benchmark main (min) PR (min) Change Samples
gemm 0.104s 0.105s +0.3% 20 / 20
gemm_oob 0.116s 0.117s +0.3% 20 / 20
indirect_load 0.022s 0.022s +0.2% 20 / 20
nested_loop 0.235s 0.235s +0.0% 20 / 20
block_pointer_loop_advance 0.125s 0.125s -0.0% 20 / 20
liger_jsd 0.139s 0.138s -0.1% 20 / 20
flaggems_layernorm 0.393s 0.394s +0.1% 20 / 20
swiglu 0.170s 0.170s +0.0% 20 / 20
cross_entropy 0.969s 0.971s +0.1% 20 / 20
fused_linear_jsd 0.209s 0.209s -0.1% 20 / 20
Total 2.483s 2.485s +0.1% N/A

Iterations: 1 warmup + 20 measured
Samples are shown as main / PR; long pytest benchmarks may use fewer samples.

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant