TileLang 编程基本知识点

TileLang 的定位可以用一句话概括：把 CUDA kernel 里「怎么切 tile、数据放哪级存储、什么时候预取」这几件事变成显式的 Python 语句，其余的寄存器分配、指令选择、同步插入交给编译器。它不是又一个 Triton–Triton 隐藏 shared memory，TileLang 让你直接写 T.alloc_shared。这个差别决定了两者能碰到的性能天花板不同。

本文整理 TileLang 的编程基本知识点：它是什么、编程模型的 5 个原语、循环原语与 T.Pipelined 的工作原理、同一份代码在不同 GPU 架构上如何 lowering、四个进阶主题（Split-K / warp specialization / autotune / Blackwell 两条路径），以及写完 kernel 之后的三个标准动作（验证 / dump 源码 / bench）。

---

一、TileLang 是什么

项目	现状（2026-08）
版本	v0.1.13（2026-08-03 发布）
GitHub	tile-ai/tilelang，7.2k stars
底层	TVM（IR 已迁移到 TIRX）
Python	≥ 3.10
主力后端	CUDA（SM70~SM120）
其他后端	ROCm/HIP、Apple Metal、LLVM CPU（实验）、CuTe DSL（实验）、WebGPU（实验）
生态后端	华为 Ascend、沐曦 MACA、摩尔线程 MUSA（独立仓库维护）

出身是学术项目：主要由 LeiWang1999、chengyupku、nox-410 在北大杨智教授指导下开发，部分工作在 MSRA 实习期间完成。2025-01 开源。

值得注意的是上游模型厂在用它：TileLang 仓库的 examples/ 里有 deepseek_mla、deepseek_v32、deepseek_v4、deepseek_mhc 四个目录，DeepSeek 系列的 MLA / 稀疏注意力 / mHC 融合 kernel 都有 TileLang 参考实现。这意味着读 TileLang examples 等于读一份最新算子的可执行论文附录–这是它相比 Triton 的一个实际优势。

一句话总结：TileLang = Pythonic 语法 + 显式 tile/memory 层级控制 + TVM 编译基础设施，目标是「写起来像 Triton，控制力接近 CUTLASS」。

---

二、编程模型：5 个原语撑起全部

TileLang 的 API 面很窄，这是刻意的。一个完整 kernel 基本只用这 5 类原语：

T.Kernel(grid_x, grid_y, threads=N)     ← ① 定义 grid / block，拿到 block index
    │
    ├── T.alloc_shared(shape, dtype)    ← ② 显式声明 shared memory buffer
    ├── T.alloc_fragment(shape, dtype)  ← ②' 显式声明 register fragment（累加器）
    │
    └── for k in T.Pipelined(N, num_stages=3):   ← ③ 软件流水（自动双/三缓冲）
            T.copy(global_tile, shared)  ← ④ 数据搬运（按架构 lower 到 cp.async / TMA）
            T.gemm(A_s, B_s, C_frag)     ← ⑤ tile 级 MMA（映射到 Tensor Core）

完整的 FP16 GEMM + ReLU（基于官方 quickstart）：

import torch
import tilelang
import tilelang.language as T


@tilelang.jit
def matmul(A, B, block_M: int, block_N: int, block_K: int):
    M, N, K = T.const("M, N, K")
    dtype = T.float16
    accum_dtype = T.float32
    A: T.Tensor((M, K), dtype)
    B: T.Tensor((K, N), dtype)
    C = T.empty((M, N), dtype)

    # grid: (N 方向块数, M 方向块数)，每 block 128 线程
    with T.Kernel(T.ceildiv(N, block_N), T.ceildiv(M, block_M), threads=128) as (bx, by):
        A_shared = T.alloc_shared((block_M, block_K), dtype)          # shared tile
        B_shared = T.alloc_shared((block_K, block_N), dtype)
        C_local = T.alloc_fragment((block_M, block_N), accum_dtype)   # 寄存器累加器，fp32

        T.clear(C_local)

        # 三级软件流水：搬 k+1 块的同时算 k 块
        for ko in T.Pipelined(T.ceildiv(K, block_K), num_stages=3):
            T.copy(A[by * block_M, ko * block_K], A_shared)   # global -> shared
            T.copy(B[ko * block_K, bx * block_N], B_shared)
            T.gemm(A_shared, B_shared, C_local)               # tensor core，fp32 累加

        # epilogue：ReLU
        for i, j in T.Parallel(block_M, block_N):
            C_local[i, j] = T.max(C_local[i, j], 0)

        T.copy(C_local, C[by * block_M, bx * block_N])        # fragment -> global

    return C


