附录 B 图形与计算 GPU (Graphics and Computing GPUs)

核心结论

GPU 历史：从固定管线（1990s）→ 可编程顶点 / 像素着色器（2000s）→ 通用计算（GPGPU，2006+）→ 现代 SIMT 架构（2010+）。
统一架构：现代 GPU 用统一着色器核阵列同时处理图形与计算；CUDA / OpenCL 提供编程接口。
SIMT 执行：32 个线程组成 warp 同步执行；分支发散时串行执行（性能损失）。
内存层级：寄存器 / 共享内存 / L1 / L2 / GDDR（HBM / GDDR6X）；高带宽（~TB/s）是 GPU 优势核心。
CUDA 编程模型：host（CPU）+ device（GPU）；grid → block → thread；global / device / shared 等修饰符。
GPU vs CPU：GPU 数千小核、SIMT、高带宽、低延迟差；CPU 少数大核、超标量、大 cache。

本章主旨

附录 B 是 GPU 专章——把第 6 章的并行处理器扩展到 GPU。GPU 是现代图形 / ML / 科学计算的核心硬件；理解 GPU 架构（SIMT、内存层级、CUDA 模型）才能写出高性能 GPU 代码。NVIDIA GPU 是商业标准（Volta / Ampere / Hopper 架构）；AMD / Intel / Apple 也用类似 SIMT 模型。

一、核心概念

本章围绕 6 个核心概念展开：GPU 历史 → 统一架构 → SIMT → 内存层级 → CUDA 模型 → GPU vs CPU。

概念	定义 + 重要性	实现提示
GPU 历史	固定管线 → 可编程顶点 → 可编程像素 → 统一着色器 → SIMT 通用计算。	§B.1；理解 GPU 演进理解现代 SIMT 的设计选择。
统一着色器架构	顶点 / 像素 / 计算用同一组着色器核；CUDA / OpenCL 把 GPU 暴露为通用计算平台。	§B.2；现代 GPU 都是统一架构（2006+）。
SIMT 执行	32 线程组成 warp；同步执行同一条指令；分支发散时 warp 串行。	§B.3；CUDA 编程必须考虑 warp 对齐。
内存层级	寄存器 / shared / L1 / L2 / HBM / GDDR；HBM 带宽 ~3 TB/s（远超 CPU DRAM）。	§B.5；GPU 性能 = 计算 + 带宽，必须平衡。
CUDA 编程模型	host (CPU) + device (GPU)；grid → block → thread；global 等修饰符。	§B.3-§B.5；CUDA 是 GPU 编程事实标准。
GPU vs CPU	GPU 数千小核、高吞吐、低延迟；CPU 少数大核、低延迟、高单核性能。	§B.6；异构计算 = CPU 串行 + GPU 并行。

概念

定义 + 重要性

实现提示

GPU 历史

固定管线 → 可编程顶点 → 可编程像素 → 统一着色器 → SIMT 通用计算。

§B.1；理解 GPU 演进理解现代 SIMT 的设计选择。

统一着色器架构

顶点 / 像素 / 计算用同一组着色器核；CUDA / OpenCL 把 GPU 暴露为通用计算平台。

§B.2；现代 GPU 都是统一架构（2006+）。

SIMT 执行

32 线程组成 warp；同步执行同一条指令；分支发散时 warp 串行。

§B.3；CUDA 编程必须考虑 warp 对齐。

内存层级

寄存器 / shared / L1 / L2 / HBM / GDDR；HBM 带宽 ~3 TB/s（远超 CPU DRAM）。

§B.5；GPU 性能 = 计算 + 带宽，必须平衡。

CUDA 编程模型

host (CPU) + device (GPU)；grid → block → thread；global 等修饰符。

§B.3-§B.5；CUDA 是 GPU 编程事实标准。

GPU vs CPU

GPU 数千小核、高吞吐、低延迟；CPU 少数大核、低延迟、高单核性能。

§B.6；异构计算 = CPU 串行 + GPU 并行。

二、详细笔记

2.1 GPU 简史 (Brief History)

What：GPU 从固定功能到可编程到通用计算的演进。

Why：理解 GPU 为何能"通用计算"——可编程 shader 的副产品。

How：

GPU 演进（§B.1）：

1990s：固定管线（顶点变换 → 光栅化 → 像素着色），无编程性。
2000s 早期：可编程顶点 / 像素着色器（汇编级）。
2000s 中后期：高级着色语言（Cg / HLSL / GLSL）。
2006+：GPGPU 兴起（CUDA 1.0 / OpenCL 1.0）。
2010+：统一着色器架构（CUDA core / Stream processor）。

GPGPU 的兴起

2006 年 NVIDIA 发布 CUDA 1.0——第一个广泛使用的 GPGPU 平台。
关键洞察：GPU 的高吞吐 + 高带宽适合"数据并行 + 计算密集"任务（如 ML / 物理仿真）。
早期"显卡"变"加速器"——NVIDIA 股价 2006-2024 涨 100 倍。

