DAY51：阅读Warp Shuffle Functions

B.15. Warp Shuffle Functions

__shfl_sync, __shfl_up_sync, __shfl_down_sync, and __shfl_xor_sync exchange a variable between threads within a warp.

Supported by devices of compute capability 3.x or higher.

Deprecation Notice: __shfl, __shfl_up, __shfl_down, and __shfl_xor have been deprecated as of CUDA 9.0.

B.15.1. Synopsis

T can be int, unsigned int, long, unsigned long, long long, unsigned long long, float or double. With the cuda_fp16.h header included, T can also be __half or __half2.

B.15.2. Description

The __shfl_sync() intrinsics permit exchanging of a variable between threads within a warp without use of shared memory. The exchange occurs simultaneously for all active threads within the warp (and named in mask), moving 4 or 8 bytes of data per thread depending on the type.

Threads within a warp are referred to as lanes, and may have an index between 0 and warpSize-1 (inclusive). Four source-lane addressing modes are supported:

• __shfl_sync()
• Direct copy from indexed lane
• __shfl_up_sync()
• Copy from a lane with lower ID relative to caller
• __shfl_down_sync()
• Copy from a lane with higher ID relative to caller
• __shfl_xor_sync()
• Copy from a lane based on bitwise XOR of own lane ID

Threads may only read data from another thread which is actively participating in the __shfl_sync() command. If the target thread is inactive, the retrieved value is undefined.

All of the __shfl_sync() intrinsics take an optional width parameter which alters the behavior of the intrinsic. width must have a value which is a power of 2; results are undefined if width is not a power of 2, or is a number greater than warpSize.

__shfl_sync() returns the value of var held by the thread whose ID is given by srcLane. If width is less than warpSize then each subsection of the warp behaves as a separate entity with a starting logical lane ID of 0. If srcLane is outside the range [0:width-1], the value returned corresponds to the value of var held by the srcLane modulo width (i.e. within the same subsection).

__shfl_up_sync() calculates a source lane ID by subtracting delta from the caller's lane ID. The value of var held by the resulting lane ID is returned: in effect, var is shifted up the warp bydelta lanes. If width is less than warpSize then each subsection of the warp behaves as a separate entity with a starting logical lane ID of 0. The source lane index will not wrap around the value of width, so effectively the lower delta lanes will be unchanged.

__shfl_down_sync() calculates a source lane ID by adding delta to the caller's lane ID. The value of var held by the resulting lane ID is returned: this has the effect of shifting var down the warp by delta lanes. If width is less than warpSize then each subsection of the warp behaves as a separate entity with a starting logical lane ID of 0. As for __shfl_up_sync(), the ID number of the source lane will not wrap around the value of width and so the upper delta lanes will remain unchanged.

__shfl_xor_sync() calculates a source line ID by performing a bitwise XOR of the caller's lane ID with laneMask: the value of var held by the resulting lane ID is returned. If width is less than warpSize then each group of width consecutive threads are able to access elements from earlier groups of threads, however if they attempt to access elements from later groups of threads their own value of var will be returned. This mode implements a butterfly addressing pattern such as is used in tree reduction and broadcast.

The new *_sync shfl intrinsics take in a mask indicating 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 non-exited threads named in mask must execute the same intrinsic with the same mask, or the result is undefined.

B.15.3. Return Value

All __shfl_sync() intrinsics return the 4-byte word referenced by var from the source lane ID as an unsigned integer. If the source lane ID is out of range or the source thread has exited, the calling thread's own var is returned.

B.15.4. Notes

Threads may only read data from another thread which is actively participating in the __shfl_sync() command. If the target thread is inactive, the retrieved value is undefined.

width must be a power-of-2 (i.e., 2, 4, 8, 16 or 32). Results are unspecified for other values.

B.15.5. Examples

B.15.5.1. Broadcast of a single value across a warp

B.15.5.2. Inclusive plus-scan across sub-partitions of 8 threads

B.15.5.3. Reduction across a warp

本文备注/经验分享：

