DAY50：阅读Warp Vote Functions

B.13. Warp Vote Functions

Deprecation notice: __any, __all, and __ballot have been deprecated as of CUDA 9.0.

The warp vote functions allow the threads of a given warp to perform a reduction-and-broadcast operation. These functions take as input an integer predicate from each thread in the warp and compare those values with zero. The results of the comparisons are combined (reduced) across the active threads of the warp in one of the following ways, broadcasting a single return value to each participating thread:

__all_sync(unsigned mask, predicate):

Evaluate predicate for all non-exited threads in mask and return non-zero if and only if predicate evaluates to non-zero for all of them.

__any_sync(unsigned mask, predicate):

Evaluate predicate for all non-exited threads in mask and return non-zero if and only if predicate evaluates to non-zero for any of them.

__ballot_sync(unsigned mask, predicate):

Evaluate predicate for all non-exited threads in mask and return an integer whose Nth bit is set if and only if predicate evaluates to non-zero for the Nth thread of the warp and the Nth thread is active.

__activemask():

Returns a 32-bit integer mask of all currently active threads in the calling warp. The Nth bit is set if the Nth lane in the warp is active when __activemask() is called. Inactive threads are represented by 0 bits in the returned mask. Threads which have exited the program are always marked as inactive. Note that threads that are convergent at an __activemask() call are not guaranteed to be convergent at subsequent instructions unless those instructions are synchronizing warp-builtin functions.

Notes

For __all_sync, __any_sync, and __ballot_sync, a mask must be passed that specifies the threads participating in the call. A bit, representing the thread's lane ID, must be set for each participating thread to ensure they are properly converged before the intrinsic is executed by the hardware. All active threads named in mask must execute the same intrinsic with the same mask, or the result is undefined.

本文备注/经验分享：

今天的章节是warp级别的一些数据交换操作.也就是所谓的warp vote系列内置函数.它们具有两个功能: (1)1-bit数据的交换 （2）1-bit数据的规约. 这些都是在warp的级别上完成的.请注意在7.0+的计算能力上, 它们还具有让warp一致执行步伐(临时)的效果，需要说明的是, warp vote不是你写程序的时候必须的,没有它们, 你依然可以写出逻辑完备的程序.但有了它们, 可以让你活的更舒服一点.(例如, 你总是可以通过shared memory进行数据交换, 然后手工进行统计. 这样费时费力(无论从写代码的人的角度, 或者从硬件执行的角度), 直接使用它们, 不仅仅写的时候省事了, 还不容易出错, 还能更快速的执行)。 在目前的版本上, CUDA C一共提供了3个(4个)相关的函数, 即: __all_sync, 也就是CUDA 9之前的无sync结尾的__all； __any_sync, 同理, 之前叫做__ayn， 以及__ballot_sync(), 以前叫做__ballot， 你可以看到它们从CUDA 9开始, 变得不兼容了, 之前的版本均没有尾巴. 现在却多了一个尾巴(warp同步用的后缀)， 以及, 还多了一个参数(第一个参数, 用来选择你需要那些warp内部的线程们(lanes)进行同步). 目前在新版本的CUDA上使用老写法(无后缀, 无第一个参数)会得到一个警告. 所以用户应当尽快升级代码到CUDA 9， 大致上, 以前版本的__all(p)等于现在的__all_sync(0xffff, p); 其他的另外几个函数替换办法类似. 这是对于新CUDA版本+老卡(Pascal或者更低计算能力)说的. 对于新卡, 请将0xffff部分替换成你需要的其他值. 然后我们说一下这些函数的大致用途. (1)__all_sync/__any_sync这两个可以用来快速的将1-bit的值在warp内部进行规约，每个线程提供1个bit的predicate(请注意这里的predicate参数是整数类型的, 硬件实际上需要的是1-bit类型(P寄存器), 编译器会在需要的时候自动为你转换(从常规寄存器中的值到P寄存器)),all版本判断它们是否全部为真，而any版本则判断它们是否至少有1个为真.很多代码需要这种操作的, 例如可以快速的判断warp整体(或者某些部分)是否需要进行某种操作,如果整体需要或者不需要, 可以快速的越过一些片段, 而不需要在片段内部逐步处理之类的.而ballot则将来自最多这32个线程的1-bit值拼凑起来,得到1个32-bit的值,(注意上面的1和32的顺序)，然后每个线程都会得到同样的32-bit值, 也就是说, 可以快速的交换得到邻近线程们的值.很多操作也需要这种的.例如一些前序和操作(prefix-sum, 一种很重要的操作, 国内部分公司总是叫这个为scan操作)，例如说, 我们可以通过__ballot_sync,配合__popc(), __clz()或者__ffs()之类的函数,在进行某些数据结构上的插入之类的操作的时候, 快速判断warp整体需要多少个空间, 而每个具体的warp内部的线程又在什么位置上需要操作.这种非常方便.再例如说, 像是昨天的章节中,所提到的快速聚合原子操作(1个block或者warp整体对同样的地址上进行原子操作),NV曾经推荐过的手工操作, 现在变成编译器自动展开进行了,也是利用了这点.先选出当前有效的线程(__activemask, 这个还没说道), __popc统计全部需要进行的原子操作数量,执行1次总体原子操作, 将原始值传播回来(shuffle), 然后继续大家上__ffs之类的确定自己的最终模拟原子操作后的位置. （2）本章节还提供了一个__activemask, 这个是新增的，大致等效于以前的__ballot(1)，实际上这是以前快速判断当前还有多少线程存活(warp内部), (例如在很多层的if或者while之类的嵌套里面)，不需要人工跟踪那些线程已经退出/结束了, 而那些线程还在执行.直接可以快速调用它得到. 这里提到的__popc()操作(从1个32-bit值中确定1的个数)，例如__popc配合__ballot_sync可以快速确定满足条件的线程的数量, 而不是mask(也叫掩码, 从每个1-bit代表1个lane而得来)，例如__ffs配合, 从__ballot_sync得来, 然后and掉小于或者大于自己的线程的bit位置的结果, 可以快速进行自己的位置统计(快速prefix-sum)——这些都是很有用的小技巧. 没有这些技巧， 你依然可以通过正常的shared memory操作+手工统计得到,但有了这些技巧可以活的更好一点.这章你可以完全不会的. 完全不影响你正常在GPU软件开发工作上班.但是有了, 你可以, 例如减少一些加班.

