第 02 章背景、设计思想与架构25 分钟13 习题8 自测题

从 PyTorch 到 kernel：算子在栈里的位置

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

本章进度状态未读完习题 0 / 13自测未作答

学习目标

读完本章，你应当能够：

画出一次训练迭代（前向 / 反向 / 优化器 / 通信）与一次推理请求（prefill / decode / KV cache）的算子地图，说清每个位置上算子扮演什么角色。
对任意一个算子做「读多少字节、算多少次」的推导，算出算术强度，并与机器平衡点比较，判断它是 memory-bound 还是 compute-bound。
能手算 LayerNorm、GEMM、attention（prefill 的两种写法、decode）这四类算子的账，并解释为什么融合能把 attention 的中间量访存从 $O (S^{2})$ 降回 $O (S)$ 。
说清一行 PyTorch 在 eager 模式下到底发生了什么（dispatch → 逐算子 kernel），以及 torch.compile / TorchInductor 在哪些层面替你做了决定、在哪些层面做不了。
用一页判据清单判断什么时候需要自己写 kernel，并知道「语法正确、数值正确、但比库慢几百倍」这种风险的来源。
说清本章每一个数字的出处与口径：哪些是厂商规格、哪些是第三方实测、哪些是本书的公式推导。

前置知识

第 1 章：算术强度、机器平衡点、存储层级、tile 的定义。本章会直接用到「平衡点」这个词，如果忘了它可以先回去看一眼 1.3 节。
会用 PyTorch：写过 torch.nn.Module、loss.backward()、optimizer.step()，读过 F.scaled_dot_product_attention 或 HuggingFace 的建模代码。
知道矩阵乘法与 attention 的公式即可。不需要 CUDA、不需要性能分析经验，本章所有推理都是从「一个算子要读多少字节、算多少次」这种方法出发的。

本章回答一个问题：当你写下 y = model(x) 这一行时，从你的 Python 代码到硬件上真正运行的 kernel，中间隔了几层；每一层的代价是什么。 它是第 3 章（TileLang 的设计哲学）、第 4 章（编译全流程）和后面所有算子章节的共同前提。

本章的代码全部是纯 Python 的账本脚本，不需要 GPU 也能跑（code/）。写作环境没有 NVIDIA GPU，也没有昇腾 NPU，因此本章没有任何一条性能数字来自本机实测；每个数字都标了来源与测量条件。

1. 问题的由来：算子在一次迭代里出现在哪些位置

一次 loss.backward(); optimizer.step() 里算子出现的四个位置，以及每个阶段在 roofline 上的典型判据。

前向的完整算子：norm（RMSNorm / LayerNorm）→ 投影 GEMM（QKV / O / gate / up / down）→ attention（QKᵀ → softmax → PV）→ 逐元素与融合（RoPE、残差加、SiLU、mask）→ loss（cross entropy / logits）。
反向：dY = dOut · Wᵀ（每个前向 GEMM 的 2 倍 FLOPs）、dW = Xᵀ · dOut、norm / 激活 / mask 的反向、attention 反向（需要重算或保存 softmax 中间量）。
优化器是完整的算子而不是一行 Python：每参数读 grad / param / m / v、写回 param / m / v，合计 ±26 字节；通信没有算术，ring all-reduce 搬的是梯度的 2(N−1)/N 倍。
各阶段的判据：前向与反向由算力决定（长序列下 FLOPs 远大于字节数），优化器由带宽决定，通信由网络带宽决定。
占比由 code/02-step-budget/step_budget.py 的粗模型复算（8B 级模型、seq_len=4096、batch=1、8 卡、H200 标定参数），属公式推导，不是实测。

同一个 attention 算子，在 prefill 与 decode 两段的判据完全相反：一个要喂满张量核，一个要少读字节。

prefill 的细节：embedding / 位置编码 → 每一层（norm → QKV GEMM → attention → O GEMM → MLP，其中 MLP 有 3 个 GEMM）→ 采样前处理（norm + lm_head GEMM，只对最后一个 token 也需要）。
decode 段是一个循环：采样出下一个 token 后回到 decode 的第一步，直到结束。
prefill（整段 prompt）：66.1 TFLOP / 28.5 GB → compute-bound；decode（每个 token）：16.1 GFLOP / 14.5 GB → memory-bound。
decode 每产出一个 token 就要把整套权重读一遍——这就是 batch、量化、MQA/GQA、投机解码要解决的问题。
数字由 code/02-step-budget/step_budget.py 复算，是本书口径的下界推导，不是实测。

1.1 训练一步：四个阶段，四种瓶颈

把一次训练迭代拆开，算子出现在四个位置（图 1）：

一次训练迭代里算子出现的四个位置：前向、反向、优化器与通信，四者的瓶颈各不相同。

前向与反向的下界由算力决定，优化器由带宽决定，通信由网络带宽决定。
含每层内部结构的完整版本见 figures/01-training-step-operator-map.md。

图 1 的完整版本（含每层内部结构）见 图：01-training-step-operator-map。

这四个阶段的特征完全不同：

（a）前向：主体的算子是三类——归一化（RMSNorm / LayerNorm）、矩阵乘（QKV 投影、输出投影、MLP 的三个矩阵）、attention（ $Q K^{⊤}$ 、softmax、 $P V$ ）。此外还有一批容易被忽略的逐元素算子：RoPE、残差相加、SiLU/GELU、mask、类型转换。在长序列、小 batch 的训练里，前向的瓶颈几乎总是算力：FLOPs 远大于必须搬动的字节数。

（b）反向：与前向逐层镜像。这里有一个值得记住的精确关系：对一个线性层 $C = A B$ （ $A \in R^{M \times K}$ ， $B \in R^{K \times N}$ ），前向是 $2 M N K$ FLOP；反向要算两个梯度 $d A = d C B^{⊤}$ 和 $d B = A^{⊤} d C$ ，各 $2 M N K$ FLOP。所以

反向的 FLOPs = 2 × 前向的 FLOPs（对 GEMM 而言，精确成立）

对 GEMM 主导的模型，反向的总量因此约为前向的两倍；attention 的反向还要额外处理 softmax 的中间量（要么保存，要么重算），通常比这更贵。本章的账本脚本（code/02-step-budget/step_budget.py）就用这个 2 倍关系估算反向。

（c）优化器：AdamW 不是「一行 Python 代码」，它是一个完整的算子。昇腾 CANN 的官方算子文档把它的输入输出讲得很清楚：varRef（参数 $θ$ ）、 mRef、vRef 三个张量既是输入也是输出，另外读入 grad 与学习率、权重衰减、 $β_{1}^{t}$ 、 $β_{2}^{t}$ 等标量，公式为 $m_{t} = β_{1} m_{t - 1} + (1 - β_{1}) g_{t}$ 、 $v_{t} = β_{2} v_{t - 1} + (1 - β_{2}) g_{t}^{2}$ 、 $θ_{t + 1} = θ_{t} - \frac{η}{v ^ _{t} + ϵ} m^_{t} - η λ θ_{t - 1}$ （来源：hiascend.com · canncommercial_900_API_aolapi_context_ops-nn_aclnnApplyAdamW_md 第 25–60 行）。

数一下它的访存：每个参数要读 grad、param、m、v，再写回 param、m、v。如果状态是 fp32、梯度是 bf16，就是 $2 + 3 \times 4 + 3 \times 4 = 26$ 字节/参数，而算术只有十几个 FLOP。算术强度约 0.6 FLOP/Byte——比任何 GEMM 都低两个数量级。一个 70 亿参数的模型，优化器一步要搬 182 GB 数据（本书推导，见 code/01-op-traffic-model/op_traffic.py 的 adamw 一行）。这是训练里最纯粹的带宽阶段，也是 ZeRO、分片优化器状态这些技术的动机。

（d）通信：数据并行时每个 rank 都算出完整梯度，然后用 all-reduce 汇总。 ring all-reduce 搬运的数据量是梯度本身的 $2 (N - 1) / N$ 倍（ $N$ 为卡数）， 一个 FLOP 都不做。对 8B 模型、bf16 梯度、8 卡，这一项就是每步约 28 GB 的搬运量。

1.2 一次推理请求：prefill 与 decode 是两件事

推理的算子地图不一样（图 2，完整版见 图：02-inference-request-operator-map）：

prefill 一次把整个 prompt 算完（compute-bound），decode 每步只出一个 token 却要读全部权重与 KV cache（memory-bound）。

prefill（整段 prompt）：66.1 TFLOP / 28.5 GB → compute-bound；decode（每个 token）：16.1 GFLOP / 14.5 GB → memory-bound。
decode 段是一个循环：每产出一个 token 后回到下一步，直到结束。

prefill 的形状与训练前向很像： $S$ 个 token 一起过网络，矩阵乘的 $S$ 维足够大，是算力密集的。decode 完全不同：每一步只有 1 个 query token，但要与整段 KV cache 做 attention，还要把整个模型的权重读一遍来决定这一个 token。用本章的口径算（8B 级模型、32 层、KV 长度 4096、bf16）：

prefill：66.1 TFLOP，28.5 GB    → compute-bound
decode ：16.1 GFLOP，14.5 GB    → memory-bound（每产出一个 token 读 14.5 GB）

同一个 attention 算子，在两段的判据是相反的。这一条会在第 15 章（attention 家族）和第 18 章（推理与部署）反复出现。

