GPGPU中的分歧 [英] Divergence in GPGPU

问题描述

我知道，在出现分歧的情况下，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

这将被执行为：

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的性能会减半。

这篇关于GPGPU中的分歧的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持IT屋！