束表决函数作用

warp vote是2.0+的一个特性，用来在warp内交换1-bit的信息的。

有三种，

一种是直接将32个1-bit的位，映射成1个32位整数。warp中的线程0，将占据这个映射结果的最低0位，线程1占据这结果的次低位，线程31占据这个32位整数结果的最高位。将这个结果复制32份，每个线程都将得到一份。

其他两种是这个操作的扩展，一个是and操作，将这个结果的32个bit，进行and, 得到最终的1-bit结果（所谓的__all())。而另外一种是or操作，将这个结果的32-bit，进行or，得到最终的1-bit结果（所谓的__any())操作。在底层硬件实现上，实际上只有一种操作的，

这个操作同时返回1个32位整数（Regular Register常规寄存器中），和一个1-bit的结果（Predicate Register中）。但是C语言不支持多返回值。所以只能在_ballot(), __any()或者__all()中选一种，选择要那个32位（4个字节的整数）值，还是那个1位值。

predicate 是“谓词”的意思。在fermi/kepler体系中，有8个(7个)特殊的寄存器，叫P寄存器，也叫“布尔寄存器”，每个里面只能保存1位。例如,bool a= c>d,这里的a就保存在“布尔寄存器”中。所有的条件想用，（例如if (...) {}),都必须放置到P寄存器里。而除了if (...)外，像是__ballot()这种里面的条件，也是在P寄存器里的。（自动安排的，无需，也不能工作安排）。

unsigned int result = __ballot(threadIdx.x < 16);如果你的block形状是(32,8)的话，你在每个warp中，得到32个：0x0000ffff。这样如果任何一个线程手中的苹果是红色的，那么warp内的其他所有线程都知道了。这种交换，相比传统的通过Shared memory的交换，效率极高。但只能交换1-bit。__ballot()在维护缓冲区的原子性追加上，是一个常见的优化。可以提高缓冲区的原子追加效率32X。

----------------------------------------------

以下为CUDA编程指南5.0中 内容

B.12 束表决（warp vote）函数
只有计算能力为1.2 或更高的设备支持束表决（参见4.1了解束的定义）函
数。
附录B C语言扩展143
int all (int predicate) ;
为束内的所有线程计算predicate，当且仅当所有线程的predicate均非零时
返回非零值。
int any(int predicate) ;
为束内的所有线程计算predicate，当且仅当任意线程的predicate 非零时返
回非零值。
unsigned int ballot ( int predicate) ;
为束内所有线程计算predicate值，并返回一个整数，如果束内第N个线程
的predicate值为非零，则该整数的第N位为1。计算能力为2.x及以上的设备支
持此函数。