Lecture 1 · 导论：效率、能量与 roofline · Introduction: Efficiency, Energy, and Rooflines

对应 Maps to：Stanford CS349E / EE339 Efficient ML Inference at Scale（Autumn 2026；Zain Asgar、Mark Horowitz、Fredrik Kjolstad、Priyanka Raina），Lecture 1 · Tue Sep 22 · Introduction（Mark Horowitz）。阅读材料：J. Austin et al., How to Scale Your Model（Google DeepMind, 2025）的 Chapter 0（Introduction）和 Chapter 1（All About Rooflines）；NVIDIA 技术博客 Inside NVIDIA Blackwell Ultra: The Chip Powering the AI Factory Era（2025-08）。 依据说明：课程网站（cs349e.github.io）只公开了课表、阅读材料和作业安排，没有 slides 或录像。本讲依据两份阅读材料来写，另外补充了三份公开资料：Mark Horowitz 的 ISSCC 2014 主题演讲论文 Computing’s Energy Problem (and what we can do about it)（§2 的能量表出自这里，这是 Horowitz 最有代表性的观点）、X. Ma & D. Patterson 的 Challenges and Research Directions for Large Language Model Inference Hardware（IEEE Computer, 2026）、以及 Gimlet Labs 的两篇技术博客（课程算力由 Gimlet 提供，Zain Asgar 是其联合创始人）。这些补充材料是我选的，不代表课堂上实际讲了什么。 配套代码 Code：code/lecture01/（标准库 + numpy；5 个脚本 + 4 个共享模块；每个脚本都在几秒内跑完） 预计阅读时间：3–4 小时（含运行代码）

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

0.1 学习目标 Learning goals

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

1. 说清楚为什么”推理效率”值得开一整门课：推理的成本随使用量增长，而延迟、吞吐、成本、能耗四个目标互相牵制，没有一个”最好”的工作点，只有一条前沿（Pareto frontier）。
2. 用能量的视角解释今天的 AI 硬件为什么长成这样：Dennard scaling 结束之后，芯片的功耗预算固定，性能 = 能效 × 功率；而一次运算的能量远小于一次数据搬运，所以专用化和数据复用是提高能效的两条主路。
3. 用 roofline 模型估计任意一段计算在任意一块芯片上的时间下界：会算算术强度、临界强度、临界 batch；会判断 compute-bound / memory-bound / communication-bound；知道片上存储、网络、容量各自也是一层”屋顶”。
4. 把 roofline 用到 LLM 推理上：手算权重、每 token 运算量、每 token KV cache；解释为什么 decode 受限于带宽和容量、prefill 受限于算力；解释为什么长上下文时 decode 批再大也到不了 compute-bound。
5. 用 roofline 和能量的眼光读一份硬件发布材料（Blackwell Ultra）：看出哪些设计是在补哪块短板（NVFP4、288 GB HBM、softmax 指数单元、NVLink），也知道厂商数字要怎么”打折”读。
6. 描述”异构、解耦”推理的基本逻辑：一个请求里的不同阶段需要不同的硬件；近存（SRAM）和远存（HBM）是一对取舍。
7. 画出这门课的全栈地图（模型 → 服务系统 → 硬件 → 内核 → 编译器 → 让 AI 写内核），并知道后面每一讲在图上的位置。

0.2 先修与本讲范围 Prerequisites & scope

原课是研究生课程，默认你学过计算机体系结构、编译原理和一点并行编程。本讲义为没有这些背景的读者另写了几篇前置讲（见 SYLLABUS）。本讲只用到很少的硬件概念，用到时会从零讲。

0.3 本讲地图 Map

§1  为什么推理效率值得一门课：一个 token 的旅程；训练花一次、推理花一辈子；四个目标与前沿；"at scale"的两层意思
§2  能量视角：运算几乎免费；Dennard scaling 与它的终结；性能 = 能效 × 功率；指令开销；专用化；低精度   ← 代码 01
§3  Roofline：算 / 搬 / 存；上下界；算术强度与临界强度；矩阵乘法与临界 batch；低精度；
    批处理救不了的情况；分块与片上存储；网络 roofline；容量；在你自己的电脑上                         ← 代码 02、03
§4  LLM 推理上的 roofline：模型账本；decode；KV cache；容量；延迟–吞吐–能耗前沿；prefill；近存与远存   ← 代码 04
§5  案例：读懂 Blackwell Ultra：双 die；内存墙；NVFP4；softmax 的 exp；288 GB；NVLink；怎样读厂商数字   ← 代码 05
§6  异构与解耦：一个请求里的不同"性格"；近存 vs 远存；几种解耦的切法；代价
§7  全栈地图：这门课怎么走
§8  现状（2025–2026）
§9–13  小结、常见误区、练习（含答案）、延伸阅读、下一讲预告

0.4 记号 Notation（全课统一）

单位约定：FLOP 是”一次浮点运算”（乘法、加法各算一次），FLOPs 是”多少次运算”，FLOP/s 是”每秒多少次”。字节用十进制（1 GB = $10^9$ 字节），和厂商标的显存、带宽一致；需要 1024 进制时写 KiB / MiB / GiB。带宽默认指单向。

0.5 关于这门课 · About CS349E

CS349E（与电子工程系的 EE339 交叉列出）是 Stanford 2026 年秋季学期新开的研究生课，3 学分，每周二、四上课。课程目录里的描述是：

一门研究生课程，关注设计和构建大规模、解耦式（disaggregated）AI 推理系统的原理与技术。课程探讨架构、算法、系统和内核实现层面的技术，目标是在异构机器上以低延迟、高吞吐、低成本服务模型。主题包括模型优化技术、分布式服务、AI 硬件与性能建模、内核编程和 AI 编译器。（意译）

四位授课教师的背景（公开资料）恰好覆盖了这条技术栈的不同层次：

• Mark Horowitz：电子工程与计算机科学教授，研究 VLSI 电路和高能效计算。他在 ISSCC 2014 的主题演讲 Computing’s Energy Problem 是”为什么要做专用硬件”这个话题被引用最多的文献之一。他讲第 1 讲（导论）和第 4 讲（分布式推理）。
• Fredrik (Fred) Kjolstad：计算机科学系教授，研究编译器（代表作是稀疏张量代数编译器 TACO）。他讲 Transformer、KV cache 管理、内核开发、ML 编译器和”让 AI 写内核”。
• Priyanka Raina：电子工程系教授，研究深度学习加速器和硬件设计方法。她讲精度/量化/稀疏、硬件性能建模、硬件引擎和内核优化。
• Zain Asgar：Gimlet Labs 的联合创始人。Gimlet 做的是”把一个推理工作负载切开、分别放到 GPU、SRAM 芯片、CPU 等不同硬件上”的异构推理云，并为本课提供算力。他讲解耦式推理和 decode 优化。

课程形式和评分（课程网站）：

• 作业（40%）：四次，都在前五周：(1) LLM 推理；(2) 解耦式推理；(3) 张量并行；(4) 写内核。
• 课程项目（40%）：第 5 周项目提案（pitch），第 10 周两次课做展示。
• 课堂参与（20%）。
• 全课 19 次课：15 个技术主题、1 次客座讲座、1 次项目提案、2 次项目展示。

本讲义把 15 个技术主题整理成 16 讲（外加前置讲）。主要调整是把”KV cache 管理”挪到”服务系统”那一部分，把”硬件引擎”放在”硬件性能建模”之前（先认识硬件，再给它建模）。逐场对应见 SYLLABUS。

和 CS349D 的关系。如果你读过本仓库的 CS349D（AI Inference Infrastructure） 讲义，会发现两门课有一部分阅读材料重合（Orca、投机解码、DistServe、PagedAttention、Mooncake、FlashAttention、Pope et al.）。两门课的视角不同：CS349D 是一门系统论文研讨课，重点在调度、显存管理、集群；CS349E 由硬件和编译器方向的教师主讲，会一路往下走到硬件、内核和编译器。本讲义是自洽的，不要求读过 CS349D；重合的部分会换一个更偏硬件的角度来讲，并给出 CS349D 对应讲次作为延伸阅读。

0.6 关于数字 · About the numbers

本讲有三类数字，可信程度不同：

1. 公式算出来的数字（代码 01、02、04、05）：由公开规格（NVIDIA 博客和数据手册、How to Scale Your Model 的表格）、公开的模型结构（Meta Llama 3.1），按 roofline 公式计算。它们是理想下界：峰值按 100% 算，多芯片不计通信。数量级和趋势可信，但不是任何真实系统的实测值。每次运行完全相同。
2. 在 CPU 上实测的数字（代码 03）：来自编写本讲时的一台 Apple M4 Pro 笔记本（numpy 2.5，底层 BLAS 是 Apple Accelerate）。你的机器会给出不同的数字，只看趋势。
3. 别人报告的数字：论文、厂商博客里的数字一律写成”某某报告/宣称”，表示我没有复现。

1. 为什么推理效率值得一门课 · Why a whole course on inference efficiency

1.1 一个 token 的旅程

先看一个最普通的场景：你在聊天框里输入一个问题，按下回车。服务器那边发生了什么？

1. 分词（tokenize）：你的文字被切成一串整数（token id）。一个英文单词大约 1–2 个 token，一个汉字大约 1 个。
2. 排队：服务器上同时有很多人在提问，你的请求先进一个队列，等调度器把它安排到某块（或某几块）加速器上。
3. prefill：模型把你的整段提示一次性”读”进去。所有提示 token 可以同时计算，这一步算完时，模型吐出第一个输出 token。你等待的这段时间叫 TTFT（time to first token，首 token 延迟）。
4. decode 循环：之后每一步只生成一个 token：把上一步生成的 token 喂回模型，算出下一个。因为第 $t+1$ 个 token 依赖第 $t$ 个，这些步只能一步一步地做。相邻两个 token 之间的时间叫 TPOT（time per output token）或 ITL（inter-token latency）。
5. 流式返回：每生成一个 token 就（解码成文字）推给你，直到模型生成结束符或达到长度上限。

这里有一个后面会反复出现的”怪事”：decode 每一步只为每个用户生成 1 个 token，却要把整个模型的权重从显存里读一遍。一个 80 亿参数的模型，BF16 存储约 16 GB；生成一个 token 的运算只有约 160 亿次——在一块每秒能做 $10^{15}$ 次运算的 GPU 上，这只要十几微秒，可把 16 GB 从显存搬到计算单元却要好几毫秒。**芯片的大部分时间在等数据，而不是在算。**这一讲的大部分内容，就是把这件”怪事”讲透。

1.2 训练花一次，推理花一辈子

为什么推理的效率特别重要？看看钱是怎么花的：

• 训练是一次性的投入：花几周、几千到几万块芯片，得到一份权重。
• 推理是持续的开销：只要有人在用，就一直在花钱，而且和用量成正比。模型越成功，推理账单越大。

