__hfma
产品支持情况
产品 |
是否支持 |
|---|---|
Atlas 350 加速卡 |
√ |
x |
|
x |
|
x |
|
x |
|
x |
|
x |
功能说明
对输入数据x、y、z,计算x与y相乘加上z的结果。

函数原型
1 | inline bfloat16_t __hfma(bfloat16_t x, bfloat16_t y, bfloat16_t z) |
参数说明
参数名 |
输入/输出 |
描述 |
|---|---|---|
x |
输入 |
源操作数。 |
y |
输入 |
源操作数。 |
z |
输入 |
源操作数。 |
返回值说明
x * y+ z的值。
- x为±inf,y为±0,返回nan。
- x为±0,y为±inf,返回nan。
- x*y为inf,z为-inf,返回nan。
- x*y为-inf,z为inf,返回nan。
- x*y+z超出对应类型范围的最大值,返回inf。
- x*y+z小于对应类型范围的最小值,返回-inf。
- x、y、z任意一个为nan,返回nan。
约束说明
无
需要包含的头文件
使用bfloat16_t类型接口需要包含"simt_api/asc_bf16.h"头文件。
1 | #include "simt_api/asc_bf16.h" |
调用示例
- SIMT编程场景:
1 2 3 4
__global__ __launch_bounds__(1024) void KernelFma(bfloat16_t* dst, bfloat16_t* x, bfloat16_t* y, bfloat16_t* z){ int idx = threadIdx.x + blockIdx.x * blockDim.x; dst[idx] = __hfma(x[idx], y[idx], z[idx]); }
- SIMD与SIMT混合编程场景:
1 2 3 4
__simt_vf__ __launch_bounds__(1024) inline void KernelFma(__gm__ bfloat16_t* dst, __gm__ bfloat16_t* x, __gm__ bfloat16_t* y, __gm__ bfloat16_t* z){ int idx = threadIdx.x + blockIdx.x * blockDim.x; dst[idx] = __hfma(x[idx], y[idx], z[idx]); }
父主题: bfloat16类型算术函数