KV cache 本身也要算子来维护。CANN 里有一个算子 aclnnKvRmsNormRopeCache，功能是「对输入张量（kv）的尾轴，拆分出左半边用于 rms_norm 计算，右半边用于 RoPE 计算，再将计算结果分别 scatter 到两块 cache 中」（来源： hiascend.com · canncommercial_900_API_aolapi_context_ops-transformer_aclnnKvRmsNormRopeCache_md 第 25 行）。 TileLang 的 paged MLA decode kernel 里则能看到对应的接口形态：block_table 描述这个 sequence 的 KV 落在哪些块上，cache_seqlens 描述长度（_vendor/tilelang-v0.1.15/examples/deepseek_mla/example_mla_decode_paged.py:30-38）。

1.3 所有算子最终都变成 kernel

把视角拉远一点。一个 PyTorch 模型从你的代码到硬件，中间是这样几层：

从 Python 模型到硬件：中间隔着算子图与一个个 kernel，TileLang 负责的是最下面那一层。

第一跳（可选）由 torch.compile / TorchDynamo 捕获计算图；第二跳由 TorchInductor 做图级融合、调度与 kernel 生成。
TileLang 的位置在这一层：它不管「整张图怎么切」，只负责「一个算子怎么写快」。

TileLang 的位置在图的中间偏下：它不管「整张图怎么切」，只负责「一个算子怎么写快」。本书的合成文档把这句话写得很直接：

「TileLang 不是 PyTorch 的替代品，也不是完整的图编译器。它解决的是「一个算子内部怎么写快」，位置在 Inductor 之下、CUDA/Ascend C 之上。」 —— TileLang 知识体系与学习路径（编者笔记）第 44–48 行

在 TileLang 的代码里，这个位置对应几个具体入口：装饰器 @tilelang.jit （_vendor/tilelang-v0.1.15/tilelang/jit/__init__.py:571-600）、把 TIR 变成可调用对象的 tilelang.compile（同文件 :94-140）、包装好的 JITKernel （_vendor/tilelang-v0.1.15/tilelang/jit/kernel.py:41，它的 __call__ 在 :206， get_kernel_source() 在 :532，get_profiler() 在 :516）。compile 的签名里有一项 execution_backend，取值包括 "cython"、"tvm_ffi"、"nvrtc"、"torch"、"cutedsl" ——这就是「怎么和在 PyTorch 里跑的模型接起来」的那个接缝，第 26 章会展开。

这一层的存在，是本章后面所有讨论的前提：既然算子最终都要变成 kernel，那么「谁来写这个 kernel、它读多少字节、算多少次」就是一个可以被计算、被比较的问题。

2. 概念与模型：一个算子要读多少字节、算多少次

横轴是算术强度（FLOP/Byte，对数），脊点在 AI = 162：左边是 memory-bound 区，右边是 compute-bound 区。

峰值算力天花板 729 TFLOPS（fp16 张量核，标定后）；带宽天花板 y = AI × 4502 GB/s；脊点 AI = 162（H200 标定参数）。
同一个 attention 同时落在脊的两侧：prefill 落 S×S 掉进 memory-bound（AI≈42），flash 风格回到 compute-bound（AI≈2068）；decode 因为每读一个 KV 元素只做一次乘加，永远在 memory-bound 区（AI≈1）。
条形按算术强度线性绘制，跨度四个数量级，所以最左边的几条几乎看不见；精确值以右侧数字为准。
所有数字来自 code/01-op-traffic-model/op_traffic.py 的输出，属公式推导，不是实测。

2.1 记账法：四条规则

性能建模最容易被误解的一点是：「访存量」不是「你读了多少数据」，而是「算法最少必须搬多少数据」。本书采用的口径与 TileLang 生态里 TileOPs 仓库的 roofline 模型一致（TileOPs/docs/design/roofline.md 第 16–66 行）：

每个输入存储读一次，每个输出存储写一次；被读写的中间张量不单独计。
一个基本运算（add / sub / mul / div / max）记 1 FLOP；一个超越函数（exp / log / tanh / sqrt / rsqrt / cos / sin…）在账本上也记 1 FLOP。
矩阵乘的一次乘加记 2 FLOP。
一次「算子调用」的代价就是 $(f l o p s, b y t es)$ 这两个整数；由此得到

算术强度 AI = flops / bytes                   （FLOP/Byte）
访存时间   = bytes / 带宽
计算时间   = flops / 峰值算力
时间下界   = max(访存时间, 计算时间)

这套口径给的是下界：真实实现只会更慢，因为还有 cache 未命中、kernel 启动、尾部效应、同步开销。下界的用处不是预测，而是判断优化方向和验收（离下界还有多远）。

2.2 机器平衡点：把「算力」和「带宽」放在同一把尺子上

定义机器平衡点：

平衡点 = 峰值算力 / 峰值带宽     （FLOP/Byte）

于是：AI < 平衡点 时机器先被带宽卡住（memory-bound），AI > 平衡点 时先被算力卡住（compute-bound）。

用两个公开的机器参数各算一次（都是本书的推导，公式就是上面的比值）：

机器	参数	出处与条件	平衡点
H100 SXM5	1000 TFLOPS fp16 dense / over 3 TB/s，HBM3	NVIDIA 官方博客，厂商公开的 preliminary specs，未实测（`NVIDIA Hopper Architecture In-Depth 第 141 行`、`:334`）	≈ 333 FLOP/Byte
H200	989.5 TFLOPS fp16 dense / 4800 GB/s 规格值；仓库自带标定后折算成 729 TFLOPS / 4502 GB/s	`TileOPs/src/tileops/perf/profiles/h200.yaml 第 15–38 行`（标定条件写在该文件注释里：2048 MB working set、5 次 ×200 轮、显存 3201 MHz、算力受 700 W 限制）	162 FLOP/Byte（标定后）

请注意第二行里藏着一个重要的习惯：「厂商规格」和「能跑到的」不是一个数。 H200 的标定文件里，fp16 张量核的 calibration 是 0.7370（burst），持续值只有 0.6684 （受 700 W 限制）；显存带宽的 calibration 是 0.9380，而纯 copy 混合只有 0.8183。本章后面的所有例子都用标定后的 729 TFLOPS / 4502 GB/s，因为它更接近真实可用的水平。

2.3 四个可以手算的例子

下面四个例子的完整输出由 code/01-op-traffic-model/op_traffic.py 复算，数值与正文一致。约定：fp16 输入输出（每个元素 2 字节），机器用 H200 标定参数（729.3 TFLOPS / 4502 GB/s / 平衡点 162）。

例 1：逐元素链。 先看最简单的： $y = sin (cos (x))$ ， $x$ 有 16.8M 个元素（即 $4096 \times 4096$ ）。

eager（两个 kernel）：读 x(33.6MB) + 写 a(33.6MB) + 读 a(33.6MB) + 写 b(33.6MB) = 134.2 MB
融合（一个 kernel）：读 x + 写 b                                          =  67.1 MB
算术强度（融合后）：33.6 MFLOP / 67.1 MB ≈ 0.5 FLOP/Byte

结论：远低于平衡点 162，是典型的 memory-bound；融合把访存量减半，因此时间下界也从 29.8 µs 降到 14.9 µs。这就是 PyTorch 官方文档里那个例子的完整算术 （"we can turn 2 reads (x, a) and 2 writes (a, b) into 1 read (x) and 1 write (b)"，来源见 4.3 节）。

例 2：LayerNorm。 $M = 8192$ 行、 $H = 4096$ 列，fp16。

融合实现：读 x + 写 y                    = 134.2 MB
用基础算子手工拼（均值一趟、方差一趟、归一化+仿射一趟）：
          读 x 三次 + 写 y             = 268.4 MB
运算量（本书口径，每元素约 8 FLOP）：0.268 GFLOP

算术强度只有 1.0–2.0 FLOP/Byte，比平衡点低两个数量级。归一化类算子的优化空间 完全在访存上：能少读一遍就少读一遍。这也是「x.mean()、(x-m).pow(2).mean()、 rsqrt、*w+b 各写一行」这种写法在性能上不可接受的原因——eager 会为每一步各发一个 kernel，中间张量在显存里来回；而 PyTorch 自带的 F.layer_norm 是一个融合 kernel。

例 3：GEMM。 $409 6^{3}$ ，fp16。

运算量：2MNK = 1.37e11 FLOP = 137 GFLOP
访存量：读 A(33.6MB) + 读 B(33.6MB) + 写 C(33.6MB) = 100.7 MB
算术强度：1365 FLOP/Byte   ← 远高于平衡点 162

GEMM 天生是 compute-bound 的。有意思的是它可以翻转：同样按「A、B、C 各读/写一遍」记账，边长 1024 时算术强度降到 341，边长 512 时降到 171——已经贴着平衡点（这两个数字由脚本复算，也是第 1 章习题里出现过的现象）。所以「小矩阵乘」和「大矩阵乘」在 roofline 上是两种东西：大 GEMM 的技巧是「把数据留在片上喂满张量核」，小 GEMM 的技巧是「减少启动与访存」。

例 4：attention。 这是本章最有信息量的一个例子。取 $B = 1$ 、32 个 head、 head_dim $d = 128$ 、序列长度 $S = 4096$ ，fp16。

prefill（标准写法：把 $S \times S$ 的分数矩阵落到显存）

运算量：32 × (4·S²·d + 5·S²) ≈ 278 GFLOP
访存：Q/K/V/O 各一趟 134.2 MB
      + 分数矩阵被写一次、被 softmax 读一次写一次、被 PV 读一次
        = 6 × 32 × 4096² × 2 字节 ≈ 6.44 GB   ← 主体是它，不是 Q/K/V
算术强度：42.2 FLOP/Byte  →  memory-bound

prefill（flash 风格：分数只在片上存活）