Ma & Patterson（2026）开篇就说”LLM inference is a crisis”：推理芯片的年销售额预计在 5–8 年内增长 4–6 倍；训练展示了能力，推理成本决定了能不能赚钱。几个新趋势让推理越来越贵：

• 推理模型（reasoning models）：在回答之前先生成一长串”思考”token，输出长度动辄几千到几万。
• 智能体（agents）：一个任务里反复调用模型、调用工具、读写长上下文。Gimlet Labs 的博客称智能体工作负载产生的 token 数约是普通聊天的 15 倍（这是它的说法，具体倍数因应用而异）。
• MoE、长上下文、多模态、RAG：每一项都在加大对显存容量、带宽和互连的需求（Ma & Patterson 表 2）。

所以”同样的硬件，每秒多产出 2 倍的 token”直接等于”成本减半”。这就是为什么系统研究者、芯片公司、云厂商都在抢这块地。

1.3 四个目标，互相牵制

衡量一个推理系统，至少有四个维度：

麻烦在于它们互相牵制。§4.5 会用数字说明：把更多用户的请求凑成一批一起算（批大小 $B$ 变大），每块芯片的产出、每兆瓦的产出都上去了，每个 token 的能耗也下来了，可每个用户等每个 token 的时间变长了。没有哪个 $B$ 在所有维度上都最好。把所有”最优折中”画在一张图上（横轴每用户速度，纵轴每芯片或每兆瓦产出），得到的曲线叫 Pareto 前沿（Pareto frontier）。NVIDIA 的 Blackwell Ultra 博客图 6 就是这样一张图：每一代硬件的意义，就是把这条前沿往右上方推。

停下来想一想：如果有人只告诉你”新系统的吞吐提高了 3 倍”，你最想追问什么？（答案：在什么延迟下？是不是同一个前沿上的点？见 §5.8。）

1.4 “at Scale” 的两层意思

课名里的 “at Scale” 至少有两层意思：

1. 模型太大，一块芯片放不下。一个 4050 亿参数的模型，BF16 存储要 812 GB，而一块 H100 只有 80 GB 显存。模型必须切开、分到多块芯片上，芯片之间就要通信（第 4 讲）。
2. 请求太多，要成千上万块芯片。这时问题变成：怎样把请求分到不同的机器上、怎样让不同阶段跑在最合适的硬件上（第 5、7 讲）。

而且这些芯片不一定是同一种：数据中心里有不同代际的 GPU、TPU、各种推理专用芯片、CPU。**异构（heterogeneous）**是这门课的一个关键词——课程描述里明确写着”on heterogeneous machines”。§6 会解释为什么异构不是”凑合着用”，而是有物理上的道理。

1.5 这门课的视角：一路往下走

一个 LLM 推理请求的速度，由一整条”栈”上的每一层共同决定：

  模型：结构、精度、稀疏                  ← 第 2、3 讲
    ↓
  服务系统：批处理、调度、KV cache、解耦  ← 第 4–7 讲
    ↓
  硬件：GPU / TPU / SRAM 芯片的结构       ← 第 8、9 讲
    ↓
  内核：每个算子在硬件上怎样执行          ← 第 10、11 讲
    ↓
  编译器：怎样自动生成好的内核            ← 第 12、13 讲
    ↓
  让 AI 来写内核、用工具                  ← 第 14、15 讲

有了这张图，你会发现不同层次的优化其实在和同几个物理限制打交道：算力、带宽、容量、通信、能量。这一讲的任务，就是把这几个限制和衡量它们的工具（roofline）一次讲清楚，后面每一讲都会回来用它。

2. 能量视角：硬件为什么长成今天这样 · The energy view

这一节主要依据 Horowitz（ISSCC 2014）。他是本课导论的主讲人，这一节代表的是他那一派体系结构研究者看问题的方式：一切从能量出发。

2.1 一个反直觉的事实：运算几乎是免费的

我们习惯说”这个模型计算量很大”，好像”算”是最贵的事。从能量的角度看，事实正好相反。下表是 Horowitz 论文图 1.1.9 给出的 45 nm 工艺下各种操作的能量（01_energy_and_scaling.py 演示 1）：

三个结论：

1. 数据搬运比运算贵几百到几万倍。从 DRAM 读 8 个字节，够做几百次 32 位浮点乘法、几万次 8 位加法。
2. 存储越大越贵。8 KB 的 SRAM 读一次 10 pJ，1 MB 的 100 pJ：存储越大，信号要走的线越长，能量越高。这就是为什么所有处理器都有一个存储层次（memory hierarchy）：寄存器 → 小而快的 SRAM → 大一些的 SRAM（cache） → 片外 DRAM。常用的数据放在近处，用的时候便宜。
3. “管理”比”干活”贵。一条最简单的指令（比如一次加法）总共要 70 pJ，其中真正做加法的只有零点几个 pJ；剩下的花在取指令（25 pJ）、读写寄存器、控制逻辑上。

打个比方：你在厨房做饭，切菜（运算）只要几秒钟，可如果每切一刀都要去一趟超市买菜（DRAM），饭就永远做不完。好厨师会一次把菜买够、放在手边的案板上（片上存储），切很多刀才去一次超市。这一讲后面所有的”复用""分块""批处理”，本质上都是在做这件事。

这些是 45 nm（大约 2008 年）的数字。今天的 3–5 nm 工艺里，每种操作的绝对能量都小了几倍，但”DRAM ≫ SRAM ≫ 运算”的比例关系基本没变——甚至更悬殊，因为运算单元随工艺缩小得更快，而导线和片外接口缩小得慢。要记住的是比例，不是绝对值。

能量版的”算术强度”。既然搬运贵、运算便宜，自然要问：每从某一级存储读 1 字节，至少要做多少次运算，运算本身的能量才”追上”搬运的能量？设一次运算 $e_\text{op}$ pJ、读 1 字节 $e_\text{byte}$ pJ，答案是

代码 01 演示 2 算出来：对 INT8 乘加，从 DRAM 读数据时 $I^*_\text{energy}$ 是 1,413–2,826 次运算/字节；FP16 乘加是 217–433；FP32 乘加是 71–141。Horowitz 论文里有一句话正好对应：要达到每次运算约 1 pJ 的能效，应用必须用 8–16 位的短整数，每次读本地存储要完成几十次运算，每次读 DRAM 要完成大约一千次运算——“这看起来不太可能，其实很常见，卷积就是这样”。矩阵乘法也是。

注意精度越低，$I^*_\text{energy}$ 越高：低精度让”算”变便宜了，但”搬”没变，所以数据复用变得更加关键。这句话在第 3 讲（量化）还会回来。

2.2 Dennard scaling：过去 30 年的”免费午餐”

为什么”能量”这么重要？要从晶体管说起。

一块芯片的动态功耗大约是

其中 $C$ 是要充放电的电容，$V$ 是电压，$f$ 是频率，$\alpha$ 是”有多少比例的电路在翻转”。每一代新工艺把线宽缩小到原来的 $1/S$（$S \approx 1.4$，面积减半）。1974 年 Robert Dennard 等人指出，如果电压也按 $1/S$ 缩小，那么：

• 每个晶体管的电容 $\propto 1/S$，电压 $\propto 1/S$，频率可以 $\propto S$（晶体管更小、更快）；
• 每个晶体管的功耗 $\propto \frac{1}{S} \cdot \frac{1}{S^2} \cdot S = \frac{1}{S^2}$；
• 同样面积里晶体管数 $\times S^2$，所以单位面积的功耗不变。

这就是 Dennard scaling：每一代晶体管翻倍、频率提高 40%，芯片却不会更热。性能是工艺”白送”的，软件什么都不用改，下一代 CPU 自然就快了。这顿”免费午餐”吃了大约 30 年（和 Moore 定律”晶体管数每两年翻倍”一起）。

2.3 免费午餐的结束：功耗墙与暗硅

大约 2005 年前后，电压降不下去了：电压再低，晶体管”关不严”，漏电功耗急剧上升。电压不能降，上面的账就变了（代码 01 演示 3）：

• 即使频率不再提高，每个晶体管的功耗也只按 $1/S$ 下降，而单位面积晶体管数 $\times S^2$，功耗密度每代 $\times S \approx 1.4$；
• 如果还想让频率继续上涨，功耗密度每代 $\times S^2 \approx 2$。

芯片能散掉的热量（每平方厘米多少瓦）是有上限的。于是出现了三个后果：

1. 频率停滞：CPU 主频从 2005 年左右就停在 3–5 GHz。
2. 转向多核：既然单核不能更快，就在一块芯片上放更多核。但这要求软件能并行。
3. 暗硅（dark silicon）：晶体管还在变多，但不能同时打开。按演示 3 的玩具模型，6 代之后同样的散热能力只够让 13% 的晶体管同时工作。

暗硅听起来是坏事，但它指出了一条路：既然不能同时点亮所有晶体管，最划算的做法是把它们做成”只干一类活、但干得极省电”的专用电路，需要哪种就开哪种。今天的 SoC（手机芯片）里有 CPU、GPU、图像处理器、视频编解码器、神经网络加速器，就是这个思路。GPU 的 Tensor Core、TPU 的脉动阵列，也是这条路上的产物。

2.4 性能 = 能效 × 功率

现在可以说出这一节最重要的一句话。芯片的功耗有上限（一块 H100 最多 700 W，一块 Blackwell Ultra 最多 1400 W），数据中心的电力也有上限（以兆瓦计，而且越来越成为建数据中心的瓶颈）。在功耗固定的前提下：

想要更快，就只能让每次运算花更少的能量。这就是为什么 NVIDIA 现在用”每兆瓦每秒产出多少 token”（TPS/MW）来宣传新一代系统，Ma & Patterson 也把”功耗和碳排放”列为一阶优化目标：在电力受限的世界里，能效就是性能。

2.5 能量去哪了：指令开销

回到”一条指令 70 pJ，其中加法只占 0.03 pJ”这件事。处理器是可编程的：每做一次运算，都要先取一条指令、译码、读寄存器、写寄存器、更新控制状态。这些”管理”工作的能量是固定的，不管这条指令干的活多大。

所以提高能效的第一个办法是：让一条指令干更多活，把固定开销摊薄（代码 01 演示 4，运算能量统一按 FP16 乘加计）：

• SIMD（single instruction, multiple data）：一条指令同时对一组数据做同样的运算。
• GPU 的 SIMT（single instruction, multiple threads）：32 个线程组成一个 warp，共享一条指令流（前置讲 P2 §4）。
• Tensor Core：一条指令完成一个小矩阵乘法（比如 $16 \times 8 \times 16$，2048 次乘加）。Hopper 的 wgmma 让 4 个 warp 组成的”warpgroup”一条指令做 $64 \times 256 \times 16$ 的矩阵乘法。

