9.2 GPU 内存层级

在上一节中，我们概览了 GPU 的整体架构——大量流式多处理器（SM）协同工作，追求极致的吞吐量。然而，再强大的计算单元，如果数据供给不上，也只能空转等待。GPU 内存层级（Memory Hierarchy）正是解决"如何高效地把数据喂给计算单元"这一核心问题的硬件基础。

本节将从物理介质出发，逐层剖析 GPU 的内存结构，并引入 Roofline 模型，为后续的性能优化技术（如分块、算子融合）提供理论框架。

9.2.1 片上内存与片下内存：SRAM vs HBM

GPU 的存储系统在物理上可以划分为两大类：片上内存（On-chip Memory）和片下内存（Off-chip Memory）。两者在制造工艺、性能特征和设计目标上存在根本差异。

片上内存（SRAM）

片上内存使用静态随机存取存储器（SRAM）工艺制造，直接集成在 GPU 芯片的硅片上。SRAM 的每个存储单元由 6 个晶体管构成，不需要刷新操作，因此访问速度极快、延迟极低。但正因为每个存储单元占用的晶体管数量多，SRAM 的面积成本和功耗成本都很高，无法做到大容量。在 GPU 中，寄存器文件、共享内存、L1 缓存和 L2 缓存都属于片上 SRAM。

片下内存（HBM / GDDR）

片下内存位于 GPU 芯片外部，是独立的 DRAM 芯片。在高端 GPU（如 NVIDIA A100、H100）中，片下内存采用高带宽内存（High Bandwidth Memory, HBM）技术；在消费级 GPU 中，则通常采用 GDDR6/GDDR6X。

HBM 的核心设计思路与传统 DRAM 截然不同：

垂直堆叠（3D Stacking）：多层 DRAM 芯片通过硅通孔（TSV）技术垂直堆叠，缩短信号传输距离，降低功耗。
超宽总线：单个 HBM 堆栈的总线宽度达 1024 位，而传统 DDR4 仅 64 位。一块 A100 GPU 集成 6 个 HBM2e 堆栈，总线宽度达 6144 位。
以宽度换频率：HBM 的时钟频率相对较低，但凭借极宽的总线，总带宽仍远超传统方案。

尽管 HBM 的带宽已达 TB/s 量级，但与片上 SRAM 相比，其延迟仍高出一到两个数量级。这一差距构成了 GPU 编程中最核心的矛盾——计算能力的增长速度远超内存带宽的增长速度，即所谓的"内存墙"（Memory Wall）问题。

特性	片上 SRAM	片下 HBM/GDDR
制造工艺	6T SRAM，集成于芯片	DRAM，独立芯片
访问延迟	1-30 个时钟周期	200-800 个时钟周期
带宽量级	数十 TB/s（聚合）	1-3 TB/s
典型容量	KB-MB 级	数十 GB 级
功耗/bit	高	低
面积/bit	大	小

9.2.2 逻辑层次：从寄存器到全局内存

从程序员和编译器的视角来看，GPU 的内存系统呈现为一个多层金字塔结构。越靠近金字塔顶端，速度越快、容量越小、离计算单元越近；越靠近底部，则相反。

第一层：寄存器（Registers）

寄存器是 GPU 中速度最快的存储单元，访问延迟仅约 1 个时钟周期。

物理位置：位于每个 SM 内部。
作用域：每个线程拥有私有的寄存器组，用于存放局部变量和中间计算结果。
容量：单个 SM 的寄存器文件总量通常为 256 KB（如 A100），由该 SM 上运行的所有线程瓜分。单个线程可用的寄存器数量取决于 SM 同时驻留的线程数。
特点：寄存器由编译器自动分配，程序员一般不直接管理。当线程使用的寄存器数超过上限时，数据会被"溢出"（spill）到较慢的局部内存中，导致性能下降。

第二层：共享内存 / L1 缓存（Shared Memory / L1 Cache）

共享内存和 L1 缓存共享同一块片上 SRAM 硬件，访问延迟约 20-30 个时钟周期。

