CPU vs GPU内存层次：寄存器、L1/L2缓存、全局内存、共享内存的延迟对比
做算子优化，说白了就是在跟内存打交道。
我刚开始接触CUDA时，总觉得把数据搬到GPU上就完事了。结果呢？性能惨不忍睹。后来才明白——不理解内存层次，就别谈优化。今天咱们就掰开揉碎，看看CPU和GPU这两套内存体系到底差在哪。
一、CPU的内存层次：老大哥的套路
CPU的内存设计，核心思路是「又快又准」。它要处理各种 unpredictable 的任务，所以缓存策略非常激进。
| 层级 | 典型大小 | 延迟（周期） | 带宽 |
|---|---|---|---|
| 寄存器 | 几百字节 | 0-1 cycle | 极高 |
| L1 缓存 | 32KB-64KB | 3-5 cycles | ~1TB/s |
| L2 缓存 | 256KB-1MB | 10-20 cycles | ~500GB/s |
| L3 缓存 | 8MB-32MB | 30-50 cycles | ~200GB/s |
| 主存（DDR） | 16GB-512GB | 100-300 cycles | ~50GB/s |
关键点：CPU的缓存是透明的。你写代码时不用管数据在哪层缓存，硬件自动帮你搬。但代价是——你没法精确控制。
我记得有一次优化一个图像处理算子，CPU上跑得飞快，但一上GPU就卡住。后来发现是CPU的L2缓存把热点数据全兜住了，而GPU那边……嗯，后面再说。
二、GPU的内存层次：并行机器的玩法
GPU的设计哲学完全不同。它假设你有成千上万个线程，每个线程干的事很简单。所以它的内存层次更「显式」——你得自己管。
| 层级 | 典型大小 | 延迟（周期） | 带宽 |
|---|---|---|---|
| 寄存器 | 每个SM 64K-256K | 0-1 cycle | 极高 |
| L1/共享内存 | 每个SM 48KB-128KB | ~30 cycles | ~5TB/s |
| L2 缓存 | 4MB-8MB | ~200 cycles | ~2TB/s |
| 全局内存（HBM/GDDR） | 8GB-80GB | 400-800 cycles | ~1TB/s |
注意：GPU的L1和共享内存是共用同一块物理存储的。你可以通过配置决定分多少给L1、多少给共享内存。我曾经因为没配好这个，导致共享内存不够用，L1又浪费了一大片……血的教训。
三、延迟对比：数字会说话
咱们把两组数据放一起看看，差距一目了然：
| 操作 | CPU延迟 | GPU延迟 | 差距倍数 |
|---|---|---|---|
| 寄存器访问 | 0-1 cycle | 0-1 cycle | ~1x |
| L1 命中 | 3-5 cycles | ~30 cycles | ~6-10x |
| L2 命中 | 10-20 cycles | ~200 cycles | ~10-20x |
| 主存/全局内存 | 100-300 cycles | 400-800 cycles | ~2-4x |
你可能会问：为什么GPU的L1延迟比CPU高那么多？
原因很简单——GPU的L1要服务几百个线程。CPU的L1只服务1-2个核心，冲突少、命中率高。GPU的L1被大量线程争抢，延迟自然上去了。
我的建议：在GPU上，尽量把热点数据塞进寄存器。如果寄存器不够用，再考虑共享内存。全局内存是最后的选择——除非你不得不访问它。
四、一张图看懂内存层次
下面这张SVG图，把CPU和GPU的内存层次画在一起，方便对比：
五、实战中的选择策略
知道了延迟差距，怎么用？我分享几个经验：
- 能放寄存器就别放共享内存——寄存器延迟是0，共享内存要30个周期。我见过有人把循环变量放共享内存，白白浪费性能。
- 共享内存要手动管理——CPU的L1是自动的，但GPU的共享内存你得自己搬数据。用
__syncthreads()保证同步，否则数据乱掉。 - 全局内存要合并访问——连续线程访问连续地址，带宽才能跑满。我曾经因为没注意对齐，带宽利用率只有10%。
- L2缓存是共享的——所有SM共用L2，所以别指望它存太多数据。热点数据尽量往共享内存放。
核心原则：在GPU上，显式控制比自动缓存更可靠。CPU你可以偷懒，GPU不行。
六、避坑指南
我曾经犯过一个经典错误：在GPU kernel里频繁访问全局内存，以为L2能兜住。结果呢？L2被大量线程冲垮，延迟飙升到800 cycles。后来改成共享内存做缓存，性能提升了5倍。
所以我的建议是：别指望GPU的L2能救你。它只是辅助，主力还得是寄存器和共享内存。
另一个坑是寄存器溢出。GPU每个线程的寄存器数量有限（通常32-64个），如果你声明太多局部变量，编译器会把它们 spill 到全局内存——那延迟直接爆炸。我一般会检查编译报告，看看寄存器使用量，超过阈值就优化代码。
好了，关于内存层次的延迟对比就聊到这。记住一句话：数据在哪，性能就在哪。