Lecture 9 · 硬件性能建模：循环嵌套、数据流与 Interstellar · Modeling Hardware Performance

对应 Maps to：Stanford CS349E / EE339 Efficient ML Inference at Scale（Autumn 2026），原课第 7 次课 · Tue Oct 13 · Modelling Hardware Performance（Priyanka Raina）。本讲义把它放在”硬件引擎”（原第 8 次课，本讲义第 8 讲）之后：先认识硬件，再给它建模（见 SYLLABUS §0.1）。阅读材料：X. Yang, M. Gao, Q. Liu, J. Setter, J. Pu, A. Nayak, S. Bell, K. Cao, H. Ha, P. Raina, C. Kozyrakis, M. Horowitz, Interstellar: Using Halide’s Scheduling Language to Analyze DNN Accelerators（ASPLOS 2020）；AWS Neuron 文档 NKI Performance Guide（v2.26）。 依据说明：课程网站只公开了课表和阅读材料，没有 slides。本讲以两篇阅读材料为主线；Interstellar 的作者里有本课的两位老师（Raina、Horowitz）。另外参考了 Timeloop（Parashar et al. 2019）的映射分析思路。补充材料是我选的，不代表课堂上实际讲了什么。 配套代码 Code：code/lecture09/（numpy；一个共享模块 + 5 个脚本，最长的约 11 秒） 先修：前置讲 P4（循环嵌套、分块、“算法与调度分离”）、第 8 讲（脉动阵列、Tensor Core、存储层次）。 预计阅读时间：3 小时（含运行代码）

0. 本讲导航 · How to read this lecture

0.1 学习目标 Learning goals

学完这一讲，你应该能够：

1. 把卷积、全连接、Transformer 的线性层写成同一个循环嵌套（矩阵乘法 $C[M,N] \mathrel{+}= A[M,K] \cdot B[K,N]$），说出每个张量最多能被复用多少次。
2. 解释”映射”（mapping）的三个组成部分——分块、循环顺序、空间展开——以及资源分配；用”驻留规则”算出每个张量在每一级存储被装进来几次。
3. 用 Interstellar 的能量模型 $E = \sum_i \#\text{acc}_i \times e_i$ 估算一个映射的能量，并说出每一级存储每次访问的大致代价。
4. 复述 Interstellar 的两个观察：(1) 只要分块和复制用得好，不同数据流的能效相近——分块比数据流重要；(2) 好的设计里总能量不应该被某一级存储主导——每一级的大小要平衡。
5. 用 NKI 性能指南的几条规则估算真实芯片上的代价：指令固定开销、分区利用率、驻留 / moving 的选择、DMA 的传输大小、融合与片上容量。

0.2 本讲地图 Map

§1  为什么要建模：设计空间太大，不能每种都造一块芯片
§2  一切都是循环嵌套                                                   ← 代码 01 演示 1
§3  映射：分块、循环顺序、空间展开；驻留规则                             ← 代码 01
§4  数据流的分类：U|V、复制；Halide 的调度语言                           ← 代码 03 演示 3
§5  能量模型：每一级的访问次数 × 每次的能量                             ← 代码 02
§6  观察 1：数据流没那么重要，分块很重要                                 ← 代码 03
§7  观察 2：每一级存储该多大                                           ← 代码 02、04
§8  从模型到真实芯片：AWS NKI 性能指南                                  ← 代码 05
§9  建模的层次与陷阱
§10 现状（2025–2026）
§11–15  小结、常见误区、练习（含答案）、延伸阅读、下一讲预告

0.3 本讲的代码

loopnest.py 实现了一个和 Interstellar / Timeloop 同一思路的解析模型：给定矩阵乘法的形状、硬件（PE 阵列、寄存器堆 RF、全局缓冲 GB）和一个映射，算出每个张量在每一级的读写次数、PE 之间的传递次数，再乘以 Interstellar 表 3 的能量。它还能枚举全部合法的映射（几万到十几万个）找最优的，并用”逐次迭代、真的数块换了几次”的暴力计数验证解析公式。05_nki_perf_model.py 是一个按 NKI 指南写的 NeuronCore 代价模型。

1. 为什么要建模 · Why model hardware performance

第 8 讲看了三种矩阵引擎：TPU 的权重驻留脉动阵列、GPU 的 Tensor Core、SRAM 芯片。加速器的论文还有很多：DianNao 的内积单元、Eyeriss 的”行驻留”、NVDLA、各种 FPGA 设计……每篇都说自己比通用处理器好，可它们的结构差别很大。哪种最好？

Interstellar 的作者发现，各篇论文只探索了设计空间的一小块，还常常把几个独立的参数绑在一起改，所以”最优参数”的报告互相矛盾。要公平比较，需要：

