CANN/asc-devkit协作组shfl函数

shfl

【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言，原生支持C和C++标准规范，主要由类库和语言扩展层构成，提供多层级API，满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit

产品支持情况

功能说明

coalesced_group组内线程的数据交换接口，不通过共享内存实现直接读取组内指定线程的寄存器中的值。

函数原型

template <typename T> T shfl(T var, int src_rank) const

参数说明

表 1参数说明

返回值说明

coalesced_group组内指定线程输入的var值。

约束说明

无

调用示例

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

图 1shfl结果示意图

• SIMT编程场景：

  using namespace cooperative_groups; __global__ void simt_kernel(...) { ... if (threadIdx.x % 2 == 0) { coalesced_group active = coalesced_threads(); uint32_t result = active.shfl(threadIdx.x + 100, 3); } ... }

• 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(threadIdx.x + 100, 3); } ... }

【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言，原生支持C和C++标准规范，主要由类库和语言扩展层构成，提供多层级API，满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit