第 3 讲 · GPU 架构与 CUDA 编程（一）· GPU Architecture & CUDA Programming (1)

对应 Maps to：CMU 15-779（Fall 2025，Zhihao Jia），Lecture 3 · Fri 9/5 · GPU Architecture and CUDA Programming (1)。 阅读材料：原课这一讲没有指定必读论文。推荐：NVIDIA CUDA C++ Programming Guide 第 1–5 章（编程模型、硬件实现）；NVIDIA GTX 980（Maxwell）与 H100（Hopper）架构白皮书。 依据说明：原课第 3、4 讲共用一份 slides（03-CUDA-programming.pdf，69 页，“Lecture 3&4: GPU Architecture & CUDA Programming”）。本讲对应其前半部分（第 1–47 页）：SIMD 与大规模并行的计算单元；CUDA 语言与线程层次；grid / block / thread；host 与 device 代码的分离；显式的线程数与越界保护；条件执行与掩码、一致执行与发散执行；CUDA 内存模型与 cudaMemcpy；设备内存的三种类型；一维卷积的两个版本；同步原语；CUDA 编译；线程块调度与 GPU 工作调度器；SMM；warp；GTX 980 到 H100 的变化；Tensor Core。后半部分（矩阵乘法与并行归约两个案例）留给第 4 讲。gpusim.py 模拟器和全部实验是本讲义自己设计的，不代表课堂上讲了什么。 配套代码 Code：code/lecture03/（5 个脚本 + 共享模块 gpusim.py；只需 numpy；每个脚本 ≤ 5 秒） 预计阅读时间：3–4 小时（含运行代码）

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

0.1 学习目标

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

1. 解释 GPU 为什么采用 SIMD / SIMT 设计，以及这种设计在”控制逻辑”和”计算单元”之间做了什么取舍。
2. 读懂并写出简单的 CUDA 内核：__global__、<<<grid, block>>>、blockIdx / blockDim / threadIdx、越界保护。
3. 解释分支发散的成因与代价，知道怎样通过重排数据或任务让 warp 保持一致执行。
4. 说出 CUDA 的内存模型：主机与设备两个地址空间；设备上的寄存器（线程私有）、共享内存（块内共享）、全局内存（所有线程）；知道共享内存为什么存在、__syncthreads() 为什么必须有。
5. 描述一个内核从启动到执行完的过程：编译出的资源需求 → 工作调度器按资源把块分配到 SM → SM 以 warp 为单位发射指令；能计算占用率，解释延迟隐藏和波次 / 尾部效应。
6. 算出 GTX 980 和 H100 的峰值算力，说明 Tensor Core 为什么让峰值涨了一个数量级、又为什么需要 FP32 累加。

0.2 先修与本讲范围

0.3 本讲地图

§1  为什么 GPU 长这样：SIMD 与"把晶体管花在计算上"
§2  CUDA 编程模型：线程层次、host / device、越界保护                ← 代码 01（演示 1–2）
§3  条件执行：掩码、一致执行与发散执行                                ← 代码 02
§4  CUDA 内存模型：两个地址空间、三种设备内存                          ← 代码 01（演示 3）
§5  案例：一维卷积，从全局内存到共享内存                              ← 代码 03
§6  同步原语：__syncthreads、原子操作、内核边界
§7  一个内核怎样跑起来：编译、调度、SM、占用率、延迟隐藏、波次          ← 代码 04
§8  warp：块的形状与车道利用率                                        ← 代码 01（演示 4）
§9  从 GTX 980 到 H100：Tensor Core                                   ← 代码 05
§10 小结   §11 常见误区   §12 练习   §13 延伸阅读   §14 下一讲预告

0.4 本讲的模拟器 gpusim.py

没有 NVIDIA GPU 也能学 CUDA：gpusim.py 用 NumPy 执行 CUDA 风格的内核，并把性能相关的事件数出来。它采用 slides 第 2–3 页的 SIMD 视角——一个线程块里的所有线程一起执行每一条语句——所以 t.threadIdx.x 是一个数组（块内每个线程一个元素），i = t.blockIdx.x * t.blockDim.x + t.threadIdx.x 一次算出块内每个线程的全局下标。CUDA 与模拟器的对照：

它统计：warp 指令的发射次数与 SIMD 效率、分支发散次数、全局内存的请求 / 32 字节事务 / 访问过的不同段数、共享内存的 bank 冲突、__syncthreads 次数、共享内存的数据竞争。它和真 GPU 的主要差别：块内线程严格同步执行（所以专门做了竞争检测）；不模拟缓存（同时报告”无缓存”和”理想缓存”两种估算）。

