sched2tlx-perf-testing

name	sched2tlx-perf-testing
description	Run the sched2tlx perf/correctness harness over the modulo-scheduling example corpus (case1-8: GEMM, persistent GEMM, FA fwd/bwd, addmm+bias, LayerNorm, wgrad+bias, multiphase GEMM). Use when the user asks to benchmark generated-vs-handwritten kernels, check corpus correctness, compare emitter revisions, or regenerate schedule_graph.json fixtures. Never run perf unless explicitly asked.
disable-model-invocation	true

sched2tlx Perf & Correctness Harness

Never run performance tests unless the user explicitly asks.

Harness: third_party/tlx/tools/sched2tlx/examples/testing/perf_regression/perf_harness.py Corpus: third_party/tlx/tools/sched2tlx/examples/case*/

Environment (this cluster — critical)

The login shell's lmod modules break both build and runtime. Prefix EVERY python/triton-opt invocation with env -u LD_LIBRARY_PATH and use the repo venv python ($REPO/.venv/bin/python). Ignore the lua/posix noise every command prints. C++ changes require rebuild first: env -u LD_LIBRARY_PATH PATH="$REPO/.venv/bin:$HOME/.local/bin:/usr/local/bin:/usr/bin:/bin" VIRTUAL_ENV=$REPO/.venv CC=/usr/bin/gcc CXX=/usr/bin/g++ MAX_JOBS=14 uv pip install -e . --no-build-isolation

Subcommands

Run from examples/testing/perf_regression/:

• bench — committed generated.py vs handwritten.py, per-case table (shape | gen | hw | gen/hw | rel | ok; correctness is checked before timing). Auto-discovers cases with a bench_spec.py (currently case1,2,3,5,6,7): env -u LD_LIBRARY_PATH $REPO/.venv/bin/python perf_harness.py bench [--cases case6_layernorm,...]
• compare — one row per case, two columns: --rev's committed generated.py (default origin/main) vs the working tree's, cells = per-shape gen/handwritten ratios. Scans EVERY case*/ in examples/ (cases without bench_spec are listed as such); byte-identical kernels are measured once and shown "unchanged". Both columns run under the CURRENT build — it compares fixtures, not toolchains (use regression/e2e for that): ... perf_harness.py compare [--rev origin/main] [--cases ...]
• regression — before/after emitter check for a diff that touches sched2tlx (re-emits every case both sides; only true regressions fail): ... perf_harness.py regression [--diff REV | --before X --after Y] [--perf-tol 0.05]
• e2e — full pipeline from TTGIR (dump → emit → bench) with the current build: ... perf_harness.py e2e --skip-master --cases case6_layernorm. The master-comparison mode (--master-repo) needs a separately built master checkout and is buck/fbsource-oriented.

Cases the harness does NOT cover

• case4_FA_bwd: cd examples/case4_FA_bwd && ... python perf_generated.py (gen vs no-WS baseline vs handwritten WS); correctness via run_generated.py.
• case8_multiphase_gemm: cd examples/case8_multiphase_gemm && ... python bench_general.py (correctness + pool-vs-sum A/B; no handwritten reference).
• Correctness-only for any case: cd examples/<case> && ... python run_generated.py.

Comparing against another revision's kernel

run_bench(case_dir, generated_path) takes the generated.py path separately, so to bench another revision's kernel under identical methodology: git show <rev>:<path>/generated.py > /tmp/gen.py, then in python: perf_harness._print_table(perf_harness.run_bench(CASE_DIR, Path("/tmp/gen.py"))). Cheap shortcut when only fixtures differ: kernels whose generated.py is byte-identical across revisions need no re-measurement.

Regenerating fixtures

TRITON_MODULO_DUMP_SCHEDULE=<case>/schedule_graph.json \
  build/cmake.*/bin/triton-opt -allow-unregistered-dialect \
  --nvgpu-modulo-schedule <case>/<kernel>_pre_modulo.ttgir -o /dev/null
env -u LD_LIBRARY_PATH PYTHONPATH=third_party/tlx/tools/sched2tlx \
  $REPO/.venv/bin/python -m sched2tlx <case>/schedule_graph.json -o <case>/generated.py

JSON op ids are pointer-derived and never byte-stable — regen always churns schedule_graph.json; the meaningful diff signal is generated.py. Known: case3 may need TRITON_MODULO_SELECT_VARIANT=2; case2 fixtures are ancient (fresh dumps differ, pre-existing).

Methodology warnings

• bench/run_bench use triton.testing.do_bench (cold-L2). Ad-hoc CUDA-event loops (e.g. perf_generated_vs_handwritten.py) are hot-L2 and read higher. NEVER cross-compare numbers from the two.
• Check nvidia-smi first; if a run hangs for minutes, run third_party/tlx/killgpu.sh.
• One bench at a time — timing runs must not share the GPU.