访存合并:带宽的开关

Lesson 03 · CUDA 性能调优 · 访存瓶颈的头号武器

你已经能判断一个 kernel 是访存瓶颈了。现在的问题是:同样搬一样多的数据,为什么有的 kernel 带宽利用 90%,有的只有 15%? 最常见的答案是访存合并(coalescing)。这一节,你将第一次真正动手把一个 kernel 改快。

本节目标(一个可带走的胜利): 你能看出一个 kernel 的访存是否合并, 理解 warp 与 128 字节事务的关系,并知道如何把"跨步访问"改成"合并访问"。

1. 关键事实:GPU 按 warp、按整块搬数据

GPU 不是一个线程一个线程地读显存。它以 warp(32 个线程) 为单位执行, 而显存访问以 固定大小的事务(transaction) 为单位——通常是 128 字节对齐的一段。[1]

当一个 warp 的 32 个线程同时读内存时,硬件会看它们的地址:

如果 32 个地址恰好落在连续的一段(线程 i 读第 i 个 float)—— 硬件用 少数几个 128B 事务就全搬回来了。这叫 合并访问。
如果 32 个地址东一个西一个——硬件被迫发很多个事务, 每个只用到一点点,大量带宽浪费在没用的字节上。这叫 非合并访问。

2. 看图:同一个 warp,两种命运

设每个线程读一个 4 字节 float,一个 warp 32 线程,一次能用一个 128B 事务全部覆盖(32×4=128)。

✅ 合并:thread i 读 data[i]

addr: 04812…124
→ 32 线程落在同一个 128B 段 → 1 个事务,100% 字节有用

❌ 非合并:thread i 读 data[i * stride]

事务1 t0··· 事务2 t1··· 事务3 t2·…
→ 每个线程独占一个事务 → 最多 32 个事务,大部分字节白搬

同样是读 32 个数,合并访问 1 个事务搞定,非合并要 32 个——带宽有效利用率可能差 10 倍以上。这就是为什么访存瓶颈的 kernel,第一件事就是检查合并。[1]

2.5 「事务」到底是什么

合并的一切都绕着"事务"这个词。说透它:GPU 读显存不是一个字节一个字节读,而是一整块一整块读。那一块就是一次事务。

打个比方:去仓库取货,规定每次必须开一整个货箱,哪怕你只要箱里一颗螺丝。货箱 = 事务。A100 上货箱固定 32 字节(一个 sector),硬件常把 4 个凑成 128 字节一起处理。

为什么这样设计?显存(HBM)的物理特性:按大块读极快,按零散小块读极慢——和机械硬盘顺序读快、随机读慢同理。所以硬件干脆规定:不管你要几个字节,都按固定大小的块搬。

关键就在一句:一次事务搬回 128 字节,但其中有多少是你真正要用的?

情况	开几个货箱	有效利用
合并:32 线程的数据挤在一个 128B 块里	1 次事务	128/128 = 100% ✅
非合并:32 线程数据分散(stride)	最多 32 次事务	每箱只用 4B,4/128 ≈ 3% ❌

同样取 128 字节有用数据,合并开 1 个箱,非合并开 32 个箱。带宽是固定的"开箱速度", 非合并白白浪费 31/32 的箱子——这就是"差 10 倍以上"的来源。

把报告里的指标对上号: request(请求)= 一个 warp 的一条访存指令(它要的那 128B); sector = 32B 的货箱。
完全合并 → 128÷32 = 4 sectors per request(理想); 完全分散 → 每线程拖一个 = 32 sectors per request(最糟)。所以这个数越接近 4 越好,越大越不合并。

3. 代码:一个经典的"踩坑→修复"

最常见的非合并来自按行处理矩阵。假设一个 M×N 矩阵按行主序存放,让每个线程处理一整列求和:

❌ 慢:每个线程沿列走,相邻线程地址相隔 N

// thread t 处理第 t 列;一个 warp 的 32 个线程
// 同一时刻访问 data[0*N+t], data[0*N+t+1]... 看似连续?
// 但下一步它们一起跳到 row=1: 地址各 +N → 跨步!
int col = blockIdx.x * blockDim.x + threadIdx.x;
float sum = 0;
for (int row = 0; row < N; row++)
    sum += data[row * N + col];   // 每轮 warp 内连续 ✓,但…见下

其实这一版反而是合并的!同一轮里,warp 的线程 t 读 row*N + t, 地址连续。真正的坑是反过来写:

❌ 真正的坑:每个线程处理一行

