asc_shfl【免费下载链接】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 推理系列产品xAtlas 200I/500 A2 推理产品xAtlas 推理系列产品AI CorexAtlas 推理系列产品Vector CorexAtlas 训练系列产品x功能说明Warp Shfl类接口主要实现Warp级数据交换能够实现直接读取某个线程的数据而不需要通过共享内存。这类接口主要通过Warp分组实现组内线程间的数据交换操作。Warp分组Warp内的线程可分为多个组用户通过参数width配置分组宽度分组的线程数分组内的线程可进行数据交换组内线程通过相对组内起始线程位置来标识索引称为逻辑LaneId。数据交换本接口主要是获取分组内指定线程持有的var值用户通过参数src_lane指定线程。如果src_lane大于等于width指定线程的逻辑LaneId是src_lane%width。主要使用场景数据分发将固定位置的线程数据广播给其他线程动态数据交换每个线程从不同的源线程读取数据例如Warp内32个活跃线程调用asc_shfl(LaneId, 5, 16)接口每个线程的返回值为当前线程所在分组内线程编号为5的var值。图 1asc_shfl结果示意图 函数原型inline int32_t asc_shfl(int32_t var, int32_t src_lane, int32_t width warpSize)inline uint32_t asc_shfl(uint32_t var, int32_t src_lane, int32_t width warpSize)inline float asc_shfl(float var, int32_t src_lane, int32_t width warpSize)inline int64_t asc_shfl(int64_t var, int32_t src_lane, int32_t width warpSize)inline uint64_t asc_shfl(uint64_t var, int32_t src_lane, int32_t width warpSize)inline half asc_shfl(half var, int32_t src_lane, int32_t width warpSize)inline half2 asc_shfl(half2 var, int32_t src_lane, int32_t width warpSize)参数说明表 1参数说明参数名输入/输出描述var输入线程用于交换的输入操作数。src_lane输入期望获取的var值所在线程的LaneId。width输入Warp内参与交换的线程的分组宽度默认值为32。width的取值范围为(0, 32]width必须是2的倍数。返回值说明Warp内指定线程的var值。约束说明如果目标线程是非活跃状态获取到寄存器中未初始化的值。若入参width不是2的倍数或超出32返回值异常。需要包含的头文件使用除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编程场景__global__ __launch_bounds__(1024) void KernelShfl(int32_t* dst) { int idx threadIdx.x blockIdx.x * blockDim.x; int32_t laneId idx % 32; // 0-15线程返回值为116-31线程返回值为17 int32_t result asc_shfl(laneId, 1, 16); dst[idx] result; }SIMD与SIMT混合编程场景__simt_vf__ __launch_bounds__(1024) void KernelShfl(__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线程返回值为116-31线程返回值为17 int32_t result asc_shfl(laneId, 1, 16); dst[idx] result; }【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
