4、内存带宽与延迟：为什么全局内存这么慢？合并访问（Coalesced Access）的原理

4.1 全局内存——GPU的“慢速大仓库”

说实话，我刚接触CUDA时，最困惑的就是：明明GPU有几千个核心，为什么程序跑不快？

后来我才明白，瓶颈往往不在计算，而在内存。全局内存，就是那个拖后腿的“慢速大仓库”。

为什么慢？两个核心原因：带宽有限和延迟巨大。

• 延迟：全局内存访问一次，大约需要400~800个时钟周期。相比之下，寄存器访问只要1个周期。差了三个数量级。
• 带宽：虽然理论带宽很高（比如A100有2TB/s），但实际利用率往往不到30%。

你想想看，如果每个线程都乱访问内存，那带宽就被白白浪费了。我见过不少项目，计算部分优化得再好，内存访问一塌糊涂，整体性能就是上不去。

4.2 合并访问——让内存“批量发货”

全局内存的硬件设计，其实很像一个“快递分拣中心”。

每次内存请求，不是只取一个字节，而是取一个内存事务（Memory Transaction）。比如128字节或32字节这样一大块。

如果同一个warp（32个线程）访问的地址是连续的，并且对齐到128字节边界——那硬件只需要一次事务，就能满足所有线程。这就是合并访问。

反过来，如果每个线程访问的地址东一个西一个，那硬件就得发起多次事务。带宽利用率直线下降。

4.3 合并访问的三种模式

我习惯把合并访问分成三种情况，方便理解：

模式	访问方式	带宽利用率	典型场景
完美合并	连续地址，对齐到128B	~100%	数组顺序访问
部分合并	连续地址，但未对齐	50%~80%	结构体数组（AoS）
无合并	随机地址或步长过大	<20%	链表、哈希表

说白了，你要让warp里的32个线程，访问的内存地址尽量“排好队”。

4.4 代码示例：合并 vs 非合并

来看一个最经典的例子——向量加法。

// ❌ 非合并访问（步长访问）
__global__ void bad_access(float *a, float *b, float *c, int stride) {
    int idx = threadIdx.x * stride;  // 每个线程跳着访问
    c[idx] = a[idx] + b[idx];
}

// ✅ 合并访问（连续访问）
__global__ void good_access(float *a, float *b, float *c) {
    int idx = threadIdx.x;  // 连续地址
    c[idx] = a[idx] + b[idx];
}

你想想看，第一个核函数里，如果stride=2，那线程0访问地址0，线程1访问地址2……中间地址1、3被跳过了。硬件每次取128字节，结果只用了其中一半。浪费啊！

第二个核函数，线程0~31访问地址0~31，刚好在一个128字节块内。一次事务搞定。

4.5 合并访问的底层原理——内存事务

为什么硬件要这么设计？

因为全局内存的物理结构是DRAM芯片 + 内存控制器。每次读取，控制器会从DRAM中取出一个“行”（row），再从中选出一列（column）。这个“行”的大小，就是内存事务的大小。

在CUDA中，不同计算能力（Compute Capability）的事务大小不同：

• CC 2.x~5.x：128字节事务
• CC 6.x+：32字节或64字节事务（更灵活）

但不管哪种，核心思想不变：一次请求，服务尽量多的线程。

我个人习惯用一张图来理解这个过程：

4.6 实战技巧：如何写出合并访问的代码

我总结了几个实用原则：

1. 让threadIdx.x对应最内层维度——这是最常用的技巧。比如二维数组，让x方向是连续内存。
2. 避免使用AoS（数组 of 结构体）——改用SoA（结构体 of 数组）。
3. 注意对齐——每个warp访问的起始地址，最好是128字节的整数倍。
4. 使用共享内存做中转——如果必须随机访问，先加载到共享内存，再重新组织。

4.7 总结

全局内存慢，不是因为它本身慢，而是因为我们访问的方式不对。

合并访问，说白了就是让warp里的线程“结伴而行”，一起访问连续的内存地址。这样硬件可以用最少的事务，服务最多的线程。

记住一句话：线程是并行的，内存访问也要并行。

嗯，下一节我们会聊聊共享内存——那个比全局内存快两个数量级的“小快件仓库”。