系统第 5 讲 / 共 17 讲

第 5 讲：GPU 架构与编程模型

1. GPU vs CPU 架构差异

要理解为什么 LLM 训练必须使用 GPU，首先需要理解 GPU 和 CPU 在设计哲学上的根本差异。

1.1 设计哲学

CPU 的设计目标是延迟优化（latency-optimized）：尽快完成单个任务。它有强大的控制逻辑（分支预测、乱序执行）、大容量的 cache、少量但很快的核心。

GPU 的设计目标是吞吐量优化（throughput-optimized）：同时处理大量简单任务。它有数千个小核心、相对较小的 cache，但拥有极高的内存带宽和计算吞吐量。

特性	CPU（如 Intel Xeon）	GPU（如 NVIDIA H100）
核心数	16-128	~16,000 CUDA cores
时钟频率	3-5 GHz	1.5-2 GHz
设计目标	低延迟	高吞吐
Cache 大小	~100 MB	~50 MB (L2)
内存带宽	~100 GB/s	~3 TB/s (HBM3)
内存容量	~1 TB	80 GB (HBM)
BF16 峰值	~1 TFLOPS	~989 TFLOPS (Tensor Core)

为什么深度学习适合 GPU

深度学习的核心运算是矩阵乘法，它具有高度的并行性：计算输出矩阵的每个元素可以独立进行。这种"大量独立的相同计算"正是 GPU 的强项。GPU 的数千个核心可以同时计算数千个输出元素，而 CPU 的十几个核心只能同时处理十几个。

2. GPU 硬件组成

2.1 Streaming Multiprocessor (SM)

SM（Streaming Multiprocessor）是 GPU 的基本处理单元。一个 GPU 包含多个 SM，每个 SM 包含多个计算核心和共享资源。

A100：108 个 SM
H100：132 个 SM

每个 SM 内部包含：

多组 CUDA cores（通用浮点计算）
Tensor Cores（专用矩阵运算加速器）
Shared Memory / L1 Cache（片上高速存储）
Register File（寄存器文件）
Warp Scheduler（线程调度器）

2.2 CUDA Cores vs Tensor Cores

CUDA Cores 是通用浮点计算单元，执行标量的加法和乘法运算。它们可以处理任意的浮点运算，但每个 cycle 只完成一次标量操作。

Tensor Cores 是专为矩阵运算设计的加速单元。一个 Tensor Core 在一个 cycle 内可以完成一次小矩阵乘加运算（如 4×4 的 FP16 矩阵乘法 + FP32 累加），吞吐量远高于 CUDA Cores。

Tensor Core 单次操作: D = A × B + C
A: 4×4 FP16, B: 4×4 FP16, C: 4×4 FP32 → D: 4×4 FP32

Tensor Core 的重要性

在 LLM 训练中，几乎所有的 FLOPs 来自矩阵乘法。Tensor Core 的矩阵乘法吞吐量比 CUDA Core 高出一个数量级。这意味着：如果你的代码没有使用 Tensor Core（例如没有使用 BF16/FP16 精度），你的 GPU 利用率会非常低。使用 PyTorch 的 autocast 或显式转换为 BF16 可以自动利用 Tensor Core。

3. Memory Hierarchy

GPU 有一个多级内存层次结构，理解它是优化 GPU 性能的关键。从快到慢（从小到大）：

层级	每 SM 大小	带宽	延迟	用途
Registers	256 KB	~20 TB/s	~1 cycle	每个线程的局部变量
Shared Memory / L1 Cache	128-228 KB	~19 TB/s	~30 cycles	同一 Block 内线程共享数据
L2 Cache	全局 40-50 MB	~12 TB/s	~200 cycles	所有 SM 共享的缓存
HBM (Global Memory)	80 GB	~2-3 TB/s	~400 cycles	模型参数、激活值等

可以看到，Registers 和 Shared Memory 的带宽比 HBM 高出约 10 倍，延迟低约 100 倍。这就是为什么减少 HBM 访问是 GPU 优化的核心策略。

# Memory hierarchy visualization (approximate)
#
# Registers:  ████████████████████████████  ~20 TB/s
# Shared Mem: ███████████████████           ~19 TB/s
# L2 Cache:   ████████████                  ~12 TB/s
# HBM:        ████                          ~2-3 TB/s
#
# The gap between compute speed and HBM bandwidth
# is the fundamental bottleneck in GPU programming!

4. 内存带宽与计算带宽

GPU 性能受两个"带宽"限制：计算带宽（compute bandwidth）和内存带宽（memory bandwidth）。对于任何给定的操作，瓶颈是其中较慢的那个。

4.1 Compute-bound vs Memory-bound

Compute-bound（计算受限）：操作的计算量很大， GPU 的计算单元成为瓶颈。典型例子：大矩阵乘法。

