GPGPU 的分歧答案

【问题标题】：Divergence in GPGPUGPGPU 的分歧
【发布时间】：2015-11-12 04:39:31
【问题描述】：

我知道，如果每个工作项在 GPU 上执行 if 和 else 都存在分歧，最后在掩码的帮助下我们选择了一个。但我不明白如果我们都在执行，那么它如何增加执行单元的空闲。

我已经解决了几个关于 stackoverflow 的问题，但它们与它如何影响执行单元的空闲状态无关。

谁能清楚地解释一下这个概念？分歧如何增加执行单元的停顿或空闲？

【问题讨论】：

如果两个分支都必须执行，那么为什么线程不会同时使用同一个分支
好的，让我打破我的问题。如果存在分歧，则对同一线程执行 if 和 else 。据我所知，两者都是按顺序执行的。如果您认为我错了，请在这里纠正我
哦。您的解释是使用扭曲。谢谢你。我可以针对单个线程说 if 和 else 都执行了吗？

标签： cuda opencl gpgpu

【解决方案1】：

与@麦龙东讨论后的几点：执行单元空闲时间：是执行单元没有做任何有用工作的时间。

在每个周期，warp 最多可以对最多 32 个数据元素（或任何 warp 宽度）执行一条指令。如果您需要执行两条不同的指令（如经纱分歧分支的情况），您需要在两个周期内发出这些指令。

当我们使用 if 和 else 时，两个分支都会在整个 warp 中执行。但是，如果第一个线程采用 if 则它被禁用，而其余的 warp 采用 else。在其余的经线完成 else 之前，它不会做任何事情。这可以被认为是对执行单元的浪费（与少分支代码相比）。当第一个线程将执行时，其他线程将被禁用。这会导致有用工作量减少，从而增加执行单元的空闲时间。

如果一个 warp 中的所有 32 个项目都在执行仅采用 if 条件路径或 else 条件路径，则执行单元的空闲时间不会增加。可以参考here的讨论

【讨论】：

【解决方案2】：

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

让我们看一个 8 宽 SIMD 单元的小例子（真正的 GPU SIMD 单元是 32 (NVidia) 或 64-ops (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

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

【讨论】：