CUDA合并访存优化：原理与实践指南

1. CUDA合并访存的核心概念

第一次接触CUDA合并访存这个概念时，我正为一个图像处理内核的性能瓶颈发愁。当时的内核运行速度比预期慢了近3倍，通过Nsight工具分析才发现问题出在全局内存访问模式上。合并访存（Coalesced Memory Access）是CUDA编程中提升内存带宽利用率的关键技术，它允许一个线程束（warp）中的32个线程对全局内存的访问合并为少数几个内存事务。

在GPU架构中，全局内存访问是以内存事务（memory transaction）为单位进行的。每个事务可以一次性读取32字节、64字节或128字节的数据。当线程束中的线程访问连续且对齐的内存地址时，这些访问可以被硬件自动合并为一个或几个内存事务。反之，如果线程访问的内存地址分散或不连续，就可能需要发起多个内存事务，造成带宽浪费。

关键理解：合并访存不是某种特殊的API调用，而是通过合理安排线程内存访问模式来触发硬件的自动优化机制。

2. 合并访存的工作原理与硬件基础

2.1 GPU内存子系统架构

现代NVIDIA GPU的内存子系统由多个内存控制器组成，每个控制器管理一块内存分区。以Ampere架构的A100为例，它有8个内存控制器，每个控制器对应512位（64字节）的访问宽度。这意味着最理想的情况下，每个内存事务应该正好是64字节的连续数据。

当线程束发出内存请求时，内存访问单元会检查这些请求的目标地址。如果地址落在同一个对齐的64字节段内（即地址的低6位相同），这些请求就可以被合并。合并的粒度取决于计算能力版本：

• 计算能力3.x/5.x：默认128字节事务
• 计算能力6.x及以上：32/64/128字节可选
• Volta及以后架构：还支持更灵活的访问模式

2.2 合并访存的条件

要实现完美的合并访存，需要满足以下条件：

1. 地址连续性：线程束中的线程按线程ID顺序访问连续的内存地址。即threadIdx.x为n的线程访问地址base_addr + n*sizeof(type)
2. 对齐要求：首地址必须对齐到事务大小的整数倍（32/64/128字节）
3. 访问粒度：所有线程访问的数据类型大小相同，建议使用4字节或8字节类型

违反这些条件会导致事务数量增加。最坏情况下，可能需要32个独立事务（每个线程一个），带宽利用率仅为1/32。

3. 合并访存示例代码解析

3.1 基础示例：向量加法

让我们从一个简单的向量加法内核开始，展示合并与非合并访问的差异：

c复制// 合并访问版本
__global__ void vectorAdd_coalesced(float* A, float* B, float* C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        C[i] = A[i] + B[i];  // 合并访问：线程连续访问连续地址
    }
}

// 非合并访问版本
__global__ void vectorAdd_noncoalesced(float* A, float* B, float* C,