Horowitz 在论文里这样描述 GPU 的选择：浮点运算大约只有一条简单指令能量的 1/10，所以做一个大约 10 lane 的 SIMD 引擎，指令的能量就被运算主导了——“这就是 GPU 走的路”。Tensor Core 把这条路走到了头：管理开销几乎为零。

代价是灵活性：一条指令必须对一大块规整的数据做同样的事。神经网络的矩阵乘法恰好满足这个条件，这是 AI 硬件能如此高效的根本原因。反过来，凡是不规整的计算（稀疏、条件分支、逐元素的特殊函数），在这种硬件上都会吃亏——§5.5 的 softmax 就是一个例子。

演示 4 是一个简化：它把”70 pJ 每条指令”套在所有架构上，只为看数量级。GPU 真实的指令开销和 CPU 不同；而且操作数从寄存器送到计算单元也要花能量（Tensor Core 靠矩阵内部的数据复用，把这部分也摊薄了）。

2.6 专用化：两三个数量级从哪里来

把运算、指令、片上存储、DRAM 四项加起来，看一次 INT8 乘加在不同架构上的总能量（代码 01 演示 5）：

“复用因子”的意思是”每做多少次乘加才需要一次这种开销”。这些复用因子是我取的有代表性的假设，只用来说明数量级；真实芯片的能效还取决于工艺、电压、利用率等很多因素。但结论很稳健：专用硬件比通用处理器高效约两个数量级，和 Horowitz 论文图 1.1.8 “专用硬件比基于处理器的方案高效 2–3 个数量级”一致。

这两个数量级来自两件事，记住它们，后面每一讲都会用到：

1. 摊薄控制开销：一条指令（或者干脆没有指令，像脉动阵列那样数据自己流过固定的电路）完成成千上万次运算。
2. 数据复用（locality）：一个数从远处（DRAM）搬进来以后，在近处（寄存器、小 SRAM）被用很多次。

注意专用阵列剩下的能量里，DRAM 仍然占了很大一部分（0.42 pJ 里有 0.16 pJ 是 DRAM）。复用因子决定一切——这正是第 9 讲（硬件性能建模，Interstellar）要系统研究的问题：一个矩阵乘法有很多种”循环顺序”和”分块方式”，怎样安排才能让数据在便宜的存储里被复用最多次？

专用化的代价。专用芯片有两个风险：一是设计和流片很贵（动辄上亿美元）；二是押错了算法就全赔了。How to Scale Your Model 第 0 章讲了 TPU 的故事：Google 在 2010 年代中期押注”神经网络 = 矩阵乘法”，设计了以脉动阵列为核心的 TPU，结果押对了，性价比远超当时的 GPU；但”可以想象，如果神经网络没有火起来，或者变成了 TPU（比 GPU 更不灵活）处理不了的样子，代价会有多大”。硬件设计者要提前 2–3 年押注算法会长成什么样，这叫软硬件协同设计（co-design）——§5 会用 Blackwell Ultra 看一个具体的例子。

2.7 低精度：最便宜的那一刀

能量表里还有一个规律：位宽越小越便宜。32 位整数乘法 3.1 pJ，8 位只要 0.2 pJ，差 15 倍。乘法器的面积和能量大致随位宽的平方增长（$n$ 位乘法要处理 $n \times n$ 个部分积）。

低精度同时省两样东西：

• 运算更便宜、更快（同样的芯片面积能放更多低精度乘法器：H100 的 FP8 算力是 BF16 的 2 倍，Blackwell Ultra 的 NVFP4 是 FP8 的 3 倍）；
• 搬运更少（每个数的字节数减半，读权重的时间也减半）。

代价是精度：数能表示的范围和精细程度变小，模型的输出可能变差。怎样在”省”和”准”之间取舍，是第 3 讲（精度、量化与稀疏）的主题。§5.4 会先用 NVFP4 做一个小实验。

3. Roofline：三种资源与一个比值 · Rooflines

这一节依据 How to Scale Your Model 第 1 章，代码是 roofline.py（公式）、02_roofline_basics.py（例题）和 03_measure_your_machine.py（实测）。

3.1 时间去哪了：算、搬、存

一段计算为什么要 50 ms，而不是 5 ms 或 50 s？How to Scale Your Model 第 1 章开门见山：在硬件上运行一个算法，受三样东西限制——

1. 算得多快：芯片每秒能做多少次运算（FLOP/s）。
2. 搬得多快：每秒能搬多少字节（byte/s）。搬运可以发生在芯片内部（显存 HBM → 计算单元），也可以发生在芯片之间（NVLink、TPU 的 ICI、数据中心网络）。
3. 存得下多少：一共能放多少字节（显存容量）。

前两样是”速率”，决定时间；第三样是”容量”，决定能不能做、要切成几份做。于是

例子：H100 的 BF16 稠密峰值约 $9.89 \times 10^{14}$ FLOP/s，所以 $10^{12}$ 次运算至少要 1.01 ms；它的 HBM 带宽是 3.35 TB/s，所以搬 16 GB 至少要 4.8 ms。

厂商标的是峰值：How to Scale Your Model 提到，H100、B200 在实际的大矩阵乘法上通常只能跑到标称算力的 80%–85%，TPU 能接近 95%。所以 roofline 给出的是下界，真实时间会长一些。

3.2 重叠：下界与上界

“算”和”搬”可以同时进行吗？大多数情况下可以：芯片可以一边把下一块数据搬进来，一边算当前这一块（这叫流水或双缓冲，第 10 讲会写代码实现）。于是：

• 完美重叠时，总时间由较慢的那一个决定，这是下界。
• 完全不重叠时，总时间是两者之和，这是上界。上下界最多差 2 倍。

哪一项更大，我们就说这段计算是 compute-bound（受限于算力，芯片满负荷在算）还是 memory-bound / communication-bound（受限于带宽，计算单元有一部分时间在等数据）。在 memory-bound 的情况下，你买的算力有一部分被浪费了。

3.3 算术强度与临界强度

判断瓶颈不需要算出时间，只要比较两个比值：

• 算术强度（arithmetic intensity）$I$ 是这段计算的属性：每搬 1 字节，做几次运算。
• 临界强度（critical intensity，也叫 ridge point，屋脊点）$I^*$ 是这块芯片的属性：峰值算力除以带宽。H100 BF16：$9.89 \times 10^{14} / 3.35 \times 10^{12} \approx 295$ FLOP/byte；TPU v5e：$1.97 \times 10^{14} / 8.2 \times 10^{11} \approx 240$。

打个比方：芯片是一个厨师（算力 $\pi$，每秒切几刀）和一辆送菜的小货车（带宽 $\beta$，每秒送几斤菜）。一道菜每斤原料要切多少刀，就是这道菜的”强度”。切丝的菜（强度高）厨师忙不过来，货车闲着；凉拌黄瓜拍一拍就行（强度低），厨师在等货车。临界强度就是”厨师和货车刚好一样忙”的那道菜。

芯片设计者的选择就是这个比值。§5.3 会看到，每一代芯片的算力都比带宽涨得快，临界强度越来越高——同样的程序，在新芯片上会”更 memory-bound”。

3.4 读一张 roofline 图

把”可达算力”画成”算术强度”的函数：

在 log-log 坐标下，这是一条先上升、后水平的折线，像一个屋顶（roofline 由此得名）：

• 斜坡（左半边）：强度每翻一倍，可达算力翻一倍。这里的计算受限于带宽。斜坡的高度由带宽决定：带宽越大，斜坡越高。
• 平顶（右半边）：到达峰值算力，再提高强度也没用。
• 屋脊点：斜坡和平顶的交点，横坐标正好是 $I^*$。

代码 02 演示 4 画了一张 H100 BF16 的字符画版 roofline（加 --png 参数还会存一张 PNG），把各种运算放在图上。读图的方法：找到你的运算的强度，往上看它碰到的是斜坡还是平顶。

一块芯片其实不止一个屋顶：

• 不同精度有不同的平顶（FP8 比 BF16 高一倍）。
• 不同的计算单元有不同的平顶：矩阵乘法跑在 Tensor Core（TPU 叫 MXU）上，逐元素运算跑在普通向量单元上（TPU 叫 VPU），后者的峰值低得多。H100 的 Tensor Core BF16 约 989 TFLOP/s，CUDA core 的 FP32 只有约 67 TFLOP/s。
• 不同级别的存储有不同的斜坡：从 HBM 读、从 L2 cache 读、从片上共享内存读，带宽差好几倍（§3.9）。
• 芯片之间的网络也是一个斜坡（§3.10）。

3.5 例子：逐元素运算和点积

最简单的例子（代码 02 演示 1，全部 BF16，每个数 2 字节）：

• 向量加法 $z = x + y$，长度 $N$：$N$ 次加法；读 $2N$ 个数、写 $N$ 个数，共 $6N$ 字节。强度 $= N / 6N = 1/6$。
• 点积 $x \cdot y$：$N$ 次乘法 + $N-1$ 次加法；读 $4N$ 字节、写 2 字节。强度 $= (2N-1)/(4N+2) \to 1/2$。

这两个强度不随 $N$ 变化：数据再多，每个数也只被用一次。在 H100 上，它们只能用到峰值算力的 0.1% 和 0.2%。在任何 AI 芯片上，逐元素运算（激活函数、归一化、残差相加……）都是 memory-bound 的。

怎么让逐元素运算快？不是”算得更快”，而是”少搬”：把好几个逐元素运算融合（fuse）成一个内核，中间结果留在片上，不写回显存。这是第 11 讲（内核优化）和第 12 讲（编译器的图级优化）的主题。

3.6 例子：矩阵乘法与”临界 batch”

矩阵乘法 $X[B, D] \cdot W[D, F] \to Z[B, F]$（BF16）：

• 运算：$B \cdot F$ 个输出，每个是长度 $D$ 的点积，共 $2BDF$ 次运算。
• 搬运（下界，每个数只搬一次）：读 $X$ 是 $2BD$ 字节，读 $W$ 是 $2DF$，写 $Z$ 是 $2BF$。

这个近似非常重要：矩阵乘法的强度约等于 $B$。直觉是：权重 $W$ 从显存搬进来一次，被 $B$ 行输入各用一次，复用了 $B$ 次。所以：

• $B = 1$（矩阵-向量乘）：强度 ≈ 1，和点积一样糟。
• $B$ 涨到几百以上，才越过 H100 的屋脊点 295，变成 compute-bound。

代码 02 演示 1 的数字（$D = F = 8192$，H100）：

注意时间那一列：$B$ 从 1 到 256，时间几乎不变（都在读 $W$），多算 255 行输入几乎是免费的。越过屋脊点之后，时间才开始和 $B$ 成正比。这就是”批处理”在推理里如此重要的原因（§4.2）。

