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

说到CUDA编程，内存模型这块儿，我敢说——谁搞懂了内存，谁就掌握了性能的钥匙。

你想想看，GPU里有成千上万个线程同时跑。它们怎么读写数据？数据存在哪儿？读写快不快？这些问题，全跟内存模型有关。我刚开始学CUDA那会儿，就踩过一个坑：写了个核函数，逻辑完全正确，结果跑起来比CPU还慢。后来一查，问题出在全局内存的访问模式上。嗯，从那以后，我再也不敢轻视内存模型了。

今天我们就来拆解CUDA的六大内存类型。我会结合项目经验，告诉你每种内存该怎么用、什么时候用、千万别怎么用。

3.1 全局内存（Global Memory）

全局内存是GPU里最大的存储空间，所有线程都能访问。说白了，它就是GPU的“主存”。

特点：

• 容量大：一般几个GB到几十个GB
• 延迟高：访问一次大概400-800个时钟周期
• 所有线程（包括CPU）都能读写

我个人的习惯是：能用共享内存解决的问题，绝不用全局内存。但数据量太大时，全局内存是唯一的选择。

// 好的访问模式：合并访问
__global__ void coalesced_kernel(float *data) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    data[tid] = tid * 2.0f;  // 相邻线程访问相邻地址
}

// 坏的访问模式：非合并访问
__global__ void non_coalesced_kernel(float *data) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    data[tid * 32] = tid * 2.0f;  // 相邻线程访问间隔32的地址
}

我在项目中遇到过：一个图像处理核函数，原本用非合并访问，带宽利用率只有15%。改成合并访问后，直接飙到85%。性能提升了5倍多。你想想看，就改了一行代码的事。

3.2 共享内存（Shared Memory）

共享内存是CUDA编程的“杀手锏”。它位于芯片内部，延迟极低（大约20-30个时钟周期），比全局内存快一个数量级。

特点：

• 容量小：每个block通常只有几十KB
• 速度快：接近寄存器速度
• 同一block内的线程共享
• 需要手动管理数据搬入搬出

说白了，共享内存就是程序员手里的“缓存”。你得自己决定把哪些数据放进去，什么时候放进去。

// 使用共享内存做矩阵转置
__global__ void transpose_kernel(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();  // 同步！确保所有线程都加载完毕
    
    // 从共享内存写回全局内存（转置）
    out[x * width + y] = tile[threadIdx.y][threadIdx.x];
}

3.3 寄存器（Register）

寄存器是GPU里最快的内存，没有之一。每个线程都有自己的寄存器，其他线程访问不到。

特点：

• 速度：零延迟（一个时钟周期内完成）
• 容量：每个线程最多255个寄存器（不同架构有差异）
• 私有：只有当前线程能访问

你写的局部变量，只要编译器觉得合适，就会放到寄存器里。但要注意：寄存器数量是有限的。如果每个线程用的寄存器太多，就会发生“寄存器溢出”——数据被挤到本地内存里，性能骤降。

3.4 本地内存（Local Memory）

本地内存听起来像“本地”的，其实它跟全局内存是同一块物理存储。只是编译器把它伪装成“每个线程私有”的样子。

什么时候会用本地内存？

• 寄存器不够用的时候（寄存器溢出）
• 访问大型数组或结构体
• 编译器无法确定索引的数组访问

嗯，这里要注意：本地内存的延迟跟全局内存一样高。所以尽量避免使用它。怎么避免？少用局部数组，多用标量变量。

3.5 常量内存（Constant Memory）

常量内存是专门存“只读数据”的。它有两大特点：一是只读，二是带缓存。

特点：

• 容量：64KB（所有架构都一样）
• 速度：如果所有线程访问同一个地址，几乎跟寄存器一样快
• 缓存：硬件会自动缓存常量数据

// 声明常量内存
__constant__ float filter[256];

// 主机端初始化
cudaMemcpyToSymbol(filter, host_filter, sizeof(float) * 256);

// 核函数中使用
__global__ void apply_filter_kernel(float *data) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    data[tid] = data[tid] * filter[tid % 256];  // 所有线程访问filter
}

我建议：当你有只读数据，而且所有线程都访问相同地址时，优先考虑常量内存。比如卷积核的权重、查找表等。

3.6 纹理内存（Texture Memory）

纹理内存最初是为图形处理设计的，但在通用计算里也很有用。它也有缓存，而且支持一些特殊功能。

特点：

• 有缓存：对空间局部性好的访问模式很友好
• 支持插值：硬件自动做线性插值
• 支持边界处理：可以设置“超出边界返回0”等模式
• 适合二维/三维数据访问

说实话，纹理内存在通用计算里用得不多。但在图像处理、科学计算（比如网格插值）这些场景下，它比全局内存快不少。

3.7 内存模型全景图

说了这么多，我们来画一张图，把六种内存的关系理清楚。

3.8 总结与选择策略

好了，六种内存都讲完了。怎么选？我总结了一个简单的决策流程：

1. 数据是只读的吗？ → 是 → 所有线程访问相同地址？→ 是 → 用常量内存
2. 数据是只读的吗？ → 是 → 需要插值或边界处理？→ 是 → 用纹理内存
3. 数据需要被同一Block的线程共享吗？ → 是 → 用共享内存
4. 数据是临时变量，且很小？ → 用寄存器（注意不要溢出）
5. 其他情况 → 用全局内存，但注意合并访问

最后说一句：内存模型是CUDA性能优化的基石。你花一周时间搞懂它，后面优化代码时能省下一个月的时间。我当年就是吃了这个亏，走了不少弯路。希望你能少踩这些坑。