全局写入中的内存合并答案

【问题标题】：Memory coalescing in global writes全局写入中的内存合并
【发布时间】：2019-07-09 22:04:46
【问题描述】：

在 CUDA 设备中，全局内存写入中的合并是否与全局内存读取中的合并一样重要？如果是，如何解释？在这个问题上，早期的 CUDA 设备和最新的设备之间是否存在差异？

【问题讨论】：

在 CUDA C 编程指南（第 5.3.2 节）和 CUDA C 最佳实践指南（第 9.2.1 节）中广泛讨论了合并问题。这两个指南还涵盖了不同架构的合并问题。为避免复制材料，如果您查看这些文件并张贴晦涩难懂且需要澄清的要点，这将更具建设性。

【解决方案1】：

合并的写入（或缺少）会影响性能，就像合并的读取（或缺少）一样。

当一个读取请求被一个 warp 指令触发时，就会发生合并读取，例如：

int i = my_int_data[threadIdx.x+blockDim.x*blockIdx.x];

可以通过内存控制器中的单个 read 事务来满足（这实际上是说所有单独的线程读取都来自单个缓存行。）

当warp指令触发写入请求时发生合并写入，例如：

my_int_data[threadIdx.x+blockDim.x*blockIdx.x] = i;

可以通过内存控制器中的单个写入事务来满足。

对于上面我展示的例子，代际上没有区别。

但还有其他类型的读取或写入可以在后来的设备中合并（即合并为单个内存控制器事务），但在早期的设备中则不然。一个例子是“广播阅读”：

int i = my_int_data[0];

在上面的例子中，所有线程都从同一个全局位置读取。在较新的设备中，这样的读取将在单个事务中“广播”到所有线程。在一些早期的设备中，这将导致线程的序列化服务。这样的示例可能没有写入的必然结果，因为多个线程写入单个位置会产生未定义的行为。但是，“加扰”写入可能会在较新的设备上合并，但不会在较旧的设备上合并：

my_int_data[(threadIdx.x+5)%32] = i;

请注意，以上所有写入都是唯一的（在 warp 内）并且属于单独的缓存行，但它们不满足 1.0 或 1.1 设备上的合并要求，但应该在较新的设备上。

如果您阅读global memory access description for devices of cc 1.0 and 1.1，并与较新的设备进行比较，您会发现在较早的设备上合并的一些要求已在较新的设备上放宽。

【讨论】：

谢谢。您能否解释一下在写入时如何涉及缓存？您在合并的读取事务中指出，“所有单独的线程读取都来自单个高速缓存行。”所以在写入的情况下，非合并写入占用了几条 L2 缓存行，对吧？
是的，非合并内存事务跨越多个缓存行，无论是读取还是写入。缓存本身在这里不是问题。缓存线是内存控制器强制执行的基本交换量。

【解决方案2】：

我们在我开设的课程中做了这个实验。事实证明，合并在写入中比在读取中更重要，这可能是因为 L1 和 L2 缓存存储了一些未使用的数据以供以后使用。

【讨论】：