# 第 1 讲 · LLM 导论与 LLM 系统的挑战 · Introduction to LLMs & the Challenges of LLM Systems

> **对应 Maps to**：CMU 11-868 / 11-968 *Large Language Model Systems*（Spring 2026，Lei Li），Lecture 1 · Mon 1/12 · **Introduction to LLM**。
> **阅读材料**：原课第 1 讲没有指定阅读材料。本讲涉及的公式出自：Brown et al. (2020) *Language Models are Few-Shot Learners*（GPT-3）；Kaplan et al. (2020) *Scaling Laws for Neural Language Models*；Hoffmann et al. (2022) *Training Compute-Optimal Large Language Models*（Chinchilla）；Rajbhandari et al. (2020) *ZeRO*（16 字节/参数的记账）；Korthikanti et al. (2022) *Reducing Activation Recomputation in Large Transformer Models*（激活值公式）。
> **依据说明**：本讲以原课第 1 讲公开的 slides（llmsys-01-intro.pdf，51 页）为主线：规模曲线与学习目标、LLM 的能力、数学基础（下一个 token 的概率、三类语言模型、学习框架、交叉熵与困惑度、评估指标）、LLM 系统的挑战（核心问题、抽象与权衡、底层算子、不同层次的挑战、协同设计、计算 vs 数据搬运、编程模型）和课程安排。slides 把"训练一个 100B 模型要多少资源"列为学习目标，但没有在第 1 讲里算出来；§5 的完整估算、§0.6 的"城市语言"和全部实验都是本讲义自己设计的，**不代表课堂上讲了什么**。
> **配套代码 Code**：[`code/lecture01/`](../code/lecture01/)（5 个脚本 + 5 个模块；`01` 只用标准库，其余用 numpy 和 PyTorch CPU 版；`03` 第一次约 75 秒，其余每个脚本 1–5 秒）
> **预计阅读时间**：4–5 小时（含运行代码）

---

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

### 0.1 学习目标 Learning goals

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

1. 说出 LLM 规模增长的大致轨迹，并解释为什么"规模"让系统问题变成了 LLM 研究的核心问题之一。
2. 从**系统负载**的角度看 LLM 的各种能力：翻译、对话、少样本提示、推理、写代码并调用工具，分别对输入长度、输出长度、调用次数和时延提出了什么要求。
3. 写出语言模型的数学定义（下一个 token 的条件概率、链式法则），区分 **encoder-only、encoder-decoder、decoder-only** 三类模型，并能用实验说明：它们的差别只在注意力掩码，系统代价却天差地别（为什么只有因果模型能用 **KV 缓存**）。
4. 解释**交叉熵损失**和**困惑度**，知道"理论最优困惑度"的含义；能区分 n-gram 的**近似误差**和**估计误差**，理解神经语言模型为什么能同时利用长上下文和有限数据。
5. 说出 LLM 系统的**核心问题**，把一个 Transformer 层拆成**矩阵乘法、逐元素映射、归约、数据搬运**四类底层运算，并用**算术强度**判断一个运算是计算受限还是访存受限。
6. 用几条公式（参数量 ≈ 12·L·d²、训练 FLOPs ≈ 6ND、训练显存 ≈ 16 字节 × N、KV 缓存 = 2·L·头数·头维度·字节数）**估算**训练和服务一个 100B 模型所需的 GPU、时间、钱、显存和带宽，并知道每条公式在什么情况下不准。

### 0.2 先修与本讲范围 Prerequisites & scope

原课的先修是线性代数、微积分、概率统计，会 Python 和 C / C++ / Java 之一，最好学过机器学习（不强制）。本讲是导论，需要的背景都就地补上，细节留给后面的讲次。

| 默认你已经会 Assumed | 本讲只给直觉，后面再展开 Previewed only |
|---|---|
| 条件概率、期望、对数 | GPU 的结构与 CUDA 编程（第 2–4 讲） |
| 矩阵乘法；导数与链式法则 | 自动微分与深度学习框架（第 5 讲） |
| Python 基础（不熟悉的写法代码注释里都有解释） | Transformer 的每个部件（第 6 讲）；LLaMA 式结构、GQA（第 7 讲） |
| 知道"训练一个模型就是最小化损失"的大意（不熟悉的话，[SYLLABUS](../SYLLABUS.md) 里规划了前置讲 P3 来补） | 分布式训练、ZeRO、量化、FlashAttention、服务系统（第 14–27 讲） |

### 0.3 本讲地图 Map

```
§1  为什么要专门学 LLM 系统：规模曲线；三个学习目标；总题目"100B 模型要多少资源"
§2  LLM 能做什么——以及每种能力给系统带来的负载
§3  数学基础：下一个 token 的概率；链式法则；三类语言模型；学习框架；      ← 代码 01、02、03
    交叉熵与困惑度；评估指标；规模与数据
§4  LLM 系统的挑战：核心问题与指标；抽象与权衡；底层算子；不同层次的挑战；  ← 代码 05
    协同设计；计算 vs 数据搬运；编程模型
§5  回答总题目：参数量、FLOPs、时间、钱、显存、通信、推理                    ← 代码 04
§6  课程地图   §7  怎样学这门课
§8  小结   §9  常见误区   §10 练习   §11 延伸阅读   §12 下一讲预告
```

建议的阅读顺序：先读 §1–§2（不需要代码）；读 §3 时依次运行 `01`、`02`、`03`；读 §4 时运行 `05`；最后读 §5 并运行 `04`。§5 是本讲的"总账"，前面每一节都会为它准备一个零件。

### 0.4 记号 Notation（全课统一）

| 记号 | 意思 |
|---|---|
| $x_t$，$x_{1:T}$，$x_{<t}$ | 第 $t$ 个 token；长度为 $T$ 的序列；$x_t$ 之前的所有 token（"前文"） |
| $\mathcal{V}$，$V = \lvert\mathcal{V}\rvert$ | 词表及其大小 |
| $p_\theta(\cdot)$ | 参数为 $\theta$ 的模型给出的概率 |
| $N$ | 模型参数量（parameters） |
| $D$ | 训练用的 token 数（data） |
| $L$，$d$，$H$，$d_h$ | 层数、隐藏维度（`d_model`）、注意力头数、每个头的维度（$d = H \cdot d_h$） |
| $T$ | 序列长度 / 上下文长度 |
| $B$，$b$ | 批大小（batch size） |
| FLOPs | 浮点运算次数（一次乘法或一次加法各算 1 次） |
| FLOP/s | 每秒浮点运算次数（算力） |
| 算术强度 arithmetic intensity | 一个运算的 FLOPs ÷ 它必须读写的字节数，单位 FLOP/字节 |
| MFU | Model FLOPs Utilization：实际达到的有效算力 ÷ 硬件峰值算力 |
| nat | 以 $e$ 为底的信息量单位（以 2 为底时叫 bit）；1 nat ≈ 1.443 bit |
| KB / MB / GB / TB | 以 1000 为进位（$10^3$、$10^6$、$10^9$、$10^{12}$ 字节）；KiB / MiB / GiB 以 1024 为进位 |

### 0.5 关于这门课 · About 11-868

**是什么。** CMU 语言技术研究所（LTI）的研究生课，2024 年春季第一次开，2025、2026 年春季都有课程主页（Spring 2026 的 slides 里还提到了一个秋季学期的版本）。主讲 Lei Li 的背景很能说明这门课的味道：他在 CMU 读完博士后在工业界做了很多年大规模 NLP 系统，主导过字节跳动的机器翻译系统 VolcTrans 和写作机器人 Xiaomingbot（slides 上写的是服务超过十亿用户），做过 Transformer 加速库 LightSeq（第 10–11 讲的主角），2023 年回到 CMU 任教；他的 VOLT（第 8 讲）获得了 ACL 2021 最佳论文。所以这门课一半是 NLP，一半是系统，而且非常强调"能动手实现"。

**学什么。** slides 列出三个学习目标：

1. **理解**现代 LLM 系统的关键技术——例如，训练一个 100B 模型需要多少资源？
2. **工程能力**：亲手实现 LLM 系统的关键部件——快速的 CUDA 内核、可扩展的训练系统、高效的推理。
3. **创新**：发现机会，解决 LLM 系统研究里新的关键挑战。

**不适合谁。** slides 很坦率：如果你想学怎样搭深度学习模型、怎样用 PyTorch，去上 11-685/11-785（深度学习导论）；如果你只想学怎样**使用** LLM、对底层系统不感兴趣，去上 11-667；如果你没法来上课、完成作业和项目，也别选——这门课要求真的写系统代码。

**怎么考核**（Spring 2026）：5 次必做 + 2 次选做的编程作业共 44%（选做另加 5%），课堂小测 10%，参与 2%，小组项目 44%（开题 2%、中期 2%、展示 20%、终期报告 20%）。作业个人完成，主线是**从零扩展一个教学用深度学习框架 miniTorch**——写 CUDA 内核（HW1）、自动微分（HW2）、Transformer（HW3）、融合内核（HW4），之后转向分布式训练、DeepSpeed、SGLang 和 RLHF（HW5–7）；要写 Python 和 C++ / CUDA，事先不要求会 CUDA。项目 2–3 人一组，二选一：给 miniTorch 实现一篇 LLM 系统论文里的组件（不许调用外部的 PyTorch / TensorFlow 代码），或者做一个以投稿系统 / ML 会议为目标的研究项目；最后在课上做海报展示，最好带演示。算力由匹兹堡超算中心（PSC Bridges-2）提供。

**原课作业 vs 本讲义代码。** 原课作业在 GPU 集群上写 CUDA。本讲义假设你只有一台普通电脑：GPU 内核用 Python 写的模拟器执行并逐项计数（第 2 讲起），性能用公开规格和 roofline 模型估算，多卡用"列表里的多个设备"模拟，所有模型都小到能在 CPU 上几十秒训练完。详见 [SYLLABUS.md](../SYLLABUS.md) 的 §0.4。

### 0.6 本讲的玩具语料 · The city language

讲语言模型时，我们总想回答一个问题：**这个模型离"最好"还差多远？** 对真实英语这个问题没有答案，因为没人知道英语的真实分布。所以代码 [`city_corpus.py`](../code/lecture01/city_corpus.py) 造了一门小语言，它的每一句话都由几次"掷骰子"生成，骰子的概率都写在代码里。四种句型：

| 句型 | 模板 | 例子 | 概率 |
|---|---|---|---|
| T1 | `<城市> is a city of <特色> .` | pittsburgh is a city of bridges . | 0.25 |
| T2 | `<城市> has very nice <好东西> .` | santa barbara has very nice beaches . | 0.20 |
| T3 | `the <人群> from <城市> who <活动> <时间> love the <特色> .` | the students from pittsburgh who swim at night love the bridges . | 0.35 |
| T4 | `in <城市> it is often <天气> .` | in seattle it is often rainy . | 0.20 |

10 个城市各有自己的"特色""好东西""天气"分布（例如 pittsburgh 的特色：bridges 0.60、rivers 0.20、steel 0.12、hills 0.06、corn 0.02——这正是 slides 上 "Pittsburgh is a city of ___" 的例子）。T3 是专门设计的**长程依赖**：句末的特色只取决于 7–8 个 token 之前的城市，中间隔着人群、活动、时间这些和答案无关的"填充物"。这些概率都是为教学编的，不是真实统计。

因为一切都是已知的，我们能**精确**算出：任何句子的概率、任何前缀之后的下一个 token 分布、测试集上任何模型都不可能超过的"理论最优困惑度"。运行 `python3 city_corpus.py` 可以看到这门语言的概况：9127 个不同的句子（其中 9000 个是 T3），词表 80 个 token，熵率 0.7506 nat/token，理论最优困惑度 2.118。

---

## 1. 为什么要专门学"LLM 系统" · Why LLM systems

### 1.1 一条陡峭的曲线

slides 的第二、三页是两张时间线：从 2000 年前后的神经网络语言模型，到 2018 年的 BERT，再到 2020–2024 年的一系列大模型；以及一张对数坐标的"LLM 规模"图。把图上的模型和它们公开的参数量列出来：

| 年份 | 模型 | 参数量 | 备注 |
|---|---|---|---|
| 2017 | Transformer（原始论文） | base 约 65M，big 约 213M | 机器翻译 |
| 2018 | GPT-1 | 0.117B | 第一个"预训练 + 微调"的 GPT |
| 2019 | GPT-2 | 1.5B | |
| 2020 | GPT-3 | 175B | 少样本提示、上下文学习 |
| 2021 | Gopher（DeepMind） | 280B | |
| 2022 | PaLM（Google） | 540B | |
| 2023 | GPT-4 | 未公开 | 原课第 14 讲的同一张图上标的是 1.8T，这是外界流传的估计 |
| 2024 | LLaMA-3.1 | 405B | 开放权重里最大的稠密模型之一 |
| 2024 | DeepSeek-V3 | 671B（每个 token 激活 37B） | 混合专家（MoE） |
| 2025 | Qwen3 | 235B（激活 22B） | MoE |
| 2025 | Kimi K2 | 约 1T（激活 32B） | MoE |

两个趋势一眼就能看出来：

1. **七年涨了一万倍。** 从 GPT-1 的 1 亿参数到万亿参数。同一时期，单块 GPU 的显存只从十几 GB 涨到一两百 GB（V100 16/32 GB、A100 80 GB、H100 80 GB、B200 约 192 GB），16 位精度的峰值算力涨了十几倍。模型增长的速度远远超过了单块硬件增长的速度——**差出来的部分，只能靠系统来补**：多卡并行、显存优化、压缩、更好的内核。
2. **2024 年以后，最大的模型几乎都是 MoE。** "总参数"和"每个 token 实际用到的参数"分开了：DeepSeek-V3 有 671B 参数，但每个 token 只激活 37B。这本身就是一个"为了系统而改模型"的例子——在计算量几乎不变的前提下增加容量，代价是更复杂的通信和负载均衡（第 17 讲）。

### 1.2 本课的总题目

原课第一个学习目标后面跟着一个具体问题：**训练一个 1000 亿参数的模型，需要多少资源？**

这个问题看起来只要一个数字，其实要拆成一串问题：一个 token 要算多少次乘法？一块 GPU 一秒能做多少次？要多少块 GPU、多少天、多少钱？显存装得下吗？GPU 之间要传多少数据、传得过来吗？训练完了，服务一个用户要多少显存、一秒能生成多少个 token？

本讲 §5 会用几条公式把这串问题全部回答一遍（结果是：约 1.2×10²⁴ 次浮点运算、1024 块 H100 一个多月、上百万美元、1.6 TB 的训练状态、每步 200 GB 的梯度通信……）。更重要的是，每一条公式都对应着这门课后面的一个单元——**这门课的结构，就是这本账的结构**（§6）。

### 1.3 为什么系统成了 LLM 研究的核心

slides 后面有一句话可以概括这门课的立场：**"Scaling is all you need"——既要往上扩（scale up），也要往下缩（scale down）。**

- **往上扩**：更大的模型、更多的数据、更长的上下文。每往上一步，都会撞上一堵新的墙：一块卡放不下（→ 模型并行、ZeRO）、算不完（→ 数据并行、更快的内核）、传不动（→ 通信优化）。
- **往下缩**：同样的能力用更少的资源。量化（第 19–20 讲）、参数高效微调（第 23 讲）、更好的推理系统（第 22–27 讲）都属于这一类。一个模型训练一次，却要被调用几十亿次，**推理的成本最终往往超过训练**。

在这两个方向上，决定成败的常常不是"模型设计得多巧妙"，而是"系统能不能把硬件用满"。§5.8 会看到，同样的模型和数据，MFU 从 20% 提到 60%，训练时间能差 3 倍。

---

## 2. LLM 能做什么——以及每种能力给系统带来的负载 · Capabilities

### 2.1 一组例子

slides 用了一连串例子展示 LLM 的能力（下面是转述）：

- **写宣传文案**：让模型参考课程主页，写一条宣传这门课的"爆款推文"。
- **翻译**：传统机器翻译系统 vs ChatGPT / GPT-4 的对比。
- **模仿一个人聊天**：一个 CMU 学生项目（WhatsApp-Llama）用自己的聊天记录微调模型，大约 10 段聊天记录、几个小时的训练，就能生成模仿本人风格的回复。
- **按例子改写风格**：给几个"把句子改成尤达大师的说话方式"的例子，模型就能改写一个新句子——这是**上下文学习**（in-context learning）：不改参数，只靠提示里的例子学会新任务。
- **常识推理**：Winograd 式的代词消解（"Steven 对 Robert 做猪肉很反感，因为 ___ 是素食者"，答案是 Steven）。
- **数学推理**：一道小学应用题（网球的个数），模型一步一步写出推理过程再给出答案。
- **考试**：一项研究发现 ChatGPT 在沃顿商学院的 MBA 运营管理考试里能拿到 B− 到 B。
- **写代码并执行**：一道工厂流程题（七道工序各有产能，问瓶颈在哪、整条线产能多少）。模型先写出分析，再**生成一段 Python 代码、交给 Python 执行器运行**，用运行结果来验证答案；另一个例子是让模型写一个找"山峰"位置的函数。
- **画图**：按描述画一幅"羊驼沿着 CMU 校园里的雕塑《Walking to the Sky》往上爬"的图。

这些例子在原课里是为了说明"LLM 已经很强、很通用"。但对一门**系统**课来说，更值得问的是：**每一种用法，对系统提出了什么要求？**

### 2.2 换一个角度：能力 → 负载

一个 LLM 请求的"形状"可以用四个量刻画：**输入多长**（提示词，决定预填充 prefill 的计算量）、**输出多长**（决定逐 token 解码 decode 的步数）、**一个任务要调用模型几次**、**用户能等多久**。

| 用法 | 输入 | 输出 | 调用次数 | 时延要求 | 主要的系统挑战 | 相关讲次 |
|---|---|---|---|---|---|---|
| 翻译、改写 | 短 | 短 | 1 | 中 | 大批量吞吐；把很多短请求凑成批 | 9、22 |
| 多轮对话 | 越聊越长 | 中 | 每轮 1 次 | 高（要"秒回"） | 历史对话每轮都要重新输入 → 复用前缀的 KV 缓存 | 22、24、27 |
| 少样本提示 | 长（例子很多） | 短 | 1 | 中 | 预填充很重；很多请求共用同一段例子 → 前缀缓存 | 21、22 |
| 一步一步推理 | 中 | **很长** | 1 | 低–中 | 解码步数多；每步都要读一遍全部权重 → 访存受限 | 9、19、24 |
| 写代码 + 调用工具 / 智能体 | 越来越长 | 中 | **很多次** | 高 | 一次任务是一串调用；中间结果和上下文要反复使用 | 22、26、27 |
| 用自己的数据微调 | — | — | — | — | 训练显存；只训练少量参数 | 18、23 |
| 画图 | 短 | 图像 | 1 | 中 | 另一类模型（扩散模型），不是本课重点 | — |

这张表不是原课 slides 上的，但它解释了这门课后半部分的选题：几乎每一讲服务系统的论文，都是在针对上表的某一行做优化。

### 2.3 一个贯穿全课的比喻：瓶颈

slides 里那道工厂题值得多看一眼。七道工序的产能（吨/小时）分别是 60、55、56（14 吨 ÷ 0.25 小时）、50（200 吨 ÷ 4 小时）、62.5、59、82.5（3 台 × 27.5）。**整条生产线的产能等于最慢的那一道工序：50 吨/小时**，瓶颈是第二个反应器。

这正是系统优化的第一原则：**先找瓶颈**。把一道不是瓶颈的工序加快十倍，整条线一点都不会变快。05 号脚本第 4 部分把同样的道理用在一个训练步上：如果"计算"要 6.0 秒、"梯度通信"要 8.1 秒，那么不论把计算优化得多快，只要两者不能重叠，一步至少要 8.1 秒。后面几乎每一讲都会先问：**这里的瓶颈是计算、访存、通信，还是调度？**

---

## 3. 数学基础 · Mathematical Foundations

### 3.1 语言模型 = 下一个 token 的概率模型

**直觉。** 给你半句话 "Pittsburgh is a city of ___"，你脑子里会冒出几个候选：bridges 很可能，rivers 也说得通，corn 就有点奇怪但不是不可能。语言模型做的就是这件事，只不过它给出的不是一个答案，而是**整个词表上的一个概率分布**：

$$
P(x_t \mid \text{提示词},\ x_{1:t-1})
$$

> **为什么说 "token" 而不是 "词"？** 模型的词表不是字典里的词，而是分词器（tokenizer）切出来的片段：常见的词是一个 token，罕见的词会被切成好几个子词，空格、标点、数字也各有 token。城市语言里我们偷懒，直接把一个词当一个 token（"new york" 是两个 token）。分词的方式会影响序列长度、词表大小、嵌入矩阵的参数量和 softmax 的计算量——这是第 8 讲的内容。本讲说"词"的地方，严格地说都应该是"token"。

slides 举了两个例子："Santa Barbara has very nice ___"（beach 0.5、weather 0.4、snow 0.01）和 "Pittsburgh is a city of ___"（bridges 0.6、corn 0.02）。我们的城市语言里嵌入了同样的概率表。01 号脚本第 2 部分**精确**算出了几个前缀之后的真实分布：

| 前缀 | 真实的下一个 token 分布（前几项） |
|---|---|
| pittsburgh is a city of | bridges 0.600 · rivers 0.200 · steel 0.120 · hills 0.060 · corn 0.020 |
| santa barbara has very nice | beaches 0.500 · weather 0.400 · views 0.060 · food 0.030 · snow 0.010 |
| the students from pittsburgh who swim at night love the | bridges 0.600 · rivers 0.200 · steel 0.120 · hills 0.060 · corn 0.020 |
| the chefs from seattle who read on weekends love the | coffee 0.500 · rain 0.250 · music 0.150 · mountains 0.070 · bridges 0.030 |
| 只看最近两个 token "love the"（最好的 trigram） | bridges 0.150 · beaches 0.098 · music 0.083 · coffee 0.081 · museums 0.060 · parks 0.052 |

两点值得注意：

1. **答案是分布，不是一个词。** corn 只有 0.02 的概率，但它不是 0。"按分布采样"时偶尔会生成它——这就是为什么同一个提示词问两次，LLM 会给出不同的回答。
2. **看多长的上下文，决定了能答多好。** 同样以 "love the" 结尾，pittsburgh 的人最可能爱 bridges（0.60），seattle 的人最可能爱 coffee（0.50）。如果只看最近两个 token，就只能给出所有城市的平均：bridges 降到 0.150。**能利用多长的上下文，直接决定模型质量**——而长上下文正是 LLM 系统最昂贵的需求之一（KV 缓存随上下文线性增长，注意力的计算随它平方增长）。

### 3.2 链式法则：恒等式，而不是假设

一句话的概率怎样由"下一个 token 的概率"得到？slides 写的是概率论里的**链式法则**：

$$
P(x_{1:T}) = \prod_{t=1}^{T} P(x_t \mid x_{1:t-1})
$$

例如 $P(\text{Pittsburgh is a city of bridges}) = P(\text{Pittsburgh}) \cdot P(\text{is} \mid \text{Pittsburgh}) \cdot P(\text{a} \mid \text{Pittsburgh is}) \cdots$。

**为什么这样做？** 直接给每个完整句子一个概率是不可能的——可能的句子有无穷多个。链式法则把"给句子打分"变成了"一次只预测一个 token"，而后者可以用同一个函数（神经网络）反复调用来实现。

**关键的一点：链式法则是恒等式，它不做任何假设。** 只要每一步都以**全部**前文为条件，连乘的结果一定等于句子的真实概率。01 号脚本第 3 部分验证了这一点：对 "the students from pittsburgh who swim at night love the bridges ."，12 个真实条件概率连乘得到 $\log p = -8.3630425116$，直接查表也是 $-8.3630425116$，只差 $5.3 \times 10^{-15}$（浮点舍入）。

近似发生在**模型**这一步：用什么东西去估计每一个 $P(x_t \mid x_{<t})$。最古老的办法是 n-gram 的**马尔可夫假设**——只看最近 $n-1$ 个 token。把上面那句话换成"只看最近两个 token"的最好估计，前 11 个位置都和真实值一样（那些位置本来就只依赖最近的词），唯独最后的 bridges 从 0.600 变成 0.150，整句的概率被低估了 4.00 倍——它依赖的 pittsburgh 在窗口之外。slides 说得很简洁：用神经网络（Transformer、CNN、RNN）来预测每一个条件概率。各种语言模型的区别，**只在于怎样近似这一串条件概率**。

### 3.3 为什么 ChatGPT 改变了 AI 的格局

slides 给了三条原因：

1. **在海量原始文本上预训练（约 300B token）+ 少量人类反馈。** 预训练的数据不需要人工标注，互联网上有的是；人类反馈只需要相对很少的量，就能把一个"续写机器"变成一个"助手"。
2. **能听懂指令（instruction following）。** 用户用自然语言描述任务即可，不需要为每个任务训练一个模型。
3. **上下文学习（in-context learning）。** 在提示里给几个例子，模型就能泛化到新任务（§2.1 的尤达风格改写）。

从系统的角度看，这三条有一个共同的后果：**一个模型、一种接口（续写文本）服务所有任务**。这让"把这一个模型训练好、服务好"的价值变得极高，也让这门课的所有优化都有了用武之地。

### 3.4 三类语言模型：差别只在"谁能看谁"

slides 把语言模型分成三类：

- **encoder-only**（BERT）：**掩码语言模型**（masked LM）。随机遮住一些 token，让模型根据**两边**的上下文猜被遮住的词，建模 $P(x_{\text{mask}} \mid x_{\text{其余}})$。它主要用来做理解类任务；也可以用来生成（例如非自回归翻译 NAT，以及 Lei Li 组的 REDER），但不常见。
- **encoder-decoder**（原始 Transformer、T5、翻译模型）：编码器读完整个输入 $x$，解码器一个一个地生成输出 $y$，建模 $P(y \mid x) = \prod_t P(y_t \mid y_{<t}, x)$。编码器可以是双向 LSTM、Transformer 或 CNN；解码器可以是 LSTM、Transformer、非自回归 Transformer 或 CNN。
- **decoder-only**（GPT 系列、LLaMA 及今天几乎所有 LLM）：**因果语言模型**（causal LM），直接建模 $P(x) = \prod_t P(x_t \mid x_{<t})$。可以用 Transformer 或单向 LSTM 实现。slides 的评价是：最流行的 LLM 结构。

用 Transformer 实现时，三类模型用的是**同一种注意力运算**，区别只在**注意力掩码**（attention mask）——位置 $i$ 可以"看"哪些位置 $j$。02 号脚本第 1 部分把它们画了出来（■ 允许，· 禁止）：

```
 (a) encoder-only：全部可见          (b) decoder-only：只看自己和过去     (c) encoder-decoder（拼接后的等价掩码）
     1  2  3  4  5  6                    1  2  3  4  5  6                    源1 源2 源3 目1 目2 目3
  1  ■  ■  ■  ■  ■  ■                 1  ■  ·  ·  ·  ·  ·              源1  ■   ■   ■   ·   ·   ·
  2  ■  ■  ■  ■  ■  ■                 2  ■  ■  ·  ·  ·  ·              源2  ■   ■   ■   ·   ·   ·
  3  ■  ■  ■  ■  ■  ■                 3  ■  ■  ■  ·  ·  ·              源3  ■   ■   ■   ·   ·   ·
  4  ■  ■  ■  ■  ■  ■                 4  ■  ■  ■  ■  ·  ·              目1  ■   ■   ■   ■   ·   ·
  5  ■  ■  ■  ■  ■  ■                 5  ■  ■  ■  ■  ■  ·              目2  ■   ■   ■   ■   ■   ·
  6  ■  ■  ■  ■  ■  ■                 6  ■  ■  ■  ■  ■  ■              目3  ■   ■   ■   ■   ■   ■
```

（真正的 encoder-decoder 模型里编码器和解码器各有一套参数，目标语言通过"交叉注意力"看源语言；(c) 是把它画成一个拼接序列后的等价掩码。只用一套参数、按 (c) 的掩码训练的模型叫 prefix-LM。）

**为什么今天的 LLM 几乎都是 decoder-only？** 从建模的角度，原因之一是 §3.3：所有任务都能写成"续写"，一个因果模型就够了。但从**系统**的角度，还有三条更硬的理由，02 号脚本逐条做了实验：

**(1) 训练可以完全并行，而且每个 token 都有用。** 因果掩码意味着：改动第 $j$ 个 token，只有位置 $\ge j$ 的输出会变。02 号脚本第 2 部分把 "pittsburgh" 换成 "seattle"，因果模型里前 4 个位置的隐藏状态完全不变（`○○○○●●●●●●●●●`），双向模型里所有位置都变了（`●●●●●●●●●●●●●`）。所以一个长度为 $T$ 的序列，一次前向就能同时得到 $T$ 个"预测下一个词"的训练样本，而且每个预测都看不到答案。BERT 式的掩码语言模型只对被遮住的约 15% 的位置计算损失（原始 BERT 的设定）；encoder-decoder 只对目标语言那一半计算损失（第 5 部分）：

| 模型 / 训练目标 | 处理的 token | 有监督的预测 | 每个 token 的训练信号 |
|---|---|---|---|
| decoder-only（下一个 token） | 512 | 512 | 1.00 |
| encoder-only（遮住 15%） | 512 | 76 | 0.15 |
| encoder-decoder（源、目标各一半） | 512 | 256 | 0.50 |

**(2) 生成时可以缓存过去的计算（KV 缓存）。** 自回归生成时，每生成一个新 token 都要再跑一次模型。因果模型里，旧位置看不到新 token，所以它们在每一层算出的键 $k$ 和值 $v$ **永远不会变**——算一次存起来，下一步直接复用。用一层注意力画出来：

```
 预填充 prefill：一次算完提示词的 5 个位置，把每个位置的 k、v 存进缓存
   位置      0     1     2     3     4
   缓存 K   k0    k1    k2    k3    k4
   缓存 V   v0    v1    v2    v3    v4
                                     └─ 位置 4 的输出 → 采样出第 1 个新 token x5

 解码 decode 第 1 步：只算新 token x5 的 q5、k5、v5
   q5 和缓存里的 k0..k5 做点积 → 权重 → 对 v0..v5 加权求和 → 输出 → 采样出 x6
   缓存追加：K = [k0 … k5]，V = [v0 … v5]      （k0..k4、v0..v4 一个都没重算）

 解码第 2 步：只算 x6 的 q6、k6、v6，和 k0..k6 做注意力……如此往复
```

每层都有这样一份缓存。每一步只有"新 token 的那一行"需要计算，但它要**读**整个缓存——所以缓存越长，每一步读的字节越多（§5.7）。02 号脚本第 3 部分验证了这一点：追加一个 token 后，因果模型旧位置的隐藏状态最多只变 $1.19 \times 10^{-6}$（浮点舍入量级），双向模型变了 $0.817$；用 KV 缓存逐个 token 计算出的 logits 和整句一次计算的结果最多差 $2.09 \times 10^{-7}$。双向模型每来一个新 token，所有位置都要重算——没有东西可缓存。

**(3) 缓存省下的计算随长度平方增长。** 02 号脚本第 4 部分用 PyTorch 的 FLOP 计数器数了一下（提示词 5 个 token，贪心生成 G 个）：

| 生成 token 数 G | 位置计算次数（无缓存 / 有缓存） | MFLOPs（无缓存 / 有缓存） | 倍数 | 结果相同 |
|---|---|---|---|---|
| 8 | 68 / 12 | 14.4 / 2.5 | 5.7× | 是 |
| 32 | 656 / 36 | 144.0 / 7.8 | 18.5× | 是 |
| 59 | 2006 / 63 | 458.6 / 14.1 | 32.6× | 是 |

