从 TVM 到 TileLang为什么Tile 化正在重写 AI 编译器的底层叙事【免费下载链接】tilelangDomain-specific language designed to streamline the development of high-performance GPU/CPU/Accelerators kernels项目地址: https://gitcode.com/GitHub_Trending/ti/tilelang十年前深度学习编译器给出的叙事是一次编写处处优化开发者写出高层算子描述编译器通过 schedule 自动变换来榨取硬件性能。十年后这条叙事的根基正在松动。以 TileLang 为代表的新一代 Tile-centric 编程模型把分块Tiling从编译器内部的一次调度动作升格为编程语言的一等公民——它要求开发者像 Kernel 工程师一样思考同时让编译器接管那些曾经让 CUTLASS 模板写起来痛不欲生的机械细节。本文结合 TileLang 源码与社区实践拆解这场范式迁移的内在逻辑。一、TVM 时代Tiling 只是调度手段TVM 的经典架构是分层抽象声明式的 TETensor Expression负责描述计算命令式的 TIR 承载底层 IR而连接两者的核心概念是schedule。在 TVM 的世界里tile、split、reorder、fuse都是一系列程序变换原语——编译器把开发者写的计算图切碎、重排、映射到线程和内存上。这种设计的隐含假设是编译器而非开发者掌握性能的关键决策。schedule 提供的是搜索空间tile 的大小、顺序只是搜索空间里的坐标。但这一假设在 GPU 上逐渐失效现代 GPU 的内存层次全局内存、Shared Memory、寄存器/register fragment、Tensor Core 的 MMA 布局极其细碎仅凭高层 schedule 无法表达这个累加器必须放在寄存器 fragment 里并按 wgmma 的 64×N 布局切分调度变换是非局部的一个cache_read或split会连锁影响后续所有 pass 的行为开发者难以预测最终生成的代码当 TVM 面对 FlashAttention 这类算子时人们发现真正决定性能的不是 loop 顺序而是数据在每一层内存里以什么形状、被哪些线程以什么布局持有——而这恰恰是 TE/TIR 时代没有建模的东西。于是社区里出现了一个有趣的反向运动与其让编译器在 opaque 的调度空间里猜不如把硬件的关键事实搬到语言表面让懂硬件的人直接说出来。TileLang 是这一运动最彻底的代表之一。二、TileLang 的转向Shared Memory、Fragment、Pipeline 成为一等公民TileLang 建立于 TVM 的编译基础设施之上其 Python 语言层直接构建在 TVM 的 TIRX 之上但在编程模型上做了根本性偏移。看 examples/gemm/example_gemm.py 里一个标准的 tiled GEMMwith T.Kernel(T.ceildiv(N, block_N), T.ceildiv(M, block_M), threads128) as (bx, by): A_shared T.alloc_shared((block_M, block_K), dtype) B_shared T.alloc_shared((block_K, block_N), dtype) C_local T.alloc_fragment((block_M, block_N), accum_dtype) T.clear(C_local) for k in T.Pipelined(T.ceildiv(K, block_K), num_stages3): T.copy(A[by * block_M, k * block_K], A_shared) T.copy(B[k * block_K, bx * block_N], B_shared) T.gemm(A_shared, B_shared, C_local) T.copy(C_local, C[by * block_M, bx * block_N])这段代码里没有任何一处写blockIdx、threadIdx、同步或 double buffering但每一行都在描述硬件结构T.alloc_shared/T.alloc_fragment直接把缓冲区钉在 Shared Memory 和寄存器 fragment 上。在 tilelang/language/allocate.py 中这些 API 是对 TVM buffer allocation 的薄封装alloc_shared默认落在shared.dynscopealloc_fragment落在local.fragmentscopealloc_var落在local.var——scope 本身就是语言类型系统的一部分T.Pipelined(..., num_stages3)声明软件流水线生产者T.copy与消费者T.gemm的重叠由编译器编排见 tilelang/language/loop.py 与 docs/programming_guides/software_pipeline.md 中关于stage/order的说明T.gemm一个 tile 级别的算子调用。在 tilelang/language/gemm_op.py 中它负责校验 A/B/C 的形状语义并把transpose、warp policy、accumulator 清零等意图编码进 13 个槽位的 tile-op 协议中最终由后端 lowering 到 Tensor Core 指令。值得强调的是显式建模并不意味着用户要手写一切。恰恰相反TileLang 的策略是用户声明数据在哪一层内存、以什么 tile 计算编译器负责把这份声明一致地落到每个 warp 头上。这正是它与先写高层计算、再靠调度变换的 TVM 旧路线的分水岭——前者把硬件事实作为输入后者把硬件事实当作需要推导的结论。下图展示了 TileLang 文档中对 GEMM 多级 tiling 的说明数据流经全局内存 → Shared Memory → register fragment每一级都有显式的分配原语对应三、布局推断编译器最值钱的部分被保留下来如果一切都要开发者亲手指定TileLang 就和手写 CUDA 没有区别。它的高明之处在于把布局推断Layout Inference留给编译器用户只给出 fragment 的逻辑形状编译器推导每个线程实际持有哪些寄存器、Shared Memory 如何 swizzle、warpgroup 之间如何切分。这一点在 examples/deepseek_mla/README.md 的 MLADeepSeek 多头潜在注意力案例中体现得淋漓尽致。MLA 的难点在于 head dim 高达 576acc_o太大单个 warpgroup 会寄存器溢出必须把acc_o沿 dim 切给两个 warpgroup而两个 warpgroup 又都需要完整的acc_s作为后续acc_s V的输入。这段逻辑听起来极其绕但 TileLang 用户只需要给T.gemm一个policyT.GemmWarpPolicy.FullCol注解T.gemm(Q_shared, K_shared, acc_s, transpose_BTrue, policyT.GemmWarpPolicy.FullRow)随后布局推断自动推导出QK 阶段每个 warpgroup 持有acc_s_0: [block_M, block_N/2]而 PV 阶段需要完整acc_s: [block_M, block_N]进而把中间的S_shared也推导为[block_M, block_N]。同时warp specialization谁做 TMA producer、谁做 consumer、mbarrier 同步、threadblock swizzle、Shared Memory 的 bank-conflict 规避全部由编译器在幕后完成——用户可见的只是T.use_swizzle(panel_size10)这样的单行声明。这就是Kernel 工程师的直觉更值钱的真正含义直觉没有被机器取代而是被语言接纳了。一个懂 GPU 内存层次、懂 bank conflict、懂 register pressure 的工程师现在可以把这些直觉直接写成语言结构而不是费力地把它们编码成一系列 schedule 变换的组合。而编译器负责的是把直觉转换成正确且一致的低层实现并暴露出足够多的调试工具如 docs/tools/lower_trace.md 中按 pass 展示 IR 变化的 Lower Trace、docs/tools/layout_visualization.md 中的布局可视化让工程师验证自己的判断。结果层面这种人机协同在社区基准里得到了验证TileLang 用约 80 行 Python 实现的 MLA decode在 H100 上与 DeepSeek 官方 FlashMLA基于 CUTLASS 模板性能相当并显著超过 FlashInfer 与 Triton 的实现在 benchmark/matmul/README.md 记录的 FP16 GEMM 基准中H800 上 8192×8192 规模吞吐可达 758 TFLOPsK8192逼近硬件理论峰值。四、对 Triton 等 DSL 的冲击与启发Triton 用block 级编程 自动调优重新定义了 GPU DSL 的体验开发者写 tile 计算编译器负责线程映射与共享内存管理。TileLang 走的是另一条路——不隐藏硬件而是显式建模硬件。这两条路线在社区讨论中被反复对照Triton 凭借 Python 接口和自动调优更适合动态 shape、快速迭代场景而 TileLang 在结构化稀疏等对内存布局敏感的算子上一度取得领先性能因为它允许开发者精确控制 Shared Memory 的 swizzle 布局与 warp 划分。这种差异不是偶然。稀疏 GEMM 需要在数据进入 Shared Memory 之前就完成压缩与元数据编排见 examples/gemm_sp/example_gemm_sp.py 中A_sparse、E两路并行载入再交给T.gemm_sp这类算子的性能密码藏在数据以什么形状进入什么内存里仅靠高层语义与自动调度很难表达清楚。TileLang 的实践表明当显式硬件建模带来的收益超过自动化的便利时性能叙事的天平就会倾斜。更具启发意义的是 TileLang 正在验证的另一个命题同一个 tile 程序可以横跨异构后端。仓库 tilelang/backend/README.md 描述了一套后端即编译器垂直切片的架构每个目标后端拥有自己的语言方言、pass pipeline 与 codegen而共享的 TileLang 前端与编译入口保持后端无关。目前主仓库已覆盖 NVIDIA CUDASM70–SM120、AMD ROCm/HIP、Apple Metal、Huawei Ascend 950 与 LLVM CPU 后端另有昇腾 A2/A3、摩尔线程 MUSA、海光 HCU、MetaX、天狼星等生态适配。这一点在社区舆情中引发了连锁反应——围绕 TileLang 的讨论已经从又一个 GPU DSL升级为国产芯片的确定性算力编译器当 CUDA 的垄断叙事被打破一个以 tile 为第一等公民、不绑定任何单家硬件的编程模型自然成为连接大模型算子与多样算力的公共语言。对 Triton 及整个 DSL 生态而言TileLang 的启示是深层的tile 不只是一个性能优化技巧而是一个完整的抽象层。谁能把tile 内存层次 流水线这套组合设计得既足够显式以容纳硬件直觉又足够自动以卸载机械劳动谁就更可能成为下一代 AI 编译器叙事的执笔者。结语叙事重写的本质回到开头的追问为什么Tile 化正在重写 AI 编译器的底层叙事因为 AI 编译器的重心正在从自动搜索最优调度迁移到让懂硬件的人把硬件知识写进语言让编译器把它变成正确高效的机器码。TVM 时代tile 是调度手段TileLang 时代tile 是编程模型本身——Shared Memory、register fragment、pipeline 这些硬件实体第一次以一等公民的身份进入 DSL而布局推断、warp specialization、TMA 编排这些曾经只属于 CUTLASS 专家的手艺变成了编译器的日常职责。这条路线在 GEMM、FlashAttention、稀疏注意力、MLA 等大模型核心算子上的反复验证以及它在国产芯片生态中的快速蔓延共同指向一个结论下一轮 AI 编译器的竞争比的不是谁能藏起硬件而是谁能更好地教开发者与硬件对话。而这场对话的第一门语言就是 tile。【免费下载链接】tilelangDomain-specific language designed to streamline the development of high-performance GPU/CPU/Accelerators kernels项目地址: https://gitcode.com/GitHub_Trending/ti/tilelang创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考