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

tlx-amd-testing

스타178

포크61

업데이트2026년 6월 24일 06:59

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).

설치

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

Manus에서 실행

출처

facebookexperimental

facebookexperimental/triton

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

다운로드

Manus에서 실행

관련 직업SOC

SOC 직업 분류 기준

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

SKILL.md

readonly

이 저장소의 다른 Skills

같은 저장소

sched2tlx-perf-testing

facebookexperimental/triton

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.

2026-07-16178

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

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

tlx-amd-testing

description

TLX-AMD Tutorial Kernel Testing

AMD tutorial kernels follow the same standardized layout as the NVIDIA (Hopper/Blackwell) reference: one shared correctness file (test_correctness.py, arch-gated) and one perf script per (op, arch). Each kernel is an importable module under third_party/tlx/tutorials/.

Kernel inventory

Kernel module	Op	Correctness test(s)	Perf script	Arch gate
`amd_gemm_warp_pipeline.py`	GEMM	`test_amd_gemm_warp_pipeline`	`test_amd_gemm_perf.py` (`warp_pipeline`)	`is_hip_cdna4`
`amd_gemm_pipelined.py`	GEMM (LDS pipeline)	`test_amd_gemm_pipelined`	`test_amd_gemm_perf.py` (`pipelined`)	`is_hip`
`amd_fa_pipelined.py`	Flash Attention	`test_amd_fa_pipelined`	`test_amd_fa_perf.py` (`simple`, `prefetch`)	`is_hip_cdna4`
`amd_fa_persistent.py`	Flash Attention (persistent)	`test_amd_fa_persistent`, `test_amd_fa_persistent_cross_attention`	`test_amd_fa_perf.py` (`persistent`)	`is_hip_cdna4`
`amd_addmm_glu.py`	addmm + GLU (gated linear unit, not GELU)	`test_amd_addmm_glu`	`test_amd_addmm_glu_perf.py`	`is_hip_cdna4`
`ikbo/ikbo_fa_triton.py`	IKBO Flash Attention	`test_ikbo_fa`	`test_amd_ikbo_fa_perf.py`	none (any HIP/CUDA)
`ikbo/ikbo_lce_triton.py`	IKBO LCE (logit cross-entropy — not attention)	`test_ikbo_lce`	`test_amd_ikbo_lce_perf.py`	none (any HIP/CUDA)
`amd_tdm_gemm_pipelined.py`	GEMM (TDM)	`test_amd_tdm_gemm_pipelined`	—	`is_hip_gfx1250`
`amd_mxfp_gemm_tdm_pipelined.py`	GEMM (MXFP, TDM)	`test_amd_mxfp_gemm_tdm_pipelined`	`test_amd_mxfp_gemm_perf.py`	`is_hip_gfx1250`

gfx950 = CDNA4 = MI350-class (is_hip_cdna4()). gfx1250 is a separate, newer target (is_hip_gfx1250()). On gfx950, the gfx1250-only GEMM tests auto-skip.

Correctness

All AMD correctness lives in the single shared file; tests self-gate via @pytest.mark.skipif, so only the relevant cases run per GPU.

# All AMD + IKBO (gfx1250-only cases auto-skip on gfx950):
pytest third_party/tlx/tutorials/testing/test_correctness.py -v -k "amd or ikbo"

# Whole file — Hopper/Blackwell cases auto-skip on AMD (what CI runs, no -k):
pytest third_party/tlx/tutorials/testing/test_correctness.py -v

-k "amd" alone does not select the IKBO tests (test_ikbo_* has no "amd" in its node id) — use -k "amd or ikbo".

Perf

Never run perf unless explicitly asked. Use the kernel-perf-testing skill for run mechanics. denoise.sh works on AMD (it runs the command with NUMA pinning); its GPU clock/power lock is nvidia-smi-based and is skipped on AMD, so expect slightly higher run-to-run variance. Pick a free GPU with rocm-smi.

CI

.github/workflows/mi350.yml runs on a gfx950 (MI350/CDNA4) runner and mirrors .github/workflows/h100.yml:

mi350-tlx-test — TLX unit tests (python/test/unit/language/test_tlx_*.py)
- the tutorial correctness suite (test_correctness.py). AMD/IKBO run; Hopper/Blackwell and gfx1250 cases auto-skip.
mi350-meta-triton-test — TritonBench perf coverage (perf-regression lives here, not in the perf scripts above).

Nightly failures are filed as issues via report-nightly-failure.yml.

Local run note

After any C++ change (or a stale checkout), the in-tree libtriton.so can lag the Python source and every AMD kernel fails at compile with AttributeError: module '...amd.passes.ttgpuir' has no attribute '<pass>'. Fix: rebuild with make dev-install-llvm. If GPU tests hang, run third_party/tlx/killgpu.sh.