Implement cuda::std::numeric_limits for __half and __nv_bfloat16

libcudacxx/include/cuda/std/limits

@@ -22,7 +22,10 @@
 #endif // no system header
 
 #include <cuda/std/__bit/bit_cast.h>
-#include <cuda/std/__type_traits/is_arithmetic.h>
+#include <cuda/std/__type_traits/integral_constant.h>
+#include <cuda/std/__type_traits/is_extended_floating_point.h>
+#include <cuda/std/__type_traits/is_floating_point.h>
+#include <cuda/std/__type_traits/is_integral.h>
 #include <cuda/std/climits>
 #include <cuda/std/version>
 
@@ -46,7 +49,46 @@ enum float_denorm_style
   denorm_present       = 1
 };
 
-template <class _Tp, bool = is_arithmetic<_Tp>::value>
+enum class __numeric_limits_type
+{
+  __integral,
+  __bool,
+  __floating_point,
+  __other,
+};
+
+template <class _Tp>
+_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 _Tp, __numeric_limits_type = __make_numeric_limits_type<_Tp>()>
 class __numeric_limits_impl
 {
 public:
@@ -135,7 +177,7 @@ struct __int_min<_Tp, __digits, false>
 };
 
 template <class _Tp>
-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<bool, true>
+class __numeric_limits_impl<bool, __numeric_limits_type::__bool>
 {
 public:
   using type = bool;
@@ -286,7 +328,7 @@ public:
 };
 
 template <>
-class __numeric_limits_impl<float, true>
+class __numeric_limits_impl<float, __numeric_limits_type::__floating_point>
 {
 public:
   using type = float;
@@ -381,7 +423,7 @@ public:
 };
 
 template <>
-class __numeric_limits_impl<double, true>
+class __numeric_limits_impl<double, __numeric_limits_type::__floating_point>
 {
 public:
   using type = double;
@@ -476,7 +518,7 @@ public:
 };
 
 template <>
-class __numeric_limits_impl<long double, true>
+class __numeric_limits_impl<long double, __numeric_limits_type::__floating_point>
 {
 #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 _Tp>
 class numeric_limits : public __numeric_limits_impl<_Tp>
 {};

libcudacxx/test/libcudacxx/std/containers/views/mdspan/my_int.hpp

@@ -1,6 +1,9 @@
 #ifndef _MY_INT_HPP
 #define _MY_INT_HPP
 
+#include <cuda/std/limits>
+#include <cuda/std/type_traits>
+
 #include "test_macros.h"
 
 struct my_int_non_convertible;
@@ -22,6 +25,10 @@ template <>
 struct cuda::std::is_integral<my_int> : cuda::std::true_type
 {};
 
+template <>
+class cuda::std::numeric_limits<my_int> : public cuda::std::numeric_limits<int>
+{};
+
 // Wrapper type that's not implicitly convertible
 
 struct my_int_non_convertible
@@ -43,6 +50,10 @@ template <>
 struct cuda::std::is_integral<my_int_non_convertible> : cuda::std::true_type
 {};
 
+template <>
+class cuda::std::numeric_limits<my_int_non_convertible> : public cuda::std::numeric_limits<int>
+{};
+
 // Wrapper type that's not nothrow-constructible
 
 struct my_int_non_nothrow_constructible
@@ -62,4 +73,8 @@ template <>
 struct cuda::std::is_integral<my_int_non_nothrow_constructible> : cuda::std::true_type
 {};
 
+template <>
+class cuda::std::numeric_limits<my_int_non_nothrow_constructible> : public cuda::std::numeric_limits<int>
+{};
+
 #endif

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<long double>();
 #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<cuda::std::complex<double>>::is_specialized,
                 "!cuda::std::numeric_limits<cuda::std::complex<double> >::is_specialized");
 

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 <cuda/std/limits>
+
+template <class T>
+__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

libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/const_data_members.pass.cpp

@@ -42,173 +42,80 @@ template <class T>
 __host__ __device__ void test(T)
 {}
 
