Support FP16 traits on CTK 12.0

NVIDIA · Jan 27, 2025 · 8653d48 · 8653d48
1 parent abfb7b4
commit 8653d48
Show file tree

Hide file tree

Showing 37 changed files with 152 additions and 152 deletions.
diff --git a/libcudacxx/include/cuda/std/__type_traits/is_extended_floating_point.h b/libcudacxx/include/cuda/std/__type_traits/is_extended_floating_point.h
@@ -53,7 +53,7 @@ _CCCL_INLINE_VAR constexpr bool __is_extended_floating_point_v
 #  endif // !_CCCL_NO_INLINE_VARIABLES
 #endif // !_CCCL_NO_VARIABLE_TEMPLATES
 
-#if defined(_LIBCUDACXX_HAS_NVFP16)
+#if defined(_CCCL_HAS_FP16)
 template <>
 struct __is_extended_floating_point<__half> : true_type
 {};
@@ -62,9 +62,9 @@ struct __is_extended_floating_point<__half> : true_type
 template <>
 _CCCL_INLINE_VAR constexpr bool __is_extended_floating_point_v<__half> = true;
 #  endif // !_CCCL_NO_INLINE_VARIABLES
-#endif // _LIBCUDACXX_HAS_NVFP16
+#endif // _CCCL_HAS_FP16
 
-#if defined(_LIBCUDACXX_HAS_NVBF16)
+#if defined(_CCCL_HAS_BF16)
 template <>
 struct __is_extended_floating_point<__nv_bfloat16> : true_type
 {};
@@ -73,7 +73,7 @@ struct __is_extended_floating_point<__nv_bfloat16> : true_type
 template <>
 _CCCL_INLINE_VAR constexpr bool __is_extended_floating_point_v<__nv_bfloat16> = true;
 #  endif // !_CCCL_NO_INLINE_VARIABLES
-#endif // _LIBCUDACXX_HAS_NVBF16
+#endif // _CCCL_HAS_BF16
 
 #if _CCCL_HAS_NVFP8()
 template <>

diff --git a/libcudacxx/include/cuda/std/limits b/libcudacxx/include/cuda/std/limits
@@ -608,7 +608,7 @@ public:
 #endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE
 };
 
-#if defined(_LIBCUDACXX_HAS_NVFP16)
+#if defined(_CCCL_HAS_FP16)
 template <>
 class __numeric_limits_impl<__half, __numeric_limits_type::__floating_point>
 {
@@ -681,9 +681,9 @@ public:
   static constexpr bool tinyness_before          = false;
   static constexpr float_round_style round_style = round_to_nearest;
 };
-#endif // _LIBCUDACXX_HAS_NVFP16
+#endif // _CCCL_HAS_FP16
 
-#if defined(_LIBCUDACXX_HAS_NVBF16)
+#if defined(_CCCL_HAS_BF16)
 template <>
 class __numeric_limits_impl<__nv_bfloat16, __numeric_limits_type::__floating_point>
 {
@@ -756,7 +756,7 @@ public:
   static constexpr bool tinyness_before          = false;
   static constexpr float_round_style round_style = round_to_nearest;
 };
-#endif // _LIBCUDACXX_HAS_NVBF16
+#endif // _CCCL_HAS_BF16
 
 #if _CCCL_HAS_NVFP8()
 #  if defined(_CCCL_BUILTIN_BIT_CAST) || _CCCL_STD_VER >= 2014

diff --git a/...est/libcudacxx/libcxx/utilities/meta/meta.unary/meta.unary.cat/is_floating_point.pass.cpp b/...est/libcudacxx/libcxx/utilities/meta/meta.unary/meta.unary.cat/is_floating_point.pass.cpp
@@ -80,12 +80,12 @@ int main(int, char**)
   test_is_floating_point<float>();
   test_is_floating_point<double>();
   test_is_floating_point<long double>();
-#ifdef _LIBCUDACXX_HAS_NVFP16
+#ifdef _CCCL_HAS_FP16
   test_is_floating_point<__half>();
