__hdiv

产品支持情况

产品	是否支持
Atlas 350 加速卡	√
Atlas A3 训练系列产品/Atlas A3 推理系列产品	x
Atlas A2 训练系列产品/Atlas A2 推理系列产品	x
Atlas 200I/500 A2 推理产品	x
Atlas 推理系列产品AI Core	x
Atlas 推理系列产品Vector Core	x
Atlas 训练系列产品	x

功能说明

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

函数原型

bfloat16_t __hdiv(const bfloat16_t x, const bfloat16_t y)

参数说明

表1 参数说明
参数名	输入/输出	描述
x	输入	源操作数。
y	输入	源操作数。

返回值说明

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

x值	y值	非饱和模式返回值	饱和模式返回值
±0	±0	nan	0
±inf	±inf	nan	0
有限值	±inf	符号由x和y的符号异或决定，值为0	符号由x和y的符号异或决定，值为0
inf	正有限值	inf	ASCRT_MAX_NORMAL_BF16
inf	负有限值	-inf	-ASCRT_MAX_NORMAL_BF16
-inf	正有限值	-inf	-ASCRT_MAX_NORMAL_BF16
-inf	负有限值	inf	ASCRT_MAX_NORMAL_BF16
>0	0	inf	ASCRT_MAX_NORMAL_BF16
<0	0	-inf	-ASCRT_MAX_NORMAL_BF16
>0	-0	-inf	-ASCRT_MAX_NORMAL_BF16
<0	-0	inf	ASCRT_MAX_NORMAL_BF16
±0	非0	符号由x和y的符号异或决定，值为0	符号由x和y的符号异或决定，值为0
x，y任意一个为nan		nan	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]);
}

父主题： bfloat16类型算术函数