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

Posted 2023-02-16

【中文标题】分支分歧真的那么糟糕吗？

【英文标题】：Is branch divergence really so bad?

【发布时间】：2013-06-20 20:58:30

【问题描述】：

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

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

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

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

foo(A)
bar(B)

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

在第一种情况下，线程 1 和 2 串行执行。调用这两个时钟周期。

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

【问题讨论】：

简短的回答是，只有当一个 warp 中的所有线程不遵循同一路径通过分支时，分支才会有问题。发生这种情况时，您将获得指令重放，这会降低指令吞吐量和性能。 但是编译器足够聪明，可以将“次要”分支转换为条件执行，这基本上是免费的。所以这通常是一个被大大夸大的问题。

@talonmies 我不知道其中一些术语。指令回放？我假设这意味着在同一个 warp 中的不同线程上串行执行相同的指令，但我发现的唯一 definition 似乎表明指令重放是由 not 执行从发送的所有指令引起的主持人，在这种情况下，我不确定您的意思。至于条件执行，我知道的唯一定义是分支分歧的同义词。

分支分歧有不同程度的开销。在许多情况下，编译器可以使用预测来消除分支或更改为扭曲活动掩码。在这些情况下，您只会受到执行条件测试和设置谓词的指令执行延迟的影响。如果代码确实存在分歧，那么问题是执行了多少附加指令来执行多个代码路径。分支本身和为分歧而记账可能会增加开销并导致扭曲停止等待分支地址解析和获取。这可以通过更高的入住率来隐藏。

我只是想指出您的问题中可能存在的错误。您提到“不同之处在于，在前一种情况下，如果 foo(A) 涉及从内存读取，我想线程 2 可以在该延迟期间开始执行”但我认为这不会发生，因为线程 1 和线程 2 是在同一个 warp 中，这意味着它们总是执行相同的指令（尽管其中一个可能被屏蔽而另一个没有）

有一篇关于减少分支分歧的有趣论文：eecis.udel.edu/~cavazos/cisc879/papers/a3-han.pdf。它很好地解释了当两个分支不是真正不同时可以使用哪些类型的优化。我不相信真正的独特性会经常发生。

【参考方案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 像素的块连续： BL0 BL1 BL2 ... GR0 GR1 GR2 ...

我可以写类似的代码：

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

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

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

【讨论】：

所以基本上你所说的并不是分支分歧一般来说不好，而是当有机会利用 GPU 的 SIMD 架构时应该消除它（如上面的示例）有空吗？

是的。分支分歧在任何地方都不好（在某种程度上）。 CPU 制造商在 CPU 推测执行和分支预测方面投入了大量精力，以解决分支分歧的负面影响。在 GPU 上，它的效果非常明显。如果你可以通过巧妙地重新设计你的算法来避免它，那很好。但如果不是这样，机器至少可以让你灵活地以这种方式编码，这样程序员的工作就更容易接受了。机器可以处理代码中的分支这一事实优于一些更严格的架构。

好的，您正在进入分支预测，我认为您的上述帖子涵盖了与分支分歧有关的 GPU 特定问题，并且我对您的理解正确。非常感谢！

我建议分支成本高昂（无论是 CPU 还是 GPU）。 CPU 制造商已经通过某些方法（例如分支预测、推测执行）在一定程度上解决了这个问题。在 GPU 方面，没有分支预测或推测执行。但无论如何，无论是在 GPU 代码还是 CPU 代码中，分歧都会付出代价。

啊，刚刚阅读了指令流水线。我明白你现在的意思了。谢谢你没有让我沉迷于我的无知。 :D