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

Question

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

据我了解：


内存合并是内存控制器的运行时优化（在硬件中实现）。在运行时确定需要多少个内存事务来完成扭曲的加载/存储。除非完全合并，否则经纱的加载/存储指令可能为issued repeatedly。
内存向量化是编译时的优化。矢量化加载/存储的内存事务数是固定的。每个向量加载/存储指令仅发出一次。
可合并的GPU加载/存储指令比SSE矢量加载/存储指令更具表现力。例如，一个st.global.s32 PTX指令可以存储到32个任意的存储位置（warp大小为32），而movdqa SSE指令只能存储到一个连续的存储器块中。
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在指令调度/乱序执行机制上花费了很多晶体管，因此仅使其扩展以能够并行运行更多微指令是不可行的。 （https://electronics.stackexchange.com/questions/443186/why-not-make-one-big-cpu-core）。

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

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

要利用此优势，不仅需要在加载/存储时进行矢量化，还需要在ALU计算中进行矢量化。如果您的数据不连续，则必须使用标量负载+混洗或使用AVX2 / AVX512将数据收集到SIMD向量中，这些负载采用基地址+（比例）索引的向量。



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

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




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


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

在Intel x86 CPU上，未对齐的加载/存储会有一些细微的折皱：在高速缓存行拆分中，必须重播uop（从预留站）以进行访问的另一部分（到另一高速缓存行）。

用Intel的术语来说，分担负载的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?（自Nehalem以来为Intel，i3 / i5 / i7系列的开始，AMD自推土机以来，我认为为他们引入了L3缓存。）

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

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