CUDA内存模型全景：从全局到寄存器，每一级都有它的脾气
做GPU编程这些年，我见过太多人一上来就写cudaMalloc，然后抱怨程序跑得慢。其实啊，很多时候问题不在算法，而在内存访问上。今天我们就来聊聊CUDA的内存模型——说白了，就是数据在GPU里是怎么存放、怎么流动的。
我个人习惯把CUDA内存想象成一栋楼。全局内存是地下室，大但慢；共享内存是一楼大厅，大家都能快速访问；寄存器呢，就是每个人的私人办公室，最快但也最有限。你想想看，如果程序频繁往地下室跑，那效率能高吗？
一、全局内存：GPU的“大仓库”
全局内存是GPU里容量最大的存储空间，通常有几十GB。所有线程都能读写它，但代价是访问延迟很高——大概400-800个时钟周期。
关键点：全局内存的访问必须合并。什么意思呢？就是相邻的线程最好访问相邻的地址。我在项目中遇到过，有人写了个随机访问的kernel，结果带宽利用率不到10%。
举个例子：
// 好的访问模式：合并访问
float value = data[threadIdx.x]; // 相邻线程访问相邻地址
// 坏的访问模式：非合并访问
float value = data[threadIdx.x * 64]; // 跨度太大，带宽浪费二、共享内存：线程块的“公共客厅”
共享内存是每个线程块私有的，容量很小（通常48KB-164KB），但速度比全局内存快得多——延迟大约20-30个时钟周期。
我记得有一次做矩阵乘法优化，用共享内存做数据分块，性能直接提升了5倍。共享内存最大的坑是什么？bank conflict。说白了就是多个线程同时访问同一个bank，导致串行化。
避坑指南：我曾经在共享内存上踩过一个大坑。当时做归约操作，没注意bank conflict，结果性能还不如直接用全局内存。后来改成__shared__ float sdata[32];并做padding，才把性能提上来。
// 共享内存声明
__shared__ float shared_data[256];
// 使用共享内存做数据缓存
if (threadIdx.x < N) {
shared_data[threadIdx.x] = global_data[blockIdx.x * blockDim.x + threadIdx.x];
}
__syncthreads(); // 同步！别忘了三、寄存器：线程的“私人保险箱”
寄存器是GPU里最快的存储，延迟几乎为0。每个线程能用的寄存器数量有限，通常是32-255个（取决于计算能力）。
你想想看，如果一个线程用了太多寄存器，编译器就会把多余的数据“溢出”到本地内存——那速度就慢下来了。我一般建议，kernel函数里尽量少声明局部变量，能用const就用const。
注意：寄存器溢出是个隐形杀手。我曾经优化一个kernel，发现占用率只有25%，查了半天才发现是寄存器用太多，数据被挤到本地内存去了。减少几个临时变量后，占用率升到67%，性能翻倍。
四、本地内存：寄存器的“替补队员”
本地内存听起来像“本地”的，其实它物理上就在全局内存里。当一个线程的寄存器不够用时，编译器会自动把数据放到本地内存。访问延迟和全局内存一样高。
怎么避免？嗯，我有个小技巧：编译时加上--ptxas-options=-v参数，看看每个线程用了多少寄存器。如果超过32个，就要小心了。
// 编译时查看寄存器使用量
nvcc --ptxas-options=-v mykernel.cu五、常量内存：只读的“公告栏”
常量内存是只读的，容量64KB，但有个好处——它有专门的缓存。如果所有线程都访问同一个地址，那速度非常快。但如果每个线程访问不同地址，那就惨了，会串行化。
我在项目中常用常量内存来存查找表、滤波器系数这类数据。比如做图像处理时，把卷积核放在常量内存里，效果很好。
// 声明常量内存
__constant__ float filter[256];
// 主机端复制数据
cudaMemcpyToSymbol(filter, host_filter, sizeof(float) * 256);六、纹理内存：为空间局部性而生
纹理内存也是只读的，但它有硬件支持的插值、边界处理等功能。特别适合图像处理、随机访问模式不规则的场景。
说实话，纹理内存现在用得少了。因为全局内存的缓存机制越来越强，纹理内存的优势不那么明显了。但如果你做的是图像处理或者需要硬件插值，纹理内存还是首选。
七、延迟对比一览表
| 内存类型 | 容量 | 延迟（时钟周期） | 作用域 | 可读写 |
|---|---|---|---|---|
| 寄存器 | 每个线程最多255个 | ~1 | 单个线程 | 读/写 |
| 共享内存 | 每个线程块48-164KB | ~20-30 | 线程块内 | 读/写 |
| 本地内存 | 每个线程最多512KB | ~400-800 | 单个线程 | 读/写 |
| 常量内存 | 64KB | ~200-400（有缓存） | 全局 | 只读 |
| 纹理内存 | 取决于设备 | ~200-400（有缓存） | 全局 | 只读 |
| 全局内存 | 几十GB | ~400-800 | 全局 | 读/写 |
八、一张图看懂CUDA内存层级
下面这张SVG图，是我自己画的。它展示了CUDA内存的层级结构，从最底层的全局内存到最顶层的寄存器。你看一眼就明白了。
从这张图可以看得很清楚：越往上速度越快，容量越小；越往下速度越慢，容量越大。这就是典型的“金字塔”结构。
九、实战建议：怎么选？
说了这么多，到底怎么用？我总结了几条经验：
- 频繁访问的数据：尽量放到共享内存或寄存器里。比如卷积操作，把输入数据块缓存到共享内存。
- 只读且所有线程访问相同地址：用常量内存。比如数学常数、查找表。
- 不规则访问模式：考虑纹理内存。比如光线追踪中的随机纹理采样。
- 大容量数据：只能用全局内存，但要注意合并访问。
- 临时变量：尽量用寄存器，别声明太多数组。
我的习惯：每次写kernel之前，先画个内存访问图。哪些数据放哪里，数据怎么流动，想清楚了再动手。这比写完了再优化要省事得多。
好了，CUDA内存模型就聊到这里。记住一句话：数据在哪里，性能就在哪里。把数据放到离计算最近的地方，你的程序自然就快了。