Move definition of floating point trait functions to their own file

NVIDIA · Nov 17, 2024 · 1836891 · 1836891
1 parent 31d53db
commit 1836891
Show file tree

Hide file tree

Showing 7 changed files with 1,033 additions and 288 deletions.
diff --git a/libcudacxx/include/cuda/std/__cccl/builtin.h b/libcudacxx/include/cuda/std/__cccl/builtin.h
@@ -210,6 +210,33 @@
 #  undef _CCCL_BUILTIN_IS_CONSTANT_EVALUATED
 #endif // _CCCL_STD_VER < 2014 && _CCCL_CUDA_COMPILER_NVCC
 
+#if _CCCL_CHECK_BUILTIN(builtin_isfinite) || defined(_CCCL_COMPILER_GCC) || defined(_CCCL_COMPILER_NVRTC)
+#  define _CCCL_BUILTIN_ISFINITE(...) __builtin_isfinite(__VA_ARGS__)
+#endif // _CCCL_CHECK_BUILTIN(isfinite)
+
+// Below 11.7 nvcc treats the builtin as a host only function
+#if _CCCL_CUDACC_BELOW(11, 7)
+#  undef _CCCL_BUILTIN_ISFINITE
+#endif // _CCCL_CUDACC_BELOW(11, 7)
+
+#if _CCCL_CHECK_BUILTIN(builtin_isinf) || defined(_CCCL_COMPILER_GCC)
+#  define _CCCL_BUILTIN_ISINF(...) __builtin_isinf(__VA_ARGS__)
+#endif // _CCCL_CHECK_BUILTIN(isinf)
+
+// Below 11.7 nvcc treats the builtin as a host only function
+#if _CCCL_CUDACC_BELOW(11, 7)
+#  undef _CCCL_BUILTIN_ISINF
+#endif // _CCCL_CUDACC_BELOW(11, 7)
+
+#if _CCCL_CHECK_BUILTIN(builtin_isnan) || defined(_CCCL_COMPILER_GCC)
+#  define _CCCL_BUILTIN_ISNAN(...) __builtin_isnan(__VA_ARGS__)
+#endif // _CCCL_CHECK_BUILTIN(isnan)
+
+// Below 11.7 nvcc treats the builtin as a host only function
+#if _CCCL_CUDACC_BELOW(11, 7)
+#  undef _CCCL_BUILTIN_ISNAN
+#endif // _CCCL_CUDACC_BELOW(11, 7)
+
 #if (_CCCL_CHECK_BUILTIN(builtin_launder) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 70000))
 #  define _CCCL_BUILTIN_LAUNDER(...) __builtin_launder(__VA_ARGS__)
 #endif // _CCCL_CHECK_BUILTIN(builtin_launder) && gcc >= 7
@@ -321,6 +348,15 @@
 #  define _CCCL_BUILTIN_OPERATOR_NEW(...)    __builtin_operator_new(__VA_ARGS__)
 #endif // _CCCL_CHECK_BUILTIN(__builtin_operator_new) && _CCCL_CHECK_BUILTIN(__builtin_operator_delete)
 
+#if _CCCL_CHECK_BUILTIN(builtin_signbit) || defined(_CCCL_COMPILER_GCC)
+#  define _CCCL_BUILTIN_SIGNBIT(...) __builtin_signbit(__VA_ARGS__)
+#endif // _CCCL_CHECK_BUILTIN(builtin_signbit)
+
+// Below 11.7 nvcc treats the builtin as a host only function
+#if _CCCL_CUDACC_BELOW(11, 7)
+#  undef _CCCL_BUILTIN_SIGNBIT
+#endif // _CCCL_CUDACC_BELOW(11, 7)
+
 #if _CCCL_HAS_BUILTIN(__decay) && defined(_CCCL_CUDA_COMPILER_CLANG)
 #  define _CCCL_BUILTIN_DECAY(...) __decay(__VA_ARGS__)
 #endif // _CCCL_HAS_BUILTIN(__decay) && clang-cuda
@@ -593,18 +629,6 @@
 #  define _CCCL_BUILTIN_IS_VOLATILE(...) __is_volatile(__VA_ARGS__)
 #endif // _CCCL_HAS_BUILTIN(__is_volatile)
 
-#if _CCCL_CHECK_BUILTIN(isfinite)
-#  define _CCCL_BUILTIN_ISFINITE(...) __builtin_isfinite(__VA_ARGS__)
-#endif // _CCCL_CHECK_BUILTIN(isfinite)
-
-#if _CCCL_CHECK_BUILTIN(isinf)
-#  define _CCCL_BUILTIN_ISINF(...) __builtin_isinf(__VA_ARGS__)
-#endif // _CCCL_CHECK_BUILTIN(isinf)
-
-#if _CCCL_CHECK_BUILTIN(isnan)
-#  define _CCCL_BUILTIN_ISNAN(...) __builtin_isnan(__VA_ARGS__)
-#endif // _CCCL_CHECK_BUILTIN(isnan)
-
 #if _CCCL_CHECK_BUILTIN(make_integer_seq) || (defined(_CCCL_COMPILER_MSVC) && _CCCL_MSVC_VERSION >= 1923)
 #  define _CCCL_BUILTIN_MAKE_INTEGER_SEQ(...) __make_integer_seq<__VA_ARGS__>
 #endif // _CCCL_CHECK_BUILTIN(make_integer_seq)