运算量：同样的 278 GFLOP
访存：只有 Q/K/V/O 四趟 = 134.2 MB
算术强度：2068 FLOP/Byte  →  compute-bound
时间下界从 1.461 ms 降到 380.6 µs（约 3.8 倍）

decode（KV 长度 32768，每步 1 个 token）

运算量：32 × (4·S·d + 5·S) ≈ 0.54 GFLOP
访存：读 K/V cache = 2 × 32 × 32768 × 128 × 2 字节 ≈ 537 MB
算术强度：1.01 FLOP/Byte  →  极端 memory-bound

把三个数字放在一起，你会看到同一个 attention 的三种面貌：

写法	算术强度	判据	瓶颈
prefill，落 $S \times S$	42.2	memory-bound	中间量在 HBM 往返
prefill，flash 风格	2068	compute-bound	张量核利用率
decode	1.01	memory-bound	KV cache 的读取

FlashAttention 那篇论文的核心论证正是第一条与第二条的对比：把「HBM 的读写次数」当成一等公民来优化，"IO-aware" 指的就是这件事（论文摘要见 FlashAttention: Fast and Memory-Efficient Exact Att…；其报告的训练加速为 BERT-large 15%、GPT-2 3×、LRA 2.4×，测量条件见论文正文）。而 decode 之所以是 memory-bound，是因为每读一个 KV 元素只做一次乘加——这正是 MQA/GQA、KV cache 量化、投机解码这些技术要解决的问题（第 18 章）。

「prefill 是 compute-bound」并不意味着它一定能跑满算力。FlashAttention-3 的论文摘要给了一个公开的数据点：FlashAttention-2 在 H100 上只做到 35% 的利用率；他们用 warp specialization 把 Tensor Core 与 TMA 的异步性重叠起来、并引入 FP8，才把 FP16 推到 740 TFLOPs/s（75% 利用率）、FP8 接近 1.2 PFLOPs/s （来源：FlashAttention-3: Fast and Accurate Attention with …，论文摘要，测量条件见论文正文；这是论文自述的数字，不是本章的实测）。换句话说：判断「它该是 compute-bound」只是第一步，做到那一步是另一回事—— 这正是第 15 章要讲的 attention 演进史。

2.4 把算子拼成模型：两张账本

单个算子的账再清楚，也不足以说明「训练慢在哪」。code/02-step-budget/step_budget.py 把它拼成一个 8B 级模型（32 层、hidden 4096、ffn 14336、32 个 query head / 8 个 KV head、 seq_len 4096、batch 1、8 卡、bf16 权重、fp32 优化器状态）：

训练一步（下界，来自脚本实际输出）
  前向（全部层）    66.1 TFLOP    28.45 GB    90.575 ms   27.9%
  反向（全部层）   132.1 TFLOP    56.91 GB   181.151 ms   55.9%
  优化器            0.12 TFLOP   208.78 GB    46.375 ms   14.3%
  通信                  0         28.11 GB     6.243 ms    1.9%

推理
  prefill（整段 prompt） 66.1 TFLOP   28.45 GB    90.575 ms
  decode（每个 token）    0.016 TFLOP  14.50 GB     3.220 ms

这张表是本书口径的粗模型（假设写在脚本注释里：反向按 2 倍前向 FLOPs、激活只算残差流与 MLP 中间量、通信按 ring all-reduce），它的价值在于给出数量级关系：

训练一步里，算力集中在前向 + 反向（合计约 84%），优化器是纯带宽阶段（120 GFLOP 的算术，却要搬 208.78 GB）；
prefill 与 decode 的判据完全相反，因此优化手段也完全不同——一个要喂满张量核，一个要少读字节。

2.5 这个模型的边界

三件事必须说清楚，否则上面的数字会被误用：

它是下界，不是预测。 真实 kernel 一定更慢。TileOPs 的 roofline 文档甚至定义了一类专门的失败模式：当 sol_time 和实测时间都低于某个阈值时，说明 kernel 太小、启动开销占主导，模型对这种情况「没有抓手」，只能标成 latency-bound （TileOPs/docs/design/roofline.md 第 40–42 行）。
cache 与 locality 被忽略了。 模型假设每次访问都直接去 HBM。真实程序有 L2、有片上复用，实际访存可能远低于（也可能高于）这个下界。
超越函数的记账是「账本上记 1 FLOP」，不是硬件成本。exp、tanh 在硬件上走 SFU，吞吐远低于 FMA。所以 softmax、SiLU 密集的算子，实际时间会比这个模型给出的更多， TileOPs 的文档也明确写了「不保证 SFU-bound 的 kernel 就在极限」（同文件 :44-52）。

3. 用法：一行 PyTorch 到底做了什么

两条路径最终都落到 kernel 上，差别只在「谁来决定有几个 kernel、每个 kernel 里放多少算子」。

eager 的代价：每个 Tensor 运算都要经过 dispatcher、选中一个 aten 算子，再展开成一到多个 kernel launch，中间结果写回显存。
eager：读 x、写 a、读 a、写 b = 2 读 + 2 写；Inductor 融合后：读 x、写 b = 1 读 + 1 写（PyTorch 官方文档）。
Inductor 生成的 kernel 名字形如 triton_poi_fused_<aten算子>_<序号>：poi 表示 pointwise，fused 表示多个算子被融进同一个 kernel。
图中「谁来决定 kernel 边界」这一层就是第 26 章要讲的集成点：是让 Inductor 生成 kernel，还是把一个 TileLang 算子挂在同一个位置。

3.1 eager 模式：一行代码 = 一串 kernel launch

PyTorch 默认是 eager（即时执行） 模式。y = torch.sin(torch.cos(x)) 的执行过程是：

Python 解释器逐行执行，先算 torch.cos(x)，再算 torch.sin(...)；
每个 Tensor 上的运算都经过 dispatcher，选中一个 aten 算子（如 aten::cos）；
这个算子被展开成一次（有时是多次）kernel launch，写回一张新分配的显存张量。

所以 cos 之后接 sin，在 eager 下就是两个 kernel，中间结果 a 必须落到显存再读回来。 PyTorch 官方文档对这件事的说法是：

"Pointwise ops in eager mode are suboptimal because each one would need to read a tensor from the memory, make some changes, and then write back those changes." —— PyTorch 官方文档「Getting Started with torch.compile」， https://docs.pytorch.org/docs/2.14/user_guide/torch_compiler/torch.compiler_get_started.html （2026-10-01 核验）

除了访存，eager 还有第二个代价：启动开销。每个 kernel 都要从 Python 侧发起，小算子上这笔开销可能比计算本身还大。PyTorch 的性能文档里有一个很直观的口径：在 densenet121 这种小算子的模型上，小 batch 时 "Percent of time when GPU is busy" 只有 32.69%，也就是说大部分墙钟时间花在 CPU 侧的准备与启动上（https://docs.pytorch.org/docs/2.14/user_guide/torch_compiler/torch.compiler_inductor_profiling.html，2026-10-01 核验；测量条件见该页所引的 benchmark 脚本）。

3.2 你会看到什么：kernel 表

判断「我的模型里有多少个 kernel、每个叫什么」，标准做法是看 profiler。用 TORCHINDUCTOR_UNIQUE_KERNEL_NAMES=1 之后，Inductor 生成的 kernel 会带上语义化名字，形如 triton_poi_fused_cat_155：poi 表示 pointwise，fused 表示它把多个算子融进了同一个 kernel，cat 是被融合的 ATen 算子的名字（来源同上）。

这个命名规则本身就是一条很好的证据链：Inductor 的 kernel 边界是「融合后的算子组」，不是「一个 ATen 算子一个 kernel」。 名字里的 cat 告诉你哪些 eager 下会各占一个 kernel 的算子在编译后合并了。

3.3 用本章的脚本复算

python3 code/01-op-traffic-model/op_traffic.py          # 单算子账本
python3 code/02-step-budget/step_budget.py              # 一次迭代/一次请求

两个脚本只用标准库，输出就是本章的表格。建议按下面顺序做一次：

先用默认参数跑一遍，确认能复现正文里的数字；
把 --machine h100-spec 换成 H100，看平衡点从 162 变成 333，哪些算子的判据会翻转；
把 --seq-len 从 4096 改到 32768，看训练一步里 attention 的占比怎么变；
把 --params 改成 70e9、--world 改成 64，看优化器与通信的占比怎么变。

这四步做完，你就有了「不看 profiler 也能对量级有判断」的能力——这正是后面第 28 章（性能模型与算子设计方法论）的起点。

4. 实现机制：融合是默认答案，torch.compile 做到哪一步

融合不是「少写几行代码」，而是把中间张量的读写从显存里省掉。

eager 的结构：read x → cos 写 a → read a → sin 写 b；融合的结构：read x → cos→sin（中间值只在寄存器）→ write b。
第三行是量级上的差别，也是 FlashAttention 把「读写 HBM 的次数」当一等公民来优化的核心论证。
同样的动机写在昇腾 CANN 的算子文档里（如 aclnnAddRmsNormDynamicQuant「减少搬入搬出操作」）；表格中访存量由 code/01-op-traffic-model/op_traffic.py 输出。

4.1 逐算子执行的两个代价

把 2.3 节的例子总结成一条规律：

每一个「读入 → 计算 → 写出」的算子，都要为它读写的每个字节付一次 HBM 的钱。
链上第 k 个算子读的是第 k-1 个算子刚写下去的中间张量——这次往返本可以完全避免。

于是有两类代价：

访存代价：中间张量的写 + 读，等于给数据多买了两次票（例 1、例 2）；
启动与调度代价：kernel 数量多、每个都很小的时候，启动开销占比升高（3.1 节）。

