第四课：SM 内部结构与指令执行过程

导语：这一部分进入 SM 内部，重点讲清楚 Warp Scheduler、发射、执行单元、Scoreboard 和流水线之间的关系。

前面已经知道：

Thread → 组成 Warp
Warp → 分配到 SM
SM → 执行 Warp 的指令

现在继续向 SM 内部深入。

先建立一个简化模型：

图示：第四课：SM 内部结构与指令执行过程

不同架构，例如 Volta、Ampere、Hopper、Blackwell，具体数量和组织方式不同，但基本思想相近。

---

一、SM 不是一个巨大的计算器

SM 更像一个小型多线程处理器。

它需要完成几件事情：

1. 保存很多 Warp 的状态；
2. 判断哪些 Warp 当前可以执行；
3. 从可执行 Warp 中选择一个或多个；
4. 读取线程寄存器；
5. 把指令发送到合适的执行单元；
6. 等待结果完成；
7. 继续执行后续指令。

可以把 SM 想象成一个工厂：

Warp Scheduler：生产调度员
Scoreboard：依赖状态表
Register File：原料仓库
CUDA Core：普通算术生产线
INT Unit：整数生产线
LSU：访存生产线
SFU：特殊函数生产线
Tensor Core：矩阵生产线

每种指令需要送到对应的生产线。

---

二、一个 Warp 包含哪些状态

一个 Warp 不只是 32 个线程编号。

硬件还需要保存它的执行状态，例如：

图示：二、一个 Warp 包含哪些状态

其中最重要的是：

• PC；
• Active Mask；
• 寄存器；
• 等待状态。

---

1. 程序计数器 PC

PC 表示 Warp 当前执行到哪条指令。

例如程序为：

x = A[i];
y = B[i];
z = x + y;
C[i] = z;

对应的简化指令序列可能是：

PC 0：LOAD A[i]
PC 1：LOAD B[i]
PC 2：ADD x, y
PC 3：STORE C[i]

Warp 当前执行完 PC 1 后，下一步通常准备执行 PC 2。

---

2. Active Mask

一个 Warp 有 32 个线程，但并不代表每条指令都一定有 32 个线程参与。

硬件会用一个 32 位掩码表示哪些线程当前有效。

例如：

Active Mask = 11111111111111111111111111111111

表示 32 个线程全部活跃。

如果只有偶数线程活跃：

Active Mask = 01010101010101010101010101010101

发射同一条指令时，只有掩码中为 1 的线程执行。

这就是分支发散时屏蔽部分线程的基础机制。

---

三、Warp Scheduler 做什么

Warp Scheduler，也就是 Warp 调度器，负责从 SM 中已经驻留的 Warp 里选择可以执行的 Warp。

假设一个 SM 中有多个 Warp：

Warp 0：等待显存
Warp 1：可以执行
Warp 2：等待前一条乘法结果
Warp 3：可以执行
Warp 4：等待同步
Warp 5：可以执行

调度器会从 Warp 1、Warp 3、Warp 5 中选择一个或多个发射。

调度器不会随意选择，它通常要考虑：

• 指令操作数是否准备好；
• 对应执行单元是否空闲；
• Warp 是否在等待内存；
• Warp 是否在等待同步；
• Warp 是否结束；
• 指令是否存在数据依赖。

---

四、Eligible Warp 和 Ready Warp

可以把 Warp 的状态粗略分为两类。

1. 不可发射 Warp

例如：

等待 Global Memory 数据
等待 Shared Memory 数据
等待前一条指令结果
等待 __syncthreads()
等待原子操作

这类 Warp 当前不能继续。

2. 可发射 Warp

它满足：

下一条指令已知
操作数已经准备好
对应执行单元可用
没有同步阻塞

这种 Warp 可以称为 Eligible Warp，即符合发射条件的 Warp。

调度器希望每个周期都有足够多的 Eligible Warp。

如果所有 Warp 都在等待，SM 的执行单元就会空闲。

---

五、Instruction Issue 和 Dispatch

这两个词经常一起出现。

Issue：发射

调度器决定：

让哪个 Warp 的哪条指令开始执行。

