一、CUDA内存模型全景：从全局内存到寄存器，一张图看懂GPU存储层次

做GPU编程这些年，我见过太多人一上来就写cudaMalloc，然后抱怨程序跑得慢。其实啊，很多时候问题不在算法，而在内存——你根本没用对存储层次。

今天我们就来彻底搞懂CUDA的内存模型。说白了，GPU的存储结构就像一座金字塔，从顶层的寄存器到底层的全局内存，速度差了上百倍。不理解这个，你的优化就无从谈起。

1.1 GPU存储层次概览

先看这张图，我亲手画的，把CUDA里所有内存类型都串起来了：

这张图我每次培训都会拿出来讲。你看，越靠近顶层的存储，速度越快、容量越小、每个线程的私有程度越高。底层正好相反。

1.2 全局内存（Global Memory）

全局内存是GPU里最大的存储区域，通常有4GB到80GB不等。所有线程都能读写它，但代价是——慢。

我在项目中遇到过最典型的坑：新手把频繁访问的变量也放在全局内存里。你想想看，每次访问都要等几百个周期，性能能好才怪。

1.3 共享内存（Shared Memory）

共享内存是CUDA编程的精华所在。它位于芯片上，延迟只有全局内存的十分之一左右。同一个线程块里的所有线程都能访问它。

属性	值
大小	每个SM 48KB~164KB（取决于架构）
延迟	约20-30个时钟周期
作用域	线程块内共享
生命周期	线程块执行期间

共享内存最经典的用法就是矩阵分块。把数据先搬到共享内存里，让块内线程反复使用，避免反复访问全局内存。

// 共享内存声明示例
__global__ void kernel() {
    __shared__ float cache[256];  // 每个块独享
    
    int tid = threadIdx.x;
    cache[tid] = global_data[tid];
    __syncthreads();  // 同步！确保所有线程都写完了
    
    // 现在可以快速访问cache了
    float val = cache[tid];
}

1.4 寄存器（Registers）

寄存器是GPU里最快的存储，没有之一。每个线程私有的寄存器，延迟几乎为零。但数量有限——每个线程最多255个（不同架构略有差异）。

编译器会自动把局部变量分配到寄存器。但如果你声明了太多变量，就会发生「寄存器溢出」，变量被挤到局部内存里——那速度就一落千丈了。

1.5 常量内存与纹理内存

这两个属于「特殊用途」的内存类型，用对了很香。

• 常量内存（Constant Memory）：只读，64KB，有缓存。适合存放系数、查找表这类数据。所有线程读同一地址时效率极高。
• 纹理内存（Texture Memory）：有硬件插值支持，适合图像处理。二维局部性访问时表现优异。

我记得有个物理模拟项目，把粒子碰撞的系数表放到常量内存里，配合广播机制，性能提升了40%。

1.6 局部内存（Local Memory）

局部内存听起来像「局部」的，其实它物理上就在全局内存里。只是每个线程有自己的私有空间。编译器在寄存器不够用时会自动把变量 spill 到这里。

1.7 缓存体系：L1和L2

L1缓存每个SM独有，L2缓存所有SM共享。它们对程序员是透明的，但理解它们的行为有助于写出缓存友好的代码。

缓存	大小	作用域	用途
L1	128KB/SM（可配置）	每个SM	缓存全局/局部内存访问
L2	1~6MB	所有SM共享	缓存全局内存访问

我个人习惯在访问全局内存时，尽量让相邻线程访问相邻地址。这样能充分利用缓存行（cache line），减少实际的内存请求次数。

1.8 一张表总结所有内存类型

内存类型	位置	缓存	访问权限	生命周期
寄存器	片上	无	单个线程读写	线程内
局部内存	片外DRAM	L1/L2	单个线程读写	线程内
共享内存	片上	无	块内所有线程读写	块内
全局内存	片外DRAM	L2	所有线程+主机	应用程序
常量内存	片外DRAM	专用常量缓存	所有线程只读	应用程序
纹理内存	片外DRAM	专用纹理缓存	所有线程只读	应用程序

1.9 实战建议：怎么选内存？

说白了，选内存就三条原则：

1. 能放寄存器就别放共享内存——但别溢出
2. 能放共享内存就别放全局内存——尤其数据要复用的时候
3. 只读数据优先考虑常量内存——64KB够用的话

嗯，CUDA的内存模型就这些。别看种类多，其实核心就一句话：把数据放在离计算最近的地方。理解了这一点，你的优化方向就不会跑偏。

---

想要更多资料，联系可可