__hfma

Applicability

Function Usage

Computes the result of x multiplied by y plus z for the input data x, y, and z.

Prototype

1

inline half __hfma(half x, half y, half z)

1

inline bfloat16_t __hfma(bfloat16_t x, bfloat16_t y, bfloat16_t z)

Parameters

Returns

Value of x * y + z.

• If x is ±inf and y is ±0, the return value is nan.
• If x is ±0 and y is ±inf, the return value is nan.
• If x * y is inf and z is -inf, the return value is nan.
• If x * y is -inf and z is inf, the return value is nan.
• If x * y + z exceeds the maximum value of the corresponding type, the return value is inf.
• If x * y + z is less than the minimum value of the corresponding type, the return value is -inf.
• If any of x, y, or z is nan, the return value is nan.

Restrictions

For SIMT programming, this API is not supported.

Header File to Be Included

To use the half API, the simt_api/asc_fp16.h header file must be included. To use the bfloat16_t API, the simt_api/asc_bf16.h header file must be included.

1

#include "simt_api/asc_fp16.h"

1

#include "simt_api/asc_bf16.h"

Examples

1
2
3
4

__simt_vf__ __launch_bounds__(1024) inline void KernelFma(__gm__ half* dst, __gm__ half* x, __gm__ half* y, __gm__ half* z){
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    dst[idx] = __hfma(x[idx], y[idx], z[idx]);
}