TileLang 入门详细教程

更新时间：2026-05-25

这份教程面向已经学过 CUDA、Triton、MLIR 或高性能算子的人。目标不是背 API，而是建立一个能写、能调、能在面试里讲清楚的 TileLang 心智模型。

资料核验：

截至 2026-05-25，TileLang 文档站页面显示版本为 0.1.10，GitHub releases 页面最新稳定标签为 v0.1.9。学习时把文档版本和 release 标签分开看。

1. TileLang 是什么

TileLang 是一个面向高性能 kernel 的 Python DSL。你用 Python 写 tile 级别的程序结构，TileLang 把它 lowering 到 TVM IR，再生成目标后端代码并运行。

用一句话记：

TileLang 让你用接近算法 tile 的方式描述 GPU kernel，同时保留 shared memory、thread、pipeline、tensor core 等性能控制点。

它适合学习和实现这些东西：

  • GEMM、conv、reduction、elementwise fusion。
  • FlashAttention、linear attention、MoE、quantization kernel。
  • 需要从算法 tile 推到 GPU memory hierarchy 的自定义算子。
  • 想理解 TVM/TIR 但又不想一开始就被 schedule API 淹没的人。

2. 先学什么

建议顺序：

  1. 复习 CUDA block、warp、thread、shared memory、register、coalesced load。
  2. 跑通 TileLang 安装和官方 examples。
  3. 先读 VectorAdd，再读 GEMM，再读 FlashAttention 类例子。
  4. 自己写一个 tiled GEMM 骨架。
  5. 加 pipeline、autotune、profile。
  6. 回头看生成的 IR 和 CUDA 代码。
TileLang 入门学习路线
从能跑示例到能讲清 GPU kernel DSL 的完整路径
0 背景
1 安装
2 读例子
复习 CUDA block warp shared memory
理解 TIR 和 tile 的关系
pip 或源码安装
确认 Python CUDA glibc 要求
VectorAdd GEMM FlashAttention
先看 program 再看 generated code
3 写 kernel
4 调优
5 面试输出
用 T.Kernel 定义网格
用 T.Parallel 和 memory scope 表达并行
tile size num threads pipeline stages
autotune 加 profile 验证
能解释 TileLang vs Triton CUDA TVM
能画出 GEMM tile 数据流
每天怎么学
第一天 跑通安装与 VectorAdd
第二天 拆 T.Kernel 和 T.Parallel
第三天 手写 tiled GEMM 骨架
第四天 加 shared memory 和 pipeline
第五天 autotune 并看 profiler
第六天 对比 Triton CUDA TVM
第七天 复盘面试题
目标不是背 API
而是能解释编程模型和性能路径

学习路线 · Excalidraw 源文件

3. 安装和环境

官方文档给出的关键要求包括：

  • Python 3.10 或更高。
  • Linux glibc 2.28 或更高。
  • CUDA 要求分两类：宿主机 CUDA 10 或更高，或者 pip 提供的 CUDA toolchain 13 或更高。
  • Windows 源码构建需要 Python、CMake、Visual Studio Build Tools 和 MSVC C++ toolchain。

常见安装路径：

pip install tilelang
pip install git+https://github.com/tile-ai/tilelang.git
git clone --recursive https://github.com/tile-ai/tilelang.git
cd tilelang
pip install . -v
python -c "import tilelang; print(tilelang.__version__)"

本仓库当前没有把 TileLang 作为依赖安装；这份教程提供学习路径和图解，不声明本机已经成功运行 TileLang kernel。

4. 编程模型

TileLang 程序通常分三层理解：

  • Python 函数层：用 @tilelang.jit 包住一个返回 T.prim_func 的函数。
  • Tile DSL 层：用 T.Kernel、T.Parallel、T.copy、T.gemm、memory scope 表达 tile 级计算。
  • 编译运行层：TileLang lowering 到 TVM IR，生成目标代码，再绑定运行时。
