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.

Jump to install

Repository: facebookexperimental/triton
Last source activity: February 12, 2026 at 16:14
Detected SKILL.md language: English
Stars: 187
Forks: 67

Install options

The review-first prompt is selected by default. You can switch to a direct command or download a local copy.

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.

A direct command skips the review prompt. Inspect the source before running it.

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

The command stays on one line. Scroll horizontally to inspect it before copying.

Prefer a local copy? Download the files currently available to SkillsMP.

Review the source files

Read SKILL.md and any companion files shown by SkillsMP before deciding whether to install.

SKILL.md

Source instructions · Read-only preview

Back to search Back to installation

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

View all from this repository

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

Source facts

Install options

Review the source files

More from this repository

More from this repository

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