企业官网建设流程全解析

TileLang IR Lower Trace：逐 Pass 可视化 TIR 变化，并支持生成代码编辑-重编译工作流

【免费下载链接】tilelangDomain-specific language designed to streamline the development of high-performance GPU/CPU/Accelerators kernels项目地址: https://gitcode.com/GitHub_Trending/ti/tilelang

IR Lower Trace 是 TileLang 内置的零侵入编译调试工具，它透明地捕获编译流水线中每一个Pass 前后的 TIR——包括最终生成 C/CUDA/HIP 源码的 codegen 步骤——并在终端和/或自包含 HTML 页面中渲染人类可读的 diff 报告。本文完整覆盖该工具的环境变量用法、输出产物结构、两层 Python API、HTML 报告能力，以及“编辑生成源码再重新编译”的工作流，并结合 tilelang/tools/lower_trace 的源码实现解释其 per-compilation 会话隔离、增量落盘与三方比较（baseline/working/latest）等底层机制。读完本文后，你可以在不改动任何 TileLang 程序的前提下打开TL_LOWER_TRACE观察完整降级流水线，并学会用lower_trace()/enable()API 做精细控制。

1. 它是做什么的：比 pass_diff 更强的 Pass 观测器

IR Lower Trace 展示 TileLang 的 TIR 随编译 Pass 推进而发生的变化。启用单个环境变量即可，无需修改任何代码；它记录每个 Pass 前后的 IR，并对 codegen 产出的设备源码一并落盘。

它是旧工具 TILELANG_PASS_DIFF 的官方继任者（后者文档开头已标注Superseded，仅为向后兼容保留）。相比旧工具，它新增/强化了六项能力：

• Phase 上下文—— 每个 Pass 都带流水线阶段标签（如pipeline_c、phase1_...），可以判断某 Pass 属于哪个后端阶段；
• Codegen 捕获—— 记录最终的 TIR → 源码降级过程，并把生成的 C/CUDA/HIP 代码写到磁盘供查看或编辑；
• 编辑-重编译工作流—— 直接编辑磁盘上的生成源码后重跑，你的修改会被注入回编译过程（带冲突检测）；
• Per-compilation 隔离—— 每次内核编译拥有独立的记录、Pass 编号、流水线作用域和报告生命周期；并发内核不会写入同一份共享内存 trace；
• 原始.tirdump—— 每个 Pass 前后的 IR 按 phase 与 Pass 索引写到磁盘；
• 崩溃安全的增量 HTML—— 报告在每个 Pass 之后都会刷新落盘，进程崩溃时部分结果依然存活；
• 增强 HTML 报告—— 侧边栏 Pass 导航、状态圆点、phase 页签、j/k键盘导航、Shift+E全局展开、F7手动对齐、明暗主题。

2. 快速上手：用环境变量观察降级流水线

最简单的方式是在运行脚本前设置TL_LOWER_TRACE环境变量：

# HTML 报告（设为 1/on/true/yes 时的默认输出） TL_LOWER_TRACE=1 python3 my_script.py # 仅在终端打印彩色 diff TL_LOWER_TRACE=terminal python3 my_script.py # 终端输出 + HTML 报告同时启用 TL_LOWER_TRACE=both python3 my_script.py # 禁用（默认状态——零开销，不安装任何 patch） python3 my_script.py

启用时机与生效范围。TileLang 在tilelang被导入时读取该设置，并通过工具工厂（tool factory）为未来的编译会话注册捕获器；它不会 patch TVM 的默认PassContext或全局 codegen 函数。这一点与旧版 pass_diff（直接 wrapPass.__call__的全进程 hook）有本质区别。从源码看，环境变量值由 env.get_lower_trace_mode() 解析：0/false/no/off返回禁用，1/true/yes/on/html归一为"html"，terminal/both原样保留，无法识别的真值回退到"html"；TL_LOWER_TRACE_DIR由 env.get_lower_trace_dir() 读取，空值时回退到./tmp/lower_trace_dir。若要改变导入之后的设置，请调用enable()（见第 4 节）。

HTML 模式启用后，脚本目录下会维护一个稳定的符号链接<script_dir>/report.html，始终指向最新一次运行的报告，直接用浏览器打开即可：

# 典型位置 open tmp/lower_trace_dir/my_script/report.html

符号链接的维护逻辑见 LowerTraceSession.update_html_link()：先原子替换report.html指向本次运行的report.html，若文件系统不支持 symlink（如某些 Windows 场景）则退化为文件复制，保证链接可用。

2.1 输出目录结构

一次运行会在TL_LOWER_TRACE_DIR下产生如下布局：

