三、全局内存访问模式：合并访问与非合并访问、对齐要求、内存事务优化

好，咱们来聊聊全局内存。这是GPU里最基础、也是最容易踩坑的地方。我见过太多同学，核函数写得飞快，结果一跑性能惨不忍睹。问题出在哪？十有八九是全局内存访问没搞好。

说白了，GPU的全局内存带宽很大，但延迟也很高。如果你访问方式不对，带宽利用率可能连10%都不到。这就像你开着一辆法拉利，却在早晚高峰的市区里堵着——空有马力，跑不起来。

3.1 合并访问：GPU性能的命门

什么是合并访问？简单说，就是一个warp里的32个线程，访问全局内存时，如果能合并成少数几个内存事务，那就是合并访问。反之，如果每个线程各访各的，那就是非合并访问。

为什么会这样？因为GPU的内存控制器是按128字节或32字节的粒度来发起事务的。你想想看，如果32个线程访问的地址是连续的，那一个128字节的事务就能搞定。但如果地址是分散的，那就得发起多个事务，甚至每个线程一个事务。

我在项目中遇到过这样一个案例：一个矩阵转置的核函数，最初写的时候没注意访问模式，结果带宽利用率只有20GB/s左右。后来改成合并访问，直接飙到140GB/s。差距就是这么大。

3.1.1 合并访问的条件

要达成合并访问，需要满足几个条件：

• 对齐：访问的起始地址必须是32字节或128字节的倍数
• 连续性：线程0访问地址A，线程1访问A+4，线程2访问A+8...以此类推
• 数据类型一致：所有线程访问的数据类型和大小相同

嗯，这里要注意，对齐和连续性是两个不同的概念。对齐说的是地址本身，连续性说的是线程间的地址关系。两者缺一不可。

3.2 非合并访问的典型场景

非合并访问最常见的情况就是按列访问。比如你有一个二维数组，按行存储，但你的线程是按列来读的。那每个线程访问的地址之间就隔了一行的距离，根本连不起来。

// 非合并访问：按列读取
__global__ void col_read(float *data, int width, int height) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    float val = data[col * width + row];  // 注意这里：col * width + row
    // ...
}

// 合并访问：按行读取
__global__ void row_read(float *data, int width, int height) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    float val = data[row * width + col];  // row * width + col
    // ...
}

你看，只是把索引顺序换了一下，性能天差地别。我曾经帮一个团队优化过类似的代码，他们原本的核函数跑一次要12毫秒，改成合并访问后直接降到1.5毫秒。8倍的提升，就改了一行代码。

3.3 对齐要求

对齐要求说白了就是：你访问的起始地址，最好是32字节或128字节的倍数。为什么？因为GPU的内存事务就是按这个粒度来的。

举个例子，如果你要读一个float4（16字节），起始地址是64的倍数，那一个事务就能搞定。但如果起始地址是60，那就跨了两个128字节的边界，需要两个事务。

数据类型	推荐对齐	说明
char/short	32字节	一个warp访问32字节
int/float	128字节	一个warp访问128字节
double/float2	128字节	一个warp访问256字节
float4/double2	256字节	一个warp访问512字节

我个人习惯在分配全局内存时，用cudaMalloc就够了，它会自动对齐到256字节。但如果你自己管理内存池，那就得小心对齐问题了。

3.4 内存事务优化

内存事务，就是GPU从全局内存读取数据的最小单元。每个事务的大小可以是32字节、128字节，或者在某些架构上是64字节。优化的目标就是：用最少的事务，传输最多的有用数据。

3.4.1 事务浪费的典型情况

你想想看，如果一个warp的32个线程，每个线程只读4字节，总共需要128字节。如果地址是连续的且对齐的，那一个128字节的事务就够了。但如果地址是分散的，比如每个线程读一个int，但地址间隔很大，那可能就需要多个事务。

更糟糕的是，即使地址是连续的，但如果起始地址没对齐，也可能导致事务浪费。比如起始地址是60，那第一个事务读0-127字节，第二个事务读128-255字节。但实际上你只需要60-188字节的数据，第一个事务浪费了60字节，第二个事务浪费了68字节。利用率只有50%。

3.4.2 利用向量化数据类型

一个很实用的技巧是使用向量化数据类型，比如float4、int4、double2。这样每个线程一次读16字节，32个线程就是512字节。虽然事务大小变大了，但事务数量减少了，整体效率反而更高。

// 使用float4进行合并访问
__global__ void vec_add(float4 *a, float4 *b, float4 *c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx].x = a[idx].x + b[idx].x;
        c[idx].y = a[idx].y + b[idx].y;
        c[idx].z = a[idx].z + b[idx].z;
        c[idx].w = a[idx].w + b[idx].w;
    }
}

我在做图像处理时特别喜欢用uchar4，因为一个像素正好是RGBA四个通道，一个线程处理一个像素，完美对齐。

3.4.3 共享内存做中转

如果实在无法避免非合并访问，比如矩阵转置这种操作，那可以用共享内存做中转。先把数据从全局内存按合并方式读到共享内存，再从共享内存按任意方式读取。因为共享内存的延迟低，而且支持广播和随机访问。

// 矩阵转置：用共享内存做中转
__global__ void transpose(float *in, float *out, int width, int height) {
    __shared__ float tile[TILE_DIM][TILE_DIM];
    
    int x = blockIdx.x * TILE_DIM + threadIdx.x;
    int y = blockIdx.y * TILE_DIM + threadIdx.y;
    
    // 合并写入共享内存
    if (x < width && y < height) {
        tile[threadIdx.y][threadIdx.x] = in[y * width + x];
    }
    
    __syncthreads();
    
    // 合并写回全局内存（转置后）
    x = blockIdx.y * TILE_DIM + threadIdx.x;
    y = blockIdx.x * TILE_DIM + threadIdx.y;
    
    if (x < height && y < width) {
        out[y * height + x] = tile[threadIdx.x][threadIdx.y];
    }
}

这个模式在矩阵运算中非常常见。我记得有一次优化一个卷积层，就是用共享内存做数据重排，把非合并访问变成了合并访问，性能提升了3倍多。

3.5 性能分析工具的使用

光靠猜是不行的。我建议你用nvidia-smi或者Nsight Compute来实际看看内存事务的情况。Nsight Compute里有个Memory Workload Analysis的section，会告诉你每个warp平均发了多少个事务，以及事务的利用率。

如果看到Sector Utilization低于50%，那基本可以断定你的访问模式有问题。这时候就该回头检查一下地址对齐和连续性了。

好了，这一章的内容就到这里。下一章我们会聊聊共享内存的bank conflict问题，那又是一个容易踩坑的地方。