1. 为什么 GPU 长这样：SIMD · Why GPUs look the way they do

1.1 一个指令，多份数据

slides 从一张对比图开始：传统的处理器核是”单指令、单数据”（SISD）——一套取指 / 译码逻辑，驱动一个 ALU；现代处理器核是”单指令、多数据”（SIMD）——同一套取指 / 译码逻辑，驱动很多个 ALU，同一条指令广播给所有 ALU，每个 ALU 处理自己的那份数据。要提高算力，只需要加 ALU，而不需要复制昂贵的控制逻辑。

为什么这对 GPU 是对的选择？ 一个处理器芯片的晶体管可以花在两类地方：

• 让单个线程跑得快：大缓存、分支预测、乱序执行……这些是 CPU 的做法，目标是降低单个任务的延迟；
• 让很多线程一起跑：更多的 ALU，目标是提高整体的吞吐。

图形渲染和深度学习的计算有一个共同特点：同样的操作施加在海量的数据上（每个像素、每个矩阵元素做同样的运算）。对这类负载，控制逻辑是可以被大量数据共享的，于是 GPU 把绝大多数晶体管花在 ALU 上，并且把它们组织成 SIMD 的形式。slides 第 4 页那张”海量并行计算单元”的图就是结果：上万个 ALU，每 32 个共享一套指令流。

1.2 代价：SIMD 要求”大家做同样的事”

SIMD 的效率建立在一个假设上：同一组 ALU 上的数据要执行同样的指令。一旦它们要做不同的事（§3 的分支发散），或者要的数据不在手边（§7 的访存延迟），这套设计就会浪费。GPU 编程的大部分技巧，都是在维护这个假设。

2. CUDA 编程模型 · The CUDA programming model

2.1 CUDA 是什么

slides 介绍：CUDA 于 2007 年随 NVIDIA Tesla 架构推出，是一种”类 C”的 GPU 编程语言。它的设计目标是保持很低的”抽象距离”——CUDA 的抽象紧贴现代 GPU 的能力和性能特征。好处是程序员能精确控制硬件；代价是程序员必须理解硬件（这一讲和下一讲的全部内容），而且硬件变了，程序往往也要改（第 1 讲 §4.1）。

2.2 线程层次：grid、block、thread

一个 CUDA 程序由线程的层次结构组成（slides 第 7–9 页）：

• 一次内核启动产生一个网格（grid），网格由若干线程块（block）组成，块由若干线程（thread）组成；
• 网格和块都可以是 1–3 维的，便于对应向量、矩阵、体数据；
• 每个线程通过四个内置变量知道自己是谁：gridDim（网格有多大）、blockIdx（我的块在网格里的位置）、blockDim（块有多大）、threadIdx（我在块里的位置）。

slides 的例子：

const int Nx = 12, Ny = 6;
dim3 threadsPerBlock(4, 3, 1);
dim3 numBlocks(Nx / threadsPerBlock.x, Ny / threadsPerBlock.y, 1);   // (3, 2, 1)
matrixAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);                  // 6 个块 × 12 个线程 = 72 个线程

__global__ void matrixAdd(float A[Ny][Nx], float B[Ny][Nx], float C[Ny][Nx]) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    C[j][i] = A[j][i] + B[j][i];
}

代码 01 的演示 1 在模拟器上跑了这个内核，并画出了每个元素由哪个块计算：

b0 b0 b0 b0 b1 b1 b1 b1 b2 b2 b2 b2
b0 b0 b0 b0 b1 b1 b1 b1 b2 b2 b2 b2
b0 b0 b0 b0 b1 b1 b1 b1 b2 b2 b2 b2
b3 b3 b3 b3 b4 b4 b4 b4 b5 b5 b5 b5
...

每个块负责一个 4×3 的小矩形；块内线程按 x 最快、再 y 的顺序编号（块 4 里 t0–t3 是第一行、t4–t7 是第二行）。这个编号顺序很重要：它决定了线程怎样被分进 warp（§8），进而决定访存能否合并（第 4 讲）。

为什么要有”块”这一层？ 如果只有一个扁平的线程数组，硬件就无法知道哪些线程需要互相协作。块是协作的单位：同一个块的线程被放在同一个 SM 上，可以通过共享内存交换数据、用 __syncthreads() 同步；不同块之间不能直接协作（§6）。这让硬件可以自由地把块分到任意 SM 上、按任意顺序执行（§7）——同一个程序因此能在 6 个核和 16 个核的 GPU 上都跑起来（slides 第 26 页）。

2.3 host 代码与 device 代码

slides 强调 CUDA 程序被静态地分成两部分：