1. 一个能描述所有这些设计的统一语言；
2. 一个能快速评估一个设计的模型（造芯片要几个月到几年；RTL 仿真一个设计要几小时）；
3. 这个模型要足够准（Interstellar 的解析模型和真正综合出来的电路误差小于 2%）。

建模的工具链从粗到细：roofline（第 1 讲，几秒钟、只看算力和带宽）→ 访问次数的解析模型（本讲，毫秒级评估一个映射，能搜索几十万个）→ 周期级模拟（第 8 讲的脉动阵列模拟、GPU 模拟器）→ RTL 仿真和综合（最准、最慢）。好的流程是先用粗模型大范围搜索，再用细模型验证少数候选。

2. 一切都是循环嵌套 · Everything is a loop nest

代码：01_loop_nests.py 演示 1

卷积层是 7 层循环（Interstellar 算法 1）：

for b in [0, B):            # 批
  for k in [0, K):          # 输出通道
    for c in [0, C):        # 输入通道
      for y in [0, Y):      # 输出高
        for x in [0, X):    # 输出宽
          for fy in [0, FY):    # 卷积核高
            for fx in [0, FX):  # 卷积核宽
              O[b][k][x][y] += I[b][c][x+fx][y+fy] * W[k][c][fx][fy]

全连接层只剩 B、K、C 三层（其余上界为 1）。Transformer 的线性层也一样。把卷积用 im2col（把每个输出位置要看的那一小块输入展开成一行）改写，所有这些都变成一个矩阵乘法 $C[M,N] \mathrel{+}= A[M,K] \cdot B[K,N]$。本讲的代码只处理矩阵乘法，讲到 Interstellar 的卷积时会说明对应关系。

矩阵乘法有 $MNK$ 次乘加，却只有 $MK + KN + MN$ 个数要搬：A 的每个数最多被用 $N$ 次，B 的每个数最多被用 $M$ 次，C 的每个数累加 $K$ 次。代码 01 演示 1：

“乘加 / 最少要搬的数”是复用的上限，由层的形状决定。卷积和 prefill 有上百到上千次；批 1 的全连接层只有 1 次——每个权重只被用一次，怎么安排都省不掉从 DRAM 读权重。这就是第 1 讲 roofline 的”算术强度”在循环层面的样子。能不能达到这个上限，由映射决定。

3. 映射：分块、循环顺序、空间展开 · Mappings

代码：01_loop_nests.py 演示 2–4、loopnest.py

3.1 三个维度的设计空间

Interstellar 把加速器的设计空间描述成三个正交的维度（论文图 1）：

1. 分块（loop blocking）：每个循环切成几段，最内的段在最小的存储（寄存器堆 RF）里转，往外依次是片上缓冲（GB）、DRAM。“切成几段、每段多大”决定了每一级存储里放多大的块。
2. 数据流（dataflow）：哪些循环在 PE 阵列上空间展开（spatial unrolling）——并行地放在不同的 PE 上。
3. 资源分配：PE 有多少、每一级存储多大。

（论文把循环顺序归在”分块”里：split 和 reorder 一起叫 loop blocking。）

在我们的代码里，一个映射就是：每一维 $d \in \{M, N, K\}$ 的大小拆成四个因子 $d = d_\text{rf} \times d_\text{sp} \times d_\text{gb} \times d_\text{dram}$（RF 里的时间循环、阵列上的空间展开、GB 层的时间循环、DRAM 层的时间循环），再加上 GB 层和 DRAM 层的循环顺序。

3.2 驻留规则：循环顺序决定谁不用重新装

假设每一级存储为每个张量放一个块。一个张量的块什么时候要重新装？看它上面的所有时间循环（从外到内）：

从最外层开始把上界乘起来，乘到”最内层的那个和这个张量有关、上界大于 1 的循环”为止。它里面那些和这个张量无关的循环转一圈，这个张量的块不变——不用重新装。这个张量在这一级”驻留”（stationary）。

代码 01 演示 2（$M = N = K = 8$，切成 4，外层每维循环 2 次，共 8 次迭代）：

所以 weight / output / input stationary 这些名字，本质上只是循环顺序的不同：最内层的循环和谁无关，谁就驻留。第 8 讲 TPU 的脉动阵列是权重驻留：权重装进 PE 以后，一批输入流过去。

代码 01 演示 4 和 loopnest.py 的自检把这条规则和”逐次迭代整个循环嵌套、记下每个张量的块下标什么时候变”的暴力计数逐个比较：28 个随机映射（含空间展开）全部一致。有了这条规则，访问次数不用模拟就能算出来，评估一个映射只要微秒级的时间。

3.3 分块大小决定 DRAM 流量

代码 01 演示 3（$256 \times 512 \times 512$，循环顺序 M → N → K，只看 GB 和 DRAM 两级）：

