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

Question

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

我的理解是：

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

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

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

Answer 1

警告：我不太了解/理解 GPU 的架构/微架构。这种理解的一部分是从问题+其他人在此处的评论/答案中所写的内容拼凑而成的。

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

顺便说一句，在实际的 DRAM 控制器介入之前，CPU 有缓存来吸收对同一缓存线的多次访问。当然，GPU 也有缓存。

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

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

通用 CPU 在指令调度/乱序执行机制上花费了大量晶体管，因此仅仅让它更宽以能够并行运行更多的 uops 是不可行的。（https://electronics.stackexchange.com/questions/443186/why-not-make-one-big-cpu-core）。

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

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

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

但 GPU SIMD 不同。它适用于大规模并行问题，您对每个元素都做同样的事情。“管道”可以非常轻量级，因为它不需要支持乱序执行或寄存器重命名，或者特别是分支和异常。这使得仅具有标量执行单元而无需处理来自连续地址的固定块中的数据变得可行。

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

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

是的，这在逻辑上是正确的。在实践中，内部结构可能稍微复杂一些，例如 AMD Ryzen 将 256 位向量操作拆分为 128 位一半，或者英特尔 Sandybridge/IvB 在具有 256 位宽 FP ALU 的情况下仅用于加载 + 存储。

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

在英特尔术语中，拆分负载的 uop 被分派两次，但只发出一次 + 退出。

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

但这些东西纯粹是在 CPU 内核内部，用于访问 L1d 缓存。 CPU <-> 内存事务在整个缓存行中，在现代 x86 CPU 上具有回写 L1d / L2 私有缓存和共享 L3 -英特尔酷睿 i7 处理器中使用了哪种缓存映射技术？ （英特尔自 Nehalem 开始，i3/i5/i7 系列的开始，AMD 自推土机以来我认为为他们引入了 L3 缓存。）

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

SIMD 的帮助是在 CPU 内部完成更多工作，以跟上更快的内存。或者对于数据适合 L2 或 L1d 缓存的问题，可以非常快速地处理该数据。