How to Scale Your Model 把这条结论写成了一句口诀：BF16 矩阵乘法要 compute-bound，在大多数 TPU 上需要 $B > 240$，在 GPU 上大约 $B > 300$。这里的 $B$ 是 token 数，不是序列数：512 个长度 4096 的序列分到 2048 块芯片上，每块芯片每步处理的是 $512 \times 4096 / 2048 = 1024$ 个 token。

用精确公式求临界 $B$（roofline.critical_batch，代码 02 演示 2），会比 $I^*$ 略大一点，因为分母里还有读 $X$、写 $Z$ 的字节：

两个值得注意的点：H200 和 H100 是同一颗计算芯片，只是 HBM 更快（4.8 vs 3.35 TB/s），屋脊点从 295 降到 206——带宽变大，更容易 compute-bound。NVFP4 则把屋脊点推到了 1,875（下一小节解释）。

3.7 低精度怎样改变强度

How to Scale Your Model 第 1 章的练习 1、2 讲的是同一件事的两面（代码 02 演示 3，TPU v5e）：

• 练习 1：int8 × int8，用 int8 算力。int8 峰值是 bf16 的 2 倍，但每个数只有 1 字节，强度 $\approx 2B$。两边同时翻倍，临界 $B$ 不变（书中近似 240，精确值 256）。低精度让”算”和”搬”同时变快，不会自动改变瓶颈。
• 练习 2：int8 权重 × bf16 激活，用 bf16 算力。只有读权重的字节减半，算力不变，强度 $\approx 2B$，临界 $B$ 减半（书中近似 120，精确值 128）。这叫仅权重量化（weight-only quantization），在 decode（$B$ 小、读权重是瓶颈）时特别有效。

NVFP4 是一个有意思的例外。Blackwell Ultra 的 NVFP4 算力是 FP8 的 3 倍（15 vs 5 PFLOP/s），而每个数的字节只少了 1.78 倍（NVFP4 平均 4.5 bit，见 §5.4）。算力涨得比字节省得多，屋脊点跳到 1,875，临界 $B$ 涨到约 750。也就是说：在 NVFP4 下，这块芯片”算力相对更富余”，要更大的批才能喂饱。

3.8 批处理救不了的情况：“每个样本一个矩阵”

练习 4 问了一个看似刁钻的问题：如果批里每个样本都有自己的矩阵，$x[b] \cdot W[b]$，强度是多少？

运算量还是 $2BDF$，可搬运量变成了 $BD + \mathbf{BDF} + BF$（每个样本的矩阵都要搬一遍）。强度 $\approx 2BDF / (2 \cdot BDF) = 1$（BF16），和 $B$ 无关（代码 02 演示 3：$B$ 从 1 到 4096，强度都是 1.000）。

这为什么重要？因为 decode 阶段的注意力正是这个形状。每个序列都有自己的 KV cache（之前所有 token 的 key 和 value），批里的序列再多，也不能共享同一份 K、V。算一下每层每个序列的注意力（BF16，ctx = 4096，32 个 query 头，每头 128 维）：

强度 = query 头数 / KV 头数，和 $B$、ctx 都无关。批处理救不了 decode 的注意力，能救它的是模型结构：让多个 query 头共享同一组 K、V（GQA/MQA），每读一次 KV 就能多用几次。这是 GQA、MQA（以及 DeepSeek 的 MLA）被发明出来的系统动机之一，第 2 讲会展开。

3.9 分块与片上存储：屋顶不止一层

§3.6 的搬运量是”每个数只搬一次”的下界，前提是片上存储大到能放下整个矩阵。真实的芯片片上存储很小（H100 每个 SM 只有约 228 KB 的共享内存），大矩阵乘法必须分块（tiling）：把输出切成 $b_m \times b_n$ 的小块，每块沿 $k$ 方向分若干步，每步把输入的一小块搬进片上存储，算完再搬下一块。

How to Scale Your Model 给出了分块后的搬运量公式（$(m, k) \cdot (k, n)$，块大小 $b_m, b_k, b_n$，$t_m = m / b_m$ 等）：

块越大，每个搬进来的数被复用越多次，强度越高（代码 02 演示 5：$32 \times 32$ 的块强度 16，$256 \times 256$ 的块强度 128）。那块能有多大？输出块的累加器（FP32）要一直待在片上，直到 $k$ 方向全部算完。一个 SM 的寄存器是 256 KiB，最多放下 $256 \times 256$ 的 FP32 累加器，此时强度 128——还不到 H100 屋脊点 295 的一半。要到 295，输出块得约 $590 \times 590$，累加器 1.3 MiB。

所以单个 SM “自己搬自己算”是到不了屋脊点的。真实 GPU 用两种办法补上差距：

1. L2 cache 共享：H100 有 50 MB、B200 有 126 MB 的 L2 cache，被所有 SM 共享。多个 SM 需要的同一块输入，只从 HBM 读一次，其余从 L2 读。所以”每个 SM 看到的强度”是相对 L2 带宽的，而 HBM 实际流量要少得多。roofline 在这里多了一层屋顶。
2. 硬件为更大的块而改：Blackwell 给每个 SM 加了 256 KB 的 Tensor Memory（TMEM）专门放累加器，并支持两个 SM 合作做一次矩阵乘法、共享操作数（Blackwell Ultra 博客：“dual-thread-block MMA, where paired SMs cooperate on a single MMA operation, sharing operands and reducing redundant memory traffic”）。两者都是为了做更大的块、提高强度。

这是”硬件为了 roofline 而设计”的一个具体例子。第 8 讲（硬件引擎）会讲 SM 的内部结构，第 10–11 讲会亲手写分块的矩阵乘法内核。

3.10 网络 roofline

上面的屋顶都在一块芯片内部。模型一大，就要切到多块芯片上，芯片之间的网络成了新的”带宽”。How to Scale Your Model 举了一个例子：把 $X[B, D] \cdot W[D, F]$ 沿 $D$ 维切成两半，放在两块芯片上：

• 芯片 0 算 $X[:, :D/2] \cdot W[:D/2, :]$，芯片 1 算另一半，每块芯片的运算是 $BDF$；
• 两边各得到一份完整形状 $[B, F]$ 的”部分和”，要发给对方相加，发送 $2BF$ 字节（BF16）。

$B$ 在比值里约掉了：算的量和发的量都正比于 $B$。网络强度只和”被切开的那一维”$D$ 有关。临界条件是 $D/2 > \pi / \beta_\text{link}$（代码 02 演示 6）：

最后两行藏着一个趋势：Blackwell Ultra 的 BF16 算力（按 FP8 的一半推算）是 H100 的 2.5 倍，NVLink 只涨了 2 倍，网络临界强度反而变高了。算力增长快于通信，是贯穿全课的趋势。第 4 讲（分布式推理）会用这套方法系统地决定”模型该怎么切”。

3.11 第三种资源：容量

容量和前两种资源不一样：它不是速率，而是一道硬约束。放不下就是放不下，不是”慢一点”。容量不够时只有几个选择：

• 切分：把模型或 KV cache 分到多块芯片上（代价是通信）；
• 压缩：低精度存储（第 3 讲）、更省 KV 的模型结构（GQA/MLA，第 2 讲）；
• 换出（offload）：把暂时不用的数据放到更大更慢的存储里（CPU 内存、SSD），用的时候再搬回来（代价是带宽和延迟，第 7 讲的 Mooncake）。

§4.4 会看到，在 LLM 推理里，容量常常是比带宽更早碰到的天花板。

3.12 在你自己的电脑上量一条 roofline

roofline 是一个通用模型，不只适用于 GPU。代码 03 在笔记本 CPU 上用 numpy 量了三样东西。编写本讲时一台 Apple M4 Pro 的结果（每次运行会有几个百分点的波动）：

• 带宽：复制一个 268 MB 的数组（单核）约 130 GB/s；对 $8192 \times 8192$ 的 FP32 矩阵做矩阵-向量乘（多线程 BLAS）约 189 GB/s。
• 峰值算力：$4096 \times 4096$ 的 FP32 矩阵乘法约 3.3 TFLOP/s。
• 屋脊点：$3.3 \times 10^{12} / 1.89 \times 10^{11} \approx 17$ FLOP/byte。比 H100（295）低得多：CPU 的算力和 GPU 差得多，带宽差得没那么多。

然后扫描 $X[B, 8192] \cdot W[8192, 8192]$ 的 $B$：

大的趋势和 roofline 一致：小 $B$ 时耗时几乎不变、FLOP/s 随 $B$ 线性上涨；大 $B$ 时趋于峰值、耗时开始正比增长。但有一个意外：$B = 2$ 到 32 时，耗时（约 5.5 ms）是 $B = 1$（1.33 ms）的 4 倍，等效带宽只有约 50 GB/s，只到 roofline 预测的 25%。

原因不在 roofline，而在实现：numpy 对 $B = 1$ 调用专门的矩阵-向量乘（GEMV）代码，$B \ge 2$ 改走通用矩阵乘法（GEMM），而这台机器上的 GEMM 对”很瘦”的输入（只有几行）没有把内存带宽用满。这给了我们两条重要的教训：

1. roofline 只给出上界。能不能贴近上界，取决于代码（内核）写得好不好。
2. LLM decode 恰好就是”很瘦”的矩阵乘法（$B$ 是几到几十个序列）。通用的矩阵乘法库不一定为它优化过，推理引擎要专门为 decode 写内核（第 10–11 讲）。

（表里 $B = 1$ 的”实测/预测”略超过 100%，是因为作为预测基准的带宽本身也是实测的、带着几个百分点的噪声。）

4. 把 roofline 用到 LLM 推理上 · LLM inference on the roofline

代码：llm_math.py（推理账本）和 04_llm_inference_on_rooflines.py。这一节的数字全部是理想下界：峰值按 100% 算，多芯片不计通信，不算激活读写和每步固定开销。

4.1 一个 token 要做多少事：模型账本

先建立一本”账”。一个 Llama 风格的 decoder-only Transformer，每一层有注意力（$W_q, W_k, W_v, W_o$ 四个矩阵）和 MLP（SwiGLU 的 gate、up、down 三个矩阵），外加词嵌入表和输出层 lm_head。第 2 讲会逐个矩阵推导，这里先记住三条近似：

1. 每个 token 的运算量 ≈ 2 × 参数量。每个参与矩阵乘法的参数，对每个 token 做一次乘法、一次加法。（词嵌入只是”查表”，不算。）
2. 注意力打分的额外运算 ∝ 上下文长度：每个 token 要和上下文里的每个位置做两次点积（$QK^\top$ 和 $PV$），每层 $4 \cdot n_\text{heads} \cdot d_\text{head} \cdot \text{ctx}$ 次。
3. 每个 token 的 KV cache $= 2 \times n_\text{layers} \times n_\text{kv heads} \times d_\text{head} \times$ 每个数的字节数。

用公开的模型结构算出来（代码 04 演示 1；参数量和 Meta 公布的 8.03B、70.6B、405B 相差不到 0.3%）：

