TileLang 命令大全:compile_only 编译、性能分析、pass 调试与环境变量速查
TileLang 的命令分四族:编译构建、性能分析、pass 调试与构建期环境变量。最常用的两条是——不接 GPU 也能编 kernel 的 python -I -m tilelang.tools.compile_only --target cuda --output_file out.cu example.py,和静态估算开销的 tilelang.tools.Analyzer.analysis(tir, device);编译产物缓存在 ~/.tilelang/cache。
本文是速查表,按「编译构建 → 性能与调试 → 环境变量与缓存」三族整理,每条都给命令、用途与预期输出。装环境的过程见《TileLang 怎么安装》。
TileLang 的编译与构建命令有哪些?
TileLang 的编译命令分两类:把 kernel 脚本编译成后端源码的 compile_only(python -I -m tilelang.tools.compile_only --target cuda --output_file out.cu example.py),以及把 TileLang 本体从源码构建安装的 pip / 脚本命令(来源)。
- 离线编译 kernel →
compile_only。执行python -I -m tilelang.tools.compile_only --target cuda --output_file out.cu example.py。预期结果:out.cu里出现生成的 CUDA 源码;-I隔离环境变量,--target必须是cuda、rocm这类具体后端,传auto会被拒绝(来源)。 - 失败要能卡住流水线。同一命令编译失败时退出码为 1。预期结果:脚本或 CI 依据退出码判定失败,可直接接进自动化校验。
- 构建 TileLang 本体(NVIDIA 侧)。源码目录下按需设置
USE_CUDA=ON,再用 pip 或python -m build -w产出 wheel。预期结果:dist/下生成 wheel,构建开关与你的工具链对应(来源)。 - 构建 TileLang 本体(昇腾侧)。
./build_wheel_ascend.sh [--enable-llvm],或直接bash install_ascend.sh装进当前环境。预期结果:前者在dist/产出 wheel,后者完成就地安装(来源)。
TileLang 的性能分析与调试命令怎么用?
TileLang 提供了不跑 kernel 就能估开销的 Analyzer,以及若干个只用于开发期定位问题的调试入口——布局可视化 plot_layout、自动并行化调试 python -m tilelang.autodd、lowering 跟踪 TL_LOWER_TRACE 与 pass 对比 TILELANG_PASS_DIFF(来源)。
- 估 FLOPs 与访存 →
Analyzer.analysis。调用tilelang.tools.Analyzer.analysis(tir, device)。预期结果:得到该 kernel 的 FLOPs、全局访存量与 roofline 时间估算,用来判断瓶颈在算力还是带宽(来源)。 - 看布局与时序 →
plot_layout。调用tilelang.tools.plot_layout。预期结果:输出 tiling 与内存层次的布局视图,用于确认分块是否按预期落到各级存储。 - 调自动并行化 →
python -m tilelang.autodd。执行该模块。预期结果:打印自动并行化 / 自动微分相关的推导与决策信息,便于定位并行方案为什么不理想(来源)。 - 看 lowering 全过程 →
TL_LOWER_TRACE。设置该环境变量后重新编译。预期结果:终端输出 TIR 到后端的逐步 lowering 跟踪;需要比对某个 pass 前后的差异时,改用TILELANG_PASS_DIFF。
这些命令属于开发期工具:Analyzer 用于调优前定方向,plot_layout / autodd / TL_LOWER_TRACE 用于定位编译期问题,都不参与生产推理链路。
TileLang 的构建期环境变量与编译缓存
TileLang 的构建期环境变量按硬件路线分:NVIDIA 侧用 USE_CUDA 打开 CUDA 后端,USE_ROCM / USE_METAL 对应 ROCm 与 Metal,USE_LLVM、TVM_ROOT 决定编译栈与 TVM 源码位置,WITH_PIP_CUDA_TOOLCHAIN 与 NO_VERSION_LABEL 控制工具链来源与版本标签;昇腾侧由 ASCEND_HOME_PATH 定位昇腾库。编译产物统一缓存在 ~/.tilelang/cache(来源)。
| 变量 / 路径 | 作用 | 用在哪条路线 |
|---|---|---|
USE_CUDA | 打开 CUDA 后端 | NVIDIA |
USE_ROCM | 打开 ROCm 后端 | AMD |
USE_METAL | 打开 Metal 后端 | Apple |
USE_LLVM | 启用 LLVM 后端 | 通用 / 昇腾自建 |
TVM_ROOT | 指向 TVM 源码位置 | 源码构建 |
WITH_PIP_CUDA_TOOLCHAIN | 使用 pip 提供的 CUDA 工具链 | NVIDIA |
NO_VERSION_LABEL | 关闭 wheel 的版本标签 | 构建打包 |
ASCEND_HOME_PATH | 指向昇腾工具链根目录 | 昇腾 |
~/.tilelang/cache | 编译产物缓存目录 | 两条路线 |
- 改设置后要重编。构建期变量在编译时读取。预期结果:设完变量重新构建才生效,只 export 不重编不会改变已装产物。
- 清缓存要连目录一起删。执行
rm -rf ~/.tilelang/cache。预期结果:下次编译重新生成产物;升级版本或改动 kernel 后仍命中旧结果时,这是第一步排查手段。 - 昇腾侧变量来自 CANN。
ASCEND_HOME_PATH通常由set_env.sh设置。预期结果:echo $ASCEND_HOME_PATH有值,编译器才能找到昇腾库。
TileLang 命令速查表
| 目的 | 命令 | 预期输出 |
|---|---|---|
| 离线编译 kernel | python -I -m tilelang.tools.compile_only --target cuda --output_file out.cu example.py | 生成 out.cu;失败退出码 1 |
| 估算 FLOPs / 访存 | tilelang.tools.Analyzer.analysis(tir, device) | FLOPs、全局访存量、roofline 时间 |
| 查看布局 | tilelang.tools.plot_layout | tiling 与内存层次布局 |
| 自动并行化调试 | python -m tilelang.autodd | 并行化决策信息 |
| lowering 跟踪 | TL_LOWER_TRACE=1(环境变量) | TIR 逐步 lowering 日志 |
| pass 前后对比 | TILELANG_PASS_DIFF=1(环境变量) | pass 差异输出 |
| 升版本 | pip install -U tilelang | 装上更新的正式版 |
| 装 nightly | pip install tilelang -f https://tile-ai.github.io/whl/nightly | 装上当天预览版 |
| 卸载 | pip uninstall tilelang | 移除包(昇腾侧为 tilelang-ascend) |
TileLang 命令使用注意事项
--target不接受auto。compile_only要求显式后端,写成cuda/rocm这类具体值。- 调试变量只用在排查期。
TL_LOWER_TRACE、TILELANG_PASS_DIFF会产生大量日志,确认无关后记得取消设置。 - 版本与缓存要一起管。升级 TileLang 前先看《TileLang 怎么更新》,它同时讲了钉版本与清
~/.tilelang/cache的时机。 - 卸载别误伤依赖。
pip uninstall的边界与残留清理,见《TileLang 怎么卸载》。
来源:TileLang Tools、TileLang compile_only tool、TileLang Analyzer、TileLang Installation Guide
常见问题
TileLang 可以在没有 GPU 的机器上完成编译:执行 python -I -m tilelang.tools.compile_only --target cuda --output_file out.cu example.py,它会输出生成的 CUDA 源码而不执行。注意 target 必须显式写成 cuda、rocm 这类具体后端,传 auto 会被拒绝;编译失败时进程退出码为 1,可用于 CI 校验。
TileLang 的性能分析入口是 tilelang.tools.Analyzer.analysis(tir, device),它在不实际运行 kernel 的前提下静态估算 FLOPs、全局访存量以及 roofline 时间,适合在调优前判断瓶颈在算力还是带宽。把目标设备传进 device 参数,昇腾与 NVIDIA 都能用。
TileLang 的调试分两层:想看布局与时序可以用 tilelang.tools.plot_layout 画布局图,想定位自动并行化决策可以跑 python -m tilelang.autodd。要观察 TIR 到后端的每一步 lowering,则设置环境变量 TL_LOWER_TRACE 打开跟踪,需要比对 pass 前后差异时用 TILELANG_PASS_DIFF。
TileLang 默认把编译产物缓存在 ~/.tilelang/cache,同一个 kernel 再次编译会命中缓存直接复用,省掉重编时间。改动 kernel 源码后仍命中旧产物时,删掉该目录即可强制重编;升级版本后也建议清一次,避免新旧产物混用。
TileLang 的构建期环境变量按路线划分:NVIDIA 侧用 USE_CUDA 打开 CUDA 后端,ROCM 用 USE_ROCM、Apple 用 USE_METAL,TVM_ROOT 指向 TVM 源码位置,WITH_PIP_CUDA_TOOLCHAIN 与 NO_VERSION_LABEL 控制工具链来源和版本标签;昇腾侧则由 CANN 的 set_env.sh 设置 ASCEND_HOME_PATH 供编译器定位昇腾库。
相关术语
- compile_only
- compile_only 是 TileLang 自带的离线编译工具(python -m tilelang.tools.compile_only),把 TileLang 脚本编译成目标后端的源码输出到文件,不需要 GPU 在场,常用于 CI 与代码审查。— TileLang compile_only tool
- Analyzer
- Analyzer 是 TileLang 的静态性能分析接口,tilelang.tools.Analyzer.analysis(tir, device) 在不运行 kernel 的情况下估算 FLOPs、全局访存量与 roofline 时间。— TileLang Analyzer
- TL_LOWER_TRACE
- TL_LOWER_TRACE 是 TileLang 的调试环境变量,打开后会在 lowering 过程中输出每一步 TIR 变换的跟踪日志,用于定位编译期问题。— TileLang Tools
- 编译缓存
- TileLang 的编译缓存位于 ~/.tilelang/cache,保存已编译的 kernel 产物;再次编译相同 kernel 时直接复用,改代码或升级版本后应清理。— TileLang Installation Guide
来源
- TileLang Tools· TileLang (tile-ai)
- TileLang compile_only tool· TileLang (tile-ai)
- TileLang Analyzer· TileLang (tile-ai)
- TileLang Installation Guide· TileLang (tile-ai)