企业官网建设流程全解析

TileLang Compile-Only 工具实战：无 GPU 环境下检查 lower() 生成的内核源码

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

本文围绕 TileLang 的 Compile-Only 工具（python -m tilelang.tools.compile_only）展开：它能把一个 TileLang 程序降级（lower）为可直接查看的内核源码，全程不触碰 GPU、CUDA 驱动或设备端编译。读完本文，你可以掌握该工具的完整命令行用法与退出码语义、各 target 取值规则（含sm_80固定与 JSON 形式）、#line指令映射机制，以及compile_kernel_source编程接口的调用方式，并能对照仓库源码理解每一条行为约束背后的实现依据。

1. 工具定位：只降级、不运行

Compile-Only 工具的完整用法一句话概括：

python -m tilelang.tools.compile_only --output_file out.c example.py

它的核心动作是调用tilelang.lower(..., enable_device_compile=False)停止在"生成内核源码"这一步，然后把产物写入--output_file指定的文件——不再往nvcc送任何东西，不产生 PTX/SASS，也不需要机器上有 GPU 或 CUDA 驱动。这个特性使它成为 TileLang 接入在线 Compiler Explorer 的入口：在线环境天然没有设备，只能依赖这种纯代码生成路径；同时在本地调试时，它也用来检查某个 kernel 经lower()之后到底长什么样。

从源码结构看，"只降级"这一行为对应 engine/lower.py 中lower_with_context的enable_device_compile参数：其文档注释明确说明该参数"控制设备代码是否在本阶段编译，默认禁用，因为 JIT 适配器通常单独处理设备端编译"。Compile-Only 工具正是显式传入enable_device_compile=False，复用同一条降级管线却跳过了设备编译与 host 运行时模块的构建，因此无需张量实参、也不会触碰 GPU。

2. 环境要求与 wheel 前置条件

该工具随 TileLang 标准安装一起发布，没有额外可选依赖。两个与运行环境相关的前提值得注意：

• 默认ctarget 全平台可用：CPU C 源码生成走的是 TVM C codegen，任何机器（包括纯 CPU 主机）都能跑通。
• --target cuda依赖 wheel 中是否编译进 CUDA codegen FFI：Linux wheel 内置了该 FFI，可以在没有 GPU 的机器上输出 CUDA 源码；macOS Metal wheel 没有编译进该 FFI，请求 CUDA target 时会得到一个清晰的CUDA codegen FFI missing诊断，而不是AttributeError之类的裸异常。

这个能力探测在源码里只有一个函数 cuda_codegen_available：

def cuda_codegen_available() -> bool: """Return whether this wheel can lower CUDA (``AnnotateDeviceBoundTmaCopies``).""" try: from tilelang.cuda import _ffi_api return hasattr(_ffi_api, "AnnotateDeviceBoundTmaCopies") except Exception: return False

它以 CUDA 变换层中AnnotateDeviceBoundTmaCopies这个 FFI 符号是否存在作为探针：只要 wheel 能完成 TMA 相关设备端变换，就认为具备完整的 CUDA 降级能力。这个设计让不同平台的 wheel 对同一命令行表现出可预期的差异化行为，也解释了上文"软失败"（fail softly）措辞的来源。

3. Quick Start：一个自包含示例与两种 target 输出

3.1 编写只定义、不启动的示例

示例文件必须是"compile-only"形态——定义 kernel，但绝不启动它、也不在 import 时分配张量：

# example.py import tilelang import tilelang.language as T @tilelang.jit def add(A, B): N = 64 A: T.Tensor((N,), T.float16) B: T.Tensor((N,), T.float16) C = T.empty((N,), T.float16) with T.Kernel(1, threads=64): for i in T.Parallel(N): C[i] = A[i] + B[i] return C

3.2 编译为 CPU C 源码

python -I -m tilelang.tools.compile_only --output_file out.c example.py

out.c中会得到生成的内核源码（文档给出的真实产物形态）：

// tilelang target: {"kind":"c","tag":"","keys":["cpu"],"host":{"kind":"c","tag":"","keys":["cpu"]}} #include <tl_templates/cpp/common.h> #ifdef __cplusplus extern "C" #endif int32_t add_kernel(half* A, half* B, half* C); #ifdef __cplusplus extern "C" #endif int32_t add_kernel(half* A, half* B, half* C) { for (int32_t i = 0; i < 8; ++i) { *(half8*)(C + (i * 8)) = (*(half8*)(A + (i * 8)) + *(half8*)(B + (i * 8))); } return 0; }

几个可以直观验证的细节：首行注释携带了完整的 target JSON（便于工具链消费）；half8说明 64 元素的T.Parallel循环在降级过程中被向量化为 8 路half8访问；kernel 被暴露为extern "C"的int32_t函数，这是 TileLang CPU 端 host codegen 的固定签名风格。

3.3 编译为 CUDA 源码（同样无需 GPU）

python -I -m tilelang.tools.compile_only --target cuda --output_file out.cu example.py