M = N = K = 1024
# 用静态 shape 编译出可执行 kernel
matmul_kernel = matmul.compile(M=M, N=N, K=K, block_M=128, block_N=128, block_K=32)

a = torch.randn(M, K, device="cuda", dtype=torch.float16)
b = torch.randn(K, N, device="cuda", dtype=torch.float16)
c = matmul_kernel(a, b)

30 行写完一个带 fused epilogue 的 Tensor Core GEMM。几个关键点：

1. T.const("M, N, K") 声明符号化 shape，.compile(...) 按实际 shape 特化出可执行 kernel（也可以直接调用让 @tilelang.jit 在首次调用时惰性编译）。
2. 累加器 dtype 和存储 dtype 分离：C_local 是 fp32 fragment，写回时才降到 fp16。这是数值稳定的标准配方。
3. epilogue 用 T.Parallel 表达，不需要另起 kernel–省一次 HBM 往返。
4. 没有一行同步代码。__syncthreads() 由编译器根据 T.Pipelined 的依赖关系自动插入。

一句话总结：TileLang 的 API 面窄到只有 5 类原语，但这 5 类恰好覆盖了 GPU kernel 性能的全部决定因素–tile 怎么切、数据放哪级存储、什么时候预取。

---

三、循环原语：从 T.serial 到 T.Pipelined

GEMM 例子里已经出现过两种循环（T.Pipelined 和 T.Parallel）。TileLang 的循环构造一共就这四个，按暴露给编译器的并行度递进，先顺次过一遍：

原语	语义	典型用途
T.serial	普通 for 循环，迭代之间有依赖	递推、边界处理
T.unroll	要求编译器完全展开	小循环，省掉分支和循环开销
T.Parallel	嵌套并行循环，所有迭代互相独立	elementwise、epilogue
T.Pipelined	软件流水，生产者-消费者跨迭代重叠	GEMM / attention 主循环

for i in T.serial(N):                        # 串行：下一轮依赖上一轮的结果
    ...

for k in T.unroll(K_TILE):                   # 展开：编译期摊平
    acc += a[k] * b[k]

for i, j in T.Parallel(M, N):                # 并行：迭代独立，映射到线程
    C[i, j] = A[i, j] + B[i, j]

for ko in T.Pipelined(iters, num_stages=3):  # 流水：copy 与 compute 时间重叠
    ...

补充几点：

• T.serial 支持三参数形式 T.serial(0, N, 2)（起点、终点、步长）；
• T.Parallel 可以加 coalesced_width= 提示控制访存合并宽度，loop_layout= 挂 fragment layout 标注；
• 另有一个高级构造 T.Persistent，表达 persistent thread-block 风格的循环（5.1 提到的 stream-K 变体就靠它）；
• Python 原生的 if/else、while、break/continue 都可用，条件是 TIR 表达式即可；潜在的越界访问由 LegalizeSafeMemoryAccess pass 自动加 guard（见第七节）。

前三个原语都好理解，真正值得单独一节展开的是最后一个。

T.Pipelined 做了什么

上面例子里最「魔法」的一行是 for ko in T.Pipelined(...)。num_stages=3 不是「循环展开 3 次」，而是建立 3 级软件流水：

时间 ->
iter 0:  [copy k=0]
iter 1:  [copy k=1] [gemm k=0]    ← 进入稳态：搬运与计算重叠
iter 2:  [copy k=2] [gemm k=1]
...
iter N-1:           [gemm k=N-2]  ← epilogue：只剩计算
         └─ HBM 延迟被后续 iter 的计算隐藏