没有缓存时，第 $k$ 步要把 $P+k$ 个位置全部重算，总量随 $G$ 平方增长；有缓存时每步只算一个新位置。但要注意两点：有缓存时每步的计算并不是常数——新 token 仍要和所有旧的 $k$、$v$ 做注意力，这一项随上下文线性增长；而且缓存要占**显存**，每个 token、每一层都要存一个 $k$ 和一个 $v$。§5.7 会算出，对一个 100B 模型，一条 2048 token 的序列的 KV 缓存就有 6.7 GB。**KV 缓存把"重算过去"换成了"存储过去"**——第 22、24、26、27 讲整整四讲都在处理由它引出的问题。

（⏱ 同一张表里本机的实测耗时只差 1.6 倍左右：这个模型太小，时间主要花在 Python 和内核调用的固定开销上。"FLOPs 少了"不等于"同样比例地变快"——这是 §4.7 的主题。）

### 3.5 LLM 的学习框架：预训练 → 监督微调 → 强化学习

slides 把 LLM 的训练画成三个阶段：

1. **预训练**（pre-training）：在海量原始文本（维基百科、书籍、Reddit 等）上做下一个 token 预测。数据越多越好、通常从网上爬取，但**质量很重要**（维基百科、书籍、过滤过的 Reddit 帖子）。模型结构上的选择（稀疏注意力、混合专家等）也在这一阶段决定。
2. **监督微调**（supervised fine-tuning, SFT）：用"问题 → 好的回答"这样的成对数据继续训练，让模型学会按指令回答。
3. **强化学习**（RL，例如 RLHF）：用偏好数据（同一个问题的两个回答，哪个更好）训练，让模型的回答更符合人的偏好。

从系统的角度看，三个阶段的负载完全不同：预训练是**超大规模、长时间、吞吐优先**的训练（第 14–18 讲）；SFT 的数据量小得多，常常用参数高效微调（第 23 讲）；RLHF 要同时运行好几个模型、在"生成"和"训练"之间来回切换（第 25 讲）。

### 3.6 训练目标：交叉熵

预训练的损失是**下一个 token 预测的交叉熵**（cross-entropy loss）：

$$
\mathcal{L}_{\text{CE}} = -\frac{1}{T} \sum_{t=1}^{T} \log P_\theta(x_t \mid x_{<t})
$$

**直觉。** 模型在第 $t$ 步给真正出现的那个 token 分配了概率 $P_\theta(x_t \mid x_{<t})$；分配得越高，$-\log$ 就越小。把所有位置平均起来，就是"平均每个 token 有多惊讶"。

**为什么是交叉熵？** 最小化交叉熵等价于**最大似然估计**：让模型给训练数据分配的总概率最大。它还有一个漂亮的分解：

$$
\underbrace{\mathbb{E}\big[-\log P_\theta\big]}_{\text{模型的交叉熵}} = \underbrace{H(P_{\text{真实}})}_{\text{不可约：语言本身的随机性}} + \underbrace{\mathrm{KL}(P_{\text{真实}} \,\|\, P_\theta)}_{\text{可约：模型还没学到的}}
$$

KL 散度永远 $\ge 0$，所以任何模型的交叉熵都不可能低于真实分布的熵。在城市语言里，我们能精确算出第一项（熵率 0.7506 nat/token），于是每个模型"还差多少"就是第二项——本讲后面的表格里"超出下界"的那一列就是它（单位：毫 nat/token）。

**训练是怎样并行的？** 因为因果掩码（§3.4），一个长度为 $T$ 的窗口一次前向就给出 $T$ 个位置的预测，损失是它们的平均——这叫**教师强制**（teacher forcing）：训练时每个位置的"前文"用的是真实数据，而不是模型自己生成的。

**数据怎样喂进去：打包（packing）。** 句子长短不一，最直接的做法是把每句补齐（padding）到同样长度——但补齐的位置算了也白算。03 号脚本第 1 部分在 20000 句训练集上算了一下：

| 做法 | 需要计算的位置数 | 浪费在填充上 |
|---|---|---|
| 每句补齐到全局最长 | 280,000 | 30.2% |
| 随机每 32 句一批、补齐到批内最长 | 278,784 | 29.9% |
| 先按长度排序再分批（分桶） | 195,424 | 0.1% |
| 首尾相接打包成一条 token 流 | 195,308 | 0.0% |

随机分批几乎没有改善：每 32 句里几乎总有一句最长的 T3。分桶能消除填充，但同一批里几乎全是同一种句型，批次不再是整体分布的随机样本。LLM 预训练几乎都用**打包**：把文档首尾相接（中间放一个分隔符，这里是 `<bos>`），按固定长度的窗口切开，每个位置都是真实 token。

### 3.7 怎样衡量一个语言模型

**困惑度（perplexity）** 是语言模型最常用的指标：

$$
\mathrm{PPL} = \exp\!\Big(-\frac{1}{T}\sum_{t=1}^{T} \log P_\theta(x_t \mid x_{<t})\Big) = \exp(\mathcal{L}_{\text{CE}})
$$

**直觉。** 困惑度可以理解成"模型平均在几个候选之间犹豫"。如果模型对每个 token 都在 $V$ 个词里均匀地猜，困惑度正好是 $V$；如果它每次都确定地猜对，困惑度是 1。城市语言的理论最优困惑度是 2.118：即使知道真实分布，平均每个位置也还有约 2 个"同样合理"的选择。

01 号脚本第 4 部分在同一个测试集（2000 句、17577 个 token）上比较了各种 n-gram，并把误差拆成两种：

| n（看 n−1 个 token） | 1 | 2 | 3 | 4 | 5 | 6 | 8 | 10 | 14 |
|---|---|---|---|---|---|---|---|---|---|
| 无限数据（精确计算） | 44.987 | 3.031 | 2.480 | 2.412 | 2.357 | 2.260 | **2.123** | 2.123 | 2.123 |
| 2000 句训练数据 | 45.039 | 3.076 | 2.501 | 2.439 | 2.410 | **2.370** | 2.457 | 2.608 | 2.934 |
| 20000 句训练数据 | 44.995 | 3.036 | 2.482 | 2.415 | 2.362 | 2.276 | **2.170** | 2.192 | 2.312 |

（真实分布下界 2.1234；Witten–Bell 插值平滑；词表大小 80，unigram 的 44.99 已经远好于均匀猜测的 80。）

- **第一行是近似误差。** 数据无限多时，窗口越长越好，n = 8（看 7 个 token）就到达下界：这门语言最远的依赖正好是 7 个 token。这一行和下界的差距完全来自"模型结构看不到"。
- **后两行加上了估计误差。** 数据有限时，窗口太长反而变差：2000 句时最好的是 n = 6，n = 14 退到 2.934。长上下文的组合太多，大部分只见过一两次，估不准。极端情况是不平滑的"∞-gram"（等于背下训练集）：测试集里有 238 句从没在训练集里出现过，它们的概率是 0，困惑度是**无穷大**。

n-gram 只能在这两种误差之间取舍，因为它把 "pittsburgh who swim at night love the" 和 "pittsburgh who sing every day love the" 当成两个毫无关系的上下文。神经语言模型的优势恰恰在这里：它能学会"中间那些词不重要"，在不同上下文之间**共享统计信息**，于是能同时用长上下文和有限数据（§3.8）。

**生成的质量也会受影响。** 01 号脚本第 5 部分从各个 n-gram 里采样了 1000 句：bigram 只有 0.117 是这门语言里真实可能出现的句子，trigram 0.381，8-gram 0.966。trigram 生成的 T3 句子有 0.799 是城市和特色对不上的，例如 "the engineers from chicago who run on weekends love the coffee ."——生成到句末时，城市早已滑出窗口，模型"忘了"自己前面说过什么。

**其他指标。** 困惑度适合开放式生成，但它只告诉你模型对数据有多"不惊讶"。slides 还列了两类指标：

- **基于参考答案的指标**：BLEU、ROUGE（数 n-gram 的重合，老但经典）；COMET、BLEURT、SEScore2（用一个训练好的模型打分）；InstructScore（能给出解释的打分，Lei Li 组的工作）。
- **下游任务指标**：命名实体识别看 F1，问答看准确率或匹配率，代码生成看 pass@k（生成 k 个程序至少有一个通过测试的概率），检索看 nDCG，摘要看 ROUGE，翻译看 COMET / SEScore2 / BLEU。

对系统研究来说，这些指标的用处是**守住底线**：量化、剪枝、推测解码、KV 缓存压缩这些优化都可能悄悄损害模型质量，必须用困惑度和下游指标证明"变快了、但没变差"。城市语言的好处是能精确地量出"变差了多少"——第 19–20 讲（量化）会直接用这把尺子。

### 3.8 实验：在 CPU 上预训练一个小 GPT

03 号脚本用 [`tiny_gpt.py`](../code/lecture01/tiny_gpt.py) 里的小 GPT（2 层、d = 64、4 个头、上下文 64，共 109,312 个参数）在 20000 句城市语言上预训练：2000 步 × 每批 32 个窗口 × 每个窗口 64 个 token = 409.6 万个 token（约 21 遍训练集），⏱ 单线程 CPU 上约 22 秒。Transformer 内部每个部件为什么这样设计，第 6 讲会细讲；这里把它当作"给定前文、输出下一个 token 分布"的函数。

**训练曲线**（验证集困惑度；验证集的理论下界 2.1139）：第 1 步 80.59（随机初始化 ≈ 均匀猜测 = 词表大小）→ 250 步 2.399 → 500 步 2.152 → 1000 步 2.141 → 2000 步 2.119（只超出下界 2.6 毫 nat/token）。

**和 n-gram 比**（同一个测试集、同样的 20000 句训练数据）：

| 模型 | 困惑度 | 超出下界（毫 nat/token） |
|---|---|---|
| 均匀猜测 | 80.000 | — |
| 1-gram | 44.995 | 3053.6 |
| 3-gram | 2.482 | 156.0 |
| 8-gram（n-gram 里最好的） | 2.170 | 21.8 |
| **小 GPT（109,312 参数）** | **2.130** | **3.1** |
| 真实分布（理论下界） | 2.123 | 0 |

小 GPT 的可约损失只有最好的 n-gram 的 14%。

**它学到的分布像不像？** 03 号脚本第 3 部分把小 GPT 和 3-gram 的下一个 token 分布与真实分布比较（每格依次是 真实 / 小 GPT / 3-gram）：

| 前缀 | 前两个 token | KL(真实‖小 GPT) | KL(真实‖3-gram) |
|---|---|---|---|
| pittsburgh is a city of | bridges 0.60/0.62/0.15 · rivers 0.20/0.19/0.05 | 0.0063 | 1.361 |
| santa barbara has very nice | beaches 0.50/0.47/0.09 · weather 0.40/0.42/0.11 | 0.0082 | 1.210 |
| the students from pittsburgh who swim at night love the | bridges 0.60/0.61/0.15 · rivers 0.20/0.18/0.04 | 0.0051 | 1.382 |
| the chefs from seattle who read on weekends love the | coffee 0.50/0.48/0.09 · rain 0.25/0.30/0.04 | 0.0169 | 1.405 |
| the doctors from miami who sing in winter love the | beaches 0.60/0.62/0.10 · music 0.20/0.16/0.08 | 0.0235 | 1.465 |

小 GPT 几乎把真实的条件概率表"背"了下来，包括长程的 T3。一个我起初没想到的结果：3-gram 在 T1、T2 上**也**很差（KL 1.36、1.21）——"a city of ___" 和 "very nice ___" 里都没有城市，城市在 3 个 token 之外。按句型拆开看（第 4 部分），3-gram 在四种句型上都超出至少 98 毫 nat/token（T1 最多，236.3）；8-gram 在 T1、T2、T4 上几乎没有差距，但在 T3 上还差 44.8，而小 GPT 只差 2.2：9000 种 T3 组合里大部分在训练集里只出现过零到一两次，n-gram 只能靠平滑"猜"，小 GPT 却学会了"填充词与答案无关"。

**采样。** 03 号脚本第 5 部分从小 GPT 采样 1000 句：

| 温度 | 合法句子 | T3 中不合法 | T3 的比例（真实 0.35） | 含 pittsburgh（真实 0.20） | 不同句子个数 |
|---|---|---|---|---|---|
| 1.0 | 0.989 | 0.019 | 0.371 | 0.197 | 470 |
| 0.5 | 1.000 | 0.000 | 0.653 | 0.311 | 601 |

温度 1.0 时 98.9% 是合法句子（8-gram 是 96.6%）。降低温度（把概率分布"变尖"）后，每一次"掷骰子"都更偏向最可能的选项：T3 的比例从 0.371 升到 0.653，含 pittsburgh 的句子从 0.197 升到 0.311。**一个反直觉的结果**：不同句子的个数反而从 470 个**增加**到 601 个。我原本以为"低温 = 更单调"，但最可能的句型 T3 恰好是组合最多的（9000 种），而 T3 里的人群、活动、时间是均匀分布——温度改变不了均匀分布。"低温更单调"只在每一步的分布都有明显高峰时成立；衡量多样性要看具体用什么指标。第 9 讲（解码）会系统地讨论这类取舍。

### 3.9 规模与数据都重要

slides 引用 GPT-3 的结果（在 300B token 上训练，不同大小模型的困惑度曲线）得出结论：**模型规模和数据规模都重要**。03 号脚本第 6 部分在城市语言上做了一个迷你版：4 种宽度（d = 8、16、32、64，2 层）× 3 种数据量，每个模型都训练 1000 步（步数相同，所以计算量只随模型大小变化）。表中是测试集超出下界的损失（括号里是训练集上的），单位毫 nat/token：

| 训练数据 | d=8（1,760 参数） | d=16（6,592） | d=32（25,472） | d=64（100,096） |
|---|---|---|---|---|
| 500 句 | 137.6（112.4） | 101.8（−5.9） | **85.6**（−13.0） | 178.3（−56.9） |
| 2000 句 | 78.3（73.9） | 23.6（12.4） | 32.5（20.3） | **16.5**（−0.3） |
| 20000 句 | 93.4（92.7） | 12.7（12.2） | 14.1（14.4） | **6.1**（6.4） |

（参数数不含嵌入。）

- **数据少时，大模型会过拟合。** 500 句时最好的是 d = 32；d = 64 在训练集上比真实分布还"好"（−56.9，意思是它把训练句背了下来、给它们的概率比真实概率还高），测试集却最差（178.3）。
- **数据多时，大模型最好。** 20000 句时 d = 64 最好（6.1）。
- **模型太小，数据再多也没用。** d = 8 无论给多少数据都停在 78 以上——它的容量装不下这门语言的规律。

所以只加一样，迟早会碰到另一样的天花板。（这张表每格只训练了一次，十几个毫 nat 以内的差异——比如 d = 16 和 d = 32 谁更好、d = 8 在 2000 句和 20000 句之间的起伏——可能只是随机波动或优化没收敛；上面的结论看的是几十到上百毫 nat 的大趋势。）

规模定律（scaling laws）把这个观察变成了定量的公式：Kaplan 等人（2020）发现损失随参数量、数据量、计算量都呈幂律下降；Hoffmann 等人（2022，Chinchilla）进一步指出，在固定的计算预算下，参数量和训练 token 数应该**同比例**增长，经验上大约是"每个参数 20 个 token"。§5 会用这个经验法则来确定 100B 模型的训练数据量。

---

## 4. LLM 系统的挑战 · Challenges in LLM Systems

### 4.1 现代 LLM 是什么样的系统

slides 用三个词概括现代 LLM：

- **通才（generalist）**：把任何 AI 任务都表述成"基于 token 的序列生成"。
- **可指令（instructibility）**：一个交互式系统，能接受自然语言的、几乎无穷多样的命令。
- **能行动（agentic）**：调用外部工具、执行动作、接收反馈。

对系统来说，这意味着负载**既大又杂**：同一个模型要同时服务短的翻译请求、长的推理请求、成串的智能体调用（§2.2 的表）。

### 4.2 核心系统问题与衡量指标

slides 把关键的系统问题写成一句话：**用更少的资源（GPU、显存、电），更快地在更大的数据集上计算（训练 / 推理）更大的 LLM。** 这句话里每个词都对应一类指标：

| 维度 | 训练时关心 | 推理时关心 |
|---|---|---|
| 快 | 吞吐量（token/秒）、训练总时长 | 首 token 时延（TTFT）、token 间时延（TPOT）、吞吐量 |
| 少用资源 | GPU 数、显存峰值、MFU（有效算力 / 峰值算力） | 每 GPU 能服务的并发数、显存占用 |
| 便宜 | GPU 小时、美元 | 每百万 token 的成本 |
| 省电 | 总耗电（MWh） | 每 token 的能耗 |
| 质量不变 | 收敛后的损失 / 下游指标 | 困惑度 / 下游指标不下降（§3.7） |

这些指标经常互相冲突：把更多请求凑成一批能提高吞吐，却会增加每个请求的时延；量化能省显存，却可能损害质量。系统设计的大部分工作，就是在这些冲突之间做取舍。

### 4.3 正确的抽象与权衡

slides 给出两条系统设计的原则：

- **正确的抽象（right abstraction）**：设计有用的"积木"，把复杂性藏起来，让上层的开发者不用操心。例如写模型的人调用 `torch.softmax`，不用知道它在 GPU 上怎样实现；用 DDP 的人只写单卡的训练代码，不用自己写梯度同步。
- **权衡（tradeoffs）**：先搞清楚什么是根本的约束、什么是主要的成功指标。例如推理服务里，显存容量是硬约束，时延是指标，吞吐是成本。

一个好的抽象值多少？05 号脚本第 3 部分对一个 256×256 的矩阵按行做 softmax，三种写法算的是同一个函数（结果只差 1.5×10⁻⁸），速度却差了一个多数量级（⏱ 本机）：纯 Python 循环约 3.4 毫秒；NumPy（max、exp、sum、除法四次遍历）约 0.11 毫秒，快约 30 倍；PyTorch 的 `torch.softmax`（一个专门的内核）约 0.065 毫秒，快约 53 倍。纯 Python 慢在每个元素都要经过解释器；NumPy 和 PyTorch 的差距则来自内核本身的写法——NumPy 的写法会把中间结果写出来再读回去。**写模型的人只需要选对积木，做积木的人负责把它做快**——这门课两边都要学。

### 4.4 语言模型里的计算：从网络层到底层算子

slides 列出了语言模型里常见的网络层：多头注意力、LayerNorm、dropout、线性层、非线性激活（tanh、ReLU、GELU）、softmax。它们最终都归结为四类**底层运算**：

| 底层运算 | 例子 | 原课作业 |
|---|---|---|
| 矩阵 / 张量乘法 | 线性层、注意力的 Q·Kᵀ 和 权重·V | HW1 的 MatrixMultiply 内核 |
| 归约（reduction） | 求和、求平均、求最大值（softmax 的分母、LayerNorm 的均值和方差） | HW1 的 reduce 内核 |
| 逐元素映射（map / zip） | 加偏置、残差相加、GELU、dropout、缩放 | HW1 的 map、zip 内核 |
| 数据搬运 | 拷贝、转置、拼接、在 CPU 和 GPU 之间 / GPU 和 GPU 之间传输 | 贯穿全课 |

这四类运算的**性格**完全不同。衡量它的工具是**算术强度**：一个运算的 FLOPs ÷ 它必须读写的字节数。硬件也有一个对应的数——**平衡点**（ridge point）= 峰值算力 ÷ 显存带宽。H100 的 BF16 张量核心标称 989 TFLOP/s、显存带宽 3.35 TB/s，平衡点约 295 FLOP/字节：算术强度低于 295 的运算，GPU 还没算完就在等数据，叫**访存受限**（memory-bound）；高于 295 的叫**计算受限**（compute-bound）。理想耗时用 **roofline 模型**估计：$t \approx \max(\text{FLOPs} / \text{峰值算力},\ \text{字节数} / \text{带宽})$。画成图（对数坐标）就是一条"屋顶线"：

```
 能达到的算力 FLOP/s（对数）
      │                    ┌─────────────────────────────  峰值算力（H100 BF16 约 989 TFLOP/s）
      │                   ╱       计算受限区：再提高算术强度也不会更快
      │                 ╱    ● 训练时的大矩阵乘法（强度 ~1600）
      │               ╱
      │             ╱  ● 训练时的 Q·Kᵀ（强度 ~114）
      │           ╱
      │         ╱   访存受限区：斜率 = 显存带宽（3.35 TB/s），能达到的算力 = 强度 × 带宽
      │       ╱
      │     ╱ ● 解码时的矩阵乘法（强度 ~1）
      │   ╱ ● softmax / GELU / LayerNorm（强度 0.2–2）
      └──────────────┼──────────────────────────────────  算术强度 FLOP/字节（对数）
                     平衡点 ≈ 295
```

一个运算落在斜线上，说明它被带宽卡住了：想让它更快，要么少读写字节（融合、量化、FlashAttention），要么提高它的算术强度（批处理）；提升 GPU 的峰值算力对它毫无帮助。

05 号脚本第 1 部分把一层 GPT 式 Transformer 拆成 12 个算子，用公开规格算了一遍（假设每个算子单独执行、中间结果都写回显存）。"假想 100B"模型（d = 10240）对一条 2048 token 的序列做训练前向时：

| 算子 | 类型 | 算术强度 FLOP/字节 | 瓶颈 | 占 FLOPs | 占理想时间 |
|---|---|---|---|---|---|
| QKV 投影 | 矩阵乘法 | 1617 | 计算 | 24.18% | 20.7% |
| Q·Kᵀ | 矩阵乘法 | 114 | **访存** | 1.61% | 3.6% |
| 缩放 + 掩码 + softmax | 归约+逐元素 | 1.2 | 访存 | 0.03% | 6.4% |
| 注意力权重·V | 矩阵乘法 | 114 | 访存 | 1.61% | 3.6% |
| 输出投影 | 矩阵乘法 | 1463 | 计算 | 8.06% | 6.9% |
| MLP 升维 / 降维 | 矩阵乘法 | 1638 | 计算 | 各 32.24% | 各 27.6% |
| GELU | 逐元素 | 2.0 | 访存 | 0.01% | 1.6% |
| LayerNorm ×2、残差相加 ×2 | 归约 / 逐元素 | 0.2–2.0 | 访存 | ≈0 | 合计 2.0% |

汇总起来：

| 场景 | 矩阵乘法占 FLOPs | 矩阵乘法占理想时间 | 一层的理想时间 |
|---|---|---|---|
| 假想 100B，训练 / prefill（1 × 2048 token） | 99.95% | 90.0% | 6.288 ms |
| GPT-2 small（d = 768），训练（8 × 1024 token） | 99.43% | 57.3% | 0.440 ms |
| 假想 100B，解码（1 条序列生成 1 个 token） | 99.95% | 100.0% | 0.777 ms |

三个结论，分别是后面三讲的出发点：

1. **"小"算子不做运算，却占时间。** softmax、GELU、LayerNorm、残差相加几乎不贡献 FLOPs，却要把整块数据读一遍写一遍。模型越小越明显：GPT-2 small 上矩阵乘法只占 57.3% 的时间。把几个小算子合成一个内核、中间结果不落显存，叫**算子融合**（kernel fusion）——第 10–11 讲（LightSeq）的核心。
2. **注意力里的矩阵乘法也是访存受限的。** Q·Kᵀ 要把"头数 × T × T"的分数矩阵写回显存，softmax 再读出来、写回去。**FlashAttention**（第 21 讲）的核心就是不把这个矩阵写出去。
3. **解码时一切都是访存受限的。** 每个 token 只用一次每个权重，算术强度约为 1。一层 0.777 毫秒几乎全是"读权重"的时间——这就是 LLM 推理服务的根本矛盾（§4.7、§5.7）。

### 4.5 不同抽象层次上的挑战

slides 用一张图把系统挑战按抽象层次分开：

| 层次 | 要解决的问题 | 对应讲次 |
|---|---|---|
| 分布式 / 并行系统 | 巨大的模型、巨大的数据集、很长的序列；怎样切分、调度、通信 | 第 14–18 讲（训练）、第 22–27 讲（服务） |
| 深度学习框架 | 让开发者容易写、容易改模型和算法（自动微分、模块化） | 第 5 讲（miniTorch）、第 12 讲（JAX） |
| 数据块上的算子 | 快速的 CUDA / TPU 内核；数据与模型的压缩 | 第 2–4、10–11、13、19–21 讲 |

后面几乎每一讲都可以放进这张表的某一格。

### 4.6 模型–算法–系统协同设计

slides 强调：LLM 需要**模型、算法、系统一起设计**——模型结构、训练 / 推理算法、软件优化（切分、调度、数据搬运、隐藏延迟）、硬件加速（专用指令）。只优化其中一层往往不够。几个本课会讲到的例子：

| 例子 | 改了哪一层 | 换来了什么 | 讲次 |
|---|---|---|---|
| GQA（多个查询头共用一组 K、V） | 模型结构 | KV 缓存缩小 H / H_kv 倍（§5.7：假想 100B 从 3.28 MB/token 降到 328 KB/token） | 7 |
| FlashAttention | 算法 + 内核 | 数学上完全相同，但不再把 T×T 矩阵写回显存 | 21 |
| 推测解码 | 解码算法 | 输出分布不变，一次验证多个 token，缓解访存受限 | 9 |
| 量化（INT8 / INT4） | 数值表示 + 内核 | 权重字节数减半或减到 1/4：既省显存又少读字节 | 19–20 |
| MoE | 模型结构 + 并行策略 | 参数涨十倍、每个 token 的计算几乎不变，代价是 all-to-all 通信 | 17 |
| PagedAttention | 系统（内存管理） | KV 缓存不再碎片化，同样的显存装下更大的批 | 24 |

### 4.7 计算 vs 数据搬运

slides 的原话很直白：**光让计算变快是不够的，搬数据也要时间。** 它列了三种数据搬运：大模型有大量参数，在设备 / 节点之间传参数和梯度可能比计算还慢；批量的样本和单条序列的负载不同；长上下文的 LLM 需要很大的工作内存。

**本机实测（⏱，05 号脚本第 2 部分）。** 单线程下，2048×2048 的矩阵乘法约 3.1–3.3 TFLOP/s（Apple 芯片有专门的矩阵单元），而两个 64 MiB（1677 万个 float32）的数组逐元素相加只有约 9.7 GFLOP/s——慢了 300 多倍。加法每做 1 次运算要搬 12 字节，它的速度由内存带宽（约 116 GB/s）决定，和 CPU 算得多快无关。本机的平衡点约 28 FLOP/字节，H100 是 295：GPU 的算力相对带宽更"富余"，访存受限的问题在 GPU 上**更**严重。（一个例外：GELU 在本机只有约 9 GB/s，比加法慢约 9 倍——CPU 上 erf 的计算本身就很贵，所以它在 CPU 上其实是计算受限的。同一个算子在不同硬件上瓶颈可以不同。）

**搬数据要多久（05 号脚本第 4 部分，公开规格）：**

| 数据 | 大小 | PCIe 4.0 x16（32 GB/s） | InfiniBand NDR（50 GB/s） | NVLink 4（450 GB/s，单向） | H100 显存（3.35 TB/s） |
|---|---|---|---|---|---|
| 100B 模型的 BF16 权重 | 202 GB | 6.33 s | 4.05 s | 0.45 s | 0.060 s |
| 一层的 BF16 权重 | 2.52 GB | 79 ms | 50 ms | 5.6 ms | 0.75 ms |
| 一条 2048 token 序列的 KV 缓存（MHA） | 6.71 GB | 0.21 s | 0.13 s | 15 ms | 2.0 ms |

对照：这一层对 2048 个 token 做前向，在 H100 上理想耗时只有 5.21 毫秒，而经 PCIe 把这一层的权重搬过来要 79 毫秒（15 倍）。所以"把权重放在 CPU 内存里、用时再搬"（卸载 offload）的方案必须提前预取、和计算重叠（第 18 讲的 ZeRO-Offload、附加讲 A2 的 KTransformers）。不同层次的带宽差两个数量级，**"数据放在哪里"决定了能跑多快**。

**批处理为什么能救推理（⏱，05 号脚本第 5 部分）。** 解码一个 token 时，权重矩阵的每个元素只用一次；如果同时解码 b 个序列，读一次权重就能用 b 次。在本机上对一个 4096×4096 的 float32 权重（64 MiB，远大于缓存）做 $y = x W^\top$：b 从 8 增加到 64，处理的 token 多了 8 倍，每次调用的耗时只从约 0.9 毫秒变到约 1.2 毫秒；b = 256 时每个 token 的成本约是 b = 1 的 1/26–1/32，并且达到约 2.6 TFLOP/s，接近矩阵乘法的峰值（变成计算受限，再加大批就没用了）。

这里还出现了一个**意外**：b = 2、4 时每次调用（约 1.7–1.9 毫秒）比 b = 1（约 0.34–0.42 毫秒）还慢好几倍，甚至比 b = 8 慢。原因是数学库对 b = 1 走专门的矩阵-向量内核（有效带宽约 160–200 GB/s），对 b ≥ 2 换成通用的矩阵乘法内核，后者对很窄的输入并不高效。同一个库对不同形状选不同实现，"理论上的平滑趋势"会被实现细节打断——这也是 GPU 上要为特定形状专门写内核、做自动调优的原因（第 4、10 讲）。

### 4.8 编程模型的三个层次

slides 最后问：什么是构建 AI 应用的好抽象？它分了三层：

- **上层**：把模型集成进产品系统，持续追踪和改进产品质量（附加讲 A1 的应用栈）。
- **中层**：构建模型训练和推理软件、流式数据流（第 5、12、14–18、22–27 讲）。
- **底层**：GPU 内核、编译器等（第 2–4、10–13、21 讲）。

一个好的抽象，**让程序员不再操心一个或多个问题（比如性能或故障），同时还能支撑上面各种各样的应用**。这门课的大部分讲次在中层和底层；但每讲完一个技术，都值得问一句：它把什么藏了起来？代价是什么？

---

## 5. 回答总题目：训练一个 100B 模型需要多少资源？

这一节用 04 号脚本和 [`llm_budget.py`](../code/lecture01/llm_budget.py) 把 §1.2 的问题算到底。原则是**每条公式先验证、再使用**。用来估算的"假想 100B"是 GPT-3 式结构：80 层、d = 10240、80 个头（每头 128 维）、词表 50257、上下文 2048。

### 5.1 参数量：N ≈ 12·L·d²

**推导。** 一层 GPT 式 Transformer 里有：注意力的 Q、K、V、输出 4 个 $d \times d$ 矩阵（$4d^2$），MLP 的两个 $d \times 4d$ 矩阵（$8d^2$），合计 $12d^2$；再加上偏置和两个 LayerNorm 共 $13d$ 个参数。$L$ 层就是 $12Ld^2 + 13Ld$，再加上词嵌入 $Vd$ 和位置嵌入 $Td$。当 $d$ 上万时，一次项可以忽略：**$N \approx 12 L d^2$ + 嵌入**。

**验证 1：和真实建出来的模型比。** 04 号脚本第 1 部分对 4 种配置（含不共享输出层、自定义 MLP 宽度）用 PyTorch 真的建出模型数参数，和 `tiny_gpt.param_formula`、`llm_budget.param_count` 两套公式**逐个相等**（例如 109,312 = 109,312 = 109,312）。

**验证 2：和公开模型的官方数字比**（第 2 部分）：

