AI/rocm-systems
2
0
Çatalla 0
Kod Konular Değişiklik İstekleri İşlemler Paketler Projeler Sürüm Viki Aktivite

SWDEV-395559 - Move unsigned short up.

Kaynağa Gözat 
Change-Id: I3373a7bc22e632dc10b94b93ff023444d456b2b7


[ROCm/clr commit: cbd91bdc19]
Bu işleme şunda yer alıyor:
Jaydeep Patel
2023-04-27 09:04:58 +00:00
işlemeyi yapan: Jaydeepkumar Patel
ebeveyn 44512e3228
işleme fa88970f6a
1 değiştirilmiş dosya ile 32 ekleme ve 28 silme
Yama Dosyasını İndir Diff Dosyasını İndir Tüm dosyaları genişlet Tüm dosyaları daralt
projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h
+32 -28
Dosyayı Görüntüle
@@ -66,11 +66,11 @@ THE SOFTWARE.
union {
static_assert(sizeof(_Float16_2) == sizeof(unsigned short[2]), "");
_Float16_2 data;
struct {
unsigned short x;
unsigned short y;
};
_Float16_2 data;
};
};
@@ -353,11 +353,11 @@ THE SOFTWARE.
static_assert(
sizeof(_Float16_2) == sizeof(unsigned short[2]), "");
_Float16_2 data;
struct {
unsigned short x;
unsigned short y;
};
_Float16_2 data;
};
// CREATORS
@@ -443,7 +443,11 @@ THE SOFTWARE.
__HOST_DEVICE__
operator decltype(data)() const { return data; }
__HOST_DEVICE__
operator __half2_raw() const { return __half2_raw{data}; }
operator __half2_raw() const {
__half2_raw r;
r.data = data;
return r;
}
// ACCESSORS - DEVICE ONLY
#if !defined(__HIP_NO_HALF_OPERATORS__)
@@ -587,7 +591,7 @@ THE SOFTWARE.
__HOST_DEVICE__
__half2 __high2half2(__half2 x)
{
return __half2_raw{
return __half2{
_Float16_2{
static_cast<__half2_raw>(x).data.y,
static_cast<__half2_raw>(x).data.y}};
@@ -597,7 +601,7 @@ THE SOFTWARE.
__HOST_DEVICE__
__half2 __lows2half2(__half2 x, __half2 y)
{
return __half2_raw{
return __half2{
_Float16_2{
static_cast<__half2_raw>(x).data.x,
static_cast<__half2_raw>(y).data.x}};
@@ -607,7 +611,7 @@ THE SOFTWARE.
__HOST_DEVICE__
__half2 __highs2half2(__half2 x, __half2 y)
{
return __half2_raw{
return __half2{
_Float16_2{
static_cast<__half2_raw>(x).data.y,
static_cast<__half2_raw>(y).data.y}};
@@ -617,7 +621,7 @@ THE SOFTWARE.
__HOST_DEVICE__
__half2 __lowhigh2highlow(__half2 x)
{
return __half2_raw{
return __half2{
_Float16_2{
static_cast<__half2_raw>(x).data.y,
static_cast<__half2_raw>(x).data.x}};
@@ -710,7 +714,7 @@ THE SOFTWARE.
__HOST_DEVICE__
__half2 __float2half2_rn(float x)
{
return __half2_raw{
return __half2{
_Float16_2{
static_cast<_Float16>(x), static_cast<_Float16>(x)}};
}
@@ -718,7 +722,7 @@ THE SOFTWARE.
__HOST_DEVICE__
__half2 __floats2half2_rn(float x, float y)
{
return __half2_raw{_Float16_2{
return __half2{_Float16_2{
static_cast<_Float16>(x), static_cast<_Float16>(y)}};
}
inline
@@ -1406,7 +1410,7 @@ THE SOFTWARE.
__HOST_DEVICE__
__half2 __hadd2(__half2 x, __half2 y)
{
return __half2_raw{
return __half2{
static_cast<__half2_raw>(x).data +
static_cast<__half2_raw>(y).data};
}
@@ -1414,14 +1418,14 @@ THE SOFTWARE.
__HOST_DEVICE__
__half2 __habs2(__half2 x)
{
return __half2_raw{
return __half2{
__ocml_fabs_2f16(static_cast<__half2_raw>(x).data)};
}
inline
__HOST_DEVICE__
__half2 __hsub2(__half2 x, __half2 y)
{
return __half2_raw{
return __half2{
static_cast<__half2_raw>(x).data -
static_cast<__half2_raw>(y).data};
}
@@ -1429,7 +1433,7 @@ THE SOFTWARE.
__HOST_DEVICE__
__half2 __hmul2(__half2 x, __half2 y)
{
return __half2_raw{
return __half2{
static_cast<__half2_raw>(x).data *
static_cast<__half2_raw>(y).data};
}
@@ -1464,7 +1468,7 @@ THE SOFTWARE.
__HOST_DEVICE__
__half2 __hfma2(__half2 x, __half2 y, __half2 z)
{
return __half2_raw{__ocml_fma_2f16(x, y, z)};
return __half2{__ocml_fma_2f16(x, y, z)};
}
inline
__HOST_DEVICE__
@@ -1479,7 +1483,7 @@ THE SOFTWARE.
__HOST_DEVICE__
__half2 __h2div(__half2 x, __half2 y)
{
return __half2_raw{
return __half2{
static_cast<__half2_raw>(x).data /
static_cast<__half2_raw>(y).data};
}
@@ -1622,61 +1626,61 @@ THE SOFTWARE.
__HOST_DEVICE__
__half2 h2trunc(__half2 x)
{
return __half2_raw{__ocml_trunc_2f16(x)};
return __half2{__ocml_trunc_2f16(x)};
}
inline
__HOST_DEVICE__
__half2 h2ceil(__half2 x)
{
return __half2_raw{__ocml_ceil_2f16(x)};
return __half2{__ocml_ceil_2f16(x)};
}
inline
__HOST_DEVICE__
__half2 h2floor(__half2 x)
{
return __half2_raw{__ocml_floor_2f16(x)};
return __half2{__ocml_floor_2f16(x)};
}
inline
__HOST_DEVICE__
__half2 h2rint(__half2 x)
{
return __half2_raw{__ocml_rint_2f16(x)};
return __half2{__ocml_rint_2f16(x)};
}
inline
__HOST_DEVICE__
__half2 h2sin(__half2 x)
{
return __half2_raw{__ocml_sin_2f16(x)};
return __half2{__ocml_sin_2f16(x)};
}
inline
__HOST_DEVICE__
__half2 h2cos(__half2 x)
{
return __half2_raw{__ocml_cos_2f16(x)};
return __half2{__ocml_cos_2f16(x)};
}
inline
__HOST_DEVICE__
__half2 h2exp(__half2 x)
{
return __half2_raw{__ocml_exp_2f16(x)};
return __half2{__ocml_exp_2f16(x)};
}
inline
__HOST_DEVICE__
__half2 h2exp2(__half2 x)
{
return __half2_raw{__ocml_exp2_2f16(x)};
return __half2{__ocml_exp2_2f16(x)};
}
inline
__HOST_DEVICE__
__half2 h2exp10(__half2 x)
{
return __half2_raw{__ocml_exp10_2f16(x)};
return __half2{__ocml_exp10_2f16(x)};
}
inline
__HOST_DEVICE__
__half2 h2log2(__half2 x)
{
return __half2_raw{__ocml_log2_2f16(x)};
return __half2{__ocml_log2_2f16(x)};
}
inline
__HOST_DEVICE__
@@ -1701,7 +1705,7 @@ THE SOFTWARE.
__half2 __hisinf2(__half2 x)
{
auto r = __ocml_isinf_2f16(x);
return __half2_raw{_Float16_2{
return __half2{_Float16_2{
static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
}
inline
@@ -1709,14 +1713,14 @@ THE SOFTWARE.
__half2 __hisnan2(__half2 x)
{
auto r = __ocml_isnan_2f16(x);
return __half2_raw{_Float16_2{
return __half2{_Float16_2{
static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
}
inline
__HOST_DEVICE__
__half2 __hneg2(__half2 x)
{
return __half2_raw{-static_cast<__half2_raw>(x).data};
return __half2{-static_cast<__half2_raw>(x).data};
}
} // Anonymous namespace.