分支差异真的那么糟糕吗？

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.所以现在我的warp需要两倍的执行时间.

但是,如果我重新排列我的像素数据,使得颜色分量可能是32像素的块:BL0 BL1 BL2 ... GR0 GR1 GR2 ......

我可以写类似的代码:

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

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

这是一个简单的例子,可能是愚蠢的,但它说明了可能有办法解决扭曲分歧,而不涉及运行所有不同路径的所有代码.