第3章内存模型与数据管理

第3章：内存模型与数据管理：全局内存、共享内存、寄存器、本地内存、常量内存与纹理内存、cudaMalloc/cudaMemcpy实战

好，咱们进入正题。这一章，我打算聊聊CUDA的内存模型。

说实话，很多初学者写CUDA程序，上来就是cudaMalloc、cudaMemcpy，跑通了就完事了。但性能嘛，往往惨不忍睹。为什么？说白了，就是没搞懂数据到底在哪儿、怎么流动的。

我在项目中遇到过好几次，代码逻辑完全正确，但就是跑不快。最后定位下来，全是内存访问模式惹的祸。你想想看，GPU有几千个核心，如果都在等数据，那再多的计算单元也是白搭。

3.1 全局内存：最常用，但也最慢

全局内存，就是你在主机端用cudaMalloc分配的那块显存。所有线程都能访问，容量最大，但延迟也最高——大概几百个时钟周期。

嗯，这里要注意：全局内存的访问，一定要满足合并访问（coalesced access）条件。什么意思？就是相邻的线程要访问相邻的内存地址。

核心原则：让线程ID与内存地址的映射关系保持连续。

举个例子，假设你有一个float数组，每个线程读一个元素：

// 好的写法：合并访问
int idx = threadIdx.x + blockIdx.x * blockDim.x;
float val = data[idx];  // 相邻线程访问相邻地址

// 坏的写法：非合并访问
int idx = threadIdx.x + blockIdx.x * blockDim.x;
float val = data[idx * 32];  // 每个线程跳着读，带宽浪费严重

我曾经在一个图像处理项目里，就因为没注意这个，带宽利用率只有10%左右。改完合并访问后，直接提升了4倍。你说这差距大不大？

3.2 共享内存：片上的“高速缓存”

共享内存是GPU芯片上的一块SRAM，速度比全局内存快两个数量级。但它有个限制：同一个线程块内的线程才能共享。

我个人习惯，只要遇到数据复用的情况，第一反应就是往共享内存里放。比如矩阵分块乘法、归约操作，这些都是共享内存的经典应用场景。

避坑指南：我曾经因为没做同步，导致共享内存里的数据被多个线程乱读写，结果算出来的结果全是错的。记住，用共享内存一定要加__syncthreads()。

共享内存的声明方式很简单：

__shared__ float shared_data[256];  // 每个block独享一份

它的容量有限，一般也就几十KB。所以别想着把所有数据都塞进去，得精打细算。

3.3 寄存器：最快的存储，但数量有限

寄存器是每个线程私有的，速度最快，零延迟。但每个线程能用的寄存器数量是有限的——一般是255个左右（取决于GPU架构）。

如果你在kernel里声明了太多局部变量，编译器就会把多余的变量“溢出”到本地内存。本地内存虽然逻辑上是私有的，但物理上其实在全局内存里，速度一下子就掉下来了。

注意：可以用--maxrregcount编译选项限制寄存器数量，但别设得太低，否则溢出更严重。

3.4 本地内存：寄存器的“备胎”

本地内存听起来挺高级，其实就是全局内存的一部分。当寄存器不够用时，编译器会自动把变量放到这里。每个线程有自己的本地内存空间，但访问延迟和全局内存一样高。

怎么避免？少声明大数组，少用递归。我见过有人把整个结构体塞进kernel的局部变量里，结果性能惨不忍睹。

3.5 常量内存与纹理内存

常量内存适合存储只读数据，而且所有线程访问同一个地址时，硬件会广播，效率极高。纹理内存则适合二维空间局部性访问，比如图像处理中的插值操作。

我记得有一次做光线追踪，把场景数据放到常量内存里，性能提升了30%。但要注意，常量内存只有64KB，别塞太多。

__constant__ float const_data[1024];  // 声明常量内存

// 主机端复制
cudaMemcpyToSymbol(const_data, host_data, sizeof(float) * 1024);

3.6 cudaMalloc / cudaMemcpy 实战

这两个API是CUDA编程的基石。cudaMalloc负责在设备端分配全局内存，cudaMemcpy负责在主机和设备之间搬运数据。

我建议你养成一个好习惯：每次分配内存后都检查返回值，每次拷贝后都同步一下。别嫌麻烦，调试的时候能省你半天时间。

float *d_data;
cudaMalloc(&d_data, N * sizeof(float));

cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice);

// 别忘了释放
cudaFree(d_data);

小技巧：可以用cudaMallocPitch分配二维数组，它会自动对齐行宽，提高访问效率。

3.7 本章小结

这一章的内容比较多，但核心就一句话：数据在哪儿，决定了你的程序能跑多快。

全局内存：容量大，速度慢，注意合并访问
共享内存：速度快，容量小，记得同步
寄存器：最快，但别用太多
本地内存：寄存器的溢出区，尽量避免
常量/纹理内存：特定场景下的利器

下一章，我会聊聊线程束（warp）的调度机制。到时候你会发现，很多性能问题其实都跟线程束的执行方式有关。咱们下章见。

📚 课程目录