4

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

我的理解是：

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

它是否正确？我是否错过了一个重要方面（也许是线程屏蔽）？

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

4

2 回答 2

5

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

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 的情况下进行加载+存储。

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

在英特尔术语中，用于拆分负载的 uop 被分派了两次，但只有 issue + retires 一次。

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

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

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

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

于 2019-07-10T08:46:44.693 回答
0

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

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

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

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

但是，没有什么类似于 Intel CPU 内核中的合并。由于缓存一致性，在优化并行内存访问方面，您能做的最好的事情就是让每个核心访问独立的内存区域。

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

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

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

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

于 2019-07-10T09:53:18.347 回答