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

Question

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

Answer 1

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

合并读取发生在由 warp 指令触发的读取请求时，例如:

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

可以通过内存 Controller 中的单个读取事务来满足(这实际上是说所有单独的线程读取都来自单个缓存行。)

当由 warp 指令触发的写请求时发生合并写，例如:

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

可以通过存储 Controller 中的单个写入事务来满足。

对于我展示的上述示例，代际没有差异。

但是还有其他类型的读取或写入可以在较晚的设备中合并(即崩溃到单个内存 Controller 事务)，但不能在较早的设备中合并。一个例子是“广播阅读”:

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 ，并与后来的设备进行比较，您会看到在较早的设备上进行合并的一些要求已在较晚的设备上放宽。