ワンクリックでManusで任意のスキルを実行

ir-debugging

スター173

フォーク55

更新日2026年2月12日 16:14

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.

インストール

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

Manusで実行

ソース

facebookexperimental

facebookexperimental/triton

GitHub リポジトリを開く Creator のリポジトリを見る

ダウンロード

Manusで実行

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

このリポジトリの他の Skills

同じリポジトリ

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

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-06-22173

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

running-with-buck

facebookexperimental/triton

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-06-16173

debug-failing-gpu

facebookexperimental/triton

Recover from GPU-busy / GPU-unavailable failures. Use when a command (pytest, python, a TLX/Triton kernel run, a benchmark) fails with errors indicating the GPU is busy, out of memory, or unavailable — e.g. "CUDA error: out of memory", "all CUDA-capable devices are busy or unavailable", "CUDA-capable device(s) is/are busy or unavailable", "RuntimeError: No CUDA GPUs are available", "device-side assert", or a hang on the first CUDA call. Runs find_working_gpu.sh to locate a healthy GPU and re-runs the failed command pinned to it via CUDA_VISIBLE_DEVICES.

2026-06-15173

tlx-api-reference

facebookexperimental/triton

TLX DSL API reference for low-level GPU primitives. Use when writing or modifying TLX kernel code that uses barriers (mbarrier, named barriers), memory allocation (local_alloc, SMEM, TMEM), TMA operations, warp specialization (async_tasks, async_task), CLC (cluster launch control), or wgmma instructions. Covers Hopper and Blackwell hardware differences.

2026-06-08173

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

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