When：选 GPU / 选加速器（CUDA / ROCm / oneAPI）；理解 GPU 编程模型。

Example：NVIDIA H100（2023）有 132 SM × 128 CUDA 核 = 16896 核；FP16 算力 1979 TFLOPS。

2.2 统一着色器架构 (Unified Shader Architecture)

What：现代 GPU 用同一组着色器核处理顶点 / 像素 / 计算任务。

Why：硬件利用率最高——不再有"顶点单元空闲而像素单元忙"。

How：

着色器核：执行标量 / 向量指令；NVIDIA 用 CUDA core，AMD 用 Stream processor。
统一分配器：硬件动态调度线程到空闲核。
SM（Streaming Multiprocessor）：一组 CUDA core + 共享内存 + 寄存器堆。

统一架构的工程优势

硬件利用率：负载均衡（任何任务都可占用所有核）。
编程模型：CUDA 屏蔽图形与计算的差异。
功耗：减少专用单元数量；降低晶体管冗余。

When：GPU 选型；理解 GPU 性能模型。

Example：NVIDIA Volta 架构（2017）首次引入独立 Tensor Core 用于 AI 加速；与 CUDA core 并存。

2.3 SIMT 执行 (SIMT Execution)

What：32 线程组成 warp，同步执行同一指令；分支发散时串行。

Why：理解 SIMT 才能写出 warp 高效的 CUDA 代码（避免分支发散）。

How：

SIMT 模型（§B.3）：

Warp：32 线程同步执行。
SIMD：单条指令处理 32 数据。
Branch divergence：warp 内线程走不同分支时串行执行（性能损失）。

性能优化：

避免 warp 内分支发散（用 __shfl_sync 同步或重排数据）。
warp 大小对齐（典型 32 / 64 / 128）。
occupancy 优化（active warps / max warps 比例）。

warp divergence 的性能影响

例：if-else 中，warp 内一半线程 if 分支、一半 else 分支。
硬件：先执行 if 分支（mask out else），再执行 else 分支（mask out if）。
代价：原本 1 个 cycle 的分支变成 2 个 cycle——性能减半。

When：CUDA kernel 优化；shader 性能分析。

Example：if (threadIdx.x % 2 == 0) 在 warp 内分支发散；改用 if (threadIdx.x < 16) 时所有线程同分支。

2.4 GPU 内存层级 (GPU Memory Hierarchy)

What：GPU 多级内存——寄存器 / shared / L1 / L2 / HBM；高带宽是 GPU 优势。

Why：GPU 性能 = 计算 + 带宽；理解内存层级才能写出 bandwidth-bound / compute-bound 代码。

How：

GPU 内存层级（§B.5）：

寄存器：每线程私有；最快（1 cycle）；瓶颈是寄存器数（限制 occupancy）。
Shared memory：每 block 共享；on-chip；显式管理（shared）。
L1 cache：每 SM 私有；自动管理。
L2 cache：所有 SM 共享；自动管理。
HBM / GDDR：device memory；带宽 1-3 TB/s（HBM3 / GDDR6X）。

GPU 内存的关键差异

带宽：HBM3 ~3 TB/s（远超 CPU DRAM 50-100 GB/s）——这是 GPU 优势核心。
延迟：HBM ~400 cycles（vs DRAM 100-200 cycles）——略高于 CPU。
容量：H100 80 GB（HBM3）；远小于 CPU 系统内存（TB 级）。

When：GPU 编程；内存优化（coalesced access）；带宽 vs 计算权衡。

Example：矩阵乘法若访问非 coalesced，GPU 带宽利用率可能 < 10%；改 coalesced 后接近 90%。

2.5 CUDA 编程模型 (CUDA Programming Model)

What：host (CPU) + device (GPU)；grid → block → thread；CUDA C 扩展 C 语言。

Why：CUDA 是 GPU 编程事实标准；理解模型才能写高性能 CUDA 代码。

How：

CUDA 模型：

Grid：所有线程（一维 / 二维 / 三维）。
Block：一组线程（典型 128-1024）；可同步 / 共享内存。
Thread：基本执行单元；每个线程有自己的寄存器 / PC。

函数修饰符：

global：kernel 函数，从 host 调用，在 device 上执行。
device：device 函数，从 kernel 调用。
host：普通 host 函数（默认）。

CUDA 编程的关键决策

Block 大小：32 的倍数（warp 对齐）；通常 128-256。
Grid 大小：(数据量 + block 大小 - 1) / block 大小。
Shared memory：用于 block 内线程通信；典型 16-48 KB。
同步：__syncthreads() 同步 block 内所有线程。

When：写 CUDA kernel；性能调优；多 GPU 编程。

Example：

__global__ void saxpy(float *x, float *y, float a, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) y[i] = a * x[i] + y[i];
}
// 调用：saxpy<<<(n+255)/256, 256>>>(x, y, 2.0f, n);
----