注意这两类代价都是在算术强度很低的时候才致命。GEMM 融合与否，差别有限（省下的只是输出矩阵 C 的写出与再读入，而它的访存是被 A、B 主导的）；逐元素链和归一化融合与否，差别是 2 倍到几十倍。融合的收益正比于「省下的字节数占原本访存量的比例」，而不是正比于「省下了几个算子」。

4.2 融合的账本与厂商证据

图 5（图：05-fusion-ledger）给出了本章三个融合例子的账本：

例子	融合前访存	融合后访存	下界之比
$sin (cos (x))$ ，16.8M 元素	134.2 MB	67.1 MB	2.0×
LayerNorm（手工拆三趟）	268.4 MB	134.2 MB	2.0×
attention prefill（落 $S \times S$ ）	6.58 GB	134.2 MB	49×

融合不只是「编译器的小技巧」，它是生产算子库的组织方式。三个来自昇腾 CANN 官方文档的证据：

aclnnAddRmsNormDynamicQuant：「将RmsNorm前的Add算子和RmsNorm归一化输出给到的1个或2个 DynamicQuant算子融合起来，减少搬入搬出操作」（hiascend.com · canncommercial_900_API_aolapi_context_ops-nn_aclnnAddRmsNormDynamicQuant_md 第 25 行）；
aclnnKvRmsNormRopeCache：切分 + RMSNorm + RoPE + 写入两块 cache，一个算子；
aclnnApplyAdamW：整个优化器更新，一个算子。

TileLang 侧的对应物是 examples 里的融合算子：examples/norm/rms_norm.py 把「平方求和 → 归约 → rsqrt → 缩放」写在一个 kernel 里（_vendor/tilelang-v0.1.15/examples/norm/rms_norm.py:22-30 里的 T.reduce_sum 与 T.rsqrt 就在同一个 T.Kernel 作用域内）；examples/elementwise/example_elementwise_add.py 是逐元素融合的最小例子（_vendor/tilelang-v0.1.15/examples/elementwise/example_elementwise_add.py:11-33）。

4.3 torch.compile / TorchInductor 能做到什么

torch.compile 想做的事，就是把 4.1 节的两类代价交给编译器去消除。它的组件与分工（来源：PyTorch 官方 torch.compiler 文档， https://docs.pytorch.org/docs/2.14/user_guide/torch_compiler/torch.compiler.html，2026-10-01 核验）：

组件	做什么
TorchDynamo	用 CPython 的 Frame Evaluation API 捕获 PyTorch 计算图，得到一张 Aten IR 图；遇到无法捕获的代码就产生 graph break，回退到 eager
AOTAutograd	「captures not only the user-level code, but also backpropagation」——提前把反向也捕获成图，于是前向和反向都能被编译
TorchInductor	默认后端：先做图级优化（融合、内存复用、乱序执行、工作放置），再生成代码；在 NVIDIA/AMD/Intel GPU 上「leverages OpenAI Triton as the key building block」

Inductor 自己把「最重要的优化」说得很清楚：

"The single most important optimization that inductor performs is fusion." —— PyTorch 官方文档「Getting Started with torch.compile」（URL 见上；2026-10-01 核验）

除了融合，官方 FAQ 还列了另外三条通用手段：水平融合（把多个矩阵乘排在一起执行，例如 grouped GEMM）、乱序执行（按数据依赖重排节点、复用缓冲区）、自动工作放置 （把节点映射到合适的硬件与存储资源）。FAQ 里还有一句很重要的话：Inductor "first takes care of fusing whatever it can and only then generates Triton kernels" ——先决定融合边界，再生成 kernel。（https://docs.pytorch.org/docs/2.14/user_guide/torch_compiler/torch.compiler_faq.html，2026-10-01 核验）

在工程上还有两件值得知道的事：

自动调参：TORCHINDUCTOR_MAX_AUTOTUNE 会让 Inductor 多试几组 Triton 配置并选实测最快的一组；这对应本书第 11 章的自动调优话题。
CUDA Graph：Inductor 可以自动把一串 kernel 用 CUDA Graph 包起来，消除逐 kernel 的启动开销（这就是 3.1 节里「GPU busy 只有 32%」这类问题的解法）。

4.4 边界在哪

torch.compile 覆盖不到的地方，恰好就是你要自己写 kernel 的地方。四类边界：

（a）图捕获本身会断。 Dynamo 遇到它无法捕获的代码就 graph break，回退到 eager； fullgraph=True 会把 graph break 变成错误而不是静默回退。graph break 的代价是「这段代码不参与融合」，而且反向在 break 处无法继续（「We cannot propagate gradients through a graph break」，同上 FAQ 页）。

（b）融合能力受限于「它能看见什么」。 Inductor 的融合是图级的：它能把相邻的逐元素算子合并，但它不会为你发明一个代价模型驱动的、跨层重排的数据流。当一个算子的性能取决于「tile 多大、数据放在哪一级存储、几条流水怎样重叠」时（attention、MLA、 MoE、paged decode），这些决策不在它的融合抽象里。

（c）后端目标有限。 Inductor 的主要 GPU 后端建立在 Triton 之上；昇腾 NPU 不在 Triton 的目标列表里，这也是昇腾侧需要 TileLang 与 CANN 算子库的原因之一（这是本书基于上面官方文档措辞的判断，不是文档原话）。

（d）自定义算子的接缝有明确约束。 把 TileLang（或其他手写）算子挂进 torch.compile 的图里，需要走 torch.library.custom_op 这类注册接口。TileOPs 项目把这些约束写得很具体（TileOPs/src/tileops/ops/compile_boundary.py 第 1–25 行）：被 Dynamo 追踪的 forward 不能构造 kernel、不能进入 TileLang builder——缓存查找、 kernel 构造与启动都必须发生在未追踪的一侧；传给自定义算子的 key 必须是 str，因为「dynamo treats string operator arguments as compile-time constants, while an int is generalized to an unhashable SymInt」；而且 Inductor 会把 fake tensor 的输出形状烤进磁盘缓存，所以同一个 key 绝不能在不同实例之间复用。同一套设计文档还把一条纪律写死了：「PyTorch is the spec oracle, not a runtime path」——算子层不允许在 forward 里直接调用 torch.sum、F.softmax 这类高层算子，否则「库里测的和发的其实是 PyTorch」（TileOPs/docs/design/layer-boundaries.md）。

本机限制：本章写作环境没有 GPU/NPU，上面提到的 TORCH_COMPILE_DEBUG、 TORCHINDUCTOR_UNIQUE_KERNEL_NAMES、Nsight Compute 都没有在本机跑过； 4.3 节的所有说法都直接来自 PyTorch 官方文档的原文，4.4 节的约束来自 TileOPs 源码快照。它们的可复现步骤见 资料库索引 的「在真实机器上怎么接着往下做」。

5. 进阶与陷阱：什么时候你需要自己写 kernel

前面四节其实已经在回答这个问题了。把判据整理成一页清单（图 6， 图：06-when-to-write-a-kernel）：

前置动作：先测量。 用 profiler 找出占 GPU 时间最多的几个算子，先试 torch.compile 和厂商库里的现成融合算子。不做这一步就直接写 kernel，是资源浪费。 下面四条是「试过之后仍然要做」的理由。

判据 1：融合缺失，且省下的访存占比很大。 典型是「一串低算术强度的算子被拆成很多小 kernel」——norm + 残差 + RoPE + 量化、mask + softmax 的组合。判断方式就是本章的账本：算一遍融合前后的 bytes，比值在 2 倍以上就值得做。

判据 2：库里根本没有这个算子。 新论文里的算子（MLA、NSA、DSA indexer、稀疏 attention、新的位置编码）在库更新之前都不存在。生产里能看到这类算子的 TileLang 实现，比如 SeerAttention-R 报告「Using TileLang, we develop a highly optimized sparse decoding kernel」（SeerAttention-R: Sparse Attention Adaptation for Lo…），以及把 NSA 的设计动机直接写成「arithmetic intensity-balanced algorithm design」的论文（Native Sparse Attention: Hardware-Aligned and Nativ…）。

判据 3：需要库不支持的低精度或布局。 fp8 block scaling、int4 权重、 NZ/fractal 布局、按 block 的 scale 与反量化——这些是「数据怎么摆」的问题，库要么不支持，要么不支持你要的那种组合。第 27 章会专门讲数值与低精度。

判据 4：单个算子已经贴着下界，但整段流水还在等。 这时要动的是跨算子的数据流：多缓冲、流水线、persistent kernel、把多个算子放进一次 launch。TileLang 把这类能力做成了语言里的显式概念（T.Pipelined(num_stages=...)、存储作用域），第 11、22 章会展开；昇腾 README 里那张「你写 TileLang，Ascend C 里可以省掉什么」的表，其中「Synchronization | None (Just write operations in order!)」一行说的就是这件事（_vendor/tilelang-v0.1.15/tilelang/ascend/README.md:141-152）。

然后是风险：写得对≠写得快。 一篇 2026 年的实证研究在 NVIDIA A100 与 GH200 上检查了 22 个 Triton/TileLang kernel，结论是：「an idiomatic TileLang LayerNorm kernel passes KernelBench's correctness check while running more than 300× slower than the PyTorch baseline」；而且慢的原因在不同算子族里不一样——归一化/归约类的慢「mainly repairable authoring defects, such as sequential reductions and unnecessary dtype conversions」，卷积与大 GEMM 的差距则来自代码生成与调优覆盖不足。这篇论文提出的两个轻量验收判据正是本章的工具：库相对效率（和 PyTorch/厂商库比）与 roofline 利用率 （和本章算出的下界比）——「together they flag every functionally valid but inefficient kernel in our suite」（Correct but Slow: An Empirical Study of the GPU Ker…）。