物理位置：位于每个 SM 内部。
作用域：同一线程块（Thread Block）内的所有线程可以共享访问。
容量：每个 SM 通常拥有 128-228 KB 的可配置空间（以 A100 为例为 192 KB），可在共享内存和 L1 缓存之间灵活划分。
共享内存的关键特征：
- 程序员显式管理：开发者可以通过 CUDA 代码明确控制哪些数据被加载到共享内存、何时加载、如何读取。这种显式控制是分块（Tiling）技术的基础。
- 线程间通信：同一线程块内的线程可以通过共享内存快速交换数据，配合同步原语（如 __syncthreads()）实现协作计算。
L1 缓存的特征：
- 硬件自动管理：对全局内存的访问会被硬件自动缓存到 L1，对程序员透明。
- 与 CPU 的 L1 缓存功能类似，但因 GPU 更依赖大规模并行来隐藏延迟，L1 的重要性相对不同。

第三层：L2 缓存（L2 Cache）

物理位置：位于 GPU 芯片上，但在所有 SM 之外，是一个全局共享资源。
作用域：被所有 SM 共享，对程序员透明。
容量：现代 GPU 的 L2 缓存通常在数 MB 到数十 MB 量级（A100 为 40 MB，H100 为 50 MB）。
延迟：约 200-300 个时钟周期，介于 L1 和全局内存之间。
作用：充当全局内存的最后一级缓存，减少对 HBM 的直接访问。当多个 SM 访问相同的全局内存数据时，L2 缓存可以有效降低重复读取的开销。

第四层：全局内存（Global Memory）

物理位置：位于 GPU 芯片外部的 HBM 或 GDDR 芯片。
作用域：可被所有 SM 上的所有线程访问，也可被主机 CPU 通过 PCIe/NVLink 读写。
容量：GB 级别（A100 为 80 GB HBM2e，H100 为 80 GB HBM3）。
延迟：200-800 个时钟周期（取决于是否命中 L2 缓存）。
带宽：虽然单次访问延迟高，但 HBM 的聚合带宽可达 TB/s 级别（A100 为 2 TB/s，H100 为 3.35 TB/s）。
作用：存储模型参数、输入数据、输出结果等所有大规模数据。所有跨线程块的数据通信必须通过全局内存完成。

下表汇总了典型 GPU（以 NVIDIA A100 为参考）各层内存的关键参数：

层级	类型	每 SM 容量	全芯片容量	访问延迟（周期）	带宽量级	管理方式
寄存器	SRAM	256 KB	~27 MB	~1	~数十 TB/s	编译器
共享内存/L1	SRAM	192 KB	~20 MB	20-30	~19 TB/s	程序员/硬件
L2 缓存	SRAM	—	40 MB	200-300	~5 TB/s	硬件
全局内存 (HBM)	DRAM	—	80 GB	200-800	2 TB/s	程序员

从表中可以清晰看到：从寄存器到全局内存，容量增大了约 3000 倍，而延迟也增大了 200-800 倍。这种巨大的速度-容量权衡，决定了高性能 GPU 编程的核心策略——尽一切可能让数据留在离计算单元最近的高速存储中。

9.2.3 内存带宽与算术强度

理解了内存层级之后，一个自然的问题是：对于一个具体的计算任务，它的性能瓶颈到底在哪里？是计算单元不够快，还是数据搬运不够快？回答这个问题需要引入算术强度（Arithmetic Intensity）的概念。

算术强度的定义

算术强度衡量的是一个计算任务中"计算量"与"内存访问量"的比值：

Arithmetic Intensity = \frac{FLOPs（浮点运算次数）}{Bytes（内存读写的字节数）}

单位为 FLOPs/Byte。这个比值刻画了"每从内存搬运一个字节的数据，能执行多少次浮点运算"。

直觉理解

