tilelang-ascend: Skill for Claude Code

TileLang-Ascend Pass 错误诊断与修复技能

本技能提供 TileLang-Ascend 框架中 Pass 层面错误的完整诊断和修复能力，支持从前端报错、Core Dump 崩溃追溯到后端 C++ Pass 的修复，并集成 AOT 目标代码验证机制。

使用场景

• 遇到 Core Dump 或 Segmentation Fault 等严重崩溃，且无其他明确错误信息
• TileLang 编译报错或生成的 Ascend 目标代码结果不符合预期（精度/性能问题）
• 编写或修改 src/transform/*.cc 下的新 Pass 时出现逻辑异常
• 需要使用 AOT 提前编译模式进行算子调试与倒推

触发机制

当用户输入包含以下关键字时自动触发：

• "Core Dump"、"Segment Fault"、"段错误"
• "TileLang 编译报错"、"Pass 生成代码错误"
• "排查精度问题，分析 IR"、"使用 AOT 调试这个算子"
• "修复 TileLang Pass"

工作流程

步骤 1：环境与状态检查

在进行任何 Pass 调试前，必须确保底层环境正常。

1. 执行环境检查脚本验证 TL_ROOT, PYTHONPATH, ASCEND_HOME_PATH, LD_LIBRARY_PATH 等变量。
2. 确认 C++ 工程是否处于最新编译状态。

步骤 2：多维度错误信息获取 (Dump & Catch)

根据报错类型，采取不同的信息获取策略：

1. 严重崩溃 (Core Dump / Segment Fault):

   • 现象：脚本运行直接崩溃，没有任何 Python 异常堆栈。
   • 动作：使用 GDB 挂载 Python 进程进行排查。

     gdb --args python xxx.py
     # 在 GDB 命令行输入：
     (gdb) run
     # 崩溃后输入 bt 查看堆栈，精确定位是 C++ 源码中哪个函数/Pass 引发了内存违规
     

2. C++ 编译报错 (Build Error):

   • 读取 /tmp/tmp*.cpp (获取最新的生成的失败代码) 定位语法或底层 API 错误。

3. 逻辑/精度异常 (IR Dump):

   • 追踪 Pass 阶段性变换。在 tilelang/engine/lower.py 中的 lower() 函数内，添加 print(mod) 以查看最终 IR 图：

     def lower(...):
         ...
         mod = OptimizeForTarget(mod, target)
         print(mod) # <--- 在此处打印 IR，排查 Pass 优化后的结果
         codegen_mod = device_codegen(mod, target)
     

4. 目标代码分析 (Target Dump):

   • 通过 @tilelang.jit(debug_root_path=...) 或 .get_kernel_source() 抓取最终生成的 C++ 核函数。

步骤 3：核心错误原因分析 (TileLang 专属模式)

分析获取的 IR、目标代码或 GDB 堆栈，重点排查 TileLang-Ascend 特有的雷区：

1. 循环嵌套与次数匹配：对比 Cube 和 Vec 部分的循环 extent，诊断类似 "Vec loop times is not enough" 的报错。
2. 跨核同步点 (Sync Points)：追踪 ascend_auto_set_cross_flag / wait_flag，排查死锁或数据竞争。
3. 工作空间 (Workspace) 访问：检查张量切片范围，确认访问是否越界。
4. 数据依赖与流水线 (Pipeline)：分析 T.Pipelined 的 num_stages，检查读写重叠冲突。

步骤 4：修复方案制定（必要时使用 AOT 隔离）

根据步骤 3 的证据提出最小修复。需要区分 TileLang lowering 与目标代码行为时，可用 AOT 构造隔离实验：

方案 A：直接修复 C++ Pass（适用于 GDB 已明确指出问题代码行的情况）

• 直接定位 src/transform/ 或 src/target/ 下的 .cc 文件，提供代码修改建议。