CUDA 编译器如何知道扭曲的发散行为？

Question

CUDA 编程指南 (v4.1) 在第 5.4.2 节中描述了关于谓词指令的内容：

编译器将分支指令替换为谓词 指令只有当指令的数量由控制 分支条件小于或等于某个阈值：如果 编译器确定该条件可能产生许多 发散warp，这个阈值为7，否则为4。

1. 一个条件如何产生许多发散的扭曲？给定条件可以 只将经线分成两部分。 很多在这里是什么意思？
2. 即使上面说得通，编译器怎么能知道运行时 经线的发散行为？

Answer 1

Warp 永远不会“分裂”。它们要么需要“条件执行”（因此在屏蔽不参与线程的情况下执行）来服务有条件的发散代码路径，要么不需要。

关于一个条件如何产生多个不同的扭曲，请考虑以下人为设计的示例：

if (threadIdx.x < 128) {
   // Only first four warps process here
   int modthirtytwo = threadIdx.x % 32;

   if (modthirtytwo == 0) {
      // Action A only first thread in the warp
   } else {
      // Action B for the other threads in the warp
   }
}

在这里，代码可以产生多个不同的扭曲，编译器应该能够在编译时对行为进行建模。如果为内核的编译器指定了启动边界，那就更好了。将此情况与仅使用一个扭曲的共享内存减少进行比较。

if (threadIdx.x < 32) {
   if (threadIdx.x < 16)  shm[threadIdx.x] += shm[threadIdx.x+16];
   if (threadIdx.x < 8)   shm[threadIdx.x] += shm[threadIdx.x+8];
   if (threadIdx.x < 4)   shm[threadIdx.x] += shm[threadIdx.x+4];
   if (threadIdx.x < 2)   shm[threadIdx.x] += shm[threadIdx.x+2];
   if (threadIdx.x == 0)  shm[0] += shm[1];
}

这里的分歧被限制在每块一个单一的扭曲。这段文字的意思是这两种情况下的编译器行为可能不同。

似乎“新”编译器（它已在 OpenCL 中使用了几年）对在分支变得更经济之前应该使用多少谓词指令具有启发性。而且指令管道中的很多分支似乎对性能不利，所以当编译器可以计算出代码会产生更高的“分支密度”时，它会更喜欢更多的谓词指令而不是分支。