块越大，复用越多，DRAM 流量越接近”每个数只搬一次”；但要放得下。同样 64 KiB，块的形状不同，流量也差 1.5 倍。分块、顺序、每一级的容量要一起选——这就是映射的搜索空间。第 8 讲 GPU 的数据块、TPU 的 VMEM、NKI 的 SBUF 分块，都是在做这件事。

4. 数据流的分类 · A taxonomy of dataflows

4.1 U|V 记号

Interstellar 用”哪些循环展开到阵列的哪个方向”来精确地定义数据流：$U|V$ 表示循环 $U$ 展开到阵列的竖直方向、$V$ 展开到水平方向。以前论文里的名字都能写成这种形式（论文表 1）：

$C \mid K$（输入通道 | 输出通道）是最常见的一种：它也能直接映射 MLP 和 LSTM 的矩阵乘法，TPU 就是它。注意同样叫”权重驻留”，$C \mid K$ 和 $FX \mid FY$ 的数据复用方式完全不同——U|V 记号比名字精确。

4.2 复制：维度太小时把阵列填满

如果展开的维度比阵列小，一部分 PE 就空着。论文图 2：输入通道 $C = 3$ 展开到 16 行，只有 3 行在干活。复制（replication）：在同一个方向上再展开另一个循环（比如 $X$ 展开 5 份），$3 \times 5 = 15$ 行在干活，利用率从 3/16 回到 15/16。记号写成 $UW \mid V$。代码 03 演示 3 复现了这个例子（$K = 3$ 的矩阵乘法）：

复制把时间缩短了 5 倍，能量几乎不变——利用率影响的是吞吐，不是能量。复制要求阵列的互连足够灵活（不同”组”之间的数据要传更远，论文图 3），Eyeriss v2、CGRA、FPGA 支持。

4.3 Halide 的调度语言就是加速器的描述语言

Interstellar 标题里的 Halide 是一个图像处理的领域专用语言，核心思想是算法与调度分离（前置讲 P4 §9.1）：算法只说”算什么”，调度说”按什么顺序、在哪里算”。论文发现，描述加速器需要的变换几乎都是 Halide 调度原语的子集（论文表 2）：

一个卷积层的例子（论文清单 1）：split 把 x、y 切成 8×8 的块、reorder 调整顺序、in().compute_at(output, xo) 为每个块的输入和权重开缓冲、unroll(xi, 4).systolic({xi}) 把内层的 xi 展开成 4 个 PE 的脉动阵列。他们改了 Halide 编译器，让它从这样的调度直接生成硬件（经 Catapult 高层综合到 28 nm 的 RTL）。同一个算法，换一个调度就是另一块加速器——公平比较由此成为可能。第 12、13 讲（编译器）会回到”调度语言”这个想法。

5. 能量模型 · The energy model

代码：02_energy_model.py

5.1 公式

$e_i$ 是第 $i$ 级存储访问一次的能量；$\#\text{acc}_i$ 是它被访问的次数，由各级的复用次数 $RT_j$ 决定（数据在这一级存在的期间，被下一级（更便宜的那一级）访问了几次）。PE 之间的传递当作额外的一级，距离不同代价不同。$e_i$ 由资源分配决定；$RT_j$ 由数据流和分块决定（§3 的驻留规则）。所以整个设计空间的优化就是：在三个维度上最小化 $E$。

5.2 能量表

Interstellar 表 3（28 nm、每次 16 位访问；代码 02 演示 1）：

两条规律：存储越大每次访问越贵（寄存器堆和大小成正比，SRAM 每翻一倍 ×1.5）；DRAM 比一次乘加贵三个多数量级。和第 1 讲 Horowitz 的 45 nm 表比，存储和运算的比例更悬殊（工艺更新、存储分了大量 bank）。

5.3 一个映射的能量花在哪

代码 02 演示 2：一个小 Transformer 层的投影（256 个 token、512×512），在 Eyeriss 式的配置（16×16 PE、512 B 寄存器堆、128 KB SRAM）上搜索最好的映射：

• 每次乘加要从寄存器堆读 A、B、读写 C 共 4 次；512 B 的寄存器堆每次 0.96 pJ——寄存器堆一项占一半。
• DRAM 是第二大头：128 KB 的 SRAM 放不下足够大的块。
• 运算本身只占约 1%。和第 8 讲一样：能量几乎全花在搬数据上，而且最多的花在离运算最近、被访问次数最多的那一级。

5.4 模型的可信度

论文用 Halide → 高层综合 → 28 nm 综合出三个真实的设计验证解析模型，能量误差小于 2%（论文表 4、图 7），也能复现 Eyeriss 论文的结果。我们的实现是它的简化版：每一级为每个张量放一个块（另一份用于双缓冲只占容量）、输出的部分和第一次从 0 开始、多播和空间归约用 hop 计、不模拟带宽和停顿。所以它适合比较能量和访问次数，不能预测时间。

6. 观察 1：数据流没那么重要，分块很重要 · Observation 1

