Merge pull request #466 from ROCm-Developer-Tools/feature_use_Float16
Feature use _Float16 and match CUDA __half behaviour.
Cette révision appartient à :
Maneesh Gupta
GitHub
+1612
-415
Fichier diff supprimé car celui-ci est trop grand
Voir la Diff
@@ -0,0 +1,257 @@
|
||||
#pragma once
|
||||
|
||||
#if defined(__cplusplus)
|
||||
#include <cstring>
|
||||
#endif
|
||||
|
||||
struct __half_raw {
|
||||
unsigned short x;
|
||||
};
|
||||
|
||||
struct __half2_raw {
|
||||
unsigned short x;
|
||||
unsigned short y;
|
||||
};
|
||||
|
||||
#if defined(__cplusplus)
|
||||
struct __half;
|
||||
|
||||
__half __float2half(float);
|
||||
float __half2float(__half);
|
||||
|
||||
// BEGIN STRUCT __HALF
|
||||
struct __half {
|
||||
protected:
|
||||
unsigned short __x;
|
||||
public:
|
||||
// CREATORS
|
||||
__half() = default;
|
||||
__half(const __half_raw& x) : __x{x.x} {}
|
||||
#if !defined(__HIP_NO_HALF_CONVERSIONS__)
|
||||
__half(float x) : __x{__float2half(x).__x} {}
|
||||
__half(double x) : __x{__float2half(x).__x} {}
|
||||
#endif
|
||||
__half(const __half&) = default;
|
||||
__half(__half&&) = default;
|
||||
~__half() = default;
|
||||
|
||||
// MANIPULATORS
|
||||
__half& operator=(const __half&) = default;
|
||||
__half& operator=(__half&&) = default;
|
||||
__half& operator=(const __half_raw& x) { __x = x.x; return *this; }
|
||||
#if !defined(__HIP_NO_HALF_CONVERSIONS__)
|
||||
__half& operator=(float x)
|
||||
{
|
||||
__x = __float2half(x).__x;
|
||||
return *this;
|
||||
}
|
||||
__half& operator=(double x)
|
||||
{
|
||||
return *this = static_cast<float>(x);
|
||||
}
|
||||
#endif
|
||||
|
||||
// ACCESSORS
|
||||
operator float() const { return __half2float(*this); }
|
||||
operator __half_raw() const { return __half_raw{__x}; }
|
||||
};
|
||||
// END STRUCT __HALF
|
||||
|
||||
// BEGIN STRUCT __HALF2
|
||||
struct __half2 {
|
||||
protected:
|
||||
__half x;
|
||||
__half y;
|
||||
public:
|
||||
// CREATORS
|
||||
__half2() = default;
|
||||
__half2(const __half2_raw& ix)
|
||||
:
|
||||
x{reinterpret_cast<const __half&>(ix.x)},
|
||||
y{reinterpret_cast<const __half&>(ix.y)}
|
||||
{}
|
||||
__half2(const __half& ix, const __half& iy) : x{ix}, y{iy} {}
|
||||
__half2(const __half2&) = default;
|
||||
__half2(__half2&&) = default;
|
||||
~__half2() = default;
|
||||
|
||||
// MANIPULATORS
|
||||
__half2& operator=(const __half2&) = default;
|
||||
__half2& operator=(__half2&&) = default;
|
||||
__half2& operator=(const __half2_raw& ix)
|
||||
{
|
||||
x = reinterpret_cast<const __half_raw&>(ix.x);
|
||||
y = reinterpret_cast<const __half_raw&>(ix.y);
|
||||
return *this;
|
||||
}
|
||||
|
||||
// ACCESSORS
|
||||
operator __half2_raw() const
|
||||
{
|
||||
return __half2_raw{
|
||||
reinterpret_cast<const unsigned short&>(x),
|
||||
reinterpret_cast<const unsigned short&>(y)};
|
||||
}
|
||||
};
|
||||
// END STRUCT __HALF2
|
||||
|
||||
namespace
|
||||
{
|
||||
inline
|
||||
unsigned short __internal_float2half(
|
||||
float flt, unsigned int& sgn, unsigned int& rem)
|
||||
{
|
||||
unsigned int x{};
|
||||
std::memcpy(&x, &flt, sizeof(flt));
|
||||
|
||||
unsigned int u = (x & 0x7fffffffU);
|
||||
sgn = ((x >> 16) & 0x8000U);
|
||||
|
||||
// NaN/+Inf/-Inf
|
||||
if (u >= 0x7f800000U) {
|
||||
rem = 0;
|
||||
return static_cast<unsigned short>(
|
||||
(u == 0x7f800000U) ? (sgn | 0x7c00U) : 0x7fffU);
|
||||
}
|
||||
// Overflows
|
||||
if (u > 0x477fefffU) {
|
||||
rem = 0x80000000U;
|
||||
return static_cast<unsigned short>(sgn | 0x7bffU);
|
||||
}
|
||||
// Normal numbers
|
||||
if (u >= 0x38800000U) {
|
||||
rem = u << 19;
|
||||
u -= 0x38000000U;
|
||||
return static_cast<unsigned short>(sgn | (u >> 13));
|
||||
}
|
||||
// +0/-0
|
||||
if (u < 0x33000001U) {
|
||||
rem = u;
|
||||
return static_cast<unsigned short>(sgn);
|
||||
}
|
||||
// Denormal numbers
|
||||
unsigned int exponent = u >> 23;
|
||||
unsigned int mantissa = (u & 0x7fffffU);
|
||||
unsigned int shift = 0x7eU - exponent;
|
||||
mantissa |= 0x800000U;
|
||||
rem = mantissa << (32 - shift);
|
||||
return static_cast<unsigned short>(sgn | (mantissa >> shift));
|
||||
}
|
||||
|
||||
inline
|
||||
__half __float2half(float x)
|
||||
{
|
||||
__half_raw r;
|
||||
unsigned int sgn{};
|
||||
unsigned int rem{};
|
||||
r.x = __internal_float2half(x, sgn, rem);
|
||||
if (rem > 0x80000000U || (rem == 0x80000000U && (r.x & 0x1))) ++r.x;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
inline
|
||||
__half __float2half_rn(float x) { return __float2half(x); }
|
||||
|
||||
inline
|
||||
__half __float2half_rz(float x)
|
||||
{
|
||||
__half_raw r;
|
||||
unsigned int sgn{};
|
||||
unsigned int rem{};
|
||||
r.x = __internal_float2half(x, sgn, rem);
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
inline
|
||||
__half __float2half_rd(float x)
|
||||
{
|
||||
__half_raw r;
|
||||
unsigned int sgn{};
|
||||
unsigned int rem{};
|
||||
r.x = __internal_float2half(x, sgn, rem);
|
||||
if (rem && sgn) ++r.x;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
inline
|
||||
__half __float2half_ru(float x)
|
||||
{
|
||||
__half_raw r;
|
||||
unsigned int sgn{};
|
||||
unsigned int rem{};
|
||||
r.x = __internal_float2half(x, sgn, rem);
|
||||
if (rem && !sgn) ++r.x;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
inline
|
||||
__half2 __float2half2_rn(float x)
|
||||
{
|
||||
return __half2{__float2half_rn(x), __float2half_rn(x)};
|
||||
}
|
||||
|
||||
inline
|
||||
__half2 __floats2half2_rn(float x, float y)
|
||||
{
|
||||
return __half2{__float2half_rn(x), __float2half_rn(y)};
|
||||
}
|
||||
|
||||
inline
|
||||
float __internal_half2float(unsigned short x)
|
||||
{
|
||||
unsigned int sign = ((x >> 15) & 1);
|
||||
unsigned int exponent = ((x >> 10) & 0x1f);
|
||||
unsigned int mantissa = ((x & 0x3ff) << 13);
|
||||
|
||||
if (exponent == 0x1fU) { /* NaN or Inf */
|
||||
mantissa = (mantissa ? (sign = 0, 0x7fffffU) : 0);
|
||||
exponent = 0xffU;
|
||||
} else if (!exponent) { /* Denorm or Zero */
|
||||
if (mantissa) {
|
||||
unsigned int msb;
|
||||
exponent = 0x71U;
|
||||
do {
|
||||
msb = (mantissa & 0x400000U);
|
||||
mantissa <<= 1; /* normalize */
|
||||
--exponent;
|
||||
} while (!msb);
|
||||
mantissa &= 0x7fffffU; /* 1.mantissa is implicit */
|
||||
}
|
||||
} else {
|
||||
exponent += 0x70U;
|
||||
}
|
||||
unsigned int u = ((sign << 31) | (exponent << 23) | mantissa);
|
||||
float f;
|
||||
memcpy(&f, &u, sizeof(u));
|
||||
|
||||
return f;
|
||||
}
|
||||
|
||||
inline
|
||||
float __half2float(__half x)
|
||||
{
|
||||
return __internal_half2float(static_cast<__half_raw>(x).x);
|
||||
}
|
||||
|
||||
inline
|
||||
float __low2float(__half2 x)
|
||||
{
|
||||
return __internal_half2float(static_cast<__half2_raw>(x).x);
|
||||
}
|
||||
|
||||
inline
|
||||
float __high2float(__half2 x)
|
||||
{
|
||||
return __internal_half2float(static_cast<__half2_raw>(x).y);
|
||||
}
|
||||
} // Anonymous namespace.
|
||||
|
||||
#if !defined(HIP_NO_HALF)
|
||||
using half = __half;
|
||||
using half2 = __half2;
|
||||
#endif
|
||||
#endif // defined(__cplusplus)
|
||||
@@ -0,0 +1,76 @@
|
||||
/*
|
||||
Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
|
||||
// /*
|
||||
// Half Math Functions
|
||||
// */
|
||||
|
||||
extern "C"
|
||||
{
|
||||
__attribute__((const)) _Float16 __ocml_ceil_f16(_Float16);
|
||||
_Float16 __ocml_cos_f16(_Float16);
|
||||
__attribute__((pure)) _Float16 __ocml_exp_f16(_Float16);
|
||||
__attribute__((pure)) _Float16 __ocml_exp10_f16(_Float16);
|
||||
__attribute__((pure)) _Float16 __ocml_exp2_f16(_Float16);
|
||||
__attribute__((const)) _Float16 __ocml_floor_f16(_Float16);
|
||||
__attribute__((const))
|
||||
_Float16 __ocml_fma_f16(_Float16, _Float16, _Float16);
|
||||
__attribute__((const)) int __ocml_isinf_f16(_Float16);
|
||||
__attribute__((const)) int __ocml_isnan_f16(_Float16);
|
||||
__attribute__((pure)) _Float16 __ocml_log_f16(_Float16);
|
||||
__attribute__((pure)) _Float16 __ocml_log10_f16(_Float16);
|
||||
__attribute__((pure)) _Float16 __ocml_log2_f16(_Float16);
|
||||
__attribute__((const)) _Float16 __llvm_amdgcn_rcp_f16(_Float16);
|
||||
__attribute__((const)) _Float16 __ocml_rint_f16(_Float16);
|
||||
__attribute__((const)) _Float16 __ocml_rsqrt_f16(_Float16);
|
||||
_Float16 __ocml_sin_f16(_Float16);
|
||||
__attribute__((const)) _Float16 __ocml_sqrt_f16(_Float16);
|
||||
__attribute__((const)) _Float16 __ocml_trunc_f16(_Float16);
|
||||
|
||||
typedef _Float16 __2f16 __attribute__((ext_vector_type(2)));
|
||||
typedef short __2i16 __attribute__((ext_vector_type(2)));
|
||||
|
||||
__attribute__((const)) __2f16 __ocml_ceil_2f16(__2f16);
|
||||
__2f16 __ocml_cos_2f16(__2f16);
|
||||
__attribute__((pure)) __2f16 __ocml_exp_2f16(__2f16);
|
||||
__attribute__((pure)) __2f16 __ocml_exp10_2f16(__2f16);
|
||||
__attribute__((pure)) __2f16 __ocml_exp2_2f16(__2f16);
|
||||
__attribute__((const)) __2f16 __ocml_floor_2f16(__2f16);
|
||||
__attribute__((const)) __2f16 __ocml_fma_2f16(__2f16, __2f16, __2f16);
|
||||
__attribute__((const)) __2i16 __ocml_isinf_2f16(__2f16);
|
||||
__attribute__((const)) __2i16 __ocml_isnan_2f16(__2f16);
|
||||
__attribute__((pure)) __2f16 __ocml_log_2f16(__2f16);
|
||||
__attribute__((pure)) __2f16 __ocml_log10_2f16(__2f16);
|
||||
__attribute__((pure)) __2f16 __ocml_log2_2f16(__2f16);
|
||||
inline
|
||||
__2f16 __llvm_amdgcn_rcp_2f16(__2f16 x) // Not currently exposed by ROCDL.
|
||||
{
|
||||
return __2f16{__llvm_amdgcn_rcp_f16(x.x), __llvm_amdgcn_rcp_f16(x.y)};
|
||||
}
|
||||
__attribute__((const)) __2f16 __ocml_rint_2f16(__2f16);
|
||||
__attribute__((const)) __2f16 __ocml_rsqrt_2f16(__2f16);
|
||||
__2f16 __ocml_sin_2f16(__2f16);
|
||||
__attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16);
|
||||
__attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16);
|
||||
}
|
||||
Référencer dans un nouveau ticket
Bloquer un utilisateur