__hif822half2
Applicability
Product |
Supported |
|---|---|
Atlas 350 Accelerator Card |
√ |
x |
|
x |
|
x |
|
x |
|
x |
|
x |
Function Usage
Converts the two components of hifloat8x2_t data into half-precision floating-point data, with the result returned as the half2 type.
Prototype
1 | inline half2 __hif822half2(const hifloat8x2_t x) |
Parameters
Parameter |
Input/Output |
Description |
|---|---|---|
x |
Input |
Source operand. |
Returns
half2 data converted from the hifloat8x2_t input
Restrictions
For SIMT programming, this API is not supported.
Header Files 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
For SIMD and SIMT programming:
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_hif822half2(__gm__ hifloat8x2_t* input, __gm__ half2* output, uint32_t input_total_length) { uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x; // Each thread processes one piece of hifloat8x2_t data, that is, two pieces of hifloat8_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] = __hif822half2(input[idx]); } __global__ __vector__ void cast_kernel(__gm__ uint8_t* input, __gm__ half* output, uint32_t input_total_length) { asc_vf_call<simt_hif822half2>(dim3(1024), (__gm__ hifloat8x2_t*)input, (__gm__ half2*)output, input_total_length); } |
Parent topic: Type Conversion