diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16_gcc.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16_gcc.h index 9b31f9e3ce..480fd8120a 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16_gcc.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16_gcc.h @@ -95,160 +95,157 @@ struct __half2_raw { }; // END STRUCT __HALF2 - namespace + inline + unsigned short __internal_float2half( + float flt, unsigned int& sgn, unsigned int& rem) { - inline - unsigned short __internal_float2half( - float flt, unsigned int& sgn, unsigned int& rem) - { - unsigned int x{}; - std::memcpy(&x, &flt, sizeof(flt)); + unsigned int x{}; + std::memcpy(&x, &flt, sizeof(flt)); - unsigned int u = (x & 0x7fffffffU); - sgn = ((x >> 16) & 0x8000U); + unsigned int u = (x & 0x7fffffffU); + sgn = ((x >> 16) & 0x8000U); - // NaN/+Inf/-Inf - if (u >= 0x7f800000U) { - rem = 0; - return static_cast( - (u == 0x7f800000U) ? (sgn | 0x7c00U) : 0x7fffU); + // NaN/+Inf/-Inf + if (u >= 0x7f800000U) { + rem = 0; + return static_cast( + (u == 0x7f800000U) ? (sgn | 0x7c00U) : 0x7fffU); + } + // Overflows + if (u > 0x477fefffU) { + rem = 0x80000000U; + return static_cast(sgn | 0x7bffU); + } + // Normal numbers + if (u >= 0x38800000U) { + rem = u << 19; + u -= 0x38000000U; + return static_cast(sgn | (u >> 13)); + } + // +0/-0 + if (u < 0x33000001U) { + rem = u; + return static_cast(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(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 */ } - // Overflows - if (u > 0x477fefffU) { - rem = 0x80000000U; - return static_cast(sgn | 0x7bffU); - } - // Normal numbers - if (u >= 0x38800000U) { - rem = u << 19; - u -= 0x38000000U; - return static_cast(sgn | (u >> 13)); - } - // +0/-0 - if (u < 0x33000001U) { - rem = u; - return static_cast(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(sgn | (mantissa >> shift)); + } else { + exponent += 0x70U; } + unsigned int u = ((sign << 31) | (exponent << 23) | mantissa); + float f; + memcpy(&f, &u, sizeof(u)); - 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 f; + } - return r; - } + inline + float __half2float(__half x) + { + return __internal_half2float(static_cast<__half_raw>(x).x); + } - inline - __half __float2half_rn(float x) { return __float2half(x); } + inline + float __low2float(__half2 x) + { + return __internal_half2float(static_cast<__half2_raw>(x).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. + inline + float __high2float(__half2 x) + { + return __internal_half2float(static_cast<__half2_raw>(x).y); + } #if !defined(HIP_NO_HALF) using half = __half;