计算机系统基础-02：GPU

GPU 分类

独立 GPU 常见于桌面和服务器。特点：

• GPU 是单独芯片。
• 拥有独立显存 VRAM。
• 通过 PCIe 与 CPU 连接。
• 计算能力和内存带宽通常更高。

集成 GPU 常见于手机、笔记本等移动设备。特点：

• 与 CPU 集成在同一系统中。
• 与 CPU 共享系统内存 DRAM。
• 功耗较低。
• 计算能力和内存带宽通常低于高端独立 GPU。

NVIDIA GPU 架构

NVIDIA GPU 由多个 Streaming Multiprocessor, SM 组成。多个 SM 共享全局 L2 Cache，并访问 GPU 显存。

可以将 SM 理解为 GPU 的基本计算单元，类似 CPU 中的核心，但其设计目标不同：SM 更接近一个向量处理器，内部包含多个执行单元。

一个 SM 通常包含：

• 多个 CUDA Core。
• Load/Store 单元。
• Special Function Unit, SFU。
• 寄存器文件。
• warp scheduler。
• 片上共享内存和 L1 Cache。

CUDA Core 中包含整数 ALU 和浮点 ALU，用于执行算术操作。

SIMT

GPU 使用 SIMT Single Instruction Multiple Threads 执行模型。

SIMT 的基本含义：

• 一条指令只解码一次。
• 同组线程执行相同指令。
• 每个线程使用自己的寄存器和数据。
• NVIDIA GPU 中一个 SIMT 组称为 warp，通常包含 32 个线程。

warp 调度

一个 SM 可以同时驻留多个 warp。例如 Fermi GPU 每个 SM 有 2 个 warp scheduler。SM 可在多个 warp 之间切换：

• 若当前 warp 因内存访问等原因阻塞，可切换到已就绪的 warp。
• 每个线程的状态持久保存在寄存器中。
• warp 切换不需要像 CPU 线程切换那样保存和恢复上下文，因此成本很低。

这种机制通过并发驻留多个 warp 来隐藏内存访问延迟。

分支发散

SIMT 要求同一 warp 中线程执行相同指令。若线程在分支处走向不同路径，会出现 thread divergence。

若同一 warp 中部分线程执行 A()，部分线程执行 B()，GPU 会分阶段执行不同分支，并在某些阶段禁用不属于当前分支的线程。被禁用的核心在这些周期内不产生有效工作。

分支发散会降低 warp 内执行单元利用率。GPU 程序应尽量让同一 warp 内线程执行相同控制流。

GPU 内存模型

GPU 层次

GPU 具有嵌套层次：

thread -> warp/SIMT group -> SM -> GPU

GPU 包含多种内存。不同内存在容量、延迟、带宽、读写权限和共享范围上不同。

类型	读取延迟	大小	读写属性	共享范围
Register	1 cycle	每个 SM 约 128 KB	R/W	单个线程
Local Memory	400 ~ 600 cycles	-	R/W	单个线程
L1 Cache	1 ~ 32 cycles	每个 SM 约 192 KB	透明缓存	SM
Shared Memory	1 ~ 32 cycles	每个 SM 约 192 KB	R/W	SIMT groups
L2 Cache	32 ~ 64 cycles	约 6 MB	透明缓存	GPU
Global Memory	400 ~ 600 cycles	GB 级	R/W	GPU
Constant Memory	100 ~ 200 cycles	KB 级	只读	GPU
Texture Memory	400 ~ 600 cycles	-	只读	GPU

• Register 延迟最低，只属于单个线程。
• Local Memory 类似 CPU 的栈，但物理上通常位于 global memory。
• Shared Memory 是可编程片上内存，常用于同一 SM 内线程协作。
• Global Memory 容量大延迟高，是 GPU 数据的主要存放位置。
• Constant Memory 位于全局内存体系中，但支持专门的读取路径和缓存机制。
• L1/L2 Cache 类似 CPU cache，对程序通常是透明的。