所以正确的顺序是：

测量 → 算账（本章）→ 先试 torch.compile / 库算子 → 仍然要写，就用 roofline 利用率验收

四条判据分别对应四种「试过 torch.compile 与厂商库之后仍然要写」的情况；起点永远是先测量。

判据 1 的完整表述：相邻的逐元素 / 归一化算子被拆成很多个小 kernel → memory-bound × 多趟访存。
判据 2：库算子覆盖不到这个融合模式（MLA / NSA / indexer / RoPE + cache 写入）。
判据 3：需要库不支持的低精度 / 布局（fp8 block scaling、int4 权重、NZ / fractal 布局）。
判据 4：单个算子已经贴着下界，但整段流水仍有串行等待（流水、persistent、多缓冲）。
决策树本身是本书的归纳，判据在正文 5.1–5.4 逐条给出证据；反例来自 A100 / GH200 上 22 个 kernel 的实测（正文 5.5 节）。
正确的顺序：测量 → 算账 → 先试 torch.compile / 库算子 → 仍然要写，就用 roofline 利用率与库相对效率验收。

5.1 本章的三条常见陷阱

把厂商规格当基线。 1000 TFLOPS 与 729 TFLOPS 不是一个东西；calibration 这类标定值连测量条件（显存频率、功耗上限、working set）都会改变。引用数字时必须带条件。
只看 FLOPs 不看字节。 本章一半的例子（LayerNorm、decode attention、AdamW）算术强度不超过 2 FLOP/Byte，它们的性能与 FLOPs 基本无关。
把「算子数量」当瓶颈。 融合的收益取决于省下的字节占比，不是省下的算子个数。一条 $k$ 个算子的逐元素链，理论上最多把这条链的访存降到 $1/ k$ ；但如果这条链的访存只占整个模型的百分之一，端到端收益仍然不到百分之一。

5.2 本章不做什么

不教你写 TileLang kernel：第 7–9 章从安装与第一个 kernel 开始。
不做真实的性能测量与调优：第 12 章（调试与性能分析）、第 11 章（流水线与调优）。
不展开各类算子的完整实现：第 13–16 章。
不展开编译器内部：第 4 章给全流程，第 20–26 章深入 pass 与后端。

6. 本章小结

算子在地图上的位置决定了它的命运：训练一步里，前向/反向是算力阶段，优化器是纯带宽阶段（每参数 ±26 字节），通信是纯搬运（ring all-reduce 搬 $2 (N - 1) / N$ 倍梯度）；推理请求里，prefill 是算力阶段，decode 是带宽阶段（每 token 读全套权重 + 整段 KV cache）。
账本方法：数出 $(f l o p s, b y t es)$ → 算算术强度 → 和机器平衡点比 → 得到 memory-bound / compute-bound 与时间下界。四个手算例子：逐元素链（0.5 F/B）、 LayerNorm（1–2 F/B）、GEMM 4096³（1365 F/B）、attention（落 $S^{2}$ 42 F/B、 flash 风格 2068 F/B、decode 1 F/B）。
Eager 的代价是「每个算子都要为它读写的字节付一次 HBM 的钱」，外加逐 kernel 的启动开销；cos 接 sin 这种两算子链，融合后访存量减半。
torch.compile 的分工：Dynamo 捕获图（graph break 就回退）、AOTAutograd 把反向也捕获、Inductor 先融合再生成 Triton kernel，另外还有 CUDA Graph 与自动调参。它覆盖「整图怎么切、相邻算子怎么并」，不覆盖「单个复杂算子的数据流怎么排」。
什么时候自己写 kernel：融合缺失且省下的访存占比大、库里没有这个算子、需要不被支持的低精度/布局、需要跨算子的数据流重组。先测量，再用 roofline 利用率验收，因为「正确但慢 300 倍」是真实发生过的事。
TileLang 的位置：Inductor 之下、CUDA/Ascend C 之上，负责「一个算子内部怎么写快」。这正是第 3 章要开始讲的东西。

习题

13 题 · 答案在提交后展开

先在纸上或输入框里写下你的答案，再展开参考答案对照。客观判断请自己打勾：这三颗按钮只影响本机保存的学习进度。

E1综合难度 ★无需硬件

把一次训练迭代的四个阶段——前向、反向、优化器、通信——分别归类到「compute-bound / memory-bound / 纯搬运」，并对每一个给出一句话的理由。对优化器那一段，请给出「每个参数要搬多少字节」的来源（哪份文档、哪个张量）。

思路：回到「算术强度 vs 平衡点」这一个判据，逐阶段看它主要花在哪。

参考答案

阶段	归类	理由
前向	compute-bound	主体是 GEMM 与 attention，长序列下 FLOPs 远大于必须搬动的字节数（8B 模型、seq 4096：66.1 TFLOP 对 28.45 GB）
反向	compute-bound	每个前向 GEMM 的反向是它的 2 倍 FLOPs（ $d A = d C B^{⊤}$ 、 $d B = A^{⊤} d C$ 各 $2 M N K$ ），合计约 3 倍前向 FLOPs
优化器	memory-bound	每参数读 grad/param/m/v、写 param/m/v；bf16 梯度 + fp32 状态是 26 字节/参数，算术只有十几 FLOP
通信	纯搬运（memory/网络 bandound）	ring all-reduce 搬 $2 (N - 1) / N$ 倍梯度数据，0 FLOP

来源：优化器的张量与公式见 hiascend.com · canncommercial_900_API_aolapi_context_ops-nn_aclnnApplyAdamW_md 第 25–60 行 （varRef/mRef/vRef 三个输入输出张量 + grad）。

常见错误：把反向归成 memory-bound（只想到「要读激活」，忽略了它同样是 2 倍 FLOPs 的 GEMM 主场）；或者忘记「反向的 FLOPs 是前向的 2 倍」是对 GEMM 精确、对 attention 偏乐观。

E2综合难度 ★无需硬件

同一个 attention 算子，在 prefill 与 decode 两段里的判据相反。请用算术强度的语言解释这件事：prefill 里每读一个字节的数据做多少次乘加？decode 里呢？（提示：decode 的 query 只有 1 个 token。）

E3综合难度 ★★无需硬件

本章说「厂商规格和能跑到的不是一个数」。请举出本章用到的两个具体事实支持这句话，并说明它们各自的测量/标定条件。如果用 989.5 TFLOPS、4800 GB/s（厂商规格）而不是 729.3 TFLOPS、4502 GB/s（标定值）来算平衡点，会得到什么值？这个差别会改变本章哪些算子的判据、不会改变哪些？

E4综合难度 ★★无需硬件

一条 $k$ 个算子的逐元素链，融合后这条链的访存从 $2 k$ 趟降到 $2$ 趟。那么是不是 $k$ 越大收益越大？请说明「这条链本身的收益」和「端到端收益」分别由什么决定，并给出一个「 $k$ 很大、端到端收益却很小」的场景。

计算题

E5综合难度 ★★无需硬件

手算 LayerNorm， $M = 4096$ 、 $H = 4096$ 、fp16：

融合实现的访存量与算术强度；
「均值一趟、方差一趟、归一化+仿射一趟」这种朴素拆分的访存量与算术强度；
在 H200 标定参数下，两者的时间下界各是多少；
如果把 $H$ 从 4096 改成 512，结论会不会变？为什么？

E6综合难度 ★★无需硬件

用 H100 的厂商规格（1000 TFLOPS fp16 / 3.0 TB/s，平衡点 333）：

手算 $204 8^{3}$ 的 fp16 GEMM 的算术强度与判据；
边长是多少时它刚好落在脊点上（算术强度等于平衡点）？给出表达式而不是只给数字；
用本章脚本验证你的两个答案。

思路：先写通用式，再代数字。

参考答案

$N = 2048$ ：FLOPs $= 2 N^{3} = 1.718 \times 1 0^{10}$ ；bytes $= 3 N^{2} b = 3 \cdot 204 8^{2} \cdot 2 = 25.17$ MB；算术强度 $= 682.7$ FLOP/Byte，compute-bound（>333）。
通用式： $AI = \frac{2 N ^{3}}{3 N ^{2} b} = \frac{2 N}{3 b} = \frac{N}{3}$ （fp16， $b = 2$ ）。令 $N /3 = 333 \Rightarrow N \approx 999$ 。也就是说只有边长小于约 1000 的 GEMM 才会掉进 memory-bound 区（这里用的是「A/B/C 各读或写一遍」的记账口径，真实实现因为有 cache 与片上复用，边界会移动）。

复算：

python3 -c "import op_traffic as t; \
print(t.judge(t.gemm(2048,2048,2048), t.MACHINES['h100-spec']))"

（在该目录下运行）会得到 compute-bound；把 2048 换成 512 会得到 memory-bound。

常见错误：把 $b$ 写成 1（fp8）或 4（fp32）却不说明；或者忘记 $3 N^{2}$ 里的系数 3 （A、B、C 三块）。

E7综合难度 ★★★无需硬件

$attention$ prefill， $B = 1$ 、32 head、 $d = 128$ 、 $S = 8192$ 、fp16：

手算「把 $S \times S$ 的分数矩阵落到显存」这一版的总访存量（写出每一项是什么）；
手算 flash 风格（分数不落显存）版本的访存量；
两者的比值是多少？如果把 $S$ 从 8192 改成 16384，这个比值会怎么变？由此解释为什么这个中间量被称为 $O (S^{2})$ 问题。

代码阅读题

E8综合难度 ★★无需硬件

阅读 code/01-op-traffic-model/op_traffic.py 里的 attention_prefill 函数，回答：