| 模型 | 结构 | 公式 | 官方 | 12·L·d² |
|---|---|---|---|---|
| GPT-3 175B | 96 层、d = 12288 | 174.60B | 175B | 173.9B |
| Llama-3 8B | 32 层、d = 4096、GQA 32/8、SwiGLU | 8.03B | 8.03B | 6.4B |
| Llama-3 70B | 80 层、d = 8192 | 70.55B | 70.6B | 64.4B |
| Llama-3.1 405B | 126 层、d = 16384 | 405.85B | 405B | 405.9B |
| 假想 100B | 80 层、d = 10240 | 101.21B | — | 100.7B |

两点提醒：**公式要跟着结构走**——LLaMA 用 SwiGLU（3 个矩阵）、GQA（K、V 投影小得多）、不共享的大词表，所以 12·L·d² 对 Llama-3 8B 明显偏低，嵌入和输出层在它的 8B 参数里占了 13.1%；405B 的 12·L·d² 碰巧几乎等于真实值，是因为 SwiGLU 的约 9.75d² 加上 GQA 注意力的约 2.1d² 又凑成了约 12d²——**巧合不能当规律用**。

### 5.2 训练 token 数

GPT-3 175B 训练了 300B token，约每个参数 1.7 个 token。Chinchilla（2022）的规模定律指出，在固定计算预算下更好的做法是每个参数约 20 个 token。我们取后者：**D = 20 × 101.2B ≈ 2.02T token**（这是一个假设，§5.8 看它的影响）。

### 5.3 训练 FLOPs：6ND

**推导。** 前向传播时，一个 token 经过一个权重，要做一次乘法和一次加法，即 2 FLOPs，所以前向约 $2N$ FLOPs / token。反向传播对每个矩阵乘法 $Y = XW$ 要算两个梯度：对输入 $\partial L/\partial X = (\partial L/\partial Y) W^\top$（传给前一层）和对权重 $\partial L/\partial W = X^\top (\partial L/\partial Y)$（交给优化器），每个都和前向一样大，所以反向约 $4N$。合计：

$$
C_{\text{train}} \approx 6 N D
$$

**验证。** 04 号脚本第 3 部分用 PyTorch 的 `FlopCounterMode` 数了三个小模型一次前向 + 反向的 FLOPs，和精确公式 $3 \times (2W + 4LTd) \times \text{token 数}$（$W$ 是参与矩阵乘法的权重数，$4LTd$ 是注意力的 Q·Kᵀ 和 权重·V）**逐个相等**（例如 368,050,176 = 368,050,176）。

但 6ND 本身是个近似，它忽略了注意力项 $4LTd$。这一项相对于 $2N \approx 24Ld^2$ 的比例大约是 $T/(6d)$ 的量级：

| 模型 | 训练上下文 T | 注意力项占前向 FLOPs | 6·N·D 比精确值 |
|---|---|---|---|
| GPT-3 175B | 2048 | 2.7% | 偏低约 3% |
| 假想 100B | 2048 | 3.2% | 偏低约 4% |
| Llama-3 8B | 8192 | 22.2% | 偏低约 28% |
| Llama-3.1 405B | 8192 | 7.7% | 偏低约 8% |

所以 **6ND 适合做数量级估算**；上下文越长、模型越"窄"，越要把注意力项加回来。（这里的注意力按"算满 T×T"计；真实的注意力内核利用因果掩码只算一半，第 21 讲。）

对假想 100B：$C \approx 6 \times 101.2\text{B} \times 2.02\text{T} \approx 1.23 \times 10^{24}$ FLOPs（加上注意力项多 3.3%）。

### 5.4 时间与钱

$$
\text{天数} = \frac{C}{\text{GPU 数} \times \text{每卡峰值算力} \times \text{MFU}} \div 86400
$$

H100 的 BF16 标称峰值是 989 TFLOP/s，大规模训练的 MFU 通常在 30%–50%，我们取 40%（假设）。04 号脚本第 4 部分：

| GPU 数 | 天数 | GPU 小时 | 成本（按假设的 2 美元 / GPU 小时） | GPU 满载耗电 |
|---|---|---|---|---|
| 1 | 35,963（约 99 年） | 86.3 万 | 约 170 万美元 | 604 MWh |
| 1024 | 35.1 | 86.3 万 | 约 170 万美元 | 604 MWh |
| 4096 | 8.8 | 86.3 万 | 约 170 万美元 | 604 MWh |
| 16384 | 2.2 | 86.3 万 | 约 170 万美元 | 604 MWh |

**第一个结论：必须分布式。** 一块 H100 要算 99 年。理想情况下 GPU 数加倍、天数减半，总 GPU 小时和钱不变——前提是并行效率不打折扣，而这正是第 14–18 讲要争取的。（价格和耗电都是粗估：云上价格随时间和合同差别很大；耗电只算了 GPU 的满载功耗 700 W，不含 CPU、网络和制冷。）

### 5.5 显存

混合精度 + Adam 训练时，每个参数要存（ZeRO 论文的记账方式）：BF16 权重 2 字节、BF16 梯度 2 字节、FP32 主权重 4 字节、Adam 的一阶矩 4 字节、二阶矩 4 字节，**共 16 字节**。对 101.2B 参数：

| 项目 | 每参数字节 | 总量 |
|---|---|---|
| BF16 权重 | 2 | 202 GB |
| BF16 梯度 | 2 | 202 GB |
| FP32 主权重 | 4 | 405 GB |
| FP32 Adam m | 4 | 405 GB |
| FP32 Adam v | 4 | 405 GB |
| **合计（不含激活值）** | **16** | **1.62 TB** |

至少要 21 块 80 GB 的 H100 才放得下这些"训练状态"。还没算**激活值**：反向传播要用到前向时各层的中间结果。按 Korthikanti 等人（2022）的公式（每层约 $s \cdot b \cdot d \cdot (34 + 5 a s / d)$ 字节，$s$ 是序列长度、$b$ 是批大小、$a$ 是头数），一条 2048 token 的序列、不做任何重计算时要 191 GB——是一块 H100 显存的 2.4 倍。

**第二个结论：必须把模型和状态切开。** 张量并行和流水线并行切模型（第 16 讲），ZeRO 切优化器状态、梯度和参数（第 18 讲），激活重计算用"多算一遍"换显存（第 18 讲），FlashAttention 不保存 T×T 的注意力矩阵（第 21 讲）。

### 5.6 通信

假想一个"纯数据并行"的方案：1024 块卡各存一份完整模型（实际放不下，但先这样算），每块卡算自己那份数据的梯度，再用**环形 all-reduce**（第 14 讲推导）把梯度平均。每块卡要发送（也接收）约 $2 \times$ 梯度字节数：BF16 梯度 202 GB，走每卡一个 400 Gb/s 的 InfiniBand 网卡（50 GB/s）要 **8.09 秒**；如果都能走机内的 NVLink（450 GB/s 单向）也要 0.90 秒。而每步（假设全局批 400 万 token）的计算只要 **6.00 秒**。

**第三个结论：通信可能比计算还慢。** 串行执行时一步要 14.1 秒，就算通信和计算完全重叠，也要 8.1 秒——瓶颈在通信（§2.3 的工厂比喻）。所以要让通信和计算重叠（第 15 讲 DDP 的分桶）、让流量大的通信走机内 NVLink、机间只同步切过片的梯度（第 14–16 讲）、用 MoE 时还要处理 all-to-all（第 17 讲）。

### 5.7 推理

**权重。** BF16 权重 202 GB，至少要 3 块 H100 才放得下。假设用一台 8 卡 H100 服务器（共 640 GB），放完权重还剩 438 GB 给 KV 缓存。

**KV 缓存。** 每个 token、每一层要存一个键向量和一个值向量：$2 \times L \times H_{kv} \times d_h \times 2$ 字节。

| 模型 | 每 token 的 KV 缓存 | 剩余显存能存多少 token | 相当于多少条 2048 token 的序列 |
|---|---|---|---|
| 假想 100B（MHA，80 个 KV 头） | 3.28 MB | 133,539 | 65 |
| 假想 100B + GQA（8 个 KV 头） | 328 KB | 1,335,391 | 652 |

GQA 让 KV 缓存缩小到 1/10，能同时服务的序列数多 10 倍——"模型结构为系统让路"的典型例子。

**解码速度的上限。** 每解码一步（每个序列生成 1 个 token），要把全部权重读一遍（整批共享），再读每条序列自己的 KV 缓存；roofline 给出理想耗时的下限（04 号脚本第 5 部分，8 卡理想均分，上下文 2048）：

| 模型 | 批大小 | KV 缓存 | 放得下 | 每步毫秒 | 瓶颈 | token/秒（全批合计） |
|---|---|---|---|---|---|---|
| MHA | 1 | 6.7 GB | ✓ | 7.80 | 访存 | 128 |
| MHA | 16 | 107 GB | ✓ | 11.56 | 访存 | 1,384 |
| MHA | 64 | 430 GB | ✓ | 23.58 | 访存 | 2,714 |
| MHA | 256 | 1.72 TB | ✗ | — | — | — |
| GQA-8 | 1 | 0.67 GB | ✓ | 6.45 | 访存 | 155 |
| GQA-8 | 16 | 10.7 GB | ✓ | 6.83 | 访存 | 2,344 |
| GQA-8 | 64 | 43 GB | ✓ | 8.03 | 访存 | 7,971 |
| GQA-8 | 256 | 172 GB | ✓ | 12.84 | 访存 | 19,943 |

- 批大小 1 时每步约 7.8 毫秒、每秒只有 128 个 token，而算力几乎闲着：生成每个 token 都要把 202 GB 权重读一遍。
- 批越大，同一次读权重被越多序列分摊，吞吐上升（GQA 从 155 到 19,943 token/秒）。但每条序列自己的 KV 缓存也要读：MHA 批 64 时 KV 缓存（430 GB）已经比权重还大，吞吐只比批 16 多 2.0 倍；批 256 干脆放不下。
- 表里**所有情况都是访存受限**。"怎样在显存里塞下更大的批、怎样少读一些字节"，正是第 19–20 讲（量化）、第 21 讲（FlashAttention）、第 22–24 讲（连续批处理、PagedAttention）的核心。

### 5.8 敏感性：哪个假设影响最大

04 号脚本第 6 部分改变三个假设（1024 卡），看训练天数怎样变：

| GPU | MFU | 训练 token 数 | 天数 |
|---|---|---|---|
| A100 | 20% | 20×N = 2.02T | 222.7 |
| A100 | 40% | 20×N | 111.3 |
| H100 | 20% | 20×N | 70.2 |
| H100 | 40% | 20×N | 35.1 |
| H100 | 60% | 20×N | 23.4 |
| H100 | 40% | 3×N = 0.30T（GPT-3 式） | 5.3 |

- 训练 token 数从约 3×N 到 20×N：计算量差将近 7 倍——这是**模型 / 数据**的选择。
- A100 换成 H100：快约 3 倍——这是**硬件**。
- MFU 从 20% 提到 60%：快 3 倍——这是**系统**。同样的模型、同样的数据、同样的硬件，系统做得好不好，训练时间能差几倍、成本能差上百万美元。

**一页纸的总账（假想 100B，Chinchilla 式训练）**：

| 问题 | 答案 | 关键公式 | 后续讲次 |
|---|---|---|---|
| 多少参数 | 101.2B | 12·L·d² + 嵌入 | 6、7 |
| 多少数据 | 2.02T token | 20 × N（假设） | 7 |
| 多少计算 | 1.23×10²⁴ FLOPs | 6ND（+注意力项） | 6、14 |
| 多少时间 / 钱 | 1024 块 H100 约 35 天，约 86 万 GPU 小时 | C / (GPU 数 × 峰值 × MFU) | 2–4、10–13 |
| 多少显存 | 训练状态 1.62 TB + 激活值 | 16 字节 × N | 16、18、21 |
| 多少通信 | 每步 all-reduce 约 200 GB 梯度 | 2 × 梯度字节 / 带宽 | 14–17 |
| 服务一个用户 | 至少 3 块 H100；批 1 约 128 token/秒；每 token KV 3.28 MB | 权重字节 / 带宽 | 9、19–27 |

---

## 6. 课程地图 · The course map

本讲义一讲对应原课一讲，完整大纲见 [SYLLABUS.md](../SYLLABUS.md)。按单元看：

| 单元 | 讲次 | 回答 §5 总账里的哪一行 |
|---|---|---|
| ① LLM 与 GPU 编程基础 | 1 导论；2–3 GPU 编程基础；4 GPU 加速（分块）；5 深度学习框架与自动微分 | "一块 GPU 能算多快"；"框架怎样把模型变成算子" |
| ② Transformer 与 LLM 基础 | 6 Transformer；7 预训练 LLM（T5、GPT-3、LLaMA）；8 分词与嵌入；9 生成与推测解码 | "多少参数、多少计算"；"服务时每步做什么" |
| ③ 在加速器上加速 Transformer | 10–11 LightSeq（推理 / 训练的算子融合）；12 TPU 与 JAX；13 Pallas 与 Splash Attention | "怎样把 MFU 提上去" |
| ④ 分布式训练与大模型 | 14 分布式训练与 all-reduce；15 DDP；16 流水线 / 张量并行；17 MoE；18 ZeRO | "多少显存、多少通信"；"怎样用一千块 GPU" |
| ⑤ 压缩与高效注意力 | 19–20 量化（GPTQ）；21 FlashAttention | "怎样少读字节、少占显存" |
| ⑥ 服务、微调与 RL 系统 | 22 SGLang；23 LoRA / QLoRA；24 vLLM；25 RLHF 系统；26 预填充 / 解码分离；27 KV 缓存 | "怎样服务成千上万个用户"；"怎样便宜地微调和对齐" |

原课的 7 次作业沿着同一条线：HW1 写 map / zip / reduce / 矩阵乘法的 CUDA 内核（§4.4 那张表的四类运算），HW2 写自动微分，HW3 搭一个 decoder-only Transformer，HW4 写融合的 softmax 和 LayerNorm 内核（§4.4 的结论 1），HW5–6 做数据并行、流水线并行、ZeRO + LoRA 和 SGLang 推理，HW7 做 RLHF。

本讲的概念在后面哪里展开：

| 本讲的概念 | 在哪里展开 |
|---|---|
| 注意力掩码、KV 缓存 | 第 6 讲（实现）、第 9 讲（解码）、第 22、24、27 讲（管理） |
| 打包、数据并行的大批量 | 第 14–15 讲 |
| 困惑度作为"质量底线" | 第 19–20 讲（量化后的困惑度）、第 23 讲 |
| 算术强度、roofline | 第 2–4 讲（GPU）、第 10、21 讲（融合、FlashAttention） |
| 16 字节 / 参数、激活值 | 第 18 讲（ZeRO、重计算） |
| 环形 all-reduce | 第 14 讲 |
| GQA、SwiGLU、RoPE | 第 7 讲 |
| 批处理与访存受限的解码 | 第 22、24、26 讲 |

---

## 7. 怎样学这门课 · How to approach this course

1. **先估算，再动手。** 遇到任何一个系统问题，先用 §5 的那几条公式算一下理想情况下的上下限。估算错了也没关系，错在哪里本身就很有信息量（本讲的 04 号脚本里，"6ND 在大模型上只差几个百分点"就被 Llama-3 8B 的数字推翻了）。
2. **测量，而不是猜测。** 本讲 05 号脚本里"批大小 2 比批大小 1 还慢"这种现象，不测是想不到的。看到反常的数字，先确认测量方法没错，再去找原因。
3. **找瓶颈。** 优化不是瓶颈的部分毫无意义（§2.3）。每一讲都问：瓶颈在计算、访存、通信还是调度？
4. **没有 GPU 也能学。** 原课作业在 GPU 集群上完成；本讲义的代码全部在 CPU 上运行：GPU 内核用模拟器执行并逐项计数，性能用公开规格估算，多卡用"列表里的多个设备"模拟。原理是一样的，而且每一个"为什么这样更快"都能被数出来。
5. **读原始论文。** 这门课的阅读材料几乎都是近几年的系统论文（LightSeq、ZeRO、Megatron、GPTQ、FlashAttention、vLLM、SGLang、DistServe……）。每讲的"延伸阅读"会列出它们。

---

## 8. 小结 · Summary

- LLM 的规模七年涨了一万倍，远快于单块硬件的增长；**差出来的部分只能靠系统补**。2024 年以后最大的模型大多是 MoE，总参数和每个 token 的激活参数分开了。
- 每种 LLM 能力都对应一种**负载形状**（输入长度、输出长度、调用次数、时延要求），服务系统的优化都是针对这些形状做的。
- 语言模型 = **下一个 token 的条件概率**；链式法则是恒等式，模型之间的区别只在怎样近似每个条件概率。看多长的上下文、能不能在上下文之间共享统计信息，决定了模型质量。
- 三类语言模型的差别只在**注意力掩码**；因果掩码让 decoder-only 模型训练时每个 token 都是训练信号、生成时能用 **KV 缓存**（我们逐位验证了缓存的正确性，并测出生成 59 个 token 时省下 32.6 倍的 FLOPs）。
- 预训练最小化**交叉熵**；交叉熵 = 不可约的熵 + 可约的 KL。困惑度 = exp(交叉熵)。在城市语言上，小 GPT 的可约损失只有最好 n-gram 的 14%。**规模和数据都重要**：数据少时大模型过拟合，模型太小时数据再多也没用。
- LLM 系统的核心问题：**用更少的资源、更快地、在更大的数据上计算更大的模型**。网络层归结为**矩阵乘法、归约、逐元素映射、数据搬运**四类底层运算；用**算术强度**和 **roofline** 判断瓶颈。"小"算子不贡献 FLOPs 却占时间；注意力在训练时也访存受限；**解码时一切都是访存受限的**。
- 训练一个 100B 模型：约 1.2×10²⁴ FLOPs，1024 块 H100（MFU 40%）约 35 天、86 万 GPU 小时；训练状态 1.62 TB；纯数据并行时梯度通信比计算还慢。服务它：至少 3 块 H100，批 1 时每秒约 128 个 token，每个 token 的 KV 缓存 3.28 MB（GQA 降到 1/10）。**必须分布式、必须切分、必须重叠通信、必须批处理**——这就是这门课后面所有内容的出发点。

## 9. 常见误区 · Common pitfalls

1. **"困惑度越低越好，最低能到 1。"** 只有在数据完全确定时才能到 1。语言本身有随机性，任何模型的交叉熵都不可能低于真实分布的熵（城市语言里困惑度的下界是 2.118）。比较困惑度时还要保证**词表和测试集相同**——换一个分词器，困惑度就不可比了（第 8 讲）。
2. **"链式法则是一种近似。"** 链式法则是恒等式；近似发生在用模型估计每个条件概率的时候（例如 n-gram 只看最近 n−1 个 token）。
3. **"KV 缓存是一种近似 / 会改变输出。"** 对因果模型，KV 缓存是**精确**的：02 号脚本里缓存与不缓存的 logits 只差浮点舍入（2×10⁻⁷），贪心生成的结果完全相同。它只适用于因果注意力。
4. **"FLOPs 少了就会同样比例地变快。"** 02 号脚本里 FLOPs 少了 32.6 倍，⏱ 实测只快了约 1.6 倍；05 号脚本里 softmax 只占 0.03% 的 FLOPs，却占 6.4% 的理想时间。时间取决于瓶颈：计算、访存、开销中最慢的那个。
5. **"6ND 是精确公式。"** 它忽略了注意力项，长上下文、窄模型时会明显偏低（Llama-3 8B 在 8K 上下文时偏低约 28%）；也不含逐元素运算和通信。
6. **"训练显存 = 参数量 × 2 字节（BF16）。"** 混合精度 + Adam 训练要 16 字节/参数，还要加上激活值；2 字节/参数只是推理时权重的大小。
7. **"GPU 数加倍，训练就快一倍。"** 只有在计算之外的开销（通信、同步、负载不均）都能被重叠或摊薄时才成立。§5.6 的例子里，通信本身就比计算慢。
8. **"推理比训练便宜，不需要优化。"** 单次推理便宜，但模型会被调用几十亿次；而且解码是访存受限的，单个请求只能用上 GPU 算力的一小部分（§5.7 的批 1：每秒 128 个 token）。
9. **"低温采样 = 更单调。"** 只有在每一步分布都有明显高峰时才成立。03 号脚本里，低温反而让不同句子从 470 个增加到 601 个，因为它把句型推向了组合最多的 T3。

## 10. 练习 · Exercises

**练习 1（困惑度）。** (a) 证明：一个对 $V$ 个 token 均匀猜测的模型，困惑度正好是 $V$。(b) 城市语言的熵率是 0.7506 nat/token，换成 bit/token 是多少？对应的困惑度是多少？

> **参考答案**：(a) 每个位置 $P = 1/V$，交叉熵 $= -\log(1/V) = \log V$，困惑度 $= \exp(\log V) = V$。(b) $0.7506 / \ln 2 \approx 1.083$ bit/token；困惑度 $= e^{0.7506} \approx 2.118$（用 bit 算是 $2^{1.083}$，结果一样——困惑度与对数的底无关）。

**练习 2（链式法则与马尔可夫假设）。** 01 号脚本第 3 部分里，trigram 近似把 "the students from pittsburgh who swim at night love the bridges ." 的概率低估了 4.00 倍。只看那张表，说明这个 4.00 是怎么来的。

> **参考答案**：12 个位置里只有第 11 个（bridges）的两个概率不同：真实 0.6000，trigram 0.1501；其余位置相同，连乘时相互抵消。比值 0.6000 / 0.1501 ≈ 4.00。

**练习 3（n-gram 的两种误差）。** 在 01 号脚本第 4 部分的表里，为什么 2000 句数据时最好的 n 是 6，而 20000 句时是 8？如果有 200 万句数据，你预计最好的 n 和它的困惑度会怎样？

> **参考答案**：n 越大近似误差越小（无限数据时 n = 8 到达下界），但需要估计的上下文组合越多、估计误差越大。数据越多，估计误差越小，最优的 n 就往大的方向移。有 200 万句时，n = 8 的估计误差会继续缩小，困惑度会更接近 2.123；n > 8 不再降低近似误差（已经是 0），只会增加估计误差，所以最优的 n 大概率仍是 8（或者和 n = 10 几乎打平）。写作时我用 80000 句试过：n = 8、10、14 的困惑度分别是 2.132、2.136、2.183，趋势与此一致。你可以把 01 号脚本里 `make_splits` 的 `n_train` 改大来验证（注意运行时间会变长）。

**练习 4（KV 缓存为什么只适用于因果模型）。** 用一两句话解释：为什么 BERT 式的双向模型在逐词生成时不能用 KV 缓存？02 号脚本里哪个数字直接说明了这一点？

> **参考答案**：双向注意力里，旧位置能"看到"新加入的 token，所以每加入一个 token，所有旧位置每一层的隐藏状态（以及由它们算出的 $k$、$v$）都会改变，之前存下的结果就作废了。02 号脚本第 3 部分：追加一个 token 后，双向模型旧位置的隐藏状态最大变化 0.817（因果模型只有 1.19×10⁻⁶）。

**练习 5（参数量）。** 一个 GPT 式模型：24 层、d = 2048、16 个头、词表 50257、上下文 2048、输出层与嵌入共享。用 §5.1 的公式算出它的参数量，并说出 12·L·d² 占了多少。

> **参考答案**：$12Ld^2 = 12 \times 24 \times 2048^2 = 1{,}207{,}959{,}552$；一次项 $13Ld = 638{,}976$；词嵌入 $Vd = 102{,}926{,}336$；位置嵌入 $Td = 4{,}194{,}304$；最后的 LayerNorm $2d = 4{,}096$。合计 **1,315,723,264**（约 1.32B），其中 12·L·d² 占约 91.8%，词嵌入约 7.8%，位置嵌入约 0.3%。（可以用 `llm_budget.param_count(ModelShape("ex", 24, 2048, 16, 50257, 2048))` 验证。）

**练习 6（训练计算量）。** Llama-3 8B（8.03B 参数）用了约 15T token 预训练。(a) 用 6ND 估算训练 FLOPs。(b) 按 H100 峰值 989 TFLOP/s、MFU 40%，要多少 GPU 小时？(c) 它的训练上下文是 8192，用精确公式（含注意力项）重算 (b)。

> **参考答案**：(a) $6 \times 8.03 \times 10^9 \times 15 \times 10^{12} \approx 7.23 \times 10^{23}$ FLOPs。(b) $7.23 \times 10^{23} / (989 \times 10^{12} \times 0.4) / 3600 \approx 50.7$ 万 GPU 小时。(c) 精确公式给出 $8.69 \times 10^{23}$ FLOPs（多约 20%），约 61.0 万 GPU 小时。（这两个数是在假设的 MFU 下的理想估算，不是 Meta 报告的实际用量。）

**练习 7（KV 缓存的显存）。** Llama-3 70B：80 层、8 个 KV 头、每头 128 维、BF16。(a) 每个 token 的 KV 缓存是多少字节？(b) 同时服务 32 条长度 8192 的序列，KV 缓存一共多大？和它的 BF16 权重（约 141 GB）比呢？

> **参考答案**：(a) $2 \times 80 \times 8 \times 128 \times 2 = 327{,}680$ 字节 = 320 KiB。(b) $327{,}680 \times 8192 \times 32 \approx 85.9$ GB，约是权重的 61%。如果它用 64 个 KV 头的标准多头注意力，这个数会变成 8 倍（约 687 GB）——GQA 的价值一目了然。

**练习 8（解码速度的上限）。** 在一块 H100（显存带宽 3.35 TB/s）上用 BF16 跑 Llama-3 8B（权重约 16.06 GB），批大小 1。(a) 忽略 KV 缓存，每秒最多能生成多少个 token？(b) 为什么提高 GPU 的算力对 (a) 没有帮助？

> **参考答案**：(a) 每生成一个 token 要把全部权重读一遍：$16.06 \text{ GB} / 3.35 \text{ TB/s} \approx 4.8$ 毫秒，即每秒最多约 209 个 token。（加上上下文 8192 的 KV 缓存，`decode_step_seconds` 给出约 196 token/秒。）(b) 批 1 时算术强度约为 1 FLOP/字节，远低于 H100 的平衡点 295，计算单元大部分时间在等数据；瓶颈是显存带宽，不是算力。要提速只能少读字节（量化）、或者让一次读取服务更多 token（批处理、推测解码）。

**练习 9（算术强度）。** 两个各有 $10^9$ 个元素的 BF16 张量逐元素相加。(a) 它的算术强度是多少？(b) 在 H100 上按 roofline 估算耗时，计算和访存各要多久？

> **参考答案**：(a) 每个元素 1 次加法，读 2 个数、写 1 个数，共 6 字节：$1/6 \approx 0.17$ FLOP/字节。(b) 计算：$10^9 / 989 \times 10^{12} \approx 1.0$ 微秒；访存：$6 \times 10^9 / 3.35 \times 10^{12} \approx 1.79$ 毫秒。访存比计算慢约 1800 倍——逐元素运算是典型的访存受限运算，这正是算子融合（第 10 讲）要解决的问题。

**练习 10（瓶颈）。** §2.3 的工厂里，如果把二号反应器换成一个 400 吨的（仍然要 4 小时），整条线的产能变成多少？瓶颈换成了哪道工序？这对系统优化有什么启示？

> **参考答案**：二号反应器的产能变成 100 吨/小时，新的瓶颈是锁斗（55 吨/小时），整条线从 50 提高到 55 吨/小时——把瓶颈的产能翻倍，整体只提高了 10%。启示：消除一个瓶颈之后，瓶颈会转移到下一个最慢的环节；每次优化后都要重新测量、重新找瓶颈。

**练习 11（规模与数据）。** 在 §3.9 的表里，d = 64 在 500 句数据上训练集"超出下界"是负数（−56.9），这意味着什么？为什么在 20000 句数据上就不再是负数了？

> **参考答案**：负数意味着模型给训练句子分配的概率比真实概率还高——它记住了这些具体的句子（过拟合），而这些概率在没见过的测试句上就不成立了，所以测试集损失最高（178.3）。20000 句数据里，同样大小的模型已经无法逐句记忆（每个具体句子出现得太少、需要记的太多），只能学习一般规律，于是训练集和测试集的损失都接近下界（6.4 和 6.1）。

**练习 12（思考题：协同设计）。** §5.7 的表里，GQA-8 在批 256 时仍然是访存受限的。设想你可以改动模型、算法、系统三者中的任何一个，列出至少三种能提高这台 8 卡服务器解码吞吐的办法，并说出每种办法改的是哪一层、代价是什么。

> **参考答案（不唯一）**：① 权重量化到 INT4（模型 / 数值表示）：读权重的字节数降到 1/4，代价是可能损失精度（第 19–20 讲）。② KV 缓存量化或压缩（系统 / 数值表示）：降低每条序列要读的字节、能装下更大的批，代价是精度和编解码开销（第 27 讲）。③ 推测解码（算法）：用一个小模型先猜几个 token、大模型一次验证，一次读权重产生多个 token，代价是草稿模型的计算和猜错时的浪费（第 9 讲）。④ PagedAttention（系统）：消除 KV 缓存碎片，同样的显存装下更多序列（第 24 讲）。⑤ 更少的 KV 头 / MQA，或 MoE（模型结构）：分别缩小 KV 缓存或每个 token 的计算，代价是可能损失质量或增加通信（第 7、17 讲）。

## 11. 延伸阅读 · Further reading

- **原课材料**：第 1 讲 slides（llmsys-01-intro.pdf）；课程主页 https://llmsystem.github.io/llmsystem2026spring/ ；作业网站 https://llmsystem.github.io/llmsystemhomework/ 。
- **规模与训练计算量**：Brown et al. (2020) GPT-3（表 2.1 的模型形状、训练 token 数）；Kaplan et al. (2020) *Scaling Laws for Neural Language Models*（2.1 节的 FLOPs 推导）；Hoffmann et al. (2022) Chinchilla（计算最优的参数 / 数据配比）。
- **显存与激活值**：Rajbhandari et al. (2020) *ZeRO*（16 字节/参数的记账，第 18 讲的阅读材料）；Korthikanti et al. (2022) *Reducing Activation Recomputation in Large Transformer Models*（激活值公式）。
- **roofline 模型**：Williams, Waterman & Patterson (2009) *Roofline: An Insightful Visual Performance Model for Multicore Architectures*。
- **硬件规格**：NVIDIA H100 / A100 的数据手册（本讲用的是 SXM 版本的 dense BF16 标称值）。
- **系统化的估算练习**：Austin et al. (2025) *How to Scale Your Model*（Google DeepMind 的在线书，对 TPU 上的同类估算讲得很细，和第 12–13 讲互补）。

## 12. 下一讲预告 · Next lecture

slides 的最后一页写着下一讲的两个主题：**GPU 内部是什么样的，以及怎样给 GPU 编程**。第 2 讲会从本讲 §4.4 那张"底层运算"的表出发：一个文本分类网络怎样拆成矩阵乘法、逐元素运算和归约；一台 GPU 服务器由哪些部件组成、数据在 CPU、内存、PCIe、NVLink 和 GPU 之间怎样流动；GPU 由哪些 SM 组成、为什么能同时跑上万个线程；一个 CUDA 内核是怎样被切成网格、线程块和线程束来执行的。我们会写一个 Python 的 SIMT 模拟器，让每一个 CUDA 概念都能被运行、被计数——这也是原课 HW1（map、zip、reduce、矩阵乘法内核）的起点。
# 第 2 讲 · GPU 编程基础（一）：GPU 服务器、GPU 架构与 CUDA 执行模型 · GPU Programming Basics 1

