全局写操作中的内存合并

Question

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

在CUDA设备中，全局内存写入的合并和全局内存读取的合并同样重要。这是因为当多个线程尝试同时写入全局内存时，合并操作可以减少访问内存的次数，并提高内存带宽利用率。在早期的CUDA设备中，由于硬件限制，全局内存写入的合并可能会受到更大的影响。但是，在最新的CUDA设备中，这个问题已经得到了改善。

- Farzad

---

在CUDA C编程指南（第5.3.2节）和CUDA C最佳实践指南（第9.2.1节）中广泛讨论了协同合并问题。这两个指南还涵盖了不同架构的协同合并问题。为了避免重复材料，如果您查看这些文件并发布需要澄清的模糊点，将会更具建设性。 - Vitality

Answer 1

Coalesced writes（或缺乏）会影响性能，就像 coalesced reads（或缺乏）一样。当由warp指令触发的读取请求时，会发生coalesced read，例如：

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

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

当由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设备上的合并要求，但应该能够满足较新设备上的要求。

如果您阅读cc 1.0和1.1设备的全局内存访问描述，并与较新设备进行比较，您将看到在早期设备上用于合并的某些要求已经在后来的设备上放宽了。

- Robert Crovella

---

谢谢。您能否进一步解释在写入情况下缓存是如何参与的？您指出在合并读取事务中，“所有单个线程读取都来自单个缓存行。”那么在写入情况下，非合并写入会占用多个L2缓存行，对吗？ - Farzad

是的，非合并内存事务跨越一个以上的高速缓存行，无论是读取还是写入。这里并不涉及高速缓存本身的问题。高速缓存行是由内存控制器强制执行的基本交换量子。 - Robert Crovella

Answer 2

我们在我所开设的课程中进行了这个实验。聚合在写入操作中比读取操作中略微更重要，可能是因为L1和L2缓存会存储一些未使用的数据以备后用。

- Levi Barnes