<TL_LOWER_TRACE_DIR>/ └── <script_name>/ # 由 sys.argv[0] 推导，例如 "my_script" ├── report.html # 符号链接 → 最新一次运行的报告 ├── codegen.cpp # 生成的 codegen 源码（可编辑，见第 5 节） ├── codegen.cpp.original # 编辑/重编译工作流的基线快照 ├── codegen.cpp.latest # 最近一次运行的真实 codegen 输出 └── .run_records/ └── run_<YYYYMMDD_HHMMSS_ffffff>_<pid>_<unique-id>/ ├── report.html # 本次运行的完整报告 ├── pipeline_c/ # 每个 phase 一个子目录（示例） │ ├── 00_BindTarget_before.tir │ ├── 00_BindTarget_after.tir │ ├── 01_Simplify_before.tir │ └── 01_Simplify_after.tir ├── phase2_optimize/ # 另一个 phase（示意） │ └── ... ├── codegen/ │ ├── 42_codegen_before.tir │ └── 42_codegen_after.cpp └── unscoped/ # 不在任何 pipeline 窗口内的 Pass └── ...

几个可以从源码确认的实现细节：

• 脚本名取自sys.argv[0]的主文件名（去扩展名，空时回退"kernel"），运行目录名格式为run_<时间戳>_<pid>_<uuid8>，见 LowerTraceSession.ensure_run_dir()；
• 每个 phase 拥有独立子目录（phase 名经 文件名安全化 处理）；Pass 编号全局递增、零填充两位（00_、01_…），排序无歧义；
• codegen 记录的after产物写成.cpp（生成的源码），普通 Pass 两侧均为.tir；落盘实现见 LowerTraceSession.save_raw_files()——注意其中对 IO 失败只打印 WARNING，不会让追踪故障拖垮编译本身。

3. 编程 API 第一层：一次性lower_trace()

当只需要对一条固定的 Pass 链做对比、而不想安装任何全局 hook 时，使用 one-shot API（实现见 lower_trace()）：

from tilelang.tools import lower_trace as lt from tilelang import tvm import tilelang.transform as transform # 对比单个 Pass results = lt.lower_trace(func, transform.Simplify(), mode="terminal") # 对比命名 Pass 链，并写出 HTML 报告 results = lt.lower_trace( func, [ ("Annotate", tvm.tirx.transform.AnnotateDeviceRegions()), ("Split", tvm.tirx.transform.SplitHostDevice()), ("ThreadSync", transform.ThreadSync("shared")), ], mode="both", html_path="my_diff.html", )

返回值是list[dict]，每个 Pass 步一个条目，字段为name、before_script、after_script、diff_lines、insertions、deletions、changed。从源码看还有两个值得注意的行为：

• 异常即报告：某个 Pass 抛错时，该步会被记录为error条目（含捕获到的beforeIR），terminal 模式下先打印[FAILED]段（异常信息 + 失败前 IR）再向上重抛，HTML 中该 Pass 显示为失败状态并附错误框（init.py 异常处理）；
• HTML 写入是“尽力而为”：报告在finally中生成，写报告失败绝不会被当作降低级的真实错误抛出，只在降级成功时才会掩盖问题（best-effort 注释）。

4. 编程 API 第二层：enable()/disable()编译级插桩

要追踪真实内核的完整编译流水线（即环境变量的效果，但以编程方式控制），用enable()（实现见 enable()）：

from tilelang.tools import lower_trace as lt # 为进程内后续编译启用追踪 lt.enable(mode="both") # ... 执行 tilelang.compile() / 内核编译 ...

lt.enable()的三个参数全部可选——mode、trace_dir、codegen_output在省略时回退到TL_LOWER_TRACE/TL_LOWER_TRACE_DIR环境变量（或合理默认值）：

幂等性与快照语义：enable()可安全重复调用。它的取值在编译开始时被快照进独立的LowerTraceSession——enable()实际做的事只是调用register_pass_instrumentation_tool("lower_trace", 工厂)注册一个配置工厂（core.py），因此中途修改配置不会污染一个正在编译的内核。

4.1 何时调用disable()

from tilelang.tools import lower_trace as lt lt.enable(mode="both") # 第一个内核——被追踪 kernel1 = tilelang.compile(func_a) # 第二个内核——自动进入独立会话和报告，同样被追踪 kernel2 = tilelang.compile(func_b) # 可选：禁用之后所有编译的追踪 lt.disable()

注意：若不调用lt.disable()，追踪会对之后的编译保持开启。每份 HTML 报告在其所属编译退出时定稿——包括该编译抛出异常的情况（LowerTraceSession.finish() 在异常路径也会 flush 并打印最终报告路径）。