例如：

选择 Warp 3
发射一条 FP32 ADD 指令

Dispatch：分派

把这条指令送到正确的执行单元。

例如：

图示：Dispatch：分派

简化过程如下：

图示：Dispatch：分派

---

六、一个 Warp 的指令是否一拍执行完

这是一个非常重要的问题。

假设一个 Warp 有 32 个线程，执行：

z = x + y;

从编程角度看，32 个线程都执行加法。

但物理上是否一个周期内有 32 个加法器同时完成，要看具体微架构。

可能存在以下情况：

情况一：物理执行宽度等于 32

如果有足够的执行通道，Warp 的 32 个线程可以在一个发射步骤中覆盖全部通道。

情况二：物理执行宽度小于 32

例如只有 16 条执行通道，则可能分两拍执行：

第 1 拍：Thread 0～15
第 2 拍：Thread 16～31

情况三：执行单元按更复杂方式分区

现代 GPU 中，一个 SM 往往被划分成若干处理分区，每个分区有自己的调度器、寄存器访问通路和执行单元。

因此不能简单说：

一个 Warp 一定在一个周期内全部算完。

更准确的说法是：

Warp 是调度和指令发射单位，但实际物理执行宽度由具体架构决定。

---

七、延迟和吞吐率不是一回事

GPU 架构中经常同时出现两个概念：

• Latency，延迟；
• Throughput，吞吐率。

假设一条浮点加法指令需要 4 个周期才产生结果。

这叫：

指令延迟 = 4 cycles

但如果流水线每个周期都可以接收一条新指令，那么：

吞吐率 = 每周期 1 条

示意如下：

图示：七、延迟和吞吐率不是一回事

虽然每条指令都要经过多个周期，但流水线填满以后，可以连续产生结果。

---

八、什么是流水线

假设一次乘法需要经过四个阶段：

阶段 1：读取操作数
阶段 2：部分乘积计算
阶段 3：加法归约
阶段 4：结果写回

如果完全串行：

指令 A 完成后
指令 B 才开始

效率很低。

流水线方式是：

图示：八、什么是流水线

流水线提高的是吞吐率，而不是单条指令本身的延迟。

---

九、数据依赖是什么

考虑下面的代码：

a = b + c;
d = a * e;

第二条指令使用第一条指令的结果 a。

这叫做数据依赖：

ADD 产生 a
MUL 消费 a

在 a 产生之前，乘法不能开始。

简化为：

Instruction 1：ADD  R1, R2, R3
Instruction 2：MUL  R4, R1, R5

第二条指令依赖寄存器 R1。

---

十、Scoreboard 是什么

Scoreboard 可以理解为硬件中的“寄存器依赖状态表”。

它记录：

• 哪个寄存器正在被某条指令写入；
• 哪个结果还没有完成；
• 下一条指令所需的数据是否准备好。

例如：

R1：等待 ADD 结果
R2：Ready
R3：Ready
R4：等待 LOAD 结果

Warp 下一条指令是：

MUL R5, R1, R2

因为 R1 还没有准备好，所以这条指令不能发射。

Scoreboard 会让该 Warp 暂时等待。

此时调度器可以选择其他 Warp。

---

十一、Scoreboard 如何帮助隐藏延迟

假设 Warp 0 执行：

LOAD R1, [A]
ADD  R2, R1, R3

LOAD 发出后，R1 尚未返回。

Scoreboard 标记：

R1 = Not Ready

因此 Warp 0 的 ADD 不能执行。

但其他 Warp 可能可以执行：

Warp 1：FP32 ADD
Warp 2：INT ADD
Warp 3：Tensor Core MMA

于是调度器切换到其他 Warp。

等内存返回后：

R1 = Ready

Warp 0 再次变成 Eligible Warp。

因此 GPU 隐藏延迟的完整逻辑是：

图示：十一、Scoreboard 如何帮助隐藏延迟

---

十二、线程切换为什么很快

CPU 线程切换通常可能需要保存和恢复：

• 寄存器；
• 程序计数器；
• 栈指针；
• 操作系统调度状态。

