GPGPU中的发散

Question

我知道，在出现发散的情况下，在GPU上对每个工作项执行if和each，最后在掩码的帮助下，我们选择了一个。但我不明白，如果我们两个都被执行，那么它如何增加执行单位的懒惰。

我已经经历了一些关于堆栈溢出的问题，但它们与它如何影响执行单元的闲置无关。

有人能把这个概念解释清楚吗？分歧如何增加执行单位的停滞或懒惰？

Answer 1

在执行两个代码路径以选择一个结果之后，不使用掩码，但在执行过程中使用不同的掩码，以便只启用当前代码路径中活动的当前翘曲中的线程。

让我们看一个8宽的SIMD单元(实际GPU单元为32 (NVidia)或64 ops (AMD )宽)的小例子：

if ((threadIdx.x % 2) == 0)   
{
   a = b+c;   // Even threads
} else
   a = b*c;   // Odd threads
}
d = a*2

这将以以下方式执行：

1. `if ((threadIdx.x % 2) == 0) // MASK: 11111111 all threads enabled`
2. `a = b+c; // MASK: 10101010, only even threads enabled, odd threads idle`
3. `a = b*c; // MASK: 01010101, only odd threads enabled, even threads idle`
4. `d = a*2 // MASK: 11111111, all threads are enabled again`

第4行被称为“恢复点”，因为在第1行分裂的控件流在这里合并。在第2行和第3行中，只使用了一半的执行单元，而其他执行单元保持空闲状态。因此，在执行这些行时，GPU的性能减半。

Answer 2

几点后，与@Mai龙洞:空闲时间的执行单位:是时间，执行单位没有做任何有用的工作。

在每个周期中，一个翘曲最多只能对32个数据元素执行一条指令(或任何翘曲宽度)。如果您需要执行两个不同的指令(如翘曲发散分支的情况)，则需要在两个周期内发出这些指令。

当我们使用if和the时，对整个翘曲执行两个分支。但是，如果第一个线程使用if，那么它将被禁用，而其余的翘曲则会被禁用。它不会做任何事情，直到其余的经纱完成其余的。这可以被认为是对执行单元的浪费(与分支代码相比)。当第一个线程执行时，其他线程将被禁用。这导致了有用工作量的减少，从而增加了执行单元的空闲时间。

如果翘曲中的所有32项都在执行，则执行单元的空闲时间没有增加。你可以参考here的讨论