TileLang Skill:80行Python替代500行CUDA,GPU内核开发的范式革命

TileLang Skill:80行Python替代500行CUDA,GPU内核开发的范式革命

TileLang 是一个基于 TVM 编译器基础设施的领域特定语言(DSL),专门用于简化 GPU/CPU/加速器高性能内核的开发。它用 Pythonic 语法描述内核逻辑,由编译器自动完成线程分配、内存管理、Warp 特化等底层工作,让开发者专注于数据搬运和计算逻辑本身,而非硬件细节。TileLang 最初从北京大学实验室走出,2026 年 1 月被 ICLR 接收为 Oral Paper,同年 DeepSeek V3.2-Exp 在训练中原生采用,GitHub star 数已突破 7,900,是今年 GPU 编程领域最值得关注的新兴基础设施项目之一。

TileLang 的设计原则是”让正确性证明和性能优化同源”。它集成 Z3 定理证明器在编译器级别做形式化验证,同时通过 Tile 级别的抽象让编译器自动推导出内存布局(Layout)、流水线(Pipeline)和 Warp 特化策略,开发者无需手动调优即可接近硬件峰值性能。

截至 2026-10-01,TileLang 累计约 7,928 Stars(近 30 天增长 +618),Fork 799,社区活跃维护(最新 commit 2026-10-01),GitHub trending 全球排名约 #10(本周),已获 ICLR 2026 Oral 论文认可,被 DeepSeek V3.2-Exp 正式用于训练内核编写。

GitHub 链接:https://github.com/tile-ai/tilelang

TileLang 项目主要开发者为北京大学杨智教授课题组的王磊(LeiWang1999),核心团队成员还包括 chengyupku、nox-410 等,部分工作在微软研究院完成。LeiWang 同时是 BitBLAS(混合精度计算库)的作者,TileLang 是其从 2024 年开始的核心项目,2025 年 1 月正式开源。

TileLang 是一种嵌入 Python 的 DSL,其编译器后端基于 Apache TVM。用户用 Python 书写内核描述,TileLang 将其转换为 TVM IR,再通过 TVM 的优化和代码生成管道输出目标平台(CUDA、ROCm、Metal、WebGPU)的机器码。

架构上分为三层:前端语法层(Pythonic Tile 描述语言)→ 中间表示层(TVM IR)→ 后端代码生成层(多后端代码输出)。开发者用”方块”视角描述数据在全局内存→共享内存→寄存器的搬运路径,告诉编译器”把这块数据放到共享内存”、”在寄存器里做矩阵乘法”,具体的线程分配和内存排列由编译器自动完成。

与 CUDA/Triton 的直观对比:

  • CUDA:需要手动管理线程、Warp、共享内存、同步、指令排程,几乎全部由开发者控制
  • Triton:Block 级别运算逻辑由开发者定义,编译器负责部分内存管理和调度
  • TileLang:开发者只需描述 Tile 级数据流和内存放置(Layout),编译器自动推导 Layout、流水线、Warp 特化——相当于”带主动安全辅助的自动挡”,同时支持随时切回手动模式

核心特点:

  • Pythonic 语法,门槛低:约 80 行 TileLang 代码即可实现与 500 行手写 CUDA 相当的 FlashMLA 性能
  • 多后端支持:CUDA、ROCm、Metal、WebGPU,以及华为昇腾等国产 GPU
  • 形式化验证:集成 Z3 定理证明器,在编译器级别做正确性约束
  • 2:4 稀疏张量核心支持:原生支持结构性稀疏加速
  • ICLR Oral 认证:2026 年 1 月被 ICLR 接收,学术背书扎实
  • DeepSeek 生产验证:DeepSeek V3.2-Exp 在训练中原生采用,已进入顶级模型基础设施

安装与基本使用:

# 克隆仓库
git clone https://github.com/tile-ai/tilelang.git
cd tilelang

# 安装依赖(需要 TVM,可参考 README 中的 docker 方式)
pip install tilelang

# 编写一个简单的 GEMM 内核(TileLang Pythonic 语法)
import tilelang as tl

@tl.tile(shape=[128, 128], threads=128)
def gemm(A, B, C, M, N, K):
    ...

# 编译生成 CUDA 代码
gemm.compile(target="cuda")

TileLang 提供 Beginner / Developer / Expert 三种编程模式,新手从 Beginner 入手可直接写出可用内核,进阶后可深入底层调优。

使用场景与目标用户:

TileLang 适合需要开发高性能 GPU 内核的团队和个人,具体包括:

  • AI 框架开发者:需要自定义 Attention Kernel、GEMM、MLA 等计算密集算子
  • 高性能计算工程师:在 GPU 上做科学计算、信号处理、图像处理等
  • 硬件厂商工具链团队:快速在国产 GPU(昇腾、摩尔线程 MUSA 等)上构建算子库
  • 学术研究者:用 TileLang 研究新算子架构,TileLang Puzzles 提供了交互式学习工具

对 AI 开发者而言,TileLang 的核心价值在于:当你想优化某个 Attention 变体时,不需要手写 500 行 CUDA——用 TileLang 描述清楚数据流,编译器帮你把性能做到接近专家级。

输入与输出案例:

输入:用 TileLang 写一个 FlashMLA kernel(简化示意)

import tilelang as tl

@tl.tile(shape=[128, 64], threads=128, dtype="float16")
def flash_mla(Q, K, V, O, ...):
    # 描述 tile 级别的数据搬运和注意力计算
    q_tile = Q.load(tile_id)
    k_tile, v_tile = K.load(tile_id), V.load(tile_id)
    attn_score = tl.matmul(q_tile, k_tile.transpose())
    attn_score = tl.softmax(attn_score)
    O.store(tile_id, tl.matmul(attn_score, v_tile))

输出:编译器自动生成接近手写 CUDA 的高效机器码,在 A100 上的实际测试中,80 行 TileLang 达到了与 500 行手写 CUDA 相当的性能(FlashAttention 专家库级别),同时保留了可读性和可维护性。

另一个案例:DeepSeek V3.2-Exp 训练流程中,其核心 Attention Kernel 使用 TileLang 编写,在华为昇腾和英伟达 GPU 上均达到硬件峰值性能,开发效率提升数倍。


GitHub: https://github.com/tile-ai/tilelang

评论区

0 条评论

登录后可评论。

Skill超级捕获手 919 阅读