• host 代码：在 CPU 上串行执行的普通 C/C++，负责分配内存、拷贝数据、启动内核；
• device 代码：在 GPU 上以 SIMD 方式执行。__global__ 标记内核（从 host 启动、在 device 上运行），__device__ 标记只能在 device 上调用的函数（例如 slides 里的 doubleValue）。

kernel<<<numBlocks, threadsPerBlock>>>(args) 是”批量启动”：一次性创建整个网格的线程。slides 原文的意思是：调用在所有线程结束后才返回（从 host 的角度看是这样；实际上 CUDA 的启动是异步的，host 要在需要结果时同步——第 2 讲 PyTorch 论文的”控制流与数据流分离”正是利用了这种异步）。

2.4 线程数是程序员显式给的：越界保护

slides 第 12 页指出：线程的数量由启动配置决定，而不是由数据大小决定。当 Nx = 11, Ny = 5（不是 4 和 3 的整数倍）而网格仍是 (3, 2) 时，仍然会启动 72 个线程，其中一部分落在数组之外。所以内核要写越界保护：

if (i < Nx && j < Ny) C[j][i] = A[j][i] + B[j][i];

代码 01 演示 2：去掉保护，模拟器立刻报告”线程 8（块 (0,1)）访问 A[55] 越界（大小 55）“；加上保护，结果正确，72 个线程里有 17 个空转（24%），并且有 4 个 warp 在这个 if 上发生了发散（§3）。在真实规模（$N = 10^6$）下，浪费的最多是最后一个块，可以忽略；但越界 bug 是 CUDA 里最常见的错误之一：真 GPU 上越界不一定报错，可能悄悄读写到别的数据。

3. 条件执行：掩码与发散 · Conditional execution & divergence

3.1 掩码：丢弃 ALU 的输出

slides 第 13–16 页用一个内核说明 SIMD 怎样处理 if：

float x = A[i];
if (x > 0) { x = 2.0f * x; }          // 便宜的分支
else       { x = exp(x, 5.0f); }      // 昂贵的分支
A[i] = x;

一个 warp 的 32 个线程共用一条指令流，但它们的 x 不同，有的要走 then、有的要走 else。硬件的做法是：两个分支都执行。执行 then 时，走 else 的线程被屏蔽（mask）——它们的 ALU 照样在运转，但输出被丢弃；执行 else 时反过来。分支结束后，所有线程恢复全速执行。

slides 给出两个术语：

• 一致执行（coherent execution）：同样的指令序列作用于所有元素——高效使用 GPU 的必要条件；
• 发散执行（divergent execution）：缺乏一致性——在 CUDA 程序里应该尽量减少。

3.2 发散的代价有多大（代码 02）

代码 02 把这个内核在 4096 个元素上跑了 6 种数据分布（假设 then 分支 1 条指令、else 分支 16 条——指数函数的真实指令数取决于实现，这里只取量级）：

三个结论：

1. 随机正负时，几乎和”全部为负”一样慢（4.2× vs 4.0×）：每个发散的 warp 都要把两个分支走一遍，便宜分支省下的时间全部浪费了。
2. 同样的数据，换个排列就快 1.7 倍：把数据排好序，让同一个 warp 里的元素同号，发散几乎消失（128 → 1 次）。做的有用功完全一样，差别只在”谁和谁分在一个 warp”。
3. 罕见的昂贵分支代价惊人：只有 5% 的元素走 else，但一个 warp 里”至少有一个线程走 else”的概率是 $1 - 0.95^{32} \approx 79\%$（实测 101/128 个 warp 发散），指令数是全正数的 3.5 倍；如果每个线程能独立执行，平均只需要 1.72 条。SIMD 把 32 个线程的命运捆在了一起。

代码 02 的演示 3 是多路分支的版本：switch (i % 4) 让每个 warp 里四种情况各占 8 个线程，四个分支都要执行，SIMD 效率只有 37.5%；换成 switch ((i / 32) % 4)（按 warp 号选择），每个 warp 只走一个分支，效率 100%。

实践中的对策：把罕见情况挪出热循环单独处理；先按类别重排数据（例如 MoE 里先把 token 按专家排序，第 21 讲）；或者干脆用无分支的写法（两边都算，再用 select 挑一个——当两个分支都很便宜时这往往更快）。

4. CUDA 内存模型 · The CUDA memory model

4.1 两个地址空间

slides 第 18–20 页：host（CPU）和 device（GPU）有各自独立的地址空间。device 代码不能访问 host 内存，host 代码也不能直接访问 device 内存。数据要显式地搬：

