3、内存访问模式优化：合并访问（Coalesced Access）原理与实践、对齐与跨步访问的影响
说到GPU性能优化，内存访问模式绝对是个绕不开的话题。我见过太多同学，核函数写得挺漂亮，计算逻辑也没问题，一跑起来就是慢。为什么？八成是内存访问没搞好。
说白了，GPU的算力再强，如果数据喂不进去，那也是白搭。今天我们就来聊聊这个核心问题——合并访问。
3.1 合并访问：GPU内存的“团购”机制
先问一个问题：你知道GPU一次能从显存里读多少数据吗？
嗯，不是1个float，也不是4个float。GPU的最小内存事务单位是32字节，也就是8个float。更关键的是，一个warp（32个线程）执行同一指令时，硬件会尽量把这32个线程的内存请求合并成尽可能少的内存事务。
这就是合并访问（Coalesced Access）的本质。
核心概念：当warp中32个线程访问的内存地址是连续且对齐的，硬件可以用最少的内存事务完成数据加载。反之，地址分散会导致大量冗余事务，带宽利用率急剧下降。
我在项目中遇到过这样一个案例：一个简单的向量加法，合并访问版本跑出800 GB/s的带宽，非合并版本直接掉到120 GB/s。同样的计算，差了将近7倍。你想想看，这差距有多大。
3.2 对齐访问：地址要对齐到128字节
对齐，说白了就是你的数据起始地址要“规整”。GPU的L2缓存和DRAM控制器喜欢128字节对齐的访问。
什么意思？举个例子：
// 对齐访问：地址是128字节的整数倍
float* d_data;
cudaMalloc(&d_data, N * sizeof(float)); // 默认对齐到256字节
// 非对齐访问：偏移了4个字节
float* d_unaligned = d_data + 1; // 地址偏移了4字节为什么对齐这么重要？因为GPU的内存事务是固定大小的。如果地址不对齐，一次事务可能跨越两个128字节的缓存行，导致额外的事务开销。
个人习惯：我一般用cudaMalloc分配内存，它默认会做对齐。但如果自己管理内存池，一定要用cudaMalloc或posix_memalign确保128字节对齐。
3.3 跨步访问：最隐蔽的性能杀手
跨步访问（Strided Access）是很多新手容易踩的坑。你以为自己在顺序访问，其实线程之间在“跳着走”。
看个典型例子：
// 跨步访问：每个线程访问间隔为N
__global__ void strided_kernel(float* data, int N) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
for (int i = 0; i < N; i++) {
data[tid + i * gridDim.x * blockDim.x] = 1.0f; // 跨步！
}
}
// 合并访问：连续地址
__global__ void coalesced_kernel(float* data, int N) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
for (int i = 0; i < N; i++) {
data[tid * N + i] = 1.0f; // 连续！
}
}为什么会这样？因为warp内32个线程的threadIdx.x是连续的。在跨步版本中，线程0访问地址0，线程1访问地址N，线程2访问地址2N……这些地址之间隔了N个元素，完全不连续。硬件没法合并，只能一个一个事务去读。
我曾经接手过一个矩阵转置的优化项目，原版代码用了跨步访问，带宽利用率只有15%。改成合并访问后，直接飙到85%。嗯，有时候性能问题就是这么简单粗暴。
3.4 结构体数组 vs 数组结构体
这个经典问题，我再说一遍：
| 布局方式 | 内存模式 | 适用场景 |
|---|---|---|
| AoS（结构体数组） | 跨步访问 | CPU缓存友好 |
| SoA（数组结构体） | 合并访问 | GPU内存友好 |
说白了，GPU上尽量用SoA。比如你要处理粒子的位置和速度：
// 不推荐：AoS
struct Particle {
float x, y, z;
float vx, vy, vz;
};
Particle particles[N];
// 推荐：SoA
struct Particles {
float* x, *y, *z;
float* vx, *vy, *vz;
};SoA的好处是，当warp里的线程各自处理一个粒子时，所有线程的x坐标是连续存储的，一次内存事务就能搞定。
3.5 知识体系总览
下面这张图，是我自己总结的内存访问优化核心逻辑：
3.6 实战建议
最后，给你几个我踩坑踩出来的建议：
- 用profiler看内存事务：ncu（Nsight Compute）里有个指标叫L1/TEX Cache Throughput，如果实际事务数远大于理论最小值，说明访问模式有问题。
- 尽量用一维索引：二维和三维block虽然方便，但容易搞乱内存访问顺序。我习惯先把多维索引转成一维，再算地址。
- 注意shared memory的bank conflict：虽然今天讲的是全局内存，但shared memory也有类似问题，原理相通。
- 别迷信“连续”：有时候为了连续访问，把数据结构改得乱七八糟，反而得不偿失。先profiling，再优化。
一个小技巧：如果你不确定自己的访问模式是否合并，可以在核函数里加一句__syncthreads()，然后用ncu看内存事务数。如果事务数等于ceil(32 * sizeof(type) / 128)，那就是完美合并。
嗯，内存访问模式这块，说白了就是“让数据排好队，让GPU一口吃掉”。你想想看，GPU有几千个核心，如果每个核心都在等数据，那得浪费多少时间。合并访问就是让数据流动起来的关键。
下次写核函数的时候，多想想你的数据在显存里是怎么放的。相信我，这个习惯能帮你省下大把的优化时间。