第三课：GPU 的存储层次

导语：这一部分讲 GPU 的存储层次，以及为什么同一个程序只是改变数据访问方式，性能就可能相差数倍。

GPU 性能并不只取决于“有多少个 CUDA Core”。

很多 GPU 程序真正的瓶颈不是计算，而是：

数据能不能及时送到计算单元。

因此，学习 GPU 架构必须理解它的存储层次。

可以先把 GPU 的主要存储结构理解成：

速度快、容量小
        ↑
      寄存器
    Shared Memory
       L1 Cache
       L2 Cache
   Global Memory 显存
        ↓
速度慢、容量大

整体规律是：

• 越靠近计算单元，速度越快；
• 越靠近计算单元，容量越小；
• 越远离计算单元，访问延迟越高；
• 越远离计算单元，容量通常越大。

---

一、先理解“计算”和“取数据”的关系

考虑一条简单语句：

C[i] = A[i] + B[i];

表面上看，只需要一次加法。

但硬件实际需要完成：

1. 从显存读取 A[i]
2. 从显存读取 B[i]
3. 执行加法
4. 把结果写回显存 C[i]

真正的计算只有一次加法，但却有三次显存访问。

因此，它很可能不是“算力受限”，而是“访存受限”。

可以把执行过程类比为工厂：

CUDA Core / ALU：工人
寄存器：工人手边的工作台
Shared Memory：车间内部仓库
L2 Cache：整个工厂的公共仓库
Global Memory：远处的大型总仓库

如果工人每做一次加法，都要跑到远处仓库取两个数据，计算单元再多也可能一直在等待。

---

二、GPU 中的主要存储结构

先看总体关系：

图示：二、GPU 中的主要存储结构

其中：

• 寄存器、Shared Memory 和 L1 通常靠近或位于 SM 内；
• L2 Cache 是所有 SM 共享的；
• Global Memory 是 GPU 芯片外部的显存，例如 HBM 或 GDDR。

---

三、寄存器 Register

寄存器是 GPU 中最快的存储资源之一。

每个线程都有自己的寄存器。

例如：

float x = A[i];
float y = B[i];
float z = x + y;

通常：

• x 放在当前线程的寄存器中；
• y 放在当前线程的寄存器中；
• z 也放在当前线程的寄存器中。

线程之间不能直接读取对方的寄存器。

可以理解为：

图示：三、寄存器 Register

虽然每个线程都有独立寄存器，但物理上它们存放在 SM 内部的大型 Register File 中。

---

1. 寄存器为什么快

因为寄存器非常靠近执行单元。

执行：

z = x + y;

时，数据可以直接从寄存器送到 ALU，不需要访问片外显存。

所以通常希望：

需要反复使用的数据尽可能保存在寄存器中。

---

2. 寄存器也不是越多越好

每个 SM 的寄存器总容量是有限的。

假设一个 SM 有固定数量的寄存器。

如果每个线程只使用少量寄存器：

每个线程使用 32 个寄存器
→ SM 可以同时容纳较多线程

如果每个线程使用很多寄存器：

每个线程使用 128 个寄存器
→ SM 同时能容纳的线程数量下降

例如，假设 SM 一共有 65536 个寄存器。

每线程使用 32 个寄存器时，理论上可供：

$$65536 \div 32 = 2048$$

个线程使用。

如果每线程使用 128 个寄存器，则只能供：

$$65536 \div 128 = 512$$

个线程使用。

因此：

图示：2. 寄存器也不是越多越好

这就是为什么寄存器使用量会影响 Occupancy。

---

3. 寄存器溢出 Register Spilling

如果一个线程需要的寄存器太多，硬件无法全部提供，编译器可能会把部分变量放到 Local Memory。

这叫：

Register Spilling，寄存器溢出。

虽然名字叫 Local Memory，但它通常并不是真正位于线程旁边的高速存储，而是位于显存地址空间中。

因此可能出现：

原本：寄存器访问
后来：显存访问

性能会明显下降。

所以在 CUDA 性能分析中，常常需要检查：

• 每个线程用了多少寄存器；
• 是否发生寄存器溢出；
• 是否出现 Local Memory 访问。

---

四、Local Memory 是什么

Local Memory 很容易产生误解。

“Local”并不代表它一定很快，也不代表它就在 SM 内部。

它的“Local”主要是指：

每个线程私有，其他线程不能直接访问。

下面这些情况可能使用 Local Memory：

• 大型线程私有数组；
• 寄存器不够导致溢出；
• 编译器无法放入寄存器的局部变量；
• 动态索引访问的线程私有数据。

例如：

