内存合并与向量化内存访问
Memory Coalescing vs. Vectorized Memory Access
我正在尝试了解 NVIDIA GPUs/CUDA 上的 内存合并 和 [=43 上的 向量化内存访问 之间的关系=]++。
据我了解：
- 内存合并是内存控制器的运行时间优化（在硬件中实现）。在 运行 时间确定需要多少内存事务来完成扭曲的 load/store。一个 load/store 的 warp 指令可能是 issued repeatedly 除非有完美的合并。
- 内存矢量化是一种编译时优化。矢量化 load/store 的内存事务数是固定的。每个向量 load/store 指令恰好发出一次。
- 可合并 GPU load/store 指令比 SSE 矢量 load/store 指令更具表现力。例如，
st.global.s32 PTX 指令可以存储到 32 个任意内存位置（warp 大小 32），而 movdqa SSE 指令只能存储到连续的内存块中。
- CUDA 中的内存合并似乎可以保证高效的向量化 内存访问（当访问可合并时），而在 x86-SSE 上，我们不得不希望编译器实际上对代码进行向量化（它可能会失败）或使用 SSE 内在函数手动矢量化代码，这对程序员来说更难。
这是正确的吗？我是否错过了一个重要的方面（线程屏蔽，也许）？
现在，为什么 GPU 有 运行 时间合并？这可能需要额外的硬件电路。与 CPU 中的编译时合并相比，主要优势是什么？是否有 applications/memory 访问模式由于缺少 运行 时间合并而更难在 CPU 上实现？
警告：我不太了解/理解 GPU 的架构/微架构。其中一些理解是从问题+其他人在此处的评论/答案中写的内容拼凑而成的。
GPU 让一条指令对多个数据进行操作的方式与 CPU SIMD 非常 不同。这就是为什么他们需要对内存合并的特殊支持。 CPU-SIMD 无法以需要它的方式进行编程。
顺便说一句，CPUs 具有缓存以在实际 DRAM 控制器介入之前吸收对同一缓存行的多次访问。 GPU当然也有缓存。
是的，memory-coalescing 基本上在 运行 时间做 short-vector CPU SIMD 在编译时做的事情，在单个 "core" 内。 CPU-SIMD 等效项将是 gather/scatter loads/stores，它可以优化对相邻索引的单个宽缓存访问。 现有 CPUs 不要这样做：每个元素在一个集合中分别访问缓存。如果您知道许多索引将相邻，则不应使用收集加载；将 128 位或 256 位块洗牌到位会更快。对于所有数据都是连续的常见情况，您只需使用普通向量加载指令而不是收集加载。
现代 short-vector CPU SIMD 的重点是通过 fetch/decode/exec 管道 提供更多工作，而无需 它在每个时钟周期必须解码 + 跟踪 + 执行更多 CPU 指令方面更宽。 使 CPU 管道更宽快速减少 returns 对于大多数 [=92] =], 因为大多数代码没有很多 ILP。
A general-purpose CPU 在 instruction-scheduling / out-of-order 执行机制上花费了很多晶体管，所以只需使其更宽即可 运行并行更多的 uops 是不可行的。 (https://electronics.stackexchange.com/questions/443186/why-not-make-one-big-cpu-core).
为了获得更多的吞吐量，我们可以提高频率，提高 IPC，并使用 SIMD 为每个 out-of-order 机器必须跟踪的 instruction/uop 做更多的工作。 （而且我们可以在单个芯片上构建多个内核，但是 cache-coherent 它们之间的互连 + L3 缓存 + 内存控制器很难）。现代 CPUs 使用了所有这些东西，所以我们得到频率 * IPC * SIMD 的总吞吐量能力，如果我们多线程，则乘以核心数。它们不是彼此可行的替代品，它们是正交的事情，您必须 所有 来驱动大量 FLOP 或通过 CPU 管道进行整数工作。
这就是为什么 CPU SIMD 具有广泛的 fixed-width 执行单元，而不是每个标量操作的单独指令。没有一种机制可以将一条标量指令灵活地馈送到多个执行单元。
利用这一点需要在编译时进行矢量化，不仅是您的加载/存储，还有您的 ALU 计算。如果您的数据不连续，则必须使用标量加载+随机播放或使用 AVX2/AVX512 收集采用基地址+（缩放）索引向量的加载将其收集到 SIMD 向量中。
但 GPU SIMD 不同。它适用于 massively 并行问题，您对每个元素执行相同的操作。 "pipeline" 可以非常轻量级，因为它不需要支持 out-of-order 执行或寄存器重命名，尤其是分支和异常。这使得仅具有标量执行单元而无需处理来自连续地址的固定块中的数据变得可行。
这是两种截然不同的编程模型。它们都是 SIMD，但是运行它们的硬件细节非常不同。
Each vector load/store instruction is issued exactly once.
是的，这在逻辑上是正确的。实际上，内部结构可能稍微复杂一些，例如AMD Ryzen 将 256 位矢量运算拆分为 128 位的一半，或者英特尔 Sandybridge/IvB 在具有 256 位宽 FP ALU 的情况下仅针对加载+存储执行此操作。
在 Intel x86 CPUs 上 loads/stores 有轻微的错位：在 cache-line 拆分中，uop 必须重播（从保留站）以执行访问的其他部分（到其他缓存行）。
在 Intel 术语中，拆分负载的 uop 分派 两次，但只发布 + 退役一次。
Aligned loads/stores like movdqa, 或者 movdqu 当内存恰好在 运行 时对齐，都只是对L1d缓存的一次访问（假设一个缓存命中）。除非你在 CPU 上将向量指令解码为两半，例如 AMD 的 256 位向量。
但是那东西纯粹在 CPU 核心内部，用于访问 L1d 缓存。 CPU <-> 内存事务在整个缓存行中，具有 write-back L1d / L2 私有缓存，并在现代 x86 CPUs 上共享 L3 - （Intel 自 Nehalem 以来，i3/i5/i7 系列的开始，AMD 自 Bulldozer 以来我认为为它们引入了 L3 缓存。）
在 CPU 中，write-back L1d 缓存基本上将事务合并为一个整体缓存行，无论您是否使用 SIMD。
SIMD 的帮助是在 CPU 中完成更多工作，以跟上更快的内存。或者对于数据适合 L2 或 L1d 缓存的问题，真正 快速处理该数据。
内存合并与parallel访问有关：当SM中的每个核心将访问后续内存位置时，内存访问被优化。
Viceversa，SIMD 是单核优化：当向量寄存器充满操作数并执行 SSE 操作时，并行性在 CPU 核内，对每个内部逻辑执行一个操作每个时钟周期的单位。
但是你是对的：coalesced/uncoalesced内存访问是一个运行时方面。 SIMD运算都编译进去了，我觉得没法比。
如果要进行并行处理，我会将 GPU 中的合并与 CPU 中的内存 预取 进行比较。这也是一个非常重要的运行时优化 - 我相信它在使用 SSE 的幕后也是活跃的。
然而，在 Intel CPU 内核中没有类似于合并的东西。由于高速缓存一致性，在优化并行内存访问方面，您可以做的最好的事情就是让每个内核访问独立的内存区域。
Now, why do GPUs have run-time coalescing?
图形处理针对在相邻元素上并行执行单个任务进行了优化。
例如，考虑对图像的每个像素执行操作，将每个像素分配给不同的核心。现在很明显，您希望有一个最佳路径来加载将一个像素传播到每个核心的图像。
这就是内存合并深深埋藏在 GPU 架构中的原因。
我正在尝试了解 NVIDIA GPUs/CUDA 上的 内存合并 和 [=43 上的 向量化内存访问 之间的关系=]++。
据我了解：
- 内存合并是内存控制器的运行时间优化（在硬件中实现）。在 运行 时间确定需要多少内存事务来完成扭曲的 load/store。一个 load/store 的 warp 指令可能是 issued repeatedly 除非有完美的合并。
- 内存矢量化是一种编译时优化。矢量化 load/store 的内存事务数是固定的。每个向量 load/store 指令恰好发出一次。
- 可合并 GPU load/store 指令比 SSE 矢量 load/store 指令更具表现力。例如，
st.global.s32PTX 指令可以存储到 32 个任意内存位置（warp 大小 32），而movdqaSSE 指令只能存储到连续的内存块中。 - CUDA 中的内存合并似乎可以保证高效的向量化 内存访问（当访问可合并时），而在 x86-SSE 上，我们不得不希望编译器实际上对代码进行向量化（它可能会失败）或使用 SSE 内在函数手动矢量化代码，这对程序员来说更难。
这是正确的吗？我是否错过了一个重要的方面（线程屏蔽，也许）？
现在，为什么 GPU 有 运行 时间合并？这可能需要额外的硬件电路。与 CPU 中的编译时合并相比，主要优势是什么？是否有 applications/memory 访问模式由于缺少 运行 时间合并而更难在 CPU 上实现？
警告：我不太了解/理解 GPU 的架构/微架构。其中一些理解是从问题+其他人在此处的评论/答案中写的内容拼凑而成的。
GPU 让一条指令对多个数据进行操作的方式与 CPU SIMD 非常 不同。这就是为什么他们需要对内存合并的特殊支持。 CPU-SIMD 无法以需要它的方式进行编程。
顺便说一句，CPUs 具有缓存以在实际 DRAM 控制器介入之前吸收对同一缓存行的多次访问。 GPU当然也有缓存。
是的，memory-coalescing 基本上在 运行 时间做 short-vector CPU SIMD 在编译时做的事情，在单个 "core" 内。 CPU-SIMD 等效项将是 gather/scatter loads/stores，它可以优化对相邻索引的单个宽缓存访问。 现有 CPUs 不要这样做：每个元素在一个集合中分别访问缓存。如果您知道许多索引将相邻，则不应使用收集加载；将 128 位或 256 位块洗牌到位会更快。对于所有数据都是连续的常见情况，您只需使用普通向量加载指令而不是收集加载。
现代 short-vector CPU SIMD 的重点是通过 fetch/decode/exec 管道 提供更多工作，而无需 它在每个时钟周期必须解码 + 跟踪 + 执行更多 CPU 指令方面更宽。 使 CPU 管道更宽快速减少 returns 对于大多数 [=92] =], 因为大多数代码没有很多 ILP。
A general-purpose CPU 在 instruction-scheduling / out-of-order 执行机制上花费了很多晶体管，所以只需使其更宽即可 运行并行更多的 uops 是不可行的。 (https://electronics.stackexchange.com/questions/443186/why-not-make-one-big-cpu-core).
为了获得更多的吞吐量，我们可以提高频率，提高 IPC，并使用 SIMD 为每个 out-of-order 机器必须跟踪的 instruction/uop 做更多的工作。 （而且我们可以在单个芯片上构建多个内核，但是 cache-coherent 它们之间的互连 + L3 缓存 + 内存控制器很难）。现代 CPUs 使用了所有这些东西，所以我们得到频率 * IPC * SIMD 的总吞吐量能力，如果我们多线程，则乘以核心数。它们不是彼此可行的替代品，它们是正交的事情，您必须 所有 来驱动大量 FLOP 或通过 CPU 管道进行整数工作。
这就是为什么 CPU SIMD 具有广泛的 fixed-width 执行单元，而不是每个标量操作的单独指令。没有一种机制可以将一条标量指令灵活地馈送到多个执行单元。
利用这一点需要在编译时进行矢量化，不仅是您的加载/存储，还有您的 ALU 计算。如果您的数据不连续，则必须使用标量加载+随机播放或使用 AVX2/AVX512 收集采用基地址+（缩放）索引向量的加载将其收集到 SIMD 向量中。
但 GPU SIMD 不同。它适用于 massively 并行问题，您对每个元素执行相同的操作。 "pipeline" 可以非常轻量级，因为它不需要支持 out-of-order 执行或寄存器重命名，尤其是分支和异常。这使得仅具有标量执行单元而无需处理来自连续地址的固定块中的数据变得可行。
这是两种截然不同的编程模型。它们都是 SIMD，但是运行它们的硬件细节非常不同。
Each vector load/store instruction is issued exactly once.
是的，这在逻辑上是正确的。实际上，内部结构可能稍微复杂一些，例如AMD Ryzen 将 256 位矢量运算拆分为 128 位的一半，或者英特尔 Sandybridge/IvB 在具有 256 位宽 FP ALU 的情况下仅针对加载+存储执行此操作。
在 Intel x86 CPUs 上 loads/stores 有轻微的错位：在 cache-line 拆分中，uop 必须重播（从保留站）以执行访问的其他部分（到其他缓存行）。
在 Intel 术语中，拆分负载的 uop 分派 两次，但只发布 + 退役一次。
Aligned loads/stores like movdqa, 或者 movdqu 当内存恰好在 运行 时对齐，都只是对L1d缓存的一次访问（假设一个缓存命中）。除非你在 CPU 上将向量指令解码为两半，例如 AMD 的 256 位向量。
但是那东西纯粹在 CPU 核心内部，用于访问 L1d 缓存。 CPU <-> 内存事务在整个缓存行中，具有 write-back L1d / L2 私有缓存，并在现代 x86 CPUs 上共享 L3 -
在 CPU 中，write-back L1d 缓存基本上将事务合并为一个整体缓存行，无论您是否使用 SIMD。
SIMD 的帮助是在 CPU 中完成更多工作，以跟上更快的内存。或者对于数据适合 L2 或 L1d 缓存的问题，真正 快速处理该数据。
内存合并与parallel访问有关：当SM中的每个核心将访问后续内存位置时，内存访问被优化。
Viceversa，SIMD 是单核优化：当向量寄存器充满操作数并执行 SSE 操作时，并行性在 CPU 核内，对每个内部逻辑执行一个操作每个时钟周期的单位。
但是你是对的：coalesced/uncoalesced内存访问是一个运行时方面。 SIMD运算都编译进去了，我觉得没法比。
如果要进行并行处理，我会将 GPU 中的合并与 CPU 中的内存 预取 进行比较。这也是一个非常重要的运行时优化 - 我相信它在使用 SSE 的幕后也是活跃的。
然而，在 Intel CPU 内核中没有类似于合并的东西。由于高速缓存一致性，在优化并行内存访问方面，您可以做的最好的事情就是让每个内核访问独立的内存区域。
Now, why do GPUs have run-time coalescing?
图形处理针对在相邻元素上并行执行单个任务进行了优化。
例如，考虑对图像的每个像素执行操作，将每个像素分配给不同的核心。现在很明显，您希望有一个最佳路径来加载将一个像素传播到每个核心的图像。
这就是内存合并深深埋藏在 GPU 架构中的原因。