适用读者：写过或要写 GPU 算子的人；被 thread/block/grid、tensor core、shared memory 调优折磨过的工程师。
前置：S23-J01（prefill/decode 的算力/带宽账）、基础的 CUDA 编程模型概念。
覆盖旧文：A-1988 → 15 点（方向 02：面向 Tile 的开发语言、三层抽象、矩阵乘示例、V3.2 采用）。
扩展点：E-S23-008（三层抽象的选择判据）。
上一篇：S23-J03 · 训推分离与协同设计：从"优化"到"设计"
下一篇：S23-A01 · Agent Infra 与超节点。
1. 一句话结论
算子优化的痛点在"数据流"和"调度"缠在一起：你想表达的是"这块数据怎么流动、怎么被复用"，实际写的却是"第几个线程、放哪级缓存、怎么流水"。TileLang 的做法是把两者彻底解耦——你写数据流，编译器管调度。
| 传统 CUDA 写法 | TileLang 写法 |
|---|---|
| 手动定义 thread/block/grid | 只描述 tile 与计算 |
| 手动指定 shared memory/寄存器 | 编译器分配 |
| 手动软件流水（double buffer） | 由 pipeline 声明推导 |
| 优化靠经验 + profile 反复调 | 优化靠编译器策略 + 少量提示 |
2. 三层抽象（§2）
原文给出的三个层级（可混合使用）：
| 层级 | 你写什么 | 编译器做什么 | 适合 |
|---|---|---|---|
| 入门版 | 数据块 + 计算逻辑 | 全部硬件细节与优化 | 新算子快速落地、教学 |
| 进阶版 | 调用封装的算子（如 GEMM） | 组合与调度 | 复用成熟内核、兼顾效率 |
| 精细版 | 线程级细节 | 少量介入 | 极致性能、特定硬件特性 |
"可混合"是关键：同一个 kernel 里可以对不同部分用不同层级——大部分用入门版，热点段落降到精细版。
2.1 一段 TileLang 的形态（示意）
# 形态示意（非可运行代码；真实 API 以项目文档为准）
@tilelang.jit
def matmul(A, B, C, M, N, K, block_M, block_N, block_K):
with tilelang.Kernel([M, N]) as (i, j):
with tilelang.Tile([block_M, block_N]) as (bi, bj):
# 只描述数据流：分块、搬运、累加
As = tilelang.alloc_shared([block_M, block_K], dtype)
Bs = tilelang.alloc_shared([block_K, block_N], dtype)
acc = tilelang.alloc_fragment([block_M, block_N], dtype)
tilelang.clear(acc)
for k in tilelang.Pipelined(K // block_K): # 流水由编译器展开
tilelang.copy(A[bi, k * block_K], As)
tilelang.copy(B[k * block_K, bj], Bs)
tilelang.gemm(As, Bs, acc) # 计算交给内核/编译器
tilelang.copy(acc, C[bi, bj])
说明：这段用于说明"声明数据流 + 交给编译器"的形态。具体 API 名称、装饰器与内存作用域以 TileLang 项目文档为准；本机无 GPU，未实测。
判据：同一段逻辑，如果换成 CUDA 要写 3 倍以上的调度代码，就值得用 TileLang。
3. 为什么"面向 Tile"能提升效率（§2）
原文的实验结论：TileLang 用少量代码即可获得更优的模型效率；并提到 DeepSeek V3.2 官方文档说明部分算子用 TileLang 重构（该说法需按官方文档核对，本系列标 needs_verification）。
机制上有三点：
| 机制 | 说明 | 对性能的影响 |
|---|---|---|
| Tile 级表达 | 以"块"为单位描述搬运与计算 | 编译器可自动做 tiling 与融合 |
| 调度与数据流解耦 | 同一份数据流可换不同调度策略 | 便于搜索（autotune） |
| 流水线声明式 | Pipelined 由编译器展开 double buffer |
减少手写同步错误 |
反例：把所有东西都写成"精细版"，等于回到 CUDA 的复杂度，只是换了语法——失去了这门语言的意义。
4. 与 CUDA / Triton 的关系
| 维度 | CUDA | Triton | TileLang |
|---|---|---|---|
| 抽象层级 | 线程级 | 块级（block-level DSL） | Tile 级 + 三层可下探 |
| 调度控制 | 全手动 | 部分自动 | 声明式，可下探 |
| 学习曲线 | 陡 | 中 | 中（入门版平缓） |
| 适合 | 极致定制、硬件特性 | 快速写高效 kernel | 数据流复杂、需要复用与调优搜索 |
判据：先问"这个 kernel 会不会被反复改"。一次性 kernel → Triton/CUDA 更快出活；会被反复调优的（attention 变体、MoE 路由）→ TileLang 的抽象收益更大。
5. 前提与边界
| 结论 | 成立条件 | 不成立时 |
|---|---|---|
| 少量代码获得更优效率 | 编译器策略与目标硬件匹配 | 冷门硬件/新指令集 → 需要精细版甚至退回 CUDA |
| 三层可混用 | 同一 kernel 内作用域清晰 | 混用过度 → 可读性下降、调优无从下手 |
| 生态可用 | 有对应后端与算子库 | 依赖的库缺失 → 收益打折 |
本机无 GPU，未运行任何 TileLang 示例；本节所有性能结论均来自原文转述，标为
needs_verification。
6. 常见误区
| 误区 | 纠正 |
|---|---|
| "有了 TileLang 就不用懂硬件" | 精细版仍要懂；且不懂 roofline 就选不对策略 |
| "抽象越高越慢" | 抽象影响的是可达性能上限，不是必然更慢；取决于编译器 |
| "写一次就完事" | Tile 参数的搜索（block 大小、stage 数）才是调优主战场 |
| "能跑通就是优化好了" | 要用 roofline 判断是算力瓶颈还是带宽瓶颈（J03 §3） |
7. 自测
- 你最近写的一个 kernel，调度代码占比多少？超过一半吗？
- 你能说出它的算术强度与 roofline 哪一侧是瓶颈吗？
- 如果要把它迁到 TileLang，哪一段必须保留精细版？
8. 下一篇
S23-A01 · Agent Infra 与超节点：把六个方向落到"该不该投" —— 沙箱的四个硬要求、Workflow 与 Agent 的分界、超节点带宽与上下文长度的关系。