__global__ void kernel()
{
    float array[100];
}

如果编译器无法将这个数组完全放在寄存器中，它可能被放入 Local Memory。

逻辑上：

Thread 0 的 Local Memory
Thread 1 的 Local Memory
Thread 2 的 Local Memory

但是物理上，这些数据可能位于 Global Memory 中，并通过缓存访问。

因此需要记住：

Local Memory 是线程私有的，但不一定是低延迟的。

---

五、Shared Memory 共享内存

Shared Memory 是 GPU 中非常重要的片上存储资源。

它的特点是：

• 位于 SM 附近或 SM 内部；
• 延迟比 Global Memory 低；
• 同一个 Block 内的线程可以共享；
• 不同 Block 之间不能直接共享；
• 容量有限；
• 由程序员显式管理。

例如：

__shared__ float buffer[256];

这个数组属于当前 Block。

同一个 Block 中的线程都可以访问：

buffer[threadIdx.x]

可以理解为：

图示：五、Shared Memory 共享内存

每个 Block 都有自己独立的 Shared Memory 空间。

---

1. Shared Memory 为什么重要

假设同一个数据需要被一个 Block 中的多个线程重复使用。

如果每个线程都从 Global Memory 读取：

图示：1. Shared Memory 为什么重要

会产生大量片外访存。

更好的方法是：

1. 少量线程把数据从 Global Memory 搬到 Shared Memory
2. Block 内线程同步
3. 所有线程反复访问 Shared Memory

示意如下：

图示：1. Shared Memory 为什么重要

这叫做：

数据复用。

Shared Memory 最重要的作用，就是减少对 Global Memory 的重复访问。

---

2. Shared Memory 与 Block 的关系

一个 Block 被分配到一个 SM 后，它所需要的 Shared Memory 会从该 SM 的 Shared Memory 中分配。

假设一个 SM 有 64 KB Shared Memory。

如果每个 Block 使用 16 KB：

$$64\text{ KB}\div16\text{ KB}=4$$

理论上最多可以同时容纳 4 个这样的 Block。

如果每个 Block 使用 32 KB：

$$64\text{ KB}\div32\text{ KB}=2$$

只能容纳 2 个 Block。

所以：

每个 Block 使用的 Shared Memory 越多
                ↓
一个 SM 能同时驻留的 Block 越少
                ↓
同时驻留的 Warp 可能越少

Shared Memory 和寄存器一样，都会限制 Occupancy。

---

六、为什么 Shared Memory 需要同步

考虑下面的过程：

Thread 0 把数据写入 Shared Memory
Thread 1 读取这个数据

如果 Thread 1 读取时，Thread 0 还没有写完，就会读到错误数据。

所以通常需要使用同步指令：

__syncthreads();

例如：

shared_data[threadIdx.x] = global_data[i];

__syncthreads();

float x = shared_data[other_index];

它的含义可以简单理解为：

当前 Block 中的线程都执行到这里以后，才允许继续向后执行。

执行过程：

图示：六、为什么 Shared Memory 需要同步

注意，__syncthreads() 只能同步同一个 Block 内的线程。

它不能直接同步不同 Block。

---

七、Shared Memory Bank

Shared Memory 并不是单独的一整块存储器，而是被划分为多个 Bank。

可以简单理解为：

图示：七、Shared Memory Bank

多个 Bank 可以并行访问。

假设一个 Warp 中的 32 个线程分别访问不同 Bank：

图示：七、Shared Memory Bank

这些访问可以高效并行完成。

---

1. 什么是 Bank Conflict

如果多个线程同时访问同一个 Bank 中的不同地址，就可能产生 Bank Conflict。

例如：

Thread 0 → Bank 0，地址 A
Thread 1 → Bank 0，地址 B
Thread 2 → Bank 0，地址 C

由于同一个 Bank 不能同时完成这些不同地址的访问，硬件需要分多次处理。

这称为：

Shared Memory Bank Conflict，共享内存 Bank 冲突。

结果是：

原本希望一次完成
实际拆成多次完成

性能下降。

---

2. 同地址访问不一定冲突

如果多个线程同时读取同一个 Shared Memory 地址，硬件通常可以进行广播。

例如：

Thread 0 → 地址 X
Thread 1 → 地址 X
Thread 2 → 地址 X

这种情况可能由一次读取广播给多个线程，并不一定产生严重冲突。

Bank Conflict 主要发生在：

同一个 Bank 中的多个不同地址被同时访问。

---

八、L1 Cache

L1 Cache 是靠近 SM 的一级缓存。

与 Shared Memory 不同：

• Shared Memory 通常由程序员主动管理；
• L1 Cache 通常由硬件自动管理。

