1、全局内存合并访问：为什么合并访问如此重要？从硬件架构看内存事务
各位同学，咱们今天聊聊全局内存合并访问。
说实话，我刚开始学CUDA那会儿，觉得这玩意儿就是个“锦上添花”的优化。后来在项目里被狠狠教育了一顿——一个核函数跑得比CPU还慢，查了半天，罪魁祸首就是内存访问没合并。嗯，从那以后，我再也不敢小看它了。
1.1 先说说内存事务是什么
要理解合并访问，得先搞清楚GPU怎么读内存。
GPU里的全局内存，说白了就是显存。它和CPU的内存不一样，CPU一次读一个字节、四个字节，但GPU不一样。GPU一次读一个“事务”（transaction），这个事务的大小通常是32字节、64字节，或者128字节。
为什么会这样？
因为GPU的设计哲学是“吞吐优先”。它不在乎单次访问的延迟，在乎的是单位时间内能搬多少数据。所以它把内存访问打包成固定大小的“包裹”，一次搬一大块。
我打个比方：
- CPU像快递员，一次送一个小包裹，但送得很快
- GPU像卡车，一次拉一大车货，但启动慢
你想想看，如果你让卡车每次只拉一个快递盒，那效率得多低？
1.2 合并访问到底是个啥？
合并访问，就是让同一个warp里的32个线程，访问的内存地址是连续的、对齐的。
举个例子：
// 合并访问：好！
int tid = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[tid]; // 每个线程访问连续地址
// 非合并访问：坏！
int tid = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[tid * 32]; // 每个线程跳着访问第一种写法，线程0读data[0]，线程1读data[1]……线程31读data[31]。这些地址是连续的，GPU可以一次事务就把这32个float（128字节）全部搬过来。
第二种写法，线程0读data[0]，线程1读data[32]……每个线程之间隔了128字节。GPU得发多个事务才能搞定，效率直接打骨折。
核心结论：合并访问时，一次内存事务服务整个warp。非合并访问时，一个warp可能需要多个事务，甚至几十个事务。
1.3 从硬件角度看，不合并有多痛？
我记得有一次调优，一个warp里32个线程，每个线程读一个float。按理说，一次128字节的事务就够了。
但我的代码写成了这样：
float val = data[threadIdx.x * 2]; // 步长为2结果呢？
每个线程读的地址是0, 2, 4, 6……这些地址分布在两个不同的128字节段里。GPU不得不发两个事务，每个事务只用了其中一半的数据。
浪费啊！
更惨的情况是，如果步长是32：
float val = data[threadIdx.x * 32]; // 步长为32每个线程读的地址分布在不同的128字节段里。32个线程，可能需要32个事务！
你算算，本来1个事务能搞定的事，现在要32个。带宽利用率从100%掉到3%左右。这性能能好才怪。
| 访问模式 | 事务数 | 带宽利用率 |
|---|---|---|
| 连续、对齐（合并） | 1 | ~100% |
| 步长为2 | 2 | ~50% |
| 步长为4 | 4 | ~25% |
| 随机访问 | 最多32 | ~3% |
注意：这里说的“对齐”是指起始地址要是32字节、64字节或128字节的整数倍。我曾经踩过一个坑——结构体数组（AoS）的访问模式，天然就是不对齐的。后来改成数组结构体（SoA），性能直接翻倍。
1.4 为什么GPU设计成这样？
你可能会问：为什么GPU不设计成每个线程独立访问内存？那样不就不用操心合并了吗？
嗯，这个问题我当年也想过。
答案其实很简单：成本与功耗。
如果每个线程都有独立的内存通道，那芯片面积会爆炸，功耗也压不住。GPU选择了“共享内存通道”的方案——32个线程共享一个内存控制器。这样硬件简单、功耗低，但代价就是要求程序员配合，让访问模式“合并”。
说白了，GPU把一部分优化责任交给了你。你写得好，它就飞起来；你写得不好，它就趴窝。
1.5 我个人的习惯
写CUDA代码时，我一般会先问自己三个问题：
- 这个warp里的线程，访问的地址是不是连续的？
- 起始地址是不是对齐的？
- 有没有办法把随机访问变成顺序访问？
如果答案都是“是”，那内存这块基本稳了。
如果答案有“否”，我会考虑：
- 能不能用共享内存做中转？
- 能不能调整数据布局（比如AoS转SoA）？
- 能不能用向量化加载（比如float4）？
小技巧：用 nvprof 或 Nsight Compute 查看 global_load_transactions 和 global_store_transactions 指标。如果事务数远大于理论最小值，说明你的访问模式有问题。
1.6 总结一下
合并访问，说白了就是让GPU的“卡车”一次拉满货，别空跑。
从硬件角度看，一次内存事务的成本是固定的。你让事务多干活，效率就高；你让事务少干活，效率就低。
我见过太多新手，核函数写得花里胡哨，结果内存访问一塌糊涂。性能瓶颈往往不在计算，而在内存。记住：GPU算得再快，数据喂不进去也是白搭。
下一节，咱们聊聊具体的合并访问模式，以及怎么用代码实现它。