手写 CUDA 要实现同样效果，需要自己管理 stage 数组下标、cp.async 的 commit/wait group、以及每级之间的 barrier。TileLang 把这压缩成一个参数。具体地，编译器在这一个循环上做了四件事：

1. 循环重写：把源代码里的单层循环拆成 prologue（预取前 N-1 轮）/ 稳态 body（搬第 k+N-1 块的同时算第 k 块）/ epilogue（算完尾部） 三段。稳态时 copy 和 compute 在时间上重叠。
2. 共享内存多缓冲：num_stages=3 意味着 A_shared / B_shared 会被自动复制成 3 份（三缓冲），生产者写第 k+2 块、消费者读第 k 块，互不冲突。你在源码里写的是一份 buffer，编译器做的是 buffer 乘法。
3. 异步拷贝插入：流水线里的 T.copy 会 lower 成异步拷贝（Ampere+ 上是 cp.async，Hopper+ 上是 TMA 的 cp.async.bulk），并自动配好 commit_group / wait_group 的配对。
4. 同步插入：编译器的 PipelinePlanning / InjectSoftwarePipeline / InjectTmaBarrier 等一系列 pass 负责推导生产者-消费者依赖，在正确的地方插入 __syncthreads()（或 mbarrier）。这就是为什么源代码里一行同步都没有。

注意 T.copy 本身的语义是同步的–语句结束后 dst 就可读，如果 lower 到了异步指令，编译器会补上 wait 保证这一点。想手动控制异步，用 T.async_copy（不自动插 wait，需要自己写 T.ptx_wait_group）。

手动标注 stage / order

常规 GEMM 形态的流水线，num_stages=N 就够了，编译器自己推断谁是生产者谁是消费者。当循环体顺序不寻常（比如想让「下一轮的 copy」排在「这一轮的 compute」之前发射）时，可以显式标注：

for ko in T.Pipelined(
    num_tiles,
    stage=[0, 1],     # copy 是 stage 0，gemm 是 stage 1
    order=[1, 0],     # 发射顺序上 gemm 先、copy 后
):
    T.copy(A[ko * BK], A_shared)
    T.gemm(A_shared, B_shared, C_local)

规则：

• stage / order 与循环体内的可调度语句（copy、gemm、reduction、store、同步）按源码顺序一一对齐；
• 流水线深度由 max(stage) + 1 推断，此时不要再传 num_stages；
• 循环体里的标量别名（base = ko * BK 这类 Bind 语句）不占标注位–它们没有副作用，编译器会在每个消费者处按需重放；
• 编译器会校验依赖：生产者的 stage 必须不晚于消费者，同 stage 内 order 必须生产者在消费者前。

一句话总结：T.Pipelined(num_stages=N) = 循环三分重写 + shared memory N 重缓冲 + 异步拷贝 + 自动同步，这是 GEMM/attention 类 kernel 能打到带宽/算力天花板的全部前提。

---

四、同一份代码，不同架构的 lowering

TileLang 源码里只有 T.copy 和 T.gemm 这两个「意图」，落到哪条指令由 target 决定。这是它和直接写 CUDA 最大的分工差异–你描述数据流和 tile 结构，编译器按架构选指令：

架构	异步拷贝（流水线内的 T.copy）	MMA 指令（T.gemm）
SM70~75（V100/T4）	SIMT ld.global + st.shared	mma.sync（fp16）
SM80~89（A100/3090/4090/Ada）	cp.async 多级流水	mma.sync
SM90a（H100/H200）	TMA（cp.async.bulk.tensor）+ mbarrier	wgmma（warp group MMA）
SM100a（B100/B200）	TMA + mbarrier	tcgen05.mma（TMEM 累加）
SM120（RTX 50 / RTX PRO，消费级 Blackwell）	没有 TMA，走 cp.async（LDGSTS）	普通 mma（第五代 tensor core）

