__float22hif82_rna

Applicability

Product

Supported

Atlas 350 Accelerator Card

Atlas A3 training product/Atlas A3 inference product

x

Atlas A2 training product/Atlas A2 inference product

x

Atlas 200I/500 A2 inference product

x

Atlas inference product AI Core

x

Atlas inference product Vector Core

x

Atlas training product

x

Function Usage

Converts the two components of float2 data to hifloat8_t precision in line with the CAST_ROUND mode and returns the converted hifloat8x2_t data.

Prototype

1
inline hifloat8x2_t __float22hif82_rna(const float2 x)

Parameters

Table 1 Parameters

Parameter

Input/Output

Description

x

Input

Source operand.

Returns

The hifloat8x2_t data converted from the two input components in line with the CAST_ROUND mode.

Restrictions

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

For SIMD and SIMT programming:
 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_float22hif82_rna(__gm__ float2* input, __gm__ hifloat8x2_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] = __float22hif82_rna(input[idx]);
}
__global__ __vector__ void cast_kernel(__gm__ float* input, __gm__ uint8_t* output, uint32_t input_total_length)
{
    asc_vf_call<simt_float22hif82_rna>(dim3(1024), (__gm__ float2*)input, (__gm__ hifloat8x2_t*)output, input_total_length);
}