diff --git a/libcudacxx/include/cuda/std/__cmath/nvbf16.h b/libcudacxx/include/cuda/std/__cmath/nvbf16.h
@@ -76,46 +76,6 @@ _LIBCUDACXX_HIDE_FROM_ABI __nv_bfloat16 sqrt(__nv_bfloat16 __x)
 }
 
 // floating point helper
-_LIBCUDACXX_HIDE_FROM_ABI bool signbit(__nv_bfloat16 __v)
-{
-  return ::signbit(::__bfloat162float(__v));
-}
-
-_LIBCUDACXX_HIDE_FROM_ABI bool __constexpr_isnan(__nv_bfloat16 __x) noexcept
-{
-  return ::__hisnan(__x);
-}
-
-_LIBCUDACXX_HIDE_FROM_ABI bool isnan(__nv_bfloat16 __v)
-{
-  return _CUDA_VSTD::__constexpr_isnan(__v);
-}
-
-_LIBCUDACXX_HIDE_FROM_ABI bool __constexpr_isinf(__nv_bfloat16 __x) noexcept
-{
-#  if _CCCL_STD_VER >= 2020 && _CCCL_CUDACC_BELOW(12, 3)
-  // this is a workaround for nvbug 4362808
-  return !::__hisnan(__x) && ::__hisnan(__x - __x);
-#  else // ^^^ C++20 && below 12.3 ^^^ / vvv C++17 or 12.3+ vvv
-  return ::__hisinf(__x) != 0;
-#  endif // _CCCL_STD_VER <= 2017 || _CCCL_CUDACC_BELOW(12, 3)
-}
-
-_LIBCUDACXX_HIDE_FROM_ABI bool isinf(__nv_bfloat16 __v)
-{
-  return _CUDA_VSTD::__constexpr_isinf(__v);
-}
-
-_LIBCUDACXX_HIDE_FROM_ABI bool __constexpr_isfinite(__nv_bfloat16 __x) noexcept
-{
-  return !_CUDA_VSTD::__constexpr_isnan(__x) && !_CUDA_VSTD::__constexpr_isinf(__x);
-}
-
-_LIBCUDACXX_HIDE_FROM_ABI bool isfinite(__nv_bfloat16 __v)
-{
-  return _CUDA_VSTD::__constexpr_isfinite(__v);
-}
-
 _LIBCUDACXX_HIDE_FROM_ABI __nv_bfloat16 __constexpr_copysign(__nv_bfloat16 __x, __nv_bfloat16 __y) noexcept
 {
   return __float2bfloat16(::copysignf(__bfloat162float(__x), __bfloat162float(__y)));

diff --git a/libcudacxx/include/cuda/std/__cmath/nvfp16.h b/libcudacxx/include/cuda/std/__cmath/nvfp16.h
@@ -141,46 +141,6 @@ _LIBCUDACXX_HIDE_FROM_ABI __half sqrt(__half __x)
 }
 
 // floating point helper
-_LIBCUDACXX_HIDE_FROM_ABI bool signbit(__half __v)
-{
-  return ::signbit(::__half2float(__v));
-}
-
-_LIBCUDACXX_HIDE_FROM_ABI bool __constexpr_isnan(__half __x) noexcept
-{
-  return ::__hisnan(__x);
-}
-
-_LIBCUDACXX_HIDE_FROM_ABI bool isnan(__half __v)
-{
-  return _CUDA_VSTD::__constexpr_isnan(__v);
-}
-
-_LIBCUDACXX_HIDE_FROM_ABI bool __constexpr_isinf(__half __x) noexcept
-{
-#  if _CCCL_STD_VER >= 2020 && _CCCL_CUDACC_BELOW(12, 3)
-  // this is a workaround for nvbug 4362808
-  return !::__hisnan(__x) && ::__hisnan(__x - __x);
-#  else // ^^^ C++20 && below 12.3 ^^^ / vvv C++17 or 12.3+ vvv
-  return ::__hisinf(__x) != 0;
-#  endif // _CCCL_STD_VER <= 2017 || _CCCL_CUDACC_BELOW(12, 3)
-}
-
-_LIBCUDACXX_HIDE_FROM_ABI bool isinf(__half __v)
-{
-  return _CUDA_VSTD::__constexpr_isinf(__v);
-}
-
-_LIBCUDACXX_HIDE_FROM_ABI bool __constexpr_isfinite(__half __x) noexcept
-{
-  return !_CUDA_VSTD::__constexpr_isnan(__x) && !_CUDA_VSTD::__constexpr_isinf(__x);
-}
-
-_LIBCUDACXX_HIDE_FROM_ABI bool isfinite(__half __v)
-{
-  return _CUDA_VSTD::__constexpr_isfinite(__v);
-}
-
 _LIBCUDACXX_HIDE_FROM_ABI __half __constexpr_copysign(__half __x, __half __y) noexcept
 {
   return __float2half(::copysignf(__half2float(__x), __half2float(__y)));