第7章: GPU设计中的流式多处理器设计
流式多处理器（SM）是NVIDIA GPU体系结构的基本组成单元。每个SM包含一组CUDA核心,以SIMT（单指令,多线程）的方式执行指令。SM负责管理和调度线程束,处理分支发散,并提供对共享内存和缓存的快速访问。在本章中,我们将探讨SM的微体系结构,包括其流水线、线程束调度机制、寄存器文件设计以及共享内存和L1缓存组织。
SM微体系结构和流水线
SM是一种高度并行和流水化的处理器,旨在高效地并发执行数百个线程。图7.1显示了NVIDIA Volta架构中SM的简化框图。
指令缓存
|
v
线程束调度器
|
v
调度单元 (4个线程束)
| | | |
v v v v
CUDA核心 (FP64/FP32/INT)
CUDA核心 (FP64/FP32/INT)
CUDA核心 (FP64/FP32/INT)
...
张量核心
张量核心
...
加载/存储单元
加载/存储单元
...
特殊功能单元
^
|
寄存器文件 (64 KB)
^
```分享内存/L1缓存(96 KB)
图7.1: NVIDIA Volta架构中SM的简化框图。
SM的主要组件包括:
-
指令高速缓存: 存储频繁访问的指令,以减少延迟并提高吞吐量。
-
Warp调度器: 选择准备就绪的warps并将它们派送到可用的执行单元。
-
派送单元: 每个周期从最多4个warps中提取和解码指令,并将它们派送到适当的执行单元。
-
CUDA核心: 可编程的执行单元,支持广泛的整数和浮点运算。每个Volta SM包含64个CUDA核心。
-
张量核心: 专为加速深度学习和AI工作负载而设计的专用执行单元。每个Volta SM包含8个张量核心。
-
加载/存储单元: 处理内存操作,包括对全局内存、共享内存和缓存的加载和存储。
-
特殊功能单元: 执行超越和其他复杂的数学运算。
-
寄存器文件: 提供对线程私有寄存器的快速访问。每个Volta SM有64 KB的寄存器文件。
-
共享内存/L1缓存: 一个可配置的内存空间,可用作软件管理的缓存(共享内存)或硬件管理的L1数据缓存。
SM管道被设计为最大化吞吐量,允许多个warps并发执行并隐藏内存延迟。图7.2展示了SM管道的简化视图。
指令获取
|
v
指令解码
|
v
操作数集合
|
v
执行(CUDA核心、张量核心、加载/存储单元、特殊功能单元)
|
v
写回图7.2: 简化的SM管道。
管道阶段如下:
-
指令获取: Warp调度器选择一个准备就绪的warp以下是该 Markdown 文件的中文翻译。对于代码部分,我只翻译了注释,代码本身未做翻译。
-
指令获取(Instruction Fetch): SM 从指令缓存中获取下一条指令。
-
指令解码(Instruction Decode): 获取的指令被解码,以确定操作类型、操作数和目标寄存器。
-
操作数收集(Operand Collection): 从寄存器文件或共享内存中收集指令所需的操作数。
-
执行(Execution): 指令在适当的执行单元(CUDA 核心、张量核心、加载/存储单元或特殊功能单元)上执行。
-
写回(Writeback): 执行结果被写回寄存器文件或共享内存。
为了实现高性能,SM 采用了几种技术来最大化资源利用率和隐藏延迟:
- 双发射(Dual-Issue): SM 可以在单个周期内发射两个独立的指令,从而增加指令级并行度。
- 流水线执行单元(Pipelined Execution Units): 执行单元采用流水线设计,允许 SM 在上一操作完成之前开始新的操作。
- 延迟隐藏(Latency Hiding): SM 可以在循环周期中切换 warp,从而隐藏内存访问和长延迟操作的延迟,执行其他 warp 的指令。
示例 7.1 展示了一个简单的 CUDA 内核,用于对两个向量执行元素级加法。
__global__ void vectorAdd(int *a, int *b, int *c, int n) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < n) {
c[tid] = a[tid] + b[tid];
}
}示例 7.1: 用于向量加法的 CUDA 内核。
在此示例中,内核中的每个线程计算输入向量 a 和 b 中对应元素的和,并将结果存储在输出向量 c 中。SM 通过将每个线程分配到一个 CUDA 核心并调度 warp 线程在可用核心上执行来执行此内核。加载/存储单元用于从全局内存中获取输入数据并将结果写回。
Warp 调度和分支发散处理
Ef这个 Markdown 文件的中文翻译如下:
对于 SM (流处理器) 来说, 高效的 warp 调度是最大化性能的关键。warp 调度器负责选择就绪的 warp 并将它们分派到可用的执行单元上。warp 调度器的主要目标是保持执行单元繁忙,确保总是有可执行的 warp。
SM 采用了两级 warp 调度机制:
-
Warp 调度: warp 调度器根据调度策略(如轮转或最老优先)选择就绪的 warp,然后将它们分派到可用的执行单元上。
-
指令调度: 在每个 warp 内部,SM 根据指令依赖关系和执行单元的可用性调度指令。SM 可以在单个周期内发射多条来自同一 warp 的独立指令,以最大化指令级并行性。
图 7.3 说明了两级 warp 调度机制。
Warp 池
Warp 1 (就绪)
Warp 2 (等待)
Warp 3 (就绪)
...
Warp N (就绪)
|
v
Warp 调度器
|
v
分派单元
|
v
执行单元图 7.3: 两级 warp 调度机制。
warp 调度中的一个关键挑战是处理分支发散。在 SIMT 执行模型中,一个 warp 中的所有线程以锁步的方式执行相同的指令。但是,当一个 warp 遇到分支指令(如 if-else 语句)时,有些线程可能会执行 if 分支,而其他线程会执行 else 分支。这种情况称为分支发散。
为了处理分支发散,SM 使用一种称为谓词化的技术。当 warp 遇到发散分支时,SM 会顺序执行分支的两个路径,并对不属于每个路径的线程进行屏蔽。然后使用谓词寄存器将结果合并,确保每个线程都获得正确的结果。
示例 7.2 展示了一个包含发散分支的 CUDA 内核。这是中文翻译:
__global__ void 分支内核(int *数据, int *结果) {
int 线程ID = 块索引.x * 块维度.x + 线程索引.x;
if (数据[线程ID] > 0) {
结果[线程ID] = 数据[线程ID] * 2;
} else {
结果[线程ID] = 数据[线程ID] * 3;
}
}例7.2: 具有分支的CUDA内核。
在此示例中,分支条件 数据[线程ID] > 0 可能导致某些线程在一个warp中采取if路径,而其他线程采取else路径。SM通过按顺序执行两个路径并屏蔽每个路径中非活动线程来处理此分支。
图7.4说明了具有分支线程的warp的谓词化过程。
Warp (32个线程)
线程1: 数据[1] = 5, 结果[1] = 10
线程2: 数据[2] = -3, 结果[2] = -9
...
线程32: 数据[32] = 7, 结果[32] = 14
分支:
if (数据[线程ID] > 0) {
结果[线程ID] = 数据[线程ID] * 2;
} else {
结果[线程ID] = 数据[线程ID] * 3;
}
谓词化:
步骤1: 执行if路径并设置掩码
线程1: 结果[1] = 10
线程2: (被屏蔽)
...
线程32: 结果[32] = 14
步骤2: 执行else路径并设置掩码
线程1: (被屏蔽)
线程2: 结果[2] = -9
...
线程32: (被屏蔽)
最终结果:
线程1: 结果[1] = 10
线程2: 结果[2] = -9
...
线程32: 结果[32] = 14图7.4: 具有分支线程的warp的谓词化过程。
通过使用谓词化,SM可以处理分支分歧,而无需显式分支指令或控制流分歧。但是,分支分歧仍然会影响性能,因为SM必须按顺序执行两个路径,从而降低了有效并行度。
寄存器文件和操作数收集器
寄存器文件是SM的关键组成部分,提供对线程专有寄存器的快速访问。每个SM都有一个大型寄存器文件,以支持许多活动线程并实现warp之间的高效上下文切换。在NVIDIA Volta架构中，每个SM有一个64 KB的寄存器文件，组织为32个2 KB的银行。寄存器文件被设计为提供高带宽和低延迟访问,以支持大量并发线程。
为了最小化银行冲突并提高性能,SM采用了一种称为操作数收集的技术。操作数收集器是专门的单元,它从寄存器文件银行收集操作数,并将它们传送到执行单元。通过使用操作数收集器,SM可以减少银行冲突的影响,并提高执行单元的利用率。
图7.5显示了SM中寄存器文件和操作数收集器的简化图。
寄存器文件 (64 KB)
银行 1 (2 KB)
银行 2 (2 KB)
...
银行 32 (2 KB)
|
v
操作数收集器
|
v
执行单元图7.5: SM中的寄存器文件和操作数收集器。
操作数收集器的工作原理是从多个指令和多个Warp中收集操作数,允许SM在单个周期内从不同的Warp发出指令到执行单元。这有助于隐藏寄存器文件访问的延迟,并提高SM的总体吞吐量。
示例7.3显示了一个执行两个向量点积的CUDA内核。
__global__ void dotProduct(float *a, float *b, float *result, int n) {
// 为每个块创建一个共享内存数组来存储部分和
__shared__ float partialSum[256];
// 获取当前线程的索引
int tid = threadIdx.x;
// 计算当前线程处理的元素索引
int i = blockIdx.x * blockDim.x + threadIdx.x;
// 初始化当前线程的部分和为0
partialSum[tid] = 0;
// 计算点积的部分和
while (i < n) {
partialSum[tid] += a[i] * b[i];
i += blockDim.x * gridDim.x;
}
// 等待所有线程完成部分和计算
__syncthreads();
// 使用树状归约方式计算最终的点积
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
if (tid < s) {
partialSum[tid] += partialSum[tid + s];
}
__syncthreads();
}
// 将最终结果写入输出数组
if (tid == 0) {
result[blockIdx.x] = partialSum[0];
}
}在此示例中,每个线程使用分配给它的元素计算点积的部分和。以下是该 Markdown 文件的中文翻译版本。对于代码部分,我只翻译了注释,而没有翻译代码本身。
从输入向量中提取元素。部分和存储在共享内存数组 partialSum 中。当所有线程完成计算部分和后,进行并行归约以求出最终的点积结果。
操作数收集器在此示例中起着关键作用,通过有效收集共享内存访问和算术运算的操作数,避免了银行冲突并提高了执行单元的利用率。
结论
流式多处理器是现代 GPU 架构中的核心计算单元。其设计着眼于通过细粒度多线程、SIMT 执行和高效的操作数收集来最大化吞吐量并隐藏内存延迟。
SM 的关键组件包括:选择要执行的 warp 的 warp 调度器、处理分支分散和聚合的 SIMT 堆栈、为线程私有寄存器提供快速访问的寄存器文件和操作数收集器,以及支持低延迟数据共享和重用的共享内存和 L1 缓存。
随着 GPU 架构的不断发展,在分支分散处理、warp 调度和寄存器文件设计等领域的研究将对提高未来 GPU 的性能和效率至关重要。诸如动态 warp 形成、线程块压缩和操作数重用缓存等新技术有望显著增强 SM 的功能,为并行计算工作负载带来新的性能水平。