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 a la instalación

Repositorio: facebookexperimental/triton
Última actividad en el origen: 12 de febrero de 2026 a las 16:14
Idioma detectado de SKILL.md: inglés
Estrellas: 187
Forks: 67

Opciones de instalación

De forma predeterminada está seleccionado el prompt que primero revisa el origen. Puedes cambiar a un comando directo o descargar una copia local.

Instalar con Codex o Claude Copia este prompt, pégalo en Codex, Claude u otro asistente, y deja que revise la página de la skill y la instale por ti.

Un comando directo omite el prompt de revisión. Revisa el origen antes de ejecutarlo.

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

El comando permanece en una sola línea. Desplázate horizontalmente para revisarlo antes de copiarlo.

¿Prefieres una copia local? Descarga los archivos que SkillsMP tiene disponibles ahora.

Revisa los archivos de origen

Lee SKILL.md y los archivos complementarios que muestra SkillsMP antes de decidir si quieres instalarlo.

SKILL.md

Instrucciones de origen · Vista previa de solo lectura

Volver a la búsqueda Volver a la instalación

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 todos los de este repositorio

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

Datos de origen

Opciones de instalación

Revisa los archivos de origen

Más de este repositorio

Más de este repositorio

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