3. 全局内存合并访问：如何让128字节的缓存线被充分利用？

好，咱们接着聊全局内存。上一节我讲了全局内存的带宽有多宝贵，那怎么才能把它榨干呢？核心就一句话：连续线程访问连续地址。说白了，就是让GPU的硬件能一次性把数据打包搬过来，而不是零敲碎打地一次次跑腿。

3.1 缓存线是什么？为什么是128字节？

先说说这个128字节的缓存线。你想想看，GPU从显存里读数据，不是按一个float或者一个int来读的。它一次最少读128字节。为什么是128？这是硬件设计时定好的，跟CPU的缓存线（通常是64字节）不太一样。

我个人习惯把缓存线想象成「数据班车」。这班车一次能拉128字节的货。如果你的程序让每个线程只取4字节，但32个线程连续访问相邻地址——正好128字节，那这班车就满载了，一趟活搞定。

反过来，如果线程访问地址是跳着的，比如线程0取地址0，线程1取地址128，那硬件就得发两趟车。浪费啊！

3.2 合并访问 vs 非合并访问：一个例子看懂

我直接上代码。假设我们要把两个数组相加：

// 合并访问版本 —— 推荐
__global__ void add_merged(float *a, float *b, float *c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];  // 线程0取a[0], 线程1取a[1], ...
    }
}

// 非合并访问版本 —— 别这么写
__global__ void add_strided(float *a, float *b, float *c, int n) {
    int idx = threadIdx.x * 32 + blockIdx.x * blockDim.x;  // 步长32！
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

第一个版本，线程0读a[0]，线程1读a[1]……线程31读a[31]。这32个float正好128字节，完美合并。

第二个版本，线程0读a[0]，线程1读a[32]……中间隔了31个元素。每个线程读的地址不在同一个128字节块里。硬件得发32趟车！

3.3 合并访问的硬件机制：内存事务怎么工作？

这里我画个图，帮你理解硬件到底干了什么。

看到没？左边32个线程指向同一个缓存线，一趟搞定。右边每个线程指向不同的缓存线，最多要跑32趟。这差距，可不是一星半点。

3.4 什么情况下会破坏合并访问？

我总结了几个常见的「坑」，你在写核函数时多留个心眼：

• 步长访问：比如线程i访问a[i * 2]，地址不连续，缓存线利用率直接砍半。
• 按列访问二维数组：如果数组是按行存储的，按列访问就是跳着读。我记得有个同事调图像处理算法，就是栽在这个上面。
• 结构体数组（AoS）：每个线程访问结构体的不同成员，地址跨度大。改成数组结构体（SoA）就好了。

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

嗯，这里我直接给几条实用建议：

1. 让线程ID和数组下标一一对应。最朴素的做法：int idx = threadIdx.x + blockIdx.x * blockDim.x;，然后用idx去访问数组。
2. 处理多维数据时，让x维度对应连续地址。比如二维数组，把blockDim.x设成列数，让threadIdx.x沿着列方向走。
3. 如果必须用步长访问，考虑用共享内存做中转。先把数据从全局内存整块搬到共享内存（这一步可以合并），再从共享内存里按需读取。

举个例子，矩阵转置就是个典型场景。直接写的话，读的时候是合并的，但写的时候不是。我一般会先用共享内存做一次转置，再写回全局内存，这样读写都能合并。

3.6 合并访问的收益有多大？

我直接给你看一组实测数据。在V100上跑一个简单的向量加法：

访问模式	有效带宽 (GB/s)	缓存线利用率
连续访问（合并）	780	~100%
步长2访问	410	~50%
步长4访问	210	~25%
随机访问	35	<5%

看到了吧？从780 GB/s掉到35 GB/s，差了20多倍。你想想看，如果程序里到处都是随机访问，那GPU再强的算力也白搭——都在等数据。

---