> **对应 Maps to**：CMU 11-868 / 11-968 *Large Language Model Systems*（Spring 2026，Lei Li），Lecture 2 · Wed 1/14 · **GPU Programming Basics 1**。
> **阅读材料**：Hwu, Kirk & El Hajj, *Programming Massively Parallel Processors*（PMPP，第 4 版）第 2 章（异构数据并行计算）、第 4 章（计算架构与调度）；NVIDIA *CUDA C++ Programming Guide* 的 "Programming Model" 与 "Hardware Implementation" 两章；NVIDIA A100 / H100 架构白皮书。
> **依据说明**：本讲以原课第 2 讲公开的 slides（llmsys-02-gpu-programming.pdf，36 页）为主线：从神经网络层到底层算子、GPU 服务器的组成与通信、GPU 产品线与 SM 结构、CPU vs GPU、主机与设备、SIMT、线程块与线程束、CUDA 内核与编译。slides 里的 GPU 对比表、服务器配置、A6000 / H100 的 SM 参数都来自原课；**§0.4 的 SIMT 模拟器、全部实验、占用率 / 波次 / 延迟隐藏的计算、"为什么是 32"的玩具模型都是本讲义自己设计的，不代表课堂上讲了什么**。
> **配套代码 Code**：[`code/lecture02/`](../code/lecture02/)（5 个脚本 + `simt.py` 模拟器 + `gpu_model.py` 纸面 GPU；每个脚本 1 秒左右，输出重复运行逐位相同）
> **预计阅读时间**：4 小时（含运行代码）

---

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

### 0.1 学习目标 Learning goals

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

1. 把一个神经网络（以 slides 的影评分类网络为例）拆成**五种底层运算**：查表搬运、矩阵乘法、逐元素映射（map / zip）、归约（reduce），并数出每一种的 FLOPs 和读写字节，判断谁是瓶颈。
2. 说出一台 **GPU 服务器**由哪些部件组成（CPU、内存、NVMe、GPU、PCIe、NVLink），估算数据在它们之间流动的时间，并解释**为什么数据并行训练时梯度走哪条路很重要**。
3. 读懂 GPU 规格表：区分 **dense 和稀疏**的峰值、算出**平衡点**；从"SM 数 × 每 SM 核心数 × 2 × 频率"推出峰值算力；说出一个 **SM** 里有什么（4 个分区、CUDA 核心、张量核心、寄存器、L1 / 共享内存）。
4. 解释 **CPU 和 GPU 的设计取向**为什么不同（延迟优先 vs 吞吐优先）。
5. 描述 CUDA 的**执行模型**：主机与设备、内核、网格、线程块、线程、**线程束（warp）**、**SIMT**；写出 `i = blockIdx.x * blockDim.x + threadIdx.x` 并知道它的边界情况。
6. 用**占用率**、**波次**和**延迟隐藏**解释"为什么 GPU 要同时驻留几万个线程"，并用 Little 定律算出"需要多少个线程束才能用满带宽"。
7. 识别**线程束分化**和**非合并访存**，知道它们为什么让内核变慢、怎样避免。

### 0.2 先修与本讲范围 Prerequisites & scope

| 默认你已经会 Assumed | 本讲只给直觉，后面再展开 Previewed only |
|---|---|
| C 的指针、数组、行优先存储（[前置讲 P1](prereq1_python_to_c.md)，特别是 §7 逐行读 CUDA 程序） | `cudaMalloc` / `cudaMemcpy` 的细节、多维网格、HW1 的四个内核（第 3 讲） |
| 存储层次、Little 定律、并发与竞争（[前置讲 P4](prereq4_systems_crash_course.md)） | 共享内存与分块、合并访存与 bank 冲突的优化（第 4 讲） |
| 算术强度、平衡点、roofline（[第 1 讲](lecture01_intro_llm_systems.md) §4.4） | 张量核心的编程、算子融合（第 10–11、21 讲） |
| 线性层、ReLU、softmax（[前置讲 P3](prereq3_neural_network_crash_course.md)） | 多卡通信与 all-reduce 的推导（第 14–15 讲） |

### 0.3 本讲地图 Map

```
§1  从网络层到底层运算：一个影评分类网络 = 五种运算                     ← 代码 01
§2  GPU 服务器：部件、带宽、为什么梯度怎么走很重要                      ← 代码 02 第 5 部分
§3  GPU 架构：规格表、SM、张量核心、CPU vs GPU                         ← 代码 02 第 1–4 部分
§4  程序在 GPU 上怎样执行：主机与设备、内核、网格 / 块 / 线程、线程束、  ← 代码 03、04、05
    分化、占用率、波次、延迟隐藏、编译
§5  小结   §6  常见误区   §7  练习   §8  延伸阅读   §9  下一讲预告
```

建议的阅读顺序：§1 和 §2–§3 可以先读，不依赖模拟器；读 §4 之前先看一遍 §0.4，再依次运行 `03`、`05`、`04`。

### 0.4 本讲的工具：一个 Python 写的 SIMT 模拟器

原课的作业要在 GPU 集群上编译运行 CUDA。本讲义假设你只有一台普通电脑，所以写了一个**SIMT 模拟器** [`simt.py`](../code/lecture02/simt.py)：内核写成 Python 函数，模拟器按 CUDA 的规则为网格里的每个线程调用它，并且把每一次显存访问都记下来。它比真 GPU 慢上百万倍，但它能**数**——这正是学 GPU 编程最需要的：

| CUDA C | 模拟器里的写法 |
|---|---|
| `__global__ void k(float* A, int n)` | `def k(t, A, n):`（第一个参数 `t` 提供内置变量） |
| `threadIdx.x`、`blockIdx.y`、`blockDim.x`、`gridDim.x` | `t.threadIdx.x`、`t.blockIdx.y`、`t.blockDim.x`、`t.gridDim.x` |
| `k<<<grid, block>>>(dA, n);` | `launch(k, grid, block, dA, n)`，返回统计 `LaunchStats` |
| `cudaMalloc` / `cudaMemcpy` / `cudaFree` | `dev.malloc` / `dev.memcpy` / `dev.free`（或便捷的 `dev.to_device`、`dev.to_host`） |
| `__shared__ float As[16][16];` | `As = t.shared("As", (16, 16))` |
| `__syncthreads();` | `yield t.syncthreads()` |
| `atomicAdd(&C[i], v)` | `C.atomic_add(i, v)` |

它统计的东西都是"为什么这样写更快"的答案：每个线程束的每条访存指令产生几个 32 字节的**内存事务**（合并访存）、共享内存有没有 **bank 冲突**、线程束里有多少线程在**空等**（分化）、有没有**数据竞争**。它**不**模拟时间——GPU 上的耗时要用 [`gpu_model.py`](../code/lecture02/gpu_model.py) 的纸面模型（公开规格 + roofline）去估。讲义里同时给出等价的 CUDA C，以后有 GPU（比如 Colab 的免费 T4）可以对照运行。

---

## 1. 从网络层到底层运算 · Layers to low-level operators

### 1.1 为什么从这里开始

第 1 讲的结论是：LLM 系统的问题，归根结底是**算得快**和**搬得少**。可是"一个 Transformer 层"太复杂，没法直接讨论怎样在硬件上算快。slides 的做法是先往下拆：**任何神经网络层，最后都落到几种底层运算上**；只要把这几种运算在 GPU 上做快，所有网络就都快了。这正是深度学习框架的分层（第 5 讲）：上层写模型的人只拼积木，下层写内核的人只管把几块积木做到极致。

原课 HW1 要写的四个 CUDA 内核——map、zip、reduce、矩阵乘法——就是这几块积木。

### 1.2 一个影评分类网络

slides 用的例子是影评情感分类（正面、负面、中性）。网络从下往上是：

```
"It is a good movie"
   → Embedding（查表：每个词 → 一个 d 维向量）
   → Linear（d → h）→ ReLU → Linear（h → C）     对每个 token 各做一遍
   → Avg（对句子里的所有 token 取平均）
   → Softmax（C 个类别的概率）
```

01 号脚本用 [P3 的影评小语料](../code/prereq3/nn_lab.py)（`ReviewWorld`）真的跑了一遍，并把每一步都改写成对底层运算的调用：

| 层 | 底层运算 | 类型 |
|---|---|---|
| Embedding | gather（按编号取出表里的行） | 数据搬运，0 次运算 |
| Linear | matmul（x · W）+ zip(add)（加偏置，广播） | 矩阵乘法 + 逐元素 |
| ReLU | map(relu) | 逐元素 |
| Avg（带掩码） | zip(mul)（乘掩码）+ reduce(sum)（沿长度求和）+ reduce(sum)（数有效 token）+ zip(div) | 逐元素 + 归约 |
| Softmax | reduce(max) + zip(sub) + map(exp) + reduce(sum) + zip(div) | 归约 + 逐元素 |

一次前向一共 15 次底层运算，**种类只有五种**。拼出来的概率和 PyTorch 的 `nn.Embedding`、`nn.Linear`、`torch.softmax` 相比，最大差 3.0×10⁻⁸——两边做的是同一组运算，只在求和顺序等细节上有 float32 的舍入差别。

### 1.3 谁算得多，谁搬得多

玩具网络太小，看不出性格。脚本第 2 部分把尺寸换成接近 GPT-2 small 的设置（32 条 × 128 token，d = 768，隐藏层 h = 3072，3 个类别），并按 H100 的 FP32 规格（67 TFLOPS、3.35 TB/s，平衡点 20 FLOP/字节）用 roofline 估算每个运算的理想耗时（假设每个运算单独执行、输入输出都经过显存）：

| 层 | 运算 | FLOPs | 读写字节 | 算术强度 FLOP/字节 | H100 理想耗时 |
|---|---|---|---|---|---|
| Embedding | gather | 0 | 25.2 MB | 0 | 7.5 μs |
| Linear1 | matmul | 193 亿 | 72.4 MB | 267 | 288.5 μs |
| Linear1 | zip(add) | 1258 万 | 100.7 MB | 0.12 | 30.1 μs |
| ReLU | map(relu) | 1258 万 | 100.7 MB | 0.12 | 30.0 μs |
| Linear2 | matmul | 7550 万 | 50.4 MB | 1.50 | 15.1 μs |
| 其余 10 个小运算 | | 合计约 4.1 万 | 约 0.28 MB | 0.10–0.25 | 合计约 0.08 μs |

按运算类型汇总：

| 运算类型 | 次数 | 占 FLOPs | 占字节 | 占理想时间 |
|---|---|---|---|---|
| 矩阵乘法 | 2 | 99.87% | 35.1% | 81.8% |
| 逐元素 | 8 | 0.13% | 57.7% | 16.2% |
| 数据搬运 | 1 | 0% | 7.2% | 2.0% |
| 归约 | 4 | 0.00% | 0.02% | 0.01% |

这张表就是第 1 讲 §4.4 的结论在一个具体网络上的样子：

- **矩阵乘法几乎包揽了全部运算**，而且算术强度高（267 远大于平衡点 20），是计算受限的。要让它快，得用好 GPU 的算力——第 3、4 讲的主题。
- **逐元素运算几乎不做运算，却读写了一半以上的字节**。它们的算术强度约 0.12，速度完全由显存带宽决定，占了 16% 的时间。要让它们快，只能少搬数据——把几个小运算合成一个内核（算子融合，第 10–11 讲）。
- Linear2 也是矩阵乘法，但它的输出只有 3 列，算术强度只有 1.5，同样是访存受限的。**"是不是矩阵乘法"不决定瓶颈，算术强度才决定。**

### 1.4 运算的顺序：一个代数改写

看一眼网络结构：Linear2 是 $x W + b$（仿射），Avg 是加权求和（线性，权重之和为 1）。**两者可以交换顺序**：先对 token 取平均、再做 Linear2，数学上完全相同。脚本第 3 部分验证了这一点：

| 顺序 | Linear2 的 FLOPs | 与原顺序的最大差 |
|---|---|---|
| 原顺序：… → ReLU → Linear2 → Avg | 75,509,760 | 0 |
| 改写：… → ReLU → Avg → Linear2 | 589,920 | 8.9×10⁻⁸ |
| 错误改写：… → Avg → ReLU → Linear2 | 589,920 | 0.077 |

Linear2 的计算量少了 128 倍（T = 128：原来对每个 token 各做一次，现在每条影评只做一次），结果只差 float32 的舍入。但**整个网络只省了 0.26%**——大头是 Linear1。如果 Avg 能再往前挪到 Linear1 之前，Linear1 也能省 128 倍；可惜中间隔着 ReLU，$\max(0,\cdot)$ 不是线性的，"平均之后再 ReLU"和"ReLU 之后再平均"不相等（概率差 0.077）。

两个启发：

1. **"不改变结果、只改变计算量"的改写，是编译器和框架做图优化的基本手段**（第 5 讲的计算图、第 12 讲的 XLA）。前提是先证明它在数学上等价，这需要知道每个运算的性质：线性、可交换、可结合。
2. **模型结构本身就是计算量的设计**。原课 HW2 的网络把顺序定成"先对嵌入取平均，再过 Linear → ReLU → Linear"——那是另一个模型（它丢掉了"每个 token 先各自非线性变换"的能力），但便宜 T 倍。第 1 讲 §4.6 说的"模型–算法–系统协同设计"，在这么小的网络上就已经出现了。

---

## 2. GPU 服务器 · Components of a GPU server

### 2.1 一台服务器里有什么

slides 给出了作者实验室里一台真实服务器的配置：一个 4U 机箱里装着 2 颗 AMD EPYC 9354 CPU、16 条 64 GB 的 DDR5 内存（共 1 TB）、4 块 15.36 TB 的 PCIe 4.0 NVMe 固态硬盘、8 块 NVIDIA RTX A6000（每块 48 GB 显存），以及 4 个双槽 NVLink 桥（把 8 块卡两两连起来）。

把它画成数据流的样子：

```
   NVMe SSD ×4 ──PCIe──┐                          ┌── GPU0 ═NVLink═ GPU1
   （每块约 7 GB/s）      │                          ├── GPU2 ═NVLink═ GPU3
                       CPU ×2 ── PCIe 4.0 x16 ─────┤
   DDR5 内存 1 TB ──────┘    （每卡约 32 GB/s）      ├── GPU4 ═NVLink═ GPU5
   （每颗 CPU 约 300 GB/s）                         └── GPU6 ═NVLink═ GPU7
                                                     每块 GPU 自己的显存：768 GB/s
```

slides 标的两个关键带宽：GPU 之间的 **NVLink 112.5 GB/s**（A6000 的 NVLink 桥，这是双向之和，单向约 56 GB/s），以及 **PCIe 4.0 x16 32 GB/s**（16 条通道 × 每条约 2 GB/s，单向）。另外，CPU 内存的带宽 slides 画的是 64–512 GB/s，A6000 自己的显存是 768 GB/s。

**带宽差了一个数量级以上**：显存 768 GB/s，NVLink 56 GB/s，PCIe 32 GB/s，NVMe 7 GB/s。数据放在哪一层、从哪条路走，决定了能跑多快——这是 P4 存储层次那一课在服务器尺度上的重演。

### 2.2 为什么和我们有关：梯度怎么走

slides 在这里抛出一个问题：考虑数据并行训练时，把梯度从一块 GPU 搬到另一块 GPU——服务器的结构为什么和这件事有关？02 号脚本第 5 部分算了一遍。场景：数据并行训练一个 13 亿参数的模型，每张卡每步处理 8 条 × 2048 = 16,384 个 token。

| 数据流 | 大小 | 时间 |
|---|---|---|
| 从 NVMe 读一步的训练数据（token 编号） | 66 KB | 0.019 ms |
| CPU 内存 → GPU（PCIe）送这批 token | 66 KB | 0.012 ms |
| 开训前：从 4 块 NVMe 读 BF16 权重 | 2.6 GB | 93 ms（只读一次） |
| GPU0 → GPU1 发一份 FP32 梯度（NVLink） | 5.2 GB | 92.5 ms |
| GPU0 → GPU2 发一份 FP32 梯度（PCIe） | 5.2 GB | 162.5 ms |
| 8 卡环形 all-reduce 梯度（环上最慢的一段是 PCIe） | 5.2 GB | 284 ms |
| 假想：8 卡全部用 NVLink 连成环 | 5.2 GB | 162 ms |

对照：每张卡这一步的计算量 $6 \cdot N \cdot \text{tokens} = 1.28 \times 10^{14}$ FLOPs，按 A6000 BF16 dense 峰值（154.8 TFLOPS）的 40% 算要 2,064 ms。

三个结论：

1. **训练时 CPU 送给 GPU 的数据很少**。一步只有 66 KB 的 token 编号——权重、优化器状态、激活值全都常驻显存。新手常以为"CPU ↔ GPU 拷贝"是训练的瓶颈，对 LLM 训练来说通常不是。（推理服务和"把权重放在 CPU 内存里"的卸载方案是另一回事，见第 18 讲和附加讲 A2。）
2. **真正大的流量在 GPU 之间**。数据并行的每一步都要把和模型一样大的梯度在卡之间交换一遍（all-reduce，第 14 讲）。这里是 284 ms，占计算时间的 14%。
3. **拓扑决定速度**。这台机器只有两两之间有 NVLink 桥，环形 all-reduce 必然要经过 PCIe，被最慢的一段卡住。如果所有卡都用 NVLink 互连（数据中心的 DGX / HGX 服务器用 NVSwitch 做到全互连），同样的通信只要 162 ms。

第 14–15 讲会讲怎样把这笔通信和反向计算重叠起来（P4 §3 的重叠、第 15 讲的分桶），把它藏住。

### 2.3 计算设备：CPU、GPU、FPGA

slides 并列了三种计算设备：CPU（AMD EPYC 9754：128 核 256 线程，2.25 GHz，256 MB L3；Intel Xeon 8593Q：64 核 128 线程，320 MB L3）、GPU（RTX A6000：10,752 个核心，48 GB，38.7 TFLOPS——slides 特意注明它比 2001 年排名第一的超级计算机还强）、FPGA（可重新配置的电路，适合定制的低延迟流水线）。LLM 系统几乎都在 GPU（以及 Google 的 TPU，第 12 讲）上运行。为什么？§3.4 回答。

---

## 3. GPU 架构 · GPU architecture

### 3.1 读懂规格表：稀疏 vs dense、平衡点

slides 给出 NVIDIA 三代数据中心 GPU 的对比（单位 TFLOPS，照抄）：

| | A100（Ampere） | H100（Hopper） | B200（Blackwell） |
|---|---|---|---|
| FP64 | 9.7 | 34 | 37 |
| FP32 | 19.5 | 67 | 75 |
| TF32 张量核心 | 312 | 989 | 2,200 |
| FP16 / BF16 张量核心 | 624 | 1,979 | 4,500 |
| FP8 张量核心 | — | 3,958 | 9,000 |
| FP4 张量核心 | — | — | 18,000 |
| 显存 | 80 GB HBM2e | 80 GB HBM3 | 192 GB HBM3e |
| 显存带宽 | 2 TB/s | 3.2 TB/s | 最高 7.7 TB/s |

读这张表有两个要点。

**第一，张量核心一栏是"含稀疏加速"的数字。** NVIDIA 产品页上张量核心的峰值默认带一个脚注：使用 **2:4 结构化稀疏**（每 4 个连续的权重里恰好有 2 个是 0，硬件跳过它们）。普通的稠密矩阵乘法只能拿到一半：H100 的 BF16 dense 是 989.4 TFLOPS，A100 是 312（第 1 讲用的也是 dense 值）。读任何规格表都要先看脚注——这是做"信封背面计算"的第一课。

**第二，平衡点在变高。** 02 号脚本第 1 部分用 dense 值算了平衡点（峰值算力 ÷ 显存带宽）：

| GPU | FP32 | BF16 张量核心 | 带宽 | FP32 平衡点 | BF16 平衡点 |
|---|---|---|---|---|---|
| RTX A6000 | 38.7 | 154.8 | 0.768 TB/s | 50 | 202 |
| A100 SXM 80GB | 19.5 | 312 | 2.04 TB/s | 10 | 153 |
| H100 SXM 80GB | 67 | 989.4 | 3.35 TB/s | 20 | 295 |
| B200 | 75 | 2,250 | 7.7 TB/s | 10 | 292 |

- 张量核心比普通 FP32 核心快一个数量级以上（H100：989 vs 67），所以训练都用 BF16 / FP8 的矩阵乘法；但平衡点也从约 20 FLOP/字节升到约 300——**越快的算力越难喂饱**。
- **每一代的算力涨得比带宽快**：A100 → B200，BF16 算力涨了 7.2 倍，显存带宽只涨了 3.8 倍。访存受限的问题一代比一代严重，这就是为什么 FlashAttention（第 21 讲）、量化（第 19–20 讲）、算子融合（第 10–11 讲）这些"少搬数据"的技术越来越重要。

### 3.2 从芯片到 SM

slides 以 RTX A6000（GA102 芯片）为例：芯片上有 **84 个 SM**（Streaming Multiprocessor，流式多处理器），12 个 32 位的显存控制器（合计 384 位宽的显存总线），所有 SM 共享 6 MB 的 **L2 缓存**。H100 的完整芯片更大，SXM 版本启用 132 个 SM（PCIe 版 114 个），配 80 GB HBM3 显存。

**SM 是 GPU 的基本计算单元**，可以理解为"GPU 里的一个核"——但它和 CPU 的核很不一样。slides 列出的 SM 内部结构（Ampere GA102）：

- 每个 SM 分成 **4 个分区**（partition / processing block），每个分区有 32 个 FP32 核心，对应一个线程束的 32 个线程；整个 SM 共 128 个核心，一个周期能做 128 次 FP32 运算（一次乘加 FMA 算作一条指令）。
- 每个分区有 **64 KB 的寄存器文件**（16,384 个 32 位寄存器，访问最快），每个 SM 共 256 KB。
- 每个 SM 有 **128 KB 的 L1 缓存 / 共享内存**（两者共用一块 SRAM，比例可配置）。H100 的 SM 增加到 256 KB。
- 每个分区还有一个**张量核心**（Tensor Core）：一条指令完成一个小矩阵块的乘加（例如 BF16 输入、FP32 累加），是矩阵乘法的主力。H100 的张量核心新增了 **FP8** 支持（slides 第 20 页），配合软件里的"Transformer Engine"按层选择精度——第 19 讲会讲 FP8 训练。
- 每个分区有一个**线程束调度器**（warp scheduler），每个周期从它管理的线程束里挑一个、发射一条指令。

### 3.3 峰值算力从哪里来

有了这些数，就能自己算峰值。02 号脚本第 2 部分：

$$
\text{FP32 峰值} = \text{SM 数} \times \text{每 SM 的 FP32 核心数} \times 2 \times \text{频率}
$$

| GPU | SM 数 | 每 SM 核心 | 核心总数 | boost 频率 | 算出来 | 规格表 |
|---|---|---|---|---|---|---|
| RTX A6000 | 84 | 128 | 10,752 | 1.80 GHz | 38.71 | 38.7 |
| A100 SXM | 108 | 64 | 6,912 | 1.41 GHz | 19.49 | 19.5 |
| H100 SXM | 132 | 128 | 16,896 | 1.98 GHz | 66.91 | 67 |

"2"来自乘加指令：$a \cdot b + c$ 一条指令算两次运算。A6000 的 10,752 个核心就是 slides 写的数字。A100 每个 SM 只有 64 个 FP32 核心（每个分区 16 个，一条线程束指令要分两个周期执行），它把芯片面积主要花在了张量核心上。

一个容易误解的地方：**"CUDA 核心"不是 CPU 意义上的核**。CPU 的每个核都能独立取指令、译码、执行；GPU 的"核心"只是一条 32 宽的向量流水线里的一个**车道**（lane），它不取指令，只执行调度器发给整个线程束的那条指令。真正能独立调度的单位是每个 SM 的 4 个线程束调度器——H100 全卡一共 132 × 4 = 528 个。

### 3.4 CPU vs GPU：延迟优先 vs 吞吐优先

slides 用一张表对比了 AMD EPYC 9754 和 RTX A6000。02 号脚本第 3 部分把它补全：

| | AMD EPYC 9754 | RTX A6000 |
|---|---|---|
| 硬件线程 / 核心 | 256（128 核 × 2） | 10,752 个 FP32 核心；最多同时驻留 129,024 个线程 |
| 频率 | 2.25 GHz（基础） | 1.8 GHz（boost） |
| FP32 峰值（slides 的算法：每核每周期 1 次乘加） | 576 GFLOPS | 38,700 GFLOPS |
| FP32 峰值（用满 CPU 的 SIMD：每核每周期 32 次） | 9.2 TFLOPS | 38.7 TFLOPS |
| 张量核心（BF16，dense） | — | 154.8 TFLOPS |
| 内存带宽 | 461 GB/s（12 通道 DDR5-4800） | 768 GB/s |
| 最后一级缓存 | 256 MB L3 | 6 MB L2 |
| 功耗 | 360 W | 300 W |

slides 的 576 GFLOPS 是按"每核每周期一次乘加"算的，而现代 CPU 也有很宽的 SIMD（AVX-512：一条指令处理 16 个 float）。同口径比较，A6000 的 FP32 只快约 4.2 倍；加上张量核心，差距是 17 倍。差距没有"576 vs 38,700"那么夸张——但这恰好说明了**真正的差别不在核心数，而在设计取向**：

- **CPU 是延迟优先的**：大缓存（256 MB L3）、乱序执行、分支预测、高频率，目的是让**少数几个线程的每一步都尽快完成**。适合逻辑复杂、分支多、前后依赖强的程序（操作系统、数据库、Python 解释器）。
- **GPU 是吞吐优先的**：缓存小（6 MB L2），把更多晶体管花在算术单元和寄存器上；单个线程很慢，但同时驻留十几万个线程，**靠数量换吞吐**。一个线程在等数据时，别的线程接着算（§4.7）。适合"同一段程序在海量数据上各算各的"——神经网络的矩阵乘法和逐元素运算正好是这样。

这也解释了 §2.2 的分工：CPU 负责读数据、调度、发射内核，GPU 负责算。

### 3.5 寄存器和共享内存：为什么是稀缺资源

02 号脚本第 4 部分把每个 SM 的资源列了出来：

| GPU | 每 SM 寄存器 | 每 SM L1+共享 | 其中共享内存最多 | 每 SM 最多线程 | 满驻留时每线程寄存器 | 全卡寄存器 |
|---|---|---|---|---|---|---|
| RTX A6000 | 256 KB | 128 KB | 100 KB | 1536 | 42 | 21.0 MB |
| A100 SXM | 256 KB | 192 KB | 164 KB | 2048 | 32 | 27.0 MB |
| H100 SXM | 256 KB | 256 KB | 228 KB | 2048 | 32 | 33.0 MB |

H100 全卡的寄存器加起来约 33 MB，和全卡的 L1 + 共享内存一样大，是 50 MB L2 的 66%——寄存器是 GPU 上容量最大的片上存储之一，而 CPU 每个核心只有几百个寄存器。为什么要这么多？因为**每个驻留的线程都有自己的一份寄存器，常驻不动**。slides 说"线程束之间的切换是瞬间的"，原因就在这里：CPU 切换线程要把寄存器保存到内存、再把另一个线程的读回来，要花几微秒；GPU 上所有驻留线程束的寄存器都在寄存器文件里，调度器每个周期都可以换一个线程束发射，不需要任何保存和恢复。

代价是寄存器要分着用：H100 一个 SM 驻留满 2048 个线程时，每个线程平均只有 32 个寄存器。内核用的寄存器越多，能驻留的线程就越少——§4.6 的占用率。

---

## 4. 程序在 GPU 上怎样执行 · Program execution on GPU

### 4.1 主机与设备

slides 给出 CUDA 的编程模型：

- **CPU 是主机（host）**，运行普通的 C++ 程序；
- **GPU 是设备（device）**，运行 CUDA **内核**（kernel）；
- 一个 CUDA 程序一部分在 CPU 上跑、一部分在 GPU 上跑；
- **数据要在系统内存和显存之间搬运**（经过 PCIe，slides 的图：PCIe 32 GB/s，系统内存 64–512 GB/s，A6000 显存 768 GB/s）。

主机和设备有各自的内存，主机代码不能直接读写显存（反过来也一样，除非用统一内存等机制）。在模拟器里，主机代码直接读写 `DeviceArray` 会报错，必须用 `dev.memcpy`——第 3 讲会把 `cudaMalloc` → `cudaMemcpy` → 发射内核 → `cudaMemcpy` → `cudaFree` 这五步走一遍，并算清拷贝的时间账。

### 4.2 内核：一份串行的程序，几万份同时运行

slides 对内核的描述只有四句，每一句都重要：

1. 每个内核是一个在 GPU 上运行的函数（程序）；
2. **程序本身是串行的**——内核代码里没有任何"并行"的写法；
3. 可以同时运行很多（上万个）线程；
4. **每个线程用自己的线程编号，去处理数据的不同部分**。

这就是 **SIMT**（Single Instruction, Multiple Threads，单指令多线程）：写一份串行代码，硬件把它在几万个线程上同时执行，每个线程看到的唯一区别是内置变量 `threadIdx`、`blockIdx` 的值。03 号脚本第 1 部分的向量加法：

```c
__global__ void vecAdd(const float* A, const float* B, float* C, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;   // 我是第几个线程？
    if (i < n) C[i] = A[i] + B[i];                   // 我负责第 i 个元素
}
vecAdd<<<(n + 255) / 256, 256>>>(dA, dB, dC, n);    // 发射：⌈n/256⌉ 个块，每块 256 个线程
```

n = 10、每块 4 个线程时，网格有 ⌈10/4⌉ = 3 个块、12 个线程。模拟器把每个线程算出的 i 记了下来：

| i | blockIdx.x | threadIdx.x | 算式 |
|---|---|---|---|
| 0–3 | 0 | 0–3 | 0 × 4 + 0..3 |
| 4–7 | 1 | 0–3 | 1 × 4 + 0..3 |
| 8–9 | 2 | 0–1 | 2 × 4 + 0..1 |

另外 2 个线程（块 2 的 threadIdx 2、3）算出的 i ≥ 10，被 `if (i < n)` 挡在外面。结果和 NumPy 的 `a + b` 逐位相同。

这个编号公式和 P1 §5 的行优先下标一模一样：把线程看成一个 `gridDim × blockDim` 的二维数组，i 就是它展平后的位置。

### 4.3 网格、线程块、线程

为什么要分两级（块和线程），而不是直接发射 n 个线程？slides 的图给出了答案：**线程块被整体分配到某一个 SM 上执行**，块内的线程可以通过共享内存协作、用 `__syncthreads()` 同步；不同块之间没有任何执行顺序的保证，也不能互相等待。

```
   网格 Grid（一次内核发射）
   ┌─────────────┬─────────────┬─────────────┬───
   │ 线程块 0     │ 线程块 1     │ 线程块 2     │ …   ← 每个块整体分到一个 SM
   │ 256 个线程   │ 256 个线程   │ 256 个线程   │
   └─────────────┴─────────────┴─────────────┴───
           │ 块内再按 32 个一组切成线程束（warp）
           ▼
   线程束 0（线程 0–31）、线程束 1（32–63）、……、线程束 7（224–255）
```

这个两级结构是 CUDA 设计里最聪明的一点：

