在CUDA 9中以`_sync()`命名的一些内部元素；语义相同吗？

Question

在CUDA 9中，nVIDIA似乎有了“合作团体”的新概念；由于某些原因，__ballot()现在(= CUDA 9)被废弃，转而支持__ballot_sync()。那是别名还是语义改变了？

..。其他内置程序也有类似的问题，它们现在已经将__sync()添加到它们的名称中。

Answer 1

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

首先，为了奠定基础工作，有必要认识到Volta 介绍可能性 for 独立线程调度，通过引入一个每线程程序计数器和其他更改。因此，Volta有可能在长时间内表现为非曲同步行为，而在执行期间，以前的体系结构可能仍然是扭曲同步的。

大多数翘曲本质只为实际参与的线程提供预期的结果(即，在该周期中，实际对该指令的问题是活动的)。现在，程序员可以通过新的mask参数明确说明预期将参与哪些线程。然而，也有一些需求，特别是对Pascal和以前的体系结构。来自程序设计指南

但是，请注意，对于Pascal和更早的体系结构，mask中的所有线程必须在收敛中执行相同的翘曲内禀指令，而掩码中所有值的合并必须等于warp的活动掩码。

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

这相关问题讨论mask参数。这个答案并不是为了解决独立线程调度可能产生的所有可能问题以及对翘曲级本质的影响。为此，我鼓励大家阅读编程指南。