CANN Ascend C SIMT协作组shfl_up函数

shfl_up【免费下载链接】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 训练系列产品不支持功能说明获取coalesced_group组内当前线程向前偏移delta的线程的数据。函数原型template typename T T shfl_up(T var, unsigned int delta) const参数说明表1参数说明参数名输入/输出描述var输入线程用于交换的输入操作数。支持的数据类型为half、int32_t、uint32_t、float、half2、int64_t、uint64_t、bfloat16_t、bfloat16x2_t。delta输入期望获取的var值所在线程在组内相对当前线程的向前偏移值。返回值说明coalesced_group组内当前线程向前偏移delta的线程输入的var值。若偏移后超出组范围返回当前线程输入的var值。约束说明无调用示例示例代码中的条件分支将一个warp中所有线程id是偶数的线程组成coalesced_group协作组组内各线程shfl_up接口返回结果如下图所示。图1shfl_up结果示意图 SIMT编程场景using namespace cooperative_groups; __global__ void simt_kernel(...) { ... if (threadIdx.x % 2 0) { coalesced_group active coalesced_threads(); uint32_t result active.shfl_up(threadIdx.x 100, 2); } ... }SIMD与SIMT混合编程场景using namespace cooperative_groups; __simt_vf__ inline void simt_kernel(...) { ... if (threadIdx.x % 2 0) { coalesced_group active coalesced_threads(); uint32_t result active.shfl_up(threadIdx.x 100, 2); } ... }【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考