- **块内**：线程在同一个 SM 上，可以共享一块高速的共享内存、可以同步——这给了"协作"的能力（第 4 讲的分块矩阵乘法靠它）。
- **块间**：完全独立——这给了"扩展"的能力。同一个网格，在 84 个 SM 的 A6000 上和 132 个 SM 的 H100 上都能跑，SM 多就同时跑更多块（§4.6 的波次）。程序不需要知道 GPU 有多少个 SM。

网格和块都可以是一到三维的（`dim3`），第 3 讲处理矩阵时会用到二维的网格和块。

### 4.4 线程束：真正的执行单位

slides 接着往下拆：SM 把每个线程块切成**线程束**（warp），**线程束是 GPU 创建、管理、调度、执行线程的基本单位**。每个线程束 32 个线程：

- 从同一个程序地址开始；
- 各有自己的程序计数器和寄存器；
- **每个周期执行一条共同的指令**；
- 可以分支、独立执行（代价见 §4.5）。

线程束是按块内的线性编号切的：线程 0–31 是第 0 个线程束，32–63 是第 1 个……每个分区的调度器管理一部分线程束，每个周期选一个"准备好了"（下一条指令的操作数都已就绪）的线程束发射。slides 的说法是：每个 SM 同一时刻可以并行执行 4 个线程束（每个分区一个）。

**块大小要取 32 的倍数。** 03 号脚本第 2 部分用 n = 1000 个 float 试了不同的块大小：

| 块大小 | 块数 | 发射的线程 | 线程束 | 线程束里的空位 | 读请求 | 读事务 |
|---|---|---|---|---|---|---|
| 32 | 32 | 1024 | 32 | 0 | 64 | 250 |
| 100 | 10 | 1000 | 40 | 280 | 80 | 290 |
| 128 | 8 | 1024 | 32 | 0 | 64 | 250 |
| 256 | 4 | 1024 | 32 | 0 | 64 | 250 |

（"读请求"是一个线程束的一条读指令；"读事务"是它产生的 32 字节内存事务个数。）

块大小 100 时，每个块被切成 4 个线程束，最后一个只有 4 个线程，其余 28 个位置空着——它们照样占用调度机会和寄存器。读事务也多了 16%：一是这些残缺的线程束要单独发请求（读请求从 64 个变成 80 个）；二是奇数号块从第 400、1200……字节开始（400 = 12 × 32 + 16），没有对齐到 32 字节边界，这些块里一个满线程束读的 128 字节要跨 5 个 sector 而不是 4 个。

块大小是 32 的倍数时，读事务正好是 1000 × 4 字节 × 2 个数组 ÷ 32 字节 = 250 个：线程束里 32 个线程读相邻的 4 字节，合起来正好是 4 个完整的 32 字节 sector，一个字节都不浪费。这就是**合并访存**（coalesced access）——第 4 讲的重点，这里先记住它的形状："**一个线程束在同一条指令里，一起访问一段连续的内存**"。

### 4.5 线程束分化

"每个周期执行一条共同的指令"和"可以独立分支"怎么同时成立？答案是：**线程束里的线程走了不同的分支时，线程束把两条路依次都走一遍**；走 if 那条路时，走 else 的线程被屏蔽（masked），原地空等，反之亦然。这叫**线程束分化**（warp divergence）。

05 号脚本第 1 部分用模拟器的指令计数（每执行一行 Python 代码算一条指令；同一线程束里执行同一行的线程合起来算一次"发射"）量了一下。一个块 128 个线程，一半走 if、一半走 else，两条路各循环 20 次：

| 分法 | 线程执行的指令 | 线程束发射次数 | SIMT 效率 |
|---|---|---|---|
| 所有线程都走 if（无分化） | 6144 | 192 | 100% |
| 按线程束分：线程束 0、2 走 if，1、3 走 else | 6144 | 192 | 100% |
| 按奇偶分：每个线程束里一半一半 | 6016 | 352 | 53.4% |

（SIMT 效率 = 线程实际执行的指令数 ÷ (发射次数 × 32)。）

同样是"一半线程走 if、一半走 else"，**按线程束对齐**时没有任何代价——分化只发生在线程束之间，而线程束之间本来就是独立调度的。**按奇偶分**时，每个线程束都要把两条路走一遍，发射次数变成 1.83 倍，效率 53%（不到正好 50%，是因为分支前后的公共代码大家一起执行）。线程做的事几乎一样多，时间却几乎翻倍。

**工作量不均是同一个问题。** 05 号脚本第 2 部分让 128 个线程各循环 0–31 次中的某个次数：

| 分配方式 | 总循环次数 | Σ 每个线程束里的最大次数 | 发射次数 | SIMT 效率 |
|---|---|---|---|---|
| 随机排列 | 2030 | 122 | 260 | 55.0% |
| 按工作量排序后再分配 | 2030 | 79 | 174 | 82.1% |
| 每个线程都循环平均次数 | 2048 | 64 | 144 | 100% |

线程束的循环要转到它里面次数最多的那个线程做完为止。把工作量相近的线程放进同一个线程束，总工作量不变，效率从 55% 升到 82%。LLM 系统里"长短不一"到处都是：一个批里的序列长度不同、稀疏矩阵每一行的非零元个数不同（第 4 讲的 CSR）。按长度排序或分桶、把长任务拆碎，是 GPU 上常见的对策；第 22 讲的连续批处理是同一个问题在服务系统层面的版本。

> **关于"独立的程序计数器"。** slides 说每个线程"有自己的程序计数器"——这是 NVIDIA 从 Volta 架构（2017）开始的"独立线程调度"：分化的两条路可以交错推进，不再必须先走完一条。但同一时刻一个线程束仍然只执行一条指令，分化的代价并没有消失；它主要改变的是正确性语义（例如同一线程束内的线程也可能互相等待，需要 `__syncwarp()`）。

### 4.6 线程块怎样分到 SM 上：占用率与波次

slides 说"一个 SM 上能放多少个线程束，取决于它们申请的资源和 SM 剩余的资源"。具体地，一个 SM 能同时驻留的块数受四个上限约束，取最小者：

1. 每个 SM 最多驻留的线程束数（A100 / H100：64 个，即 2048 个线程；A6000：48 个）；
2. 每个 SM 最多驻留的块数（A100 / H100：32；A6000：16）；
3. 寄存器：每个线程用的寄存器 × 线程数 不能超过 64K；
4. 共享内存：每块申请的共享内存 × 块数 不能超过 SM 的共享内存。

驻留的线程束数 ÷ 最大线程束数，叫**占用率**（occupancy）。04 号脚本第 1 部分（[`gpu_model.occupancy`](../code/lecture02/gpu_model.py) 是简化版的占用率计算器）：

| 每块线程数 | RTX A6000（最多 48 束、16 块） | H100（最多 64 束、32 块） |
|---|---|---|
| 32 | 16 块 / 16 束 = 33%（块数上限） | 32 块 / 32 束 = 50%（块数上限） |
| 64 | 16 块 / 32 束 = 67%（块数上限） | 32 块 / 64 束 = 100% |
| 128–512 | 100% | 100% |
| 768 | 2 块 / 48 束 = 100% | 2 块 / 48 束 = 75% |
| 1024 | 1 块 / 32 束 = 67% | 2 块 / 64 束 = 100% |

块太小会被"每个 SM 最多几个块"卡住；块太大可能放不下整数个（A6000 每个 SM 最多 1536 个线程，1024 个线程的块只放得下 1 个）。128–512 是常见的选择。

资源用量的影响更隐蔽（H100，每块 256 个线程）：每线程寄存器从 32 涨到 40，驻留块数就从 8 掉到 6（占用率 75%）；64 个寄存器时 50%。每块申请 48 KB 共享内存时只能驻留 4 块（50%），160 KB 时只有 1 块（12%）。写内核时多开几个局部变量、多申请一块共享内存，都可能让占用率掉一截。

占用率不是越高越好——第 4 讲的分块矩阵乘法就是故意用更多共享内存、更低的占用率换更少的显存访问。但占用率太低时，没有足够的线程束来掩盖访存延迟（§4.7）。

**slides 的问题**：`fullKernel<<<dim3(2,3,4), dim3(2,4,8)>>>` 一共 24 个块 × 64 个线程 = 1536 个线程，能在 A6000 上同时运行吗？04 号脚本第 2 部分：每个 SM 最多驻留 16 个这样的块，全卡 84 × 16 = 1344 个，24 个块当然能同时驻留——实际上只用到 24 个 SM（每个 1 块），另外 60 个 SM 空着。但"能同时驻留"不等于"把 GPU 用满了"：每个 SM 上只有 2 个线程束，4 个调度器有一半没事干；全卡能驻留 129,024 个线程，这个网格只用了 1.2%。**GPU 要喂饱，网格里的线程数要比全卡能驻留的多好几倍。**

**波次与尾部效应。** 块数超过全卡能同时驻留的数目时，块会一波一波地执行：前一波的块做完，空出位置，下一波的块补上。04 号脚本第 3 部分假设 H100 每个 SM 同时驻留 1 个块、每个块耗时相同：

| 块数 | 波数 | SM 的平均利用率 |
|---|---|---|
| 132 | 1 | 100% |
| 133 | 2 | 50.4% |
| 264 | 2 | 100% |
| 265 | 3 | 66.9% |
| 1000 | 8 | 94.7% |
| 10000 | 76 | 99.7% |

133 个块要 2 波，第二波只有 1 个块在算，其余 131 个 SM 干等——利用率和只用一半 SM 差不多。块数很多时尾巴被摊薄。一个矩阵乘法切成多少块，会因为这种"波次量化"（wave quantization）影响速度，cuBLAS 会针对矩阵形状挑选分块方式（第 4 讲）。

### 4.7 为什么要这么多线程：延迟隐藏

slides 第 4 讲会给出一组数字：读寄存器约 1 个周期、读共享内存约 5 个周期、**读全局内存（显存）约 500 个周期**。一个线程发出读请求之后，要等 500 个周期数据才回来——这段时间它什么也做不了。CPU 用大缓存和乱序执行来对付这个问题；GPU 的办法是：**让别的线程束接着算**。

04 号脚本第 4 部分做了一个线程束调度器的逐周期模拟（[`gpu_model.simulate_scheduler`](../code/lecture02/gpu_model.py)）：每个线程束反复做"读 128 字节 → c 条依赖这次数据的计算指令"，访存延迟 500 个周期；调度器每个周期最多发射一条指令。另外加上带宽的限制：H100 的 3.35 TB/s 分到每个 SM 的每个调度器，约 3.2 字节 / 周期，即每 40 个周期才能读一次 128 字节。结果（调度器发射了指令的周期占比；括号里是公式 $\min(1,\ W(1+c)/(L+c),\ (1+c)/40)$ 的值）：

| 每次读之后的计算 c | 1 个线程束 | 2 | 4 | 8 | 12 | 16 |
|---|---|---|---|---|---|---|
| 4 | 1%（1%） | 2%（2%） | 4%（4%） | 8%（8%） | 12%（12%） | 12%（12%） |
| 16 | 3%（3%） | 7%（7%） | 13%（13%） | 26%（26%） | 39%（40%） | 42%（42%） |
| 64 | 11%（12%） | 22%（23%） | 45%（46%） | 85%（92%） | 99%（100%） | 99%（100%） |

- 只有 1 个线程束时，调度器绝大部分时间在等数据。**线程束越多，一个在等的时候另一个就能发射**，利用率几乎随线程束数线性上升——这就是 slides 说的"用大量线程隐藏访存延迟"。
- 计算少的内核（c = 4）在 12–13 个线程束时撞上带宽上限（12%），再多线程束也没用：它是**访存受限**的，瓶颈是显存带宽本身。计算多的内核（c = 64）在 16 个线程束时接近 100%：**计算受限**。

**用 Little 定律算"需要多少线程"**（P4 §2）：在路上的请求数 = 发射速率 × 延迟。模拟在 c = 4、16 个线程束时测得平均 12.4 个读请求在路上，发射速率 0.025 个 / 周期，× 500 = 12.5，吻合。要用满带宽，每个调度器需要同时有 500 ÷ 40 ≈ 12.5 个读请求在路上，整个 SM 约 50 个，全卡约 6,600 个——约 0.8 MiB 的数据时刻"在路上"。H100 每个 SM 最多驻留 64 个线程束（每个调度器 16 个），和需要的 12.5 个是同一个量级。**GPU 是按"带宽 × 延迟"来配线程数的。** 这也是为什么占用率太低时，访存受限的内核连带宽都吃不满。

（这个模型很简化：真实的线程束可以连发几条互不依赖的读指令，即指令级并行，也能帮助隐藏延迟；缓存命中时延迟更短。它用来看趋势，不代表精确数字。）

### 4.8 三个常见错误，以及"一个线程处理多个元素"

03 号脚本第 3 部分在模拟器上犯了三个最常见的错误：

| 错误 | 模拟器的反应 | 真 GPU 上 |
|---|---|---|
| 忘了 `if (i < n)` | `IllegalAddress`：blockIdx=(3,0,0) threadIdx=(232,0,0) 越界访问 A[1000] | 不一定报错：可能悄悄读写到别的数组；越过已分配的区域时才报 "an illegal memory access" |
| 块数写成 `n / 256`（向下取整） | 不报错；只发射了 768 个线程，最后 232 个元素还是垃圾值（NaN） | 同样不报错——最危险 |
| 一个块 2048 个线程 | `LaunchError`：block.x = 2048 超过上限 1024 | 内核根本不运行；要主动调用 `cudaGetLastError()` 才看得到 |

教训：块数永远向上取整（`(n + 255) / 256`）；测试时一定要用**不是块大小整数倍**的 n；发射之后检查错误。

**元素比线程多怎么办？** 一个线程处理多个元素有两种分法。03 号脚本第 4 部分：n = 4096，只发射 4 × 64 = 256 个线程，每个线程处理 16 个元素：

| 分法 | 线程读次数 | 读请求 | 读事务 | 每请求事务 | 不同的 sector |
|---|---|---|---|---|---|
| 跨步循环：`for (i = g; i < n; i += 总线程数)` | 8192 | 256 | 1024 | 4.0 | 1536 |
| 每人一段：线程 g 处理 `[16g, 16g+16)` | 8192 | 256 | 8192 | 32.0 | 1536 |

两种分法读的元素完全一样，用到的不同 sector 也一样多，但每人一段的读事务是跨步循环的 **8 倍**。原因在"同一时刻"：线程束里的 32 个线程是一起执行同一条指令的。跨步循环时它们读 32 个相邻元素（4 个 sector）；每人一段时相邻线程隔 16 个元素（64 字节），32 个线程落在 32 个不同的 sector 里。

**CPU 上正好相反**：一个 CPU 线程顺着自己那一段往下读，缓存行和预取器都很高兴（P4 §1），所以 CPU 的并行库通常按段切分。在 GPU 上思考访存时，要想的是**"一个线程束在同一条指令里一起访问了哪些地址"**，而不是"一个线程依次访问了哪些地址"。这是从 CPU 编程转到 GPU 编程最需要扭过来的直觉。

### 4.9 编译 CUDA 代码

slides 的最后一页技术内容是一条编译命令，原课 HW1 每次改完 `combine.cu` 都要运行它：

```bash
nvcc -o output.so --shared src.cu -Xcompiler -fPIC
```

- `nvcc` 是 NVIDIA 的 CUDA 编译器驱动：它把 `.cu` 文件拆成两部分，**设备代码**（`__global__` 函数）编译成 GPU 的指令（中间表示 PTX，再到具体架构的机器码 SASS），**主机代码**交给系统的 C++ 编译器（gcc / clang）。
- `--shared -o output.so`：生成一个共享库（P1 §1.1 的"链接"那一步的产物），而不是可执行文件。
- `-Xcompiler -fPIC`：把 `-fPIC`（生成位置无关代码，共享库必须这样编译）转交给主机编译器。

生成的 `.so` 再由 Python 用 ctypes 加载、调用（P1 §6.2 用同样的方法调用过自己写的 C 函数）。miniTorch 的 CUDA 后端就是这样把 `combine.so` 里的内核接到 Python 的张量运算上的——第 3 讲会把这条链路完整地走一遍。

---

## 5. 小结 · Summary

1. **神经网络 = 五种底层运算**：gather、matmul、map、zip、reduce。矩阵乘法包揽了几乎全部 FLOPs（真实尺寸下 99.87%）而且是计算受限的；逐元素运算几乎不做运算，却读写了一半以上的字节，是访存受限的。**算术强度**决定瓶颈，不是运算的名字。
2. **GPU 服务器的带宽分层**：显存（几百 GB/s 到几 TB/s）≫ NVLink（几十到几百 GB/s）> PCIe（32 GB/s）> NVMe（7 GB/s）。训练时 CPU → GPU 的数据很少，**GPU 之间的梯度**才是大流量，拓扑决定它有多快。
3. **读规格表先看脚注**：张量核心的峰值默认含 2:4 稀疏，dense 是一半。平衡点 = 峰值算力 ÷ 带宽，一代比一代高。
4. **GPU 由几十到上百个 SM 组成**；每个 SM 有 4 个分区，每个分区一个线程束调度器、32（或 16）个 FP32 车道、一个张量核心、64 KB 寄存器；SM 内共享一块 L1 / 共享内存，全卡共享 L2 和显存。峰值 = SM 数 × 车道数 × 2 × 频率。
5. **CPU 延迟优先，GPU 吞吐优先**：GPU 用海量常驻线程（每个都有自己的寄存器）换吞吐，线程束切换不花时间。
6. **执行模型**：主机发射内核 → 网格由线程块组成 → 块整体分到一个 SM → 块切成 32 个线程的**线程束** → 调度器每周期给一个线程束发射一条指令（SIMT）。块内可以协作和同步，块间完全独立，所以同一个程序能在不同大小的 GPU 上扩展。
7. **三个效率问题**：线程束**分化**（同一束走不同分支 → 依次执行）、**非合并访存**（同一条指令访问分散的地址 → 事务变多）、**占用率 / 波次**（资源用量和块数决定 SM 有没有被喂饱）。
8. **延迟隐藏**：一次读显存约 500 个周期，靠别的线程束填补；按 Little 定律，用满带宽需要"带宽 × 延迟"那么多数据在路上，这决定了 GPU 要驻留多少线程。

## 6. 常见误区 · Common pitfalls

- **"CUDA 核心就是 GPU 的核，A6000 有一万个核。"** CUDA 核心只是向量流水线的车道，不能独立取指令；能独立调度的是每个 SM 的 4 个线程束调度器。
- **"GPU 比 CPU 快 67 倍。"** 那是拿 GPU 的峰值和 CPU 的标量峰值比。同口径（都用 SIMD）A6000 的 FP32 只快约 4 倍；真正的数量级差距来自张量核心和更高的显存带宽，而且只对数据并行的程序成立。
- **"规格表上 H100 是 1979 TFLOPS。"** 那是含 2:4 稀疏的数字，普通矩阵乘法的上限是 989。
- **"线程越多越好 / 占用率越高越好。"** 线程足够隐藏延迟之后，再多也没用；有时故意降低占用率换取更多共享内存和寄存器反而更快（第 4 讲）。
- **"一个线程顺序处理一段连续数据最快。"** 这是 CPU 的直觉。在 GPU 上要让同一个线程束里相邻的线程访问相邻的地址（跨步循环），否则事务数会成倍增加。
- **"块数写成 n / 256 就行。"** 整数除法向下取整，最后不满一块的元素没人算，而且不报错。永远写 `(n + 255) / 256`，并保留 `if (i < n)`。
- **"if / else 在 GPU 上和 CPU 上一样便宜。"** 同一个线程束里走不同分支时，两条路都要走一遍。按线程束对齐的分支则没有代价。
- **"训练慢是因为 CPU 往 GPU 拷数据慢。"** 对 LLM 训练，每步送进 GPU 的只是 token 编号（几十 KB）；大流量是 GPU 之间的梯度同步。
- **"在模拟器上结果对，内核就没有竞争。"** 模拟器按固定顺序执行线程，真 GPU 上顺序不确定。这就是模拟器另外做竞争检测的原因（第 4 讲会看到漏写 `__syncthreads()` 的例子）。

## 7. 练习 · Exercises

1. 一块假想的 GPU 有 100 个 SM，每个 SM 有 128 个 FP32 核心，boost 频率 1.5 GHz。它的 FP32 峰值是多少？
2. 某 GPU 的规格表写着"BF16 Tensor Core：1979 TFLOPS\*（\*with sparsity）"。一个 4096 × 4096 × 4096 的稠密矩阵乘法，按 60% 的 dense 峰值要多久？
3. n = 1,000,000，每块 256 个线程，需要几个块？一共发射多少个线程？多少个线程什么也不做？
4. 每块 256 个线程时，处理第 1000 个元素（下标从 0 开始）的是哪个线程？把答案和 03 号脚本第 3 部分错误 1 的报错信息对照。
5. A100 上每块 256 个线程，内核每个线程用 48 个寄存器，不用共享内存。每个 SM 能驻留几个块？占用率是多少？瓶颈是什么？
6. H100（132 个 SM），每个 SM 能同时驻留 2 个块，网格有 300 个块，每个块耗时相同。要几波？SM 的平均利用率是多少？
7. 一个线程束的 32 个线程分别读 `A[32 * lane]`（float，lane = 0..31）。这条指令产生几个 32 字节事务？搬运的字节里有多少是真正用到的？
8. 块大小 128，比较两种写法的分化：(a) `if (threadIdx.x < 16) f(); else g();` (b) `if (threadIdx.x < 64) f(); else g();`。哪一种有分化？为什么？
9. A100 的显存带宽 2.039 TB/s，频率 1.41 GHz，假设访存延迟 500 个周期。按 Little 定律，全卡需要多少字节"在路上"才能用满带宽？平均每个 SM 是多少个 128 字节的线程束读请求？和每个 SM 最多 64 个线程束相比如何？
10. 在 slides 那台服务器上，把 10 GB 的梯度从 GPU0 发到 GPU1（有 NVLink 桥，单向 56.25 GB/s）和发到 GPU2（只能走 PCIe，32 GB/s），各要多久？
11. 向量加法 `C[i] = A[i] + B[i]`（float32）的算术强度是多少？在 H100 上它最多能达到多少 GFLOPS？占 FP32 峰值的百分之几？
12. 判断下面的改写是否保持结果不变，并说明理由：(a) Embedding → Avg → Linear 改成 Embedding → Linear → Avg；(b) Linear → GELU → Avg → Linear 改成 Linear → Avg → GELU → Linear；(c) Linear → Avg → Softmax 改成 Linear → Softmax → Avg。

### 参考答案

1. 100 × 128 × 2 × 1.5×10⁹ = 3.84×10¹³，即 **38.4 TFLOPS**。
2. dense 峰值 = 1979 / 2 ≈ 989.5 TFLOPS。FLOPs = 2 × 4096³ ≈ 1.374×10¹¹。时间 = 1.374×10¹¹ / (989.4×10¹² × 0.6) ≈ **0.23 毫秒**。如果误用 1979，会低估一半。
3. ⌈1,000,000 / 256⌉ = **3907** 个块，发射 3907 × 256 = **1,000,192** 个线程，其中 **192** 个线程的 i ≥ n，被越界检查挡住。
4. 1000 = 3 × 256 + 232：**blockIdx.x = 3，threadIdx.x = 232**——正是 03 号脚本错误 1 报告的那个越界线程（n = 1000 时它是第一个 i ≥ n 的线程）。
5. 每个线程束占 ⌈48 × 32 / 256⌉ × 256 = 1536 个寄存器，每块 8 个线程束 = 12,288 个；65,536 / 12,288 = 5.33 → **5 个块** = 40 个线程束，占用率 40 / 64 = **62.5%**，瓶颈是**寄存器**（线程束上限允许 8 块，块数上限允许 32 块）。
6. 位置数 = 132 × 2 = 264，⌈300 / 264⌉ = **2 波**，利用率 = 300 / (2 × 264) = **56.8%**。
7. 相邻线程相隔 32 个 float = 128 字节，每个线程落在不同的 sector：**32 个事务**，搬运 32 × 32 = 1024 字节，真正用到 32 × 4 = 128 字节，**利用率 12.5%**。
8. (a) **有分化**：每个线程束（0–31、32–63……）里，只有线程束 0 的前 16 个线程走 f，后 16 个走 g；线程束 0 两条路都要走（其余线程束全走 g，没有分化）。(b) **没有分化**：64 是 32 的倍数，线程束 0、1 全走 f，线程束 2、3 全走 g。
9. 延迟 = 500 / 1.41 GHz ≈ 354.6 ns。在路上的字节 = 2.039×10¹² × 354.6×10⁻⁹ ≈ **723 KB**；每个 SM ≈ 723 KB / 108 ≈ 6.7 KB ≈ **52 个**128 字节的读请求。每个 SM 最多驻留 64 个线程束，够用但余量不大——如果每个线程束同时只有一个读请求在路上，占用率要达到约 81%（52 / 64）才能用满带宽（或者靠指令级并行让每个线程束同时发多个读请求）。
10. NVLink：10 / 56.25 ≈ **178 ms**；PCIe：10 / 32 ≈ **313 ms**。
11. 每个元素 1 次加法，读 8 字节、写 4 字节，算术强度 1/12 ≈ **0.083 FLOP/字节**。上限 = 0.083 × 3.35×10¹² ≈ **279 GFLOPS**，是 FP32 峰值 67 TFLOPS 的 **0.42%**。向量加法在任何 GPU 上都是访存受限的。
12. (a) **保持**：Avg 和 Linear 都是线性 / 仿射的，可以交换（第 1 部分的同一个道理），后者的 Linear 计算量多 T 倍。(b) **不保持**：GELU 是非线性的，平均之后再 GELU ≠ GELU 之后再平均。(c) **不保持**：softmax 是非线性的，"平均的 softmax"≠"softmax 的平均"。

## 8. 延伸阅读 · Further reading

- **PMPP 第 4 版**第 2 章（CUDA 程序结构、向量加法）、第 4 章（GPU 架构、块的调度、线程束与 SIMD 硬件、控制分化、延迟隐藏、资源划分与占用率）。本讲 §4 的大部分概念都能在第 4 章找到更系统的推导。
- **NVIDIA CUDA C++ Programming Guide**：第 2 章 "Programming Model"（线程层次、内存层次、异构编程）、第 4 章 "Hardware Implementation"（SIMT 架构、硬件多线程）。附录 "Compute Capabilities" 有每一代 GPU 的全部上限（本讲占用率计算用到的数字都在那里）。
- **NVIDIA A100 / H100 架构白皮书**：SM 的结构图、张量核心、2:4 结构化稀疏、TMA、FP8 Transformer Engine 的原始说明。
- **Volkov (2010) *Better Performance at Lower Occupancy***（GTC 演讲）：用指令级并行代替线程级并行、低占用率也能跑满的经典论证，是 §4.6"占用率不是越高越好"的出处。
- **Lindholm et al. (2008) *NVIDIA Tesla: A Unified Graphics and Computing Architecture***（IEEE Micro）：SIMT 和线程束的最初描述。

## 9. 下一讲预告 · Next lecture

slides 的下一页写着：**写 GPU 程序**。第 3 讲把本讲的执行模型变成真正的代码：一个 CUDA 程序的完整五步（`cudaMalloc`、`cudaMemcpy`、发射内核、拷回、`cudaFree`）以及它们的时间账——为什么 PCIe 拷贝可能比内核本身慢一百倍；二维、三维的网格和线程编号，以及 slides 里那个矩阵加法内核为什么"对但慢"；最后在模拟器上实现原课 HW1 的四个内核（map、zip、reduce、矩阵乘法），处理任意步长和广播，并和 NumPy 逐位对比。
# 第 3 讲 · GPU 编程基础（二）：CUDA 操作与矩阵 / 张量计算 · GPU Programming Basics 2

> **对应 Maps to**：CMU 11-868 / 11-968 *Large Language Model Systems*（Spring 2026，Lei Li），Lecture 3 · Wed 1/21 · **GPU Programming Basics 2**。
> **阅读材料**：PMPP 第 4 版第 2 章（CUDA 程序结构）、第 3 章（多维网格与数据）；原课 HW1 说明（CUDA Programming：map、zip、reduce、矩阵乘法内核及其接入）；Mark Harris, *Optimizing Parallel Reduction in CUDA*（NVIDIA 教程）。
> **依据说明**：本讲以原课第 3 讲公开的 slides（llmsys-03-gpu-programming2.pdf，32 页）为主线：CUDA 的基本操作（`cudaMalloc`、`cudaMemcpy`、发射内核、`cudaFree`）、GPU 的三种内存、H100 的 TMA 与异步执行、`__global__` / `__device__` / `__host__`、`dim3` 与内置变量、VecAdd / MatAdd / fullKernel 三个例子。§3 的四类内核按原课 HW1 的说明组织（slides 只给了 Colab 笔记本的链接）。**模拟器上的全部实验、时间账、§4 的并行归约都是本讲义自己设计的，不代表课堂上讲了什么**。
> **配套代码 Code**：[`code/lecture03/`](../code/lecture03/)（4 个脚本 + `kernels.py`；复用第 2 讲的 `simt.py`、`gpu_model.py` 和 P1 的 `tensor_layout.py`；每个脚本 4 秒以内，输出重复运行逐位相同）
> **预计阅读时间**：4 小时（含运行代码）

---

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

### 0.1 学习目标 Learning goals

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

1. 写出一个完整 CUDA 程序的**五步**（分配显存、拷入、发射内核、拷出、释放），解释每个 API 的参数（特别是 `cudaMalloc(&dA, size)` 为什么要传 `&dA`），并能算出**拷贝 vs 内核**的时间账，判断一个运算值不值得搬到 GPU 上。
2. 说出 GPU 的三种内存（寄存器、共享内存、全局内存）各自的**作用域和生命周期**，以及 `__global__`、`__device__`、`__host__` 的含义。
3. 用 `dim3` 配置一维 / 二维 / 三维的网格和线程块，写出任意维度下的**全局线程编号**，处理**越界**，并让线程束的访问**合并**（threadIdx.x 对应内存里连续的那一维）。
4. 用"**存储 + 形状 + 步长**"实现原课 HW1 的四类内核：**map、zip（带广播）、reduce、矩阵乘法**，并把它们接到主机端的张量上。
5. 理解**并行归约**：关键路径、树形归约的三种写法（分化与 bank 冲突的差别）、原子操作与数据竞争，以及为什么 GPU 上的浮点归约"每次可能不一样"。

### 0.2 先修与本讲范围 Prerequisites & scope

| 默认你已经会 Assumed | 本讲只给直觉，后面再展开 Previewed only |
|---|---|
| 第 2 讲：SIMT、线程束、合并访存的定义、SIMT 模拟器的用法（§0.4） | 共享内存分块、bank 冲突的系统讲解、转置（第 4 讲） |
| P1 §3.5（`void**`）、§5（行优先、步长、视图、HW1 的下标函数） | 异步拷贝、流、TMA 的编程（第 10–11、21 讲） |
| P2 §3（浮点加法不满足结合律） | 自动微分、框架的分层（第 5 讲） |

### 0.3 本讲地图 Map

```
§1  一个 CUDA 程序的五步：分配、拷入、发射、拷出、释放；三种内存；时间账；异步   ← 代码 01
§2  多维网格与线程编号：dim3、fullKernel、矩阵加法的 x / y、越界、块的形状       ← 代码 02
§3  张量计算：HW1 的 map、zip、reduce、矩阵乘法（步长与广播）；主机端封装      ← 代码 03
§4  并行归约：关键路径、树形归约的三种写法、原子操作与不确定性                   ← 代码 04
§5  小结   §6  常见误区   §7  练习   §8  延伸阅读   §9  下一讲预告
```