float* deviceA;
cudaMalloc(&deviceA, bytes);                                   // 在 device 地址空间分配
cudaMemcpy(deviceA, A, bytes, cudaMemcpyHostToDevice);         // host → device
// 注意：在 host 代码里写 deviceA[i] 是非法的

模拟器里，在主机代码中写 A[3] 会得到 HostAccessError。代码 01 演示 3 还算了一笔账：100 万个 float 的向量加法，拷贝 12.6 MB 走 PCIe 5.0（标称约 64 GB/s 单向）约 197 µs；内核本身在 H100 上受 HBM 带宽限制，约 3.8 µs——拷贝是计算的 52 倍。这就是为什么深度学习框架让张量常驻显存、整个训练循环都在 GPU 上进行。（现代系统有统一内存、零拷贝、GPUDirect 等机制模糊了这条边界，但”跨 PCIe 搬数据很贵”这件事没有变。）

4.2 设备上的三种内存

slides 第 21 页列出 device 内核可用的三种内存：

slides 提的问题是：为什么要有共享内存？ 答案是：让同一个块的线程能够协作——一个线程从全局内存搬进来的数据，块里的其他线程也能用。这是”复用”的基础，第 4 讲的分块矩阵乘法完全建立在它上面。

5. 案例：一维卷积 · Case study: 1D convolution

5.1 两个版本

slides 第 22–24 页的例子：output[i] = (input[i] + input[i+1] + input[i+2]) / 3，$N = 1024 \times 1024$，每块 128 个线程。

版本 1：每个线程直接从全局内存读 3 个数。

int index = blockIdx.x * blockDim.x + threadIdx.x;
float result = 0.0f;
for (int i = 0; i < 3; i++) result += input[index + i];
output[index] = result / 3.f;

共享内存版：块里的线程先合作把这个块需要的 130 个数（128 个 + 右边 2 个”光环”元素）搬进共享内存，同步之后再从共享内存读：

__shared__ float support[THREADS_PER_BLK + 2];
support[threadIdx.x] = input[index];
if (threadIdx.x < 2) support[THREADS_PER_BLK + threadIdx.x] = input[index + THREADS_PER_BLK];
__syncthreads();                                         // 等块里所有线程都搬完
float result = 0.0f;
for (int i = 0; i < 3; i++) result += support[threadIdx.x + i];
output[index] = result / 3.f;

slides 的结论：每块从全局内存读 130 次，而不是 3 × 128 = 384 次。

5.2 数一数（代码 03）

几个值得细看的地方：

• 事务数降了 3.29 倍，而不是 3 倍。 版本 1 读 input[i+1]、input[i+2] 的指令地址没有对齐到 32 字节，一个 warp 读的 128 字节横跨 5 个段，所以三条读指令一共 4 + 5 + 5 = 14 个事务。对齐也会影响访存效率。
• 两个版本访问的”不同段数”完全一样。 版本 1 多出来的读取，全都是重复读邻居已经读过的数据。真 GPU 有 L1 / L2 缓存，会挡住其中很大一部分——所以在这个例子里，共享内存在真 GPU 上的收益介于两种估算之间，不会有 3 倍那么多。
• 那共享内存的价值到底在哪？ 在于可控的复用：程序员明确地决定哪些数据放在片上、被谁用、用多少次，而不是寄希望于缓存的替换策略。复用的数据量越大、访问越不规则，缓存越靠不住。第 4 讲的矩阵乘法里，每个数据要被复用几十上百次，共享内存的作用会大得多。
• 没有 bank 冲突：一个 warp 的 32 个线程读连续的 32 个 float，正好落在 32 个不同的 bank（第 4 讲细讲 bank）。

5.3 忘了 __syncthreads() 会怎样？

代码 03 演示 3 把 __syncthreads() 删掉。在模拟器上数值仍然是对的（它让块内线程严格同步执行），但竞争检测报告了 98,304 次竞争（每块 12 次）：每个 warp 的最后两个线程要读下一个 warp 搬进共享内存的数据，而在真 GPU 上，不同 warp 的执行进度是不确定的——那个 warp 可能还没搬完。这类 bug 的可怕之处在于：它可能在你的测试里一直”恰好正确”，换一块卡、换一个负载就出错。

6. 同步原语 · Synchronization primitives

slides 第 25 页列出 CUDA 的同步手段：

1. __syncthreads()：块内栅栏，等块里所有线程都到达这一点。注意它只能同步一个块内的线程；而且所有线程都必须执行到它——把它放在只有部分线程进入的 if 里是未定义行为，可能死锁（模拟器会报 BarrierError）。
2. 原子操作：例如 atomicAdd(float* addr, float amount)，在全局内存和共享内存上都可以用。多个线程同时加到同一个地址，结果不会丢。代价是同一地址上的原子操作要串行化。
3. host / device 同步：一个内核返回时，所有线程都已结束——内核边界是一个隐式的全局栅栏。

