问关于warp投票函数

EN

在CUDA-histogram和CUDA NPP库中使用__ballot来快速生成位掩码，并将其与__popc内部特性相结合，以非常有效地实现布尔约简。

在引入__ballot之前，__all和__any被用在了reduction中，尽管我想不出还有其他的用法。

__ballot的原型如下

unsigned int __ballot(int predicate);

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

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

实际上，atomicOr的原型是

int atomicOr(int* address, int val);

atomicOr读取address指向的值，使用val执行逐位OR操作，然后将该值写回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]));

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

有关更多详细信息，请参阅Morgan Kaufmann CUDA Programming的Shane Cook

作为使用__ballot应用编程接口的算法的一个例子，我会提到D.M.Hughes等人的内核内流压缩。它在流压缩的前缀sum部分中使用，以计算(每扭曲)通过谓词的元素的数量。

Here the paper. In-k Stream Compaction