GPU 内核访存优化：coalescing、共享内存与 bank conflict
说明：本篇为综述/无第一手实验数据，采用公开实践与 CUDA 编程模型路线。
实验目标
很多 GPU 内核跑得慢不是"算不动"而是"取数慢"——访存带宽没用满。GPU 的并行度全靠"大量线程同时访存"，而访存效率由三个机制决定：全局内存的合并访问（coalescing）、共享内存的正确使用、bank conflict 的规避。本篇目标：用可复现代码拆解这三个机制——为什么连续访问快、共享内存怎么用、bank conflict 怎么避免，以及向量化加载。读完能回答：这个内核的访存效率在哪丢了、怎么改。
代码逐段拆解
第一段：合并访问 vs 非合并访问
cuda
// 坏例子：线程按列取数（非合并访问，带宽浪费 8-32 倍）
__global__ void bad(float* a, float* b, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
b[i] = a[i * 32]; // 线程 i 取 a[i*32]：间隔 32 个元素
}
// 好例子：线程连续取数（合并访问，一次事务取满）
__global__ void good(float* a, float* b, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
b[i] = a[i]; // 相邻线程取相邻地址：完美合并
}要点：
- 合并访问的定义：同一 warp（32 线程）同时访问连续地址——硬件把一次 warp 访存合并成最少的内存事务（128B/事务）。非连续（跨步/转置）访问则每次只用了事务的一部分，带宽利用率掉到 1/8-1/32。
- 90% 的访存问题是数据布局与线程映射不匹配：数组布局按"线程 i 取 a[i]"设计，而不是按"逻辑索引"设计——转置、按列遍历、结构体数组（AoS）都是经典坑。
- 验证方法：Nsight Compute 的 memory workload 分析（08 章 nsight 篇）直接给"合并度/带宽利用率"——先测后改。
第二段：共享内存与 bank conflict
cuda
// 共享内存：块内共享、比全局内存快 ~100 倍，但按 bank 组织
__global__ void smem(float* in, float* out, int n) {
__shared__ float tile[32][33]; // 注意：第 2 维故意 +1（33 而非 32）
int i = threadIdx.x;
tile[i][0] = in[blockIdx.x * 32 + i]; // 写入
__syncthreads();
// 读 tile[i][i]：如果第 2 维是 32，同一 bank 冲突；33 则错开
out[blockIdx.x * 32 + i] = tile[i][i];
}要点：
- 共享内存按 bank 组织：32 个 bank（每 bank 4B），同一 warp 访问同一 bank 的不同地址 = bank conflict（串行化，慢 2-32 倍）；访问不同 bank = 并行。
tile[32][33]的 +1 是经典技巧：二维数组按行存储，行宽 32 时第 i 行第 i 列全落同一 bank（对角访问冲突）；行宽 +1（33）把每行错开，对角访问落入不同 bank——"padding 一列"是规避 bank conflict 最常用的手段。__syncthreads()是共享内存的纪律：写后必须同步再读，否则块内线程读到未写完的数据。
原理讲解
三层访存层次与优化顺序
全局内存（几百 GB，慢） → 共享内存（每 SM 几十 KB，快 100 倍）
→ 寄存器（最快，但私有）优化顺序：① 先保证全局内存合并访问（数据布局）→ ② 数据复用多时搬进共享内存 → ③ 循环内尽量用寄存器暂存 → ④ 最后处理 bank conflict。共享内存只在该数据被块内多线程复用（如矩阵分块、卷积滑动窗）时才值得搬——一次读就用的数据放共享内存反而亏。
向量化加载（float4）
cuda
// 用 float4 一次读 16B：减少指令数、提升带宽利用（对齐要求）
float4 v = reinterpret_cast<float4*>(a)[i];- 向量化与合并访问配合：每个线程读 float4（16B），warp 一次读 512B——比逐 float 读少 4 倍指令，带宽更满。
- 对齐要求：float4 要求 16B 对齐——数据布局要按 16B 对齐分配，否则 illegal address。
访存模式 vs 计算模式匹配
| 计算模式 | 访存模式 | 优化 |
|---|---|---|
| 逐元素（加/乘） | 连续流 | 合并访问 + 向量化 |
| 分块（矩阵乘/卷积） | 分块复用 | 共享内存 + padding |
| 归约（求和/找极值） | 跨线程 | 共享内存分步归约 |
核心认知：访存优化不是"加缓存"，而是让访问模式匹配硬件的并行访存机制（合并、bank、向量化）——先按计算模式设计数据布局，再按访存模式选优化手段。
实验结果
（估算口径，非实测；带宽受限内核）
| 优化 | 带宽利用提升 | 代价 |
|---|---|---|
| 合并访问（布局修正） | 2-8 倍 | 改数据布局/索引 |
| 共享内存 + padding | 1.5-4 倍（复用场景） | 代码复杂度 |
| float4 向量化 | 1.3-2 倍 | 对齐约束 |
工程启示
- 先布局后优化：数据布局按"线程 i 取 a[i]"设计，合并访问是访存优化的地基——布局错了后面都白搭。
- 共享内存在"复用"时才用：矩阵分块/卷积等复用场景才搬共享内存；一次读场景直接全局读。
- padding 一列治 bank conflict：对角/跨步访问冲突时，行宽 +1 是最快解法。
- Nsight 验证合并度：memory workload 指标（08 章 nsight 篇）量化"合并度/带宽利用"，优化前后对比以数据验收。
下一篇预告：访存优化讲完，08 编译器章与 04 存储章本轮都补到位；下一批可回 11 训练章或 13 应用章继续，或转 SEO 收尾。由你裁量。