--- title: TileLang 入门教程与 Excalidraw 图解 jupyter: python3 --- 更新时间：2026-05-25 本目录是一套 TileLang 入门到面试复盘教程，包含详细 Markdown 教程和 6 张标准 `.excalidraw` 图。 ## 文件 1. [`TUTORIAL`](TUTORIAL.qmd)：完整文字教程。 2. `00_tilelang_learning_roadmap.excalidraw`：学习路线。 3. `01_tilelang_programming_model.excalidraw`：编程模型。 4. `02_tilelang_kernel_anatomy.excalidraw`：kernel 解剖。 5. `03_tilelang_gemm_tile_flow.excalidraw`：GEMM tile 数据流。 6. `04_tilelang_tuning_debug_workflow.excalidraw`：调优与调试。 7. `05_tilelang_interview_map.excalidraw`：面试速查。 8. `generate_tilelang_tutorial.py`：重新生成教程和图。 ## 资料来源 - 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。 ## 练习题 1. 把同一个 GEMM tile 分别用 CUDA thread/block、Triton program、TileLang `T.Kernel` 三种语言描述，比较抽象粒度。 2. 修改 `BM/BN/BK/num_threads/num_stages`，记录 shared memory、寄存器压力和 latency 的变化。 3. 给一个非整除矩阵 shape，补齐边界保护，并用小 shape 对齐 PyTorch baseline。 4. 读生成代码或 IR，标出 global -> shared -> fragment -> global 的每一次数据移动。 ## 面试题 1. TileLang 为什么强调 tiled dataflow，而不是让用户直接写每个 thread 的计算？ 2. `T.copy`、`T.gemm`、`T.Pipelined` 分别对应 GPU kernel 里的哪些性能关键路径？ 3. TileLang 和 Triton 都能写 GEMM，二者在抽象层、编译基础设施和调优方式上有什么不同？ 4. 如果 TileLang kernel 性能不如预期，你会先检查 tile shape、memory scope、pipeline、生成代码还是 profiler 证据？ ## 图示 ```{=html}
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 源文件

```{=html}
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 源文件

```{=html}
TileLang Kernel 解剖图
把一段 TileLang 程序拆成 shape、grid、loop、memory、primitive
程序骨架
def kernel(M, N, K, dtype):
@T.prim_func
def main(A, B, C):
with T.Kernel(grid, threads):
形状和 dtype
决定 buffer 维度
决定 index 计算边界
grid 和 threads
决定 CTA 数量
决定每个 CTA 的并行资源
GPU 映射
Block 或 CTA
处理一个 tile
例如 C 的 BM x BN 子矩阵
Threads
T.Parallel 负责并行索引
threadIdx 负责合作搬运或计算
Memory scopes
global 到 shared 到 fragment
减少 HBM 往返
Primitives
copy gemm reduce clear
表达常见 tile 操作
写程序时按这个顺序检查
1
每个 CTA 负责哪块输出
2
输入 tile 从哪搬到 shared
3
K 维怎么循环和 pipeline
4
fragment 上怎么累加
5
边界和 dtype 是否正确
6
调优参数是否可枚举
```

kernel 解剖 · Excalidraw 源文件

```{=html}
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 源文件

```{=html}
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 源文件

```{=html}
TileLang 面试速查图
把学习内容收束成可讲清楚的问答框架
一句话定义
TileLang 是面向高性能算子的 Python DSL
核心是用 tile 级抽象表达 GPU kernel
再经由 TVM 体系 lowering 到目标代码
和其他工具比较
vs CUDA
CUDA 控制最细 但样板多
TileLang 更接近算法 tile 表达
vs Triton
Triton 编程体验更成熟
TileLang 更强调 TVM IR 和调度体系衔接
必须能画出的图
GEMM tile
A/B global 到 shared
fragment 累加
C 写回
FlashAttention tile
Q block 扫 K/V tile
online softmax 更新 m l O
调优循环
baseline correctness
autotune profile 固化
编译路径
Python DSL 到 TIR
到 CUDA source
到 runtime
常见追问
为什么要 shared memory
减少 global memory 访问
提升 tile 复用
pipeline 解决什么
隐藏 global load 延迟
让搬运和计算重叠
autotune 调什么
tile size threads stages
目标是延迟和资源平衡
怎么验证
小 shape 对齐 torch
大 shape profile
再看生成代码
回答模板
先讲算法 tile
再讲 GPU 映射
最后讲调优和验证
不要一上来背 API
```

面试速查 · Excalidraw 源文件