Run any Skill in Manus with one click

sched2tlx-perf-testing

Stars178

Forks61

UpdatedJuly 16, 2026 at 13:47

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.

Installation

Install with Codex or Claude Copy this prompt, paste it into Codex, Claude, or another assistant, and let it review the skill page and install it for you.

Run Skill in Manus

Source

facebookexperimental

facebookexperimental/triton

View GitHub Repository View Creator Repositories

Download

Run Skill in Manus

Related occupationsSOC

Based on SOC occupation classification

Software DevelopersComputer and Mathematical Occupations·SOC 15-1252

SKILL.md

readonly

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.

More from this repository

same repository

barrier-visualization

facebookexperimental/triton

Produce a structured barrier report for AutoWS (automatic warp specialization) IR. Use when the user wants to visualize, audit, or debug barrier usage across warp-specialized partitions, or when debugging a GPU kernel hang (deadlock). For hangs, first dump IR using the ir-debugging skill, then run this barrier analysis to find the barrier that actually deadlocks -- reasoning with the mbarrier phase model (NOT raw arrive/wait counts, which give false positives), plus missing backward barriers and other synchronization issues. Covers mbarriers, named barriers, tcgen05 commit, TMA-implicit arrives, Aref-based synchronization, and producer/consumer barrier patterns.

2026-07-01178

compute-sanitizer

facebookexperimental/triton

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.

2026-06-26178

kernel-perf-testing

facebookexperimental/triton

Run TLX kernel performance benchmarks on Hopper, Blackwell, and AMD (gfx950/CDNA4, gfx1250) GPUs. Use when user asks to benchmark, profile, or measure performance of any TLX kernel (GEMM, Flash Attention, addmm+GLU, IKBO variants). Handles GPU selection, denoise wrapping (NVIDIA only), and version flags. Never run unless explicitly asked.

2026-06-24178

tlx-amd-testing

facebookexperimental/triton

Test and run TLX-AMD tutorial kernels (gfx950/CDNA4 and gfx1250) and understand their CI. Use when working on AMD TLX tutorial kernels — GEMM (warp-pipeline, LDS-pipelined, TDM, MXFP), Flash Attention (simple, prefetch, persistent), addmm+GLU, or IKBO (FA, LCE) — running their correctness or perf, checking arch gating (gfx950 vs gfx1250), or the MI350 CI workflow. Covers the standardized layout (one correctness file, one perf file per op×arch).

2026-06-24178

autows-authoring

facebookexperimental/triton

Author Triton kernels with automatic warp specialization (AutoWS). Use when writing new AutoWS kernels, adding warp_specialize=True to tl.range loops, choosing tl.range kwargs and JIT options, debugging why WS was not applied, or structuring a kernel to work with both Meta WS and upstream OAI Triton. Covers GEMM and Flash Attention patterns on Hopper and Blackwell.

2026-06-23178

ir-override-ablation

facebookexperimental/triton

Design and run Triton TTGIR debugging ablations using ir_override. Use when reducing a provided or dumped TTGIR, trying user-provided or agent-generated ablation/oblation ideas, updating a test harness around ir_override, or preserving a compile/runtime failure while simplifying IR to expose a fundamental compiler or lowering gap.

2026-06-18178

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)

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

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

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.