__floats2bfloat162_rn

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 input data x and y to the bfloat16 type in line with the CAST_RINT mode, fills both the front and back parts of bfloat16x2, and returns the converted bfloat16x2 data.

Prototype

1
inline bfloat16x2_t __floats2bfloat162_rn(const float x, const float y)

Parameters

Table 1 Parameters

Parameter

Input/Output

Description

x

Input

Source operand.

y

Input

Source operand.

Returns

Converted data returned after this API converts the input float data to the bfloat16 type in line with the CAST_RINT mode and fills both the front and back parts of bfloat16x2

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.

1
#include "simt_api/asc_bf16.h"

Examples

For SIMD and SIMT programming:
 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
// Use small vectors to improve data movement efficiency.
__simt_vf__ __launch_bounds__(1024) inline void simt_floats2bfloat162_rn(__gm__ float* input1, __gm__ float* input2, __gm__ bfloat16x2_t* output, uint32_t input_total_length)
{
    uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx > input_total_length) {
        return;
    }
    output[idx] = __floats2bfloat162_rn(input1[idx], input2[idx]);
}
__global__ __vector__ void cast_kernel(__gm__ float* input1,  __gm__ float* input2, __gm__ bfloat16_t* output, uint32_t input_total_length)
{
    asc_vf_call<simt_floats2bfloat162_rn>(dim3(1024), input1, input2, (__gm__ bfloat16x2_t*)output, input_total_length);
}