ir-debugging

Debug Triton compilation by dumping IR at each stage (TTIR, TTGIR, LLVM, PTX). Use when investigating compilation failures, kernel performance, register spills, or when user asks to inspect IR output. Covers TRITON_KERNEL_DUMP, MLIR_ENABLE_DUMP, LLVM_IR_ENABLE_DUMP, TRITON_DUMP_PTXAS_LOG, and related env vars.

インストールへ移動

リポジトリ: facebookexperimental/triton
ソースの最終更新活動: 2026年2月12日 16:14
検出された SKILL.md の言語: 英語
スター: 187
フォーク: 67

インストール方法

デフォルトでは、最初にソースを確認する Prompt が選択されています。直接コマンドに切り替えるか、ローカルコピーをダウンロードすることもできます。

Codex または Claude でインストールこの Prompt をコピーして Codex、Claude、または他のアシスタントに貼り付けると、Skill ページを確認してインストールできます。

直接コマンドでは確認用 Prompt が省略されます。実行前にソースを確認してください。

npx skills add https://github.com/facebookexperimental/triton --skill ir-debugging

コマンドは1行のまま表示されます。コピー前に横へスクロールして全体を確認してください。

ローカルで確認しますか？SkillsMP が現在取得できるファイルをダウンロードできます。

ソースファイルを確認

インストールを決める前に、SKILL.md と SkillsMP に表示されている付属ファイルをお読みください。

SKILL.md

ソースの指示 · 読み取り専用プレビュー

検索に戻るインストールに戻る

running-with-buck

How to build and run GPU targets under Buck in fbcode. Use when invoking buck2 run / buck2 build for any GPU benchmark, test, or kernel — selecting the GPU architecture and CUDA version, using @mode/opt and the beta Triton modifier, passing environment variables through, and running from the right directory. Covers the general requirements plus the B200/GB200 (b200a, CUDA >= 12.8) and GB300 (b300a, CUDA >= 13.0) hardware requirements.

2026-08-03

sched2tlx-perf-testing

Run the sched2tlx perf/correctness harness over the modulo-scheduling example corpus (case1-9: GEMM, persistent GEMM, FA fwd/bwd, addmm+bias, LayerNorm, wgrad+bias, multiphase GEMM, scaled_mm). 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-27

barrier-visualization

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-01

このリポジトリのすべてを見る

IR Debugging

Environment variables

Env var

What it does

TRITON_KERNEL_DUMP=1

Dump IR at every compilation stage to ~/.triton/dump/

TRITON_PRINT_AUTOTUNING=1

Use human-readable per-config subdirectories instead of hashes (combine with KERNEL_DUMP)

TRITON_KERNEL_DUMP_BEST_CONFIG=1

Dump IR only for the winning autotuned config (re-compiles with dumping, avoids noise)

MLIR_ENABLE_DUMP=1

Dump MLIR IR during pass execution (filter by kernel: MLIR_ENABLE_DUMP=_kernel)

LLVM_IR_ENABLE_DUMP=1

Dump LLVM IR (print-after-all)

NVPTX_ENABLE_DUMP=1

Dump NVPTX backend IR

TRITON_DUMP_PTXAS_LOG=1

Dump ptxas assembler logs (register usage, spills)

TRITON_INTERPRET=1

Run kernels in interpreter mode (no GPU needed)

TRITON_ALWAYS_COMPILE=1

Bypass cache, force recompilation

TRITON_DUMP_TTGIR_TO_TLX=1

Dump TTGIR back to TLX Python (reverse-engineer IR)

Decision tree: what are you debugging?

"Kernel produces wrong results" → TRITON_INTERPRET=1 to run on CPU, or TRITON_KERNEL_DUMP=1 to inspect IR at each stage

"Kernel is slow / register spills" → TRITON_DUMP_PTXAS_LOG=1 to check register usage and spills

"Which autotuned config won and why?" → TRITON_KERNEL_DUMP_BEST_CONFIG=1 TRITON_PRINT_AUTOTUNING=1

