From 4f2efaf9d57ed297aa1e16c8e2207c96cac3d7ca Mon Sep 17 00:00:00 2001 From: David Bayer <48736217+davebayer@users.noreply.github.com> Date: Tue, 14 Jan 2025 09:01:58 +0100 Subject: [PATCH 1/3] Implement `cuda::std::numeric_limits` for `__half` and `__nv_bfloat16` (#3361) * implement `cuda::std::numeric_limits` for `__half` and `__nv_bfloat16` --- libcudacxx/include/cuda/std/limits | 206 +++++++++++++++- .../std/containers/views/mdspan/my_int.hpp | 15 ++ .../limits/is_specialized.pass.cpp | 7 + .../limits/numeric.limits.members/common.h | 41 ++++ .../const_data_members.pass.cpp | 225 +++++------------- .../denorm_min.pass.cpp | 15 +- .../numeric.limits.members/digits.pass.cpp | 7 +- .../numeric.limits.members/digits10.pass.cpp | 6 + .../numeric.limits.members/epsilon.pass.cpp | 15 +- .../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 | 26 +- .../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 | 16 +- .../numeric.limits.members/max.pass.cpp | 15 +- .../max_digits10.pass.cpp | 6 + .../max_exponent.pass.cpp | 6 + .../max_exponent10.pass.cpp | 6 + .../numeric.limits.members/min.pass.cpp | 15 +- .../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 | 15 +- .../round_style.pass.cpp | 6 + .../signaling_NaN.pass.cpp | 10 +- .../tinyness_before.pass.cpp | 6 + .../numeric.limits.members/traps.pass.cpp | 6 + 36 files changed, 563 insertions(+), 201 deletions(-) create mode 100644 libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/common.h diff --git a/libcudacxx/include/cuda/std/limits b/libcudacxx/include/cuda/std/limits index 98c63813b7b..ad529f2082b 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,46 @@ enum float_denorm_style denorm_present = 1 }; -template ::value> +enum class __numeric_limits_type +{ + __integral, + __bool, + __floating_point, + __other, +}; + +template +_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)) + { + 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 ()> class __numeric_limits_impl { public: @@ -135,7 +177,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 +254,7 @@ public: }; template <> -class __numeric_limits_impl +class __numeric_limits_impl { public: using type = bool; @@ -286,7 +328,7 @@ public: }; template <> -class __numeric_limits_impl +class __numeric_limits_impl { public: using type = float; @@ -381,7 +423,7 @@ public: }; template <> -class __numeric_limits_impl +class __numeric_limits_impl { public: using type = double; @@ -476,7 +518,7 @@ public: }; template <> -class __numeric_limits_impl +class __numeric_limits_impl { #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE @@ -551,6 +593,156 @@ public: #endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE }; +#if defined(_LIBCUDACXX_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 // _LIBCUDACXX_HAS_NVFP16 + +#if defined(_LIBCUDACXX_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 // _LIBCUDACXX_HAS_NVBF16 + template class numeric_limits : public __numeric_limits_impl<_Tp> {}; 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 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..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,6 +68,13 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16>(); +#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/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/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..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 @@ -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 defined(_LIBCUDACXX_HAS_NVFP16) + test_type<__half>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test_type<__nv_bfloat16>(); +#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 730adc30d36..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**) @@ -65,6 +66,12 @@ int main(int, char**) test(LDBL_TRUE_MIN); # endif #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half>(__double2half(5.9604644775390625e-08)); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + 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 #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..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,6 +55,11 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif - +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, 11>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, 8>(); +#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 3295686ea49..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,6 +59,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, 3>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, 2>(); +#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 5bc22e7f5f2..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**) @@ -56,6 +57,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(LDBL_EPSILON); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half>(__double2half(0.0009765625)); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + 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/has_denorm.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/has_denorm.pass.cpp index e62208d7e3b..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,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, cuda::std::denorm_present>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, cuda::std::denorm_present>(); +#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 3a665fe2c9b..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,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, false>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, false>(); +#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 be41dabb02c..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,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, true>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, true>(); +#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 2d13db35438..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,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, true>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, true>(); +#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 d5cf5096bb7..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,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, true>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, true>(); +#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 2d1c29f6f31..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**) @@ -62,6 +64,12 @@ int main(int, char**) # ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(1. / 0.); # endif +# if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half>(__double2half(1.0 / 0.0)); +# endif // _LIBCUDACXX_HAS_NVFP16 +# if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16>(__double2bfloat16(1.0 / 0.0)); +# endif // _LIBCUDACXX_HAS_NVBF16 // MSVC has issues with producing INF with divisions by zero. #else test(INFINITY); @@ -69,11 +77,13 @@ int main(int, char**) # ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(INFINITY); # endif +# if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half>(__double2half(INFINITY)); +# endif // _LIBCUDACXX_HAS_NVFP16 +# if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16>(__double2bfloat16(INFINITY)); +# endif // _LIBCUDACXX_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..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,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, true>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, true>(); +#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 be45efae70c..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,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, false>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, false>(); +#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 6221cd6ed59..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,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif // _LIBCUDACXX_HAS_NO_LONG_DOUBLE +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, true>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, true>(); +#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 3d166f31f28..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,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, false>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, false>(); +#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 7b1adabf0c7..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,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, false>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, false>(); +#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 d7f98766343..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,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, true>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, true>(); +#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 6fec93e4a3d..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); } @@ -35,6 +36,7 @@ int main(int, char**) { test(false); test(CHAR_MIN); + test(SCHAR_MIN); test(0); #ifndef TEST_COMPILER_NVRTC @@ -64,6 +66,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(-LDBL_MAX); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half>(__double2half(-65504.0)); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + 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 67c94051729..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); } @@ -64,6 +65,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(LDBL_MAX); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half>(__double2half(65504.0)); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + 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_digits10.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max_digits10.pass.cpp index cd5892e6c8c..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,6 +55,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, 5>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, 4>(); +#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 aeb9189d315..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,6 +55,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, 16>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, 128>(); +#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 ca0eb2917f6..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,6 +55,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, 4>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, 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 53d196d2a51..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); } @@ -65,6 +66,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(LDBL_MIN); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half>(__double2half(6.103515625e-05)); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + 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/min_exponent.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min_exponent.pass.cpp index b075bcff87d..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,6 +55,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, -13>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, -125>(); +#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 c787cf4caab..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,6 +55,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, -4>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, -37>(); +#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 ce38b3ed60d..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 @@ -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 defined(_LIBCUDACXX_HAS_NVFP16) + test<__half>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16>(); +#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 5a709b3aefc..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,6 +55,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, FLT_RADIX>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, FLT_RADIX>(); +#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 01d10e80fb9..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**) @@ -56,6 +57,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(0.5); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half>(__double2half(0.5)); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16>(__double2bfloat16(0.5)); +#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 7a7099662f0..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,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, cuda::std::round_to_nearest>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, cuda::std::round_to_nearest>(); +#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 164d54c5741..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 @@ -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 defined(_LIBCUDACXX_HAS_NVFP16) + test<__half>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16>(); +#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 70d832dc547..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,6 +54,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, false>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, false>(); +#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 5c66acb56ce..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,6 +60,12 @@ int main(int, char**) #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE test(); #endif +#if defined(_LIBCUDACXX_HAS_NVFP16) + test<__half, false>(); +#endif // _LIBCUDACXX_HAS_NVFP16 +#if defined(_LIBCUDACXX_HAS_NVBF16) + test<__nv_bfloat16, false>(); +#endif // _LIBCUDACXX_HAS_NVBF16 return 0; } From daab0a4c1cc01fe631e41602ba211e0b59b8c59c Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Tue, 14 Jan 2025 09:39:24 +0100 Subject: [PATCH 2/3] Cleanup util_arch (#2773) --- cub/cub/util_arch.cuh | 31 ++++++++++++++++--------------- 1 file changed, 16 insertions(+), 15 deletions(-) diff --git a/cub/cub/util_arch.cuh b/cub/cub/util_arch.cuh index b1da6a03b5d..3c6aea5cc5c 100644 --- a/cub/cub/util_arch.cuh +++ b/cub/cub/util_arch.cuh @@ -47,6 +47,10 @@ #include #include +#include +#include +#include + // Legacy include; this functionality used to be defined in here. #include @@ -113,27 +117,24 @@ namespace detail static constexpr ::cuda::std::size_t max_smem_per_block = 48 * 1024; } // namespace detail -template +template struct RegBoundScaling { - enum - { - ITEMS_PER_THREAD = CUB_MAX(1, NOMINAL_4B_ITEMS_PER_THREAD * 4 / CUB_MAX(4, sizeof(T))), - BLOCK_THREADS = CUB_MIN(NOMINAL_4B_BLOCK_THREADS, - ((cub::detail::max_smem_per_block / (sizeof(T) * ITEMS_PER_THREAD)) + 31) / 32 * 32), - }; + static constexpr int ITEMS_PER_THREAD = + ::cuda::std::max(1, Nominal4ByteItemsPerThread * 4 / ::cuda::std::max(4, int{sizeof(T)})); + static constexpr int BLOCK_THREADS = + ::cuda::std::min(Nominal4ByteBlockThreads, + ::cuda::ceil_div(int{detail::max_smem_per_block} / (int{sizeof(T)} * ITEMS_PER_THREAD), 32) * 32); }; -template +template struct MemBoundScaling { - enum - { - ITEMS_PER_THREAD = - CUB_MAX(1, CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(T), NOMINAL_4B_ITEMS_PER_THREAD * 2)), - BLOCK_THREADS = CUB_MIN(NOMINAL_4B_BLOCK_THREADS, - ((cub::detail::max_smem_per_block / (sizeof(T) * ITEMS_PER_THREAD)) + 31) / 32 * 32), - }; + static constexpr int ITEMS_PER_THREAD = ::cuda::std::max( + 1, ::cuda::std::min(Nominal4ByteItemsPerThread * 4 / int{sizeof(T)}, Nominal4ByteItemsPerThread * 2)); + static constexpr int BLOCK_THREADS = + ::cuda::std::min(Nominal4ByteBlockThreads, + ::cuda::ceil_div(int{detail::max_smem_per_block} / (int{sizeof(T)} * ITEMS_PER_THREAD), 32) * 32); }; #endif // Do not document From 0e635521ff2eeafdd4a89e8a991ceb6a393b2566 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Tue, 14 Jan 2025 09:39:41 +0100 Subject: [PATCH 3/3] Deprecate thrust::null_type (#3367) --- thrust/testing/tuple.cu | 2 ++ thrust/thrust/tuple.h | 8 +++++++- 2 files changed, 9 insertions(+), 1 deletion(-) diff --git a/thrust/testing/tuple.cu b/thrust/testing/tuple.cu index 4be47b1de81..cccffec0081 100644 --- a/thrust/testing/tuple.cu +++ b/thrust/testing/tuple.cu @@ -529,6 +529,7 @@ DECLARE_UNITTEST(TestTupleCTAD); #endif // _CCCL_STD_VER >= 2017 // Ensure that we are backwards compatible with the old thrust::tuple implementation +_CCCL_SUPPRESS_DEPRECATED_PUSH static_assert( thrust::tuple_size>::value == 9, ""); static_assert(thrust::tuple_size>::value == 10, ""); +_CCCL_SUPPRESS_DEPRECATED_POP diff --git a/thrust/thrust/tuple.h b/thrust/thrust/tuple.h index f54e814baa6..ff57db3fd97 100644 --- a/thrust/thrust/tuple.h +++ b/thrust/thrust/tuple.h @@ -44,9 +44,11 @@ THRUST_NAMESPACE_BEGIN // define null_type for backwards compatibility -struct null_type +struct CCCL_DEPRECATED_BECAUSE("Please remove null_type from parameters to tuple<...>") null_type {}; +_CCCL_SUPPRESS_DEPRECATED_PUSH + _CCCL_HOST_DEVICE inline bool operator==(const null_type&, const null_type&) { return true; @@ -77,6 +79,8 @@ _CCCL_HOST_DEVICE inline bool operator>(const null_type&, const null_type&) return false; } +_CCCL_SUPPRESS_DEPRECATED_POP + /*! \addtogroup utility * \{ */ @@ -176,6 +180,7 @@ using _CUDA_VSTD::tie; THRUST_NAMESPACE_END _LIBCUDACXX_BEGIN_NAMESPACE_STD +_CCCL_SUPPRESS_DEPRECATED_PUSH template <> struct tuple_size> {}; +_CCCL_SUPPRESS_DEPRECATED_POP _LIBCUDACXX_END_NAMESPACE_STD