GPU 访存合并：同一份代码，带宽能差 10 倍

新手写 CUDA 常遇到怪事：算法一样、线程数一样，换个数据排布方式，性能差出一个数量级。十有八九是**访存合并（memory coalescing）**的问题——GPU 内存系统的第一性原理。

## GPU 读内存是"按批"读的

GPU 的全局内存控制器一次事务读取 **32/64/128 字节**（一个 cache line 对齐的段）。一个 warp（32 个线程）发起访存时，硬件会看这 32 个线程要的地址：

- **合并访问**：线程 i 访问地址 `base + i`（连续），32 个请求合并成 **1~2 个**内存事务，带宽利用率接近 100%
- **跨步访问**：线程 i 访问 `base + i * stride`（比如按列访问行优先存储的矩阵），32 个请求分散到 32 个 cache line，变成 **32 个**事务，实际有用数据只占很小一部分

实测中，跨步访问的等效带宽经常只有合并访问的 **1/10 甚至更低**。注意：你的代码"逻辑"完全正确，只是慢——这是最隐蔽的性能坑。

## 典型场景与改法

```c
// 坏：按列访问行优先矩阵（stride = N）
float x = A[col * N + row]; // 32 个线程各读各的 cache line

// 好：转置后按行访问，或用 shared memory 做分块转置
float x = A_T[row * N + col]; // 连续访问，完美合并
```

工程上常用的三板斧：

1. **数据布局优先**：SoA（结构体数组）通常比 AoS（数组结构体）更容易合并；矩阵运算先想清楚行列优先
2. **Shared memory 做中转**：先把一块数据合并读入 shared memory，转置/重排后再计算。CUDA 矩阵乘法的经典分块优化，本质就是在解决合并问题
3. **向量化加载**：`float4` 一次读 16 字节，减少事务数；配合 `__ldg()` 只读缓存，效果叠加

## 怎么验证

- **Nsight Compute**：看 `dram__throughput` 和 `l1tex__t_bytes`，再看 "Memory Workload Analysis" 里的合并效率指标，一目了然
- **简单估算**：算一下你的 kernel 的算术强度（FLOP/byte）。如果远低于 GPU 的 ridge point（峰值算力/峰值带宽），那你就是访存 bound，优化合并访问是第一优先级

记住这个顺序：**先保证访存合并，再谈 shared memory，再谈指令级优化**。很多 CUDA 性能问题，根子上都是第一步没做好。

补充：除了合并访问，还要注意共享内存的 bank conflict 和数据对齐，常见做法是用共享内存做转置、结构体数组（AoS）改成数组结构体（SoA）。用 Nsight Compute 可以直接看全局内存的事务效率。你们在优化 kernel 时，是先看 memory throughput 指标还是先看 occupancy？