TileLang 编程模型总览
Python DSL 前端到 TVM IR 再到可运行 kernel
用户写的部分
@tilelang.jit
Python 函数返回 T.prim_func
Tile 级 DSL
循环
T.Parallel T.serial T.Pipelined T.unroll
编译和运行
Lowering 到 TVM IR
检查结构是否符合硬件映射
T.Tensor 和 T.Buffer
描述形状 dtype 与输入输出
存储层级
global shared local fragment
生成目标代码
CUDA HIP CPU 或其他后端
T.Kernel grid threads
描述 CTA 网格和线程块
算子原语
T.copy T.gemm T.clear T.reduce_sum
Runtime 调用
传入 torch tensor 并执行
三层心智模型
Beginner
用模板写出正确 kernel
先让数据路径可解释
Developer
控制 tile 形状、线程绑定、共享内存、pipeline
开始做性能调优
Expert
检查 IR、生成代码、tensorize、layout
把 DSL 映射到具体 GPU 指令

编程模型 · Excalidraw 源文件

你读一个 TileLang kernel 时，先找这五个东西：

  1. 输入输出 tensor 的 shape 和 dtype。
  2. T.Kernel 的 grid 和 threads。
  3. 每个 CTA 负责哪一块输出 tile。
  4. global、shared、local、fragment 的数据搬运路径。
  5. 计算 primitive 和 loop pipeline。

5. 最小 VectorAdd 思路

VectorAdd 的学习价值不是性能，而是确认 DSL 的基本形状：

Code 
import tilelang
import tilelang.language as T
from tilelang import jit


@jit
def vector_add(n, block_size=256, dtype="float32"):
    @T.prim_func
    def main(a: T.Tensor((n,), dtype),
             b: T.Tensor((n,), dtype),
             c: T.Tensor((n,), dtype)):
        with T.Kernel(T.ceildiv(n, block_size), threads=block_size) as bx:
            for tx in T.Parallel(block_size):
                i = bx * block_size + tx
                if i < n:
                    c[i] = a[i] + b[i]
    return main


kernel = vector_add(1024)

这段代码的重点：

  • T.Kernel 建立 block 维度。
  • T.Kernel(..., threads=block_size) 指定每个 block 的线程数，T.Parallel 声明 block 内的并行循环，由编译器映射到线程。
  • i < n 是边界保护。
  • 先跑小 shape 对齐 torch，再扩到大 shape。

6. GEMM 怎么学

GEMM 是 TileLang 入门的核心练习，因为它把所有重要概念都串起来：

  • A 和 B 从 global memory 搬到 shared memory。
  • 每个 CTA 负责 C 的一个 BM x BN tile。
  • K 维按 BK 切块循环。
  • fragment accumulator 累加。
  • pipeline 尝试隐藏 global load 延迟。
  • autotune 搜索 BM、BN、BK、threads、stages。
Tiled GEMM 数据流
用一张图理解 TileLang 为什么要显式表达 tile、shared memory 和 pipeline
全局内存
A matrix
M x K
B matrix
K x N
C matrix
M x N
一个 CTA 的工作
A tile
BM x BK
shared
B tile
BK x BN
shared
Accumulator
BM x BN
fragment
K 维主循环
for ko in K tiles
load A tile and B tile
pipeline stages
当前 tile 计算
下一 tile 预取
store C tile
写回 global
性能要点
tile size 太小
访存复用不足
launch 和同步开销占比高
tile size 太大
shared memory 或 register 压力升高
occupancy 下降
pipeline 合理
隐藏 global memory 延迟
提高 tensor core 利用
边界处理清晰
非整除 shape 也能跑
先正确再调优

GEMM tile 数据流 · Excalidraw 源文件

写 GEMM 时先画数据流，再写代码。不要先陷入 API 细节：

  1. 画出 C tile。
  2. 反推这个 C tile 需要哪块 A tile 和 B tile。
  3. 决定 A/B tile 怎么搬到 shared。
  4. 决定 K loop 怎么展开或 pipeline。
  5. 决定 accumulator dtype。
  6. 最后才调 tile size。

7. 调优与调试

TileLang 的调优不应该从随机改参数开始。推荐 workflow：

  1. 写 baseline，只求正确。
  2. 用小 shape 和 PyTorch 结果对齐。
  3. 打印或检查生成代码，确认索引和边界。
  4. 加 shared memory 复用。
  5. 加 pipeline。
  6. 用 autotune 枚举关键参数。
  7. 用 profiler 看瓶颈。
  8. 固化 shape family 和最佳配置。
