__asc_cvt_float2_to_fp8x2

Applicability

Function Usage

Converts two components of float2 data to 8-bit floating-point data of the specified type (float8_e4m3x2_t and float8_e5m2x2_t) in line with the CAST_RINT mode, and performs overflow processing based on the specified saturation behavior (saturation or non-saturation mode). The conversion result is stored as the __asc_fp8x2_storage_t type in bit-level packing format. This type is a 16-bit unsigned integer (unsigned short int) used to store data of the float8_e4m3x2_t or float8_e5m2x2_t type.

Prototype

1

__asc_fp8x2_storage_t __asc_cvt_float2_to_fp8x2(const float2 x, const __asc_saturation_t saturate, const __asc_fp8_interpretation_t fp8_interpretation)

Parameters

Returns

The __asc_fp8x2_storage_t data converted from the two input components based on the specified 8-bit floating-point number type and saturation behavior in line with the CAST_RINT mode.

• The float8_e4m3_t data type does not have the inf value. In non-saturation mode, if an input is greater than the range that can be represented by this type, the overflow result is nan. In saturation mode, the overflow result is the maximum or minimum value that can be represented by this type.
• The float8_e5m2_t data type has the inf value. In non-saturation mode, if an input is greater than the range that can be represented by this type, the conversion result is the inf value of the corresponding sign. In saturation mode, the overflow result is the maximum or minimum value that can be represented by this type.
• In saturation mode, for the float8_e5m2_t and float8_e4m3_t data types, if bit[50] of the CTRL register is 0, the nan value is converted to 0; if bit[50] of the CTRL register is 1, the nan value is converted to the nan value of the fp8 type. The value of the CTRL register can be obtained by referring to GetCtrlSpr(ISASI).

Restrictions

Before using this API, set bit[60] of the CTRL register to 0. Otherwise, the saturation behavior does not take effect. For details, see Methods of Controlling Saturation Behaviors.

For SIMT programming, this API is not supported.

Header File to Be Included

To use this API, the simt_api/asc_fp8.h header file must be included.

1

#include "simt_api/asc_fp8.h"

Examples

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14

// Use small vectors to improve data movement efficiency.
__simt_vf__ __launch_bounds__(1024) inline void simt_asc_cvt_float2_to_fp8x2(__gm__ float2* input, __gm__ uint16_t* output, uint32_t input_total_length)
{
    uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
    // Each thread processes one piece of float2 data, that is, two pieces of float data. Therefore, threads with idx greater than or equal to input_total_length/2 do not process data.
    if (idx >= input_total_length / 2) {
        return;
    }
    output[idx] = __asc_cvt_float2_to_fp8x2(input[idx], __ASC_NOSAT, __ASC_E4M3);
}
__global__ __vector__ void cast_kernel(__gm__ float* input, __gm__ uint16_t* output, uint32_t input_total_length)
{
    asc_vf_call<simt_asc_cvt_float2_to_fp8x2>(dim3(1024), (__gm__ float2*)input, output, input_total_length);
}