算术强度低的操作（如逐元素加法、ReLU 激活函数）：每个元素只做 1 次运算，却需要从内存读入再写出，算术强度约为 0.25 FLOPs/Byte（以 FP32 为例，读 4 字节 + 写 4 字节 = 8 字节，做 1 次运算）。这类操作的性能完全受限于内存带宽。
算术强度高的操作（如矩阵乘法）：对于两个 $n \times n$ 的矩阵相乘，总计算量为 $O (n^{3})$ ，而内存访问量为 $O (n^{2})$ ，算术强度为 $O (n)$ 。当 $n$ 足够大时，算术强度可以非常高，此时性能受限于计算能力而非内存带宽。

典型操作的算术强度

操作	计算量	内存访问量	算术强度
逐元素加法 (FP32)	$n$	$12 n$ 字节 (读 2 写 1)	~0.08 FLOPs/Byte
向量点积 (FP32)	$2 n$	$8 n$ 字节	~0.25 FLOPs/Byte
矩阵乘法 $n \times n$ (FP32)	$2 n^{3}$	$12 n^{2}$ 字节	$n / 6$ FLOPs/Byte
批量矩阵乘法（大 $n$ ）	$2 n^{3}$	$12 n^{2}$ 字节	$n / 6$ FLOPs/Byte

矩阵乘法之所以在 GPU 上能够高效执行，正是因为其算术强度随问题规模线性增长，能够充分利用 GPU 强大的计算能力。而逐元素操作无论问题多大，算术强度始终很低，注定受限于内存带宽。

9.2.4 Roofline 模型

有了算术强度的概念，我们就可以引入 Roofline 模型——一个简洁而强大的性能分析框架，用于判断任何计算任务的性能瓶颈。

模型定义

Roofline 模型将程序的可达性能（Attainable Performance）表示为其算术强度的函数：

P = min (π, β \cdot I)

其中：

$P$ 为程序的可达性能（单位：FLOP/s）
$π$ 为硬件的峰值计算性能（单位：FLOP/s），例如 A100 的 FP32 峰值性能为 19.5 TFLOP/s
$β$ 为硬件的峰值内存带宽（单位：Byte/s），例如 A100 的 HBM 带宽为 2 TB/s
$I$ 为程序的算术强度（单位：FLOP/Byte）

图形解读

在以算术强度 $I$ 为横轴（对数刻度）、可达性能 $P$ 为纵轴（对数刻度）的坐标系中，Roofline 模型呈现为两条线段的组合——形如一个"屋顶"：

斜线部分（屋檐）： $P = β \cdot I$ 。在这个区域，性能随算术强度线性增长，斜率等于内存带宽 $β$ 。程序处于**内存受限（Memory-bound）**状态——计算单元在等待数据。
水平线部分（屋顶）： $P = π$ 。一旦算术强度超过某个阈值，性能不再增长，被峰值计算性能封顶。程序处于**计算受限（Compute-bound）**状态——数据供给充足，但计算单元已满载运行。

两条线的交点处的算术强度称为脊点（Ridge Point）：

I^{*} = \frac{π}{β}

脊点的物理意义是：当程序的算术强度恰好等于 $I^{*}$ 时，计算能力和内存带宽同时被完全利用，硬件资源得到最大化发挥。

以 A100 为例：

I^{*} = \frac{19.5 TFLOP/s}{2 TB/s} = 9.75 FLOP/Byte

这意味着，在 A100 上，一个程序的算术强度至少需要达到约 10 FLOP/Byte，才能不被内存带宽拖后腿。

用 Roofline 模型诊断性能

将具体操作标注在 Roofline 图上，可以直观地看出优化方向：

逐元素操作（算术强度 ~0.1）：远在脊点左侧，深度内存受限。即使 GPU 的计算能力再强 10 倍也不会变快——瓶颈在带宽。优化方向是减少内存访问（算子融合、减少中间结果的读写）。
小规模矩阵乘法（ $n = 64$ ，算术强度 ~10）：接近脊点，开始过渡到计算受限。
大规模矩阵乘法（ $n = 4096$ ，算术强度 ~680）：远在脊点右侧，完全计算受限。优化方向是提升计算效率（使用 Tensor Core、低精度计算）。

Roofline 模型的实践意义

Roofline 模型揭示了 GPU 性能优化的两条基本路径：