需要补充的事:(1)硬件实际上(5.X/6.X/7.X)总是将这些固定的规约操作和按位统计一体完成的. 也就是说, ballot和any/all这些, 实际上编译出来的都是基本一样的指令. 前者等于不要统计结果/规约结果的后者. 后者等于不要按位结果的前者. 但是两个结果的延迟稍微不同. 32-bit统计的值的结果先出来. 规约操作的值(P寄存器)的结果需要较长的延迟才能出来.但这个延迟往往不会造成困扰.因为比global memoy的访存之类的还是要快很多的(ballot操作的32-bit基本上和普通的加法的延迟差不多，规约统计的值的延迟大约和普通的判断(if之类)的结果出来延迟差不多)。均不需要过度忧虑(给某些总是喜欢忧虑的人说的) (2)还有另外一个没有公开的warp vote函数：__eq操作，该操作可以快速统计参与eq统计的线程们的结果是否一致.硬件支持该操作, 但没有导出到CUDA C中(幸运的是, 导出到了PTX).需要的用户可以单独按照本手册的说法, 能够通过内嵌PTX的方式来使用它.此外, 还需要说明的是warp vote这些, 在竞争对手A家的卡中, 都是免费操作.因为有个单独的每个CU(相当于SM, 叫法不同而已), 里面都有一个"标量运算单元"，你可以理解成每个CU里面除了SP外, 都赠送了一个免费的CPU(很弱不过)，这些操作都可以通过该免费CPU来完成.相比N卡, 等于0代价.但是很遗憾的是, 都7年过去了.至今AMD没有将它们导出到OpenCL C中.所以你依然只能看, 不能用.所以选择N卡是一个明智的选择. 再好的硬件, 不能给用户利用, 不能保护用户的智力投资, 等于是0.

有不明白的地方，请在本文后留言

或者在我们的技术论坛bbs.gpuworld.cn上发帖