-#endif // _LIBCUDACXX_HAS_NVFP16
-#ifdef _LIBCUDACXX_HAS_NVBF16
+#endif // _CCCL_HAS_FP16
+#ifdef _CCCL_HAS_BF16
   test_is_floating_point<__nv_bfloat16>();
-#endif // _LIBCUDACXX_HAS_NVBF16
+#endif // _CCCL_HAS_BF16
 #if _CCCL_HAS_NVFP8()
   test_is_floating_point<__nv_fp8_e4m3>();
   test_is_floating_point<__nv_fp8_e5m2>();

diff --git a/...udacxx/test/libcudacxx/std/language.support/support.limits/limits/is_specialized.pass.cpp b/...udacxx/test/libcudacxx/std/language.support/support.limits/limits/is_specialized.pass.cpp
@@ -68,12 +68,12 @@ int main(int, char**)
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double>();
 #endif
-#if defined(_LIBCUDACXX_HAS_NVFP16)
+#if defined(_CCCL_HAS_FP16)
   test<__half>();
-#endif // _LIBCUDACXX_HAS_NVFP16
-#if defined(_LIBCUDACXX_HAS_NVBF16)
+#endif // _CCCL_HAS_FP16
+#if defined(_CCCL_HAS_BF16)
   test<__nv_bfloat16>();
-#endif // _LIBCUDACXX_HAS_NVBF16
+#endif // _CCCL_HAS_BF16
 
   static_assert(!cuda::std::numeric_limits<cuda::std::complex<double>>::is_specialized,
                 "!cuda::std::numeric_limits<cuda::std::complex<double> >::is_specialized");

diff --git a/...est/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/common.h b/...est/libcudacxx/std/language.support/support.limits/limits/numeric.limits.members/common.h
@@ -51,18 +51,18 @@ __host__ __device__ inline bool float_eq(__nv_fp8_e5m2 x, __nv_fp8_e5m2 y)
 }
 #endif // _CCCL_HAS_NVFP8
 
-#if defined(_LIBCUDACXX_HAS_NVFP16)
+#if defined(_CCCL_HAS_FP16)
 __host__ __device__ inline bool float_eq(__half x, __half y)
 {
   return __heq(x, y);
 }
-#endif // _LIBCUDACXX_HAS_NVFP16
+#endif // _CCCL_HAS_FP16
 
-#if defined(_LIBCUDACXX_HAS_NVBF16)
+#if defined(_CCCL_HAS_BF16)
 __host__ __device__ inline bool float_eq(__nv_bfloat16 x, __nv_bfloat16 y)
 {
   return __heq(x, y);
 }
-#endif // _LIBCUDACXX_HAS_NVBF16
+#endif // _CCCL_HAS_BF16
 
 #endif // NUMERIC_LIMITS_MEMBERS_COMMON_H
diff --git a/...language.support/support.limits/limits/numeric.limits.members/const_data_members.pass.cpp b/...language.support/support.limits/limits/numeric.limits.members/const_data_members.pass.cpp
@@ -110,12 +110,12 @@ int main(int, char**)
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test_type<long double>();
 #endif // _LIBCUDACXX_HAS_NO_LONG_DOUBLE
-#if defined(_LIBCUDACXX_HAS_NVFP16)
+#if defined(_CCCL_HAS_FP16)
   test_type<__half>();
-#endif // _LIBCUDACXX_HAS_NVFP16
-#if defined(_LIBCUDACXX_HAS_NVBF16)
+#endif // _CCCL_HAS_FP16
+#if defined(_CCCL_HAS_BF16)
   test_type<__nv_bfloat16>();
-#endif // _LIBCUDACXX_HAS_NVBF16
+#endif // _CCCL_HAS_BF16
 #if _CCCL_HAS_NVFP8()
   test_type<__nv_fp8_e4m3>();
   test_type<__nv_fp8_e5m2>();

