L21：图形处理器（GPUs）

MIT 6.5900 Fall 2024 · Joel Emer（部分幻灯片致谢 Srini Devadas）主题：SIMT 执行模型、流式多处理器（SM）、GPU 存储层级、分支/访存发散、CUDA 编程模型

一、为什么研究 GPU

非常成功的商品化加速器/协处理器；
GPU 结合两种提效策略：大规模并行（massive parallelism）与专门化（specialization）；
体现了加速器中性能与可编程性之间的张力；
在可编程性范畴内，体现了"让常见情况快"（make the common case fast）的原则。

GPU 最初是为 3D 渲染设计的：

输入：场景描述（3D 表面几何，如三角网格；表面材质、光照、相机等）。输出：场景图像。渲染的简单定义：计算 3D 网格中每个三角形对图像中每个像素外观的贡献。

二、图形处理器时间线

90 年代中期以前：大部分图形处理在 CPU；VGA 控制器加速部分显示功能；
90 年代中—2000 年代中：用于 2D/3D 图形的固定功能加速器（三角形装配与光栅化、纹理映射与着色），编程用 OpenGL / DirectX API；
现代 GPU
：
- 一部分固定功能硬件（纹理、光栅操作、光线追踪……）；
- 加上可编程的数据并行多处理器；
- 编程：OpenGL/DirectX，外加更通用的语言（CUDA、OpenCL 等）。

GPU 在现代系统中的两种形态

形态	特征
独立 GPU（Discrete）	与 CPU 分离，经 PCIe 或专用链路连接，独立显存
集成 GPU（Integrated）	CPU 与 GPU 同片，共享主存与末级缓存

权衡：独立 GPU 带宽高但需跨总线拷贝数据；集成 GPU 无需拷贝，但 CPU 与 GPU 争用同一内存。

三、SIMT 执行模型（Single Instruction Multiple Thread）

SIMT 的核心概念（绿色为 Nvidia 术语）：

许多线程，每个线程有私有的体系结构状态（如寄存器）；
一起发射的线程组称为 warp；
一起发射的所有线程执行同一条指令；
整条流水线称为一个流式多处理器（Streaming Multiprocessor, SM）；
每个执行单元称为一条 lane。

多线程 + SIMT

通过在多个 warp 之间切换来隐藏延迟：单个 PC/取指/译码逻辑被一个 warp 内的多条 lane 共享，每条 lane 有自己的 GPR 与功能单元。

流式多处理器（SM）概览

每个 SM 支持数十个 warp（如 Kepler 为 64 个），每 warp 32 线程；
每周期取 1 条指令；
每周期发射 1 条就绪指令；
- 简单记分牌（scoreboarding）：warp 内全部元素都就绪才发射；
指令广播到所有 lane；
多线程是主要的延迟隐藏机制。

Little's Law（再次出现）

吞吐量 T = \frac{在飞数 N}{延迟 L}

示例：64 个 warp（在飞指令数），期望吞吐 1 指令/周期，则平均指令延迟须 $< 64$ 周期。

上下文大小 vs 上下文数目

SM 可根据所需寄存器（与共享内存）支持可变数目的上下文：

少量大上下文 → 寄存器溢出（register spill）更少；
大量小上下文 → 延迟容忍更强；
选择权交给编译器。

Kepler 示例：最多 64 warp @ ≤32 寄存器/线程；最少 8 warp @ 256 寄存器/线程。

四、GPU 的多种存储类型

存储类型	特征
每线程私有内存（Per Thread Memory）	无跨线程共享，小且固定大小，可用于常量，多体实现（可位于全局内存）
共享暂存内存（Shared Scratchpad）	线程共享数据，每 SM 16K–64K，分体以获得高带宽，由地址合并单元（ACU）+ 交叉开关馈送
全局内存（Global Memory）	大容量共享，同样受访存发散影响，配 Cache 层级

访存发散（Memory Access Divergence）

所有 load 都是 gather，所有 store 都是 scatter；
地址合并单元（Address Coalescing Unit, ACU）检测顺序/步长模式并合并访存请求，但复杂模式会产生多次低带宽请求（即访存发散）；
编程要点：编程模型允许任意访问模式，但写高效 GPU 代码需让大多数访问不冲突。

Cache 设计

加 Cache 以节省访存带宽，并可做数据压缩/解压；
串行化 Cache 访问：先查 tag store 再访 data store，用延迟换功耗/灵活性——只访问含数据的那个 bank，并便于更复杂的 Cache 组织（如更高相联度）。

五、分支与发散

分支发散（Branch Divergence）处理

与向量处理器类似，但掩码在内部处理——每个 warp 维护一个栈，保存未走路径的 PC 与掩码。

条件分支上：

把当前掩码压栈；
把未走路径的掩码与 PC 压栈；
设置已走路径的掩码。

已走路径结束：弹出未走路径的掩码与 PC 并执行。未走路径结束：弹出分支前的原始掩码。若某掩码全 0，则跳过该块。

