关于warp投票功能

Question

CUDA编程指南介绍了warp投票功能的概念,"_ all"," _ any"和"__ballot".

我的问题是:哪些应用程序将使用这3个功能？

Answer 1

原型__ballot如下

unsigned int __ballot(int predicate);

如果predicate非零，则__ballot返回一个N设置了第 th 位的值，其中N是线程索引。

结合atomicOrand __popc，它可用于累积每个经线中具有真谓词的线程数。

确实，原型atomicOr是

int atomicOr(int* address, int val);

并atomicOr读取 指向的值address，使用 执行按位OR运算val，然后将值写回address并将其旧值作为返回参数返回。

另一方面，__popc返回使用32-bit 参数设置的位数。

因此，指令

volatile __shared__ u32 warp_shared_ballot[MAX_WARPS_PER_BLOCK];

const u32 warp_sum = threadIdx.x >> 5;

atomicOr(&warp_shared_ballot[warp_num],__ballot(data[tid]>threshold));

atomicAdd(&block_shared_accumulate,__popc(warp_shared_ballot[warp_num]));

可用于计算谓词为真的线程数。

有关更多详细信息，请参阅 Shane Cook、CUDA 编程、Morgan Kaufmann

Answer 2

__ballot用于CUDA-histogram和 CUDA NPP 库中，用于快速生成位掩码，并将其与__popc内在函数相结合，以非常有效地实现布尔归约。

__all并且__any在引入 之前被用于reduce __ballot，虽然我想不出它们的任何其他用途。