penCL 2.0 通过扩展cl_khr_subgroups提供一些基础子组操作支持包括获取子组 ID、组内 ID 等基本功能组内断言(any/all)、广播(broadcast)、归约(reduce)、扫描(scan)等基本操作同时允许一些可选扩展支持更丰富的子组操作比如洗牌(shuffle)、投票(ballot)等。其中子组内洗牌用于直接交换子组内寄存器值而不借助全局或本地内存类比 CUDA 束内洗牌经常用于一些算法优化是比较有用的子组函数但是大多数 GPU 都不支持官方的cl_khr_subgroup_shuffle扩展。当然一些厂商也可以通过厂商扩展、内联汇编(Inline Assembly)、内置函数(Intrinsics/Builtins)等途径支持子组洗牌。Nvidia通过内置函数__nvvm_shfl_mode_type(sm_70 以前)、__nvvm_shfl_sync_mode_type(sm_70, ptx60[1])支持。CUDA 中的线程束洗牌原语也是通过这些内置函数实现的。内置函数功能gentype __nvvm_shfl_mode_type(gentype val, int offset, int clamp)按照mode定义规则将offset处车道的val传递到本车道对clamp内的目标车道有效gentype __nvvm_shfl_sync_mode_type(uint mask, gentype val, int offset, int clamp)按照mode定义规则将offset处车道的val传递到本车道并按照mask同步线程束对clamp内的目标车道有效其中type可取i32、f32对应gentype为int、float。其中mode取值如下mode目标车道对标 OpenCL 子组函数对应 OpenCL 子组扩展idxoffsetsub_group_shufflecl_khr_subgroup_shufflebflylaneid ^ offsetsub_group_shuffle_xorcl_khr_subgroup_shuffleuplaneid - offsetsub_group_shuffle_upcl_khr_subgroup_shuffle_relativedownlaneid offsetsub_group_shuffle_downcl_khr_subgroup_shuffle_relativeNvidia 的 OpenCL 驱动支持 PTX 内联汇编因此可直接调用shfl.mode.type和shfl.sync.mode.type等汇编指令。Intel提供扩展cl_intel_subgroups支持多种子组洗牌操作。目前 Intel 是对 OpenCL 支持比较好的厂商该扩展自 OpenCL 1.2 加入并且直接实现了多种子组洗牌、投票函数。扩展函数功能gentype intel_sub_group_shuffle(gentype data, uint sub_group_local_id)将sub_group_local_id子组工作项的数据值传递给本子组工作项。gentype intel_sub_group_shuffle_down(gentype curr, gentype, gentype next, uint delta)将sub_group_local_iddelta作为目的子组索引。若子组索引大于等于 0 且小于max_sub_group_size则传递curr。若子组索引大于等于max_sub_group_size且小于max_sub_group_size* 2则将目的子组索引减max_sub_group_size传递next。gentype intel_sub_group_shuffle_up(gentype curr, gentype, gentype next, uint delta)将sub_group_local_id-delta作为目的子组索引。若子组索引大于等于 0 且小于max_sub_group_size则传递curr。若子组索引大于等于负max_sub_group_size且小于 0则将目的子组索引加max_sub_group_size传递next。gentype intel_sub_group_shuffle_xor(gentype data, uint val)将sub_group_local_idXORval子组工作项的数据值传递给本子组工作项。AMD ROCm有前缀为__builtin_amdgcn_的内置函数ds_permute、ds_bpermute前者将数据从当前车道推送(Push)到目标车道而后者从目标车道拉取(Pull)到当前车道与 OpenCL 中的sub_group_shuffle定义类似。内置函数功能int __builtin_amdgcn_ds_permute(int offset, int val)将本车道的数据值传输到索引为offset的车道int __builtin_amdgcn_ds_bpermute(int offset, int val)将索引为offset的车道的数据值传输到本车道另外AMD ROCm 还提供ds_swizzle(字节单位交换)、dpp_mov(数据并行移动)等内部函数实现更丰富的子组洗牌操作。AMD ROCm 的 OpenCL 驱动同样支持内联汇编可以直接调用ds.permute、ds.bpermute等 AMDGCN 指令。AMD PAL/HSAILROCm 不支持 Windows 集显其默认驱动为 AMD PAL/HSAIL。AMD PAL/HSAIL 提供内置函数gentype __hsail_activelanepermute_wavewidth_type(gentype src, uint lid, gentype ival, bool useival)实现活跃车道交换不同步[2]。该函数将目的索引lid的车道的数据值交换到本车道。若lid非法当useival为真时返回ival否则未定义。其中w可取b1、b32、b64表示该指令进行交换的数据宽度对应gentype为uchar、uint、ulong。使用__hsail_activelanepermute_wavewidth_type时需要在代码中声明该函数否则编译时会报错找不到符号。参见HSAFoundation/builtins[3]。另外sub_group_broadcast是用该指令与__hsail_wavebarrier()实现的因此也可用子组广播进行洗牌。Qualcommon Adreno提供扩展cl_qcom_subgroup_shuffle支持子组洗牌操作允许全组洗牌、宽 4 洗牌、宽 8 洗牌。扩展函数gentype qcom_sub_group_shuffle_op(gentype source_value,uint offset,qcom_sub_group_shuffle_width_modes_t width,gentype default_value);其中gentype可为uchar、char、ushort、short、uint、int、ulong、long以及float如果启用cl_khr_fp16扩展则也可取half。其中op可为up、down、rotate_up、rotate_down、xor。qcom_sub_group_shuffle_width_modes_t为枚举类型取值如下子组洗牌宽度模式含义CLK_SUB_GROUP_SHUFFLE_WIDTH_WAVE_SIZE_QCOM子组内所有工作项参与洗牌CLK_SUB_GROUP_SHUFFLE_WIDTH_W4_QCOM子组内每4个工作项之间洗牌CLK_SUB_GROUP_SHUFFLE_WIDTH_W8_QCOM子组内每8个工作项之间洗牌