CANN/asc-devkit：asc_shfl_xor线程交换API

asc_shfl_xor【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit产品支持情况产品是否支持Ascend 950PR/Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品xAtlas A2 训练系列产品 / Atlas A2 推理系列产品x功能说明获取Warp内当前线程LaneId与输入lane_mask做异或操作LaneId^lane_mask得到的dstLaneId对应线程输入的用于交换的var值如果目标线程是非活跃状态获取到寄存器中未初始化的值。其中参数width用于划分Warp内线程的分组。参数width设置参与交换的32个线程的分组宽度默认值为32即所有线程分为1组。在多个分组场景width小于32下每个线程获取位于本组内或线程编号更小的组内的dstLaneId对应线程的var值也就是说如果dstLaneId小于当前线程所在分组的起始LaneIddstLaneId对应的线程位于线程编号更小的组内则可以获取该dstLaneId线程的var值如果dstLaneId大于当前线程所在分组的最大LaneId则返回当前线程的var值。例如Warp内32个活跃线程调用asc_shfl_xor(LaneId, 1, 16)接口每个线程的返回值为当前线程LaneId^1对应线程的var值。图 1asc_shfl_xor结果示意图函数原型__simt_callee__ inline int32_t asc_shfl_xor(int32_t var, int32_t lane_mask, int32_t width warpSize)__simt_callee__ inline uint32_t asc_shfl_xor(uint32_t var, int32_t lane_mask, int32_t width warpSize)__simt_callee__ inline float asc_shfl_xor(float var, int32_t lane_mask, int32_t width warpSize)__simt_callee__ inline int64_t asc_shfl_xor(int64_t var, int32_t lane_mask, int32_t width warpSize)__simt_callee__ inline uint64_t asc_shfl_xor(uint64_t var, int32_t lane_mask, int32_t width warpSize)__simt_callee__ inline half asc_shfl_xor(half var, int32_t lane_mask, int32_t width warpSize)__simt_callee__ inline half2 asc_shfl_xor(half2 var, int32_t lane_mask, int32_t width warpSize)参数说明表 1参数说明参数名输入/输出描述var输入线程用于交换的输入操作数。lane_mask输入与当前线程LaneId做异或运算的操作数。取值范围为[0, 32)且小于width。width输入Warp内参与交换的线程的分组宽度默认值为32。width的取值范围为(0, 32]width必须是2的倍数。返回值说明Warp内指定线程的var值未初始化undefined的值约束说明无需要包含的头文件使用除half、half2类型之外的接口需要包含simt_api/device_warp_functions.h头文件使用half和half2类型接口需要包含simt_api/asc_fp16.h头文件。#include simt_api/device_warp_functions.h#include simt_api/asc_fp16.h调用示例__simt_vf__ __launch_bounds__(1024) inline void KernelShflXor(__gm__ int32_t* dst) { // asc_vf_call参数dim3{1024, 1, 1} int idx threadIdx.x blockIdx.x * blockDim.x; int32_t laneId idx % 32; // 0-15线程返回值分别为{1,0,3,2,5,4,7,6,9,8,11,10,13,12,15,14} // 16-31线程返回值为{17,16,19,18,21,20,23,22,25,24,27,26,29,28,31,30} int32_t result asc_shfl_xor(laneId, 1, 16); dst[idx] result; } // asc_shfl_xor实现reducesum __simt_vf__ __launch_bounds__(1024) inline void KernelShflXorReduceSum(__gm__ int32_t* dst) { int idx threadIdx.x blockIdx.x * blockDim.x; int32_t laneId idx % 32; int32_t value laneId; value asc_shfl_xor(value, 1, 31); // 1 value asc_shfl_xor(value, 2, 31); // 2 value asc_shfl_xor(value, 4, 31); // 4 value asc_shfl_xor(value, 8, 31); // 8 value asc_shfl_xor(value, 16, 31); // 16 dst[idx] value; }【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考