cuda - 指令级并行和线程级并行如何在 GPU 上工作？

Question

假设我正在尝试对数组大小 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 的通用性，那么当您知道每个周期会触发多少个添加项时，为什么还要在树中减少呢？

Answer 1

一个线程无法让 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 都有缓存，并且缓存会在某种程度上弥补这两种方法之间的差距，但它们绝不会弥补所有经线成员发出组合请求与单个经线之间的巨大差异成员按顺序发出一组请求。

您可能需要阅读 GPU 内存合并。由于您的问题似乎以 OpenCL 为中心，您可能对 this document 感兴趣.