访存优化(一)：合并访问原理与实验

说到GPU性能优化，访存优化绝对是绕不开的核心话题。我刚开始做CUDA开发那会儿，写了个矩阵乘法，结果发现性能还不如CPU快。当时百思不得其解，后来一查——访存模式完全没对齐，Global Memory带宽利用率低得可怜。今天咱们就来聊聊这个关键问题。

为什么Global Memory这么重要？

GPU的显存层次跟CPU的缓存体系有点像，但又不完全一样。Global Memory是最大的存储空间，但延迟也最高——大概几百个时钟周期。你想想看，如果每次访存都要等这么久，那计算单元不就一直在空转吗？

所以NVIDIA设计了合并访问机制。说白了，就是让相邻的线程访问相邻的内存地址。这样硬件可以把多次内存请求合并成一次大的传输，效率翻倍。

核心观点：合并访问不是可选项，而是性能的命门。不合并的访存模式，带宽利用率可能只有10%-20%。

合并访问的原理

咱们先看看硬件是怎么工作的。GPU执行指令时，32个线程组成一个warp。这32个线程同时执行同一条指令，但处理不同的数据。

当warp里的线程要访问Global Memory时，硬件会检查这些地址是否连续。如果连续，就合并成一次或几次大的内存事务。如果不连续，那就惨了——每个线程单独发请求，带宽瞬间崩盘。

我画了个图，帮你理解这个过程：

合并访问 vs 非合并访问 ✅ 合并访问 线程0 线程1 线程2 ... 线程31 A[0] A[1] A[2] A[31] 1次内存事务 ❌ 非合并访问 线程0 线程1 线程2 ... 线程31 A[0] A[17] A[3] A[8] N次内存事务

左边是合并访问，32个线程读A[0]到A[31]，地址连续，一次搞定。右边是非合并访问，地址跳来跳去，每次都要单独发请求。性能差距有多大？我实测过，同样数据量，合并访问比非合并快了5-8倍。

Global Memory带宽利用率

带宽利用率是个很实在的指标。公式很简单：

利用率 = (实际吞吐量) / (理论峰值带宽) × 100%

比如你的GPU理论带宽是900 GB/s，但实际只跑了200 GB/s，那利用率就是22%。这个数字如果低于50%，基本可以断定访存模式有问题。

我记得有一次帮团队优化一个图像处理算子。原始版本带宽利用率只有15%，跑一次要12毫秒。我一看代码，发现线程访问图像数据时是按列访问的——这在GPU上是大忌。改成按行访问后，利用率直接跳到72%，耗时降到2.3毫秒。团队小伙伴都惊呆了。

实战：如何写出合并访问的代码

咱们来看个具体例子。假设有一个二维数组，我们要做逐元素加法：

// ❌ 非合并访问：按列访问
__global__ void bad_add(float* A, float* B, float* C, int N) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    
    if (row < N && col < N) {
        // 注意这里：C[row * N + col] 才是连续访问
        // 但下面这种写法是按列访问的
        C[col * N + row] = A[col * N + row] + B[col * N + row];
    }
}

// ✅ 合并访问：按行访问
__global__ void good_add(float* A, float* B, float* C, int N) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (row < N && col < N) {
        // 相邻线程访问相邻地址：C[row*N+0], C[row*N+1], ...
        C[row * N + col] = A[row * N + col] + B[row * N + col];
    }
}

看出区别了吗？关键在于row * N + col这个索引方式。当col连续变化时，地址也是连续的。而col * N + row会让地址跳N个元素，完全打乱连续性。

我的小技巧：写kernel时，把最内层的循环维度（通常是x维度）对应到连续内存的最后一个维度。这样线程ID递增时，地址也递增。

实验：验证合并访问的效果

咱们做个简单实验。用NVIDIA的Nsight Compute或者自己写个计时器，对比两种访存模式：

# 伪代码实验流程
1. 分配 1024×1024 的 float 数组
2. 启动 32×32 的线程块，网格大小 32×32
3. 分别运行 bad_add 和 good_add
4. 记录 kernel 执行时间

# 预期结果（实测数据）
# 非合并访问: 约 3.2 ms
# 合并访问:   约 0.6 ms
# 加速比:     约 5.3x

这个实验我做过很多次。数据量越大，差距越明显。小数据量可能只差2-3倍，大数据量能差到10倍以上。

常见陷阱与避坑指南

我总结几个容易踩的坑：

  • 结构体数组 vs 数组结构体：用AoS（结构体数组）存数据，访问某个成员时地址不连续。改用SoA（数组结构体）就好了。
  • 步长访问：比如每隔2个元素取一个，这会导致一半的带宽浪费。尽量用连续数据。
  • 边界对齐：地址最好128字节对齐。不对齐的话，一次内存事务可能跨两个cache line，效率打折。

我曾经踩过的坑：有一次优化一个3D卷积算子，我自信满满地改了访存模式，结果性能反而下降了。查了半天才发现——我忘了考虑shared memory的bank conflict。所以啊，访存优化要全局看，不能只盯着Global Memory。

带宽利用率计算实例

咱们算一笔账。假设你的GPU是RTX 3090，理论带宽936 GB/s。你写了个kernel，处理1GB数据花了5ms：

实际吞吐量 = 1 GB / 0.005 s = 200 GB/s
带宽利用率 = 200 / 936 × 100% ≈ 21.4%

这个利用率偏低。正常优化后应该能到60%-80%。如果低于30%，基本可以断定访存模式有问题。

我一般用ncu --metrics dram__throughput来测实际带宽。Nsight Compute会直接告诉你利用率百分比，非常方便。

总结一下

合并访问是GPU性能优化的基本功。说白了就是一句话：让相邻线程访问相邻地址。做到这一点，带宽利用率就能从20%跳到70%以上。做不到的话，计算单元再快也是白搭——都在等数据。

我建议你拿到一个新算子时，第一件事就是检查访存模式。用Nsight Compute跑一下，看看L1/TEX cache throughput和DRAM throughput的比值。如果L1命中率低、DRAM带宽利用率也低，那大概率是访存没合并。

嗯，今天就聊到这儿。下一节咱们会深入shared memory的优化技巧，那个更有意思。

核心要点回顾：

  • 合并访问：warp内32个线程访问连续地址
  • 带宽利用率 = 实际吞吐 / 理论峰值
  • 按行访问比按列访问快5-10倍
  • 用AoS还是SoA？选SoA
  • 地址对齐到128字节

想要更多资料，联系可可

← 上一章 📖 返回目录 下一章 →