几个直觉：

• 8K 上下文时，8B 模型的注意力打分多出约 29% 的运算——长上下文时注意力的运算不能忽略。
• KV cache 每 token 只有几百 KiB，看起来很小；但一个 8K 上下文的序列就要 1–4 GB，几十个序列就和权重一样大了。
• 三个模型的 KV / token 增长得比参数慢得多（8 组 KV 头的 GQA 设计）。

4.2 decode：每一步都要把权重读一遍

decode 的每一步，批里 $B$ 个序列各生成 1 个 token。这一步要：

• 读全部权重一次（所有序列共享，被 $B$ 个 token 复用）；
• 读每个序列各自的 KV cache（不能共享）；
• 做 $B \times$ 每 token 运算量次运算。

代码 04 演示 2：Llama-3.1-8B，一块 H100，BF16，每个序列上下文 4096：

先看 $B = 1$：一步至少 4.64 ms，几乎全花在读 15 GB 的权重上；计算只要 17 µs，H100 的算力用了不到 0.4%。这就是 §1.1 说的”怪事”的精确版本。

再看批处理：读权重的时间不随 $B$ 变化。$B$ 从 1 涨到 16，每块芯片的产出涨了 10.5 倍，每个用户的速度只降了 34%。批处理是推理的第一杠杆：同一次权重读取，服务更多的人。第 6 讲（Orca 的连续批处理）会讲怎样在请求随时到达、随时结束的情况下维持一个大批。

4.3 KV cache：批处理摊不薄的那部分

但表里的”读 KV”随 $B$ 线性增长：$B = 256$ 时读 KV 是读权重的 9.2 倍。再加大 $B$，芯片产出几乎不涨（5,626 → 5,917），每个用户却越来越慢。整张表里瓶颈始终是 memory——为什么到不了 compute-bound？

这正是 §3.8 的”每个样本一个矩阵”。设每个序列每一步的计算时间是 $f/\pi$（$f$ 是每 token 运算量），读自己 KV 的时间是 $k/\beta$（$k$ = ctx × 每 token KV 字节），读权重的时间是 $W/\beta$。要 compute-bound，需要

只有当每多一个序列带来的计算时间大于它带来的读 KV 时间（$f/\pi > k/\beta$）时，加大 $B$ 才能追上固定的读权重时间。代码 04 演示 3：

ctx = 0（只看权重）时，$B^* \approx 296$，和矩阵乘法的临界 batch 一致。但上下文只要超过几百个 token，批再大也到不了 compute-bound：强度上限 = 每 token 运算量 / 每序列 KV 字节，随 ctx 下降，低于屋脊点 295。而聊天、推理、智能体的上下文通常是几千到几十万个 token。

结论：长上下文 decode 本质上是 KV cache 的”带宽 + 容量”问题，不是算力问题。减小 KV（GQA/MQA/MLA、KV 量化）、少读 KV（稀疏注意力）、把 KV 放到合适的地方（KV cache 管理、解耦），都是在攻这一点（第 2、3、5、7 讲）。

4.4 容量：批大小的天花板

§4.2 说批越大越划算（至少到某一点为止）。可批能有多大？显存要放下权重 + 所有序列的 KV cache：

代码 04 演示 4（预留 10% 给激活和工作区，KV 用 BF16）：

• 70B 的 BF16 权重（141 GB）放不进一块 H100：要么量化，要么切到多块芯片（第 4 讲）。
• 128K 上下文时，即使 8B 模型、一块 H100 也只能放 3 个序列——远低于”高效”需要的几百。容量直接把批大小压在了低效区。
• 这就是为什么 Blackwell Ultra 把显存加到 288 GB（H100 的 3.6 倍）：更大的显存是推理芯片最重要的卖点之一，不亚于算力。再配合 NVFP4 把权重压小，腾出的空间全部给 KV cache。

4.5 延迟、吞吐、能耗的前沿

现在把 §1.3 的四个目标放到一起。Llama-3.1-70B，8 块芯片理想切分（算力和带宽 ×8，不计通信），ctx = 4096，扫描 $B$（代码 04 演示 5，节选）：

（8×H100 FP8 在这个上下文下最多只能放 376 个序列，$B = 512$ 放不下。）

怎么读这张表：

• 每个部署从上到下就是沿前沿移动：用户速度一路下降；芯片产出、每兆瓦产出一路上升；每 token 能耗、每百万 token 的芯片·时一路下降。选哪个点取决于产品：聊天要 TPOT 在几十毫秒以内，离线批处理只看成本。
• 批处理同时是”省钱”和”省电”的杠杆：$B = 1$ 时每个 token 要花 7–15 焦耳（8 块芯片满功率等一个 token），$B = 256$ 时降到 0.26–0.34 焦耳。再往后降得很慢——KV 读取不能被摊薄。
• 新硬件把整条前沿往右上推：同样 8 块芯片，Blackwell Ultra 在每个 $B$ 上都更快（带宽 8 vs 3.35 TB/s），显存还能多放几倍的序列。
• NVFP4 在小 $B$ 时收益最大（权重主导：1.11 → 0.63 ms），大 $B$ 时收益变小（KV 主导，而这里的 KV 仍是 BF16）。

“芯片·时 / 百万 token”乘以你的每卡每小时价格，就是每百万 token 的芯片成本。比如 0.13 芯片·时 × 每小时若干美元。这只是芯片租金，没算 CPU、网络、电费、空闲时间。

能耗按最大功耗算，是偏保守的估计：memory-bound 的芯片通常达不到最大功耗。另外这些都是不计通信的理想值；8 卡张量并行每层要做两次 all-reduce，小 $B$ 时通信延迟占比很高（第 4 讲）。

4.6 prefill：另一种”性格”

prefill 一次处理整段提示：$n$ 个 token 共享同一次权重读取，强度 ≈ $n$。代码 04 演示 6（Llama-3.1-8B，一块 H100，BF16）：

• 几百个 token 就越过了屋脊点：prefill 是 compute-bound 的，和 decode 正好相反。同一个模型的两个阶段，“性格”完全不同。
• 提示很长时，注意力的运算按 $n^2$ 增长：128K 的提示里注意力占了 70% 的运算，TTFT 下界涨到 6.5 秒（而且这是 100% 峰值的理想值）。

prefill 吃算力，decode 吃带宽和容量。如果把它们放在同一块 GPU 上，长 prefill 会卡住正在 decode 的用户（生成停顿），decode 又让 GPU 的算力闲着。把它们放到不同的机器、甚至不同种类的芯片上，就是第 5 讲”解耦式推理”（DistServe）的出发点。

4.7 近存与远存：SRAM 芯片

到目前为止，“主存”都是 HBM：一种堆叠在芯片旁边的 DRAM，容量大（几十到几百 GB），带宽几 TB/s。还有另一条路：把权重直接放在芯片上的 SRAM 里。SRAM 就在计算单元旁边，带宽高得多，但每比特面积大得多（6 个晶体管存 1 bit，DRAM 只要 1 个晶体管加 1 个电容），所以容量小得多。Groq、Cerebras、d-Matrix 都走这条路。

一个有用的比值是主存读遍/秒 = 带宽 / 容量：每秒能把整块主存从头到尾读几遍。如果一个模型恰好塞满主存，batch = 1 时每个用户每秒最多生成这么多个 token。代码 04 演示 7（Llama-3.1-70B，FP8 权重 70.6 GB，batch = 1）：

（Groq 3 LPU 的数字是 NVIDIA 公布的托盘规格除以 16 颗芯片；Corsair 是 Gimlet 博客给的每卡数字。）

• HBM 芯片每秒只能把显存读几十遍；SRAM 芯片能读几万到几十万遍。
• 但 SRAM 很小：放一个 70B 的 FP8 模型要一百多块 Groq 3 LPU。瓶颈从”带宽”换成了”容量”（要多少芯片、多少钱）和”互连”（跨这么多芯片的延迟）。
• “一步下界”那一列用了一个粗略的假设：每一层的激活要从一块芯片传到下一块，每次至少 1 µs（How to Scale Your Model 给 TPU 的每跳延迟估计），80 层就是 80 µs。在这个假设下，SRAM 方案的一步时间由通信决定，而不是由读权重决定。这和 Ma & Patterson 的判断一致：LLM 推理的主要挑战是内存和互连（尤其是互连延迟），而不是算力。

所以没有一种芯片在所有阶段都最好。§6 会接着讲这对”推理系统该怎么搭”意味着什么。

5. 案例：读懂 Blackwell Ultra · Case study: co-design in Blackwell Ultra

依据 NVIDIA 技术博客 Inside NVIDIA Blackwell Ultra（2025-08，2025-09-24 更新）。代码：05_codesign_blackwell_ultra.py。厂商博客是营销材料，这一节的目的正是练习用 roofline 和能量的眼光去读它：每个设计决定在补哪块短板？哪些数字要打折？

5.1 先看规格表

博客表 2 的主要数字（Blackwell 列对应 GB200 中的 B200）：

每个 SM 有 128 个 CUDA core、4 个第五代 Tensor Core、256 KB 的 TMEM 和若干特殊函数单元（SFU），全芯片 160 个 SM。

5.2 为什么是两块 die

芯片制造用光刻机把电路”印”在硅片上，一次曝光的最大面积（reticle limit，光罩极限）大约是 26 mm × 33 mm ≈ 858 mm²。H100 已经接近这个极限。想要更多晶体管，又不能做更大的单片 die，怎么办？

Blackwell 的答案是两块光罩极限大小的 die，用一种叫 NV-HBI 的芯片间接口连起来，带宽 10 TB/s，对软件来说表现为一块 GPU（共享 L2 cache、内存访问完全一致）。总晶体管数 2080 亿，是 Hopper 的 2.6 倍。

为什么 10 TB/s 很重要？因为如果两块 die 之间的带宽不够，它们就会像两块独立的 GPU，程序员要自己处理”数据在哪块 die 上”。10 TB/s 比 HBM 带宽（8 TB/s）还高，让”一块 GPU”的抽象成立。这是一个为了保持编程模型不变而付出的硬件代价——CUDA 的生态（二十年的软件）是 NVIDIA 最大的资产之一。

5.3 算力 vs 带宽：内存墙

把四代芯片并排（代码 05 演示 1，相对 H100）：

两件事：

1. “7.5 倍”里有一大半来自换了数值格式。博客说 Blackwell Ultra 的 NVFP4 算力是 Hopper 的 7.5 倍——但 Hopper 不支持 FP4，这是拿 NVFP4 和 FP8 比。同为 FP8，只有 2.5 倍。读厂商数字时，第一件事是确认比的是不是同一种精度。
2. 算力涨得比带宽快得多。从 H100 到 Blackwell Ultra，算力（按厂商口径）涨了 7.6 倍，带宽只涨了 2.4 倍，屋脊点从 591 涨到 1,875。芯片越来越”算力过剩、带宽不足”——这就是常说的内存墙（memory wall）。Ma & Patterson 给过一个更长期的数字：2012–2022 年 NVIDIA GPU 的 64 位算力涨了 80 倍，带宽只涨了 17 倍。