-#define TEST_NUMERIC_LIMITS(type)                           \
-  test(cuda::std::numeric_limits<type>::is_specialized);    \
-  test(cuda::std::numeric_limits<type>::digits);            \
-  test(cuda::std::numeric_limits<type>::digits10);          \
-  test(cuda::std::numeric_limits<type>::max_digits10);      \
-  test(cuda::std::numeric_limits<type>::is_signed);         \
-  test(cuda::std::numeric_limits<type>::is_integer);        \
-  test(cuda::std::numeric_limits<type>::is_exact);          \
-  test(cuda::std::numeric_limits<type>::radix);             \
-  test(cuda::std::numeric_limits<type>::min_exponent);      \
-  test(cuda::std::numeric_limits<type>::min_exponent10);    \
-  test(cuda::std::numeric_limits<type>::max_exponent);      \
-  test(cuda::std::numeric_limits<type>::max_exponent10);    \
-  test(cuda::std::numeric_limits<type>::has_infinity);      \
-  test(cuda::std::numeric_limits<type>::has_quiet_NaN);     \
-  test(cuda::std::numeric_limits<type>::has_signaling_NaN); \
-  test(cuda::std::numeric_limits<type>::has_denorm);        \
-  test(cuda::std::numeric_limits<type>::has_denorm_loss);   \
-  test(cuda::std::numeric_limits<type>::is_iec559);         \
-  test(cuda::std::numeric_limits<type>::is_bounded);        \
-  test(cuda::std::numeric_limits<type>::is_modulo);         \
-  test(cuda::std::numeric_limits<type>::traps);             \
-  test(cuda::std::numeric_limits<type>::tinyness_before);   \
-  test(cuda::std::numeric_limits<type>::round_style);
+template <class T>
+__host__ __device__ void test_type_helper()
+{
+  test(cuda::std::numeric_limits<T>::is_specialized);
+  test(cuda::std::numeric_limits<T>::digits);
+  test(cuda::std::numeric_limits<T>::digits10);
+  test(cuda::std::numeric_limits<T>::max_digits10);
+  test(cuda::std::numeric_limits<T>::is_signed);
+  test(cuda::std::numeric_limits<T>::is_integer);
+  test(cuda::std::numeric_limits<T>::is_exact);
+  test(cuda::std::numeric_limits<T>::radix);
+  test(cuda::std::numeric_limits<T>::min_exponent);
+  test(cuda::std::numeric_limits<T>::min_exponent10);
+  test(cuda::std::numeric_limits<T>::max_exponent);
+  test(cuda::std::numeric_limits<T>::max_exponent10);
+  test(cuda::std::numeric_limits<T>::has_infinity);
+  test(cuda::std::numeric_limits<T>::has_quiet_NaN);
+  test(cuda::std::numeric_limits<T>::has_signaling_NaN);
+  test(cuda::std::numeric_limits<T>::has_denorm);
+  test(cuda::std::numeric_limits<T>::has_denorm_loss);
+  test(cuda::std::numeric_limits<T>::is_iec559);
+  test(cuda::std::numeric_limits<T>::is_bounded);
+  test(cuda::std::numeric_limits<T>::is_modulo);
+  test(cuda::std::numeric_limits<T>::traps);
+  test(cuda::std::numeric_limits<T>::tinyness_before);
+  test(cuda::std::numeric_limits<T>::round_style);
+}
+
+template <class T>
+__host__ __device__ void test_type()
+{
+  test_type_helper<T>();
+  test_type_helper<const T>();
+  test_type_helper<volatile T>();
+  test_type_helper<const volatile T>();
+}
 
 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<bool>();
