CANN/asc-devkit SIMT线程数据交换

shfl【免费下载链接】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 训练系列产品不支持功能说明thread_block_tile组内线程的数据交换接口不通过共享内存实现直接读取组内指定线程的寄存器中的值。函数原型template typename T T shfl(T var, int src_rank) const参数说明表1参数说明参数名输入/输出描述var输入线程用于交换的输入操作数。支持的数据类型为half、int32_t、uint32_t、float、half2、int64_t、uint64_t、bfloat16_t、bfloat16x2_t。src_rank输入期望获取的var值所在的线程在组内的排名。当src_rank大于等于组内线程数时获取src_rank % Size对应rank线程的var值。返回值说明thread_block_tile组内指定线程输入的var值。约束说明无调用示例示例代码中以4个线程为一组划分线程块获取组内排名为2的线程输入的var值。图1shfl接口返回结果示意图SIMT编程场景using namespace cooperative_groups; __global__ void simt_kernel(...) { ... thread_block block this_thread_block(); auto tile4 tiled_partition4(block); uint32_t result tile4.shfl(threadIdx.x 100, 2); ... }SIMD与SIMT混合编程场景using namespace cooperative_groups; __simt_vf__ inline void simt_kernel(...) { ... thread_block block this_thread_block(); auto tile4 tiled_partition4(block); uint32_t result tile4.shfl(threadIdx.x 100, 2); ... }【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考