---

## 1. 一个 CUDA 程序的五步 · Basic CUDA operations

### 1.1 回顾：内核在哪里跑，数据在哪里放

第 2 讲讲了内核是怎样执行的：网格 → 线程块 → 线程束 → 线程，每个线程用 `blockIdx`、`threadIdx` 算出自己负责的数据。但还有一个问题没回答：**数据是怎么到 GPU 上的？** 第 2 讲 §4.1 说过，主机（CPU）和设备（GPU）有各自的内存，内核只能读写显存。slides 第 11 页的图再一次提醒：CPU 的系统内存和 GPU 的显存之间隔着一条 PCIe（32 GB/s），而 GPU 访问自己显存的带宽是它的几十倍。

所以每个 CUDA 程序都要明确地安排数据的搬运。slides 把它总结成五个操作：

1. CPU 在 GPU 上**分配显存**：`cudaMalloc`
2. CPU 把数据**拷到显存**（host to device）：`cudaMemcpy`
3. CPU **发射内核**
4. CPU 把结果**拷回主机**（device to host）：`cudaMemcpy`
5. **释放显存**：`cudaFree`

### 1.2 五个 API

**分配**：`cudaError_t cudaMalloc(void** devPtr, size_t size)`。

```c
int *dA;
cudaMalloc(&dA, n * sizeof(int));
```

为什么传 `&dA`？因为 `cudaMalloc` 要**修改 `dA` 本身**——把新分配的显存地址写进去；而它的返回值已经被用来返回错误码了。C 只有按值传参，要修改调用者的变量，就得传它的地址（P1 §3.4–§3.5）。`dA` 的类型是 `int*`，它的地址是 `int**`，统一转成 `void**` 传进去。slides 第 13 页画的正是这件事：CPU 内存里的变量 `dA`（地址 306）里存着一个显存地址（12）。

注意 `dA` 本身是**主机上的一个变量**，它的**值**是一个显存地址。主机代码不能解引用它（`*dA` 会读主机内存的地址 12，通常段错误）；只能把它传给内核，或传给 `cudaMemcpy`。

**释放**：`cudaFree(void* devPtr)`。slides 写着"不要忘了！"——显存不会因为函数返回而自动释放。

**拷贝**：`cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind)`。参数顺序是**先目的、后来源**（和 `memcpy` 一样），`count` 是**字节数**，`kind` 说明方向（`cudaMemcpyHostToDevice`、`cudaMemcpyDeviceToHost`）。

**发射**：`VecAddKernel<<<num_blocks, threads_per_block>>>(dA, dB, dC, n);`，其中 `num_blocks = (n + threads_per_block - 1) / threads_per_block`（向上取整，第 2 讲 §4.8）。

01 号脚本第 1 部分把 slides 第 29 页的 `VecAddCUDA` 一行行搬到模拟器上，每一步之后打印显存的状态：

| 步骤 | 显存里 | 已拷贝 |
|---|---|---|
| ① cudaMalloc × 3 | dA、dB、dC（各 1000 个 int32，内容是垃圾值） | 0 次 |
| ② cudaMemcpy(HostToDevice) × 2 | 同上，dA、dB 有了数据 | 2 次，8,000 字节 |
| ③ VecAddKernel<<<4, 256>>> | 同上，dC 有了结果 | 2 次 |
| ④ cudaMemcpy(DeviceToHost) | 同上 | 3 次，12,000 字节 |
| ⑤ cudaFree × 3 | （空） | 3 次 |

结果和 NumPy 的 `A + B` 逐位相同，没有泄漏。

### 1.3 GPU 的三种内存

slides 第 15 页列出了内核能用的三种内存：

| 内存 | 谁能看见 | 活多久 | 在哪里 | 速度 |
|---|---|---|---|---|
| **寄存器**（内核里的局部变量） | 一个线程私有 | 线程结束就没了 | SM 的寄存器文件 | 最快（约 1 个周期） |
| **共享内存**（`__shared__`） | 同一个线程块的所有线程 | 块结束就没了 | SM 上的 SRAM（和 L1 共用） | 很快（约 5 个周期） |
| **全局内存**（`cudaMalloc` 得到的） | 所有线程，以及主机（通过 cudaMemcpy） | **跨内核发射一直存在**，直到 cudaFree | 显存（HBM / GDDR） | 慢（约 500 个周期） |

最后一行的"跨内核发射一直存在"很关键：一个内核写进全局内存的结果，下一个内核可以直接读。这让"数据留在显存里，被一串内核接力处理"成为可能（§1.5）。共享内存的用法是第 4 讲的主题。

### 1.4 时间账：什么样的运算值得搬到 GPU 上

01 号脚本第 2 部分按公开规格算了一笔账（每次拷贝的固定开销取 10 μs、每次发射内核取 5 μs，都是量级）。向量加法：拷入两个数组、内核读 8 字节写 4 字节、拷出一个数组：

| GPU | 元素数 | 拷贝 | 内核 | 拷贝占比 |
|---|---|---|---|---|
| RTX A6000（PCIe 4.0，32 GB/s） | 1,000 | 0.030 ms | 0.0050 ms | 85.8% |
| RTX A6000 | 10,000,000 | 3.780 ms | 0.161 ms | 95.9% |
| RTX A6000 | 100,000,000 | 37.5 ms | 1.57 ms | 96.0% |
| H100（PCIe 5.0，64 GB/s） | 100,000,000 | 18.8 ms | 0.363 ms | 98.1% |

**拷贝占了九成以上**。原因很简单：内核和拷贝搬的是同样多的字节，而显存带宽是 PCIe 的 24 倍（A6000）到 52 倍（H100）。只用一次的数据，专门搬到 GPU 上做一次加法，不如在 CPU 上直接做。

矩阵乘法呢？拷贝 $3n^2$ 个 float，计算 $2n^3$ 次——拷贝的每个字节对应约 $n/6$ 次运算。在 A6000 上按 FP32 峰值算：

| n | 拷贝 | 计算 | 计算 ÷ 拷贝 |
|---|---|---|---|
| 1024 | 0.42 ms | 0.06 ms | 0.14 |
| 4096 | 6.3 ms | 3.6 ms | 0.56 |
| 16384 | 101 ms | 227 ms | 2.26 |

即使是算术强度很高的矩阵乘法，n = 4096 时计算也只有拷贝的一半多；到 n = 16384 才是两倍多。用张量核心（BF16 快约 4 倍），拷贝的占比还会更大。**PCIe 太慢了。**

结论不是"GPU 没用"，而是：**数据要留在显存里，被很多个内核反复使用**。训练时权重、梯度、优化器状态全部常驻显存，每步从主机送进来的只有一小批 token 编号（第 2 讲 §2.2：66 KB）——正是这个道理。

### 1.5 数据留在显存里

01 号脚本第 3 部分做了一个对照：连续做 8 个逐元素运算 $y = a_j x + b_j$。

| 写法 | 拷贝次数 | 结果与 NumPy 逐位相同 | n = 1 亿时 A6000 上的纸面时间 |
|---|---|---|---|
| 每个运算都"拷入 → 发射 → 拷出" | 16 | 是 | 209 ms |
| 拷入一次，8 个内核在显存里接力（交换输入输出指针） | 2 | 是 | 33 ms |

结果完全相同，时间差 6.2 倍。PyTorch 的设计就是第二种：张量一旦 `.to('cuda')`，之后所有运算的输入输出都在显存里；只有 `.item()`、`.cpu()`、`print(tensor)` 才会拷回主机——而且是**同步**的（见 §1.6），所以训练循环里随手写的 `loss.item()` 会让 CPU 停下来等 GPU。

再往前一步：8 个内核之间，中间结果还是要写回显存再读出来。如果把 8 个运算合成一个内核，每个元素只读一次、写一次——这就是**算子融合**（第 10–11 讲）。

### 1.6 异步执行与 TMA

上面的时间账假设每一步做完才开始下一步。真实的 CUDA 并不完全是这样：

- **内核发射对主机是异步的**：`kernel<<<...>>>(...)` 只是把内核放进一个队列（流，stream）就立即返回，CPU 可以接着做别的事（比如准备下一批数据）。
- **`cudaMemcpy` 是同步的**：它会等之前队列里的内核全部完成、拷贝结束才返回。异步版本 `cudaMemcpyAsync` 配合多个流，可以让"拷贝下一批"和"计算这一批"重叠（P4 §3 的重叠，第 15 讲会用到）。

slides 第 17–18 页介绍了 Hopper（H100）的新硬件：**张量内存加速器（Tensor Memory Accelerator，TMA）**。在 H100 之前，把一块数据从全局内存搬到共享内存，要由线程们一条条地发读指令、再一条条地写共享内存；TMA 让一个线程发一条指令（`cuda::memcpy_async(dst, src, shape, barrier)`），由专门的硬件把一整块多维数据搬过去，搬完时通知一个**屏障**（`barrier.wait()`）。线程们在等待期间可以继续算上一块——搬运和计算在硬件层面异步重叠。FlashAttention-3（第 21 讲）就重度依赖 TMA 和这种异步流水线。本讲只需要记住它的动机：**数据搬运太贵，要让专门的硬件去搬，算的单元专心算**。

### 1.7 显存管理的错误，以及 `__global__` / `__device__` / `__host__`

01 号脚本第 4 部分在模拟器上犯了五个错误：

| 错误 | 模拟器的反应 |
|---|---|
| 循环里 cudaMalloc，忘了 cudaFree（3 次 × 4 MiB） | 程序结束时还有 3 块、12 MiB 没释放 |
| 同一块显存 cudaFree 两次 | 报错：cudaFree 了一个无效的指针 |
| 拷贝方向写反（`cudaMemcpy(host, dev, …, HostToDevice)`） | 报错：方向要求 dst 是显存、src 是主机内存 |
| cudaFree 之后又在内核里用 | 报错：访问已经 cudaFree 的数组 |
| `__global__` 函数返回一个值 | 报错：内核不能有返回值 |

真 CUDA 里，这些错误大多**不会立刻报出来**：重复释放、释放后使用的后果往往是之后某个不相干的地方崩溃；方向写反时两个指针在 C 里都是 `float*`，编译器分辨不出。所以真实的 CUDA 代码会检查每一个 API 的返回值（`cudaError_t`），发射内核后调用 `cudaGetLastError()`。

slides 第 19 页的三个函数修饰符：

| 修饰符 | 谁调用 | 在哪里执行 | 备注 |
|---|---|---|---|
| `__global__` | 主机（CPU） | 设备（GPU） | 内核；返回 `void`；用 `<<<…>>>` 发射 |
| `__device__` | 设备 | 设备 | GPU 上的辅助函数，可以有返回值 |
| `__host__` | 主机 | 主机 | 普通的 CPU 函数（不写修饰符时的默认） |

主机代码和设备代码写在同一个 `.cu` 文件里，`nvcc` 按修饰符把它们分开编译（第 2 讲 §4.9）。在模拟器里，`__device__` 函数就是被内核调用的普通 Python 函数（脚本里的 `device_helper`）；HW1 的下标函数 `to_index`、`index_to_position`、`broadcast_index` 在 CUDA 版本里就是 `__device__` 函数。

---

## 2. 多维网格与线程编号 · Multi-dimensional grids

### 2.1 dim3 与内置变量

slides 第 21–22 页：主机在发射时用 `dim3` 指定网格和线程块的大小，每一维都可以是 1–3 维：

```c
dim3 Dg(4, 2, 1);           // 网格：4 × 2 × 1 = 8 个块
dim3 Db(8, 8, 1);           // 线程块：8 × 8 × 1 = 64 个线程（≤ 1024）
kernelFuncName<<<Dg, Db>>>(args);
```

内核里有四个内置变量，都有 `.x .y .z` 三个字段：`gridDim`（网格大小）、`blockIdx`（块在网格里的编号）、`blockDim`（块大小）、`threadIdx`（线程在块里的编号）。限制：每块最多 1024 个线程，`blockDim.x`、`.y` 最多 1024、`.z` 最多 64；网格的 x 方向最多 $2^{31}-1$，y、z 方向最多 65535。

### 2.2 fullKernel：三维编号的公式

slides 第 27–28 页的例子：`blocks_per_grid(2, 3, 4)`、`threads_per_block(2, 4, 8)`，一共 24 个块 × 64 个线程 = 1536 个线程。每个线程算出全局编号：

```c
int block_id = blockIdx.x + blockIdx.y * gridDim.x + blockIdx.z * gridDim.x * gridDim.y;
int block_offset = block_id * blockDim.x * blockDim.y * blockDim.z;
int thread_offset = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.x * blockDim.y;
int tid = block_offset + thread_offset;
```

02 号脚本第 1 部分让每个线程对 `COUNT[tid]` 做一次原子加，并打开了全局竞争检测：**1536 个 tid 每个恰好被写了一次，0 次竞争**——公式是一一对应的。它就是把 `(blockIdx.z, .y, .x, threadIdx.z, .y, .x)` 这六个数当作一个六维下标、按行优先展平（P1 §5）：x 变化最快，z 最慢。例如 blockIdx = (1, 2, 3)、threadIdx = (1, 3, 7) 的线程：block_id = 1 + 2·2 + 3·6 = 23，thread_offset = 1 + 3·2 + 7·8 = 63，tid = 23 × 64 + 63 = 1535，正是最后一个。

线程束按**块内编号**（thread_offset）每 32 个切一刀：块 (0,0,0) 的第 0 个线程束是 thread_offset 0–31，它们的 threadIdx 是 (0..1, 0..3, 0..3)——x 方向只有 2，所以一个线程束跨了 y 和 z 方向。**块的哪一维放在 x 上，决定了哪些线程在同一个线程束里一起访存**。

### 2.3 矩阵加法：x 对应行还是对应列？

slides 第 26 页的矩阵加法内核：

```c
__global__ void MatAddKernel(float* A, float* B, float* C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    C[i * N + j] = A[i * N + j] + B[i * N + j];
}
```

它让 x 方向的编号 i 当**行号**。02 号脚本第 2 部分把它和"行号来自 y、列号来自 x"的写法放在一起（N = 32，一个 32 × 32 的块）：

| 写法 | 结果正确 | 读事务 | 读事务 / 读请求 | 写事务 | 搬运的字节 |
|---|---|---|---|---|---|
| slides：i（行）= x 方向 | 是 | 2048 | 32.0 | 1024 | 98,304 |
| 改写：行 = y、列 = x | 是 | 256 | 4.0 | 128 | 12,288 |

两种写法结果都对，但 slides 的写法搬运的字节是 **8 倍**。原因：一个线程束是 threadIdx.x = 0..31、threadIdx.y 相同的 32 个线程。slides 的写法里 x 是行号，这 32 个线程访问 32 个不同的行，地址相隔 N × 4 = 128 字节，每个线程独占一个 sector；改写后它们访问同一行的 32 个相邻元素，4 个 sector 就够了。

**规则**：行优先存储的矩阵，让 **threadIdx.x**（线程束里变化最快的那一维）对应**列号**（内存里连续的那一维）。这条规则在 HW1 的矩阵乘法提示和第 4 讲里都会反复出现。

### 2.4 越界：N 不是块大小的整数倍

slides 第 30 页的 `MatAddCUDA` 用 `THREADS = 32`、`BLOCKS = (n + THREADS - 1) / THREADS`，网格是 BLOCKS × BLOCKS。但 `MatAddKernel` 没有越界检查——slides 的例子 N = 32 正好整除。02 号脚本第 3 部分取 N = 40：网格 2 × 2 个块、4096 个线程，矩阵只有 1600 个元素。slides 的内核立刻越界（blockIdx=(1,0,0)、threadIdx=(8,0,0) 的线程访问 A[1600]）；加上 `if (row < N && col < N)` 之后结果正确，4096 个线程里有 2496 个（61%）什么也不做。

二维时要对行、列**分别**检查。多出来的线程是"向上取整"的代价，N 远大于块大小时可以忽略；矩阵很小或形状很"瘦"时（例如 1 × 4096 的矩阵用 32 × 32 的块）浪费会很大，这时应该换用一维的块，或者换一种分工。

### 2.5 二维线程块的形状

同样每块 256 个线程，块的形状不同，合并程度也不同。02 号脚本第 4 部分（64 × 64 的矩阵加法，行 = y、列 = x）：

| blockDim.x × blockDim.y | 读事务 | 读事务 / 读请求 |
|---|---|---|
| 32 × 8、64 × 4 | 1024 | 4.0 |
| 16 × 16、8 × 32 | 1024 | 4.0 |
| 4 × 64 | 2048 | 8.0 |
| 1 × 256 | 8192 | 32.0 |

在"按 32 字节 sector 计数"的模型里，x 方向至少 8 个 float（32 字节）时每个 sector 都被用满：blockDim.x = 16 时一个线程束跨两行、每行读 64 字节（2 个 sector），总数和 32 × 8 一样。x 方向只有 4 时每行只读 16 字节（半个 sector），事务翻倍；只有 1 时每个线程独占一个 sector。实践中 x 方向通常取 16 或 32（16 × 16、32 × 8、32 × 32）：x = 32 时一个线程束只碰一行、地址完全连续，对 128 字节的缓存行和显存的突发传输也最友好。另外，`dim3(64, 32)` = 2048 个线程会直接发射失败。

---

## 3. 张量计算：原课 HW1 的四类内核 · Tensor computation on GPU

### 3.1 张量 = 存储 + 形状 + 步长

第 2 讲 §1 说过，神经网络最后只剩几种底层运算。原课 HW1 要求在 `combine.cu` 里写其中四种的 CUDA 内核，并接到 miniTorch 的张量上：

| 内核 | 做什么 | 例子 |
|---|---|---|
| `mapKernel` | 对每个元素做一元函数 | $f(x) = x^2$：[1, 2, 3] → [1, 4, 9] |
| `zipKernel` | 对两个张量对应位置做二元函数（可以广播） | 加法：[1, 2, 3] + [4, 5, 6] → [5, 7, 9] |
| `reduceKernel` | 沿某一维归约 | 沿第 1 维求和：[[1,2,3],[4,5,6]] → [6, 15] |
| `MatrixMultiplyKernel` | （批量）矩阵乘法 | — |

难点不在运算本身，而在**下标**：miniTorch 的张量和 NumPy、PyTorch 一样，是"一维存储 + 形状 + 步长"（HW1 说明的 "Data layout and strides" 一节，P1 §5）：

$$
\text{位置} = \sum_k \text{index}[k] \times \text{strides}[k]
$$

转置、切片、广播都只改形状和步长，不搬数据（视图）。所以内核拿到的输入可能是任意步长的，**同一个内核必须对连续的张量、转置之后的视图、广播之后的张量都算对**。

[`kernels.py`](../code/lecture03/kernels.py) 用模拟器把这四个内核写了一遍（对应 HW1 的 `combine.cu`），并写了主机端的封装（对应 HW1 的 `cuda_kernel_ops.py`）。下标换算直接用 P1 的 [`tensor_layout.py`](../code/prereq1/tensor_layout.py)：`to_index`（第 i 个元素 → 多维下标）、`index_to_position`（多维下标 → 一维位置）、`broadcast_index`（输出的下标 → 被广播的输入的下标）。

### 3.2 map：每个线程一个输出元素

```python
def map_kernel(t, out, out_shape, out_strides, out_size, a, a_shape, a_strides, fn):
    i = t.blockIdx.x * t.blockDim.x + t.threadIdx.x
    if i < out_size:
        out_index = to_index(i, out_shape)                          # 我负责的输出的多维下标
        a_index = broadcast_index(out_index, out_shape, a_shape)    # 对应的输入下标（广播）
        out[index_to_position(out_index, out_strides)] = fn(a[index_to_position(a_index, a_strides)])
```

思路是 HW1 的提示："每个线程处理输出的一个元素；用线程和块的编号算出全局编号；做越界检查；考虑基于步长的下标"。03 号脚本第 1 部分用同一个内核处理三种输入（32 × 32）：

| 输入 | 步长 | 与 NumPy 逐位相同 | 读事务 / 读请求 |
|---|---|---|---|
| 连续的 | (32, 1) | 是 | 4.0 |
| 转置视图 `X.permute(1, 0)` | (1, 32) | 是 | 32.0 |
| (1, 32) 的一行广播成 (32, 32) | (32, 1) | 是 | 4.0 |

内核只看 shape 和 strides，三种都算对了，没有一次额外的拷贝。但视图有代价：输出按行优先连续写，转置视图的输入却是**按列读**的，相邻线程读的地址相隔 128 字节，每个请求 32 个事务。PyTorch 里对转置后的张量调用 `.contiguous()`，就是先付一次拷贝（把数据按新的顺序重排），换后续运算的合并访存。什么时候值得？取决于这个张量之后会被读几次。

### 3.3 zip：两个输入各自广播

zip 的输出形状由广播规则决定（`shape_broadcast`：从最后一维对齐，每一维要么相等、要么有一个是 1），每个线程把自己的输出下标分别"缩"回 a 和 b 的形状。03 号脚本第 2 部分：

| 情况 | 输出形状 | 与 NumPy 逐位相同 |
|---|---|---|
| 同形状 (4, 8) + (4, 8) | (4, 8) | 是 |
| 加偏置 (16, 8) + (8,) | (16, 8) | 是 |
| 外积式 (5, 1) × (1, 7) | (5, 7) | 是 |
| 三维 (2, 1, 8) + (3, 1) | (2, 3, 8) | 是 |

逐元素运算没有求和顺序的问题，所以和 NumPy 逐位相同。广播的那个输入里，同一个元素被很多线程读到，但内存里只有一份——这就是"加偏置"不需要先把偏置复制成整个矩阵的原因。

### 3.4 reduce：一个线程一个输出，还是一个块一个输出？

归约要把一整条"纤维"（沿 dim 的所有元素）合成一个数。有两种分工：

- **一个线程算一个输出**：线程 i 负责第 i 个输出，自己沿 dim 从头加到尾。简单，不需要同步。
- **一个线程块算一个输出**（HW1 的提示）：块号 = 第几个输出；块内线程沿 dim 跨步各加一部分，写进共享内存，再做**树形归约**（§4.2），最后线程 0 写结果。

03 号脚本第 3 部分对形状 (8, 16, 64) 的张量分别沿第 0 维和第 2 维求和：

| 沿哪一维 | 分工 | 块数 | 线程数 | 与同顺序参考逐位相同 | 与 float64 精确和的最大差 | 读事务 / 读请求 |
|---|---|---|---|---|---|---|
| 0（不连续） | 一个线程 / 输出 | 32 | 1024 | 是 | 9.3e-07 | 4.0 |
| 0 | 一个块 / 输出 | 1024 | 32768 | 是 | 6.7e-07 | 8.0 |
| 2（连续） | 一个线程 / 输出 | 4 | 128 | 是 | 4.4e-06 | 32.0 |
| 2 | 一个块 / 输出 | 128 | 4096 | 是 | 2.0e-06 | 4.0 |

两点观察：

1. **数值**：所有写法都和"按同样顺序计算"的参考实现逐位相同，和 float64 的精确和相差 1e-6 量级。不同的分工就是不同的求和顺序，float32 的舍入就不同（P2 §3）。所以**测试归约内核要用容差比较**——原课 HW1 的测试也是这样。
2. **合并访存取决于"沿哪一维归约"和"怎样分工"的搭配**。沿不连续的第 0 维时，"一个线程 / 输出"让相邻线程算相邻的输出、读相邻的地址，合并得很好；"一个块 / 输出"让块内线程沿第 0 维读，地址相隔 4 KB，每个线程独占一个 sector（这一维只有 8 个元素，32 个线程里还只有 8 个有活干）。沿连续的第 2 维时**正好反过来**。没有一种分工在所有情况下都最好——PyTorch 的归约内核会根据归约的维度和形状选择不同的策略。

同一个内核换一个 `fn` 和初始值就是别的归约：求最大值用 $-\infty$（与 NumPy 逐位相同），求积用 1。

### 3.5 矩阵乘法：一个线程一个输出元素

HW1 的"简单并行"提示：每个线程算输出矩阵的一个元素，读 A 的一行和 B 的一列做内积。`kernels.matmul_kernel` 用三维网格：`blockIdx.z` 是第几个批，x 方向对应输出的列、y 方向对应行（§2.3 的规则），批维长度为 1 时步长当作 0（广播）。03 号脚本第 4 部分：

| 情况 | 输出形状 | 与同顺序参考逐位相同 | 与 NumPy 逐位相同 | 与 NumPy 的最大差 | 线程读次数 |
|---|---|---|---|---|---|
| (24, 16) @ (16, 20) | (24, 20) | 是 | 否 | 1.9e-06 | 15,360 |
| (3, 12, 16) @ (3, 16, 10) | (3, 12, 10) | 是 | 否 | 1.9e-06 | 11,520 |
| (4, 12, 16) @ (1, 16, 10)（批维广播） | (4, 12, 10) | 是 | 否 | 1.4e-06 | 15,360 |
| (24, 16) @ Bᵀ（B 是转置视图） | (24, 20) | 是 | 否 | 1.4e-06 | 15,360 |

- **和 NumPy 不是逐位相同**，差在最后几位：NumPy 调用 BLAS，BLAS 会分块、换求和顺序、用 FMA（一次舍入的乘加）。我们的内核按 k 从 0 到 m−1 依次"先乘、再加"（两次舍入），和同顺序的参考逐位相同。两者都是正确的 float32 结果。（真 GPU 上 nvcc 默认会把 `a*b + c` 编译成 FMA，所以真 GPU 的结果又会差最后一位。）
- **线程读次数 = 2·n·p·m**：每个输出元素读 m 个 A 的元素和 m 个 B 的元素。每做一次乘加（2 FLOPs）就要读两个 float（8 字节）——**算术强度 0.25 FLOP/字节**。按第 2 讲的平衡点（A100 FP32 约 10），这个内核只能用到 FP32 峰值的约 2.5%。为什么？怎么办？——这正是第 4 讲的开场。

### 3.6 主机端：把内核接到张量上

HW1 每道题的 Part B 是"集成"：在 `cuda_kernel_ops.py` 里写 `CudaKernelOps.map(fn)`、`zip(fn)`、`reduce(fn, start)`、`matrix_multiply(a, b)`，让 miniTorch 的张量运算调用你的内核。`kernels.py` 里的 `tensor_map` 等函数做的是同样的事：

```
tensor_map(dev, fn, a):
    建一个连续的输出张量（形状由输入或广播决定）
    把 a 的存储 cudaMemcpy 到显存，为输出 cudaMalloc
    块数 = ⌈输出元素数 / 每块线程数⌉
    发射 map_kernel，参数是两个设备指针 + 两组 (shape, strides) + fn
    把输出 cudaMemcpy 回来，cudaFree 两块显存
```

注意它和 §1.4–§1.5 的矛盾：每次调用都拷入拷出，正是"每次都拷来拷去"的写法。原课的 miniTorch 为了教学清晰也是这样做的（所以作业里的 CUDA 后端在小张量上不一定比 CPU 快）；PyTorch 则让张量常驻显存，只在需要时拷回。第 5 讲会把这几个内核接成迷你框架的一个"后端"，看清"前端—自动微分—后端"的分层。

---

## 4. 并行归约 · Parallel reduction

### 4.1 关键路径：并行算法和串行算法的根本区别

把 16384 个数加成一个数，04 号脚本第 1 部分比较了三种分工：

| 分工 | 线程数 | 栅栏次数 | 原子操作 | 关键路径（必须依次完成的步数） | 与精确和之差 |
|---|---|---|---|---|---|
| 一个线程从头加到尾 | 1 | 0 | 0 | 16,384 | +3.76e-04 |
| 一个块（256 线程）+ 树形归约 | 256 | 9 | 0 | 64 + 8 = 72 | −5.16e-06 |
| 16 个块 × 256 线程 + 树形归约 + atomicAdd | 4096 | 144 | 16 | 4 + 8 + 16 = 28 | −5.16e-06 |

三种写法读的数一样多（16,384 次），总的加法次数也差不多，差别全在**关键路径**：每个线程先依次累加自己那份（n / 线程数步），树形归约每轮一步（log₂256 = 8 轮），打在同一个地址上的原子操作只能排队（16 步）。只要线程足够多，**GPU 的时间由关键路径决定**——这是写并行算法时首先要问的问题。

树形归约的思想来自结合律：$(a+b)+(c+d)$ 和 $((a+b)+c)+d$ 数学上相等，但前者的两次加法可以同时做。n 个数的树有 $\log_2 n$ 层，每层内的加法互不依赖。顺便，它在数值上也更好：和 P2 的成对求和一样，树形求和的误差（5.16e-06）比从头加到尾（3.76e-04）小两个数量级。

### 4.2 树形归约的三种写法

块内树形归约有几种经典写法，原课 HW1 的可选提示给出的是第一种（这几种写法出自 Mark Harris 的教程 *Optimizing Parallel Reduction in CUDA*）：

```c
// interleaved（HW1 的 reduce0）：第 s 轮，tid 是 2s 的倍数的线程干活
for (s = 1; s < blockDim.x; s *= 2) { if (tid % (2*s) == 0) sdata[tid] += sdata[tid + s]; __syncthreads(); }
// strided：第 s 轮，前 blockDim/(2s) 个线程干活，各算 index = 2·s·tid
for (s = 1; s < blockDim.x; s *= 2) { int index = 2*s*tid; if (index < blockDim.x) sdata[index] += sdata[index + s]; __syncthreads(); }
// sequential：步长从 blockDim/2 开始减半，前 s 个线程干活
for (s = blockDim.x/2; s > 0; s >>= 1) { if (tid < s) sdata[tid] += sdata[tid + s]; __syncthreads(); }
```

interleaved 和 strided 合并的是**完全相同的配对**（第 s 轮把位置 2s·k 和 2s·k + s 合并），只是换了"由哪个线程来做"；sequential 换了一种配对（位置 k 和 k + s）。04 号脚本第 2 部分（8 行 × 1024 列，每行一个 256 线程的块，开了指令计数）：

| 写法 | 与同顺序参考逐位相同 | 线程束发射次数 | SIMT 效率 | 共享内存请求 | 共享内存轮数 | 每请求轮数 |
|---|---|---|---|---|---|---|
| interleaved | 是 | 6824 | 87.3% | 1200 | 1200 | 1.00 |
| strided | 是 | 6840 | 95.5% | 360 | 1200 | 3.33 |
| sequential | 是 | 6328 | 95.2% | 360 | 360 | 1.00 |

- **interleaved**：每一轮都让**每个线程束里的一部分线程**干活（s = 1 时每个线程束里偶数号的 16 个线程）。这是线程束分化（第 2 讲 §4.5），SIMT 效率 87%；而且每个线程束每一轮都要发共享内存请求（1200 个，另两种只有 360 个）。
- **strided**：让编号连续的线程干活，分化少了，请求也降到 360 个。但线程 tid 访问 `sdata[2s·tid]`：s = 1 时一个线程束访问第 0, 2, 4, …, 62 号字，两两落在同一个 **bank** 上（共享内存有 32 个 bank，第 k 个字在第 k mod 32 号 bank；同一 bank 的不同字要排队），s 越大冲突越严重。平均每请求 3.33 轮，总轮数还是 1200——省下来的全被 **bank 冲突**吃回去了；它的发射次数也没比 interleaved 少（多算了一行 index）。
- **sequential**：干活的线程编号连续（分化只在最后几轮 s < 32 时出现），访问的地址也连续（没有 bank 冲突）。发射次数最少，共享内存总轮数只有另两种的 30%。