今天的主要内容是warp shuffle, 这是一个从计算能力3.0(Kepler)开始提供的特性.如同之前的warp vote一样, 同样可以进行warp内部的数据交换.但有几个区别: (1)没有数据的规约处理功能. (2)交换的数据是32-bit的(4B), 比warp vote的1-bit要大很多. 较新版本的CUDA(例如CUDA 9+), 具有增强版本的warp shuffle功能,例如可以交换64-bit的数据, 此时编译器将自动为你拆分成2个32-bit的shfl指令,但并不排除将来的硬件, 会直接实现64-bit的shfl版本.用户也可以无视本章节, 因为和warp vote一样, 这不是一个必选的特性.用户可以完全无视warp shuffle, 而写出功能完备的代码.不使用warp shuffle的时候, 需要通过shared memory进行数据交换. 后者这种交换不如warp shuffle高效.换句话说, warp shuffle有如下特色(相比shared memory上的交换): (1)不需要为参与数据交换的warp(和内部的线程们--lanes)分配shared memory.这样可以减少shared memory的使用. 甚至有些代码, 使用满了48KB的每个block最大的shared memory, 此时无法继续分配空间进行数据交换使用.但依然可以进行warp shuffle.(请注意, 你也可以临时将shared memory中的内容交换到寄存器, 然后将空出来的shared memory用作数据交换. 然后交换完成后, 再将被破坏掉的shared memory中的内容保存回去. 但这样比较低效)。 所以完全不占用空间, 相比shared memory上的方式, 是一个很大的特色. (2)不需要计算地址, 用户可以直接按值目标线程的位置之类的信息(例如, lane id, lane id是warp中的分量线程的编号, 0-31一般)进行交换. 而传统的shared memory上的方式则需要先计算写入地址(位置, 例如线程编号 *4 + 基地址, 如果要每个线程交换4B大小的话), 然后写入，写入完成后需要继续计算读取地址(同上)，然后读取，这样的话需要较高的SP来进行辅助. 而shfl可以直接计算常见类型的地址, 等于免费了(例如向上shuffle 1个4B位置的时候, 该计算可以免费)。很多时候, 访问shared memory密集的代码并没有卡在shared memory上(可以通过profiler看), 而是卡在地址计算上. 这有的时候非常尴尬.而warp shuffle不仅仅节省了空间, 还节省了计算量. (3)点则是, shuffle的本质依然是使用shared memory, 它被编译后, 生成的指令被GPU执行的时候, 依然是提交一条指令给shared memory,只是shared memory不进行任何操作, 就地将数据按照一定的方式打乱(shuffle么. 类似洗牌), 然后直接返回.相比普通的shared memory上的1次写入+1次读取, shuffle可以一条指令内直接完成，这样就算SP上的地址计算不是一个瓶颈, 当shared memory成为瓶颈的时候, 通过shuffle可以等效的提高一倍效率. (4)点则是, 你现在不需要同步了. 以前通过shared memory上的交换, 除非使用locked-step的warp之类的 + volatile关键字之类的技巧, 你需要block同步的, 这样的代价就很大.而warp shuffle可以直接交换, 不需要block级别的同步, 也没有之前的技巧容易导致的BUG(例如用户忘记了volatile), 非常易用.正因为至少有这4方面的好处, 你能用warp shuffle的地址就应当使用. 除非你用不了(例如, 数据的交换需要在block大范围内, 而不是warp范围内, 则此时你只能走shared memory,这也是我们之前说过的shared memory的三大用途之一)。 回到具体的CUDA 9+, warp shuffle从这个版本起,引入了不兼容的改变(多了_sync后缀和需要warp内部同步的线程的掩码, 因为从计算能力7.0起, warp不一定必须完全步伐一致的执行). 对于新的CUDA 9.2 + 老卡(例如GTX1080), 用户可以指定0xffff之类的mask,(为了配合新硬件, 老硬件总会慢慢老去并从市面上消失的, 将来只会有新卡）此时将可以模拟原本的无后缀, 无第一个参数的老shuffle系列函数行为.或者用户依然可以维持老用法, 但会得到一个警告, 建议用户尽快迁移. 本章节所提供的shfl系列函数, 具有4个主要版本. 分别是up/down, 即目标线程的编号可以通过简单的加减法来确定的时候.还有一种是叫xor, 这个是将线程编号进行异或操作。另外一个则是直接索引(这个是没有后缀的那个), 用户可以直接要求指定一个线程编号.这4种. 其中的xor版本需要注意一下, 这也叫buffterfly型操作(蝴蝶)，因为蝴蝶的翅膀是类似X形状交错的。网上有一些教程, 对xor版本里面的线程编号变化情况有图,你会看到真的很像蝴蝶. 像是很多操作, 例如前序和，像是常见的一些规约操作，都可以全部或者部分通过shuffle来进行.手册上的一些例子, 包括本章节最后的:那个规约求和的例子, 可以适合用在一个长的block内部的规约的最后(当缩减到warp规模的时候).实际上手册之前这个例子有个对比的.是一个通过shared memory上的规约，后者版本需要多次的shared memory的读写.而warp shuffle版本没有这个要求.

继续补充一点:以前很多代码需要 #if __CUDA_ARCH < 300 进行shared memory上的空间分配, 以及, 后续的shared memory上的交换 #else 直接warp shuffle #endif

现在这种代码已经基本上不需要了. 这是为了Kepler之前的卡进行的. 现在从CUDA 9开始, 已经不再支持对这些卡进行开发了(但可以运行)， 用户看一些书的时候, 应当直接看#else的部分, 而不需要再考虑进行前者了. 此外, 还需要注意的是, 从计算能力3.0(最初支持shuffle的版本)到计算能力7.2， warp shuffle的速率均是32条/SM/clock， 但是因为实际上SM里面的SP数量是变化的. 例如从Kepler的192个, 到Maxwell/Pascal(6.1)的128个, 到Pascal(6.0)和GV100(7.0)的64个SP， 实际上warp shuffle的等效速率是从1/6到1/4到1/2发生变化的，换句话说, 越来越快了。用户应当考虑在不同的硬件上, 进行warp shuffle, 和就地直接重新计算数据(如果能通过直接计算算出来的话), 两者之间的代价权衡. 此外, 关于32这点, 也可以看出的确是走shared memory的,(shared memory这些年一直是32个banks，之前的32是按线程单位提供的---请注意profiler的指令单位是warp单位, 也就是1条/SM/周期)，该指令具有典型的shared memory的周期,当用户的代码卡shared memory操作,或者对延迟非常敏感而无法掩盖的时候, 应当注意shuffle本身的延迟.(比计算指令高不少, 而且可变)但是一般来说, 除非特殊情况, 使用shuffle总是有正面效果的. 大致这样.几个例子上的应用, 用户可以看一下. 都比较简单。

这里需要说一下. 所有的本章节的shfl函数都接受一个width参数, 可以在warp进步进一步的切分大小. 很多时候, 例如在邻近的16个线程或者4个线程之类的场合, 直接指定大小可以等效的减轻用户计算线程ID的计算量.此时则应当考虑使用. 此外, 为了未来着想, 用户当应当只需要warp内部的较小的范围的时候, 应当直接指定较小的范围, 例如4, 8, 16这种. 未来的一些硬件可能对这些有特殊的优化. 例如A家的一些硬件, 当在特定的sub-wave的范围内进行交换的时候, 可以通过某些特殊的设计, 直接将额外的指令执行消除, 同时减少了指令量, 也消除了延迟. 例如连续的16个wave内的某些交换操作, 可以直接不走LDS(等于N卡的shared memory), 直接0成本.虽然目前N卡没有对特定的范围情况进行优化处理。 但用户能这样写就应当保持这样写, 以取得未来的更好效果.此外, 几乎如同一个惯例,虽然A家的硬件如此优秀, 不仅仅提供了wave内的shuffle功能, 还提供了特殊情况下的硬件优化,但很遗憾的是, 无论是这种优化, 还是普通的shuffle功能, 至今依然不能使用，AMD至今没有将它们导出到OpenCL C种.因此虽然N卡的warp shuffle比AMD的wave shuffle更往后推出,但上来就可以用的。 选择N卡, 的确是你明智的选择. 小声说一句： CUDA C版本的warp shuffle虽然已经很强了, 但PTX版本的功能更强.PTX版本是双返回值的.除了CUDA C这里能返回交换后的数据外,还能返回是否真的参与了交换(例如因为越界)，需要额外功能的shuffle支持, 应当考虑PTX嵌入.

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

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