5. HTML 报告特性

HTML 报告是单一自包含文件（无外部资源），提供：

• 侧边栏：逐 Pass 导航、状态圆点（● 有变更 / ○ 无操作 / ✕ 失败 / ◆ codegen）、+/−行数统计；可折叠、可拖拽调宽；
• Phase 页签：按流水线阶段过滤 Pass；
• Summary 栏：可点击的过滤徽章（changed / failed / codegen）；
• 并排 diff：GitHub 风格配色、词级行内高亮、可折叠上下文（↑↓ Expand按钮展开隐藏的一致行）；
• 键盘导航：j/k在 Pass 间移动，Shift+E展开/折叠全部，F7类 Beyond-Compare 手动对齐；
• 明/暗主题切换，通过localStorage持久化；
• 复制按钮：复制任意 Pass 的 before/after IR；
• 错误框：失败的 Pass 在其崩溃前的 IR 旁展示异常信息。

报告渲染逻辑集中在 tilelang/tools/lower_trace/html.py（约 1800 行，含内嵌 JS/CSS），增量刷新由 flush_html() 驱动。

6. Codegen 源码捕获与编辑-重编译工作流

追踪启用时，最终 codegen 步骤（TIR → C/CUDA/HIP/…）会被拦截，生成源码写入<script_dir>/codegen.cpp（或codegen_output=指定的路径），你可以查看——甚至直接编辑——真正要被编译的代码。

为支持“编辑生成代码后带着修改重跑”，IR Lower Trace 维护三个协作文件：

每次运行时，三方比较（baseline / working copy / 当前 codegen 输出）决定如何处理。决策表的实现即 _resolve_codegen_edit_locked()：

关于.bak文件：CONFLICT与INIT-BACKUP产生的备份目的不同。INIT-BACKUP是在工具接管前保存来路不明的既有codegen.cpp；CONFLICT是在 codegen 变化覆盖用户编辑之前保留用户改动。恢复CONFLICT编辑可用diff codegen.cpp.original.bak codegen.cpp.bak。

6.1 典型工作流

1. 查看—— 用TL_LOWER_TRACE=1跑一次，打开codegen.cpp阅读生成源码；
2. 编辑—— 修改codegen.cpp（例如加printf、调整循环）。不要动.original；
3. 重跑—— 再跑一次。由于codegen.cpp与.original不同而 codegen 输出未变，你会看到PATCHED from …/codegen.cpp，编辑过的源码被编译；
4. 迭代—— 持续编辑、重跑。每次运行都重新注入你的工作副本；
5. 如果 codegen 本身变了（比如你改了 TileLang 程序）——两种结果：
   • 你的编辑恰好与新 codegen 输出一致 →SYNCED（基线前移，你的编辑被保留）；
   • 你的编辑和 codegen 都变了且二者不同 →CONFLICT。工作副本备份到codegen.cpp.bak，旧基线备份到codegen.cpp.original.bak。用diff codegen.cpp.original.bak codegen.cpp.bak找回你的编辑，再将其重新应用到刚重新生成的codegen.cpp上。

6.2 编辑-重编译的后端要求

编辑-重编译工作流要求使用源码编译型执行后端——nvrtc、cython或cutedsl。这些后端使用*_without_compilecodegen FFI 产生仅含源码的模块，然后在运行时经 NVRTC / Cython / CuTeDSL 编译（被编辑过的）源串。从源码看，工具维护了一份 “源码型 FFI”白名单_SOURCE_ONLY_CODEGEN_FFIS，包括target.build.tilelang_cuda_without_compile、target.build.tilelang_cutedsl_without_compile、target.build.tilelang_hip_without_compile、target.build.tilelang_metal_without_compile、target.build.tilelang_c等条目；当命中的 FFI 在该白名单内，run_codegen()会用 TVM 的CSourceModuleCreate以用户编辑的源码重建一个新的CSourceModule返回给下游（_make_patched_source_module()、run_codegen()）。

而默认的tvm_ffi后端在 codegen 阶段就已从 TIR 预编译出设备二进制（PTX/hsaco）。当该后端激活时你编辑codegen.cpp，会看到一条NOTE：你的编辑已被记录进 trace 供 diff 查看，但并未被重新编译（NOTE 与后端提示逻辑——CUDA 目标提示execution_backend='nvrtc'，HIP 目标提示execution_backend='cython'）。要用上编辑-重编译，请切换到源码编译型后端：

# CUDA 目标： tilelang.compile(..., execution_backend="nvrtc") # HIP 目标： tilelang.compile(..., execution_backend="cython")

