Kernel 调试

本指南涵盖为昇腾 NPU 编译的 Triton kernel 的调试技术，包括提取中间表示（IR）、分析编译器变换以及诊断同步问题。

提取中间表示

Triton 会缓存编译后的 kernel 和中间产物。通过检查缓存目录，您可以检查各个编译阶段的 IR。

默认缓存位置

默认情况下，Triton 将编译后的 kernel 缓存在 ~/.triton/cache/ 中。每个 kernel 存储在以基于 kernel 签名和编译选项的 MD5 哈希命名的子目录中。

# 列出缓存的 kernel
ls ~/.triton/cache/

对于昇腾 NPU 编译，每个 kernel 的缓存目录包含：

kernel.ttir —— 带有分布式原语的高级 Triton IR
kernel.ttadapter （release/3.2.2）或 kernel.bcmlir （master 分支）—— 用作 bishengir-compile 输入的适配器 IR
编译后的二进制文件 —— NPU 可执行文件（.o）

编译流水线：

[Triton Python Kernel]
     ↓ (triton.compile)
[kernel.ttir]
     ↓ (昇腾后端适配)
[kernel.ttadapter]  (release/3.2.2)
或
[kernel.bcmlir]     (master 分支)
     ↓ (bishengir-compile)
[kernel.o]  (NPU 可执行文件)

启用调试转储

要在编译期间转储 IR 文件（在命中缓存时很有用），启用调试模式：

export TRITON_DEBUG=1
python your_script.py

这会在 ~/.triton/dump/ 中创建 IR 转储，包含以下文件：

kernel.ttir.mlir
kernel.ttadapter.mlir 或 kernel.bcmlir.mlir

强制重新编译（绕过缓存）：

export TRITON_ALWAYS_COMPILE=1
export TRITON_DEBUG=1
python your_script.py

检查缓存的 IR

# 查找最新的 kernel 哈希
ls -lt ~/.triton/cache/ | head

# 检查特定 kernel 的 IR
cd ~/.triton/cache/<hash>
cat kernel.ttir        # 高级 Triton IR
cat kernel.ttadapter   # 适配器 IR（release/3.2.2）
# 或
cat kernel.bcmlir      # BCM IR（master 分支）

用于验证：

正确的 kernel 参数和张量形状
预期的分布式原语（wait、notify、symm_at）
内存分配和布局变换

编译期间打印 IR

昇腾后端使用 bishengir-compile 将适配器 IR 降低为 NPU 二进制文件。您可以向 bishengir-compile 传递 MLIR 调试标志以在每个编译器 pass 后打印 IR。

使用调试标志手动调用

从缓存中提取适配器 IR，然后使用调试标志手动调用 bishengir-compile：

# 首先，运行脚本以生成缓存的 IR
export TRITON_DEBUG=1
python your_script.py

# 查找缓存的适配器 IR
cd ~/.triton/cache/<hash>

# 使用 IR 打印手动编译
bishengir-compile kernel.ttadapter --target=Ascend910B3 \
  --mlir-print-ir-after-all \
  --enable-auto-multi-buffer=True \
  --enable-hfusion-compile=true \
  --enable-hivm-compile=true \
  --enable-triton-kernel-compile=true

这会产生详细的输出，显示每个 pass 后的 IR 变换：

// -----// IR Dump After SomePass //----- //
module {
  func.func @kernel(...) {
    ...
  }
}

// -----// IR Dump After AnotherPass //----- //
module {
  func.func @kernel(...) {
    ...
  }
}

有用的 MLIR 调试标志

传递给 bishengir-compile 的常用标志：

# 在所有 pass 后打印 IR
--mlir-print-ir-after-all

# 仅在特定 pass 后打印 IR
--mlir-print-ir-after=pass-name

# 在所有 pass 之前打印 IR
--mlir-print-ir-before-all

# 仅在 IR 更改时打印（减少噪音）
--mlir-print-ir-after-change

# 禁用多线程以获得确定性输出
--mlir-disable-threading

用于调试的其他 bishengir-compile 标志：

# 启用调试信息生成
--enable-debug-info=true

# 在特定 HIVM pass 后打印 IR
--hivm-compile-args=bishengir-print-ir-after=hivm-inject-sync