这里有个容易踩的坑：不能按 SM 版本号大小推测能力。SM120 数字上大于 SM90，但它没有 TMA、没有 tcgen05，异步拷贝走的是 Ampere 时代引入的 cp.async。在 SM120 卡上写完全相同的 T.Pipelined 代码，编译器会自动退回 cp.async 风格的流水–语义不变，只是底层搬运指令不同。这正是 T.Pipelined 这个抽象的价值：你表达的是「我要 N 级预取流水」这个意图，而不是「我要发 cp.async.bulk.tensor」这个指令。

target 通过三种方式指定：

kernel = tilelang.compile(func, target={"kind": "cuda", "arch": "sm_90"})

@tilelang.jit(target="cuda")          # 或裸字符串
def factory(...): ...

# 或环境变量（适合整台机器固定 GPU 型号的场景）
# export TILELANG_DEFAULT_TARGET='{"kind": "cuda", "arch": "sm_90"}'

auto（默认）按 CUDA → HIP → Metal 顺序探测。arch 直接对应 NVCC 的 -arch=sm_XX；需要一份代码出多个 SASS 时用 code 列表（fatbin）。跨厂商同理：HIP 配 mcpu="gfx90a"，Metal / LLVM CPU / WebGPU 各有对应 kind。

判断「我写的代码在这张卡上到底变成了什么」，最直接的办法还是后面第六节的 dump 源码–看生成的 CUDA 里是 cp.async、TMA descriptor 还是 wgmma，比读文档更可靠。

---

五、进阶主题

5.1 Split-K：grid 第三维 + atomic_add

M、N 小而 K 很大时，M×N 的 tile 数不够填满 SM，就把 K 切开分给多个 block 并行累加。TileLang 里不需要新原语–T.Kernel 的第三个 grid 维就是 split 因子，最后用 T.atomic_add 归约（来自 examples/gemm_splitk）：

splitK = K // split_k

with T.Kernel(T.ceildiv(N, block_N), T.ceildiv(M, block_M), split_k, threads=128) as (bx, by, bz):
    ...
    T.clear(C_local)
    for ko in T.Pipelined(T.ceildiv(splitK, block_K), num_stages=0):
        T.copy(A[by * block_M, bz * splitK + ko * block_K], A_shared)   # K 维偏移带上 bz
        T.copy(B[bz * splitK + ko * block_K, bx * block_N], B_shared)
        T.gemm(A_shared, B_shared, C_local)

    for i, j in T.Parallel(block_M, block_N):
        T.atomic_add(C[by * block_M + i, bx * block_N + j], C_local[i, j])

要点：bz 只出现在 K 维索引里；累加结束整体做一次 atomic_add（而不是每个元素多次原子写）；输出 C 必须先清零；num_stages=0 是官方示例关掉了自动流水（生产代码里该开的还是要开）。注意 atomic_add 的归约顺序不定，数值不可复现–对复现性有要求的场合要改成两阶段确定性归约（partial 先写回 workspace 再单独 reduce），本博客 DeepSeek-V4 mHC Pre-Block 融合 Kernel 详解 里有完整分析。同一目录下还有 stream-K 变体（gemm_streamk，把尾部 wave 按 K 拆给 peer block 再 fixup），是 persistent kernel 的入门样本。

5.2 warp specialization：T.ws + mbarrier

Hopper 之后，生产者和消费者可以拆成不同 warp group，各自跑各自的「循环」，靠 mbarrier 握手。TileLang 用 T.ws(i) 划分角色、T.alloc_barrier 建握手信号（来自 examples/warp_specialize）：