例如线程执行：

x = A[i];

它逻辑上访问 Global Memory，但数据可能已经存在 L1 Cache 中。

如果命中 L1：

线程请求数据
    ↓
L1 Cache 找到
    ↓
直接返回

如果没有命中：

线程请求数据
    ↓
L1 Cache 未找到
    ↓
继续访问 L2 Cache

缓存的基本作用是保存最近使用的数据，从而降低重复访问显存的开销。

---

Shared Memory 和 L1 Cache 的区别

二者都接近 SM，但管理方式不同。

对比	Shared Memory	L1 Cache
管理者	程序员	硬件
是否显式读写	是	否
线程范围	同一 Block 共享	SM 上的访问可受益
数据是否确定存在	程序员控制	不保证命中
主要用途	显式数据复用、线程通信	自动缓存访存数据

可以理解为：

Shared Memory：
程序员指定“把这个数据放在这里”

L1 Cache：
硬件自己判断“这个数据可能还会用到”

---

九、L2 Cache

L2 Cache 是 GPU 中所有 SM 共享的缓存。

结构大致如下：

SM 0 ─┐
SM 1 ─┤
SM 2 ─┼── L2 Cache ── Global Memory
SM 3 ─┘

它的作用包括：

• 缓存 Global Memory 数据；
• 减少显存访问；
• 在不同 SM 访问相同数据时提供一定的数据复用；
• 连接片上计算结构和显存控制器；
• 支持一些原子操作和一致性机制。

L2 比 L1 更大，但一般也更慢。

---

十、Global Memory

Global Memory 通常就是 GPU 的片外显存，例如：

• GDDR6；
• GDDR7；
• HBM2e；
• HBM3；
• HBM3e。

它的特点是：

• 容量大；
• 所有 SM 都可以访问；
• CPU 可以通过数据传输接口向其中复制数据；
• 带宽很高；
• 访问延迟也很高。

需要区分两个概念：

带宽

带宽表示单位时间内能传输多少数据。

例如：

$$1\text{ TB/s}$$

表示每秒理论上可以传输约 1 TB 数据。

延迟

延迟表示发出一次访问请求后，需要等待多久才能得到结果。

GPU 显存可以同时具有：

高带宽
+
高延迟

这并不矛盾。

可以类比为高速货运列车：

• 一次可以运输很多货物，所以带宽高；
• 货物从远方仓库送来仍然需要时间，所以延迟高。

---

十一、GPU 为什么能够容忍高显存延迟

GPU 的办法不是让每次访存都立刻返回，而是同时运行大量 Warp。

例如：

Warp 0：等待 Global Memory
Warp 1：执行计算
Warp 2：访问 Shared Memory
Warp 3：执行整数运算

调度器暂时不执行 Warp 0，而去执行其他已经准备好的 Warp。

所以：

图示：十一、GPU 为什么能够容忍高显存延迟

这也是 GPU 需要大量线程的重要原因。

---

十二、显存访问合并 Coalescing

这是 GPU 访存中非常重要的概念。

假设一个 Warp 有 32 个线程，每个线程读取一个 float。

一个 float 占 4 字节。

如果线程连续访问：

Thread 0 读取 A[0]
Thread 1 读取 A[1]
Thread 2 读取 A[2]
...
Thread 31 读取 A[31]

总数据量为：

$$32\times4=128\,\mathrm{B}$$

这些地址连续，硬件可以把它们组合成少量内存事务。

这叫：

Coalesced Memory Access，合并访存。

对应代码：

float x = A[global_thread_id];

通常是比较理想的访问方式。

---

1. 不连续访问

假设线程这样访问：

Thread 0 读取 A[0]
Thread 1 读取 A[100]
Thread 2 读取 A[200]
...
Thread 31 读取 A[3100]

每个线程访问的位置相距很远。

硬件可能需要发起大量独立内存事务。

一次 Warp 访存
        ↓
被拆成许多显存请求
        ↓
实际传输了大量无用数据
        ↓
有效带宽下降

即使显存理论带宽很高，程序也无法充分使用。

---

2. 合并访存的本质

合并访存并不是要求所有线程读取同一个地址，而是希望：

同一 Warp 中的线程访问相邻、对齐的数据地址。

理想情况：

Thread 0  → 地址 0
Thread 1  → 地址 4
Thread 2  → 地址 8
Thread 3  → 地址 12

较差情况：

图示：2. 合并访存的本质

因此，GPU 编程不仅要考虑算什么，还要考虑线程如何排列数据地址。

---

十三、二维图像访问示例

假设有一张二维图像：

image[row][column]