代码：03_dataflow_vs_blocking.py

6.1 每种数据流配上最好的分块

代码 03 演示 1（16×16 PE、64 B 寄存器堆、128 KB SRAM；每种数据流都搜索它自己最好的分块和循环顺序）：

（pJ/乘加。）最好和最差的数据流只差 1%–2%。这就是 Interstellar 的观察 1：

同样的硬件资源下，只要用了合适的分块和复制，很多不同的数据流都能达到相近的、接近最优的能效。

论文在 AlexNet、GoogLeNet 的各层、三种硬件配置、批大小 1 和 16 上都看到同样的结论（图 8）。事后看并不奇怪：卷积和矩阵乘法里每个数都有很大的复用潜力，好的分块让绝大多数访问发生在最小的那一级（论文：最优分块下 98% 的访问发生在寄存器堆）；数据在 PE 之间怎么传，只占总能量的很小一部分（演示 2 里 array 一项不到 1%）。复用很少的层（批 1 的全连接），数据必须从 DRAM 来，片上的数据流也影响不大。

6.2 分块之间差别很大

代码 03 演示 2：同一种数据流（$K \mid N$）、同一块硬件，113,796 个合法映射（分块 × 两级循环顺序）的能量分布：

只有 0.2% 的映射在最低能量的 1.25 倍以内（论文图 10：AlexNet CONV3 上是 30%；我们的搜索空间包括很多糟糕的循环顺序，分布更宽）。随手挑一个映射，比最好的差好几倍。所以决定能效的是分块和循环顺序，不是数据流的名字——加速器离不开一个好的映射搜索 / 编译器。

6.3 数据流影响的是利用率

§4.2 已经看到：数据流和复制决定了 PE 阵列能不能填满，也就是吞吐。论文图 9：不支持复制时各种数据流的利用率差别很大；支持复制以后差别大多消失。论文最后选 $C \mid K$，理由是通道维通常最大、最容易填满固定大小的阵列。

7. 观察 2：每一级存储该多大 · Observation 2

代码：02_energy_model.py 演示 3–4、04_memory_hierarchy.py

7.1 缩小寄存器堆

既然数据流不太重要，硬件设计者能决定的就是资源分配。§5.3 的分解显示 512 B 的寄存器堆独占一半能量——每次乘加访问它 4 次，它每次太贵了。代码 04 演示 1 把寄存器堆和全局缓冲的大小一起扫一遍（一个 $d = 512$ 的小 Transformer 层在 256 个 token 上的全部线性层，pJ/乘加）：

• 寄存器堆从 512 B 缩到 32–64 B，能量降约 40%（Eyeriss 式的 8.82 → 5.00）；论文在 AlexNet 上测到最多 2.6 倍（图 12）。
• 但不能太小：16 B 装不下足够大的块，比 32 B 还差——存在最优点。
• 全局缓冲越大越好、收益递减：256 KB → 512 KB 只再降约 6%。论文：“超过 256 KB 以后收益可以忽略”。
• 最好的配置（64 B、512 KB）比 Eyeriss 式的好 2.04 倍。

7.2 能量应该平衡

代码 04 演示 2：

寄存器堆缩小以后不再是大头，DRAM 变成了主导；再加大全局缓冲，DRAM 降下来，各级更均衡。这就是论文的观察 2：

一个高效系统的总能量不应该被存储层次中的某一级主导。

某一级太大太贵，就把它缩小、把复用推到别的级去；某一级访问太多，就把它上面那一级加大。这也解释了为什么一些输出驻留、权重驻留的设计表现不好：它们在寄存器堆这一级抓不住足够的复用，结果 DRAM 主导了能量。

例外：如果 DRAM 主导、但 DRAM 的访问次数已经是最少的（每个输入读一次、每个输出写一次），系统就是受限于内存的——Amdahl 定律说，别的地方再优化也没用。代码 02 演示 4：批 1 的 4096×4096 全连接层，DRAM 约 200.8 pJ/乘加，下限是 200.1——几乎全是躲不掉的”每个权重读一次”。论文专门指出批 1、MLP、LSTM 属于这一类。这和 TPU 论文（第 8 讲 §4）、LLM 的 decode（第 1 讲 §4.2）是同一个结论。

7.3 相邻两级的比例，以及阵列变大时

论文还发现：在原来的层次里再加一级小寄存器堆（16 B → 256 B → 256 KB），AlexNet 又省 25%；经验规则是相邻两级存储的容量之比在 4–16 之间时最好。原因是每一级既要大到能抓住足够的复用，又要小到每次访问便宜；比例太大，下一级被访问得太多。

代码 04 演示 3（阵列大小变化时的最优配置）：