"Need to see MLIR passes" → MLIR_ENABLE_DUMP=1 (optionally filter: MLIR_ENABLE_DUMP=_my_kernel)

"Need to see final PTX/LLVM" → LLVM_IR_ENABLE_DUMP=1 and/or NVPTX_ENABLE_DUMP=1

"Cached result is stale" → TRITON_ALWAYS_COMPILE=1 to force recompilation

Common combos

# Full dump of best config with readable directory names TRITON_KERNEL_DUMP_BEST_CONFIG=1 TRITON_PRINT_AUTOTUNING=1 python my_kernel.py # Debug register pressure TRITON_DUMP_PTXAS_LOG=1 TRITON_ALWAYS_COMPILE=1 python my_kernel.py # Inspect MLIR passes for a specific kernel MLIR_ENABLE_DUMP=_my_kernel TRITON_ALWAYS_COMPILE=1 python my_kernel.py # Full IR pipeline dump TRITON_KERNEL_DUMP=1 TRITON_ALWAYS_COMPILE=1 python my_kernel.py

Reference files

Full Python knobs: python/triton/knobs.py

C++ env vars: include/triton/Tools/Sys/GetEnv.hpp

IR Debugging

Environment variables

Env var

What it does

TRITON_KERNEL_DUMP=1

Dump IR at every compilation stage to ~/.triton/dump/

TRITON_PRINT_AUTOTUNING=1

Use human-readable per-config subdirectories instead of hashes (combine with KERNEL_DUMP)

TRITON_KERNEL_DUMP_BEST_CONFIG=1

Dump IR only for the winning autotuned config (re-compiles with dumping, avoids noise)

MLIR_ENABLE_DUMP=1

Dump MLIR IR during pass execution (filter by kernel: MLIR_ENABLE_DUMP=_kernel)

LLVM_IR_ENABLE_DUMP=1

Dump LLVM IR (print-after-all)

NVPTX_ENABLE_DUMP=1

Dump NVPTX backend IR

TRITON_DUMP_PTXAS_LOG=1

Dump ptxas assembler logs (register usage, spills)

TRITON_INTERPRET=1

Run kernels in interpreter mode (no GPU needed)

TRITON_ALWAYS_COMPILE=1

Bypass cache, force recompilation

TRITON_DUMP_TTGIR_TO_TLX=1

Dump TTGIR back to TLX Python (reverse-engineer IR)

Decision tree: what are you debugging?

"Kernel produces wrong results" → TRITON_INTERPRET=1 to run on CPU, or TRITON_KERNEL_DUMP=1 to inspect IR at each stage

"Kernel is slow / register spills" → TRITON_DUMP_PTXAS_LOG=1 to check register usage and spills

"Which autotuned config won and why?" → TRITON_KERNEL_DUMP_BEST_CONFIG=1 TRITON_PRINT_AUTOTUNING=1

"Need to see MLIR passes" → MLIR_ENABLE_DUMP=1 (optionally filter: MLIR_ENABLE_DUMP=_my_kernel)

"Need to see final PTX/LLVM" → LLVM_IR_ENABLE_DUMP=1 and/or NVPTX_ENABLE_DUMP=1

"Cached result is stale" → TRITON_ALWAYS_COMPILE=1 to force recompilation

Common combos

Reference files

Full Python knobs: python/triton/knobs.py

C++ env vars: include/triton/Tools/Sys/GetEnv.hpp

name	ir-debugging
description	Debug Triton compilation by dumping IR at each stage (TTIR, TTGIR, LLVM, PTX). Use when investigating compilation failures, kernel performance, register spills, or when user asks to inspect IR output. Covers TRITON_KERNEL_DUMP, MLIR_ENABLE_DUMP, LLVM_IR_ENABLE_DUMP, TRITON_DUMP_PTXAS_LOG, and related env vars.

ir-debugging

ソース情報

インストール方法

ソースファイルを確認

このリポジトリの他の Skills

IR Debugging

Environment variables

Decision tree: what are you debugging?

Common combos

Reference files

IR Debugging

Environment variables

Decision tree: what are you debugging?

Common combos

Reference files

このリポジトリの他の Skills