从零到一：TileLang 系统教程

从「为什么 PyTorch 里一个算子会慢」讲起：存储层级的不对称、算术强度与机器平衡点，然后沿着 SIMT → 张量核 → 异步搬运三代硬件演进，说明为什么 CUDA、Triton、TileLang 会依次出现。

• 学习目标
• 前置知识
• 1. 问题的由来：为什么 PyTorch 里一个算子会慢
• 2. 概念与模型

把「一次训练迭代」和「一次推理请求」拆成算子的清单与账本：每个算子要搬多少字节、算多少次、是 memory-bound 还是 compute-bound，以及 torch.compile 的边界在哪、什么时候必须自己写 kernel。

• 学习目标
• 前置知识
• 1. 问题的由来：算子在一次迭代里出现在哪些位置
• 2. 概念与模型：一个算子要读多少字节、算多少次

TileLang 的核心主张：把调度空间（线程绑定、layout、tensorize、流水）从数据流里解耦出来。讲清 tile 的定义、beginner/developer/expert 三层接口，以及可组合（composable）到底是什么意思。

• 学习目标
• 前置知识
• 1. 问题的由来：调度为什么会变成负担
• 2. 概念与模型

给出一张可随时重画的仓库结构图与一次编译的完整路径：前端 → TIR → pass pipeline → device codegen → host wrapper，并对比 CUDA 与 Ascend 两条管线。

• 学习目标
• 前置知识
• 1. 问题的由来
• 2. 概念与模型

用「用户必须提供的硬件相关决策数量」这把尺子，把 CUDA、Triton、CUTLASS-CuTe、TVM、cuTile、Ascend C、PTO-DSL 放到同一条纵轴上，再用同一个带融合 epilogue 的 GEMM 逐层对照。

• 学习目标
• 前置知识
• 1. 问题的由来：为什么需要一张坐标系
• 2. 概念与模型：三层抽象，七个对象

用存储层级、执行单元、数据搬运、同步四个维度描述任意加速器，再分别落到 NVIDIA（grid/warp/Tensor Core/TMA/cluster）与昇腾（Cube/Vector/MTE/UB/L1/L0）上，并解释为什么「GPU 是 SIMT、NPU 是 SIMD」已经过时。

• 学习目标
• 前置知识
• 1. 问题的由来：为什么必须有「片上存储」和「分块搬运」
• 2. 概念与模型：描述任何加速器的四个维度

装好 TileLang 并写出第一个 kernel：安装的三条路径、kernel 的四个部件、如何把生成的 CUDA 源码逐行读回来，以及 JIT、编译缓存与 kernel 复用的行为。

• 学习目标
• 前置知识
• 1. 问题的由来
• 2. 概念与模型：一个 kernel 的四个部件

把第 7 章的生成源码往回拆：内核定义与网格、dtype 与内存作用域、T.Parallel/循环/控制流各自对应 IR 里什么、代价是什么，以及 Python 兼容性的边界在哪。

• 学习目标
• 前置知识
• 1. 问题的由来：kernel 里到底要"表达"什么
• 2. 概念与模型：六个构件

逐个讲透核心原语：T.copy 的搬运与越界语义、T.gemm 如何选指令、reduce/atomic/fill 的用法与坑，以及昇腾专属的 T.dual_copy、T.SimtVF/T.SimdVF 为什么必须存在。

• 学习目标
• 前置知识
• 1. 问题的由来：原语是「接口」，不是「实现」
• 2. 概念与模型

解释「同一个张量为什么会有摆法问题」：从 fragment 与线程的对应、bank conflict、swizzle 到昇腾的 fractal 与 ND→NZ，并给出把布局画出来观察的可复现步骤。

• 学习目标
• 前置知识
• 1. 问题的由来：同一个张量，为什么会有「摆法」问题
• 2. 概念与模型

把搬运与计算重叠起来：T.Pipelined 的 stage/order 语义、多缓冲的容量代价、如何判断一份配置还放不放得下，以及 autotuner 的搜索空间怎么定义、结果怎么读。

• 学习目标
• 前置知识
• 1. 问题的由来：不做流水会怎样
• 2. 概念与模型

把 kernel 问题分成生成、正确性、性能、偶发四类，给出对应的工具链：编译期日志与 IR Lower Trace、运行期打印与异常 dump、对拍容差的来源，以及 Nsight Compute / msProf 的读法。

• 学习目标
• 前置知识
• 1. 问题的由来
• 2. 概念与模型

从最不起眼的逐元素算子讲起：用访存账本说明融合的价值，再逐步实现 scale+add、LayerNorm/RMSNorm 与 RoPE，最后合成 RMSNorm+RoPE+量化的融合算子。

• 学习目标
• 前置知识
• 1. 问题的由来：最"简单"的算子也值得认真写
• 2. 概念与模型

GEMM 的性能账本与五级实现阶梯：从朴素分块到 shared memory、软件流水、swizzle 注解与官方 intrinsic 版本，再到 split-K/stream-K/grouped 与低精度变体的选择依据。