2.6 GPU vs CPU (GPU vs CPU)

What：GPU 与 CPU 在架构 / 用途上互补。

Why：现代计算 = CPU 串行 + GPU 并行；理解两者差异才能选对硬件。

How：

GPU vs CPU 对比（§B.6）：

维度	CPU	GPU
核数	4-32 大核	1000-10000 小核
单核频率	3-5 GHz	1-2 GHz
单核性能	高	低
吞吐量	中	极高
延迟	低	高
内存带宽	50-100 GB/s	1-3 TB/s
内存容量	16-128 GB	8-80 GB
适用	串行 / 控制流	数据并行
编程	C/C++/Java	CUDA / OpenCL

维度

CPU

GPU

核数

4-32 大核

1000-10000 小核

单核频率

3-5 GHz

1-2 GHz

单核性能

高

低

吞吐量

中

极高

延迟

低

高

内存带宽

50-100 GB/s

1-3 TB/s

内存容量

16-128 GB

8-80 GB

适用

串行 / 控制流

数据并行

编程

C/C++/Java

CUDA / OpenCL

异构计算的工程模式

CPU + GPU 异构：CPU 处理控制流 / I/O / 串行逻辑；GPU 处理数据并行。
数据传输开销：CPU ↔ GPU 数据传输带宽受 PCIe / NVLink 限制（数十 GB/s）——批处理大块传输更高效。
AMD APU / Apple M1：CPU + GPU 共享统一内存——传输开销消失。

When：硬件选型；异构编程；性能调优。

Example：ML 训练用 GPU（数据并行计算密集）；推理用 CPU（延迟敏感）+ GPU（吞吐密集）。

三、关键图表

视觉图表

Figure 1. 图 B-1：GPU 架构（SM / 内存层级）

Figure 2. 图 B-2：CUDA 编程模型（grid / block / thread）

Figure 3. 图 B-3：warp 分支发散

Figure 4. 图 B-4：GPU vs CPU roofline

非可视化条目

非可视化条目（表 / 算法）

编号

内容摘要

表 B.1

GPU 架构演进（固定管线 → 顶点 → 像素 → 统一）。

表 B.2

NVIDIA GPU 性能对比（V100 / A100 / H100）。

式 B-1 至 B-5

occupancy / roofline / 带宽计算。

列表 B.1-B.x

CUDA 示例代码（saxpy / matrix multiply / reduction）。

核心公式对照表

概念

公式

Occupancy

\(\text{Occupancy} = \frac{\text{Active warps}}{\text{Max warps per SM}}\)

Roofline

\(\text{Perf} = \min(\text{Peak FLOPS}, \text{Bandwidth} \times \text{AI})\)

加速比（GPU vs CPU）

\(\text{Speedup} = \frac{N_{\text{core GPU}}}{N_{\text{core CPU}}} \times \frac{f_{\text{GPU}}}{f_{\text{CPU}}} \times \text{utilization}\)

内存带宽

\(\text{Bandwidth} = \text{Bus width} \times \text{Rate}\)

四、思维导图

mindmap
  root((附录 B GPU))
    历史
      固定管线
      可编程着色器
      通用计算CUDA
    统一架构
      顶点像素计算
      共享着色器核
      动态调度
    SIMT
      32线程warp
      分支发散
      同步执行
    内存层级
      寄存器shared
      L1L2cache
      HBM高带宽
    CUDA
      gridblockthread
      hostdevice
      修饰符
    GPUvsCPU
      数千小核vs少数大核
      高吞吐vs低延迟
      数据并行vs串行

五、重点与易错点

warp 分支发散是性能杀手：避免 if (threadIdx.x % 2 == 0) 这类发散；用 if (threadIdx.x < 16) 等对齐分支。
Coalesced memory access：连续线程访问连续内存 → 合并为一次大带宽传输；非对齐 → 多次小带宽传输。
Occupancy 优化：每 SM active warps 越多，GPU 利用率越高；但太多反而 register pressure。
Host ↔ Device 数据传输开销：PCIe 带宽有限；尽量减少传输、批处理大块数据。
GPU 不适合分支密集任务：控制流分支多 → warp 发散严重 → 性能差。
CUDA 不是唯一 GPU 编程模型：OpenCL、ROCm（AMD）、Vulkan Compute、Metal（Apple）都是替代。
Tensor Core：NVIDIA Volta+ 的专用矩阵计算单元；ML 训练 / 推理必备；与 CUDA core 不同。
GPU 内存容量远小于 CPU：H100 80 GB vs 服务器 1+ TB；大数据训练需 CPU + NVMe offload。
跨章衔接：第 3 章 SIMD 是 GPU SIMT 的前身；第 6 章并行处理器是 GPU 的理论基础；附录 D 给出 ARM 等嵌入式 ISA 与 GPU 配合。