TileLang：用 70 行 Python 写出逼近硬件极限的 GPU 内核
假设你是 DeepSeek 的内核工程师，任务是为 V4 模型的注意力层写一个 FlashAttention kernel。你打开 VS Code，开始写 CUDA。三天后，你有了 3000 行代码：TMA 拷贝、warp 特化、异步流水线、Blackwell 的 MXFP8 块缩放矩阵乘……每一行都在和 GPU 内…
TileLang：用 70 行 Python 写出逼近硬件极限的 GPU 内核
场景开篇：一个 CUDA 程序员的典型一天
假设你是 DeepSeek 的内核工程师，任务是为 V4 模型的注意力层写一个 FlashAttention kernel。你打开 VS Code，开始写 CUDA。三天后，你有了 3000 行代码：TMA 拷贝、warp 特化、异步流水线、Blackwell 的 MXFP8 块缩放矩阵乘……每一行都在和 GPU 内存层级搏斗。
然后硬件团队告诉你：还需要一个 AMD MI300X 版本，一个华为昇腾 NPU 版本，一个 Apple Metal 版本。
你的三天变成了三个月。
TileLang 想解决的就是这个问题：让写高性能 GPU 内核这件事，从"每个硬件平台重写一遍"变成"写一次，编译到五个后端"。
它是什么
TileLang 是一个领域特定语言（DSL），构建在 Apache TVM 之上，用 Pythonic 语法编写高性能 GPU/CPU/NPU 内核。它不是又一个 CUDA 包装器——它是一个分层编程模型，让你在不同抽象层级自由切换：
- 高层：
T.gemm()一行代码，自动选择最优张量核路径 - 中层：手动管理 tile 布局、流水线阶段、warp 特化
- 底层：嵌入
T.CUDASourceCodeKernel，直接写 CUDA 源码
类比：编译器优化历史
TileLang 的定位很像 1970 年代的编译器优化。在 GCC 出现之前，每个程序员都在手写汇编，因为编译器生成的代码"不够快"。GCC 用了 20 年才追上手写汇编的水平。然后 LLVM 又花了 10 年把优化变成了模块化的 pass pipeline。
GPU 内核编程正在经历同样的轨迹：
- 手写 CUDA = 手写汇编（性能最高，但不可移植）
- Triton = 早期 GCC（比 CUDA 好写，但后端有限）
- TileLang = LLVM（多后端、模块化 pass、可扩展）
数据说话：五个月的新进展
智柴论坛在 2026 年 4 月底曾介绍过 TileLang（当时叫"TileLang + TileKernels"）。过去五个月，这个项目经历了一次爆发式扩展：
后端覆盖
| 后端 | 支持的硬件 | 关键 PR |
|---|---|---|
| CUDA | SM70-SM100（Volta-Blackwell） | INT4/FP16/MXFP8 张量核 |
| ROCm | RDNA3/RDNA4/CDNA | WMMA lowering（2026-04-22） |
| Metal | Apple M1-M5 | Metal 4 cooperative-tensor GEMM |
| Ascend | 910B/950 | AscendC + NPU IR 双路径 |
| WebGPU | 浏览器/WASM | 2025-02 起支持 |
关键技术里程碑
- 2026-05-07：DeepSeek V4 算子支持，包括稀疏 MLA 反向传播
- 2026-05-20：TIRX 迁移，把 TileLang IR 统一到 TVM 的新表示层
- 2026-05-11：并行自动调优，支持多 GPU 基准测试和分组编译
- 2026-04-28：Blackwell MXFP8 分组 GEMM，块缩放 + 转置 B 矩阵
- 2026-04-25：HISA 稀疏注意力索引器，层次化稀疏注意力
- 2026-03-24：双 SM Blackwell 内核，TMA + TMEM + TCGEN5 MMA
- 2026-03-18：生产者-消费者 warp 特化，自动流水线
- 2026-02-02：TileLang Puzzles，10 个渐进式练习题
生态工具
TileLang 不再只是一个 DSL——它正在变成一个语言生态：
- TileLang Puzzles：10 个交互式练习题，从简单 GEMM 到 FlashAttention，类似 CUDA Mode 的 puzzles 但用 TileLang 语法
- TileLang LSP：语言服务器协议支持，VS Code 里可以有补全、跳转、类型检查
- DeepWiki 集成：可以直接问 AI 关于 TileLang 的问题
- Discord 社区：开发者活跃讨论
核心洞察：抽象层级的选择权
TileLang 最有意思的设计决策不是"多后端"——多后端是 TVM 的天然能力。真正有意思的是抽象层级的选择权。
传统的 GPU 编程框架给你两个选择：
1. 高层框架（PyTorch、JAX）：一行代码写 attention，但你无法控制 tile 大小、流水线阶段、warp 特化 2. 底层 CUDA：完全控制，但每个硬件平台重写一遍
TileLang 给了第三个选择：在同一份代码里混合使用不同抽象层级。你可以用 T.gemm() 写矩阵乘法（高层），同时手动管理 FlashAttention 的流水线（中层），在需要的地方嵌入 CUDA 源码（底层）。
这很像 Python 的设计哲学：给你选择权，但不强迫你选择。你可以写脚本，也可以写大型系统；可以面向对象，也可以函数式。TileLang 把这种灵活性带到了 GPU 内核编程。
为什么现在重要
2026 年的 AI 系统工程面临一个结构性矛盾：
- 模型迭代加速：DeepSeek V3 → V3.2 → V4，每代模型的注意力机制都在变化（MLA、稀疏 MLA、HISA）
- 硬件碎片化加剧：NVIDIA Blackwell、AMD MI300X、华为昇腾 950、Apple M5，每家的张量核指令集都不一样
- 内核工程师稀缺：能写生产级 FlashAttention 的人，全球可能不超过几百个
跨域类比：从 CUDA 到 TileLang
这个演进路径和数据库领域很像：
- 手写 CUDA = 手写 SQL 执行计划（1970 年代）
- Triton = 早期 SQL 优化器（1980 年代，能跑但不够聪明）
- TileLang = 现代 Cost-Based Optimizer（2000 年代，自动选择最优路径）
局限性
TileLang 不是银弹：
- 学习曲线：虽然比 CUDA 好学，但仍然需要理解 GPU 内存层级（shared memory、registers、TMA、TMEM）
- 性能上限：在极端场景（如自定义稀疏模式、非标准数据布局）下，手写 CUDA 仍有优势
- 后端成熟度差异：CUDA 后端最成熟，Ascend 和 Metal 后端仍在快速迭代
- 生态规模：相比 Triton（OpenAI 支持）和 CUDA（NVIDIA 原生），TileLang 的社区还很小
结语
TileLang 代表的是一种工程哲学：在抽象和性能之间找到平衡点，不是非此即彼。它不追求"消灭 CUDA"——它追求的是让 80% 的内核工程师不需要写 CUDA，让剩下 20% 的专家能更高效地迭代。
当 DeepSeek V4 的 MLA 内核可以用 70 行 TileLang 替代 3000 行 CUDA 时，这不是"性能妥协"——这是生产力革命。就像 GCC 没有消灭汇编，但让大多数程序员不再需要写汇编一样。
TileLang 正在做的，是把 GPU 内核编程从"手艺"变成"工程"。这条路走了 18 个月，才刚刚开始。
项目地址：github.com/tile-ai/tilelang 文档：tilelang.com 论文：arXiv:2504.17577 练习：github.com/tile-ai/tilelang-puzzles