From 5b63fe0423412a4a5663f70e0a11650eb75d1d2f Mon Sep 17 00:00:00 2001 From: David Bayer Date: Sun, 12 Jan 2025 18:02:22 +0100 Subject: [PATCH 1/7] implement `cuda::std::numeric_limits` for `__half` and `__nv_bfloat16` --- libcudacxx/include/cuda/std/limits | 196 ++++++++++++++- .../limits/is_specialized.pass.cpp | 7 + .../const_data_members.pass.cpp | 225 +++++------------- .../denorm_min.pass.cpp | 6 + .../numeric.limits.members/digits.pass.cpp | 7 +- .../numeric.limits.members/digits10.pass.cpp | 6 + .../numeric.limits.members/epsilon.pass.cpp | 6 + .../has_denorm.pass.cpp | 6 + .../has_denorm_loss.pass.cpp | 6 + .../has_infinity.pass.cpp | 6 + .../has_quiet_NaN.pass.cpp | 6 + .../has_signaling_NaN.pass.cpp | 6 + .../numeric.limits.members/infinity.pass.cpp | 16 +- .../is_bounded.pass.cpp | 6 + .../numeric.limits.members/is_exact.pass.cpp | 6 + .../numeric.limits.members/is_iec559.pass.cpp | 6 + .../is_integer.pass.cpp | 6 + .../numeric.limits.members/is_modulo.pass.cpp | 6 + .../numeric.limits.members/is_signed.pass.cpp | 6 + .../numeric.limits.members/lowest.pass.cpp | 7 + .../numeric.limits.members/max.pass.cpp | 6 + .../max_digits10.pass.cpp | 6 + .../max_exponent.pass.cpp | 6 + .../max_exponent10.pass.cpp | 6 + .../numeric.limits.members/min.pass.cpp | 6 + .../min_exponent.pass.cpp | 6 + .../min_exponent10.pass.cpp | 6 + .../numeric.limits.members/quiet_NaN.pass.cpp | 10 +- .../numeric.limits.members/radix.pass.cpp | 6 + .../round_error.pass.cpp | 6 + .../round_style.pass.cpp | 6 + .../signaling_NaN.pass.cpp | 10 +- .../tinyness_before.pass.cpp | 6 + .../numeric.limits.members/traps.pass.cpp | 6 + 34 files changed, 461 insertions(+), 173 deletions(-) diff --git a/libcudacxx/include/cuda/std/limits b/libcudacxx/include/cuda/std/limits index 98c63813b7b..abed2d1ad68 100644 --- a/libcudacxx/include/cuda/std/limits +++ b/libcudacxx/include/cuda/std/limits @@ -22,7 +22,10 @@ #endif // no system header #include -#include +#include +#include +#include +#include #include #include @@ -46,7 +49,36 @@ enum float_denorm_style denorm_present = 1 }; -template ::value> +enum class __numeric_limits_type +{ + __integral, + __bool, + __floating_point, + __other, +}; + +template +struct __numeric_limits_type_selector : integral_constant<__numeric_limits_type, __numeric_limits_type::__other> +{}; + +template +struct __numeric_limits_type_selector<_Tp, enable_if_t<_CCCL_TRAIT(is_integral, _Tp)>> + : integral_constant<__numeric_limits_type, __numeric_limits_type::__integral> +{}; + +template <> +struct __numeric_limits_type_selector + : integral_constant<__numeric_limits_type, __numeric_limits_type::__bool> +{}; + +template +struct __numeric_limits_type_selector< + _Tp, + enable_if_t<(_CCCL_TRAIT(is_floating_point, _Tp) || _CCCL_TRAIT(__is_extended_floating_point, _Tp))>> + : integral_constant<__numeric_limits_type, __numeric_limits_type::__floating_point> +{}; + +template ::value> class __numeric_limits_impl { public: @@ -135,7 +167,7 @@ struct __int_min<_Tp, __digits, false> }; template -class __numeric_limits_impl<_Tp, true> +class __numeric_limits_impl<_Tp, __numeric_limits_type::__integral> { public: using type = _Tp; @@ -212,7 +244,7 @@ public: }; template <> -class __numeric_limits_impl +class __numeric_limits_impl { public: using type = bool; @@ -286,7 +318,7 @@ public: }; template <> -class __numeric_limits_impl +class __numeric_limits_impl { public: using type = float; @@ -381,7 +413,7 @@ public: }; template <> -class __numeric_limits_impl +class __numeric_limits_impl { public: using type = double; @@ -476,7 +508,7 @@ public: }; template <> -class __numeric_limits_impl +class __numeric_limits_impl { #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE @@ -551,6 +583,156 @@ public: #endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE }; +#if _CCCL_HAS_NVFP16 +template <> +class __numeric_limits_impl<__half, __numeric_limits_type::__floating_point> +{ +public: + using type = __half; + + static constexpr bool is_specialized = true; + + static constexpr bool is_signed = true; + static constexpr int digits = 11; + static constexpr int digits10 = 3; + static constexpr int max_digits10 = 5; + _LIBCUDACXX_HIDE_FROM_ABI static constexpr type min() noexcept + { + return type(__half_raw{0x0400u}); + } + _LIBCUDACXX_HIDE_FROM_ABI static constexpr type max() noexcept + { + return type(__half_raw{0x7bffu}); + } + _LIBCUDACXX_HIDE_FROM_ABI static constexpr type lowest() noexcept + { + return type(__half_raw{0xfbffu}); + } + + static constexpr bool is_integer = false; + static constexpr bool is_exact = false; + static constexpr int radix = __FLT_RADIX__; + _LIBCUDACXX_HIDE_FROM_ABI static constexpr type epsilon() noexcept + { + return type(__half_raw{0x1400u}); + } + _LIBCUDACXX_HIDE_FROM_ABI static constexpr type round_error() noexcept + { + return type(__half_raw{0x3800u}); + } + + static constexpr int min_exponent = -13; + static constexpr int min_exponent10 = -4; + static constexpr int max_exponent = 16; + static constexpr int max_exponent10 = 4; + + static constexpr bool has_infinity = true; + static constexpr bool has_quiet_NaN = true; + static constexpr bool has_signaling_NaN = true; + static constexpr float_denorm_style has_denorm = denorm_present; + static constexpr bool has_denorm_loss = false; + _LIBCUDACXX_HIDE_FROM_ABI static constexpr type infinity() noexcept + { + return type(__half_raw{0x7c00u}); + } + _LIBCUDACXX_HIDE_FROM_ABI static constexpr type quiet_NaN() noexcept + { + return type(__half_raw{0x7e00u}); + } + _LIBCUDACXX_HIDE_FROM_ABI static constexpr type signaling_NaN() noexcept + { + return type(__half_raw{0x7d00u}); + } + _LIBCUDACXX_HIDE_FROM_ABI static constexpr type denorm_min() noexcept + { + return type(__half_raw{0x0001u}); + } + + static constexpr bool is_iec559 = true; + static constexpr bool is_bounded = true; + static constexpr bool is_modulo = false; + + static constexpr bool traps = false; + static constexpr bool tinyness_before = false; + static constexpr float_round_style round_style = round_to_nearest; +}; +#endif // _CCCL_HAS_NVFP16 + +#if _CCCL_HAS_NVBF16 +template <> +class __numeric_limits_impl<__nv_bfloat16, __numeric_limits_type::__floating_point> +{ +public: + using type = __nv_bfloat16; + + static constexpr bool is_specialized = true; + + static constexpr bool is_signed = true; + static constexpr int digits = 8; + static constexpr int digits10 = 2; + static constexpr int max_digits10 = 4; + _LIBCUDACXX_HIDE_FROM_ABI static constexpr type min() noexcept + { + return type(__nv_bfloat16_raw{0x0080u}); + } + _LIBCUDACXX_HIDE_FROM_ABI static constexpr type max() noexcept + { + return type(__nv_bfloat16_raw{0x7f7fu}); + } + _LIBCUDACXX_HIDE_FROM_ABI static constexpr type lowest() noexcept + { + return type(__nv_bfloat16_raw{0xff7fu}); + } + + static constexpr bool is_integer = false; + static constexpr bool is_exact = false; + static constexpr int radix = __FLT_RADIX__; + _LIBCUDACXX_HIDE_FROM_ABI static constexpr type epsilon() noexcept + { + return type(__nv_bfloat16_raw{0x3c00u}); + } + _LIBCUDACXX_HIDE_FROM_ABI static constexpr type round_error() noexcept + { + return type(__nv_bfloat16_raw{0x3f00u}); + } + + static constexpr int min_exponent = -125; + static constexpr int min_exponent10 = -37; + static constexpr int max_exponent = 128; + static constexpr int max_exponent10 = 38; + + static constexpr bool has_infinity = true; + static constexpr bool has_quiet_NaN = true; + static constexpr bool has_signaling_NaN = true; + static constexpr float_denorm_style has_denorm = denorm_present; + static constexpr bool has_denorm_loss = false; + _LIBCUDACXX_HIDE_FROM_ABI static constexpr type infinity() noexcept + { + return type(__nv_bfloat16_raw{0x7f80u}); + } + _LIBCUDACXX_HIDE_FROM_ABI static constexpr type quiet_NaN() noexcept + { + return type(__nv_bfloat16_raw{0x7fc0u}); + } + _LIBCUDACXX_HIDE_FROM_ABI static constexpr type signaling_NaN() noexcept + { + return type(__nv_bfloat16_raw{0x7fa0u}); + } + _LIBCUDACXX_HIDE_FROM_ABI static constexpr type denorm_min() noexcept + { + return type(__nv_bfloat16_raw{0x0001u}); + } + + static constexpr bool is_iec559 = true; + static constexpr bool is_bounded = true; + static constexpr bool is_modulo = false; + + static constexpr bool traps = false; + static constexpr bool tinyness_before = false; + static constexpr float_round_style round_style = round_to_nearest; +}; +#endif // _CCCL_HAS_NVBF16 + template class numeric_limits : public __numeric_limits_impl<_Tp> {}; diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/is_specialized.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/is_specialized.pass.cpp index 2ecd59004bb..18dd0e5b03d 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/is_specialized.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/is_specialized.pass.cpp @@ -68,6 +68,13 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16>(); +#endif // _CCCL_HAS_NVBF16 + static_assert(!cuda::std::numeric_limits>::is_specialized, "!cuda::std::numeric_limits >::is_specialized"); diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/const_data_members.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/const_data_members.pass.cpp index 8db1a9f5f0c..e435ed91129 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/const_data_members.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/const_data_members.pass.cpp @@ -42,173 +42,80 @@ template __host__ __device__ void test(T) {} -#define TEST_NUMERIC_LIMITS(type) \ - test(cuda::std::numeric_limits::is_specialized); \ - test(cuda::std::numeric_limits::digits); \ - test(cuda::std::numeric_limits::digits10); \ - test(cuda::std::numeric_limits::max_digits10); \ - test(cuda::std::numeric_limits::is_signed); \ - test(cuda::std::numeric_limits::is_integer); \ - test(cuda::std::numeric_limits::is_exact); \ - test(cuda::std::numeric_limits::radix); \ - test(cuda::std::numeric_limits::min_exponent); \ - test(cuda::std::numeric_limits::min_exponent10); \ - test(cuda::std::numeric_limits::max_exponent); \ - test(cuda::std::numeric_limits::max_exponent10); \ - test(cuda::std::numeric_limits::has_infinity); \ - test(cuda::std::numeric_limits::has_quiet_NaN); \ - test(cuda::std::numeric_limits::has_signaling_NaN); \ - test(cuda::std::numeric_limits::has_denorm); \ - test(cuda::std::numeric_limits::has_denorm_loss); \ - test(cuda::std::numeric_limits::is_iec559); \ - test(cuda::std::numeric_limits::is_bounded); \ - test(cuda::std::numeric_limits::is_modulo); \ - test(cuda::std::numeric_limits::traps); \ - test(cuda::std::numeric_limits::tinyness_before); \ - test(cuda::std::numeric_limits::round_style); +template +__host__ __device__ void test_type_helper() +{ + test(cuda::std::numeric_limits::is_specialized); + test(cuda::std::numeric_limits::digits); + test(cuda::std::numeric_limits::digits10); + test(cuda::std::numeric_limits::max_digits10); + test(cuda::std::numeric_limits::is_signed); + test(cuda::std::numeric_limits::is_integer); + test(cuda::std::numeric_limits::is_exact); + test(cuda::std::numeric_limits::radix); + test(cuda::std::numeric_limits::min_exponent); + test(cuda::std::numeric_limits::min_exponent10); + test(cuda::std::numeric_limits::max_exponent); + test(cuda::std::numeric_limits::max_exponent10); + test(cuda::std::numeric_limits::has_infinity); + test(cuda::std::numeric_limits::has_quiet_NaN); + test(cuda::std::numeric_limits::has_signaling_NaN); + test(cuda::std::numeric_limits::has_denorm); + test(cuda::std::numeric_limits::has_denorm_loss); + test(cuda::std::numeric_limits::is_iec559); + test(cuda::std::numeric_limits::is_bounded); + test(cuda::std::numeric_limits::is_modulo); + test(cuda::std::numeric_limits::traps); + test(cuda::std::numeric_limits::tinyness_before); + test(cuda::std::numeric_limits::round_style); +} + +template +__host__ __device__ void test_type() +{ + test_type_helper(); + test_type_helper(); + test_type_helper(); + test_type_helper(); +} struct other {}; int main(int, char**) { - // bool - TEST_NUMERIC_LIMITS(bool) - TEST_NUMERIC_LIMITS(const bool) - TEST_NUMERIC_LIMITS(volatile bool) - TEST_NUMERIC_LIMITS(const volatile bool) - - // char - TEST_NUMERIC_LIMITS(char) - TEST_NUMERIC_LIMITS(const char) - TEST_NUMERIC_LIMITS(volatile char) - TEST_NUMERIC_LIMITS(const volatile char) - - // signed char - TEST_NUMERIC_LIMITS(signed char) - TEST_NUMERIC_LIMITS(const signed char) - TEST_NUMERIC_LIMITS(volatile signed char) - TEST_NUMERIC_LIMITS(const volatile signed char) - - // unsigned char - TEST_NUMERIC_LIMITS(unsigned char) - TEST_NUMERIC_LIMITS(const unsigned char) - TEST_NUMERIC_LIMITS(volatile unsigned char) - TEST_NUMERIC_LIMITS(const volatile unsigned char) - - // wchar_t - TEST_NUMERIC_LIMITS(wchar_t) - TEST_NUMERIC_LIMITS(const wchar_t) - TEST_NUMERIC_LIMITS(volatile wchar_t) - TEST_NUMERIC_LIMITS(const volatile wchar_t) - -#if TEST_STD_VER > 2017 && defined(__cpp_char8_t) - // char8_t - TEST_NUMERIC_LIMITS(char8_t) - TEST_NUMERIC_LIMITS(const char8_t) - TEST_NUMERIC_LIMITS(volatile char8_t) - TEST_NUMERIC_LIMITS(const volatile char8_t) -#endif - - // char16_t - TEST_NUMERIC_LIMITS(char16_t) - TEST_NUMERIC_LIMITS(const char16_t) - TEST_NUMERIC_LIMITS(volatile char16_t) - TEST_NUMERIC_LIMITS(const volatile char16_t) - - // char32_t - TEST_NUMERIC_LIMITS(char32_t) - TEST_NUMERIC_LIMITS(const char32_t) - TEST_NUMERIC_LIMITS(volatile char32_t) - TEST_NUMERIC_LIMITS(const volatile char32_t) - - // short - TEST_NUMERIC_LIMITS(short) - TEST_NUMERIC_LIMITS(const short) - TEST_NUMERIC_LIMITS(volatile short) - TEST_NUMERIC_LIMITS(const volatile short) - - // int - TEST_NUMERIC_LIMITS(int) - TEST_NUMERIC_LIMITS(const int) - TEST_NUMERIC_LIMITS(volatile int) - TEST_NUMERIC_LIMITS(const volatile int) - - // long - TEST_NUMERIC_LIMITS(long) - TEST_NUMERIC_LIMITS(const long) - TEST_NUMERIC_LIMITS(volatile long) - TEST_NUMERIC_LIMITS(const volatile long) - -#ifndef _LIBCUDACXX_HAS_NO_INT128 - TEST_NUMERIC_LIMITS(__int128_t) - TEST_NUMERIC_LIMITS(const __int128_t) - TEST_NUMERIC_LIMITS(volatile __int128_t) - TEST_NUMERIC_LIMITS(const volatile __int128_t) -#endif - - // long long - TEST_NUMERIC_LIMITS(long long) - TEST_NUMERIC_LIMITS(const long long) - TEST_NUMERIC_LIMITS(volatile long long) - TEST_NUMERIC_LIMITS(const volatile long long) - - // unsigned short - TEST_NUMERIC_LIMITS(unsigned short) - TEST_NUMERIC_LIMITS(const unsigned short) - TEST_NUMERIC_LIMITS(volatile unsigned short) - TEST_NUMERIC_LIMITS(const volatile unsigned short) - - // unsigned int - TEST_NUMERIC_LIMITS(unsigned int) - TEST_NUMERIC_LIMITS(const unsigned int) - TEST_NUMERIC_LIMITS(volatile unsigned int) - TEST_NUMERIC_LIMITS(const volatile unsigned int) - - // unsigned long - TEST_NUMERIC_LIMITS(unsigned long) - TEST_NUMERIC_LIMITS(const unsigned long) - TEST_NUMERIC_LIMITS(volatile unsigned long) - TEST_NUMERIC_LIMITS(const volatile unsigned long) - - // unsigned long long - TEST_NUMERIC_LIMITS(unsigned long long) - TEST_NUMERIC_LIMITS(const unsigned long long) - TEST_NUMERIC_LIMITS(volatile unsigned long long) - TEST_NUMERIC_LIMITS(const volatile unsigned long long) - + test_type(); + test_type(); + test_type(); + test_type(); + test_type(); +#if TEST_STD_VER >= 2020 && defined(__cpp_char8_t) + test_type(); +#endif // TEST_STD_VER >= 2020 && defined(__cpp_char8_t) + test_type(); + test_type(); + test_type(); + test_type(); + test_type(); + test_type(); + test_type(); + test_type(); + test_type(); + test_type(); #ifndef _LIBCUDACXX_HAS_NO_INT128 - TEST_NUMERIC_LIMITS(__uint128_t) - TEST_NUMERIC_LIMITS(const __uint128_t) - TEST_NUMERIC_LIMITS(volatile __uint128_t) - TEST_NUMERIC_LIMITS(const volatile __uint128_t) -#endif - - // float - TEST_NUMERIC_LIMITS(float) - TEST_NUMERIC_LIMITS(const float) - TEST_NUMERIC_LIMITS(volatile float) - TEST_NUMERIC_LIMITS(const volatile float) - - // double - TEST_NUMERIC_LIMITS(double) - TEST_NUMERIC_LIMITS(const double) - TEST_NUMERIC_LIMITS(volatile double) - TEST_NUMERIC_LIMITS(const volatile double) - + test_type<__int128_t>(); +#endif // _LIBCUDACXX_HAS_NO_INT128 + test_type(); + test_type(); #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE - // long double - TEST_NUMERIC_LIMITS(long double) - TEST_NUMERIC_LIMITS(const long double) - TEST_NUMERIC_LIMITS(volatile long double) - TEST_NUMERIC_LIMITS(const volatile long double) -#endif - - // other - TEST_NUMERIC_LIMITS(other) - TEST_NUMERIC_LIMITS(const other) - TEST_NUMERIC_LIMITS(volatile other) - TEST_NUMERIC_LIMITS(const volatile other) + test_type(); +#endif // _LIBCUDACXX_HAS_NO_LONG_DOUBLE +#if _CCCL_HAS_NVFP16 + test_type<__half>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test_type<__nv_bfloat16>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/denorm_min.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/denorm_min.pass.cpp index 730adc30d36..673dd998224 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/denorm_min.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/denorm_min.pass.cpp @@ -65,6 +65,12 @@ int main(int, char**) test(LDBL_TRUE_MIN); # endif #endif +#if _CCCL_HAS_NVFP16 + test<__half>(CUDART_MIN_DENORM_FP16); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16>(CUDART_MIN_DENORM_BF16); +#endif // _CCCL_HAS_NVBF16 #if !defined(__FLT_DENORM_MIN__) && !defined(FLT_TRUE_MIN) # error Test has no expected values for floating point types #endif diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/digits.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/digits.pass.cpp index 63ecf93515f..8748040024e 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/digits.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/digits.pass.cpp @@ -55,6 +55,11 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif - +#if _CCCL_HAS_NVFP16 + test<__half, 11>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, 8>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/digits10.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/digits10.pass.cpp index 3295686ea49..603472b2fe1 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/digits10.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/digits10.pass.cpp @@ -59,6 +59,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half, 3>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, 2>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/epsilon.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/epsilon.pass.cpp index 5bc22e7f5f2..1e572f5a416 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/epsilon.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/epsilon.pass.cpp @@ -56,6 +56,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(LDBL_EPSILON); #endif +#if _CCCL_HAS_NVFP16 + test<__half>(0.0009765625); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16>(0.0078125); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_denorm.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_denorm.pass.cpp index e62208d7e3b..cb7048710e5 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_denorm.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_denorm.pass.cpp @@ -54,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half, cuda::std::denorm_present>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, cuda::std::denorm_present>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_denorm_loss.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_denorm_loss.pass.cpp index 3a665fe2c9b..0ba20db0291 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_denorm_loss.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_denorm_loss.pass.cpp @@ -54,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half, false>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, false>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_infinity.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_infinity.pass.cpp index be41dabb02c..ce21ccdd24c 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_infinity.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_infinity.pass.cpp @@ -54,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half, true>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, true>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_quiet_NaN.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_quiet_NaN.pass.cpp index 2d13db35438..323526ac0ca 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_quiet_NaN.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_quiet_NaN.pass.cpp @@ -54,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half, true>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, true>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_signaling_NaN.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_signaling_NaN.pass.cpp index d5cf5096bb7..7dd2825f3bd 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_signaling_NaN.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_signaling_NaN.pass.cpp @@ -54,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half, true>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, true>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/infinity.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/infinity.pass.cpp index 2d1c29f6f31..36e1a01f960 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/infinity.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/infinity.pass.cpp @@ -62,6 +62,12 @@ int main(int, char**) # ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(1. / 0.); # endif +# if _CCCL_HAS_NVFP16 + test<__half>(1.0 / 0.0); +# endif // _CCCL_HAS_NVFP16 +# if _CCCL_HAS_NVBF16 + test<__nv_bfloat16>(1.0 / 0.0); +# endif // _CCCL_HAS_NVBF16 // MSVC has issues with producing INF with divisions by zero. #else test(INFINITY); @@ -69,11 +75,13 @@ int main(int, char**) # ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(INFINITY); # endif +# if _CCCL_HAS_NVFP16 + test<__half>(INFINITY); +# endif // _CCCL_HAS_NVFP16 +# if _CCCL_HAS_NVBF16 + test<__nv_bfloat16>(INFINITY); +# endif // _CCCL_HAS_NVBF16 #endif return 0; } - -#ifndef TEST_COMPILER_NVRTC -float zero = 0; -#endif diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_bounded.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_bounded.pass.cpp index 2dd4bd94fbc..e3311f600f1 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_bounded.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_bounded.pass.cpp @@ -54,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half, true>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, true>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_exact.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_exact.pass.cpp index be45efae70c..2975cf0a3a6 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_exact.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_exact.pass.cpp @@ -54,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half, false>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, false>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_iec559.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_iec559.pass.cpp index 6221cd6ed59..4ce90bfeee4 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_iec559.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_iec559.pass.cpp @@ -54,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif // _LIBCUDACXX_HAS_NO_LONG_DOUBLE +#if _CCCL_HAS_NVFP16 + test<__half, true>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, true>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_integer.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_integer.pass.cpp index 3d166f31f28..abe1115c975 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_integer.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_integer.pass.cpp @@ -54,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half, false>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, false>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_modulo.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_modulo.pass.cpp index 7b1adabf0c7..b8ab8352a61 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_modulo.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_modulo.pass.cpp @@ -54,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half, false>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, false>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_signed.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_signed.pass.cpp index d7f98766343..e46eba94577 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_signed.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_signed.pass.cpp @@ -54,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half, true>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, true>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/lowest.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/lowest.pass.cpp index 6fec93e4a3d..43d1070446d 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/lowest.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/lowest.pass.cpp @@ -35,6 +35,7 @@ int main(int, char**) { test(false); test(CHAR_MIN); + test(SCHAR_MIN); test(0); #ifndef TEST_COMPILER_NVRTC @@ -64,6 +65,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(-LDBL_MAX); #endif +#if _CCCL_HAS_NVFP16 + test<__half>(-65504.0); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16>(-3.3895313892515355e+38); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max.pass.cpp index 67c94051729..ac0c0116835 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max.pass.cpp @@ -64,6 +64,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(LDBL_MAX); #endif +#if _CCCL_HAS_NVFP16 + test<__half>(65504.0); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16>(3.3895313892515355e+38); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_digits10.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_digits10.pass.cpp index cd5892e6c8c..b22863f831a 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_digits10.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_digits10.pass.cpp @@ -55,6 +55,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half, 5>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, 4>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_exponent.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_exponent.pass.cpp index aeb9189d315..fd5c08c95f0 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_exponent.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_exponent.pass.cpp @@ -55,6 +55,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half, 16>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, 128>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_exponent10.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_exponent10.pass.cpp index ca0eb2917f6..b5b3d2d1db9 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_exponent10.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_exponent10.pass.cpp @@ -55,6 +55,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half, 4>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, 38>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min.pass.cpp index 53d196d2a51..6c7a56da6b5 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min.pass.cpp @@ -65,6 +65,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(LDBL_MIN); #endif +#if _CCCL_HAS_NVFP16 + test<__half>(6.103515625e-05); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16>(1.17549435082228750796873653722e-38); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min_exponent.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min_exponent.pass.cpp index b075bcff87d..5a857c117c7 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min_exponent.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min_exponent.pass.cpp @@ -55,6 +55,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half, -13>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, -125>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min_exponent10.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min_exponent10.pass.cpp index c787cf4caab..566f7d0a2ef 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min_exponent10.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min_exponent10.pass.cpp @@ -55,6 +55,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half, -4>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, -37>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/quiet_NaN.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/quiet_NaN.pass.cpp index ce38b3ed60d..872d8ade1aa 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/quiet_NaN.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/quiet_NaN.pass.cpp @@ -38,7 +38,9 @@ __host__ __device__ void test_imp(cuda::std::false_type) template __host__ __device__ inline void test() { - test_imp(cuda::std::is_floating_point()); + constexpr bool is_float = cuda::std::is_floating_point::value || cuda::std::__is_extended_floating_point::value; + + test_imp(cuda::std::integral_constant{}); } int main(int, char**) @@ -72,6 +74,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/radix.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/radix.pass.cpp index 5a709b3aefc..436838a9df7 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/radix.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/radix.pass.cpp @@ -55,6 +55,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half, FLT_RADIX>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, FLT_RADIX>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/round_error.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/round_error.pass.cpp index 01d10e80fb9..6c774728388 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/round_error.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/round_error.pass.cpp @@ -56,6 +56,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(0.5); #endif +#if _CCCL_HAS_NVFP16 + test<__half>(0.5); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16>(0.5); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/round_style.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/round_style.pass.cpp index 7a7099662f0..c9d9a06b595 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/round_style.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/round_style.pass.cpp @@ -54,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half, cuda::std::round_to_nearest>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, cuda::std::round_to_nearest>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/signaling_NaN.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/signaling_NaN.pass.cpp index 164d54c5741..2ad1327643e 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/signaling_NaN.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/signaling_NaN.pass.cpp @@ -38,7 +38,9 @@ __host__ __device__ void test_imp(cuda::std::false_type) template __host__ __device__ inline void test() { - test_imp(cuda::std::is_floating_point()); + constexpr bool is_float = cuda::std::is_floating_point::value || cuda::std::__is_extended_floating_point::value; + + test_imp(cuda::std::integral_constant{}); } int main(int, char**) @@ -72,6 +74,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/tinyness_before.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/tinyness_before.pass.cpp index 70d832dc547..e99510cc508 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/tinyness_before.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/tinyness_before.pass.cpp @@ -54,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half, false>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, false>(); +#endif // _CCCL_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/traps.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/traps.pass.cpp index 5c66acb56ce..12808cec121 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/traps.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/traps.pass.cpp @@ -60,6 +60,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if _CCCL_HAS_NVFP16 + test<__half, false>(); +#endif // _CCCL_HAS_NVFP16 +#if _CCCL_HAS_NVBF16 + test<__nv_bfloat16, false>(); +#endif // _CCCL_HAS_NVBF16 return 0; } From aca1e83965a78267458e51293d3b03c600919fa2 Mon Sep 17 00:00:00 2001 From: David Bayer Date: Mon, 13 Jan 2025 08:43:40 +0100 Subject: [PATCH 2/7] use `_LIBCUDACXX_HAS_type` instead of `_CCCL_HAS_type` --- libcudacxx/include/cuda/std/limits | 8 ++++---- .../limits/is_specialized.pass.cpp | 8 ++++---- .../const_data_members.pass.cpp | 8 ++++---- .../numeric.limits.members/denorm_min.pass.cpp | 8 ++++---- .../numeric.limits.members/digits.pass.cpp | 8 ++++---- .../numeric.limits.members/digits10.pass.cpp | 8 ++++---- .../numeric.limits.members/epsilon.pass.cpp | 8 ++++---- .../numeric.limits.members/has_denorm.pass.cpp | 8 ++++---- .../has_denorm_loss.pass.cpp | 8 ++++---- .../numeric.limits.members/has_infinity.pass.cpp | 8 ++++---- .../has_quiet_NaN.pass.cpp | 8 ++++---- .../has_signaling_NaN.pass.cpp | 8 ++++---- .../numeric.limits.members/infinity.pass.cpp | 16 ++++++++-------- .../numeric.limits.members/is_bounded.pass.cpp | 8 ++++---- .../numeric.limits.members/is_exact.pass.cpp | 8 ++++---- .../numeric.limits.members/is_iec559.pass.cpp | 8 ++++---- .../numeric.limits.members/is_integer.pass.cpp | 8 ++++---- .../numeric.limits.members/is_modulo.pass.cpp | 8 ++++---- .../numeric.limits.members/is_signed.pass.cpp | 8 ++++---- .../numeric.limits.members/lowest.pass.cpp | 8 ++++---- .../limits/numeric.limits.members/max.pass.cpp | 8 ++++---- .../numeric.limits.members/max_digits10.pass.cpp | 8 ++++---- .../numeric.limits.members/max_exponent.pass.cpp | 8 ++++---- .../max_exponent10.pass.cpp | 8 ++++---- .../limits/numeric.limits.members/min.pass.cpp | 8 ++++---- .../numeric.limits.members/min_exponent.pass.cpp | 8 ++++---- .../min_exponent10.pass.cpp | 8 ++++---- .../numeric.limits.members/quiet_NaN.pass.cpp | 8 ++++---- .../limits/numeric.limits.members/radix.pass.cpp | 8 ++++---- .../numeric.limits.members/round_error.pass.cpp | 8 ++++---- .../numeric.limits.members/round_style.pass.cpp | 8 ++++---- .../signaling_NaN.pass.cpp | 8 ++++---- .../tinyness_before.pass.cpp | 8 ++++---- .../limits/numeric.limits.members/traps.pass.cpp | 8 ++++---- 34 files changed, 140 insertions(+), 140 deletions(-) diff --git a/libcudacxx/include/cuda/std/limits b/libcudacxx/include/cuda/std/limits index abed2d1ad68..b258cf7af04 100644 --- a/libcudacxx/include/cuda/std/limits +++ b/libcudacxx/include/cuda/std/limits @@ -583,7 +583,7 @@ public: #endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE }; -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) template <> class __numeric_limits_impl<__half, __numeric_limits_type::__floating_point> { @@ -656,9 +656,9 @@ public: static constexpr bool tinyness_before = false; static constexpr float_round_style round_style = round_to_nearest; }; -#endif // _CCCL_HAS_NVFP16 +#endif // _LIBCUDACXX_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#if defined(_LIBCUDACXX_HAS_NVBF16) template <> class __numeric_limits_impl<__nv_bfloat16, __numeric_limits_type::__floating_point> { @@ -731,7 +731,7 @@ public: static constexpr bool tinyness_before = false; static constexpr float_round_style round_style = round_to_nearest; }; -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 template class numeric_limits : public __numeric_limits_impl<_Tp> diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/is_specialized.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/is_specialized.pass.cpp index 18dd0e5b03d..7113c0e2772 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/is_specialized.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/is_specialized.pass.cpp @@ -68,12 +68,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 static_assert(!cuda::std::numeric_limits>::is_specialized, "!cuda::std::numeric_limits >::is_specialized"); diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/const_data_members.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/const_data_members.pass.cpp index e435ed91129..769080cff83 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/const_data_members.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/const_data_members.pass.cpp @@ -110,12 +110,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test_type(); #endif // _LIBCUDACXX_HAS_NO_LONG_DOUBLE -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test_type<__half>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test_type<__nv_bfloat16>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/denorm_min.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/denorm_min.pass.cpp index 673dd998224..182ae69d327 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/denorm_min.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/denorm_min.pass.cpp @@ -65,12 +65,12 @@ int main(int, char**) test(LDBL_TRUE_MIN); # endif #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half>(CUDART_MIN_DENORM_FP16); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16>(CUDART_MIN_DENORM_BF16); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 #if !defined(__FLT_DENORM_MIN__) && !defined(FLT_TRUE_MIN) # error Test has no expected values for floating point types #endif diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/digits.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/digits.pass.cpp index 8748040024e..efce1ccf678 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/digits.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/digits.pass.cpp @@ -55,11 +55,11 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, 11>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, 8>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/digits10.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/digits10.pass.cpp index 603472b2fe1..32990ece4b1 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/digits10.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/digits10.pass.cpp @@ -59,12 +59,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, 3>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, 2>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/epsilon.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/epsilon.pass.cpp index 1e572f5a416..63b0bc7b8b4 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/epsilon.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/epsilon.pass.cpp @@ -56,12 +56,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(LDBL_EPSILON); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half>(0.0009765625); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16>(0.0078125); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_denorm.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_denorm.pass.cpp index cb7048710e5..5a0a05ab73b 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_denorm.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_denorm.pass.cpp @@ -54,12 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, cuda::std::denorm_present>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, cuda::std::denorm_present>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_denorm_loss.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_denorm_loss.pass.cpp index 0ba20db0291..450e51b8111 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_denorm_loss.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_denorm_loss.pass.cpp @@ -54,12 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, false>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, false>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_infinity.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_infinity.pass.cpp index ce21ccdd24c..646f5e20160 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_infinity.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_infinity.pass.cpp @@ -54,12 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, true>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, true>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_quiet_NaN.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_quiet_NaN.pass.cpp index 323526ac0ca..626b4110695 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_quiet_NaN.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_quiet_NaN.pass.cpp @@ -54,12 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, true>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, true>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_signaling_NaN.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_signaling_NaN.pass.cpp index 7dd2825f3bd..20cd04d107e 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_signaling_NaN.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_signaling_NaN.pass.cpp @@ -54,12 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, true>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, true>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/infinity.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/infinity.pass.cpp index 36e1a01f960..a13ff0692df 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/infinity.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/infinity.pass.cpp @@ -62,12 +62,12 @@ int main(int, char**) # ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(1. / 0.); # endif -# if _CCCL_HAS_NVFP16 +# if defined(_LIBCUDACXX_HAS_NVFP16) test<__half>(1.0 / 0.0); -# endif // _CCCL_HAS_NVFP16 -# if _CCCL_HAS_NVBF16 +# endif // _LIBCUDACXX_HAS_NVFP16 +# if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16>(1.0 / 0.0); -# endif // _CCCL_HAS_NVBF16 +# endif // _LIBCUDACXX_HAS_NVBF16 // MSVC has issues with producing INF with divisions by zero. #else test(INFINITY); @@ -75,12 +75,12 @@ int main(int, char**) # ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(INFINITY); # endif -# if _CCCL_HAS_NVFP16 +# if defined(_LIBCUDACXX_HAS_NVFP16) test<__half>(INFINITY); -# endif // _CCCL_HAS_NVFP16 -# if _CCCL_HAS_NVBF16 +# endif // _LIBCUDACXX_HAS_NVFP16 +# if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16>(INFINITY); -# endif // _CCCL_HAS_NVBF16 +# endif // _LIBCUDACXX_HAS_NVBF16 #endif return 0; diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_bounded.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_bounded.pass.cpp index e3311f600f1..9e671c5d905 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_bounded.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_bounded.pass.cpp @@ -54,12 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, true>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, true>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_exact.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_exact.pass.cpp index 2975cf0a3a6..cfc9a6cab90 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_exact.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_exact.pass.cpp @@ -54,12 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, false>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, false>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_iec559.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_iec559.pass.cpp index 4ce90bfeee4..945347ff4b5 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_iec559.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_iec559.pass.cpp @@ -54,12 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif // _LIBCUDACXX_HAS_NO_LONG_DOUBLE -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, true>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, true>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_integer.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_integer.pass.cpp index abe1115c975..65dd98fdb04 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_integer.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_integer.pass.cpp @@ -54,12 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, false>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, false>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_modulo.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_modulo.pass.cpp index b8ab8352a61..6d82269e1c8 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_modulo.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_modulo.pass.cpp @@ -54,12 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, false>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, false>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_signed.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_signed.pass.cpp index e46eba94577..eb39869bf24 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_signed.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_signed.pass.cpp @@ -54,12 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, true>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, true>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/lowest.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/lowest.pass.cpp index 43d1070446d..60cec9e8533 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/lowest.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/lowest.pass.cpp @@ -65,12 +65,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(-LDBL_MAX); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half>(-65504.0); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16>(-3.3895313892515355e+38); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max.pass.cpp index ac0c0116835..fa06dab8963 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max.pass.cpp @@ -64,12 +64,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(LDBL_MAX); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half>(65504.0); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16>(3.3895313892515355e+38); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_digits10.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_digits10.pass.cpp index b22863f831a..92b3d13ea61 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_digits10.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_digits10.pass.cpp @@ -55,12 +55,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, 5>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, 4>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_exponent.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_exponent.pass.cpp index fd5c08c95f0..81d5ae07795 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_exponent.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_exponent.pass.cpp @@ -55,12 +55,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, 16>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, 128>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_exponent10.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_exponent10.pass.cpp index b5b3d2d1db9..4c426b37460 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_exponent10.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_exponent10.pass.cpp @@ -55,12 +55,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, 4>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, 38>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min.pass.cpp index 6c7a56da6b5..14ebad929c3 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min.pass.cpp @@ -65,12 +65,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(LDBL_MIN); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half>(6.103515625e-05); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16>(1.17549435082228750796873653722e-38); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min_exponent.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min_exponent.pass.cpp index 5a857c117c7..e3150f8dc8e 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min_exponent.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min_exponent.pass.cpp @@ -55,12 +55,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, -13>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, -125>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min_exponent10.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min_exponent10.pass.cpp index 566f7d0a2ef..cbca8e04171 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min_exponent10.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min_exponent10.pass.cpp @@ -55,12 +55,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, -4>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, -37>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/quiet_NaN.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/quiet_NaN.pass.cpp index 872d8ade1aa..74e7f427941 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/quiet_NaN.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/quiet_NaN.pass.cpp @@ -74,12 +74,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/radix.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/radix.pass.cpp index 436838a9df7..9765db6f760 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/radix.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/radix.pass.cpp @@ -55,12 +55,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, FLT_RADIX>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, FLT_RADIX>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/round_error.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/round_error.pass.cpp index 6c774728388..ffc5d738109 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/round_error.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/round_error.pass.cpp @@ -56,12 +56,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(0.5); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half>(0.5); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16>(0.5); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/round_style.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/round_style.pass.cpp index c9d9a06b595..3fb436381a7 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/round_style.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/round_style.pass.cpp @@ -54,12 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, cuda::std::round_to_nearest>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, cuda::std::round_to_nearest>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/signaling_NaN.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/signaling_NaN.pass.cpp index 2ad1327643e..69ba66038de 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/signaling_NaN.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/signaling_NaN.pass.cpp @@ -74,12 +74,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/tinyness_before.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/tinyness_before.pass.cpp index e99510cc508..70cde2711a1 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/tinyness_before.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/tinyness_before.pass.cpp @@ -54,12 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, false>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, false>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/traps.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/traps.pass.cpp index 12808cec121..7dd7eee68cc 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/traps.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/traps.pass.cpp @@ -60,12 +60,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif -#if _CCCL_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVFP16) test<__half, false>(); -#endif // _CCCL_HAS_NVFP16 -#if _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) test<__nv_bfloat16, false>(); -#endif // _CCCL_HAS_NVBF16 +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } From a9e4daf25eec298e4397467df947fb8d4af9c845 Mon Sep 17 00:00:00 2001 From: David Bayer Date: Mon, 13 Jan 2025 13:00:41 +0100 Subject: [PATCH 3/7] fix review --- libcudacxx/include/cuda/std/limits | 50 ++++++++++++++++++------------ 1 file changed, 30 insertions(+), 20 deletions(-) diff --git a/libcudacxx/include/cuda/std/limits b/libcudacxx/include/cuda/std/limits index b258cf7af04..79fa35bf860 100644 --- a/libcudacxx/include/cuda/std/limits +++ b/libcudacxx/include/cuda/std/limits @@ -57,28 +57,38 @@ enum class __numeric_limits_type __other, }; -template -struct __numeric_limits_type_selector : integral_constant<__numeric_limits_type, __numeric_limits_type::__other> -{}; - template -struct __numeric_limits_type_selector<_Tp, enable_if_t<_CCCL_TRAIT(is_integral, _Tp)>> - : integral_constant<__numeric_limits_type, __numeric_limits_type::__integral> -{}; - -template <> -struct __numeric_limits_type_selector - : integral_constant<__numeric_limits_type, __numeric_limits_type::__bool> -{}; - -template -struct __numeric_limits_type_selector< - _Tp, - enable_if_t<(_CCCL_TRAIT(is_floating_point, _Tp) || _CCCL_TRAIT(__is_extended_floating_point, _Tp))>> - : integral_constant<__numeric_limits_type, __numeric_limits_type::__floating_point> -{}; +constexpr __numeric_limits_type __make_numeric_limits_type() +{ +#if !defined(_CCCL_NO_IF_CONSTEXPR) + _CCCL_IF_CONSTEXPR (_CCCL_TRAIT(is_same, _Tp, bool)) + { + return __numeric_limits_type::__bool; + } + else _CCCL_IF_CONSTEXPR (_CCCL_TRAIT(is_integral, _Tp)) + { + return __numeric_limits_type::__integral; + } + else _CCCL_IF_CONSTEXPR (_CCCL_TRAIT(is_floating_point, _Tp) || _CCCL_TRAIT(__is_extended_floating_point, _Tp)) + { + return __numeric_limits_type::__floating_point; + } + else + { + return __numeric_limits_type::__other; + } +#else // ^^^ !_CCCL_NO_IF_CONSTEXPR ^^^ // vvv _CCCL_NO_IF_CONSTEXPR vvv + return _CCCL_TRAIT(is_same, _Tp, bool) + ? __numeric_limits_type::__bool + : (_CCCL_TRAIT(is_integral, _Tp) + ? __numeric_limits_type::__integral + : (_CCCL_TRAIT(is_floating_point, _Tp) || _CCCL_TRAIT(__is_extended_floating_point, _Tp) + ? __numeric_limits_type::__floating_point + : __numeric_limits_type::__other)); +#endif // _CCCL_NO_IF_CONSTEXPR +} -template ::value> +template ()> class __numeric_limits_impl { public: From 9d59ab9ed4d8937090005072c753aab1b571172d Mon Sep 17 00:00:00 2001 From: David Bayer Date: Mon, 13 Jan 2025 13:12:05 +0100 Subject: [PATCH 4/7] fix mdspan test --- .../std/containers/views/mdspan/my_int.hpp | 15 +++++++++++++++ 1 file changed, 15 insertions(+) diff --git a/libcudacxx/test/libcudacxx/std/containers/views/mdspan/my_int.hpp b/libcudacxx/test/libcudacxx/std/containers/views/mdspan/my_int.hpp index 4f27784cd61..df34fa1d42e 100644 --- a/libcudacxx/test/libcudacxx/std/containers/views/mdspan/my_int.hpp +++ b/libcudacxx/test/libcudacxx/std/containers/views/mdspan/my_int.hpp @@ -1,6 +1,9 @@ #ifndef _MY_INT_HPP #define _MY_INT_HPP +#include +#include + #include "test_macros.h" struct my_int_non_convertible; @@ -22,6 +25,10 @@ template <> struct cuda::std::is_integral : cuda::std::true_type {}; +template <> +class cuda::std::numeric_limits : public cuda::std::numeric_limits +{}; + // Wrapper type that's not implicitly convertible struct my_int_non_convertible @@ -43,6 +50,10 @@ template <> struct cuda::std::is_integral : cuda::std::true_type {}; +template <> +class cuda::std::numeric_limits : public cuda::std::numeric_limits +{}; + // Wrapper type that's not nothrow-constructible struct my_int_non_nothrow_constructible @@ -62,4 +73,8 @@ template <> struct cuda::std::is_integral : cuda::std::true_type {}; +template <> +class cuda::std::numeric_limits : public cuda::std::numeric_limits +{}; + #endif From e7071c5db906e5c56f4cde3cd9a70cc8ab44cf1e Mon Sep 17 00:00:00 2001 From: David Bayer Date: Mon, 13 Jan 2025 13:15:39 +0100 Subject: [PATCH 5/7] fix denorm_min test for nvrtc --- .../limits/numeric.limits.members/denorm_min.pass.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/denorm_min.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/denorm_min.pass.cpp index 182ae69d327..36c92718ec7 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/denorm_min.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/denorm_min.pass.cpp @@ -66,10 +66,10 @@ int main(int, char**) # endif #endif #if defined(_LIBCUDACXX_HAS_NVFP16) - test<__half>(CUDART_MIN_DENORM_FP16); + test<__half>(5.9604644775390625e-08); #endif // _LIBCUDACXX_HAS_NVFP16 #if defined(_LIBCUDACXX_HAS_NVBF16) - test<__nv_bfloat16>(CUDART_MIN_DENORM_BF16); + test<__nv_bfloat16>(9.18354961579912115600575419705e-41); #endif // _LIBCUDACXX_HAS_NVBF16 #if !defined(__FLT_DENORM_MIN__) && !defined(FLT_TRUE_MIN) # error Test has no expected values for floating point types From a512c51e858d07b9689ef94a82749c14282fd474 Mon Sep 17 00:00:00 2001 From: David Bayer Date: Mon, 13 Jan 2025 14:07:34 +0100 Subject: [PATCH 6/7] fix `__make_numeric_limits_type` --- libcudacxx/include/cuda/std/limits | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libcudacxx/include/cuda/std/limits b/libcudacxx/include/cuda/std/limits index 79fa35bf860..ad529f2082b 100644 --- a/libcudacxx/include/cuda/std/limits +++ b/libcudacxx/include/cuda/std/limits @@ -58,7 +58,7 @@ enum class __numeric_limits_type }; template -constexpr __numeric_limits_type __make_numeric_limits_type() +_LIBCUDACXX_HIDE_FROM_ABI constexpr __numeric_limits_type __make_numeric_limits_type() { #if !defined(_CCCL_NO_IF_CONSTEXPR) _CCCL_IF_CONSTEXPR (_CCCL_TRAIT(is_same, _Tp, bool)) From 1e8a58302ea47631b5786a41f0fe71f0e791828b Mon Sep 17 00:00:00 2001 From: David Bayer Date: Mon, 13 Jan 2025 15:44:59 +0100 Subject: [PATCH 7/7] disable operators and conversions for extended floating point types --- .../limits/numeric.limits.members/common.h | 41 +++++++++++++++++++ .../denorm_min.pass.cpp | 13 +++--- .../numeric.limits.members/epsilon.pass.cpp | 13 +++--- .../numeric.limits.members/infinity.pass.cpp | 18 ++++---- .../numeric.limits.members/lowest.pass.cpp | 13 +++--- .../numeric.limits.members/max.pass.cpp | 13 +++--- .../numeric.limits.members/min.pass.cpp | 13 +++--- .../round_error.pass.cpp | 13 +++--- 8 files changed, 93 insertions(+), 44 deletions(-) create mode 100644 libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/common.h diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/common.h b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/common.h new file mode 100644 index 00000000000..15b48836839 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/common.h @@ -0,0 +1,41 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef NUMERIC_LIMITS_MEMBERS_COMMON_H +#define NUMERIC_LIMITS_MEMBERS_COMMON_H + +// Disable all the extended floating point operations and conversions +#define __CUDA_NO_HALF_CONVERSIONS__ 1 +#define __CUDA_NO_HALF_OPERATORS__ 1 +#define __CUDA_NO_BFLOAT16_CONVERSIONS__ 1 +#define __CUDA_NO_BFLOAT16_OPERATORS__ 1 + +#include + +template +__host__ __device__ bool float_eq(T x, T y) +{ + return x == y; +} + +#if defined(_LIBCUDACXX_HAS_NVFP16) +__host__ __device__ inline bool float_eq(__half x, __half y) +{ + return __heq(x, y); +} +#endif // _LIBCUDACXX_HAS_NVFP16 + +#if defined(_LIBCUDACXX_HAS_NVBF16) +__host__ __device__ inline bool float_eq(__nv_bfloat16 x, __nv_bfloat16 y) +{ + return __heq(x, y); +} +#endif // _LIBCUDACXX_HAS_NVBF16 + +#endif // NUMERIC_LIMITS_MEMBERS_COMMON_H diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/denorm_min.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/denorm_min.pass.cpp index 36c92718ec7..cc64ed14686 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/denorm_min.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/denorm_min.pass.cpp @@ -14,15 +14,16 @@ #include #include +#include "common.h" #include "test_macros.h" template __host__ __device__ void test(T expected) { - assert(cuda::std::numeric_limits::denorm_min() == expected); - assert(cuda::std::numeric_limits::denorm_min() == expected); - assert(cuda::std::numeric_limits::denorm_min() == expected); - assert(cuda::std::numeric_limits::denorm_min() == expected); + assert(float_eq(cuda::std::numeric_limits::denorm_min(), expected)); + assert(float_eq(cuda::std::numeric_limits::denorm_min(), expected)); + assert(float_eq(cuda::std::numeric_limits::denorm_min(), expected)); + assert(float_eq(cuda::std::numeric_limits::denorm_min(), expected)); } int main(int, char**) @@ -66,10 +67,10 @@ int main(int, char**) # endif #endif #if defined(_LIBCUDACXX_HAS_NVFP16) - test<__half>(5.9604644775390625e-08); + test<__half>(__double2half(5.9604644775390625e-08)); #endif // _LIBCUDACXX_HAS_NVFP16 #if defined(_LIBCUDACXX_HAS_NVBF16) - test<__nv_bfloat16>(9.18354961579912115600575419705e-41); + test<__nv_bfloat16>(__double2bfloat16(9.18354961579912115600575419705e-41)); #endif // _LIBCUDACXX_HAS_NVBF16 #if !defined(__FLT_DENORM_MIN__) && !defined(FLT_TRUE_MIN) # error Test has no expected values for floating point types diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/epsilon.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/epsilon.pass.cpp index 63b0bc7b8b4..fa42c5e8fe6 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/epsilon.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/epsilon.pass.cpp @@ -14,15 +14,16 @@ #include #include +#include "common.h" #include "test_macros.h" template __host__ __device__ void test(T expected) { - assert(cuda::std::numeric_limits::epsilon() == expected); - assert(cuda::std::numeric_limits::epsilon() == expected); - assert(cuda::std::numeric_limits::epsilon() == expected); - assert(cuda::std::numeric_limits::epsilon() == expected); + assert(float_eq(cuda::std::numeric_limits::epsilon(), expected)); + assert(float_eq(cuda::std::numeric_limits::epsilon(), expected)); + assert(float_eq(cuda::std::numeric_limits::epsilon(), expected)); + assert(float_eq(cuda::std::numeric_limits::epsilon(), expected)); } int main(int, char**) @@ -57,10 +58,10 @@ int main(int, char**) test(LDBL_EPSILON); #endif #if defined(_LIBCUDACXX_HAS_NVFP16) - test<__half>(0.0009765625); + test<__half>(__double2half(0.0009765625)); #endif // _LIBCUDACXX_HAS_NVFP16 #if defined(_LIBCUDACXX_HAS_NVBF16) - test<__nv_bfloat16>(0.0078125); + test<__nv_bfloat16>(__double2bfloat16(0.0078125)); #endif // _LIBCUDACXX_HAS_NVBF16 return 0; diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/infinity.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/infinity.pass.cpp index a13ff0692df..34527e300c5 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/infinity.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/infinity.pass.cpp @@ -14,6 +14,8 @@ #include #include +#include "common.h" + // MSVC has issues with producing INF with divisions by zero. #if defined(_MSC_VER) # include @@ -24,10 +26,10 @@ template __host__ __device__ void test(T expected) { - assert(cuda::std::numeric_limits::infinity() == expected); - assert(cuda::std::numeric_limits::infinity() == expected); - assert(cuda::std::numeric_limits::infinity() == expected); - assert(cuda::std::numeric_limits::infinity() == expected); + assert(float_eq(cuda::std::numeric_limits::infinity(), expected)); + assert(float_eq(cuda::std::numeric_limits::infinity(), expected)); + assert(float_eq(cuda::std::numeric_limits::infinity(), expected)); + assert(float_eq(cuda::std::numeric_limits::infinity(), expected)); } int main(int, char**) @@ -63,10 +65,10 @@ int main(int, char**) test(1. / 0.); # endif # if defined(_LIBCUDACXX_HAS_NVFP16) - test<__half>(1.0 / 0.0); + test<__half>(__double2half(1.0 / 0.0)); # endif // _LIBCUDACXX_HAS_NVFP16 # if defined(_LIBCUDACXX_HAS_NVBF16) - test<__nv_bfloat16>(1.0 / 0.0); + test<__nv_bfloat16>(__double2bfloat16(1.0 / 0.0)); # endif // _LIBCUDACXX_HAS_NVBF16 // MSVC has issues with producing INF with divisions by zero. #else @@ -76,10 +78,10 @@ int main(int, char**) test(INFINITY); # endif # if defined(_LIBCUDACXX_HAS_NVFP16) - test<__half>(INFINITY); + test<__half>(__double2half(INFINITY)); # endif // _LIBCUDACXX_HAS_NVFP16 # if defined(_LIBCUDACXX_HAS_NVBF16) - test<__nv_bfloat16>(INFINITY); + test<__nv_bfloat16>(__double2bfloat16(INFINITY)); # endif // _LIBCUDACXX_HAS_NVBF16 #endif diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/lowest.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/lowest.pass.cpp index 60cec9e8533..e3b832dfd9b 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/lowest.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/lowest.pass.cpp @@ -16,18 +16,19 @@ #include #include +#include "common.h" #include "test_macros.h" template __host__ __device__ void test(T expected) { - assert(cuda::std::numeric_limits::lowest() == expected); + assert(float_eq(cuda::std::numeric_limits::lowest(), expected)); assert(cuda::std::numeric_limits::is_bounded); - assert(cuda::std::numeric_limits::lowest() == expected); + assert(float_eq(cuda::std::numeric_limits::lowest(), expected)); assert(cuda::std::numeric_limits::is_bounded); - assert(cuda::std::numeric_limits::lowest() == expected); + assert(float_eq(cuda::std::numeric_limits::lowest(), expected)); assert(cuda::std::numeric_limits::is_bounded); - assert(cuda::std::numeric_limits::lowest() == expected); + assert(float_eq(cuda::std::numeric_limits::lowest(), expected)); assert(cuda::std::numeric_limits::is_bounded); } @@ -66,10 +67,10 @@ int main(int, char**) test(-LDBL_MAX); #endif #if defined(_LIBCUDACXX_HAS_NVFP16) - test<__half>(-65504.0); + test<__half>(__double2half(-65504.0)); #endif // _LIBCUDACXX_HAS_NVFP16 #if defined(_LIBCUDACXX_HAS_NVBF16) - test<__nv_bfloat16>(-3.3895313892515355e+38); + test<__nv_bfloat16>(__double2bfloat16(-3.3895313892515355e+38)); #endif // _LIBCUDACXX_HAS_NVBF16 return 0; diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max.pass.cpp index fa06dab8963..7ba6dabb1d2 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max.pass.cpp @@ -16,18 +16,19 @@ #include #include +#include "common.h" #include "test_macros.h" template __host__ __device__ void test(T expected) { - assert(cuda::std::numeric_limits::max() == expected); + assert(float_eq(cuda::std::numeric_limits::max(), expected)); assert(cuda::std::numeric_limits::is_bounded); - assert(cuda::std::numeric_limits::max() == expected); + assert(float_eq(cuda::std::numeric_limits::max(), expected)); assert(cuda::std::numeric_limits::is_bounded); - assert(cuda::std::numeric_limits::max() == expected); + assert(float_eq(cuda::std::numeric_limits::max(), expected)); assert(cuda::std::numeric_limits::is_bounded); - assert(cuda::std::numeric_limits::max() == expected); + assert(float_eq(cuda::std::numeric_limits::max(), expected)); assert(cuda::std::numeric_limits::is_bounded); } @@ -65,10 +66,10 @@ int main(int, char**) test(LDBL_MAX); #endif #if defined(_LIBCUDACXX_HAS_NVFP16) - test<__half>(65504.0); + test<__half>(__double2half(65504.0)); #endif // _LIBCUDACXX_HAS_NVFP16 #if defined(_LIBCUDACXX_HAS_NVBF16) - test<__nv_bfloat16>(3.3895313892515355e+38); + test<__nv_bfloat16>(__double2bfloat16(3.3895313892515355e+38)); #endif // _LIBCUDACXX_HAS_NVBF16 return 0; diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min.pass.cpp index 14ebad929c3..c24c3fde869 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min.pass.cpp @@ -16,18 +16,19 @@ #include #include +#include "common.h" #include "test_macros.h" template __host__ __device__ void test(T expected) { - assert(cuda::std::numeric_limits::min() == expected); + assert(float_eq(cuda::std::numeric_limits::min(), expected)); assert(cuda::std::numeric_limits::is_bounded || !cuda::std::numeric_limits::is_signed); - assert(cuda::std::numeric_limits::min() == expected); + assert(float_eq(cuda::std::numeric_limits::min(), expected)); assert(cuda::std::numeric_limits::is_bounded || !cuda::std::numeric_limits::is_signed); - assert(cuda::std::numeric_limits::min() == expected); + assert(float_eq(cuda::std::numeric_limits::min(), expected)); assert(cuda::std::numeric_limits::is_bounded || !cuda::std::numeric_limits::is_signed); - assert(cuda::std::numeric_limits::min() == expected); + assert(float_eq(cuda::std::numeric_limits::min(), expected)); assert(cuda::std::numeric_limits::is_bounded || !cuda::std::numeric_limits::is_signed); } @@ -66,10 +67,10 @@ int main(int, char**) test(LDBL_MIN); #endif #if defined(_LIBCUDACXX_HAS_NVFP16) - test<__half>(6.103515625e-05); + test<__half>(__double2half(6.103515625e-05)); #endif // _LIBCUDACXX_HAS_NVFP16 #if defined(_LIBCUDACXX_HAS_NVBF16) - test<__nv_bfloat16>(1.17549435082228750796873653722e-38); + test<__nv_bfloat16>(__double2bfloat16(1.17549435082228750796873653722e-38)); #endif // _LIBCUDACXX_HAS_NVBF16 return 0; diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/round_error.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/round_error.pass.cpp index ffc5d738109..ba5049fc49f 100644 --- a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/round_error.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/round_error.pass.cpp @@ -14,15 +14,16 @@ #include #include +#include "common.h" #include "test_macros.h" template __host__ __device__ void test(T expected) { - assert(cuda::std::numeric_limits::round_error() == expected); - assert(cuda::std::numeric_limits::round_error() == expected); - assert(cuda::std::numeric_limits::round_error() == expected); - assert(cuda::std::numeric_limits::round_error() == expected); + assert(float_eq(cuda::std::numeric_limits::round_error(), expected)); + assert(float_eq(cuda::std::numeric_limits::round_error(), expected)); + assert(float_eq(cuda::std::numeric_limits::round_error(), expected)); + assert(float_eq(cuda::std::numeric_limits::round_error(), expected)); } int main(int, char**) @@ -57,10 +58,10 @@ int main(int, char**) test(0.5); #endif #if defined(_LIBCUDACXX_HAS_NVFP16) - test<__half>(0.5); + test<__half>(__double2half(0.5)); #endif // _LIBCUDACXX_HAS_NVFP16 #if defined(_LIBCUDACXX_HAS_NVBF16) - test<__nv_bfloat16>(0.5); + test<__nv_bfloat16>(__double2bfloat16(0.5)); #endif // _LIBCUDACXX_HAS_NVBF16 return 0;