__highs2bfloat162

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

Extracts the high 16 bits from each of the two bfloat162 inputs, fills them into bfloat162, and returns the result.

Prototype

inline bfloat16x2_t __highs2bfloat162(const bfloat16x2_t x, const bfloat16x2_t y)

Parameters

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

Returns

Result returned after the high 16 bits are extracted from the two bfloat162 inputs and filled into bfloat162

Restrictions

For SIMT programming, this API is not supported.

Header Files to Be Included

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

#include "simt_api/asc_bf16.h"

Examples

For SIMD and SIMT programming:

// Use small vectors to improve data movement efficiency.
__simt_vf__ __launch_bounds__(1024) inline void simt_highs2bfloat162(__gm__ bfloat16x2_t* input1, __gm__ bfloat16x2_t* input2, __gm__ bfloat16x2_t* 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] = __highs2bfloat162(input1[idx], input2[idx]);
}
__global__ __vector__ void cast_kernel(__gm__ bfloat16_t* input1,  __gm__ bfloat16_t* input2, __gm__ bfloat16_t* output, uint32_t input_total_length)
{
    asc_vf_call<simt_highs2bfloat162>(dim3(1024), (__gm__ bfloat16x2_t*)input1, (__gm__ bfloat16x2_t*)input2, (__gm__ bfloat16x2_t*)output, input_total_length);
}

Parent topic: Type Conversion