CUDA内存模型：全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存

说到CUDA编程，内存模型这块儿，我敢说是最容易被忽视、却又最影响性能的环节。很多人一上来就写kernel，跑得慢就怪显卡不行。其实呢，很多时候是内存访问没搞对。今天咱们就把CUDA的内存体系彻底捋一遍。

内存层次结构：一张图看懂全局

先看这张图，我亲手画的。它把CUDA里所有内存类型和它们的关系讲得明明白白。

这张图我每次培训都会拿出来。你看，从外到内，速度越来越快，容量越来越小。全局内存在最外层，是所有线程都能访问的「大仓库」。但访问一次要几百个时钟周期，慢得让人心疼。

全局内存：最常用也最容易被坑

全局内存就是显存。你在CPU端用cudaMalloc分配的那块空间，就是它。所有线程都能读写，但延迟很高。

合并访问是全局内存性能的关键。什么意思呢？就是让同一个warp里的32个线程，访问连续的内存地址。硬件会把多个访问合并成一次或几次事务。

// ✅ 合并访问：线程0访问data[0]，线程1访问data[1]...
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[idx];  // 连续访问，效率高

// ❌ 非合并访问：线程0访问data[0]，线程1访问data[stride]...
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[idx * stride];  // 跨步访问，效率低

我有个项目做矩阵乘法，一开始没注意合并访问，性能只有理论峰值的10%。后来把数据布局从行主序改成列主序，配合正确的索引方式，性能直接翻了6倍。嗯，合并访问就是这么重要。

共享内存：程序员手中的利器

共享内存是片上的，速度比全局内存快一个数量级。但它需要你手动管理——说白了就是你自己负责把数据从全局内存搬进来，算完再搬回去。

共享内存按bank组织。每个bank每周期只能服务一个地址。如果多个线程访问同一个bank的不同地址，就发生bank conflict，访问会串行化。

// 共享内存声明
__shared__ float s_data[256];

// 避免bank conflict的技巧
// 如果每个线程访问s_data[threadIdx.x]，没有冲突
// 但如果访问s_data[threadIdx.x * 2]，就可能产生冲突
// 可以加padding：__shared__ float s_data[256 + 1];

寄存器与本地内存

寄存器是CUDA里最快的存储，没有之一。每个线程的私有变量，只要放得下，编译器都会优先分配到寄存器。

但寄存器数量有限。每个SM的寄存器总数是固定的（比如Kepler架构是65536个）。如果每个线程用的寄存器太多，就会降低occupancy（驻留的warp数量）。

本地内存是个坑。它名义上是「本地」的，实际上存储在全局内存里。当寄存器不够用，或者访问数组时索引不是编译期常量，编译器就会把变量「溢出」到本地内存。你想想看，本来以为很快的变量，突然变成了几百个周期的访问，性能能不崩吗？

常量内存与纹理内存

常量内存容量只有64KB，但有专门的缓存。当所有线程访问同一个常量地址时，速度极快——只需要一个周期。但如果每个线程访问不同的地址，那就惨了，会串行化。

纹理内存呢，我主要用在图像处理上。它支持硬件插值，而且对二维空间局部性有优化。比如做图像缩放、矩阵转置，用纹理内存往往比全局内存快不少。

内存访问模式：别让硬件白干活

最后聊一下访问模式。GPU是SIMT架构，一个warp里的32个线程执行同一条指令。如果它们访问的内存地址是连续的、对齐的，硬件就能用最少的事务搞定。

我个人习惯在写kernel之前，先画个内存访问图。把线程ID和地址的对应关系画出来，一眼就能看出是不是合并访问。这招虽然土，但特别管用。

内存模型这块儿，说白了就是「数据在哪、怎么搬、怎么访问」三个问题。搞明白了，你的CUDA程序性能至少能翻一倍。我见过太多人在这上面栽跟头了，希望你能少走弯路。

想要更多资料，联系可可