Kernel 调试
===========

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

提取中间表示
------------

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

默认缓存位置
^^^^^^^^^^^^

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

.. code-block:: bash

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

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

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

编译流水线：

.. code-block:: text

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

启用调试转储
^^^^^^^^^^^^

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

.. code-block:: bash

   export TRITON_DEBUG=1
   python your_script.py

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

- ``kernel.ttir.mlir``
- ``kernel.ttadapter.mlir`` 或 ``kernel.bcmlir.mlir``

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

.. code-block:: bash

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

检查缓存的 IR
^^^^^^^^^^^^^

.. code-block:: bash

   # 查找最新的 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``：

.. code-block:: bash

   # 首先，运行脚本以生成缓存的 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 变换：

.. code-block:: text

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

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

有用的 MLIR 调试标志
^^^^^^^^^^^^^^^^^^^^

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

.. code-block:: bash

   # 在所有 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`` 标志：

.. code-block:: bash

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

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

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

.. code-block:: bash

   bishengir-compile --help

用于 IR 转储的环境变量
^^^^^^^^^^^^^^^^^^^^^^

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

.. code-block:: bash

   # 在每个优化 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``：

.. code-block:: bash

   # 提取缓存的适配器 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 产生不确定或间歇性错误结果
- 想要隔离问题是在同步逻辑还是计算逻辑中

**工作流程：**

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

示例调试会话
------------

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

.. code-block:: bash

   # 步骤 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：

.. code-block:: bash

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

或手动调用 ``bishengir-compile`` 时：

.. code-block:: bash

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

Kernel 参数验证
^^^^^^^^^^^^^^^

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

.. code-block:: python

   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 的调试：

.. code-block:: bash

   export TRITON_INTERPRET=1
   python your_script.py

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

启用设备打印
^^^^^^^^^^^^

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

.. code-block:: bash

   export TRITON_DEVICE_PRINT=1
   python your_script.py

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

常见陷阱
--------

缓存的 Kernel
^^^^^^^^^^^^^

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

.. code-block:: bash

   rm -rf ~/.triton/cache/*

或强制重新编译：

.. code-block:: bash

   export TRITON_ALWAYS_COMPILE=1

非确定性输出
^^^^^^^^^^^^

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

.. code-block:: bash

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

详细日志
^^^^^^^^

启用详细的 Triton 日志记录：

.. code-block:: bash

   export TRITON_DEBUG=1
   python your_script.py

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

.. code-block:: bash

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

报告问题
--------

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

1. **最小复现器** —— 触发问题的最小 kernel
2. **缓存的 IR** —— 附加失败 kernel 的 ``~/.triton/cache/<hash>/`` 内容
3. **编译器输出** —— 使用 ``--mlir-print-ir-after-all`` 或相关环境变量的完整日志
4. **环境信息** —— CANN 版本、Triton-Ascend 版本、NPU 型号、分支（release/3.2.2 vs master）
5. **解决方法状态** —— ``--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>`_

总结
----

.. list-table::
   :header-rows: 1
   :widths: 30 70

   * - 技术
     - 命令
   * - 启用 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`` 诊断同步错误