（这里的计数包含了每个线程先累加 4 个数的阶段，三种写法完全相同，所以整体的发射次数差别看起来不大。）这个例子把第 2 讲的分化和第 4 讲要正式讲的 bank 冲突放到了一起：**同样的计算，"让哪个线程做哪件事"决定了效率**。

### 4.3 原子操作、竞争与不确定性

多级归约的最后一步，是把每个块的部分和加到同一个位置上。04 号脚本第 3 部分（2048 个数）：

| 写法 | 原子操作次数 | 检测到的竞争 | 模拟器上的结果 |
|---|---|---|---|
| `OUT[0] = OUT[0] + X[i]`（不用原子操作） | 0 | 4094 | −95.087730 |
| `atomicAdd(&OUT[0], X[i])`（每个线程一次） | 2048 | 0 | −95.087730 |
| 块内树形归约 + 每块一次 atomicAdd（8 个块） | 8 | 0 | −95.087700 |

第一行最有教育意义：**不用原子操作的写法在模拟器上"算对了"**——因为模拟器让线程一个接一个地执行。真 GPU 上，两个线程可能同时读到旧值、各自加完再写回，其中一个的加法就丢了（P4 §3 的"读-加-写"竞争）。这正是"测试通过不代表没有竞争"的典型例子，模拟器的竞争检测器把它抓了出来。

`atomicAdd` 把"读-加-写"变成硬件保证不可分割的一步，结果正确；代价是打在同一个地址上的原子操作只能排队。先在块内归约、每块只做一次原子操作，原子操作从 2048 次降到 8 次。

但原子操作只保证"不丢"，**不保证顺序**：真 GPU 上每次运行，各个块完成的先后不同，加法的顺序就不同。float32 不满足结合律，所以结果每次可能差最后几位。脚本对照：同样 2048 个 float32 按 20 种随机顺序依次相加，得到 15 个不同的结果（范围 −95.087830 到 −95.087662）。这是 GPU 训练"不能逐位复现"的常见来源之一。需要确定性时，要用固定顺序的归约（PyTorch 的 `torch.use_deterministic_algorithms(True)` 会换用确定性的内核，通常更慢）。

---

## 5. 小结 · Summary

1. **CUDA 程序的五步**：cudaMalloc（传 `&dA`，因为要修改指针本身）→ cudaMemcpy 拷入（先目的、后来源，字节数）→ 发射内核 → cudaMemcpy 拷出 → cudaFree。
2. **PCIe 比显存慢几十倍**：只用一次的数据不值得搬到 GPU 上；正确的用法是数据常驻显存、被一串内核反复使用（再进一步是算子融合）。
3. **三种内存**：寄存器（线程私有）、共享内存（块内共享、块结束即消失）、全局内存（所有线程可见、跨内核发射一直存在）。`__global__` 由主机调用在设备执行，`__device__` 由设备调用在设备执行。
4. **多维编号**就是把 (blockIdx, threadIdx) 按行优先展平，x 变化最快；线程束按块内编号切分。行优先矩阵要让 **threadIdx.x 对应列号**，否则事务数成倍增加；二维也要对行、列分别做越界检查。
5. **HW1 的四类内核**：每个线程（或每个块）负责一个输出位置，用 `to_index` → `broadcast_index` → `index_to_position` 找到输入。同一个内核对连续张量、转置视图、广播都正确，但视图会改变访存模式。
6. **归约**：一个线程 / 输出 vs 一个块 / 输出，哪个合并取决于沿哪一维归约。浮点归约的结果依赖顺序，测试要用容差或同顺序参考。
7. **朴素矩阵乘法**的算术强度只有 0.25 FLOP/字节——第 4 讲的问题。
8. **并行归约**看关键路径：树形归约把 n 步变成 log₂n 层；sequential 写法同时避开了分化和 bank 冲突；跨块汇总用原子操作，它保证正确但不保证顺序。

## 6. 常见误区 · Common pitfalls

- **"`cudaMalloc(dA, size)` 就行。"** 必须传 `&dA`：函数要把地址写回你的变量。传 `dA` 只是把一个未初始化的值复制进去。
- **"`cudaMemcpy(src, dst, …)`。"** 顺序是先目的、后来源，和 `memcpy` 一样；`count` 是字节数，不是元素个数。
- **"把计算搬到 GPU 上总会更快。"** 对只用一次的数据，PCIe 拷贝比计算贵得多（向量加法拷贝占 96%）。
- **"`i = blockIdx.x * blockDim.x + threadIdx.x` 当行号很自然。"** 行优先矩阵里 x 方向应该对应列号，否则一个线程束访问 32 个不同的行。
- **"例子里 N = 32 没有越界检查，我的代码也不用。"** slides 的例子正好整除；换一个 N 就越界。
- **"转置视图是免费的。"** 不拷贝数据是免费的，但之后的读取变成了按列读，事务可能多 8 倍。
- **"归约结果和 NumPy 不完全一样，所以我的内核错了。"** 浮点加法不满足结合律，换一种分工就换了一种顺序；要和同顺序参考比，或用容差。
- **"树形归约里写 `if (tid < s) { …; __syncthreads(); }`。"** 栅栏不能放在只有部分线程执行的分支里（第 2 讲 `BarrierError`）；`if` 只包住"加"那一行。
- **"不用原子操作也测对了。"** 在模拟器上或碰巧的时序下"算对"不代表没有竞争。多线程往同一个位置写，要么用原子操作，要么先归约。
- **"用了原子操作，结果就可复现。"** 原子操作不保证顺序，浮点结果每次可能差最后几位。

## 7. 练习 · Exercises

1. 写出 `cudaMalloc` 的原型，解释为什么它的第一个参数是 `void**`，以及调用时为什么写 `&dA`。
2. 在 A6000 上（PCIe 32 GB/s、显存 768 GB/s，每次拷贝 + 10 μs、每次发射 + 5 μs）对 10⁷ 个 float32 做向量加法，拷贝和内核各要多久？把答案和 01 号脚本的表对照。
3. `dim3 Dg(4, 2, 1); dim3 Db(8, 8, 1);` 一共多少个块、多少个线程？每个块有几个线程束？
4. 在 slides 的 fullKernel（网格 (2,3,4)、块 (2,4,8)）里，blockIdx = (1, 2, 3)、threadIdx = (1, 3, 7) 的线程算出的 tid 是多少？
5. 1000 × 1000 的矩阵加法用 16 × 16 的块，网格多大？一共发射多少个线程？多少个线程什么也不做？
6. slides 的 MatAddKernel（x 方向当行号）用在 1024 × 1024 的矩阵上、块 32 × 32。一个线程束的一条读指令产生几个 32 字节事务？改成 x 方向当列号呢？
7. 一个形状 (2, 3, 4) 的连续张量做 `permute(2, 0, 1)`，新的形状和步长是什么？新张量的元素 [3, 1, 2] 在存储里的位置是多少？
8. a 的形状 (3, 1)、b 的形状 (1, 4)，`zip(add)` 的输出形状是什么？输出下标 (2, 3) 对应 a 和 b 的哪个下标？
9. 对形状 (8, 16, 64) 的张量沿第 2 维求和，用"一个块 / 输出、每块 32 个线程、sequential 树形归约"。一共几个块？每个线程先累加几个元素？树形归约几轮？每个块几次栅栏？全部块一共几次？和 03 号脚本的表对照。
10. 朴素矩阵乘法 C (n×p) = A (n×m) · B (m×p)，每个线程算一个输出。一共读多少次全局内存？做多少 FLOPs？算术强度（按线程的读取字节算）是多少？
11. interleaved 树形归约，块大小 256。第 s = 1 轮每个线程束有几个线程干活？s = 16 轮呢？s = 32 轮有几个线程束里有线程干活？
12. 2048 个线程各对同一个地址做一次 atomicAdd，和"8 个块各自树形归约、再各做一次 atomicAdd"相比，原子操作那一段的关键路径各是多少步？为什么后者在真 GPU 上快得多？

### 参考答案

1. `cudaError_t cudaMalloc(void** devPtr, size_t size)`。返回值被错误码占用了，新分配的显存地址只能通过参数"带回来"；C 是按值传参，要修改调用者的指针变量 `dA`，就要传它的地址 `&dA`（类型 `float**`，转成通用的 `void**`）。
2. 拷贝：3 × (10 μs + 4×10⁷ 字节 / 32×10⁹) = 3 × (0.01 + 1.25) ms = **3.78 ms**；内核：5 μs + 12×10⁷ / 768×10⁹ ≈ 0.005 + 0.156 = **0.161 ms**。拷贝占 95.9%，与脚本一致。
3. 块：4 × 2 × 1 = **8**；每块 8 × 8 = 64 个线程，共 **512** 个线程；每块 64 / 32 = **2** 个线程束。
4. block_id = 1 + 2 × 2 + 3 × 2 × 3 = 23；thread_offset = 1 + 3 × 2 + 7 × 2 × 4 = 63；tid = 23 × 64 + 63 = **1535**（最后一个线程）。
5. ⌈1000/16⌉ = 63，网格 **63 × 63**；线程 63² × 256 = **1,016,064**；空转 1,016,064 − 1,000,000 = **16,064**。
6. 一个线程束是 threadIdx.x = 0..31（同一个 threadIdx.y）：x 当行号时 32 个线程访问 32 个不同的行，地址相隔 4096 字节，**32 个事务**；x 当列号时访问同一行的 32 个相邻 float（128 字节），**4 个事务**。
7. 原步长 (12, 4, 1)；新形状 **(4, 2, 3)**、新步长 **(1, 12, 4)**；位置 = 3 × 1 + 1 × 12 + 2 × 4 = **23**。
8. 输出 **(3, 4)**；a 的下标 **(2, 0)**（第 1 维大小为 1，取 0），b 的下标 **(0, 3)**。
9. 输出 8 × 16 × 1 = **128** 个块；每个线程累加 64 / 32 = **2** 个元素；log₂32 = **5** 轮；每块栅栏 1（写共享内存之后）+ 5 = **6** 次；共 128 × 6 = **768** 次，与脚本表中的 768 一致。
10. 读 **2·n·p·m** 次（每个输出 m 次 A、m 次 B）；FLOPs **2·n·p·m**；算术强度 = 2npm / (8npm) = **0.25 FLOP/字节**。
11. s = 1：tid 是 2 的倍数的线程，每个线程束 **16** 个；s = 16：tid 是 32 的倍数，每个线程束 **1** 个（SIMT 效率 1/32）；s = 32：tid 是 64 的倍数，**8 个线程束里只有 4 个**有线程干活（线程束 0、2、4、6），每个只有 1 个线程。
12. 前者 **2048 步**（同一地址的原子操作串行）；后者树形归约 8 轮 + **8 步**原子操作。真 GPU 上同一地址的原子操作要在 L2 里排队处理，2048 次的排队远比 8 次长；而且块内归约用的是快得多的共享内存。

## 8. 延伸阅读 · Further reading

- **PMPP 第 4 版**第 2 章（向量加法、cudaMalloc / cudaMemcpy、错误检查）、第 3 章（多维网格、彩色转灰度与图像模糊、矩阵乘法）、第 10 章（归约：树形归约、分化、合并、线程粗化、原子操作）。
- **Mark Harris, *Optimizing Parallel Reduction in CUDA***（NVIDIA 教程幻灯片）：从 interleaved 到 sequential，再到循环展开、线程粗化，一步步把归约做到接近带宽上限。§4.2 的三种写法出自这里。
- **原课 HW1 说明**（https://llmsystem.github.io/llmsystemhomework/ ）与 **CUDA_Code_Examples.ipynb**（原课 llmsys_code_examples 仓库的 simple_cuda_demo，slides 推荐在 Colab 上运行）：本讲的四类内核在真 GPU 上的版本。
- **NVIDIA CUDA C++ Programming Guide**："Programming Interface → CUDA Runtime"（设备内存、异步并发执行、流）；"Asynchronous Data Copies"与 Hopper 调优指南里的 TMA。
- **PyTorch 文档 *Reproducibility***：哪些 CUDA 运算是非确定性的，以及 `use_deterministic_algorithms` 的代价。

## 9. 下一讲预告 · Next lecture

第 3 讲留下了一个数字：朴素矩阵乘法的算术强度只有 **0.25 FLOP/字节**。原课第 4 讲从这里出发：在 A100 上，这个内核受显存带宽限制，只能用到 FP32 峰值的 **2.48%**。怎样把它提高几十倍？答案是**分块**（tiling）：让一个线程块把 A 和 B 的一小块搬进共享内存，块内的线程反复使用它们。之后是**内存并行**：合并访存、转置的三个版本、共享内存的 **bank 冲突**和"+1 填充"的技巧；再之后是**稀疏矩阵**（CSR 格式与 SpMV 内核）和 **cuBLAS**（为什么通常不该自己写 GEMM，以及它"列优先"的接口怎么用）。模拟器会把每一步省下的访存都数出来。
# 第 4 讲 · GPU 加速：分块、内存并行、稀疏矩阵与 cuBLAS · GPU Acceleration

> **对应 Maps to**：CMU 11-868 / 11-968 *Large Language Model Systems*（Spring 2026，Lei Li），Lecture 4 · Mon 1/26 · **GPU Acceleration**。
> **阅读材料**：PMPP 第 4 版第 5 章（存储架构与数据局部性：分块）、第 6 章（性能考虑：合并访存、线程粗化）、第 14 章（稀疏矩阵计算）；NVIDIA *CUDA Programming Guide*；cuBLAS 文档。
> **依据说明**：本讲以原课第 4 讲公开的 slides（llmsys-04-gpu-acceleration.pdf，50 页）为主线：计算–全局内存访问比与 A100 上 2.48% 的计算、存储层次的延迟、CUDA 的内存模型、分块矩阵乘法的步骤与 MatMulTiledKernel、内存的限制、DRAM 突发与合并访存、矩阵转置的三个版本与 bank 冲突、CSR 与 SpMV 内核、cuBLAS 的接口。**模拟器上的全部计数、TILE 的扫描、漏掉栅栏 / 边界处理的反例、线程粗化、"一个线程束一行"的 SpMV、2:4 稀疏、行优先交换技巧和 roofline 阶梯都是本讲义自己设计的，不代表课堂上讲了什么**。
> **配套代码 Code**：[`code/lecture04/`](../code/lecture04/)（5 个脚本 + `accel.py`；复用第 2 讲的 `simt.py`、`gpu_model.py`；每个脚本 6 秒以内，输出重复运行逐位相同）
> **预计阅读时间**：4–5 小时（含运行代码）

---

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

### 0.1 学习目标 Learning goals

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

1. 算出一个内核的**计算–全局内存访问比**，并用它和 GPU 的平衡点解释"朴素矩阵乘法为什么只能用到 A100 峰值的 2.48%"。
2. 说出 CUDA 的**内存模型**（寄存器、本地、共享、全局、常量）各自的延迟、作用域、生命周期，以及变量声明怎样决定它放在哪里。
3. 写出**分块矩阵乘法**，解释它为什么把全局内存读取减少到 1/TILE、两个 `__syncthreads()` 各防什么、边界怎样处理、TILE 受哪些资源限制，以及**线程粗化**再省在哪里。
4. 用"一个线程束在同一条指令里访问了哪些地址"判断**合并访存**；理解矩阵转置的三个版本，算出共享内存的 **bank 冲突**，并解释"+1 填充"为什么有效。
5. 用 **CSR** 表示稀疏矩阵，写出 SpMV 内核，判断多稀疏才划算，并理解非结构化稀疏和 **2:4 结构化稀疏**的区别。
6. 正确调用 **cuBLAS**：列优先、leading dimension、transa / transb，以及让行优先张量直接调用 cuBLAS 的"交换技巧"；知道为什么几乎总是用库而不是自己写 GEMM。

### 0.2 先修与本讲范围 Prerequisites & scope

| 默认你已经会 Assumed | 本讲只给直觉，后面再展开 Previewed only |
|---|---|
| 第 2 讲：SM、线程束、合并访存的定义、占用率、延迟隐藏 | 张量核心的编程（mma 指令）、Triton / CUTLASS（第 10–13 讲） |
| 第 3 讲：五步、二维线程编号（x 对应列）、HW1 的朴素矩阵乘法、树形归约与 bank 冲突的初见 | FlashAttention 的分块（第 21 讲：同样的思想用在注意力上） |
| 第 1 讲 §4.4：算术强度、平衡点、roofline | 稀疏的 MoE（第 17 讲）、量化后的 GEMM（第 19–20 讲） |

### 0.3 本讲地图 Map

```
§1  访存效率决定一切：计算–全局访存比、2.48%、延迟、内存模型、重复读取    ← 代码 01
§2  分块：共享内存里的复用；两个栅栏；边界；资源限制；线程粗化            ← 代码 02
§3  内存并行：DRAM 突发与合并访存；转置的三个版本；bank 冲突与 +1         ← 代码 03
§4  稀疏矩阵：CSR、SpMV、多稀疏才划算、2:4 结构化稀疏                     ← 代码 04
§5  cuBLAS：接口、列优先、leading dimension、行优先技巧、roofline 阶梯    ← 代码 05
§6  小结   §7  常见误区   §8  练习   §9  延伸阅读   §10 下一讲预告
```

slides 把本讲概括为四种 GPU 加速技术：**分块**（PMPP 第 5 章）、**内存并行**（第 5、6 章）、**稀疏矩阵乘法**、**cuBLAS**。前三种有一个共同点：**计算本身一点没变，变的只是"哪个线程在什么时候从哪里取数"**。

---

## 1. 访存效率决定一切 · Memory access efficiency is critical

### 1.1 计算–全局内存访问比

第 3 讲的朴素矩阵乘法内核，每个线程算 C 的一个元素：

```c
__global__ void MatMulKernel(float *a, float *b, float *c, int N) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row >= N || col >= N) return;
    float Pvalue = 0.0;
    for (int k = 0; k < N; k++)
        Pvalue += a[row * N + k] * b[k * N + col];     // 1 次乘 + 1 次加；读全局内存 2 次（各 4 字节）
    c[row * N + col] = Pvalue;
}
```

slides 定义了一个量：**计算–全局内存访问比**（compute-to-global-memory-access ratio）= 每从全局内存读一个字节，做多少次浮点运算。循环体做 2 次运算、读 8 字节，比值是 **0.25 FLOP/字节**。它就是第 1 讲的"算术强度"，只是专门按全局内存的访问来数。

### 1.2 A100 上的 2.48%

slides 用 A100 80GB PCIe 的规格算了一笔账（01 号脚本第 1 部分复现）：

| GPU | 显存带宽 | FP32 | TF32 张量核心 | 带宽能喂饱的算力 | 占 FP32 峰值 | 占 TF32 峰值 |
|---|---|---|---|---|---|---|
| A100 80GB PCIe | 1,935 GB/s | 19.5 TFLOPS | 156 TFLOPS | 1935 × 0.25 = 483.75 GFLOPS | **2.48%** | 0.31% |
| H100 SXM | 3,350 GB/s | 67 TFLOPS | 494.7 TFLOPS | 837.5 GFLOPS | 1.25% | 0.17% |

显存每秒只能送来 1,935 GB 数据，每字节只能支撑 0.25 次运算，所以无论计算单元多快，这个内核最多跑 483.75 GFLOPS——**访存受限**（memory-bound）。要让 FP32 核心跑满，比值要提高到平衡点 10.1；要让 TF32 张量核心跑满，要 81。换句话说：**每从显存读一个字节，要设法让它参与几十次运算**。这一讲的前三节都是在想办法做到这一点。

### 1.3 读数据比计算慢得多

slides 用一条加法说明为什么"读"这么贵。`C[i] = A[i] + B[i]` 编译成 PTX（GPU 的中间汇编）是四条指令：

| 指令 | 做什么 | 延迟 |
|---|---|---|
| `ld.global.f32 %f1, [%rd1];` | 读 A[i] | 约 500 个周期 |
| `ld.global.f32 %f2, [%rd2];` | 读 B[i] | 约 500 个周期 |
| `add.f32 %f3, %f1, %f2;` | 加法 | 约 1 个周期 |
| `st.global.f32 [%rd3], %f3;` | 写 C[i] | （发出去就行，不需要等） |

slides 给出的 CUDA 存储层次延迟（量级）：**寄存器约 1 个周期；共享内存、L1 缓存、常量缓存约 5 个周期；全局内存约 500 个周期**。共享内存比全局内存快约 100 倍，而且就在 SM 上。第 2 讲 §4.7 讲过 GPU 用大量线程束掩盖这 500 个周期；但掩盖得再好，带宽也是硬上限（§1.2）。真正的出路是**少读全局内存**。

### 1.4 CUDA 的内存模型

slides 第 11 页给出变量声明和内存的对应（01 号脚本第 3 部分）：

| 变量声明 | 内存 | 作用域 | 生命周期 |
|---|---|---|---|
| `int var;`（内核里的局部变量） | 寄存器 | 线程 | 内核这一次执行 |
| `int varArr[N];`（局部数组） | 本地内存（local） | 线程 | 内核这一次执行 |
| `__device__ __shared__ int SharedVar;` | 共享内存 | 线程块 | 内核这一次执行 |
| `__device__ int GlobalVar;` | 全局内存 | 网格（所有线程） | 整个应用程序 |
| `__device__ __constant__ int constVar;` | 常量内存 | 网格 | 整个应用程序 |

两个容易踩的坑：**局部数组**通常放不进寄存器，会被放到"本地内存"——名字叫本地，其实在显存里，和全局内存一样慢；**共享内存是程序员手动管理的缓存**——搬什么、什么时候搬、什么时候可以覆盖，都要自己决定（§2.5 的两个栅栏）。

### 1.5 机会在哪里：重复读取

slides 第 12–13 页指出了优化的机会：**线程块里的线程会用到相同的数据**。C 的同一行的所有元素都要 A 的同一行，同一列的都要 B 的同一列。01 号脚本第 4 部分在模拟器上把朴素矩阵乘法（N = 32）的字节数按三种口径数了一遍：

| 口径 | 字节数 | FLOP / 字节 |
|---|---|---|
| 线程要求读写的字节（每次访问 4 字节） | 266,240 | 0.25 |
| 合并之后实际搬运的 32 字节事务 | 135,168 | 0.48 |
| 数据本身（A、B、C 各 N² 个 float） | 12,288 | 5.33 |

A 的每个元素被读了 N = 32 次。按 slides 的口径（每次读都去显存），比值是 0.25；如果每个数只从显存读一次，比值是 N/6 = 5.33，而且 N 越大越高。"要求读的"和"数据本身"差了 22 倍——这就是**重复读取**，也是优化的空间。真实 GPU 的 L1 / L2 缓存会自动接住一部分重复读取，但容量有限、命中与否不受控制；**分块**用共享内存把"重复使用"变成程序里明确的、有保证的事情。

---

## 2. 分块 · Tiling

### 2.1 直觉：把要反复用的书搬到书桌上

想象你要写一份报告，需要反复查阅图书馆书库里的几十本书。每查一个数就跑一趟书库（全局内存，500 个周期），太慢了。更好的办法：先把接下来一段时间要用的几本书一次搬到书桌上（共享内存，5 个周期），查完这几本再换下一批。书桌只有这么大（共享内存容量有限），所以要分批（分块）；几个人共用一张书桌（线程块共享），换书之前要确认大家都查完了（栅栏）。

### 2.2 分块矩阵乘法的步骤

slides 第 14–18 页一步步画出了这个过程。C 被分成 TILE × TILE 的小块，每个线程块负责一块，每个线程负责其中一个元素。要算这一块，需要 A 的一条"横带"（TILE 行）和 B 的一条"竖带"（TILE 列）；把它们沿 k 方向切成 TILE × TILE 的小块，一相（phase）处理一对：

1. **（同时）** 把 A、B 的第一对小块搬进共享内存：块内**每个线程搬一个元素**；`__syncthreads()` 等大家搬完。
2. 每个线程用共享内存里的小块算出部分和（As 的一行 · Bs 的一列，TILE 次乘加）；`__syncthreads()` 等大家算完。
3. **（同时）** 搬下一对小块……
4. 继续累加部分和……直到走完整条横带和竖带。

slides 第 36 页的内核（原课的现场编程环节）：

```c
#define TILE_WIDTH 2
__global__ void MatMulTiledKernel(float* d_A, float* d_B, float* d_C, int N) {
    __shared__ float As[TILE_WIDTH][TILE_WIDTH];
    __shared__ float Bs[TILE_WIDTH][TILE_WIDTH];
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    float Cvalue = 0;
    for (int ph = 0; ph < N / TILE_WIDTH; ++ph) {
        As[threadIdx.y][threadIdx.x] = d_A[row * N + ph * TILE_WIDTH + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = d_B[(ph * TILE_WIDTH + threadIdx.y) * N + col];
        __syncthreads();
        for (int k = 0; k < TILE_WIDTH; ++k)
            Cvalue += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        __syncthreads();
    }
    d_C[row * N + col] = Cvalue;
}
```

为什么能省？块里同一行的 TILE 个线程都要 A 的同一个元素——现在由**一个**线程把它搬进共享内存，**TILE 个**线程从共享内存里读。每个从全局内存读进来的数被用了 TILE 次，全局内存读取减少到 1/TILE，计算–全局访存比提高到 **0.25 × TILE**。

### 2.3 实验：TILE 从 1 到 32

02 号脚本第 1 部分（N = 32，[`accel.matmul_tiled`](../code/lecture04/accel.py)）：

| 写法 | 线程读全局内存 | 全局读事务 | 栅栏次数 | 与朴素版逐位相同 | 计算–全局访存比 | A100 FP32 峰值上限 |
|---|---|---|---|---|---|---|
| 朴素 | 65,536 | 4,096 | 0 | — | 0.25 | 2.5% |
| TILE = 2 | 32,768 | 16,384 | 8,192 | 是 | 0.5 | 5.0% |
| TILE = 4 | 16,384 | 4,096 | 1,024 | 是 | 1 | 9.9% |
| TILE = 8 | 8,192 | 1,024 | 128 | 是 | 2 | 19.8% |
| TILE = 16 | 4,096 | 512 | 16 | 是 | 4 | 39.7% |
| TILE = 32 | 2,048 | 256 | 2 | 是 | 8 | 79.4% |

- 线程读全局内存的次数正好是 **2N³ / TILE**。运算总量没变，读的地方变了：TILE = 16 时共享内存的读取 65,536 次（和朴素版的全局读取一样多），外加搬进来时写共享内存的 4,096 次。
- **结果和朴素版逐位相同**：每个线程仍然按 k = 0, 1, …, N−1 的顺序累加。分块是纯粹的性能优化，不改变数值。
- **一个反常的数**：TILE = 2 时线程读全局内存的次数减半，全局读事务却是朴素版的 **4 倍**。2 × 2 的块只有 4 个线程，凑不满一个线程束，每条读指令只取 2 个数，合并访存完全失效。块不能太小——一行至少要有 8 个线程（第 3 讲 §2.5）。
- 按 slides 的口径，TILE = 16 时能用到 A100 FP32 峰值的约 40%，TILE = 32 时约 80%。要超过平衡点还需要更多手段（§2.8、§5.5）。

### 2.4 两个栅栏各防什么

02 号脚本第 2 部分故意去掉其中一个栅栏（N = 32，TILE = 8）：

| 写法 | 错误的元素 | 其中 NaN | 检测到的竞争 |
|---|---|---|---|
| 两个栅栏都在 | 0 / 1024 | 0 | 0 |
| 漏掉第二个（算完之后） | 1008 / 1024 | 0 | 26,880 |
| 漏掉第一个（搬完之后） | 1008 / 1024 | 1008 | 35,840 |

- **第一个栅栏**（搬完之后）保证"**读之前大家都写完了**"（read-after-write）。漏掉它，线程会读到同伴还没搬进来的位置。模拟器的共享内存初始是 NaN，所以结果里出现了 NaN；真 GPU 上读到的是上一相留下的旧值，错得更隐蔽。
- **第二个栅栏**（算完之后）保证"**写之前大家都读完了**"（write-after-read）。漏掉它，跑得快的线程会进入下一相、覆盖 As / Bs，而别的线程还在用旧的那一块。结果是错的数，而不是 NaN。
- 1024 个元素里有 16 个"碰巧对了"：模拟器按编号依次执行线程，每个块（共 16 个）里编号最大的那个线程运行时，它要的数都已经被前面的线程搬好了。真 GPU 上的执行顺序不确定，同样的 bug 可能时有时无——这正是并发 bug 难调的原因，也是模拟器另外做竞争检测的原因。

### 2.5 边界：N 不是 TILE 的倍数

slides 的内核用 `ph < N / TILE_WIDTH`（整数除法），而且不检查越界——slides 的例子 N 正好是 TILE 的倍数。02 号脚本第 3 部分取 N = 10、TILE = 4：

| 写法 | 结果 |
|---|---|
| slides 的写法，网格 ⌈10/4⌉ = 3 | 越界：blockIdx=(0,2,0)、threadIdx=(0,2,0) 的线程访问 A[100] |
| slides 的写法，网格 10 // 4 = 2 | 不报错；只算了左上角 8 × 8，每个元素只加了前 8 项（左上角最大差 5.59），其余 36 个元素是垃圾值 |
| 修正：相数 ⌈N/TILE⌉，越界位置往共享内存写 0，只让矩阵内的线程写 C | 与 NumPy 差 4.8e-07 |

修正的要点：相数向上取整；搬运时如果行或列超出 N，就往共享内存写 0（0 乘任何数都是 0）；最后只让矩阵内的线程写 C。**越界的线程不能提前 `return`**——它们还要参与搬运和栅栏（第 2 讲的 `BarrierError`：`__syncthreads()` 必须被块内所有线程执行到）。

### 2.6 TILE 能有多大：资源限制

slides 第 20–21 页列出了限制：H100 每个 SM 有 256 KB 的 L1 / 共享内存、64K 个寄存器；A100 是 192 KB 和 64K（slides 的表里 A100 写的是 104 个 SM，NVIDIA 白皮书给出的是 108 个）。分块用的共享内存 As + Bs = 2 × TILE² × 4 字节：TILE = 32 时 As、Bs 各 32 × 32 × 4 = 4 KB（slides 的例子），合计 8 KB，离每块的上限（A100 163 KB、H100 227 KB）还很远。02 号脚本第 4 部分：

| TILE | 每块线程 | 每块共享内存 | A100 每 SM 驻留 | H100 每 SM 驻留 | 计算–访存比 |
|---|---|---|---|---|---|
| 8 | 64 | 0.5 KB | 32 块 / 64 束 = 100% | 同左 | 2 |
| 16 | 256 | 2 KB | 8 块 / 64 束 = 100% | 同左 | 4 |
| 32 | 1024 | 8 KB | 2 块 / 64 束 = 100% | 同左 | 8 |
| 64 | 4096 | 32 KB | 超过每块 1024 个线程，无法发射 | | |

真正先撞上的是**每块最多 1024 个线程**：只要还是"一个线程一个输出"，TILE 就不能超过 32，比值就停在 8。要让每个块覆盖更大的输出块，必须让**每个线程算好几个输出**，用寄存器再做一层复用（§2.7、§5.5）。

> **关于 slides 第 21 页的寄存器算术。** slides 写"1024 个线程时，每个线程只能用 16384 / 1024 = 16 个寄存器"。16384 是**一个分区**的寄存器数（64 KB）；整个 SM 有 4 个分区、65,536 个寄存器，1024 个线程的块分布在 4 个分区上，所以每个线程最多能用 64 个；SM 驻留满 2048 个线程时每个 32 个（第 2 讲 §3.5）。结论的方向不变：线程越多，每个线程能用的寄存器越少。