为什么带宽涨得慢？Gimlet 的 SRAM 博客给了一个几何上的解释：片外内存的带宽受芯片周长的限制（信号要从芯片边缘的引脚出去），而算力随芯片面积增长。面积长得比周长快，所以每单位算力分到的片外带宽越来越少。HBM 靠把内存颗粒堆叠起来、用几千根极短的连线和芯片相连来缓解这个问题：Ma & Patterson 的表 1 显示，从 HBM3 到 HBM4，每个堆栈的引脚数从 1024 翻倍到 2048。这正是 Rubin 带宽一下涨到 22 TB/s 的原因——当 decode 受限于带宽，下一代最该补的就是带宽。

5.4 NVFP4：两级缩放

4 位浮点数（FP4，格式 E2M1：1 位符号、2 位指数、1 位尾数）只能表示 ±{0, 0.5, 1, 1.5, 2, 3, 4, 6} 这 15 个数。要用它存一组任意大小的权重，必须先除以一个缩放因子（scale），把数值范围”挪”到 [−6, 6] 里。问题是缩放因子要按最大的那个数来选：如果一组数里有几个特别大的”离群值”（outlier），缩放因子就会被拉得很大，其余的普通数除完之后都挤在 0 附近，被舍入成 0。

解决办法是块缩放（block scaling）：每一小块数用自己的缩放因子。博客描述 NVFP4 用两级缩放：

• 每 16 个数共享一个 FP8（E4M3）格式的缩放因子；
• 整个张量再有一个 FP32 的缩放因子（把块缩放因子挪到 E4M3 能表示的范围里）。

平均每个数 $4 + 8/16 = 4.5$ bit。另一个常见格式 MXFP4（OCP 的 Microscaling 规范）是每 32 个数一个缩放因子，而且缩放因子只能是 2 的幂（E8M0），平均 4.25 bit。

代码 05 演示 2 真的用 numpy 把 100 万个”像权重的”随机数（Student-t 分布，有重尾和离群值）量化了一遍：

整个张量一个缩放因子的 FP4，82.5% 的数被量化成了 0，基本不能用。块缩放把误差降到 9–13%：离群值只”连累”同一块里的十几个邻居。NVFP4 的块更小、缩放因子更精细，比 MXFP4 再好一些，代价是每个数多 0.25 bit。博客说 NVFP4 “markedly lower error than standard FP4”——在这组数据上可以复现这个方向。

两个提醒：

• 博客说 NVFP4 的模型精度”通常和 FP8 相差不到约 1%“——那是模型任务上的精度，不是这里的逐元素误差。逐元素误差 9% 和任务精度差 1% 之间是什么关系，要靠第 3 讲的实验来理解。注意，误差不会在点积里抵消：P3 §6.6 验证了矩阵乘法输出的相对误差和权重的相对误差几乎一样（9.23% → 9.25%）；差别来自模型本身对噪声的容忍度，以及不同层、不同张量的敏感度很不一样。
• 存储账：3000 亿参数 × 4.5 bit = 169 GB，放进 288 GB 还剩 119 GB 给 KV cache。这就是博客所说”300B+ 参数的模型可以完整放在一块 GPU 上”的来历。

5.5 被遗忘的单元：softmax 的指数运算

博客里有一个看起来不起眼的改进：“Accelerated softmax in the attention layer”——把特殊函数单元（SFU）计算指数 exp 的吞吐翻了一倍。为什么要专门改这个？

注意力的每个分数 $s = q \cdot k$ 都要做一次 exp（softmax），同时伴随两次长度为 $d_\text{head}$ 的点积（$QK^\top$ 一次，$PV$ 一次），共 $4 \cdot d_\text{head}$ 次运算。点积在 Tensor Core 上做，exp 在 SFU 上做。两者的速度之比决定了谁是瓶颈：

代码 05 演示 3（大于 100% 表示 exp 是瓶颈）：

故事就在这张表里：Blackwell 的 Tensor Core 比 Hopper 快了约 2.5 倍，exp 吞吐却只从 4.5 涨到 5 T/s。FP8 注意力、$d_\text{head} = 128$ 时，exp 需要的时间是矩阵乘法的约 2 倍——注意力的速度被 SFU 卡住了，再快的 Tensor Core 也没用。Blackwell Ultra 把 exp 吞吐翻倍到 10.7 T/s，正好把这个比值拉回 100% 以下。

这是 Amdahl 定律的一个活例子：一个系统的速度受限于最慢的那一环；把一部分加速得越多，剩下没加速的部分就越显眼。一个原本”不重要”的单元，因为别的单元变快了，就成了瓶颈。

软件也在同一个瓶颈上使劲：据公开介绍，FlashAttention-4 等面向 Blackwell 的注意力内核把一部分 exp 改用普通 FMA 单元上的多项式近似来算，把负载从 SFU 上分出去；NVIDIA 的 Rubin 博客则说 Rubin 继续把指数吞吐提高到 Blackwell 的 2 倍（FP32）和 4 倍（BF16/FP16）。硬件和软件一起围着同一个瓶颈转——这就是协同设计。

还有一个细节：$d_\text{head}$ 越小，每次 exp 配的矩阵乘法越少，越容易被 exp 卡住（$d_\text{head} = 64$ 在三代芯片的 FP8 下都超过 100%）。模型结构的一个”小参数”，直接决定了硬件的哪个单元成为瓶颈。

（这张表只比较”exp 吞吐 vs 矩阵乘法吞吐”；softmax 还要求最大值、减法、乘法等普通运算，没有计入，所以结论偏乐观。）

5.6 288 GB：为 KV cache 而生

§4.4 已经说明了容量的意义。博客列出的好处正是”三件事”：完整放下 3000 亿以上参数的模型、更长的上下文不用把 KV cache 换出到 CPU、更高的并发。博客正文给出的 HBM 配置是”8 个 12 层（12-Hi）堆栈、16 个 512 位控制器（总位宽 8192 位）“，这组数字是自洽的：每个 12 层堆栈 36 GB，8 个正好 288 GB；每个堆栈 1024 位接口，8 个正好 8192 位。可是文末的更新说明又说已经把图和这一节改成”proper 12 HBM stacks instead of 8”，而正文并没有变——“12”也许指的是 12 层。厂商材料也会有含糊和前后不一致的地方，引用时要回到能自洽验证的数字上。

5.7 NVLink 5 与 NVL72：把 72 块 GPU 当一块用

NVLink 5 每块 GPU 1.8 TB/s 双向（18 条链路 × 100 GB/s），是 Hopper 的 2 倍。GB300 NVL72 把 72 块 GPU（36 个 Grace Blackwell Ultra 超级芯片）用 NVLink Switch 全互连在一个机柜里，聚合带宽 130 TB/s，最多可以扩展到 576 块 GPU 的无阻塞网络。

为什么推理需要这么大的”scale-up 域”？回到 §3.10 的网络 roofline：模型切得越细（比如把一个万亿参数的 MoE 模型切到几十块 GPU 上），芯片间通信越多；只有在高带宽、低延迟的 NVLink 域内，这种切分才不会被通信拖垮。第 4 讲会算具体的账。

5.8 怎样读厂商的数字

最后，把这一节的经验总结成一张检查清单。看到”新硬件/新系统快了 X 倍”时，问：

1. 同一种精度吗？ NVFP4 vs FP8 天然差 2–3 倍。
2. 稠密还是稀疏？ “with sparsity”的数字通常是稠密的 2 倍，而且只在权重满足 2:4 结构化稀疏时才能达到（How to Scale Your Model 练习 5 专门提醒了这一点）。
3. 单芯片还是整个系统？ “GB300 NVL72 比 Hopper 平台的 AI 工厂产出高 50 倍”这类数字是机柜级的，把芯片、互连、软件、精度的提升全乘在一起。
4. 在什么延迟下？ 吞吐要和延迟一起报，最好是整条前沿（§4.5）。
5. “up to”是什么意思？ 最好情况，通常是特定模型、特定配置。Rubin 的 “up to 50 PFLOPS NVFP4” 就没有说明是否稠密。
6. 峰值还是实测？ 峰值算力在真实矩阵乘法里通常只能跑到 80% 左右，memory-bound 的部分更远。

这不是说厂商在骗人——这些数字通常都有据可查——而是说数字的含义需要你自己弄清楚。roofline 给了你一把尺子。

6. 异构与解耦：没有一种芯片适合所有阶段 · Heterogeneous, disaggregated inference

这一节是对课程主题”disaggregated AI inference systems on heterogeneous machines”的预告，依据 Gimlet Labs 的两篇博客和 NVIDIA 的 Groq 3 LPX 博客。第 5 讲会深入。

6.1 一个请求里的不同”性格”

§4 已经看到，同一个 LLM 请求里藏着好几种完全不同的计算：

Gimlet 的博客引用了 Wang et al. 的测量：在一个 8B 的 Llama 模型上，prefill 和 decode 的各个内核在 roofline 图上散落在完全不同的位置，甚至同一个阶段里的不同算子强度也不同。在一种硬件上跑所有阶段，必然有一些阶段跑在”不合适”的硬件上。

6.2 近存 vs 远存

Gimlet 的 SRAM 博客把问题归结为一个更深的取舍：近存（near-compute memory）还是远存（far-compute memory）。

• SRAM：每 bit 6 个晶体管，和计算单元用同一种工艺做在同一块芯片上，读取约 1 ns，带宽随芯片面积增长；但密度低，容量受芯片面积限制，而且每一平方毫米给了 SRAM 就不能给计算单元（零和）。
• DRAM（HBM）：每 bit 1 个晶体管加 1 个电容，要用专门的工艺、做在芯片外面，读取约 10–15 ns；密度高、容量大；但带宽受芯片周长限制。

GPU 的办法是多级存储层次：HBM 放全部数据，L2（H100 50 MB、B200 126 MB）、共享内存、寄存器层层缓存，靠数据复用（分块）减少对 HBM 的访问。这在复用率高的计算上非常有效（训练、prefill、大批量）。SRAM 芯片的办法是把整个工作集放在片上（Groq 3 LPU 每颗 500 MB、d-Matrix Corsair 每卡 2 GB、Cerebras 的整片晶圆 44 GB），不需要复用也能跑得快——正好适合复用率低的计算（小批量 decode、小模型）。

Gimlet 的结论值得记住：工作集大小决定算术强度，算术强度决定该用近存还是远存。“SRAM vs HBM”只是这个更深层决定的一种实现。

6.3 几种解耦的切法

“解耦”（disaggregation）就是把一个推理工作负载切开，不同部分放到不同的机器上。目前有几种切法：

