Warp 匹配函数：

__match_any_sync和__match_all_sync对warp内的线程之间的变量执行广播和比较操作。由计算能力7.x或更高的设备支持。Synopsys：

unsigned int __match_any_sync(unsigned mask, T value);unsigned int __match_all_sync(unsigned mask, T value, int    *pred)

T可以是int，unsigned int，long，unsigned long，long long，unsigned long long，float或double。

描述：__match_sync（）内部函数允许在同步线程中命名的线程之后，对warp中的线程之间的值进行广播和比较。__match_any_sync：返回掩码中具有相同值的线程的掩码。__match_all_sync:如果掩码中的所有线程都具有相同的值，则返回掩码; 否则返回0。 如果掩码中的所有线程都具有相同的值值，则Predicate pred将设置为true; 否则谓词设置为false。新的* _sync匹配内部函数接收指示参与调用的线程的掩码。 表示线程的通道ID的位必须为每个参与线程设置，以确保它们在内部由硬件执行之前正确收敛。 在掩码中命名的所有非退出线程都必须使用相同的掩码执行相同的内部属性，否则结果未定义。 Warp Shuffle函数：__shfl_sync，__shfl_up_sync，__shfl_down_sync和__shfl_xor_sync在变形内的线程之间交换变量。由计算能力3.x或更高的设备支持。弃用声明：自CUDA 9.0起，__shfl，__shfl_up，__shfl_down和__shfl_xor已被弃用。概要：

T __shfl_sync(unsigned mask, T var, int srcLane, int width = warpSize);T __shfl_up_sync(unsigned mask, T var, unsigned int delta, int width = warpSize);T __shfl_down_sync(unsigned mask, T var, unsigned int delta, int    width = warpSize);T __shfl_xor_sync(unsigned mask, T var, int laneMask, int width = warpSize)

T可以是int，unsigned int，long，unsigned long，long long，unsigned long long，float或double。 通过包含cuda_fp16.h头文件，T也可以是__half或__half2。

描述：__shfl_sync（）内部函数允许在不使用共享内存的情况下交换变形内的线程之间的变量。 对于warp中的所有活动线程（并在mask中命名），交换同时发生，根据类型，每个线程移动4或8个字节的数据。变形中的线程称为通道，可能具有介于0和warpSize-1（含）之间的索引。 支持四种源通道寻址模式：__shfl_sync()：从索引lane直接复制。__shfl_up_sync():从ID较低的lane复制相对于呼叫者。__shfl_down_sync():从ID较高的lane复制相对于呼叫者。__shfl_xor_sync()：基于自己lane ID的按位XOR从车道复制线程只能从另一个主动参与__shfl_sync（）命令的线程读取数据。 如果目标线程处于非活动状态，则检索到的值不确定。所有的__shfl_sync（）内部函数都有一个可选的宽度参数，它可以改变内部的行为。 宽度必须是2的幂的值; 如果宽度不是2的幂，或者是大于warpSize的数字，结果是不确定的。__shfl_sync（）返回由srcLane给出ID的线程持有的var的值。 如果宽度小于warpSize，那么warp的每个子部分都会作为一个开始逻辑通道ID为0的独立实体运行。如果srcLane超出范围[0：width-1]，则返回的值对应于var的值 由srcLane模宽度（即在同一小节内）。__shfl_up_sync（）通过从呼叫者的laneID减去delta来计算源lane ID。 返回结果lane标识所保存的var的值：实际上，var通过三角形lane向上移动变形。 如果宽度小于warpSize，那么warp的每个子部分都表现为一个开始逻辑通道ID为0的独立实体。源通道索引不会环绕宽度值，因此实际上较低delta通道将保持不变。__shfl_down_sync（）通过将delta添加到呼叫者的lane ID来计算源lane ID。 返回结果lane 标识所保存的var的值将被返回：这具有通过三角形lane 将变形向下移动var的效果。 如果宽度小于warpSize，那么warp的每个子部分都表现为一个独立的实体，其起始逻辑通道ID为0.对于__shfl_up_sync（），源通道的ID号不会环绕width的值，因此 上三角洲lane 将保持不变。__shfl_xor_sync（）通过调用laneMask执行呼叫者lane ID的按位XOR来计算源线路ID：返回由生成的lane ID保存的var值。 如果width小于warpSize，那么每组宽度连续的线程都可以访问先前线程组中的元素，但是如果它们尝试从后面的线程组访问元素，则会返回它们自己的var值。 该模式实现了一种蝴蝶寻址模式，例如用于树状缩减和广播。新的* _sync shfl内部函数接收一个指示参与调用的线程的掩码。 表示线程的通道ID的位必须为每个参与线程设置，以确保它们在内部由硬件执行之前正确收敛。 在掩码中命名的所有非退出线程都必须使用相同的掩码执行相同的内部属性，否则结果未定义。返回值：所有的__shfl_sync（）内在函数都会将源通道ID中由var引用的4字节字作为无符号整数返回。 如果源通道标识超出范围或源线程退出，则返回调用线程自己的变量。注意：线程只能从另一个主动参与__shfl_sync（）命令的线程读取数据。 如果目标线程处于非活动状态，则检索到的值不确定。宽度必须是2的幂（即2，4，8，16或32）。 其他值的结果未指定。