### 2.7 线程粗化：一个块算好几个输出块

在分块版本里，横向相邻的两个块要的是**同一条 A 的横带**，却各自从全局内存读一遍。**线程粗化**（thread coarsening，PMPP 6.3 节）让一个块负责横向相邻的 COARSE 个输出块、每个线程算 COARSE 个输出：A 的小块每相只读一次，被 COARSE 个输出块共用。02 号脚本第 5 部分（N = 32，TILE = 8）：

| COARSE | 块数 | 每线程的累加器 | 读 A | 读 B | 全局读合计 | 计算–全局访存比 |
|---|---|---|---|---|---|---|
| 1 | 16 | 1 | 4,096 | 4,096 | 8,192 | 2.00 |
| 2 | 8 | 2 | 2,048 | 4,096 | 6,144 | 2.67 |
| 4 | 4 | 4 | 1,024 | 4,096 | 5,120 | 3.20 |

A 的读取降到 1/COARSE，结果仍与朴素版逐位相同。代价是每个线程要 COARSE 个累加器（更多寄存器 → 占用率可能下降），块数减少（并行度降低；矩阵小时 GPU 可能喂不饱）。粗化是"用更少的线程做更多的事"，在并行度过剩时才划算。把它推到底——块在共享内存里分块、线程在寄存器里再分块（每个线程算 8 × 8 之类的小块）——就是 cuBLAS 这类库的做法（§5.5）。

---

## 3. 内存并行：合并访存与 bank 冲突 · Memory parallelism

### 3.1 DRAM 突发与合并访存

分块解决的是"读多少次"，这一节解决"怎么读"。slides 第 22 页：显存每次至少传 32 字节；**一个线程束里连续的访问会被合并（coalesced）在一起**；C 和 CUDA 的多维数组按行优先存放，这让 DRAM 的突发传输（burst）成为可能，比一个个单独访问快得多。

03 号脚本第 1 部分让一个线程束的 32 个线程按不同的步长和起点各读一个 float：

| 步长（元素） | 起点偏移 | 事务（32 字节） | 有用字节占比 |
|---|---|---|---|
| 1 | 0 | 4 | 100% |
| 1 | 1 | 5 | 80% |
| 1 | 8 | 4 | 100% |
| 2 | 0 | 8 | 50% |
| 4 | 0 | 16 | 25% |
| 8、16、32 | 0 | 32 | 12.5% |
| 0（都读同一个数） | 0 | 1 | 12.5% |

- 步长 1、起点对齐：32 个 float = 128 字节 = 4 个 sector，一个字节都不浪费（slides 第 23 页"一次读取服务一个线程束的多个线程"）。
- 起点偏移 1 个元素就跨了 5 个 sector；偏移 8 个元素（32 字节）又对齐了。`cudaMalloc` 返回的地址至少 256 字节对齐，所以"从数组开头按线程编号读"总是对齐的。
- 步长越大越浪费（slides 第 24 页"相隔 32 字节，两个线程要两次读取"）；步长 ≥ 8 后每个线程独占一个 sector。所有线程读同一个数是广播，1 个事务。

### 3.2 矩阵转置：读和写总有一边是按列走的

slides 用矩阵转置作为例子。朴素版本（第 26 页）：

```c
#define INDX(row, col, ld) (((row) * (ld)) + (col))       // 行优先的二维下标
__global__ void naive_cuda_transpose(int m, float *a, float *c) {
    int myCol = blockDim.x * blockIdx.x + threadIdx.x;
    int myRow = blockDim.y * blockIdx.y + threadIdx.y;
    if (myRow < m && myCol < m)
        c[INDX(myCol, myRow, m)] = a[INDX(myRow, myCol, m)];
}
```

slides 第 27–28 页问：原数组里的位置是多少？转置后的位置是多少？访问合并吗？答案：读 `a[myRow * m + myCol]` 时，一个线程束（threadIdx.x = 0..31）读同一行的 32 个相邻元素——**合并**；写 `c[myCol * m + myRow]` 时，myCol 随 threadIdx.x 变化，相邻线程写**同一列**、相隔 m 个元素——**不合并**。转置的本质矛盾：读和写，总有一边是按列走的。

### 3.3 用共享内存把"按列走"挪进片上

slides 第 29–30 页的关键思路：先把一个小块按行（合并地）读进共享内存，`__syncthreads()`，再按行（合并地）写出它的转置——"按列走"的那一步发生在共享内存里，而共享内存不需要合并。

```c
__global__ void smem_cuda_transpose(int m, float *a, float *c) {
    __shared__ float smemArray[THREADS_PER_BLOCK_X][THREADS_PER_BLOCK_Y];
    const int tileCol = blockDim.x * blockIdx.x;
    const int tileRow = blockDim.y * blockIdx.y;
    smemArray[threadIdx.x][threadIdx.y] = a[INDX(tileRow + threadIdx.y, tileCol + threadIdx.x, m)];
    __syncthreads();
    c[INDX(tileCol + threadIdx.y, tileRow + threadIdx.x, m)] = smemArray[threadIdx.y][threadIdx.x];
}
```

03 号脚本第 2 部分（64 × 64 矩阵，块 32 × 32）：

| 内核 | 结果正确 | 读 a：事务/请求 | 写 c：事务/请求 | 全局事务合计 | 共享内存请求 | 共享内存轮数 |
|---|---|---|---|---|---|---|
| naive | 是 | 4 | 32 | 4,608 | 0 | 0 |
| 共享内存 | 是 | 4 | 4 | 1,024 | 256 | 4,224 |
| 共享内存，数组声明成 [32][33] | 是 | 4 | 4 | 1,024 | 256 | 256 |

全局事务从 4,608 降到 1,024——读写都合并了。但第二行出现了新问题：共享内存 256 个请求却要 4,224 轮。

### 3.4 bank 冲突与 "+1"

slides 第 31 页：**共享内存分成 32 个 bank，每个 bank 宽 4 字节；访问同一个 bank 的不同元素要依次进行**。连续的 4 字节字依次落在 bank 0、1、……、31、0、1……，即第 k 个字在第 k mod 32 号 bank。一个线程束的 32 个线程如果访问 32 个不同 bank 的字，一轮完成；落在同一个 bank 的不同字要排队（读同一个字是广播，不排队）。03 号脚本第 3 部分把 bank 编号算了出来：

| 访问 | 线程 0–5 访问的字号 | 对应的 bank | 不同 bank 数 | 需要几轮 |
|---|---|---|---|---|
| `smemArray[32][32]`，写 `[tx][ty]` | 0, 32, 64, 96, 128, 160… | 0, 0, 0, 0, 0, 0… | 1 | **32** |
| `smemArray[32][33]`，写 `[tx][ty]` | 0, 33, 66, 99, 132, 165… | 0, 1, 2, 3, 4, 5… | 32 | 1 |
| `smemArray[32][32]`，读 `[ty][tx]` | 0, 1, 2, 3, 4, 5… | 0, 1, 2, 3, 4, 5… | 32 | 1 |

写 `smemArray[threadIdx.x][threadIdx.y]` 时，一个线程束写的是同一**列**的 32 个位置：线程 tx 访问第 tx × 32 号字，全都是 bank 0——**32 路冲突**，要排 32 轮。slides 第 33 页的修正只改了一个数：把数组声明成 `[THREADS_PER_BLOCK_X][THREADS_PER_BLOCK_Y + 1]`，多开一列不用。现在线程 tx 访问第 tx × 33 号字，bank = tx × 33 mod 32 = tx，32 个线程正好错开到 32 个 bank。多花 128 字节的共享内存，换掉了 32 路冲突（slides 第 34 页画出了加一列之后的新排布）。

规律和全局内存的合并访存是对应的：**线程束里相邻的线程访问相邻的字，两种内存都最高兴**。不同的是，全局内存怕"跨太多 sector"，共享内存怕"挤在同一个 bank"。

### 3.5 放到真实尺寸上

03 号脚本第 4 部分把模拟出的事务数等比放大到 8192 × 8192，按 A100 的 1,935 GB/s 估算（不计缓存、不计 bank 冲突）：

| 版本 | 搬运的字节 | 时间 | 有效带宽（读 + 写的数据量 ÷ 时间） |
|---|---|---|---|
| naive | 2,416 MB | 1.249 ms | 430 GB/s |
| 共享内存 | 537 MB | 0.277 ms | 1,935 GB/s |
| 共享内存 + 1 | 537 MB | 0.277 ms | 1,935 GB/s |

转置没有任何运算，数据量是读写各一遍 537 MB；naive 版本要搬 4.5 倍的字节，有效带宽只有峰值的 22%。两个共享内存版本在这个纸面模型里一样快——差别在 bank 冲突，模型没有计入；真实 GPU 上带冲突的版本会因为 32 路排队明显变慢。真实的 L2 缓存会接住 naive 版本的一部分写，所以实测差距通常没有 4.5 倍那么大，但结论不变：**纯搬运的内核，性能完全取决于访存模式**。

---

## 4. 稀疏矩阵 · Sparse matrix multiplication

### 4.1 为什么关心稀疏

很多矩阵大部分元素是 0：图的邻接矩阵、推荐系统的用户–物品矩阵、剪枝后的神经网络权重、MoE 里"每个 token 只用几个专家"（第 17 讲）。把 0 也存下来、也拿来乘，既浪费内存又浪费带宽。稀疏格式只存非零元。

### 4.2 CSR：压缩稀疏行

slides 第 41 页的 CSR（Compressed Sparse Row）用三个数组。04 号脚本第 1 部分：

```
M = [[3, 0, 1, 0],          data      = [3, 1, 2, 4, 1, 1, 1]   非零元的值，按行优先顺序
     [0, 0, 0, 0],          col_index = [0, 2, 1, 2, 3, 0, 3]   每个非零元在第几列
     [0, 2, 4, 1],          row_ptr   = [0, 2, 2, 5, 7]         第 r 行的非零元是 data[row_ptr[r] : row_ptr[r+1]]
     [1, 0, 0, 1]]
```

row_ptr 比行数多一个元素。第 1 行是空行，row_ptr[1] = row_ptr[2] = 2，内层循环一次也不执行——空行不需要特殊处理。

### 4.3 SpMV：串行与 CUDA

slides 第 42–43 页的稀疏矩阵-向量乘法 y = M · x：

```c
// 串行
for (int row = 0; row < n; row++) {
    float dot = 0;
    for (int el = row_ptr[row]; el < row_ptr[row + 1]; el++)
        dot += x[col_index[el]] * data[el];
    y[row] += dot;
}
// CUDA：一个线程负责一行
__global__ void SpMVCSRKernel(float *data, int *col_index, int *row_ptr, float *x, float *y, int num_rows) {
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < num_rows) { /* 同上的内层循环 */ }
}
```

x = [1, 2, 3, 4] 时，串行循环、模拟器上的内核、稠密乘法都得到 [6, 0, 20, 5]。注意 `x[col_index[el]]`：x 是按 col_index **间接**取的（gather）。这是稀疏计算和稠密计算最大的不同——访问哪里由数据决定，编译器和硬件都没法提前知道，合并访存和缓存都难以保证。

### 4.4 多稀疏才划算

04 号脚本第 2 部分（n = 4096，A100 纸面估算）：

| 非零元比例 | CSR / 稠密的存储 | 稠密 matvec | CSR SpMV（x 全命中缓存） | CSR SpMV（每次 gather 落空，多搬一个 sector） |
|---|---|---|---|---|
| 0.1% | 0% | 34.7 μs | 0.1 μs | 0.4 μs |
| 1% | 2% | 34.7 μs | 0.7 μs | 3.5 μs |
| 10% | 20% | 34.7 μs | 7.0 μs | 34.7 μs |
| 25% | 50% | 34.7 μs | 17.4 μs | 86.7 μs |
| 50% | 100% | 34.7 μs | 34.7 μs | 173.4 μs |

CSR 每个非零元要存值和列号两样（各 4 字节），所以非零元比例超过 **50%** 时比稠密存储还大。SpMV 和稠密 matvec 都是访存受限的（每个矩阵元素只用一次），时间约等于搬运的字节 ÷ 带宽：x 能全部留在缓存里时，分界点在 50% 左右；gather 次次落空时，分界点降到 **10%** 左右。**稀疏要足够稀疏才值得**：图和推荐矩阵常常在 0.1% 以下，而剪枝到 50% 的神经网络权重用 CSR 反而更慢。

### 4.5 一个线程一行，还是一个线程束一行？

slides 的内核"一个线程一行"有两个问题：相邻线程同时读各自那一行的第 k 个非零元，位置相隔一整行——**不合并**；各行长度不同时，线程束要等它里面**最长**的那一行——**分化**（第 2 讲 §4.5）。PMPP 第 14 章的改进是"一个线程束一行"：32 个线程读同一行连续的 32 个非零元，最后在共享内存里树形归约。04 号脚本第 3 部分（256 行 × 1024 列）：

| 矩阵 | 内核 | 读 data + col_index 的事务 | 全部读事务 | 线程束发射次数 | SIMT 效率 |
|---|---|---|---|---|---|
| 每行 8 个 | 一个线程一行 | 4,096 | 5,741 | 168 | 100% |
| 每行 8 个 | 一个线程束一行 | 512 | 3,268 | 11,904 | 72% |
| 每行 64 个 | 一个线程一行 | 32,768 | 41,648 | 1,064 | 100% |
| 每行 64 个 | 一个线程束一行 | 4,096 | 18,131 | 12,416 | 77% |
| 多数 2 个、少数 200 个 | 一个线程一行 | 3,468 | 5,526 | 1,656 | 10% |
| 多数 2 个、少数 200 个 | 一个线程束一行 | 966 | 3,095 | 12,000 | 72% |

结果不是"一个线程束一行更好"那么简单：

- **访存**：一个线程束一行读矩阵本身的事务少 8 倍（每行 64 个时 32,768 vs 4,096）；算上 x 的 gather（两种写法都是随机访问），全部读事务少 2.3 倍。
- **指令**：一个线程束一行的发射次数多一个数量级——每一行都要一个线程束、都要做一次 5 轮的树形归约；行短时大部分车道闲着（每行 8 个只用 8 个车道读数）。
- **分化**：行长不均匀时，一个线程一行的 SIMT 效率只有 10%。

SpMV 是访存受限的，真 GPU 上通常是访存那一项说了算，所以长行一般用一个线程束一行；行很短时，事务只省了 1.8 倍却要多付几十倍的指令，常用"几个线程一行"的折中。实际的库（cuSPARSE）会按行长分组、对不同的组用不同的内核，或者换别的格式（ELL、COO、混合格式）。**没有一种分工在所有矩阵上都最好**——这句话在第 3 讲的归约里已经出现过一次。

### 4.6 2:4 结构化稀疏：规格表上"含稀疏"的数字

第 2 讲读规格表时说，张量核心的峰值默认"含 2:4 结构化稀疏"。04 号脚本第 4 部分演示了它：**每 4 个连续的权重里只保留绝对值最大的 2 个**，其余置 0；压缩存储时只存这 2 个值，外加每个值在组内的位置（2 位）。用压缩形式算 W·x 和用剪枝后的稠密矩阵算，最大差 2.4e-07。FP16 权重压缩后每个原始权重占 9 位，是稠密的 56.2%。

2:4 稀疏的非零元比例固定是 50%——按 §4.4 的算账，CSR 在这个比例下毫无优势。但它的模式是**固定的**：位置只要 2 位就能记下，硬件（Ampere 起的张量核心）可以直接跳过一半的乘法，同样的时间做 2 倍的有效运算。代价是模型要按这个模式剪枝（通常还要微调恢复精度），而且只对权重适用。

**稀疏的两条路**：非结构化稀疏（CSR 等）适合极稀疏的数据，靠省存储和省访存取胜；结构化稀疏（2:4、块稀疏）适合神经网络，靠和硬件配合取胜。

---

## 5. cuBLAS · CUDA Basic Linear Algebra Subroutines

### 5.1 为什么用库

slides 第 45 页：cuBLAS 是 CUDA 的基础线性代数子程序库，一个专门做通用矩阵乘法（GEMM）等运算的轻量库。§1–§3 把朴素矩阵乘法从 2.5% 提到了约 80%（FP32）；要用上张量核心、在每一代 GPU、每一种数据类型、每一种矩阵形状上都接近峰值，需要的东西远不止分块（§5.5）。实践中，**矩阵乘法几乎总是调用库**；自己写 GEMM 是为了理解它。

### 5.2 接口

slides 第 46–48 页：

```c
cublasStatus_t cublasCreate(cublasHandle_t *handle);         // 使用前必须调用：创建库的上下文
cublasStatus_t cublasDestroy(cublasHandle_t handle);         // 使用后必须调用
cublasStatus_t cublasSdot(handle, int n, const float *x, int incx, const float *y, int incy, float *result);
cublasStatus_t cublasSgemv(handle, cublasOperation_t trans, int m, int n, const float *alpha,
                           const float *A, int lda, const float *x, int incx, const float *beta, float *y, int incy);
                           // y = α · op(A) · x + β · y
cublasStatus_t cublasSgemm(handle, cublasOperation_t transa, cublasOperation_t transb, int m, int n, int k,
                           const float *alpha, const float *A, int lda, const float *B, int ldb,
                           const float *beta, float *C, int ldc);
                           // C = α · op(A) · op(B) + β · C
```

函数名里的 S 表示单精度（float），D 双精度，H 半精度；`op(X)` 是 X（`CUBLAS_OP_N`）或 Xᵀ（`CUBLAS_OP_T`）。`handle` 保存库的状态（使用的流、工作空间等），创建一次、反复使用。

### 5.3 列优先与 leading dimension

BLAS 诞生于 Fortran 的年代，Fortran 的数组是**列优先**的，所以 cuBLAS 也按列优先解释矩阵。05 号脚本第 1 部分：缓冲区 [1, 2, 3, 4, 5, 6] 按行优先读是 [[1,2,3],[4,5,6]]，按列优先读成 3 × 2 是 [[1,4],[2,5],[3,6]]——**同一段内存，行优先看是 X，列优先看就是 Xᵀ**。

**leading dimension**（lda、ldb、ldc）是"相邻两列的起点相隔几个元素"。整块矩阵时 ld 等于行数；对一个大矩阵的**子块**运算时，指针指向子块的起点，ld 给大矩阵的行数——不用拷贝。例如 5 × 4 列优先矩阵从 (1, 1) 开始的 3 × 2 子块：指针偏移 1 + 1 × 5 = 6，ld = 5。这和 P1 §5 的步长是同一个思想（列优先时行步长 1、列步长 ld）；incx、incy 则是向量的步长。BLAS 的接口从一开始就是"指针 + 步长"。

05 号脚本第 2 部分用 NumPy 按 BLAS 的语义实现了 sgemm / sgemv / sdot：四种 transa / transb 组合、子块（偏移 13、lda = 6）、incx = incy = 2 都与 NumPy 一致（最大差 ≤ 2.4e-07）；lda 写得比行数小时报参数错误。传错一个整数，库只会报"第几个参数非法"，或者更糟——静默地算错。

### 5.4 行优先张量怎样调用列优先的库

我们的张量（NumPy、PyTorch、miniTorch）都是行优先的。想算 C = A · B，又不想做任何转置拷贝，怎么办？05 号脚本第 3 部分：

- 行优先存储的 A（M × K）在列优先的库眼里是 Aᵀ（K × M），B 是 Bᵀ；我们想要行优先的 C，在库眼里就是 Cᵀ。
- 而 Cᵀ = (A·B)ᵀ = **Bᵀ · Aᵀ**。

所以只要**把 A、B 的位置交换、m 和 n 交换**：`sgemm('N', 'N', m = N, n = M, k = K, B, ldb = N, A, lda = K, C, ldc = N)`，结果缓冲区按行优先读出来就是 A · B（与 NumPy 差 4.6e-08）；不交换直接调用则错（差 2.98）。PyTorch 在 GPU 上调用 cuBLAS 时就是这样处理行优先张量的；张量是转置视图时，改用 'T' 参数，同样不拷贝。

### 5.5 从朴素到库：roofline 阶梯

05 号脚本第 4 部分把本讲的各种写法放在同一张表里（A100 80GB PCIe，FP32 数据）：

| 写法 | 计算–全局访存比 | FP32 峰值上限 | TF32 峰值上限 |
|---|---|---|---|
| 朴素 | 0.25 | 2.5% | 0.3% |
| 共享内存分块 TILE = 16 | 4 | 39.7% | 5.0% |
| 共享内存分块 TILE = 32 | 8 | 79.4% | 9.9% |
| 块分块 64 × 64 + 寄存器分块 | 16 | 100% | 19.8% |
| 块分块 128 × 128 + 寄存器分块 | 32 | 100% | 39.7% |
| 块分块 256 × 128 + 寄存器分块 | 42.7 | 100% | 52.9% |

块分块的比值：一个块算 BM × BN 的输出，每一步读 A 的 BM × BK 和 B 的 BK × BN，做 2·BM·BN·BK 次运算，比值 = 2·BM·BN / ((BM + BN) × 4 字节)，和 BK 无关——块越大、越接近正方形，比值越高。每个线程用寄存器算一个 8 × 8 之类的小块，块就能突破"每块 1024 个线程"的限制（§2.6）。到 128 × 128，FP32 已经计算受限。

但张量核心的平衡点是 81：同样的分块只能用到约 40%–53%。要把张量核心喂饱，cuBLAS（以及 CUTLASS、Triton 生成的内核）还要做很多事：用张量核心的矩阵指令；为每一代 GPU、每种数据类型、每类矩阵形状挑选分块大小（第 2 讲的波次量化也在考虑之内）；用双缓冲和异步拷贝（Hopper 上是 TMA，第 3 讲 §1.6）让搬运和计算重叠；让 L2 缓存在块之间复用数据；把偏置、激活函数融合进结果的写回（cuBLASLt 的 epilogue，第 10–11 讲的算子融合）。这就是 slides 总结里那句"现成可用的向量、矩阵-向量、矩阵-矩阵运算"背后的工作量。

---

## 6. 小结 · Summary

1. **计算–全局内存访问比**决定访存受限的内核能跑多快。朴素矩阵乘法 0.25 FLOP/字节，在 A100 上只能用到 FP32 峰值的 2.48%。
2. **全局内存约 500 个周期，共享内存约 5 个，寄存器约 1 个**。局部数组会落到"本地内存"（其实在显存里）。
3. **分块**：线程块把 A、B 的小块搬进共享内存，块内反复使用，全局读取降到 1/TILE、比值升到 0.25 × TILE；结果与朴素版逐位相同。两个 `__syncthreads()` 分别防"读之前没写完"和"写之前没读完"；边界要向上取整、补 0、不能提前 return；块不能太小（合并失效）也不能超过 1024 个线程。
4. **线程粗化 / 寄存器分块**：让每个线程算多个输出，再省一层读取，突破线程数的限制。
5. **合并访存**：一个线程束在同一条指令里访问连续、对齐的地址。转置的读和写总有一边按列走，用共享内存把它挪到片上。
6. **bank 冲突**：共享内存 32 个 bank，第 k 个字在 k mod 32 号；同一 bank 的不同字要排队。按列访问 [32][32] 是 32 路冲突，声明成 [32][33] 即可消除。
7. **稀疏**：CSR 用 data / col_index / row_ptr；比例超过约 50%（gather 落空时约 10%）就不划算。一个线程一行 vs 一个线程束一行是访存和指令的权衡。2:4 结构化稀疏靠固定模式和张量核心配合。
8. **cuBLAS**：列优先、leading dimension、op(X)；行优先张量用"交换 A、B"的技巧直接调用。生产中的 GEMM 几乎总是用库。

## 7. 常见误区 · Common pitfalls

- **"分块减少了计算量。"** 运算总量一次没少，减少的是全局内存的读取次数；共享内存的读取和原来的全局读取一样多。
- **"TILE 越大越好。"** 一个线程一个输出时 TILE 不能超过 32（每块 1024 个线程）；共享内存、寄存器用多了占用率会掉；TILE 太小（例如 2）反而让合并访存失效。
- **"一个 `__syncthreads()` 就够了。"** 搬完之后和算完之后各需要一个，分别防两种竞争；漏掉任何一个，结果都会错，而且在真 GPU 上可能时对时错。
- **"越界的线程直接 return 就好。"** 在用栅栏的内核里，越界的线程也要参与搬运（写 0）和栅栏。
- **"共享内存很快，怎么访问都行。"** 同一个 bank 的不同字要排队；按列访问 32 × 32 的数组是 32 路冲突。
- **"合并访存只看一个线程依次访问的地址是否连续。"** 看的是一个线程束在**同一条指令**里一起访问的地址（第 2 讲 §4.8）。
- **"稀疏总比稠密快。"** 非零元比例要低于约 50%（gather 不命中时约 10%）才划算；50% 稀疏的权重要靠 2:4 这样的结构化稀疏和专门的硬件。
- **"cuBLAS 和我的张量一样是行优先的。"** cuBLAS 是列优先的；直接把行优先缓冲区传进去会算错，要用交换技巧或 transa / transb。
- **"自己写的分块 GEMM 能比 cuBLAS 快。"** 在一般形状上几乎不可能；库用了张量核心、异步流水线、按形状调优等一整套技术。

## 8. 练习 · Exercises

1. 向量加法 C = A + B 和 AXPY（y = a·x + y）的计算–全局内存访问比各是多少（float32）？
2. 在 H100 SXM（3.35 TB/s、FP32 67 TFLOPS）上，朴素矩阵乘法最多能达到 FP32 峰值的多少？
3. N = 1024 的矩阵乘法，朴素版和 TILE = 16 的分块版各读全局内存多少次（按线程计）？
4. TILE = 32 时 As、Bs 合计多少字节共享内存？如果输入改成 FP16 呢？
5. 用自己的话解释分块矩阵乘法里两个 `__syncthreads()` 各防的是什么竞争。
6. N = 100、TILE = 16 时，正确的相数是多少？slides 的 `ph < N / TILE_WIDTH` 会做几相、漏掉几项？
7. 一个线程束的线程 tx 读 `A[2 * tx]`（float，A 从对齐的地址开始），产生几个 32 字节事务？读 `A[tx + 3]` 呢？
8. `__shared__ float s[32][32];` 一个线程束（threadIdx.x = 0..31）读 `s[threadIdx.x][0]` 几路冲突？读 `s[0][threadIdx.x]` 呢？如果声明成 `s[32][33]`，读 `s[threadIdx.x][0]` 呢？
9. `__shared__ float s[64];` 线程 tx（0..31）读 `s[2 * tx]`，有 bank 冲突吗？几路？
10. 写出矩阵 [[0, 5, 0], [7, 0, 0], [0, 0, 0], [1, 2, 3]] 的 CSR 三个数组。
11. 如果值用 8 字节的 double、列号用 4 字节的 int，CSR 在非零元比例多少以下才比稠密存储省空间？
12. A 是行优先的 M × K，B 是行优先的 N × K，想得到行优先的 C = A · Bᵀ（M × N）。sgemm 的参数应该怎样填？

### 参考答案

1. 向量加法：1 次运算，读 8 字节 + 写 4 字节 → **1/12 ≈ 0.083 FLOP/字节**；AXPY：1 次乘 + 1 次加 = 2 次运算，读 x、y 写 y 共 12 字节 → **1/6 ≈ 0.17**。两者都严重访存受限。
2. 3.35×10¹² × 0.25 = 837.5 GFLOPS，占 67 TFLOPS 的 **1.25%**（H100 的算力涨得比带宽快，比例比 A100 更低）。
3. 朴素：2N³ = **2,147,483,648** 次；TILE = 16：2N³ / 16 = **134,217,728** 次。
4. 2 × 32 × 32 × 4 = **8 KB**；FP16（2 字节）时 **4 KB**。
5. 第一个（搬完之后）：保证块内所有线程都把自己负责的元素写进共享内存之后，才有人开始读——防"写后读"。第二个（算完之后）：保证所有线程都用完这一相的小块之后，才有人开始写下一相——防"读后写"。
6. 正确的相数 ⌈100/16⌉ = **7**；slides 的写法做 100 // 16 = **6** 相，漏掉最后 100 − 96 = **4** 项（而且网格向上取整时会越界）。
7. `A[2·tx]`：32 个 float 跨 256 字节 → **8 个**事务（有用 50%）。`A[tx + 3]`：起点偏移 12 字节，128 字节跨 **5 个** sector。
8. `s[tx][0]`：字号 tx × 32，全在 bank 0 → **32 路冲突**。`s[0][tx]`：连续的 32 个字 → **无冲突**。`s[32][33]` 时 `s[tx][0]` 的字号 tx × 33，bank = tx → **无冲突**。
9. 字号 0, 2, 4, …, 62 → bank 0, 2, …, 30, 0, 2, …, 30：每个偶数 bank 被两个不同的字占用 → **2 路冲突**。
10. data = **[5, 7, 1, 2, 3]**，col_index = **[1, 0, 0, 1, 2]**，row_ptr = **[0, 1, 2, 2, 5]**（第 2 行是空行）。
11. 每个非零元 8 + 4 = 12 字节，稠密每个元素 8 字节：12·d < 8 → **d < 2/3 ≈ 66.7%**（忽略 row_ptr）。
12. 想要的 Cᵀ = (A·Bᵀ)ᵀ = B·Aᵀ。库眼里 A 的缓冲区是 Aᵀ（K × M，op = 'N'），B 的缓冲区是 Bᵀ（K × N），而我们需要 B（N × K），所以 op = 'T'：**sgemm('T', 'N', m = N, n = M, k = K, B, ldb = K, A, lda = K, C, ldc = N)**。（用 05 号脚本的 sgemm 验证过，最大差 1.2e-07。）

## 9. 延伸阅读 · Further reading

- **PMPP 第 4 版**第 5 章（存储类型、分块、边界检查、存储用量对占用率的影响）、第 6 章（合并访存、隐藏延迟、线程粗化、优化清单）、第 14 章（COO、CSR、ELL、JDS 等稀疏格式与它们的 SpMV 内核）。
- **NVIDIA CUDA C++ Best Practices Guide**："Memory Optimizations"（合并访存、共享内存与 bank 冲突、矩阵转置的例子）。
- **Mark Harris, *An Efficient Matrix Transpose in CUDA C/C++***（NVIDIA 技术博客）：转置三个版本的实测带宽。
- **Simon Boehm, *How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance***（博客）：从朴素到接近 cuBLAS 的十步，每一步都有实测数字，是 §5.5 阶梯的实践版。
- **NVIDIA CUTLASS** 文档：GEMM 的分层分块（线程块 → 线程束 → 指令）在模板库里的实现。
- **Mishra et al. (2021) *Accelerating Sparse Deep Neural Networks***（NVIDIA）：2:4 结构化稀疏的硬件支持与"剪枝—微调"的流程。

## 10. 下一讲预告 · Next lecture

前四讲都在"底层"：GPU 怎样执行、内核怎样写快。第 5 讲往上走一层：**深度学习框架与自动微分**。`loss.backward()` 是怎样对任意网络算出每个参数的梯度的？slides 从通用的迭代学习算法（梯度下降）出发，讲计算图、拓扑排序、前向与反向模式的自动微分、向量–雅可比积，再到一个框架的组成（TensorFlow 式的占位符、变量、运算节点、会话）。我们会自己写一个迷你框架：标量和张量两层的自动微分（原课 HW2 的 `topological_sort` 和 `backpropagate`）、静态图和动态图两种风格，并把第 3 讲的 SIMT 内核接成它的一个"后端"，最后用它训练 P3 的影评分类网络。
