CUDA全局内存合并访问原理与优化实践

1. 全局内存合并访问的核心原理

在CUDA编程中，全局内存合并访问是提升GPU性能的关键技术。理解其工作原理需要从硬件层面入手：现代GPU的显存控制器以32字节为基本单位（称为"扇区"）进行数据传输。当一个线程束（32个线程）发出内存请求时，硬件会尝试将这些请求合并为最少次数的32字节事务。

合并访问的理想情况是：线程束中连续的线程访问连续的32位数据（如float类型）。例如，线程0访问地址0x0000，线程1访问0x0004，依此类推。这种情况下，硬件可以将整个线程束的访问合并为单个128字节（32线程×4字节）的事务，正好对应4个32字节扇区。

关键点：合并不是指线程访问相同地址，而是指访问连续的地址空间。完全相同的地址反而会导致性能下降（bank conflict）。

2. 合并访问的判断方法

2.1 代码层面的判断

通过分析内核代码的内存访问模式可以初步判断是否满足合并条件。以下是一个典型示例：

cpp复制__global__ void checkCoalescing(float* data, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        // 合并访问：连续线程访问连续地址
        float value = data[tid];  // tid递增 ⇒ 地址连续
        
        // 非合并访问示例（注释掉）：
        // float value = data[tid * 32 % n]; // 跨步访问
    }
}

判断标准：

1. 每个线程访问的数据大小是否为4字节（32位）
2. 线程ID与访问地址是否呈线性关系
3. 访问的起始地址是否32字节对齐

2.2 使用Nsight Compute工具验证

代码分析只能提供初步判断，实际需要使用性能分析工具验证。NVIDIA Nsight Compute提供了详细的内存访问统计：

bash复制ncu --metrics l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,l1tex__t_requests_pipe_lsu