GPU 的 Warp 切换不是传统意义上的操作系统上下文切换。

因为多个 Warp 的状态已经同时保存在 SM 内部，例如：

Warp 0 的寄存器状态
Warp 1 的寄存器状态
Warp 2 的寄存器状态
Warp 3 的寄存器状态

调度器只需要选择另一个 Warp 的状态来发射。

所以切换通常不需要把完整上下文搬到内存中。

这也是 GPU 能够频繁切换 Warp 的原因。

---

十三、寄存器文件如何工作

每个线程都有自己的逻辑寄存器。

假设一个 Warp 中，每个线程都执行：

z = x + y;

那么硬件需要同时读取：

32 个线程的 x
32 个线程的 y

然后写回：

32 个线程的 z

因此 Register File 必须提供很高的带宽。

可以简单理解为：

图示：十三、寄存器文件如何工作

指令发射时，根据 Warp ID 和寄存器编号读取对应数据。

---

寄存器文件为什么很大

假设一个 SM 同时驻留 2048 个线程，每个线程使用 64 个 32 位寄存器：

$$2048\times64\times32\text{ bit}$$

换算成字节：

$$2048\times64\times4=524288\text{ bytes}$$

也就是约 512 KB。

因此 GPU 的寄存器文件往往非常大。

但它仍然属于宝贵资源，因为：

• 容量大；
• 端口多；
• 带宽要求高；
• 功耗较高。

---

十四、CUDA Core 到底是什么

CUDA Core 通常指执行普通标量浮点运算的算术单元。

典型操作包括：

FP32 ADD
FP32 MUL
FP32 FMA

其中 FMA 是：

$$D=A\times B+C$$

一次 FMA 通常可计为两个浮点操作：

• 一次乘法；
• 一次加法。

所以厂商计算峰值算力时，经常使用：

$$P_{\text{peak}}= N_{\text{cores}}\times f\times\text{ops per cycle}$$

如果每个核心每周期执行一个 FMA，则通常按 2 FLOPs 计算。

---

十五、INT Unit 做什么

整数单元执行：

• 整数加减；
• 位运算；
• 比较；
• 地址计算；
• 循环计数；
• 索引计算。

例如：

int i = blockIdx.x * blockDim.x + threadIdx.x;

这里的乘法、加法和索引计算主要是整数运算。

神经网络看起来主要是浮点或低精度矩阵计算，但程序中仍有大量整数操作，例如：

• 地址生成；
• 循环控制；
• 边界判断；
• 张量索引。

某些架构可以让整数指令和浮点指令并行执行，从而提高吞吐率。

---

十六、Load/Store Unit

Load/Store Unit，简称 LSU，主要处理：

• 地址计算；
• 内存读请求；
• 内存写请求；
• Shared Memory 访问；
• Global Memory 访问；
• Local Memory 访问；
• 某些原子操作。

例如：

x = A[i];

大致需要：

1. 计算 A + i × sizeof(float)
2. 形成内存请求
3. 合并 Warp 内多个线程请求
4. 查询缓存
5. 必要时访问显存
6. 返回数据
7. 写入线程寄存器

LSU 不只是简单“读内存”，还承担地址和访问组织工作。

---

十七、SFU 是什么

SFU 是 Special Function Unit，即特殊函数单元。

它主要执行普通加法器和乘法器不擅长的运算，例如：

• 倒数；
• 平方根；
• 正弦；
• 余弦；
• 指数；
• 对数；
• 插值近似。

这些运算通常通过近似、多项式、查表或专用硬件完成。

例如神经网络中的：

exp(x)
1 / sqrt(x)

可能会使用 SFU 或一系列组合指令实现。

SFU 数量通常比普通 FP32 执行单元少，因此大量特殊函数可能成为瓶颈。

---

十八、Tensor Core 是什么

Tensor Core 是专门执行小矩阵乘加运算的单元。

普通 CUDA Core 可能逐个执行：

$$a\times b+c$$

Tensor Core 则一次处理矩阵块：

$$D=A\times B+C$$

例如：

小块矩阵 A
×
小块矩阵 B
+
累加矩阵 C

