cann/asc-devkit bfloat16除法函数

__hdiv

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

功能说明

计算两个bfloat16类型数据相除的结果,并遵循CAST_RINT模式对结果进行舍入。

函数原型

bfloat16_t __hdiv(const bfloat16_t x, const bfloat16_t y)

参数说明

表1参数说明

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

返回值说明

输入数据相除的结果。本接口受全局饱和模式影响,特殊值如下:

x值y值非饱和模式返回值饱和模式返回值
±0±0nan0
±inf±infnan0
有限值±inf符号由x和y的符号异或决定,值为0符号由x和y的符号异或决定,值为0
inf正有限值infASCRT_MAX_NORMAL_BF16
inf负有限值-inf-ASCRT_MAX_NORMAL_BF16
-inf正有限值-inf-ASCRT_MAX_NORMAL_BF16
-inf负有限值infASCRT_MAX_NORMAL_BF16
>00infASCRT_MAX_NORMAL_BF16
<00-inf-ASCRT_MAX_NORMAL_BF16
>0-0-inf-ASCRT_MAX_NORMAL_BF16
<0-0infASCRT_MAX_NORMAL_BF16
±0非0符号由x和y的符号异或决定,值为0符号由x和y的符号异或决定,值为0
x,y任意一个为nannan0

约束说明

针对Ascend 950PR/Ascend 950DT,本接口不支持Subnormal场景:处于Subnormal范围内的输入和输出值,都会被刷新为保留符号的0。

需要包含的头文件

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

#include "simt_api/asc_bf16.h"

调用示例

  • SIMT编程场景:

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

    __simt_vf__ __launch_bounds__(1024) inline void KernelHdiv(__gm__ bfloat16_t* dst, __gm__ bfloat16_t* x, __gm__ bfloat16_t* y) { int idx = threadIdx.x + blockIdx.x * blockDim.x; dst[idx] = __hdiv(x[idx], y[idx]); }

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

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