diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h index 6fa86e94b9..670cd475c2 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h @@ -529,35 +529,35 @@ THE SOFTWARE. } inline - __device__ + __host__ __device__ __half __low2half(__half2 x) { return __half{__half_raw{static_cast<__half2_raw>(x).data.x}}; } inline - __device__ + __host__ __device__ __half __high2half(__half2 x) { return __half{__half_raw{static_cast<__half2_raw>(x).data.y}}; } inline - __device__ + __host__ __device__ __half2 __half2half2(__half x) { return __half2{x, x}; } inline - __device__ + __host__ __device__ __half2 __halves2half2(__half x, __half y) { return __half2{x, y}; } inline - __device__ + __host__ __device__ __half2 __low2half2(__half2 x) { return __half2{ @@ -567,7 +567,7 @@ THE SOFTWARE. } inline - __device__ + __host__ __device__ __half2 __high2half2(__half2 x) { return __half2_raw{ @@ -577,7 +577,7 @@ THE SOFTWARE. } inline - __device__ + __host__ __device__ __half2 __lows2half2(__half2 x, __half2 y) { return __half2_raw{ @@ -587,7 +587,7 @@ THE SOFTWARE. } inline - __device__ + __host__ __device__ __half2 __highs2half2(__half2 x, __half2 y) { return __half2_raw{ @@ -597,7 +597,7 @@ THE SOFTWARE. } inline - __device__ + __host__ __device__ __half2 __lowhigh2highlow(__half2 x) { return __half2_raw{ @@ -1046,16 +1046,16 @@ THE SOFTWARE. __half __ldcs(const __half* ptr) { return *ptr; } inline - __device__ + __host__ __device__ __half2 __ldg(const __half2* ptr) { return *ptr; } inline - __device__ + __host__ __device__ __half2 __ldcg(const __half2* ptr) { return *ptr; } inline - __device__ + __host__ __device__ __half2 __ldca(const __half2* ptr) { return *ptr; } inline - __device__ + __host__ __device__ __half2 __ldcs(const __half2* ptr) { return *ptr; } // Relations @@ -1121,7 +1121,7 @@ THE SOFTWARE. bool __hgtu(__half x, __half y) { return __hgt(x, y); } inline - __device__ + __host__ __device__ __half2 __heq2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(x).data == @@ -1129,7 +1129,7 @@ THE SOFTWARE. return __builtin_convertvector(-r, _Float16_2); } inline - __device__ + __host__ __device__ __half2 __hne2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(x).data != @@ -1137,7 +1137,7 @@ THE SOFTWARE. return __builtin_convertvector(-r, _Float16_2); } inline - __device__ + __host__ __device__ __half2 __hle2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(x).data <= @@ -1145,7 +1145,7 @@ THE SOFTWARE. return __builtin_convertvector(-r, _Float16_2); } inline - __device__ + __host__ __device__ __half2 __hge2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(x).data >= @@ -1153,7 +1153,7 @@ THE SOFTWARE. return __builtin_convertvector(-r, _Float16_2); } inline - __device__ + __host__ __device__ __half2 __hlt2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(x).data < @@ -1161,7 +1161,7 @@ THE SOFTWARE. return __builtin_convertvector(-r, _Float16_2); } inline - __device__ + __host__ __device__ __half2 __hgt2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(x).data > @@ -1169,83 +1169,83 @@ THE SOFTWARE. return __builtin_convertvector(-r, _Float16_2); } inline - __device__ + __host__ __device__ __half2 __hequ2(__half2 x, __half2 y) { return __heq2(x, y); } inline - __device__ + __host__ __device__ __half2 __hneu2(__half2 x, __half2 y) { return __hne2(x, y); } inline - __device__ + __host__ __device__ __half2 __hleu2(__half2 x, __half2 y) { return __hle2(x, y); } inline - __device__ + __host__ __device__ __half2 __hgeu2(__half2 x, __half2 y) { return __hge2(x, y); } inline - __device__ + __host__ __device__ __half2 __hltu2(__half2 x, __half2 y) { return __hlt2(x, y); } inline - __device__ + __host__ __device__ __half2 __hgtu2(__half2 x, __half2 y) { return __hgt2(x, y); } inline - __device__ + __host__ __device__ bool __hbeq2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(__heq2(x, y)); return r.data.x != 0 && r.data.y != 0; } inline - __device__ + __host__ __device__ bool __hbne2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(__hne2(x, y)); return r.data.x != 0 && r.data.y != 0; } inline - __device__ + __host__ __device__ bool __hble2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(__hle2(x, y)); return r.data.x != 0 && r.data.y != 0; } inline - __device__ + __host__ __device__ bool __hbge2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(__hge2(x, y)); return r.data.x != 0 && r.data.y != 0; } inline - __device__ + __host__ __device__ bool __hblt2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(__hlt2(x, y)); return r.data.x != 0 && r.data.y != 0; } inline - __device__ + __host__ __device__ bool __hbgt2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(__hgt2(x, y)); return r.data.x != 0 && r.data.y != 0; } inline - __device__ + __host__ __device__ bool __hbequ2(__half2 x, __half2 y) { return __hbeq2(x, y); } inline - __device__ + __host__ __device__ bool __hbneu2(__half2 x, __half2 y) { return __hbne2(x, y); } inline - __device__ + __host__ __device__ bool __hbleu2(__half2 x, __half2 y) { return __hble2(x, y); } inline - __device__ + __host__ __device__ bool __hbgeu2(__half2 x, __half2 y) { return __hbge2(x, y); } inline - __device__ + __host__ __device__ bool __hbltu2(__half2 x, __half2 y) { return __hblt2(x, y); } inline - __device__ + __host__ __device__ bool __hbgtu2(__half2 x, __half2 y) { return __hbgt2(x, y); } // Arithmetic @@ -1334,7 +1334,7 @@ THE SOFTWARE. } inline - __device__ + __host__ __device__ __half2 __hadd2(__half2 x, __half2 y) { return __half2_raw{ @@ -1342,14 +1342,14 @@ THE SOFTWARE. static_cast<__half2_raw>(y).data}; } inline - __device__ + __host__ __device__ __half2 __habs2(__half2 x) { return __half2_raw{ __ocml_fabs_2f16(static_cast<__half2_raw>(x).data)}; } inline - __device__ + __host__ __device__ __half2 __hsub2(__half2 x, __half2 y) { return __half2_raw{ @@ -1357,7 +1357,7 @@ THE SOFTWARE. static_cast<__half2_raw>(y).data}; } inline - __device__ + __host__ __device__ __half2 __hmul2(__half2 x, __half2 y) { return __half2_raw{ @@ -1365,7 +1365,7 @@ THE SOFTWARE. static_cast<__half2_raw>(y).data}; } inline - __device__ + __host__ __device__ __half2 __hadd2_sat(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(__hadd2(x, y)); @@ -1374,7 +1374,7 @@ THE SOFTWARE. __clamp_01(__half_raw{r.data.y})}; } inline - __device__ + __host__ __device__ __half2 __hsub2_sat(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(__hsub2(x, y)); @@ -1383,7 +1383,7 @@ THE SOFTWARE. __clamp_01(__half_raw{r.data.y})}; } inline - __device__ + __host__ __device__ __half2 __hmul2_sat(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(__hmul2(x, y)); @@ -1392,13 +1392,13 @@ THE SOFTWARE. __clamp_01(__half_raw{r.data.y})}; } inline - __device__ + __host__ __device__ __half2 __hfma2(__half2 x, __half2 y, __half2 z) { return __half2_raw{__ocml_fma_2f16(x, y, z)}; } inline - __device__ + __host__ __device__ __half2 __hfma2_sat(__half2 x, __half2 y, __half2 z) { auto r = static_cast<__half2_raw>(__hfma2(x, y, z)); @@ -1407,7 +1407,7 @@ THE SOFTWARE. __clamp_01(__half_raw{r.data.y})}; } inline - __device__ + __host__ __device__ __half2 __h2div(__half2 x, __half2 y) { return __half2_raw{ @@ -1550,82 +1550,82 @@ THE SOFTWARE. } inline - __device__ + __host__ __device__ __half2 h2trunc(__half2 x) { return __half2_raw{__ocml_trunc_2f16(x)}; } inline - __device__ + __host__ __device__ __half2 h2ceil(__half2 x) { return __half2_raw{__ocml_ceil_2f16(x)}; } inline - __device__ + __host__ __device__ __half2 h2floor(__half2 x) { return __half2_raw{__ocml_floor_2f16(x)}; } inline - __device__ + __host__ __device__ __half2 h2rint(__half2 x) { return __half2_raw{__ocml_rint_2f16(x)}; } inline - __device__ + __host__ __device__ __half2 h2sin(__half2 x) { return __half2_raw{__ocml_sin_2f16(x)}; } inline - __device__ + __host__ __device__ __half2 h2cos(__half2 x) { return __half2_raw{__ocml_cos_2f16(x)}; } inline - __device__ + __host__ __device__ __half2 h2exp(__half2 x) { return __half2_raw{__ocml_exp_2f16(x)}; } inline - __device__ + __host__ __device__ __half2 h2exp2(__half2 x) { return __half2_raw{__ocml_exp2_2f16(x)}; } inline - __device__ + __host__ __device__ __half2 h2exp10(__half2 x) { return __half2_raw{__ocml_exp10_2f16(x)}; } inline - __device__ + __host__ __device__ __half2 h2log2(__half2 x) { return __half2_raw{__ocml_log2_2f16(x)}; } inline - __device__ + __host__ __device__ __half2 h2log(__half2 x) { return __ocml_log_2f16(x); } inline - __device__ + __host__ __device__ __half2 h2log10(__half2 x) { return __ocml_log10_2f16(x); } inline - __device__ + __host__ __device__ __half2 h2rcp(__half2 x) { return __llvm_amdgcn_rcp_2f16(x); } inline - __device__ + __host__ __device__ __half2 h2rsqrt(__half2 x) { return __ocml_rsqrt_2f16(x); } inline - __device__ + __host__ __device__ __half2 h2sqrt(__half2 x) { return __ocml_sqrt_2f16(x); } inline - __device__ + __host__ __device__ __half2 __hisinf2(__half2 x) { auto r = __ocml_isinf_2f16(x); @@ -1633,7 +1633,7 @@ THE SOFTWARE. static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}}; } inline - __device__ + __host__ __device__ __half2 __hisnan2(__half2 x) { auto r = __ocml_isnan_2f16(x); @@ -1641,7 +1641,7 @@ THE SOFTWARE. static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}}; } inline - __device__ + __host__ __device__ __half2 __hneg2(__half2 x) { return __half2_raw{-static_cast<__half2_raw>(x).data};