1. prefill / decode 解耦：prefill 和 decode 放在不同的 GPU 池里，互不干扰（DistServe、Splitwise，第 5 讲）。NVIDIA 的 Rubin CPX 是一款面向 prefill 的芯片，据报道用 GDDR7 而不是 HBM——prefill 是 compute-bound 的，用不着 HBM 那么高的带宽。
2. 注意力 / FFN 解耦：NVIDIA 的 Groq 3 LPX 博客描述了这样一种方案：Rubin GPU 负责 prefill 和 decode 的注意力（需要大容量放 KV cache），Groq 3 LPU 负责 decode 中对延迟敏感的 FFN 和 MoE 专家计算（权重放在 SRAM 里，读得极快），由 NVIDIA Dynamo 在两种后端之间调度。NVIDIA 宣称这一组合相比 GB200 NVL72 有最高 35 倍的每兆瓦吞吐（厂商数字，见 §5.8 的检查清单）。
3. 把投机解码的草稿模型放到 SRAM 芯片上：Gimlet 用 gpt-oss-120b 作为目标模型、一个 16 亿参数的草稿模型放在两张 d-Matrix Corsair 卡上，prefill 和验证放在 GPU 上。它报告在相同能效（每千瓦吞吐）下，端到端请求延迟比”草稿模型也放在 GPU 上”快 2–10 倍（结果是测量和建模结合得到的，草稿接受率按代码场景设定）。第 5、6 讲会讲投机解码本身和这个实验。

6.4 解耦的代价

解耦不是免费的。把阶段放到不同的机器上，就要在机器之间搬数据：prefill 算出的 KV cache 要传给 decode 机器，注意力的输出要传给 FFN 机器，每一步都要传。回到 §2 的能量视角和 §3.10 的网络 roofline：

• 搬的数据量和带宽决定了时间开销（通信能不能和计算重叠？）；
• 跨芯片、跨机器的搬运是能量表里最贵的那类操作；
• 调度变复杂了：要决定每个请求、每个阶段去哪台机器，还要处理负载不均、故障、尾延迟。

所以解耦值不值得，要算账：换到更合适的硬件上省下的时间和能量，要大于搬运和协调的代价。这正是这门课要训练的能力。

6.5 Ma & Patterson 的四个研究方向

Ma & Patterson 认为 LLM decode 的主要瓶颈是内存容量、内存带宽和互连延迟，而不是算力，并提出了四个硬件研究方向：

1. 高带宽闪存（High Bandwidth Flash, HBF）：像 HBM 那样堆叠闪存颗粒，容量约是 HBM 的 10 倍，带宽相近；但写入寿命有限、读取延迟是微秒级，所以只适合放不常改动的数据（推理时的权重、变化很慢的上下文语料），不能放 KV cache。
2. 近存计算（Processing-Near-Memory）：把计算单元放到内存颗粒旁边。
3. 3D 计算-逻辑堆叠：把 DRAM 直接堆在计算芯片上，获得比 HBM 更高的带宽（d-Matrix 的下一代芯片 Raptor 据称走这条路）。
4. 低延迟互连：推理的通信消息小而频繁，延迟比带宽更重要；而为训练设计的超级计算机网络主要优化的是带宽。

这四个方向都在攻 §4 里看到的那几堵墙。

7. 全栈地图：这门课怎么走 · The course map

7.1 每一层都在和 roofline 打交道

7.2 课程作业与本讲义代码

原课的四次作业正好对应栈的几层：

我们没有 GPU，所以本讲义的代码遵循一个原则：正确性用真实计算验证（numpy 实现，结果要和参考实现一致），性能用模型估计（这一讲的 roofline、第 9 讲的循环嵌套能量模型），能在 CPU 上实测的就实测（像代码 03 那样，只看趋势）。

7.3 本讲义的路线图

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

截至 2026 年秋天，几条值得关注的动向（都是厂商或第三方的公开说法，我没有独立验证）：

• NVIDIA Rubin（2026 年公布）：HBM4，每 GPU 最高 288 GB、22 TB/s（Blackwell 的 2.75 倍）；“up to” 50 PFLOPS NVFP4；NVLink 6 每 GPU 3.6 TB/s；指数吞吐是 Blackwell 的 2 倍（FP32）/ 4 倍（BF16）；还在注意力里引入了把中间分数压缩成 2:4 稀疏的硬件支持。带宽涨幅远大于前一代——印证了”decode 受限于带宽”。
• SRAM 芯片进入主流：据 Gimlet 的博客，NVIDIA 在 2025 年 12 月以约 200 亿美元授权了 Groq 的技术，Cerebras 宣布了一份为 OpenAI 提供 750 MW 推理算力的协议。NVIDIA 随后公布了 Groq 3 LPX（256 颗 LPU 的机柜，128 GB SRAM、40 PB/s 的片上带宽），定位为 Vera Rubin 平台中专门负责低延迟 decode 的部分。
• 异构推理云：Gimlet Labs 等公司把一个模型的不同阶段放到不同厂商的硬件上运行；NVIDIA 也在自家产品线里做异构（Rubin CPX 做 prefill、LPX 做低延迟 decode）。
• Google TPU7x（Ironwood）：How to Scale Your Model 的表格给出每芯片 192 GB HBM、7.4 TB/s、约 4.6 PFLOPS（int8/FP8）。
• 低精度成为默认：NVFP4 / MXFP4 在 SGLang、TensorRT-LLM、vLLM 中都有支持；4 位权重 + 8 位或 4 位激活的推理越来越常见（第 3 讲）。
• 内存越来越贵：Ma & Patterson 指出，2023–2025 年 HBM 每 GB 容量和每 GB/s 带宽的价格都上涨了约 1.35 倍，而普通 DDR 内存在下降；他们认为这种分化会长期存在。
• “每兆瓦 token 数”成为衡量标准：电力成为数据中心扩张的硬约束，厂商开始用 TPS/MW 而不是单纯的 FLOPS 来宣传。

9. 小结 · Summary

1. 推理是持续的开销，和用量成正比；推理模型、智能体、长上下文让每个请求的 token 数和内存需求不断增长。延迟、吞吐、成本、能耗四个目标互相牵制，只有前沿，没有最优点。
2. 能量视角：Dennard scaling 在 2005 年前后结束，芯片功耗密度不再随工艺下降，功耗预算成了硬约束，性能 = 能效 × 功率。一次运算只要零点几到几个 pJ，一次 DRAM 读取要上千 pJ，一条指令的管理开销 70 pJ。所以提高能效的两条主路是专用化（摊薄控制开销）和数据复用（少从远处搬），低精度则同时让运算和搬运都变便宜。
3. roofline：时间下界 $= \max(T_\text{math}, T_\text{comms})$；比较算术强度 $I$ 和临界强度 $I^* = \pi / \beta$ 就知道瓶颈。逐元素运算 $I \approx$ 常数（总是 memory-bound），矩阵乘法 $I \approx B$（要几百个 token 才 compute-bound），“每个样本一个矩阵”$I \approx 1$（批处理救不了）。屋顶不止一层：精度、计算单元、存储层级、网络各有各的屋顶；容量是硬约束。
4. LLM 推理：decode 每一步都要读全部权重和每个序列的 KV cache。批处理摊薄了权重读取，却摊不薄 KV 读取；上下文超过几百个 token 后，decode 批再大也到不了 compute-bound；显存容量又把批大小压在低效区。prefill 正好相反，是 compute-bound 的。
5. 协同设计：Blackwell Ultra 的每个改进都能用这套工具解释——NVFP4 用块缩放换来可用的 4 位精度、288 GB HBM 为 KV cache 腾地方、加倍的 exp 吞吐修补 Tensor Core 变快后暴露的 softmax 瓶颈、NVLink 扩大 scale-up 域。算力增长快于带宽和互连，是长期趋势（内存墙）。
6. 异构与解耦：一个请求里不同阶段的强度和瓶颈不同；近存（SRAM）带宽高、容量小，远存（HBM）容量大、带宽受限；把不同阶段放到合适的硬件上，代价是搬数据和协调。

10. 常见误区 · Common pitfalls

1. “LLM 计算量大，所以推理慢是因为算力不够。” decode 时 H100 的算力用了不到 1%，时间几乎全花在读权重和 KV 上。先算强度，再下结论。
2. “换成低精度，瓶颈就会变。” 如果权重、激活、算力都换成同一种低精度，强度和屋脊点同比例变化，瓶颈不变；只有”仅权重量化”这种不对称的做法才会改变临界 batch。（NVFP4 在 Blackwell Ultra 上是个例外，因为它的算力涨得比字节省得多。）
3. “批大小 $B$ 指的是序列数。” roofline 里的 $B$ 是这一步一起处理的 token 数。decode 时恰好等于序列数，prefill 时是提示长度之和。
4. “roofline 算出的时间就是实际时间。” 它是下界：峰值通常只能跑到 80% 左右；而且内核写得不好时（比如代码 03 里 numpy 的 GEMM 处理瘦矩阵），可能连下界的一半都到不了。
5. “KV cache 每 token 才几百 KiB，可以忽略。” 长上下文时它比权重还大，而且是 decode 到不了 compute-bound 的根本原因。
6. “芯片越多越快，8 块芯片就快 8 倍。” 切分要通信。网络 roofline 说明通信量和计算量的比值只和被切开的维度有关；小批量时通信延迟（不只是带宽）往往是主要开销。
7. “屋脊点越高的芯片越先进。” 屋脊点是算力和带宽的比值。屋脊点越高，越难喂饱；对 memory-bound 的 decode 来说，带宽本身更重要。H200 的屋脊点比 H100 低，decode 却更快。
8. “SRAM 芯片比 GPU 快，所以会取代 GPU。” SRAM 芯片在带宽上占优，但容量小、需要很多芯片、互连延迟成为新瓶颈；GPU 在复用率高的计算（prefill、训练、大批量）上占优。更可能的结局是两者分工。
9. “厂商说快了 7.5 倍，那就是 7.5 倍。” 先问：同一种精度吗？稠密还是稀疏？单芯片还是整个机柜？在什么延迟下？（§5.8）
10. “能耗只要看芯片的功耗就行。” 每个 token 的能耗 = 功率 × 时间 / 这一步产出的 token 数，强烈依赖批大小（$B = 1$ 和 $B = 256$ 差 40 倍以上）。而且按最大功耗算是偏保守的。

11. 练习 · Exercises

带 💻 的题可以用本讲的模块验证。参考答案在每题后面（先自己做再看）。

练习 1（能量账）💻。在 45 nm 工艺下，读一遍 Llama-3.1-8B 的 BF16 权重（16.06 GB）要多少能量（按 DRAM 每 64 bit 1.3–2.6 nJ）？生成一个 token 的矩阵乘法（约 75 亿次乘加，按 FP16 乘加 1.5 pJ）要多少能量？batch 至少要多大，运算能量才追上读权重的能量？

