 生成的内核源码)
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_compileFalse)停止在生成内核源码这一步然后把产物写入--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_compileFalse复用同一条降级管线却跳过了设备编译与 host 运行时模块的构建因此无需张量实参、也不会触碰 GPU。2. 环境要求与 wheel 前置条件该工具随 TileLang 标准安装一起发布没有额外可选依赖。两个与运行环境相关的前提值得注意默认ctarget 全平台可用CPU C 源码生成走的是 TVM C codegen任何机器包括纯 CPU 主机都能跑通。--target cuda依赖 wheel 中是否编译进 CUDA codegen FFILinux wheel 内置了该 FFI可以在没有 GPU 的机器上输出 CUDA 源码macOS Metal wheel 没有编译进该 FFI请求 CUDA target 时会得到一个清晰的CUDA codegen FFI missing诊断而不是AttributeError之类的裸异常。这个能力探测在源码里只有一个函数 cuda_codegen_availabledef 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, threads64): for i in T.Parallel(N): C[i] A[i] B[i] return C3.2 编译为 CPU C 源码python -I -m tilelang.tools.compile_only --output_file out.c example.pyout.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 源码同样无需 GPUpython -I -m tilelang.tools.compile_only --target cuda --output_file out.cu example.pyextern 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, threads64)的线程数一一对应。4. 输入文件的处理机制工具对输入文件的处理有三条规则每一条都能在 compile_only.py 中找到对应实现输入文件作为普通 Python 模块被 import。其顶层代码会真实执行所以示例必须保持 compile-only 形态只定义 kernel不 launch、不在 import 时分配张量。推荐用python -Iisolated 模式运行 CLICompiler Explorer 正是这样调用的。实现见 load_example它用importlib以模块名tilelang_compile_only_input加载并执行该文件而不是当 launch 脚本跑。选取模块顺序中第一个tilelang.jit函数或PrimFunc。discover_prim_func 按vars(module)的定义顺序遍历命中JITImpl就取其get_tir()命中裸PrimFunc就直接返回。由于遍历的是定义顺序一个先于tilelang.jitkernel 定义的PrimFunc会抢走编译机会——这一点有专门测试 test_discover_prim_func_respects_module_order 锁定行为。kernel 只被降级绝不被设备端编译或执行因此不需要任何张量实参也不会触碰 GPU。5. CLI 参考与退出码语义完整命令形态python -I -m tilelang.tools.compile_only [--target TARGET] --output_file OUTPUT input_file参数含义input_filecompile-only 示例文件。被 import绝不 launch。--output_file必需。生成内核源码的落盘位置。不得与输入文件同名指向输入文件的符号链接同样被拒绝。--target显式 target缺省为c。退出码与产物管理语义由 cli_main 实现成功退出码0源码写入--output_file。任何失败退出码1向 stderr 打印单行tilelang compile-only error: ...并且不留下陈旧的输出文件——编译开始前会先删除--output_file处的旧产物output.unlink(missing_okTrue)见 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--target取值行为c默认CPU C 源码。所有 wheel 可用无需 GPU。cuda固定到sm_80的 CUDA 源码不发生设备探测。cuda -archsm_90指定架构的 CUDA 源码。选项必须形如-keyvalue。{kind: cuda, arch: sm_90}JSON 形式 target。裸{kind: cuda}会被补全为sm_80{kind: auto}被拒绝。auto拒绝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 解析每个选项必须形如-archsm_90否则抛出CUDA target options must look like -archsm_90, or use JSON ...的指引性错误。JSON 形式会先json.loads再校验必须是带kind的对象kind为auto时按 auto 规则拒绝。在缺少 CUDA codegen FFI 的 wheel 上如 macOS Metal wheel所有CUDA target 形式cuda、cuda -archsm_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.pyTL_EMIT_LINE_DIRECTIVES tl.emit_line_directives文档说明其作用是从 TIR spans 在生成的 C 系源码中输出#line指令把生成的语句映射回 Python 源码行配合 nvcc 常开的-lineinfoPTX.loc条目会指向 Python 源码默认 False。工具侧由 _pass_context_config 处理版本兼容它用getattr(PassConfigKey, TL_EMIT_LINE_DIRECTIVES, None)探测 key 是否存在——旧版构建不认识这个 key直接硬编码传入会在 PassContext 里报错所以探测失败就返回空 dict。compile_kernel_source 在PassContext(opt_level3, config_pass_context_config())上下文中执行tilelang.lower(func, targetresolved, enable_device_compileFalse)取回CompiledArtifact.kernel_source若源码为空则抛出lower produced empty kernel_source。注意opt_level3与目标 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, targetc)接受PrimFunc例如来自JITImpl.get_tir()或 IRModule以enable_device_compileFalse降级并返回非空的内核源码字符串。target若是字符串会先经过resolve_target。resolve_target(target)把 CLI 风格的 target 字符串映射为传给lower()的显式 target字符串或 dict执行第 6 节表格中的全部规则。cuda_codegen_available()报告当前安装的 wheel 能否降级 CUDA target供调用方自行决定分支。10. 使用限制只编译模块中第一个 kerneltilelang.jit或PrimFunc按定义顺序取第一个同一文件中的其余 kernel 被忽略。产物是降级后的内核源码不是可执行二进制不会有任何东西被送到nvccPTX/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.pylower()与enable_device_compile参数tilelang/engine/lower.py如果需要在观察生成源码之外进一步定位性能问题或 IR 演化可以继续沿 docs/tools/index.md 的工具体系查阅 Analyzerroofline 估算、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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
TileLang Compile-Only 工具实战：无 GPU 环境下检查 lower() 生成的内核源码
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_compileFalse)停止在生成内核源码这一步然后把产物写入--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_compileFalse复用同一条降级管线却跳过了设备编译与 host 运行时模块的构建因此无需张量实参、也不会触碰 GPU。2. 环境要求与 wheel 前置条件该工具随 TileLang 标准安装一起发布没有额外可选依赖。两个与运行环境相关的前提值得注意默认ctarget 全平台可用CPU C 源码生成走的是 TVM C codegen任何机器包括纯 CPU 主机都能跑通。--target cuda依赖 wheel 中是否编译进 CUDA codegen FFILinux wheel 内置了该 FFI可以在没有 GPU 的机器上输出 CUDA 源码macOS Metal wheel 没有编译进该 FFI请求 CUDA target 时会得到一个清晰的CUDA codegen FFI missing诊断而不是AttributeError之类的裸异常。这个能力探测在源码里只有一个函数 cuda_codegen_availabledef 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, threads64): for i in T.Parallel(N): C[i] A[i] B[i] return C3.2 编译为 CPU C 源码python -I -m tilelang.tools.compile_only --output_file out.c example.pyout.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 源码同样无需 GPUpython -I -m tilelang.tools.compile_only --target cuda --output_file out.cu example.pyextern 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, threads64)的线程数一一对应。4. 输入文件的处理机制工具对输入文件的处理有三条规则每一条都能在 compile_only.py 中找到对应实现输入文件作为普通 Python 模块被 import。其顶层代码会真实执行所以示例必须保持 compile-only 形态只定义 kernel不 launch、不在 import 时分配张量。推荐用python -Iisolated 模式运行 CLICompiler Explorer 正是这样调用的。实现见 load_example它用importlib以模块名tilelang_compile_only_input加载并执行该文件而不是当 launch 脚本跑。选取模块顺序中第一个tilelang.jit函数或PrimFunc。discover_prim_func 按vars(module)的定义顺序遍历命中JITImpl就取其get_tir()命中裸PrimFunc就直接返回。由于遍历的是定义顺序一个先于tilelang.jitkernel 定义的PrimFunc会抢走编译机会——这一点有专门测试 test_discover_prim_func_respects_module_order 锁定行为。kernel 只被降级绝不被设备端编译或执行因此不需要任何张量实参也不会触碰 GPU。5. CLI 参考与退出码语义完整命令形态python -I -m tilelang.tools.compile_only [--target TARGET] --output_file OUTPUT input_file参数含义input_filecompile-only 示例文件。被 import绝不 launch。--output_file必需。生成内核源码的落盘位置。不得与输入文件同名指向输入文件的符号链接同样被拒绝。--target显式 target缺省为c。退出码与产物管理语义由 cli_main 实现成功退出码0源码写入--output_file。任何失败退出码1向 stderr 打印单行tilelang compile-only error: ...并且不留下陈旧的输出文件——编译开始前会先删除--output_file处的旧产物output.unlink(missing_okTrue)见 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--target取值行为c默认CPU C 源码。所有 wheel 可用无需 GPU。cuda固定到sm_80的 CUDA 源码不发生设备探测。cuda -archsm_90指定架构的 CUDA 源码。选项必须形如-keyvalue。{kind: cuda, arch: sm_90}JSON 形式 target。裸{kind: cuda}会被补全为sm_80{kind: auto}被拒绝。auto拒绝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 解析每个选项必须形如-archsm_90否则抛出CUDA target options must look like -archsm_90, or use JSON ...的指引性错误。JSON 形式会先json.loads再校验必须是带kind的对象kind为auto时按 auto 规则拒绝。在缺少 CUDA codegen FFI 的 wheel 上如 macOS Metal wheel所有CUDA target 形式cuda、cuda -archsm_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.pyTL_EMIT_LINE_DIRECTIVES tl.emit_line_directives文档说明其作用是从 TIR spans 在生成的 C 系源码中输出#line指令把生成的语句映射回 Python 源码行配合 nvcc 常开的-lineinfoPTX.loc条目会指向 Python 源码默认 False。工具侧由 _pass_context_config 处理版本兼容它用getattr(PassConfigKey, TL_EMIT_LINE_DIRECTIVES, None)探测 key 是否存在——旧版构建不认识这个 key直接硬编码传入会在 PassContext 里报错所以探测失败就返回空 dict。compile_kernel_source 在PassContext(opt_level3, config_pass_context_config())上下文中执行tilelang.lower(func, targetresolved, enable_device_compileFalse)取回CompiledArtifact.kernel_source若源码为空则抛出lower produced empty kernel_source。注意opt_level3与目标 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, targetc)接受PrimFunc例如来自JITImpl.get_tir()或 IRModule以enable_device_compileFalse降级并返回非空的内核源码字符串。target若是字符串会先经过resolve_target。resolve_target(target)把 CLI 风格的 target 字符串映射为传给lower()的显式 target字符串或 dict执行第 6 节表格中的全部规则。cuda_codegen_available()报告当前安装的 wheel 能否降级 CUDA target供调用方自行决定分支。10. 使用限制只编译模块中第一个 kerneltilelang.jit或PrimFunc按定义顺序取第一个同一文件中的其余 kernel 被忽略。产物是降级后的内核源码不是可执行二进制不会有任何东西被送到nvccPTX/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.pylower()与enable_device_compile参数tilelang/engine/lower.py如果需要在观察生成源码之外进一步定位性能问题或 IR 演化可以继续沿 docs/tools/index.md 的工具体系查阅 Analyzerroofline 估算、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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
