第三章 内存访问优化（一）：Global Memory合并访问

各位同学，今天我们来聊聊CUDA优化里最基础、也最关键的一环——全局内存的合并访问。说实话，我见过太多新手写出来的kernel，计算逻辑没问题，但跑起来就是慢得离谱。一查性能分析器，好家伙，内存带宽利用率不到10%。

为什么会这样？说白了，就是没搞懂GPU怎么读内存。

3.1 合并访问（Coalesced Access）的本质

先问大家一个问题：当一个warp（32个线程）同时执行一条加载指令时，GPU是怎么把数据取回来的？

你可能会想，每个线程独立去读自己的地址，32个线程就是32次内存访问。错！GPU没那么傻。实际上，硬件会把这32个线程的请求合并成尽可能少的内存事务（memory transaction）。

我举个例子。假设每个线程要读一个float（4字节），32个线程就是128字节。如果这128字节在内存里是连续的，并且对齐到128字节边界，那么一次128字节的事务就能搞定。这就是最理想的合并访问。

反过来，如果线程0读地址A，线程1读地址A+100，线程2读地址A+200……那硬件就没办法合并了。32个线程可能需要32次独立的事务，带宽利用率直接掉到3%。

3.2 对齐与跨步访问的影响

聊完合并，咱们再深入一点——对齐和跨步。

对齐（Alignment），指的是访问的起始地址是不是某个边界值的整数倍。对于全局内存，我建议你始终保证128字节对齐。为什么是128？因为现代GPU的L2 cache line大小就是128字节。不对齐的话，一次事务可能跨两个cache line，白白浪费带宽。

我在项目中遇到过一件事：有个同事写了一个简单的向量加法，性能死活上不去。我一看代码，数据是用malloc分配的，地址没对齐。换成cudaMalloc后，性能直接翻倍。嗯，cudaMalloc默认返回256字节对齐的地址，这就是原因。

跨步访问（Strided Access），说白了就是线程访问的地址之间有间隔。比如每个线程读一个float，但步长是2（线程0读A[0]，线程1读A[2]，线程2读A[4]……）。

跨步访问有多伤？我们来算笔账：

步长	每个warp需要的128B事务数	带宽利用率
1（连续）	1	100%
2	2	50%
4	4	25%
32	32	3.1%

你看，步长一变大，性能就断崖式下跌。所以写代码时，一定要避免跨步访问。如果数据结构本身就有跨步，考虑用共享内存做一次转置，把跨步访问变成连续访问。

3.3 实战：矩阵转置优化

理论讲完了，咱们来点实战。矩阵转置是经典的合并访问反面教材——读的时候是连续的，写的时候就是跨步的。

先看一个最朴素的实现：

__global__ void transpose_naive(float *in, float *out, 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) {
        out[x * height + y] = in[y * width + x];
    }
}

这个kernel有什么问题？你想想看：读in的时候，y * width + x，x是连续变化的，所以读是合并的。但写out的时候，x * height + y，x是连续变化的，但y是固定的——这意味着相邻线程写的是不相邻的地址。写操作完全不合并。

怎么优化？经典做法是用共享内存做分块转置：

#define TILE_SIZE 32

__global__ void transpose_coalesced(float *in, float *out, int width, int height) {
    __shared__ float tile[TILE_SIZE][TILE_SIZE];
    
    int x = blockIdx.x * TILE_SIZE + threadIdx.x;
    int y = blockIdx.y * TILE_SIZE + threadIdx.y;
    
    if (x < width && y < height) {
        tile[threadIdx.y][threadIdx.x] = in[y * width + x];
    }
    
    __syncthreads();
    
    int x_out = blockIdx.y * TILE_SIZE + threadIdx.x;
    int y_out = blockIdx.x * TILE_SIZE + threadIdx.y;
    
    if (x_out < height && y_out < width) {
        out[y_out * height + x_out] = tile[threadIdx.x][threadIdx.y];
    }
}

这个版本的思路是：先把一块数据读到共享内存里（读是合并的），然后在共享内存里完成转置（共享内存没有合并访问问题），最后再写回全局内存（写也是合并的）。

我曾经在一个图像处理项目里用这个优化，把转置的吞吐从20 GB/s提升到了180 GB/s（在V100上）。嗯，将近9倍的提升。你想想看，就改了几行代码，效果立竿见影。

3.4 避坑指南

最后，分享几个我踩过的坑：

• 别用二维数组当参数传进kernel。二维数组在内存里不一定是连续的，用一维数组加索引计算最安全。
• 小心struct of arrays（SoA）和array of structs（AoS）的选择。如果每个线程要访问struct里的同一个字段，用SoA布局，保证合并访问。
• 别忘了检查地址对齐。用cudaMalloc没问题，但如果你自己管理内存池，一定要保证每个分配都是128字节对齐的。
• 跨步访问不一定都能避免。如果实在避免不了，考虑用cudaMemcpy2D或者texture memory来缓解。

好了，这一章的内容就到这里。记住一句话：合并访问是GPU性能的基石，不合并的代码不配叫CUDA优化。