3#if defined(__cplusplus)
16#if defined(__cplusplus)
19 __half __float2half(
float);
20 float __half2float(__half);
30 #if !defined(__HIP_NO_HALF_CONVERSIONS__)
31 __half(
float x) : __x{__float2half(x).__x} {}
32 __half(
double x) : __x{__float2half(x).__x} {}
34 __half(
const __half&) =
default;
35 __half(__half&&) =
default;
39 __half& operator=(
const __half&) =
default;
40 __half& operator=(__half&&) =
default;
41 __half& operator=(
const __half_raw& x) { __x = x.x;
return *
this; }
42 #if !defined(__HIP_NO_HALF_CONVERSIONS__)
43 __half& operator=(
float x)
45 __x = __float2half(x).__x;
48 __half& operator=(
double x)
50 return *
this =
static_cast<float>(x);
55 operator float()
const {
return __half2float(*
this); }
70 x{reinterpret_cast<const __half&>(ix.x)},
71 y{reinterpret_cast<const __half&>(ix.y)}
73 __half2(
const __half& ix,
const __half& iy) : x{ix}, y{iy} {}
74 __half2(
const __half2&) =
default;
75 __half2(__half2&&) =
default;
79 __half2& operator=(
const __half2&) =
default;
80 __half2& operator=(__half2&&) =
default;
92 reinterpret_cast<const unsigned short&
>(x),
93 reinterpret_cast<const unsigned short&
>(y)};
99 unsigned short __internal_float2half(
100 float flt,
unsigned int& sgn,
unsigned int& rem)
103 std::memcpy(&x, &flt,
sizeof(flt));
105 unsigned int u = (x & 0x7fffffffU);
106 sgn = ((x >> 16) & 0x8000U);
109 if (u >= 0x7f800000U) {
111 return static_cast<unsigned short>(
112 (u == 0x7f800000U) ? (sgn | 0x7c00U) : 0x7fffU);
115 if (u > 0x477fefffU) {
117 return static_cast<unsigned short>(sgn | 0x7bffU);
120 if (u >= 0x38800000U) {
123 return static_cast<unsigned short>(sgn | (u >> 13));
126 if (u < 0x33000001U) {
128 return static_cast<unsigned short>(sgn);
131 unsigned int exponent = u >> 23;
132 unsigned int mantissa = (u & 0x7fffffU);
133 unsigned int shift = 0x7eU - exponent;
134 mantissa |= 0x800000U;
135 rem = mantissa << (32 - shift);
136 return static_cast<unsigned short>(sgn | (mantissa >> shift));
140 __half __float2half(
float x)
145 r.x = __internal_float2half(x, sgn, rem);
146 if (rem > 0x80000000U || (rem == 0x80000000U && (r.x & 0x1))) ++r.x;
152 __half __float2half_rn(
float x) {
return __float2half(x); }
155 __half __float2half_rz(
float x)
160 r.x = __internal_float2half(x, sgn, rem);
166 __half __float2half_rd(
float x)
171 r.x = __internal_float2half(x, sgn, rem);
172 if (rem && sgn) ++r.x;
178 __half __float2half_ru(
float x)
183 r.x = __internal_float2half(x, sgn, rem);
184 if (rem && !sgn) ++r.x;
190 __half2 __float2half2_rn(
float x)
192 return __half2{__float2half_rn(x), __float2half_rn(x)};
196 __half2 __floats2half2_rn(
float x,
float y)
198 return __half2{__float2half_rn(x), __float2half_rn(y)};
202 float __internal_half2float(
unsigned short x)
204 unsigned int sign = ((x >> 15) & 1);
205 unsigned int exponent = ((x >> 10) & 0x1f);
206 unsigned int mantissa = ((x & 0x3ff) << 13);
208 if (exponent == 0x1fU) {
209 mantissa = (mantissa ? (sign = 0, 0x7fffffU) : 0);
211 }
else if (!exponent) {
216 msb = (mantissa & 0x400000U);
220 mantissa &= 0x7fffffU;
225 unsigned int u = ((sign << 31) | (exponent << 23) | mantissa);
227 memcpy(&f, &u,
sizeof(u));
233 float __half2float(__half x)
235 return __internal_half2float(
static_cast<__half_raw>(x).x);
241 return __internal_half2float(
static_cast<__half2_raw>(x).x);
247 return __internal_half2float(
static_cast<__half2_raw>(x).y);
250 #if !defined(HIP_NO_HALF)
252 using half2 = __half2;
__HOST_DEVICE__ float __high2float(const __hip_bfloat162 a)
Converts high 16 bits of __hip_bfloat162 to float and returns the result.
Definition amd_hip_bf16.h:260
__HOST_DEVICE__ float __low2float(const __hip_bfloat162 a)
Converts low 16 bits of __hip_bfloat162 to float and returns the result.
Definition amd_hip_bf16.h:289
Definition hip_fp16_gcc.h:7
Definition hip_fp16_gcc.h:11