系统 第 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 GHz1.5-2 GHz
设计目标低延迟高吞吐
Cache 大小~100 MB~50 MB (L2)
内存带宽~100 GB/s~3 TB/s (HBM3)
内存容量~1 TB80 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 包含多个计算核心和共享资源。

每个 SM 内部包含:

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 的性能上限可视化为两条线的组合:

# 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 的因素:

7. GPU 代际对比

规格A100 (80GB SXM)H100 (SXM)
架构AmpereHopper
工艺7nm4nm
SM 数量108132
HBM 容量80 GB (HBM2e)80 GB (HBM3)
HBM 带宽2.0 TB/s3.35 TB/s
L2 Cache40 MB50 MB
FP32 峰值156 TFLOPS495 TFLOPS
BF16 Tensor Core 峰值312 TFLOPS989 TFLOPS
FP8 Tensor Core 峰值—1979 TFLOPS
GPU 间互联NVLink 3 (600 GB/s)NVLink 4 (900 GB/s)
TDP400W700W

7.1 关键趋势

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

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 倍。

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. 本讲总结