PE 多了 16 倍，最优的全局缓冲只大了 2 倍（32×32 碰到了扫描上限，真正的最优可能更大）：存储随 PE 数次线性增长——每一级越大每次越贵，不值得为每个 PE 保留同样多的数据（论文图 13）。每个 PE 的寄存器堆大致不变。能量随阵列变大而下降（更多 PE 共享同一份数据）；论文说下降得”很少”，我们的模型降得多一些（约三分之一），因为它的 hop 很便宜、也没算大阵列更长的导线。我们只有两级片上存储，比值落在 16–64，比论文的经验规则偏大。

7.4 一个快速的优化器

利用这两个观察，Interstellar 把搜索剪枝：固定数据流为 $C \mid K$（观察 1），只评估满足”相邻两级比例 4–16”的存储配置（观察 2）。在四个 CNN、三个 LSTM、两个 MLP 上，优化后的设计比原始的 Eyeriss 配置能效高 4.2 倍（CNN）、1.6 倍（LSTM）、1.8 倍（MLP）。LSTM 和 MLP 提升得少，正是因为它们受限于 DRAM。

8. 从模型到真实芯片：AWS NKI 性能指南 · The NKI performance guide

代码：05_nki_perf_model.py

8.1 NeuronCore 是什么样的

AWS 的 Trainium / Inferentia2 芯片由 NeuronCore 组成。每个 NeuronCore 有四个计算引擎和两级软件管理的片上存储：

• Tensor 引擎（TensorE）：128×128 的脉动阵列，做矩阵乘法，占 NeuronCore 九成以上的 FLOPs。一次矩阵乘法分两步：LoadStationary 把一个张量装进阵列（驻留），MultiplyMoving 让另一个张量流过去——就是第 8 讲的权重驻留阵列。
• Vector 引擎（VectorE）、Scalar 引擎（ScalarE）、GpSimd 引擎：逐元素运算、归约、激活函数、掩码等。
• SBUF：片上的主缓冲（按 128 个分区组织，每个分区对应计算引擎的一条并行通道）；PSUM：矩阵乘法的累加缓冲。
• DMA 引擎（16 个）：在设备内存（HBM）和 SBUF 之间搬数据。

NKI（Neuron Kernel Interface）让程序员直接写这一层：用 nl.load / nl.store 搬数据、用 nisa.nc_matmul 等指令调用引擎。这和 Interstellar 的世界很像：存储是软件管理的，映射由程序员（或编译器）决定。

指南给的目标：受限于计算时，瓶颈引擎 90% 以上的时间在忙；受限于内存时，设备内存带宽利用率（MBU）60% 以上就算好。十条优化分三类：提高算术强度（Opt #1–2）、提高计算效率（#3–8）、提高数据搬运效率（#9–10）。

8.2 提高算术强度：别重复加载、融合以避免溢出

• Opt #1 利用时间局部性：同一个输入被反复从 HBM 加载（profiler 会标出来），就调整循环顺序让它留在 SBUF 里——正是 §3 的驻留规则。但留得太多会增加 SBUF 的压力、导致溢出，要权衡。
• Opt #2 融合：一串算子分开做时，中间结果放不下 SBUF 就要溢出（spill）到 HBM 再读回来。把它们融合成一个循环、一块一块地做。

代码 05 演示 5：一个注意力头（$d = 128$），分成三个内核写时 S×S 的分数矩阵要在 HBM 进出四次，算术强度只有约 60 FLOP/字节，在我们假设的 NeuronCore（拐点约 114）上只能跑到峰值的一半左右；融合成一个内核，HBM 流量只剩 Q、K、V、O，强度随序列长度线性增长（$S = 4096$ 时 2048 FLOP/字节）。$S = 4096$ 时分数矩阵有 32 MiB，放不下 24 MiB 的 SBUF——融合的内核必须分块、只在片上保留分数矩阵的一小块。这就是 FlashAttention（第 11 讲）。

8.3 提高计算效率

• Opt #3、#4 流水：引擎之间、DMA 和计算之间按块流水，别让一个引擎等另一个（第 8 讲 §8.4 的多级流水线）。注意力里矩阵乘法在 TensorE、指数在 ScalarE、求和在 VectorE、掩码在 GpSimdE，要排成一条流水线。
• Opt #5a 自由维要够大：每条指令约 100 个周期的固定开销。代码 05 演示 1（处理 128 × 4096 个元素）：

指南的 Mamba 例子：前缀和写成 4096 条单元素指令（每条间隔 189 ns）要约 774 µs，换成一条 tensor_tensor_scan 指令只要约 3 µs。指南的规则：每个分区至少 128 个元素。

• Opt #5b 分区要用满：只用 64 个分区的指令和用 128 个一样慢。把两个 64 分区的结果写到 PSUM 的不同分区、用一条 128 分区的指令处理，快 2 倍（代码 05 演示 2）——和 Interstellar 的”复制”是同一个问题。
• Opt #6 合并指令：Scalar 引擎一条 activation 指令能同时做乘、加、exp（快 3 倍），因为数据只读一次、在引擎内部流水。
• Opt #7 矩阵 × 向量时，短的做 moving：decode 时激活只有 1 列。代码 05 演示 3（M 个 token 乘一块 128 × 4096 的权重）：