为什么 CUDA 没有块之间的全局同步？第 4 讲的并行归约会详细讨论：一是为成千上万个核构建全局同步硬件很昂贵；二是如果块数超过 GPU 能同时驻留的块数，先运行的块在栅栏处等待还没被调度的块，而后者永远等不到空出来的 SM——死锁。所以需要全局同步时，就把计算拆成多个内核，用”内核边界”当全局栅栏。

7. 一个内核怎样跑起来 · How a kernel runs

7.1 编译：二进制里写着资源需求

slides 第 26–27 页：CUDA 的目标是让同一个程序在不同规模的 GPU 上运行（6 个核的中端卡、16 个核的高端卡）。编译出的设备二进制除了指令，还写明了每个块需要的资源：例如卷积内核每块 128 个线程、每线程 8 字节局部数据、每块 130 个 float（520 字节）共享内存。

7.2 工作调度器：按资源动态分配

slides 第 28 页的两个要点：

• CUDA 的核心假设：线程块之间没有依赖，可以按任意顺序执行；
• GPU 用动态调度策略把块映射到核（SM）上，同时满足资源约束。

slides 第 31–40 页用一个玩具 GPU 一步步演示：两个核，每个核 1.5 KB 共享内存；卷积内核每块需要 520 B，所以每个核最多放 2 个块（3 × 520 B > 1.5 KB）；一共 1000 个块。代码 04 演示 1 复现了这个过程：

t=0.00  开始  块 0  核 0  上下文 0-127
t=0.00  开始  块 1  核 1  上下文 0-127
t=0.00  开始  块 2  核 0  上下文 128-255
t=0.00  开始  块 3  核 1  上下文 128-255
t=0.80  结束  块 0  核 0  上下文 0-127
t=0.80  开始  块 4  核 0  上下文 0-127        ← 块 4 接管块 0 空出来的执行上下文
t=0.90  结束  块 2  核 0  上下文 128-255
t=0.90  开始  块 5  核 0  上下文 128-255

谁先结束谁先接活，块的执行顺序和编号无关——这正是 CUDA 必须要求”块之间不能有依赖”的原因。

7.3 SM：warp 的执行引擎

slides 第 29–30 页以 GTX 980 的 SMM（Maxwell 架构的 SM）为例：

• 资源上限：最多 64 个 warp 的执行上下文（64 × 32 = 2048 个线程）、96 KB 共享内存；
• 4 个 warp 调度器，每个驱动一组 32 路的 SIMD 乘加单元；
• 每个时钟周期：从驻留的 64 个 warp 里选出最多 4 个可运行的 warp（线程级并行），每个 warp 最多发射 2 条可运行的指令（指令级并行）。

一个 128 线程的卷积块，由 4 个 warp 执行。

7.4 占用率：一个 SM 上能驻留多少 warp？

一个 SM 能同时驻留的块数，由最紧张的那种资源决定：warp 槽位、块数上限、共享内存、寄存器。驻留的 warp 数 ÷ 最大 warp 数，叫占用率（occupancy）。代码 04 演示 2：

同一份代码，在共享内存更大的 H100 上占用率从 38% 升到 88%：硬件资源变了，最优的内核参数也会变——这是第 9 讲”自动调优”的动机之一。

7.5 为什么占用率重要：延迟隐藏

GPU 的访存延迟很长（从 HBM 读数据要几百个时钟周期），而且 GPU 没有 CPU 那样的大缓存和乱序执行来降低它。GPU 的办法是同时驻留很多 warp：一个 warp 在等数据时，调度器就去发射别的 warp 的指令。只要就绪的 warp 足够多，发射槽就一直有活干，延迟就被”藏”起来了。

代码 04 演示 3 用一个极简的调度器模拟验证这一点：每个 warp 反复执行”$k$ 条算术指令 + 1 条访存”，访存要等 400 个周期。要让发射槽每个周期都有活干，需要的 warp 数是 $\lceil (k + 1 + L)/(k + 1) \rceil$（一个 warp 每 $k+1+L$ 个周期只占用 $k+1$ 个发射周期）：

每个 SM 有 4 个调度器、最多 64 个 warp，所以每个调度器最多轮换 16 个 warp。访存密集的内核需要 81 个 warp 才能藏住延迟，硬件根本放不下——这类内核注定受访存限制，这和第 1 讲的 roofline 结论一致（也呼应了第 1 讲 §4.4：“单线程喂不满内存带宽，要足够多的并发请求”）。计算密集的内核只要 5 个 warp。

