diff --git a/libcudacxx/include/cuda/std/limits b/libcudacxx/include/cuda/std/limits
index ad529f2082b..f242dd19270 100644
--- a/libcudacxx/include/cuda/std/limits
+++ b/libcudacxx/include/cuda/std/limits
@@ -29,6 +29,21 @@
 #include <cuda/std/climits>
 #include <cuda/std/version>
 
+#if defined(_LIBCUDACXX_HAS_NVFP16)
+#  include <cuda_fp16.h>
+#endif // _LIBCUDACXX_HAS_NVFP16
+
+#if defined(_LIBCUDACXX_HAS_NVBF16)
+_CCCL_DIAG_PUSH
+_CCCL_DIAG_SUPPRESS_CLANG("-Wunused-function")
+#  include <cuda_bf16.h>
+_CCCL_DIAG_POP
+#endif // _LIBCUDACXX_HAS_NVBF16
+
+#if _CCCL_HAS_NVFP8()
+#  include <cuda_fp8.h>
+#endif // _CCCL_HAS_NVFP8()
+
 _CCCL_PUSH_MACROS
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
@@ -743,6 +758,183 @@ public:
 };
 #endif // _LIBCUDACXX_HAS_NVBF16
 
+#if _CCCL_HAS_NVFP8()
+#  if defined(_CCCL_BUILTIN_BIT_CAST) || _CCCL_STD_VER >= 2014
+#    define _LIBCUDACXX_CONSTEXPR_FP8_LIMITS constexpr
+#  else // ^^^ _CCCL_BUILTIN_BIT_CAST || _CCCL_STD_VER >= 2014 ^^^ // vvv !_CCCL_BUILTIN_BIT_CAST && _CCCL_STD_VER <
+        // 2014 vvv
+#    define _LIBCUDACXX_CONSTEXPR_FP8_LIMITS
+#  endif // ^^^ !_CCCL_BUILTIN_BIT_CAST && _CCCL_STD_VER < 2014 ^^^
+
+template <>
+class __numeric_limits_impl<__nv_fp8_e4m3, __numeric_limits_type::__floating_point>
+{
+  _LIBCUDACXX_HIDE_FROM_ABI static _LIBCUDACXX_CONSTEXPR_FP8_LIMITS __nv_fp8_e4m3 __make_value(__nv_fp8_storage_t __val)
+  {
+#  if defined(_CCCL_BUILTIN_BIT_CAST)
+    return _CUDA_VSTD::bit_cast<__nv_fp8_e4m3>(__val);
+#  else // ^^^ _CCCL_BUILTIN_BIT_CAST ^^^ // vvv !_CCCL_BUILTIN_BIT_CAST vvv
+    __nv_fp8_e4m3 __ret{};
+    __ret.__x = __val;
+    return __ret;
+#  endif // ^^^ !_CCCL_BUILTIN_BIT_CAST ^^^
+  }
+
+public:
+  using type = __nv_fp8_e4m3;
+
+  static constexpr bool is_specialized = true;
+
+  static constexpr bool is_signed   = true;
+  static constexpr int digits       = 3;
+  static constexpr int digits10     = 0;
+  static constexpr int max_digits10 = 2;
+  _LIBCUDACXX_HIDE_FROM_ABI static _LIBCUDACXX_CONSTEXPR_FP8_LIMITS type min() noexcept
+  {
+    return __make_value(static_cast<__nv_fp8_storage_t>(0x08u));
+  }
+  _LIBCUDACXX_HIDE_FROM_ABI static _LIBCUDACXX_CONSTEXPR_FP8_LIMITS type max() noexcept
+  {
+    return __make_value(static_cast<__nv_fp8_storage_t>(0x7eu));
+  }
+  _LIBCUDACXX_HIDE_FROM_ABI static _LIBCUDACXX_CONSTEXPR_FP8_LIMITS type lowest() noexcept
+  {
+    return __make_value(static_cast<__nv_fp8_storage_t>(0xfeu));
+  }
+
+  static constexpr bool is_integer = false;
+  static constexpr bool is_exact   = false;
+  static constexpr int radix       = __FLT_RADIX__;
+  _LIBCUDACXX_HIDE_FROM_ABI static _LIBCUDACXX_CONSTEXPR_FP8_LIMITS type epsilon() noexcept
+  {
+    return __make_value(static_cast<__nv_fp8_storage_t>(0x20u));
+  }
+  _LIBCUDACXX_HIDE_FROM_ABI static _LIBCUDACXX_CONSTEXPR_FP8_LIMITS type round_error() noexcept
+  {
+    return __make_value(static_cast<__nv_fp8_storage_t>(0x30u));
+  }
+
+  static constexpr int min_exponent   = -6;
+  static constexpr int min_exponent10 = -2;
+  static constexpr int max_exponent   = 8;
+  static constexpr int max_exponent10 = 2;
+
+  static constexpr bool has_infinity             = false;
+  static constexpr bool has_quiet_NaN            = true;
+  static constexpr bool has_signaling_NaN        = false;
+  static constexpr float_denorm_style has_denorm = denorm_present;
+  static constexpr bool has_denorm_loss          = false;
+  _LIBCUDACXX_HIDE_FROM_ABI static _LIBCUDACXX_CONSTEXPR_FP8_LIMITS type infinity() noexcept
+  {
+    return type{};
+  }
+  _LIBCUDACXX_HIDE_FROM_ABI static _LIBCUDACXX_CONSTEXPR_FP8_LIMITS type quiet_NaN() noexcept
+  {
+    return __make_value(static_cast<__nv_fp8_storage_t>(0x7fu));
+  }
+  _LIBCUDACXX_HIDE_FROM_ABI static _LIBCUDACXX_CONSTEXPR_FP8_LIMITS type signaling_NaN() noexcept
+  {
+    return type{};
+  }
+  _LIBCUDACXX_HIDE_FROM_ABI static _LIBCUDACXX_CONSTEXPR_FP8_LIMITS type denorm_min() noexcept
+  {
+    return __make_value(static_cast<__nv_fp8_storage_t>(0x01u));
+  }
+
+  static constexpr bool is_iec559  = false;
+  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;
+};
+
+template <>
+class __numeric_limits_impl<__nv_fp8_e5m2, __numeric_limits_type::__floating_point>
+{
+  _LIBCUDACXX_HIDE_FROM_ABI static _LIBCUDACXX_CONSTEXPR_FP8_LIMITS __nv_fp8_e5m2 __make_value(__nv_fp8_storage_t __val)
+  {
+#  if defined(_CCCL_BUILTIN_BIT_CAST)
+    return _CUDA_VSTD::bit_cast<__nv_fp8_e5m2>(__val);
+#  else // ^^^ _CCCL_BUILTIN_BIT_CAST ^^^ // vvv !_CCCL_BUILTIN_BIT_CAST vvv
+    __nv_fp8_e5m2 __ret{};
+    __ret.__x = __val;
+    return __ret;
+#  endif // ^^^ !_CCCL_BUILTIN_BIT_CAST ^^^
+  }
+
+public:
+  using type = __nv_fp8_e5m2;
+
+  static constexpr bool is_specialized = true;
+
+  static constexpr bool is_signed   = true;
+  static constexpr int digits       = 2;
+  static constexpr int digits10     = 0;
+  static constexpr int max_digits10 = 2;
+  _LIBCUDACXX_HIDE_FROM_ABI static _LIBCUDACXX_CONSTEXPR_FP8_LIMITS type min() noexcept
+  {
+    return __make_value(static_cast<__nv_fp8_storage_t>(0x04u));
+  }
+  _LIBCUDACXX_HIDE_FROM_ABI static _LIBCUDACXX_CONSTEXPR_FP8_LIMITS type max() noexcept
+  {
+    return __make_value(static_cast<__nv_fp8_storage_t>(0x7bu));
+  }
+  _LIBCUDACXX_HIDE_FROM_ABI static _LIBCUDACXX_CONSTEXPR_FP8_LIMITS type lowest() noexcept
+  {
+    return __make_value(static_cast<__nv_fp8_storage_t>(0xfbu));
+  }
+
+  static constexpr bool is_integer = false;
+  static constexpr bool is_exact   = false;
+  static constexpr int radix       = __FLT_RADIX__;
+  _LIBCUDACXX_HIDE_FROM_ABI static _LIBCUDACXX_CONSTEXPR_FP8_LIMITS type epsilon() noexcept
+  {
+    return __make_value(static_cast<__nv_fp8_storage_t>(0x34u));
+  }
+  _LIBCUDACXX_HIDE_FROM_ABI static _LIBCUDACXX_CONSTEXPR_FP8_LIMITS type round_error() noexcept
+  {
+    return __make_value(static_cast<__nv_fp8_storage_t>(0x38u));
+  }
+
+  static constexpr int min_exponent   = -15;
+  static constexpr int min_exponent10 = -5;
+  static constexpr int max_exponent   = 15;
+  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 _LIBCUDACXX_CONSTEXPR_FP8_LIMITS type infinity() noexcept
+  {
+    return __make_value(static_cast<__nv_fp8_storage_t>(0x7cu));
+  }
+  _LIBCUDACXX_HIDE_FROM_ABI static _LIBCUDACXX_CONSTEXPR_FP8_LIMITS type quiet_NaN() noexcept
+  {
+    return __make_value(static_cast<__nv_fp8_storage_t>(0x7eu));
+  }
+  _LIBCUDACXX_HIDE_FROM_ABI static _LIBCUDACXX_CONSTEXPR_FP8_LIMITS type signaling_NaN() noexcept
+  {
+    return __make_value(static_cast<__nv_fp8_storage_t>(0x7du));
+  }
+  _LIBCUDACXX_HIDE_FROM_ABI static _LIBCUDACXX_CONSTEXPR_FP8_LIMITS type denorm_min() noexcept
+  {
+    return __make_value(static_cast<__nv_fp8_storage_t>(0x01u));
+  }
+
+  static constexpr bool is_iec559  = false;
+  static constexpr bool is_bounded = true;
+  static constexpr bool is_modulo  = false;
+
+  static constexpr bool traps                    = false;
+  static constexpr bool tinyness_before          = false;
+  static constexpr float_round_style round_style = round_to_nearest;
+};
+#endif // _CCCL_HAS_NVFP8()
+
 template <class _Tp>
 class numeric_limits : public __numeric_limits_impl<_Tp>
 {};
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
index 15b48836839..8400071611c 100644
--- 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
@@ -11,6 +11,7 @@
 #define NUMERIC_LIMITS_MEMBERS_COMMON_H
 
 // Disable all the extended floating point operations and conversions
