在CUDA Developer Blog上阅读了这篇文章后,我努力了解何时可以安全\正确使用__activemask()代替__ballot_sync()。
__activemask()
__ballot_sync()
在活动掩码查询部分中,作者写道:
这是不正确的,因为它将导致部分和而不是总计。
之后,在“ 机会扭曲级编程”部分中,他们使用此函数__activemask()是因为:
如果要在库函数中使用扭曲级编程,但是不能更改函数接口,则可能会很困难。
cuda gpu-warp
cuda ×1
gpu-warp ×1