（周期。）权重用”快速装载”（比 MultiplyMoving 快 4 倍）做驻留、短的激活做 moving，MultiplyMoving 只付 64 个周期的最小开销；反过来，权重要一列一列流过阵列，慢 2 倍。token 多到 128 以上两者一样——资源用满时”数据流”无所谓（§6）。但 $M = 1$ 时不管怎么映射，TensorE 的利用率只有约 1.6%：decode 的矩阵 × 向量本来就受限于内存。

• Opt #8 减少转置：TensorE 的时间里有多少花在了”交换分区维和自由维”的转置上；调整驻留 / moving 的选择、换一个引擎做，能省掉一部分。

8.4 提高数据搬运效率

• Opt #9 每次 DMA 要够大：每次传输有固定开销。指南：Trainium / Inferentia2 上每次传输要大于 32 KiB 才能达到理想带宽。代码 05 演示 4（假设每次开销 0.33 µs、每个引擎 25 GB/s）：4 B 一次几乎全是开销，4 KiB 约 33%，32 KiB 约 80%，128 KiB 约 94%。
• Opt #10 少用 DMA 转置：搬运时顺便转置会大大降低 DMA 带宽；受限于内存时，宁可正常加载再在 TensorE 上转置。

这些规则和第 8 讲 GPU 的那些（大块的 TMA、合并访存、足够多的字节在路上）本质相同：固定开销要被足够大的工作量摊薄；片上容量里要把复用做大。

9. 建模的层次与陷阱 · Levels of modeling and their pitfalls

几个常见的陷阱：

• 模型要验证：Interstellar 用综合出来的电路验证（误差 < 2%）；TPU 论文的性能模型和硬件计数器差 8%（第 8 讲 §4.5）。没有验证的模型只能看趋势。
• 模型的假设决定结论：我们的模型不计导线长度，所以”阵列越大越省能量”比论文更明显（§7.3）；不计带宽，所以看不出”利用率高但等数据”的情况。
• 测量要靠计数器和 profiler：NKI 指南的每一条优化都从 neuron-profile 里的一个”症状”开始（重复加载的警告、溢出的字节数、引擎空闲时在等哪个信号量）。TPU 论文的”陷阱”之一就是”性能计数器不能事后再加”。
• 先找瓶颈：指南的流程是先判断受限于计算还是内存，再针对瓶颈优化——优化不是瓶颈的引擎没有意义（Amdahl 定律）。

10. 现状（2025–2026）· State of the field

• 映射分析工具：Timeloop / Accelergy（NVIDIA 与 MIT）、MAESTRO、ZigZag、Interstellar 的调度器等，已经成为加速器研究的标准工具；新的工作把它们扩展到稀疏、融合的多层（层间流水）、Transformer 的注意力和 KV cache。
• 编译器里的代价模型：TVM / Ansor、Triton 的自动调优、XLA 都用代价模型挑选分块和循环顺序；越来越多地用学习到的代价模型（第 13 讲）。
• 厂商的”程序员可见的映射”：NKI、CUTLASS / CuTe（GPU 的分块抽象）、Pallas（TPU / GPU 的内核语言）把”分块、驻留、流水”直接交给内核作者；配套的 profiler（neuron-profile、Nsight Compute）就是给内核作者的性能模型。
• AI 帮忙搜索：映射空间巨大而评估便宜，正适合搜索和学习。第 14 讲的 AlphaEvolve、KernelBench 就是在这类空间里让 LLM 提方案、让评估器打分。

11. 小结 · Summary

1. 卷积、全连接、Transformer 的线性层都是同一个循环嵌套；复用的上限由形状决定（批 1 的全连接只有 1 次，prefill 上千次）。
2. 映射 = 分块 + 循环顺序 + 空间展开（+ 资源分配）。驻留规则：最内层的相关循环之外的部分决定一个张量被装几次；weight / output / input stationary 只是循环顺序的不同。
3. 能量 = Σ 每一级的访问次数 × 每次的能量；存储越大越贵，DRAM 是乘加的 2,667 倍。好的映射里运算本身只占约 1%。
4. 观察 1：每种数据流配上最好的分块，能量只差百分之几；分块之间差几倍到几十倍——分块比数据流重要。数据流和复制决定的是利用率（吞吐）。
5. 观察 2：总能量不应该被某一级主导。缩小寄存器堆（512 B → 32–64 B）、加大全局缓冲到约 256–512 KB，能量好 2 倍左右；相邻两级容量之比约 4–16；存储随 PE 数次线性增长。例外：DRAM 主导而且已经到下限时（批 1），系统受限于内存。
6. 在真实芯片上（NKI）：指令固定开销要用够大的自由维摊薄、分区要用满、decode 时短的做 moving、DMA 每次至少 32 KiB、融合以免溢出——都是”摊薄固定开销、在有限容量里把复用做大”。