int row = blockIdx.x * blockDim.x + threadIdx.x;
float sum = 0;
for (int c = 0; c < N; c++)
    sum += data[row * N + c];     // warp 内线程地址相隔 N → 非合并!

这里同一轮里线程 t 读 t*N + c,相邻线程地址差 N(一整行),完全分散。修复方法是让相邻线程访问相邻地址:换数据布局(转置)、或换线程到数据的映射方式。

判断口诀: 盯住 warp 内相邻线程(threadIdx.x 相差 1)在同一条指令里访问的地址。 地址相差 1 个元素 → 合并;相差一个大 stride → 非合并。索引里 threadIdx.x 应当乘在最低维(变化最快的维)上。

4. 在报告里怎么看出来

非合并访问在 Nsight Compute 里有明确信号。除了 SOL 显示访存瓶颈外, Memory Workload Analysis 段会出现:

[Warning] Uncoalesced global accesses. The memory access
pattern has low efficiency. X sectors per request 。

看 sectors per request(上一节已解释:理想≈4,越大越不合并)。如果是 16、32,就说明严重不合并——每个请求拖回一堆没用的 sector。[1]

4.5 延伸一:对齐(alignment)

"合并的 warp 开 1 个事务"有个隐藏前提:起始地址要对齐到事务边界。 128B 事务只能从 0、128、256… 这些 128 的倍数开始——把显存想成划好格子的货架, 事务只能整格取,不能跨格错位取。

于是即使数据完全连续、索引完美,只要起点偏了,事务数也会翻倍:

warp 要的 128B 区间	跨几格	事务数
`[0, 128)` 对齐	1 格	1 ✅
`[32, 160)` 偏了 32B	横跨 2 格	2 ❌(各有浪费)

最常见来源:非零偏移,如 kernel(data + 1, ...),或结构体字段没对齐。
CUDA 的保护:cudaMalloc 返回的指针保证至少 256B 对齐,从数组开头规矩访问通常没事。坑出在手动加偏移、自管内存时。
修复:用对齐分配、避免奇怪偏移;向量类型 float4(16B)天然带对齐;必要时 __align__(128)。[1]

合并要求"相邻线程读相邻地址";对齐额外要求"这一段的起点踩在事务边界上"。两个都满足,才真正 1 个事务。

4.6 延伸二:L1 / L2 缓存如何补救非合并

前面把模型简化成"warp 直接读显存"。真相是中间隔着两层片上缓存:

warp 访存请求
   ↓
L1(每 SM 一个,A100 上与 shared memory 共用 SRAM)
   ↓ miss
L2(全 GPU 共享,A100 有 40MB,很大)
   ↓ miss
HBM 显存(真正的"远路")

关键:事务发生在缓存这一层,不一定直达 HBM。两个后果:

后果 1:L2 能摊销部分非合并的代价

非合并、地址分散,但数据若之前被读过、还在 L2 里,这次就不必走 HBM。

非合并 + 缓存命中 → 没那么惨,L2 带宽远高于 HBM。
非合并 + 缓存未命中 → 全额惩罚,每个分散访问拖一趟 HBM,最慢。

这解释了一个困惑:同样非合并的代码,为何有时慢得要死、有时还能忍? 取决于数据是否恰好被缓存住。A100 的 40MB L2 常能兜住一部分。

后果 2:这正是报告分两层指标的原因

回想第 2 节那两行,现在能读出差异含义:

Memory %	DRAM %	含义
高	高	真在猛读显存,缓存没兜住 → 典型访存瓶颈
高	低	流量多被 L2 接住 → 瓶颈在缓存/数据复用,不在 HBM

这是判断"该优化访存模式,还是该优化数据复用(下一课 shared memory)"的关键线索。

别依赖缓存救场。 数据在不在缓存里是不可控的运气,取决于运行时历史。正确做法仍是从一开始就写合并访问——缓存是安全网,不是设计依据。数据集变大、缓存装不下时,寄希望于"L2 会兜住"一定会崩。

5. 练习:这段访存合并吗?

每题盯住 warp 内相邻线程(threadIdx.x 差 1)的地址差。点选项看反馈。

💬 随时问我。 为什么是 128 字节?二维 blockIdx 怎么算合并?转置具体怎么改? 把 kernel 贴给我,我帮你逐行看访存模式。我是你的老师。

主源推荐(本节精读)

📘 CUDA C++ Best Practices Guide — Coalesced Access to Global Memory ——这一小节用图解讲透了合并规则,是本节最高信任的来源,务必读一遍。