在CUDA 9中,nVIDIA似乎有了这种"合作团体"的新概念; 由于某些原因我不完全清楚,__ballot()现在(= CUDA 9)被弃用赞成__ballot_sync().这是别名还是语义改变了？
...对于其他内置组件的类似问题,现在已__sync()添加到他们的名字中.
没有语义不一样.函数调用本身是不同的,一个不是另一个的别名,新功能已经暴露,现在Volta架构和以前的架构之间的实现行为是不同的.
首先,设置地工作,有必要要弄清的是沃尔的事实推出的可能性为独立的线程调度,通过引入每个线程的程序计数器和其他变化.因此,Volta可能会在非warp同步行为中延长一段时间,并且在执行期间,以前的架构可能仍然是warp-synchronous.
大多数warp内在函数只通过为实际参与的线程提供预期结果(即在该周期中实际上对该指令的发布有效).程序员现在可以通过新mask参数明确指出哪些线程可以参与.但是有一些要求,特别是在Pascal和以前的架构上.从编程指南:
但是请注意,对于Pascal和早期的体系结构,所有线程
mask必须在收敛中执行相同的warp内部指令,并且mask中所有值的并集必须等于warp的活动掩码.
然而,在Volta上,warp执行引擎将在掩码中的指示线程之间产生必要的同步/参与,以便使期望/指示的操作有效(假设使用了适当_sync版本的instrinsic).为了清楚起见,warp执行引擎将重新收集在volta上分歧的线程以匹配掩码,但是它不会克服程序员引发的错误,例如阻止线程_sync()通过条件语句参与内部函数.
这个相关问题讨论了mask参数.此答案并非旨在解决独立线程调度可能产生的所有可能问题以及对warp级内在函数的影响.为此,我鼓励阅读编程指南.