if 语句中的 CUDA 分歧，没有 else 分支？

Question

在 CUDA 设备代码中，以下 if-else 语句将导致 warp 线程之间的分歧，从而导致 SIMD 硬件两次通过。假设Vs 是共享内存中的一个位置。

if (threadIdx.x % 2) {
  Vs[threadIdx.x] = 0;
} else {
  Vs[threadIdx.x] = 1;
}

我相信当我们有一个if 语句时也会有两次通过，没有 else 分支。为什么会这样？

if (threadIdx.x % 2) {
  Vs[threadIdx.x] = 0;
}

下面的if 语句会在 3 遍中完成吗？

if        (threadIdx.x < 10) {
  Vs[threadIdx.x] = 0;
} else if (threadIdx.x < 20) {
  Vs[threadIdx.x] = 1;
} else {
  Vs[threadIdx.x] = 2;
}

Answer 1

在 GPU 上，很可能只有一个带有 if-else 语句的遍 - 一个谓词遍。该条件只会在“then”块期间为一半线程打开“什么都不做”位，并为“else”块关闭另一半的“什么也不做”位。

然而，正如@njuffa 指出的那样，这取决于目标架构等参数。

更多详情见：

Branch predication on GPU

---

对于if 主体的第一个特定示例，编译器甚至可能不需要谓词传递，因为它可以重写为

Vs[threadIdx.x] = (threadIdx.x % 2 ? 0 : 1);

这在你的经线上完全一致。对于您的最后一个示例-它确实取决于，但理论上它可以再次由编译器优化为单个非谓词传递，并且您也可能会有一个谓词单路径，在三个中的每一个中都有不同的谓词范围。