Skip to content

Refactor fp masks #4246

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

Merged
miscco merged 2 commits into from
Mar 25, 2025
Merged

Refactor fp masks #4246

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
2e8feae
Refactor fp masks
davebayer Mar 24, 2025
2d8fb68
make msvc happy
davebayer Mar 24, 2025
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
4 changes: 2 additions & 2 deletions libcudacxx/include/cuda/std/__cmath/abs.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -87,8 +87,8 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI long double fabsl(long double __x) noe
template <class _Tp>
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __fabs_impl(_Tp __x) noexcept
{
const auto __val = _CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_v<_Tp>;
return _CUDA_VSTD::__fp_from_storage<_Tp>(static_cast<__fp_storage_t<_Tp>>(__val));
const auto __val = _CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_of_v<_Tp>;
return _CUDA_VSTD::__fp_from_storage<_Tp>(static_cast<__fp_storage_of_t<_Tp>>(__val));
}

#if _CCCL_HAS_NVFP16()
Expand Down
6 changes: 3 additions & 3 deletions libcudacxx/include/cuda/std/__cmath/copysign.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -90,9 +90,9 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __copysign_impl(_Tp __x,
{
if constexpr (numeric_limits<_Tp>::is_signed)
{
const auto __val = (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_v<_Tp>)
| (_CUDA_VSTD::__fp_get_storage(__y) & __fp_sign_mask_v<_Tp>);
return _CUDA_VSTD::__fp_from_storage<_Tp>(static_cast<__fp_storage_t<_Tp>>(__val));
const auto __val = (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_of_v<_Tp>)
| (_CUDA_VSTD::__fp_get_storage(__y) & __fp_sign_mask_of_v<_Tp>);
return _CUDA_VSTD::__fp_from_storage<_Tp>(static_cast<__fp_storage_of_t<_Tp>>(__val));
}
else
{
Expand Down
6 changes: 3 additions & 3 deletions libcudacxx/include/cuda/std/__cmath/fpclassify.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -86,9 +86,9 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr int __fpclassify_impl(_Tp __
else
{
const auto __storage = _CUDA_VSTD::__fp_get_storage(__x);
if ((__storage & __fp_exp_mask_v<_Tp>) == 0)
if ((__storage & __fp_exp_mask_of_v<_Tp>) == 0)
{
return (__storage & __fp_mant_mask_v<_Tp>) ? FP_SUBNORMAL : FP_ZERO;
return (__storage & __fp_mant_mask_of_v<_Tp>) ? FP_SUBNORMAL : FP_ZERO;
}
return FP_NORMAL;
}
Expand Down Expand Up @@ -166,7 +166,7 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr int fpclassify(__nv_fp8_e5m2
#if _CCCL_HAS_NVFP8_E8M0()
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr int fpclassify(__nv_fp8_e8m0 __x) noexcept
{
return ((__x.__x & __fp_exp_mask_v<__nv_fp8_e8m0>) == __fp_exp_mask_v<__nv_fp8_e8m0>) ? FP_NAN : FP_NORMAL;
return ((__x.__x & __fp_exp_mask_of_v<__nv_fp8_e8m0>) == __fp_exp_mask_of_v<__nv_fp8_e8m0>) ? FP_NAN : FP_NORMAL;
}
#endif // _CCCL_HAS_NVFP8_E8M0()

Expand Down
14 changes: 7 additions & 7 deletions libcudacxx/include/cuda/std/__cmath/isfinite.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -56,7 +56,7 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isfinite(float __x) noe
return ::isfinite(__x);
}
# if _LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST()
return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mask_v<float>) != __fp_exp_mask_v<float>;
return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mask_of_v<float>) != __fp_exp_mask_of_v<float>;
# else // ^^^ _LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST() ^^^ / vvv !_LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST() vvv
return _CUDA_VSTD::__isfinite_impl(__x);
# endif // ^^^ !_LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST() ^^^
Expand All @@ -73,7 +73,7 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isfinite(double __x) no
return ::isfinite(__x);
}
# if _LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST()
return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mask_v<double>) != __fp_exp_mask_v<double>;
return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mask_of_v<double>) != __fp_exp_mask_of_v<double>;
# else // ^^^ _LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST() ^^^ / vvv !_LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST() vvv
return _CUDA_VSTD::__isfinite_impl(__x);
# endif // ^^^ !_LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST() ^^^
Expand All @@ -94,35 +94,35 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isfinite(long double __
#if _CCCL_HAS_NVFP16()
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isfinite(__half __x) noexcept
{
return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mask_v<__half>) != __fp_exp_mask_v<__half>;
return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mask_of_v<__half>) != __fp_exp_mask_of_v<__half>;
}
#endif // _CCCL_HAS_NVFP16()

