Triton调试：Triton调试从入门到裂开再到起飞：一套工具链吃透MLIR全流程
写Triton kernel最崩溃的不是算法写不对——是编译出来结果错了、精度炸了、性能崩了，而你盯着满屏Python traceback一脸懵：“IR到底在哪一层出了问题？”
这事我踩过太多坑了。跑一个matmul，autotune选出来的配置是快了，但结果错的。加了个tl.dot，CUDA报错misaligned address。不是说Triton是"写起来像Python、跑起来像CUDA"吗？那调试为什么像拆盲盒？
别急。本文基于杜玉博老师的分享，系统梳理Triton调试的完整工具链——从标准的PDB/GDB，到官方环境变量三板斧，再到MLIR层面的triton-opt和triton-translate，让你在每一层IR都能精准定位问题。
一、先别急着掏MLIR：Python/C++层面的基础调试
很多人一上来就想dump IR，但往往问题出在Python层或是C++ runtime层。基础工具能解决的，就别上大炮。
1.1 PDB：Python kernel调试
PDB是Python内置的调试器，不需要安装任何东西。当你怀疑是kernel参数传错、grid配置不对、或者自动微分出了问题，PDB就是第一选择。
常用PDB命令：
| 命令 | 作用 |
|---|---|
b <行号> | 在指定行设置断点 |
c | 继续执行到下一个断点 |
n | 逐行执行（不进入函数内部） |
s | 步入函数内部 |
p <变量名> | 打印变量值 |
l | 查看当前行附近的代码 |
q | 退出调试 |
用法很简单，在你想打断点的地方插入一行：
importpdb;pdb.set_trace()或者启动时直接：
python-mpdb your_triton_script.py1.2 GDB：当问题沉到C++层
Triton编译流程走到后面全是C++——MLIR的Pass、LLVM的codegen、ptxas的汇编。这时候PDB就管不到了，得上GDB。
关键差别记牢了：
- PDB：用Python写的kernel逻辑调试，断点打在
@triton.jit装饰的函数里 - GDB：调试Triton编译器的C++代码，断点打在
lib/Conversion/、lib/Dialect/等源码路径
用GDB调试Triton程序的启动方式：
# Triton程序本质是python脚本，所以要这样起gdb-argspython your_script.py进GDB后：
(gdb) b TritonToTritonGPUPass.cpp:123 # 在某个Pass的特定行打断点 (gdb) r # 运行 (gdb) bt # 崩溃了？看调用栈 (gdb) p *someValue # 查看变量二、理解Triton编译流水线：你不知道数据流经了哪些层，怎么定位？
调试之前必须搞清楚Triton代码到底走过了哪些阶段。否则dump出来的IR你看都看不懂。
Triton的编译架构分三大块：
Frontend Optimizer Backend ┌─────────┐ ┌──────────────┐ ┌──────────────┐ │ Python │ → │ Triton IR │ → │ LLVM IR │ │ Kernel │ │ → TTGIR │ │ → PTX │ │ Code │ │ (各种Pass) │ │ → cubin │ └─────────┘ └──────────────┘ └──────────────┘- Frontend：把你写的
@triton.jitPython代码转成Triton IR（.ttir），同时负责kernel launch的runtime逻辑 - Optimizer：通过一系列MLIR Pass把Triton IR逐步转换为TritonGPU IR（
.ttgir）——这里会做shared memory分配、warp调度、coalesce等优化 - Backend：把TritonGPU IR转成LLVM IR（
.llir），最后通过ptxas编译成GPU能执行的cubin
问题往往出在Optimizer和Backend的交界处：你的算法逻辑是对的，但某个Pass把memory layout搞乱了，或者thread placement不合理。下面就是对付这种场景的武器。
三、环境变量三连：不写一行代码就能dump全流程IR
Triton官方提供了几个环境变量，设了就自动dump IR。不需要改任何代码。
3.1MLIR_ENABLE_DUMP=1
这个可能是最常用的调试手段。设了这个变量后，Triton会在每个MLIR Pass前后自动把IR打印到终端。
MLIR_ENABLE_DUMP=1python your_triton_script.py你能看到类似这样的输出流：
// *** IR Dump Before TritonToTritonGPUPass *** tt.func @kernel(...) { ... } // *** IR Dump After TritonToTritonGPUPass *** ttg.func @kernel(...) { ... }⚠️ 一个坑：如果设了环境变量但没看到输出，先清理triton cache：
rm-rf~/.triton/cache3.2LLVM_IR_ENABLE_DUMP=1
类似上面，但在LLVM IR层面打印每个pass前后的IR。当你怀疑是LLVM codegen出了问题（比如地址计算错误、向量化异常），用这个。
LLVM_IR_ENABLE_DUMP=1python your_triton_script.py3.3TRITON_PRINT_AUTOTUNING=1
这个不是用来debug correctness的，而是用于性能调优。autotune选出来一堆配置，想知道哪个最优、耗时多少？
TRITON_PRINT_AUTOTUNING=1python your_triton_script.py输出类似：
Triton Autotuning: best config = {'BLOCK_M': 128, 'BLOCK_N': 128, ...}, time = 12.3 us四、triton-opt：逐Pass拆解IR变换的显微镜
环境变量dump的是全量流水线，信息太多容易看花眼。当你想精细控制某个特定Pass的输入输出时，就要用triton-opt。
triton-opt是Triton源码编译后产出的工具，需要先把它的路径加到环境变量里。它的核心用法是：接收一个IR文件，跑指定的Pass，输出变换后的IR文件。
基本用法
# .ttir → .ttgir（Triton IR 降级到 TritonGPU IR）triton-opt XX.ttir -convert-triton-to-tritongpu&>XX.ttgir# 在 .ttgir 层面执行特定优化triton-opt XX.ttgir --tritongpu-coalesce --tritongpu-pipeline&>XX-opt.ttgir配合GDB使用才是完全体：
gdb triton-opt(gdb)b TritonToTritonGPUPass.cpp:123(gdb)r XX.ttir -convert-triton-to-tritongpu&>XX.ttgir这样你可以在特定Pass的C++源码中打断点，配合输入IR和输出IR做对比分析。
实战场景：分阶段lowering定位bug
这是最有效的方法——不要一次跑完整个pipeline，分阶段来：
# 第一步：Triton IR → TritonGPU IRtriton-opt model.ttir -convert-triton-to-tritongpu-omodel.ttgir# 第二步：TritonGPU IR → LLVM IRtriton-opt model.ttgir -convert-tritongpu-to-llvm-omodel.llir哪一步输出的IR不对，问题就在那一步。
五、triton-translate：从IR到PTX的最后一公里
triton-translate的职责是把MLIR转换成最终的可执行产物——PTX或LLVM IR。它在pipeline中处于最末端。
# ttgir → llirtriton-translate XX.ttgir--target=llvmir&>XX.llir# ttgir → ptx（如果需要直接看PTX汇编）triton-translate XX.ttgir --mlir-to-ptx-oXX.ptx同样配合GDB：
gdb triton-translate(gdb)b SomeBackendPass.cpp:456(gdb)r XX.ttgir--target=llvmir&>XX.llirtriton-translate和triton-opt的关系简单说：
triton-opt：管IR的变换（优化Pass）triton-translate：管IR到目标代码的翻译
六、实战推荐组合
根据问题类型，我整理了以下调试策略：
| 问题现象 | 首选工具 | 策略 |
|---|---|---|
| Python层报错、参数问题 | PDB | import pdb; pdb.set_trace() |
| 编译时crash、C++ segfault | GDB + MLIR_ENABLE_DUMP | 先设环境变量看crash在哪个pass，然后GDB打断点 |
| 结果数值不对 | MLIR_ENABLE_DUMP | dump全量IR，找关键op（tt.dot、tt.load）的数值 |
| 性能不达预期 | triton-opt 分阶段 | 看ttgir层的shared memory layout和coalesce |
| PTX生成异常 | triton-translate | 直接看生成的PTX是否符合预期 |
| autotune选配置不满意 | TRITON_PRINT_AUTOTUNING | 打印所有候选配置及耗时 |
终极组合技：
# 先dump全量IR看问题在哪一层MLIR_ENABLE_DUMP=1python your_script.py2>&1|grep-A50"error\|warning\|misaligned"# 定位到具体pass后用triton-opt精细调试triton-opt model.ttir -convert-triton-to-tritongpu -mlir-print-ir-after-all&>debug.log# 在关键pass打断点gdb triton-opt(gdb)b problematic_pass.cpp:xxx(gdb)r model.ttir -convert-triton-to-tritongpu七、总结
Triton调试的核心思路就一句话：分层定位，逐级缩小。
不要一上来就盯着MLIR看——先从Python层排除参数问题、再从autotune排除配置问题、然后用环境变量dump IR定位到具体Pass、最后用triton-opt/triton-translate精细拆解。
工具都在这里了，下次跑Triton kernel出问题的时候，别再只会print了。
本文整理自杜玉博老师的《Triton调试方法及工具》分享，结合个人实战经验补充了大量操作细节和避坑指南。
官方文档参考：Triton - Tips for Hacking