第 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/（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

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：内核写成 Python 函数，模拟器按 CUDA 的规则为网格里的每个线程调用它，并且把每一次显存访问都记下来。它比真 GPU 慢上百万倍，但它能数——这正是学 GPU 编程最需要的：

它统计的东西都是”为什么这样写更快”的答案：每个线程束的每条访存指令产生几个 32 字节的内存事务（合并访存）、共享内存有没有 bank 冲突、线程束里有多少线程在空等（分化）、有没有数据竞争。它不模拟时间——GPU 上的耗时要用 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 的影评小语料（ReviewWorld）真的跑了一遍，并把每一步都改写成对底层运算的调用：

一次前向一共 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 估算每个运算的理想耗时（假设每个运算单独执行、输入输出都经过显存）：

按运算类型汇总：

这张表就是第 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 的计算量少了 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。

对照：每张卡这一步的计算量 $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，照抄）：

读这张表有两个要点。

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

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

• 张量核心比普通 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 部分：

“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 部分把它补全：

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 的资源列了出来：

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 部分的向量加法：

__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 记了下来：

另外 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 字节内存事务个数。）

块大小 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 效率 = 线程实际执行的指令数 ÷ (发射次数 × 32)。）

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

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

线程束的循环要转到它里面次数最多的那个线程做完为止。把工作量相近的线程放进同一个线程束，总工作量不变，效率从 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 是简化版的占用率计算器）：

块太小会被”每个 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 个块、每个块耗时相同：

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）：每个线程束反复做”读 128 字节 → c 条依赖这次数据的计算指令”，访存延迟 500 个周期；调度器每个周期最多发射一条指令。另外加上带宽的限制：H100 的 3.35 TB/s 分到每个 SM 的每个调度器，约 3.2 字节 / 周期，即每 40 个周期才能读一次 128 字节。结果（调度器发射了指令的周期占比；括号里是公式 $\min(1,\ W(1+c)/(L+c),\ (1+c)/40)$ 的值）：

• 只有 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 部分在模拟器上犯了三个最常见的错误：

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

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

两种分法读的元素完全一样，用到的不同 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 都要运行它：

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 逐位对比。