goto 指令对 CUDA 代码中经线内发散的影响

Posted 2023-02-16

【中文标题】goto 指令对 CUDA 代码中经线内发散的影响

【英文标题】：The impact of goto instruction at intra-warp divergence in CUDA code

【发布时间】：2014-05-30 00:12:58

【问题描述】：

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

if( threadIdx.x < 16 ) 
    A:
    // do something.
 else 
    B:
    // do something else.

C:
// rest of code.

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

我的问题是 SM 是否能够像上面一样正确处理包含goto 指令的代码？还是不能保证选择的再收敛点是最佳的？ 例如，如果我在使用 goto 实现的 CUDA 代码中具有以下控制流

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 指令引起的经线内分歧的重新收敛点？

【参考方案1】：

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

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