并行编程实战——CUDA编程的其它Warp函数-育师

一、说明

在前面将束内原语的Vote和Shuffle进行了分析和说明，基本明白了二者的功能和用途。这时候可能就会想到，会不会还有其它的束内函数呢？那自然是有的。下面将对其它的几个束内函数进行分析和说明，不过，重点只是进行功能的分析说明。更多的细节需要思考和查找相关资料（特别是官网上的文档说明）。

二、Wrap Reduce、Warp Match和Warp Matrix

在CUDA中，还提供了Wrap Reduce、Warp Match和Warp Matrix函数即线程束归约函数、线程束匹配函数和线程束矩阵运算操作函数。
Wrap Reduce用于在线程束内计算所有线程数据的单一聚合值，如极大极小值等。主要应用场景为并行归约、求和和极值运算等
Warp Match用于在线程束内查找与当前线程拥有相同值的其它线程，如进行快速的数据查找和比较，应用场景为数据的过滤、去重及构建索引等。
Warp Matrix则是利用GPU的Tensor Core硬件加速单元，在线程束内进行小规模的矩阵乘累加操作（D = A * B + C），核心应用就是硬件加速矩阵运算。应用场景为深度学习、科学计算中的矩阵运算等。
需要注意的是，后二者都需要GPU架构的支持，在实际应用中需要先确定硬件和CUDA版本的具体支持情况，发现问题要首先确认是否是硬件或软件版本支持的问题。

三、相关的函数接口

这三个束内函数包括:

//Warp 匹配函数unsigned__match_any_sync(unsignedmask,T value);unsigned__match_all_sync(unsignedmask,T value,int*pred);//Warp 归约函数T__reduce_add_sync(unsignedmask,T value);T__reduce_min_sync(unsignedmask,T value);T__reduce_max_sync(unsignedmask,T value);unsigned__reduce_and_sync(unsignedmask,unsignedvalue);unsigned__reduce_or_sync(unsignedmask,unsignedvalue);unsigned__reduce_xor_sync(unsignedmask,unsignedvalue);//Warp矩阵运算处理函数template<typename Use,intm,intn,intk,typename T,typename Layout=void>class fragment;voidload_matrix_sync(fragment<...>&a,constT*mptr,unsignedldm);voidload_matrix_sync(fragment<...>&a,constT*mptr,unsignedldm,layout_tlayout);voidstore_matrix_sync(T*mptr,constfragment<...>&a,unsignedldm,layout_tlayout);voidfill_fragment(fragment<...>&a,constT&v);voidmma_sync(fragment<...>&d,constfragment<...>&a,constfragment<...>&b,constfragment<...>&c,bool satf=false);

上面的接口相对来说看函数名称即可明白相关的操作。后面的矩阵操作进行一下简单的说明：
矩阵的运算处理只针对算力在7.0以上的设备上，它们都定义在nvcuda::wmma命名空间中。fragment是一个重载的类，包含分布在warp所有线程中的矩阵片段。矩阵元素到fragment内部存储的映射关系未指定，且在未来的架构中可能发生变化。它仅允许特定的模板参数组合。
load_matrix_sync和store_matrix_sync接口类似于C++中的原子操作中的相关操作，前者等待所有warp线程到达 load_matrix_sync，然后从内存加载矩阵片段；而后者等待所有warp线程到达 store_matrix_sync，然后将矩阵片段 a 存储到内存。
fill_fragment用常量值v填充矩阵片段。由于矩阵元素到每个片段的映射未指定，此函数通常由warp中所有线程以共同的v值调用。
mma_sync等待所有warp线程到达mma_sync，然后执行warp同步的矩阵乘加操作D = A * B + C。

四、Warp Active Mask

在开始分析束内函数时，就在官方文档中看到相关活动掩码控制的函数接口即：

unsigned__activemask();

其主要的功能是动态获取当前活动线程的掩码，配合前面的束内原语函数进行工作。但对它的使用要非常谨慎，因为新的GPU（如Volta架构）后，由于引入了独立线程调度，即线程的活跃状态可能随时发生变化，因为由此函数获取的掩码可能不再准确。通常还是建议直接在代码逻辑中指定掩码。

五、接口更新

在最新的NVIDIA13.1的官方文档中，在这些函数进行了不同程度的更新和警告说明。比如对于上面的Shuffle和匹配函数，推荐使用libcu++中的相应接口。如“libcu++中的cuda::device::warp_match_all()函数作为__match_all_sync函数”的安全且通用的替代方案。又比如推荐使用CUB库中的归约库函数等等。
对于矩阵的处理，则会不断的发展变化，所以在官方文档中说明了“Sub-byte operations are considered preview, i.e. the data structures and APIs for them are subject to change and may not be compatible with future releases”意思说对子字节（低于一字节精度如int4,int1等）的矩阵运算在本阶段为预览版，在未来有可能发生变化并不保证兼容。
这些在新版本的开发中，需要明白这些最新的变化。

六、例程

下面看一个例程：

#include"cuda_runtime.h"#include"device_launch_parameters.h"#include<stdio.h>__global__voidWarpReduceDemo(){inttid=threadIdx.x;intretAdd=__reduce_add_sync(0xffffffff,tid);intretMin=__reduce_min_sync(0xffffffff,tid);intretMax=__reduce_max_sync(0xffffffff,tid);unsignedintopTid=tid&1;unsignedintretAnd=__reduce_and_sync(0xffffffff,opTid);unsignedintretOr=__reduce_or_sync(0xffffffff,opTid);unsignedintretXor=__reduce_xor_sync(0xffffffff,opTid);printf("threadId: %d reduce add: %d reduce min: %d reduce max: %d reduce and: %x reduce or: %x reduce xor: %x\n",threadIdx.x,retAdd,retMin,retMax,retAnd,retOr,retXor);}intmain(){WarpReduceDemo<<<1,32>>>();return0;}

由于环境中支持的架构算力太低，GTX960是5.2，所以默认的编译是无法通过的。利用在前面的CDP动态并行性中的方法将compute_52,sm_52修改为compute_85,sm85编译报不支持此架构，再次修改为compute_80,sm_80，编译通过，但运行无结果。这也符合前面提到的，如果硬件支持不够，则行为未定义。这也意味着，编译与实际运行是两个范畴。开发者可以自行指定编译的算力大小，从而决定运行的环境。
大家可以根据自己情况进行处理。