scores_bytes 的三行分别代表哪三次访存？为什么第一行是写、后两行是读+写+读？
flops 表达式里的 4.0 * s * s * d 和 5.0 * s * s 分别是什么？
如果要模拟「softmax 与 PV 已经融合、但分数矩阵仍然落显存」的实现， scores_bytes 应该怎么改？

E9综合难度 ★★无需硬件

阅读 _vendor/tilelang-v0.1.15/tilelang/jit/__init__.py 中 compile 的签名（约 94–130 行）：

execution_backend 有哪些合法取值？out_idx 的作用是什么？
target 参数接受哪些形式的输入？docstring 给的例子是什么？
从「算子在栈里的位置」这个角度看，这个签名说明了 TileLang 把自己定位在哪一层？用一句话回答，并说明你的依据。

E10综合难度 ★★★无需硬件

阅读 TileOPs/src/tileops/ops/compile_boundary.py 开头的模块文档字符串（约 1–25 行）：

为什么传给自定义算子的 key 必须是 str 而不是 int？原文怎么说的？
文档里说「inductor bakes the fake's output shape into an artifact it caches on disk」。这条事实对「同一个 key 不能复用」这个约束起了什么作用？
结合这段文字，解释为什么「被 Dynamo 追踪的 forward 不能构造 kernel」。

动手题

E11综合难度 ★★无需硬件

运行：

python3 code/01-op-traffic-model/op_traffic.py
python3 code/01-op-traffic-model/op_traffic.py --machine h100-spec

列出两份输出中判据（memory-bound / compute-bound）发生翻转的算子，以及没有翻转的算子。对每一类各挑一个，解释翻转或不翻转的原因。再回答：为什么「逐元素加」不管用哪台机器都不可能变成 compute-bound？

思路：跑两组参数，逐行比对判据列。

参考答案

python3 code/01-op-traffic-model/op_traffic.py                      # H200 标定，平衡点 162
python3 code/01-op-traffic-model/op_traffic.py --machine h100-spec  # H100 规格，平衡点 333

默认表格里的算子（逐元素加、逐元素链 ×2、LayerNorm ×2、GEMM 4096³、attention prefill ×2、 attention decode、AdamW、all-reduce，算术强度从 0.17 到 2068）没有一个翻转：它们要么远低于 162/333，要么远高于。翻转只出现在边界算子上，例如 fp16 GEMM 边长 512（算术强度 170.7）：在平衡点 162 下是 compute-bound，在 333（H100）或 206.1（H200 厂商规格）下是 memory-bound。

为什么逐元素加永远不可能变成 compute-bound？它的访存是 $3 N b$ 字节（读 $x$ 、读 $b$ 、写 $y$ ），算术只有 $N$ FLOP，算术强度恒为 $1/ (3 b) = 1/6$ FLOP/Byte（fp16），与规模无关。要让它变成 compute-bound，机器的平衡点得低于 1/6，也就是带宽比算力贵 6 倍—— 现实中不存在这样的加速器。

常见错误：以为「大一点就变成 compute-bound」。算术强度是比例，对逐元素、归约这类算子与规模无关；只有 GEMM、attention 这类「计算量随某个维度超线性增长」的算子，算术强度才会随规模增长。

E12综合难度 ★★★无需硬件

先用默认参数跑 code/02-step-budget/step_budget.py，再把 --seq-len 依次改成 512、 4096、32768、131072：

记录每个序列长度下「前向」「反向」「优化器」「通信」的占比；
找出占比变化最快的那一段，并解释为什么是它；
结论里必须提到：这四个阶段里，哪一段的绝对时间几乎不随序列长度变化。

思路：固定模型与机器，只改序列长度。

参考答案（python3 code/02-step-budget/step_budget.py --seq-len <S>）

seq_len	前向	反向	优化器	通信
512	12.1%	24.2%	56.2%	7.6%
4096	27.9%	55.9%	14.3%	1.9%
32768	32.9%	65.8%	1.1%	0.1%
131072	33.3%	66.6%	0.1%	0.0%

占比变化最快的是优化器：从 56.2% 掉到 0.1%。
原因：优化器的访存量（8B 参数约 208.78 GB）与"每个参数"有关，与序列长度无关；而前向/反向的 FLOPs 随序列长度（attention 部分甚至是 $S^{2}$ ）增长。分母长大，分子不变，占比自然被稀释。
几乎不随序列长度变化的绝对时间是优化器（46.375 ms）与通信（6.243 ms）—— 注意通信的占比在变，是因为它在总时间里的份额变了，它自己的绝对量没变。

顺带一个结论：短序列（batch 小、序列短）的训练里，训练时间的相当一部分花在 "与序列长度无关"的优化器与通信上，这时优化算子的收益上限就是那 40%–60% 之外的余量。

常见错误：把「占比下降」当成「绝对时间下降」。绝对时间没变，变的是分母。

E13综合难度 ★★★无需硬件

给下面两个算子各写一个账本函数（可以照 op_traffic.py 里 layernorm 的写法），加进脚本并跑通：

SiLU 激活： $y = x \cdot σ (x)$ ， $N$ 个元素；
RoPE 的旋转部分（不考虑 norm，不考虑写 KV cache）：对 $B \times H \times S \times d$ 的张量，按 dim 两两配对，每对做 $y_{1} = x_{1} cos θ - x_{2} sin θ$ 、 $y_{2} = x_{2} cos θ + x_{1} sin θ$ （ $θ$ 已预计算）。

对每个算子：写出 flops（要给出你采用的每元素 FLOP 计数与理由）、bytes、算术强度、在 H200 标定参数下的时间下界与判据。最后回答：如果把它们和「残差相加」融合进一个 kernel，访存量会怎么变？

思路：照 layernorm 的写法，先定「读什么、写什么」，再定「每元素几个 FLOP」，最后交给 judge。

参考答案（示例代码，可直接接在 op_traffic.py 后面）

def silu(n: float, dtype: str = "fp16") -> Cost:
    elem = ELEM_BYTES[dtype]
    # sigmoid = neg + exp + add + recip = 4 FLOP；再乘 x = 5 FLOP/元素
    # （计数口径见 TileOPs roofline.md：超越函数在账本上记 1 FLOP）
    return Cost(
        name=f"SiLU ({n/1e6:.1f}M 元素, {dtype})",
        flops=5.0 * n,
        nbytes=2.0 * n * elem,          # 读 x，写 y
        assumption="每元素 5 FLOP（sigmoid 4 + 乘 1）；读 x、写 y",
    )


def rope(b: float, heads: float, s: float, d: float, dtype: str = "fp16") -> Cost:
    elem = ELEM_BYTES[dtype]
    n = b * heads * s * d
    # 每两个元素配成一对：各 2 次乘 + 1 次加/减 = 每对 6 FLOP，即每元素 3 FLOP
    return Cost(
        name=f"RoPE 旋转部分 ({b:.0f}x{heads:.0f}x{s:.0f}x{d:.0f}, {dtype})",
        flops=3.0 * n,
        nbytes=2.0 * n * elem,          # 读 x、写 y；cos/sin 预先算好，忽略
        assumption="每对元素 4 乘 2 加减 = 6 FLOP，即每元素 3 FLOP；cos/sin 预计算",
    )

运行结果（H200 标定参数， $N = 16.777$ M 元素，两者形状取 $B = 1$ 、32 head、 $S = 4096$ 、 $d = 128$ ，即 $N = 16.777$ M）：

算子	FLOPs	访存	算术强度	时间下界	判据
SiLU	83.9 MFLOP	67.11 MB	1.25	14.9 µs	memory-bound
RoPE（旋转部分）	50.3 MFLOP	67.11 MB	0.75	14.9 µs	memory-bound

两个都是 memory-bound，时间下界完全由「读一遍、写一遍」决定。如果把它们与残差相加融合进同一个 kernel，访存量不变（还是读入一次、写出一次）—— 融合的收益体现在「省掉了它们各自读写的中间张量」，而不是省掉这两次读写本身。换句话说：这几个算子单独看都是 2 趟访存；把它们串起来才会变成 2k 趟。

常见错误：

忘记超越函数的计数规则（TileOPs 的口径是账本上记 1 FLOP，不是硬件成本）；
把 RoPE 的 cos/sin 也算进访存（它们是预计算的，与 token 数同阶但小得多，可忽略——但要在答案里说明这个假设）；
认为融合会减少「输入→输出」的读写次数。融合减少的是中间张量的往返。

本章代码

5 个文件

正文里提到的每个示例都能在这里直接阅读（带语法高亮），也可以整章浏览。

浏览本章全部代码 →

本章引用

源码链接指向锁定 tag

本章写到的每个源码位置与外部资料都在下面列出：源码链接指向锁定 tag 的 GitHub 文件，资料链接指向原文。查看全部资料索引 →

展开引用清单

第 02 章参考资料

本章所有结论都基于下表中的本地资料或锁定版本源码。版本以 版本锁定 为准：TileLang v0.1.15 （commit a35f8ddf45eba16c21211ec8822d56ce5363036f）、TileLang-Ascend TileLang-ascend-v0.1.2.000-release（commit 89a9364e47a341fffdd65ed94786b1d4f9b6a78c）。文中不出现「最新版」这类表述。

写作环境：macOS (Apple Silicon)、Python 3.11；没有 NVIDIA GPU，也没有昇腾 NPU。因此本章没有任何性能数字来自本机实测；code/ 下两个示例都能在 CPU 上离线运行，验证级别见 资料库索引。

锁定的源码引用（`_vendor/tilelang-v0.1.15/`）