12. 常见误区 · Common pitfalls

1. “我们的数据流比别人的好”——同样的资源下，配上各自最好的分块，数据流之间只差百分之几（§6.1）；比较时如果分块不同，结论就是在比分块。
2. “寄存器堆越大越好，能抓住更多复用”——每次乘加都要访问它好几次，它越大每次越贵；最优点在 32–64 B（§7.1）。
3. “片上缓冲越大越好”——收益递减，还要花面积；超过 256 KB 收益很小（§7.1）。
4. “能量低就快”——能量模型不管时间。复制把利用率从 19% 提到 94%，能量几乎不变、时间少了 5 倍（§4.2）。
5. “weight / output stationary 是硬件结构”——它们只是循环顺序；同一块硬件可以用不同的顺序（§3.2）。
6. “优化片上数据流能救批 1 的推理”——批 1 时 DRAM 已经到下限，受限于内存，片上怎么排都差不多（§7.2）。
7. “一条指令只处理几个元素也没关系，反正总量一样”——每条指令约 100 个周期的固定开销，单元素指令效率只有 1%（§8.3）。
8. “解析模型算出来的就是真实性能”——它不计带宽争用、停顿和控制开销；必须用综合结果或 profiler 验证（§9）。

13. 练习 · Exercises

练习 1（驻留规则） 矩阵乘法 $C[M,N] \mathrel{+}= A[M,K] \cdot B[K,N]$，外层循环从外到内是 N（上界 4）、K（2）、M（3）。A、B、C 的块各装几次？谁驻留？

答案：A（M、K）：最内层的相关循环是 M（第 3 个），$4 \times 2 \times 3 = 24$ 次。B（K、N）：最内层的相关循环是 K（第 2 个），$4 \times 2 = 8$ 次。C（M、N）：最内层的相关循环是 M，24 次。B 驻留（权重驻留）——最内层的 M 和 B 无关。

练习 2（DRAM 流量） $M = N = K = 512$，GB 里的块 128×128×128，DRAM 层循环 M → N → K。(a) A、B 各从 DRAM 读多少个数？(b) C 读写多少？(c) 和下限比是几倍？(d) 三个块（16 位）占多少字节？

答案：(a) A 的最内层相关循环是 K（最内层），64 个版本 × 16,384 = 1,048,576；B 同样 1,048,576。(b) C 的最内层相关循环是 N，16 个版本 × 16,384 = 262,144 次写；K 在最内层，部分和一直留在 GB 里，不用读回。(c) 共 2,359,296，下限 $3 \times 512^2 = 786{,}432$，3 倍。(d) $3 \times 128^2 \times 2 = 96$ KiB。（loopnest.access_counts 算出同样的数。）

练习 3（寄存器堆的能量下限） 每次乘加访问寄存器堆 4 次。512 B 和 64 B 的寄存器堆，光这一项每次乘加多少能量？是一次乘加本身的几倍？

答案：512 B：$4 \times 0.96 = 3.84$ pJ，51 倍；64 B：$4 \times 0.12 = 0.48$ pJ，6.4 倍。所以缩小寄存器堆是最直接的省能量办法（§7.1）。

练习 4（批 1 的 DRAM 能量） 4096×4096 的全连接层（16 位权重），每次 DRAM 访问 200 pJ。批 1 时每次乘加至少花多少 DRAM 能量？批 16 呢？

答案：批 1：每个权重只用一次，每次乘加至少读 1 个权重，≥ 200 pJ（代码 02 演示 4：200.8，下限 200.1，多出来的是激活和输出）。批 16：每个权重用 16 次，≥ 12.5 pJ（加上激活约 12.6）。

练习 5（复制） 16×16 的阵列，$K = 3$ 展开到行、$N = 64$ 展开到列。利用率多少？再把 $M$ 展开 5 份到行上呢？

答案：$3/16 \times 16/16 = 19\%$；$15/16 = 94\%$。能量几乎不变，时间少约 5 倍（代码 03 演示 3）。

练习 6（SRAM 的能量外推） Interstellar 表 3：32 KB 的 SRAM 每次 6 pJ，大小每翻一倍 ×1.5。1 MB 的 SRAM 每次多少？

答案：1 MB = 32 KB × $2^5$，$6 \times 1.5^5 \approx 45.6$ pJ。

练习 7（指令固定开销） 每条指令固定开销 100 个周期、之后每周期每个分区处理 1 个元素。自由维多大时效率达到 90%？

答案：$F / (100 + F) \ge 0.9 \Rightarrow F \ge 900$。指南的”至少 128”只够约 56%；要接近满负荷需要上千。

