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 (NVIDIA#2046)

The issue we have is that our tests rely extensively on those conversions which makes it incredibly painfull to test
  • Loading branch information
@miscco @pciolkosz
miscco authored and pciolkosz committed Aug 4, 2024
1 parent f2335fc commit 6dffe96
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 6dffe96

Please sign in to comment.