__half22hif82_rh
Applicability
Product |
Supported |
|---|---|
Atlas 350 Accelerator Card |
√ |
x |
|
x |
|
x |
|
x |
|
x |
|
x |
Function Usage
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(const half2 x) |
Parameters
Parameter |
Input/Output |
Description |
|---|---|---|
x |
Input |
Source operand. |
Returns
The hifloat8x2_t data converted from the two input components in line with the CAST_HYBRID mode.
Restrictions
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
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_half22hif82_rh(__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(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>(dim3(1024), (__gm__ half2*)input, (__gm__ hifloat8x2_t*)output, input_total_length); } |
Parent topic: Type Conversion