CUDA

CUDA 即 Compute Unified Device Architecture，风格接近 C/C++，用于表达在 GPU 上运行的并行程序。

术语	含义
Host	CPU
Host memory	CPU 内存
Host code	在 CPU 上运行的代码
Device	GPU
Device memory	GPU 内存，独立 GPU 中通常指显存
Device code	在 GPU 上运行的代码

CPU 程序的一般流程：

定义函数->分配 CPU 内存->加载数据->在 CPU 上计算->保存结果->释放 CPU 内存

独立 GPU 程序通常需要额外处理主机与设备之间的数据移动：

将 CPU/GPU 代码分别加载到 CPU/GPU 内存->为数据分配 CPU/GPU 内存->将输入数据加载到 CPU 内存->将数据从 CPU 内存复制到 GPU 内存->启动 GPU kernel->将结果从 GPU 内存复制回 CPU 内存

对独立 GPU 而言，CPU 内存和 GPU 显存是分离的。数据传输通常通过 PCIe 完成，传输开销可能影响整体性能。

CUDA Kernel

GPU kernel 是运行在 GPU 上的函数。__global__ 表示函数：

• 在 device 上执行。
• 从 host code 中调用。
• 以单个线程的视角编写。

向量加法 kernel：

__global__ void vecAdd(float *a, float *b, float *c, size_t length)
{
    /* blockIdx.x 是 block 编号，threadIdx.x 是 block 内线程编号。 */
    size_t idx = blockDim.x * blockIdx.x + threadIdx.x;
    /* 总线程数可能大于向量长度，需要边界检查避免越界访问。 */
    if (idx < length) {
        /* 每个 GPU 线程负责计算一个输出元素。 */
        c[idx] = a[idx] + b[idx];
    }
}

该函数中每个线程负责一个元素。idx 根据 block 编号和 thread 编号计算全局线程编号。

CUDA 使用 SPMD Single Program Multiple Data 模型：

• 所有处理单元执行同一个程序。
• 每个线程处理不同数据。
• 在 CUDA 中，程序通常是 kernel，处理单元是 CUDA core 上执行的线程。

• SPMD 描述程序员视角：同一个 kernel 被许多线程执行。
• SIMT 描述硬件执行方式：线程被组织成 warp，同一 warp 中线程同步执行指令。

CUDA 内存管理

Host 指针和 device 指针指向不同地址空间：

float *h_A;  // host pointer，指向 CPU 主存
float *d_A;  // device pointer，指向 GPU 显存

/* 在 CPU 侧分配输入或输出数组。 */
h_A = (float *)malloc(size);

/* 在 GPU 侧分配同样大小的显存空间，地址写入 d_A。 */
cudaMalloc((void **)&d_A, size);

• Device pointer 指向 GPU 内存。
• Device pointer 可在 host code 中传递，但不能在 host code 中解引用。
• Host pointer 指向 CPU 内存。
• Host pointer 可传给 device code，但不能在 device code 中直接解引用。
• cudaMalloc() 和 cudaFree() 类似 malloc() 和 free()，用于管理设备内存。

CPU 与 GPU 之间通过 cudaMemcpy() 复制数据，并通过方向参数指定复制方向：

/* 将输入数组从 CPU 主存复制到 GPU 显存，供 kernel 读取。 */
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

/* kernel 执行后，将结果从 GPU 显存复制回 CPU 主存。 */
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

• cudaMemcpyHostToDevice：CPU 到 GPU。
• cudaMemcpyDeviceToHost：GPU 到 CPU。

Kernel Launch

CUDA 使用 <<< >>> 从 host code 启动 device code，该操作称为 kernel launch。

/* 启动 16 个 block，每个 block 64 个 thread。 */
/* d_A、d_B、d_C 是 device pointer，kernel 在 GPU 上解引用。 */
vecAdd<<<16, 64>>>(d_A, d_B, d_C, length);

