第三章 存储层次结构：寄存器文件、共享内存、L1/L2缓存、全局内存与HBM
各位同学，今天我们来聊聊GPU里最绕不开的话题——存储层次结构。
说实话，我见过太多做AI加速的工程师，算法调得飞起，结果一上板卡，性能直接腰斩。为什么？八成是内存访问模式出了问题。你想想看，计算单元再快，数据喂不进去，那就是白搭。
3.1 寄存器文件：最快但最稀缺
寄存器文件，说白了就是离计算单元最近的家当。在GPU里，每个线程都有自己的私有寄存器。我习惯把它比作「桌面上的工具」——伸手就能拿到，但桌面就这么大。
关键参数：
- 访问延迟：1个时钟周期
- 容量：每个线程最多255个32位寄存器（NVIDIA架构）
- 作用域：线程私有
实战经验：我在做某个卷积算子优化时，发现寄存器溢出（register spilling）导致性能掉了40%。当时我死活想不通，明明算法没问题啊。后来一查，是编译器把局部变量挤到全局内存里去了。嗯，这里要注意——寄存器不够用的时候，编译器会悄悄把数据挪到L1缓存甚至全局内存，你根本不知道。
避坑指南：我曾经在写矩阵乘法时，为了省寄存器，把中间结果反复从全局内存读取。结果呢？比直接用寄存器慢了8倍。所以我的建议是：宁可少开几个线程，也要保证每个线程有足够的寄存器。
3.2 共享内存：可编程的L1缓存
共享内存是个好东西。它位于芯片上，由程序员显式控制。说白了，就是让你手动管理数据复用。
为什么需要它？因为全局内存太慢了。你想想看，一个线程从全局内存读数据要几百个周期，而共享内存只要几十个周期。如果能让一个线程块里的所有线程共享同一份数据，那性能提升是立竿见影的。
典型用法：
// 声明共享内存
__shared__ float tile[BLOCK_SIZE][BLOCK_SIZE];
// 从全局内存加载到共享内存
tile[threadIdx.y][threadIdx.x] = global[base + threadIdx.y * WIDTH + threadIdx.x];
// 同步，确保所有线程都加载完毕
__syncthreads();
// 从共享内存读取，进行矩阵乘法计算
float sum = 0.0f;
for (int k = 0; k < BLOCK_SIZE; ++k) {
sum += tile[threadIdx.y][k] * tile[k][threadIdx.x];
}注意：共享内存有bank冲突。我在项目中遇到过，明明用了共享内存，性能却没提升多少。一分析，发现是bank冲突导致串行化访问。解决办法？要么padding，要么重新组织数据布局。
3.3 L1/L2缓存：硬件自动管理
L1和L2缓存是硬件自动管理的，程序员没法直接控制。但你可以通过数据访问模式来「暗示」缓存。
L1缓存：
- 每个SM独立拥有
- 容量：通常16KB-128KB
- 延迟：约30个时钟周期
- 特点：对局部性敏感
L2缓存：
- 所有SM共享
- 容量：通常几MB
- 延迟：约200个时钟周期
- 特点：缓解全局内存带宽压力
我个人习惯把L1缓存当作「意外之喜」。什么意思呢？就是不要依赖它，但如果数据访问模式好，它能帮你白捡性能。比如，连续访问内存地址，L1缓存命中率就高。反之，随机访问？那L1基本等于摆设。
3.4 全局内存：大容量但慢
全局内存是GPU里最大的存储空间，通常有几十GB。但它的延迟也是最高的——几百个时钟周期。
访问模式决定性能：
| 访问模式 | 带宽利用率 | 说明 |
|---|---|---|
| 合并访问（coalesced） | 接近100% | 相邻线程访问相邻地址 |
| 非合并访问 | 10%-50% | 线程访问分散地址 |
| 随机访问 | <10% | 基本等于没优化 |
实战案例：我记得有一次优化一个图像处理算子，原始代码是这么写的：
// 糟糕的访问模式
for (int i = 0; i < N; ++i) {
output[threadIdx.x] += input[threadIdx.x * N + i];
}改成合并访问后：
// 好的访问模式
for (int i = 0; i < N; ++i) {
output[threadIdx.x] += input[i * N + threadIdx.x];
}性能直接提升了5倍。你想想看，就改了一行代码，效果天差地别。
3.5 HBM：高带宽内存
HBM（High Bandwidth Memory）是近年来GPU的标配。它通过3D堆叠技术，把多个DRAM die堆在一起，中间用硅通孔（TSV）连接。
HBM的优势：
- 带宽极高：HBM2e可达1.6TB/s，HBM3超过3TB/s
- 功耗低：相比GDDR6，能效比提升约3倍
- 容量适中：单颗HBM2e通常8GB-16GB
说白了，HBM就是给GPU修了一条超宽的高速公路。但要注意，带宽再高，也架不住你乱访问。我见过有人用HBM跑随机访问，结果带宽利用率不到5%。
我的建议：用HBM时，一定要做带宽测试。我曾经在项目里用了一个第三方库，结果它内部做了大量小粒度随机访问，把HBM的带宽全浪费了。后来我们自己重写了数据加载部分，带宽利用率从15%提升到了85%。
3.6 存储层次总结
好了，我们来捋一捋整个存储层次：
| 存储层次 | 容量 | 延迟 | 管理方式 | 作用域 |
|---|---|---|---|---|
| 寄存器文件 | 几十KB/SM | 1周期 | 编译器 | 线程私有 |
| 共享内存 | 几十KB/SM | ~30周期 | 程序员 | 线程块内 |
| L1缓存 | 16-128KB/SM | ~30周期 | 硬件 | SM内 |
| L2缓存 | 几MB | ~200周期 | 硬件 | 全局 |
| 全局内存 | 几十GB | ~400周期 | 程序员 | 全局 |
| HBM | 8-80GB | ~400周期 | 硬件+程序员 | 全局 |
核心原则：数据越热，越要放在靠近计算单元的地方。说白了，就是「热数据上移，冷数据下移」。我见过太多人把热数据放在全局内存里反复读，结果性能惨不忍睹。
最后说一句，存储层次结构不是死的。不同的GPU架构，参数会有差异。但核心思想不变——理解数据流动，优化访问模式。你把这个搞明白了，AI加速就成功了一半。
下一章，我们来聊聊指令流水线和Warp调度。嗯，那又是另一个有意思的话题了。