最后一行是一个意外的发现：同样的内核，如果调度器每个周期轮换一个 warp（轮转），所有 warp 会齐头并进、几乎同时发出访存、又同时等待（“车队效应”），8 个 warp 只有 69%；而”贪心后最老”（一直发射同一个 warp 直到它停下）让各个 warp 的访存时刻自然错开，8 个 warp 就达到 100%。真实 GPU 的调度器策略是保密的，但这说明调度策略本身也影响延迟隐藏。

7.6 波次与尾部效应

如果每个块耗时相同，一个 GPU 一次能并行执行的块数（SM 数 × 每 SM 驻留的块数）就是一”波”（wave）。代码 04 演示 4：H100 有 132 个 SM，假设每个 SM 驻留 2 个块，一波 264 个块：

只多 1 个块，就要多跑一整波，而那一波里只有 1 个块在干活——这叫波次量化（wave quantization）。块越多，尾部的浪费被摊得越薄。矩阵乘法库会根据矩阵形状挑选分块大小，部分原因就是为了避免难看的尾部（第 4、9 讲）。

8. warp：块的形状与车道利用率 · Warps

slides 第 41 页：warp 是 NVIDIA GPU 的实现细节——一个线程块里每 32 个线程组成一个 warp，用 32 宽的 SIMD 一起执行。CUDA 的编程模型里没有 warp，但几乎所有性能现象都以它为单位：分支发散看的是 warp 内部是否一致（§3），合并访存看的是一个 warp 的 32 个地址是否连续（第 4 讲），bank 冲突看的是一个 warp 的 32 个共享内存地址落在哪些 bank（第 4 讲）。

块内线程按 x 最快、y 其次、z 最后的顺序编号，然后每 32 个一组切成 warp。代码 01 演示 4 列出了几种块形状的车道利用率：

slides 的 (4, 3) 块是为了画图方便；实际内核的块大小应该是 32 的倍数（常见 128 或 256）。

9. 从 GTX 980 到 H100：Tensor Core · From GTX 980 to H100

9.1 峰值算力怎么算

slides 第 43–44 页：GTX 980（2014）有 16 个 SMM，每个 SMM 4 组 32 路乘加单元，时钟约 1.1 GHz，峰值 4.6 TFLOPs；整卡最多驻留 16 × 64 = 1024 个 warp（32,768 个线程）。到 H100（2022），warp 宽度（32）和每个 SM 最多驻留的 warp 数（64）都没变，SM 从 16 个增加到 132 个，共享内存从 96 KB 增加到 H100 每个 SM 256 KB 的 L1 / 共享内存（其中最多 228 KB 可配成共享内存），峰值涨到约 1000 TFLOPs——slides 说”主要是因为 Tensor Core”。

代码 05 演示 1 把这些数字拆开：

从 GTX 980 到 H100 的 Tensor Core 峰值，215 倍 = SM 个数 ×8.25 × 每 SM 每周期 ×16 × 时钟 ×1.63。注意中间那个 16：只看 CUDA core，每个 SM 每周期的运算数十年没变；H100 的 CUDA core 比 GTX 980 快 15 倍，全靠更多的 SM 和更高的时钟。真正让每个 SM 的算力翻 16 倍的，是 Tensor Core。

9.2 Tensor Core：一条指令做一个小矩阵乘法

slides 第 46–47 页：Tensor Core 是 SM 里的矩阵乘法单元。CUDA core 的一条 FFMA 指令让 32 个线程各做 1 次乘加；Tensor Core 的一条 mma 指令让一个 warp 一起完成一个小矩阵乘加，例如 mma.sync.m16n8k16：$16 \times 8 \times 16 = 2048$ 次乘加。代码 05 演示 2：算一个 $64 \times 64 \times 64$ 的矩阵块（262,144 次乘加），

• CUDA core：8,192 条 warp 级 FFMA 指令；
• Tensor Core mma.m16n8k16：128 条（少 64 倍）；
• Hopper 的 wgmma.m64n64k16（一个 warpgroup = 4 个 warp 一起发射）：4 条。

指令发射槽本身就是有限资源（§7.5 的调度器每周期只能发一条），Tensor Core 把”一条指令做的事”变大，才可能把算力喂满。代价是编程更难：操作数必须按特定的”片段”布局放在寄存器或共享内存里——这是第 6 讲 warp 特化、第 8 讲 Triton 要处理的细节。

9.3 为什么累加器要用 FP32

Tensor Core 的标准用法是低精度输入、FP32 累加。代码 05 演示 3 模拟了 FP16 输入的点积（参考值用 FP64 计算，20 次取平均）：