extern "C" __global__ void __launch_bounds__(64, 1) add_kernel(const half_t* __restrict__ A, const half_t* __restrict__ B, half_t* __restrict__ C) { C[((int)threadIdx.x)] = (A[((int)threadIdx.x)] + B[((int)threadIdx.x)]); }

CUDA 路径与 CPU 路径的差别在于：架构（arch）是固定（pinned）的而不是探测出来的——默认固定到sm_80，因此整个流程不发生任何设备检测。__launch_bounds__(64, 1)与T.Kernel(1, threads=64)的线程数一一对应。

4. 输入文件的处理机制

工具对输入文件的处理有三条规则，每一条都能在 compile_only.py 中找到对应实现：

1. 输入文件作为普通 Python 模块被 import。其顶层代码会真实执行，所以示例必须保持 compile-only 形态：只定义 kernel，不 launch、不在 import 时分配张量。推荐用python -I（isolated 模式）运行 CLI，Compiler Explorer 正是这样调用的。实现见 load_example：它用importlib以模块名tilelang_compile_only_input加载并执行该文件，而不是当 launch 脚本跑。
2. 选取"模块顺序中第一个"@tilelang.jit函数或PrimFunc。discover_prim_func 按vars(module)的定义顺序遍历，命中JITImpl就取其get_tir()，命中裸PrimFunc就直接返回。由于遍历的是定义顺序，一个先于@tilelang.jitkernel 定义的PrimFunc会"抢走"编译机会——这一点有专门测试 test_discover_prim_func_respects_module_order 锁定行为。
3. kernel 只被降级，绝不被设备端编译或执行，因此不需要任何张量实参，也不会触碰 GPU。

5. CLI 参考与退出码语义

完整命令形态：

python -I -m tilelang.tools.compile_only [--target TARGET] --output_file OUTPUT input_file

退出码与产物管理语义（由 cli_main 实现）：

• 成功：退出码0，源码写入--output_file。
• 任何失败：退出码1，向 stderr 打印单行tilelang compile-only error: ...，并且不留下陈旧的输出文件——编译开始前会先删除--output_file处的旧产物（output.unlink(missing_ok=True)，见 cli_main 的注释说明这是对齐 Compiler Explorer wrapper 的做法：CE 会复用同一个--output_file，失败的运行不能把上一次的产物留给消费者去读）。
• 输出路径保护：若--output_file与输入文件是同一文件（含经符号链接指向同一文件的情形），直接报错--output_file must not be the input file，输入文件内容保持不变。_paths_are_same同时用Path.resolve()与samefile双重判定。测试 test_compile_only_cli_keeps_input_when_output_is_symlink 验证了符号链接情形下输入文件原样保留。

对应地，"失败不留陈旧产物"的行为由 test_compile_only_cli_unlinks_stale_output_on_failure 专门回归：先在输出路径写入一个"昨天的汇编"占位，再让编译失败，断言占位内容不再存在。

6. Target 规则：必须显式，auto被拒绝

Target 必须显式给出。auto会触发设备探测，因此在纯字符串和 JSON 两种形式下都被拒绝：

$ python -I -m tilelang.tools.compile_only --target auto --output_file out.c example.py tilelang compile-only error: target must be explicit; do not use auto

这些规则集中实现在 resolve_target，几个源码级细节：

• 模块顶部有常量_PINNED_CUDA_TARGET = {"kind": "cuda", "arch": "sm_80"}（第 40 行）：裸cuda字符串、不带arch的 JSON{"kind": "cuda"}都会归一到它，保证"无设备检测"这一不变量。
• 带选项的 CUDA 字符串由 _parse_cuda_cli_options 解析：每个选项必须形如-arch=sm_90，否则抛出CUDA target options must look like -arch=sm_90, or use JSON ...的指引性错误。
• JSON 形式会先json.loads再校验必须是带kind的对象，kind为auto时按 auto 规则拒绝。

在缺少 CUDA codegen FFI 的 wheel 上（如 macOS Metal wheel），所有CUDA target 形式（cuda、cuda -arch=sm_90、JSON 形式）都会软失败为：

tilelang compile-only error: CUDA codegen FFI missing (e.g. macOS Metal wheel); use default --target c

该软失败在 compile_kernel_source 中实现：目标解析为 CUDA 而cuda_codegen_available()为假时，抛出带上述文案的RuntimeError，避免AnnotateDeviceBoundTmaCopies缺失时泄漏裸AttributeError。测试 test_compile_only_cli_cuda_soft_fails_without_ffi 覆盖了全部三种 CUDA target 形式的软失败路径（在无 CUDA FFI 的环境上才执行）。

7. 错误报告：诊断而非 traceback

编译问题以 stderr 上的单行诊断和退出码1报告，而不是围绕"缺 GPU/缺驱动"的 traceback：

$ python -I -m tilelang.tools.compile_only --output_file out.c broken.py tilelang compile-only error: no @tilelang.jit kernel or PrimFunc found

