显存体系详解：全局显存、共享显存、L1/L2缓存、寄存器文件的带宽与延迟特性
做GPU编程这些年，我最大的感触就是——搞懂显存体系，你就掌握了性能优化的半壁江山。很多人一上来就写核函数，结果跑起来发现带宽利用率不到20%，还一脸懵。说白了，你连数据在哪儿、怎么搬的都没搞清楚，怎么可能写出高性能代码？
今天咱们就把GPU的存储体系掰开揉碎了讲。我会结合自己踩过的坑，带你彻底搞懂每一级存储的特性。
一、GPU存储层次全景图
先看一张我手绘的存储层次图。这张图我每次培训都会拿出来讲，因为它太重要了。
这张图你看懂了吗？从上到下，速度越来越慢，容量越来越大。我当年刚接触CUDA时，总觉得全局显存够快了啊，结果一测延迟，好家伙，800个时钟周期！从那以后，我再也不敢随便乱访存了。
二、寄存器文件：最快的存储，但也是最稀缺的
寄存器文件是离计算单元最近的存储。它的速度有多快？一个时钟周期就能完成读写。说白了，这就是GPU的"贴身侍卫"。
关键数据：
- 带宽：约20 TB/s（每个SM内部）
- 延迟：1个时钟周期
- 容量：每个SM约256KB（以A100为例）
- 分配方式：每个线程私有，编译器自动分配
我在项目中遇到过一个问题：一个简单的向量加法，性能死活上不去。后来一查，每个线程用了32个寄存器，SM的occupancy被压到了25%。你想想看，本来能跑2048个线程的SM，现在只能跑512个，性能能好吗？
我的经验：用 --ptxas-options=-v 编译选项查看寄存器使用量。如果每个线程超过32个寄存器，就要考虑优化了。我曾经把一个kernel的寄存器从48压到24，性能直接翻倍。
三、L1缓存与共享显存：可编程的"快车道"
L1缓存和共享显存在物理上是同一块硬件，但用法完全不同。L1是自动的，你只管读数据，它自己缓存。共享显存则需要你手动管理——显式地把数据从全局显存搬进来。
我个人习惯把共享显存当作"程序员可控的缓存"。为什么这么说？因为它的延迟只有20-30个周期，比全局显存快了一个数量级。
| 特性 | L1缓存 | 共享显存 |
|---|---|---|
| 控制方式 | 硬件自动 | 程序员手动 |
| 延迟 | ~30 cycles | ~20 cycles |
| 带宽 | ~10 TB/s | ~10 TB/s |
| 容量配置 | 可配（如48KB L1 + 16KB SM） | 可配（如16KB L1 + 48KB SM） |
| 典型用途 | 全局显存访问的自动缓存 | 线程块内数据共享、归约 |
嗯，这里要注意：共享显存虽然快，但用不好就是灾难。我曾经见过一个同事，把整个矩阵都塞进共享显存，结果每个SM只能跑一个线程块，occupancy直接归零。
避坑指南：共享显存不是越大越好。我曾经为了省事，把共享显存配到最大（48KB），结果L1只剩16KB，全局显存的缓存命中率暴跌。后来我学乖了，根据访存模式动态调整：
- 如果数据重用率高 → 多分给共享显存
- 如果随机访问多 → 多分给L1缓存
四、L2缓存：芯片的"中转站"
L2缓存是所有SM共享的。它的带宽大概在2-4 TB/s，延迟200-400个周期。虽然比不上L1，但比全局显存还是快不少。
我记得有一次做矩阵乘法优化，发现性能瓶颈不在计算，而在L2缓存的bank冲突。每个SM都在抢L2的同一块数据，结果带宽利用率只有30%。后来我调整了数据布局，让每个SM访问不同的L2 slice，性能提升了2倍。
L2缓存的关键特性：
- 芯片全局共享，所有SM都能访问
- 容量：4MB-40MB（取决于GPU型号）
- 缓存行大小：128字节（和全局显存一致）
- 写策略：写回（write-back）
五、全局显存：容量最大，但也是最慢的
全局显存就是GPU的"主存"。HBM2e的带宽能做到2 TB/s，GDDR6X也能到1 TB/s。但延迟嘛...400-800个周期。你想想看，CPU的L1缓存延迟才几个周期，这差距有多大。
但是！全局显存有一个隐藏特性：合并访问。如果相邻线程访问相邻地址，硬件会把多次访问合并成一次。我见过太多人没利用好这个特性，导致带宽利用率不到10%。
// ❌ 错误示范：非合并访问
__global__ void bad_access(float *data) {
int tid = threadIdx.x;
float val = data[tid * 1024]; // 每个线程间隔1024个元素
}
// ✅ 正确示范：合并访问
__global__ void good_access(float *data) {
int tid = threadIdx.x;
float val = data[tid]; // 相邻线程访问相邻地址
}我的建议：写全局显存访问代码时，脑子里要有一张"线程-地址"映射图。保证同一个warp内的32个线程访问的地址是连续的、对齐的。我曾经用这个技巧，把带宽利用率从15%拉到了85%。
六、各存储层次的实战对比
说了这么多，咱们来个实战对比。假设你要处理一个1GB的数据集：
| 存储层次 | 理论带宽 | 实际可用带宽 | 处理1GB耗时 | 适用场景 |
|---|---|---|---|---|
| 寄存器 | 20 TB/s | ~15 TB/s | ~0.07ms | 局部变量、中间结果 |
| 共享显存 | 10 TB/s | ~8 TB/s | ~0.13ms | 线程块内数据共享 |
| L2缓存 | 3 TB/s | ~2 TB/s | ~0.5ms | 数据重用、缓存友好 |
| 全局显存 | 1.5 TB/s | ~1 TB/s | ~1ms | 大容量数据存储 |
看到差距了吧？寄存器比全局显存快了200多倍。所以优化的核心思路就是：让数据尽量待在上层存储中。
七、总结：我的"三板斧"
做GPU性能优化这么多年，我总结了三句话：
- 能放寄存器的，绝不放共享显存——编译器会自动优化，但你要注意寄存器压力
- 能放共享显存的，绝不放全局显存——手动管理虽然麻烦，但收益巨大
- 访问全局显存时，一定要合并——这是最容易被忽视的性能杀手
说白了，GPU的存储体系就像一座金字塔。你越靠近塔尖，速度越快，但容量越小。作为工程师，你的任务就是在容量和速度之间找到最佳平衡点。
好了，这一章的内容就到这里。记住我今天讲的这些数字和特性，下次写kernel的时候，多想想你的数据到底在哪儿、怎么搬的。相信我，性能提升会让你惊喜的。