__float2bfloat162_rn
Applicability
Product |
Supported |
|---|---|
Atlas 350 Accelerator Card |
√ |
x |
|
x |
|
x |
|
x |
|
x |
|
x |
Function Usage
Converts float data to the bfloat16 type in line with the CAST_RINT mode, fills both the front and back parts of bfloat16x2, and returns the filled bfloat16x2 data.
Prototype
1 | inline bfloat16x2_t __float2bfloat162_rn(const float x) |
Parameters
Parameter |
Input/Output |
Description |
|---|---|---|
x |
Input |
Source operand. |
Returns
Filled data returned after this API converts the input 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
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_float2bfloat162_rn(__gm__ float* input, __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] = __float2bfloat162_rn(input[idx]); } __global__ __vector__ void cast_kernel(__gm__ float* input, __gm__ bfloat16_t* output, uint32_t input_total_length) { asc_vf_call<simt_float2bfloat162_rn>(dim3(1024), input, (__gm__ bfloat16x2_t*)output, input_total_length); } |