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

Question

CUDA Programming Guide (v4.1)在 Sec 5.4.2 中描述了关于谓词指令的内容：

仅当分支条件控制的指令数小于或等于某个阈值时，编译器才将分支指令替换为谓词指令：如果编译器确定该条件可能产生许多发散扭曲，则此阈值为 7，否则是4。

1. 一个条件如何产生许多不同的扭曲？给定条件只能将经线分成两部分。许多在这里是什么意思？
2. 即使上述内容是有道理的，编译器如何知道扭曲的运行时发散行为？

Answer 1

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

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

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 中使用了几年）在分支变得更经济之前应该使用多少谓词指令具有启发性。而且指令流水线中的很多分支似乎对性能不利，所以当编译器可以计算出代码会产生更高的“分支密度”时，它会更喜欢谓词指令而不是分支。