CUDA内存模型：全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存
聊到CUDA编程，内存模型是绕不开的核心话题。我刚开始接触CUDA时，总觉得把数据扔到GPU上就能自动跑得快——结果被现实狠狠教育了一顿。后来才明白，不理解内存层次结构，写出来的核函数性能可能还不如CPU串行。
说白了，GPU的内存不是一块铁板。它有多种类型，每种都有不同的速度、容量和用途。你想想看，如果所有数据都放在同一个地方，那访存瓶颈会卡死整个程序。所以英伟达设计了这套分层的内存体系。
全局内存（Global Memory）
全局内存是GPU中容量最大的存储空间，通常有数GB甚至数十GB。所有线程都能访问它，但代价是——慢。延迟通常在400-800个时钟周期。
我在项目中遇到过这样一个坑：有个矩阵乘法的核函数，数据全放在全局内存里，结果性能还不如CPU的OpenBLAS。后来一分析，访存带宽根本没跑满，因为全局内存的访问模式不对。
关键点：全局内存的访问必须满足合并访问（Coalesced Access）。也就是说，相邻线程要访问相邻的地址。这样才能让一个内存事务把多个数据一次性搬进来。
// 好的合并访问：线程0访问data[0]，线程1访问data[1]...
int idx = threadIdx.x + blockIdx.x * blockDim.x;
float val = data[idx];
// 差的合并访问：线程0访问data[0]，线程1访问data[1000]...
int idx = threadIdx.x * 1000 + blockIdx.x;
float val = data[idx];嗯，这里要注意：全局内存默认是没有缓存的（除非你用L1/L2缓存，但那是另一回事）。每次读写都直接走DRAM，所以一定要尽量减少全局内存的访问次数。
共享内存（Shared Memory）
共享内存是我个人最喜欢的存储类型。它位于每个线程块内部，所有线程块内的线程都能访问。速度极快，延迟只有几个时钟周期，比全局内存快两个数量级。
但共享内存的容量很小——每个SM（流式多处理器）通常只有几十到一百多KB。所以你得精打细算。
我习惯把共享内存当作手动管理的缓存。比如做矩阵分块时，先把数据从全局内存搬到共享内存，然后让线程块内的所有线程反复使用这些数据。
__shared__ float tile[BLOCK_SIZE][BLOCK_SIZE];
// 从全局内存加载到共享内存
tile[threadIdx.y][threadIdx.x] = A[row * WIDTH + col];
__syncthreads(); // 同步！确保所有线程都加载完毕
// 然后从共享内存中读取，速度飞快
float val = tile[threadIdx.y][threadIdx.x];注意：共享内存使用不当会导致Bank Conflict。共享内存被分成32个Bank，如果多个线程同时访问同一个Bank的不同地址，就会发生冲突，导致串行化。我曾经因为这个，一个核函数性能直接腰斩。
避免Bank Conflict的技巧：适当填充数据，或者调整访问模式。比如二维数组每行多加一个元素做padding，就能让相邻线程访问不同的Bank。
寄存器（Register）
寄存器是GPU中最快的存储单元，没有之一。每个线程私有的，延迟为0（一个时钟周期内完成）。
但寄存器数量极其有限。每个SM的寄存器总数是固定的（比如Kepler架构是65536个），分配到每个线程的数量取决于线程数。如果每个线程用太多寄存器，就会降低SM上同时运行的线程块数量——这叫寄存器压力。
我遇到过这种情况：一个核函数性能上不去，用nvcc的--ptxas-options=-v一看，每个线程用了64个寄存器。结果SM上只能同时跑很少的线程，隐藏延迟的能力大大下降。
建议：可以用__launch_bounds__来限制每个线程的最大寄存器数，让编译器帮你做寄存器溢出（spill）到本地内存。虽然本地内存慢，但有时候比降低并行度要好。
本地内存（Local Memory）
本地内存听起来像是“本地”的，其实它物理上位于全局内存中。只是每个线程私有的，编译器自动管理。
什么时候会用本地内存？当寄存器不够用的时候。比如数组索引不是编译期常量、或者寄存器溢出时，编译器会把数据放到本地内存里。
本地内存的延迟和全局内存一样高，而且没有缓存（除非L1/L2）。所以，尽量减少本地内存的使用。我个人的经验是：如果发现核函数性能异常低，先检查是不是有大量本地内存访问。
// 这个数组可能被放到本地内存（如果寄存器不够）
int temp[10];
for (int i = 0; i < 10; i++) {
temp[i] = data[i] * 2;
}你可以用cuobjdump或者NVIDIA Nsight查看生成的PTX代码，看看有没有st.local或ld.local指令。如果有，说明有本地内存访问。
常量内存（Constant Memory）
常量内存是全局内存的一个特殊区域，大小为64KB。它被所有线程共享，但只能读取，不能写入（在核函数内）。
常量内存最大的特点是：当所有线程访问同一个地址时，速度极快（相当于寄存器速度）。但如果每个线程访问不同的地址，性能会急剧下降。
我习惯把一些固定的参数放在常量内存里，比如滤波器的系数、数学常数等。用__constant__修饰符声明，然后在主机端用cudaMemcpyToSymbol复制数据。
__constant__ float filter[256];
// 主机端
float h_filter[256];
// ... 初始化h_filter ...
cudaMemcpyToSymbol(filter, h_filter, sizeof(float) * 256);适用场景：所有线程都读取相同的值。比如做图像处理时，所有像素点都用同一个卷积核。这时候常量内存比全局内存快得多。
纹理内存（Texture Memory）
纹理内存最初是为图形学设计的，但后来在通用计算中也很有用。它也是全局内存的一种特殊形式，有自己独立的缓存（纹理缓存）。
纹理内存有几个独特的能力：
- 硬件插值：可以自动做线性插值，适合图像缩放、采样
- 边界处理：支持多种寻址模式（clamp、wrap、mirror等）
- 缓存优化：对二维空间局部性有优化，适合随机访问
我在做医学图像处理时用过纹理内存。CT图像是三维的，但切片是二维的。用纹理内存做双线性插值，比手动实现快了将近一倍。
// 声明纹理引用
texture<float, 2, cudaReadModeElementType> tex;
// 绑定纹理
cudaBindTexture2D(NULL, tex, d_data, &tex_channel_desc, width, height, pitch);
// 在核函数中采样
float val = tex2D(tex, x, y);小技巧：如果你的数据访问模式是非合并的（比如随机访问），纹理内存往往比全局内存快。因为它有专门的缓存，而且对二维空间局部性更友好。
各内存类型对比
| 内存类型 | 位置 | 访问范围 | 延迟 | 容量 | 生命周期 |
|---|---|---|---|---|---|
| 寄存器 | SM内部 | 单个线程 | 0周期 | 有限（每个SM数千个） | 核函数执行期间 |
| 共享内存 | SM内部 | 线程块内 | 几个周期 | 几十KB | 线程块执行期间 |
| 本地内存 | 全局内存 | 单个线程 | 400-800周期 | 有限（寄存器溢出） | 核函数执行期间 |
| 全局内存 | 板载DRAM | 所有线程+主机 | 400-800周期 | 数GB | 程序运行期间 |
| 常量内存 | 板载DRAM | 所有线程（只读） | 几个周期（同地址） | 64KB | 程序运行期间 |
| 纹理内存 | 板载DRAM | 所有线程（只读） | 取决于缓存命中 | 同全局内存 | 程序运行期间 |
总结一下我的经验
写CUDA程序，说白了就是在管理数据流动。你要把数据从慢的存储（全局内存）搬到快的存储（共享内存、寄存器），然后尽可能多地复用。
我一般遵循这几个原则：
- 尽量减少全局内存访问——能合并就合并，能缓存就缓存
- 共享内存是手动缓存——分块、tiling、reduction都用它
- 寄存器是黄金——但别用太多，否则影响并行度
- 常量内存和纹理内存是特种兵——特定场景下比全局内存强很多
嗯，最后说一句：别光看理论。我建议你写几个小实验，用nvprof或Nsight Compute看看实际的内存访问模式。只有亲手调过，才能真正理解这些内存类型之间的差异。