Global 内存访问与 Memory Coalescing 实验解析

摘要

第 7 课围绕 CUDA 中非常关键的Memory Coalescing，即内存合并访问展开。

实验通过对比连续访问data[idx]和跨步访问data[idx * stride]，验证了 Global 内存访问模式对性能的巨大影响。

结果显示，在 Tesla T4 上，当stride=32时，非合并访问比合并访问慢约25.5 倍；
当stride=16时慢约14.6 倍；
当stride=8时慢约8.0 倍。
这说明 GPU 性能不仅取决于“算多少”，更取决于 warp 内线程是否访问连续地址。

1. 为什么要学习 Memory Coalescing？

在前几课中，我们已经学习了：

Pinned Memory：优化 CPU ↔ GPU 数据传输 Shared Memory：优化 GPU kernel 内部数据复用 CUDA Stream：让传输与计算重叠

第 7 课关注的是另一个核心问题：

同样访问 Global 内存，为什么有的访问方式很快，有的访问方式会慢几十倍？

GPU 的 Global 内存访问不是以单个线程为单位孤立执行，而是以warp为基本调度单位。

一个 warp 通常有：

32 个线程

如果这 32 个线程访问的是连续地址，GPU 可以把这些访问合并成较少的内存事务。

如果这 32 个线程访问的是分散地址，GPU 就需要发起更多内存事务，带宽利用率会急剧下降。

2. 什么是 Memory Coalescing？

Memory Coalescing 可以理解为：

warp 内线程访问连续内存地址时，GPU 将多个线程的访存请求合并成较少的内存事务，从而提高 Global 内存带宽利用率。

2.1 合并访问：Coalesced Access

例如：

thread 0 -> data[0] thread 1 -> data[1] thread 2 -> data[2] ... thread 31 -> data[31]

这些地址是连续的。

对于float来说，每个元素 4 字节，32 个线程访问的数据正好是连续的一段：

32 × 4 bytes = 128 bytes

这类访问非常适合 GPU 内存系统合并处理。

2.2 非合并访问：Uncoalesced Access

如果访问模式变成：

thread 0 -> data[0] thread 1 -> data[32] thread 2 -> data[64] thread 3 -> data[96] ... thread 31 -> data[992]

这时线程之间访问地址相隔很远。

对于stride=32：

相邻线程地址间隔 = 32 × 4 bytes = 128 bytes

也就是说，一个 warp 内的每个线程几乎都落在不同的内存区域，GPU 很难把它们合并成少数内存事务。

结果就是：

访问请求变多 带宽利用率下降 kernel 时间显著增加

3. 实验代码设计

本次实验设计两个 kernel。

3.1 Coalesced kernel

__global__voidcoalesced_access(float*data,size_t ops){size_t idx=(size_t)blockIdx.x*blockDim.x+threadIdx.x;if(idx<ops){data[idx]+=1.0f;}}

这个 kernel 的访问模式是：

data[0], data[1], data[2], ...

也就是连续访问。

3.2 Uncoalesced kernel

__global__voiduncoalesced_access(float*data,size_t ops,intstride){size_t idx=(size_t)blockIdx.x*blockDim.x+threadIdx.x;if(idx<ops){size_t access_idx=idx*(size_t)stride;data[access_idx]+=1.0f;}}

这个 kernel 的访问模式是：

data[0], data[stride], data[2 * stride], ...

也就是跨步访问。

4. 实验设置

实验设置如下：

Array size: 4096 MB block size: 256 threads 测试 stride: 8、16、32 计时方式: cudaEvent 统计范围: kernel 执行时间

代码中：

constsize_t n=1ULL<<30;constsize_t bytes=n*sizeof(float);constsize_t ops=n/stride;

这里的设计很重要：

ops = n / stride

这样可以保证 uncoalesced kernel 不越界。

同时，在同一组 stride 下：

coalesced kernel 和 uncoalesced kernel 执行相同数量的线程操作

区别只在于访问模式不同。

5. 实验结果

5.1 原始结果

这个结果符合课程预期：

stride 越大 warp 内线程访问越分散 memory coalescing 越差 性能下降越明显

6. 实验结果解读

6.1 Coalesced 访问为什么快？

以stride=32那组为例，coalesced kernel 访问的是：

data[0], data[1], data[2], ...

warp 内线程访问连续地址，GPU 可以高效合并访问。

所以它只用了：

1.09542 ms

这说明 GPU 对连续 Global 内存访问的带宽利用率很高。

6.2 Uncoalesced 访问为什么慢？

同样是stride=32，uncoalesced kernel 访问的是：

data[0], data[32], data[64], data[96], ...

一个 warp 内相邻线程间隔：

