Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Remove xfloat32 which is not used in hiptensor #169

Merged
merged 1 commit into from
Dec 12, 2023
Merged

Remove xfloat32 which is not used in hiptensor #169

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 0 additions & 5 deletions library/include/hiptensor/internal/native_types.hpp
View file
Open in desktop
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
View file
Open in desktop
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