+  test_type<char>();
+  test_type<signed char>();
+  test_type<unsigned char>();
+  test_type<wchar_t>();
+#if TEST_STD_VER >= 2020 && defined(__cpp_char8_t)
+  test_type<char8_t>();
+#endif // TEST_STD_VER >= 2020 && defined(__cpp_char8_t)
+  test_type<char16_t>();
+  test_type<char32_t>();
+  test_type<short>();
+  test_type<unsigned short>();
+  test_type<int>();
+  test_type<unsigned int>();
+  test_type<long>();
+  test_type<unsigned long>();
+  test_type<long long>();
+  test_type<unsigned long long>();
 #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<float>();
+  test_type<double>();
 #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<long double>();
+#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;
 }

libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/denorm_min.pass.cpp

@@ -14,15 +14,16 @@
 #include <cuda/std/cfloat>
 #include <cuda/std/limits>
 
+#include "common.h"
 #include "test_macros.h"
 
 template <class T>
 __host__ __device__ void test(T expected)
 {
-  assert(cuda::std::numeric_limits<T>::denorm_min() == expected);
-  assert(cuda::std::numeric_limits<const T>::denorm_min() == expected);
-  assert(cuda::std::numeric_limits<volatile T>::denorm_min() == expected);
-  assert(cuda::std::numeric_limits<const volatile T>::denorm_min() == expected);
+  assert(float_eq(cuda::std::numeric_limits<T>::denorm_min(), expected));
+  assert(float_eq(cuda::std::numeric_limits<const T>::denorm_min(), expected));
+  assert(float_eq(cuda::std::numeric_limits<volatile T>::denorm_min(), expected));
+  assert(float_eq(cuda::std::numeric_limits<const volatile T>::denorm_min(), expected));
 }
 
 int main(int, char**)
@@ -65,6 +66,12 @@ int main(int, char**)
   test<long double>(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

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<long double, LDBL_MANT_DIG>();
 #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;
 }

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<long double, LDBL_DIG>();
 #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;
 }

libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/epsilon.pass.cpp

@@ -14,15 +14,16 @@
 #include <cuda/std/cfloat>
 #include <cuda/std/limits>
 
+#include "common.h"
 #include "test_macros.h"
 
 template <class T>
 __host__ __device__ void test(T expected)
 {
-  assert(cuda::std::numeric_limits<T>::epsilon() == expected);
-  assert(cuda::std::numeric_limits<const T>::epsilon() == expected);
-  assert(cuda::std::numeric_limits<volatile T>::epsilon() == expected);
-  assert(cuda::std::numeric_limits<const volatile T>::epsilon() == expected);
+  assert(float_eq(cuda::std::numeric_limits<T>::epsilon(), expected));
+  assert(float_eq(cuda::std::numeric_limits<const T>::epsilon(), expected));
+  assert(float_eq(cuda::std::numeric_limits<volatile T>::epsilon(), expected));
+  assert(float_eq(cuda::std::numeric_limits<const volatile T>::epsilon(), expected));
 }
 
 int main(int, char**)
@@ -56,6 +57,12 @@ int main(int, char**)
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double>(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;
 }

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<long double, cuda::std::denorm_present>();
 #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;
 }

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<long double, false>();
 #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;
 }

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<long double, true>();
 #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;
 }

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<long double, true>();
 #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;
 }

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<long double, true>();
 #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;
 }

libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/infinity.pass.cpp

@@ -14,6 +14,8 @@
 #include <cuda/std/cfloat>
 #include <cuda/std/limits>
 
+#include "common.h"
+
 // MSVC has issues with producing INF with divisions by zero.
 #if defined(_MSC_VER)
 #  include <cmath>
@@ -24,10 +26,10 @@
 template <class T>
 __host__ __device__ void test(T expected)
 {
-  assert(cuda::std::numeric_limits<T>::infinity() == expected);
-  assert(cuda::std::numeric_limits<const T>::infinity() == expected);
-  assert(cuda::std::numeric_limits<volatile T>::infinity() == expected);
-  assert(cuda::std::numeric_limits<const volatile T>::infinity() == expected);
+  assert(float_eq(cuda::std::numeric_limits<T>::infinity(), expected));
+  assert(float_eq(cuda::std::numeric_limits<const T>::infinity(), expected));
+  assert(float_eq(cuda::std::numeric_limits<volatile T>::infinity(), expected));
+  assert(float_eq(cuda::std::numeric_limits<const volatile T>::infinity(), expected));
 }
 
 int main(int, char**)
