3. 显存层次结构：GPU显存体系
聊GPU显存，我习惯把它想象成一个「金字塔」。
塔尖最快，但容量最小。塔基最慢，但容量最大。你写的每一行CUDA代码，本质上都在跟这个金字塔打交道。搞不懂它，吞吐优化就是空中楼阁。
3.1 金字塔的四层结构
先看这张图，我手绘的，帮你建立整体认知：
说白了，这就是一个「快小贵」vs「慢大便宜」的博弈。你想想看，如果所有数据都能塞进寄存器，那GPU早就飞起来了。但现实是——寄存器就那么一丁点。
3.2 各层级详解
寄存器（Register）
速度：~20个时钟周期。这是GPU里最快的地方，没有之一。
容量：每个线程最多255个32位寄存器（以NVIDIA为例）。注意，是「每个线程」。一个SM里有几千个线程在跑，寄存器总量是固定的，你用得越多，能同时跑的线程就越少。
我踩过的坑：有一次我写一个矩阵乘法的kernel，为了省事把所有中间变量都声明成float，结果寄存器溢出（spill）到了local memory。性能直接掉了40%。后来我手动复用寄存器，把一些中间结果算完就扔，才把性能拉回来。
L1缓存 / Shared Memory
速度：~30个时钟周期。比寄存器慢一点点，但比L2快一个数量级。
容量：每个SM独立拥有。以A100为例，每个SM有192KB的L1+Shared Memory总和，你可以通过配置来分配多少给L1、多少给Shared Memory。
我个人习惯：如果数据需要在线程块内共享，我会优先用Shared Memory。如果只是单纯的数据重用，交给L1缓存自动管理就好。
cudaFuncSetAttribute 来设置每个线程块能用的Shared Memory上限。别设太大，否则SM能同时跑的线程块数会减少。
L2缓存
速度：~200个时钟周期。比L1慢，但比HBM快得多。
容量：所有SM共享。A100有40MB，H100有50MB。嗯，这个容量其实挺大的，很多模型参数都能塞进去。
避坑指南：我曾经以为L2缓存是「透明的」，不用管它。后来发现，如果多个SM频繁访问同一块显存地址，L2的命中率会很高。但如果每个SM都访问不同的地址，L2就变成了一个昂贵的「搬运工」。所以，数据访问模式要尽量让L2命中。
HBM / 全局显存
速度：~400~800个时钟周期。这是最慢的一层，但容量最大。
容量：几十GB。A100有40GB/80GB，H100有80GB。大模型训练全靠它撑着。
说白了，全局显存就是你的「仓库」。所有数据最终都存在这里。但每次从仓库取货，都要花几百个周期。所以，尽量减少对全局显存的直接访问，这是吞吐优化的核心原则之一。
3.3 速度与容量的量化对比
我整理了一张表，方便你直观对比：
| 层级 | 典型延迟（周期） | 典型容量（A100） | 作用域 |
|---|---|---|---|
| 寄存器 | ~20 | 每个线程 255 个 | 线程私有 |
| L1 / Shared Memory | ~30 | 每个 SM 192KB | 线程块内共享 |
| L2 缓存 | ~200 | 40 MB | 全局共享 |
| HBM 全局显存 | ~400~800 | 40 GB / 80 GB | 全局共享 |
你看这个延迟差距：寄存器到HBM，差了40倍。这意味着，如果你能从HBM里少读一次数据，换成从寄存器里读，性能就能提升40倍。当然，这是理想情况，但方向是对的。
3.4 实战中的层次协同
光知道层级没用，你得会用。我分享一个实际案例：
之前优化一个Attention kernel。原始版本直接读HBM里的Q、K、V矩阵，每次计算都从全局显存取数。性能惨不忍睹。
我的优化思路：
- 把Q矩阵的一整块数据加载到Shared Memory（L1层），让线程块内的所有线程共享。
- 把K矩阵分块加载到寄存器，每个线程只负责自己需要的部分。
- 计算过程中，中间结果全部留在寄存器里，绝不写回HBM。
- 最终结果才写回HBM。
结果呢？吞吐提升了3倍。说白了，就是让数据尽量待在金字塔的上层，少去底层「旅游」。
3.5 总结一下
显存层次结构，说白了就是一套「缓存体系」。你写的代码，决定了数据在这个金字塔里怎么流动。
- 寄存器：最快，但极小。别让变量溢出。
- L1/Shared Memory：次快，线程块内共享。手动管理能带来巨大收益。
- L2：全局共享，自动管理。注意访问模式。
- HBM：最慢，但容量大。尽量减少直接访问。
嗯，理解了这个层次结构，你再看后面的「吞吐优化」和「显存复用」，就会觉得顺理成章了。因为所有优化技巧，本质上都是在跟这个金字塔做博弈。
想要更多资料，联系可可