示例：4 线程/warp、初始掩码 1111，嵌套 if/else 通过压栈/弹栈逐层管理掩码。优化：若分支全部走同一方向，可跳过另一路径。

分支发散与加锁

if (condition[i]) {
    while (locked(map0[i])) {}
    lock(locks[map0[i]]);
} else {
    unlock(locks[map1[i]]);
}

基于 warp 的实现可能死锁：因为 warp 内线程必须一起前进，持锁线程可能无法继续，而等锁线程又卡住整个 warp。

六、GPU 编程环境与线程模型

编程语言

CUDA（仅 Nvidia）：类 C 语言，运行在 GPU 上；库有 cuDNN、cuBLAS、cuFFT；
OpenCL（开放标准）：类 C，可跑在 GPU/CPU/FPGA 上，通常优化程度不如 CUDA。

CUDA 线程模型

SPMD（Single-Program Multiple Data）模型：

每个上下文是一个 thread（有寄存器、有局部内存）；
并行线程打包成 block（有共享内存，线程用 barrier 同步，block 运行到结束或中止）；
grid 包含独立的 block（可并发执行，共享全局内存，但 block 间同步能力有限）。

与向量术语的对应

CUDA	向量
Thread	标量循环的 1 次迭代（向量循环中的 1 个元素）
Block	向量化循环体（示例中 VL=256）
Grid	可向量化的循环

DAXPY 示例：CUDA 代码每个 block 启动 256 个线程。

七、执行、调度与同步

核函数执行（Kernel Execution）

把输入数据从 CPU 内存传到 GPU 内存；
启动核函数（grid）；
等待核函数结束（若同步）；
把结果传回 CPU 内存。

数据传输可能主导执行时间；集成 GPU 配统一地址空间可免拷贝，但 CPU 与 GPU 会争用内存。

硬件调度

grid 可由 CPU 或 GPU 启动（来自多个 CPU 线程/进程的工作）；
由硬件单元把 grid 调度到各 SM 上，基于优先级；
多级调度：活跃 grid 数有限，其余排队/暂停。

同步（Synchronization）

线程块内 barrier 同步（__syncthreads()）：用 warp 分组简化跟踪，计数器记录已到达 barrier 的 warp 数；
对全局内存的原子操作（读-改-写：add、exchange、compare-and-swap……）在内存控制器或 L2 处执行；
块间同步能力有限：不能等待其他 block 完成。

八、ISA、编译与系统级问题

ISA 与编译

GPU 微架构与指令集变化非常频繁；
为兼容性：编译器生成中间伪汇编（如 Nvidia PTX），GPU 驱动在运行时 JIT 针对具体微架构生成代码；
实践中性能可移植性差——代码常针对特定 GPU 架构调优。

系统级问题

指令语义：异常处理；
调度：每个核函数不可抢占（但可中止），资源管理与调度由 GPU 驱动负责，对 OS 不透明；
内存管理：早期 GPU 无虚拟内存；近期支持基本虚拟内存（grid 间保护，无分页）；独立 GPU 需主机—设备拷贝；最新 GPU 才支持分页。

九、现代 GPU 实例与性能

多线程多核芯片实例

Nvidia Grace/Hopper GH100（2022）：144 个 SM、60MB 共享 L2、12 个内存控制器、3 TB/s（HBM3）、增强张量核（tensor core）、图形固定功能逻辑（纹理单元、光栅操作）、新层级线程块簇（thread block cluster）。

CPU vs GPU 性能

来源 Stanford CS231n：（部分优化的）CPU 与 CUDA 库（cuDNN）之比，体现 GPU 在 DNN 类负载上的显著优势。

小结

GPU 用 SIMT + 大规模多线程隐藏延迟，以 warp 为发射粒度；
存储层级（私有/共享/全局 + Cache）与访存合并是性能关键，发散是主要敌人；
CUDA 的 thread/block/grid 与向量的 element/loop-body/loop 一一对应；
ISA 频繁变化靠 PTX + JIT 维持兼容，但牺牲了性能可移植性。

下一讲：Security

L21：图形处理器（GPUs） ​

一、为什么研究 GPU ​

二、图形处理器时间线 ​

GPU 在现代系统中的两种形态 ​

三、SIMT 执行模型（Single Instruction Multiple Thread） ​

多线程 + SIMT ​

流式多处理器（SM）概览 ​

Little's Law（再次出现） ​

上下文大小 vs 上下文数目 ​

四、GPU 的多种存储类型 ​

访存发散（Memory Access Divergence） ​

Cache 设计 ​

五、分支与发散 ​

分支发散（Branch Divergence）处理 ​

分支发散与加锁 ​

六、GPU 编程环境与线程模型 ​

编程语言 ​

CUDA 线程模型 ​

与向量术语的对应 ​

七、执行、调度与同步 ​

核函数执行（Kernel Execution） ​

硬件调度 ​

同步（Synchronization） ​

八、ISA、编译与系统级问题 ​

ISA 与编译 ​

系统级问题 ​

九、现代 GPU 实例与性能 ​

多线程多核芯片实例 ​

CPU vs GPU 性能 ​

小结 ​