CANN/asc-devkit bfloat16x2倒数函数

📅 2026/7/17 8:17:05 👁️ 阅读次数 📝 编程学习
CANN/asc-devkit bfloat16x2倒数函数

h2rcp

【免费下载链接】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 训练系列产品:不支持

功能说明

获取输入数据x各元素的倒数。

函数原型

inline bfloat16x2_t h2rcp(bfloat16x2_t x)

参数说明

表1参数说明

参数名输入/输出描述
x输入源操作数,输入数据。

返回值说明

输入数据各元素的倒数。本接口受全局饱和模式影响,特殊值如下:

输入非饱和模式返回值饱和模式返回值
0infASCRT_MAX_NORMAL_BF16
-0-0-0
nannan0
inf00
-inf-0-0

约束说明

针对Ascend 950PR/Ascend 950DT,本接口不支持Subnormal场景:本接口内部实现使用到了除法运算符,由于除法运算符不支持Subnormal场景,当输入x为Subnormal数据时,会导致本接口最终结果为±inf。

需要包含的头文件

使用bfloat16x2_t类型接口需要包含"simt_api/asc_bf16.h"头文件。

#include "simt_api/asc_bf16.h"

调用示例

  • SIMT编程场景:

    __global__ __launch_bounds__(1024) void KernelRcp(bfloat16x2_t* dst, bfloat16x2_t* x) { int idx = threadIdx.x + blockIdx.x * blockDim.x; dst[idx] = h2rcp(x[idx]); }
  • SIMD与SIMT混合编程场景:

    __simt_vf__ __launch_bounds__(1024) inline void KernelRcp(__gm__ bfloat16x2_t* dst, __gm__ bfloat16x2_t* x) { int idx = threadIdx.x + blockIdx.x * blockDim.x; dst[idx] = h2rcp(x[idx]); }

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

创作声明:本文部分内容由AI辅助生成(AIGC),仅供参考