TileLang 调优与调试工作流
从正确性到性能，再回到 IR 和 profiler 定位瓶颈
写 baseline
打印和检查
加 pipeline
先写最直观 tile kernel
用小 shape 对齐 torch 结果
用 T.print 或生成代码
确认索引和边界
T.Pipelined 控制阶段
检查 shared memory 生命周期
autotune
profile
固化模板
枚举 block size num threads
记录 latency 和 best config
Nsight 看 occupancy memory tensor core
不要只看单次耗时
把 shape family 和配置记录下来
形成可复用 kernel
常见失败模式
结果错误
索引越界、mask 缺失、dtype 累加错误
性能不涨
tile 形状不匹配硬件、访存未合并、pipeline 没隐藏延迟
只会调参数
没有先画数据流
不知道瓶颈在 HBM shared 还是 register

调优与调试 · Excalidraw 源文件

常见坑：

  • tile size 太大导致 register 或 shared memory 压力过高。
  • tile size 太小导致访存复用不足。
  • pipeline stage 加了，但 load 和 compute 没真正重叠。
  • 只看 latency，不看 occupancy、memory throughput、tensor core utilization。
  • 边界 shape 没测，只测了能整除的矩阵。

8. 和 CUDA、Triton、TVM 怎么比较

面试里可以这样讲：

  • CUDA：控制最细，工程样板最多。适合极限手工优化。
  • Triton：Python 体验成熟，面向 block program，常用于 PyTorch 自定义算子。
  • TVM：编译体系完整，schedule 和后端覆盖广，但学习曲线更陡。
  • TileLang：把 tile 级 DSL、TVM IR、GPU kernel 控制点放在一起，适合用较高层表达写出可调优 kernel。

不要说 TileLang 一定替代 Triton 或 CUDA。更稳的说法是：

TileLang 适合在自定义高性能 kernel 中更清楚地表达 tile 数据流，并利用 TVM 编译基础设施做 lowering、生成和调优。

9. 七天练习计划

第一天：安装、跑 VectorAdd、理解 T.Kernel。

第二天：读 GEMM example，把 block、thread、shared、fragment 标出来。

第三天：自己写一个最小 tiled GEMM 骨架，不追求快。

第四天：加 shared memory staging，比较 global-only 与 shared 版本。

第五天：加入 pipeline，观察代码结构变化。

第六天：autotune tile 参数，记录不同 shape 的 best config。

第七天：用 05_tilelang_interview_map.excalidraw 复盘面试题。

10. 面试问答

Q：TileLang 的核心抽象是什么？

A：用 tile 级 DSL 描述 kernel 数据流和调度意图，包括 grid、thread、memory scope、copy、gemm、pipeline，再经由 TVM IR lowering 到目标代码。

Q：为什么要先画 GEMM tile 数据流？

A：因为性能来自数据复用和硬件映射。你必须先知道每个 CTA 算哪块 C、需要哪块 A/B、怎么进入 shared、在哪个 fragment 累加，才能合理调参数。

Q：autotune 调什么？

A：主要调 tile shape、threads、pipeline stages、memory layout 等参数。目标是在 occupancy、register pressure、shared memory、memory bandwidth、tensor core 利用之间平衡。

Q：什么时候不用 TileLang？

A：如果只需要一个很普通的 PyTorch op，就不需要自定义 kernel。如果团队已经有成熟 Triton/CUDA 模板，并且 TileLang 后端或部署环境不匹配，也不应强行迁移。

11. 文件清单

  • 00_tilelang_learning_roadmap.excalidraw：学习路线。
  • 01_tilelang_programming_model.excalidraw：编程模型。
  • 02_tilelang_kernel_anatomy.excalidraw：kernel 解剖。
  • 03_tilelang_gemm_tile_flow.excalidraw：GEMM tile 数据流。
  • 04_tilelang_tuning_debug_workflow.excalidraw：调优与调试。
  • 05_tilelang_interview_map.excalidraw：面试速查。
  • generate_tilelang_tutorial.py：重新生成本目录所有教程文件。