全局写操作中的内存合并

全局写操作中的内存合并

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设备的全局内存访问描述,并与较新设备进行比较,您将看到在早期设备上用于合并的某些要求已经在后来的设备上放宽了。