• 学习目标
• 前置知识
• 1. 问题的由来：GEMM 的性能账本
• 2. 五级阶梯：从朴素版到 intrinsic

从朴素 attention 的三趟访存推导到 online softmax 的一趟实现，逐级加上分块、流水与片上复用，并覆盖 causal、varlen、GQA、decode、paged KV 与 split-KV 等变体。

• 学习目标
• 前置知识
• 1. 问题的由来：attention 的时间花在哪
• 2. 概念与模型

大模型里最贵的两段结构：MoE 的路由、重排与 grouped GEMM，以及 MLA 的 KV 压缩如何改变访存与 kernel 形态；最后给出一次 decode step 的耗时占比与瓶颈类型总结。

• 学习目标
• 前置知识
• 1. 问题的由来：一个 decode step 里最贵的是什么
• 2. MoE：四段结构，两段不是矩阵乘

用 HuggingFace 的数据、tokenizer 与权重，把 TileLang 算子接进真实的训练循环：显存与算力两本账、数据管线、混合精度与 checkpoint，以及哪些热点算子值得换成自己的 kernel。

• 学习目标
• 前置知识
• 1. 问题的由来：「从零训练一个 7B 模型」这句话哪里是真的
• 2. 可行性核算：显存与算力两本账

训练完的模型怎么变成服务：prefill 与 decode 的差异、KV cache 的显存账本、paged KV 与 continuous batching，以及 decode attention、量化 GEMM 如何接进 vLLM / TileRT 这类框架。

• 学习目标
• 前置知识
• 1. 问题的由来：训练完的模型怎么变成服务
• 2. 概念与模型

在昇腾上写第一个算子：环境与 CANN 安装、GM→UB/L1/L0 的数据流、Ascend C 对照实现、msProf 调优、ACLGraph 入图与设备端调试手段。

• 学习目标
• 前置知识
• 本章的验证边界（务必先读）
• 1. 问题的由来：在昇腾上写算子，「难」在哪

把 TileLang 的编译期读成一条可验证的管线：TIR 的基本构件、CUDA 与 Ascend 两条 pass 序列的分组与职责，以及 LayoutInference、LowerTileOp、ThreadSync 这些承重 pass 各自改变了什么。

• 学习目标
• 前置知识
• 1. 问题的由来
• 2. 概念与模型

layout 推导的形式化与实现：约束从哪里来、如何在 tile op 图上传播、冲突时如何权衡转换代价，以及昇腾的 fractal 归一化、ND→NZ 与边界填充为什么是另一类问题。

• 学习目标
• 前置知识
• 1. 问题的由来：为什么 layout 必须被「推导」
• 2. 形式化：layout 代数

软件流水的算法层：依赖距离与模调度、缓冲版本数的推导、多缓冲的物化与复用，以及从读写集推导同步点（CUDA 的 ThreadSync 与昇腾的 InsertSync）的原理。

• 学习目标
• 前置知识
• 1. 问题的由来：把「重叠」变成可判定的事
• 2. 软件流水的形式化

昇腾后端最有技术含量的一步：把 kernel 拆成带资源标签的任务、估计延迟与容量、用 Z3 求解调度与多缓冲，再把解落成真实的 Set/Wait。讲清模型、约束与它失效时的后果。

• 学习目标
• 前置知识
• 1. 问题的由来：Ascend C 里最难的不是写，是排
• 2. 建模：把 kernel 变成一个调度问题

NVIDIA 后端如何把 tile 程序翻译成 CUDA 与 PTX：指令选择的谓词链、TMA 判定、warp specialization、cluster/DSM，以及 Blackwell 的 tcgen05 与 TMEM 带来的新约束。

• 学习目标
• 前置知识
• 1. 问题的由来
• 2. 概念与模型

昇腾后端的四条路线与它们各自的落点：主干 ascend 方言、适配器的 Ascend C&PTO 与 AscendNPU IR，以及 T.alloc_l1/l0a/l0b/l0c、dual_copy 在代码生成里变成了什么。

• 学习目标
• 前置知识
• 本章的验证边界（务必先读）
• 1. 问题的由来：一次 `target="ascend"` 之后到底发生了什么

从装饰器到可调用 kernel 的完整链路：JIT 生命周期、缓存键与失效、packed ABI 的参数校验、五个 adapter 的分工，以及张量转换、stream 语义与 autograd 的边界。

• 学习目标
• 前置知识
• 1. 问题的由来
• 2. 概念与模型

低精度的数值基础与工程约束：fp16/bf16/fp8/MXFP4/int4 的表示与误差、block scale 的打包与存储布局，以及如何选参考实现对拍、如何区分数值误差与逻辑错误。

• 学习目标
• 前置知识
• 1. 问题的由来
• 2. 数值基础

把前面的经验上升成方法论：先算（性能模型、roofline）、再定 tile、再定流水，最后才写代码；配一个完整的融合算子案例与一份反模式清单。