#if _CCCL_HAS_NVBF16()
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isfinite(__nv_bfloat16 __x) noexcept
{
return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mask_v<__nv_bfloat16>) != __fp_exp_mask_v<__nv_bfloat16>;
return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mask_of_v<__nv_bfloat16>) != __fp_exp_mask_of_v<__nv_bfloat16>;
}
#endif // _CCCL_HAS_NVBF16()

#if _CCCL_HAS_NVFP8_E4M3()
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isfinite(__nv_fp8_e4m3 __x) noexcept
{
return (__x.__x & __fp_exp_mant_mask_v<__nv_fp8_e4m3>) != __fp_exp_mant_mask_v<__nv_fp8_e4m3>;
return (__x.__x & __fp_exp_mant_mask_of_v<__nv_fp8_e4m3>) != __fp_exp_mant_mask_of_v<__nv_fp8_e4m3>;
}
#endif // _CCCL_HAS_NVFP8_E4M3()

#if _CCCL_HAS_NVFP8_E5M2()
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isfinite(__nv_fp8_e5m2 __x) noexcept
{
return (__x.__x & __fp_exp_mask_v<__nv_fp8_e5m2>) != __fp_exp_mask_v<__nv_fp8_e5m2>;
return (__x.__x & __fp_exp_mask_of_v<__nv_fp8_e5m2>) != __fp_exp_mask_of_v<__nv_fp8_e5m2>;
}
#endif // _CCCL_HAS_NVFP8_E5M2()

#if _CCCL_HAS_NVFP8_E8M0()
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isfinite(__nv_fp8_e8m0 __x) noexcept
{
return __x.__x != __fp_exp_mask_v<__nv_fp8_e8m0>;
return __x.__x != __fp_exp_mask_of_v<__nv_fp8_e8m0>;
}
#endif // _CCCL_HAS_NVFP8_E8M0()

Expand Down
11 changes: 6 additions & 5 deletions libcudacxx/include/cuda/std/__cmath/isinf.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -67,7 +67,7 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isinf(float __x) noexce
{
return ::isinf(__x);
}
return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_v<float>) == __fp_exp_mask_v<float>;
return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_of_v<float>) == __fp_exp_mask_of_v<float>;
#else // ^^^ _LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST() ^^^ / vvv !_LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST() vvv
return _CUDA_VSTD::__isinf_impl(__x);
#endif // ^^^ !_CCCL_BUILTIN_ISINF ^^^
Expand All @@ -89,7 +89,7 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isinf(double __x) noexc
{
return ::isinf(__x);
}
return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_v<double>) == __fp_exp_mask_v<double>;
return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_of_v<double>) == __fp_exp_mask_of_v<double>;
#else // ^^^ _LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST() ^^^ / vvv !_LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST() vvv
return _CUDA_VSTD::__isinf_impl(__x);
#endif // ^^^ !_CCCL_BUILTIN_ISINF ^^^
Expand Down Expand Up @@ -120,7 +120,7 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isinf(__half __x) noexc
# endif // _CCCL_STD_VER <= 2017 || _CCCL_CUDACC_BELOW(12, 3)
}
# endif // _LIBCUDACXX_HAS_NVFP16()
return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_v<__half>) == __fp_exp_mask_v<__half>;
return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_of_v<__half>) == __fp_exp_mask_of_v<__half>;
}
#endif // _CCCL_HAS_NVFP16()

Expand All @@ -138,7 +138,8 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isinf(__nv_bfloat16 __x
# endif // _CCCL_STD_VER <= 2017 || _CCCL_CUDACC_BELOW(12, 3)
}
# endif // _LIBCUDACXX_HAS_NVBF16()
return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_v<__nv_bfloat16>) == __fp_exp_mask_v<__nv_bfloat16>;
return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_of_v<__nv_bfloat16>)
== __fp_exp_mask_of_v<__nv_bfloat16>;
}
#endif // _CCCL_HAS_NVBF16()

Expand All @@ -152,7 +153,7 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isinf(__nv_fp8_e4m3) no
#if _CCCL_HAS_NVFP8_E5M2()
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isinf(__nv_fp8_e5m2 __x) noexcept
{
return (__x.__x & __fp_exp_mant_mask_v<__nv_fp8_e5m2>) == __fp_exp_mask_v<__nv_fp8_e5m2>;
return (__x.__x & __fp_exp_mant_mask_of_v<__nv_fp8_e5m2>) == __fp_exp_mask_of_v<__nv_fp8_e5m2>;
}
#endif // _CCCL_HAS_NVFP8_E5M2()

Expand Down
15 changes: 8 additions & 7 deletions libcudacxx/include/cuda/std/__cmath/isnan.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -85,7 +85,8 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isnan(__half __x) noexc
# endif // _LIBCUDACXX_HAS_NVFP16()

