__float22half2_rz_sat

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

In saturation mode, converts the two components of float2 data to half-precision floating-point data in line with the CAST_TRUNC mode and returns the converted half2 data.

Prototype

inline half2 __float22half2_rz_sat(const float2 x)

Parameters

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

Returns

The half2 data converted from the two input components in saturation mode in line with the CAST_TRUNC mode.

Restrictions

Before using this API, set bit[60] of the CTRL register to 0. Otherwise, the saturation mode 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_fp16.h header file must be included.

#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_float22half2_rz_sat(__gm__ float2* input, __gm__ half2* 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] = __float22half2_rz_sat(input[idx]);
}

__global__ __vector__ void cast_kernel(__gm__ float* input, __gm__ half* output, uint32_t input_total_length)
{
    asc_vf_call<simt_float22half2_rz_sat>(dim3(1024), (__gm__ float2*)input, (__gm__ half2*)output, input_total_length);
}

Parent topic: Type Conversion