全局内存合并访问：Coalescing 原理与实战

说到 GPU 显存优化，全局内存合并访问绝对是个绕不开的话题。我刚开始写 CUDA 程序那会儿，总觉得显存带宽不够用，后来才发现——不是带宽不够，是我没用好。

说白了，全局内存合并访问就是让 GPU 的一次内存请求，能尽可能多地加载有用的数据。你想想看，GPU 访问显存是以 128 字节为单位的，如果你只用了其中 4 个字节，那剩下的 124 个字节就白白浪费了。这就是带宽利用率低下的根源。

合并访问的原理

先说说硬件是怎么工作的。当一个 warp（32 个线程）执行一条加载指令时，GPU 会把这些线程的访存请求合并成尽可能少的内存事务。理想情况下，32 个线程访问连续的 128 字节，只需要一次事务就能搞定。

我遇到过最典型的反面教材：一个 warp 里的线程访问地址是 0, 32, 64, 96... 这种跨步访问。结果呢？一次请求被拆成了 32 次独立的事务，带宽利用率直接掉到 3%。

对齐与步长的影响

对齐是什么意思？举个例子：

// 对齐访问 - 好
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[idx];  // 每个线程访问连续地址

// 未对齐访问 - 差
int idx = blockIdx.x * blockDim.x + threadIdx.x + 1;
float val = data[idx];  // 偏移了 4 字节，导致跨边界

步长的影响更隐蔽。我做过一个实验，用不同的步长访问同一个数组：

为什么会这样？因为步长 2 意味着每个线程只用到 8 字节，但 GPU 还是得加载 128 字节。剩下的 120 字节？全浪费了。

矩阵转置的合并访问优化

矩阵转置是个经典案例。先看一个 naive 的实现：

__global__ void transpose_naive(float *in, float *out, int n) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (row < n && col < n) {
        out[col * n + row] = in[row * n + col];
    }
}

这个写法有个大问题：读的时候是合并的（行连续），但写的时候是跨步的（列连续）。我实测过，对于 4096x4096 的矩阵，带宽利用率只有 15% 左右。

怎么优化？用共享内存做中转：

__global__ void transpose_coalesced(float *in, float *out, int n) {
    __shared__ float tile[TILE_DIM][TILE_DIM];
    
    int row = blockIdx.y * TILE_DIM + threadIdx.y;
    int col = blockIdx.x * TILE_DIM + threadIdx.x;
    
    if (row < n && col < n) {
        // 合并读
        tile[threadIdx.y][threadIdx.x] = in[row * n + col];
    }
    
    __syncthreads();
    
    // 交换行列索引
    int new_row = blockIdx.x * TILE_DIM + threadIdx.y;
    int new_col = blockIdx.y * TILE_DIM + threadIdx.x;
    
    if (new_row < n && new_col < n) {
        // 合并写
        out[new_row * n + new_col] = tile[threadIdx.x][threadIdx.y];
    }
}

这个版本的关键在于：读和写都是合并的。读的时候按行连续，写的时候也按行连续。中间通过共享内存做了一次转置。

知识体系结构

下面这张图展示了全局内存合并访问的核心逻辑：

最后说个实战技巧。如果你不确定自己的访存模式是否合并，可以用 NVIDIA Nsight Compute 的 Memory Workload Analysis 看一下。它会告诉你每个 warp 产生了多少次内存事务。理想情况是 1 次，如果看到 4 次或更多，那就得优化了。

嗯，全局内存合并访问就讲这么多。记住一句话：让同一个 warp 的线程访问连续地址，这是 GPU 编程的基本功。

想要更多资料，联系可可