const auto __storage = _CUDA_VSTD::__fp_get_storage(__x);
return ((__storage & __fp_exp_mask_v<__half>) == __fp_exp_mask_v<__half>) && (__storage & __fp_mant_mask_v<__half>);
return ((__storage & __fp_exp_mask_of_v<__half>) == __fp_exp_mask_of_v<__half>)
&& (__storage & __fp_mant_mask_of_v<__half>);
}
#endif // _CCCL_HAS_NVFP16()

Expand All @@ -100,30 +101,30 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isnan(__nv_bfloat16 __x
# endif // _LIBCUDACXX_HAS_NVFP16()

const auto __storage = _CUDA_VSTD::__fp_get_storage(__x);
return ((__storage & __fp_exp_mask_v<__nv_bfloat16>) == __fp_exp_mask_v<__nv_bfloat16>)
&& (__storage & __fp_mant_mask_v<__nv_bfloat16>);
return ((__storage & __fp_exp_mask_of_v<__nv_bfloat16>) == __fp_exp_mask_of_v<__nv_bfloat16>)
&& (__storage & __fp_mant_mask_of_v<__nv_bfloat16>);
}
#endif // _CCCL_HAS_NVBF16()

#if _CCCL_HAS_NVFP8_E4M3()
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isnan(__nv_fp8_e4m3 __x) noexcept
{
return (__x.__x & __fp_exp_mant_mask_v<__nv_fp8_e4m3>) == __fp_exp_mant_mask_v<__nv_fp8_e4m3>;
return (__x.__x & __fp_exp_mant_mask_of_v<__nv_fp8_e4m3>) == __fp_exp_mant_mask_of_v<__nv_fp8_e4m3>;
}
#endif // _CCCL_HAS_NVFP8_E4M3()

#if _CCCL_HAS_NVFP8_E5M2()
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isnan(__nv_fp8_e5m2 __x) noexcept
{
return ((__x.__x & __fp_exp_mask_v<__nv_fp8_e5m2>) == __fp_exp_mask_v<__nv_fp8_e5m2>)
&& (__x.__x & __fp_mant_mask_v<__nv_fp8_e5m2>);
return ((__x.__x & __fp_exp_mask_of_v<__nv_fp8_e5m2>) == __fp_exp_mask_of_v<__nv_fp8_e5m2>)
&& (__x.__x & __fp_mant_mask_of_v<__nv_fp8_e5m2>);
}
#endif // _CCCL_HAS_NVFP8_E5M2()

#if _CCCL_HAS_NVFP8_E8M0()
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isnan(__nv_fp8_e8m0 __x) noexcept
{
return (__x.__x & __fp_exp_mask_v<__nv_fp8_e8m0>) == __fp_exp_mask_v<__nv_fp8_e8m0>;
return (__x.__x & __fp_exp_mask_of_v<__nv_fp8_e8m0>) == __fp_exp_mask_of_v<__nv_fp8_e8m0>;
}
#endif // _CCCL_HAS_NVFP8_E8M0()

Expand Down
2 changes: 1 addition & 1 deletion libcudacxx/include/cuda/std/__cmath/signbit.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -68,7 +68,7 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool __signbit_impl(_Tp __x)
{
if constexpr (numeric_limits<_Tp>::is_signed)
{
return _CUDA_VSTD::__fp_get_storage(__x) & __fp_sign_mask_v<_Tp>;
return _CUDA_VSTD::__fp_get_storage(__x) & __fp_sign_mask_of_v<_Tp>;
}
else
{
Expand Down
6 changes: 3 additions & 3 deletions libcudacxx/include/cuda/std/__floating_point/cccl_fp.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -35,7 +35,7 @@ class __cccl_fp
{
static_assert(_Fmt != __fp_format::__invalid);

using __storage_type = __fp_storage_t<__cccl_fp>;
using __storage_type = __fp_storage_t<_Fmt>;

__storage_type __storage_;

Expand Down Expand Up @@ -97,9 +97,9 @@ class __cccl_fp
}

template <class _Tp>
friend _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __fp_from_storage(__fp_storage_t<_Tp> __v) noexcept;
friend _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __fp_from_storage(__fp_storage_of_t<_Tp> __v) noexcept;
template <class _Tp>
friend _LIBCUDACXX_HIDE_FROM_ABI constexpr __fp_storage_t<_Tp> __fp_get_storage(_Tp __v) noexcept;
friend _LIBCUDACXX_HIDE_FROM_ABI constexpr __fp_storage_of_t<_Tp> __fp_get_storage(_Tp __v) noexcept;
};

_LIBCUDACXX_END_NAMESPACE_STD
Expand Down
Loading