2、全局内存访问模式：合并访问原理、对齐要求、步长访问对带宽的影响
好，我们直接进入正题。全局内存访问，是GPU性能优化的核心痛点之一。很多同学写CUDA代码，觉得“反正数据在显存里，怎么读都一样”。大错特错。我见过太多项目，就因为内存访问模式没调好，性能直接腰斩，甚至更惨。
说白了，全局内存的带宽虽然高，但它有脾气。你得顺着它的性子来。这个“性子”，就是合并访问（Coalesced Access）。
2.1 合并访问：到底在合并什么？
先问一个问题：一个warp（32个线程）同时发起内存访问，硬件会怎么处理？
答案是：硬件会把这些访问请求合并成尽可能少的事务（memory transaction）。每个事务的大小通常是32字节或128字节。合并得越好，事务数越少，带宽利用率越高。
我个人习惯把合并访问想象成“拼车”。32个线程要去取数据，如果它们要的数据在内存里是连续挨着的，那硬件就派一辆大车（128字节事务）一次性拉回来。如果数据是东一个西一个，那就得派多辆车，甚至每辆车只拉一点点数据。你说哪个效率高？
核心原则：一个warp内的线程，访问的地址要尽量连续，并且对齐到128字节边界。
我在项目中遇到过这样一个案例：一个图像处理核，每个线程处理一个像素，但像素在内存中是按行存储的。如果线程0访问像素0，线程1访问像素1，线程2访问像素2……这就是完美的合并访问。但如果线程0访问像素0，线程1访问像素1024（下一行的开头），那就完蛋了——每次访问都跨行，硬件根本没法合并。
2.2 对齐要求：地址要对齐到哪？
对齐，说白了就是地址要“整”。硬件喜欢地址是32字节或128字节的整数倍。
为什么？因为内存控制器是按“段”来取数据的。一个段就是32字节或128字节。如果你的起始地址不是段的起始地址，那一次请求可能跨越两个段，硬件就得发两次事务。
举个例子：
// 对齐访问（推荐）
float *data;
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[idx]; // 每个线程访问连续地址，且起始地址对齐
// 未对齐访问（性能差）
float *data;
int idx = blockIdx.x * blockDim.x + threadIdx.x + 1; // 偏移了1个float
float val = data[idx]; // 起始地址不对齐到128字节边界
嗯，这里要注意：偏移一个元素，可能就让你的带宽利用率从100%掉到25%。我曾经调试过一个模型推理的算子，性能死活上不去，最后发现是输入数据的地址偏移了4个字节。改回来之后，速度直接快了3倍。
避坑指南：我曾经在分配显存时，用了cudaMallocPitch而不是cudaMalloc，结果因为pitch不是128字节对齐的，导致后续所有访问都未对齐。调试了整整两天才发现。所以，分配内存时尽量用cudaMalloc，或者确保pitch是128的倍数。
2.3 步长访问：为什么步长越大，带宽越差？
步长访问（Strided Access），就是每个线程访问的数据之间有空隙。比如线程0访问地址0，线程1访问地址4，线程2访问地址8……步长就是4个float（16字节）。
为什么会这样？你想想看：
- 步长为1：32个线程访问连续地址，硬件发1个128字节事务就够了。
- 步长为2：32个线程访问地址0,2,4,...，数据分散在64个float范围内。硬件需要发2个128字节事务，但只用到其中一半的数据。带宽利用率50%。
- 步长为4：需要4个事务，利用率25%。
- 步长为N：利用率大约为1/N。
我整理了一个表格，方便你直观感受：
| 步长（单位：float） | 事务数（128字节） | 带宽利用率 | 典型场景 |
|---|---|---|---|
| 1 | 1 | 100% | 连续数组访问 |
| 2 | 2 | 50% | 访问二维矩阵的列（列优先存储） |
| 4 | 4 | 25% | 隔点采样 |
| 8 | 8 | 12.5% | 稀疏数据访问 |
| 随机 | 32 | ~3% | 散列访问、哈希表 |
你看，步长从1变成2，带宽利用率直接腰斩。这就是为什么我总说：能连续访问，就别跳着访问。
小技巧：如果你必须按列访问一个二维数组（比如矩阵转置），可以考虑先用共享内存做一次转置，再按行写入。虽然多了共享内存的开销，但往往比直接按列访问全局内存快得多。我自己的项目中，用这个方法把矩阵转置的速度提升了4倍。
2.4 实战：如何检查你的访问模式？
光说不练假把式。我建议你用NVIDIA的Nsight Compute工具，它可以直接告诉你每个kernel的全局内存访问效率。
具体看这几个指标：
- Global Load Efficiency：实际使用的数据量 / 总读取数据量。理想值是100%。如果低于50%，说明你的访问模式有问题。
- Global Store Efficiency：同理，看写入效率。
- Sector Utilization：每个sector（32字节）被用到的比例。如果很多sector只用了1个字节，那就是典型的未合并访问。
我记得有一次帮一个团队优化推荐系统模型，他们的Embedding lookup kernel效率只有12%。一看，每个线程访问的Embedding向量是随机分布的，步长巨大。后来我们改成了让一个warp内的线程访问连续的Embedding slot，效率直接拉到了85%以上。
2.5 总结一下
全局内存访问，说白了就三句话：
- 让一个warp内的线程访问连续地址——步长为1最好。
- 确保起始地址对齐到128字节——别小看这4个字节的偏移。
- 避免随机访问和跨行访问——能用共享内存转置就别硬来。
你想想看，GPU的显存带宽动辄几百GB/s，但如果你访问模式不对，实际能用的可能只有几十GB/s。这就像开着一辆法拉利，却在堵车的胡同里慢慢挪——车再好也白搭。
下一章我们会聊共享内存的bank conflict，那又是一个容易踩坑的地方。到时候我会分享一个我踩过的“血泪史”，保证让你印象深刻。