合并访问原则：全局内存合并访问的条件与实现方法

说到GPU显存访问，我第一个想聊的就是「合并访问」。这玩意儿说白了，就是让GPU一次内存请求能拉回尽可能多的有用数据。你想想看，显存带宽再大，如果每次访问都浪费掉，那性能肯定上不去。

我在项目中遇到过不少同学，核函数写得挺漂亮，结果一跑profiling，发现内存带宽利用率只有10%出头。嗯，十有八九是合并访问没做好。

什么是合并访问？

先看一个简单的概念。GPU的全局内存是以事务（transaction）为单位访问的。一个事务通常传输32字节或128字节的数据。当线程束（warp）里的32个线程访问全局内存时，硬件会尽量把这些访问合并成尽可能少的事务。

如果32个线程访问的地址是连续的、对齐的，那一个事务就能搞定。如果地址是散乱的，那就可能需要多个事务，甚至每个线程一个事务。这就是合并访问的核心思想。

合并访问的条件

要达成合并访问，需要满足几个条件。我按重要程度排个序：

1. 地址连续：线程束内32个线程访问的地址必须是连续的。比如线程0访问地址A，线程1访问A+4，线程2访问A+8……以此类推。
2. 地址对齐：起始地址必须是32字节或128字节的整数倍。说白了，就是地址要对齐到事务边界上。
3. 数据类型一致：所有线程访问的数据类型和大小要相同。混用int和float会导致无法合并。

为什么会这样？因为GPU的内存控制器是按固定大小的事务来工作的。你地址不对齐，一个事务可能只用到一半数据，另一半就浪费了。

代码示例：合并 vs 非合并

咱们直接看代码。假设有一个float数组，每个线程要读一个元素：

// ✅ 合并访问：连续地址
__global__ void coalesced_read(float *in, float *out, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        out[idx] = in[idx];  // 线程0读in[0], 线程1读in[1], ...
    }
}

// ❌ 非合并访问：跨步访问
__global__ void strided_read(float *in, float *out, int N, int stride) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        out[idx] = in[idx * stride];  // 线程0读in[0], 线程1读in[stride], ...
    }
}

第一个核函数，线程束内的线程访问in[0], in[1], in[2]... 地址连续，完美合并。第二个核函数，如果stride=2，线程0读in[0]，线程1读in[2]，中间跳过了in[1]。这就没法合并了。

矩阵转置的合并访问陷阱

矩阵转置是个经典例子。咱们看看两种写法：

// 按行读，按列写
__global__ void transpose_naive(float *in, float *out, int width) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < width && col < width) {
        out[col * width + row] = in[row * width + col];
    }
}

// 使用共享内存优化
__global__ void transpose_coalesced(float *in, float *out, int width) {
    __shared__ float tile[TILE_SIZE][TILE_SIZE];
    
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    
    if (row < width && col < width) {
        tile[threadIdx.y][threadIdx.x] = in[row * width + col];
    }
    
    __syncthreads();
    
    int new_row = blockIdx.x * TILE_SIZE + threadIdx.y;
    int new_col = blockIdx.y * TILE_SIZE + threadIdx.x;
    
    if (new_row < width && new_col < width) {
        out[new_row * width + new_col] = tile[threadIdx.x][threadIdx.y];
    }
}

第一个版本，读的时候是连续的（按行读），但写的时候是跨步的（按列写）。写操作没法合并。第二个版本用共享内存做中转，读和写都是连续的。性能差距能到10倍以上。

我曾经在一个图像处理项目里踩过这个坑。转置一张4K图像，naive版本跑了50多毫秒，优化后只要5毫秒。嗯，从那以后我再也不敢小看合并访问了。

结构体数组 vs 数组结构体

这也是个常见问题。看两种数据布局：

// AoS: 结构体数组
struct Particle {
    float x, y, z;
};
Particle particles[N];

// SoA: 数组结构体
struct Particles {
    float x[N], y[N], z[N];
};

如果每个线程处理一个粒子，只访问x坐标：

• AoS布局：线程0读particles[0].x，线程1读particles[1].x。地址间隔是12字节（三个float），没法合并。
• SoA布局：线程0读x[0]，线程1读x[1]。地址连续，完美合并。

合并访问的检查方法

我一般用两种方法检查合并访问情况：

1. NVIDIA Nsight Compute：看Memory Workload Analysis里的Global Load/Store Efficiency。如果效率低于80%，说明有优化空间。
2. 手动计算：算一下每个线程束实际访问的地址范围，除以事务大小，看需要多少个事务。

举个例子，如果每个线程读4字节，32个线程总共需要128字节。如果地址连续且对齐，一个128字节事务就够了，效率100%。如果地址跨步为2，那需要两个128字节事务，效率50%。

知识体系图

下面这张图总结了合并访问的核心逻辑：

总结一下

合并访问是GPU优化的基本功。说白了就一句话：让相邻线程访问相邻地址。做到这一点，你的核函数就已经赢在起跑线上了。

我个人的习惯是，每次写完核函数，先跑一下Nsight Compute看看内存效率。如果低于80%，就回头检查数据布局和访问模式。这个习惯帮我省了不少调试时间。

最后提醒一句：合并访问虽然重要，但也不是银弹。有时候为了算法简洁，牺牲一点合并访问也是可以接受的。关键是要知道自己在做什么，以及代价是什么。