__half22hif82_rh_sat
Applicability
Product |
Supported |
|---|---|
Atlas 350 Accelerator Card |
√ |
x |
|
x |
|
x |
|
x |
|
x |
|
x |
Function Usage
In saturation mode, converts the two components of half2 data to hifloat8_t precision in line with the CAST_HYBRID mode and returns the converted hifloat8x2_t data.
Prototype
1 | inline hifloat8x2_t __half22hif82_rh_sat(const half2 x) |
Parameters
Parameter |
Input/Output |
Description |
|---|---|---|
x |
Input |
Source operand. |
Returns
The hifloat8x2_t data converted from the two input components in saturation mode in line with the CAST_HYBRID 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_fp8.h header file must be included.
1 | #include "simt_api/asc_fp8.h" |
Examples
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_half22hif82_rh_sat(__gm__ half2* 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 half2 data, that is, two pieces of half 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] = __half22hif82_rh_sat(input[idx]); } __global__ __vector__ void cast_kernel(__gm__ half* input, __gm__ uint8_t* output, uint32_t input_total_length) { asc_vf_call<simt_half22hif82_rh_sat>(dim3(1024), (__gm__ half2*)input, (__gm__ hifloat8x2_t*)output, input_total_length); } |