第一章：CUDA内存模型概述

各位同学，今天我们来聊聊CUDA内存模型。说实话，这是整个CUDA编程里最基础也最关键的部分。我见过太多人一上来就写核函数，结果性能惨不忍睹，最后发现是内存访问模式出了问题。

CUDA的内存模型，说白了就是一张「存储层次图」。CPU有L1、L2、L3缓存，GPU也有类似的结构，但更复杂一些。咱们先看个总览：

内存类型	位置	访问速度	生命周期	作用域
寄存器	SM内部	最快（1周期）	线程生命周期	单个线程
共享内存	SM内部	快（~5周期）	线程块生命周期	线程块内
本地内存	显存	慢（~100周期）	线程生命周期	单个线程
全局内存	显存	慢（~400周期）	应用程序生命周期	所有线程
常量内存	显存（带缓存）	较快（缓存命中时）	应用程序生命周期	所有线程
纹理内存	显存（带缓存）	较快（空间局部性好）	应用程序生命周期	所有线程

1. 全局内存

全局内存是GPU里最大的存储区域，也是我们最常打交道的地方。它位于显存上，所有线程都能访问。但代价就是——慢。我记得第一次做矩阵乘法时，没注意全局内存的合并访问，结果性能只有理论值的十分之一。

举个例子：

// 好的访问模式：合并访问
float value = data[threadIdx.x];  // 线程0访问data[0]，线程1访问data[1]...

// 坏的访问模式：非合并访问
float value = data[threadIdx.x * 1024];  // 线程0访问data[0]，线程1访问data[1024]...

你想想看，如果每个线程都跳着访问，硬件就得为每个线程单独发一次内存请求，那效率能高吗？

2. 共享内存

共享内存是我个人最喜欢的GPU内存类型。它位于SM内部，速度比全局内存快两个数量级。但容量有限——每个SM通常只有几十KB到一百多KB。

共享内存的作用是什么？说白了就是「线程块内的数据交换站」。线程块内的所有线程都能读写它，而且速度极快。

但要注意——共享内存有bank冲突。嗯，这是个坑。我曾经在一个图像处理项目里，因为没处理好bank冲突，性能直接腰斩。

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

// 使用共享内存做归约
__syncthreads();  // 别忘了同步！

3. 寄存器

寄存器是GPU里最快的内存，没有之一。每个线程都有自己的寄存器文件，编译器会尽量把局部变量分配到寄存器里。

但寄存器数量有限。每个SM的寄存器总数是固定的，比如有些架构是65536个。如果每个线程用太多寄存器，就会导致SM能同时调度的线程数减少，这叫「寄存器压力」。

怎么查看寄存器使用量？用nvcc的编译选项：

nvcc --ptxas-options=-v mykernel.cu

它会告诉你每个线程用了多少寄存器。

4. 本地内存

本地内存听起来像「本地」的，但其实它位于显存上。当寄存器不够用时，编译器会把一些变量「溢出」到本地内存。说白了，就是寄存器的「备胎」。

为什么会发生这种情况？比如你声明了一个大数组：

float big_array[100];  // 寄存器放不下，会放到本地内存

本地内存的访问延迟和全局内存差不多，所以尽量避免。我个人的习惯是：如果发现核函数里用了很多本地内存，就重新审视算法，看看能不能减少每个线程的工作量。

5. 常量内存

常量内存有点特殊。它位于显存，但带有缓存。适合存储那些所有线程都读取的、不会变化的数据。比如物理常数、查找表等。

常量内存的缓存是广播式的——如果所有线程访问同一个地址，那速度极快。但如果每个线程访问不同的地址，那就退化成普通全局内存了。

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

// 在主机端初始化
cudaMemcpyToSymbol(coefficients, host_data, sizeof(float) * 256);

我记得有一次做光线追踪，把材质参数放在常量内存里，性能提升了30%。因为所有线程都在读同一组参数，缓存命中率极高。

6. 纹理内存

纹理内存最初是为图形学设计的，但后来也被广泛用于通用计算。它也有缓存，而且对二维空间局部性特别友好。

什么时候用纹理内存？当你需要做插值、边界处理、或者访问模式不规则时。比如图像处理中的双线性插值，纹理内存硬件直接支持。

使用纹理内存需要绑定纹理对象：

// 创建纹理对象
cudaTextureObject_t texObj;
cudaResourceDesc resDesc;
// ... 配置参数

// 在核函数中使用
float value = tex2D(texObj, x, y);

总结一下

好了，咱们把CUDA的六种内存类型都过了一遍。你想想看，其实核心就一句话：把频繁访问的数据放在快的内存里，把大块但少访问的数据放在慢的内存里。

实际项目中，我一般这样选型：

• 线程私有、频繁使用 → 寄存器
• 线程块内共享、需要协作 → 共享内存
• 所有线程都读、数据不变 → 常量内存
• 空间局部性好、需要插值 → 纹理内存
• 其他情况 → 全局内存（注意合并访问）

下一章咱们会深入讲共享内存的bank冲突和优化技巧。嗯，那才是真正考验功力的地方。

公众号:蓝海资料掘金营，微信deep3321