+#define __CUDA_NO_FP8_CONVERSIONS__      1
 #define __CUDA_NO_HALF_CONVERSIONS__     1
 #define __CUDA_NO_HALF_OPERATORS__       1
 #define __CUDA_NO_BFLOAT16_CONVERSIONS__ 1
@@ -24,6 +25,32 @@ __host__ __device__ bool float_eq(T x, T y)
   return x == y;
 }
 
+#if _CCCL_HAS_NVFP8()
+__host__ __device__ inline __nv_fp8_e4m3 make_fp8_e4m3(double x, __nv_saturation_t sat = __NV_NOSAT)
+{
+  __nv_fp8_e4m3 res;
+  res.__x = __nv_cvt_double_to_fp8(x, sat, __NV_E4M3);
+  return res;
+}
+
+__host__ __device__ inline __nv_fp8_e5m2 make_fp8_e5m2(double x, __nv_saturation_t sat = __NV_NOSAT)
+{
+  __nv_fp8_e5m2 res;
+  res.__x = __nv_cvt_double_to_fp8(x, sat, __NV_E5M2);
+  return res;
+}
+
+__host__ __device__ inline bool float_eq(__nv_fp8_e4m3 x, __nv_fp8_e4m3 y)
+{
+  return float_eq(__half{__nv_cvt_fp8_to_halfraw(x.__x, __NV_E4M3)}, __half{__nv_cvt_fp8_to_halfraw(y.__x, __NV_E4M3)});
+}
+
+__host__ __device__ inline bool float_eq(__nv_fp8_e5m2 x, __nv_fp8_e5m2 y)
+{
+  return float_eq(__half{__nv_cvt_fp8_to_halfraw(x.__x, __NV_E5M2)}, __half{__nv_cvt_fp8_to_halfraw(y.__x, __NV_E5M2)});
+}
+#endif // _CCCL_HAS_NVFP8
+
 #if defined(_LIBCUDACXX_HAS_NVFP16)
 __host__ __device__ inline bool float_eq(__half x, __half y)
 {
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 769080cff83..b095d63afcd 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
@@ -116,6 +116,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test_type<__nv_bfloat16>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test_type<__nv_fp8_e4m3>();
+  test_type<__nv_fp8_e5m2>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 cc64ed14686..475f41a3388 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
@@ -72,6 +72,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16>(__double2bfloat16(9.18354961579912115600575419705e-41));
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3>(make_fp8_e4m3(0.001953125));
+  test<__nv_fp8_e5m2>(make_fp8_e5m2(0.0000152587890625));
+#endif // _CCCL_HAS_NVFP8()
 #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 efce1ccf678..0d3c910b672 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
@@ -61,5 +61,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16, 8>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3, 3>();
+  test<__nv_fp8_e5m2, 2>();
+#endif // _CCCL_HAS_NVFP8()
+
   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 32990ece4b1..bd66aeecfeb 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
@@ -12,10 +12,25 @@
 
 #include <cuda/std/cfloat>
 #include <cuda/std/limits>
+#include <cuda/std/type_traits>
 
 #include "test_macros.h"
 
-template <class T, int expected>
+template <class T, cuda::std::enable_if_t<cuda::std::is_integral<T>::value, int> = 0>
+__host__ __device__ constexpr int make_expected_digits10()
+{
+  // digits * log10(2)
+  return static_cast<int>((cuda::std::numeric_limits<T>::digits * 30103l) / 100000l);
+}
+
+template <class T, cuda::std::enable_if_t<!cuda::std::is_integral<T>::value, int> = 0>
+__host__ __device__ constexpr int make_expected_digits10()
+{
+  // (digits - 1) * log10(2)
+  return static_cast<int>(((cuda::std::numeric_limits<T>::digits - 1) * 30103l) / 100000l);
+}
+
+template <class T, int expected = make_expected_digits10<T>()>
 __host__ __device__ void test()
 {
   static_assert(cuda::std::numeric_limits<T>::digits10 == expected, "digits10 test 1");
@@ -30,41 +45,45 @@ __host__ __device__ void test()
 
 int main(int, char**)
 {
-  test<bool, 0>();
-  test<char, 2>();
-  test<signed char, 2>();
-  test<unsigned char, 2>();
-  test<wchar_t, 5 * sizeof(wchar_t) / 2 - 1>(); // 4 -> 9 and 2 -> 4
+  test<bool>();
+  test<char>();
+  test<signed char>();
+  test<unsigned char>();
+  test<wchar_t>();
 #if TEST_STD_VER > 2017 && defined(__cpp_char8_t)
-  test<char8_t, 2>();
+  test<char8_t>();
 #endif
 #ifndef _LIBCUDACXX_HAS_NO_UNICODE_CHARS
-  test<char16_t, 4>();
-  test<char32_t, 9>();
+  test<char16_t>();
+  test<char32_t>();
 #endif // _LIBCUDACXX_HAS_NO_UNICODE_CHARS
-  test<short, 4>();
-  test<unsigned short, 4>();
-  test<int, 9>();
-  test<unsigned int, 9>();
-  test<long, sizeof(long) == 4 ? 9 : 18>();
-  test<unsigned long, sizeof(long) == 4 ? 9 : 19>();
-  test<long long, 18>();
-  test<unsigned long long, 19>();
+  test<short>();
+  test<unsigned short>();
+  test<int>();
+  test<unsigned int>();
+  test<long>();
+  test<unsigned long>();
+  test<long long>();
+  test<unsigned long long>();
 #ifndef _LIBCUDACXX_HAS_NO_INT128
-  test<__int128_t, 38>();
-  test<__uint128_t, 38>();
+  test<__int128_t>();
+  test<__uint128_t>();
 #endif
-  test<float, FLT_DIG>();
-  test<double, DBL_DIG>();
+  test<float>();
+  test<double>();
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
-  test<long double, LDBL_DIG>();
+  test<long double>();
 #endif
 #if defined(_LIBCUDACXX_HAS_NVFP16)
-  test<__half, 3>();
+  test<__half>();
 #endif // _LIBCUDACXX_HAS_NVFP16
 #if defined(_LIBCUDACXX_HAS_NVBF16)
-  test<__nv_bfloat16, 2>();
+  test<__nv_bfloat16>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3>();
+  test<__nv_fp8_e5m2>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 fa42c5e8fe6..15366bdf308 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
@@ -63,6 +63,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16>(__double2bfloat16(0.0078125));
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3>(make_fp8_e4m3(0.125));
+  test<__nv_fp8_e5m2>(make_fp8_e5m2(0.25));
+#endif // _CCCL_HAS_NVFP8()
 
   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 5a0a05ab73b..8fa506b93ce 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
@@ -60,6 +60,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16, cuda::std::denorm_present>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3, cuda::std::denorm_present>();
+  test<__nv_fp8_e5m2, cuda::std::denorm_present>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 450e51b8111..3b7722acd8b 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
@@ -60,6 +60,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16, false>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3, false>();
+  test<__nv_fp8_e5m2, false>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 646f5e20160..ebddcb4421e 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
@@ -60,6 +60,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16, true>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3, false>();
+  test<__nv_fp8_e5m2, true>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 626b4110695..908f2d7fa4a 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
@@ -60,6 +60,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16, true>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3, true>();
+  test<__nv_fp8_e5m2, true>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 20cd04d107e..62d81c8a524 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
@@ -60,6 +60,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16, true>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3, false>();
+  test<__nv_fp8_e5m2, true>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 34527e300c5..627105a4a8c 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
@@ -70,6 +70,10 @@ int main(int, char**)
 #  if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16>(__double2bfloat16(1.0 / 0.0));
 #  endif // _LIBCUDACXX_HAS_NVBF16
+#  if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3>(__nv_fp8_e4m3{});
+  test<__nv_fp8_e5m2>(make_fp8_e5m2(1.0 / 0.0));
+#  endif // _CCCL_HAS_NVFP8()
 // MSVC has issues with producing INF with divisions by zero.
 #else
   test<float>(INFINITY);
@@ -83,6 +87,10 @@ int main(int, char**)
 #  if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16>(__double2bfloat16(INFINITY));
 #  endif // _LIBCUDACXX_HAS_NVBF16
+#  if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3>(__nv_fp8_e4m3{});
+  test<__nv_fp8_e5m2>(make_fp8_e5m2(INFINITY));
+#  endif // _CCCL_HAS_NVFP8()
 #endif
 
   return 0;
diff --git a/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_bounded.pass.cpp b/libcudacxx/test/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/is_bounded.pass.cpp
index 9e671c5d905..eeb9740e4e2 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
@@ -60,6 +60,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16, true>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3, true>();
+  test<__nv_fp8_e5m2, true>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 cfc9a6cab90..c3c2e027c72 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
@@ -60,6 +60,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16, false>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3, false>();
+  test<__nv_fp8_e5m2, false>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 945347ff4b5..7bab40e8826 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
@@ -60,6 +60,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16, true>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3, false>();
+  test<__nv_fp8_e5m2, false>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 65dd98fdb04..68e7437f1e0 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
@@ -60,6 +60,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16, false>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3, false>();
+  test<__nv_fp8_e5m2, false>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 6d82269e1c8..992be2b18b7 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
@@ -60,6 +60,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16, false>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3, false>();
+  test<__nv_fp8_e5m2, false>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 eb39869bf24..be7e4f235a7 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
@@ -60,6 +60,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16, true>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3, true>();
+  test<__nv_fp8_e5m2, true>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 e3b832dfd9b..6a8b2a9c181 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
@@ -72,6 +72,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16>(__double2bfloat16(-3.3895313892515355e+38));
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3>(make_fp8_e4m3(-448.0));
+  test<__nv_fp8_e5m2>(make_fp8_e5m2(-57344.0));
+#endif // _CCCL_HAS_NVFP8()
 
   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 7ba6dabb1d2..a1582e41b22 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
@@ -71,6 +71,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16>(__double2bfloat16(3.3895313892515355e+38));
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3>(make_fp8_e4m3(448.0));
+  test<__nv_fp8_e5m2>(make_fp8_e5m2(57344.0));
+#endif // _CCCL_HAS_NVFP8()
 
   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 92b3d13ea61..d01a4aa099c 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
@@ -12,10 +12,24 @@
 
 #include <cuda/std/cfloat>
 #include <cuda/std/limits>
+#include <cuda/std/type_traits>
 
 #include "test_macros.h"
 
-template <class T, int expected>
+template <class T, cuda::std::enable_if_t<cuda::std::is_integral<T>::value, int> = 0>
+__host__ __device__ constexpr int make_expected_max_digits10()
+{
+  return 0;
+}
+
+template <class T, cuda::std::enable_if_t<!cuda::std::is_integral<T>::value, int> = 0>
+__host__ __device__ constexpr int make_expected_max_digits10()
+{
+  // std::ceil(std::numeric_limits<float>::digits * std::log10(2) + 1)
+  return static_cast<int>((cuda::std::numeric_limits<T>::digits * 30103l + 99999l) / 100000l) + 1;
+}
+
+template <class T, int expected = make_expected_max_digits10<T>()>
 __host__ __device__ void test()
 {
   static_assert(cuda::std::numeric_limits<T>::max_digits10 == expected, "max_digits10 test 1");
@@ -26,41 +40,45 @@ __host__ __device__ void test()
 
 int main(int, char**)
 {
-  test<bool, 0>();
-  test<char, 0>();
-  test<signed char, 0>();
-  test<unsigned char, 0>();
-  test<wchar_t, 0>();
+  test<bool>();
+  test<char>();
+  test<signed char>();
+  test<unsigned char>();
+  test<wchar_t>();
 #if TEST_STD_VER > 2017 && defined(__cpp_char8_t)
-  test<char8_t, 0>();
+  test<char8_t>();
 #endif
 #ifndef _LIBCUDACXX_HAS_NO_UNICODE_CHARS
-  test<char16_t, 0>();
-  test<char32_t, 0>();
+  test<char16_t>();
+  test<char32_t>();
 #endif // _LIBCUDACXX_HAS_NO_UNICODE_CHARS
-  test<short, 0>();
-  test<unsigned short, 0>();
-  test<int, 0>();
-  test<unsigned int, 0>();
-  test<long, 0>();
-  test<unsigned long, 0>();
-  test<long long, 0>();
-  test<unsigned long long, 0>();
+  test<short>();
+  test<unsigned short>();
+  test<int>();
+  test<unsigned int>();
+  test<long>();
+  test<unsigned long>();
+  test<long long>();
+  test<unsigned long long>();
 #ifndef _LIBCUDACXX_HAS_NO_INT128
-  test<__int128_t, 0>();
-  test<__uint128_t, 0>();
+  test<__int128_t>();
+  test<__uint128_t>();
 #endif
-  test<float, 2 + (FLT_MANT_DIG * 30103) / 100000>();
-  test<double, 2 + (DBL_MANT_DIG * 30103) / 100000>();
+  test<float>();
+  test<double>();
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
-  test<long double, 2 + (LDBL_MANT_DIG * 30103) / 100000>();
+  test<long double>();
 #endif
 #if defined(_LIBCUDACXX_HAS_NVFP16)
-  test<__half, 5>();
+  test<__half>();
 #endif // _LIBCUDACXX_HAS_NVFP16
 #if defined(_LIBCUDACXX_HAS_NVBF16)
-  test<__nv_bfloat16, 4>();
+  test<__nv_bfloat16>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3>();
+  test<__nv_fp8_e5m2>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 81d5ae07795..3027e9f06f5 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
@@ -12,10 +12,17 @@
 
 #include <cuda/std/cfloat>
 #include <cuda/std/limits>
+#include <cuda/std/type_traits>
 
 #include "test_macros.h"
 
-template <class T, int expected>
+template <class T, cuda::std::enable_if_t<cuda::std::is_integral<T>::value, int> = 0>
+__host__ __device__ constexpr int make_expected_max_exponent()
+{
+  return 0;
+}
+
+template <class T, int expected = make_expected_max_exponent<T>()>
 __host__ __device__ void test()
 {
   static_assert(cuda::std::numeric_limits<T>::max_exponent == expected, "max_exponent test 1");
@@ -26,29 +33,29 @@ __host__ __device__ void test()
 
 int main(int, char**)
 {
-  test<bool, 0>();
-  test<char, 0>();
-  test<signed char, 0>();
-  test<unsigned char, 0>();
-  test<wchar_t, 0>();
+  test<bool>();
+  test<char>();
+  test<signed char>();
+  test<unsigned char>();
+  test<wchar_t>();
 #if TEST_STD_VER > 2017 && defined(__cpp_char8_t)
-  test<char8_t, 0>();
+  test<char8_t>();
 #endif
 #ifndef _LIBCUDACXX_HAS_NO_UNICODE_CHARS
-  test<char16_t, 0>();
-  test<char32_t, 0>();
+  test<char16_t>();
+  test<char32_t>();
 #endif // _LIBCUDACXX_HAS_NO_UNICODE_CHARS
-  test<short, 0>();
-  test<unsigned short, 0>();
-  test<int, 0>();
-  test<unsigned int, 0>();
-  test<long, 0>();
-  test<unsigned long, 0>();
-  test<long long, 0>();
-  test<unsigned long long, 0>();
+  test<short>();
+  test<unsigned short>();
+  test<int>();
+  test<unsigned int>();
+  test<long>();
+  test<unsigned long>();
+  test<long long>();
+  test<unsigned long long>();
 #ifndef _LIBCUDACXX_HAS_NO_INT128
-  test<__int128_t, 0>();
-  test<__uint128_t, 0>();
+  test<__int128_t>();
+  test<__uint128_t>();
 #endif
   test<float, FLT_MAX_EXP>();
   test<double, DBL_MAX_EXP>();
@@ -61,6 +68,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16, 128>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3, 8>();
+  test<__nv_fp8_e5m2, 15>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 4c426b37460..5924aee173d 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
@@ -12,10 +12,17 @@
 
 #include <cuda/std/cfloat>
 #include <cuda/std/limits>
+#include <cuda/std/type_traits>
 
 #include "test_macros.h"
 
-template <class T, int expected>
+template <class T, cuda::std::enable_if_t<cuda::std::is_integral<T>::value, int> = 0>
+__host__ __device__ constexpr int make_expected_max_exponent10()
+{
+  return 0;
+}
+
+template <class T, int expected = make_expected_max_exponent10<T>()>
 __host__ __device__ void test()
 {
   static_assert(cuda::std::numeric_limits<T>::max_exponent10 == expected, "max_exponent10 test 1");
@@ -26,29 +33,29 @@ __host__ __device__ void test()
 
 int main(int, char**)
 {
-  test<bool, 0>();
-  test<char, 0>();
-  test<signed char, 0>();
-  test<unsigned char, 0>();
-  test<wchar_t, 0>();
+  test<bool>();
+  test<char>();
+  test<signed char>();
+  test<unsigned char>();
+  test<wchar_t>();
 #if TEST_STD_VER > 2017 && defined(__cpp_char8_t)
-  test<char8_t, 0>();
+  test<char8_t>();
 #endif
 #ifndef _LIBCUDACXX_HAS_NO_UNICODE_CHARS
-  test<char16_t, 0>();
-  test<char32_t, 0>();
+  test<char16_t>();
+  test<char32_t>();
 #endif // _LIBCUDACXX_HAS_NO_UNICODE_CHARS
-  test<short, 0>();
-  test<unsigned short, 0>();
-  test<int, 0>();
-  test<unsigned int, 0>();
-  test<long, 0>();
-  test<unsigned long, 0>();
-  test<long long, 0>();
-  test<unsigned long long, 0>();
+  test<short>();
+  test<unsigned short>();
+  test<int>();
+  test<unsigned int>();
+  test<long>();
+  test<unsigned long>();
+  test<long long>();
+  test<unsigned long long>();
 #ifndef _LIBCUDACXX_HAS_NO_INT128
-  test<__int128_t, 0>();
-  test<__uint128_t, 0>();
+  test<__int128_t>();
+  test<__uint128_t>();
 #endif
   test<float, FLT_MAX_10_EXP>();
   test<double, DBL_MAX_10_EXP>();
@@ -61,6 +68,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16, 38>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3, 2>();
+  test<__nv_fp8_e5m2, 4>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 c24c3fde869..15f470909df 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
@@ -72,6 +72,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16>(__double2bfloat16(1.17549435082228750796873653722e-38));
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3>(make_fp8_e4m3(0.015625));
+  test<__nv_fp8_e5m2>(make_fp8_e5m2(0.000061035));
+#endif // _CCCL_HAS_NVFP8()
 
   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 e3150f8dc8e..b63d653a7c3 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
@@ -12,10 +12,17 @@
 
 #include <cuda/std/cfloat>
 #include <cuda/std/limits>
+#include <cuda/std/type_traits>
 
 #include "test_macros.h"
 
-template <class T, int expected>
+template <class T, cuda::std::enable_if_t<cuda::std::is_integral<T>::value, int> = 0>
+__host__ __device__ constexpr int make_expected_min_exponent()
+{
+  return 0;
+}
+
+template <class T, int expected = make_expected_min_exponent<T>()>
 __host__ __device__ void test()
 {
   static_assert(cuda::std::numeric_limits<T>::min_exponent == expected, "min_exponent test 1");
@@ -26,29 +33,29 @@ __host__ __device__ void test()
 
 int main(int, char**)
 {
-  test<bool, 0>();
-  test<char, 0>();
-  test<signed char, 0>();
-  test<unsigned char, 0>();
-  test<wchar_t, 0>();
+  test<bool>();
+  test<char>();
+  test<signed char>();
+  test<unsigned char>();
+  test<wchar_t>();
 #if TEST_STD_VER > 2017 && defined(__cpp_char8_t)
-  test<char8_t, 0>();
+  test<char8_t>();
 #endif
 #ifndef _LIBCUDACXX_HAS_NO_UNICODE_CHARS
-  test<char16_t, 0>();
-  test<char32_t, 0>();
+  test<char16_t>();
+  test<char32_t>();
 #endif // _LIBCUDACXX_HAS_NO_UNICODE_CHARS
-  test<short, 0>();
-  test<unsigned short, 0>();
-  test<int, 0>();
-  test<unsigned int, 0>();
-  test<long, 0>();
-  test<unsigned long, 0>();
-  test<long long, 0>();
-  test<unsigned long long, 0>();
+  test<short>();
+  test<unsigned short>();
+  test<int>();
+  test<unsigned int>();
+  test<long>();
+  test<unsigned long>();
+  test<long long>();
+  test<unsigned long long>();
 #ifndef _LIBCUDACXX_HAS_NO_INT128
-  test<__int128_t, 0>();
-  test<__uint128_t, 0>();
+  test<__int128_t>();
+  test<__uint128_t>();
 #endif
   test<float, FLT_MIN_EXP>();
   test<double, DBL_MIN_EXP>();
@@ -61,6 +68,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16, -125>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3, -6>();
+  test<__nv_fp8_e5m2, -15>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 cbca8e04171..a6ff20e7fde 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
@@ -12,10 +12,17 @@
 
 #include <cuda/std/cfloat>
 #include <cuda/std/limits>
+#include <cuda/std/type_traits>
 
 #include "test_macros.h"
 
-template <class T, int expected>
+template <class T, cuda::std::enable_if_t<cuda::std::is_integral<T>::value, int> = 0>
+__host__ __device__ constexpr int make_expected_min_exponent10()
+{
+  return 0;
+}
+
+template <class T, int expected = make_expected_min_exponent10<T>()>
 __host__ __device__ void test()
 {
   static_assert(cuda::std::numeric_limits<T>::min_exponent10 == expected, "min_exponent10 test 1");
@@ -26,29 +33,29 @@ __host__ __device__ void test()
 
 int main(int, char**)
 {
-  test<bool, 0>();
-  test<char, 0>();
-  test<signed char, 0>();
-  test<unsigned char, 0>();
-  test<wchar_t, 0>();
+  test<bool>();
+  test<char>();
+  test<signed char>();
+  test<unsigned char>();
+  test<wchar_t>();
 #if TEST_STD_VER > 2017 && defined(__cpp_char8_t)
-  test<char8_t, 0>();
+  test<char8_t>();
 #endif
 #ifndef _LIBCUDACXX_HAS_NO_UNICODE_CHARS
-  test<char16_t, 0>();
-  test<char32_t, 0>();
+  test<char16_t>();
+  test<char32_t>();
 #endif // _LIBCUDACXX_HAS_NO_UNICODE_CHARS
-  test<short, 0>();
-  test<unsigned short, 0>();
-  test<int, 0>();
-  test<unsigned int, 0>();
-  test<long, 0>();
-  test<unsigned long, 0>();
-  test<long long, 0>();
-  test<unsigned long long, 0>();
+  test<short>();
+  test<unsigned short>();
+  test<int>();
+  test<unsigned int>();
+  test<long>();
+  test<unsigned long>();
+  test<long long>();
+  test<unsigned long long>();
 #ifndef _LIBCUDACXX_HAS_NO_INT128
-  test<__int128_t, 0>();
-  test<__uint128_t, 0>();
+  test<__int128_t>();
+  test<__uint128_t>();
 #endif
   test<float, FLT_MIN_10_EXP>();
   test<double, DBL_MIN_10_EXP>();
@@ -61,6 +68,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16, -37>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3, -2>();
+  test<__nv_fp8_e5m2, -5>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 74e7f427941..2d6d9582f5c 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
@@ -18,29 +18,63 @@
 #include "test_macros.h"
 
 template <class T>
-__host__ __device__ void test_imp(cuda::std::true_type)
+__host__ __device__ bool is_nan(T x)
 {
-  assert(cuda::std::isnan(cuda::std::numeric_limits<T>::quiet_NaN()));
-  assert(cuda::std::isnan(cuda::std::numeric_limits<const T>::quiet_NaN()));
-  assert(cuda::std::isnan(cuda::std::numeric_limits<volatile T>::quiet_NaN()));
-  assert(cuda::std::isnan(cuda::std::numeric_limits<const volatile T>::quiet_NaN()));
+  return cuda::std::isnan(x);
 }
 
+#if _CCCL_HAS_NVFP8()
+__host__ __device__ bool is_nan(__nv_fp8_e4m3 x)
+{
+  return is_nan(__half{__nv_cvt_fp8_to_halfraw(x.__x, __NV_E4M3)});
+}
+
+__host__ __device__ bool is_nan(__nv_fp8_e5m2 x)
+{
+  return is_nan(__half{__nv_cvt_fp8_to_halfraw(x.__x, __NV_E5M2)});
+}
+#endif // _CCCL_HAS_NVFP8()
+
 template <class T>
-__host__ __device__ void test_imp(cuda::std::false_type)
+__host__ __device__ void test_impl(cuda::std::true_type)
 {
-  assert(cuda::std::numeric_limits<T>::quiet_NaN() == T());
-  assert(cuda::std::numeric_limits<const T>::quiet_NaN() == T());
-  assert(cuda::std::numeric_limits<volatile T>::quiet_NaN() == T());
-  assert(cuda::std::numeric_limits<const volatile T>::quiet_NaN() == T());
+  assert(is_nan(cuda::std::numeric_limits<T>::quiet_NaN()));
+  assert(is_nan(cuda::std::numeric_limits<const T>::quiet_NaN()));
+  assert(is_nan(cuda::std::numeric_limits<volatile T>::quiet_NaN()));
+  assert(is_nan(cuda::std::numeric_limits<const volatile T>::quiet_NaN()));
 }
 
 template <class T>
-__host__ __device__ inline void test()
+__host__ __device__ bool equal_to(T x, T y)
 {
-  constexpr bool is_float = cuda::std::is_floating_point<T>::value || cuda::std::__is_extended_floating_point<T>::value;
+  return x == y;
+}
+
+#if _CCCL_HAS_NVFP8()
+__host__ __device__ bool equal_to(__nv_fp8_e4m3 x, __nv_fp8_e4m3 y)
+{
+  return x.__x == y.__x;
+}
 
-  test_imp<T>(cuda::std::integral_constant<bool, is_float>{});
+__host__ __device__ bool equal_to(__nv_fp8_e5m2 x, __nv_fp8_e5m2 y)
+{
+  return x.__x == y.__x;
+}
+#endif // _CCCL_HAS_NVFP8()
+
+template <class T>
+__host__ __device__ void test_impl(cuda::std::false_type)
+{
+  assert(equal_to(cuda::std::numeric_limits<T>::signaling_NaN(), T()));
+  assert(equal_to(cuda::std::numeric_limits<const T>::signaling_NaN(), T()));
+  assert(equal_to(cuda::std::numeric_limits<volatile T>::signaling_NaN(), T()));
+  assert(equal_to(cuda::std::numeric_limits<const volatile T>::signaling_NaN(), T()));
+}
+
+template <class T>
+__host__ __device__ inline void test()
+{
+  test_impl<T>(cuda::std::integral_constant<bool, cuda::std::numeric_limits<T>::has_quiet_NaN>{});
 }
 
 int main(int, char**)
@@ -80,6 +114,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3>();
+  test<__nv_fp8_e5m2>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 9765db6f760..7e5c87927aa 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
@@ -61,6 +61,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16, FLT_RADIX>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3, FLT_RADIX>();
+  test<__nv_fp8_e5m2, FLT_RADIX>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 ba5049fc49f..d4faf373a09 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
@@ -63,6 +63,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16>(__double2bfloat16(0.5));
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3>(make_fp8_e4m3(0.5));
+  test<__nv_fp8_e5m2>(make_fp8_e5m2(0.5));
+#endif // _CCCL_HAS_NVFP8()
 
   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 3fb436381a7..8515581d650 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
@@ -60,6 +60,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16, cuda::std::round_to_nearest>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3, cuda::std::round_to_nearest>();
+  test<__nv_fp8_e5m2, cuda::std::round_to_nearest>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 69ba66038de..19ace1b3d2c 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
@@ -18,29 +18,63 @@
 #include "test_macros.h"
 
 template <class T>
-__host__ __device__ void test_imp(cuda::std::true_type)
+__host__ __device__ bool is_nan(T x)
 {
-  assert(cuda::std::isnan(cuda::std::numeric_limits<T>::signaling_NaN()));
-  assert(cuda::std::isnan(cuda::std::numeric_limits<const T>::signaling_NaN()));
-  assert(cuda::std::isnan(cuda::std::numeric_limits<volatile T>::signaling_NaN()));
-  assert(cuda::std::isnan(cuda::std::numeric_limits<const volatile T>::signaling_NaN()));
+  return cuda::std::isnan(x);
 }
 
+#if _CCCL_HAS_NVFP8()
+__host__ __device__ bool is_nan(__nv_fp8_e4m3 x)
+{
+  return is_nan(__half{__nv_cvt_fp8_to_halfraw(x.__x, __NV_E4M3)});
+}
+
+__host__ __device__ bool is_nan(__nv_fp8_e5m2 x)
+{
+  return is_nan(__half{__nv_cvt_fp8_to_halfraw(x.__x, __NV_E5M2)});
+}
+#endif // _CCCL_HAS_NVFP8()
+
 template <class T>
-__host__ __device__ void test_imp(cuda::std::false_type)
+__host__ __device__ void test_impl(cuda::std::true_type)
 {
-  assert(cuda::std::numeric_limits<T>::signaling_NaN() == T());
-  assert(cuda::std::numeric_limits<const T>::signaling_NaN() == T());
-  assert(cuda::std::numeric_limits<volatile T>::signaling_NaN() == T());
-  assert(cuda::std::numeric_limits<const volatile T>::signaling_NaN() == T());
+  assert(is_nan(cuda::std::numeric_limits<T>::signaling_NaN()));
+  assert(is_nan(cuda::std::numeric_limits<const T>::signaling_NaN()));
+  assert(is_nan(cuda::std::numeric_limits<volatile T>::signaling_NaN()));
+  assert(is_nan(cuda::std::numeric_limits<const volatile T>::signaling_NaN()));
 }
 
 template <class T>
-__host__ __device__ inline void test()
+__host__ __device__ bool equal_to(T x, T y)
 {
-  constexpr bool is_float = cuda::std::is_floating_point<T>::value || cuda::std::__is_extended_floating_point<T>::value;
+  return x == y;
+}
+
+#if _CCCL_HAS_NVFP8()
+__host__ __device__ bool equal_to(__nv_fp8_e4m3 x, __nv_fp8_e4m3 y)
+{
+  return x.__x == y.__x;
+}
 
-  test_imp<T>(cuda::std::integral_constant<bool, is_float>{});
+__host__ __device__ bool equal_to(__nv_fp8_e5m2 x, __nv_fp8_e5m2 y)
+{
+  return x.__x == y.__x;
+}
+#endif // _CCCL_HAS_NVFP8()
+
+template <class T>
+__host__ __device__ void test_impl(cuda::std::false_type)
+{
+  assert(equal_to(cuda::std::numeric_limits<T>::signaling_NaN(), T()));
+  assert(equal_to(cuda::std::numeric_limits<const T>::signaling_NaN(), T()));
+  assert(equal_to(cuda::std::numeric_limits<volatile T>::signaling_NaN(), T()));
+  assert(equal_to(cuda::std::numeric_limits<const volatile T>::signaling_NaN(), T()));
+}
+
+template <class T>
+__host__ __device__ inline void test()
+{
+  test_impl<T>(cuda::std::integral_constant<bool, cuda::std::numeric_limits<T>::has_signaling_NaN>{});
 }
 
 int main(int, char**)
@@ -80,6 +114,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3>();
+  test<__nv_fp8_e5m2>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 70cde2711a1..38dec8c872b 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
@@ -60,6 +60,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16, false>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3, false>();
+  test<__nv_fp8_e5m2, false>();
+#endif // _CCCL_HAS_NVFP8()
 
   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 7dd7eee68cc..55d7eb990db 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
@@ -66,6 +66,10 @@ int main(int, char**)
 #if defined(_LIBCUDACXX_HAS_NVBF16)
   test<__nv_bfloat16, false>();
 #endif // _LIBCUDACXX_HAS_NVBF16
+#if _CCCL_HAS_NVFP8()
+  test<__nv_fp8_e4m3, false>();
+  test<__nv_fp8_e5m2, false>();
+#endif // _CCCL_HAS_NVFP8()
 
   return 0;
 }