@@ -62,18 +64,26 @@ int main(int, char**)
 #  ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double>(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<float>(INFINITY);
   test<double>(INFINITY);
 #  ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double>(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

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<long double, true>();
 #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;
 }

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<long double, false>();
 #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;
 }

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<long double, true>();
 #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;
 }

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<long double, false>();
 #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;
 }

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<long double, false>();
 #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;
 }

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<long double, true>();
 #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;
 }

libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/lowest.pass.cpp

@@ -16,25 +16,27 @@
 #include <cuda/std/cstdint>
 #include <cuda/std/limits>
 
+#include "common.h"
 #include "test_macros.h"
 
 template <class T>
 __host__ __device__ void test(T expected)
 {
-  assert(cuda::std::numeric_limits<T>::lowest() == expected);
+  assert(float_eq(cuda::std::numeric_limits<T>::lowest(), expected));
   assert(cuda::std::numeric_limits<T>::is_bounded);
-  assert(cuda::std::numeric_limits<const T>::lowest() == expected);
+  assert(float_eq(cuda::std::numeric_limits<const T>::lowest(), expected));
   assert(cuda::std::numeric_limits<const T>::is_bounded);
-  assert(cuda::std::numeric_limits<volatile T>::lowest() == expected);
+  assert(float_eq(cuda::std::numeric_limits<volatile T>::lowest(), expected));
   assert(cuda::std::numeric_limits<volatile T>::is_bounded);
-  assert(cuda::std::numeric_limits<const volatile T>::lowest() == expected);
+  assert(float_eq(cuda::std::numeric_limits<const volatile T>::lowest(), expected));
   assert(cuda::std::numeric_limits<const volatile T>::is_bounded);
 }
 
 int main(int, char**)
 {
   test<bool>(false);
   test<char>(CHAR_MIN);
+
   test<signed char>(SCHAR_MIN);
   test<unsigned char>(0);
 #ifndef TEST_COMPILER_NVRTC
@@ -64,6 +66,12 @@ int main(int, char**)
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double>(-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;
 }

libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/max.pass.cpp

@@ -16,18 +16,19 @@
 #include <cuda/std/cstdint>
 #include <cuda/std/limits>
 
+#include "common.h"
 #include "test_macros.h"
 
 template <class T>
 __host__ __device__ void test(T expected)
 {
-  assert(cuda::std::numeric_limits<T>::max() == expected);
+  assert(float_eq(cuda::std::numeric_limits<T>::max(), expected));
   assert(cuda::std::numeric_limits<T>::is_bounded);
-  assert(cuda::std::numeric_limits<const T>::max() == expected);
+  assert(float_eq(cuda::std::numeric_limits<const T>::max(), expected));
   assert(cuda::std::numeric_limits<const T>::is_bounded);
-  assert(cuda::std::numeric_limits<volatile T>::max() == expected);
+  assert(float_eq(cuda::std::numeric_limits<volatile T>::max(), expected));
   assert(cuda::std::numeric_limits<volatile T>::is_bounded);
-  assert(cuda::std::numeric_limits<const volatile T>::max() == expected);
+  assert(float_eq(cuda::std::numeric_limits<const volatile T>::max(), expected));
   assert(cuda::std::numeric_limits<const volatile T>::is_bounded);
 }
 
@@ -64,6 +65,12 @@ int main(int, char**)
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double>(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;
 }

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<long double, 2 + (LDBL_MANT_DIG * 30103) / 100000>();
 #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;
 }

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<long double, LDBL_MAX_EXP>();
 #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;
 }

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<long double, LDBL_MAX_10_EXP>();
 #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;
 }

libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/min.pass.cpp

