26 #if defined(__cplusplus) 28 #include <type_traits> 32 #if defined(__clang__) && (__clang_major__ > 5) 33 typedef _Float16 _Float16_2
__attribute__((ext_vector_type(2)));
37 static_assert(
sizeof(_Float16) ==
sizeof(
unsigned short),
"");
46 static_assert(
sizeof(_Float16_2) ==
sizeof(
unsigned short[2]),
"");
56 #if defined(__cplusplus) 57 #include "hip_fp16_math_fwd.h" 58 #include "hip_vector_types.h" 63 template<>
struct is_floating_point<_Float16> : std::true_type {};
66 template<
bool cond,
typename T =
void>
67 using Enable_if_t =
typename std::enable_if<cond, T>::type;
73 static_assert(
sizeof(_Float16) ==
sizeof(
unsigned short),
"");
84 #if !defined(__HIP_NO_HALF_CONVERSIONS__) 86 __half(decltype(data) x) : data{x} {}
89 Enable_if_t<std::is_floating_point<T>{}>* =
nullptr>
91 __half(T x) : data{
static_cast<_Float16
>(x)} {}
94 __half(
const __half&) =
default;
96 __half(__half&&) =
default;
101 #if !defined(__HIP_NO_HALF_CONVERSIONS__) 103 typename T, Enable_if_t<std::is_integral<T>{}>* =
nullptr>
105 __half(T x) : data{
static_cast<_Float16
>(x)} {}
110 __half& operator=(
const __half&) =
default;
112 __half& operator=(__half&&) =
default;
120 volatile __half& operator=(
const __half_raw& x)
volatile 125 volatile __half& operator=(
const volatile __half_raw& x)
volatile 135 volatile __half& operator=(
__half_raw&& x)
volatile 140 volatile __half& operator=(
volatile __half_raw&& x)
volatile 145 #if !defined(__HIP_NO_HALF_CONVERSIONS__) 148 Enable_if_t<std::is_floating_point<T>{}>* =
nullptr>
150 __half& operator=(T x)
152 data =
static_cast<_Float16
>(x);
158 #if !defined(__HIP_NO_HALF_CONVERSIONS__) 160 typename T, Enable_if_t<std::is_integral<T>{}>* =
nullptr>
162 __half& operator=(T x)
164 data =
static_cast<_Float16
>(x);
169 #if !defined(__HIP_NO_HALF_OPERATORS__) 171 __half& operator+=(
const __half& x)
177 __half& operator-=(
const __half& x)
183 __half& operator*=(
const __half& x)
189 __half& operator/=(
const __half& x)
195 __half& operator++() { ++data;
return *
this; }
197 __half operator++(
int)
204 __half& operator--() { --data;
return *
this; }
206 __half operator--(
int)
215 #if !defined(__HIP_NO_HALF_CONVERSIONS__) 219 std::is_floating_point<T>{} &&
220 !std::is_same<T, double>{}>* =
nullptr>
221 operator T()
const {
return data; }
232 #if !defined(__HIP_NO_HALF_CONVERSIONS__) 234 typename T, Enable_if_t<std::is_integral<T>{}>* =
nullptr>
236 operator T()
const {
return data; }
239 #if !defined(__HIP_NO_HALF_OPERATORS__) 241 __half operator+()
const {
return *
this; }
243 __half operator-()
const 246 tmp.data = -tmp.data;
252 #if !defined(__HIP_NO_HALF_OPERATORS__) 256 __half operator+(
const __half& x,
const __half& y)
258 return __half{x} += y;
263 __half operator-(
const __half& x,
const __half& y)
265 return __half{x} -= y;
270 __half operator*(
const __half& x,
const __half& y)
272 return __half{x} *= y;
277 __half operator/(
const __half& x,
const __half& y)
279 return __half{x} /= y;
284 bool operator==(
const __half& x,
const __half& y)
286 return x.data == y.data;
291 bool operator!=(
const __half& x,
const __half& y)
298 bool operator<(
const __half& x,
const __half& y)
300 return x.data < y.data;
305 bool operator>(
const __half& x,
const __half& y)
307 return y.data < x.data;
312 bool operator<=(
const __half& x,
const __half& y)
319 bool operator>=(
const __half& x,
const __half& y)
323 #endif // !defined(__HIP_NO_HALF_OPERATORS__) 332 sizeof(_Float16_2) ==
sizeof(
unsigned short[2]),
"");
347 __half2(decltype(data) x) : data{x} {}
349 __half2(
const __half& x,
const __half& y)
353 static_cast<__half_raw>(y).data}
356 __half2(
const __half2&) =
default;
358 __half2(__half2&&) =
default;
360 ~__half2() =
default;
364 __half2& operator=(
const __half2&) =
default;
366 __half2& operator=(__half2&&) =
default;
375 #if !defined(__HIP_NO_HALF_OPERATORS__) 377 __half2& operator+=(
const __half2& x)
383 __half2& operator-=(
const __half2& x)
389 __half2& operator*=(
const __half2& x)
395 __half2& operator/=(
const __half2& x)
401 __half2& operator++() {
return *
this += _Float16_2{1, 1}; }
403 __half2 operator++(
int)
410 __half2& operator--() {
return *
this -= _Float16_2{1, 1}; }
412 __half2 operator--(
int)
422 operator decltype(data)()
const {
return data; }
427 #if !defined(__HIP_NO_HALF_OPERATORS__) 429 __half2 operator+()
const {
return *
this; }
431 __half2 operator-()
const 434 tmp.data = -tmp.data;
440 #if !defined(__HIP_NO_HALF_OPERATORS__) 444 __half2 operator+(
const __half2& x,
const __half2& y)
446 return __half2{x} += y;
451 __half2 operator-(
const __half2& x,
const __half2& y)
453 return __half2{x} -= y;
458 __half2 operator*(
const __half2& x,
const __half2& y)
460 return __half2{x} *= y;
465 __half2 operator/(
const __half2& x,
const __half2& y)
467 return __half2{x} /= y;
472 bool operator==(
const __half2& x,
const __half2& y)
474 auto r = x.data == y.data;
475 return r.x != 0 && r.y != 0;
480 bool operator!=(
const __half2& x,
const __half2& y)
487 bool operator<(
const __half2& x,
const __half2& y)
489 auto r = x.data < y.data;
490 return r.x != 0 && r.y != 0;
495 bool operator>(
const __half2& x,
const __half2& y)
502 bool operator<=(
const __half2& x,
const __half2& y)
509 bool operator>=(
const __half2& x,
const __half2& y)
513 #endif // !defined(__HIP_NO_HALF_OPERATORS__) 521 __half2 make_half2(__half x, __half y)
523 return __half2{x, y};
528 __half __low2half(__half2 x)
535 __half __high2half(__half2 x)
542 __half2 __half2half2(__half x)
544 return __half2{x, x};
549 __half2 __halves2half2(__half x, __half y)
551 return __half2{x, y};
556 __half2 __low2half2(__half2 x)
561 static_cast<__half2_raw>(x).data.x}};
566 __half2 __high2half2(__half2 x)
571 static_cast<__half2_raw>(x).data.y}};
576 __half2 __lows2half2(__half2 x, __half2 y)
581 static_cast<__half2_raw>(y).data.x}};
586 __half2 __highs2half2(__half2 x, __half2 y)
591 static_cast<__half2_raw>(y).data.y}};
596 __half2 __lowhigh2highlow(__half2 x)
601 static_cast<__half2_raw>(x).data.x}};
607 short __half_as_short(__half x)
614 unsigned short __half_as_ushort(__half x)
621 __half __short_as_half(
short x)
629 __half __ushort_as_half(
unsigned short x)
639 __half __float2half(
float x)
645 __half __float2half_rn(
float x)
651 __half __float2half_rz(
float x)
657 __half __float2half_rd(
float x)
663 __half __float2half_ru(
float x)
669 __half2 __float2half2_rn(
float x)
673 static_cast<_Float16
>(x), static_cast<_Float16>(x)}};
677 __half2 __floats2half2_rn(
float x,
float y)
680 static_cast<_Float16
>(x), static_cast<_Float16>(y)}};
684 __half2 __float22half2_rn(float2 x)
686 return __floats2half2_rn(x.x, x.y);
692 float __half2float(__half x)
698 float __low2float(__half2 x)
704 float __high2float(__half2 x)
710 float2 __half22float2(__half2 x)
713 static_cast<__half2_raw>(x).data.x,
714 static_cast<__half2_raw>(x).data.y);
720 int __half2int_rn(__half x)
726 int __half2int_rz(__half x)
732 int __half2int_rd(__half x)
738 int __half2int_ru(__half x)
746 __half __int2half_rn(
int x)
752 __half __int2half_rz(
int x)
758 __half __int2half_rd(
int x)
764 __half __int2half_ru(
int x)
772 short __half2short_rn(__half x)
778 short __half2short_rz(__half x)
784 short __half2short_rd(__half x)
790 short __half2short_ru(__half x)
798 __half __short2half_rn(
short x)
804 __half __short2half_rz(
short x)
810 __half __short2half_rd(
short x)
816 __half __short2half_ru(
short x)
824 long long __half2ll_rn(__half x)
830 long long __half2ll_rz(__half x)
836 long long __half2ll_rd(__half x)
842 long long __half2ll_ru(__half x)
850 __half __ll2half_rn(
long long x)
856 __half __ll2half_rz(
long long x)
862 __half __ll2half_rd(
long long x)
868 __half __ll2half_ru(
long long x)
876 unsigned int __half2uint_rn(__half x)
882 unsigned int __half2uint_rz(__half x)
888 unsigned int __half2uint_rd(__half x)
894 unsigned int __half2uint_ru(__half x)
902 __half __uint2half_rn(
unsigned int x)
908 __half __uint2half_rz(
unsigned int x)
914 __half __uint2half_rd(
unsigned int x)
920 __half __uint2half_ru(
unsigned int x)
928 unsigned short __half2ushort_rn(__half x)
934 unsigned short __half2ushort_rz(__half x)
940 unsigned short __half2ushort_rd(__half x)
946 unsigned short __half2ushort_ru(__half x)
954 __half __ushort2half_rn(
unsigned short x)
960 __half __ushort2half_rz(
unsigned short x)
966 __half __ushort2half_rd(
unsigned short x)
972 __half __ushort2half_ru(
unsigned short x)
980 unsigned long long __half2ull_rn(__half x)
986 unsigned long long __half2ull_rz(__half x)
992 unsigned long long __half2ull_rd(__half x)
998 unsigned long long __half2ull_ru(__half x)
1006 __half __ull2half_rn(
unsigned long long x)
1012 __half __ull2half_rz(
unsigned long long x)
1018 __half __ull2half_rd(
unsigned long long x)
1024 __half __ull2half_ru(
unsigned long long x)
1032 __half __ldg(
const __half* ptr) {
return *ptr; }
1035 __half __ldcg(
const __half* ptr) {
return *ptr; }
1038 __half __ldca(
const __half* ptr) {
return *ptr; }
1041 __half __ldcs(
const __half* ptr) {
return *ptr; }
1045 __half2 __ldg(
const __half2* ptr) {
return *ptr; }
1048 __half2 __ldcg(
const __half2* ptr) {
return *ptr; }
1051 __half2 __ldca(
const __half2* ptr) {
return *ptr; }
1054 __half2 __ldcs(
const __half2* ptr) {
return *ptr; }
1059 bool __heq(__half x, __half y)
1062 static_cast<__half_raw>(y).data;
1066 bool __hne(__half x, __half y)
1069 static_cast<__half_raw>(y).data;
1073 bool __hle(__half x, __half y)
1076 static_cast<__half_raw>(y).data;
1080 bool __hge(__half x, __half y)
1083 static_cast<__half_raw>(y).data;
1087 bool __hlt(__half x, __half y)
1090 static_cast<__half_raw>(y).data;
1094 bool __hgt(__half x, __half y)
1097 static_cast<__half_raw>(y).data;
1101 bool __hequ(__half x, __half y) {
return __heq(x, y); }
1104 bool __hneu(__half x, __half y) {
return __hne(x, y); }
1107 bool __hleu(__half x, __half y) {
return __hle(x, y); }
1110 bool __hgeu(__half x, __half y) {
return __hge(x, y); }
1113 bool __hltu(__half x, __half y) {
return __hlt(x, y); }
1116 bool __hgtu(__half x, __half y) {
return __hgt(x, y); }
1120 __half2 __heq2(__half2 x, __half2 y)
1123 static_cast<__half2_raw>(y).data;
1125 static_cast<_Float16
>(r.x), static_cast<_Float16>(r.y)}};
1129 __half2 __hne2(__half2 x, __half2 y)
1132 static_cast<__half2_raw>(y).data;
1134 static_cast<_Float16
>(r.x), static_cast<_Float16>(r.y)}};
1138 __half2 __hle2(__half2 x, __half2 y)
1141 static_cast<__half2_raw>(y).data;
1143 static_cast<_Float16
>(r.x), static_cast<_Float16>(r.y)}};
1147 __half2 __hge2(__half2 x, __half2 y)
1150 static_cast<__half2_raw>(y).data;
1152 static_cast<_Float16
>(r.x), static_cast<_Float16>(r.y)}};
1156 __half2 __hlt2(__half2 x, __half2 y)
1159 static_cast<__half2_raw>(y).data;
1161 static_cast<_Float16
>(r.x), static_cast<_Float16>(r.y)}};
1165 __half2 __hgt2(__half2 x, __half2 y)
1168 static_cast<__half2_raw>(y).data;
1170 static_cast<_Float16
>(r.x), static_cast<_Float16>(r.y)}};
1174 __half2 __hequ2(__half2 x, __half2 y) {
return __heq2(x, y); }
1177 __half2 __hneu2(__half2 x, __half2 y) {
return __hne2(x, y); }
1180 __half2 __hleu2(__half2 x, __half2 y) {
return __hle2(x, y); }
1183 __half2 __hgeu2(__half2 x, __half2 y) {
return __hge2(x, y); }
1186 __half2 __hltu2(__half2 x, __half2 y) {
return __hlt2(x, y); }
1189 __half2 __hgtu2(__half2 x, __half2 y) {
return __hgt2(x, y); }
1193 bool __hbeq2(__half2 x, __half2 y)
1196 return r.data.x != 0 && r.data.y != 0;
1200 bool __hbne2(__half2 x, __half2 y)
1203 return r.data.x != 0 && r.data.y != 0;
1207 bool __hble2(__half2 x, __half2 y)
1210 return r.data.x != 0 && r.data.y != 0;
1214 bool __hbge2(__half2 x, __half2 y)
1217 return r.data.x != 0 && r.data.y != 0;
1221 bool __hblt2(__half2 x, __half2 y)
1224 return r.data.x != 0 && r.data.y != 0;
1228 bool __hbgt2(__half2 x, __half2 y)
1231 return r.data.x != 0 && r.data.y != 0;
1235 bool __hbequ2(__half2 x, __half2 y) {
return __hbeq2(x, y); }
1238 bool __hbneu2(__half2 x, __half2 y) {
return __hbne2(x, y); }
1241 bool __hbleu2(__half2 x, __half2 y) {
return __hble2(x, y); }
1244 bool __hbgeu2(__half2 x, __half2 y) {
return __hbge2(x, y); }
1247 bool __hbltu2(__half2 x, __half2 y) {
return __hblt2(x, y); }
1250 bool __hbgtu2(__half2 x, __half2 y) {
return __hbgt2(x, y); }
1255 __half __clamp_01(__half x)
1266 __half __hadd(__half x, __half y)
1270 static_cast<__half_raw>(y).data};
1274 __half __hsub(__half x, __half y)
1278 static_cast<__half_raw>(y).data};
1282 __half __hmul(__half x, __half y)
1286 static_cast<__half_raw>(y).data};
1290 __half __hadd_sat(__half x, __half y)
1292 return __clamp_01(__hadd(x, y));
1296 __half __hsub_sat(__half x, __half y)
1298 return __clamp_01(__hsub(x, y));
1302 __half __hmul_sat(__half x, __half y)
1304 return __clamp_01(__hmul(x, y));
1308 __half __hfma(__half x, __half y, __half z)
1311 static_cast<__half_raw>(x).data,
1312 static_cast<__half_raw>(y).data,
1313 static_cast<__half_raw>(z).data)};
1317 __half __hfma_sat(__half x, __half y, __half z)
1319 return __clamp_01(__hfma(x, y, z));
1323 __half __hdiv(__half x, __half y)
1327 static_cast<__half_raw>(y).data};
1332 __half2 __hadd2(__half2 x, __half2 y)
1336 static_cast<__half2_raw>(y).data};
1340 __half2 __hsub2(__half2 x, __half2 y)
1344 static_cast<__half2_raw>(y).data};
1348 __half2 __hmul2(__half2 x, __half2 y)
1352 static_cast<__half2_raw>(y).data};
1356 __half2 __hadd2_sat(__half2 x, __half2 y)
1365 __half2 __hsub2_sat(__half2 x, __half2 y)
1374 __half2 __hmul2_sat(__half2 x, __half2 y)
1383 __half2 __hfma2(__half2 x, __half2 y, __half2 z)
1389 __half2 __hfma2_sat(__half2 x, __half2 y, __half2 z)
1391 auto r =
static_cast<__half2_raw>(__hfma2(x, y, z));
1398 __half2 __h2div(__half2 x, __half2 y)
1402 static_cast<__half2_raw>(y).data};
1406 #if (__hcc_workweek__ >= 19015) || __HIP_CLANG_ONLY__ 1409 float amd_mixed_dot(__half2 a, __half2 b,
float c,
bool saturate) {
1410 return __ockl_fdot2(static_cast<__half2_raw>(a).data,
1411 static_cast<__half2_raw>(b).data,
1417 __half htrunc(__half x)
1420 __ocml_trunc_f16(static_cast<__half_raw>(x).data)};
1424 __half hceil(__half x)
1427 __ocml_ceil_f16(static_cast<__half_raw>(x).data)};
1431 __half hfloor(__half x)
1434 __ocml_floor_f16(static_cast<__half_raw>(x).data)};
1438 __half hrint(__half x)
1441 __ocml_rint_f16(static_cast<__half_raw>(x).data)};
1445 __half hsin(__half x)
1448 __ocml_sin_f16(static_cast<__half_raw>(x).data)};
1452 __half hcos(__half x)
1455 __ocml_cos_f16(static_cast<__half_raw>(x).data)};
1459 __half hexp(__half x)
1462 __ocml_exp_f16(static_cast<__half_raw>(x).data)};
1466 __half hexp2(__half x)
1469 __ocml_exp2_f16(static_cast<__half_raw>(x).data)};
1473 __half hexp10(__half x)
1476 __ocml_exp10_f16(static_cast<__half_raw>(x).data)};
1480 __half hlog2(__half x)
1483 __ocml_log2_f16(static_cast<__half_raw>(x).data)};
1487 __half hlog(__half x)
1490 __ocml_log_f16(static_cast<__half_raw>(x).data)};
1494 __half hlog10(__half x)
1497 __ocml_log10_f16(static_cast<__half_raw>(x).data)};
1501 __half hrcp(__half x)
1504 __llvm_amdgcn_rcp_f16(static_cast<__half_raw>(x).data)};
1508 __half hrsqrt(__half x)
1511 __ocml_rsqrt_f16(static_cast<__half_raw>(x).data)};
1515 __half hsqrt(__half x)
1518 __ocml_sqrt_f16(static_cast<__half_raw>(x).data)};
1522 bool __hisinf(__half x)
1524 return __ocml_isinf_f16(static_cast<__half_raw>(x).data);
1528 bool __hisnan(__half x)
1530 return __ocml_isnan_f16(static_cast<__half_raw>(x).data);
1534 __half __hneg(__half x)
1541 __half2 h2trunc(__half2 x)
1547 __half2 h2ceil(__half2 x)
1553 __half2 h2floor(__half2 x)
1559 __half2 h2rint(__half2 x)
1565 __half2 h2sin(__half2 x)
1571 __half2 h2cos(__half2 x)
1577 __half2 h2exp(__half2 x)
1583 __half2 h2exp2(__half2 x)
1589 __half2 h2exp10(__half2 x)
1595 __half2 h2log2(__half2 x)
1601 __half2 h2log(__half2 x) {
return __ocml_log_2f16(x); }
1604 __half2 h2log10(__half2 x) {
return __ocml_log10_2f16(x); }
1607 __half2 h2rcp(__half2 x) {
return __llvm_amdgcn_rcp_2f16(x); }
1610 __half2 h2rsqrt(__half2 x) {
return __ocml_rsqrt_2f16(x); }
1613 __half2 h2sqrt(__half2 x) {
return __ocml_sqrt_2f16(x); }
1616 __half2 __hisinf2(__half2 x)
1618 auto r = __ocml_isinf_2f16(x);
1620 static_cast<_Float16
>(r.x), static_cast<_Float16>(r.y)}};
1624 __half2 __hisnan2(__half2 x)
1626 auto r = __ocml_isnan_2f16(x);
1628 static_cast<_Float16
>(r.x), static_cast<_Float16>(r.y)}};
1632 __half2 __hneg2(__half2 x)
1638 #if !defined(HIP_NO_HALF) 1639 using half = __half;
1640 using half2 = __half2;
1642 #endif // defined(__cplusplus) 1643 #elif defined(__GNUC__) 1644 #include "hip_fp16_gcc.h" 1645 #endif // !defined(__clang__) && defined(__GNUC__)
_Float16 __2f16 __attribute__((ext_vector_type(2)))
Copies the memory address of symbol symbolName to devPtr.
Definition: hip_fp16_math_fwd.h:53
Definition: hip_fp16_gcc.h:11
#define __host__
Definition: host_defines.h:41
Definition: hip_fp16_gcc.h:7