练习 8（短的做 moving） 一块 128 × 4096 的权重乘 1 个 token。按代码 05 的模型：权重做驻留时，要装几次驻留张量、每次多少周期？反过来呢？

答案：权重驻留：$4096 / 128 = 32$ 次装载，每次 $\max(128/4, \max(1, 64)) = 64$ 周期，共 2,048；激活驻留：1 次装载、权重分 $4096 / 512 = 8$ 段流过，每段 512 周期，共 4,096。慢 2 倍。

练习 9（DMA 效率） 每次 DMA 固定开销 0.33 µs、每个引擎 25 GB/s。一次传 32 KiB 的效率多少？要达到 95% 每次要多大？

答案：32 KiB 传 $32768 / 25\text{e}9 = 1.31\ \mu s$，效率 $1.31 / (1.31 + 0.33) = 80\%$。95%：$x / 25\text{e}9 \ge 19 \times 0.33\ \mu s$，$x \ge 157$ KB（约 153 KiB）。

练习 10（融合的算术强度） 一个注意力头，$S = 4096$，$d = 128$，BF16。(a) 运算量 $4 S^2 d$？(b) 不融合时 HBM 流量（Q、K、V、O 各一次 + 分数矩阵进出四次）？(c) 两种情况的算术强度？

答案：(a) $4 \times 4096^2 \times 128 \approx 8.59$ GFLOP。(b) $4 \times 4096 \times 128 \times 2 = 4.19$ MB；分数 $4 \times 4096^2 \times 2 = 134$ MB；共 138.4 MB。(c) 不融合 62 FLOP/字节；融合（只有 4.19 MB）2048 FLOP/字节。

练习 11（为什么数据流不重要） 用 §5 的能量模型解释：为什么在卷积层上，换一种数据流（配上最好的分块）能量几乎不变，换一种分块却能差几倍？

答案：能量 = Σ 访问次数 × 每次能量。卷积 / 大矩阵乘法的复用潜力很大，好的分块能让绝大多数访问发生在最便宜的那一级（寄存器堆），DRAM 和缓冲的访问很少。数据流只改变 PE 之间怎么传数据——这一级（hop）很便宜，而且占的比例很小。分块却直接决定了数据在哪一级被复用、贵的那几级被访问多少次。

练习 12（相邻两级的比例） 为什么相邻两级存储的容量之比不能太大也不能太小？

答案：下一级（大的）每次访问的能量随大小增长；上一级（小的）每次便宜但装不下太多。比例太大：小的那一级抓不住多少复用，大量访问落到大而贵的那一级；比例太小：两级的能量差不多，多一级反而多一次搬运。Interstellar 的经验是 4–16。

14. 延伸阅读 · Further reading

阅读材料

• X. Yang et al., Interstellar: Using Halide’s Scheduling Language to Analyze DNN Accelerators, ASPLOS 2020. https://arxiv.org/abs/1809.04070 ；分析模型的代码：https://github.com/xuanyoya/Interstellar-CNN-scheduler
• AWS Neuron 文档, NKI Performance Guide. https://awsdocs-neuron.readthedocs-hosted.com/en/v2.26.0/general/nki/nki_perf_guide.html （同一文档集里的 NeuronDevice Architecture Guide 讲各个引擎的细节。）

更多

• A. Parashar et al., Timeloop: A Systematic Approach to DNN Accelerator Evaluation, ISPASS 2019；Y. N. Wu et al., Accelergy, ICCAD 2019。
• Y.-H. Chen, J. Emer, V. Sze, Eyeriss: A Spatial Architecture for Energy-Efficient Dataflow for Convolutional Neural Networks, ISCA 2016（行驻留、“数据流”这个词的出处）。
• V. Sze, Y.-H. Chen, T.-J. Yang, J. Emer, Efficient Processing of Deep Neural Networks（Morgan & Claypool, 2020）——这一领域的教材。
• J. Ragan-Kelley et al., Halide: A Language and Compiler for Optimizing Parallelism, Locality, and Recomputation in Image Processing Pipelines, PLDI 2013。

本仓库

• 前置讲 P4：循环嵌套、依赖分析、分块与”算法与调度分离”。
• 第 8 讲：脉动阵列、Tensor Core、数据块与流水线。

15. 下一讲预告 · Next lecture

这两讲站在硬件设计者的角度：给定一个计算，怎样的硬件和映射最好。接下来换到另一边：给定一块硬件（GPU），怎样写出跑满它的内核。第 10 讲从最朴素的矩阵乘法内核开始，一步步做：合并访存 → 共享内存分块 → 寄存器分块 → 双缓冲 → Tensor Core，每一步都用 roofline 和本讲的”访问次数”去分析它为什么变快——这也是原课作业 4 的内容。

→ 第 10 讲：内核开发