• <<<16, 64>>> 指定并行结构。
• 第一个参数是 block 数量。
• 第二个参数是每个 block 的 thread 数量。
• 圆括号中是 kernel 参数。

该例启动 16 个 block，每个 block 有 64 个 thread，总线程数为 $16 \times 64 = 1024$。

CUDA 线程模型

CUDA 程序包含层次化并发线程：

Grid -> Block -> Thread

基本规则：

• 一个 grid 由多个 block 组成。
• 一个 block 由多个 thread 组成。
• 启动一个 kernel 等价于启动一个 grid。
• Grid 和 block 默认是一维，也可扩展到最多三维。

常见维度变量：

变量	含义
gridDim.x/y/z	grid 在各维度上的 block 数量
blockDim.x/y/z	每个 block 在各维度上的 thread 数量
blockIdx.x/y/z	当前 block 的编号
threadIdx.x/y/z	当前 thread 在 block 内的编号

一维数组索引公式：

$$\text{thread\_id} = \text{blockDim.x} \cdot \text{blockIdx.x} + \text{threadIdx.x}$$

总线程数：

$$\text{total\_threads} = \text{gridDim.x} \cdot \text{blockDim.x}$$

例如：

gridDim.x = 16
blockDim.x = 64
blockIdx.x = 2
threadIdx.x = 1

thread_id = 64 * 2 + 1 = 129
total_threads = 16 * 64 = 1024

CUDA block 与 SM 的映射规则：

• 一个 block 中的所有 thread 在同一个 SM 上运行。
• 一个 block 不会被拆分到多个 SM。
• 一个 block 在执行期间不会迁移到其他 SM。
• 一个 SM 可同时执行多个 block。
• block 由硬件自动分配到 SM。

这种映射使 block 内线程可以通过 shared memory 协作。

GPU 性能趋势

GPU 性能趋势考虑带宽、计算能力、能效和计算/带宽比：

• GPU 显存带宽持续增长。带宽决定从 device memory 读取或写入数据的速度。对于访存密集、每字节计算量较低的程序，带宽通常成为主要瓶颈。

• GPU 峰值浮点计算能力持续增长，尤其适合独立算术操作较多的任务。对于每字节数据执行较多浮点操作的程序，计算单元吞吐可能成为主要瓶颈。

• GPU 在单位能耗上的计算能力持续提升，因此适合高并行吞吐任务。但实际能效仍取决于访存模式、分支发散、线程占用率和算法的并行度。

• 计算能力与带宽的比值常用 算术强度 Arithmetic Intensity / Operational Intensity 描述：

  $$I=\frac{\text{floating point operations}}{\text{bytes moved from device memory}}$$

  其含义是：从 device memory 移动每个字节数据所对应的浮点运算数量。

总体趋势是峰值计算能力相对于带宽增长更快，因此硬件要求程序具有更高算术强度，才能充分利用计算单元。

Roofline 模型

Roofline 模型用于判断给定硬件上的程序是 memory bandwidth bound 还是 compute bound。

核心思想：

• 横轴表示算术强度 $I$。
• 纵轴表示可达到的计算性能。
• 当 $I$ 较低时，性能受内存带宽限制。
• 当 $I$ 较高时，性能受峰值计算能力限制。

性能上界可表示为：

$$P \le \min(P_{peak},\ I \cdot B_{peak})$$

其中：

• $P$ 是实际可达性能。
• $P_{peak}$ 是硬件峰值计算性能。
• $B_{peak}$ 是硬件峰值内存带宽。
• $I$ 是算术强度。

因此优化 GPU 程序时，需要判断瓶颈位置：

瓶颈类型	特征	优化方向
带宽受限	算术强度低，访存占主导	减少 global memory 访问，合并访存，使用 shared memory，提高数据复用
计算受限	算术强度高，计算单元占主导	提高指令吞吐，减少分支发散，提高 occupancy，利用专用计算库