FP16 只有 10 位尾数：累加和变大以后，再加上一个小的乘积会被舍入掉（“大数吃小数”），误差随 K 迅速增长。矩阵乘法的 K 维往往有几千，所以累加器必须用 FP32。输入可以用低精度（FP16 / BF16 / FP8），因为它们省带宽、省芯片面积、算得快，而单个乘积的相对误差是可控的。

10. 小结 · Summary

1. GPU 把晶体管花在 ALU 上，用 SIMD 让一套控制逻辑驱动很多 ALU；代价是要求同一组线程做同样的事。
2. CUDA 的线程层次：grid → block → thread，内置变量 gridDim / blockIdx / blockDim / threadIdx；块是协作的单位，块之间独立；线程数由程序员显式给出，所以要写越界保护。
3. 分支发散：一个 warp 走了不同的分支，就要把所有分支都执行一遍。随机正负的数据让内核慢 4.2 倍；排好序后快 1.7 倍；5% 的罕见分支让 79% 的 warp 付出全部代价。
4. 内存模型：host / device 两个地址空间，拷贝很贵（PCIe 拷贝是内核本身的 52 倍）；设备上有寄存器、共享内存、全局内存三级。共享内存让块内线程协作复用数据，__syncthreads() 保证可见性，忘了它会产生（可能被测试掩盖的）数据竞争。
5. 一个内核怎样跑：编译时确定资源需求；工作调度器按资源动态分配块；SM 以 warp 为单位发射；占用率由最紧张的资源决定；足够多的驻留 warp 才能藏住访存延迟；块数不是”波”的整数倍时有尾部浪费。
6. Tensor Core 让每个 SM 每周期的运算数翻 16 倍，是 H100 比 GTX 980 快 200 倍的主要原因；它要求低精度输入、FP32 累加和特定的数据布局。

11. 常见误区 · Common misconceptions

1. “CUDA 线程和 CPU 线程一样，可以各干各的。” 同一个 warp 的 32 个线程共享指令流；它们”各干各的”（发散）时，性能按分支数成倍下降。
2. “线程越多越好 / 块越大越好。” 块太大会让一个 SM 只能放下很少的块（寄存器、共享内存不够），也会让尾部更难看；块太小（< 32 或不是 32 的倍数）浪费车道。要看占用率和波次。
3. “占用率越高越快。” 占用率只是藏延迟的手段之一；很多高性能内核故意用低占用率换取更多的寄存器和共享内存（每个线程做更多的工作、指令级并行更高，第 4 讲的寄存器分块）。
4. “在我的测试里没出错，所以没有数据竞争。” §5.3：删掉 __syncthreads() 后结果仍然可能”恰好正确”。
5. “__syncthreads() 能同步所有线程。” 只能同步一个块内的线程；块之间没有栅栏（除了内核边界）。
6. “GPU 有缓存，所以不用管访存模式。” 缓存很小、而且是共享的；可控的复用（共享内存、寄存器分块）和合并访存仍然是性能的关键（第 4 讲）。
7. “峰值 1000 TFLOPs 意味着随便写的 CUDA 代码都能快 200 倍。” 那是 Tensor Core 在低精度下的峰值；不用 Tensor Core 的 FP32 代码，峰值只有 67 TFLOPs，而且还要先解决访存问题。

12. 练习 · Exercises（含参考答案）

练习 1（线程下标）：blockDim = (16, 8)，blockIdx = (2, 3)，threadIdx = (5, 7)。这个线程负责的 $(i, j)$ 是多少？它在块内的线性编号是多少？属于块内第几个 warp？

答：$i = 2 \times 16 + 5 = 37$，$j = 3 \times 8 + 7 = 31$。线性编号 $= 5 + 7 \times 16 = 117$，属于第 $\lfloor 117/32 \rfloor = 3$ 个 warp（从 0 开始数）。

练习 2（网格大小）：要处理 $N = 1000$ 个元素，每块 256 个线程，网格应该开多大？会有多少个线程空转？

答：$\lceil 1000 / 256 \rceil = 4$ 个块，共 1024 个线程，24 个空转（都在最后一个块里）。常见写法：(N + 255) / 256。

练习 3（发散）：一个 warp 里 32 个线程执行 if (threadIdx.x < 16) A(); else B();，A 需要 10 条指令、B 需要 30 条。这个 warp 要花多少条指令的时间？SIMD 效率是多少？如果改成 if (threadIdx.x / 32 % 2 == 0) 呢？

答：两个分支都要执行：$10 + 30 = 40$ 条；有用的线程-指令 = $16 \times 10 + 16 \times 30 = 640$，效率 $640 / (40 \times 32) = 50\%$。改成按 warp 号选择后，同一个 warp 的条件相同，只执行一个分支，效率 100%（但不同的 warp 走的分支不同，总工作量取决于有多少 warp 走 B）。