with T.Kernel(..., threads=256) as (bx, by):
    A_shared = T.alloc_shared((block_M, block_K), dtype)
    ...
    data_is_ready = T.alloc_barrier(arrive_count=128)
    compute_is_done = T.alloc_barrier(arrive_count=128)

    with T.ws(1):                       # 消费者 warp group
        T.clear(C_local)

    for ko in T.Pipelined(T.ceildiv(K, block_K), num_stages=0):   # 手动流水，不走自动重写
        with T.ws(0):                   # 生产者：等上一轮算完 → TMA 搬数 → 通知就绪
            T.barrier_wait(compute_is_done, (ko + 1) % 2)
            T.tma_copy(A[by * block_M, ko * block_K], A_shared, barrier=data_is_ready)
            T.tma_copy(B[ko * block_K, bx * block_N], B_shared, barrier=data_is_ready)
            T.barrier_arrive(data_is_ready)
        with T.ws(1):                   # 消费者：等数就绪 → gemm → 通知算完
            T.barrier_wait(data_is_ready, ko % 2)
            T.gemm(A_shared, B_shared, C_local)
            T.barrier_arrive(compute_is_done)

    with T.ws(1):
        T.copy(C_local, C[by * block_M, bx * block_N])

要点：

• num_stages=0 关掉自动流水–因为流水逻辑已经由两个 warp group 的 barrier 协议手动表达了，双缓冲体现在 barrier_wait 的 % 2 相位翻转上；
• T.tma_copy(..., barrier=...) 是显式 TMA 入口，比 T.copy 更低一层；
• 这是 TileLang 里「控制力接近 CUTLASS」的具体形态：mbarrier 的 arrive_count、相位奇偶全部由你负责。examples 目录里同族还有 barrierpipe / softpipe 等多种流水协议写法可以对照。

5.3 autotune：tile 参数交给搜索

block_M / block_N / block_K / num_stages / threads 这组参数的理论最优值依赖具体 GPU 和问题规模，手调不现实。TileLang 内置 autotuner，用法是把可调参数写成带默认值的函数参数，再套一层装饰器：

def matmul_configs(M, N, K):
    return [
        dict(block_M=BM, block_N=BN, block_K=BK, num_stages=S, threads=TH)
        for BM in [64, 128]
        for BN in [64, 128]
        for BK in [32, 64]
        for S in [2, 3]
        for TH in [128, 256]
    ]

@tilelang.autotune(configs=matmul_configs, warmup=25, rep=100, timeout=60)
@tilelang.jit(out_idx=[-1])
def matmul(M: int, N: int, K: int,
           block_M: int = 128, block_N: int = 128, block_K: int = 32,
           threads: int = 128, num_stages: int = 3, ...):
    ...

with set_autotune_inputs(a, b, c):      # 固定输入，保证各 config 可比
    tuned = matmul(M, N, K)             # 编译+验证+benchmark 全部 config，返回最优

值得知道的工程细节：候选 kernel 并行编译、逐个 benchmark；每个 config 会先过正确性检查（ref_prog 或默认的 torch 对比，容差 rtol/atol 默认 1e-2）；结果缓存在 ~/.tilelang/cache/autotuner，缓存 key 包含 TileLang 版本 + 函数源码 + config 列表，改代码自动失效。调稳定后建议把最优 config 烘焙成函数默认值写回源码，autotune 只作为开发期工具。

5.4 Blackwell 两条路径

Blackwell 不是一个统一架构，写跟架构相关的 kernel 时必须分开看：

	SM100a（B100/B200，数据中心）	SM120（RTX PRO / RTX 50，消费级）
异步拷贝	TMA + mbarrier	cp.async（LDGSTS），无 TMA
MMA 指令	tcgen05 + TMEM	普通 mma
two-SM（2-CTA）kernel	✅	❌
NVFP4 block-scaled	✅ T.mma_gemm_blockscaled	✅ T.mma_gemm_blockscaled（2026-07-30 加入）
对应 example	blockscaled_gemm_sm100、gemm_tcgen05	gemm_sm120

SM100a-- tcgen05 路径。第五代 Tensor Core 的 MMA 指令 tcgen05.mma 从 shared memory 直读操作数、累加到独立的 Tensor Memory（TMEM），还支持两个 CTA 配对发射。对应到 TileLang 是一组新原语：T.alloc_tmem 分配 TMEM 累加器，T.tcgen05_gemm 发射 MMA（不带隐式等待），T.alloc_barrier + T.mbarrier_wait_parity(mbar, k % 2) 手动做相位同步。这条路径目前是实验性 preview：同步协议要自己写，官方 README 明说「manual implementation required」。较新版本提供了半自动入口 T.gemm(..., mbar=...)–发射后自动插入匹配的 mbarrier_wait_parity，并把 fence 插入交给 InjectTcgen05Fence pass。examples/gemm_tcgen05/ 下有从裸 tcgen05_gemm 到 warp-specialized persistent、再到 2-CTA stream-K 的完整梯度。

