TileLang 入门详细教程

更新时间：2026-05-25

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

资料核验：

TileLang docs: https://tilelang.com/
TileLang GitHub: https://github.com/tile-ai/tilelang
TileLang releases: https://github.com/tile-ai/tilelang/releases
TileLang paper: https://arxiv.org/abs/2504.17577
TVM docs: https://tvm.apache.org/docs/

截至 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. 先学什么

建议顺序：

复习 CUDA block、warp、thread、shared memory、register、coalesced load。
跑通 TileLang 安装和官方 examples。
先读 VectorAdd，再读 GEMM，再读 FlashAttention 类例子。
自己写一个 tiled GEMM 骨架。
加 pipeline、autotune、profile。
回头看生成的 IR 和 CUDA 代码。

对应图：00_tilelang_learning_roadmap.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。

常见安装路径：

Code

pip install tilelangpip install git+https://github.com/tile-ai/tilelang.gitgit clone --recursive https://github.com/tile-ai/tilelang.gitcd tilelangpip install . -vpython -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，生成目标代码，再绑定运行时。对应图：01_tilelang_programming_model.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 tilelangimport tilelang.language as Tfrom tilelang import jit@jitdef 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 mainkernel = vector_add(1024)

这段代码的重点：- T.Kernel 建立 block 维度。- thread_binding 建立 thread 维度。- 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。对应图：03_tilelang_gemm_tile_flow.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 和最佳配置。对应图：04_tilelang_tuning_debug_workflow.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：重新生成本目录所有教程文件。

--- title: TileLang 入门详细教程 jupyter: python3 --- 更新时间：2026-05-25 这份教程面向已经学过 CUDA、Triton、MLIR 或高性能算子的人。目标不是背 API，而是建立一个能写、能调、能在面试里讲清楚的 TileLang 心智模型。资料核验： - TileLang docs: https://tilelang.com/ - TileLang GitHub: https://github.com/tile-ai/tilelang - TileLang releases: https://github.com/tile-ai/tilelang/releases - TileLang paper: https://arxiv.org/abs/2504.17577 - TVM docs: https://tvm.apache.org/docs/ 截至 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 代码。对应图：`00_tilelang_learning_roadmap.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。常见安装路径： ```{python} #| language: bash pip install tilelangpip install git+https://github.com/tile-ai/tilelang.gitgit clone --recursive https://github.com/tile-ai/tilelang.gitcd tilelangpip install . -vpython -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，生成目标代码，再绑定运行时。对应图：`01_tilelang_programming_model.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 的基本形状： ```{python} #| language: python import tilelangimport tilelang.language as Tfrom tilelang import jit@jitdef 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 mainkernel = vector_add(1024) ``` 这段代码的重点：- `T.Kernel` 建立 block 维度。- `thread_binding` 建立 thread 维度。- `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。对应图：`03_tilelang_gemm_tile_flow.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 和最佳配置。对应图：`04_tilelang_tuning_debug_workflow.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`：重新生成本目录所有教程文件。