CUDA内存模型全景：从全局内存到寄存器，一张图看懂GPU存储层次
做GPU优化的人，心里都有一张内存层次图。我刚开始接触CUDA时，最头疼的就是搞不清各种内存的区别。全局内存、共享内存、寄存器、常量内存、纹理内存……名字一大堆，用起来更是容易踩坑。
说白了，GPU的内存模型就像一座金字塔。越往上，速度越快，容量越小。越往下，速度越慢，容量越大。今天我就带你把这个金字塔彻底搞明白。
核心观点：CUDA优化的本质，就是让数据尽可能待在金字塔的上层。数据离计算单元越近，程序跑得越快。
一、GPU存储层次全景图
先上一张我自己画的图。这张图我用了好多年，每次给团队新人培训都会拿出来讲。你把它存下来，以后写代码时对照着看，思路会清晰很多。
这张图里，我特意把寄存器放在最顶端。为什么？因为寄存器是GPU里最快的存储单元，访问延迟只有1个时钟周期。而最底层的全局内存，延迟高达400-800个周期。你想想看，这中间差了将近三个数量级。
二、逐层拆解：每层内存的脾气秉性
1. 寄存器（Register）
寄存器是每个线程私有的。说白了，你在kernel里定义的局部变量，只要编译器没优化掉，基本都放在寄存器里。
我的经验：我曾经在一个矩阵乘法的项目里，因为每个线程用了太多寄存器，导致occupancy从100%掉到了33%。性能直接腰斩。后来我手动拆分了变量，把一些中间结果放到共享内存里，occupancy才拉回来。
寄存器数量是有限的。每个SM有65536个寄存器（以A100为例），分配到每个线程最多255个。如果你每个线程用了超过这个数，寄存器就会溢出到局部内存——那玩意儿其实在全局内存里，慢得要命。
避坑指南：我曾经见过一个新手，在kernel里定义了一个巨大的结构体数组，结果编译器默默地把所有变量都 spill 到了局部内存。性能惨不忍睹。记得用 --ptxas-options=-v 编译选项查看寄存器使用情况。
2. L1缓存 / 共享内存
这两个东西在硬件上共用同一块物理存储。你可以通过API配置它们的比例。比如在Turing架构上，可以配成48KB共享内存+16KB L1，或者反过来。
共享内存是块上所有线程共享的。它的速度跟L1差不多，但它是可编程的。你可以手动控制数据什么时候加载进来，什么时候写回去。
关键区别：L1是硬件自动管理的，共享内存是程序员手动管理的。手动管理虽然麻烦，但你能精确控制数据复用，这是优化的核心。
3. L2缓存
L2缓存是所有SM共享的。它缓存全局内存和局部内存的访问。在A100上，L2缓存有40MB，已经相当大了。
但要注意，L2缓存是透明的。你没法直接控制它。不过你可以通过调整数据访问模式来利用它。比如，让同一个warp里的线程访问连续的地址，这样L2缓存的命中率会高很多。
4. 常量内存 / 纹理内存
常量内存只有64KB，但它的读取有专门的缓存和广播机制。如果同一个warp里的所有线程读取同一个常量地址，那只需要一次内存事务，速度极快。
纹理内存呢，它专门为空间局部性优化过。如果你在做图像处理或者需要插值计算，纹理内存比全局内存快不少。
我个人的习惯：凡是只读且所有线程访问相同数据的参数，我都用常量内存。比如卷积核的权重、矩阵乘法的系数。64KB虽然小，但放这些参数绰绰有余。
5. 全局内存
全局内存是GPU里最大的存储空间，也是速度最慢的。所有线程都能访问它，但每次访问都要付出几百个周期的代价。
优化全局内存访问的核心原则就一条：合并访问。也就是让相邻的线程访问相邻的地址。这样GPU可以把多次内存请求合并成一次大的事务，效率高得多。
// 好的访问模式：合并访问
__global__ void good_access(float* data) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
float val = data[idx]; // 相邻线程访问相邻地址
}
// 坏的访问模式：非合并访问
__global__ void bad_access(float* data) {
int idx = threadIdx.y * gridDim.x * blockDim.x + threadIdx.x;
float val = data[idx * 32]; // 相邻线程访问间隔32的地址
}三、一张表总结所有内存特性
| 内存类型 | 作用域 | 生命周期 | 延迟 | 容量 | 缓存 |
|---|---|---|---|---|---|
| 寄存器 | 单个线程 | kernel执行期间 | 1 cycle | 256KB/SM | 无 |
| L1缓存 | 单个SM | 自动管理 | ~30 cycles | 24-128KB/SM | 硬件控制 |
| 共享内存 | 单个Block | kernel执行期间 | ~30 cycles | 48-164KB/Block | 可编程 |
| L2缓存 | 所有SM | 自动管理 | 100-200 cycles | 40KB-40MB | 硬件控制 |
| 常量内存 | 所有线程 | 程序运行期间 | ~200 cycles | 64KB | 专用缓存 |
| 纹理内存 | 所有线程 | 程序运行期间 | ~200 cycles | 依赖设备 | 专用缓存 |
| 全局内存 | 所有线程 | 程序运行期间 | 400-800 cycles | GB级 | L1+L2 |
四、实际项目中的内存选择策略
好了，理论讲完了。咱们来点实际的。你在写一个kernel时，怎么决定数据放哪？
我一般按这个顺序思考：
- 先看数据是不是只读且所有线程共享 → 放常量内存（64KB以内）或纹理内存
- 再看数据是不是被同一个block多次复用 → 放共享内存，手动管理
- 如果数据只被单个线程使用 → 尽量用寄存器，别用局部内存
- 剩下的，放全局内存 → 但一定要保证合并访问
一个真实的案例：我之前优化过一个流体力学模拟程序。原始版本把所有数据都放在全局内存里，每个线程反复读取同一块数据。我把它改成：把每个block需要的子区域加载到共享内存，计算完再写回全局内存。结果性能提升了4倍。说白了，就是让数据在金字塔上层多待一会儿。
五、常见误区与避坑
- 误区一：寄存器越多越好 —— 错！寄存器太多会降低occupancy，反而更慢。我建议每个线程控制在32个寄存器以内。
- 误区二：共享内存越大越好 —— 错！共享内存用太多，每个SM能同时运行的block数就少了。要平衡。
- 误区三：全局内存随便访问 —— 错！不合并访问的全局内存，性能可能差10倍以上。
- 误区四：L2缓存不用管 —— 错！虽然你不能直接控制L2，但数据访问模式会影响L2命中率。尽量让数据访问有规律。
我曾经踩过的坑：有一次我为了省事，把所有数据都声明为 __device__ 全局变量。结果每个线程都去读写同一个地址，导致严重的bank conflict和缓存抖动。后来改成每个线程处理自己的数据分区，问题才解决。
嗯，到这里，CUDA内存模型的全景图你应该心里有数了。记住那张金字塔图，写代码时多想想：我的数据现在在金字塔的哪一层？能不能往上挪一挪？
想通了这个问题，你的CUDA优化之路就成功了一半。