Taichi 编译器开发实用指南:跨 Python/C++ 调试与 IR 分阶段打印
【免费下载链接】taichiProductive, portable, and performant GPU programming in Python.项目地址: https://gitcode.com/GitHub_Trending/ta/taichi
导读
本文是面向 Taichi 编译器开发者的实战指南,聚焦 Taichi 源码开发中最常用的调试手段:理解一条 kernel 从 Python 装饰器到后端机器码的完整编译流程,并在各个阶段打印 IR(中间表示)进行验证。读完本文,你将掌握ti.init()中 6 个 IR 打印参数的准确含义与适用场景,了解这些参数在编译流水线中的底层作用位置,并具备一套可复用的跨 Python/C++ 代码导航与编译调试工作流。
本文内容以仓库文档 development_tips.md 为骨架展开,并结合 taichi/program/compile_config.h、taichi/python/export_lang.cpp、python/taichi/lang/misc.py 等源码进行印证。
前置条件:先完成开发安装
本文中的一切调试技巧都要求你以开发者模式从源码构建 Taichi,而不是使用pip install taichi安装的发布版本。在继续之前,请确保已经完整走通 开发安装指南 中的流程:
- 克隆仓库(注意
--recursive,LLVM、spdlog 等外部依赖以 submodule 形式引入); - 通过
./build.py --shell进入带完整构建环境的 shell; - 执行
python3 setup.py develop以符号链接方式安装,保证 Python 侧修改即时生效。
dev_install.md还详细列出了TAICHI_CMAKE_ARGS中诸如TI_WITH_CUDA、TI_WITH_VULKAN等后端开关(默认值可见其参数表)。如果你要调试 CUDA 后端的汇编输出,需要在构建时确保 CUDA 后端已开启。
Taichi 编译器工作流速览:一条 kernel 的一生
在深入 IR 打印之前,先建立全局认知。Life of a Taichi Kernel 一文完整描述了 kernel 的五个生命周期阶段:
- Kernel 注册:执行
@ti.kernel装饰器时,函数体 AST 被记录下来,此时不做任何编译; - 模板实例化与缓存:首次调用时实例化,相同模板签名(如
add(x, 42)与add(x, 1)签名同为(x, ti.i32))复用已编译产物,ti.template()参数变化才触发重新实例化; - Python AST 变换:
ASTTransformer将函数体 AST 改写为可执行、能产出 Taichi 前端 IR 的 Python 脚本; - Taichi IR 编译、优化与可执行文件生成:前端 IR 降级为层级 SSA IR,经历循环向量化、类型推断、公共子表达式消除(CSE)、死指令消除(DIE)、常量折叠、store forwarding、访问降级、数据访问优化、自动微分、并行化与 offload、原子操作 demote 等一系列 pass,最终由 LLVM 或 Metal/OpenGL 等后端编译器生成可执行代码;
- Kernel 启动:以多线程 CPU 任务或 GPU kernel 形式执行。
理解这个流水线后,IR 打印参数的用途就一目了然:它们分别对应第 2~4 阶段的不同产物。其中第 4 阶段的 IR pass 过程主要由print_ir控制,其打印动作在 taichi/transforms/compile_to_offloads.cpp 中通过make_pass_printer为每个 pass 安装,真正实现"每个优化 pass 前后各打印一份 IR"的效果。
语言标准约定:C++17 与 Python 3.7+
参与 Taichi 编译器开发时,请遵守以下语言标准约定(原文档明确要求):
- C++ 部分:整个编译器以C++17编写,你可以放心假设 C++17 特性(如
std::optional、结构化绑定、if constexpr等)始终可用; - Python 部分:以Python 3.7+为目标,同样可以直接使用 3.7 引入的语言特性(如 dataclass、
from __future__ import annotations等)。
这一约定在 dev_install.md 的前置条件表中也有呼应:构建要求 Clang++ >= 10(推荐 Clang++ 15),并建议libstdc++-10-dev或直接安装g++。遵守语言标准不仅能避免构建问题,也让你的贡献与既有代码风格一致。
高效的 Python/C++ 跨语言代码导航
Taichi 的 Python 前端(python/taichi/lang)与 C++ 编译器核心(taichi)通过 pybind11 绑定相连。如果你在开发语言前端(Python/C++ 接口),经常需要在"Python 调用处 → C++ 绑定定义 → C++ 实现"之间跳转。
原文档推荐使用ffi-navigator这一工具:它允许你从 Python 绑定直接跳转到对应的 C++ 定义。安装与编辑器配置请按其项目 README 进行。典型场景包括:
- 在 python/taichi/lang 中看到一个 Python API,想定位其底层 C++ 实现;
- 在 C++ 侧修改了某个内部接口后,反向确认 Python 侧有哪些调用点。
一个可参考的对照入口是 taichi/python/export_lang.cpp,它用 pybind11 的def_readwrite把CompileConfig的编译配置字段逐一导出到 Python,正是"Python 配置 → C++ 配置对象"这一桥接关系的最直观体现。
用 ti.init 参数打印各阶段 IR(核心)
当你使用ti.init(arch=desired_arch, **kwargs)创建 Taichi 程序时,向kwargs传入以下参数即可让编译器打印不同阶段的 IR:
| 参数 | 作用 | 默认值 | 打印内容 |
|---|---|---|---|
print_ir=True | 打印 kernel(不含 accessor)编译的 Taichi IR 变换过程 | false | 每个优化 pass 前后的 SSA IR |
print_accessor_ir=True | 打印数据访问器(data accessor)的 IR 变换过程 | false | accessor 这种特殊简单 kernel 的 IR |
print_struct_llvm_ir=True | 保存 Taichi struct 编译器发射的 LLVM IR | false | struct 编译(如 SNode 布局相关)生成的 LLVM IR |
print_kernel_llvm_ir=True | 保存 Taichi kernel 编译器发射的 LLVM IR | false | kernel 生成的未经优化的 LLVM IR |
print_kernel_llvm_ir_optimized=True | 保存每个 kernel 优化后的 LLVM IR | false | 经优化 pass 处理后的 LLVM IR |
print_kernel_asm=True | 保存每个 kernel 的汇编代码(文档标注 CUDA only) | false | 后端最终生成的汇编 |
这些字段的真实定义与默认值可以在 taichi/program/compile_config.h(print_ir、print_accessor_ir)与 compile_config.h 的 LLVM 相关字段(print_struct_llvm_ir、print_kernel_llvm_ir、print_kernel_llvm_ir_optimized、print_kernel_asm)中查到,全部在 compile_config.cpp 中初始化为false。
下面逐个深入说明。
print_ir:观察整个 IR pass 流水线
print_ir=True是最常用、信息量最大的开关。它在kernel 编译阶段(不包含 accessor)打印整个 IR 变换过程:
- 前端 IR 被降级为层级 SSA IR 之后,每个 pass(类型推断、CSE、DIE、常量折叠、访问降级、offload 等)执行前后都会输出一份 IR 快照;
- 打印动作由 taichi/transforms/compile_to_offloads.cpp 的
make_pass_printer驱动,传入verbose(即print_ir)与print_ir_dbg_info两个配置; - LLVM 后端编译入口 taichi/codegen/llvm/kernel_compiler.cpp 同样读取
compile_config.print_ir控制输出;SPIR-V 后端则在 taichi/codegen/spirv/kernel_compiler.cpp 中传入。
典型用法:
import taichi as ti ti.init(arch=ti.cpu, print_ir=True) x = ti.field(dtype=ti.f32, shape=128) @ti.kernel def add(field: ti.template(), delta: ti.i32): for i in field: field[i] += delta add(x, 42)终端会随之输出从初始 IR 到各 pass 变换后的完整流水线记录,可用于定位某个优化是否生效、类型推断是否正确、offload 是否按预期切分。
print_accessor_ir:调试数据访问器
print_accessor_ir=True打印数据访问器(data accessor)的 IR 变换过程,原文档特别说明它"很少使用",只在调试数据访问器的编译时才需要。原因在于:
- Python 作用域中的
x[1, 2, 3] = 3实际上调用的是x的写入 accessor kernel; print(y[42])调用的则是y的读取 accessor kernel。
这些 accessor 是"特殊且简单的 kernel",其编译路径与普通 kernel 不同,因此独立出一个开关。在 Python 侧,其底层入口可以追溯到 python/taichi/lang/_ndarray.py 中NdarrayHostAccessor的getter/setter,它们最终触发对应的访问 kernel 执行。
三档 LLVM IR:struct / kernel / 优化后
print_struct_llvm_ir、print_kernel_llvm_ir、print_kernel_llvm_ir_optimized三个参数分别控制 LLVM IR 的三个产出时机:
- struct 编译器:负责 SNode 结构布局相关代码的生成,taichi/codegen/llvm/struct_llvm.cpp 中读取
print_struct_llvm_ir决定是否 dump; - kernel 编译器:负责 kernel 主体代码生成,taichi/codegen/llvm/codegen_llvm.cpp 读取
print_kernel_llvm_ir; - 优化后:kernel 经过 LLVM 优化管线后的 IR,taichi/codegen/cpu/codegen_cpu.cpp 与 DX12 后端的 dx12_global_optimize_module.cpp 均会读取
print_kernel_llvm_ir_optimized。
对比同一 kernel 的"发射 IR"与"优化后 IR",可以直观看到 LLVM 层做了哪些向量化、内联与指令合并。
print_kernel_asm:查看最终汇编
print_kernel_asm=True保存每个 kernel 生成的汇编代码。原文档标注该能力CUDA only,用于检查 GPU 侧最终指令质量(如是否产生不必要的局部内存访问、是否发生寄存器溢出)。从源码结构看,taichi/codegen/cpu/codegen_cpu.cpp 同样存在对print_kernel_asm的处理,因此实际可用性请以你构建时启用的后端与版本为准;调试 CUDA 时它是确认后端发射质量的第一手资料。
这些参数是如何生效的:配置项流转链路
理解配置项从 Python 到 C++ 的完整链路,有助于你将来新增类似调试开关:
- C++ 定义:字段声明于 compile_config.h,默认值初始化于 compile_config.cpp;
- pybind11 导出:taichi/python/export_lang.cpp 用
def_readwrite将字段暴露为 Python 属性,这是ti.init(**kwargs)能直接写print_ir=True的基础; - 前端注入:python/taichi/lang/misc.py 的
ti.init文档明确列出print_ir等高频 kwargs,并在 misc.py 中通过_EnvironmentConfigurator(kwargs, cfg)将CompileConfig的全部字段注册到配置解析器——这意味着一方面kwargs可以直接覆盖任意编译配置,另一方面这些配置也保留了通过环境变量注入的通道(具体前缀规则可查阅_EnvironmentConfigurator的实现); - 各后端消费:编译流水线中的
compile_to_offloads.cpp、LLVM/CPU/CUDA/DX12/SPIR-V 各 codegen 在对应阶段读取字段决定是否输出。
实战组合:一套推荐的调试会话
把上文参数组合起来,可以得到一个完整的"前端到汇编"调试配置:
import taichi as ti ti.init( arch=ti.cuda, # 或 ti.cpu debug=True, # 开启边界检查等调试辅助 print_ir=True, # 查看 Taichi IR pass 全过程 print_kernel_llvm_ir=True, # 查看发射的 LLVM IR print_kernel_llvm_ir_optimized=True, # 查看优化后的 LLVM IR # print_kernel_asm=True, # CUDA 下查看最终汇编 ) @ti.kernel def saxpy(x: ti.template(), y: ti.template(), a: ti.f32): for i in x: y[i] = a * x[i] + y[i] x = ti.field(dtype=ti.f32, shape=1024) y = ti.field(dtype=ti.f32, shape=1024) saxpy(x, y, 2.0)排查问题的推荐顺序:
- 先用
print_ir确认前端与优化 pass 层面是否符合预期(类型、循环结构、offload); - 若问题疑似出在 LLVM 生成/优化阶段,再打开三档 LLVM IR 开关对比;
- CUDA 下最后用
print_kernel_asm检查实际发射的 PTX/汇编指令。
数据访问器(Data Accessor)的特殊性
原文档用一个专门注释强调了 accessor 的特殊性,值得单独理解:
Data accessors in Python-scope are implemented as special Taichi kernels. For example,
x[1, 2, 3] = 3will call the writing accessor kernel ofx, andprint(y[42])will call the reading accessor kernel ofy.
这解释了三个现象:
- 为什么
print_ir打印不到它们:普通 kernel 编译与 accessor 编译走不同开关,print_ir明确排除 accessor; - 为什么要有
print_accessor_ir:当你的 bug 表现为"Python 侧读写字面索引行为异常"时,就需要用这个开关观察 accessor 的 IR; - 为什么它们也是 kernel:accessor 同样是编译产物、同样可被缓存,调试时可以把它们当作最小化的 kernel 来理解。
进一步深入
- 完整阅读 kernel 生命周期:Life of a Taichi Kernel;
- 掌握从源码构建与全部 CMake 开关:开发安装指南;
- 查看 IR 打印相关配置的全部字段与默认值:compile_config.h、compile_config.cpp;
- 观察 IR pass 打印器如何挂在编译流水线上:compile_to_offloads.cpp;
- 理解 Python 侧
ti.init如何解析这些 kwargs:misc.py。
掌握上述调试手段后,无论是为 Taichi 贡献新 pass、修复后端代码生成问题,还是深入理解自动微分与并行化细节,你都能第一时间看到编译器的真实行为,而不是在黑盒中猜测。
【免费下载链接】taichiProductive, portable, and performant GPU programming in Python.项目地址: https://gitcode.com/GitHub_Trending/ta/taichi
创作声明:本文部分内容由AI辅助生成(AIGC),仅供参考