7. 工作原理：三层协作机制

插桩会话由完整的编译器入口点拥有，例如tilelang.compile()、tilelang.lower()以及分组编译；更底层的 pipeline 与 codegen 组件只是向活动会话贡献事件，自身不会创建独立会话。IR Lower Trace 使用三个协作层：

1. TVMPassInstrument—— before/after 回调捕获str(mod)，计算+/−行数并追加一条 LowerRecord（字段含phase、name、index、before/after_text、changed、add/del_lines、status、error_msg、depth、parent_index）。共享的嵌套回调栈记录深度与父子关系，包括被其他 TVM Pass 内部调用的 Pass；不在任何 pipeline 窗口内的 Pass 被打上unscoped阶段标签。观测端实现见 _LowerTraceObserver：pass_started保存 before 文本，pass_finished生成记录并在 HTML 模式触发增量 flush；未完成的 Pass 走passes_incomplete路径记为 FAILED。
2. 显式PassPipeline.lower作用域—— 通过 pipeline_scope() 向后端流水线中的 Pass 提供上下文本地阶段标签（如pipeline_c），贡献给调用方拥有的会话，而不是 patch 类本身。这是正常的 pipeline 边界，不是 monkey-patch；多次流水线运行会得到run2_前缀以避免阶段名冲突。
3. 显式后端 codegen hook—— TileLang 的设备/主机 codegen 注册表将每个target.build.*调用路由到当前编译会话，从而捕获最终 TIR → 源码降级并驱动三文件编辑-重编译工作流，而不替换任何进程级 FFI。没有活动会话时，注册表直接调用底层 codegen（零开销路径，run_codegen 短路）。

此外还有几条关键设计决策：

• 运行时追加记录（而非预注册）：运行时被跳过的条件 Pass——例如should_force_let_inline()为False时的LetInline——干脆不出现在报告中，不会留下幽灵/跳过占位；
• 增量刷新 HTML：每个 Pass 后都刷一次（总成本 O(n)），即使SIGKILL或部分崩溃，已有结果仍留在磁盘上；
• 多 PassContext 场景：当一次逻辑编译使用多个 PassContext（例如分组 lowering 之后接设备与主机 codegen），它们从同一编译会话获取新的回调对象，记录共享该会话内一个单调递增索引；不同的编译从零索引、新记录列表开始（对应第 2 节目录结构中的 per-compilation 隔离）。

注意：enable()只影响 TileLang 管理的编译会话，绝不修改 enable 时刻碰巧处于当前状态的 PassContext。对于独立的 TVM PassContext，lt.create_pass_instrument()仍可用于手动挂接（实现）。另外，捕获 IR 以及 HTML 模式下每个 Pass 后重写报告会引入调试开销——常规构建与基准测试请保持该功能关闭。

8. 如何选择输出模式与实用技巧

模式选择：

• 少量 Pass、需要即时反馈时用terminal；
• 降级流水线产生大量变更、或需要完整对比 before/after 脚本时用html；
• 交互迭代但需要留存可分享报告时用both。

IR Lower Trace 比较的是IRModule.script()返回的文本形式。文本变化是某个 Pass 行为的有用证据，但本身不足以证明发生了语义或性能变化。

技巧清单：

• 快速检查用terminal模式—— 彩色 diff 随 Pass 执行实时打印；
• 深入分析用html模式—— 跨大量 Pass 导航、展开隐藏上下文、复制 IR 片段；
• 配合TL_LOWER_TRACE_DIR将报告指到特定位置，例如 CI 中运行或跨运行对比；
• hook 捕获降级流水线中的全部 Pass，包括由tilelang.compile()内部触发的 Pass，适合理解完整编译流程；
• 如果之前用过TILELANG_PASS_DIFF，请切换到TL_LOWER_TRACE—— 它是严格超集，也是后续功能改进的承载工具（旧文档 docs/tools/pass_diff.md 已标注被取代，仅为向后兼容保留）。

9. 参考

• 原始文档：docs/tools/lower_trace.md
• 工具实现：tilelang/tools/lower_trace/__init__.py（one-shot API）、tilelang/tools/lower_trace/core.py（会话/观测器/三文件工作流）、tilelang/tools/lower_trace/diff.py、tilelang/tools/lower_trace/html.py
• 插桩基础设施：tilelang/instrumentation
• 环境变量解析：tilelang/env.py
• 被取代的旧工具：docs/tools/pass_diff.md

【免费下载链接】tilelangDomain-specific language designed to streamline the development of high-performance GPU/CPU/Accelerators kernels项目地址: https://gitcode.com/GitHub_Trending/ti/tilelang