查看所有可用的 bishengir-compile 选项：

bishengir-compile --help

用于 IR 转储的环境变量

Triton-Ascend 提供环境变量以在编译期间控制 IR 转储：

# 在每个优化 pass 之前转储 MLIR IR
export MLIR_ENABLE_DUMP=1

# 在每个 LLVM 优化之前转储 LLVM IR
export LLVM_IR_ENABLE_DUMP=1

# 在每个编译阶段生成 MLIR 复现器文件
export TRITON_REPRODUCER_PATH=/tmp/reproducer

# 启用详细调试输出
export TRITON_DEBUG=1

# 强制重新编译（绕过缓存）
export TRITON_ALWAYS_COMPILE=1

python your_script.py

注意：如果命中缓存，MLIR_ENABLE_DUMP 可能不起作用。使用 TRITON_ALWAYS_COMPILE=1 强制重新编译，或使用 rm -rf ~/.triton/cache/ 清除缓存。

调试同步问题

分布式 kernel 使用基于数据依赖关系的细粒度同步。编译器计算同步点，但这些计算可能存在错误或边界情况。要隔离编译器同步问题，您可以强制使用全局屏障。

使用 `--enable-hivm-inject-barrier-all-sync`

此标志将细粒度同步替换为每个同步点的全局屏障。

要使用它，手动调用 bishengir-compile：

# 提取缓存的适配器 IR
export TRITON_DEBUG=1
python your_script.py

cd ~/.triton/cache/<hash>

# 使用全局屏障重新编译
bishengir-compile kernel.ttadapter --target=Ascend910B3 \
  --enable-hivm-inject-barrier-all-sync=true \
  --enable-auto-multi-buffer=True \
  --enable-hfusion-compile=true \
  --enable-hivm-compile=true \
  --enable-triton-kernel-compile=true

效果：

每个同步操作都替换为 barrier_all()
所有 PE 在每个同步点等待
性能下降，但如果问题与同步相关，正确性会提高

使用场景：

怀疑存在数据竞争或错误的同步
Kernel 产生不确定或间歇性错误结果
想要隔离问题是在同步逻辑还是计算逻辑中

工作流程：

使用细粒度同步（默认）运行 kernel —— 观察问题
使用 --enable-hivm-inject-barrier-all-sync=true 手动重新编译并替换缓存的二进制文件
如果全局屏障修复了问题，则编译器的同步分析可能存在错误 —— 使用最小复现器报告
如果问题持续存在，则问题在其他地方（计算、内存访问等）

示例调试会话

以下是产生错误结果的 kernel 的完整调试工作流程：

# 步骤 1：启用转储和调试模式
export TRITON_DEBUG=1
export TRITON_ALWAYS_COMPILE=1
python failing_kernel.py

# 步骤 2：检查转储的 IR
ls ~/.triton/dump/
cat ~/.triton/dump/kernel.ttir.mlir
cat ~/.triton/dump/kernel.ttadapter.mlir  # 或 kernel.bcmlir.mlir

# 步骤 3：查找缓存的 kernel 哈希
ls -lt ~/.triton/cache/ | head
cd ~/.triton/cache/<latest_hash>

# 步骤 4：使用详细 IR 打印手动重新编译
bishengir-compile kernel.ttadapter --target=Ascend910B3 \
  --mlir-print-ir-after-all \
  --enable-auto-multi-buffer=True \
  --enable-hfusion-compile=true \
  --enable-hivm-compile=true \
  --enable-triton-kernel-compile=true \
  2>&1 | tee compile_log.txt

# 步骤 5：在日志中搜索特定操作
grep -A 10 "notify" compile_log.txt
grep -A 10 "symm_at" compile_log.txt
grep -A 10 "inject-sync" compile_log.txt

# 步骤 6：尝试使用全局屏障
bishengir-compile kernel.ttadapter --target=Ascend910B3 \
  --enable-hivm-inject-barrier-all-sync=true \
  --enable-auto-multi-buffer=True \
  --enable-hfusion-compile=true \
  --enable-hivm-compile=true \
  --enable-triton-kernel-compile=true \
  -o kernel_fixed.o

# 步骤 7：替换缓存的二进制文件并测试
cp kernel_fixed.o <original_binary_name>
python failing_kernel.py

高级技术