Tensor Core 的优势在于：

• 一次完成大量乘加；
• 支持 FP16、BF16、TF32、FP8、INT8、FP4 等精度；
• 吞吐率远高于普通标量运算；
• 特别适合神经网络和矩阵乘法。

但它要求数据按照适合矩阵指令的格式组织。

所以有 Tensor Core 不等于程序自动获得高性能。

还需要：

• 合适的数据精度；
• 合适的矩阵维度；
• 足够的数据复用；
• 高效的内存布局；
• 足够高的 Tensor Core 利用率。

---

十九、一条指令的完整执行流程

假设 Warp 执行：

z = x + y;

可以粗略分成以下步骤。

第一步：取指

根据 Warp 的 PC 找到下一条指令。

PC → ADD R3, R1, R2

第二步：检查依赖

Scoreboard 检查：

R1 是否 Ready
R2 是否 Ready
R3 是否存在写冲突

第三步：选择 Warp

Warp Scheduler 判断该 Warp 是否 Eligible。

第四步：发射

发射 ADD 指令。

第五步：读取寄存器

从 Register File 中读取 Warp 内各线程的 R1 和 R2。

第六步：分派

把指令送到 FP32 执行流水线。

第七步：执行

执行多个线程的加法。

第八步：写回

把结果写回各线程的 R3。

第九步：更新 Scoreboard

R3 被标记为 Ready。

第十步：更新 PC

Warp 准备执行下一条指令。

---

二十、看一个完整例子

假设 Kernel 中有：

float x = A[i];
float y = B[i];
float z = x * y;
C[i] = z;

简化指令为：

I0：LOAD R1, [A+i]
I1：LOAD R2, [B+i]
I2：MUL  R3, R1, R2
I3：STORE [C+i], R3

执行过程可能是：

图示：二十、看一个完整例子

这正是 GPU 依靠多 Warp 隐藏内存延迟的过程。

---

二十一、指令级并行和线程级并行

GPU 中有两种重要的并行性。

1. Thread-Level Parallelism

大量 Warp 之间并行。

例如：

Warp 0 等待
Warp 1 执行
Warp 2 执行
Warp 3 准备

这是 GPU 最主要的并行方式之一。

2. Instruction-Level Parallelism

同一个 Warp 中，彼此独立的指令可以重叠执行。

例如：

a = b + c;
d = e * f;

第二条指令不依赖第一条。

理论上可以在不同执行管线中重叠。

而下面的代码：

a = b + c;
d = a * f;

第二条依赖第一条，不能立即执行。

因此更高的指令独立性有利于提高流水线利用率。

---

二十二、Dual Issue 是什么

部分 GPU 架构中，一个 Warp Scheduler 在合适条件下可以在一个周期发射两条指令。

这称为 Dual Issue。

但通常需要满足：

• 两条指令相互独立；
• 使用不同执行管线；
• 操作数已经准备好；
• 对应执行单元可用；
• 架构允许这样的组合。

例如可能出现：

一条 FP32 指令
+
一条 INT 指令

同时发射。

但不能简单理解为：

任意两条指令都能同时执行。

Dual Issue 的具体规则高度依赖架构。

---

二十三、结构冲突

即使数据已经准备好，指令也不一定能马上执行。

假设多个 Warp 都要使用同一种执行单元：

Warp 0：SFU 指令
Warp 1：SFU 指令
Warp 2：SFU 指令
Warp 3：SFU 指令

但 SFU 数量有限，就会发生资源竞争。

这称为结构冲突。

类似地，如果大量 Warp 同时执行：

• Tensor Core 指令；
• Load/Store 指令；
• 特殊函数；
• 原子操作；

对应管线可能成为瓶颈。

所以性能不仅取决于有多少 Warp，还取决于指令混合是否合理。

---

二十四、流水线停顿的常见原因

Warp 不能发射时，常见原因包括：

数据依赖停顿

前一条指令结果还没有生成

内存停顿

等待 L1、L2 或 Global Memory

执行单元忙

目标流水线暂时无法接收新指令

同步停顿

等待同一 Block 其他线程到达同步点