在内存中通常按行连续存放：

第 0 行：像素 0、1、2、3、4……
第 1 行：像素 0、1、2、3、4……

如果一个 Warp 中的线程沿水平方向访问：

图示：十三、二维图像访问示例

地址连续，容易合并访存。

如果沿竖直方向访问：

图示：十三、二维图像访问示例

相邻线程访问地址相距一整行，访存可能不连续。

所以矩阵转置、卷积等程序经常使用 Shared Memory 对数据重新排列。

---

十四、Shared Memory Tiling

Tiling 可以翻译为分块或分片。

这是 GPU 优化矩阵计算时最常见的方法之一。

考虑矩阵乘法：

$$C=A\times B$$

其中：

$$C_{ij}=\sum_k A_{ik}B_{kj}$$

计算不同的 $C_{ij}$ 时，会反复使用 A 和 B 中的数据。

如果每次都从 Global Memory 读取：

每做一次乘法
→ 读取一个 A
→ 读取一个 B

会产生大量重复访问。

更好的办法是把矩阵切成小块：

图示：十四、Shared Memory Tiling

示意：

图示：十四、Shared Memory Tiling

Tiling 的本质是：

用一次 Global Memory 加载，换取多次片上数据使用。

这是提高计算强度的重要方式。

---

十五、计算强度 Arithmetic Intensity

计算强度可以简单理解为：

$$\text{arithmetic intensity} = \frac{\text{operation count}} {\text{bytes accessed}}$$

例如数组加法：

C[i] = A[i] + B[i];

每个元素：

• 读取 A：4 字节；
• 读取 B：4 字节；
• 写入 C：4 字节；
• 进行 1 次加法。

大致为：

$$\frac{1\,\mathrm{op}}{12\,\mathrm{B}}$$

计算强度很低，通常受显存带宽限制。

矩阵乘法通过数据复用，可以让加载进来的数据参与大量乘加运算，计算强度较高，通常更容易充分使用 Tensor Core 或 CUDA Core。

---

十六、带宽受限和计算受限

GPU 程序的瓶颈大致可以分成两类。

1. Memory Bound：访存受限

特点：

• 计算量不大；
• 需要搬运大量数据；
• 计算单元经常等待数据；
• 提高 CUDA Core 数量帮助不大。

例如：

• 数组拷贝；
• 数组加法；
• 简单逐像素操作；
• 部分稀疏计算。

优化方向：

减少 Global Memory 访问
提高合并访存
增加数据复用
使用 Shared Memory
改善缓存命中率

2. Compute Bound：计算受限

特点：

• 每加载一份数据，会进行大量计算；
• 执行单元长期忙碌；
• 算力是主要瓶颈。

例如：

• 大规模矩阵乘法；
• 高计算量卷积；
• 部分科学计算。

优化方向：

提高执行单元利用率
使用 Tensor Core
减少指令依赖
使用合适的数据精度
增加并行度

---

十七、Constant Memory

Constant Memory 用于存放只读数据。

特点是：

• 数据通常由 CPU 写入；
• GPU 线程读取；
• Kernel 执行期间不修改；
• 有专门缓存；
• 当一个 Warp 中多个线程读取同一个地址时效率较高。

例如所有线程都使用同一个参数：

float scale = constant_data[0];

Warp 中：

Thread 0 读取地址 X
Thread 1 读取地址 X
...
Thread 31 读取地址 X

硬件可以把同一个值广播给所有线程。

适合存放：

• 固定参数；
• 小型查找表；
• 卷积常量；
• 所有线程共享的只读配置。

如果 Warp 中每个线程读取完全不同的 Constant Memory 地址，性能优势可能下降。

---

十八、Texture Memory

Texture Memory 最初主要服务于图形纹理访问。

它适合：

• 二维或三维空间局部访问；
• 只读数据；
• 图像数据；
• 坐标访问；
• 某些具有空间局部性的非规则访存。

现代 GPU 中，纹理缓存和普通只读缓存的具体实现会因架构而变化。

从入门角度，只需要理解：

Texture Memory 是针对图像和空间局部性访问优化的只读存储路径。

它不是额外的一块巨大物理内存，数据通常仍然来自 Global Memory，只是使用了特殊的访问和缓存机制。

---

十九、不同存储空间的作用范围

可以用下面的表格概括：

存储类型	谁可以访问	典型作用范围
Register	单个 Thread	线程私有
Local Memory	单个 Thread	线程私有
Shared Memory	同一 Block	Block 内共享
L1 Cache	靠近单个 SM	SM 附近
L2 Cache	全部 SM	整个 GPU
Global Memory	全部线程	整个 GPU
Constant Memory	全部线程只读	整个 GPU
Texture Memory	全部线程只读	整个 GPU

