合并写入(或缺少写入)可以影响性能,就像合并读取(或缺少读取)一样.
当由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];
在上面的示例中,所有线程都从同一个全局位置读取.在较新的设备中,这种读取将"广播"到单个事务中的所有线程.在某些早期设备中,这将导致线程的序列化服务.这样的示例在写入中可能没有必然结果,因为写入单个位置的多个线程会给出未定义的行为.但是,"scrambled"写入可能会在较新的设备上合并,但不会更旧:
my_int_data[(threadIdx.x+5)%32] = i;
请注意,上面的所有写入都是唯一的(在warp中)并且属于单个缓存行,但它们不满足1.0或1.1设备上的合并要求,但应该在较新的设备上.
如果您阅读cc 1.0和1.1设备的全局内存访问描述,并与后续设备进行比较,您将看到在以后的设备上放宽的早期设备上的一些合并要求.