3

假设我正在尝试对数组大小 n 进行简单的缩减，比如说保持在一个工作单元内……比如说添加所有元素。一般策略似乎是在每个 GPU 上生成许多工作项，从而减少树中的项。天真地这似乎需要 log n 步，但这并不是说第一波线程所有这些线程都一次性完成，是吗？他们被安排在经线中。

for(int offset = get_local_size(0) / 2;
      offset > 0;
      offset >>= 1) {
     if (local_index < offset) {
       float other = scratch[local_index + offset];
       float mine = scratch[local_index];
       scratch[local_index] = (mine < other) ? mine : other;
     }
     barrier(CLK_LOCAL_MEM_FENCE);
   }

因此并行添加了 32 个项目，然后该线程在屏障处等待。另一个 32 去，我们在障碍物处等待。再走 32 次，我们在屏障处等待，直到所有线程都完成了进入树最顶层所需的 n/2 次加法，然后我们绕过循环。凉爽的。

这看起来不错，但也许很复杂？我知道指令级并行性很重要，所以为什么不生成一个线程并执行类似的操作

while(i<array size){
    scratch[0] += scratch[i+16]
    scratch[1] += scratch[i+17]
    scratch[2] += scratch[i+17]
    ...
    i+=16
}
...
int accum = 0;
accum += scratch[0]
accum += scratch[1]
accum += scratch[2]
accum += scratch[3]
...

这样所有的添加都发生在一个扭曲中。现在你有一个线程可以让 gpu 尽可能地忙碌。

现在假设指令级并行性并不是真正的事情。下面的情况如何，工作大小设置为 32（经纱数）。

for(int i = get_local_id(0);i += 32;i++){
    scratch[get_local_id(0)] += scratch[i+get_local_id(0)]
}

然后将前 32 项相加。我想那 32 个线程会一次又一次地触发。

如果您不反对放弃 OpenCL 的通用性，那么当您知道每个循环会触发多少个添加时，为什么还要在树中减少呢？

4

1 回答 1

5

一个线程不能让 GPU 保持忙碌。这与说一个线程可以保持 8 核 CPU 繁忙大致相同。

为了最大限度地利用计算资源以及可用的内存带宽，有必要利用整个机器（即所有可以执行线程的可用资源）。

对于大多数较新的 GPU，您当然可以通过指令级并行性来提高性能，方法是让您的线程代码按顺序具有多个独立指令。但是您不能将所有这些都投入到单个线程中并期望它能够提供良好的性能。

当您有 2 个顺序指令时，如下所示：

scratch[0] += scratch[i+16]
scratch[1] += scratch[i+17]

这对 ILP 有好处，因为这两个操作完全相互独立。但是，由于 GPU 发出内存事务的方式，第一行代码将参与特定的内存事务，而第二行代码将必然参与不同的内存事务。

当我们有一个warp一起工作时，一行代码是这样的：

float other = scratch[local_index + offset];

将导致 warp 的所有成员生成请求，但这些请求将全部组合成一个或两个内存事务。这就是您可以实现完全带宽利用的方式。

尽管大多数现代 GPU 都有缓存，并且缓存会在某种程度上弥合这两种方法之间的差距，但它们绝不会弥补让所有 warp 成员发出组合请求与单个 warp 之间的事务中的巨大差异成员按顺序发出一组请求。

您可能想了解 GPU 内存合并。由于您的问题似乎是以 OpenCL 为中心的，因此您可能对本文档感兴趣。

于 2013-08-06T01:07:07.127 回答