内存模型（上）：全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存
说到CUDA的内存模型，我刚开始接触的时候也觉得头大。各种内存类型，各有各的脾气。你想想看，一个GPU上有那么多计算单元，数据怎么喂给它们，怎么存中间结果，怎么输出，这背后全是内存模型的学问。
说白了，CUDA的内存模型就是一张「数据流动的地图」。你得知道数据从哪里来，到哪里去，中间经过哪些缓存。搞不清楚这个，性能优化就是瞎忙活。
核心观点：CUDA性能优化的本质，就是让数据尽可能靠近计算单元。离计算单元越近，访问速度越快，但容量越小。这是一个经典的「空间换时间」的权衡。
1. 全局内存（Global Memory）
全局内存是GPU上最大的存储空间，所有线程都能访问。但大，就意味着慢。我记得第一次做矩阵乘法优化时，没注意全局内存的访问模式，结果性能惨不忍睹。
全局内存的延迟大概在400-800个时钟周期。什么概念？就是线程等一个数据的时间，够它做几百次加法了。
关键优化点：合并访问（Coalesced Access）
全局内存有个特性：相邻的线程访问相邻的地址，效率最高。这叫合并访问。我见过太多人踩这个坑了——线程0访问地址0，线程1访问地址100，线程2访问地址200...这样每个线程都得单独发一次内存请求，效率直接打骨折。
// 好的访问模式：合并访问
__global__ void coalesced_access(float *A, float *B, int N) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < N) {
B[idx] = A[idx] * 2.0f; // 相邻线程访问相邻地址
}
}
// 差的访问模式：非合并访问
__global__ void non_coalesced_access(float *A, float *B, int N) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < N) {
B[threadIdx.x * N + blockIdx.x] = A[idx] * 2.0f; // 跨步访问
}
}我的经验：写kernel之前，先在纸上画一下线程和数据的对应关系。确保线程索引和数据索引是线性对应的。这一步花不了几分钟，但能省下后面调优的大量时间。
2. 共享内存（Shared Memory）
共享内存是CUDA编程的「瑞士军刀」。它位于芯片内部，延迟只有几个时钟周期，比全局内存快两个数量级。但容量有限，一个block通常只有几十KB。
共享内存的作用是什么？说白了就是「线程间协作」。同一个block里的线程，可以通过共享内存交换数据，避免反复访问全局内存。
我做过一个图像滤波的算子，原始版本每个线程都要从全局内存读9个像素值。用了共享内存后，先把整个tile加载到共享内存，然后每个线程只从共享内存读数据。性能提升了5倍。
// 共享内存使用示例：矩阵分块乘法
__global__ void matmul_shared(float *A, float *B, float *C, int N) {
__shared__ float As[TILE_SIZE][TILE_SIZE];
__shared__ float Bs[TILE_SIZE][TILE_SIZE];
int row = blockIdx.y * TILE_SIZE + threadIdx.y;
int col = blockIdx.x * TILE_SIZE + threadIdx.x;
float sum = 0.0f;
for (int t = 0; t < N / TILE_SIZE; t++) {
// 协作加载数据到共享内存
As[threadIdx.y][threadIdx.x] = A[row * N + t * TILE_SIZE + threadIdx.x];
Bs[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * N + col];
__syncthreads();
// 从共享内存计算
for (int k = 0; k < TILE_SIZE; k++) {
sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
}
__syncthreads();
}
C[row * N + col] = sum;
}注意：共享内存有两个「坑」——bank conflict和同步。bank conflict会导致访问串行化，性能下降。我曾经在一个项目中，因为没处理好bank conflict，共享内存的加速效果从4倍降到了1.5倍。另外，别忘了__syncthreads()，否则会出现数据竞争。
3. 寄存器（Register）
寄存器是离计算单元最近的内存，速度最快，没有延迟。每个线程都有自己的寄存器，其他线程看不到。
但寄存器数量有限。每个SM的寄存器总数是固定的，比如1024个。如果每个线程用32个寄存器，那一个SM最多同时运行32个线程。用多了，反而会降低并行度。
我见过一个极端案例：一个kernel用了128个寄存器，结果occupancy掉到了12.5%。性能还不如用全局内存做中间存储的版本。
我的建议：用--ptxas-options=-v编译选项查看寄存器使用量。如果超过32个，就要考虑是否能把一些变量放到共享内存或本地内存里。记住：寄存器不是越多越好。
4. 本地内存（Local Memory）
本地内存听起来像是「本地」的，但实际上它存在全局内存里。只是每个线程私有的，其他线程访问不到。
什么时候会用本地内存？寄存器不够用的时候。编译器会自动把「溢出」的变量放到本地内存。但本地内存的延迟和全局内存一样，400-800个时钟周期。
怎么避免？减少每个线程的变量数量。比如把大的数组放到共享内存，或者拆分成多个小kernel。
5. 常量内存（Constant Memory）
常量内存是只读的，所有线程都能访问。它有自己的缓存，大小只有64KB。但好处是，如果所有线程访问同一个地址，缓存命中率极高，延迟几乎为零。
我常用常量内存来存一些「查表」数据。比如做颜色空间转换时，把转换矩阵放到常量内存里。所有线程读同一个矩阵，缓存效率非常高。
// 常量内存声明
__constant__ float filter_kernel[256];
// 主机端初始化
cudaMemcpyToSymbol(filter_kernel, host_kernel, sizeof(float) * 256);
// kernel中使用
__global__ void apply_filter(float *input, float *output, int N) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < N) {
// 所有线程读同一个常量内存地址，缓存命中
output[idx] = input[idx] * filter_kernel[0];
}
}注意：常量内存的读取有广播机制。如果所有线程读同一个地址，一次读取就够了。但如果每个线程读不同的地址，那就得串行读取，性能会急剧下降。所以常量内存适合「所有线程读同一个值」的场景。
6. 纹理内存（Texture Memory）
纹理内存最初是为图形处理设计的，但用在通用计算里也有奇效。它有自己的缓存，对二维空间局部性有优化。而且支持硬件插值、边界处理等功能。
我做过一个图像缩放的算子，用纹理内存的硬件双线性插值功能，比自己写插值快了3倍。而且代码还更简洁。
纹理内存适合什么场景？图像处理、数据有空间局部性、需要插值或边界处理的场景。但要注意，纹理内存是只读的，不能写。
知识体系总览
下面这张图总结了CUDA内存模型的层次结构。你可以看到，从寄存器到全局内存，速度递减，容量递增。优化的核心就是：把频繁访问的数据放到离计算单元近的地方。
总结与避坑指南
好了，我们来捋一捋。CUDA内存模型的核心就一句话：数据在哪里，决定了你的kernel能跑多快。
- 寄存器：最快，但别用太多。超过32个就要警惕。
- 共享内存：block内协作的利器。注意bank conflict和同步。
- 全局内存：大而慢。一定要做合并访问。
- 常量/纹理内存：只读场景的加速器。适合查表和图像处理。
- 本地内存：寄存器的「备胎」。能不用就不用。
我曾经踩过的坑：有一次做流体模拟，我把所有中间变量都放在寄存器里，结果occupancy掉到了6%。后来把一些不常用的变量挪到共享内存，occupancy回到了50%，性能提升了8倍。所以记住：别让寄存器成为性能瓶颈。
嗯，内存模型的上半部分就讲到这里。下一章我们会深入讨论内存访问模式、bank conflict的详细分析，以及如何用NVIDIA Nsight工具来诊断内存问题。这些内容会更贴近实战，到时候我会分享更多项目中的真实案例。
想要更多资料，联系可可