我知道，在出现分歧的情况下，IF和ELSE都是在GPU上为每个工作项执行的，最后在掩码的帮助下我们选择了其中一个。但我不能理解，如果我们两个都被执行了，那么它怎么会增加执行单元的空闲。

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

有人能给我解释清楚这个概念吗？分歧如何增加执行单元中的停滞或空闲？

推荐答案

在执行两个代码路径以选择一个结果后，不使用掩码，但在执行期间使用不同的掩码，以便仅启用在当前执行的代码路径中处于活动状态的当前Warp内的线程。

让我们看一个8宽SIMD单元的小示例(实际的GPU SIMD单元是32(Nvidia)或64-op(AMD GCN)宽)：

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

这将被执行为：

if ((threadIdx.x % 2) == 0) // MASK: 11111111 all threads enabled

a = b+c; // MASK: 10101010, only even threads enabled, odd threads idle

a = b*c; // MASK: 01010101, only odd threads enabled, even threads idle

d = a*2 // MASK: 11111111, all threads are enabled again