4. 全局内存合并访问

说到全局内存访问，我得先跟你聊聊一个常见的误区。很多新手写CUDA程序，觉得只要把数据拷贝到显存里就完事了。嗯，其实没那么简单。全局内存的访问方式，直接决定了你的核函数能跑多快。

我刚开始做GPU优化那会儿，有个项目怎么调都跑不满带宽。后来才发现，问题出在内存访问模式上。改了一行代码，性能直接翻了三倍。今天我们就来彻底搞懂这个「合并访问」。

4.1 访问粒度：32字节 vs 128字节

先说说硬件层面的东西。GPU的全局内存，不是按字节一个一个读的。它有自己的脾气——一次最少读32字节，最多读128字节。

为什么会这样？说白了，这是为了平衡带宽利用率和硬件复杂度。32字节对应一个warp里32个线程，每个线程读4字节（比如一个float）。128字节则是缓存行（cache line）的大小。

粒度	对应场景	典型用途
32字节	每个线程读4字节，32线程对齐	float数组、int数组
128字节	缓存行传输，空间局部性好	结构体、纹理内存

我个人习惯把32字节叫做「warp粒度」，128字节叫「缓存粒度」。记住一个原则：尽量让一个warp的32个线程，访问连续的内存地址。这样硬件就能用最少的次数，搬最多的数据。

4.2 合并访问的条件与判断

什么叫合并访问？我给你打个比方。想象你去图书馆借书，一次最多借32本。如果你和你的31个朋友，每人要借的书都在同一个书架上挨着放，管理员一趟就能搬完。这就是合并访问。

反过来，如果你们32个人要借的书分散在图书馆各个角落，管理员就得跑32趟。这就是非合并访问，性能惨不忍睹。

判断是否合并访问，其实就三个条件：

• 对齐：起始地址必须是32字节（或128字节）的倍数
• 连续：线程i访问地址A，线程i+1访问地址A+sizeof(type)
• 顺序：线程ID和地址顺序一致（不是必须，但推荐）

我曾经在一个图像处理项目里，因为没注意对齐，带宽利用率只有20%。后来加了个padding，直接干到90%以上。嗯，这就是对齐的重要性。

4.3 结构体数组（AoS）vs 数组结构体（SoA）

这个坑我踩过，而且踩得很深。先看两种数据布局：

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

// SoA：数组结构体
struct Particles {
    float *x, *y, *z;
    float *vx, *vy, *vz;
};
Particles ps;

你想想看，如果每个线程处理一个粒子，要读它的x坐标。在AoS布局下，一个warp的32个线程读到的地址是：

particles[0].x, particles[1].x, particles[2].x ...

这些地址之间隔了6个float（24字节）！完全不是连续的。硬件每次读32字节，里面只有4个字节有用，浪费率高达87.5%。

换成SoA呢？所有x坐标连续存储。32个线程读的就是连续的32个float，一次128字节的传输，全部用上。带宽利用率接近100%。

4.4 使用cudaMallocPitch与cudaMemcpy2D

处理二维数据时，有个很容易忽略的问题：行对齐。比如一张1024x768的图像，每行1024个像素。但硬件要求每行的起始地址对齐到128字节。

如果图像宽度不是128字节的倍数，怎么办？手动padding？太麻烦了。CUDA提供了现成的API：

size_t pitch;
float* d_data;
cudaMallocPitch(&d_data, &pitch, width * sizeof(float), height);

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

cudaMallocPitch会自动计算合适的pitch（每行实际占用的字节数），保证每行对齐。你只需要用pitch来算地址就行：

// 访问第row行第col列
float val = d_data[row * pitch / sizeof(float) + col];

4.5 实战：图像滤波的合并访问优化

光说不练假把式。我们拿一个3x3均值滤波来实战。先看原始版本：

__global__ void blur_naive(float* in, float* out, int w, int h) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    
    if (col < w && row < h) {
        float sum = 0.0f;
        for (int dy = -1; dy <= 1; dy++) {
            for (int dx = -1; dx <= 1; dx++) {
                int r = min(max(row + dy, 0), h-1);
                int c = min(max(col + dx, 0), w-1);
                sum += in[r * w + c];
            }
        }
        out[row * w + col] = sum / 9.0f;
    }
}

这个版本有什么问题？每个线程读9个像素，而且相邻线程读的像素有大量重叠。虽然访问模式是连续的，但重复读取浪费了带宽。

优化思路：利用共享内存。每个block加载一块数据到共享内存，然后所有线程从共享内存读。这样全局内存只访问一次：

__global__ void blur_shared(float* in, float* out, int w, int h, int pitch) {
    __shared__ float tile[TILE_SIZE + 2][TILE_SIZE + 2]; // 加2是边界
    
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    
    // 加载数据到共享内存（合并访问！）
    int in_idx = row * pitch + col;
    tile[threadIdx.y + 1][threadIdx.x + 1] = in[in_idx];
    
    // 处理边界...
    __syncthreads();
    
    // 从共享内存计算（无全局内存访问）
    float sum = 0.0f;
    for (int dy = 0; dy < 3; dy++) {
        for (int dx = 0; dx < 3; dx++) {
            sum += tile[threadIdx.y + dy][threadIdx.x + dx];
        }
    }
    out[row * w + col] = sum / 9.0f;
}

这个版本里，全局内存的访问是完全合并的——连续线程读连续地址。共享内存的访问虽然不合并（因为bank conflict），但共享内存的延迟低得多，影响不大。

最后送你一个检查合并访问的小技巧：在Nsight Compute里看Memory Workload Analysis，如果L2命中率低、全局内存请求次数多，基本就是非合并访问。我曾经靠这个工具，半小时定位了一个性能瓶颈。

这张图把合并访问的知识点串起来了。从核心概念出发，延伸到访问粒度、判断条件、数据布局，再到具体的优化技术。你写代码的时候，可以对照这张图检查自己的实现。

想要更多资料，联系可可