SM120-- 传统路径。没有 tcgen05、TMEM 和 TMA，走的仍是 SM80 风格的 mma.sync + cp.async 流水线，TileLang 现有代码基本直接可用。换句话说，为 Hopper 写的 kernel 迁到 RTX 50 通常只是换个 arch，迁到 B200 才需要考虑 TMEM 那套新原语。两边唯一真正共享的新能力是 NVFP4 block-scaled MMA（T.mma_gemm_blockscaled，SM120 路径 2026-07-30 加入，可对照 Colfax 那篇 NVFP4 Blockscaled GEMM on RTX Pro Blackwell (sm12x)），但底层指令并不相同。

编译 target 上，数据中心 Blackwell 需要 fatbin 时可以 {"kind": "cuda", "arch": "sm_100f", "code": ["sm_100a", "sm_103a"]} 一份代码出多个 SASS。

---

六、写完之后的三个动作

quickstart 的官方流程把「kernel 写完之后」固化成了三步，建议形成肌肉记忆：

① 验证正确性–先于一切性能讨论，且以 fp32 参考为准：

ref32 = torch.relu(a.float() @ b.float())
torch.testing.assert_close(c.float(), ref32, rtol=1e-2, atol=0.05)

为什么要跟 fp32 参考比而不是直接 torch.relu(a @ b)：kernel 内部用 fp32 累加，a @ b 则是 fp16 累加，拿后者做参考的话误差来源两边不一致–你分不清到底是自己 kernel 错了还是累加精度差异。用 a.float() @ b.float() 做基准，差异才能归因到 kernel 本身。K=1024 累加下 atol=0.05 是合理范围。更复杂的 kernel（attention、MoE）可以保留一个 PyTorch 参考实现专门做这件事；TileLang 的 autotuner 也用同样思路验证每个候选 config。

② dump 生成源码–看编译器到底做了什么：

cuda_source = matmul_kernel.get_kernel_source()
print(cuda_source)

返回的是最终生成的完整 CUDA 源码。这一步回答所有「lowering 疑问」：T.copy 变成了 cp.async 还是 TMA descriptor？T.gemm 是 mma.sync 还是 wgmma？多缓冲分配了多少 shared memory？__syncthreads() 插在了哪？调性能之前先读一遍生成代码，能省掉大量盲猜。换个 num_stages 或 block_K 再 dump 一次、diff 两份源码，比读任何文档都直观。

③ benchmark–拿可信的延迟数字：

profiler = matmul_kernel.get_profiler(tensor_supply_type=tilelang.TensorSupplyType.Normal)
latency = profiler.do_bench()      # ms

get_profiler 会自动生成输入、warmup、多次重复取统计，不用自己写 torch.cuda.synchronize() + time.time() 那套容易测错的东西。注意 TensorSupplyType.Normal 指定用正态分布造输入–对 GEMM 无影响，但对带 exp 的 attention kernel 会影响数值路径，别用全 0。没有 JITKernel 对象时也可以直接 from tilelang.profiler import do_bench 包一个 callable。有了 ① 的正确性和 ③ 的基线数字，后面任何改动（换 tile 尺寸、加 swizzle、上 warp specialization）都是可度量的。

---

七、学习路径与调试工具

环境与选卡

pip install tilelang
python -c "import tilelang; print(tilelang.__version__)"

需要 Python ≥ 3.10。想要最新特性走 nightly：pip install tilelang --find-links https://tile-ai.github.io/whl/nightly。

没有 GPU 怎么办：TileLang 的 CUDA 后端需要真卡才能编译执行，Apple Silicon 可以走 metal target 但例子覆盖有限。学 TileLang 属于典型的短时高强度用卡–跑几小时 examples 就停，不需要长期占资源，RunPod 按小时租一张卡是最省事的路子。

