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

文章目录

摘要

• 1. 为什么要学习 Memory Coalescing？
• 2. 什么是 Memory Coalescing？
  • 2.1 合并访问：Coalesced Access
  • 2.2 非合并访问：Uncoalesced Access
• 3. 实验代码设计
  • 3.1 Coalesced kernel
  • 3.2 Uncoalesced kernel
• 4. 实验设置
• 5. 实验结果
  • 5.1 原始结果
• 6. 实验结果解读
  • 6.1 Coalesced 访问为什么快？
  • 6.2 Uncoalesced 访问为什么慢？
• 7. 为什么 stride 越大，slowdown 越明显？
• 8. 为什么 coalesced 时间也会变化？
• 9. 有效带宽估算
  • 9.1 Coalesced 有效带宽
  • 9.2 Uncoalesced 有效带宽
• 10. 和矩阵、图像、深度学习有什么关系？
  • 10.1 矩阵访问
  • 10.2 图像处理
  • 10.3 深度学习 Tensor Layout
• 11. 实践优化建议
  • 11.1 让 `threadIdx.x` 对应连续内存
  • 11.2 二维数组优先保证行方向连续
  • 11.3 对非连续访问进行重排
• 本课结论

摘要

第 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__ void coalesced_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__ void uncoalesced_access(float* data, size_t ops, int stride) {
    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 执行时间

代码中：

const size_t n = 1ULL << 30;
const size_t bytes = n * sizeof(float);
const size_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 优化中，访存模式往往比算术指令本身更决定性能。