__low2float

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 low 16 bits of the input data into floating-point data and returns the result.

Prototype

inline float __low2float(const half2 x)

__simt_callee__ inline float __low2float(const bfloat16x2_t x)

Parameters

**Table 1** Parameters
Parameter	Input/Output	Description
x	Input	Source operand.

Returns

Floating-point data converted from the low 16 bits of the input data

Restrictions

For SIMT programming, this API is not supported.

Header Files to Be Included

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

#include "simt_api/asc_bf16.h"

#include "simt_api/asc_fp16.h"

Examples

For SIMD and SIMT programming:

// Use small vectors to improve data movement efficiency.
__simt_vf__ __launch_bounds__(1024) inline void simt_low2float(__gm__ bfloat16x2_t* input, __gm__ float* output, uint32_t input_total_length)
{
    uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
    // Each thread processes one piece of bfloat16x2_t data, that is, two pieces of bfloat16_t 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] = __low2float(input[idx]);
}
__global__ __vector__ void cast_kernel(__gm__ bfloat16_t* input, __gm__ float* output, uint32_t input_total_length)
{
    asc_vf_call<simt_low2float>(dim3(1024),  (__gm__ bfloat16x2_t*)input, output, input_total_length);
}

Parent topic: Type Conversion