cuda - 在 CUDA 9 中附加了一些以 `_sync()` 命名的内在函数；语义相同？

Question

在 CUDA 9 中，nVIDIA 似乎有了这种“合作组”的新概念；出于某种原因，我并不完全清楚，__ballot()现在（= CUDA 9）不推荐使用__ballot_sync(). 那是别名还是语义发生了变化？

...对于现在已__sync()添加到其名称中的其他内置函数的类似问题。

Answer 1

不，语义不一样。函数调用本身是不同的，一个不是另一个的别名，新功能已经暴露，现在 Volta 架构和以前的架构之间的实现行为不同。

首先，要打好基础，有必要认识到 Volta通过引入每线程程序计数器和其他更改引入了独立线程调度的可能性。因此，Volta 有可能在很长一段时间内以非扭曲同步行为表现，并且在以前的架构可能仍然是扭曲同步的执行期间。

大多数warp内在函数仅通过为实际参与的线程提供预期结果（即在该周期中对于该指令的发出实际上是活动的）来工作。程序员现在可以通过新mask参数明确预期哪些线程参与。但是有一些要求，特别是在 Pascal 和以前的体系结构上。从编程指南：

但是请注意，对于 Pascal 和更早的架构，所有线程在mask收敛时必须执行相同的 warp 内在指令，并且 mask 中所有值的并集必须等于 warp 的活动掩码。

然而，在 Volta 上，warp 执行引擎将在掩码中的指示线程之间带来必要的同步/参与，以使所需/指示的操作有效（假设使用了适当_sync版本的 instrinsic）。需要明确的是，warp 执行引擎将重新收敛在 volta 上发散的线程以匹配掩码，但是它不会克服程序员引起的错误，例如_sync()通过条件语句阻止线程参与内在函数。

这个相关问题讨论了mask参数。该答案并非旨在解决独立线程调度可能产生的所有可能问题以及对warp级别内在函数的影响。为此，我鼓励阅读编程指南。