MUSA Math API Reference
The MUSA Math API provides device-side numeric types, conversions, mathematical functions, and low-level intrinsics for MUSA kernels. Use this reference to choose the correct API for the data type, precision, rounding behavior, saturation behavior, and memory access pattern used by a kernel.
Most entries include the C/C++ signature, the operation performed, and the effect that matters to application code. For conversion APIs, pay attention to the destination format, rounding mode, and saturation rule. For load and store intrinsics, pay attention to the pointer type, value type, and memory access constraints documented for the target device. For packed or vector intrinsics, the operation is applied independently to each lane or packed element unless the description states otherwise.
API Category Index
- FP4 Intrinsics (10)
- FP6 Intrinsics (10)
- FP8 Intrinsics (11)
- Half Precision Intrinsics (241)
- Bfloat16 Precision Intrinsics (247)
- Single Precision Mathematical Functions (79)
- Single Precision Intrinsics (96)
- Double Precision Mathematical Functions (134)
- Double Precision Intrinsics (64)
- FP128 Quad Precision Mathematical Functions (46) [New in MUSA 5.2]
- Floating-Point Classification Intrinsics (6)
- Scalar Mathematical Helper Functions (45)
- Integer Intrinsics (361)
- SIMD Intrinsics (130)
- Type and Enum Reference
- Struct Reference
Type and Enum Reference
This section lists the public storage types, interpretation enums, saturation modes, and numeric structs used by the Math API signatures.
FP4 Types and Enums
__mt_fp4_interpretation_t
typedef enum __mt_fp4_interpretation_t {
__MT_E2M1,
} __mt_fp4_interpretation_t;
Enumerates the possible interpretations of 4-bit values when referring to them as FP4 types.
| Enumerator | Description |
|---|---|
__MT_E2M1 | Represents FP4 numbers of e2m1 kind. |
FP4 Storage Types
| Type | Definition | Description |
|---|---|---|
__mt_fp4_storage_t | typedef __mt_fp8_storage_t __mt_fp4_storage_t | 8-bit unsigned integer storage for FP4 values. |
__mt_fp4x2_storage_t | typedef __mt_fp8_storage_t __mt_fp4x2_storage_t | 8-bit unsigned integer storage for pairs of FP4 values. |
__mt_fp4x4_storage_t | typedef __mt_fp8x2_storage_t __mt_fp4x4_storage_t | 16-bit unsigned integer storage for tetrads of FP4 values. |
FP4 Struct Types
| Type | Description |
|---|---|
__mt_fp4_e2m1 | C++ structure for handling an FP4 value of e2m1 kind. |
__mt_fp4x2_e2m1 | C++ structure for handling a vector of two FP4 values of e2m1 kind. |
__mt_fp4x4_e2m1 | C++ structure for handling a vector of four FP4 values of e2m1 kind. |
FP6 Types and Enums
__mt_fp6_interpretation_t
typedef enum __mt_fp6_interpretation_t {
__MT_E2M3,
__MT_E3M2,
} __mt_fp6_interpretation_t;
Enumerates the possible interpretations of 6-bit values when referring to them as FP6 types.
| Enumerator | Description |
|---|---|
__MT_E2M3 | Represents FP6 numbers of e2m3 kind. |
__MT_E3M2 | Represents FP6 numbers of e3m2 kind. |
FP6 Storage Types
| Type | Definition | Description |
|---|---|---|
__mt_fp6_storage_t | typedef __mt_fp8_storage_t __mt_fp6_storage_t | 8-bit unsigned integer storage for FP6 values. |
__mt_fp6x2_storage_t | typedef __mt_fp8x2_storage_t __mt_fp6x2_storage_t | 16-bit unsigned integer storage for pairs of FP6 values. |
__mt_fp6x4_storage_t | typedef __mt_fp8x4_storage_t __mt_fp6x4_storage_t | 32-bit unsigned integer storage for tetrads of FP6 values. |
FP6 Struct Types
| Type | Description |
|---|---|
__mt_fp6_e2m3 | C++ structure for handling an FP6 value of e2m3 kind. |
__mt_fp6_e3m2 | C++ structure for handling an FP6 value of e3m2 kind. |
__mt_fp6x2_e2m3 | C++ structure for handling a vector of two FP6 values of e2m3 kind. |
__mt_fp6x2_e3m2 | C++ structure for handling a vector of two FP6 values of e3m2 kind. |
__mt_fp6x4_e2m3 | C++ structure for handling a vector of four FP6 values of e2m3 kind. |
__mt_fp6x4_e3m2 | C++ structure for handling a vector of four FP6 values of e3m2 kind. |
FP8 Types and Enums
__musa_saturation_t
typedef enum __musa_saturation_t {
__MT_NOSAT,
__MT_SATFINITE,
} __musa_saturation_t;
Enumerates the saturation modes used when narrowing values to FP8 destination types.
| Enumerator | Description |
|---|---|
__MT_NOSAT | Does not saturate finite results when rounding produces a value outside the destination range. For e4m3, values larger than the maximum finite value become NaN. |
__MT_SATFINITE | Saturates values larger than the maximum finite value to MAXNORM with the same sign as the input. |
__mt_fp8_interpretation_t
typedef enum __mt_fp8_interpretation_t {
__MT_E4M3,
__MT_E5M2,
} __mt_fp8_interpretation_t;
Enumerates the possible interpretations of 8-bit values when referring to them as FP8 types.
| Enumerator | Description |
|---|---|
__MT_E4M3 | Represents FP8 numbers of e4m3 kind. |
__MT_E5M2 | Represents FP8 numbers of e5m2 kind. |
FP8 Storage Types
| Type | Definition | Description |
|---|---|---|
__mt_fp8_storage_t | typedef unsigned char __mt_fp8_storage_t | 8-bit unsigned integer storage for FP8 values. |
__mt_fp8x2_storage_t | typedef unsigned short int __mt_fp8x2_storage_t | 16-bit unsigned integer storage for pairs of FP8 values. |
__mt_fp8x4_storage_t | typedef unsigned int __mt_fp8x4_storage_t | 32-bit unsigned integer storage for tetrads of FP8 values. |
FP8 Struct Types
| Type | Description |
|---|---|
__mt_fp8_e4m3 | C++ structure for handling an FP8 value of e4m3 kind. |
__mt_fp8_e5m2 | C++ structure for handling an FP8 value of e5m2 kind. |
__mt_fp8_e8m0 | C++ structure for handling an FP8 value of e8m0 kind. |
__mt_fp8x2_e4m3 | C++ structure for handling a vector of two FP8 values of e4m3 kind. |
__mt_fp8x2_e5m2 | C++ structure for handling a vector of two FP8 values of e5m2 kind. |
__mt_fp8x2_e8m0 | C++ structure for handling a vector of two FP8 values of e8m0 kind. |
__mt_fp8x4_e4m3 | C++ structure for handling a vector of four FP8 values of e4m3 kind. |
__mt_fp8x4_e5m2 | C++ structure for handling a vector of four FP8 values of e5m2 kind. |
__mt_fp8x4_e8m0 | C++ structure for handling a vector of four FP8 values of e8m0 kind. |
Half Precision Types
| Type | Description |
|---|---|
__half | C++ structure for storing a half-precision floating-point value. |
__half2 | C++ structure for storing two half-precision floating-point values. |
__half_raw | Raw storage representation of a half-precision value. |
__half2_raw | Raw storage representation of two half-precision values. |
half | Alias for __half. |
half2 | Alias for __half2. |
Bfloat16 Precision Types
| Type | Description |
|---|---|
__mt_bfloat16 | C++ structure for storing an mt_bfloat16 floating-point value. |
__mt_bfloat162 | C++ structure for storing two mt_bfloat16 floating-point values. |
__mt_bfloat16_raw | Raw storage representation of an mt_bfloat16 value. |
__mt_bfloat162_raw | Raw storage representation of two mt_bfloat16 values. |
mt_bfloat16 | Alias for __mt_bfloat16. |
mt_bfloat162 | Alias for __mt_bfloat162. |
Rounding Mode
Several conversion APIs accept enum musaRoundMode. This enum is shared with the broader MUSA runtime type definitions. Math API entries describe how the rounding mode affects each operation where the behavior is operation-specific.
Struct Reference
This section documents the numeric wrapper structures used by the Math API. These structures provide storage, constructors, conversion operators, and arithmetic or comparison operators where supported by the type.
__half
struct __half
__half data type
This structure implements the datatype for storing half-precision floating-point numbers. The
structure implements assignment, arithmetic and comparison operators, and type conversions.
16 bits are being used in total: 1 sign bit, 5 bits for the exponent, and the significand is being
stored in 10 bits. The total precision is 11 bits. There are 15361 representable numbers within the
interval [0.0, 1.0], endpoints included. On average we have log10(2**11) ~ 3.311 decimal digits.
The objective here is to provide IEEE754-compliant implementation of binary16 type and arith-
metic with limitations due to device hardware not supporting floating-point exceptions.
Public functions
__half() = default
Constructor by default.
Empty default constructor. The result is uninitialized.
__half2
struct __half2
__half2 data type
This structure implements the datatype for storing two half-precision floating-point numbers.
The structure implements assignment, arithmetic and comparison operators, and type conversions.
- NOTE: __half2 is visible to non-MUSA host compilers
Public functions
__half2() = defaultConstructor by default. Empty default constructor. The result is uninitialized.
__half2_raw
struct __half2_raw
__half2_raw data type
Type allows static initialization of half2 until it becomes a built-in type.
- Note: this initialization is as a bit-field representation of half2, and not a conversion from short2 to half2. Such representation will be deprecated in a future version of MUSA.
- Note: this is visible to non-MUSA compilers, including C-only compilations
Public members
unsigned short xStorage field contains bits of the lower half part.unsigned short yStorage field contains bits of the upper half part.
__half_raw
struct __half_raw
__half_raw data type
Type allows static initialization of half until it becomes a built-in type.
- Note: this initialization is as a bit-field representation of half, and not a conversion from short to half. Such representation will be deprecated in a future version of MUSA.
- Note: this is visible to non-MUSA compilers, including C-only compilations
Public members
unsigned short xStorage field contains bits representation of the half floating-point number.
__mt_bfloat16
struct __mt_bfloat16
mt_bfloat16 data type
This structure implements the datatype for storing mt_bfloat16 floating-point numbers. The
structure implements assignment operators and type conversions. 16 bits are being used in
total: 1 sign bit, 8 bits for the exponent, and the significand is being stored in 7 bits. The total
precision is 8 bits.
Public functions
__mt_bfloat16() = default
Constructor by default.
Empty default constructor. The result is uninitialized.
__mt_bfloat162
struct __mt_bfloat162
mt_bfloat162 data type
This structure implements the datatype for storing two mt_bfloat16 floating-point numbers. The
structure implements assignment, arithmetic and comparison operators, and type conversions.
- NOTE: __mt_bfloat162 is visible to non-MUSA host compilers
Public functions
__mt_bfloat162() = defaultConstructor by default. Empty default constructor. The result is uninitialized.
__mt_bfloat162_raw
struct __mt_bfloat162_raw
__mt_bfloat162_raw data type Type allows static initialization of mt_bfloat162 until it becomes a built-in type.
- Note: this initialization is as a bit-field representation of mt_bfloat162, and not a conversion from short2 to mt_bfloat162. Such representation will be deprecated in a future version of MUSA.
__mt_bfloat16_raw
struct __mt_bfloat16_raw
__mt_bfloat16_raw data type Type allows static initialization of mt_bfloat16 until it becomes a built-in type.
- Note: this initialization is as a bit-field representation of mt_bfloat16, and not a conversion from short to mt_bfloat16. Such representation will be deprecated in a future version of MUSA.
- Note: this is visible to non-MUSA compilers, including C-only compilations
Public members
unsigned short xStorage field contains bits representation of the mt_bfloat16 floating-point number.
__mt_fp4_e2m1
struct __mt_fp4_e2m1
__mt_fp4_e2m1 datatype This structure implements the datatype for handling fp4 floating-point numbers of e2m1 kind: with 1 sign, 2 exponent, 1 implicit and 1 explicit mantissa bits. This encoding does not support Inf/NaN. The structure implements converting constructors and operators. Public functions
__mt_fp4x2_e2m1
struct __mt_fp4x2_e2m1
__mt_fp4x2_e2m1 datatype This structure implements the datatype for handling two fp4 floating-point numbers of e2m1 kind each. The structure implements converting constructors and operators. Public functions
__mt_fp4x4_e2m1
struct __mt_fp4x4_e2m1
__mt_fp4x4_e2m1 datatype This structure implements the datatype for handling four fp4 floating-point numbers of e2m1 kind each. The structure implements converting constructors and operators. Public functions
__mt_fp6_e2m3
struct __mt_fp6_e2m3
__mt_fp6_e2m3 datatype This structure implements the datatype for storing fp6 floating-point numbers of e2m3 kind: with 1 sign, 2 exponent, 1 implicit and 3 explicit mantissa bits. This encoding does not support Inf/NaN. The structure implements converting constructors and operators. Public functions
__mt_fp6_e3m2
struct __mt_fp6_e3m2
__mt_fp6_e3m2 datatype This structure implements the datatype for handling fp6 floating-point numbers of e3m2 kind: with 1 sign, 3 exponent, 1 implicit and 2 explicit mantissa bits. This encoding does not support Inf/NaN. The structure implements converting constructors and operators. Public functions
__mt_fp6x2_e2m3
struct __mt_fp6x2_e2m3
__mt_fp6x2_e2m3 datatype This structure implements the datatype for handling two fp6 floating-point numbers of e2m3 kind each. The structure implements converting constructors and operators. Public functions
__mt_fp6x2_e3m2
struct __mt_fp6x2_e3m2
__mt_fp6x2_e3m2 datatype This structure implements the datatype for handling two fp6 floating-point numbers of e3m2 kind each. The structure implements converting constructors and operators. Public functions
__mt_fp6x4_e2m3
struct __mt_fp6x4_e2m3
__mt_fp6x4_e2m3 datatype This structure implements the datatype for handling four fp6 floating-point numbers of e2m3 kind each. The structure implements converting constructors and operators. Public functions
__mt_fp6x4_e3m2
struct __mt_fp6x4_e3m2
__mt_fp6x4_e3m2 datatype This structure implements the datatype for handling four fp6 floating-point numbers of e3m2 kind each. The structure implements converting constructors and operators. Public functions
__mt_fp8_e4m3
struct __mt_fp8_e4m3
__mt_fp8_e4m3 datatype This structure implements the datatype for storing fp8 floating-point numbers of e4m3 kind: with 1 sign, 4 exponent, 1 implicit and 3 explicit mantissa bits. The encoding doesn’t support Infinity. NaNs are limited to 0x7F and 0xFF values. The structure implements converting constructors and operators. Public functions __mt_fp8_e4m3() = default Constructor by default.
__mt_fp8_e5m2
struct __mt_fp8_e5m2
__mt_fp8_e5m2 datatype This structure implements the datatype for handling fp8 floating-point numbers of e5m2 kind: with 1 sign, 5 exponent, 1 implicit and 2 explicit mantissa bits. The structure implements converting constructors and operators. Public functions __mt_fp8_e5m2() = default Constructor by default.
__mt_fp8_e8m0
struct __mt_fp8_e8m0
__mt_fp8_e8m0 datatype This structure implements the datatype for handling 8-bit scale factors of e8m0 kind: interpreted as powers of two with biased exponent. Bias equals to 127, so numbers 0 through 254 represent 2^-127 through 2^127. Number 0xFF = 255 is reserved for NaN. The structure implements converting constructors and operators.
__mt_fp8x2_e4m3
struct __mt_fp8x2_e4m3
__mt_fp8x2_e4m3 datatype This structure implements the datatype for storage and operations on the vector of two fp8 values of e4m3 kind each: with 1 sign, 4 exponent, 1 implicit and 3 explicit mantissa bits. The encoding doesn’t support Infinity. NaNs are limited to 0x7F and 0xFF values. Public functions __mt_fp8x2_e4m3() = default Constructor by default.
__mt_fp8x2_e5m2
struct __mt_fp8x2_e5m2
__mt_fp8x2_e5m2 datatype This structure implements the datatype for handling two fp8 floating-point numbers of e5m2 kind each: with 1 sign, 5 exponent, 1 implicit and 2 explicit mantissa bits. The structure implements converting constructors and operators. Public functions __mt_fp8x2_e5m2() = default Constructor by default.
__mt_fp8x2_e8m0
struct __mt_fp8x2_e8m0
__mt_fp8x2_e8m0 datatype This structure implements the datatype for storage and operations on the vector of two scale factors of e8m0 kind each. Public functions __mt_fp8x2_e8m0() = default Constructor by default.
__mt_fp8x4_e4m3
struct __mt_fp8x4_e4m3
__mt_fp8x4_e4m3 datatype This structure implements the datatype for storage and operations on the vector of four fp8 values of e4m3 kind each: with 1 sign, 4 exponent, 1 implicit and 3 explicit mantissa bits. The encoding doesn’t support Infinity. NaNs are limited to 0x7F and 0xFF values. Public functions __mt_fp8x4_e4m3() = default Constructor by default.
__mt_fp8x4_e5m2
struct __mt_fp8x4_e5m2
__mt_fp8x4_e5m2 datatype This structure implements the datatype for handling four fp8 floating-point numbers of e5m2 kind each: with 1 sign, 5 exponent, 1 implicit and 2 explicit mantissa bits. The structure implements converting constructors and operators. Public functions __mt_fp8x4_e5m2() = default Constructor by default.
__mt_fp8x4_e8m0
struct __mt_fp8x4_e8m0
__mt_fp8x4_e8m0 datatype This structure implements the datatype for storage and operations on the vector of scale factors of e8m0 kind each. Public functions __mt_fp8x4_e8m0() = default Constructor by default.
FP4 Intrinsics
Use FP4 intrinsics when a kernel needs to pack higher-precision values into 4-bit floating-point storage or unpack FP4 values for computation. These APIs are mainly useful for low-precision data movement, model parameter storage, and kernels that explicitly manage FP4 formats.
FP4 Conversion and Data Movement
__musa_cvt_double_to_fp4
__mt_fp4_storage_t __musa_cvt_double_to_fp4(const double x, const __mt_fp4_interpretation_t fp4_interpretation, const enum musaRoundMode rounding)
Converts input double precision x to fp4 type of the requested kind using specified rounding mode and saturating the out-of-range values.
Converts input x to fp4 type of the kind specified by fp4_interpretation parameter, using rounding mode specified by rounding parameter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result in positive MAXNORM.
- The
__mt_fp4_storage_tvalue holds the result of conversion.
__musa_cvt_double2_to_fp4x2
__mt_fp4x2_storage_t __musa_cvt_double2_to_fp4x2(const double2 x, const __mt_fp4_interpretation_t fp4_interpretation, const enum musaRoundMode rounding)
Converts input vector of two double precision numbers packed in double2 x into a vector of two values of fp4 type of the requested kind using specified rounding mode and saturating the out-of-range values.
Converts input vector x to a vector of two fp4 values of the kind specified by fp4_interpretation parameter, using rounding mode specified by rounding parameter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result in positive MAXNORM.
- The
__mt_fp4x2_storage_tvalue holds the result of conversion.
__musa_cvt_float_to_fp4
__mt_fp4_storage_t __musa_cvt_float_to_fp4(const float x, const __mt_fp4_interpretation_t fp4_interpretation, const enum musaRoundMode rounding)
Converts input single precision x to fp4 type of the requested kind using specified rounding mode and saturating the out-of-range values.
Converts input x to fp4 type of the kind specified by fp4_interpretation parameter, using rounding mode specified by rounding parameter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result in positive MAXNORM.
- The
__mt_fp4_storage_tvalue holds the result of conversion.
__musa_cvt_float2_to_fp4x2
__mt_fp4x2_storage_t __musa_cvt_float2_to_fp4x2(const float2 x, const __mt_fp4_interpretation_t fp4_interpretation, const enum musaRoundMode rounding)
Converts input vector of two single precision numbers packed in float2 x into a vector of two values of fp4 type of the requested kind using specified rounding mode and saturating the out-of-range values.
Converts input vector x to a vector of two fp4 values of the kind specified by fp4_interpretation parameter, using rounding mode specified by rounding parameter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result in positive MAXNORM.
- The
__mt_fp4x2_storage_tvalue holds the result of conversion.
__musa_cvt_halfraw_to_fp4
__mt_fp4_storage_t __musa_cvt_halfraw_to_fp4(const __half_raw x, const __mt_fp4_interpretation_t fp4_interpretation, const enum musaRoundMode rounding)
Converts input half precision x to fp4 type of the requested kind using specified rounding mode and saturating the out-of-range values.
Converts input x to fp4 type of the kind specified by fp4_interpretation parameter, using rounding mode specified by rounding parameter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result in positive MAXNORM.
- The
__mt_fp4_storage_tvalue holds the result of conversion.
__musa_cvt_halfraw2_to_fp4x2
__mt_fp4x2_storage_t __musa_cvt_halfraw2_to_fp4x2(const __half2_raw x, const __mt_fp4_interpretation_t fp4_interpretation, const enum musaRoundMode rounding)
Converts input vector of two half precision numbers packed in __half2_raw x into a vector of two values of fp4 type of the requested kind using specified rounding mode and saturating the out-of-range values.
Converts input vector x to a vector of two fp4 values of the kind specified by fp4_interpretation parameter, using rounding mode specified by rounding parameter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result in positive MAXNORM.
- The
__mt_fp4x2_storage_tvalue holds the result of conversion.
__musa_cvt_bfloat16raw_to_fp4
__mt_fp4_storage_t __musa_cvt_bfloat16raw_to_fp4(const __mt_bfloat16_raw x, const __mt_fp4_interpretation_t fp4_interpretation, const enum musaRoundMode rounding)
Converts input mt_bfloat16 precision x to fp4 type of the requested kind using specified rounding mode and saturating the out-of-range values.
Converts input x to fp4 type of the kind specified by fp4_interpretation parameter, using rounding mode specified by rounding parameter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result in positive MAXNORM.
- The
__mt_fp4_storage_tvalue holds the result of conversion.
__musa_cvt_bfloat16raw2_to_fp4x2
__mt_fp4x2_storage_t __musa_cvt_bfloat16raw2_to_fp4x2(const __mt_bfloat162_raw x, const __mt_fp4_interpretation_t fp4_interpretation, const enum musaRoundMode rounding)
Converts input vector of two mt_bfloat16 precision numbers packed in __mt_bfloat162_raw x into a vector of two values of fp4 type of the requested kind using specified rounding mode and saturating the out-of-range values.
Converts input vector x to a vector of two fp4 values of the kind specified by fp4_interpretation parameter, using rounding mode specified by rounding parameter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result in positive MAXNORM.
- The
__mt_fp4x2_storage_tvalue holds the result of conversion.
__musa_cvt_fp4_to_halfraw
__half_raw __musa_cvt_fp4_to_halfraw(const __mt_fp4_storage_t x, const __mt_fp4_interpretation_t fp4_interpretation)
Converts input fp4 x of the specified kind to half precision.
Converts input x of fp4 type of the kind specified by fp4_interpretation parameter to half precision.
- The
__half_rawvalue holds the result of conversion.
__musa_cvt_fp4x2_to_halfraw2
__half2_raw __musa_cvt_fp4x2_to_halfraw2(const __mt_fp4x2_storage_t x, const __mt_fp4_interpretation_t fp4_interpretation)
Converts input vector of two fp4 values of the specified kind to a vector of two half precision values packed in __half2_raw structure.
Converts input vector x of fp4 type of the kind specified by fp4_interpretation parameter to a vector of two half precision values and returns as __half2_raw structure.
- The
__half2_rawvalue holds the result of conversion.
FP6 Intrinsics
Use FP6 intrinsics when a kernel needs a compact floating-point format with more precision or dynamic range than FP4 while still using smaller storage than FP8 or half precision. The interpretation parameter selects the FP6 layout used by the stored value.
FP6 Conversion and Data Movement
__musa_cvt_double_to_fp6
__mt_fp6_storage_t __musa_cvt_double_to_fp6(const double x, const __mt_fp6_interpretation_t fp6_interpretation, const enum musaRoundMode rounding)
Converts input double precision x to fp6 type of the requested kind using specified rounding mode and saturating the out-of-range values.
Converts input x to fp6 type of the kind specified by fp6_interpretation parameter, using rounding mode specified by rounding parameter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result in positive MAXNORM.
- The
__mt_fp6_storage_tvalue holds the result of conversion.
__musa_cvt_double2_to_fp6x2
__mt_fp6x2_storage_t __musa_cvt_double2_to_fp6x2(const double2 x, const __mt_fp6_interpretation_t fp6_interpretation, const enum musaRoundMode rounding)
Converts input vector of two double precision numbers packed in double2 x into a vector of two values of fp6 type of the requested kind using specified rounding mode and saturating the out-of-range values.
Converts input vector x to a vector of two fp6 values of the kind specified by fp6_interpretation parameter, using rounding mode specified by rounding parameter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result in positive MAXNORM.
- The
__mt_fp6x2_storage_tvalue holds the result of conversion.
__musa_cvt_float_to_fp6
__mt_fp6_storage_t __musa_cvt_float_to_fp6(const float x, const __mt_fp6_interpretation_t fp6_interpretation, const enum musaRoundMode rounding)
Converts input single precision x to fp6 type of the requested kind using specified rounding mode and saturating the out-of-range values.
Converts input x to fp6 type of the kind specified by fp6_interpretation parameter, using rounding mode specified by rounding parameter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result in positive MAXNORM.
- The
__mt_fp6_storage_tvalue holds the result of conversion.
__musa_cvt_float2_to_fp6x2
__mt_fp6x2_storage_t __musa_cvt_float2_to_fp6x2(const float2 x, const __mt_fp6_interpretation_t fp6_interpretation, const enum musaRoundMode rounding)
Converts input vector of two single precision numbers packed in float2 x into a vector of two values of fp6 type of the requested kind using specified rounding mode and saturating the out-of-range values.
Converts input vector x to a vector of two fp6 values of the kind specified by fp6_interpretation parameter, using rounding mode specified by rounding parameter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result in positive MAXNORM.
- The
__mt_fp6x2_storage_tvalue holds the result of conversion.
__musa_cvt_halfraw_to_fp6
__mt_fp6_storage_t __musa_cvt_halfraw_to_fp6(const __half_raw x, const __mt_fp6_interpretation_t fp6_interpretation, const enum musaRoundMode rounding)
Converts input half precision x to fp6 type of the requested kind using specified rounding mode and saturating the out-of-range values.
Converts input x to fp6 type of the kind specified by fp6_interpretation parameter, using rounding mode specified by rounding parameter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result in positive MAXNORM.
- The
__mt_fp6_storage_tvalue holds the result of conversion.
__musa_cvt_halfraw2_to_fp6x2
__mt_fp6x2_storage_t __musa_cvt_halfraw2_to_fp6x2(const __half2_raw x, const __mt_fp6_interpretation_t fp6_interpretation, const enum musaRoundMode rounding)
Converts input vector of two half precision numbers packed in __half2_raw x into a vector of two values of fp6 type of the requested kind using specified rounding mode and saturating the out-of-range values.
Converts input vector x to a vector of two fp6 values of the kind specified by fp6_interpretation parameter, using rounding mode specified by rounding parameter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result in positive MAXNORM.
- The
__mt_fp6x2_storage_tvalue holds the result of conversion.
__musa_cvt_bfloat16raw_to_fp6
__mt_fp6_storage_t __musa_cvt_bfloat16raw_to_fp6(const __mt_bfloat16_raw x, const __mt_fp6_interpretation_t fp6_interpretation, const enum musaRoundMode rounding)
Converts input mt_bfloat16 precision x to fp6 type of the requested kind using specified rounding mode and saturating the out-of-range values.
Converts input x to fp6 type of the kind specified by fp6_interpretation parameter, using rounding mode specified by rounding parameter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result in positive MAXNORM.
- The
__mt_fp6_storage_tvalue holds the result of conversion.
__musa_cvt_bfloat16raw2_to_fp6x2
__mt_fp6x2_storage_t __musa_cvt_bfloat16raw2_to_fp6x2(const __mt_bfloat162_raw x, const __mt_fp6_interpretation_t fp6_interpretation, const enum musaRoundMode rounding)
Converts input vector of two mt_bfloat16 precision numbers packed in __mt_bfloat162_raw x into a vector of two values of fp6 type of the requested kind using specified rounding mode and saturating the out-of-range values.
Converts input vector x to a vector of two fp6 values of the kind specified by fp6_interpretation parameter, using rounding mode specified by rounding parameter. Large out-of-range values saturate to MAXNORM of the same sign. NaN input values result in positive MAXNORM.
- The
__mt_fp6x2_storage_tvalue holds the result of conversion.
__musa_cvt_fp6_to_halfraw
__half_raw __musa_cvt_fp6_to_halfraw(const __mt_fp6_storage_t x, const __mt_fp6_interpretation_t fp6_interpretation)
Converts input fp6 x of the specified kind to half precision.
Converts input x of fp6 type of the kind specified by fp6_interpretation parameter to half precision.
- The
__half_rawvalue holds the result of conversion.
__musa_cvt_fp6x2_to_halfraw2
__half2_raw __musa_cvt_fp6x2_to_halfraw2(const __mt_fp6x2_storage_t x, const __mt_fp6_interpretation_t fp6_interpretation)
Converts input vector of two fp6 values of the specified kind to a vector of two half precision values packed in __half2_raw structure.
Converts input vector x of fp6 type of the kind specified by fp6_interpretation parameter to a vector of two half precision values and returns as __half2_raw structure.
- The
__half2_rawvalue holds the result of conversion.
FP8 Intrinsics
Use FP8 intrinsics to convert values to and from compact 8-bit floating-point formats. These APIs are useful when a kernel explicitly controls FP8 storage, scaling factors, saturation behavior, or conversion between FP8 and half/bfloat16 data paths.
FP8 Conversion and Data Movement
__musa_cvt_float_to_fp8
__mt_fp8_storage_t __musa_cvt_float_to_fp8(const float x, const __musa_saturation_t saturate, const __mt_fp8_interpretation_t fp8_interpretation)
Converts input single precision x to fp8 type of the requested kind using round-to-nearest-even rounding and requested saturation mode.
Converts input x to fp8 type of the kind specified by fp8_interpretation parameter, using round-to-nearest-even rounding and saturation mode specified by saturate parameter.
- The
__mt_fp8_storage_tvalue holds the result of conversion.
__musa_cvt_halfraw_to_fp8
__mt_fp8_storage_t __musa_cvt_halfraw_to_fp8(const __half_raw x, const __musa_saturation_t saturate, const __mt_fp8_interpretation_t fp8_interpretation)
Converts input half precision x to fp8 type of the requested kind using round-to-nearest-even rounding and requested saturation mode.
Converts input x to fp8 type of the kind specified by fp8_interpretation parameter, using round-to-nearest-even rounding and saturation mode specified by saturate parameter.
- The
__mt_fp8_storage_tvalue holds the result of conversion.
__musa_cvt_bfloat16raw_to_fp8
__mt_fp8_storage_t __musa_cvt_bfloat16raw_to_fp8(const __mt_bfloat16_raw x, const __musa_saturation_t saturate, const __mt_fp8_interpretation_t fp8_interpretation)
Converts input __mt_bfloat16_raw precision x to fp8 type of the requested kind using round-to-nearest-even rounding and requested saturation mode.
Converts input x to fp8 type of the kind specified by fp8_interpretation parameter, using round-to-nearest-even rounding and saturation mode specified by saturate parameter.
- The
__mt_fp8_storage_tvalue holds the result of conversion.
__musa_cvt_bfloat16raw_to_e8m0
__mt_fp8_storage_t __musa_cvt_bfloat16raw_to_e8m0(const __mt_bfloat16_raw x, const __musa_saturation_t saturate, const enum musaRoundMode rounding)
Converts input bfloat16 input into a scaling factor of e8m0 kind.
Input number's absolute value is rounded to the closest power of two in the direction specified via rounding parameter. Rounded results that are smaller than the smallest representable target format number 2^-127 are then clipped to 2^-127. Results that are larger than the largest representable target format number 2^127 are either clipped to 2^127 if saturate equals to __MT_SATFINITE, or convert to NaN otherwise. NaN inputs convert into NaN output, encoded as 0xFF in the target format.
- The
__mt_fp8_storage_tvalue holds the result of conversion.
__musa_cvt_bfloat162raw_to_e8m0x2
__mt_fp8x2_storage_t __musa_cvt_bfloat162raw_to_e8m0x2(const __mt_bfloat162_raw x, const __musa_saturation_t saturate, const enum musaRoundMode rounding)
Converts a pair of bfloat16 values into a pair of scaling factors of e8m0 kind.
__musa_cvt_bfloat16raw_to_e8m0() for details of conversion.
- The
__mt_fp8x2_storage_tvalue holds the result of conversion.
__musa_cvt_float_to_e8m0
__mt_fp8_storage_t __musa_cvt_float_to_e8m0(const float x, const __musa_saturation_t saturate, const enum musaRoundMode rounding)
Converts input float value into a scaling factor of e8m0 kind.
__musa_cvt_bfloat16raw_to_e8m0() for details of conversion.
- The
__mt_fp8_storage_tvalue holds the result of conversion.
__musa_cvt_float2_to_e8m0x2
__mt_fp8x2_storage_t __musa_cvt_float2_to_e8m0x2(const float2 x, const __musa_saturation_t saturate, const enum musaRoundMode rounding)
Converts a pair of float values into a pair of scaling factors of e8m0 kind.
__musa_cvt_bfloat16raw_to_e8m0() for details of conversion.
- The
__mt_fp8x2_storage_tvalue holds the result of conversion.
__musa_cvt_double_to_e8m0
__mt_fp8_storage_t __musa_cvt_double_to_e8m0(const double x, const __musa_saturation_t saturate, const enum musaRoundMode rounding)
Converts input double value into a scaling factor of e8m0 kind.
__musa_cvt_bfloat16raw_to_e8m0() for details of conversion.
- The
__mt_fp8_storage_tvalue holds the result of conversion.
__musa_cvt_double2_to_e8m0x2
__mt_fp8x2_storage_t __musa_cvt_double2_to_e8m0x2(const double2 x, const __musa_saturation_t saturate, const enum musaRoundMode rounding)
Converts a pair of double values into a pair of scaling factors of e8m0 kind.
__musa_cvt_bfloat16raw_to_e8m0() for details of conversion.
- The
__mt_fp8x2_storage_tvalue holds the result of conversion.
__musa_cvt_e8m0_to_bf16raw
__mt_bfloat16_raw __musa_cvt_e8m0_to_bf16raw(const __mt_fp8_storage_t x)
Converts input scaling factor value of e8m0 kind into bfloat16.
Input scales are exact powers of two or a NaN value, also representable in the target format.
- The
__mt_bfloat16_rawvalue holds the result of conversion.
__musa_cvt_e8m0x2_to_bf162raw
__mt_bfloat162_raw __musa_cvt_e8m0x2_to_bf162raw(const __mt_fp8x2_storage_t x)
Converts input pair of scaling factors of e8m0 kind into a pair of bfloat16 values.
- The
__mt_bfloat162_rawvalue holds the result of conversion.
Half Precision Intrinsics
Use half precision intrinsics for device code that stores or computes with 16-bit floating-point values. Scalar __half APIs operate on one half value; __half2 APIs operate on two packed half values and are useful when a kernel can process two lanes with the same operation.
Half Precision Conversion and Data
__double2half
__half __double2half(const double a)
Converts double number to half precision in round-to-nearest-even mode and returns half with converted value.
Converts double number a to half precision in round-to-nearest-even mode.
Parameters
a(in): double
Returns
- half
- a converted to half.
__float22half2_rn
__half2 __float22half2_rn(const float2 a)
Converts both components of float2 number to half precision in round-to-nearest-even mode and returns half2 with converted values.
Converts both components of float2 to half precision in round-to-nearest mode and combines the results into one half2 number. Low 16 bits of the return value correspond to a.x and high 16 bits of the return value correspond to a.y.
converted float2 components.
Parameters
a(in): float2..
Returns
- half2
- The
half2which has corresponding halves equal to the
__float2half
__half __float2half(const float a)
Converts float number to half precision in round-to-nearest-even mode and returns half with converted value.
Converts float number a to half precision in round-to-nearest-even mode.
Parameters
a(in): float
Returns
- half
- a converted to half.
__float2half2_rn
__half2 __float2half2_rn(const float a)
Converts input to half precision in round-to-nearest-even mode and populates both halves of half2 with converted value.
Converts input a to half precision in round-to-nearest-even mode and populates both halves of half2 with converted value.
float
precision number.
Parameters
a(in):
Returns
- half2
- The
half2value with both halves equal to the converted half
__float2half_rd
__half __float2half_rd(const float a)
Converts float number to half precision in round-down mode and returns half with converted value.
Converts float number a to half precision in round-down mode.
Parameters
a(in): float
Returns
- half
- a converted to half.
__float2half_rn
__half __float2half_rn(const float a)
Converts float number to half precision in round-to-nearest-even mode and returns half with converted value.
Converts float number a to half precision in round-to-nearest-even mode.
Parameters
a(in): float
Returns
- half
- a converted to half.
__float2half_ru
__half __float2half_ru(const float a)
Converts float number to half precision in round-up mode and returns half with converted value.
Converts float number a to half precision in round-up mode.
Parameters
a(in): float
Returns
- half
- a converted to half.
__float2half_rz
__half __float2half_rz(const float a)
Converts float number to half precision in round-towards-zero mode and returns half with converted value.
Converts float number a to half precision in round-towards-zero mode.
converted to half.
Parameters
a(in): float
Returns
- half
- a
__floats2half2_rn
__half2 __floats2half2_rn(const float a, const float b)
Converts both input floats to half precision in round-to-nearest-even mode and returns half2 with converted values.
Converts both input floats to half precision in round-to-nearest-even mode and combines the results into one half2 number. Low 16 bits of the return value correspond to the input a, high 16 bits correspond to the input b.
converted input floats.
Parameters
a(in): floatb(in): float
Returns
- half2
- The
half2value with corresponding halves equal to the
Half2 Arithmetic Functions
__h2div
__half2 __h2div(const __half2 a, const __half2 b)
Performs half2 vector division in round-to-nearest-even mode.
Divides half2 input vector a by input vector b in round-to-nearest mode.
Parameters
a(in): half2..b(in): half2..
Returns
- half2
- The elementwise division of
awithb.
Half Arithmetic Functions
__habs
__half __habs(const __half a)
Calculates the absolute value of input half number and returns the result.
Calculates the absolute value of input half number and returns the result.
Parameters
a(in): half
Returns
- half
- The absolute value of a.
Half2 Arithmetic Functions
__habs2
__half2 __habs2(const __half2 a)
Calculates the absolute value of both halves of the input half2 number and returns the result.
Calculates the absolute value of both halves of the input half2 number and returns the result.
Parameters
a(in): half2..
Returns
- half2
- Returns
awith the absolute value of both halves.
Half Arithmetic Functions
__hadd
__half __hadd(const __half a, const __half b)
Performs half addition in round-to-nearest-even mode.
Performs half addition of inputs a and b, in round-to-nearest-even mode.
Parameters
a(in): halfb(in): half
Returns
- half
- The sum of
aandb.
Half2 Arithmetic Functions
__hadd2
__half2 __hadd2(const __half2 a, const __half2 b)
Performs half2 vector addition in round-to-nearest-even mode.
Performs half2 vector add of inputs a and b, in round-to-nearest mode.
Parameters
a(in): half2..b(in): half2..
Returns
- half2
- The sum of vectors
aandb.
__hadd2_rn
__half2 __hadd2_rn(const __half2 a, const __half2 b)
Performs half2 vector addition in round-to-nearest-even mode. Prevents floating-point contractions of mul+add into fma.
Performs half2 vector add of inputs a and b, in round-to-nearest mode.
Parameters
a(in): half2..b(in): half2..
Returns
- half2
- The sum of
aandb, with respect to saturation.
__hadd2_sat
__half2 __hadd2_sat(const __half2 a, const __half2 b)
Performs half2 vector addition in round-to-nearest-even mode, with saturation to [0.0, 1.0].
Performs half2 vector add of inputs a and b, in round-to-nearest mode, and clamps the results to range [0.0, 1.0]. NaN results are flushed to +0.0.
Parameters
a(in): half2..b(in): half2..
Returns
- half2
- The sum of
aandb, with respect to saturation.
Half Arithmetic Functions
__hadd_rn
__half __hadd_rn(const __half a, const __half b)
Performs half addition in round-to-nearest-even mode. Prevents floating-point contractions of mul+add into fma.
Performs half addition of inputs a and b, in round-to-nearest-even mode.
Parameters
a(in): halfb(in): half
Returns
- half
- The sum of
aandb.
__hadd_sat
__half __hadd_sat(const __half a, const __half b)
Performs half addition in round-to-nearest-even mode, with saturation to [0.0, 1.0].
Performs half add of inputs a and b, in round-to-nearest-even mode, and clamps the result to range [0.0, 1.0]. NaN results are flushed to +0.0.
Parameters
a(in): halfb(in): half
Returns
- half
- The sum of
aandb, with respect to saturation.
Half Precision Conversion and Data Movement 69
__half
__half(const __half_raw &hr)
Constructor from __half_raw .
__half
__half()
Constructor from default .
__half
explicit __half(const __mt_bfloat16 f)
Construct __half from __mt_bfloat16 input using default round-to-nearest-even round-ing mode.
__half
__half(const float f)
Construct __half from float input using default round-to-nearest-even rounding mode.
__half
__half(const double f)
Construct __half from double input using default round-to-nearest-even rounding mode.
__half
__half(const short val)
Construct __half from short integer input using default round-to-nearest-even rounding mode.
__half
__half(const unsigned short val)
Construct __half from short integer input using default round-to-nearest-even rounding mode.
__half
__half(const int val)
Construct __half from integer input using default round-to-nearest-even rounding mode.
__half
__half(const unsigned int val)
Construct __half from unsigned int input using default round-to-nearest-even rounding mode.
__half
__half(const long val)
Construct __half from long input using default round-to-nearest-even rounding mode.
__half
__half(const unsigned long val)
Construct __half from unsigned long input using default round-to-nearest-even rounding mode.
__half
__half(const long long val)
Construct __half from long long input using default round-to-nearest-even rounding mode.
__half
__half(const unsigned long long val)
Construct __half from unsigned long long input using default round-to-nearest-even rounding mode.
Half Precision Conversion and Data
__half2
__half2(const __half &a, const __half &b)
Constructor from two __half variables.
__half2
__half2()
Constructor from default.
__half2
__half2(const __half2 &src)
Copy constructor.
__half2
__half2(const __half2 &&src)
Move constructor, available for C++11 and later dialects.
__half2
__half2(const __half2_raw &h2r)
Constructor from __half2_raw .
__half22float2
float2 __half22float2(const __half2 a)
Converts both halves of half2 to float2 and returns the result.
Converts both halves of half2 input a to float2 and returns the result.
Parameters
a(in): half2..
Returns
- float2
- a converted to float2.
__half2char_rz
signed char __half2char_rz(const __half h)
Convert a half to a signed char in round-towards-zero mode.
Convert the half-precision floating-point value h to a signed char integer in round-towards-zero mode. NaN inputs are converted to 0.
Parameters
h(in): half
Returns
- signed char
- h converted to a signed char using round-towards-zero mode.
__half2float
float __half2float(const __half a)
Converts half number to float.
Converts half number a to float.
Parameters
a(in): float
Returns
- float
- a converted to float.
__half2half2
__half2 __half2half2(const __half a)
Returns half2 with both halves equal to the input value.
Returns half2 number with both halves equal to the input a half number.
Parameters
a(in): half
Returns
- half2
- The vector which has both its halves equal to the input
a.
__half2int_rd
int __half2int_rd(const __half h)
Convert a half to a signed integer in round-down mode.
Convert the half-precision floating-point value h to a signed integer in round-down mode.
Parameters
h(in): half
Returns
- int
- h converted to a signed integer.
__half2int_rn
int __half2int_rn(const __half h)
Convert a half to a signed integer in round-to-nearest-even mode.
Convert the half-precision floating-point value h to a signed integer in round-to-nearest-even mode. NaN inputs are converted to 0.
Parameters
h(in): half. Is only being
Returns
- int
- h converted to a signed integer.
__half2int_ru
int __half2int_ru(const __half h)
Convert a half to a signed integer in round-up mode.
Convert the half-precision floating-point value h to a signed integer in round-up mode.
Parameters
h(in): half
Returns
- int
- h converted to a signed integer.
__half2int_rz
int __half2int_rz(const __half h)
Convert a half to a signed integer in round-towards-zero mode.
Convert the half-precision floating-point value h to a signed integer in round-towards-zero mode.
Parameters
h(in): - half. Is only being read.
Returns
- int
- h converted to a signed integer using round-towards-zero mode.
- __half2int_rz (±0) returns 0.
- __half2int_rz (+∞) returns INT_MAX = 0x7FFFFFFF.
- __half2int_rz (−∞) returns INT_MIN = 0x80000000.
- __half2int_rz(NaN) returns 0.
__half2ll_rd
long long int __half2ll_rd(const __half h)
Convert a half to a signed 64-bit integer in round-down mode.
Convert the half-precision floating-point value h to a signed 64-bit integer in round-down mode.
Parameters
h(in): half
Returns
- long long int
- h converted to a signed 64-bit integer.
__half2ll_rn
long long int __half2ll_rn(const __half h)
Convert a half to a signed 64-bit integer in round-to-nearest-even mode.
Convert the half-precision floating-point value h to a signed 64-bit integer in round-to-nearest-even mode.
being.
Parameters
h(in): half. Is only
Returns
- long long int
- h converted to a signed 64-bit integer.
__half2ll_ru
long long int __half2ll_ru(const __half h)
Convert a half to a signed 64-bit integer in round-up mode.
Convert the half-precision floating-point value h to a signed 64-bit integer in round-up mode.
Parameters
h(in): half
Returns
- long long int
- h converted to a signed 64-bit integer.
__half2ll_rz
long long int __half2ll_rz(const __half h)
Convert a half to a signed 64-bit integer in round-towards-zero mode.
Convert the half-precision floating-point value h to a signed 64-bit integer in round-towards-zero mode.
read.
Parameters
h(in): half. Is only being
Returns
- long long int
- h converted to a signed 64-bit integer.
__half2short_rd
short int __half2short_rd(const __half h)
Convert a half to a signed short integer in round-down mode.
Convert the half-precision floating-point value h to a signed short integer in round-down mode.
Parameters
h(in): half
Returns
- short int
- h converted to a signed short integer.
Half Precision Conversion and Data Movement 69
__half2short_rn
short int __half2short_rn(const __half h)
Convert a half to a signed short integer in round-to-nearest-even mode.
Convert the half-precision floating-point value h to a signed short integer in round-to-nearest-even mode.
being.
Parameters
h(in): half. Is only
Returns
- short int
- h converted to a signed short integer.
__half2short_ru
short int __half2short_ru(const __half h)
Convert a half to a signed short integer in round-up mode.
Convert the half-precision floating-point value h to a signed short integer in round-up mode.
Parameters
h(in): half
Returns
- short int
- h converted to a signed short integer.
__half2short_rz
short int __half2short_rz(const __half h)
Convert a half to a signed short integer in round-towards-zero mode.
Convert the half-precision floating-point value h to a signed short integer in round-towards-zero mode.
read.
Parameters
h(in): half. Is only being
Returns
- short int
- h converted to a signed short integer.
__half2uchar_rz
unsigned char __half2uchar_rz(const __half h)
Convert a half to an unsigned char in round-towards-zero mode.
Convert the half-precision floating-point value h to an unsigned char in round-towards-zero mode. NaN inputs are converted to 0.
Parameters
h(in): half
Returns
- unsigned char
- h converted to an unsigned char using round-towards-zero mode.
__half2uint_rd
unsigned int __half2uint_rd(const __half h)
Convert a half to an unsigned integer in round-down mode.
Convert the half-precision floating-point value h to an unsigned integer in round-down mode.
Parameters
h(in): half
Returns
- unsigned int
- h converted to an unsigned integer.
__half2uint_rn
unsigned int __half2uint_rn(const __half h)
Convert a half to an unsigned integer in round-to-nearest-even mode.
Convert the half-precision floating-point value h to an unsigned integer in round-to-nearest-even mode.
read.
Parameters
h(in): half. Is only being
Returns
- unsigned int
- h converted to an unsigned integer.
__half2uint_ru
unsigned int __half2uint_ru(const __half h)
Convert a half to an unsigned integer in round-up mode.
Convert the half-precision floating-point value h to an unsigned integer in round-up mode.
Parameters
h(in): half
Returns
- unsigned int
- h converted to an unsigned integer.
__half2uint_rz
unsigned int __half2uint_rz(const __half h)
Convert a half to an unsigned integer in round-towards-zero mode.
Convert the half-precision floating-point value h to an unsigned integer in round-towards-zero mode.
Parameters
h(in): half
Returns
- unsigned int
- h converted to an unsigned integer.
__half2ull_rd
unsigned long long int __half2ull_rd(const __half h)
Convert a half to an unsigned 64-bit integer in round-down mode.
Convert the half-precision floating-point value h to an unsigned 64-bit integer in round-down mode.
Parameters
h(in): half
Returns
- unsigned long long int
- h converted to an unsigned 64-bit integer.
__half2ull_rn
unsigned long long int __half2ull_rn(const __half h)
Convert a half to an unsigned 64-bit integer in round-to-nearest-even mode.
Convert the half-precision floating-point value h to an unsigned 64-bit integer in round-to-nearest-even mode.
being.
Parameters
h(in): half. Is only
Returns
- unsigned long long int
- h converted to an unsigned 64-bit integer.
__half2ull_ru
unsigned long long int __half2ull_ru(const __half h)
Convert a half to an unsigned 64-bit integer in round-up mode.
Convert the half-precision floating-point value h to an unsigned 64-bit integer in round-up mode.
Parameters
h(in): half
Returns
- unsigned long long int
- h converted to an unsigned 64-bit integer.
__half2ull_rz
unsigned long long int __half2ull_rz(const __half h)
Convert a half to an unsigned 64-bit integer in round-towards-zero mode.
Convert the half-precision floating-point value h to an unsigned 64-bit integer in round-towards-zero mode.
read.
Parameters
h(in): half. Is only being
Returns
- unsigned long long int
- h converted to an unsigned 64-bit integer.
__half2ushort_rd
unsigned short int __half2ushort_rd(const __half h)
Convert a half to an unsigned short integer in round-down mode.
Convert the half-precision floating-point value h to an unsigned short integer in round-down mode.
Parameters
h(in): half
Returns
- unsigned short int
- h converted to an unsigned short integer.
__half2ushort_rn
unsigned short int __half2ushort_rn(const __half h)
Convert a half to an unsigned short integer in round-to-nearest-even mode.
Convert the half-precision floating-point value h to an unsigned short integer in round-to-nearest-even mode.
being.
Parameters
h(in): half. Is only
Returns
- unsigned short int
- h converted to an unsigned short integer.
__half2ushort_ru
unsigned short int __half2ushort_ru(const __half h)
Convert a half to an unsigned short integer in round-up mode.
Convert the half-precision floating-point value h to an unsigned short integer in round-up mode.
Parameters
h(in): half
Returns
- unsigned short int
- h converted to an unsigned short integer.
__half2ushort_rz
unsigned short int __half2ushort_rz(const __half h)
Convert a half to an unsigned short integer in round-towards-zero mode.
Convert the half-precision floating-point value h to an unsigned short integer in round-towards-zero mode.
read.
Parameters
h(in): half. Is only being
Returns
- unsigned short int
- h converted to an unsigned short integer.
Half Precision Conversion and Data Movement 71
__half_as_short
short int __half_as_short(const __half h)
Reinterprets bits in a half as a signed short integer.
Reinterprets the bits in the half-precision floating-point number h as a signed short integer.
Parameters
h(in): half
Returns
- short int
- The reinterpreted value.
__half_as_ushort
unsigned short int __half_as_ushort(const __half h)
Reinterprets bits in a half as an unsigned short integer.
Reinterprets the bits in the half-precision floating-point h as an unsigned short number.
Parameters
h(in): half
Returns
- unsigned short int
- The reinterpreted value.
__halves2half2
__half2 __halves2half2(const __half a, const __half b)
Combines two half numbers into one half2 number.
Combines two input half number a and b into one half2 number. Input a is stored in low 16 bits of the return value, input b is stored in high 16 bits of the return value.
being.
Parameters
a(in): half. Is onlyb(in): half
Returns
- half2
- The half2 with one half equal to
aand the other tob.
Half2 Comparison Functions
__hbeq2
bool __hbeq2(const __half2 a, const __half2 b)
Performs half2 vector if-equal comparison and returns boolean true iff both half results are true, boolean false otherwise.
Performs half2 vector if-equal comparison of inputs a and b. The bool result is set to true only if both half if-equal comparisons evaluate to true, or false otherwise. NaN inputs generate false results.
of vectors a and b are true;
Parameters
a(in): half2..b(in): half2..
Returns
- bool
- true if both
halfresults of if-equal comparison - false otherwise.
__hbequ2
bool __hbequ2(const __half2 a, const __half2 b)
Performs half2 vector unordered if-equal comparison and returns boolean true iff both half results are true, boolean false otherwise.
Performs half2 vector if-equal comparison of inputs a and b. The bool result is set to true only if both half if-equal comparisons evaluate to true, or false otherwise. NaN inputs generate true results.
comparison of vectors a and b are true;
Parameters
a(in): half2..b(in): half2..
Returns
- bool
- true if both
halfresults of unordered if-equal - false otherwise.
__hbge2
bool __hbge2(const __half2 a, const __half2 b)
Performs half2 vector greater-equal comparison and returns boolean true iff both half results are true, boolean false otherwise.
Performs half2 vector greater-equal comparison of inputs a and b. The bool result is set to true only if both half greater-equal comparisons evaluate to true, or false otherwise. NaN inputs generate false results.
only.
comparison of vectors a and b are true;
Parameters
a(in): half2..b(in): half2. Is
Returns
- bool
- true if both
halfresults of greater-equal - false otherwise.
__hbgeu2
bool __hbgeu2(const __half2 a, const __half2 b)
Performs half2 vector unordered greater-equal comparison and returns boolean true iff both half results are true, boolean false otherwise.
Performs half2 vector greater-equal comparison of inputs a and b. The bool result is set to true only if both half greater-equal comparisons evaluate to true, or false otherwise. NaN inputs generate true results.
only.
greater-equal comparison of vectors a and b are true;
Parameters
a(in): half2..b(in): half2. Is
Returns
- bool
- true if both
halfresults of unordered - false otherwise.
__hbgt2
bool __hbgt2(const __half2 a, const __half2 b)
Performs half2 vector greater-than comparison and returns boolean true iff both half results are true, boolean false otherwise.
Performs half2 vector greater-than comparison of inputs a and b. The bool result is set to true only if both half greater-than comparisons evaluate to true, or false otherwise. NaN inputs generate false results.
only.
comparison of vectors a and b are true;
Parameters
a(in): half2..b(in): half2. Is
Returns
- bool
- true if both
halfresults of greater-than - false otherwise.
__hbgtu2
bool __hbgtu2(const __half2 a, const __half2 b)
Performs half2 vector unordered greater-than comparison and returns boolean true iff both half results are true, boolean false otherwise.
Performs half2 vector greater-than comparison of inputs a and b. The bool result is set to true only if both half greater-than comparisons evaluate to true, or false otherwise. NaN inputs generate true results.
only.
greater-than comparison of vectors a and b are true;
Parameters
a(in): half2..b(in): half2. Is
Returns
- bool
- true if both
halfresults of unordered - false otherwise.
__hble2
bool __hble2(const __half2 a, const __half2 b)
Performs half2 vector less-equal comparison and returns boolean true iff both half results are true, boolean false otherwise.
Performs half2 vector less-equal comparison of inputs a and b. The bool result is set to true only if both half less-equal comparisons evaluate to true, or false otherwise. NaN inputs generate false results.
of vectors a and b are true;
Parameters
a(in): half2..b(in): half2..
Returns
- bool
- true if both
halfresults of less-equal comparison - false otherwise.
__hbleu2
bool __hbleu2(const __half2 a, const __half2 b)
Performs half2 vector unordered less-equal comparison and returns boolean true iff both half results are true, boolean false otherwise.
Performs half2 vector less-equal comparison of inputs a and b. The bool result is set to true only if both half less-equal comparisons evaluate to true, or false otherwise. NaN inputs generate true results.
comparison of vectors a and b are true;
Parameters
a(in): half2..b(in): half2..
Returns
- bool
- true if both
halfresults of unordered less-equal - false otherwise.
__hblt2
bool __hblt2(const __half2 a, const __half2 b)
Performs half2 vector less-than comparison and returns boolean true iff both half results are true, boolean false otherwise.
Performs half2 vector less-than comparison of inputs a and b. The bool result is set to true only if both half less-than comparisons evaluate to true, or false otherwise. NaN inputs generate false results.
of vectors a and b are true;
Parameters
a(in): half2..b(in): half2..
Returns
- bool
- true if both
halfresults of less-than comparison - false otherwise.
__hbltu2
bool __hbltu2(const __half2 a, const __half2 b)
Performs half2 vector unordered less-than comparison and returns boolean true iff both half results are true, boolean false otherwise.
Performs half2 vector less-than comparison of inputs a and b. The bool result is set to true only if both half less-than comparisons evaluate to true, or false otherwise. NaN inputs generate true results.
vectors a and b are true;
Parameters
a(in): half2..b(in): half2..
Returns
- bool
- true if both
halfresults of unordered less-than comparison of - false otherwise.
__hbne2
bool __hbne2(const __half2 a, const __half2 b)
Performs half2 vector not-equal comparison and returns boolean true iff both half results are true, boolean false otherwise.
Performs half2 vector not-equal comparison of inputs a and b. The bool result is set to true only if both half not-equal comparisons evaluate to true, or false otherwise. NaN inputs generate false results.
of vectors a and b are true,
Parameters
a(in): half2..b(in): half2..
Returns
- bool
- true if both
halfresults of not-equal comparison - false otherwise.
__hbneu2
bool __hbneu2(const __half2 a, const __half2 b)
Performs half2 vector unordered not-equal comparison and returns boolean true iff both half results are true, boolean false otherwise.
Performs half2 vector not-equal comparison of inputs a and b. The bool result is set to true only if both half not-equal comparisons evaluate to true, or false otherwise. NaN inputs generate true results.
comparison of vectors a and b are true;
Parameters
a(in): half2..b(in): half2..
Returns
- bool
- true if both
halfresults of unordered not-equal - false otherwise.
Half2 Arithmetic Functions
__hcmadd
__half2 __hcmadd(const __half2 a, const __half2 b, const __half2 c)
Performs fast complex multiply-accumulate
Interprets vector half2 input pairs a, b, and c as complex numbers in half precision and performs complex multiply-accumulate operation: a*b + c
numbers a, b, and c
Parameters
a(in): half2..b(in): half2..c(in): half2..
Returns
- half2
- The result of complex multiply-accumulate operation on complex
Half Arithmetic Functions
__hdiv
__half __hdiv(const __half a, const __half b)
Performs half division in round-to-nearest-even mode.
Divides half input a by input b in round-to-nearest mode.
Parameters
a(in): halfb(in): half
Returns
- half
- The result of dividing
abyb.
Half Comparison Functions
__heq
bool __heq(const __half a, const __half b)
Performs half if-equal comparison.
Performs half if-equal comparison of inputs a and b. NaN inputs generate false results.
Parameters
a(in): halfb(in): half
Returns
- bool
- The boolean result of if-equal comparison of
aandb.
Half2 Comparison Functions
__heq2
__half2 __heq2(const __half2 a, const __half2 b)
Performs half2 vector if-equal comparison.
Performs half2 vector if-equal comparison of inputs a and b. The corresponding half results are set to 1.0 for true, or 0.0 for false. NaN inputs generate false results.
read.
Parameters
a(in): half2. Is only beingb(in): half2..
Returns
- half2
- The vector result of if-equal comparison of vectors
aandb.
Half Comparison Functions
__hequ
bool __hequ(const __half a, const __half b)
Performs half unordered if-equal comparison.
Performs half if-equal comparison of inputs a and b. NaN inputs generate true results.
b.
Parameters
a(in): halfb(in): half
Returns
- bool
- The boolean result of unordered if-equal comparison of
aand
Half2 Comparison Functions
__hequ2
__half2 __hequ2(const __half2 a, const __half2 b)
Performs half2 vector unordered if-equal comparison.
Performs half2 vector if-equal comparison of inputs a and b. The corresponding half results are set to 1.0 for true, or 0.0 for false. NaN inputs generate true results.
read.
and b.
Parameters
a(in): half2. Is only beingb(in): half2..
Returns
- half2
- The vector result of unordered if-equal comparison of vectors
a
Half Arithmetic Functions
__hfma
__half __hfma(const __half a, const __half b, const __half c)
Performs half fused multiply-add in round-to-nearest-even mode.
Performs half multiply on inputs a and b, then performs a half add of the result with c, rounding the result once in round-to-nearest-even mode.
Parameters
a(in): halfb(in): halfc(in): half
Returns
- half
- The result of fused multiply-add operation on
a,b, andc.
Half2 Arithmetic Functions
__hfma2
__half2 __hfma2(const __half2 a, const __half2 b, const __half2 c)
Performs half2 vector fused multiply-add in round-to-nearest-even mode.
Performs half2 vector multiply on inputs a and b, then performs a half2 vector add of the result with c, rounding the result once in round-to-nearest-even mode.
Parameters
a(in): half2..b(in): half2..c(in): half2..
Returns
- half2
- The result of elementwise fused multiply-add operation on vectors
a,b, andc.
__hfma2_sat
__half2 __hfma2_sat(const __half2 a, const __half2 b, const __half2 c)
Performs half2 vector fused multiply-add in round-to-nearest-even mode, with saturation to [0.0, 1.0].
Performs half2 vector multiply on inputs a and b, then performs a half2 vector add of the result with c, rounding the result once in round-to-nearest-even mode, and clamps the results to range [0.0, 1.0]. NaN results are flushed to +0.0.
Parameters
a(in): half2..b(in): half2..c(in): half2..
Returns
- half2
- The result of elementwise fused multiply-add operation on vectors
a,b, andc, with respect to saturation.
Half Arithmetic Functions
__hfma_sat
__half __hfma_sat(const __half a, const __half b, const __half c)
Performs half fused multiply-add in round-to-nearest-even mode, with saturation to [0.0, 1.0].
Performs half multiply on inputs a and b, then performs a half add of the result with c, rounding the result once in round-to-nearest-even mode, and clamps the result to range [0.0, 1.0]. NaN results are flushed to +0.0.
Parameters
a(in): halfb(in): halfc(in): half
Returns
- half
- The result of fused multiply-add operation on
a,b, andc, with respect to saturation.
Half Comparison Functions
__hge
bool __hge(const __half a, const __half b)
Performs half greater-equal comparison.
Performs half greater-equal comparison of inputs a and b. NaN inputs generate false results.
Parameters
a(in): halfb(in): half
Returns
- bool
- The boolean result of greater-equal comparison of
aandb.
Half2 Comparison Functions 115
__hge2
__half2 __hge2(const __half2 a, const __half2 b)
Performs half2 vector greater-equal comparison.
Performs half2 vector greater-equal comparison of inputs a and b. The corresponding half results are set to 1.0 for true, or 0.0 for false. NaN inputs generate false results.
read.
Parameters
a(in): half2. Is only beingb(in): half2..
Returns
- half2
- The vector result of greater-equal comparison of vectors
aandb.
__hge2_mask
unsigned __hge2_mask(const __half2 a, const __half2 b)
Performs half2 vector greater-equal comparison.
Performs half2 vector greater-equal comparison of inputs a and b. The corresponding half unsigned bits are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate true results.
only.
vectors a and b.
Parameters
a(in): half2. Isb(in): half2..
Returns
- usigned int
- The
half2vector result of greater-equal comparison of
Half Comparison Functions
__hgeu
bool __hgeu(const __half a, const __half b)
Performs half unordered greater-equal comparison.
Performs half greater-equal comparison of inputs a and b. NaN inputs generate true results.
and b.
Parameters
a(in): halfb(in): half
Returns
- bool
- The boolean result of unordered greater-equal comparison of
a
Half2 Comparison Functions 115
__hgeu2
__half2 __hgeu2(const __half2 a, const __half2 b)
Performs half2 vector unordered greater-equal comparison.
Performs half2 vector greater-equal comparison of inputs a and b. The corresponding half results are set to 1.0 for true, or 0.0 for false. NaN inputs generate true results.
read.
vectors a and b.
Parameters
a(in): half2. Is only beingb(in): half2..
Returns
- half2
- The
half2vector result of unordered greater-equal comparison of
__hgeu2_mask
unsigned __hgeu2_mask(const __half2 a, const __half2 b)
Performs half2 vector unordered greater-equal comparison.
Performs half2 vector greater-equal comparison of inputs a and b. The corresponding half unsigned bits are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate true results.
only.
vectors a and b.
Parameters
a(in): half2. Isb(in): half2..
Returns
- usigned int
- The
half2vector result of unordered greater-equal comparison of
Half Comparison Functions
__hgt
bool __hgt(const __half a, const __half b)
Performs half greater-than comparison.
Performs half greater-than comparison of inputs a and b. NaN inputs generate false results.
Parameters
a(in): halfb(in): half
Returns
- bool
- The boolean result of greater-than comparison of
aandb.
Half2 Comparison Functions 115
__hgt2
__half2 __hgt2(const __half2 a, const __half2 b)
Performs half2 vector greater-than comparison.
Performs half2 vector greater-than comparison of inputs a and b. The corresponding half results are set to 1.0 for true, or 0.0 for false. NaN inputs generate false results.
read.
Parameters
a(in): half2. Is only beingb(in): half2..
Returns
- half2
- The vector result of greater-than comparison of vectors
aandb.
__hgt2_mask
unsigned __hgt2_mask(const __half2 a, const __half2 b)
Performs half2 vector greater-than comparison.
Performs half2 vector greater-than comparison of inputs a and b. The corresponding half unsigned bits are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate false results.
only.
Parameters
a(in): half2. Isb(in): half2..
Returns
- unsigned int
- The vector result of greater-than comparison of vectors
aandb.
Half Comparison Functions
__hgtu
bool __hgtu(const __half a, const __half b)
Performs half unordered greater-than comparison.
Performs half greater-than comparison of inputs a and b. NaN inputs generate true results.
and b.
Parameters
a(in): halfb(in): half
Returns
- bool
- The boolean result of unordered greater-than comparison of
a
Half2 Comparison Functions 115
__hgtu2
__half2 __hgtu2(const __half2 a, const __half2 b)
Performs half2 vector unordered greater-than comparison.
Performs half2 vector greater-than comparison of inputs a and b. The corresponding half results are set to 1.0 for true, or 0.0 for false. NaN inputs generate true results.
read.
vectors a and b.
Parameters
a(in): half2. Is only beingb(in): half2..
Returns
- half2
- The
half2vector result of unordered greater-than comparison of
__hgtu2_mask
unsigned __hgtu2_mask(const __half2 a, const __half2 b)
Performs half2 vector unordered greater-than comparison.
Performs half2 vector greater-than comparison of inputs a and b. The corresponding half unsigned bits are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate true results.
only.
Parameters
a(in): half2. Isb(in): half2..
Returns
- unsigned int
- The vector result of unordered greater-than comparison of vectors
aandb.
Half Precision Conversion and Data Movement 71
__high2float
float __high2float(const __half2 a)
Converts high 16 bits of half2 to float and returns the result
Converts high 16 bits of half2 input a to 32-bit floating-point number and returns the result.
being.
Parameters
a(in): half2. Is only
Returns
- float
- The high 16 bits of
aconverted to float.
__high2half
__half __high2half(const __half2 a)
Returns high 16 bits of half2 input.
Returns high 16 bits of half2 input a.
Parameters
a(in): half2..
Returns
- half
- The high 16 bits of the input.
__high2half2
__half2 __high2half2(const __half2 a)
Extracts high 16 bits from half2 input.
Extracts high 16 bits from half2 input a and returns a new half2 number which has both halves equal to the extracted bits.
- half2..
Parameters
a(in):
Returns
- half2
- The half2 with both halves equal to the high 16 bits of the input.
__highs2half2
__half2 __highs2half2(const __half2 a, const __half2 b)
Extracts high 16 bits from each of the two half2 inputs and combines into one half2 number.
Extracts high 16 bits from each of the two half2 inputs and combines into one half2 number. High 16 bits from input a is stored in low 16 bits of the return value, high 16 bits from input b is stored in high 16 bits of the return value.
Parameters
a(in): half2..b(in): half2..
Returns
- half2
- The high 16 bits of
aand ofb.
Half Comparison Functions
__hisinf
int __hisinf(const __half a)
Checks if the input half number is infinite.
Checks if the input half number a is infinite.
Parameters
a(in): half
Returns
- int
- -1 iff
ais equal to negative infinity, - 1 iff
ais equal to positive infinity, - 0 otherwise.
__hisnan
bool __hisnan(const __half a)
Determine whether half argument is a NaN.
Determine whether half value a is a NaN.
Parameters
a(in): half
Returns
- bool
- true iff argument is NaN.
Half2 Comparison Functions 115
__hisnan2
__half2 __hisnan2(const __half2 a)
Determine whether half2 argument is a NaN.
Determine whether each half of input half2 number a is a NaN.
1.0 for NaN, 0.0 otherwise.
Parameters
a(in): half2..
Returns
- half2
- The half2 with the corresponding
halfresults set to
Half Comparison Functions
__hle
bool __hle(const __half a, const __half b)
Performs half less-equal comparison.
Performs half less-equal comparison of inputs a and b. NaN inputs generate false results.
Parameters
a(in): halfb(in): half
Returns
- bool
- The boolean result of less-equal comparison of
aandb.
Half2 Comparison Functions 115
__hle2
__half2 __hle2(const __half2 a, const __half2 b)
Performs half2 vector less-equal comparison.
Performs half2 vector less-equal comparison of inputs a and b. The corresponding half results are set to 1.0 for true, or 0.0 for false. NaN inputs generate false results.
read.
Parameters
a(in): half2. Is only being read.b(in): half2. Is only being read.
Returns
- half2
- The
half2result of less-equal comparison of vectorsaandb.
__hle2_mask
unsigned __hle2_mask(const __half2 a, const __half2 b)
Performs half2 vector less-equal comparison.
Performs half2 vector less-equal comparison of inputs a and b. The corresponding half unsigned bits are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate false results.
only.
Parameters
a(in): half2. Isb(in): half2..
Returns
- unsigned int
- The
half2result of less-equal comparison of vectorsaandb.
Half Comparison Functions
__hleu
bool __hleu(const __half a, const __half b)
Performs half unordered less-equal comparison.
Performs half less-equal comparison of inputs a and b. NaN inputs generate true results.
b.
Parameters
a(in): halfb(in): half
Returns
- bool
- The boolean result of unordered less-equal comparison of
aand
Half2 Comparison Functions 115
__hleu2
__half2 __hleu2(const __half2 a, const __half2 b)
Performs half2 vector unordered less-equal comparison.
Performs half2 vector less-equal comparison of inputs a and b. The corresponding half results are set to 1.0 for true, or 0.0 for false. NaN inputs generate true results.
and b.
Parameters
a(in): half2..b(in): half2..
Returns
- half2
- The vector result of unordered less-equal comparison of vectors
a
__hleu2_mask
unsigned __hleu2_mask(const __half2 a, const __half2 b)
Performs half2 vector unordered less-equal comparison.
Performs half2 vector less-equal comparison of inputs a and b. The corresponding half unsigned bits are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate true results.
and b.
Parameters
a(in): half2..b(in): half2..
Returns
- unsigned int
- The vector result of unordered less-equal comparison of vectors
a
Half Comparison Functions
__hlt
bool __hlt(const __half a, const __half b)
Performs half less-than comparison.
Performs half less-than comparison of inputs a and b. NaN inputs generate false results.
Parameters
a(in): halfb(in): half
Returns
- bool
- The boolean result of less-than comparison of
aandb.
Half2 Comparison Functions 115
__hlt2
__half2 __hlt2(const __half2 a, const __half2 b)
Performs half2 vector less-than comparison.
Performs half2 vector less-than comparison of inputs a and b. The corresponding half results are set to 1.0 for true, or 0.0 for false. NaN inputs generate false results.
read.
b.
Parameters
a(in): half2. Is only beingb(in): half2..
Returns
- half2
- The half2 vector result of less-than comparison of vectors
aand
__hlt2_mask
unsigned __hlt2_mask(const __half2 a, const __half2 b)
Performs half2 vector less-than comparison.
Performs half2 vector less-than comparison of inputs a and b. The corresponding half unsigned bits are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate false results.
only.
b.
Parameters
a(in): half2. Isb(in): half2..
Returns
- unsigned int
- The half2 vector result of less-than comparison of vectors
aand
Half Comparison Functions
__hltu
bool __hltu(const __half a, const __half b)
Performs half unordered less-than comparison.
Performs half less-than comparison of inputs a and b. NaN inputs generate true results.
b.
Parameters
a(in): halfb(in): half
Returns
- bool
- The boolean result of unordered less-than comparison of
aand
Half2 Comparison Functions 115
__hltu2
__half2 __hltu2(const __half2 a, const __half2 b)
Performs half2 vector unordered less-than comparison.
Performs half2 vector less-than comparison of inputs a and b. The corresponding half results are set to 1.0 for true, or 0.0 for false. NaN inputs generate true results.
read.
and b.
Parameters
a(in): half2. Is only beingb(in): half2..
Returns
- half2
- The vector result of unordered less-than comparison of vectors
a
__hltu2_mask
unsigned __hltu2_mask(const __half2 a, const __half2 b)
Performs half2 vector unordered less-than comparison.
Performs half2 vector less-than comparison of inputs a and b. The corresponding half unsigned bits are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate true results.
only.
and b.
Parameters
a(in): half2. Isb(in): half2..
Returns
- unsigned int
- The vector result of unordered less-than comparison of vectors
a
Half Comparison Functions
__hmax
__half __hmax(const __half a, const __half b)
Calculates half maximum of two input values.
Calculates half max(a, b) defined as a > b ? a : b. - If either of inputs is NaN, the other input is returned. - If both inputs are NaNs, then canonical NaN is returned. - If values of both inputs are 0.0, then +0.0 > -0.0.
Parameters
a(in): halfb(in): half
Returns
- half
Half2 Comparison Functions 115
__hmax2
__half2 __hmax2(const __half2 a, const __half2 b)
Calculates half2 vector maximum of two inputs.
Calculates half2 vector max(a, b). Elementwise half operation is defined as a > b ? a : b. - If either of inputs is NaN, the other input is returned. - If both inputs are NaNs, then canonical NaN is returned. - If values of both inputs are 0.0, then +0.0 > -0.0.
Parameters
a(in): half2..b(in): half2..
Returns
- half2
- The result of elementwise maximum of vectors
aandb
__hmax2_nan
__half2 __hmax2_nan(const __half2 a, const __half2 b)
Calculates half2 vector maximum of two inputs, NaNs pass through.
Calculates half2 vector max(a, b). Elementwise half operation is defined as a > b ? a : b. - If either of inputs is NaN, then canonical NaN is returned. - If values of both inputs are 0.0, then +0.0 > -0.0.
NaNs pass through
Parameters
a(in): half2..b(in): half2..
Returns
- half2
- The result of elementwise maximum of vectors
aandb, with
Half Comparison Functions
__hmax_nan
__half __hmax_nan(const __half a, const __half b)
Calculates half maximum of two input values, NaNs pass through.
Calculates half max(a, b) defined as a > b ? a : b. - If either of inputs is NaN, then canonical NaN is returned. - If values of both inputs are 0.0, then +0.0 > -0.0.
Parameters
a(in): halfb(in): half
Returns
- half
__hmin
__half __hmin(const __half a, const __half b)
Calculates half minimum of two input values.
Calculates half min(a, b) defined as a < b ? a : b. - If either of inputs is NaN, the other input is returned. - If both inputs are NaNs, then canonical NaN is returned. - If values of both inputs are 0.0, then +0.0 > -0.0.
Parameters
a(in): halfb(in): half
Returns
- half
Half2 Comparison Functions 115
__hmin2
__half2 __hmin2(const __half2 a, const __half2 b)
Calculates half2 vector minimum of two inputs.
Calculates half2 vector min(a, b). Elementwise half operation is defined as a < b ? a : b. - If either of inputs is NaN, the other input is returned. - If both inputs are NaNs, then canonical NaN is returned. - If values of both inputs are 0.0, then +0.0 > -0.0.
Parameters
a(in): half2..b(in): half2..
Returns
- half2
- The result of elementwise minimum of vectors
aandb
__hmin2_nan
__half2 __hmin2_nan(const __half2 a, const __half2 b)
Calculates half2 vector minimum of two inputs, NaNs pass through.
Calculates half2 vector min(a, b). Elementwise half operation is defined as a < b ? a : b. - If either of inputs is NaN, then canonical NaN is returned. - If values of both inputs are 0.0, then +0.0 > -0.0.
NaNs pass through
Parameters
a(in): half2..b(in): half2..
Returns
- half2
- The result of elementwise minimum of vectors
aandb, with
Half Comparison Functions
__hmin_nan
__half __hmin_nan(const __half a, const __half b)
Calculates half minimum of two input values, NaNs pass through.
Calculates half min(a, b) defined as a < b ? a : b. - If either of inputs is NaN, then canonical NaN is returned. - If values of both inputs are 0.0, then +0.0 > -0.0.
Parameters
a(in): halfb(in): half
Returns
- half
Half Arithmetic Functions
__hmul
__half __hmul(const __half a, const __half b)
Performs half multiplication in round-to-nearest-even mode.
Performs half multiplication of inputs a and b, in round-to-nearest mode.
Parameters
a(in): halfb(in): half
Returns
- half
- The result of multiplying
aandb.
Half2 Arithmetic Functions
__hmul2
__half2 __hmul2(const __half2 a, const __half2 b)
Performs half2 vector multiplication in round-to-nearest-even mode.
Performs half2 vector multiplication of inputs a and b, in round-to-nearest-even mode.
Parameters
a(in): half2..b(in): half2..
Returns
- half2
- The result of elementwise multiplying the vectors
aandb.
__hmul2_rn
__half2 __hmul2_rn(const __half2 a, const __half2 b)
Performs half2 vector multiplication in round-to-nearest-even mode. Prevents floating-point contractions of mul+add or sub into fma.
Performs half2 vector multiplication of inputs a and b, in round-to-nearest-even mode.
Parameters
a(in): half2..b(in): half2..
Returns
- half2
- The result of elementwise multiplying the vectors
aandb.
__hmul2_sat
__half2 __hmul2_sat(const __half2 a, const __half2 b)
Performs half2 vector multiplication in round-to-nearest-even mode, with saturation to [0.0, 1.0].
Performs half2 vector multiplication of inputs a and b, in round-to-nearest-even mode, and clamps the results to range [0.0, 1.0]. NaN results are flushed to +0.0.
with respect to saturation.
Parameters
a(in): half2..b(in): half2..
Returns
- half2
- The result of elementwise multiplication of vectors
aandb,
Half Arithmetic Functions
__hmul_rn
__half __hmul_rn(const __half a, const __half b)
Performs half multiplication in round-to-nearest-even mode. Prevents floating-point contractions of mul+add or sub into fma.
Performs half multiplication of inputs a and b, in round-to-nearest mode.
Parameters
a(in): halfb(in): half
Returns
- half
- The result of multiplying
aandb.
__hmul_sat
__half __hmul_sat(const __half a, const __half b)
Performs half multiplication in round-to-nearest-even mode, with saturation to [0.0, 1.0].
Performs half multiplication of inputs a and b, in round-to-nearest mode, and clamps the result to range [0.0, 1.0]. NaN results are flushed to +0.0.
half
Parameters
a(in): halfb(in):
Returns
- half
- The result of multiplying
aandb, with respect to saturation.
Half Comparison Functions
__hne
bool __hne(const __half a, const __half b)
Performs half not-equal comparison.
Performs half not-equal comparison of inputs a and b. NaN inputs generate false results.
Parameters
a(in): halfb(in): half
Returns
- bool
- The boolean result of not-equal comparison of
aandb.
Half2 Comparison Functions 115
__hne2
__half2 __hne2(const __half2 a, const __half2 b)
Performs half2 vector not-equal comparison.
Performs half2 vector not-equal comparison of inputs a and b. The corresponding half results are set to 1.0 for true, or 0.0 for false. NaN inputs generate false results.
read.
Parameters
a(in): half2. Is only beingb(in): half2..
Returns
- half2
- The vector result of not-equal comparison of vectors
aandb.
__hne2_mask
unsigned __hne2_mask(const __half2 a, const __half2 b)
Performs half2 vector not-equal comparison.
Performs half2 vector not-equal comparison of inputs a and b. The corresponding half unsigned bits are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate false results.
only.
Parameters
a(in): half2. Isb(in): half2..
Returns
- unsigned int
- The vector result of not-equal comparison of vectors
aandb.
Half Arithmetic Functions
__hneg
__half __hneg(const __half a)
Negates input half number and returns the result.
Negates input half number and returns the result.
Parameters
a(in): half
Returns
- half
- minus a
Half2 Arithmetic Functions
__hneg2
__half2 __hneg2(const __half2 a)
Negates both halves of the input half2 number and returns the result.
Negates both halves of the input half2 number a and returns the result.
half2..
Parameters
a(in):
Returns
- half2
- Returns
awith both halves negated.
Half Comparison Functions
__hneu
bool __hneu(const __half a, const __half b)
Performs half unordered not-equal comparison.
Performs half not-equal comparison of inputs a and b. NaN inputs generate true results.
b.
Parameters
a(in): halfb(in): half
Returns
- bool
- The boolean result of unordered not-equal comparison of
aand
Half2 Comparison Functions 115
__hneu2
__half2 __hneu2(const __half2 a, const __half2 b)
Performs half2 vector unordered not-equal comparison.
Performs half2 vector not-equal comparison of inputs a and b. The corresponding half results are set to 1.0 for true, or 0.0 for false. NaN inputs generate true results.
read.
and b.
Parameters
a(in): half2. Is only beingb(in): half2..
Returns
- half2
- The vector result of unordered not-equal comparison of vectors
a
__hneu2_mask
unsigned __hneu2_mask(const __half2 a, const __half2 b)
Performs half2 vector not-equal comparison.
Performs half2 vector unordered not-equal comparison of inputs a and b. The corresponding half unsigned bits are set to 0xFFFF for true, or 0x0 for false. NaN inputs generate false results.
half2..
and b.
Parameters
a(in):b(in): half2..
Returns
- unsigned int
- The vector result of unordered not-equal comparison of vectors
a
Half Arithmetic Functions
__hsub
__half __hsub(const __half a, const __half b)
Performs half subtraction in round-to-nearest-even mode.
Subtracts half input b from input a in round-to-nearest mode.
Parameters
a(in): halfb(in): half
Returns
- half
- The result of subtracting
bfroma.
Half2 Arithmetic Functions
__hsub2
__half2 __hsub2(const __half2 a, const __half2 b)
Performs half2 vector subtraction in round-to-nearest-even mode.
Subtracts half2 input vector b from input vector a in round-to-nearest-even mode.
Parameters
a(in): half2..b(in): half2..
Returns
- half2
- The subtraction of vector
bfroma.
__hsub2_rn
__half2 __hsub2_rn(const __half2 a, const __half2 b)
Performs half2 vector subtraction in round-to-nearest-even mode. Prevents floating-point contractions of mul+sub into fma.
Subtracts half2 input vector b from input vector a in round-to-nearest-even mode.
Parameters
a(in): half2..b(in): half2..
Returns
- half2
- The subtraction of vector
bfroma.
__hsub2_sat
__half2 __hsub2_sat(const __half2 a, const __half2 b)
Performs half2 vector subtraction in round-to-nearest-even mode, with saturation to [0.0, 1.0].
Subtracts half2 input vector b from input vector a in round-to-nearest-even mode, and clamps the results to range [0.0, 1.0]. NaN results are flushed to +0.0.
Parameters
a(in): half2..b(in): half2..
Returns
- half2
- The subtraction of vector
bfroma, with respect to saturation.
Half Arithmetic Functions
__hsub_rn
__half __hsub_rn(const __half a, const __half b)
Performs half subtraction in round-to-nearest-even mode. Prevents floating-point contractions of mul+sub into fma.
Subtracts half input b from input a in round-to-nearest mode.
Parameters
a(in): halfb(in): half
Returns
- half
- The result of subtracting
bfroma.
__hsub_sat
__half __hsub_sat(const __half a, const __half b)
Performs half subtraction in round-to-nearest-even mode, with saturation to [0.0, 1.0].
Subtracts half input b from input a in round-to-nearest mode, and clamps the result to range [0.0, 1.0]. NaN results are flushed to +0.0.
saturation.
Parameters
a(in): halfb(in): half
Returns
- half
- The result of subtraction of
bfroma, with respect to
Half Precision Conversion and Data Movement 71
__int2half_rd
__half __int2half_rd(const int i)
Convert a signed integer to a half in round-down mode.
Convert the signed integer value i to a half-precision floating-point value in round-down mode.
read.
Parameters
i(in): int. Is only being
Returns
- half
- i converted to half.
__int2half_rn
__half __int2half_rn(const int i)
Convert a signed integer to a half in round-to-nearest-even mode.
Convert the signed integer value i to a half-precision floating-point value in round-to-nearest-even mode.
only.
Parameters
i(in): int. Is
Returns
- half
- i converted to half.
__int2half_ru
__half __int2half_ru(const int i)
Convert a signed integer to a half in round-up mode.
Convert the signed integer value i to a half-precision floating-point value in round-up mode.
read.
Parameters
i(in): int. Is only being
Returns
- half
- i converted to half.
__int2half_rz
__half __int2half_rz(const int i)
Convert a signed integer to a half in round-towards-zero mode.
Convert the signed integer value i to a half-precision floating-point value in round-towards-zero mode.
being.
Parameters
i(in): int. Is only
Returns
- half
- i converted to half.
__ldca
__half2 __ldca(const __half2 *const ptr)
Loads a value from the specified global memory address.
Parameters
ptr(in): source memory location to load from
Returns
- The value loaded from
ptr
__ldca
__half __ldca(const __half *const ptr)
Loads a value from the specified global memory address.
Parameters
ptr(in): source memory location to load from
Returns
- The value loaded from
ptr
__ldcg
__half2 __ldcg(const __half2 *const ptr)
Loads a value from the specified global memory address.
Parameters
ptr(in): source memory location to load from
Returns
- The value loaded from
ptr
__ldcg
__half __ldcg(const __half *const ptr)
Loads a value from the specified global memory address.
Parameters
ptr(in): source memory location to load from
Returns
- The value loaded from
ptr
__ldcs
__half2 __ldcs(const __half2 *const ptr)
Loads a value from the specified global memory address.
Parameters
ptr(in): source memory location to load from
Returns
- The value loaded from
ptr
__ldcs
__half __ldcs(const __half *const ptr)
Loads a value from the specified global memory address.
Parameters
ptr(in): source memory location to load from
Returns
- The value loaded from
ptr
__ldcv
__half2 __ldcv(const __half2 *const ptr)
Loads a value from the specified global memory address.
Parameters
ptr(in): source memory location to load from
Returns
- The value loaded from
ptr
__ldcv
__half __ldcv(const __half *const ptr)
Loads a value from the specified global memory address.
Parameters
ptr(in): source memory location to load from
Returns
- The value loaded from
ptr
__ldg
__half2 __ldg(const __half2 *const ptr)
Loads a value from the specified global memory address.
Parameters
ptr(in): source memory location to load from
Returns
- The value loaded from
ptr
__ldg
__half __ldg(const __half *const ptr)
Loads a value from the specified global memory address.
Parameters
ptr(in): source memory location to load from
Returns
- The value loaded from
ptr
__ll2half_rd
__half __ll2half_rd(const long long int i)
Convert a signed 64-bit integer to a half in round-down mode.
Convert the signed 64-bit integer value i to a half-precision floating-point value in round-down mode.
only.
Parameters
i(in): long long int. Is
Returns
- half
- i converted to half.
__ll2half_rn
__half __ll2half_rn(const long long int i)
Convert a signed 64-bit integer to a half in round-to-nearest-even mode.
Convert the signed 64-bit integer value i to a half-precision floating-point value in round-to-nearest-even mode.
int
Parameters
i(in): long long
Returns
- half
- i converted to half.
Half Precision Conversion and Data Movement 73
__ll2half_ru
__half __ll2half_ru(const long long int i)
Convert a signed 64-bit integer to a half in round-up mode.
Convert the signed 64-bit integer value i to a half-precision floating-point value in round-up mode.
being.
Parameters
i(in): long long int. Is only
Returns
- half
- i converted to half.
__ll2half_rz
__half __ll2half_rz(const long long int i)
Convert a signed 64-bit integer to a half in round-towards-zero mode.
Convert the signed 64-bit integer value i to a half-precision floating-point value in round-towards-zero mode.
int
Parameters
i(in): long long
Returns
- half
- i converted to half.
__low2float
float __low2float(const __half2 a)
Converts low 16 bits of half2 to float and returns the result
Converts low 16 bits of half2 input a to 32-bit floating-point number and returns the result.
Parameters
a(in): half2..
Returns
- float
- The low 16 bits of
aconverted to float.
__low2half
__half __low2half(const __half2 a)
Returns low 16 bits of half2 input.
Returns low 16 bits of half2 input a.
Parameters
a(in): half2..
Returns
- half
- Returns
halfwhich contains low 16 bits of the inputa.
__low2half2
__half2 __low2half2(const __half2 a)
Extracts low 16 bits from half2 input.
Extracts low 16 bits from half2 input a and returns a new half2 number which has both halves equal to the extracted bits.
- half2..
Parameters
a(in):
Returns
- half2
- The half2 with both halves equal to the low 16 bits of the input.
__lowhigh2highlow
__half2 __lowhigh2highlow(const __half2 a)
Swaps both halves of the half2 input.
Swaps both halves of the half2 input and returns a new half2 number with swapped halves.
Parameters
a(in): half2..
Returns
- half2
- a with its halves being swapped.
__lows2half2
__half2 __lows2half2(const __half2 a, const __half2 b)
Extracts low 16 bits from each of the two half2 inputs and combines into one half2 number.
Extracts low 16 bits from each of the two half2 inputs and combines into one half2 number. Low 16 bits from input a is stored in low 16 bits of the return value, low 16 bits from input b is stored in high 16 bits of the return value.
Parameters
a(in): half2..b(in): half2..
Returns
- half2
- The low 16 bits of
aand ofb.
__shfl_down_sync
__half2 __shfl_down_sync(const unsigned mask, const __half2 var, const unsigned int delta, const int width)
Exchange a variable between threads within a warp. Copy from a thread with higher ID relative to the caller.
Calculates a source thread ID by adding delta to the caller's thread ID. The value of var held by the resulting thread ID is returned: this has the effect of shifting var down the warp by delta threads. If width is less than warpSize then each subsection of the warp behaves as a separate entity with a starting logical thread ID of 0. As for __shfl_up_sync(), the ID number of the source thread will not wrap around the value of width and so the upper delta threads will remain unchanged.
int
as half2. If the source thread ID is out of range or the source thread has exited, the calling thread's own var is returned.
Parameters
mask(in): unsignedvar(in): half2..delta(in): intwidth(in): int
Returns
- Returns the 4-byte word referenced by var from the source thread ID
__shfl_down_sync
__half __shfl_down_sync(const unsigned mask, const __half var, const unsigned int delta, const int width)
Exchange a variable between threads within a warp. Copy from a thread with higher ID relative to the caller.
Calculates a source thread ID by adding delta to the caller's thread ID. The value of var held by the resulting thread ID is returned: this has the effect of shifting var down the warp by delta threads. If width is less than warpSize then each subsection of the warp behaves as a separate entity with a starting logical thread ID of 0. As for __shfl_up_sync(), the ID number of the source thread will not wrap around the value of width and so the upper delta threads will remain unchanged.
int
as half2. If the source thread ID is out of range or the source thread has exited, the calling thread's own var is returned.
Parameters
mask(in): unsignedvar(in): half2..delta(in): intwidth(in): int
Returns
- Returns the 4-byte word referenced by var from the source thread ID
__shfl_sync
__half2 __shfl_sync(const unsigned mask, const __half2 var, const int delta, const int width)
Exchange a variable between threads within a warp. Direct copy from indexed thread.
Returns the value of var held by the thread whose ID is given by delta. If width is less than warpSize then each subsection of the warp behaves as a separate entity with a starting logical thread ID of 0. If delta is outside the range [0:width-1], the value returned corresponds to the value of var held by the delta modulo width (i.e. within the same subsection). width must have a value which is a power of 2; results are undefined if width is not a power of 2, or is a number greater than warpSize.
unsigned int
as half2. If the source thread ID is out of range or the source thread has exited, the calling thread's own var is returned.
Parameters
mask(in):var(in): half2..delta(in): intwidth(in): int
Returns
- Returns the 4-byte word referenced by var from the source thread ID
__shfl_sync
__half __shfl_sync(const unsigned mask, const __half var, const int delta, const int width)
Exchange a variable between threads within a warp. Direct copy from indexed thread.
Returns the value of var held by the thread whose ID is given by delta. If width is less than warpSize then each subsection of the warp behaves as a separate entity with a starting logical thread ID of 0. If delta is outside the range [0:width-1], the value returned corresponds to the value of var held by the delta modulo width (i.e. within the same subsection). width must have a value which is a power of 2; results are undefined if width is not a power of 2, or is a number greater than warpSize.
unsigned int
as half2. If the source thread ID is out of range or the source thread has exited, the calling thread's own var is returned.
Parameters
mask(in):var(in): half2..delta(in): intwidth(in): int
Returns
- Returns the 4-byte word referenced by var from the source thread ID
__shfl_up_sync
__half2 __shfl_up_sync(const unsigned mask, const __half2 var, const unsigned int delta, const int width)
Exchange a variable between threads within a warp. Copy from a thread with lower ID relative to the caller.
Calculates a source thread ID by subtracting delta from the caller's lane ID. The value of var held by the resulting lane ID is returned: in effect, var is shifted up the warp by delta threads. If width is less than warpSize then each subsection of the warp behaves as a separate entity with a starting logical thread ID of 0. The source thread index will not wrap around the value of width, so effectively the lower delta threads will be unchanged. width must have a value which is a power of 2; results are undefined if width is not a power of 2, or is a number greater than warpSize.
unsigned int
as half2. If the source thread ID is out of range or the source thread has exited, the calling thread's own var is returned.
Parameters
mask(in):var(in): half2..delta(in): intwidth(in): int
Returns
- Returns the 4-byte word referenced by var from the source thread ID
__shfl_up_sync
__half __shfl_up_sync(const unsigned mask, const __half var, const unsigned int delta, const int width)
Exchange a variable between threads within a warp. Copy from a thread with lower ID relative to the caller.
Calculates a source thread ID by subtracting delta from the caller's lane ID. The value of var held by the resulting lane ID is returned: in effect, var is shifted up the warp by delta threads. If width is less than warpSize then each subsection of the warp behaves as a separate entity with a starting logical thread ID of 0. The source thread index will not wrap around the value of width, so effectively the lower delta threads will be unchanged. width must have a value which is a power of 2; results are undefined if width is not a power of 2, or is a number greater than warpSize.
unsigned int
as half2. If the source thread ID is out of range or the source thread has exited, the calling thread's own var is returned.
Parameters
mask(in):var(in): half2..delta(in): intwidth(in): int
Returns
- Returns the 4-byte word referenced by var from the source thread ID
__shfl_xor_sync
__half2 __shfl_xor_sync(const unsigned mask, const __half2 var, const int delta, const int width)
Exchange a variable between threads within a warp. Copy from a thread based on bitwise XOR of own thread ID.
Calculates a source thread ID by performing a bitwise XOR of the caller's thread ID with mask: the value of var held by the resulting thread ID is returned. If width is less than warpSize then each group of width consecutive threads are able to access elements from earlier groups of threads, however if they attempt to access elements from later groups of threads their own value of var will be returned. This mode implements a butterfly addressing pattern such as is used in tree reduction and broadcast.
as half2. If the source thread ID is out of range or the source thread has exited, the calling thread's own var is returned.
Parameters
mask(in): unsigned intvar(in): half2..delta(in): intwidth(in): int
Returns
- Returns the 4-byte word referenced by var from the source thread ID
__shfl_xor_sync
__half __shfl_xor_sync(const unsigned mask, const __half var, const int delta, const int width)
Exchange a variable between threads within a warp. Copy from a thread based on bitwise XOR of own thread ID.
Calculates a source thread ID by performing a bitwise XOR of the caller's thread ID with mask: the value of var held by the resulting thread ID is returned. If width is less than warpSize then each group of width consecutive threads are able to access elements from earlier groups of threads, however if they attempt to access elements from later groups of threads their own value of var will be returned. This mode implements a butterfly addressing pattern such as is used in tree reduction and broadcast.
as half2. If the source thread ID is out of range or the source thread has exited, the calling thread's own var is returned.
Parameters
mask(in): unsigned intvar(in): half2..delta(in): intwidth(in): int
Returns
- Returns the 4-byte word referenced by var from the source thread ID
__short2half_rd
__half __short2half_rd(const short int i)
Convert a signed short integer to a half in round-down mode.
Convert the signed short integer value i to a half-precision floating-point value in round-down mode.
being.
Parameters
i(in): short int. Is only
Returns
- half
- i converted to half.
__short2half_rn
__half __short2half_rn(const short int i)
Convert a signed short integer to a half in round-to-nearest-even mode.
Convert the signed short integer value i to a half-precision floating-point value in round-to-nearest-even mode.
Is only.
Parameters
i(in): short int.
Returns
- half
- i converted to half.
__short2half_ru
__half __short2half_ru(const short int i)
Convert a signed short integer to a half in round-up mode.
Convert the signed short integer value i to a half-precision floating-point value in round-up mode.
being.
Parameters
i(in): short int. Is only
Returns
- half
- i converted to half.
__short2half_rz
__half __short2half_rz(const short int i)
Convert a signed short integer to a half in round-towards-zero mode.
Convert the signed short integer value i to a half-precision floating-point value in round-towards-zero mode.
only.
Parameters
i(in): short int. Is
Returns
- half
- i converted to half.
__short_as_half
__half __short_as_half(const short int i)
Reinterprets bits in a signed short integer as a half.
Reinterprets the bits in the signed short integer i as a half-precision floating-point number.
Parameters
i(in): short int
Returns
- half
- The reinterpreted value.
__stcg
void __stcg(__half2 *const ptr, const __half2 value)
Stores a value to the specified global memory address.
Parameters
ptr(out): destination memory location for the storevalue(in): value to write toptr
__stcg
void __stcg(__half *const ptr, const __half value)
Stores a value to the specified global memory address.
Parameters
ptr(out): destination memory location for the storevalue(in): value to write toptr
__stcs
void __stcs(__half2 *const ptr, const __half2 value)
Stores a value to the specified global memory address.
Parameters
ptr(out): destination memory location for the storevalue(in): value to write toptr
__stcs
void __stcs(__half *const ptr, const __half value)
Stores a value to the specified global memory address.
Parameters
ptr(out): destination memory location for the storevalue(in): value to write toptr
__stwb
void __stwb(__half2 *const ptr, const __half2 value)
Stores a value to the specified global memory address.
Parameters
ptr(out): destination memory location for the storevalue(in): value to write toptr
__stwb
void __stwb(__half *const ptr, const __half value)
Stores a value to the specified global memory address.
Parameters
ptr(out): destination memory location for the storevalue(in): value to write toptr
__stwt
void __stwt(__half2 *const ptr, const __half2 value)
Stores a value to the specified global memory address.
Parameters
ptr(out): destination memory location for the storevalue(in): value to write toptr
__stwt
void __stwt(__half *const ptr, const __half value)
Stores a value to the specified global memory address.
Parameters
ptr(out): destination memory location for the storevalue(in): value to write toptr
__uint2half_rd
__half __uint2half_rd(const unsigned int i)
Convert an unsigned integer to a half in round-down mode.
Convert the unsigned integer value i to a half-precision floating-point value in round-down mode.
being.
Parameters
i(in): unsigned int. Is only
Returns
- half
- i converted to half.
__uint2half_rn
__half __uint2half_rn(const unsigned int i)
Convert an unsigned integer to a half in round-to-nearest-even mode.
Convert the unsigned integer value i to a half-precision floating-point value in round-to-nearest-even mode.
int
Parameters
i(in): unsigned
Returns
- half
- i converted to half.
__uint2half_ru
__half __uint2half_ru(const unsigned int i)
Convert an unsigned integer to a half in round-up mode.
Convert the unsigned integer value i to a half-precision floating-point value in round-up mode.
being.
Parameters
i(in): unsigned int. Is only
Returns
- half
- i converted to half.
__uint2half_rz
__half __uint2half_rz(const unsigned int i)
Convert an unsigned integer to a half in round-towards-zero mode.
Convert the unsigned integer value i to a half-precision floating-point value in round-towards-zero mode.
Is only.
Parameters
i(in): unsigned int.
Returns
- half
- i converted to half.
__ull2half_rd
__half __ull2half_rd(const unsigned long long int i)
Convert an unsigned 64-bit integer to a half in round-down mode.
Convert the unsigned 64-bit integer value i to a half-precision floating-point value in round-down mode.
int
Parameters
i(in): unsigned long long
Returns
- half
- i converted to half.
__ull2half_rn
__half __ull2half_rn(const unsigned long long int i)
Convert an unsigned 64-bit integer to a half in round-to-nearest-even mode.
Convert the unsigned 64-bit integer value i to a half-precision floating-point value in round-to-nearest-even mode.
long long int
Parameters
i(in): unsigned
Returns
- half
- i converted to half.
__ull2half_ru
__half __ull2half_ru(const unsigned long long int i)
Convert an unsigned 64-bit integer to a half in round-up mode.
Convert the unsigned 64-bit integer value i to a half-precision floating-point value in round-up mode.
Is only.
Parameters
i(in): unsigned long long int.
Returns
- half
- i converted to half.
__ull2half_rz
__half __ull2half_rz(const unsigned long long int i)
Convert an unsigned 64-bit integer to a half in round-towards-zero mode.
Convert the unsigned 64-bit integer value i to a half-precision floating-point value in round-towards-zero mode.
long int
Parameters
i(in): unsigned long
Returns
- half
- i converted to half.
__ushort2half_rd
__half __ushort2half_rd(const unsigned short int i)
Convert an unsigned short integer to a half in round-down mode.
Convert the unsigned short integer value i to a half-precision floating-point value in round-down mode.
Is only.
Parameters
i(in): unsigned short int.
Returns
- half
- i converted to half.
__ushort2half_rn
__half __ushort2half_rn(const unsigned short int i)
Convert an unsigned short integer to a half in round-to-nearest-even mode.
Convert the unsigned short integer value i to a half-precision floating-point value in round-to-nearest-even mode.
short int
Parameters
i(in): unsigned
Returns
- half
- i converted to half.
__ushort2half_ru
__half __ushort2half_ru(const unsigned short int i)
Convert an unsigned short integer to a half in round-up mode.
Convert the unsigned short integer value i to a half-precision floating-point value in round-up mode.
only.
Parameters
i(in): unsigned short int. Is
Returns
- half
- i converted to half.
Half Precision Conversion and Data Movement 75
__ushort2half_rz
__half __ushort2half_rz(const unsigned short int i)
Convert an unsigned short integer to a half in round-towards-zero mode.
Convert the unsigned short integer value i to a half-precision floating-point value in round-towards-zero mode.
short int
Parameters
i(in): unsigned
Returns
- half
- i converted to half.
__ushort_as_half
__half __ushort_as_half(const unsigned short int i)
Reinterprets bits in an unsigned short integer as a half.
Reinterprets the bits in the unsigned short integer i as a half-precision floating-point number.
Parameters
i(in): unsigned short int
Returns
- half
- The reinterpreted value.
Half Arithmetic Functions
atomicAdd
__half2 atomicAdd(__half2 *const address, const __half2 val)
Adds val to the value stored at address in global or shared memory, and writes this value back to address .
atomicAdd
__half atomicAdd(__half *const address, const __half val)
Adds val to the value stored at address in global or shared memory, and writes this value back to address .
Half2 Math Functions
h2ceil
__half2 h2ceil(const __half2 h)
Calculate half2 vector ceiling of the input argument.
For each component of vector h compute the smallest integer value not less than h.
Parameters
h(in): half2..
Returns
- half2
- The vector of smallest integers not less than
h.
h2cos
__half2 h2cos(const __half2 a)
Calculates half2 vector cosine in round-to-nearest-even mode.
Calculates half2 cosine of input vector a in round-to-nearest-even mode.
Parameters
a(in): half2..
Returns
- half2
- The elementwise cosine on vector
a.
h2exp
__half2 h2exp(const __half2 a)
Calculates half2 vector exponential function in round-to-nearest mode.
Calculates half2 exponential function of input vector a in round-to-nearest-even mode.
Parameters
a(in): half2..
Returns
- half2
- The elementwise exponential function on vector
a.
h2exp10
__half2 h2exp10(const __half2 a)
Calculates half2 vector decimal exponential function in round-to-nearest-even mode.
Calculates half2 decimal exponential function of input vector a in round-to-nearest-even mode.
Parameters
a(in): half2..
Returns
- half2
- The elementwise decimal exponential function on vector
a.
h2exp2
__half2 h2exp2(const __half2 a)
Calculates half2 vector binary exponential function in round-to-nearest-even mode.
Calculates half2 binary exponential function of input vector a in round-to-nearest-even mode.
Parameters
a(in): half2..
Returns
- half2
- The elementwise binary exponential function on vector
a.
h2floor
__half2 h2floor(const __half2 h)
Calculate the largest integer less than or equal to h.
For each component of vector h calculate the largest integer value which is less than or equal to h.
being.
Parameters
h(in): half2. Is only
Returns
- half2
- The vector of largest integers which is less than or equal to
h.
h2log
__half2 h2log(const __half2 a)
Calculates half2 vector natural logarithm in round-to-nearest-even mode.
Calculates half2 natural logarithm of input vector a in round-to-nearest-even mode.
Parameters
a(in): half2..
Returns
- half2
- The elementwise natural logarithm on vector
a.
h2log10
__half2 h2log10(const __half2 a)
Calculates half2 vector decimal logarithm in round-to-nearest-even mode.
Calculates half2 decimal logarithm of input vector a in round-to-nearest-even mode.
Parameters
a(in): half2..
Returns
- half2
- The elementwise decimal logarithm on vector
a.
h2log2
__half2 h2log2(const __half2 a)
Calculates half2 vector binary logarithm in round-to-nearest-even mode.
Calculates half2 binary logarithm of input vector a in round-to-nearest mode.
Parameters
a(in): half2..
Returns
- half2
- The elementwise binary logarithm on vector
a.
h2rcp
__half2 h2rcp(const __half2 a)
Calculates half2 vector reciprocal in round-to-nearest-even mode.
Calculates half2 reciprocal of input vector a in round-to-nearest-even mode.
Parameters
a(in): half2..
Returns
- half2
- The elementwise reciprocal on vector
a.
h2rint
__half2 h2rint(const __half2 h)
Round input to nearest integer value in half-precision floating-point number.
Round each component of half2 vector h to the nearest integer value in half-precision floating-point format, with halfway cases rounded to the nearest even integer value.
Parameters
h(in): half2..
Returns
- half2
- The vector of rounded integer values.
h2rsqrt
__half2 h2rsqrt(const __half2 a)
Calculates half2 vector reciprocal square root in round-to-nearest mode.
Calculates half2 reciprocal square root of input vector a in round-to-nearest-even mode.
Parameters
a(in): half2..
Returns
- half2
- The elementwise reciprocal square root on vector
a.
h2sin
__half2 h2sin(const __half2 a)
Calculates half2 vector sine in round-to-nearest-even mode.
Calculates half2 sine of input vector a in round-to-nearest-even mode.
Parameters
a(in): half2..
Returns
- half2
- The elementwise sine on vector
a.
h2sqrt
__half2 h2sqrt(const __half2 a)
Calculates half2 vector square root in round-to-nearest-even mode.
Calculates half2 square root of input vector a in round-to-nearest mode.
Parameters
a(in): half2..
Returns
- half2
- The elementwise square root on vector
a.
h2trunc
__half2 h2trunc(const __half2 h)
Truncate half2 vector input argument to the integral part.
Round each component of vector h to the nearest integer value that does not exceed h in magnitude.
read.
Parameters
h(in): half2. Is only being
Returns
- half2
- The truncated
h.
Half Math Functions
hceil
__half hceil(const __half h)
Calculate ceiling of the input argument.
Compute the smallest integer value not less than h.
Parameters
h(in): half
Returns
- half
- The smallest integer value not less than
h.
hcos
__half hcos(const __half a)
Calculates half cosine in round-to-nearest-even mode.
Calculates half cosine of input a in round-to-nearest-even mode.
Parameters
a(in): half
Returns
- half
- The cosine of
a.
hexp
__half hexp(const __half a)
Calculates half natural exponential function in round-to-nearest mode.
Calculates half natural exponential function of input a in round-to-nearest-even mode.
Parameters
a(in): half
Returns
- half
- The natural exponential function on
a.
hexp10
__half hexp10(const __half a)
Calculates half decimal exponential function in round-to-nearest mode.
Calculates half decimal exponential function of input a in round-to-nearest-even mode.
Parameters
a(in): half
Returns
- half
- The decimal exponential function on
a.
hexp2
__half hexp2(const __half a)
Calculates half binary exponential function in round-to-nearest mode.
Calculates half binary exponential function of input a in round-to-nearest-even mode.
Parameters
a(in): half
Returns
- half
- The binary exponential function on
a.
hfloor
__half hfloor(const __half h)
Calculate the largest integer less than or equal to h.
Calculate the largest integer value which is less than or equal to h.
Parameters
h(in): half
Returns
- half
- The largest integer value which is less than or equal to
h.
hlog
__half hlog(const __half a)
Calculates half natural logarithm in round-to-nearest-even mode.
Calculates half natural logarithm of input a in round-to-nearest-even mode.
Parameters
a(in): half
Returns
- half
- The natural logarithm of
a.
hlog10
__half hlog10(const __half a)
Calculates half decimal logarithm in round-to-nearest-even mode.
Calculates half decimal logarithm of input a in round-to-nearest-even mode.
Parameters
a(in): half
Returns
- half
- The decimal logarithm of
a.
hlog2
__half hlog2(const __half a)
Calculates half binary logarithm in round-to-nearest-even mode.
Calculates half binary logarithm of input a in round-to-nearest-even mode.
Parameters
a(in): half
Returns
- half
- The binary logarithm of
a.
hrcp
__half hrcp(const __half a)
Calculates half reciprocal in round-to-nearest-even mode.
Calculates half reciprocal of input a in round-to-nearest-even mode.
Parameters
a(in): half
Returns
- half
- The reciprocal of
a.
hrint
__half hrint(const __half h)
Round input to nearest integer value in half-precision floating-point number.
Round h to the nearest integer value in half-precision floating-point format, with halfway cases rounded to the nearest even integer value.
Parameters
h(in): half
Returns
- half
- The nearest integer to
h.
hrsqrt
__half hrsqrt(const __half a)
Calculates half reciprocal square root in round-to-nearest-even mode.
Calculates half reciprocal square root of input a in round-to-nearest mode.
Parameters
a(in): half
Returns
- half
- The reciprocal square root of
a.
hsin
__half hsin(const __half a)
Calculates half sine in round-to-nearest-even mode.
Calculates half sine of input a in round-to-nearest-even mode.
Parameters
a(in): half
Returns
- half
- The sine of
a.
hsqrt
__half hsqrt(const __half a)
Calculates half square root in round-to-nearest-even mode.
Calculates half square root of input a in round-to-nearest-even mode.
Parameters
a(in): half
Returns
- half
- The square root of
a.
htanh
__half htanh(const __half a)
Calculates half hyperbolic tangent function in round-to-nearest-even mode.
Calculates half hyperbolic tangent function: _math_formula (a)_musa_math_formula in round-to-nearest-even mode.
Parameters
a(in): half
Returns
- half
- The hyperbolic tangent function of
a.
htanh_approx
__half htanh_approx(const __half a)
Calculates approximate half hyperbolic tangent function.
Calculates approximate half hyperbolic tangent function.
Parameters
a(in): half
Returns
- half
- The approximate hyperbolic tangent function of
a.
htrunc
__half htrunc(const __half h)
Truncate input argument to the integral part.
Round h to the nearest integer value that does not exceed h in magnitude.
Parameters
h(in): half
Returns
- half
- The truncated integer value.
Half Precision Conversion and Data Movement 75
make_half2
__half2 make_half2(const __half x, const __half y)
Vector function, combines two __half numbers into one __half2 number.
Bfloat16 Precision Intrinsics
Use bfloat16 intrinsics when the kernel needs the wider exponent range of bfloat16 with reduced mantissa precision. Scalar __mt_bfloat16 APIs operate on one value; __mt_bfloat162 APIs operate on two packed bfloat16 values.
Bfloat16 Precision Conversion and Data
__bfloat1622float2
float2 __bfloat1622float2(const __mt_bfloat162 a)
Converts both halves of mt_bfloat162 to float2 and returns the result.
Converts both halves of mt_bfloat162 input a to float2 and returns the result.
Parameters
a(in): mt_bfloat162
Returns
- float2
- a converted to float2.
__bfloat162bfloat162
__mt_bfloat162 __bfloat162bfloat162(const __mt_bfloat16 a)
Returns mt_bfloat162 with both halves equal to the input value.
Returns mt_bfloat162 number with both halves equal to the input a mt_bfloat16 number.
Parameters
a(in): mt_bfloat16
Returns
- mt_bfloat162
- The vector which has both its halves equal to the input
a.
__bfloat162char_rz
signed char __bfloat162char_rz(const __mt_bfloat16 h)
Convert a mt_bfloat16 to a signed char in round-towards-zero mode.
Convert the mt_bfloat16 floating-point value h to a signed char in round-towards-zero mode. NaN inputs are converted to 0.
Parameters
h(in): mt_bfloat16
Returns
- signed char
- h converted to an signed char.
__bfloat162float
float __bfloat162float(const __mt_bfloat16 a)
Converts mt_bfloat16 number to float.
Converts mt_bfloat16 number a to float.
Parameters
a(in): float
Returns
- float
- a converted to float.
__bfloat162int_rd
int __bfloat162int_rd(const __mt_bfloat16 h)
Convert a mt_bfloat16 to a signed integer in round-down mode.
Convert the mt_bfloat16 floating-point value h to a signed integer in round-down mode.
Parameters
h(in): mt_bfloat16
Returns
- int
- h converted to a signed integer.
__bfloat162int_rn
int __bfloat162int_rn(const __mt_bfloat16 h)
Convert a mt_bfloat16 to a signed integer in round-to-nearest-even mode.
Convert the mt_bfloat16 floating-point value h to a signed integer in round-to-nearest-even mode.
being.
Parameters
h(in): mt_bfloat16. Is only
Returns
- int
- h converted to a signed integer.
__bfloat162int_ru
int __bfloat162int_ru(const __mt_bfloat16 h)
Convert a mt_bfloat16 to a signed integer in round-up mode.
Convert the mt_bfloat16 floating-point value h to a signed integer in round-up mode.
Parameters
h(in): mt_bfloat16
Returns
- int
- h converted to a signed integer.
__bfloat162int_rz
int __bfloat162int_rz(const __mt_bfloat16 h)
Convert a mt_bfloat16 to a signed integer in round-towards-zero mode.
Convert the mt_bfloat16 floating-point value h to a signed integer in round-towards-zero mode.
read.
Parameters
h(in): mt_bfloat16. Is only being
Returns
- int
- h converted to a signed integer.
__bfloat162ll_rd
long long int __bfloat162ll_rd(const __mt_bfloat16 h)
Convert a mt_bfloat16 to a signed 64-bit integer in round-down mode.
Convert the mt_bfloat16 floating-point value h to a signed 64-bit integer in round-down mode.
Parameters
h(in): mt_bfloat16
Returns
- long long int
- h converted to a signed 64-bit integer.
__bfloat162ll_rn
long long int __bfloat162ll_rn(const __mt_bfloat16 h)
Convert a mt_bfloat16 to a signed 64-bit integer in round-to-nearest-even mode.
Convert the mt_bfloat16 floating-point value h to a signed 64-bit integer in round-to-nearest-even mode.
Parameters
h(in): mt_bfloat16
Returns
- long long int
- h converted to a signed 64-bit integer.
__bfloat162ll_ru
long long int __bfloat162ll_ru(const __mt_bfloat16 h)
Convert a mt_bfloat16 to a signed 64-bit integer in round-up mode.
Convert the mt_bfloat16 floating-point value h to a signed 64-bit integer in round-up mode.
Parameters
h(in): mt_bfloat16
Returns
- long long int
- h converted to a signed 64-bit integer.
__bfloat162ll_rz
long long int __bfloat162ll_rz(const __mt_bfloat16 h)
Convert a mt_bfloat16 to a signed 64-bit integer in round-towards-zero mode.
Convert the mt_bfloat16 floating-point value h to a signed 64-bit integer in round-towards-zero mode.
Parameters
h(in): mt_bfloat16
Returns
- long long int
- h converted to a signed 64-bit integer.
__bfloat162short_rd
short int __bfloat162short_rd(const __mt_bfloat16 h)
Convert a mt_bfloat16 to a signed short integer in round-down mode.
Convert the mt_bfloat16 floating-point value h to a signed short integer in round-down mode.
Parameters
h(in): mt_bfloat16
Returns
- short int
- h converted to a signed short integer.
__bfloat162short_rn
short int __bfloat162short_rn(const __mt_bfloat16 h)
Convert a mt_bfloat16 to a signed short integer in round-to-nearest-even mode.
Convert the mt_bfloat16 floating-point value h to a signed short integer in round-to-nearest-even mode.
Parameters
h(in): mt_bfloat16
Returns
- short int
- h converted to a signed short integer.
__bfloat162short_ru
short int __bfloat162short_ru(const __mt_bfloat16 h)
Convert a mt_bfloat16 to a signed short integer in round-up mode.
Convert the mt_bfloat16 floating-point value h to a signed short integer in round-up mode.
Parameters
h(in): mt_bfloat16
Returns
- short int
- h converted to a signed short integer.
__bfloat162short_rz
short int __bfloat162short_rz(const __mt_bfloat16 h)
Convert a mt_bfloat16 to a signed short integer in round-towards-zero mode.
Convert the mt_bfloat16 floating-point value h to a signed short integer in round-towards-zero mode.
Parameters
h(in): mt_bfloat16
Returns
- short int
- h converted to a signed short integer.
__bfloat162uchar_rz
unsigned char __bfloat162uchar_rz(const __mt_bfloat16 h)
Convert a mt_bfloat16 to an unsigned char in round-towards-zero mode.
__bfloat162uint_rd
unsigned int __bfloat162uint_rd(const __mt_bfloat16 h)
Convert a mt_bfloat16 to an unsigned integer in round-down mode.
Convert the mt_bfloat16 floating-point value h to an unsigned integer in round-down mode.
Parameters
h(in): mt_bfloat16
Returns
- unsigned int
- h converted to an unsigned integer.
__bfloat162uint_rn
unsigned int __bfloat162uint_rn(const __mt_bfloat16 h)
Convert a mt_bfloat16 to an unsigned integer in round-to-nearest-even mode.
Convert the mt_bfloat16 floating-point value h to an unsigned integer in round-to-nearest-even mode.
being.
Parameters
h(in): mt_bfloat16. Is only
Returns
- unsigned int
- h converted to an unsigned integer.
__bfloat162uint_ru
unsigned int __bfloat162uint_ru(const __mt_bfloat16 h)
Convert a mt_bfloat16 to an unsigned integer in round-up mode.
Convert the mt_bfloat16 floating-point value h to an unsigned integer in round-up mode.
Parameters
h(in): mt_bfloat16
Returns
- unsigned int
- h converted to an unsigned integer.
__bfloat162uint_rz
unsigned int __bfloat162uint_rz(const __mt_bfloat16 h)
Convert a mt_bfloat16 to an unsigned integer in round-towards-zero mode.
Convert the mt_bfloat16 floating-point value h to an unsigned integer in round-towards-zero mode.
read.
Parameters
h(in): mt_bfloat16. Is only being
Returns
- unsigned int
- h converted to an unsigned integer.
__bfloat162ull_rd
unsigned long long int __bfloat162ull_rd(const __mt_bfloat16 h)
Convert a mt_bfloat16 to an unsigned 64-bit integer in round-down mode.
Convert the mt_bfloat16 floating-point value h to an unsigned 64-bit integer in round-down mode.
read.
Parameters
h(in): mt_bfloat16. Is only being
Returns
- unsigned long long int
- h converted to an unsigned 64-bit integer.
__bfloat162ull_rn
unsigned long long int __bfloat162ull_rn(const __mt_bfloat16 h)
Convert a mt_bfloat16 to an unsigned 64-bit integer in round-to-nearest-even mode.
Convert the mt_bfloat16 floating-point value h to an unsigned 64-bit integer in round-to-nearest-even mode.
only.
Parameters
h(in): mt_bfloat16. Is
Returns
- unsigned long long int
- h converted to an unsigned 64-bit integer.
__bfloat162ull_ru
unsigned long long int __bfloat162ull_ru(const __mt_bfloat16 h)
Convert a mt_bfloat16 to an unsigned 64-bit integer in round-up mode.
Convert the mt_bfloat16 floating-point value h to an unsigned 64-bit integer in round-up mode.
read.
Parameters
h(in): mt_bfloat16. Is only being
Returns
- unsigned long long int
- h converted to an unsigned 64-bit integer.
__bfloat162ull_rz
unsigned long long int __bfloat162ull_rz(const __mt_bfloat16 h)
Convert a mt_bfloat16 to an unsigned 64-bit integer in round-towards-zero mode.
Convert the mt_bfloat16 floating-point value h to an unsigned 64-bit integer in round-towards-zero mode.
only.
Parameters
h(in): mt_bfloat16. Is
Returns
- unsigned long long int
- h converted to an unsigned 64-bit integer.
__bfloat162ushort_rd
unsigned short int __bfloat162ushort_rd(const __mt_bfloat16 h)
Convert a mt_bfloat16 to an unsigned short integer in round-down mode.
Convert the mt_bfloat16 floating-point value h to an unsigned short integer in round-down mode.
read.
Parameters
h(in): mt_bfloat16. Is only being
Returns
- unsigned short int
- h converted to an unsigned short integer.
__bfloat162ushort_rn
unsigned short int __bfloat162ushort_rn(const __mt_bfloat16 h)
Convert a mt_bfloat16 to an unsigned short integer in round-to-nearest-even mode.
Convert the mt_bfloat16 floating-point value h to an unsigned short integer in round-to-nearest-even mode.
only.
Parameters
h(in): mt_bfloat16. Is
Returns
- unsigned short int
- h converted to an unsigned short integer.
__bfloat162ushort_ru
unsigned short int __bfloat162ushort_ru(const __mt_bfloat16 h)
Convert a mt_bfloat16 to an unsigned short integer in round-up mode.
Convert the mt_bfloat16 floating-point value h to an unsigned short integer in round-up mode.
read.
Parameters
h(in): mt_bfloat16. Is only being
Returns
- unsigned short int
- h converted to an unsigned short integer.
__bfloat162ushort_rz
unsigned short int __bfloat162ushort_rz(const __mt_bfloat16 h)
Convert a mt_bfloat16 to an unsigned short integer in round-towards-zero mode.
Convert the mt_bfloat16 floating-point value h to an unsigned short integer in round-towards-zero mode.
being.
Parameters
h(in): mt_bfloat16. Is only
Returns
- unsigned short int
- h converted to an unsigned short integer.
__bfloat16_as_short
short int __bfloat16_as_short(const __mt_bfloat16 h)
Reinterprets bits in a mt_bfloat16 as a signed short integer.
Reinterprets the bits in the mt_bfloat16 floating-point number h as a signed short integer.
Parameters
h(in): mt_bfloat16
Returns
- short int
- The reinterpreted value.
__bfloat16_as_ushort
unsigned short int __bfloat16_as_ushort(const __mt_bfloat16 h)
Reinterprets bits in a mt_bfloat16 as an unsigned short integer.
Reinterprets the bits in the mt_bfloat16 floating-point h as an unsigned short number.
Parameters
h(in): mt_bfloat16
Returns
- unsigned short int
- The reinterpreted value.
__double2bfloat16
__mt_bfloat16 __double2bfloat16(const double a)
Converts double number to mt_bfloat16 precision in round-to-nearest-even mode and returns mt_bfloat16 with converted value.
Converts double number a to mt_bfloat16 precision in round-to-nearest-even mode.
Parameters
a(in): double
Returns
- mt_bfloat16
- a converted to mt_bfloat16.
__float22bfloat162_rn
__mt_bfloat162 __float22bfloat162_rn(const float2 a)
Converts both components of float2 number to mt_bfloat16 precision in round-to-nearest-even mode and returns mt_bfloat162 with converted values.
Converts both components of float2 to mt_bfloat16 precision in round-to-nearest mode and combines the results into one mt_bfloat162 number. Low 16 bits of the return value correspond to a.x and high 16 bits of the return value correspond to a.y.
being.
converted float2 components.
Parameters
a(in): float2. Is only
Returns
- mt_bfloat162
- The
mt_bfloat162which has corresponding halves equal to the
__float2bfloat16
__mt_bfloat16 __float2bfloat16(const float a)
Converts float number to mt_bfloat16 precision in round-to-nearest-even mode and returns mt_bfloat16 with converted value.
Converts float number a to mt_bfloat16 precision in round-to-nearest-even mode.
Parameters
a(in): float
Returns
- mt_bfloat16
- a converted to mt_bfloat16.
__float2bfloat162_rn
__mt_bfloat162 __float2bfloat162_rn(const float a)
Converts input to mt_bfloat16 precision in round-to-nearest-even mode and populates both halves of mt_bfloat162 with converted value.
Converts input a to mt_bfloat16 precision in round-to-nearest-even mode and populates both halves of mt_bfloat162 with converted value.
mt_bfloat16 precision number.
Parameters
a(in): float
Returns
- mt_bfloat162
- The
mt_bfloat162value with both halves equal to the converted
__float2bfloat16_rd
__mt_bfloat16 __float2bfloat16_rd(const float a)
Converts float number to mt_bfloat16 precision in round-down mode and returns mt_bfloat16 with converted value.
Converts float number a to mt_bfloat16 precision in round-down mode.
Parameters
a(in): float
Returns
- mt_bfloat16
- a converted to mt_bfloat16.
Bfloat16 Precision Conversion and Data Movement 161
__float2bfloat16_rn
__mt_bfloat16 __float2bfloat16_rn(const float a)
Converts float number to mt_bfloat16 precision in round-to-nearest-even mode and returns mt_bfloat16 with converted value.
Converts float number a to mt_bfloat16 precision in round-to-nearest-even mode.
Parameters
a(in): float
Returns
- mt_bfloat16
- a converted to mt_bfloat16.
__float2bfloat16_ru
__mt_bfloat16 __float2bfloat16_ru(const float a)
Converts float number to mt_bfloat16 precision in round-up mode and returns mt_bfloat16 with converted value.
Converts float number a to mt_bfloat16 precision in round-up mode.
Parameters
a(in): float
Returns
- mt_bfloat16
- a converted to mt_bfloat16.
__float2bfloat16_rz
__mt_bfloat16 __float2bfloat16_rz(const float a)
Converts float number to mt_bfloat16 precision in round-towards-zero mode and returns mt_bfloat16 with converted value.
Converts float number a to mt_bfloat16 precision in round-towards-zero mode.
mt_bfloat16
Parameters
a(in): float
Returns
- a converted to mt_bfloat16.
__floats2bfloat162_rn
__mt_bfloat162 __floats2bfloat162_rn(const float a, const float b)
Converts both input floats to mt_bfloat16 precision in round-to-nearest-even mode and returns mt_bfloat162 with converted values.
Converts both input floats to mt_bfloat16 precision in round-to-nearest-even mode and combines the results into one mt_bfloat162 number. Low 16 bits of the return value correspond to the input a, high 16 bits correspond to the input b.
converted input floats.
Parameters
a(in): floatb(in): float
Returns
- mt_bfloat162
- The
mt_bfloat162value with corresponding halves equal to the
Bfloat162 Arithmetic Functions
__h2div
__mt_bfloat162 __h2div(const __mt_bfloat162 a, const __mt_bfloat162 b)
Performs mt_bfloat162 vector division in round-to-nearest-even mode.
Divides mt_bfloat162 input vector a by input vector b in round-to-nearest mode.
Parameters
a(in): mt_bfloat162b(in): mt_bfloat162
Returns
- mt_bfloat162
- The elementwise division of
awithb.