全局内存高效访问：合并访问的艺术

全局内存，说白了就是GPU里那块最大的存储空间。所有线程都能访问，但速度嘛...嗯，得看你怎么用。我刚开始做CUDA开发时，总觉得全局内存就是个大仓库，随便怎么搬都行。结果一跑性能分析，发现带宽利用率低得可怜。

为什么会这样？因为全局内存有个脾气——它喜欢「合并访问」。

合并访问的核心原理

先说说什么是合并访问。当warp里的32个线程同时访问全局内存时，硬件会把它们的请求合并成尽可能少的内存事务。如果访问模式是连续的、对齐的，那一个事务就能搞定。否则，就得拆成多个事务，带宽就浪费了。

我个人习惯用一个比喻来理解：想象你去图书馆借书，32个人要借同一排书架上的书。如果大家按顺序借第1、2、3...32本，管理员一趟就能搬完。如果大家东借一本西借一本，管理员就得来回跑好几趟。

对齐要求：别小看这个细节

对齐要求其实很简单：每个内存事务的起始地址必须是事务大小的整数倍。比如128字节的事务，起始地址必须是128的倍数。

我在项目中遇到过一个问题：明明数据是连续的，但带宽就是上不去。后来一查，发现结构体没对齐。每个线程访问的是结构体里的某个成员，但结构体大小不是2的幂次，导致地址错位。

跨步访问优化：一个常见的陷阱

跨步访问，说白了就是线程访问的地址之间有间隔。比如线程0访问地址0，线程1访问地址N，线程2访问地址2N...这种模式在矩阵转置、图像处理中很常见。

你想想看，如果N=32，那每个线程访问的地址正好落在不同的内存段里，合并访问就完全失效了。每个线程都得单独发一个事务，带宽利用率直接掉到1/32。

怎么优化？我建议用共享内存做中转。先把数据按合并访问的方式读到共享内存里，再按跨步访问的方式读出来。虽然多了个步骤，但速度反而更快。

// 跨步访问的优化示例
__global__ void transpose_kernel(float *in, float *out, int width) {
    __shared__ float tile[TILE_SIZE][TILE_SIZE];
    
    int x = blockIdx.x * TILE_SIZE + threadIdx.x;
    int y = blockIdx.y * TILE_SIZE + threadIdx.y;
    
    // 合并写入共享内存
    tile[threadIdx.y][threadIdx.x] = in[y * width + x];
    __syncthreads();
    
    // 合并读出到输出
    out[x * width + y] = tile[threadIdx.x][threadIdx.y];
}

cudaMallocPitch与cudaMemcpy2D：帮你自动对齐

说实话，手动对齐二维数组挺麻烦的。好在CUDA提供了两个好用的API：cudaMallocPitch和cudaMemcpy2D。

cudaMallocPitch分配二维数组时，会自动调整每行的步长（pitch），确保每行起始地址对齐。这样你就不用操心对齐问题了。

// 使用cudaMallocPitch分配二维数组
size_t pitch;
float *d_data;
cudaMallocPitch(&d_data, &pitch, width * sizeof(float), height);

// 使用cudaMemcpy2D拷贝数据
cudaMemcpy2D(d_data, pitch, h_data, width * sizeof(float), 
             width * sizeof(float), height, cudaMemcpyHostToDevice);

// 内核中访问时，使用pitch作为步长
__global__ void kernel(float *data, size_t pitch, int width, int height) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < width && y < height) {
        float *row = (float *)((char *)data + y * pitch);
        row[x] = some_value;
    }
}

实战中的合并访问检查

怎么知道你的代码有没有实现合并访问？我推荐用NVIDIA Nsight Compute或者Visual Profiler。它们会直接告诉你每个内核的内存事务数。

如果发现事务数远大于理论最小值，那就说明合并访问没做好。这时候，我一般会检查以下几点：

• 线程索引和地址的映射关系是否连续
• 数据类型大小是否一致（别混用float和double）
• 结构体是否对齐
• 二维数组的pitch是否正确使用

知识体系总览

下面这张图总结了全局内存高效访问的核心要点，我把它画成了流程图，方便你理解整个优化思路。

嗯，全局内存访问这块，说白了就是「对齐」和「连续」四个字。只要记住这两个原则，大部分性能问题都能迎刃而解。

想要更多资料，联系可可