32 × 4 bytes = 128 bytes

这会导致：

• 原本可以合并的访问被拆散
• 内存事务数量显著增加
• cache line / memory sector 利用率下降（GPU 为了满足分散的线程访问，不得不搬来一整段内存数据，但每段里只有很少几个字节真正被线程使用；结果是内存事务变多、有效带宽下降、kernel 变慢）
• 实际带宽被浪费

所以它耗时：

27.9479 ms

相比 coalesced 慢了：

25.5133x

这是非常显著的性能差距。

7. 为什么 stride 越大，slowdown 越明显？

slowdown 是：

这个趋势很清晰：

stride=8 ：线程间地址间隔 32 bytes stride=16 ：线程间地址间隔 64 bytes stride=32 ：线程间地址间隔 128 bytes

stride 越大，warp 内线程访问越分散。

因此：

合并访问能力越差 内存事务越多 带宽浪费越严重 性能越差

所以 slowdown 从约 8 倍逐渐扩大到约 25 倍。

8. 为什么 coalesced 时间也会变化？

coalesced 的时间分别是：

你会发现：

ops 翻倍，coalesced time 也大致翻倍

原因是：

ops=n/stride;

所以 stride 越小，实际访问次数越多。

这说明 coalesced kernel 的行为比较稳定，主要受有效访问次数控制。

9. 有效带宽估算

因为：

data[idx]+=1.0f;

大致可以看作：

一次读 一次写

所以可粗略估算有效数据量：

useful_bytes = ops × sizeof(float) × 2

有效带宽：

bandwidth = useful_bytes / time

9.1 Coalesced 有效带宽

Coalesced 访问的有效带宽比较稳定，说明连续访问能充分利用 Global 内存带宽。

9.2 Uncoalesced 有效带宽

从有效带宽看，uncoalesced 的带宽利用率明显低得多。

尤其是stride=32时，有效带宽只有约：

9.6 GB/s

而 coalesced 约为：

245 GB/s

两者差距非常大。

10. 和矩阵、图像、深度学习有什么关系？

Memory Coalescing 在很多场景中都非常重要。

10.1 矩阵访问

如果矩阵按行主序存储：

A[row*N+col]

那么同一个 warp 中线程访问连续col通常比较快：

A[row][0], A[row][1], A[row][2], ...

但如果按列访问：

A[0][col], A[1][col], A[2][col], ...

在行主序内存中就是跨步访问，可能变慢。

10.2 图像处理

图像通常也是二维数组。

如果线程按行连续处理像素：

pixel[y][x], pixel[y][x+1], pixel[y][x+2]

访问通常较好。

如果线程跨列、跨通道、跨行访问，可能产生非合并访问。

10.3 深度学习 Tensor Layout

深度学习中常见 layout：

NCHW NHWC

不同 layout 会影响某些 kernel 中线程访问是否连续。

所以在高性能推理框架中，layout 选择和 memory coalescing 密切相关。

11. 实践优化建议

最重要的工程建议是：

11.1 让threadIdx.x对应连续内存

在 CUDA kernel 中，尽量让：

threadIdx.x = 0,1,2,3...

对应访问：

data[base + 0] data[base + 1] data[base + 2] data[base + 3]

而不是：

data[base + 0 * stride] data[base + 1 * stride] data[base + 2 * stride]

11.2 二维数组优先保证行方向连续

如果数据是 C/C++ 行主序：

data[row*width+col]

那么通常应该让：

threadIdx.x 对应 col

这样同一个 warp 的线程更容易访问连续地址。

11.3 对非连续访问进行重排

如果算法天然需要跨步访问，可以考虑：

1. 改变数据布局 2. 使用 shared memory 做 tile 重排 3. 让读取阶段 coalesced，计算阶段在 shared memory 中调整访问模式 4. 合并多个小访问，减少随机访问

矩阵转置就是经典例子：

Global 内存读写尽量 coalesced 中间用 Shared Memory 处理转置

本课结论

CUDA kernel 的性能不仅取决于计算量，也强烈依赖 Global 内存访问模式。

实验结果显示：

stride=8 时，uncoalesced 约慢 8 倍 stride=16 时，uncoalesced 约慢 14.6 倍 stride=32 时，uncoalesced 约慢 25.5 倍

这说明：

warp 内线程访问连续地址 → 内存事务少 → 带宽利用率高 → kernel 快 warp 内线程访问分散地址 → 内存事务多 → 带宽利用率低 → kernel 慢

一句话总结：

Memory Coalescing 的本质，是让一个 warp 内的线程尽量访问连续 Global 内存地址，从而减少内存事务、提高带宽利用率；在 CUDA 优化中，访存模式往往比算术指令本身更决定性能。