diff --git a/...cxx/std/language.support/support.limits/limits/numeric.limits.members/denorm_min.pass.cpp b/...cxx/std/language.support/support.limits/limits/numeric.limits.members/denorm_min.pass.cpp
@@ -66,12 +66,12 @@ int main(int, char**)
   test<long double>(LDBL_TRUE_MIN);
 #  endif
 #endif
-#if defined(_LIBCUDACXX_HAS_NVFP16)
+#if defined(_CCCL_HAS_FP16)
   test<__half>(__double2half(5.9604644775390625e-08));
-#endif // _LIBCUDACXX_HAS_NVFP16
-#if defined(_LIBCUDACXX_HAS_NVBF16)
+#endif // _CCCL_HAS_FP16
+#if defined(_CCCL_HAS_BF16)
   test<__nv_bfloat16>(__double2bfloat16(9.18354961579912115600575419705e-41));
-#endif // _LIBCUDACXX_HAS_NVBF16
+#endif // _CCCL_HAS_BF16
 #if _CCCL_HAS_NVFP8()
   test<__nv_fp8_e4m3>(make_fp8_e4m3(0.001953125));
   test<__nv_fp8_e5m2>(make_fp8_e5m2(0.0000152587890625));

diff --git a/...cudacxx/std/language.support/support.limits/limits/numeric.limits.members/digits.pass.cpp b/...cudacxx/std/language.support/support.limits/limits/numeric.limits.members/digits.pass.cpp
@@ -55,12 +55,12 @@ int main(int, char**)
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double, LDBL_MANT_DIG>();
 #endif
-#if defined(_LIBCUDACXX_HAS_NVFP16)
+#if defined(_CCCL_HAS_FP16)
   test<__half, 11>();
-#endif // _LIBCUDACXX_HAS_NVFP16
-#if defined(_LIBCUDACXX_HAS_NVBF16)
+#endif // _CCCL_HAS_FP16
+#if defined(_CCCL_HAS_BF16)
   test<__nv_bfloat16, 8>();
-#endif // _LIBCUDACXX_HAS_NVBF16
+#endif // _CCCL_HAS_BF16
 #if _CCCL_HAS_NVFP8()
   test<__nv_fp8_e4m3, 3>();
   test<__nv_fp8_e5m2, 2>();

diff --git a/...dacxx/std/language.support/support.limits/limits/numeric.limits.members/digits10.pass.cpp b/...dacxx/std/language.support/support.limits/limits/numeric.limits.members/digits10.pass.cpp
@@ -74,12 +74,12 @@ int main(int, char**)
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double>();
 #endif
-#if defined(_LIBCUDACXX_HAS_NVFP16)
+#if defined(_CCCL_HAS_FP16)
   test<__half>();
-#endif // _LIBCUDACXX_HAS_NVFP16
-#if defined(_LIBCUDACXX_HAS_NVBF16)
+#endif // _CCCL_HAS_FP16
+#if defined(_CCCL_HAS_BF16)
   test<__nv_bfloat16>();
-#endif // _LIBCUDACXX_HAS_NVBF16
+#endif // _CCCL_HAS_BF16
 #if _CCCL_HAS_NVFP8()
   test<__nv_fp8_e4m3>();
   test<__nv_fp8_e5m2>();

diff --git a/...udacxx/std/language.support/support.limits/limits/numeric.limits.members/epsilon.pass.cpp b/...udacxx/std/language.support/support.limits/limits/numeric.limits.members/epsilon.pass.cpp
@@ -57,12 +57,12 @@ int main(int, char**)
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double>(LDBL_EPSILON);
 #endif
-#if defined(_LIBCUDACXX_HAS_NVFP16)
+#if defined(_CCCL_HAS_FP16)
   test<__half>(__double2half(0.0009765625));
-#endif // _LIBCUDACXX_HAS_NVFP16
-#if defined(_LIBCUDACXX_HAS_NVBF16)
+#endif // _CCCL_HAS_FP16
+#if defined(_CCCL_HAS_BF16)
   test<__nv_bfloat16>(__double2bfloat16(0.0078125));
