4. 内存访问模式优化

大家好，我是你们的老朋友。今天咱们来聊聊GPU内存访问模式。说实话，这是很多人在写CUDA/HIP程序时最容易忽略、但影响最大的地方。我见过太多人把核函数写得花里胡哨，结果性能被内存访问拖垮了。嗯，咱们今天就把这块硬骨头啃下来。

1. 合并访问（Coalesced Access）—— 性能的命脉

先问大家一个问题：为什么GPU的显存带宽那么高，但你的程序跑起来还是慢？

答案很简单——你没有利用好合并访问。说白了，合并访问就是让同一个warp（32个线程）访问的显存地址，在物理上尽量连续。这样GPU可以用一次内存事务，把32个线程需要的数据一次性搬回来。

核心原则：相邻线程访问相邻地址。

我举个例子。假设你有一个float数组，每个线程要读一个元素：

// ✅ 合并访问 —— 好！
int idx = threadIdx.x + blockIdx.x * blockDim.x;
float val = data[idx];  // 线程0读data[0]，线程1读data[1]...

// ❌ 非合并访问 —— 差！
int idx = threadIdx.x + blockIdx.x * blockDim.x;
float val = data[idx * 2];  // 线程0读data[0]，线程1读data[2]... 中间跳过了

我在项目中遇到过这样一个案例：一个矩阵转置的核函数，原本用非合并方式写，带宽利用率只有20%。改成合并访问后，性能直接翻了4倍。你想想看，只是改了一下索引顺序，效果就这么明显。

2. 对齐与跨步访问 —— 细节决定成败

对齐访问，指的是每个线程访问的地址，最好是它自身数据类型大小的整数倍。比如float是4字节，地址最好是4的倍数。这听起来简单，但实际中很容易踩坑。

跨步访问（Strided Access）就更常见了。比如你有一个结构体数组：

struct Particle {
    float x, y, z;  // 每个粒子3个float
};

Particle* particles;
// 如果只取x分量：
float x = particles[threadIdx.x].x;  // 跨步为3个float

这种访问模式，每个线程只取结构体中的一个字段，导致相邻线程访问的地址间隔很大。GPU一次内存事务只能搬回连续的一段数据，结果大部分带宽都浪费了。

我的建议：遇到结构体数组，尽量拆成数组结构体（SoA）。把x、y、z分别放到三个独立的数组里。这样相邻线程访问的就是连续地址了。

我曾经在一个粒子模拟项目中，把AoS改成SoA，内存带宽利用率从15%提升到了85%。嗯，这个优化几乎不花什么代码成本，效果却立竿见影。

3. 共享内存Bank Conflict —— 隐藏的杀手

共享内存是片上的，速度很快。但它有个特点：被分成了32个Bank（存储体）。每个Bank每周期只能响应一个请求。如果同一个warp里的多个线程，同时访问同一个Bank的不同地址，就会发生Bank Conflict，导致访问串行化。

举个例子：

__shared__ float smem[32][32];

// 情况1：无冲突
float val = smem[threadIdx.x][threadIdx.y];  // 每个线程访问不同行，不同Bank

// 情况2：有冲突
float val = smem[threadIdx.y][threadIdx.x];  // 同一行的不同列，可能冲突

为什么情况2会冲突？因为共享内存的Bank是按照地址的低5位来分配的。如果多个线程访问的地址低5位相同，就会打到同一个Bank上。

注意：Bank Conflict只发生在同一个warp内。不同warp之间没有这个问题。

我记得有一次调试一个矩阵乘法的核函数，性能始终上不去。用NVIDIA的Nsight工具一看，Bank Conflict率高达80%。后来我把共享内存的声明从smem[32][32]改成smem[32][33]（加了一列padding），冲突就几乎消失了。性能提升了30%。

4. 内存事务与缓存行 —— 底层原理

GPU从显存读取数据时，不是按字节读的，而是按缓存行（Cache Line）来读。在NVIDIA的架构中，一个缓存行通常是128字节。也就是说，哪怕你只读一个float（4字节），GPU也会把周围128字节的数据都搬回来。

这就解释了为什么合并访问这么重要：

  • 如果32个线程读连续的128字节，一次内存事务就搞定了。
  • 如果32个线程读分散的地址，可能需要多次内存事务，带宽就浪费了。

内存事务的大小也有讲究。对于全局内存，GPU会尽量使用128字节的事务。但如果访问模式不对，可能会退化成32字节的事务，效率直接打折扣。

避坑指南：我曾经在写一个图像滤波的核函数时，每个线程只处理一个像素，但像素的存储格式是RGB（3个字节）。结果每个线程读3个字节，地址不对齐，导致大量内存事务浪费。后来我把数据打包成float4（16字节对齐），性能就正常了。

5. 知识体系总览

下面这张图，是我自己总结的内存访问优化知识体系。你可以把它当作一个检查清单：

GPU内存访问优化知识体系 内存访问优化 合并访问 相邻线程访问相邻地址 一次内存事务搬128字节 对齐与跨步访问 地址对齐到数据类型大小 避免跨步，使用SoA Bank Conflict 32个Bank，每周期一个请求 加padding避免冲突 内存事务与缓存行 缓存行128字节 事务大小：128B vs 32B 优化策略 数据重排、padding 使用向量化加载（float4） 核心：让内存访问模式匹配硬件的工作方式 工具：Nsight Compute / ROCProfiler 分析内存事务

6. 实战建议

说了这么多，给大家几个可以直接用的建议：

  1. 优先保证合并访问。写核函数时，先问自己：相邻线程访问的地址连续吗？如果不连续，想办法重排数据或调整索引。
  2. 用float4或int4加载数据。一次加载16字节，比4次加载4字节效率高得多。我习惯在核函数开头用float4把数据读进来，然后再拆开用。
  3. 共享内存加padding。如果声明smem[N][N]有Bank Conflict，试试smem[N][N+1]。多出来的那一列不用，但能避免冲突。
  4. 用工具说话。别猜，用Nsight Compute或ROCProfiler看内存事务数。如果实际事务数远大于理论最小值，说明访问模式有问题。

一个小技巧：当你不知道访问模式对不对时，可以算一下理论内存事务数。公式很简单：理论事务数 = (总数据量) / (缓存行大小)。如果实际事务数超过理论值的2倍，就该优化了。

好了，这一章的内容就到这里。内存访问模式是GPU优化的基本功，说白了就是让硬件干得舒服。你顺着硬件的脾气来，它就会给你好性能；你非要跟它对着干，那性能就只能看天了。

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