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.

Ir para a instalação

Repositório: facebookexperimental/triton
Última atividade na origem: 12 de fevereiro de 2026 às 16:14
Idioma detectado do SKILL.md: inglês
Estrelas: 187
Forks: 67

Opções de instalação

Por padrão, está selecionado o prompt que primeiro revisa a origem. Você pode mudar para um comando direto ou baixar uma cópia local.

Instalar com Codex ou Claude Copie este prompt, cole no Codex, Claude ou outro assistente e deixe que ele revise a página da skill e instale para você.

Um comando direto ignora o prompt de revisão. Verifique a origem antes de executá-lo.

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

O comando permanece em uma só linha. Role horizontalmente para revisá-lo antes de copiar.

Prefere uma cópia local? Baixe os arquivos disponíveis atualmente no SkillsMP.

Revise os arquivos de origem

Leia o SKILL.md e os arquivos complementares exibidos pelo SkillsMP antes de decidir se vai instalar.

SKILL.md

Instruções da origem · Visualização somente leitura

Voltar à pesquisa Voltar à instalação

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

Ver todas deste repositório

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

Informações da origem

Opções de instalação

Revise os arquivos de origem

Mais deste repositório

IR Debugging

Environment variables

Decision tree: what are you debugging?

Common combos

Reference files

Mais deste repositório

IR Debugging

Environment variables

Decision tree: what are you debugging?

Common combos

Reference files