24 #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_FP16_H 25 #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_FP16_H 27 #include <hip/hcc_detail/hip_common.h> 31 #if defined(__cplusplus) 33 #include <type_traits> 37 #if __HCC_OR_HIP_CLANG__ 38 typedef _Float16 _Float16_2 __attribute__((ext_vector_type(2)));
42 static_assert(
sizeof(_Float16) ==
sizeof(
unsigned short),
"");
51 static_assert(
sizeof(_Float16_2) ==
sizeof(
unsigned short[2]),
"");
61 #if defined(__cplusplus) 62 #include "hip_fp16_math_fwd.h" 63 #include "hip_vector_types.h" 68 template<>
struct is_floating_point<_Float16> : std::true_type {};
71 template<
bool cond,
typename T =
void>
72 using Enable_if_t =
typename std::enable_if<cond, T>::type;
78 static_assert(
sizeof(_Float16) ==
sizeof(
unsigned short),
"");
89 #if !defined(__HIP_NO_HALF_CONVERSIONS__) 91 __half(decltype(data) x) : data{x} {}
94 Enable_if_t<std::is_floating_point<T>{}>* =
nullptr>
96 __half(T x) : data{
static_cast<_Float16
>(x)} {}
99 __half(
const __half&) =
default;
101 __half(__half&&) =
default;
106 #if !defined(__HIP_NO_HALF_CONVERSIONS__) 108 typename T, Enable_if_t<std::is_integral<T>{}>* =
nullptr>
110 __half(T x) : data{
static_cast<_Float16
>(x)} {}
115 __half& operator=(
const __half&) =
default;
117 __half& operator=(__half&&) =
default;
125 volatile __half& operator=(
const __half_raw& x)
volatile 130 volatile __half& operator=(
const volatile __half_raw& x)
volatile 140 volatile __half& operator=(
__half_raw&& x)
volatile 145 volatile __half& operator=(
volatile __half_raw&& x)
volatile 150 #if !defined(__HIP_NO_HALF_CONVERSIONS__) 153 Enable_if_t<std::is_floating_point<T>{}>* =
nullptr>
155 __half& operator=(T x)
157 data =
static_cast<_Float16
>(x);
163 #if !defined(__HIP_NO_HALF_CONVERSIONS__) 165 typename T, Enable_if_t<std::is_integral<T>{}>* =
nullptr>
167 __half& operator=(T x)
169 data =
static_cast<_Float16
>(x);
174 #if !defined(__HIP_NO_HALF_OPERATORS__) 176 __half& operator+=(
const __half& x)
182 __half& operator-=(
const __half& x)
188 __half& operator*=(
const __half& x)
194 __half& operator/=(
const __half& x)
200 __half& operator++() { ++data;
return *
this; }
202 __half operator++(
int)
209 __half& operator--() { --data;
return *
this; }
211 __half operator--(
int)
220 #if !defined(__HIP_NO_HALF_CONVERSIONS__) 224 std::is_floating_point<T>{} &&
225 !std::is_same<T, double>{}>* =
nullptr>
227 operator T()
const {
return data; }
237 #if !defined(__HIP_NO_HALF_CONVERSIONS__) 239 typename T, Enable_if_t<std::is_integral<T>{}>* =
nullptr>
241 operator T()
const {
return data; }
244 #if !defined(__HIP_NO_HALF_OPERATORS__) 246 __half operator+()
const {
return *
this; }
248 __half operator-()
const 251 tmp.data = -tmp.data;
257 #if !defined(__HIP_NO_HALF_OPERATORS__) 261 __half operator+(
const __half& x,
const __half& y)
263 return __half{x} += y;
268 __half operator-(
const __half& x,
const __half& y)
270 return __half{x} -= y;
275 __half operator*(
const __half& x,
const __half& y)
277 return __half{x} *= y;
282 __half operator/(
const __half& x,
const __half& y)
284 return __half{x} /= y;
289 bool operator==(
const __half& x,
const __half& y)
291 return x.data == y.data;
296 bool operator!=(
const __half& x,
const __half& y)
303 bool operator<(
const __half& x,
const __half& y)
305 return x.data < y.data;
310 bool operator>(
const __half& x,
const __half& y)
312 return y.data < x.data;
317 bool operator<=(
const __half& x,
const __half& y)
324 bool operator>=(
const __half& x,
const __half& y)
328 #endif // !defined(__HIP_NO_HALF_OPERATORS__) 337 sizeof(_Float16_2) ==
sizeof(
unsigned short[2]),
"");
352 __half2(decltype(data) x) : data{x} {}
354 __half2(
const __half& x,
const __half& y)
358 static_cast<__half_raw>(y).data}
361 __half2(
const __half2&) =
default;
363 __half2(__half2&&) =
default;
365 ~__half2() =
default;
369 __half2& operator=(
const __half2&) =
default;
371 __half2& operator=(__half2&&) =
default;
380 #if !defined(__HIP_NO_HALF_OPERATORS__) 382 __half2& operator+=(
const __half2& x)
388 __half2& operator-=(
const __half2& x)
394 __half2& operator*=(
const __half2& x)
400 __half2& operator/=(
const __half2& x)
406 __half2& operator++() {
return *
this += _Float16_2{1, 1}; }
408 __half2 operator++(
int)
415 __half2& operator--() {
return *
this -= _Float16_2{1, 1}; }
417 __half2 operator--(
int)
427 operator decltype(data)()
const {
return data; }
432 #if !defined(__HIP_NO_HALF_OPERATORS__) 434 __half2 operator+()
const {
return *
this; }
436 __half2 operator-()
const 439 tmp.data = -tmp.data;
445 #if !defined(__HIP_NO_HALF_OPERATORS__) 449 __half2 operator+(
const __half2& x,
const __half2& y)
451 return __half2{x} += y;
456 __half2 operator-(
const __half2& x,
const __half2& y)
458 return __half2{x} -= y;
463 __half2 operator*(
const __half2& x,
const __half2& y)
465 return __half2{x} *= y;
470 __half2 operator/(
const __half2& x,
const __half2& y)
472 return __half2{x} /= y;
477 bool operator==(
const __half2& x,
const __half2& y)
479 auto r = x.data == y.data;
480 return r.x != 0 && r.y != 0;
485 bool operator!=(
const __half2& x,
const __half2& y)
492 bool operator<(
const __half2& x,
const __half2& y)
494 auto r = x.data < y.data;
495 return r.x != 0 && r.y != 0;
500 bool operator>(
const __half2& x,
const __half2& y)
507 bool operator<=(
const __half2& x,
const __half2& y)
514 bool operator>=(
const __half2& x,
const __half2& y)
518 #endif // !defined(__HIP_NO_HALF_OPERATORS__) 526 __half2 make_half2(__half x, __half y)
528 return __half2{x, y};
533 __half __low2half(__half2 x)
540 __half __high2half(__half2 x)
547 __half2 __half2half2(__half x)
549 return __half2{x, x};
554 __half2 __halves2half2(__half x, __half y)
556 return __half2{x, y};
561 __half2 __low2half2(__half2 x)
566 static_cast<__half2_raw>(x).data.x}};
571 __half2 __high2half2(__half2 x)
576 static_cast<__half2_raw>(x).data.y}};
581 __half2 __lows2half2(__half2 x, __half2 y)
586 static_cast<__half2_raw>(y).data.x}};
591 __half2 __highs2half2(__half2 x, __half2 y)
596 static_cast<__half2_raw>(y).data.y}};
601 __half2 __lowhigh2highlow(__half2 x)
606 static_cast<__half2_raw>(x).data.x}};
612 short __half_as_short(__half x)
619 unsigned short __half_as_ushort(__half x)
626 __half __short_as_half(
short x)
634 __half __ushort_as_half(
unsigned short x)
644 __half __float2half(
float x)
650 __half __float2half_rn(
float x)
656 __half __float2half_rz(
float x)
662 __half __float2half_rd(
float x)
668 __half __float2half_ru(
float x)
674 __half2 __float2half2_rn(
float x)
678 static_cast<_Float16
>(x), static_cast<_Float16>(x)}};
682 __half2 __floats2half2_rn(
float x,
float y)
685 static_cast<_Float16
>(x), static_cast<_Float16>(y)}};
689 __half2 __float22half2_rn(float2 x)
691 return __floats2half2_rn(x.x, x.y);
697 float __half2float(__half x)
703 float __low2float(__half2 x)
709 float __high2float(__half2 x)
715 float2 __half22float2(__half2 x)
718 static_cast<__half2_raw>(x).data.x,
719 static_cast<__half2_raw>(x).data.y);
725 int __half2int_rn(__half x)
731 int __half2int_rz(__half x)
737 int __half2int_rd(__half x)
743 int __half2int_ru(__half x)
751 __half __int2half_rn(
int x)
757 __half __int2half_rz(
int x)
763 __half __int2half_rd(
int x)
769 __half __int2half_ru(
int x)
777 short __half2short_rn(__half x)
783 short __half2short_rz(__half x)
789 short __half2short_rd(__half x)
795 short __half2short_ru(__half x)
803 __half __short2half_rn(
short x)
809 __half __short2half_rz(
short x)
815 __half __short2half_rd(
short x)
821 __half __short2half_ru(
short x)
829 long long __half2ll_rn(__half x)
835 long long __half2ll_rz(__half x)
841 long long __half2ll_rd(__half x)
847 long long __half2ll_ru(__half x)
855 __half __ll2half_rn(
long long x)
861 __half __ll2half_rz(
long long x)
867 __half __ll2half_rd(
long long x)
873 __half __ll2half_ru(
long long x)
881 unsigned int __half2uint_rn(__half x)
887 unsigned int __half2uint_rz(__half x)
893 unsigned int __half2uint_rd(__half x)
899 unsigned int __half2uint_ru(__half x)
907 __half __uint2half_rn(
unsigned int x)
913 __half __uint2half_rz(
unsigned int x)
919 __half __uint2half_rd(
unsigned int x)
925 __half __uint2half_ru(
unsigned int x)
933 unsigned short __half2ushort_rn(__half x)
939 unsigned short __half2ushort_rz(__half x)
945 unsigned short __half2ushort_rd(__half x)
951 unsigned short __half2ushort_ru(__half x)
959 __half __ushort2half_rn(
unsigned short x)
965 __half __ushort2half_rz(
unsigned short x)
971 __half __ushort2half_rd(
unsigned short x)
977 __half __ushort2half_ru(
unsigned short x)
985 unsigned long long __half2ull_rn(__half x)
991 unsigned long long __half2ull_rz(__half x)
997 unsigned long long __half2ull_rd(__half x)
1003 unsigned long long __half2ull_ru(__half x)
1011 __half __ull2half_rn(
unsigned long long x)
1017 __half __ull2half_rz(
unsigned long long x)
1023 __half __ull2half_rd(
unsigned long long x)
1029 __half __ull2half_ru(
unsigned long long x)
1037 __half __ldg(
const __half* ptr) {
return *ptr; }
1040 __half __ldcg(
const __half* ptr) {
return *ptr; }
1043 __half __ldca(
const __half* ptr) {
return *ptr; }
1046 __half __ldcs(
const __half* ptr) {
return *ptr; }
1050 __half2 __ldg(
const __half2* ptr) {
return *ptr; }
1053 __half2 __ldcg(
const __half2* ptr) {
return *ptr; }
1056 __half2 __ldca(
const __half2* ptr) {
return *ptr; }
1059 __half2 __ldcs(
const __half2* ptr) {
return *ptr; }
1064 bool __heq(__half x, __half y)
1067 static_cast<__half_raw>(y).data;
1071 bool __hne(__half x, __half y)
1074 static_cast<__half_raw>(y).data;
1078 bool __hle(__half x, __half y)
1081 static_cast<__half_raw>(y).data;
1085 bool __hge(__half x, __half y)
1088 static_cast<__half_raw>(y).data;
1092 bool __hlt(__half x, __half y)
1095 static_cast<__half_raw>(y).data;
1099 bool __hgt(__half x, __half y)
1102 static_cast<__half_raw>(y).data;
1106 bool __hequ(__half x, __half y) {
return __heq(x, y); }
1109 bool __hneu(__half x, __half y) {
return __hne(x, y); }
1112 bool __hleu(__half x, __half y) {
return __hle(x, y); }
1115 bool __hgeu(__half x, __half y) {
return __hge(x, y); }
1118 bool __hltu(__half x, __half y) {
return __hlt(x, y); }
1121 bool __hgtu(__half x, __half y) {
return __hgt(x, y); }
1125 __half2 __heq2(__half2 x, __half2 y)
1128 static_cast<__half2_raw>(y).data;
1129 return __builtin_convertvector(-r, _Float16_2);
1133 __half2 __hne2(__half2 x, __half2 y)
1136 static_cast<__half2_raw>(y).data;
1137 return __builtin_convertvector(-r, _Float16_2);
1141 __half2 __hle2(__half2 x, __half2 y)
1144 static_cast<__half2_raw>(y).data;
1145 return __builtin_convertvector(-r, _Float16_2);
1149 __half2 __hge2(__half2 x, __half2 y)
1152 static_cast<__half2_raw>(y).data;
1153 return __builtin_convertvector(-r, _Float16_2);
1157 __half2 __hlt2(__half2 x, __half2 y)
1160 static_cast<__half2_raw>(y).data;
1161 return __builtin_convertvector(-r, _Float16_2);
1165 __half2 __hgt2(__half2 x, __half2 y)
1168 static_cast<__half2_raw>(y).data;
1169 return __builtin_convertvector(-r, _Float16_2);
1173 __half2 __hequ2(__half2 x, __half2 y) {
return __heq2(x, y); }
1176 __half2 __hneu2(__half2 x, __half2 y) {
return __hne2(x, y); }
1179 __half2 __hleu2(__half2 x, __half2 y) {
return __hle2(x, y); }
1182 __half2 __hgeu2(__half2 x, __half2 y) {
return __hge2(x, y); }
1185 __half2 __hltu2(__half2 x, __half2 y) {
return __hlt2(x, y); }
1188 __half2 __hgtu2(__half2 x, __half2 y) {
return __hgt2(x, y); }
1192 bool __hbeq2(__half2 x, __half2 y)
1195 return r.data.x != 0 && r.data.y != 0;
1199 bool __hbne2(__half2 x, __half2 y)
1202 return r.data.x != 0 && r.data.y != 0;
1206 bool __hble2(__half2 x, __half2 y)
1209 return r.data.x != 0 && r.data.y != 0;
1213 bool __hbge2(__half2 x, __half2 y)
1216 return r.data.x != 0 && r.data.y != 0;
1220 bool __hblt2(__half2 x, __half2 y)
1223 return r.data.x != 0 && r.data.y != 0;
1227 bool __hbgt2(__half2 x, __half2 y)
1230 return r.data.x != 0 && r.data.y != 0;
1234 bool __hbequ2(__half2 x, __half2 y) {
return __hbeq2(x, y); }
1237 bool __hbneu2(__half2 x, __half2 y) {
return __hbne2(x, y); }
1240 bool __hbleu2(__half2 x, __half2 y) {
return __hble2(x, y); }
1243 bool __hbgeu2(__half2 x, __half2 y) {
return __hbge2(x, y); }
1246 bool __hbltu2(__half2 x, __half2 y) {
return __hblt2(x, y); }
1249 bool __hbgtu2(__half2 x, __half2 y) {
return __hbgt2(x, y); }
1254 __half __clamp_01(__half x)
1265 __half __hadd(__half x, __half y)
1269 static_cast<__half_raw>(y).data};
1273 __half __hsub(__half x, __half y)
1277 static_cast<__half_raw>(y).data};
1281 __half __hmul(__half x, __half y)
1285 static_cast<__half_raw>(y).data};
1289 __half __hadd_sat(__half x, __half y)
1291 return __clamp_01(__hadd(x, y));
1295 __half __hsub_sat(__half x, __half y)
1297 return __clamp_01(__hsub(x, y));
1301 __half __hmul_sat(__half x, __half y)
1303 return __clamp_01(__hmul(x, y));
1307 __half __hfma(__half x, __half y, __half z)
1310 static_cast<__half_raw>(x).data,
1311 static_cast<__half_raw>(y).data,
1312 static_cast<__half_raw>(z).data)};
1316 __half __hfma_sat(__half x, __half y, __half z)
1318 return __clamp_01(__hfma(x, y, z));
1322 __half __hdiv(__half x, __half y)
1326 static_cast<__half_raw>(y).data};
1331 __half2 __hadd2(__half2 x, __half2 y)
1335 static_cast<__half2_raw>(y).data};
1339 __half2 __hsub2(__half2 x, __half2 y)
1343 static_cast<__half2_raw>(y).data};
1347 __half2 __hmul2(__half2 x, __half2 y)
1351 static_cast<__half2_raw>(y).data};
1355 __half2 __hadd2_sat(__half2 x, __half2 y)
1364 __half2 __hsub2_sat(__half2 x, __half2 y)
1373 __half2 __hmul2_sat(__half2 x, __half2 y)
1382 __half2 __hfma2(__half2 x, __half2 y, __half2 z)
1388 __half2 __hfma2_sat(__half2 x, __half2 y, __half2 z)
1390 auto r =
static_cast<__half2_raw>(__hfma2(x, y, z));
1397 __half2 __h2div(__half2 x, __half2 y)
1401 static_cast<__half2_raw>(y).data};
1405 #if (__hcc_workweek__ >= 19015) || __HIP_CLANG_ONLY__ 1408 float amd_mixed_dot(__half2 a, __half2 b,
float c,
bool saturate) {
1409 return __ockl_fdot2(static_cast<__half2_raw>(a).data,
1410 static_cast<__half2_raw>(b).data,
1416 __half htrunc(__half x)
1419 __ocml_trunc_f16(static_cast<__half_raw>(x).data)};
1423 __half hceil(__half x)
1426 __ocml_ceil_f16(static_cast<__half_raw>(x).data)};
1430 __half hfloor(__half x)
1433 __ocml_floor_f16(static_cast<__half_raw>(x).data)};
1437 __half hrint(__half x)
1440 __ocml_rint_f16(static_cast<__half_raw>(x).data)};
1444 __half hsin(__half x)
1447 __ocml_sin_f16(static_cast<__half_raw>(x).data)};
1451 __half hcos(__half x)
1454 __ocml_cos_f16(static_cast<__half_raw>(x).data)};
1458 __half hexp(__half x)
1461 __ocml_exp_f16(static_cast<__half_raw>(x).data)};
1465 __half hexp2(__half x)
1468 __ocml_exp2_f16(static_cast<__half_raw>(x).data)};
1472 __half hexp10(__half x)
1475 __ocml_exp10_f16(static_cast<__half_raw>(x).data)};
1479 __half hlog2(__half x)
1482 __ocml_log2_f16(static_cast<__half_raw>(x).data)};
1486 __half hlog(__half x)
1489 __ocml_log_f16(static_cast<__half_raw>(x).data)};
1493 __half hlog10(__half x)
1496 __ocml_log10_f16(static_cast<__half_raw>(x).data)};
1500 __half hrcp(__half x)
1503 __llvm_amdgcn_rcp_f16(static_cast<__half_raw>(x).data)};
1507 __half hrsqrt(__half x)
1510 __ocml_rsqrt_f16(static_cast<__half_raw>(x).data)};
1514 __half hsqrt(__half x)
1517 __ocml_sqrt_f16(static_cast<__half_raw>(x).data)};
1521 bool __hisinf(__half x)
1523 return __ocml_isinf_f16(static_cast<__half_raw>(x).data);
1527 bool __hisnan(__half x)
1529 return __ocml_isnan_f16(static_cast<__half_raw>(x).data);
1533 __half __hneg(__half x)
1540 __half2 h2trunc(__half2 x)
1546 __half2 h2ceil(__half2 x)
1552 __half2 h2floor(__half2 x)
1558 __half2 h2rint(__half2 x)
1564 __half2 h2sin(__half2 x)
1570 __half2 h2cos(__half2 x)
1576 __half2 h2exp(__half2 x)
1582 __half2 h2exp2(__half2 x)
1588 __half2 h2exp10(__half2 x)
1594 __half2 h2log2(__half2 x)
1600 __half2 h2log(__half2 x) {
return __ocml_log_2f16(x); }
1603 __half2 h2log10(__half2 x) {
return __ocml_log10_2f16(x); }
1606 __half2 h2rcp(__half2 x) {
return __llvm_amdgcn_rcp_2f16(x); }
1609 __half2 h2rsqrt(__half2 x) {
return __ocml_rsqrt_2f16(x); }
1612 __half2 h2sqrt(__half2 x) {
return __ocml_sqrt_2f16(x); }
1615 __half2 __hisinf2(__half2 x)
1617 auto r = __ocml_isinf_2f16(x);
1619 static_cast<_Float16
>(r.x), static_cast<_Float16>(r.y)}};
1623 __half2 __hisnan2(__half2 x)
1625 auto r = __ocml_isnan_2f16(x);
1627 static_cast<_Float16
>(r.x), static_cast<_Float16>(r.y)}};
1631 __half2 __hneg2(__half2 x)
1637 #if !defined(HIP_NO_HALF) 1638 using half = __half;
1639 using half2 = __half2;
1641 #endif // defined(__cplusplus) 1642 #elif defined(__GNUC__) 1643 #include "hip_fp16_gcc.h" 1644 #endif // !defined(__clang__) && defined(__GNUC__) 1646 #endif // HIP_INCLUDE_HIP_HCC_DETAIL_HIP_FP16_H
Definition: hip_fp16_gcc.h:11
#define __host__
Definition: host_defines.h:41
Definition: hip_fp16_gcc.h:7