内存合并与向量化内存访问

Question

我试图了解 NVIDIA GPU/CUDA 上的内存合并与 x86-SSE/C++ 上的矢量化内存访问之间的关系。

据我了解：

• 内存合并是内存控制器的运行时优化（在硬件中实现）。完成warp的加载/存储需要多少内存事务是在运行时确定的。除非有完美的合并，否则扭曲的加载/存储指令可能是 issued repeatedly。
• 内存矢量化是一种编译时优化。矢量化加载/存储的内存事务数是固定的。每个向量加载/存储指令只发出一次。
• 可合并的 GPU 加载/存储指令比 SSE 向量加载/存储指令更具表现力。例如，st.global.s32 PTX 指令可以存储到 32 个任意内存位置（warp 大小为 32），而movdqa SSE 指令只能存储到连续的内存块中。
• CUDA 中的内存合并似乎保证了高效的矢量化内存访问（当访问可合并时），而在 x86-SSE 上，我们不得不希望编译器实际上对代码进行了矢量化（它可能无法这样做）或使用 SSE 内在函数手动矢量化代码，这对程序员来说更难。

这是正确的吗？我是否错过了一个重要方面（可能是线程屏蔽）？

现在，为什么 GPU 具有运行时合并功能？这可能需要额外的硬件电路。与 CPU 中的编译时合并相比，主要优势是什么？是否存在由于缺少运行时合并而更难在 CPU 上实现的应用程序/内存访问模式？

Answer 1

警告：我不太了解/了解 GPU 的架构/微架构。其中一些理解是从问题 + 其他人在 cmets / answers 中写的内容拼凑而成的。

GPU 让一条指令对多条数据进行操作的方式非常不同于 CPU SIMD。这就是为什么他们需要对内存合并的特殊支持。 CPU-SIMD 无法以需要的方式进行编程。

顺便说一句，在实际 DRAM 控制器参与之前，CPU 具有缓存以吸收对同一缓存行的多次访问。当然，GPU 也有缓存。

---

是的，内存合并基本上在运行时执行短向量 CPU SIMD 在编译时执行的操作，在单个“核心”内执行。 CPU-SIMD 等价物是收集/分散加载/存储，可以优化对相邻索引的单个广泛访问缓存。 现有 CPU 不这样做：每个元素分别访问缓存一次聚会。如果您知道许多索引将相邻，则不应使用收集负载；将 128 位或 256 位块随机放置到位会更快。对于所有数据都是连续的常见情况，您只需使用正常向量加载指令而不是聚集加载。

现代短向量 CPU SIMD 的重点是通过 fetch/decode/exec 管道提供更多工作而不使其在必须解码 + 跟踪 + exec 方面更宽每个时钟周期的 CPU 指令数。 使 CPU 流水线更宽会很快导致大多数用例的收益递减，因为大多数代码没有很多 ILP。

通用 CPU 在指令调度/乱序执行机制上花费了大量晶体管，因此仅将其扩大到能够并行运行更多微指令是不可行的。 (https://electronics.stackexchange.com/questions/443186/why-not-make-one-big-cpu-core)。

为了获得更多的吞吐量，我们可以提高频率、提高 IPC 并使用 SIMD 来完成乱序机器必须跟踪的每条指令/微指令的更多工作。 （我们可以在单个芯片上构建多个内核，但是它们之间的缓存一致互连 + L3 缓存 + 内存控制器很难）。现代 CPU 使用所有这些东西，因此我们获得了频率 * IPC * SIMD 的总吞吐量能力，如果我们使用多线程，则乘以内核数。它们不是彼此可行的替代品，它们是正交的事情，您必须全部通过 CPU 管道驱动大量 FLOP 或整数工作。

这就是为什么 CPU SIMD 具有宽的固定宽度执行单元，而不是为每个标量操作单独的指令。没有一种机制可以将一条标量指令灵活地馈送到多个执行单元。

利用这一点需要在编译时进行向量化，不仅是您的加载/存储，还包括您的 ALU 计算。如果您的数据不连续，则必须将其收集到 SIMD 向量中，或者使用标量加载 + 随机播放，或者使用 AVX2 / AVX512 收集采用基地址 + （缩放）索引向量的加载。

---

但 GPU SIMD 不同。它适用于 大量 并行问题，您对每个元素执行相同的操作。 “管道”可以非常轻量级，因为它不需要支持乱序执行或寄存器重命名，尤其是分支和异常。这使得只拥有标量执行单元而不需要处理来自连续地址的固定块中的数据变得可行。

这是两种截然不同的编程模型。它们都是 SIMD，但运行它们的硬件细节却大不相同。

---

每个向量加载/存储指令只发出一次。

是的，这在逻辑上是正确的。在实践中，内部结构可能会稍微复杂一些，例如AMD Ryzen 将 256 位向量操作分成 128 位两半，或者 Intel Sandybridge/IvB 仅用于加载和存储，同时具有 256 位宽 FP ALU。

在 Intel x86 CPU 上存在未对齐的加载/存储问题：在高速缓存行拆分时，必须（从保留站）重播 uop 以执行访问的另一部分（到另一高速缓存行） )。

在英特尔术语中，用于拆分负载的 uop 会调度两次，但只会发出一次 + 退出一次。

对齐的加载/存储，如movdqa 或movdqu（当内存恰好在运行时对齐时）只是对 L1d 缓存的一次访问（假设缓存命中）。除非您使用的 CPU 将向量指令解码为两半，例如 AMD 的 256 位向量。

---

但这些东西纯粹是在 CPU 内核内部，用于访问 L1d 缓存。 CPU 内存事务在整个高速缓存行中，具有回写式 L1d / L2 私有高速缓存，并在现代 x86 CPU 上共享 L3 - Which cache mapping technique is used in intel core i7 processor?（Intel 自 Nehalem，i3 的开始/i5/i7 系列，自 Bulldozer 以来的 AMD，我认为为它们引入了 L3 缓存。）

在 CPU 中，是回写式 L1d 缓存基本上将事务合并到整个缓存行中，无论您是否使用 SIMD。

SIMD 有助于在 CPU 内完成更多工作，以跟上更快的内存。或者对于数据适合 L2 或 L1d 缓存的问题，真正快速处理该数据。

Answer 2

内存合并与parallel 访问有关：当SM 中的每个内核将访问后续内存位置时，内存访问得到优化。

反之亦然，SIMD 是单核优化：当向量寄存器充满操作数并执行 SSE 操作时，并行性在 CPU 内核内部，每个时钟周期对每个内部逻辑单元执行一次操作。

但是您是对的：合并/取消合并的内存访问是运行时方面的。 SIMD 操作是编译进去的。我认为他们不能很好地比较。

如果我要进行并行处理，我会将 GPU 中的合并与 CPU 中的内存预取进行比较。这也是一个非常重要的运行时优化——我相信它在使用 SSE 的幕后也很活跃。

但是，在 Intel CPU 内核中没有任何类似于 colescing 的东西。由于缓存一致性，优化并行内存访问的最佳方法是让每个内核访问独立的内存区域。

现在，为什么 GPU 具有运行时合并功能？

图形处理针对在相邻元素上并行执行单个任务进行了优化。

例如，考虑对图像的每个像素执行操作，将每个像素分配给不同的核心。现在很明显，您希望有一个最佳路径来加载将一个像素散布到每个核心的图像。

这就是内存合并深埋在 GPU 架构中的原因。