分支停顿

执行不同分支路径或等待重新汇合

屏障停顿

等待 barrier 或 memory fence

指令取出停顿

指令缓存或前端供应不足

---

二十五、为什么高 Occupancy 不一定性能高

假设一个 SM 有很多 Warp，但这些 Warp 全都在等待同一个资源：

Warp 0：等待显存
Warp 1：等待显存
Warp 2：等待显存
Warp 3：等待显存
……

即使 Occupancy 很高，也可能没有 Warp 可以发射。

反过来，某个计算密集型程序可能 Occupancy 不高，但每个 Warp 有足够多的独立指令，可以持续使用执行单元。

因此：

高 Occupancy
≠
一定高性能

真正重要的是：

每周期是否有 Eligible Warp
执行管线是否被充分利用
内存系统是否能够持续供数

---

二十六、Latency Hiding 的两个来源

GPU 隐藏延迟主要依赖两种机制。

第一种：Warp-Level Latency Hiding

当前 Warp 等待时，执行其他 Warp。

Warp 0 等待
→ 执行 Warp 1
→ 执行 Warp 2
→ 执行 Warp 3

第二种：Instruction-Level Latency Hiding

同一个 Warp 内执行独立指令。

a = b + c;
d = e * f;

两个运算互不依赖，可以重叠。

实际 GPU 通常同时利用这两种方式。

---

二十七、SIMT 和流水线的关系

SIMT 解决的是：

一条指令如何作用于多个线程。

流水线解决的是：

多条指令如何在时间上重叠执行。

二者不是同一个概念。

可以这样理解：

SIMT：
一条 ADD 指令
→ 同时处理 Warp 中多个线程的数据

流水线：
ADD A 还没结束时
→ ADD B 已经进入下一批处理

---

二十八、从 RTL 角度看一个简化 SM

如果从数字电路设计角度抽象，一个简化 SM 可以表示为：

图示：二十八、从 RTL 角度看一个简化 SM

其中核心控制结构包括：

• Warp 状态表；
• Ready 队列；
• Scoreboard；
• 发射仲裁器；
• 执行流水线；
• 写回仲裁；
• 内存返回队列。

---

二十九、为什么 GPU 前端比 CPU 简单

高性能 CPU 通常有：

• 复杂分支预测；
• 乱序执行；
• 寄存器重命名；
• Reorder Buffer；
• 推测执行；
• 复杂异常恢复。

GPU 通常更依赖：

• 大量 Warp；
• Scoreboard；
• 相对简单的 Warp 调度；
• 编译器安排；
• 多线程延迟隐藏。

因此 GPU 往往牺牲复杂的单线程优化，把面积和功耗用于：

• 更多执行单元；
• 更大寄存器文件；
• 更多线程状态；
• 更高内存带宽；
• 更多 Tensor Core。

---

三十、本节最重要的结构图

图示：三十、本节最重要的结构图

---

本节核心结论

第一，Warp Scheduler 调度的是 Warp，而不是单个 Thread。

第二，Warp 是指令发射单位，但 32 个线程是否在一个周期内全部完成，取决于物理执行宽度。

第三，Scoreboard 用于跟踪寄存器和指令依赖。

第四，一个 Warp 等待数据时，调度器会执行其他 Warp。

第五，流水线可以提高吞吐率，但不会消除单条指令的延迟。

第六，延迟和吞吐率是两个不同概念。

第七，FP32、INT、LSU、SFU、Tensor Core 分别处理不同类型的指令。

第八，寄存器文件需要同时保存大量线程的状态，因此容量和带宽都很大。

第九，高 Occupancy 只是隐藏延迟的条件之一，并不保证高性能。

第十，GPU 的核心思想不是让单个线程特别快，而是让大量 Warp 持续填满执行流水线。

下一部分最适合学习：

第五课：分支、谓词执行与 Warp 发散

包括：

if/else 在硬件中如何执行
Active Mask 如何变化
Predicate Register 是什么
Warp 如何重新汇合
循环为什么也会发散
分支发散什么时候严重
怎样编写更适合 GPU 的控制流