cuda - 全局写入中的内存合并

Question

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

Answer 1

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

当由 warp 指令触发读取请求时，会发生合并读取，例如：

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 设备的全局内存访问描述，并与更高版本的设备进行比较，您会看到在早期设备上合并的一些要求在以后的设备上已经放宽了。

Answer 2

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