원클릭으로 Manus에서 모든 스킬 실행

sched2tlx-perf-testing

스타178

포크61

업데이트2026년 7월 16일 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.

설치

Codex 또는 Claude로 설치 이 Prompt를 복사해 Codex, Claude 또는 다른 어시스턴트에 붙여 넣으면 Skill 페이지를 검토하고 설치를 진행할 수 있습니다.

Manus에서 실행

출처

facebookexperimental

facebookexperimental/triton

GitHub 저장소 열기 Creator 저장소 보기

다운로드

Manus에서 실행

관련 직업SOC

SOC 직업 분류 기준

소프트웨어 개발자컴퓨터 및 수학직·SOC 15-1252

SKILL.md

readonly

이 저장소의 다른 Skills

같은 저장소

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)

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.