---

二十、不同存储空间的生命周期

Register

通常随线程创建和结束。

Thread 创建 → 寄存器有效
Thread 结束 → 寄存器释放

Shared Memory

通常随 Block 创建和结束。

Block 分配到 SM → Shared Memory 分配
Block 执行结束 → Shared Memory 释放

Global Memory

由程序显式分配和释放。

cudaMalloc(...);
cudaFree(...);

其生命周期通常长于单个 Kernel。

---

二十一、完整的数据访问过程

假设线程执行：

float x = A[i];

可以粗略理解为：

图示：二十一、完整的数据访问过程

然后执行：

float y = x + 1.0f;

此时 x 已经在寄存器中，加法可以直接在执行单元完成。

---

二十二、CPU Cache 和 GPU Cache 的区别

CPU 和 GPU 都有 Cache，但使用目标不同。

CPU 缓存主要服务于：

• 少量线程；
• 低延迟；
• 复杂控制流；
• 单线程快速响应。

GPU 缓存主要服务于：

• 大量线程；
• 高吞吐量；
• 大规模数据流；
• 配合线程并行隐藏延迟。

GPU 通常不能只依赖“大缓存”解决问题，因为：

• 并发线程太多；
• 工作集很大；
• 数据吞吐需求很高；
• 芯片面积需要优先留给计算和寄存器资源。

所以 GPU 更强调：

缓存
+
Shared Memory
+
合并访存
+
大量 Warp
+
高显存带宽

共同解决数据供应问题。

---

二十三、一个卷积计算中的数据流

以神经网络卷积为例：

输入特征图
权重
偏置

最简单的实现可能不断从 Global Memory 读取输入和权重。

优化后通常是：

图示：二十三、一个卷积计算中的数据流

理想的数据流是：

1. 从 Global Memory 成块读取；
2. 在 Shared Memory 中复用；
3. 把当前计算数据送入寄存器；
4. 在寄存器中完成多次累加；
5. 最后只写回一次结果。

---

二十四、需要避免的常见问题

问题一：每次都访问 Global Memory

sum += global_data[index];

如果同一数据反复使用，应该考虑 Cache、Shared Memory 或寄存器复用。

问题二：线程地址不连续

相邻线程访问相隔很远的地址

会降低访存合并效率。

问题三：Shared Memory Bank Conflict

多个线程同时访问同一 Bank 中的不同地址

可能导致访问串行化。

问题四：每线程寄存器过多

会降低 SM 可同时驻留的 Warp 数量。

问题五：Shared Memory 使用过多

会降低每个 SM 能同时驻留的 Block 数量。

问题六：频繁同步

__syncthreads() 可以保证正确性，但同步过多也会造成等待。

---

二十五、本节需要掌握的结构图

Thread 私有：
    Register
    Local Memory

Block 共享：
    Shared Memory

SM 附近：
    L1 Cache

全 GPU 共享：
    L2 Cache
    Global Memory

访问速度大致可理解为：

Register
   ↓
Shared Memory / L1
   ↓
L2
   ↓
Global Memory

数据容量大致相反：

Global Memory
   ↓
L2
   ↓
Shared Memory / L1
   ↓
Register

---

本节核心结论

第一，GPU 的算力只有在数据能够及时送到执行单元时才有意义。

第二，寄存器速度快，但每个 SM 的寄存器总量有限。

第三，Shared Memory 是 Block 内线程共享的片上存储器，主要用于线程通信和数据复用。

第四，L1 和 L2 Cache 由硬件自动管理，而 Shared Memory 通常由程序员显式管理。

第五，Global Memory 容量大、带宽高，但延迟也高。

第六，同一 Warp 的线程连续访问内存，可以形成合并访存。

第七，Shared Memory 的 Bank Conflict 会使本来并行的访问变成多次访问。

第八，使用 Tiling 可以把数据从 Global Memory 搬到 Shared Memory，进行多次复用。

第九，寄存器和 Shared Memory 使用过多都会降低 SM 上的并发线程数量。

第十，GPU 优化的核心通常不是单纯减少一条指令，而是减少昂贵的数据搬运。

下一部分应学习 SM 内部结构和指令流水线，包括：

Warp Scheduler
Instruction Dispatch
CUDA Core
INT Unit
Load/Store Unit
SFU
Tensor Core
Scoreboard
指令依赖
流水线

理解这些结构后，就能进一步解释：为什么一个 Warp 并不是一拍执行完32个线程，以及多条指令在 SM 中究竟怎样流动。