From c914642086cf619841f8658d21746f7207c2f704 Mon Sep 17 00:00:00 2001 From: "Chaudhary, Jatin Jaikishan" Date: Fri, 6 Jun 2025 16:52:37 +0100 Subject: [PATCH] SWDEV-529899 - [ABI Break] remove unnamed namespace from fp16 operations (#244) Change-Id: I4feb9089b4cee145300dc78fc43a4f42473f5b1c [ROCm/clr commit: 543925e048884e40586024034bb8e9b3b84051e5] --- .../include/hip/amd_detail/amd_hip_fp16.h | 1483 ++++++++--------- 1 file changed, 740 insertions(+), 743 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h index d8ba155081..814575ea1e 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h @@ -532,772 +532,770 @@ THE SOFTWARE. }; // END STRUCT __HALF2 - namespace + inline + __HOST_DEVICE__ + __half2 make_half2(__half x, __half y) { - inline - __HOST_DEVICE__ - __half2 make_half2(__half x, __half y) - { - return __half2{x, y}; - } + return __half2{x, y}; + } - inline - __HOST_DEVICE__ - __half __low2half(__half2 x) - { - return __half{__half_raw{static_cast<__half2_raw>(x).data.x}}; - } + inline + __HOST_DEVICE__ + __half __low2half(__half2 x) + { + return __half{__half_raw{static_cast<__half2_raw>(x).data.x}}; + } - inline - __HOST_DEVICE__ - __half __high2half(__half2 x) - { - return __half{__half_raw{static_cast<__half2_raw>(x).data.y}}; - } + inline + __HOST_DEVICE__ + __half __high2half(__half2 x) + { + return __half{__half_raw{static_cast<__half2_raw>(x).data.y}}; + } - inline - __HOST_DEVICE__ - __half2 __half2half2(__half x) - { - return __half2{x, x}; - } + inline + __HOST_DEVICE__ + __half2 __half2half2(__half x) + { + return __half2{x, x}; + } - inline - __HOST_DEVICE__ - __half2 __halves2half2(__half x, __half y) - { - return __half2{x, y}; - } + inline + __HOST_DEVICE__ + __half2 __halves2half2(__half x, __half y) + { + return __half2{x, y}; + } - inline - __HOST_DEVICE__ - __half2 __low2half2(__half2 x) - { - return __half2{ - _Float16_2{ - static_cast<__half2_raw>(x).data.x, - static_cast<__half2_raw>(x).data.x}}; - } - - inline - __HOST_DEVICE__ - __half2 __high2half2(__half2 x) - { - return __half2{ - _Float16_2{ - static_cast<__half2_raw>(x).data.y, - static_cast<__half2_raw>(x).data.y}}; - } - - inline - __HOST_DEVICE__ - __half2 __lows2half2(__half2 x, __half2 y) - { - return __half2{ - _Float16_2{ - static_cast<__half2_raw>(x).data.x, - static_cast<__half2_raw>(y).data.x}}; - } - - inline - __HOST_DEVICE__ - __half2 __highs2half2(__half2 x, __half2 y) - { - return __half2{ - _Float16_2{ - static_cast<__half2_raw>(x).data.y, - static_cast<__half2_raw>(y).data.y}}; - } - - inline - __HOST_DEVICE__ - __half2 __lowhigh2highlow(__half2 x) - { - return __half2{ - _Float16_2{ - static_cast<__half2_raw>(x).data.y, - static_cast<__half2_raw>(x).data.x}}; - } - - // Bitcasts - inline - __HOST_DEVICE__ - short __half_as_short(__half x) - { - return static_cast<__half_raw>(x).x; - } - - inline - __HOST_DEVICE__ - unsigned short __half_as_ushort(__half x) - { - return static_cast<__half_raw>(x).x; - } - - inline - __HOST_DEVICE__ - __half __short_as_half(short x) - { - __half_raw r; r.x = x; - return r; - } - - inline - __HOST_DEVICE__ - __half __ushort_as_half(unsigned short x) - { - __half_raw r; r.x = x; - return r; - } - - // float -> half | half2 - inline - __HOST_DEVICE__ - __half __float2half(float x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - inline - __HOST_DEVICE__ - __half __float2half_rn(float x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - #if !defined(__HIPCC_RTC__) - // TODO: rounding behaviour is not correct for host functions. - inline - __host__ - __half __float2half_rz(float x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - inline - __host__ - __half __float2half_rd(float x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - inline - __host__ - __half __float2half_ru(float x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - #endif - inline - __device__ - __half __float2half_rz(float x) - { - return __half_raw{__ocml_cvtrtz_f16_f32(x)}; - } - inline - __device__ - __half __float2half_rd(float x) - { - return __half_raw{__ocml_cvtrtn_f16_f32(x)}; - } - inline - __device__ - __half __float2half_ru(float x) - { - return __half_raw{__ocml_cvtrtp_f16_f32(x)}; - } - inline - __HOST_DEVICE__ - __half2 __float2half2_rn(float x) - { - return __half2{ - _Float16_2{ - static_cast<_Float16>(x), static_cast<_Float16>(x)}}; - } - inline - __HOST_DEVICE__ - __half2 __floats2half2_rn(float x, float y) - { - return __half2{_Float16_2{ - static_cast<_Float16>(x), static_cast<_Float16>(y)}}; - } - inline - __HOST_DEVICE__ - __half2 __float22half2_rn(float2 x) - { - return __floats2half2_rn(x.x, x.y); - } - - // half | half2 -> float - inline - __HOST_DEVICE__ - float __half2float(__half x) - { - return static_cast<__half_raw>(x).data; - } - inline - __HOST_DEVICE__ - float __low2float(__half2 x) - { - return static_cast<__half2_raw>(x).data.x; - } - inline - __HOST_DEVICE__ - float __high2float(__half2 x) - { - return static_cast<__half2_raw>(x).data.y; - } - inline - __HOST_DEVICE__ - float2 __half22float2(__half2 x) - { - return make_float2( + inline + __HOST_DEVICE__ + __half2 __low2half2(__half2 x) + { + return __half2{ + _Float16_2{ static_cast<__half2_raw>(x).data.x, - static_cast<__half2_raw>(x).data.y); - } + static_cast<__half2_raw>(x).data.x}}; + } - // half -> int - inline - __device__ - int __half2int_rn(__half x) - { - return static_cast<__half_raw>(x).data; - } - inline - __device__ - int __half2int_rz(__half x) - { - return static_cast<__half_raw>(x).data; - } - inline - __device__ - int __half2int_rd(__half x) - { - return static_cast<__half_raw>(x).data; - } - inline - __device__ - int __half2int_ru(__half x) - { - return static_cast<__half_raw>(x).data; - } + inline + __HOST_DEVICE__ + __half2 __high2half2(__half2 x) + { + return __half2{ + _Float16_2{ + static_cast<__half2_raw>(x).data.y, + static_cast<__half2_raw>(x).data.y}}; + } - // int -> half - inline - __HOST_DEVICE__ - __half __int2half_rn(int x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - inline - __device__ - __half __int2half_rz(int x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - inline - __device__ - __half __int2half_rd(int x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - inline - __device__ - __half __int2half_ru(int x) - { - return __half_raw{static_cast<_Float16>(x)}; - } + inline + __HOST_DEVICE__ + __half2 __lows2half2(__half2 x, __half2 y) + { + return __half2{ + _Float16_2{ + static_cast<__half2_raw>(x).data.x, + static_cast<__half2_raw>(y).data.x}}; + } - // half -> short - inline - __device__ - short __half2short_rn(__half x) - { - return static_cast<__half_raw>(x).data; - } - inline - __device__ - short __half2short_rz(__half x) - { - return static_cast<__half_raw>(x).data; - } - inline - __device__ - short __half2short_rd(__half x) - { - return static_cast<__half_raw>(x).data; - } - inline - __device__ - short __half2short_ru(__half x) - { - return static_cast<__half_raw>(x).data; - } + inline + __HOST_DEVICE__ + __half2 __highs2half2(__half2 x, __half2 y) + { + return __half2{ + _Float16_2{ + static_cast<__half2_raw>(x).data.y, + static_cast<__half2_raw>(y).data.y}}; + } - // short -> half - inline - __device__ - __half __short2half_rn(short x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - inline - __device__ - __half __short2half_rz(short x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - inline - __device__ - __half __short2half_rd(short x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - inline - __device__ - __half __short2half_ru(short x) - { - return __half_raw{static_cast<_Float16>(x)}; - } + inline + __HOST_DEVICE__ + __half2 __lowhigh2highlow(__half2 x) + { + return __half2{ + _Float16_2{ + static_cast<__half2_raw>(x).data.y, + static_cast<__half2_raw>(x).data.x}}; + } - // half -> long long - inline - __device__ - long long __half2ll_rn(__half x) - { - return static_cast<__half_raw>(x).data; - } - inline - __device__ - long long __half2ll_rz(__half x) - { - return static_cast<__half_raw>(x).data; - } - inline - __device__ - long long __half2ll_rd(__half x) - { - return static_cast<__half_raw>(x).data; - } - inline - __device__ - long long __half2ll_ru(__half x) - { - return static_cast<__half_raw>(x).data; - } + // Bitcasts + inline + __HOST_DEVICE__ + short __half_as_short(__half x) + { + return static_cast<__half_raw>(x).x; + } - // long long -> half - inline - __device__ - __half __ll2half_rn(long long x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - inline - __device__ - __half __ll2half_rz(long long x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - inline - __device__ - __half __ll2half_rd(long long x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - inline - __device__ - __half __ll2half_ru(long long x) - { - return __half_raw{static_cast<_Float16>(x)}; - } + inline + __HOST_DEVICE__ + unsigned short __half_as_ushort(__half x) + { + return static_cast<__half_raw>(x).x; + } - // half -> unsigned int - inline - __device__ - unsigned int __half2uint_rn(__half x) - { - return static_cast<__half_raw>(x).data; - } - inline - __device__ - unsigned int __half2uint_rz(__half x) - { - return static_cast<__half_raw>(x).data; - } - inline - __device__ - unsigned int __half2uint_rd(__half x) - { - return static_cast<__half_raw>(x).data; - } - inline - __device__ - unsigned int __half2uint_ru(__half x) - { - return static_cast<__half_raw>(x).data; - } + inline + __HOST_DEVICE__ + __half __short_as_half(short x) + { + __half_raw r; r.x = x; + return r; + } - // unsigned int -> half - inline - __device__ - __half __uint2half_rn(unsigned int x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - inline - __device__ - __half __uint2half_rz(unsigned int x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - inline - __device__ - __half __uint2half_rd(unsigned int x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - inline - __device__ - __half __uint2half_ru(unsigned int x) - { - return __half_raw{static_cast<_Float16>(x)}; - } + inline + __HOST_DEVICE__ + __half __ushort_as_half(unsigned short x) + { + __half_raw r; r.x = x; + return r; + } - // half -> unsigned short - inline - __device__ - unsigned short __half2ushort_rn(__half x) - { - return static_cast<__half_raw>(x).data; - } - inline - __device__ - unsigned short __half2ushort_rz(__half x) - { - return static_cast<__half_raw>(x).data; - } - inline - __device__ - unsigned short __half2ushort_rd(__half x) - { - return static_cast<__half_raw>(x).data; - } - inline - __device__ - unsigned short __half2ushort_ru(__half x) - { - return static_cast<__half_raw>(x).data; - } + // float -> half | half2 + inline + __HOST_DEVICE__ + __half __float2half(float x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + inline + __HOST_DEVICE__ + __half __float2half_rn(float x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + #if !defined(__HIPCC_RTC__) + // TODO: rounding behaviour is not correct for host functions. + inline + __host__ + __half __float2half_rz(float x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + inline + __host__ + __half __float2half_rd(float x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + inline + __host__ + __half __float2half_ru(float x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + #endif + inline + __device__ + __half __float2half_rz(float x) + { + return __half_raw{__ocml_cvtrtz_f16_f32(x)}; + } + inline + __device__ + __half __float2half_rd(float x) + { + return __half_raw{__ocml_cvtrtn_f16_f32(x)}; + } + inline + __device__ + __half __float2half_ru(float x) + { + return __half_raw{__ocml_cvtrtp_f16_f32(x)}; + } + inline + __HOST_DEVICE__ + __half2 __float2half2_rn(float x) + { + return __half2{ + _Float16_2{ + static_cast<_Float16>(x), static_cast<_Float16>(x)}}; + } + inline + __HOST_DEVICE__ + __half2 __floats2half2_rn(float x, float y) + { + return __half2{_Float16_2{ + static_cast<_Float16>(x), static_cast<_Float16>(y)}}; + } + inline + __HOST_DEVICE__ + __half2 __float22half2_rn(float2 x) + { + return __floats2half2_rn(x.x, x.y); + } - // unsigned short -> half - inline - __device__ - __half __ushort2half_rn(unsigned short x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - inline - __device__ - __half __ushort2half_rz(unsigned short x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - inline - __device__ - __half __ushort2half_rd(unsigned short x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - inline - __device__ - __half __ushort2half_ru(unsigned short x) - { - return __half_raw{static_cast<_Float16>(x)}; - } + // half | half2 -> float + inline + __HOST_DEVICE__ + float __half2float(__half x) + { + return static_cast<__half_raw>(x).data; + } + inline + __HOST_DEVICE__ + float __low2float(__half2 x) + { + return static_cast<__half2_raw>(x).data.x; + } + inline + __HOST_DEVICE__ + float __high2float(__half2 x) + { + return static_cast<__half2_raw>(x).data.y; + } + inline + __HOST_DEVICE__ + float2 __half22float2(__half2 x) + { + return make_float2( + static_cast<__half2_raw>(x).data.x, + static_cast<__half2_raw>(x).data.y); + } - // half -> unsigned long long - inline - __device__ - unsigned long long __half2ull_rn(__half x) - { - return static_cast<__half_raw>(x).data; - } - inline - __device__ - unsigned long long __half2ull_rz(__half x) - { - return static_cast<__half_raw>(x).data; - } - inline - __device__ - unsigned long long __half2ull_rd(__half x) - { - return static_cast<__half_raw>(x).data; - } - inline - __device__ - unsigned long long __half2ull_ru(__half x) - { - return static_cast<__half_raw>(x).data; - } + // half -> int + inline + __device__ + int __half2int_rn(__half x) + { + return static_cast<__half_raw>(x).data; + } + inline + __device__ + int __half2int_rz(__half x) + { + return static_cast<__half_raw>(x).data; + } + inline + __device__ + int __half2int_rd(__half x) + { + return static_cast<__half_raw>(x).data; + } + inline + __device__ + int __half2int_ru(__half x) + { + return static_cast<__half_raw>(x).data; + } - // unsigned long long -> half - inline - __device__ - __half __ull2half_rn(unsigned long long x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - inline - __device__ - __half __ull2half_rz(unsigned long long x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - inline - __device__ - __half __ull2half_rd(unsigned long long x) - { - return __half_raw{static_cast<_Float16>(x)}; - } - inline - __device__ - __half __ull2half_ru(unsigned long long x) - { - return __half_raw{static_cast<_Float16>(x)}; - } + // int -> half + inline + __HOST_DEVICE__ + __half __int2half_rn(int x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + inline + __device__ + __half __int2half_rz(int x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + inline + __device__ + __half __int2half_rd(int x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + inline + __device__ + __half __int2half_ru(int x) + { + return __half_raw{static_cast<_Float16>(x)}; + } - // Load primitives - inline - __device__ - __half __ldg(const __half* ptr) { return *ptr; } - inline - __device__ - __half __ldcg(const __half* ptr) { return *ptr; } - inline - __device__ - __half __ldca(const __half* ptr) { return *ptr; } - inline - __device__ - __half __ldcs(const __half* ptr) { return *ptr; } + // half -> short + inline + __device__ + short __half2short_rn(__half x) + { + return static_cast<__half_raw>(x).data; + } + inline + __device__ + short __half2short_rz(__half x) + { + return static_cast<__half_raw>(x).data; + } + inline + __device__ + short __half2short_rd(__half x) + { + return static_cast<__half_raw>(x).data; + } + inline + __device__ + short __half2short_ru(__half x) + { + return static_cast<__half_raw>(x).data; + } - inline - __HOST_DEVICE__ - __half2 __ldg(const __half2* ptr) { return *ptr; } - inline - __HOST_DEVICE__ - __half2 __ldcg(const __half2* ptr) { return *ptr; } - inline - __HOST_DEVICE__ - __half2 __ldca(const __half2* ptr) { return *ptr; } - inline - __HOST_DEVICE__ - __half2 __ldcs(const __half2* ptr) { return *ptr; } + // short -> half + inline + __device__ + __half __short2half_rn(short x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + inline + __device__ + __half __short2half_rz(short x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + inline + __device__ + __half __short2half_rd(short x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + inline + __device__ + __half __short2half_ru(short x) + { + return __half_raw{static_cast<_Float16>(x)}; + } - // Relations - inline - __HOST_DEVICE__ - bool __heq(__half x, __half y) - { - return static_cast<__half_raw>(x).data == - static_cast<__half_raw>(y).data; - } - inline - __HOST_DEVICE__ - bool __hne(__half x, __half y) - { - return static_cast<__half_raw>(x).data != - static_cast<__half_raw>(y).data; - } - inline - __HOST_DEVICE__ - bool __hle(__half x, __half y) - { - return static_cast<__half_raw>(x).data <= - static_cast<__half_raw>(y).data; - } - inline - __HOST_DEVICE__ - bool __hge(__half x, __half y) - { - return static_cast<__half_raw>(x).data >= - static_cast<__half_raw>(y).data; - } - inline - __HOST_DEVICE__ - bool __hlt(__half x, __half y) - { - return static_cast<__half_raw>(x).data < - static_cast<__half_raw>(y).data; - } - inline - __HOST_DEVICE__ - bool __hgt(__half x, __half y) - { - return static_cast<__half_raw>(x).data > - static_cast<__half_raw>(y).data; - } - inline __HOST_DEVICE__ - bool __hequ(__half x, __half y) { - return !(static_cast<__half_raw>(x).data < static_cast<__half_raw>(y).data) && - !(static_cast<__half_raw>(x).data > static_cast<__half_raw>(y).data); - } - inline __HOST_DEVICE__ - bool __hneu(__half x, __half y) { - return !(static_cast<__half_raw>(x).data == static_cast<__half_raw>(y).data); - } - inline __HOST_DEVICE__ - bool __hleu(__half x, __half y) { - return !(static_cast<__half_raw>(x).data > static_cast<__half_raw>(y).data); - } - inline - __HOST_DEVICE__ - bool __hgeu(__half x, __half y) { - return !(static_cast<__half_raw>(x).data < static_cast<__half_raw>(y).data); - } - inline - __HOST_DEVICE__ - bool __hltu(__half x, __half y) { - return !(static_cast<__half_raw>(x).data >= static_cast<__half_raw>(y).data); - } - inline - __HOST_DEVICE__ - bool __hgtu(__half x, __half y) { - return !(static_cast<__half_raw>(x).data <= static_cast<__half_raw>(y).data); - } + // half -> long long + inline + __device__ + long long __half2ll_rn(__half x) + { + return static_cast<__half_raw>(x).data; + } + inline + __device__ + long long __half2ll_rz(__half x) + { + return static_cast<__half_raw>(x).data; + } + inline + __device__ + long long __half2ll_rd(__half x) + { + return static_cast<__half_raw>(x).data; + } + inline + __device__ + long long __half2ll_ru(__half x) + { + return static_cast<__half_raw>(x).data; + } - inline - __HOST_DEVICE__ - __half2 __heq2(__half2 x, __half2 y) - { - auto r = static_cast<__half2_raw>(x).data == - static_cast<__half2_raw>(y).data; - return __builtin_convertvector(-r, _Float16_2); - } - inline - __HOST_DEVICE__ - __half2 __hne2(__half2 x, __half2 y) - { - auto r = static_cast<__half2_raw>(x).data != - static_cast<__half2_raw>(y).data; - return __builtin_convertvector(-r, _Float16_2); - } - inline - __HOST_DEVICE__ - __half2 __hle2(__half2 x, __half2 y) - { - auto r = static_cast<__half2_raw>(x).data <= - static_cast<__half2_raw>(y).data; - return __builtin_convertvector(-r, _Float16_2); - } - inline - __HOST_DEVICE__ - __half2 __hge2(__half2 x, __half2 y) - { - auto r = static_cast<__half2_raw>(x).data >= - static_cast<__half2_raw>(y).data; - return __builtin_convertvector(-r, _Float16_2); - } - inline - __HOST_DEVICE__ - __half2 __hlt2(__half2 x, __half2 y) - { - auto r = static_cast<__half2_raw>(x).data < - static_cast<__half2_raw>(y).data; - return __builtin_convertvector(-r, _Float16_2); - } - inline - __HOST_DEVICE__ - __half2 __hgt2(__half2 x, __half2 y) - { - auto r = static_cast<__half2_raw>(x).data > - static_cast<__half2_raw>(y).data; - return __builtin_convertvector(-r, _Float16_2); - } - inline __HOST_DEVICE__ - __half2 __hequ2(__half2 x, __half2 y) { - auto r = !(static_cast<__half2_raw>(x).data < static_cast<__half2_raw>(y).data) && - !(static_cast<__half2_raw>(x).data > static_cast<__half2_raw>(y).data); - return __builtin_convertvector(-r, _Float16_2); - } - inline - __HOST_DEVICE__ - __half2 __hneu2(__half2 x, __half2 y) { - auto r = !(static_cast<__half2_raw>(x).data == static_cast<__half2_raw>(y).data); - return __builtin_convertvector(-r, _Float16_2); - } - inline - __HOST_DEVICE__ - __half2 __hleu2(__half2 x, __half2 y) { - auto r = !(static_cast<__half2_raw>(x).data > static_cast<__half2_raw>(y).data); - return __builtin_convertvector(-r, _Float16_2); - } - inline - __HOST_DEVICE__ - __half2 __hgeu2(__half2 x, __half2 y) { - auto r = !(static_cast<__half2_raw>(x).data < static_cast<__half2_raw>(y).data); - return __builtin_convertvector(-r, _Float16_2); - } - inline - __HOST_DEVICE__ - __half2 __hltu2(__half2 x, __half2 y) { - auto r = !(static_cast<__half2_raw>(x).data >= static_cast<__half2_raw>(y).data); - return __builtin_convertvector(-r, _Float16_2); - } - inline - __HOST_DEVICE__ - __half2 __hgtu2(__half2 x, __half2 y) { - auto r = !(static_cast<__half2_raw>(x).data <= static_cast<__half2_raw>(y).data); - return __builtin_convertvector(-r, _Float16_2); - } + // long long -> half + inline + __device__ + __half __ll2half_rn(long long x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + inline + __device__ + __half __ll2half_rz(long long x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + inline + __device__ + __half __ll2half_rd(long long x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + inline + __device__ + __half __ll2half_ru(long long x) + { + return __half_raw{static_cast<_Float16>(x)}; + } - inline - __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 - __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 - __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 - __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 - __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 - __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 - __HOST_DEVICE__ - bool __hbequ2(__half2 x, __half2 y) { return __hbeq2(x, y); } - inline - __HOST_DEVICE__ - bool __hbneu2(__half2 x, __half2 y) { return __hbne2(x, y); } - inline - __HOST_DEVICE__ + // half -> unsigned int + inline + __device__ + unsigned int __half2uint_rn(__half x) + { + return static_cast<__half_raw>(x).data; + } + inline + __device__ + unsigned int __half2uint_rz(__half x) + { + return static_cast<__half_raw>(x).data; + } + inline + __device__ + unsigned int __half2uint_rd(__half x) + { + return static_cast<__half_raw>(x).data; + } + inline + __device__ + unsigned int __half2uint_ru(__half x) + { + return static_cast<__half_raw>(x).data; + } + + // unsigned int -> half + inline + __device__ + __half __uint2half_rn(unsigned int x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + inline + __device__ + __half __uint2half_rz(unsigned int x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + inline + __device__ + __half __uint2half_rd(unsigned int x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + inline + __device__ + __half __uint2half_ru(unsigned int x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + + // half -> unsigned short + inline + __device__ + unsigned short __half2ushort_rn(__half x) + { + return static_cast<__half_raw>(x).data; + } + inline + __device__ + unsigned short __half2ushort_rz(__half x) + { + return static_cast<__half_raw>(x).data; + } + inline + __device__ + unsigned short __half2ushort_rd(__half x) + { + return static_cast<__half_raw>(x).data; + } + inline + __device__ + unsigned short __half2ushort_ru(__half x) + { + return static_cast<__half_raw>(x).data; + } + + // unsigned short -> half + inline + __device__ + __half __ushort2half_rn(unsigned short x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + inline + __device__ + __half __ushort2half_rz(unsigned short x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + inline + __device__ + __half __ushort2half_rd(unsigned short x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + inline + __device__ + __half __ushort2half_ru(unsigned short x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + + // half -> unsigned long long + inline + __device__ + unsigned long long __half2ull_rn(__half x) + { + return static_cast<__half_raw>(x).data; + } + inline + __device__ + unsigned long long __half2ull_rz(__half x) + { + return static_cast<__half_raw>(x).data; + } + inline + __device__ + unsigned long long __half2ull_rd(__half x) + { + return static_cast<__half_raw>(x).data; + } + inline + __device__ + unsigned long long __half2ull_ru(__half x) + { + return static_cast<__half_raw>(x).data; + } + + // unsigned long long -> half + inline + __device__ + __half __ull2half_rn(unsigned long long x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + inline + __device__ + __half __ull2half_rz(unsigned long long x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + inline + __device__ + __half __ull2half_rd(unsigned long long x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + inline + __device__ + __half __ull2half_ru(unsigned long long x) + { + return __half_raw{static_cast<_Float16>(x)}; + } + + // Load primitives + inline + __device__ + __half __ldg(const __half* ptr) { return *ptr; } + inline + __device__ + __half __ldcg(const __half* ptr) { return *ptr; } + inline + __device__ + __half __ldca(const __half* ptr) { return *ptr; } + inline + __device__ + __half __ldcs(const __half* ptr) { return *ptr; } + + inline + __HOST_DEVICE__ + __half2 __ldg(const __half2* ptr) { return *ptr; } + inline + __HOST_DEVICE__ + __half2 __ldcg(const __half2* ptr) { return *ptr; } + inline + __HOST_DEVICE__ + __half2 __ldca(const __half2* ptr) { return *ptr; } + inline + __HOST_DEVICE__ + __half2 __ldcs(const __half2* ptr) { return *ptr; } + + // Relations + inline + __HOST_DEVICE__ + bool __heq(__half x, __half y) + { + return static_cast<__half_raw>(x).data == + static_cast<__half_raw>(y).data; + } + inline + __HOST_DEVICE__ + bool __hne(__half x, __half y) + { + return static_cast<__half_raw>(x).data != + static_cast<__half_raw>(y).data; + } + inline + __HOST_DEVICE__ + bool __hle(__half x, __half y) + { + return static_cast<__half_raw>(x).data <= + static_cast<__half_raw>(y).data; + } + inline + __HOST_DEVICE__ + bool __hge(__half x, __half y) + { + return static_cast<__half_raw>(x).data >= + static_cast<__half_raw>(y).data; + } + inline + __HOST_DEVICE__ + bool __hlt(__half x, __half y) + { + return static_cast<__half_raw>(x).data < + static_cast<__half_raw>(y).data; + } + inline + __HOST_DEVICE__ + bool __hgt(__half x, __half y) + { + return static_cast<__half_raw>(x).data > + static_cast<__half_raw>(y).data; + } + inline __HOST_DEVICE__ + bool __hequ(__half x, __half y) { + return !(static_cast<__half_raw>(x).data < static_cast<__half_raw>(y).data) && + !(static_cast<__half_raw>(x).data > static_cast<__half_raw>(y).data); + } + inline __HOST_DEVICE__ + bool __hneu(__half x, __half y) { + return !(static_cast<__half_raw>(x).data == static_cast<__half_raw>(y).data); + } + inline __HOST_DEVICE__ + bool __hleu(__half x, __half y) { + return !(static_cast<__half_raw>(x).data > static_cast<__half_raw>(y).data); + } + inline + __HOST_DEVICE__ + bool __hgeu(__half x, __half y) { + return !(static_cast<__half_raw>(x).data < static_cast<__half_raw>(y).data); + } + inline + __HOST_DEVICE__ + bool __hltu(__half x, __half y) { + return !(static_cast<__half_raw>(x).data >= static_cast<__half_raw>(y).data); + } + inline + __HOST_DEVICE__ + bool __hgtu(__half x, __half y) { + return !(static_cast<__half_raw>(x).data <= static_cast<__half_raw>(y).data); + } + + inline + __HOST_DEVICE__ + __half2 __heq2(__half2 x, __half2 y) + { + auto r = static_cast<__half2_raw>(x).data == + static_cast<__half2_raw>(y).data; + return __builtin_convertvector(-r, _Float16_2); + } + inline + __HOST_DEVICE__ + __half2 __hne2(__half2 x, __half2 y) + { + auto r = static_cast<__half2_raw>(x).data != + static_cast<__half2_raw>(y).data; + return __builtin_convertvector(-r, _Float16_2); + } + inline + __HOST_DEVICE__ + __half2 __hle2(__half2 x, __half2 y) + { + auto r = static_cast<__half2_raw>(x).data <= + static_cast<__half2_raw>(y).data; + return __builtin_convertvector(-r, _Float16_2); + } + inline + __HOST_DEVICE__ + __half2 __hge2(__half2 x, __half2 y) + { + auto r = static_cast<__half2_raw>(x).data >= + static_cast<__half2_raw>(y).data; + return __builtin_convertvector(-r, _Float16_2); + } + inline + __HOST_DEVICE__ + __half2 __hlt2(__half2 x, __half2 y) + { + auto r = static_cast<__half2_raw>(x).data < + static_cast<__half2_raw>(y).data; + return __builtin_convertvector(-r, _Float16_2); + } + inline + __HOST_DEVICE__ + __half2 __hgt2(__half2 x, __half2 y) + { + auto r = static_cast<__half2_raw>(x).data > + static_cast<__half2_raw>(y).data; + return __builtin_convertvector(-r, _Float16_2); + } + inline __HOST_DEVICE__ + __half2 __hequ2(__half2 x, __half2 y) { + auto r = !(static_cast<__half2_raw>(x).data < static_cast<__half2_raw>(y).data) && + !(static_cast<__half2_raw>(x).data > static_cast<__half2_raw>(y).data); + return __builtin_convertvector(-r, _Float16_2); + } + inline + __HOST_DEVICE__ + __half2 __hneu2(__half2 x, __half2 y) { + auto r = !(static_cast<__half2_raw>(x).data == static_cast<__half2_raw>(y).data); + return __builtin_convertvector(-r, _Float16_2); + } + inline + __HOST_DEVICE__ + __half2 __hleu2(__half2 x, __half2 y) { + auto r = !(static_cast<__half2_raw>(x).data > static_cast<__half2_raw>(y).data); + return __builtin_convertvector(-r, _Float16_2); + } + inline + __HOST_DEVICE__ + __half2 __hgeu2(__half2 x, __half2 y) { + auto r = !(static_cast<__half2_raw>(x).data < static_cast<__half2_raw>(y).data); + return __builtin_convertvector(-r, _Float16_2); + } + inline + __HOST_DEVICE__ + __half2 __hltu2(__half2 x, __half2 y) { + auto r = !(static_cast<__half2_raw>(x).data >= static_cast<__half2_raw>(y).data); + return __builtin_convertvector(-r, _Float16_2); + } + inline + __HOST_DEVICE__ + __half2 __hgtu2(__half2 x, __half2 y) { + auto r = !(static_cast<__half2_raw>(x).data <= static_cast<__half2_raw>(y).data); + return __builtin_convertvector(-r, _Float16_2); + } + + inline + __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 + __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 + __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 + __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 + __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 + __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 + __HOST_DEVICE__ + bool __hbequ2(__half2 x, __half2 y) { return __hbeq2(x, y); } + inline + __HOST_DEVICE__ + bool __hbneu2(__half2 x, __half2 y) { return __hbne2(x, y); } + inline + __HOST_DEVICE__ bool __hbleu2(__half2 x, __half2 y) { return __hble2(x, y); } inline __HOST_DEVICE__ @@ -1801,7 +1799,6 @@ THE SOFTWARE. { return __half2{-static_cast<__half2_raw>(x).data}; } - } // Anonymous namespace. #if !defined(HIP_NO_HALF) using half = __half;