CANN/asc-devkit Ascend C SIMT Warp-Shfl函数文档

asc_shfl_down【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit产品支持情况Ascend 950PR/Ascend 950DT支持Atlas A3 训练系列产品/Atlas A3 推理系列产品不支持Atlas A2 训练系列产品/Atlas A2 推理系列产品不支持Atlas 200I/500 A2 推理产品不支持Atlas 推理系列产品AI Core不支持Atlas 推理系列产品Vector Core不支持Atlas 训练系列产品不支持功能说明获取Warp内当前线程向后偏移delta当前线程Lane ID delta的线程输入的用于交换的var值如果目标线程是非活跃状态获取到寄存器中未初始化的值。其中参数width用于划分Warp内线程的分组。参数width设置参与交换的32个线程的分组宽度默认值为32即所有线程分为1组。在多个分组场景width小于32下每个分组交换操作是独立的每个线程获取本组内当前线程向后偏移delta的线程的var值。如果当前线程向后偏移delta的线程编号即Lane ID delta大于所在分组的最大Lane ID则返回当前线程的var值。例如Warp内32个活跃线程调用asc_shfl_down(LaneId, 2, 16)接口每个线程的返回值为当前线程Lane ID 2对应线程的var值或者当前线程的var值。图1asc_shfl_down结果示意图函数原型inline int32_t asc_shfl_down(int32_t var, uint32_t delta, int32_t width warpSize)inline uint32_t asc_shfl_down(uint32_t var, uint32_t delta, int32_t width warpSize)inline float asc_shfl_down(float var, uint32_t delta, int32_t width warpSize)inline int64_t asc_shfl_down(int64_t var, uint32_t delta, int32_t width warpSize)inline uint64_t asc_shfl_down(uint64_t var, uint32_t delta, int32_t width warpSize)inline half asc_shfl_down(half var, uint32_t delta, int32_t width warpSize)inline half2 asc_shfl_down(half2 var, uint32_t delta, int32_t width warpSize)inline bfloat16_t asc_shfl_down(bfloat16_t var, uint32_t delta, int32_t width warpSize)inline bfloat16x2_t asc_shfl_down(bfloat16x2_t var, uint32_t delta, int32_t width warpSize)参数说明表1参数说明参数名输入/输出描述var输入线程用于交换的输入操作数。delta输入期望获取的var值所在线程相对当前线程的向后偏移值。width输入Warp内参与交换的线程的分组宽度默认值为32。width的取值范围为(0, 32]width必须是2的倍数。返回值说明Warp内指定线程的var值。约束说明如果目标线程是非活跃状态获取到寄存器中未初始化的值。若delta width所有线程都返回当前线程的var值。width不是2的倍数或者超出32返回值异常。需要包含的头文件使用除half、half2、bfloat16_t、bfloat16x2_t类型之外的接口需要包含simt_api/device_warp_functions.h头文件使用half和half2类型接口需要包含simt_api/asc_fp16.h头文件使用bfloat16_t和bfloat16x2_t类型接口需要包含simt_api/asc_bf16.h头文件。#include simt_api/device_warp_functions.h#include simt_api/asc_fp16.h#include simt_api/asc_bf16.h调用示例完整样例请参考Sobel边缘检测样例。以下示例包含两类用法示例1使用asc_shfl_down获取Warp分组内相对当前线程向后偏移delta的线程的输入值示例2使用asc_shfl_down在每个Warp内进行归约求和。示例1SIMT编程场景__global__ __launch_bounds__(1024) void KernelShflDown(int32_t* dst) { int idx threadIdx.x blockIdx.x * blockDim.x; int32_t laneId idx % 32; // 0-15线程返回值分别为{2,3,4,5,6,7,8,9,10,11,12,13,14,15,14,15} // 16-31线程返回值为{18,19,20,21,22,23,24,25,26,27,28,29,30,31,30,31} int32_t result asc_shfl_down(laneId, 2, 16); dst[idx] result; }SIMD与SIMT混合编程场景__simt_vf__ __launch_bounds__(1024) inline void KernelShflDown(__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线程返回值分别为{2,3,4,5,6,7,8,9,10,11,12,13,14,15,14,15} // 16-31线程返回值为{18,19,20,21,22,23,24,25,26,27,28,29,30,31,30,31} int32_t result asc_shfl_down(laneId, 2, 16); dst[idx] result; }示例2SIMT编程场景__global__ __launch_bounds__(1024) void KernelShflDownReduceSum(int32_t* dst) { int idx threadIdx.x blockIdx.x * blockDim.x; int32_t laneId idx % 32; int32_t value laneId; value asc_shfl_down(value, 16, 32); value asc_shfl_down(value, 8, 32); value asc_shfl_down(value, 4, 32); value asc_shfl_down(value, 2, 32); value asc_shfl_down(value, 1, 32); dst[idx] value; // 归约求和结果位于每个Warp内laneId为0的线程 }SIMD与SIMT混合编程场景__simt_vf__ __launch_bounds__(1024) inline void KernelShflDownReduceSum(__gm__ int32_t* dst) { int idx threadIdx.x blockIdx.x * blockDim.x; int32_t laneId idx % 32; int32_t value laneId; value asc_shfl_down(value, 16, 32); value asc_shfl_down(value, 8, 32); value asc_shfl_down(value, 4, 32); value asc_shfl_down(value, 2, 32); value asc_shfl_down(value, 1, 32); dst[idx] value; // 归约求和结果位于每个Warp内laneId为0的线程 }【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考