全局内存访问优化：合并访问原理与带宽利用

好，咱们今天聊点实在的。全局内存访问，说白了就是GPU和显存之间的数据搬运。很多同学写CUDA程序，一上来就写个for循环，每个线程读一个数，觉得挺正常。结果一跑，带宽利用率不到10%。嗯，问题就出在——你没有做合并访问。

什么是合并访问？

合并访问，英文叫Coalesced Access。它的核心思想很简单：让相邻的线程访问相邻的内存地址。

你想想看，GPU的显存是挂在DRAM上的。每次读取数据，不是按字节来的，而是按“事务”来的。一个内存事务，通常32字节、64字节或者128字节。如果32个线程（一个warp）同时访问内存，并且它们访问的地址是连续的，那硬件就能用一个事务把数据全取回来。这叫合并。

反过来，如果每个线程跳着读，比如线程0读地址0，线程1读地址100，线程2读地址200……那硬件就得拆成多个事务。一次warp的内存请求，可能要发几十个事务。带宽自然就下来了。

我在项目中遇到过这样一个场景：一个图像处理核，输入是二维数组，我一开始按行存储、按列访问。结果带宽只有理论峰值的5%。后来改成按行访问，带宽直接飙到80%。差别就这么大。

对齐要求：别让你的地址“歪”了

合并访问还有一个前提——对齐。什么意思？就是warp访问的起始地址，最好是32字节、64字节或128字节的整数倍。

举个例子。假设你有一个float数组，每个元素4字节。如果warp的32个线程从地址0开始读，那正好对齐到128字节边界。硬件一个事务搞定。

但如果你的起始地址是4，也就是从第二个float开始读。那地址4到132，跨越了两个128字节的缓存行。硬件就得发两个事务。虽然还是合并的，但效率打了折扣。

我曾经调试过一个模型推理的kernel，发现带宽上不去。查了半天，发现是输入数据的地址没有做cudaMalloc的默认对齐。后来我手动用cudaMallocPitch或者cudaMallocManaged，对齐到256字节，问题就解决了。

内存事务大小对带宽的影响

好，咱们来拆解一下内存事务。GPU的内存控制器，每次读取的最小单位是32字节。但实际事务大小取决于访问模式：

为什么会这样？因为硬件是按缓存行来取数据的。一个缓存行128字节。如果warp的32个线程访问的地址都在同一个128字节内，那一个事务就够了。如果分散在两个缓存行，那就两个事务。如果每个线程都落在不同的128字节里……嗯，32个事务，带宽直接除以32。

我记得有一次做矩阵转置，naive实现就是每个线程读一个元素，写一个元素。结果读和写都不合并。带宽只有几个GB/s。后来我用共享内存做分块转置，读和写都变成合并的，带宽直接到了200多GB/s。这就是事务大小的威力。

如何写出合并访问的代码？

其实就一条原则：让线程ID和内存地址线性对应。

看个例子。假设你有一个一维数组data，长度为N。你想让每个线程处理一个元素：

// 好的写法：合并访问
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
    float val = data[idx];  // 线程0读data[0]，线程1读data[1]...
}

// 坏的写法：不合并
int idx = threadIdx.x * gridDim.x + blockIdx.x;
if (idx < N) {
    float val = data[idx];  // 线程0读data[0]，线程1读data[32]...
}

第一种写法，相邻线程访问相邻地址，完美合并。第二种写法，线程0读0，线程1读32，线程2读64……每个线程之间差了32个元素。这就不合并了。你想想看，32个线程访问的地址，分布在32个不同的128字节区间里。每个线程一个事务，效率极低。

对于二维数组，道理一样。假设你有一个float*表示的矩阵，行优先存储：

// 好的写法：按行访问
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
if (row < rows && col < cols) {
    float val = matrix[row * cols + col];  // 同一行的col连续
}

// 坏的写法：按列访问
int row = blockIdx.x * blockDim.x + threadIdx.x;
int col = blockIdx.y * blockDim.y + threadIdx.y;
if (row < rows && col < cols) {
    float val = matrix[row * cols + col];  // 同一列的row不连续
}

按行访问时，同一行的列索引是连续的，所以warp内线程访问的地址也是连续的。按列访问时，同一列的row索引不连续，每个线程之间差了cols个元素。如果cols很大，那就不合并了。

实际项目中的经验

我做过一个推理引擎的优化，里面有个卷积层，输入是NHWC格式。一开始我按H维度做并行，每个线程处理一个像素的所有通道。结果发现带宽只有30GB/s。后来我改成按C维度做并行，每个线程处理一个通道的一个像素，并且让warp内的线程访问连续的W维度。带宽直接到了250GB/s。

说白了，合并访问是CUDA优化的第一课。你写kernel之前，先画个图，看看warp内线程的地址分布。如果地址是连续的，那就对了。如果不是，赶紧改。

嗯，今天就聊到这儿。下一章咱们讲共享内存的bank conflict，那个坑更多。但只要你把合并访问搞明白了，后面的路就好走了。