implement _CCCL_HAS_INCLUDE (NVIDIA#2786)

* `implement `_CCCL_HAS_INCLUDE`

.clang-format

@@ -130,6 +130,8 @@ MaxEmptyLinesToKeep: 1
 Macros:
 - _LIBCUDACXX_TEMPLATE(...)=template<...>
 - _LIBCUDACXX_REQUIRES(...)=requires (...)
+WhitespaceSensitiveMacros:
+- _CCCL_HAS_INCLUDE
 NamespaceIndentation: None
 PackConstructorInitializers: Never
 PenaltyBreakAssignment: 30

cub/cub/detail/nvtx.cuh

@@ -47,14 +47,16 @@
 // * NVTX is not explicitly disabled (via CCCL_DISABLE_NVTX or NVTX_DISABLE)
 // * C++14 is availabl for cuda::std::optional
 // * NVTX3 uses module as an identifier, which trips up NVHPC
-#if __has_include(<nvtx3/nvToolsExt.h> ) && !defined(CCCL_DISABLE_NVTX) && !defined(NVTX_DISABLE) \
-&& _CCCL_STD_VER >= 2014 &&(!defined(_CCCL_COMPILER_NVHPC) || _CCCL_STD_VER <= 2017)
+#if _CCCL_HAS_INCLUDE(<nvtx3/nvToolsExt.h>) && !defined(CCCL_DISABLE_NVTX) && !defined(NVTX_DISABLE) \
+                      && _CCCL_STD_VER >= 2014                                                       \
+                      && (!defined(_CCCL_COMPILER_NVHPC) || _CCCL_STD_VER <= 2017)
 // Include our NVTX3 C++ wrapper if not available from the CTK
-#  if __has_include(<nvtx3/nvtx3.hpp>) // TODO(bgruber): replace by a check for the first CTK version shipping the header
+// TODO(bgruber): replace by a check for the first CTK version shipping the header
+#  if _CCCL_HAS_INCLUDE(<nvtx3/nvtx3.hpp>)
 #    include <nvtx3/nvtx3.hpp>
-#  else // __has_include(<nvtx3/nvtx3.hpp>)
+#  else // _CCCL_HAS_INCLUDE(<nvtx3/nvtx3.hpp>)
 #    include "nvtx3.hpp"
-#  endif // __has_include(<nvtx3/nvtx3.hpp>)
+#  endif // _CCCL_HAS_INCLUDE(<nvtx3/nvtx3.hpp>)
 
 // We expect the NVTX3 V1 C++ API to be available when nvtx3.hpp is available. This should work, because newer versions
 // of NVTX3 will continue to declare previous API versions. See also:
@@ -103,9 +105,9 @@ CUB_NAMESPACE_END
 #    define CUB_DETAIL_NVTX_RANGE_SCOPE_IF(condition, name)
 #    define CUB_DETAIL_NVTX_RANGE_SCOPE(name)
 #  endif // NVTX3_CPP_DEFINITIONS_V1_0
-#else // __has_include(<nvtx3/nvToolsExt.h> ) && !defined(CCCL_DISABLE_NVTX) && !defined(NVTX_DISABLE) && _CCCL_STD_VER
-      // >= 2014
+#else // _CCCL_HAS_INCLUDE(<nvtx3/nvToolsExt.h> ) && !defined(CCCL_DISABLE_NVTX) && !defined(NVTX_DISABLE) &&
+      // _CCCL_STD_VER >= 2014
 #  define CUB_DETAIL_NVTX_RANGE_SCOPE_IF(condition, name)
 #  define CUB_DETAIL_NVTX_RANGE_SCOPE(name)
-#endif // __has_include(<nvtx3/nvToolsExt.h> ) && !defined(CCCL_DISABLE_NVTX) && !defined(NVTX_DISABLE) && _CCCL_STD_VER
-       // >= 2014
+#endif // _CCCL_HAS_INCLUDE(<nvtx3/nvToolsExt.h> ) && !defined(CCCL_DISABLE_NVTX) && !defined(NVTX_DISABLE) &&
+       // _CCCL_STD_VER >= 2014

cudax/include/cuda/experimental/__async/stop_token.cuh

@@ -29,7 +29,7 @@
 #include <cuda/experimental/__async/utility.cuh>
 #include <cuda/experimental/__detail/config.cuh>
 
-#if __has_include(<stop_token>) && __cpp_lib_jthread >= 201911
+#if _CCCL_HAS_INCLUDE(<stop_token>) && __cpp_lib_jthread >= 201911
 #  include <stop_token>
 #endif
 

cudax/include/cuda/experimental/__stf/utility/cuda_safe_call.cuh

@@ -38,14 +38,14 @@
 
 #include <cuda/experimental/__stf/utility/unittest.cuh>
 
-#if __has_include(<cusolverDn.h>)
+#if _CCCL_HAS_INCLUDE(<cusolverDn.h>)
 #  include <cusolverDn.h>
 #endif
 
 namespace cuda::experimental::stf
 {
 
-#if __has_include(<cusolverDn.h>)
+#if _CCCL_HAS_INCLUDE(<cusolverDn.h>)
 // Undocumented
 inline const char* cusolverGetErrorString(const cusolverStatus_t status)
 {
@@ -115,10 +115,10 @@ public:
   {
     // All "success" statuses are zero
     static_assert(cudaSuccess == 0 && CUDA_SUCCESS == 0
-#if __has_include(<cublas_v2.h>)
+#if _CCCL_HAS_INCLUDE(<cublas_v2.h>)
                     && CUBLAS_STATUS_SUCCESS == 0
 #endif
-#if __has_include(<cusolverDn.h>)
+#if _CCCL_HAS_INCLUDE(<cusolverDn.h>)
                     && CUSOLVER_STATUS_SUCCESS == 0
 #endif
                   ,
@@ -133,7 +133,7 @@ public:
     int dev = -1;
     cudaGetDevice(&dev);
 
-#if __has_include(<cusolverDn.h>)
+#if _CCCL_HAS_INCLUDE(<cusolverDn.h>)
     if constexpr (::std::is_same_v<T, cusolverStatus_t>)
     {
       format("%s(%u) [device %d] CUSOLVER error in call %s: %s.",
@@ -144,8 +144,8 @@ public:
              cusolverGetErrorString(status));
     }
     else
-#endif // __has_include(<cusolverDn.h>)
-#if __has_include(<cublas_v2.h>)
+#endif // _CCCL_HAS_INCLUDE(<cusolverDn.h>)
+#if _CCCL_HAS_INCLUDE(<cublas_v2.h>)
       if constexpr (::std::is_same_v<T, cublasStatus_t>)
     {
       format("%s(%u) [device %d] CUBLAS error in %s: %s.",
@@ -156,7 +156,7 @@ public:
              cublasGetStatusString(status));
     }
     else
-#endif // __has_include(<cublas_v2.h>)
+#endif // _CCCL_HAS_INCLUDE(<cublas_v2.h>)
       if constexpr (::std::is_same_v<T, cudaOccError>)
       {
         format("%s(%u) [device %d] CUDA OCC error in %s: %s.",
@@ -226,7 +226,7 @@ UNITTEST("cuda_exception")
 {
   auto e = cuda_exception(CUDA_SUCCESS);
   EXPECT(e.what()[0] == 0);
-#  if __has_include(<cusolverDn.h>)
+#  if _CCCL_HAS_INCLUDE(<cusolverDn.h>)
   auto e1 = cuda_exception(CUSOLVER_STATUS_ZERO_PIVOT);
   EXPECT(strlen(e1.what()) > 0u);
 #  endif

cudax/include/cuda/experimental/__stf/utility/nvtx.cuh

@@ -26,7 +26,7 @@
 #  pragma system_header
 #endif // no system header
 
-#if __has_include(<nvtx3/nvToolsExt.h>) &&(!defined(_CCCL_COMPILER_NVHPC) || _CCCL_STD_VER <= 2017)
+#if _CCCL_HAS_INCLUDE(<nvtx3/nvToolsExt.h>) && (!defined(_CCCL_COMPILER_NVHPC) || _CCCL_STD_VER <= 2017)
 #  include <nvtx3/nvToolsExt.h>
 #endif
 
@@ -41,7 +41,7 @@ class nvtx_range
 public:
   nvtx_range(const char* message)
   {
-#if __has_include(<nvtx3/nvToolsExt.h>)&&(!defined(_CCCL_COMPILER_NVHPC) || _CCCL_STD_VER <= 2017)
+#if _CCCL_HAS_INCLUDE(<nvtx3/nvToolsExt.h>) && (!defined(_CCCL_COMPILER_NVHPC) || _CCCL_STD_VER <= 2017)
     nvtxRangePushA(message);
 #endif
   }
@@ -53,7 +53,7 @@ public:
 
   ~nvtx_range()
   {
-#if __has_include(<nvtx3/nvToolsExt.h>)&&(!defined(_CCCL_COMPILER_NVHPC) || _CCCL_STD_VER <= 2017)
+#if _CCCL_HAS_INCLUDE(<nvtx3/nvToolsExt.h>) && (!defined(_CCCL_COMPILER_NVHPC) || _CCCL_STD_VER <= 2017)
     nvtxRangePop();
 #endif
   }

libcudacxx/include/cuda/std/__cccl/assert.h

@@ -25,6 +25,7 @@
 #include <cuda/std/__cccl/attributes.h>
 #include <cuda/std/__cccl/builtin.h>
 #include <cuda/std/__cccl/execution_space.h>
+#include <cuda/std/__cccl/preprocessor.h>
 
 #if !defined(_CCCL_COMPILER_NVRTC)
 #  include <assert.h>
@@ -63,7 +64,7 @@
 //! _CCCL_ASSERT_IMPL_HOST should never be used directly
 #if defined(_CCCL_COMPILER_NVRTC) // There is no host standard library in nvrtc
 #  define _CCCL_ASSERT_IMPL_HOST(expression, message) ((void) 0)
-#elif __has_include(<yvals.h>) && defined(_CCCL_COMPILER_MSVC) // MSVC uses _STL_VERIFY from <yvals.h>
+#elif _CCCL_HAS_INCLUDE(<yvals.h>) && defined(_CCCL_COMPILER_MSVC) // MSVC uses _STL_VERIFY from <yvals.h>
 #  include <yvals.h>
 #  define _CCCL_ASSERT_IMPL_HOST(expression, message) _STL_VERIFY(expression, message)
 #else // ^^^ MSVC STL ^^^ / vvv !MSVC STL vvv

libcudacxx/include/cuda/std/__cccl/extended_floating_point.h

@@ -23,24 +23,18 @@
 #endif // no system header
 
 #include <cuda/std/__cccl/diagnostic.h>
-
-#ifndef __has_include
-#  define __has_include(x) 0
-#endif // !__has_include
+#include <cuda/std/__cccl/preprocessor.h>
 
 #if !defined(_CCCL_HAS_NVFP16)
-#  if __has_include(<cuda_fp16.h>)                                              \
-    && (defined(_CCCL_CUDA_COMPILER) || defined(LIBCUDACXX_ENABLE_HOST_NVFP16)) \
-    && !defined(CCCL_DISABLE_FP16_SUPPORT)
+#  if _CCCL_HAS_INCLUDE(<cuda_fp16.h>) && (defined(_CCCL_CUDA_COMPILER) || defined(LIBCUDACXX_ENABLE_HOST_NVFP16)) \
+                        && !defined(CCCL_DISABLE_FP16_SUPPORT)
 #    define _CCCL_HAS_NVFP16 1
 #  endif
 #endif // !_CCCL_HAS_NVFP16
 
 #if !defined(_CCCL_HAS_NVBF16)
-#  if __has_include(<cuda_bf16.h>)         \
-    && defined(_CCCL_HAS_NVFP16)           \
-    && !defined(CCCL_DISABLE_BF16_SUPPORT) \
-    && !defined(CUB_DISABLE_BF16_SUPPORT)
+#  if _CCCL_HAS_INCLUDE(<cuda_bf16.h>) && defined(_CCCL_HAS_NVFP16) && !defined(CCCL_DISABLE_BF16_SUPPORT) \
+                        && !defined(CUB_DISABLE_BF16_SUPPORT)
 #    define _CCCL_HAS_NVBF16 1
 #  endif
 #endif // !_CCCL_HAS_NVBF16

libcudacxx/include/cuda/std/__cccl/preprocessor.h

@@ -11,6 +11,12 @@
 #ifndef __CCCL_PREPROCESSOR_H
 #define __CCCL_PREPROCESSOR_H
 
+#ifdef __has_include
+#  define _CCCL_HAS_INCLUDE(_X) __has_include(_X)
+#else
+#  define _CCCL_HAS_INCLUDE(_X) 0
+#endif
+
 #define _CCCL_PP_EXPAND(...)   __VA_ARGS__
 #define _CCCL_PP_EVAL(_M, ...) _M(__VA_ARGS__)
 #define _CCCL_PP_EAT(...)

libcudacxx/include/cuda/std/__mdspan/config.h

@@ -56,10 +56,6 @@
 
 #if _CCCL_STD_VER > 2011
 
-#  ifndef __has_include
-#    define __has_include(x) 0
-#  endif
-
 #  ifdef _MSVC_LANG
 #    define __MDSPAN_CPLUSPLUS _MSVC_LANG
 #  else

libcudacxx/include/cuda/std/detail/libcxx/include/__config

@@ -98,10 +98,6 @@ extern "C++" {
 #    define __has_declspec_attribute(__x) 0
 #  endif
 
-#  ifndef __has_include
-#    define __has_include(...) 0
-#  endif
-
 #  if !defined(_CCCL_CUDA_COMPILER_NVCC) && !defined(_CCCL_COMPILER_NVRTC)
 // If NVCC is not being used <complex> can safely use `long double` without warnings
 #    define _LIBCUDACXX_HAS_COMPLEX_LONG_DOUBLE
@@ -700,7 +696,7 @@ typedef unsigned int char32_t;
     && !defined(_LIBCUDACXX_HAS_THREAD_API_WIN32) && !defined(_LIBCUDACXX_HAS_THREAD_API_EXTERNAL)
 #    if defined(__FreeBSD__) || defined(__Fuchsia__) || defined(__wasi__) || defined(__NetBSD__) || defined(__linux__) \
       || defined(__GNU__) || defined(__APPLE__) || defined(__CloudABI__) || defined(__sun__)                           \
-      || (defined(__MINGW32__) && __has_include(<pthread.h>))
+      || (defined(__MINGW32__) && _CCCL_HAS_INCLUDE(<pthread.h>))
 #      define _LIBCUDACXX_HAS_THREAD_API_PTHREAD
 #    elif defined(_LIBCUDACXX_WIN32API)
 #      define _LIBCUDACXX_HAS_THREAD_API_WIN32

libcudacxx/include/cuda/std/version

@@ -23,7 +23,7 @@
 
 // We need to define our own macros to not conflict with the host stl.
 // At the same time we want bring in all feature test macros from host
-#if __has_include(<version>) // <version> should be the smallest include possible
+#if _CCCL_HAS_INCLUDE(<version>) // <version> should be the smallest include possible
 #  include <version>
 #elif !defined(_CCCL_COMPILER_NVRTC)
 #  include <ciso646> // otherwise go for the smallest possible header

thrust/thrust/detail/config/memory_resource.h

@@ -34,10 +34,10 @@
 #define THRUST_MR_DEFAULT_ALIGNMENT alignof(THRUST_NS_QUALIFIER::detail::max_align_t)
 
 #if _CCCL_STD_VER >= 2017
-#  if __has_include(<memory_resource>)
+#  if _CCCL_HAS_INCLUDE(<memory_resource>)
 #    define THRUST_MR_STD_MR_HEADER <memory_resource>
 #    define THRUST_MR_STD_MR_NS     std::pmr
-#  elif __has_include(<experimental/memory_resource>)
+#  elif _CCCL_HAS_INCLUDE(<experimental/memory_resource>)
 #    define THRUST_MR_STD_MR_HEADER <experimental/memory_resource>
 #    define THRUST_MR_STD_MR_NS     std::experimental::pmr
 #  endif

thrust/thrust/type_traits/remove_cvref.h

@@ -31,11 +31,9 @@
 #  pragma system_header
 #endif // no system header
 
-#if _CCCL_STD_VER >= 2017
-#  if __has_include(<version>)
-#    include <version>
-#  endif
-#endif
+#if _CCCL_HAS_INCLUDE(<version>)
+#  include <version>
+#endif // _CCCL_HAS_INCLUDE(<version>)
 
 #include <type_traits>