对于内存受限的操作：目标是提高算术强度，将性能点在图上向右移动。具体手段包括：
- 分块（Tiling）：将数据加载到共享内存中反复使用，减少对全局内存的重复访问。
- 算子融合（Operator Fusion）：将多个连续的逐元素操作合并为一个 kernel，避免中间结果写回全局内存再读出。
- 重计算（Recomputation）：在某些场景下，与其把中间结果存入内存再读回，不如在需要时重新计算——用计算换内存访问。
对于计算受限的操作：目标是逼近峰值计算性能，将性能点在图上向上移动。具体手段包括：
- 利用 Tensor Core：使用专用矩阵乘法单元，比通用 CUDA Core 快数倍。
- 低精度计算：从 FP32 切换到 FP16/BF16 甚至 INT8，可成倍提升峰值 FLOP/s（同时需注意数值精度）。
- 优化指令调度：减少空闲周期，提高计算单元利用率。

Roofline 模型是一个上界分析工具——它给出的是理论上限，实际性能可能因为各种原因（缓存未命中、bank 冲突、warp 发散等）低于 Roofline 预测。但它提供了判断"应该往哪个方向优化"的清晰指南，避免在错误的方向上浪费精力。

9.2.5 内存层级视角下的性能优化原则

综合以上讨论，我们可以提炼出几条围绕内存层级的核心优化原则：

数据局部性最大化

高性能 GPU 代码的首要目标是让数据尽可能停留在靠近计算单元的高速存储中。具体而言：

能用寄存器的，不用共享内存。
能用共享内存的，不访问全局内存。
必须访问全局内存时，确保访问模式对缓存友好（连续、对齐）。

全局内存访问最小化

全局内存是 GPU 中最昂贵的操作之一。每一次不必要的全局内存读写，都意味着数百个时钟周期的浪费。优化的首要目标是最小化全局内存的读写次数。

内存访问合并

当必须访问全局内存时，应确保同一个 Warp（32 个线程）的访问地址是连续且对齐的，从而触发硬件的合并读取（Coalesced Access）机制。一次合并的 128 字节读取，远比 32 次分散的 4 字节读取高效。

计算-访存比的自觉

在设计算法和选择实现策略时，始终关注操作的算术强度。对于低算术强度的操作，优先考虑算子融合等减少内存访问的技术；对于高算术强度的操作，优先考虑计算效率的优化。

本节小结

GPU 的内存系统是一个精心设计的多级层次结构，从最快的寄存器到最慢的全局内存，容量与速度之间存在数百倍的权衡。理解这个层次结构是所有 GPU 性能优化的起点。

核心要点回顾：

物理二分法：片上 SRAM 快而小，片下 HBM 大而慢，两者之间的速度鸿沟构成了 GPU 编程的核心挑战。
四级逻辑层次：寄存器（~1 周期）→ 共享内存/L1（~20-30 周期）→ L2 缓存（~200-300 周期）→ 全局内存（~200-800 周期），每一级的容量大约增长 10-100 倍，延迟也相应增大。
算术强度是衡量操作计算密集度的核心指标，定义为每字节内存访问所对应的浮点运算次数。
Roofline 模型通过 $P = min (π, β \cdot I)$ 将硬件参数与程序特征统一在一个分析框架中，明确区分内存受限与计算受限两种瓶颈状态，为优化方向提供理论指引。
几乎所有 GPU 优化技术——分块、算子融合、重计算、低精度计算——都可以从内存层级和 Roofline 模型的角度获得统一的理解：要么减少内存访问以提高算术强度，要么提高计算效率以逼近峰值性能。

在下一节中，我们将具体讨论如何利用这些原则来优化大模型训练中的关键计算模式。

9.2 GPU 内存层级 ​

9.2.1 片上内存与片下内存：SRAM vs HBM ​

9.2.2 逻辑层次：从寄存器到全局内存 ​

9.2.3 内存带宽与算术强度 ​

9.2.4 Roofline 模型 ​

9.2.5 内存层级视角下的性能优化原则 ​