读权重：$16.06 \times 10^9 / 8 \times 1.3 \text{ nJ} \approx 2.6$ J（高值 5.2 J）。运算：$7.5 \times 10^9 \times 1.5 \text{ pJ} \approx 0.011$ J。比值约 232（用低值）：batch 要约 230 以上，运算的能量才和读权重的能量相当。有意思的是，这和 H100 的”时间”屋脊点（约 295）是同一个量级——硬件设计者在同时平衡时间账和能量账。batch = 1 时，99.5% 以上的能量花在读权重上。

练习 2（逐元素运算）。RMSNorm 对一个 $[B, d]$ 的 BF16 张量做：平方、求和、开方取倒数、逐元素相乘再乘以权重。大约每个元素 4 次运算，读 2 字节、写 2 字节。它的强度是多少？在 H100 上 $B = 64$、$d = 4096$ 时，roofline 下界是多少？你觉得实际时间会被什么主导？

强度约 4 / 4 = 1 FLOP/byte，远低于 295，memory-bound。字节数 $64 \times 4096 \times 4 \approx 1.05$ MB，下界约 $1.05 \times 10^6 / 3.35 \times 10^{12} \approx 0.31$ µs。这么短的时间，实际会被启动一个内核的固定开销（通常几微秒）主导。这正是要把 RMSNorm 和前后的运算融合到一个内核里的原因（第 11、12 讲）。

练习 3（仅权重量化）💻。在 H100 上，把 $W[8192, 8192]$ 存成 int8、激活和计算仍用 BF16，用 roofline.critical_batch 求精确的临界 $B$，并和全 BF16（319）比较。换成 H200 呢？

critical_batch(8192, 8192, H100.critical_intensity("bf16"), x_bytes=2, w_bytes=1, out_bytes=2) = 160，约为全 BF16 的一半。H200（屋脊点 206）是 109。仅权重量化让 decode 更早进入 compute-bound；更重要的是，在 memory-bound 区它直接把读权重的时间减半。

练习 4（decode 时间）💻。Llama-3.1-70B，FP8 权重，一块 Blackwell Ultra，$B = 32$，ctx = 8192，BF16 KV。一步 decode 的下界是多少？瓶颈是什么？每个用户每秒约几个 token？

decode_step(LLAMA31_70B, B300, 32, 8192, w_prec="fp8")：读权重 8.69 ms，读 KV 10.74 ms，计算 1.03 ms，下界 19.43 ms，memory-bound（而且 KV 已经比权重读得多）。每个用户约 1 / 0.0194 ≈ 51 tok/s。把 KV 也存成 FP8，读 KV 减半，下界降到约 14 ms。

练习 5（容量）💻。Llama-3.1-8B（BF16 权重）放在一块 H200（141 GB）上，ctx = 32K，预留 10%。最多能放几个序列？把 KV cache 改成 FP8（kv_elem_bytes=1）呢？

max_batch(LLAMA31_8B, H200, 32768, "bf16", reserve_frac=0.1) = 25；FP8 KV 时是 51，翻了一倍。在长上下文下，KV 的精度对并发数的影响比权重的精度大得多。

练习 6（GQA 的价值）💻。假设 Llama-3.1-8B 不用 GQA，而用普通多头注意力（32 组 KV 头）。用 ModelConfig("8B-MHA", 32, 4096, 32, 32, 128, 14336, 128256) 算：参数量、每 token KV、一块 H100 在 ctx = 8K 时最多放几个序列、$B = 32$ 时一步 decode 的下界，并和原版比较。

参数量 8.84 B（多了 K、V 的投影矩阵），每 token KV 512 KiB（原版 128 KiB，4 倍）。ctx = 8K 时最多放 12 个序列（原版 52）；$B = 32$ 时一步下界 46.0 ms（原版 14.7 ms）。GQA 几乎不影响模型质量，却让 decode 的 KV 读取和占用都降到 1/4——这就是它成为标配的原因。

练习 7（网络 roofline）💻。用 sharded_matmul_two_chips 计算：把 Llama-3.1-8B 的 MLP 上投影 $X[64, 4096] \cdot W[4096, 14336]$ 沿 $D = 4096$ 切到两块 H100 上（NVLink 450 GB/s 单向）。是计算受限还是通信受限？如果换成 $D = 16384$ 的大模型呢？

网络强度 $D/2 = 2048$，H100 的网络临界强度 2198，略低：通信受限（计算 3.80 µs，通信 4.08 µs）。$D = 16384$ 时网络强度 8192，远高于 2198，计算受限。这说明小模型不值得切太细。（注意：这里是 How to Scale Your Model 的玩具切法，真实的张量并行切法不同，第 4 讲会讲。）

练习 8（softmax 瓶颈）💻。某个模型用 $d_\text{head} = 96$。在 Blackwell Ultra 上，FP8 注意力的 exp 会不会成为瓶颈？BF16 呢？

每次 exp 配 $4 \times 96 = 384$ 次运算。FP8：需要 $5 \times 10^{15} / 384 \approx 13.0 \times 10^{12}$ 次 exp/s，而 SFU 只有 $10.7 \times 10^{12}$，比值 122%，exp 是瓶颈。BF16：61%，不是。

练习 9（读厂商数字）。博客说 Blackwell Ultra 的稠密 NVFP4 算力是”Hopper H100 和 H200 的 7.5 倍”。(a) 这个比较的两边分别是什么精度？(b) 同一种精度下是多少倍？(c) 博客表 2 里 Blackwell Ultra 的 NVFP4 稀疏算力（20 PFLOPS）只比稠密（15 PFLOPS）高 1.33 倍，而 FP8 稀疏是稠密的 2 倍。你能想到什么解释？

(a) Blackwell Ultra 的 NVFP4（15 PFLOPS）对比 Hopper 的 FP8（2 PFLOPS），因为 Hopper 不支持 FP4。(b) 同为 FP8，5 vs 2 PFLOPS，2.5 倍。(c) 一种可能的解释：Blackwell Ultra 主要提升的是稠密 NVFP4 的吞吐（从 10 提到 15），而稀疏路径没有同比例提升（Blackwell 是 10 / 20，Ultra 是 15 / 20）——稀疏的峰值可能受限于别的部分（比如数据通路或稀疏元数据的处理）。博客没有解释，这是推测。教训是：同一张表里不同列的比例关系不一定一样，不要想当然地乘 2。

练习 10（你的电脑）💻。运行 03_measure_your_machine.py，记下你机器的屋脊点。如果在这台电脑上用 FP32 权重跑 8B 模型的 decode（只考虑权重，FP32 矩阵乘法强度 ≈ $B/2$），$B$ 至少多大才能 compute-bound？实测的曲线支持这个预测吗？

编写本讲时的 M4 Pro：屋脊点约 17，所以 $B \gtrsim 2 \times 17 \approx 35$（用 critical_batch(4096, 14336, 17.4, 4, 4, 4) 精确算是 36）。但实测曲线在 $B = 128$ 才达到峰值的约 60%：因为 BLAS 的 GEMM 处理瘦矩阵时效率不高（§3.12），实际达到的带宽和算力都低于 roofline 的上界。roofline 告诉你”最多能多快”，不告诉你”代码写得好不好”。

练习 11（思考题）。H200 和 H100 的计算芯片相同，H200 的 HBM 带宽高 43%、容量大 76%。对 prefill、小批量 decode、长上下文大批量 decode 三种场景，分别预测 H200 能快多少，并说明理由。

prefill 是 compute-bound，算力不变，几乎不快。小批量 decode 受限于读权重，约快 1.43 倍（带宽之比）。长上下文大批量 decode 受限于读 KV 和容量：带宽带来约 1.43 倍，更大的显存还能放下更多序列、让批更大，每芯片吞吐的提升可能超过 1.43 倍。同一块”更好的”芯片，对不同阶段的帮助完全不同。

12. 延伸阅读 · Further reading

本讲阅读材料

• J. Austin, S. Douglas, R. Frostig, A. Levskaya, et al., How to Scale Your Model, Google DeepMind, 2025. https://jax-ml.github.io/scaling-book/ —— 第 0、1 章是本讲阅读；第 2 章（TPU）、第 3 章（切分矩阵）、第 7 章（推理）是第 4 讲的阅读；第 12 章（GPU）适合配合第 8 讲。
• NVIDIA, Inside NVIDIA Blackwell Ultra: The Chip Powering the AI Factory Era, 2025. https://developer.nvidia.com/blog/inside-nvidia-blackwell-ultra-the-chip-powering-the-ai-factory-era/

本讲补充材料

• M. Horowitz, Computing’s Energy Problem (and what we can do about it), ISSCC 2014 Plenary. —— §2 的能量表和”专用化”论证的出处。
• X. Ma and D. Patterson, Challenges and Research Directions for Large Language Model Inference Hardware, IEEE Computer, 2026（arXiv:2601.05047）。—— 推理为什么是”内存和互连”问题，以及四个硬件方向。
• S. Williams, A. Waterman, D. Patterson, Roofline: An Insightful Visual Performance Model for Multicore Architectures, Communications of the ACM, 2009. —— roofline 模型的原始论文。
• J. Hennessy and D. Patterson, A New Golden Age for Computer Architecture, Communications of the ACM, 2019. —— Dennard scaling 结束后，为什么领域专用架构（DSA）是出路。
• A. Gholami et al., AI and Memory Wall, IEEE Micro, 2024. —— 算力、内存带宽、互连带宽的长期增长速度对比。
• Gimlet Labs 博客：The emerging role of SRAM-centric chips in AI inference；Low-Latency Inference with Speculative Decoding on d-Matrix Corsair and GPU（2026）。—— §4.7、§6 的来源。
• NVIDIA 技术博客：Inside NVIDIA Rubin GPU Architecture；Inside NVIDIA Groq 3 LPX（2026）。

本仓库的相关讲义

• CS349D 第 1 讲：从系统角度讲推理的指标、批处理和 Clockwork；和本讲的 §1、§4 互补。
• CS349D 前置讲 P1：GPU 的组织、分块与复用、在 CPU 上测 roofline。

13. 下一讲预告 · Next lecture

第 2 讲（Fred Kjolstad）讲 Transformer。阅读材料是两篇经典：word2vec（Efficient Estimation of Word Representations in Vector Space）和 Attention Is All You Need。我们会从”怎样把一个词变成向量”讲起，一步步搭出 decoder-only Transformer，然后用一个能真正运行的 numpy Transformer 逐个矩阵验证本讲 llm_math.py 里的参数量、运算量和 KV cache 公式，实现带 KV cache 的 decode 循环（原课作业 1 的主题），并比较 MHA、GQA、MQA 在推理时的账本。这一讲给出的”decode 受限于读权重和 KV”的结论，会在那里变成一行行可以数的代码。