Skip to content

Commit

Permalink
Merge pull request #169 from CongMa13/remove_xfloat32
Browse files Browse the repository at this point in the history
Remove xfloat32 which is not used in hiptensor
  • Loading branch information
CongMa13 authored Dec 12, 2023
2 parents b88175f + 76de7d0 commit 53d687b
Show file tree
Hide file tree
Showing 3 changed files with 5 additions and 416 deletions.
5 changes: 0 additions & 5 deletions library/include/hiptensor/internal/native_types.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -33,8 +33,6 @@
#include <type_traits>
#include <utility>

#include "xfloat32.hpp"

namespace hiptensor
{

Expand Down Expand Up @@ -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


Expand Down
82 changes: 5 additions & 77 deletions library/include/hiptensor/internal/type_traits.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -26,9 +26,11 @@

#ifndef HIPTENSOR_TYPE_TRAITS_HPP
#define HIPTENSOR_TYPE_TRAITS_HPP
#include "native_types.hpp"
#include <cfloat>

#include "config.hpp"
#include "native_types.hpp"

namespace hiptensor
{
namespace detail
Expand Down Expand Up @@ -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)
Expand All @@ -81,10 +82,6 @@ namespace hiptensor
: f32(initVal)
{
}
constexpr Fp32Bits(xfloat32_t initVal)
: xf32(initVal)
{
}
};

} // namespace detail
Expand Down Expand Up @@ -273,68 +270,6 @@ namespace std
hiptensor::detail::Fp16Bits eps(static_cast<uint16_t>(0x7FC0));
return eps.b16;
}

///////////////////////////////////////////////////////////
/////////// std::numeric_limits<xfloat32_t> //////////////
///////////////////////////////////////////////////////////

template <>
HIPTENSOR_HOST_DEVICE constexpr hiptensor::xfloat32_t
numeric_limits<hiptensor::xfloat32_t>::epsilon() noexcept
{
hiptensor::detail::Fp32Bits eps(static_cast<float>(FLT_EPSILON));
return eps.xf32;
}

template <>
HIPTENSOR_HOST_DEVICE constexpr hiptensor::xfloat32_t
numeric_limits<hiptensor::xfloat32_t>::infinity() noexcept
{
hiptensor::detail::Fp32Bits eps(static_cast<float>(HUGE_VALF));
return eps.xf32;
}

template <>
HIPTENSOR_HOST_DEVICE constexpr hiptensor::xfloat32_t
numeric_limits<hiptensor::xfloat32_t>::lowest() noexcept
{
hiptensor::detail::Fp32Bits eps(static_cast<float>(-FLT_MAX));
return eps.xf32;
}

template <>
HIPTENSOR_HOST_DEVICE constexpr hiptensor::xfloat32_t
numeric_limits<hiptensor::xfloat32_t>::max() noexcept
{
hiptensor::detail::Fp32Bits eps(static_cast<float>(FLT_MAX));
return eps.xf32;
}

template <>
HIPTENSOR_HOST_DEVICE constexpr hiptensor::xfloat32_t
numeric_limits<hiptensor::xfloat32_t>::min() noexcept
{
hiptensor::detail::Fp32Bits eps(static_cast<float>(FLT_MIN));
return eps.xf32;
}

template <>
HIPTENSOR_HOST_DEVICE constexpr hiptensor::xfloat32_t
numeric_limits<hiptensor::xfloat32_t>::quiet_NaN() noexcept
{
hiptensor::detail::Fp32Bits eps(static_cast<uint32_t>(0x7FF80000));
return eps.xf32;
}

template <>
HIPTENSOR_HOST_DEVICE constexpr hiptensor::xfloat32_t
numeric_limits<hiptensor::xfloat32_t>::signaling_NaN() noexcept
{
hiptensor::detail::Fp32Bits eps(static_cast<uint32_t>(0x7FF00000));
return eps.xf32;
}
// @endcond

} // namespace std

namespace hiptensor
Expand Down Expand Up @@ -378,13 +313,6 @@ namespace hiptensor
// b16 mantissa is 7 bits
return ((int32_t)1 << 8);
}

template <typename T, typename std::enable_if_t<std::is_same<T, xfloat32_t>::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
Loading

0 comments on commit 53d687b

Please sign in to comment.