在特定 Pass 转储 IR

要减少输出量，仅在特定 pass 转储 IR：

export MLIR_ENABLE_DUMP=kernel_name  # 仅转储特定 kernel
python your_script.py

或手动调用 bishengir-compile 时：

bishengir-compile kernel.ttadapter --target=Ascend910B3 \
  --mlir-print-ir-after=hivm-inject-sync

Kernel 参数验证

打印运行时 kernel 参数用于调试：

import triton
import triton_dist.language as tdl

@triton.jit
def my_kernel(ptr, size, ...):
    rank = tdl.rank()
    if rank == 0:
        # 使用 device_print 进行运行时调试
        pass  # 在此处插入调试逻辑

使用 `TRITON_INTERPRET`

对于不使用 NPU 执行的基于 CPU 的调试：

export TRITON_INTERPRET=1
python your_script.py

这会在 CPU 上以解释模式运行 kernel，有助于将 kernel 逻辑问题与 NPU 特定行为隔离开来。

启用设备打印

要启用 tl.device_print() 和 tl.static_print() 进行运行时调试：

export TRITON_DEVICE_PRINT=1
python your_script.py

注意：每个线程的设备打印输出的 GM 缓冲区限制为 16 KB。

常见陷阱

缓存的 Kernel

Triton 会缓存编译后的 kernel。更改编译器标志或调试设置后，清除缓存：

rm -rf ~/.triton/cache/*

或强制重新编译：

export TRITON_ALWAYS_COMPILE=1

非确定性输出

多线程编译可能产生非确定性的 IR 排序：

bishengir-compile kernel.ttadapter --mlir-disable-threading ...

详细日志

启用详细的 Triton 日志记录：

export TRITON_DEBUG=1
python your_script.py

用于详细日志记录的其他环境变量：

export TRITON_ENABLE_LLVM_DEBUG=1  # 大量 LLVM CodeGen 日志（非常大）
export MLIR_ENABLE_TIMING=1        # MLIR pass 的时间统计
export LLVM_ENABLE_TIMING=1        # LLVM pass 的时间统计

报告问题

报告编译器错误时，请包括：

最小复现器 —— 触发问题的最小 kernel
缓存的 IR —— 附加失败 kernel 的 ~/.triton/cache/<hash>/ 内容
编译器输出 —— 使用 --mlir-print-ir-after-all 或相关环境变量的完整日志
环境信息 —— CANN 版本、Triton-Ascend 版本、NPU 型号、分支（release/3.2.2 vs master）
解决方法状态 —— --enable-hivm-inject-barrier-all-sync=true 是否修复了问题？

提交问题至：

Triton-distributed-ascend: https://gitcode.com/Ascend/Triton-distributed-ascend/issues
AscendNPU-IR: https://gitcode.com/Ascend/AscendNPU-IR/issues

总结

技术	命令
启用 IR 转储	`export TRITON_DEBUG=1`
强制重新编译	`export TRITON_ALWAYS_COMPILE=1`
转储 MLIR pass	`export MLIR_ENABLE_DUMP=1`
转储 LLVM IR	`export LLVM_IR_ENABLE_DUMP=1`
打印所有 IR pass（手动）	`bishengir-compile ... --mlir-print-ir-after-all`
仅打印更改（手动）	`bishengir-compile ... --mlir-print-ir-after-change`
强制全局屏障（手动）	`bishengir-compile ... --enable-hivm-inject-barrier-all-sync=true`
禁用多线程（手动）	`bishengir-compile ... --mlir-disable-threading`
CPU 解释	`export TRITON_INTERPRET=1`
启用设备打印	`export TRITON_DEVICE_PRINT=1`
详细 LLVM 调试	`export TRITON_ENABLE_LLVM_DEBUG=1`
清除缓存	`rm -rf ~/.triton/cache/*`

关键要点：

IR 文件：kernel.ttir → kernel.ttadapter``（release/3.2.2）或 ``kernel.bcmlir``（master）→ ``kernel.o
缓存位置：~/.triton/cache/
转储位置：~/.triton/dump/``（当 ``TRITON_DEBUG=1 时）
需要手动编译才能使用 bishengir-compile 调试标志
使用 --enable-hivm-inject-barrier-all-sync=true 诊断同步错误