内存访问优化：合并访问、对齐与跨步、Bank Conflict 避免

说到 GPU 算子优化，内存访问这块儿绝对是绕不开的硬骨头。很多人写 CUDA 代码，算力利用率上不去，一看 profiler，发现大部分时间都卡在显存读写上。说白了，GPU 的计算单元快得飞起，但显存带宽就那么宽，你访问模式不对，性能直接腰斩。

我个人习惯把内存访问优化分成三个层次：合并访问、对齐与跨步、Bank Conflict 避免。这三层是递进关系，从全局内存到共享内存，一层层抠细节。今天咱们就一个一个说清楚。

核心观点：GPU 内存访问优化的本质，就是让硬件的数据搬运效率最大化。你写代码时多花十分钟思考访问模式，可能换来几倍的性能提升。

1. 合并访问（Coalesced Access）

先问一个问题：当一个 warp（32 个线程）同时去全局内存读数据时，硬件是怎么做的？

答案是：它会把 32 个线程的请求合并成尽量少的内存事务（memory transaction）。如果 32 个线程访问的地址是连续的，那一个 128 字节的事务就能搞定。如果地址是随机分散的，那就得拆成几十个事务，带宽利用率直线下降。

我在项目中遇到过最典型的反面教材：有人写了一个向量加法，但数据布局是 AoS（Array of Structures）而不是 SoA（Structure of Arrays）。32 个线程读同一个结构体的不同字段，地址跨步很大，结果合并访问完全失效，性能只有理论峰值的 10%。

来看个例子：

// ❌ 非合并访问：每个线程读一个结构体，地址不连续
struct Particle { float x, y, z; };
Particle particles[N];
float sum = 0.0f;
for (int i = threadIdx.x; i < N; i += blockDim.x) {
    sum += particles[i].x;  // 地址跨步为 sizeof(Particle) = 12 bytes
}

// ✅ 合并访问：使用 SoA 布局，地址连续
float x[N], y[N], z[N];
float sum = 0.0f;
for (int i = threadIdx.x; i < N; i += blockDim.x) {
    sum += x[i];  // 地址连续，完美合并
}

我的经验：写 kernel 之前，先画一张内存布局图。看看同一个 warp 里的线程访问的地址是不是连续的。如果是，恭喜你，合并访问达标了。

2. 对齐与跨步（Alignment & Stride）

合并访问搞定了，下一个坑是对齐。GPU 的内存事务有固定大小：32 字节、64 字节、128 字节。如果你的数据起始地址不是 128 字节对齐的，那一个事务可能只用到一部分数据，剩下的带宽就浪费了。

举个例子：假设你有一个 float 数组，起始地址是 4 字节对齐（默认 malloc 行为）。当 warp 读取前 32 个 float 时，地址范围是 [0, 128)。硬件会发一个 128 字节的事务，刚好覆盖。但如果起始地址是 100 字节，那事务覆盖 [0, 128) 和 [128, 256) 两个区间，多读了一倍的数据。

跨步访问更隐蔽。我记得有一次调试一个矩阵转置的 kernel，性能死活上不去。后来发现是因为行主序存储，但列主序读取，导致每个线程读取的地址间隔是矩阵的列数（很大），完全无法合并。

// ❌ 跨步访问：按列读取行主序矩阵
__global__ void bad_kernel(float* mat, int rows, int cols) {
    int row = threadIdx.x;
    int col = blockIdx.x;
    float val = mat[col * rows + row];  // 跨步 = rows，很大
}

// ✅ 合并访问：按行读取
__global__ void good_kernel(float* mat, int rows, int cols) {
    int row = blockIdx.x;
    int col = threadIdx.x;
    float val = mat[row * cols + col];  // 连续地址
}

注意：跨步访问有时候不可避免，比如某些卷积操作。这时候可以考虑用共享内存做中转，先把数据从全局内存搬到共享内存（合并访问），再从共享内存里按需读取（随机访问）。

3. Bank Conflict 避免

共享内存是 GPU 上的一块高速缓存，但它有 32 个 bank（以 Fermi 架构为例）。每个 bank 每周期只能响应一个请求。如果同一个 warp 里的多个线程访问同一个 bank 的不同地址，就会发生 bank conflict，导致访问串行化。

嗯，这里要注意：bank conflict 只发生在共享内存的随机访问场景。如果你所有线程都读同一个地址（广播），那没问题，硬件会优化。但如果是不同地址落在同一个 bank 上，那就糟了。

我曾经写过一个直方图统计的 kernel，用共享内存做局部统计。32 个线程同时更新 32 个不同的 bin，结果发现性能极差。一查，原来 bin 的地址刚好映射到同一个 bank 上，所有更新都串行了。

怎么避免？核心思路是地址偏移：

// ❌ 有 bank conflict：每个线程访问 s_data[threadIdx.x]
__shared__ float s_data[32];
float val = s_data[threadIdx.x];  // 如果多个线程访问不同索引但落在同一 bank...

// ✅ 无 bank conflict：使用 padding 或重新排列
__shared__ float s_data[33];  // 多一个元素做 padding，打破 bank 对齐
float val = s_data[threadIdx.x];  // 现在 bank 映射被打散了

避坑指南：我曾经在实现一个 3x3 卷积的共享内存版本时，发现 bank conflict 导致性能只有理论值的 60%。后来加了一个元素的 padding，性能直接翻倍。记住：共享内存大小加 1 往往能解决大部分 bank conflict。

知识体系总览

下面这张图总结了本章的核心逻辑，从全局内存到共享内存，逐层优化：

实战建议

说了这么多，总结几条我自己的经验：

先看 profiler：别猜，用 nvprof 或 Nsight Compute 看 memory 相关的 metrics。重点关注 global load/store efficiency 和 shared memory bank conflict。
数据布局优先：写 kernel 之前先想好数据是 SoA 还是 AoS。我见过太多人因为偷懒用 AoS，结果性能惨不忍睹。
共享内存 padding 是银弹：遇到 bank conflict 先试 padding，简单有效。但注意 padding 会增加共享内存用量，别超过 48KB。
跨步访问用共享内存中转：如果全局内存访问模式注定是跨步的，那就先合并读到共享内存，再从共享内存里随机读。虽然多了一次拷贝，但往往比直接跨步访问快。

最后提醒：优化内存访问不是一蹴而就的。我每次写新 kernel，都会先写一个朴素版本，然后看 profiler，再针对性地优化。别一开始就想着完美，迭代才是王道。

好了，这一章的内容就到这里。记住：合并访问是基础，对齐跨步是细节，bank conflict 是进阶。把这三点吃透，你的 GPU 算子性能至少能翻一倍。

核心公式：性能 ≈ 计算量 / (访存量 × 访问效率)。访问效率就是合并、对齐、bank conflict 的综合体现。

想要更多资料，联系可可

📚 课程目录

内存访问优化：合并访问、对齐与跨步、Bank Conflict 避免

1. 合并访问（Coalesced Access）

2. 对齐与跨步（Alignment & Stride）

3. Bank Conflict 避免

知识体系总览

实战建议