Skip to content

Commit

Permalink
Do not rely on conversions between float and extended floating point …
Browse files Browse the repository at this point in the history 
…types

The issue we have is that our tests rely extensively on those conversions which makes it incredibly painfull to test
  • Loading branch information
@miscco
miscco committed Jul 24, 2024
1 parent 53fe08f commit a994cc5
Show file tree
Hide file tree
Showing 2 changed files with 20 additions and 20 deletions.
20 changes: 10 additions & 10 deletions libcudacxx/include/cuda/std/__cuda/cmath_nvbf16.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -37,47 +37,47 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD
// trigonometric functions
inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 sin(__nv_bfloat16 __v)
{
NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hsin(__v);), (return __nv_bfloat16(::sin(float(__v)));))
NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hsin(__v);), (return __float2bfloat16(::sin(__bfloat162float(__v)));))
}

inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 sinh(__nv_bfloat16 __v)
{
return __nv_bfloat16(::sinh(float(__v)));
return __float2bfloat16(::sinh(__bfloat162float(__v)));
}

inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 cos(__nv_bfloat16 __v)
{
NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hcos(__v);), (return __nv_bfloat16(::cos(float(__v)));))
NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hcos(__v);), (return __float2bfloat16(::cos(__bfloat162float(__v)));))
}

inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 cosh(__nv_bfloat16 __v)
{
return __nv_bfloat16(::cosh(float(__v)));
return __float2bfloat16(::cosh(__bfloat162float(__v)));
}

inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 exp(__nv_bfloat16 __v)
{
NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hexp(__v);), (return __nv_bfloat16(::exp(float(__v)));))
NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hexp(__v);), (return __float2bfloat16(::exp(__bfloat162float(__v)));))
}

inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 hypot(__nv_bfloat16 __x, __nv_bfloat16 __y)
{
return __nv_bfloat16(::hypot(float(__x), float(__y)));
return __float2bfloat16(::hypot(__bfloat162float(__x), __bfloat162float(__y)));
}

inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 atan2(__nv_bfloat16 __x, __nv_bfloat16 __y)
{
return __nv_bfloat16(::atan2(float(__x), float(__y)));
return __float2bfloat16(::atan2(__bfloat162float(__x), __bfloat162float(__y)));
}

inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 log(__nv_bfloat16 __x)
{
NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hlog(__x);), (return __nv_bfloat16(::log(float(__x)));))
NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hlog(__x);), (return __float2bfloat16(::log(__bfloat162float(__x)));))
}

inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 sqrt(__nv_bfloat16 __x)
{
NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hsqrt(__x);), (return __nv_bfloat16(::sqrt(float(__x)));))
NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hsqrt(__x);), (return __float2bfloat16(::sqrt(__bfloat162float(__x)));))
}

// floating point helper
Expand Down Expand Up @@ -123,7 +123,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY bool isfinite(__nv_bfloat16 __v)

inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 __constexpr_copysign(__nv_bfloat16 __x, __nv_bfloat16 __y) noexcept
{
return __nv_bfloat16(::copysignf(float(__x), float(__y)));
return __float2bfloat16(::copysignf(__bfloat162float(__x), __bfloat162float(__y)));
}

inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 copysign(__nv_bfloat16 __x, __nv_bfloat16 __y)
Expand Down
20 changes: 10 additions & 10 deletions libcudacxx/include/cuda/std/__cuda/cmath_nvfp16.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -35,7 +35,7 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD
inline _LIBCUDACXX_INLINE_VISIBILITY __half sin(__half __v)
{
NV_IF_ELSE_TARGET(NV_PROVIDES_SM_53, (return ::hsin(__v);), ({
float __vf = __v;
float __vf = __half2float(__v);
__vf = ::sin(__vf);
__half_raw __ret_repr = ::__float2half_rn(__vf);

Expand All @@ -61,7 +61,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half sin(__half __v)

inline _LIBCUDACXX_INLINE_VISIBILITY __half sinh(__half __v)
{
return __half(::sinh(float(__v)));
return __float2half(::sinh(__half2float(__v)));
}

// clang-format off
Expand All @@ -71,7 +71,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half cos(__half __v)
return ::hcos(__v);
), (
{
float __vf = __v;
float __vf = __half2float(__v);
__vf = ::cos(__vf);
__half_raw __ret_repr = ::__float2half_rn(__vf);

Expand All @@ -94,7 +94,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half cos(__half __v)

inline _LIBCUDACXX_INLINE_VISIBILITY __half cosh(__half __v)
{
return __half(::cosh(float(__v)));
return __float2half(::cosh(__half2float(__v)));
}

// clang-format off
Expand All @@ -104,7 +104,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half exp(__half __v)
return ::hexp(__v);
), (
{
float __vf = __v;
float __vf = __half2float(__v);
__vf = ::exp(__vf);
__half_raw __ret_repr = ::__float2half_rn(__vf);

Expand All @@ -127,12 +127,12 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half exp(__half __v)

inline _LIBCUDACXX_INLINE_VISIBILITY __half hypot(__half __x, __half __y)
{
return __half(::hypot(float(__x), float(__y)));
return __float2half(::hypot(__half2float(__x), __half2float(__y)));
}

inline _LIBCUDACXX_INLINE_VISIBILITY __half atan2(__half __x, __half __y)
{
return __half(::atan2(float(__x), float(__y)));
return __float2half(::atan2(__half2float(__x), __half2float(__y)));
}

// clang-format off
Expand All @@ -142,7 +142,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half log(__half __x)
return ::hlog(__x);
), (
{
float __vf = __x;
float __vf = __half2float(__x);
__vf = ::log(__vf);
__half_raw __ret_repr = ::__float2half_rn(__vf);

Expand All @@ -164,7 +164,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half log(__half __x)

inline _LIBCUDACXX_INLINE_VISIBILITY __half sqrt(__half __x)
{
NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hsqrt(__x);), (return __half(::sqrt(float(__x)));))
NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hsqrt(__x);), (return __float2half(::sqrt(__half2float(__x)));))
}

// floating point helper
Expand Down Expand Up @@ -210,7 +210,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY bool isfinite(__half __v)

inline _LIBCUDACXX_INLINE_VISIBILITY __half __constexpr_copysign(__half __x, __half __y) noexcept
{
return __half(::copysignf(float(__x), float(__y)));
return __float2half(::copysignf(__half2float(__x), __half2float(__y)));
}

inline _LIBCUDACXX_INLINE_VISIBILITY __half copysign(__half __x, __half __y)
Expand Down

0 comments on commit a994cc5

Please sign in to comment.