练习 4（卷积的复用）：把一维卷积改成 7 点（input[i..i+6]），每块 256 个线程。版本 1 和共享内存版每块各读多少次全局内存？

答：版本 1：$7 \times 256 = 1792$ 次；共享内存版：$256 + 6 = 262$ 次，约 6.8 倍。卷积窗口越大，共享内存的收益越大。

练习 5（占用率）：一个内核每块 256 个线程、每线程 64 个寄存器、每块 48 KB 共享内存。在 H100 SM（64 warp、32 块、228 KB 共享内存、65536 寄存器）上能驻留几个块？占用率多少？限制因素是什么？用 gpusim.occupancy 验证。

答：warp 限制 $64 / 8 = 8$；寄存器限制 $65536 / (64 \times 256) = 4$；共享内存限制 $228 / 48 = 4$；块数上限 32。取最小值 4 块、32 个 warp、占用率 50%，寄存器和共享内存同时卡住（occupancy 会报告其中一个）。

练习 6（延迟隐藏）：访存延迟 600 个周期，每次访存之间有 30 条算术指令。一个调度器需要多少个驻留的 warp 才能让发射槽饱和？一个 SM（4 个调度器）需要多少个？占用率至少多少？

答：$\lceil (30 + 1 + 600)/(30 + 1) \rceil = \lceil 20.35 \rceil = 21$ 个 warp / 调度器，一个 SM 需要 84 个 warp，超过上限 64，所以即使占用率 100% 也无法完全藏住延迟（最多 $16 \times 31 / 631 \approx 79\%$）。要么提高每个 warp 的指令级并行（例如一次发出多个独立的访存），要么减少访存（复用）。

练习 7（波次）：一个内核在 H100 上每 SM 驻留 3 个块，网格有 400 个块。要跑几波？最后一波的利用率是多少？如果把块大小减半（块数变成 800，每 SM 驻留 6 个块）呢？

答：一波 $132 \times 3 = 396$ 个块，400 个块要 2 波，第二波只有 4 个块，总利用率 $400/792 \approx 50.5\%$。块数 800、一波 792：仍然 2 波，第二波 8 个块，利用率 $800/1584 \approx 50.5\%$——一样糟。更好的办法是选一个让块数接近 396 整数倍的分块，或者用 persistent kernel（第 7–8 讲）。

练习 8（Tensor Core 指令数）：用 mma.sync.m16n8k16 计算 $C = AB$，$A$ 是 $128 \times 64$、$B$ 是 $64 \times 256$。需要多少条 mma 指令？用 CUDA core 的 FFMA 需要多少条？

答：乘加次数 $128 \times 256 \times 64 = 2{,}097{,}152$。mma：$2{,}097{,}152 / 2048 = 1024$ 条；FFMA：$2{,}097{,}152 / 32 = 65{,}536$ 条（多 64 倍）。

13. 延伸阅读 · Further reading

• NVIDIA, CUDA C++ Programming Guide：第 2 章（编程模型）、第 4 章（硬件实现：SIMT 架构、硬件多线程）、第 5 章（性能指南：占用率、访存）。
• NVIDIA, NVIDIA H100 Tensor Core GPU Architecture 白皮书：SM 结构、Tensor Core、线程块簇、TMA，第 6 讲会用到。
• V. Volkov, Better Performance at Lower Occupancy（GTC 2010）：为什么低占用率也能跑得快——用指令级并行代替线程级并行。
• T. G. Rogers, M. O’Connor, T. Aamodt, Cache-Conscious Wavefront Scheduling（MICRO 2012）：warp 调度策略（包括 GTO）对性能的影响，和 §7.5 的车道效应有关。
• 本仓库 CMU11-868 第 2–3 讲：同样从零模拟 SIMT 执行，侧重 CUDA 编程作业的写法，可以对照阅读。

14. 下一讲预告 · Next lecture

第 4 讲《GPU 架构与 CUDA 编程（二）》讲 slides 后半部分的两个案例：矩阵乘法（朴素版 → 线程级寄存器分块 → 线程块级共享内存分块，每一步分析全局内存和共享内存的访问量，以及合作加载）和并行归约（为什么 CUDA 没有全局同步、怎样拆成多个内核；交错寻址的发散问题；步长索引；顺序寻址与 bank 冲突；warp 级的优化）。我们会在 gpusim.py 上把每个版本都跑一遍，数出事务、bank 冲突和发散，并用 roofline 估算它们在 H100 上的速度。