id	引用位置	用途	核验方式与状态
S1	`_vendor/tilelang-v0.1.15/tilelang/jit/__init__.py:94-140`	`tilelang.compile(...)` 的签名：`out_idx`、`execution_backend` 的六种取值、`target` 支持字符串/字典/`Target`	已核验（2026-10-01）：`rg -n "^def compile" tilelang/jit/__init__.py` → 94；逐行读签名与 docstring
S2	`_vendor/tilelang-v0.1.15/tilelang/jit/__init__.py:571-600`	`@tilelang.jit` 装饰器（用户侧入口）	已核验（2026-10-01）：`rg -n "^def jit"` → 571/575/588（三个重载）
S3	`_vendor/tilelang-v0.1.15/tilelang/jit/kernel.py:41,206,516,532`	`class JITKernel`、`__call__`、`get_profiler()`、`get_kernel_source()`	已核验（2026-10-01）：`rg -n "class JITKernel\|def __call__\|def get_profiler\|def get_kernel_source"` 逐条命中
S4	`_vendor/tilelang-v0.1.15/tilelang/profiler/bench.py:31-60`	`do_bench(...)`：多种计时后端、两次运行之间清 L2、按运行时长自动定 warmup/repeat	已核验（2026-10-01）：`rg -n "def do_bench"` → 31；「Clearing L2 cache between runs」在 :48
S5	`_vendor/tilelang-v0.1.15/examples/norm/rms_norm.py:22-30`	RMSNorm 的融合写法：`T.reduce_sum` 与 `T.rsqrt` 在同一个 `T.Kernel` 作用域内	已核验（2026-10-01）：`rg -n "reduce_sum\|rsqrt"` → 25、27；文件可被 `python3 -m py_compile` 解析
S6	`_vendor/tilelang-v0.1.15/examples/elementwise/example_elementwise_add.py:11-33`	逐元素融合的最小例子（`@tilelang.jit` + `T.Parallel`）	已核验（2026-10-01）：`rg -n "@tilelang.jit\|T.Parallel"` → 11、27
S7	`_vendor/tilelang-v0.1.15/examples/deepseek_mla/example_mla_decode_paged.py:30-38`	paged decode 的真实接口形态：`block_table` 与 `cache_seqlens`	已核验（2026-10-01）：`sed -n '30,38p'` 逐行阅读；文件可被 `python3 -m py_compile` 解析
S8	`_vendor/tilelang-v0.1.15/tilelang/ascend/README.md:141-152`	「你写 TileLang，Ascend C 里可以省掉什么」对照表：缓冲与数据搬运、tiling、SIMT/SIMD、调度与流水、同步	已核验（2026-10-01）：`rg -n "Compared with writing Ascend C directly"` → 141，表格到 152

「已核验」的含义：在锁定版本的工作树上用 rg -n 定位到定义行，并阅读了完整签名或 docstring；行号区间与文件内容一致。这不是运行验证——本章没有任何 API 在真实硬件上运行过。

corpus 资料引用（`资料库索引`）

抓取日期：2026-09-30（Asia/Shanghai），见 corpus 的 资料库索引。

id	本地路径	上游链接	用途	核验状态
R1	`NVIDIA Hopper Architecture In-Depth`	https://developer.nvidia.com/blog/nvidia-hopper-architecture-in-depth/	H100 SXM5：HBM3「over 3 TB/sec」（:334）、「Peak FP16 Tensor Core 1000 TFLOPS」（:141，Table 1，dense）、shared memory 可配到 228 KB（:381）	已核验（2026-10-01）：本地为全文；厂商公开资料，Table 1 自述「preliminary performance specs」，非实测
R2	`TileOPs/src/tileops/perf/profiles/h200.yaml`	https://github.com/tile-ai/TileOPs（快照 commit `50dae81fadaf`）	H200 机器参数：fp16 989.5 TFLOPS / HBM 4800 GB/s，标定 0.7370（burst）、0.6684（sustained）、0.9380（带宽最大值）、0.8183（copy 混合）；标定条件写在文件注释里	已核验（2026-10-01）：本地为仓库快照，文件头注释给出测量条件（2048 MB working set、5 次 ×200 轮、显存 3201 MHz、700 W 限制）
R3	`TileOPs/docs/design/roofline.md`	https://github.com/tile-ai/TileOPs	记账法（:31 中间量不计）、判据（:14-27）、latency-bound 定义（:40）、FLOP 计数惯例（:47-58）、SFU 提醒（:34）	已核验（2026-10-01）：逐行阅读上述行号
R4	`TileOPs/src/tileops/ops/compile_boundary.py`	https://github.com/tile-ai/TileOPs	与 `torch.compile` 的接缝约束：被追踪的 `forward` 不能构造 kernel（:15-17）；key 必须是 `str`（:19-22）；Inductor 会把 fake 输出形状烤进磁盘缓存（:24-25）	已核验（2026-10-01）：读取 :1-30
R5	`TileOPs/docs/design/layer-boundaries.md`	https://github.com/tile-ai/TileOPs	「PyTorch is the spec oracle, not a runtime path」（:47）：算子层不允许在 forward 里直接调用高层 PyTorch 算子	已核验（2026-10-01）
R6	`TileOPs/src/tileops/perf/formulas.py`	https://github.com/tile-ai/TileOPs	逐元素与复合算子的 FLOP 计数惯例（如 `sigmoid = 4`、`silu = 5`）、attention 每分数 5 FLOP	已核验（2026-10-01）：读取模块头部常量与注释
R7	`hiascend.com · canncommercial_900_API_aolapi_context_ops-nn_aclnnApplyAdamW_md`	https://www.hiascend.com/document/detail/zh/canncommercial/latest/API/aolapi/context/ops-nn/aclnnApplyAdamW.md	优化器阶段：`varRef`/`mRef`/`vRef` 三输入输出 + `grad` + 超参；AdamW 公式（:25-60）	已核验（2026-10-01）：CANN 官方算子文档镜像
R8	`hiascend.com · canncommercial_900_API_aolapi_context_ops-transformer_aclnnKvRmsNormRopeCache_md`	https://www.hiascend.com/document/detail/zh/canncommercial/latest/API/aolapi/context/ops-transformer/aclnnKvRmsNormRopeCache.md	融合算子：切分尾轴做 RMSNorm + RoPE，再 scatter 写入两块 KV cache（:25）	已核验（2026-10-01）
R9	`hiascend.com · canncommercial_900_API_aolapi_context_ops-nn_aclnnAddRmsNormDynamicQuant_md`	https://www.hiascend.com/document/detail/zh/canncommercial/latest/API/aolapi/context/ops-nn/aclnnAddRmsNormDynamicQuant.md	融合的动机原文：「…融合起来，减少搬入搬出操作」（:25）	已核验（2026-10-01）
R10	`FlashAttention: Fast and Memory-Efficient Exact Att…`	https://arxiv.org/abs/2205.14135	FlashAttention：IO-aware 的论证；BERT-large 15%、GPT-2 3×、LRA 2.4×	已核验（2026-10-01）：本地为摘要全文；数字来自论文自述，条件见论文正文，本章未把它当作 TileLang 的性能
R11	`FlashAttention-3: Fast and Accurate Attention with …`	https://arxiv.org/abs/2407.08608	FlashAttention-3：FA2 在 H100 上 35% 利用率；FA3 FP16 达 740 TFLOPs/s（75%）、FP8 接近 1.2 PFLOPs/s	已核验（2026-10-01）：本地为摘要；论文自述数字
R12	`Correct but Slow: An Empirical Study of the GPU Ker…`	https://arxiv.org/abs/2607.04454	「Correct but Slow」：22 个 Triton/TileLang kernel 在 A100 与 GH200 上；一个 LayerNorm 比 PyTorch 基线慢 300 倍以上；库相对效率 + roofline 利用率作为筛选判据	已核验（2026-10-01）：本地为摘要全文
R13	`Native Sparse Attention: Hardware-Aligned and Nativ…`	https://arxiv.org/abs/2502.11089	NSA：「arithmetic intensity-balanced algorithm design」的原文表述	已核验（2026-10-01）
R14	`SeerAttention-R: Sparse Attention Adaptation for Lo…`	https://arxiv.org/abs/2506.08889	SeerAttention-R：「Using TileLang, we develop a highly optimized sparse decoding kernel…」（TileLang 作为研究算子的实现载体）	已核验（2026-10-01）
R15	`TileLang 知识体系与学习路径（编者笔记）第 44–48 行`	（本书 corpus 内的合成文档，无独立上游链接）	TileLang 的定位：「位置在 Inductor 之下、CUDA/Ascend C 之上」	已核验（2026-10-01）

上游链接核验说明

corpus 里的 torch.compile documentation 与 PyTorch Inductor: generated kernels 只有 11 行 / 306 字节：docs.pytorch.org 的 stable 索引页返回的是一张 302 跳转页（指向文档改版后的新路径），抓取工具把跳转页存了下来，所以 corpus 里没有可用正文。本章需要 TorchDynamo / TorchInductor / AOTAutograd / kernel 命名这些内容，因此在写作时（2026-10-01）直接核验了下面四个上游页面，并逐句引用了其中的原文：

