performance - 分支分歧真的那么糟糕吗？

Question

我在互联网上看到了许多关于分支分歧以及如何避免它的问题。然而，即使在阅读了数十篇关于 CUDA 工作原理的文章后，我似乎看不出在大多数情况下避免分支分歧有什么帮助 .在有人伸出爪子扑向我之前，请允许我描述一下我认为“大多数情况”的情况。

在我看来，大多数分支分歧实例都涉及许多真正不同的代码块。例如，我们有以下场景:

if (A):
  foo(A)
else:
  bar(B)

如果我们有两个线程遇到这种分歧，线程 1 将首先执行，走路径 A。接下来，线程 2 将走路径 B。为了消除分歧，我们可以将上面的代码块改为如下所示:

foo(A)
bar(B)

假设调用 foo(A) 是安全的在线程 2 和 bar(B) 上在线程 1 上，人们可能期望性能会有所提高。但是，这是我的看法:

在第一种情况下，线程 1 和 2 串行执行。称之为两个时钟周期。

在第二种情况下，线程 1 和 2 执行 foo(A)并行，然后执行 bar(B)在平行下。这在我看来仍然像两个时钟周期，区别在于在前一种情况下，如果 foo(A)涉及从内存读取，我想线程 2 可以在该延迟期间开始执行，这导致延迟隐藏。如果是这种情况，分支发散代码会更快。

Answer 1

您假设(至少这是您提供的示例和您提供的唯一引用)避免分支分歧的唯一方法是允许所有线程执行所有代码。

在那种情况下，我同意没有太大的区别。

但是避免分支分歧可能更多地与更高级别的算法重组有关，而不仅仅是添加或删除一些 if 语句并使代码“安全”地在所有线程中执行。

我将提供一个例子。假设我知道奇数线程需要处理像素的蓝色分量，偶数线程需要处理绿色分量:

#define N 2 // number of pixel components
#define BLUE 0
#define GREEN 1
// pixel order: px0BL px0GR px1BL px1GR ...


if (threadIdx.x & 1)  foo(pixel(N*threadIdx.x+BLUE));
else                  bar(pixel(N*threadIdx.x+GREEN));

这意味着每个备用线程都采用给定的路径，无论是 foo或 bar .所以现在我的扭曲需要两倍的时间来执行。

但是，如果我重新排列我的像素数据，使颜色分量可能以 32 个像素的 block 连续:
BL0 BL1 BL2 ... GR0 GR1 GR2 ...

我可以写类似的代码:

if (threadIdx.x & 32)  foo(pixel(threadIdx.x));
else                   bar(pixel(threadIdx.x));

看起来我仍然有分歧的可能性。但是由于分歧发生在扭曲边界上，给定扭曲执行 if路径或 else路径，因此不会发生实际的分歧。

这是一个微不足道的例子，可能很愚蠢，但它说明了可能有一些方法可以解决扭曲分歧，而不涉及运行所有分歧路径的所有代码。