输入文件的语法错误、输入文件缺失、空文件（不含任何 kernel）都走同一条报告路径。cli_main把load_example/discover_prim_func/compile_kernel_source包在统一的try/except里，任何异常都折叠为{_ERROR_PREFIX}: {exc}单行输出（前缀常量_ERROR_PREFIX = "tilelang compile-only error"，见 第 41 行）。测试侧也显式断言了这一点：test_compile_only_cli_reports_compile_error 检查错误输出中不出现CUDA与driver字样——它必须是纯编译诊断，而不是设备环境的缺失报错；test_compile_only_cli_reports_missing_input_file 与 test_compile_only_cli_reports_empty_file 分别覆盖文件缺失与空模块两种情形，且都断言不留下任何输出文件。

8.#line指令：把生成源码映射回示例文件

当当前安装的构建注册了tl.emit_line_directives这个 pass config（即PassConfigKey.TL_EMIT_LINE_DIRECTIVES）时，工具会自动启用它，生成源码中便携带#line <n> "<input_file>"指令，把每条语句映射回示例文件的原始行号。Compiler Explorer 正是用这些指令把右侧生成源码点击定位到左侧输入面板的对应行。旧版 wheel 没有这个 config key 时，工具只是不产生#line指令，两种情况下都不需要额外传 flag。

实现链路值得展开：

• 该 config key 定义在 tilelang/transform/pass_config.py：TL_EMIT_LINE_DIRECTIVES = "tl.emit_line_directives"，文档说明其作用是"从 TIR spans 在生成的 C 系源码中输出#line指令，把生成的语句映射回 Python 源码行；配合 nvcc 常开的-lineinfo，PTX.loc条目会指向 Python 源码"，默认 False。
• 工具侧由 _pass_context_config 处理版本兼容：它用getattr(PassConfigKey, "TL_EMIT_LINE_DIRECTIVES", None)探测 key 是否存在——旧版构建不认识这个 key，直接硬编码传入会在 PassContext 里报错，所以探测失败就返回空 dict。
• compile_kernel_source 在PassContext(opt_level=3, config=_pass_context_config())上下文中执行tilelang.lower(func, target=resolved, enable_device_compile=False)，取回CompiledArtifact.kernel_source；若源码为空则抛出lower produced empty kernel_source。注意opt_level=3与目标 target 上下文with ..., resolved的写法——注释说明这是为了与JITKernel._compile_artifact保持一致，因为 LayoutInference 会读取Target.current()。

#line的正确性由两个测试回归：进程内的 test_compile_kernel_source_emits_line_directives 与走 CLI 子进程的 test_compile_only_cli_emits_line_directives，它们都在源码中埋了唯一的行标记注释，然后断言#line指令中确实存在指向该标记行的条目。

9. 编程接口

同样的能力也可以直接在 Python 里调用：

from tilelang.tools.compile_only import compile_kernel_source source = compile_kernel_source(add.get_tir()) # 默认 target "c" cuda_source = compile_kernel_source(add.get_tir(), "cuda") # 固定到 sm_80

模块公开的 API（见 compile_only.py 的__all__）：

• compile_kernel_source(func, target="c")：接受PrimFunc（例如来自JITImpl.get_tir()）或 IRModule，以enable_device_compile=False降级并返回非空的内核源码字符串。target若是字符串会先经过resolve_target。
• resolve_target(target)：把 CLI 风格的 target 字符串映射为传给lower()的显式 target（字符串或 dict），执行第 6 节表格中的全部规则。
• cuda_codegen_available()：报告当前安装的 wheel 能否降级 CUDA target，供调用方自行决定分支。

10. 使用限制

• 只编译模块中第一个 kernel：@tilelang.jit或PrimFunc按定义顺序取第一个，同一文件中的其余 kernel 被忽略。
• 产物是降级后的内核源码，不是可执行二进制：不会有任何东西被送到nvcc，PTX/SASS 不在该工具的能力范围内。
• 输入文件由调用方的 Python 解释器直接执行：请像对待任何可运行脚本一样对待它；python -I能把环境与用户配置隔离，但并不构成沙箱。

11. 小结与延伸阅读

Compile-Only 工具是 TileLang 工具链中"可观察性"的一端：它以最小依赖（标准安装即自带）、最强环境约束（无 GPU、无驱动、target 必须显式）换来了在任意机器上审查lower()产物的一致性体验，并通过退出码0/1、单行 stderr 诊断与"失败不留陈旧产物"三条契约保证它可以直接被自动化系统（如 Compiler Explorer）消费。关键文件索引：

• 工具实现：tilelang/tools/compile_only.py
• 官方文档：docs/tools/compile_only.md，工具总览见 docs/tools/index.md（其中把该工具列为"无 GPU 输出内核源码"的任务入口）
• 回归测试：testing/python/tools/test_tilelang_tools_compile_only.py
• #line指令的 pass config 定义：tilelang/transform/pass_config.py
• lower()与enable_device_compile参数：tilelang/engine/lower.py

如果需要在观察生成源码之外进一步定位性能问题或 IR 演化，可以继续沿 docs/tools/index.md 的工具体系查阅 Analyzer（roofline 估算）、Layout Visualization（线程/数据映射）与 IR Lower Trace（全流水线 pass 观察）等配套工具。

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