id	URL	用途	核验状态
U1	https://docs.pytorch.org/docs/2.14/user_guide/torch_compiler/torch.compiler.html	TorchDynamo（Frame Evaluation API 捕获图）、TorchInductor（默认后端，在 NVIDIA/AMD/Intel GPU 上以 Triton 为构建块）、AOTAutograd（把反向也提前捕获）	已核验（2026-10-01）：抓取全文并逐句核对原文
U2	https://docs.pytorch.org/docs/2.14/user_guide/torch_compiler/torch.compiler_get_started.html	eager 的逐元素算子「read → modify → write back」；「The single most important optimization that inductor performs is fusion」；`cos`→`sin` 从 2 读 2 写变成 1 读 1 写；`TORCH_COMPILE_DEBUG=1`	已核验（2026-10-01）
U3	https://docs.pytorch.org/docs/2.14/user_guide/torch_compiler/torch.compiler_inductor_profiling.html	kernel 命名 `triton_poi_fused_<aten算子>_<序号>`（`poi` = pointwise）；前向图与反向图各一个编译产物；densenet121 小 batch 时「GPU busy」只有 32.69%	已核验（2026-10-01）
U4	https://docs.pytorch.org/docs/2.14/user_guide/torch_compiler/torch.compiler_faq.html	加速手段（垂直融合、水平融合、乱序执行、自动工作放置）；「Inductor first takes care of fusing whatever it can and only then generates Triton kernels」；graph break 与「We cannot propagate gradients through a graph break」	已核验（2026-10-01）

注意：这些页面是会被上游改版的在线文档，且本机没有 GPU、没有在该版本 PyTorch 上复现过任何行为。它们在本章只用于说明「torch.compile 的设计意图与接口边界」，不用于任何性能结论。若将来 corpus 重新抓取到可用正文，应把上表替换成 corpus 路径。

本章出现的每一个数字（来源与条件）

数字	出现在	来源	测量/推导条件
H100：over 3 TB/s、FP16 1000 TFLOPS	2.2	R1	厂商公开资料，H100 SXM5，Table 1 自述 preliminary，未实测
H100 平衡点 ≈ 333 FLOP/Byte	2.2	本书推导（1000/3.0）	用 R1 的两个数字做比值
H200：989.5 TFLOPS、4800 GB/s、标定 0.7370 / 0.6684 / 0.9380 / 0.8183	2.2	R2	仓库自带微基准标定，条件见该文件注释；不是本章实测
H200 平衡点 162（标定后）	2.2、2.3	本书推导	729.3 TFLOPS / 4502 GB/s
逐元素链 eager 134.2 MB / 融合 67.1 MB	2.3、4.2	本书推导（`op_traffic.py`）	16.8M 个 fp16 元素，2 算子链
`cos`→`sin` 的 2 读 2 写 → 1 读 1 写	4.2、图 3、图 5	U2	PyTorch 官方文档举例
LayerNorm 134.2 MB / 268.4 MB	2.3、4.2	本书推导（`op_traffic.py`）	8192×4096 fp16，每元素 8 FLOP（本书口径）
GEMM 4096³：137 GFLOP / 100.7 MB / AI 1365	2.3	本书推导（`op_traffic.py`）	fp16，A/B/C 各读或写一遍
attention prefill（S=4096）：42.2 与 2068 FLOP/Byte	2.3	本书推导（`op_traffic.py`）	32 head、d=128、fp16；分数矩阵 6 趟访存是假设
attention decode（KV 32768）：AI 1.01	2.3	本书推导（`op_traffic.py`）	每 token 读整段 KV
FA2 在 H100 上 35% 利用率；FA3 FP16 740 TFLOPs/s（75%）、FP8 接近 1.2 PFLOPs/s	2.3	R11	论文自述，测量条件见论文正文
FlashAttention 训练加速 15% / 3× / 2.4×	2.3	R10	论文自述（BERT-large / GPT-2 / LRA），条件见论文正文
优化器 26 字节/参数、7B 参数 182 GB、8B 参数 208.78 GB	1.1、2.4	本书推导（`op_traffic.py`、`step_budget.py`）	bf16 梯度 + fp32 状态；张量与公式来自 R7
all-reduce 搬 $2 (N - 1) / N$ 倍梯度、8 卡 8B 模型 28.11 GB	1.1、2.4	本书推导	ring all-reduce 的标准数据量；无实测
训练一步 / 推理的占比表（27.9% / 55.9% / 14.3% / 1.9% 等）	2.4、习题 E12	本书推导（`step_budget.py`）	粗模型，假设写在脚本注释与 `资料库索引`
densenet121 小 batch「GPU busy」32.69%	3.1	U3	该页所引的 benchmark 脚本，非本章实测
「正确但慢 300 倍」的 LayerNorm	5	R12	22 个 kernel、A100 与 GH200，论文自述

没有来源的、或来自厂商宣传口径的数字一律未采用。 所有需要真实加速器才能得到的数字（实测算力利用率、延迟、吞吐）在本章中均未出现，或明确标注为论文自述。

待补资料

corpus 里的 torch.compile 资料不可用（优先级：高）。 torch.compile documentation 与 PyTorch Inductor: generated kernels 抓到的都是 docs.pytorch.org 的 302 跳转页（分别 11 行 / 306 字节），没有正文。本章 4.3、4.4 节因此只能引用在线官方文档（见「上游链接核验说明」）。建议后续版本把 sources.yaml 里的 URL 改成不跳转的新路径并重抓，例如 https://docs.pytorch.org/docs/2.14/user_guide/torch_compiler/torch.compiler.html、 .../torch.compiler_get_started.html、.../torch.compiler_inductor_profiling.html、 .../torch.compiler_faq.html。
roofline 模型的原始论文缺失（优先级：中）。 本章用的是 TileLang 生态里 TileOPs 的记账口径（R3），原始出处是 Williams、Waterman、Patterson 2009 年的 "Roofline: An Insightful Visual Performance Model for Multicore Architectures" （CACM；DOI 10.1145/1498765.1498785）。corpus 里没有这条来源，第 28 章若要讲性能模型史，建议补抓。
训练侧并行与显存的一手资料缺失（优先级：中）。 本章 1.1 节关于优化器状态占多少字节、ZeRO 如何分片，只有算术推导（本书口径）；corpus 里没有 ZeRO（arXiv:1910.02054）或 Megatron-LM 的一手材料。第 17 章（端到端训练）大概率需要它们。
「反向 ≈ 2 倍前向」缺少官方表述（优先级：低）。 本章的做法是对单个 GEMM 逐项推导（ $d A$ 、 $d B$ 各 $2 M N K$ ），结论精确；但没有引用官方文档的同类表述。若第 17 章要写训练成本预算，建议补一条 PyTorch 或 Megatron 的说明。
Triton 文档快照过短（优先级：低）。 Triton documentation 只有 41 行（首页快照），不足以支撑「Triton 支持哪些后端目标」这类结论。本章 4.4(c) 因此写成了「本书基于官方文档措辞的判断」而不是事实断言。

本章共 9 个二级小节；右上角目录在桌面端会跟随滚动高亮。

从 PyTorch 到 kernel：算子在栈里的位置

学习目标

前置知识

1. 问题的由来：算子在一次迭代里出现在哪些位置

1.1 训练一步：四个阶段，四种瓶颈

1.2 一次推理请求：prefill 与 decode 是两件事

1.3 所有算子最终都变成 kernel

2. 概念与模型：一个算子要读多少字节、算多少次

2.1 记账法：四条规则

2.2 机器平衡点：把「算力」和「带宽」放在同一把尺子上

2.3 四个可以手算的例子

2.4 把算子拼成模型：两张账本

2.5 这个模型的边界

3. 用法：一行 PyTorch 到底做了什么

3.1 eager 模式：一行代码 = 一串 kernel launch

3.2 你会看到什么：kernel 表

3.3 用本章的脚本复算

4. 实现机制：融合是默认答案，torch.compile 做到哪一步

4.1 逐算子执行的两个代价

4.2 融合的账本与厂商证据

4.3 torch.compile / TorchInductor 能做到什么

4.4 边界在哪

5. 进阶与陷阱：什么时候你需要自己写 kernel

5.1 本章的三条常见陷阱

5.2 本章不做什么

6. 本章小结

延伸阅读

习题

计算题

代码阅读题

动手题

章节自测

本章代码

本章引用

第 02 章参考资料

锁定的源码引用（`_vendor/tilelang-v0.1.15/`）

corpus 资料引用（`资料库索引`）

上游链接核验说明

本章出现的每一个数字（来源与条件）

待补资料

学习目标

前置知识

1. 问题的由来：算子在一次迭代里出现在哪些位置

1.1 训练一步：四个阶段，四种瓶颈

1.2 一次推理请求：prefill 与 decode 是两件事

1.3 所有算子最终都变成 kernel

2. 概念与模型：一个算子要读多少字节、算多少次

2.1 记账法：四条规则

2.2 机器平衡点：把「算力」和「带宽」放在同一把尺子上

2.3 四个可以手算的例子

2.4 把算子拼成模型：两张账本

2.5 这个模型的边界

3. 用法：一行 PyTorch 到底做了什么

3.1 eager 模式：一行代码 = 一串 kernel launch

3.2 你会看到什么：kernel 表

3.3 用本章的脚本复算

4. 实现机制：融合是默认答案，torch.compile 做到哪一步

4.1 逐算子执行的两个代价

4.2 融合的账本与厂商证据

4.3 torch.compile / TorchInductor 能做到什么

4.4 边界在哪

5. 进阶与陷阱：什么时候你需要自己写 kernel

5.1 本章的三条常见陷阱

5.2 本章不做什么

6. 本章小结

延伸阅读

习题

计算题

代码阅读题

动手题

章节自测

本章代码

本章引用

第 02 章 参考资料

锁定的源码引用（_vendor/tilelang-v0.1.15/）

corpus 资料引用（资料库索引）

上游链接核验说明

本章出现的每一个数字（来源与条件）

待补资料

第 02 章参考资料

锁定的源码引用（`_vendor/tilelang-v0.1.15/`）

corpus 资料引用（`资料库索引`）