想象一个场景:
DeepSeek在2025年初放出了FlashMLA代码,那是他们高度优化的多头潜在注意力kernel,在H100上性能号称全球最强。工程团队花了大量时间手工调优SASS汇编、优化内存访问模式、压榨每一个CUDA指令的吞吐。
然后TileLang出现了。
他们用80行Python代码,实现了同样的MLA Decoding性能——与FlashMLA在H100上完全持平。
这不是在吹牛。代码就放在GitHub上,benchmark图也在。
今天要聊的,就是这个项目。
01
GPU kernel开发,是AI基础设施里最苦的活。
说它是"屠龙之术"也不为过——会的人极少,门槛极高,需要对硬件架构有极其深入的理解。CUDA编程、内存层级、向量化、Tensor Core调度……每一个环节都有大量可优化但也极容易出错的细节。
但它的重要性,不成正比。
一个优化良好的kernel,和一个普通kernel,在大模型推理场景下,性能差距可以达到2-5倍。换句话说,给定一块H100,前者能跑出的有效算力,是后者的两到五倍。在大模型军备竞赛的背景下,这是巨大的成本差异和竞争优势。
所以有实力的大厂,都在花重金养kernel团队。字节、阿里、DeepSeek,都有自己的kernel优化团队,他们写的代码不对外公开,是核心竞争力之一。
问题在于:这种能力,是大厂的专属。
对于中小团队、研究者、甚至是大厂的普通工程师来说,写一个高性能kernel,门槛高得离谱。你要么是NVidia/AMD官方的kernel团队,要么是经过多年踩坑磨练出来的极少数人。
TileLang想要改变这个局面。
02
TileLang的核心思路很简单:用Python的语法,做汇编级别的事。
它是一个领域特定语言(DSL),构建在TVM编译器之上。开发者用Python风格的API描述计算逻辑,TileLang再将其编译成目标硬件的高效指令。
举一个GEMM(矩阵乘法)的例子,这是AI计算的核心操作:
import tilelang as tl
@tl.kernel
def gemm(A: tl.F32[1024, 512], B: tl.F32[512, 1024], C: tl.F32[1024, 1024]):
tile_m = tl.thread_axis("blockIdx.x")
tile_n = tl.thread_axis("blockIdx.y")
tile_k = tl.thread_axis("blockIdx.z")
local_mat = tl.compute((64, 64), lambda ij: ...)
# ... 核心计算逻辑
这段代码看起来像普通的Python,但它最终会被编译成针对特定GPU优化过的机器码。
TileLang负责处理那些开发者不愿意手工处理的细节:内存访问模式如何对齐、Tensor Core如何调度、异步复制如何掩盖计算延迟……
换句话说,它把硬件的"暗知识"——那些需要多年经验才能积累的对硬件行为的直觉——变成了编译器的自动优化。
03
TileLang目前支持的硬件,已经相当广泛:
NVIDIA GPU:H100(SOTA支持,自动TMA/WGMMA)、A100、V100、RTX 4090、RTX 3090、RTX A6000 AMD GPU:MI250(自动MatrixCore)、MI300X(异步复制支持) Apple:Metal(2025年新增) 华为:AscendC和AscendNPU IR后端 其他:WebGPU(2025年2月新增)
这个覆盖范围很有意思。
CUDA生态是NVIDIA的天下,但AMD MI300X在大规模推理场景里正在快速崛起(性价比高),苹果的M系列芯片在端侧AI推理里优势明显,华为Ascend是国内绕不过去的选择。
TileLang把目光放在整个GPU生态,而不只是NVIDIA——这是一个务实的选择,也反映出项目团队的野心。
04
最能说明TileLang实力的,是几个具体的实现案例。
DeepSeek MLA Decoding:80行Python代码,在H100上达到与FlashMLA同等的性能。MLA(Multi-Head Latent Attention)是DeepSeek-V2提出的一种注意力机制变体,被认为是当前最高效的大模型注意力方案之一。用80行代码复现全球最强的kernel性能,这个故事本身就很有传播力。
FlashAttention:TileLang的FlashAttention实现在H100上的性能曲线,与官方FlashAttention基本一致。FlashAttention大家应该比较熟悉,是目前大模型训练和推理中最广泛使用的注意力优化技术。
2:4稀疏Tensor Core支持:2025年7月新增的T.gemm_sp,支持NVIDIA的2:4稀疏加速。这一特性可以让稀疏模型的计算速度提升2倍而精度损失极小。需要在硬件和算法层面同时支持,目前支持范围有限,但方向很清晰。
AMD MI300X FlashMLA:2025年4月的更新,让AMD卡上也实现了FlashMLA级别的高性能。这对AMD生态意义重大——AMD在硬件上对标NVIDIA,但软件生态一直是短板。TileLang在AMD上的优化填补了一个关键空白。
05
TileLang背后的技术底座是Apache TVM。
TVM是一个成熟的深度学习编译器栈,可以将模型 lowering 到各种硬件后端。它在业界有广泛的应用,亚马逊(MXNet时代的SageMaker)、Facebook(当时的PyTorch/JIT)、百度(XLM)都有TVM相关的投入。
TileLang在TVM之上,构建了自己的一层DSL抽象。它的优势在于:
1. 编译时优化
集成了Z3定理证明器(2025年12月),用于SMT-based符号推理,可以在编译阶段做正确性验证和更激进的优化。这是编译器层面的能力提升,不是简单的封装。
2. 多后端支持
同一个TileLang程序,可以编译到不同的硬件后端。2025年新增的CuTeDSL后端(NVIDIA CUTLASS CuTe DSL)让TileLang可以直接生成CUDA的底层代码,而不是通过TVM的标准路径。这意味着更少的抽象损耗,更直接地控制硬件。
3. 低编译开销
2025年10月迁移到apache-tvm-ffi,显著降低了CPU开销。对于需要频繁重编译kernel的开发场景,这是很重要的改进。
06
TileLang目前还年轻——v0.1.0在2025年2月才发布,当前是v0.1.6.post2。
作为一个开源项目,它的活跃度相当不错:持续的版本更新、多硬件后端支持、还有配套的TileLang Puzzles(10个循序渐进的学习谜题,帮助新手上路)。
但也有明显的挑战。
学习曲线:虽然比手写CUDA容易太多,但对没有并行计算背景的开发者来说,tile/thread/warp这些概念依然有门槛。不是"零基础"就能上手的。
调试困难:GPU kernel的调试历来是难题。TileLang提供了T.print调试工具和内存布局可视化工具,但相比传统开发方式,调试体验仍有差距。
生态建设:目前TileLang的用例主要来自官方示例和benchmark。如何吸引第三方开发者贡献kernel库、如何建立类似cuDNN的kernel生态,是决定它能走多远的关键。
07
TileLang代表了一个趋势:AI Infra的下沉。
大模型的能力在快速商品化,但AI Infra的门槛依然高企。GPU kernel开发、编译器优化、硬件适配……这些工作目前还是少数人的专利。
TileLang正在尝试用"抽象+编译器优化"的方式,把这些能力从"大厂专属"变成"开发者可及"。虽然还不能完全替代手工kernel,但它让更多人有能力去够到这个层面的优化。
未来如果TileLang能建立起活跃的社区和丰富的kernel库,AI开发者在性能优化上的选择会多很多——不必依赖大厂开源的kernel,自己也能写出高效实现。
这是有价值的。
而且,80行代码比肩FlashMLA这个故事,值得让更多人知道。
项目链接:https://github.com/tile-ai/tilelang 安装:pip install tilelang 入门:GitHub上有TileLang Puzzles,适合边学边做
全文完。
调研 & 撰写: AI(Claude) 来源: TileLang GitHub仓库、官方文档及版本更新日志 发布时间: 2026年5月5日