关于warp投票功能

Question

CUDA编程指南介绍了经线投票功能的概念，“_ all”，“_ any”和“__ ballot”。

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

Answer 1

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

__all和__any在引入__ballot之前用于减少，但我想不出任何其他用途。

Answer 2

__ballot的原型如下

unsigned int __ballot(int predicate);

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

结合atomicOr和__popc，它可以用于累积具有真实谓词的每个warp中的线程数。

的确，atomicOr的原型是

int atomicOr(int* address, int val);

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

另一方面，qazxsw poi返回qazxsw poi-bit参数中设置的位数。

据此，说明

__popc

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

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

Answer 3

作为使用32 API的算法示例，我将提到D.M Hughes等人的In-Kernel Stream Compaction。它用于前缀流压缩的总和部分，以计算传递谓词的元素数（每个warp）。

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]));

Answer 4

CUDA提供了几种warp范围的广播和缩减操作，NVIDIA的架构可以有效地支持这些操作。例如，__ ballot（谓词）指令为warp的所有活动线程计算谓词，并返回一个整数，其第N位被设置当且仅当谓词为warp的第N个线程计算为非零并且第N个线程处于活动状态时[参考：GPU架构的灵活软件分析]。