@@ -16,18 +16,19 @@
 #include <cuda/std/cstdint>
 #include <cuda/std/limits>
 
+#include "common.h"
 #include "test_macros.h"
 
 template <class T>
 __host__ __device__ void test(T expected)
 {
-  assert(cuda::std::numeric_limits<T>::min() == expected);
+  assert(float_eq(cuda::std::numeric_limits<T>::min(), expected));
   assert(cuda::std::numeric_limits<T>::is_bounded || !cuda::std::numeric_limits<T>::is_signed);
-  assert(cuda::std::numeric_limits<const T>::min() == expected);
+  assert(float_eq(cuda::std::numeric_limits<const T>::min(), expected));
   assert(cuda::std::numeric_limits<const T>::is_bounded || !cuda::std::numeric_limits<const T>::is_signed);
-  assert(cuda::std::numeric_limits<volatile T>::min() == expected);
+  assert(float_eq(cuda::std::numeric_limits<volatile T>::min(), expected));
   assert(cuda::std::numeric_limits<volatile T>::is_bounded || !cuda::std::numeric_limits<volatile T>::is_signed);
-  assert(cuda::std::numeric_limits<const volatile T>::min() == expected);
+  assert(float_eq(cuda::std::numeric_limits<const volatile T>::min(), expected));
   assert(cuda::std::numeric_limits<const volatile T>::is_bounded
          || !cuda::std::numeric_limits<const volatile T>::is_signed);
 }
@@ -65,6 +66,12 @@ int main(int, char**)
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double>(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;
 }

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<long double, LDBL_MIN_EXP>();
 #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;
 }

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<long double, LDBL_MIN_10_EXP>();
 #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;
 }

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 <class T>
 __host__ __device__ inline void test()
 {
-  test_imp<T>(cuda::std::is_floating_point<T>());
+  constexpr bool is_float = cuda::std::is_floating_point<T>::value || cuda::std::__is_extended_floating_point<T>::value;
+
+  test_imp<T>(cuda::std::integral_constant<bool, is_float>{});
 }
 
 int main(int, char**)
@@ -72,6 +74,12 @@ int main(int, char**)
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double>();
 #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;
 }

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<long double, FLT_RADIX>();
 #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;
 }

libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/round_error.pass.cpp

@@ -14,15 +14,16 @@
 #include <cuda/std/cfloat>
 #include <cuda/std/limits>
 
+#include "common.h"
 #include "test_macros.h"
 
 template <class T>
 __host__ __device__ void test(T expected)
 {
-  assert(cuda::std::numeric_limits<T>::round_error() == expected);
-  assert(cuda::std::numeric_limits<const T>::round_error() == expected);
-  assert(cuda::std::numeric_limits<volatile T>::round_error() == expected);
-  assert(cuda::std::numeric_limits<const volatile T>::round_error() == expected);
+  assert(float_eq(cuda::std::numeric_limits<T>::round_error(), expected));
+  assert(float_eq(cuda::std::numeric_limits<const T>::round_error(), expected));
+  assert(float_eq(cuda::std::numeric_limits<volatile T>::round_error(), expected));
+  assert(float_eq(cuda::std::numeric_limits<const volatile T>::round_error(), expected));
 }
 
 int main(int, char**)
@@ -56,6 +57,12 @@ int main(int, char**)
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double>(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;
 }

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<long double, cuda::std::round_to_nearest>();
 #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;
 }

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 <class T>
 __host__ __device__ inline void test()
 {
-  test_imp<T>(cuda::std::is_floating_point<T>());
+  constexpr bool is_float = cuda::std::is_floating_point<T>::value || cuda::std::__is_extended_floating_point<T>::value;
+
+  test_imp<T>(cuda::std::integral_constant<bool, is_float>{});
 }
 
 int main(int, char**)
@@ -72,6 +74,12 @@ int main(int, char**)
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double>();
 #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;
 }

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<long double, false>();
 #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;
 }

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<long double, false>();
 #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;
 }