SWDEV-236670 Address Eigen unit test failure by adding __host__ attribute to half2 functions
Change-Id: Ifdc852c30a1b3704871e0ee58cb7a55d3d37fc6e
[ROCm/clr commit: bb785840b9]
Этот коммит содержится в:
Dittakavi Satyanvesh
коммит произвёл
Satyanvesh Dittakavi
Satyanvesh Dittakavi
родитель
4206cfddc5
Коммит
a0fbc8fd08
@@ -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};
|
||||
|
||||
Ссылка в новой задаче
Block a user