c - 关于全局内存访问方式

Question

一般来说，对于GPU，哪种访问模式(从连续的全局内存块读取数据)更快？

(1) 单线程或极少数线程的for循环，从全局内存块读取数据；

(2) 让很多线程，可能来自不同的 block ，同时从全局内存中读取数据。

例如

if (threadIdx.x==0)
{
  for (int i=0; i<1000; ++i)

     buffer[i]=data[i];//data is stored in global memory
}

OR:

buffer[threadIdx.x]=data[threadIdx.x];//there are 1000 threads in this thread block

Answer 1

总之，第二个应该更快。理由如下:

并行有两种:线程级并行(TLP)和指令级并行(ILP)。您的第一个代码(循环)针对 ILP，第二个代码利用 TLP。

当 TLP 被利用时，许多内存请求是同时发出的，没有任何控制流依赖性。在这种情况下，硬件可以利用线程之间的局部性来减少总内存事务(如果可能)。此外，硬件可以通过 L2 缓存组并行性、内存 Controller 并行性、DRAM 组并行性和许多其他级别的并行性来同时处理并发请求。

但是，在 ILP 案例中，现有的控制依赖性限制了并发发出的内存请求的数量。即使在循环展开的情况下也是如此(记分板大小和指令窗口大小等硬件资源限制了未完成指令的总数)。因此，许多内存请求实际上是不必要的序列化。此外，没有利用内存访问合并的硬件能力。