Memory-bound（内存受限）：操作需要频繁读写内存，但计算量相对较少，HBM 带宽成为瓶颈。典型例子：element-wise 操作（LayerNorm, ReLU, Softmax）。

# Example: Matrix multiplication vs Element-wise addition
# Both on a (4096, 4096) tensor in BF16

# Matrix multiply: C = A @ B
# - Reads: 2 * 4096² * 2 bytes = 64 MB
# - Computes: 2 * 4096³ = 137 GFLOP
# - Arithmetic intensity: 137G / 64M = ~2000 FLOP/byte → COMPUTE-BOUND

# Element-wise add: C = A + B
# - Reads: 2 * 4096² * 2 bytes = 64 MB
# - Writes: 4096² * 2 bytes = 32 MB
# - Computes: 4096² = 16.7 MFLOP
# - Arithmetic intensity: 16.7M / 96M = ~0.17 FLOP/byte → MEMORY-BOUND

5. Arithmetic Intensity 与 Roofline Model

5.1 Arithmetic Intensity（计算强度）

Arithmetic Intensity（AI）定义为每字节内存访问对应的浮点运算次数：

AI = FLOPs / Bytes Transferred (FLOP/Byte)

这是判断一个操作是 compute-bound 还是 memory-bound 的关键指标。

5.2 Roofline Model

Roofline Model 是分析 GPU 性能上限的经典框架。它将 GPU 的性能上限可视化为两条线的组合：

Memory Bandwidth Roof：当 AI 较低时，性能受限于内存带宽，上限 = AI × memory_bandwidth
Compute Roof：当 AI 较高时，性能受限于计算能力，上限 = peak_compute
Ridge Point：两条线的交点，AI = peak_compute / memory_bandwidth

# Roofline calculation for H100
peak_compute = 989e12    # BF16: 989 TFLOPS
memory_bw = 3.35e12      # HBM3: 3.35 TB/s

ridge_point = peak_compute / memory_bw
# = 989 / 3.35 ≈ 295 FLOP/Byte

# Operations with AI > 295 are compute-bound (good for GPU!)
# Operations with AI < 295 are memory-bound (need optimization)

# Typical arithmetic intensities:
# - Large matmul (4096x4096):     ~2000 FLOP/Byte → compute-bound ✓
# - LayerNorm:                    ~5 FLOP/Byte    → memory-bound ✗
# - Softmax:                      ~5 FLOP/Byte    → memory-bound ✗
# - Element-wise activation:      ~1 FLOP/Byte    → memory-bound ✗

实践启示

Transformer 中大约 60-70% 的时间花在矩阵乘法上（compute-bound），30-40% 花在 memory-bound 操作上（LayerNorm, Softmax, activation, 数据搬运等）。优化 memory-bound 操作的关键策略是 kernel fusion——将多个操作合并为一个 kernel，减少 HBM 的读写次数。这就是 Flash Attention 等优化技术的核心思想。

6. CUDA 编程模型

CUDA（Compute Unified Device Architecture）是 NVIDIA 的 GPU 编程模型。理解 CUDA 的线程组织方式对于理解 GPU 性能优化至关重要。

6.1 线程层次结构

CUDA 将线程组织为三个层次：

层次	说明	硬件映射
Thread	最小执行单元，每个 thread 执行相同的代码（SIMT）	CUDA Core / Thread Lane
Block（Thread Block）	一组 threads（最多 1024），共享 Shared Memory	一个 Block 在一个 SM 上执行
Grid	一组 Blocks，组成完整的并行任务	分布在所有 SM 上

6.2 Warp：真正的执行单位

在硬件层面，Warp 才是真正的执行单位。一个 Warp 包含 32 个 threads，它们在同一时刻执行完全相同的指令（SIMT 模型）。如果 Warp 中的 threads 走了不同的分支路径（branch divergence），两个分支会被串行执行，导致性能损失。

# CUDA kernel example: vector addition
# (Python pseudocode showing the CUDA mental model)

# Grid:  [Block 0] [Block 1] [Block 2] ...
# Block: [Thread 0] [Thread 1] ... [Thread 255]
# Warp:  [Thread 0..31] [Thread 32..63] ...

def vector_add_kernel(A, B, C, N):
    """Each thread computes one element."""
    # Global thread index
    idx = blockIdx.x * blockDim.x + threadIdx.x
    if idx < N:
        C[idx] = A[idx] + B[idx]

# Launch configuration
N = 1_000_000
block_size = 256           # 256 threads per block
grid_size = (N + block_size - 1) // block_size  # ceil(N / 256)

vector_add_kernel<<>>(A, B, C, N)

6.3 Occupancy

