tilelang-error-fixer

# 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 进程进行排查。 ```bash 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 图： ```python 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` 文件，提供代码修改建议。 **方案 B：AOT counterfactual** - 复用 `examples/gemm_aot` 的 `engine.lower` → `LibraryGenerator` → ctypes 流程；不要维护 手写 Bisheng include/flag 命令。 - 对生成 source 的实验性修改只能放在 disposable 副本中。实验通过只能定位问题层，最终 修复仍应落到产生该 source 的 frontend/pass/codegen，并用原始重现验证。 ### 步骤 5：执行修复与 C++ 重新编译 1. 展示修复方案（C++ 代码 Diff），等待用户确认。 2. 备份涉及的原始 C++ 文件。 3. **强制执行编译命令**：修改 C++ Pass 后，必须在 `build` 目录下执行 `make -j$(nproc)`。 ### 步骤 6：修复验证与回滚 1. 重新执行用户的重现脚本。 2. 对比修复前后的 IR 图（`diff before.tir after.tir`）或目标代码，确认修改符合预期。 3. **失败处理**：如果编译失败或运行不通过，恢复备份代码，清理 `/tmp` 缓存，重新进入步骤 3。 --- ## 输出格式 ### TileLang Pass 诊断与修复报告 ```markdown ## TileLang Pass 分析报告 ### 基本信息 - **触发报错模块**：[如：Segmentation Fault / CombineCV Pass] - **涉及算子/脚本**：`example_flash_attention.py` - **获取的调试信息**：[如：GDB Backtrace 指向 `src/transform/combine_cv.cc:145` / `lower.py` 打印的 IR 图] ### 错误原因分析 - **错误分类**：[如：空指针解引用 / 循环迭代不匹配] - **详细分析**：通过 GDB 堆栈发现，在处理 `T.Pipelined` 节点时，试图访问不存在的 `extent` 属性，导致 Core Dump。 - **定位代码**：`src/transform/cross_core_pipeline.cc: 89行` ### 修复方案建议 (需用户确认) **建议直接修复 Pass C++ 源码** 1. 修改 `src/transform/cross_core_pipeline.cc`，在访问节点前增加类型检查... 2. 代码 Diff: ```cpp // Before auto extent = op->extent.as<IntImmNode>()->value; // After if (const auto* imm = op->extent.as<IntImmNode>()) { auto extent = imm->value; // ... } ``` 3. 修复后将自动执行 `make -j8`。 *(生成逻辑错误可参考 `examples/gemm_aot` 做隔离 AOT 实验；不得把生成 source 当作长期修复。)* ### 执行与验证结果 - 编译状态：✅ 成功 - 测试脚本验证：✅ 通过 - Core Dump 状态：已解决 ```