-#endif // _LIBCUDACXX_HAS_NVBF16
+#endif // _CCCL_HAS_BF16
 #if _CCCL_HAS_NVFP8()
   test<__nv_fp8_e4m3>(make_fp8_e4m3(0.125));
   test<__nv_fp8_e5m2>(make_fp8_e5m2(0.25));

diff --git a/...cxx/std/language.support/support.limits/limits/numeric.limits.members/has_denorm.pass.cpp b/...cxx/std/language.support/support.limits/limits/numeric.limits.members/has_denorm.pass.cpp
@@ -54,12 +54,12 @@ int main(int, char**)
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double, cuda::std::denorm_present>();
 #endif
-#if defined(_LIBCUDACXX_HAS_NVFP16)
+#if defined(_CCCL_HAS_FP16)
   test<__half, cuda::std::denorm_present>();
-#endif // _LIBCUDACXX_HAS_NVFP16
-#if defined(_LIBCUDACXX_HAS_NVBF16)
+#endif // _CCCL_HAS_FP16
+#if defined(_CCCL_HAS_BF16)
   test<__nv_bfloat16, cuda::std::denorm_present>();
-#endif // _LIBCUDACXX_HAS_NVBF16
+#endif // _CCCL_HAS_BF16
 #if _CCCL_HAS_NVFP8()
   test<__nv_fp8_e4m3, cuda::std::denorm_present>();
   test<__nv_fp8_e5m2, cuda::std::denorm_present>();

diff --git a/...td/language.support/support.limits/limits/numeric.limits.members/has_denorm_loss.pass.cpp b/...td/language.support/support.limits/limits/numeric.limits.members/has_denorm_loss.pass.cpp
@@ -54,12 +54,12 @@ int main(int, char**)
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double, false>();
 #endif
-#if defined(_LIBCUDACXX_HAS_NVFP16)
+#if defined(_CCCL_HAS_FP16)
   test<__half, false>();
-#endif // _LIBCUDACXX_HAS_NVFP16
-#if defined(_LIBCUDACXX_HAS_NVBF16)
+#endif // _CCCL_HAS_FP16
+#if defined(_CCCL_HAS_BF16)
   test<__nv_bfloat16, false>();
-#endif // _LIBCUDACXX_HAS_NVBF16
+#endif // _CCCL_HAS_BF16
 #if _CCCL_HAS_NVFP8()
   test<__nv_fp8_e4m3, false>();
   test<__nv_fp8_e5m2, false>();

diff --git a/...x/std/language.support/support.limits/limits/numeric.limits.members/has_infinity.pass.cpp b/...x/std/language.support/support.limits/limits/numeric.limits.members/has_infinity.pass.cpp
@@ -54,12 +54,12 @@ int main(int, char**)
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double, true>();
 #endif
-#if defined(_LIBCUDACXX_HAS_NVFP16)
+#if defined(_CCCL_HAS_FP16)
   test<__half, true>();
-#endif // _LIBCUDACXX_HAS_NVFP16
-#if defined(_LIBCUDACXX_HAS_NVBF16)
+#endif // _CCCL_HAS_FP16
+#if defined(_CCCL_HAS_BF16)
   test<__nv_bfloat16, true>();
-#endif // _LIBCUDACXX_HAS_NVBF16
+#endif // _CCCL_HAS_BF16
 #if _CCCL_HAS_NVFP8()
   test<__nv_fp8_e4m3, false>();
   test<__nv_fp8_e5m2, true>();

diff --git a/.../std/language.support/support.limits/limits/numeric.limits.members/has_quiet_NaN.pass.cpp b/.../std/language.support/support.limits/limits/numeric.limits.members/has_quiet_NaN.pass.cpp
@@ -54,12 +54,12 @@ int main(int, char**)
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double, true>();
 #endif
-#if defined(_LIBCUDACXX_HAS_NVFP16)
+#if defined(_CCCL_HAS_FP16)
   test<__half, true>();