• 学习目标
• 前置知识
• 1. 问题的由来：为什么必须「先算再写」
• 2. 第一步：把一个算子变成 `(flops, bytes)`

把「读编译器」变成「改编译器」：C++ 与 Python pass 的双语结构、写一个新 pass 的完整流程、IR diff 与性能回归怎么测，以及评审时会被问什么。

• 学习目标
• 前置知识
• 1. 问题的由来
• 2. 概念与模型

区分「已发布 / 只在 main / 只在未合入 PR」三种状态，逐条核对 TVM 迁移、CuteDSL 后端、Ascend 950 与 PTO、持久化 kernel、AutoSimtVF 等方向的现状，并给出持续跟踪上游的方法。

• 学习目标
• 前置知识
• 1. 问题的由来：release 之外还有两个世界
• 2. 概念与模型：三种状态、四步判定

把改动变成上游的一部分：写算子、改编译器、加后端三条路径各自的门槛与测试要求，issue→RFC→PR→CI 的流程，以及昇腾双仓库的协作方式与自检清单。

• 学习目标
• 前置知识
• 1. 问题的由来
• 2. 概念与模型：三种路径、四道门

2026-09-30 DeepSeek 面向昇腾开源了六个组件。本章给出全景地图：每个组件解决什么问题、与 NVIDIA 侧怎么一一对应、在一次训练 step 和一次推理请求里各占哪一段，以及怎么自己核验版本、硬件与可用状态。

• 学习目标
• 前置知识
• 本章的验证边界（务必先读）
• 1. 问题的由来：为什么需要先读一张地图

TileKernels 是用 TileLang 写的生产算子库，同一份算法提供 CUDA 与 Ascend 两个后端。以 top-k gate 为例逐行对照两种写法，讲清哪些部分平台无关、哪些必须重写，以及量化、RoPE、Rand 这些算子族的接口。

• 学习目标
• 前置知识
• 本章的验证边界（务必先读）
• 1. 问题的由来：为什么是 TileKernels

另一类算子：不是矩阵乘，而是访存与小算子为主的结构性算子。以 Engram 论文与 mHC 论文（arXiv:2512.24880）为定义、以 TileKernels 的 mhc/ 与 engram/ 源码为实现，讲透门控融合、Sinkhorn 归一化，以及「除 prenorm GEMM 外都用 TileLang」这条边界。

• 学习目标
• 前置知识
• 本章的验证边界（务必先读）
• 1. 问题的由来：为什么需要「不是矩阵乘」的算子

当 TileLang 不够用（或不该用）时，专家怎么在昇腾上把 GEMM 做到 99% 硬件利用率。讲清 MAD 原语抽象包了什么、JIT 与启发式怎么选配置、MegaMoE 为什么能把通信和两个 grouped GEMM 融进一个 kernel，以及哪一类算子反而适合交给 TileLang。

• 学习目标
• 前置知识
• 本章的验证边界（务必先读）
• 1. 问题的由来：TileLang 什么时候不够用

官方为昇腾稀疏注意力写了一篇技术报告。本章把它讲成一条可跟着走的设计过程：从算存比推出「64 head/AI Core、1C2V、accumulator 放 UB」，再讲 KV 的 ND2NZ 时机、双流水排布、gather2 与 skip-scale，并划清哪些结论能用 TileLang 表达、哪些必须下钻。

• 学习目标
• 前置知识
• 本章的验证边界（务必先读）
• 1. 问题的由来：稀疏注意力为什么要在昇腾上重写一遍

算法先行，硬件随后。用阈值扫描 + 随机块序 + RadixSelect 把 TopK 的额外计算量压到 O(k log(N/k))，并给出期望复杂度的证明；再看 CUDA 侧怎么用 mask 与 PTX 把常数压到极致。它同时回答了「什么类型的算子该手写而不是用 TileLang」。

• 学习目标
• 前置知识
• 本章的验证边界（务必先读）
• 1. 问题的由来：TopK 为什么在 DeepSeek 里是个瓶颈

MoE 的专家并行会引入 all-to-all，而通信很容易把算力省下的东西吃回去。本章讲清 EP dispatch/combine 的语义、FP8 dispatch 与延迟 epilogue、HCCL/UBMEM/URMA 的分工、DeepJIT 为什么要在运行时编译，以及 EP8–EP128 的带宽数字怎么读。

• 学习目标
• 前置知识
• 本章的验证边界（务必先读）
• 1. 问题的由来：省下来的算力，被 all-to-all 吃回去

收口章：把六个组件拼回一条完整链路，分别走一遍训练的一步与推理的 prefill/decode，给出算力、片内带宽、卡间带宽三类资源的预算账与瓶颈定位清单，并把前面所有「未实现 / 实验性」的结论集中成一份诚实的现状清单。

• 学习目标
• 前置知识
• 本章的验证边界（务必先读）
• 1. 收口：这一章要回答两个问题