2、全局内存（Global Memory）：带宽瓶颈与合并访问（Coalescing）原理

好，咱们来聊聊全局内存。说实话，这是GPU编程里最容易踩坑的地方，也是性能提升空间最大的地方。我见过太多人，核函数写得飞快，结果一跑起来，慢得像蜗牛。为什么？十有八九是全局内存访问出了问题。

2.1 全局内存的“慢”是相对的

全局内存是GPU里最大的存储空间，所有线程都能访问。但它也是离计算单元最远的——你想想看，从芯片到显存，要走很长一段路。延迟大概在400-800个时钟周期。相比之下，寄存器几乎零延迟。

那怎么办？靠带宽来弥补。现代GPU的全局内存带宽很高，比如H100能达到3.35TB/s。但问题是，你得会用。用不好，实际带宽可能连理论值的10%都不到。

2.2 合并访问——GPU的“团购”机制

为什么会有合并访问这个概念？这得从GPU的内存控制器说起。

当32个线程（一个warp）同时发起内存访问时，硬件会把这些请求合并成尽可能少的“大交易”。比如，如果32个线程访问的是连续地址，那硬件可能只需要一次或两次内存事务就能搞定。如果地址是随机分散的，那就得32次甚至更多。

说白了，合并访问就是让GPU像“团购”一样，一次买一大块数据，而不是每个人单独跑一趟。我刚开始写CUDA时，没注意这个，结果一个简单的向量加法，跑出来的带宽只有理论值的5%。后来一查，发现地址没对齐，全乱套了。

2.3 什么才算“合并”的？

要满足合并访问，有几个条件：

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

举个例子，下面这个访问模式就是合并的：

// 合并访问示例
__global__ void coalesced_copy(float *in, float *out, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        out[idx] = in[idx];  // 每个线程访问连续地址
    }
}

而下面这个，就是典型的“反例”：

// 非合并访问——千万别这么写！
__global__ void strided_copy(float *in, float *out, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        out[idx] = in[idx * 32];  // 每个线程间隔32个元素
    }
}

第二种写法，32个线程访问的地址之间差了128字节，完全没法合并。我曾经在一个项目里看到有人这么写，结果带宽只有理论值的2%。改过来之后，性能直接提升了20倍。

2.4 合并访问的“粒度”

这里有个细节要注意。不同架构的合并粒度不一样：

架构	合并粒度	说明
Fermi/Kepler	128字节	一个warp的请求合并成128字节的事务
Maxwell/Pascal	32字节	粒度更细，但合并要求更灵活
Volta/Turing/Ampere	32字节	支持更复杂的合并模式

嗯，这里要注意：虽然新架构的合并粒度更细了，但并不意味着你可以随便写。最好的实践仍然是让warp内的线程访问连续地址。

2.5 常见的“坑”与避坑指南

我总结了几种常见的非合并访问模式，你写代码时可以对照检查：

• 步长访问：比如每个线程跳过N个元素。解决办法是调整数据布局，或者用共享内存做转置。
• 按列访问二维数组：在CUDA里，二维数组默认是按行存储的。按列访问会导致非合并。我曾经在一个矩阵转置项目里踩过这个坑，后来用共享内存做了分块转置，才把带宽跑满。
• 结构体数组（AoS） vs 数组结构体（SoA）：如果你用AoS，比如每个线程访问一个结构体的不同成员，那地址是分散的。改成SoA，就能实现合并访问。

2.6 实战：如何写出合并访问的代码

假设我们要做一个向量加法，数据是结构体数组：

// 不推荐：AoS布局
struct Vec3 {
    float x, y, z;
};
Vec3 *data;

// 每个线程访问同一个结构体的不同成员
float val = data[idx].x;  // 非合并！

改成SoA布局：

// 推荐：SoA布局
float *x_data, *y_data, *z_data;

// 每个线程访问连续地址
float val = x_data[idx];  // 合并！

如果数据布局已经固定，没法改怎么办？那就用共享内存做“中转”。先把数据从全局内存读到共享内存（这一步可能非合并），然后在共享内存里重新组织，再让线程访问。虽然多了一步拷贝，但整体性能往往更好。

2.7 总结

全局内存的合并访问，说白了就是让warp内的线程“步调一致”地访问连续地址。这是GPU编程的基本功，也是性能优化的第一道关卡。我个人习惯在写任何核函数之前，先画一下内存访问模式图，看看是不是合并的。

最后送你一句话：合并访问不是银弹，但不合并访问一定是毒药。 大多数情况下，只要注意数据布局和线程索引的对应关系，就能轻松实现合并访问。剩下的，就是靠共享内存和寄存器来进一步优化了——那是后面章节的内容。