-#endif // _LIBCUDACXX_HAS_NVFP16
-#if defined(_LIBCUDACXX_HAS_NVBF16)
+#endif // _CCCL_HAS_FP16
+#if defined(_CCCL_HAS_BF16)
   test<__nv_bfloat16, true>();
-#endif // _LIBCUDACXX_HAS_NVBF16
+#endif // _CCCL_HAS_BF16
 #if _CCCL_HAS_NVFP8()
   test<__nv_fp8_e4m3, true>();
   test<__nv_fp8_e5m2, true>();

diff --git a/.../language.support/support.limits/limits/numeric.limits.members/has_signaling_NaN.pass.cpp b/.../language.support/support.limits/limits/numeric.limits.members/has_signaling_NaN.pass.cpp
@@ -54,12 +54,12 @@ int main(int, char**)
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double, true>();
 #endif
-#if defined(_LIBCUDACXX_HAS_NVFP16)
+#if defined(_CCCL_HAS_FP16)
   test<__half, true>();
-#endif // _LIBCUDACXX_HAS_NVFP16
-#if defined(_LIBCUDACXX_HAS_NVBF16)
+#endif // _CCCL_HAS_FP16
+#if defined(_CCCL_HAS_BF16)
   test<__nv_bfloat16, true>();
-#endif // _LIBCUDACXX_HAS_NVBF16
+#endif // _CCCL_HAS_BF16
 #if _CCCL_HAS_NVFP8()
   test<__nv_fp8_e4m3, false>();
   test<__nv_fp8_e5m2, true>();

diff --git a/...dacxx/std/language.support/support.limits/limits/numeric.limits.members/infinity.pass.cpp b/...dacxx/std/language.support/support.limits/limits/numeric.limits.members/infinity.pass.cpp
@@ -64,12 +64,12 @@ int main(int, char**)
 #  ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double>(1. / 0.);
 #  endif
-#  if defined(_LIBCUDACXX_HAS_NVFP16)
+#  if defined(_CCCL_HAS_FP16)
   test<__half>(__double2half(1.0 / 0.0));
-#  endif // _LIBCUDACXX_HAS_NVFP16
-#  if defined(_LIBCUDACXX_HAS_NVBF16)
+#  endif // _CCCL_HAS_FP16
+#  if defined(_CCCL_HAS_BF16)
   test<__nv_bfloat16>(__double2bfloat16(1.0 / 0.0));
-#  endif // _LIBCUDACXX_HAS_NVBF16
+#  endif // _CCCL_HAS_BF16
 #  if _CCCL_HAS_NVFP8()
   test<__nv_fp8_e4m3>(__nv_fp8_e4m3{});
   test<__nv_fp8_e5m2>(make_fp8_e5m2(1.0 / 0.0));
@@ -81,12 +81,12 @@ int main(int, char**)
 #  ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double>(INFINITY);
 #  endif
-#  if defined(_LIBCUDACXX_HAS_NVFP16)
+#  if defined(_CCCL_HAS_FP16)
   test<__half>(__double2half(INFINITY));
-#  endif // _LIBCUDACXX_HAS_NVFP16
-#  if defined(_LIBCUDACXX_HAS_NVBF16)
+#  endif // _CCCL_HAS_FP16
+#  if defined(_CCCL_HAS_BF16)
   test<__nv_bfloat16>(__double2bfloat16(INFINITY));
-#  endif // _LIBCUDACXX_HAS_NVBF16
+#  endif // _CCCL_HAS_BF16
 #  if _CCCL_HAS_NVFP8()
   test<__nv_fp8_e4m3>(__nv_fp8_e4m3{});
   test<__nv_fp8_e5m2>(make_fp8_e5m2(INFINITY));

diff --git a/...cxx/std/language.support/support.limits/limits/numeric.limits.members/is_bounded.pass.cpp b/...cxx/std/language.support/support.limits/limits/numeric.limits.members/is_bounded.pass.cpp
@@ -54,12 +54,12 @@ int main(int, char**)
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double, true>();
 #endif
