第3章内存访问模式

3. 内存访问模式：合并访问与非合并访问、对齐访问与未对齐访问、空间局部性

好，咱们今天聊点实在的。显存带宽这东西，说白了就是GPU的命脉。你算力再强，数据喂不进去，那也是白搭。我这些年调优下来，发现很多性能问题，根源就出在内存访问模式上。

你想想看，GPU里有成百上千个核心，它们同时去抢显存带宽。如果访问模式乱七八糟，那带宽利用率能高才怪。今天我就把这块掰开了揉碎了讲清楚。

3.1 合并访问 vs 非合并访问

这是CUDA编程里最基础、也最重要的概念。我刚开始学的时候，觉得不就是读数据嘛，有啥区别？后来被现实狠狠教育了一回。

什么是合并访问？

当一个warp（32个线程）同时访问显存时，硬件会把它们的请求合并成尽可能少的内存事务。如果这32个线程访问的地址是连续的，那就能合并成一次或几次大的内存事务。这就是合并访问。

非合并访问呢？

就是线程访问的地址乱七八糟，东一个西一个。硬件没法合并，只能拆成很多小事务。带宽利用率直线下降。

核心要点：合并访问能让显存带宽利用率接近100%，非合并访问可能连10%都不到。

我举个例子。假设每个线程要读一个float（4字节），一个warp有32个线程，总共需要128字节。如果是合并访问，硬件一次就能搞定。如果是非合并访问，可能得拆成32次，每次4字节。这差距，你品，你细品。

// 合并访问示例
__global__ void coalesced_kernel(float* data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        float val = data[idx];  // 连续地址，完美合并
        // 处理数据...
    }
}

// 非合并访问示例
__global__ void non_coalesced_kernel(float* data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        float val = data[idx * 32];  // 间隔32个元素，无法合并
        // 处理数据...
    }
}

我的经验：我曾经优化过一个矩阵转置的kernel。原始版本是非合并访问，带宽利用率只有15%。改成合并访问后，直接飙到85%。就改了一行代码，性能提升了5倍多。这种优化，性价比极高。

3.2 对齐访问 vs 未对齐访问

对齐访问，指的是访问的起始地址是某个对齐边界（通常是32字节或128字节）的整数倍。未对齐访问就是地址不在边界上。

为什么会这样？因为GPU的内存控制器是按段（segment）来读取数据的。每个段通常是32字节或128字节。如果地址对齐，一个段就能覆盖你需要的所有数据。如果不对齐，可能就需要跨两个段去读。

注意：未对齐访问不一定会报错，但会额外消耗带宽。每个跨段访问，都会多读一部分你用不到的数据。

我记得有一次调试一个性能问题，发现kernel跑得特别慢。查了半天，发现是结构体成员没对齐。每个线程读一个结构体，结构体大小是36字节，不是32的倍数。结果每次访问都要跨段，带宽浪费了将近一半。

解决办法？重新排列结构体成员，或者用__align__关键字强制对齐。

// 未对齐的结构体
struct Misaligned {
    float a;  // 4字节
    float b;  // 4字节
    float c;  // 4字节
    float d;  // 4字节
    float e;  // 4字节
    // 总共20字节，不是32的倍数
};

// 对齐的结构体
struct __align__(32) Aligned {
    float a;  // 4字节
    float b;  // 4字节
    float c;  // 4字节
    float d;  // 4字节
    float e;  // 4字节
    float pad[3];  // 填充12字节，凑够32字节
};

3.3 空间局部性

空间局部性，说白了就是：你访问了一个地址，接下来很可能访问它附近的地址。这个原理在CPU缓存里很常见，在GPU里同样重要。

GPU的L1缓存和共享内存，都是利用空间局部性来提升性能的。如果一个warp里的线程访问的数据在空间上很接近，那缓存命中率就会很高。反之，如果数据分散在显存各处，缓存就形同虚设。

我给你们画个图，看看这几种访问模式的区别。

嗯，这张图应该很直观了。合并访问和对齐访问，都是让硬件能高效工作的关键。非合并和未对齐，就是性能杀手。

3.4 实战建议

说了这么多理论，来点实际的。我总结了几条经验，你们写kernel的时候可以对照着检查。

让线程访问连续地址——这是合并访问的基础。用threadIdx.x作为最内层索引。
注意数据类型的对齐——结构体尽量用__align__对齐到32字节或128字节。
利用共享内存做转置——如果必须非合并访问，先读到共享内存里，再按合并方式读出来。
避免跨段访问——每个线程读的数据量最好是32字节的整数倍。
用profiler检查——NVIDIA Nsight Compute里有个指标叫"Global Load Efficiency"，低于80%就要警惕了。

避坑指南：我曾经在一个项目里，用二维数组存矩阵。按行访问没问题，按列访问就炸了。后来改成用共享内存做转置，性能直接翻倍。记住：行优先存储的数组，按行访问才是合并的。

最后说一句，这些优化技巧，说白了就是让硬件干得舒服。你让硬件舒服了，它才能让你舒服。别跟硬件对着干，顺着它的脾气来，性能自然就上去了。

一句话总结：合并访问 + 对齐访问 + 良好的空间局部性 = 显存带宽利用率最大化。

公众号:蓝海资料掘金营，微信deep3321

📚 课程目录

3. 内存访问模式：合并访问与非合并访问、对齐访问与未对齐访问、空间局部性

3.1 合并访问 vs 非合并访问

3.2 对齐访问 vs 未对齐访问

3.3 空间局部性

3.4 实战建议