我试图了解下一段代码中翘曲发散的惩罚是什么。我了解原则上如何处理扭曲分歧（小分支的预测指令，大分支的扭曲投票和分支 - 如果所有扭曲都同意，否则预测指令并且没有分支，与小分支相同）。但是，我不了解具体细节 - 如何处理带有中断/继续的 while 循环。
在下面的示例中，当通道 X 的 scrapEverythingCondition() 计算结果为 true 时，将发生以下哪种情况：
- 评估在内部循环中进行，直到 n == N_N，通道 X 始终评估 nops，i 递增，所有通道现在一起工作。
- 除了车道 X 之外的每个人都执行 someMoreWork() 而车道 X 评估 nops，车道 X 执行 someCostlyInitialization() 然后 n=0 而其他所有人都评估 nops，所有车道继续一起评估内部循环（显然具有不同的 n 值）。
- 还有一些我没有想到的。
代码：
__global__ void chainKernel() {
int i = threadIdx.x + blockIdx.x * blockDim.x;
while (i < N_I) {
someCostlyInitialization();
for(int n = 0; n < N_N; ++n) {
someStatisticsComputations(n);
if (scrapEverythingCondition(n)) {
// Everything we did for current i is no good. Scrap and begin again
i -= BLOCKS*THREADS;
break;
}
someMoreWork();
}
i += BLOCKS*THREADS;
}
}
我尝试编译到 PTX 并查看生成的代码，但它对我来说太复杂了:(
编辑：感谢 Maku 的回答。我还能够使用散布在代码周围的老式 printf() 来验证答案。我能够看到哪些线程以什么顺序到达哪里，并且确实选项 1 是正确的（通道 X 被暂停，直到内部 for 循环耗尽）。