SWDEV-501779 - add correct function qualifier for fp16 functions
Some functions were __device__ only, but should be __host__ and
__device__, changed them to __HOST_DEVICE__.
Some functions were __HOST_DEVICE__ but were using ockl functions,
changed them to __device__ only.
Change-Id: Ife9e7abe60415bda68f5f9a101e6e7c39ad51064
[ROCm/clr commit: 5122b8c999]
This commit is contained in:
Jatin Chaudhary
committed by
Jatin Jaikishan Chaudhary
Jatin Jaikishan Chaudhary
parent
cf25aaddcf
commit
163f810941
@@ -41,6 +41,14 @@ THE SOFTWARE.
|
||||
#endif
|
||||
#endif // !defined(__HIPCC_RTC__)
|
||||
|
||||
#define HIPRT_INF_FP16 __ushort_as_half((unsigned short)0x7C00U)
|
||||
#define HIPRT_MAX_NORMAL_FP16 __ushort_as_half((unsigned short)0x7BFFU)
|
||||
#define HIPRT_MIN_DENORM_FP16 __ushort_as_half((unsigned short)0x0001U)
|
||||
#define HIPRT_NAN_FP16 __ushort_as_half((unsigned short)0x7FFFU)
|
||||
#define HIPRT_NEG_ZERO_FP16 __ushort_as_half((unsigned short)0x8000U)
|
||||
#define HIPRT_ONE_FP16 __ushort_as_half((unsigned short)0x3C00U)
|
||||
#define HIPRT_ZERO_FP16 __ushort_as_half((unsigned short)0x0000U)
|
||||
|
||||
#if defined(__clang__) && defined(__HIP__)
|
||||
typedef _Float16 _Float16_2 __attribute__((ext_vector_type(2)));
|
||||
|
||||
@@ -182,42 +190,42 @@ THE SOFTWARE.
|
||||
#endif
|
||||
|
||||
#if !defined(__HIP_NO_HALF_OPERATORS__)
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half& operator+=(const __half& x)
|
||||
{
|
||||
data += x.data;
|
||||
return *this;
|
||||
}
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half& operator-=(const __half& x)
|
||||
{
|
||||
data -= x.data;
|
||||
return *this;
|
||||
}
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half& operator*=(const __half& x)
|
||||
{
|
||||
data *= x.data;
|
||||
return *this;
|
||||
}
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half& operator/=(const __half& x)
|
||||
{
|
||||
data /= x.data;
|
||||
return *this;
|
||||
}
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half& operator++() { ++data; return *this; }
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half operator++(int)
|
||||
{
|
||||
__half tmp{*this};
|
||||
++*this;
|
||||
return tmp;
|
||||
}
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half& operator--() { --data; return *this; }
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half operator--(int)
|
||||
{
|
||||
__half tmp{*this};
|
||||
@@ -250,9 +258,9 @@ THE SOFTWARE.
|
||||
#endif
|
||||
|
||||
#if !defined(__HIP_NO_HALF_OPERATORS__)
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half operator+() const { return *this; }
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half operator-() const
|
||||
{
|
||||
__half tmp{*this};
|
||||
@@ -265,70 +273,70 @@ THE SOFTWARE.
|
||||
#if !defined(__HIP_NO_HALF_OPERATORS__)
|
||||
friend
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half operator+(const __half& x, const __half& y)
|
||||
{
|
||||
return __half{x} += y;
|
||||
}
|
||||
friend
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half operator-(const __half& x, const __half& y)
|
||||
{
|
||||
return __half{x} -= y;
|
||||
}
|
||||
friend
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half operator*(const __half& x, const __half& y)
|
||||
{
|
||||
return __half{x} *= y;
|
||||
}
|
||||
friend
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half operator/(const __half& x, const __half& y)
|
||||
{
|
||||
return __half{x} /= y;
|
||||
}
|
||||
friend
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool operator==(const __half& x, const __half& y)
|
||||
{
|
||||
return x.data == y.data;
|
||||
}
|
||||
friend
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool operator!=(const __half& x, const __half& y)
|
||||
{
|
||||
return !(x == y);
|
||||
}
|
||||
friend
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool operator<(const __half& x, const __half& y)
|
||||
{
|
||||
return x.data < y.data;
|
||||
}
|
||||
friend
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool operator>(const __half& x, const __half& y)
|
||||
{
|
||||
return y.data < x.data;
|
||||
}
|
||||
friend
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool operator<=(const __half& x, const __half& y)
|
||||
{
|
||||
return !(y < x);
|
||||
}
|
||||
friend
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool operator>=(const __half& x, const __half& y)
|
||||
{
|
||||
return !(x < y);
|
||||
@@ -385,42 +393,42 @@ THE SOFTWARE.
|
||||
|
||||
// MANIPULATORS - DEVICE ONLY
|
||||
#if !defined(__HIP_NO_HALF_OPERATORS__)
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half2& operator+=(const __half2& xx)
|
||||
{
|
||||
data += xx.data;
|
||||
return *this;
|
||||
}
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half2& operator-=(const __half2& xx)
|
||||
{
|
||||
data -= xx.data;
|
||||
return *this;
|
||||
}
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half2& operator*=(const __half2& xx)
|
||||
{
|
||||
data *= xx.data;
|
||||
return *this;
|
||||
}
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half2& operator/=(const __half2& xx)
|
||||
{
|
||||
data /= xx.data;
|
||||
return *this;
|
||||
}
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half2& operator++() { return *this += _Float16_2{1, 1}; }
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half2 operator++(int)
|
||||
{
|
||||
__half2 tmp{*this};
|
||||
++*this;
|
||||
return tmp;
|
||||
}
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half2& operator--() { return *this -= _Float16_2{1, 1}; }
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half2 operator--(int)
|
||||
{
|
||||
__half2 tmp{*this};
|
||||
@@ -441,9 +449,9 @@ THE SOFTWARE.
|
||||
|
||||
// ACCESSORS - DEVICE ONLY
|
||||
#if !defined(__HIP_NO_HALF_OPERATORS__)
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half2 operator+() const { return *this; }
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half2 operator-() const
|
||||
{
|
||||
__half2 tmp{*this};
|
||||
@@ -456,35 +464,35 @@ THE SOFTWARE.
|
||||
#if !defined(__HIP_NO_HALF_OPERATORS__)
|
||||
friend
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half2 operator+(const __half2& xx, const __half2& yy)
|
||||
{
|
||||
return __half2{xx} += yy;
|
||||
}
|
||||
friend
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half2 operator-(const __half2& xx, const __half2& yy)
|
||||
{
|
||||
return __half2{xx} -= yy;
|
||||
}
|
||||
friend
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half2 operator*(const __half2& xx, const __half2& yy)
|
||||
{
|
||||
return __half2{xx} *= yy;
|
||||
}
|
||||
friend
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half2 operator/(const __half2& xx, const __half2& yy)
|
||||
{
|
||||
return __half2{xx} /= yy;
|
||||
}
|
||||
friend
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool operator==(const __half2& xx, const __half2& yy)
|
||||
{
|
||||
auto r = xx.data == yy.data;
|
||||
@@ -492,14 +500,14 @@ THE SOFTWARE.
|
||||
}
|
||||
friend
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool operator!=(const __half2& xx, const __half2& yy)
|
||||
{
|
||||
return !(xx == yy);
|
||||
}
|
||||
friend
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool operator<(const __half2& xx, const __half2& yy)
|
||||
{
|
||||
auto r = xx.data < yy.data;
|
||||
@@ -507,21 +515,21 @@ THE SOFTWARE.
|
||||
}
|
||||
friend
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool operator>(const __half2& xx, const __half2& yy)
|
||||
{
|
||||
return yy < xx;
|
||||
}
|
||||
friend
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool operator<=(const __half2& xx, const __half2& yy)
|
||||
{
|
||||
return !(yy < xx);
|
||||
}
|
||||
friend
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool operator>=(const __half2& xx, const __half2& yy)
|
||||
{
|
||||
return !(xx < yy);
|
||||
@@ -619,21 +627,21 @@ THE SOFTWARE.
|
||||
|
||||
// Bitcasts
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
short __half_as_short(__half x)
|
||||
{
|
||||
return static_cast<__half_raw>(x).x;
|
||||
}
|
||||
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
unsigned short __half_as_ushort(__half x)
|
||||
{
|
||||
return static_cast<__half_raw>(x).x;
|
||||
}
|
||||
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half __short_as_half(short x)
|
||||
{
|
||||
__half_raw r; r.x = x;
|
||||
@@ -641,7 +649,7 @@ THE SOFTWARE.
|
||||
}
|
||||
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half __ushort_as_half(unsigned short x)
|
||||
{
|
||||
__half_raw r; r.x = x;
|
||||
@@ -1091,72 +1099,72 @@ THE SOFTWARE.
|
||||
|
||||
// Relations
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool __heq(__half x, __half y)
|
||||
{
|
||||
return static_cast<__half_raw>(x).data ==
|
||||
static_cast<__half_raw>(y).data;
|
||||
}
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool __hne(__half x, __half y)
|
||||
{
|
||||
return static_cast<__half_raw>(x).data !=
|
||||
static_cast<__half_raw>(y).data;
|
||||
}
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool __hle(__half x, __half y)
|
||||
{
|
||||
return static_cast<__half_raw>(x).data <=
|
||||
static_cast<__half_raw>(y).data;
|
||||
}
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool __hge(__half x, __half y)
|
||||
{
|
||||
return static_cast<__half_raw>(x).data >=
|
||||
static_cast<__half_raw>(y).data;
|
||||
}
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool __hlt(__half x, __half y)
|
||||
{
|
||||
return static_cast<__half_raw>(x).data <
|
||||
static_cast<__half_raw>(y).data;
|
||||
}
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool __hgt(__half x, __half y)
|
||||
{
|
||||
return static_cast<__half_raw>(x).data >
|
||||
static_cast<__half_raw>(y).data;
|
||||
}
|
||||
inline __device__
|
||||
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 __device__
|
||||
inline __HOST_DEVICE__
|
||||
bool __hneu(__half x, __half y) {
|
||||
return !(static_cast<__half_raw>(x).data == static_cast<__half_raw>(y).data);
|
||||
}
|
||||
inline __device__
|
||||
inline __HOST_DEVICE__
|
||||
bool __hleu(__half x, __half y) {
|
||||
return !(static_cast<__half_raw>(x).data > static_cast<__half_raw>(y).data);
|
||||
}
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool __hgeu(__half x, __half y) {
|
||||
return !(static_cast<__half_raw>(x).data < static_cast<__half_raw>(y).data);
|
||||
}
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool __hltu(__half x, __half y) {
|
||||
return !(static_cast<__half_raw>(x).data >= static_cast<__half_raw>(y).data);
|
||||
}
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool __hgtu(__half x, __half y) {
|
||||
return !(static_cast<__half_raw>(x).data <= static_cast<__half_raw>(y).data);
|
||||
}
|
||||
@@ -1307,36 +1315,45 @@ THE SOFTWARE.
|
||||
__HOST_DEVICE__
|
||||
bool __hbgtu2(__half2 x, __half2 y) { return __hbgt2(x, y); }
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool __hisnan(__half x)
|
||||
{
|
||||
__half_raw hr = x;
|
||||
return (hr.x & 0x7FFFU) > 0x7C00u;
|
||||
}
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__half __hmax(const __half x, const __half y) {
|
||||
return __half_raw{__ocml_fmax_f16(static_cast<__half_raw>(x).data,
|
||||
static_cast<__half_raw>(y).data)};
|
||||
if (__hisnan(x) && !__hisnan(y)) return y;
|
||||
if (!__hisnan(x) && __hisnan(y)) return x;
|
||||
if (__hisnan(x) && __hisnan(y)) return HIPRT_NAN_FP16;
|
||||
if (static_cast<__half_raw>(x).data > static_cast<__half_raw>(y).data)
|
||||
return __half_raw{static_cast<__half_raw>(x).data};
|
||||
return __half_raw{static_cast<__half_raw>(y).data};
|
||||
}
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half __hmax_nan(const __half x, const __half y) {
|
||||
if(__ocml_isnan_f16(static_cast<__half_raw>(x).data)) {
|
||||
return x;
|
||||
} else if (__ocml_isnan_f16(static_cast<__half_raw>(y).data)) {
|
||||
return y;
|
||||
}
|
||||
return __hmax(x, y);
|
||||
if (__hisnan(x)) return x;
|
||||
if (__hisnan(y)) return y;
|
||||
return __hmax(x, y);
|
||||
}
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half __hmin(const __half x, const __half y) {
|
||||
return __half_raw{__ocml_fmin_f16(static_cast<__half_raw>(x).data,
|
||||
static_cast<__half_raw>(y).data)};
|
||||
if (__hisnan(x) && !__hisnan(y)) return y;
|
||||
if (!__hisnan(x) && __hisnan(y)) return x;
|
||||
if (__hisnan(x) && __hisnan(y)) return HIPRT_NAN_FP16;
|
||||
if (static_cast<__half_raw>(x).data > static_cast<__half_raw>(y).data)
|
||||
return __half_raw{static_cast<__half_raw>(y).data};
|
||||
return __half_raw{static_cast<__half_raw>(x).data};
|
||||
}
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half __hmin_nan(const __half x, const __half y) {
|
||||
if(__ocml_isnan_f16(static_cast<__half_raw>(x).data)) {
|
||||
return x;
|
||||
} else if (__ocml_isnan_f16(static_cast<__half_raw>(y).data)) {
|
||||
return y;
|
||||
}
|
||||
return __hmin(x, y);
|
||||
if (__hisnan(x)) return x;
|
||||
if (__hisnan(y)) return y;
|
||||
return __hmin(x, y);
|
||||
}
|
||||
|
||||
// Arithmetic
|
||||
@@ -1352,22 +1369,27 @@ THE SOFTWARE.
|
||||
}
|
||||
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half __hadd(__half x, __half y)
|
||||
{
|
||||
return __half_raw{
|
||||
static_cast<__half_raw>(x).data +
|
||||
static_cast<__half_raw>(y).data};
|
||||
}
|
||||
inline
|
||||
__device__
|
||||
__half __habs(__half x)
|
||||
{
|
||||
return __half_raw{
|
||||
__ocml_fabs_f16(static_cast<__half_raw>(x).data)};
|
||||
}
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__half __habs(__half x)
|
||||
{
|
||||
static_assert(sizeof(_Float16) == sizeof(unsigned short));
|
||||
union {
|
||||
_Float16 fp16;
|
||||
unsigned short us;
|
||||
} u{static_cast<__half_raw>(x).data};
|
||||
u.us &= 0x7FFFu;
|
||||
return __half_raw{u.fp16};
|
||||
}
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half __hsub(__half x, __half y)
|
||||
{
|
||||
return __half_raw{
|
||||
@@ -1375,7 +1397,7 @@ THE SOFTWARE.
|
||||
static_cast<__half_raw>(y).data};
|
||||
}
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half __hmul(__half x, __half y)
|
||||
{
|
||||
return __half_raw{
|
||||
@@ -1383,19 +1405,19 @@ THE SOFTWARE.
|
||||
static_cast<__half_raw>(y).data};
|
||||
}
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half __hadd_sat(__half x, __half y)
|
||||
{
|
||||
return __clamp_01(__hadd(x, y));
|
||||
}
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half __hsub_sat(__half x, __half y)
|
||||
{
|
||||
return __clamp_01(__hsub(x, y));
|
||||
}
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half __hmul_sat(__half x, __half y)
|
||||
{
|
||||
return __clamp_01(__hmul(x, y));
|
||||
@@ -1416,7 +1438,7 @@ THE SOFTWARE.
|
||||
return __clamp_01(__hfma(x, y, z));
|
||||
}
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half __hdiv(__half x, __half y)
|
||||
{
|
||||
return __half_raw{
|
||||
@@ -1432,13 +1454,12 @@ THE SOFTWARE.
|
||||
static_cast<__half2_raw>(x).data +
|
||||
static_cast<__half2_raw>(y).data};
|
||||
}
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__half2 __habs2(__half2 x)
|
||||
{
|
||||
return __half2{
|
||||
__ocml_fabs_2f16(static_cast<__half2_raw>(x).data)};
|
||||
}
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__half2 __habs2(__half2 x)
|
||||
{
|
||||
return __half2{__habs(x.x), __habs(x.y)};
|
||||
}
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__half2 __hsub2(__half2 x, __half2 y)
|
||||
@@ -1483,13 +1504,13 @@ THE SOFTWARE.
|
||||
__clamp_01(__half_raw{r.data.y})};
|
||||
}
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__device__
|
||||
__half2 __hfma2(__half2 x, __half2 y, __half2 z)
|
||||
{
|
||||
return __half2{__ocml_fma_2f16(x, y, z)};
|
||||
}
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__device__
|
||||
__half2 __hfma2_sat(__half2 x, __half2 y, __half2 z)
|
||||
{
|
||||
auto r = static_cast<__half2_raw>(__hfma2(x, y, z));
|
||||
@@ -1655,104 +1676,99 @@ THE SOFTWARE.
|
||||
__ocml_sqrt_f16(static_cast<__half_raw>(x).data)};
|
||||
}
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
bool __hisinf(__half x)
|
||||
{
|
||||
return __ocml_isinf_f16(static_cast<__half_raw>(x).data);
|
||||
// +Inf/-Inf
|
||||
return x == HIPRT_INF_FP16 || x == __ushort_as_half((unsigned short)0xFC00U);
|
||||
}
|
||||
inline
|
||||
__device__
|
||||
bool __hisnan(__half x)
|
||||
{
|
||||
return __ocml_isnan_f16(static_cast<__half_raw>(x).data);
|
||||
}
|
||||
inline
|
||||
__device__
|
||||
__HOST_DEVICE__
|
||||
__half __hneg(__half x)
|
||||
{
|
||||
return __half_raw{-static_cast<__half_raw>(x).data};
|
||||
}
|
||||
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__device__
|
||||
__half2 h2trunc(__half2 x)
|
||||
{
|
||||
return __half2{__ocml_trunc_2f16(x)};
|
||||
}
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__device__
|
||||
__half2 h2ceil(__half2 x)
|
||||
{
|
||||
return __half2{__ocml_ceil_2f16(x)};
|
||||
}
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__device__
|
||||
__half2 h2floor(__half2 x)
|
||||
{
|
||||
return __half2{__ocml_floor_2f16(x)};
|
||||
}
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__device__
|
||||
__half2 h2rint(__half2 x)
|
||||
{
|
||||
return __half2{__ocml_rint_2f16(x)};
|
||||
}
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__device__
|
||||
__half2 h2sin(__half2 x)
|
||||
{
|
||||
return __half2{__ocml_sin_2f16(x)};
|
||||
}
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__device__
|
||||
__half2 h2cos(__half2 x)
|
||||
{
|
||||
return __half2{__ocml_cos_2f16(x)};
|
||||
}
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__device__
|
||||
__half2 h2exp(__half2 x)
|
||||
{
|
||||
return __half2{__ocml_exp_2f16(x)};
|
||||
}
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__device__
|
||||
__half2 h2exp2(__half2 x)
|
||||
{
|
||||
return __half2{__ocml_exp2_2f16(x)};
|
||||
}
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__device__
|
||||
__half2 h2exp10(__half2 x)
|
||||
{
|
||||
return __half2{__ocml_exp10_2f16(x)};
|
||||
}
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__device__
|
||||
__half2 h2log2(__half2 x)
|
||||
{
|
||||
return __half2{__ocml_log2_2f16(x)};
|
||||
}
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__device__
|
||||
__half2 h2log(__half2 x) { return __ocml_log_2f16(x); }
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__device__
|
||||
__half2 h2log10(__half2 x) { return __ocml_log10_2f16(x); }
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__device__
|
||||
__half2 h2rcp(__half2 x) {
|
||||
return _Float16_2{
|
||||
_Float16_2{static_cast<_Float16>(1.0f), static_cast<_Float16>(1.0f)} / x.data};
|
||||
}
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__device__
|
||||
__half2 h2rsqrt(__half2 x) { return __ocml_rsqrt_2f16(x); }
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__device__
|
||||
__half2 h2sqrt(__half2 x) { return __ocml_sqrt_2f16(x); }
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
__device__
|
||||
__half2 __hisinf2(__half2 x)
|
||||
{
|
||||
auto r = __ocml_isinf_2f16(x);
|
||||
@@ -1763,9 +1779,8 @@ THE SOFTWARE.
|
||||
__HOST_DEVICE__
|
||||
__half2 __hisnan2(__half2 x)
|
||||
{
|
||||
auto r = __ocml_isnan_2f16(x);
|
||||
return __half2{_Float16_2{
|
||||
static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
|
||||
return __half2{_Float16_2{static_cast<_Float16>(__hisnan(x.x) ? 1.0f : 0.0f),
|
||||
static_cast<_Float16>(__hisnan(x.y) ? 1.0f : 0.0f)}};
|
||||
}
|
||||
inline
|
||||
__HOST_DEVICE__
|
||||
|
||||
Reference in New Issue
Block a user