选卡要看想学什么（原因见第四节）：

学习目标	需要的卡
基础 tile / T.Pipelined / autotune	任意 SM80+，一张 4090 就够
TMA、T.tma_copy、warp specialization	必须 SM90a（H100/H200）
tcgen05 MMA、TMEM、two-SM kernel	必须 SM100a（B100/B200）
SM120 NVFP4 block-scaled	RTX PRO 6000 / RTX 50 系

别拿 SM120 的卡去学 TMA–消费级 Blackwell 没有这个硬件单元。

调试工具箱

TileLang 在调试工具链上比 Triton 强，别浪费：

• T.print(buffer, msg=...)：kernel 内部打印 shared/fragment buffer，TileLang 自动只从单个线程打印避免刷屏；配合 if i == 0: 谓词用。
• T.device_assert(cond, msg)：device 侧断言，CUDA target 上生效，排查越界和 NaN 比注释掉半段代码快得多。
• Pass Visualizer（2026-07 加入）：结构树浏览器，看每个编译 pass 对 IR 做了什么。
• IR Lower Trace（2026-07 加入）：逐 pass dump IR，定位「我写的 layout 到哪一步被改掉了」。
• TileLang LSP（2026-08 开源）：VSCode 里显示 buffer 的 shape / dtype / scope / 推断 layout 的 inlay hint，写 kernel 时不用反复回头查 shape，第一天就该装上。
• layout 可视化：把 fragment layout 画出来，检查 bank conflict。
• get_kernel_source()：见第六节，最强的「调试器」其实是读生成的 CUDA。
• 缓存目录 ~/.tilelang/cache/：autotuner 产物、编译出的 .so / cubin 都在这；改了代码行为不对时先怀疑缓存，TILELANG_DISABLE_CACHE=1 一键排除。
• 越界防护：LegalizeSafeMemoryAccess pass 会在可能越界的访问处自动插 guard（证明安全则自动消除），所以边界处理很多情况下不用手写 if–但自定义边界逻辑仍建议显式写。

推荐顺序

阶段	材料	目标
①	TileLang Puzzles 10 题	建立 tile 思维，比读文档有效得多
②	quickstart.py + elementwise	摸清 5 个原语
③	语言基础文档	补齐 layout / 内存作用域概念
④	examples/gemm	swizzle、autotune、架构特化
⑤	examples/flash_attention	online softmax + 双 GEMM 融合
⑥	examples/gemm_fp8 / blockscaled_gemm_sm100	量化 GEMM，per-block scale
⑦	examples/gemm_splitk / warp_specialize	本文第五节的出处
⑧	examples/deepseek_mla / deepseek_v4 / deepseek_mhc	真实生产算子

Puzzles 优先这点要强调：TileLang 的文档偏参考手册风格，直接读容易只学到 API 名字、学不到「为什么这么切」。10 个难度递增的 puzzle 会强迫你自己想清楚 tile 划分。第 ⑧ 步的收益也不在语法，而在于看懂「一个生产级 LLM 算子是如何被拆成 tile 数据流的」。

一句话总结：语法半天就能过完，TileLang 的学习成本在「tile 级思维」–而这只能靠 Puzzles 和 examples 里那些真实 kernel 攒出来。

---

参考：

• TileLang GitHub（v0.1.13，2026-08）
• 官方文档 · 语言基础 · 控制流 · 软件流水 · autotuning · target 指南
• quickstart.py：本文 GEMM + ReLU 例子与「三个动作」的出处
• gemm_splitk · warp_specialize · gemm_tcgen05：第五节各示例的出处
• TileLang Puzzles：10 个交互式练习
• TileLang LSP：编辑器支持
• tilelang-benchmark：官方 benchmark 脚本
• Colfax Research: NVFP4 Blockscaled GEMM on RTX Pro Blackwell (sm12x)：SM120 block-scaled MMA 路径
• TVM：底层编译基础设施