4. CUDA内存模型：全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存的特点与使用场景
说到CUDA编程，内存模型是绕不开的核心话题。我刚开始学的时候，觉得不就是几种内存嘛，用就是了。结果写出来的程序跑得比CPU还慢，被导师骂了一顿。后来才明白——不理解内存模型，你写的CUDA代码就是碰运气。
CUDA的内存种类挺多的，但别怕。咱们一个一个捋清楚。每种内存都有它的脾气，用对了地方，性能起飞；用错了，嗯，你会怀疑人生的。
一句话总结：CUDA内存优化的本质，就是把数据放在离计算单元最近的地方，尽量减少慢速内存的访问次数。
4.1 全局内存（Global Memory）
全局内存是GPU里最大的存储空间，所有线程都能访问。说白了，它就像你的硬盘——容量大，但慢。
特点：
- 容量最大，通常几个GB到几十个GB
- 延迟最高，大约400-800个时钟周期
- 所有线程（包括CPU）都能读写
- 没有缓存（或者说L2缓存很小）
使用场景：
- 存放输入数据和最终结果
- 线程间需要交换数据时（但效率低）
- 数据量太大，放不进其他内存时
我的经验：全局内存访问要尽量合并。什么意思呢？就是相邻线程访问相邻地址。我见过一个项目，就因为没注意合并访问，性能差了10倍。后来改成连续访问，速度立马就上来了。
// 好的访问模式：合并访问
__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.x * 100 + blockIdx.x;
float val = data[idx]; // 线程之间隔了100个元素
}4.2 共享内存（Shared Memory）
共享内存是CUDA编程的杀手锏。它位于芯片内部，速度极快，但容量很小（通常几十KB）。同一个线程块内的线程可以共享这块内存。
特点：
- 延迟极低，大约1-2个时钟周期
- 容量小，每个SM通常只有几十KB
- 同一个block内的线程可以读写
- 需要手动管理数据加载和同步
使用场景：
- 数据重用：多个线程需要反复访问同一块数据
- 线程间协作：比如做归约、扫描等操作
- 全局内存的缓存：先把数据从全局内存搬到共享内存
核心技巧：共享内存最怕的就是bank conflict。我踩过这个坑——明明用了共享内存，性能反而下降了。后来一查，是bank conflict导致串行化访问。解决办法就是加padding或者调整数据布局。
// 共享内存做矩阵转置的经典用法
__global__ void transpose(float *in, float *out, int width) {
__shared__ float tile[TILE_SIZE][TILE_SIZE];
int x = blockIdx.x * TILE_SIZE + threadIdx.x;
int y = blockIdx.y * TILE_SIZE + threadIdx.y;
// 加载到共享内存
tile[threadIdx.y][threadIdx.x] = in[y * width + x];
__syncthreads(); // 同步！别忘了
// 从共享内存读出并转置
x = blockIdx.y * TILE_SIZE + threadIdx.x;
y = blockIdx.x * TILE_SIZE + threadIdx.y;
out[y * width + x] = tile[threadIdx.x][threadIdx.y];
}4.3 寄存器（Register）
寄存器是GPU里最快的内存，没有之一。每个线程独享，编译器自动分配。但数量有限——一个SM的寄存器总数是固定的，线程越多，每个线程分到的就越少。
特点：
- 零延迟，一个时钟周期就能访问
- 每个线程私有，编译器自动管理
- 数量有限，每个线程最多255个（不同架构有差异）
- 寄存器溢出时会使用本地内存（性能骤降）
使用场景：
- 局部变量、循环计数器、临时计算结果
- 不需要在线程间共享的数据
注意：我曾经犯过一个错误——在kernel里声明了太多局部变量，导致寄存器溢出。编译器默默地把这些变量放到了本地内存里，性能直接腰斩。后来我用--maxrregcount编译选项限制寄存器数量，反而让occupancy上去了，性能提升了30%。
4.4 本地内存（Local Memory）
本地内存听起来像是片上的，其实不是。它和全局内存一样在显存里，只是每个线程私有的。编译器在寄存器不够用的时候，会把变量溢出到这里。
特点：
- 延迟高，和全局内存差不多
- 每个线程私有，其他线程看不到
- 编译器自动使用，开发者一般不需要手动管理
- 通常用于数组索引访问或者寄存器溢出
使用场景：
- 大型局部数组（比如声明了一个
float arr[100]） - 编译器判断寄存器不够用时
避坑指南：怎么知道有没有用本地内存？编译时加--ptxas-options=-v选项，看输出里的lmem字段。如果数值很大，说明寄存器溢出了，得优化代码。
4.5 常量内存（Constant Memory）
常量内存是全局内存的一种特殊形式，但它有硬件缓存（常量缓存）。适合存放那些不会改变、且所有线程都需要的只读数据。
特点：
- 有专门的常量缓存，访问速度快
- 容量小，通常64KB
- 只读，所有线程访问同一地址
- 同一个warp内的线程访问同一地址时，速度最快
使用场景：
- 查找表、系数、参数等只读数据
- 所有线程都需要访问的固定数据
// 常量内存声明
__constant__ float coeff[256];
// 主机端初始化
float h_coeff[256] = { ... };
cudaMemcpyToSymbol(coeff, h_coeff, sizeof(h_coeff));
// kernel中使用
__global__ void apply_filter(float *data) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
data[idx] = data[idx] * coeff[idx % 256]; // 所有线程访问同一地址
}我的建议：常量内存最适合warp内所有线程访问同一地址的场景。如果每个线程访问不同地址，常量缓存反而会串行化，性能还不如全局内存。我做过一个滤波器，一开始用常量内存存系数，结果每个线程访问不同系数，慢得要命。后来改成共享内存，速度翻倍。
4.6 纹理内存（Texture Memory）
纹理内存最初是为图形处理设计的，但在通用计算里也有妙用。它有自己的缓存，支持硬件插值和边界处理。
特点：
- 有纹理缓存，对空间局部性好的访问模式友好
- 支持硬件线性插值（做图像处理时很爽）
- 支持边界处理（clamp、wrap等模式）
- 只读
使用场景：
- 图像处理、计算机视觉
- 需要硬件插值的科学计算
- 非对齐访问模式（纹理缓存能处理）
我记得：有一次做医学图像重建，数据访问模式非常随机，全局内存的缓存命中率极低。换成纹理内存后，利用它的2D空间局部性缓存，性能提升了3倍。纹理内存对那种"邻居访问"模式特别友好。
4.7 各内存对比一览
| 内存类型 | 位置 | 延迟 | 容量 | 作用域 | 生命周期 |
|---|---|---|---|---|---|
| 寄存器 | 片上 | ~1 cycle | ~256个/线程 | 线程私有 | kernel执行期间 |
| 共享内存 | 片上 | ~2 cycles | ~48KB/SM | block内共享 | block执行期间 |
| 本地内存 | 显存 | ~500 cycles | 受显存限制 | 线程私有 | kernel执行期间 |
| 全局内存 | 显存 | ~500 cycles | GB级别 | 所有线程+CPU | 程序运行期间 |
| 常量内存 | 显存+缓存 | ~5 cycles（缓存命中） | 64KB | 所有线程只读 | 程序运行期间 |
| 纹理内存 | 显存+缓存 | ~100 cycles（缓存命中） | 受显存限制 | 所有线程只读 | 程序运行期间 |
4.8 内存选择策略
说了这么多，到底怎么选？我总结了一个简单的决策流程：
- 数据是否只读且所有线程访问同一地址？ → 常量内存
- 数据是否只读且访问模式不规则？ → 纹理内存
- 数据是否需要在线程块内共享？ → 共享内存
- 数据是否线程私有且量小？ → 寄存器
- 其他情况？ → 全局内存
最后提醒一句：别为了用共享内存而用共享内存。我见过有人把只读一次的数据也搬到共享内存里，结果多了两次拷贝，反而更慢。共享内存的精髓在于数据重用——同一个数据被多个线程多次访问，这时候用共享内存才划算。
好了，CUDA的六种内存就讲到这里。每种内存都有它的脾气，用对了就是利器，用错了就是坑。多写代码多踩坑，慢慢就找到感觉了。
公众号:蓝海资料掘金营，微信deep3321