GPU全局内存合并访问优化指南

1. 全局内存合并访问的核心概念

在GPU编程中，全局内存合并访问（Coalesced Memory Access）是一个直接影响程序性能的关键技术点。简单来说，它描述的是当多个线程同时访问全局内存时，这些访问请求能否被合并成更少的内存事务（memory transaction）来处理。

想象一下这样的场景：你带着32个朋友去餐厅吃饭，如果每个人都单独点餐，服务员需要来回跑32次；但如果大家能协调好，把相似的菜品合并下单，可能只需要3-4次就能完成点餐。全局内存合并访问的原理与此类似——当相邻线程访问连续的内存地址时，GPU可以将这些访问合并为一次内存事务，大幅减少内存带宽的浪费。

提示：在CUDA架构中，合并访问的最小单位是32字节（对应一个warp的32个线程），如果线程访问的内存地址落在连续的32字节对齐块内，就能实现完美合并。

2. 合并访问的硬件实现原理

2.1 GPU内存子系统的工作机制

现代GPU的内存子系统采用分层的设计架构。以NVIDIA的GPU为例，全局内存访问需要经过以下关键路径：

1. 内存控制器：负责与显存物理交互
2. L2缓存：所有SM共享的二级缓存
3. L1缓存/共享内存：每个SM独有的高速缓存
4. 寄存器文件：线程私有的最快存储

当warp中的线程发出内存请求时，内存控制器会根据访问模式决定是否合并。合并访问的理想情况是：warp中所有线程访问连续的32/64/128字节对齐的内存块（具体取决于架构）。

2.2 合并访问的判定条件

一个内存访问是否被合并，主要看三个维度：

1. 地址连续性：线程访问的地址是否连续递增
2. 对齐要求：起始地址是否符合硬件要求的对齐边界
3. 访问粒度：是否匹配内存控制器的传输单元大小

以Volta架构为例，其内存事务支持32、64和128字节的传输粒度。如果warp中所有线程访问的地址落在同一个128字节的段内，且地址是连续递增的，则只需要1次128字节的事务即可完成。

3. 合并访问的典型模式分析

3.1 理想合并访问模式

c复制// 示例1：完美合并访问
__global__ void vectorAdd(float* A, float* B, float* C) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    C[i] = A[i] + B[i];  // 所有线程访问连续地址
}

在这个经典向量加法示例中：

• 每个线程访问的A[i]、B[i]、C[i]地址都是连续且对齐的
• 一个warp的32个线程访问32个连续的float（共128字节）
• 只需1次128字节的内存事务即可满足整个warp的需求

3.2 非合并访问的典型场景

c复制// 示例2：跨步访问导致非合并
__global__ void strideAccess(float* input, float* output, int stride) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    output[tid] = input[tid * stride];  // stride>1时破坏连续性
}

当stride>1时会出现：

• 线程0访问input[0]
• 线程1访问input[stride]
• 线程2访问input[2*stride