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

Question

在 CUDA 9 中，nVIDIA 似乎有了“协作组”这个新概念；由于某些我不完全清楚的原因，__ballot() 现在(= CUDA 9)已弃用，取而代之的是 __ballot_sync()。这是别名还是语义发生了变化？

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

Answer 1

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

首先，要做好基础工作，必须认识到 Volta introduced the possibility对于independent thread scheduling ，通过引入每线程程序计数器和其他更改。因此，Volta 有可能在很长一段时间内以非扭曲同步行为运行，并且在以前的架构可能仍然是扭曲同步的执行期间。

大多数 warp 内在函数仅通过为实际参与的线程提供预期结果来工作(即在该周期中对于该指令的发出实际上是活跃的)。程序员现在可以通过新的 mask 参数明确哪些线程应该参与。然而，有一些要求，特别是在 Pascal 和以前的体系结构上。来自 the programming guide :

Note, however, that for Pascal and earlier architectures, all threads in mask must execute the same warp intrinsic instruction in convergence, and the union of all values in mask must be equal to the warp's active mask.

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

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