在CUDA 9中,nVIDIA似乎有了“合作组”这个新概念;由于某种我不太清楚的原因,__ballot()
现在(= CUDA 9)已弃用,取而代之的是__ballot_sync()
。这是别名还是语义发生了变化?
...其他内置函数现在有类似的问题__sync()
添加到他们的名字中。
不,语义不一样。函数调用本身不同,一个不是另一个的别名,新功能已经公开,并且 Volta 架构和以前的架构之间的实现行为现在有所不同。
首先,为了奠定基础,有必要认识到 Volta介绍了可能性 for 独立线程调度,通过引入每线程程序计数器和其他更改。因此,Volta 可能会在较长时间内以非扭曲同步行为运行,并且在执行期间(以前的架构可能仍然是扭曲同步的)。
大多数 warp 内在函数的工作方式是仅为实际参与的线程提供预期结果(即,在该周期中实际发出该指令时处于活动状态)。程序员现在可以通过新的方法明确哪些线程应该参与mask
范围。然而,有一些要求,特别是对 Pascal 和以前的架构。从编程指南:
但请注意,对于 Pascal 和更早的体系结构,所有线程mask
必须在收敛时执行相同的 warp 内在指令,并且 mask 中所有值的并集必须等于 warp 的活动 mask。
然而,在 Volta 上,warp 执行引擎将在掩码中指示的线程之间实现必要的同步/参与,以便使所需/指示的操作有效(假设适当的_sync
使用内在的版本)。需要明确的是,warp 执行引擎将重新聚合在 volta 上发散的线程以匹配掩码,但是它不会克服程序员引起的错误,例如阻止线程参与_sync()
通过条件语句内在的。
This相关问题讨论了mask
范围。此答案并非旨在解决独立线程调度可能出现的所有可能问题以及对扭曲级别内在函数的影响。为此,我鼓励阅读编程指南。
本文内容由网友自发贡献,版权归原作者所有,本站不承担相应法律责任。如您发现有涉嫌抄袭侵权的内容,请联系:hwhale#tublm.com(使用前将#替换为@)