__byte_perm【免费下载链接】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 训练系列产品不支持功能说明由输入的两个4字节的uint32_t类型数据组成一个8个字节的64比特位的整数通过选择器s指定选取其中的4个字节将这4个字节从低位到高位拼成一个uint32_t类型的整数。具体实现逻辑如下// 以下为C表示的BytePerm(x, y, s)计算逻辑 uint64_t tmp64 ((uint64_t)y 32) | x; // x,y拼接成uint64整数 uint8_t selector0 (s 0) 0x7; // selector0取值范围在[0, 7]之间 uint8_t selector1 (s 4) 0x7; uint8_t selector2 (s 8) 0x7; uint8_t selector3 (s 12) 0x7; uint8_t byte0 (tmp64 (selector0 * 8)) 0xFF; // 选取tmp64中的第selector0个字节 uint8_t byte1 (tmp64 (selector1 * 8)) 0xFF; uint8_t byte2 (tmp64 (selector2 * 8)) 0xFF; uint8_t byte3 (tmp64 (selector3 * 8)) 0xFF; // result为BytePerm返回值其结果由对应字节按照顺序拼接而成 uint32_t result byte0 | (byte1 8) | (byte2 16) | (byte3 24);函数原型unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s)参数说明表1参数说明参数名输入/输出描述x输入源操作数uint32_t类型与y拼接成64比特位的整数该整数的[0:31]位为x。y输入源操作数uint32_t类型与x拼接成64比特位的整数该整数的[32:63]位为y。s输入选择器uint32_t类型用于指定如何从x和y组成的8个字节64比特位的整数中提取4字节数据。具体为s[0:3]、s[4:7]、s[8:11]、s[12:15]表示的数据指定选取的字节在8个字节整数中的索引值0到7。返回值说明通过选择器s选取出的uint32_t类型的整数。当x为0y为0s为0时返回值为0。当x为1y为1s为1时返回值为16843008。约束说明无需要包含的头文件使用该接口需要包含simt_api/device_functions.h头文件。#include simt_api/device_functions.h调用示例SIMT编程场景__global__ __launch_bounds__(1024) void KernelByte_perm(unsigned int* dst, unsigned int* x, unsigned int* y, unsigned int* s) { int idx threadIdx.x blockIdx.x * blockDim.x; dst[idx] __byte_perm(x[idx], y[idx], s[idx]); }SIMD与SIMT混合编程场景__simt_vf__ __launch_bounds__(1024) inline void KernelByte_perm(__gm__ unsigned int* dst, __gm__ unsigned int* x, __gm__ unsigned int* y, __gm__ unsigned int* s) { int idx threadIdx.x blockIdx.x * blockDim.x; dst[idx] __byte_perm(x[idx], y[idx], s[idx]); }【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考