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