goto 指令对 CUDA 代码中扭曲内发散的影响

Question

对于CUDA中简单的warp内线程发散，我所知道的是SM选择一个重新收敛点（PC地址），并在两个/多个路径中执行指令，同时禁用未采取该路径的线程的执行效果。
例如，在下面的代码中：

if( threadIdx.x < 16 ) {
    A:
    // do something.
} else {
    B:
    // do something else.
}
C:
// rest of code.

C是重新收敛点，warp 调度程序在 和 处调度指令A，B同时禁用A上半 warp 处的指令，并禁用B下半 warp 处的指令。当它达到 时C，将为经线内的所有线程启用指令。

我的问题是 SM 是否能够goto像上面那样正确处理包括指令在内的代码？或者不能保证所选择的重收敛点是最佳的？
例如，如果我的 CUDA 代码中有以下控制流，使用goto

A:
// some code here.
B:
// some code here too.
if( threadIdx.x < 16 ) {
    C:
    // do something.
    goto A;
}
// do something else.
goto B;

SM是否足够聪明，可以决定B作为由指令引起的扭曲内发散的重新收敛点if？

Answer 1

一般来说，goto非结构化控制流会干扰许多编译器优化，无论平台如何。CUDA C 编译器应以goto功能正确的方式处理代码，但性能可能不是最佳的。

性能欠佳的部分原因可能是编译器对收敛点的放置。您可以使用 来检查生成的机器代码 (SASS) 中的收敛点cuobjdump --dump-sass。一条SSY指令记录了一个收敛点，.S指令上的后缀表示控制权转移到最后记录的收敛点。