-#if defined(_LIBCUDACXX_HAS_NVFP16)
+#if defined(_CCCL_HAS_FP16)
   test<__half, true>();
-#endif // _LIBCUDACXX_HAS_NVFP16
-#if defined(_LIBCUDACXX_HAS_NVBF16)
+#endif // _CCCL_HAS_FP16
+#if defined(_CCCL_HAS_BF16)
   test<__nv_bfloat16, true>();
-#endif // _LIBCUDACXX_HAS_NVBF16
+#endif // _CCCL_HAS_BF16
 #if _CCCL_HAS_NVFP8()
   test<__nv_fp8_e4m3, true>();
   test<__nv_fp8_e5m2, true>();

diff --git a/...dacxx/std/language.support/support.limits/limits/numeric.limits.members/is_exact.pass.cpp b/...dacxx/std/language.support/support.limits/limits/numeric.limits.members/is_exact.pass.cpp
@@ -54,12 +54,12 @@ int main(int, char**)
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double, false>();
 #endif
-#if defined(_LIBCUDACXX_HAS_NVFP16)
+#if defined(_CCCL_HAS_FP16)
   test<__half, false>();
-#endif // _LIBCUDACXX_HAS_NVFP16
-#if defined(_LIBCUDACXX_HAS_NVBF16)
+#endif // _CCCL_HAS_FP16
+#if defined(_CCCL_HAS_BF16)
   test<__nv_bfloat16, false>();
-#endif // _LIBCUDACXX_HAS_NVBF16
+#endif // _CCCL_HAS_BF16
 #if _CCCL_HAS_NVFP8()
   test<__nv_fp8_e4m3, false>();
   test<__nv_fp8_e5m2, false>();

diff --git a/...acxx/std/language.support/support.limits/limits/numeric.limits.members/is_iec559.pass.cpp b/...acxx/std/language.support/support.limits/limits/numeric.limits.members/is_iec559.pass.cpp
@@ -54,12 +54,12 @@ int main(int, char**)
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double, true>();
 #endif // _LIBCUDACXX_HAS_NO_LONG_DOUBLE
-#if defined(_LIBCUDACXX_HAS_NVFP16)
+#if defined(_CCCL_HAS_FP16)
   test<__half, true>();
-#endif // _LIBCUDACXX_HAS_NVFP16
-#if defined(_LIBCUDACXX_HAS_NVBF16)
+#endif // _CCCL_HAS_FP16
+#if defined(_CCCL_HAS_BF16)
   test<__nv_bfloat16, true>();
-#endif // _LIBCUDACXX_HAS_NVBF16
+#endif // _CCCL_HAS_BF16
 #if _CCCL_HAS_NVFP8()
   test<__nv_fp8_e4m3, false>();
   test<__nv_fp8_e5m2, false>();

diff --git a/...cxx/std/language.support/support.limits/limits/numeric.limits.members/is_integer.pass.cpp b/...cxx/std/language.support/support.limits/limits/numeric.limits.members/is_integer.pass.cpp
@@ -54,12 +54,12 @@ int main(int, char**)
 #ifndef _LIBCUDACXX_HAS_NO_LONG_DOUBLE
   test<long double, false>();
 #endif
-#if defined(_LIBCUDACXX_HAS_NVFP16)
+#if defined(_CCCL_HAS_FP16)
   test<__half, false>();
-#endif // _LIBCUDACXX_HAS_NVFP16
-#if defined(_LIBCUDACXX_HAS_NVBF16)
+#endif // _CCCL_HAS_FP16
+#if defined(_CCCL_HAS_BF16)
   test<__nv_bfloat16, false>();
-#endif // _LIBCUDACXX_HAS_NVBF16
+#endif // _CCCL_HAS_BF16
 #if _CCCL_HAS_NVFP8()
   test<__nv_fp8_e4m3, false>();
   test<__nv_fp8_e5m2, false>();