Occupancy 表示每个 SM 上活跃的 Warp 数占最大支持 Warp 数的比例。高 occupancy 有助于隐藏内存延迟：当一个 Warp 等待内存数据时， GPU 可以切换到另一个活跃的 Warp 继续执行。

影响 occupancy 的因素：

每个 thread 的 register 使用量：使用太多 register 会减少每个 SM 能容纳的 Warp 数
每个 Block 的 Shared Memory 使用量：同理
Block size：太小或太大都可能降低 occupancy

7. GPU 代际对比

规格	A100 (80GB SXM)	H100 (SXM)
架构	Ampere	Hopper
工艺	7nm	4nm
SM 数量	108	132
HBM 容量	80 GB (HBM2e)	80 GB (HBM3)
HBM 带宽	2.0 TB/s	3.35 TB/s
L2 Cache	40 MB	50 MB
FP32 峰值	156 TFLOPS	495 TFLOPS
BF16 Tensor Core 峰值	312 TFLOPS	989 TFLOPS
FP8 Tensor Core 峰值	—	1979 TFLOPS
GPU 间互联	NVLink 3 (600 GB/s)	NVLink 4 (900 GB/s)
TDP	400W	700W

7.1 关键趋势

从 A100 到 H100 的演进反映了以下趋势：

计算速度增长快于内存带宽：BF16 计算能力增长了 3.2 倍，但 HBM 带宽只增长了 1.7 倍。这意味着 memory-bound 操作的重要性越来越高
低精度支持：H100 新增 FP8 支持，在支持的场景下再翻倍吞吐
互联带宽增长：NVLink 从 600 GB/s 到 900 GB/s，对多 GPU 训练至关重要
功耗增长：从 400W 到 700W，散热和供电成为数据中心的挑战

Ridge Point 的变化

A100 的 ridge point ≈ 312T/2.0T = 156 FLOP/Byte，H100 的 ridge point ≈ 989T/3.35T = 295 FLOP/Byte。H100 的 ridge point 几乎是 A100 的 2 倍。这意味着在 H100 上，更多的操作变成了 memory-bound，kernel fusion 等内存优化技术变得更加重要。

8. 实际编程考量

8.1 数据传输

CPU 和 GPU 之间的数据传输（通过 PCIe 或 NVLink）是一个潜在的瓶颈。 PCIe 4.0 x16 的理论带宽约 32 GB/s，而 HBM 带宽是 2-3 TB/s——差了约 100 倍。

尽量减少 CPU-GPU 数据传输
使用 pin_memory=True 加速 data loading
使用 non_blocking=True 实现异步传输
数据预处理尽量在 GPU 上完成

8.2 Kernel Launch Overhead

每次启动一个 CUDA kernel 都有约 5-10 微秒的 overhead。如果你的程序由大量小 kernel 组成，launch overhead 可能成为瓶颈。这也是为什么 kernel fusion（合并多个小操作为一个大 kernel）很重要。

# Bad: many small kernels
x = x + bias       # kernel 1
x = layer_norm(x)  # kernel 2
x = relu(x)        # kernel 3
x = dropout(x)     # kernel 4

# Better: fused kernel
x = fused_bias_layernorm_relu_dropout(x, bias)  # single kernel

8.3 PyTorch 中的 GPU 编程

# Check GPU availability
print(torch.cuda.is_available())
print(torch.cuda.device_count())
print(torch.cuda.get_device_name(0))

# Move data to GPU
device = torch.device('cuda:0')
tensor = torch.randn(1024, 1024, device=device)

# Monitor memory
print(f"Allocated: {torch.cuda.memory_allocated()/1e9:.2f} GB")
print(f"Max allocated: {torch.cuda.max_memory_allocated()/1e9:.2f} GB")

# Synchronize for accurate timing
torch.cuda.synchronize()
start = time.time()
result = model(input)
torch.cuda.synchronize()  # Wait for GPU to finish
elapsed = time.time() - start

异步执行

GPU 操作是异步的——torch.cuda.synchronize() 会阻塞 CPU 直到所有排队的 GPU 操作完成。在做性能测量时必须调用 synchronize()，否则测到的只是 CPU 端排队的时间，而非 GPU 实际执行的时间。

9. 本讲总结

GPU 是吞吐量优化的处理器，拥有数千个核心和极高的内存带宽，非常适合深度学习的大规模并行计算
GPU 由多个 SM 组成，每个 SM 包含 CUDA Cores（通用计算）和 Tensor Cores（矩阵加速）
Memory Hierarchy：Registers → Shared Memory → L2 Cache → HBM，带宽依次递减约 10 倍
Arithmetic Intensity 决定了操作是 compute-bound 还是 memory-bound
Roofline Model 提供了性能分析的理论框架
CUDA 的线程层次：Thread → Warp (32 threads) → Block → Grid
从 A100 到 H100，计算能力增长快于内存带宽，使得内存优化越来越重要