diff --git a/library/include/hiptensor/internal/native_types.hpp b/library/include/hiptensor/internal/native_types.hpp index 6c9dbee8..69ce706f 100644 --- a/library/include/hiptensor/internal/native_types.hpp +++ b/library/include/hiptensor/internal/native_types.hpp @@ -33,8 +33,6 @@ #include #include -#include "xfloat32.hpp" - namespace hiptensor { @@ -84,9 +82,6 @@ namespace hiptensor #if !HIPTENSOR_NO_HALF using hfloat16_t = __half; #endif // !HIPTENSOR_NO_HALF - - using xfloat32_t = hiptensor_xfloat32; - // clang-format off diff --git a/library/include/hiptensor/internal/type_traits.hpp b/library/include/hiptensor/internal/type_traits.hpp index 3867839d..48566051 100644 --- a/library/include/hiptensor/internal/type_traits.hpp +++ b/library/include/hiptensor/internal/type_traits.hpp @@ -26,9 +26,11 @@ #ifndef HIPTENSOR_TYPE_TRAITS_HPP #define HIPTENSOR_TYPE_TRAITS_HPP -#include "native_types.hpp" #include +#include "config.hpp" +#include "native_types.hpp" + namespace hiptensor { namespace detail @@ -69,9 +71,8 @@ namespace hiptensor { union { - uint32_t i32; - float32_t f32; - xfloat32_t xf32; + uint32_t i32; + float32_t f32; }; constexpr Fp32Bits(uint32_t initVal) : i32(initVal) @@ -81,10 +82,6 @@ namespace hiptensor : f32(initVal) { } - constexpr Fp32Bits(xfloat32_t initVal) - : xf32(initVal) - { - } }; } // namespace detail @@ -273,68 +270,6 @@ namespace std hiptensor::detail::Fp16Bits eps(static_cast(0x7FC0)); return eps.b16; } - - /////////////////////////////////////////////////////////// - /////////// std::numeric_limits ////////////// - /////////////////////////////////////////////////////////// - - template <> - HIPTENSOR_HOST_DEVICE constexpr hiptensor::xfloat32_t - numeric_limits::epsilon() noexcept - { - hiptensor::detail::Fp32Bits eps(static_cast(FLT_EPSILON)); - return eps.xf32; - } - - template <> - HIPTENSOR_HOST_DEVICE constexpr hiptensor::xfloat32_t - numeric_limits::infinity() noexcept - { - hiptensor::detail::Fp32Bits eps(static_cast(HUGE_VALF)); - return eps.xf32; - } - - template <> - HIPTENSOR_HOST_DEVICE constexpr hiptensor::xfloat32_t - numeric_limits::lowest() noexcept - { - hiptensor::detail::Fp32Bits eps(static_cast(-FLT_MAX)); - return eps.xf32; - } - - template <> - HIPTENSOR_HOST_DEVICE constexpr hiptensor::xfloat32_t - numeric_limits::max() noexcept - { - hiptensor::detail::Fp32Bits eps(static_cast(FLT_MAX)); - return eps.xf32; - } - - template <> - HIPTENSOR_HOST_DEVICE constexpr hiptensor::xfloat32_t - numeric_limits::min() noexcept - { - hiptensor::detail::Fp32Bits eps(static_cast(FLT_MIN)); - return eps.xf32; - } - - template <> - HIPTENSOR_HOST_DEVICE constexpr hiptensor::xfloat32_t - numeric_limits::quiet_NaN() noexcept - { - hiptensor::detail::Fp32Bits eps(static_cast(0x7FF80000)); - return eps.xf32; - } - - template <> - HIPTENSOR_HOST_DEVICE constexpr hiptensor::xfloat32_t - numeric_limits::signaling_NaN() noexcept - { - hiptensor::detail::Fp32Bits eps(static_cast(0x7FF00000)); - return eps.xf32; - } - // @endcond - } // namespace std namespace hiptensor @@ -378,13 +313,6 @@ namespace hiptensor // b16 mantissa is 7 bits return ((int32_t)1 << 8); } - - template ::value, int> = 0> - constexpr auto maxExactInteger() -> int32_t - { - // xf32 mantissa is 7 bits - return ((int32_t)1 << 8); - } } // namespace hiptensor #endif // HIPTENSOR_TYPE_TRAITS_HPP diff --git a/library/include/hiptensor/internal/xfloat32.hpp b/library/include/hiptensor/internal/xfloat32.hpp deleted file mode 100644 index 6e9168cf..00000000 --- a/library/include/hiptensor/internal/xfloat32.hpp +++ /dev/null @@ -1,334 +0,0 @@ -/* ************************************************************************ - * Copyright (C) 2016-2024 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 cop- - * ies 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 IM- - * PLIED, 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 CONNE- - * CTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. - * - * ************************************************************************ */ - -/*!\file - * \brief xfloat32.h provides struct for hiptensor_xfloat32 typedef - */ - -#ifndef HIPTENSOR_XFLOAT32_HPP -#define HIPTENSOR_XFLOAT32_HPP - -#if __cplusplus < 201103L || (!defined(__HCC__) && !defined(__HIPCC__)) - -// If this is a C compiler, C++ compiler below C++11, or a host-only compiler, we only -// include a minimal definition of hiptensor_xfloat32 - -#include -typedef struct -{ - float data; -} hiptensor_xfloat32; - -#else // __cplusplus < 201103L || (!defined(__HCC__) && !defined(__HIPCC__)) - -#include -#include -#include -#include -#include -#include - -#include "config.hpp" - -struct hiptensor_xfloat32 -{ - float data; - - enum round_t - { - round_up - }; - - HIPTENSOR_HOST_DEVICE hiptensor_xfloat32() = default; - - // round upper 19 bits of IEEE float to convert to xfloat32 - explicit HIPTENSOR_HOST_DEVICE hiptensor_xfloat32(float f, round_t) - : data(float_to_xfloat32(f)) - { - } - - explicit HIPTENSOR_HOST_DEVICE hiptensor_xfloat32(float f) - : data(truncate_float_to_xfloat32(f)) - { - } - - // zero extend lower 13 bits of xfloat32 to convert to IEEE float - HIPTENSOR_HOST_DEVICE operator float() const - { - return data; - } - - explicit HIPTENSOR_HOST_DEVICE operator bool() const - { - union - { - float fp32; - uint32_t int32; - } u = {data}; - return u.int32 & 0x7fffe000; - } - - explicit HIPTENSOR_HOST_DEVICE operator uint32_t() const - { - return uint32_t(float(*this)); - } - - explicit HIPTENSOR_HOST_DEVICE operator long() const - { - return long(float(*this)); - } - - explicit HIPTENSOR_HOST_DEVICE operator double() const - { - return double(float(*this)); - } - -private: - static HIPTENSOR_HOST_DEVICE float float_to_xfloat32(float f) - { - union - { - float fp32; - uint32_t int32; - } u = {f}; - if(~u.int32 & 0x7f800000) - { - // When the exponent bits are not all 1s, then the value is zero, normal, - // or subnormal. We round the xfloat32 mantissa up by adding 0xFFF, plus - // 1 if the least significant bit of the xfloat32 mantissa is 1 (odd). - // This causes the xfloat32's mantissa to be incremented by 1 if the 13 - // least significant bits of the float mantissa are greater than 0x1000, - // or if they are equal to 0x1000 and the least significant bit of the - // xfloat32 mantissa is 1 (odd). This causes it to be rounded to even when - // the lower 13 bits are exactly 0x1000. If the xfloat32 mantissa already - // has the value 0x3ff, then incrementing it causes it to become 0x00 and - // the exponent is incremented by one, which is the next higher FP value - // to the unrounded xfloat32 value. When the xfloat32 value is subnormal - // with an exponent of 0x00 and a mantissa of 0x3FF, it may be rounded up - // to a normal value with an exponent of 0x01 and a mantissa of 0x00. - // When the xfloat32 value has an exponent of 0xFE and a mantissa of 0x3FF, - // incrementing it causes it to become an exponent of 0xFF and a mantissa - // of 0x00, which is Inf, the next higher value to the unrounded value. - - u.int32 += 0xfff + ((u.int32 >> 13) & 1); // Round to nearest, round to even - } - else if(u.int32 & 0x1fff) - { - // When all of the exponent bits are 1, the value is Inf or NaN. - // Inf is indicated by a zero mantissa. NaN is indicated by any nonzero - // mantissa bit. Quiet NaN is indicated by the most significant mantissa - // bit being 1. Signaling NaN is indicated by the most significant - // mantissa bit being 0 but some other bit(s) being 1. If any of the - // lower 13 bits of the mantissa are 1, we set the least significant bit - // of the xfloat32 mantissa, in order to preserve signaling NaN in case - // the xfloat32's mantissa bits are all 0. - u.int32 |= 0x2000; // Preserve signaling NaN - } - - u.int32 &= 0xffffe000; - return u.fp32; - } - - // Truncate instead of rounding - static HIPTENSOR_HOST_DEVICE float truncate_float_to_xfloat32(float f) - { - union - { - float fp32; - uint32_t int32; - } u = {f}; - - u.int32 = u.int32 & 0xffffe000; - return u.fp32; - } -}; - -typedef struct -{ - float data; -} hiptensor_xfloat32_public; - -static_assert(std::is_standard_layout{}, - "hiptensor_xfloat32 is not a standard layout type, and thus is " - "incompatible with C."); - -static_assert(std::is_trivial{}, - "hiptensor_xfloat32 is not a trivial type, and thus is " - "incompatible with C."); - -static_assert(sizeof(hiptensor_xfloat32) == sizeof(hiptensor_xfloat32_public) - && offsetof(hiptensor_xfloat32, data) - == offsetof(hiptensor_xfloat32_public, data), - "internal hiptensor_xfloat32 does not match public hiptensor_xfloat32"); - -inline std::ostream& operator<<(std::ostream& os, const hiptensor_xfloat32& xf32) -{ - return os << float(xf32); -} - -inline HIPTENSOR_HOST_DEVICE hiptensor_xfloat32 operator+(hiptensor_xfloat32 a) -{ - return a; -} -inline HIPTENSOR_HOST_DEVICE hiptensor_xfloat32 operator-(hiptensor_xfloat32 a) -{ - union - { - float fp32; - uint32_t int32; - } u = {a.data}; - u.int32 ^= 0x80000000; - return hiptensor_xfloat32(u.fp32); -} -inline HIPTENSOR_HOST_DEVICE hiptensor_xfloat32 operator+(hiptensor_xfloat32 a, - hiptensor_xfloat32 b) -{ - return hiptensor_xfloat32(float(a) + float(b)); -} -inline HIPTENSOR_HOST_DEVICE hiptensor_xfloat32 operator-(hiptensor_xfloat32 a, - hiptensor_xfloat32 b) -{ - return hiptensor_xfloat32(float(a) - float(b)); -} -inline HIPTENSOR_HOST_DEVICE hiptensor_xfloat32 operator*(hiptensor_xfloat32 a, - hiptensor_xfloat32 b) -{ - return hiptensor_xfloat32(float(a) * float(b)); -} -inline HIPTENSOR_HOST_DEVICE hiptensor_xfloat32 operator/(hiptensor_xfloat32 a, - hiptensor_xfloat32 b) -{ - return hiptensor_xfloat32(float(a) / float(b)); -} -inline HIPTENSOR_HOST_DEVICE bool operator<(hiptensor_xfloat32 a, hiptensor_xfloat32 b) -{ - return float(a) < float(b); -} -inline HIPTENSOR_HOST_DEVICE bool operator==(hiptensor_xfloat32 a, hiptensor_xfloat32 b) -{ - return float(a) == float(b); -} -inline HIPTENSOR_HOST_DEVICE bool operator>(hiptensor_xfloat32 a, hiptensor_xfloat32 b) -{ - return b < a; -} -inline HIPTENSOR_HOST_DEVICE bool operator<=(hiptensor_xfloat32 a, hiptensor_xfloat32 b) -{ - return !(a > b); -} -inline HIPTENSOR_HOST_DEVICE bool operator!=(hiptensor_xfloat32 a, hiptensor_xfloat32 b) -{ - return !(a == b); -} -inline HIPTENSOR_HOST_DEVICE bool operator>=(hiptensor_xfloat32 a, hiptensor_xfloat32 b) -{ - return !(a < b); -} -inline HIPTENSOR_HOST_DEVICE hiptensor_xfloat32& operator+=(hiptensor_xfloat32& a, - hiptensor_xfloat32 b) -{ - return a = a + b; -} -inline HIPTENSOR_HOST_DEVICE hiptensor_xfloat32& operator-=(hiptensor_xfloat32& a, - hiptensor_xfloat32 b) -{ - return a = a - b; -} -inline HIPTENSOR_HOST_DEVICE hiptensor_xfloat32& operator*=(hiptensor_xfloat32& a, - hiptensor_xfloat32 b) -{ - return a = a * b; -} -inline HIPTENSOR_HOST_DEVICE hiptensor_xfloat32& operator/=(hiptensor_xfloat32& a, - hiptensor_xfloat32 b) -{ - return a = a / b; -} -inline HIPTENSOR_HOST_DEVICE hiptensor_xfloat32& operator++(hiptensor_xfloat32& a) -{ - return a += hiptensor_xfloat32(1.0f); -} -inline HIPTENSOR_HOST_DEVICE hiptensor_xfloat32& operator--(hiptensor_xfloat32& a) -{ - return a -= hiptensor_xfloat32(1.0f); -} -inline HIPTENSOR_HOST_DEVICE hiptensor_xfloat32 operator++(hiptensor_xfloat32& a, int) -{ - hiptensor_xfloat32 orig = a; - ++a; - return orig; -} -inline HIPTENSOR_HOST_DEVICE hiptensor_xfloat32 operator--(hiptensor_xfloat32& a, int) -{ - hiptensor_xfloat32 orig = a; - --a; - return orig; -} - -namespace std -{ - constexpr HIPTENSOR_HOST_DEVICE bool isinf(hiptensor_xfloat32 a) - { - union - { - float fp32; - uint32_t int32; - } u = {a.data}; - return !(~u.int32 & 0x7f800000) && !(u.int32 & 0x7fe000); - } - constexpr HIPTENSOR_HOST_DEVICE bool isnan(hiptensor_xfloat32 a) - { - union - { - float fp32; - uint32_t int32; - } u = {a.data}; - return !(~u.int32 & 0x7f800000) && +(u.int32 & 0x7fe000); - } - constexpr HIPTENSOR_HOST_DEVICE bool iszero(hiptensor_xfloat32 a) - { - union - { - float fp32; - uint32_t int32; - } u = {a.data}; - return (u.fp32 == 0.0f); - } - - HIPTENSOR_HOST_DEVICE inline hiptensor_xfloat32 sin(hiptensor_xfloat32 a) - { - return hiptensor_xfloat32(sinf(float(a))); - } - HIPTENSOR_HOST_DEVICE inline hiptensor_xfloat32 cos(hiptensor_xfloat32 a) - { - return hiptensor_xfloat32(cosf(float(a))); - } - - HIPTENSOR_HOST_DEVICE constexpr hiptensor_xfloat32 real(const hiptensor_xfloat32& a) - { - return a; - } -} - -#endif // __cplusplus < 201103L || (!defined(__HCC__) && !defined(__HIPCC__)) - -#endif // HIPTENSOR_XFLOAT32_HPP