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 训练系列产品：不支持
功能说明
获取coalesced_group组内当前线程向后偏移delta的线程的数据。
函数原型
template <typename T> T shfl_down(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_down接口返回结果如下图所示。
图1shfl_down结果示意图
![](https://raw.gitcode.com/cann/asc-devkit/raw/1680a8274bf627fb844cc7e09f9e1a9e99e24425/docs/zh/api/figures/coalesced_group_shfl_down.png "shfl_down结果示意图"?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_down(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_down(threadIdx.x + 100, 2); } ... }
【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言，原生支持C和C++标准规范，主要由类库和语言扩展层构成，提供多层级API，满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit
创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
