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参数说明

返回值说明

coalesced_group组内当前线程向前偏移delta的线程输入的var值。若偏移后超出组范围，返回当前线程输入的var值。

约束说明

无

调用示例

示例代码中的条件分支将一个warp中所有线程id是偶数的线程组成coalesced_group协作组，组内各线程shfl_up接口返回结果如下图所示。

图1shfl_up结果示意图 ![](https://raw.gitcode.com/cann/asc-devkit/raw/1680a8274bf627fb844cc7e09f9e1a9e99e24425/docs/zh/api/figures/coalesced_group_shfl_up.png "shfl_up结果示意图"?utm_source=gitcode_repo_files)

• 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