Deprecate Thrust's cpp_compatibility.h macros

docs/repo.toml

@@ -271,6 +271,7 @@ doxygen_predefined = [
   "_CCCL_DIAG_SUPPRESS_GCC(x)=",
   "_CCCL_DIAG_SUPPRESS_MSVC(x)=",
   "_CCCL_DIAG_SUPPRESS_NVHPC(x)=",
+  "_CCCL_GLOBAL_CONSTANT=constexpr",
   "CUDASTF_HOST=",
   "CUDASTF_DEVICE=",
   "CUDASTF_HOST_DEVICE=",

thrust/thrust/async/copy.h

@@ -117,7 +117,7 @@ struct copy_fn final
 
 } // namespace copy_detail
 
-THRUST_INLINE_CONSTANT copy_detail::copy_fn copy{};
+_CCCL_GLOBAL_CONSTANT copy_detail::copy_fn copy{};
 
 /*! \endcond
  */

thrust/thrust/async/for_each.h

@@ -94,7 +94,7 @@ struct for_each_fn final
 
 } // namespace for_each_detail
 
-THRUST_INLINE_CONSTANT for_each_detail::for_each_fn for_each{};
+_CCCL_GLOBAL_CONSTANT for_each_detail::for_each_fn for_each{};
 
 /*! \endcond
  */

thrust/thrust/async/reduce.h

@@ -160,7 +160,7 @@ struct reduce_fn final
 
 } // namespace reduce_detail
 
-THRUST_INLINE_CONSTANT reduce_detail::reduce_fn reduce{};
+_CCCL_GLOBAL_CONSTANT reduce_detail::reduce_fn reduce{};
 
 ///////////////////////////////////////////////////////////////////////////////
 
@@ -296,7 +296,7 @@ struct reduce_into_fn final
 
 } // namespace reduce_into_detail
 
-THRUST_INLINE_CONSTANT reduce_into_detail::reduce_into_fn reduce_into{};
+_CCCL_GLOBAL_CONSTANT reduce_into_detail::reduce_into_fn reduce_into{};
 
 /*! \endcond
  */

thrust/thrust/async/scan.h

@@ -179,7 +179,7 @@ struct inclusive_scan_fn final
 
 } // namespace inclusive_scan_detail
 
-THRUST_INLINE_CONSTANT inclusive_scan_detail::inclusive_scan_fn inclusive_scan{};
+_CCCL_GLOBAL_CONSTANT inclusive_scan_detail::inclusive_scan_fn inclusive_scan{};
 
 namespace exclusive_scan_detail
 {
@@ -288,7 +288,7 @@ struct exclusive_scan_fn final
 
 } // namespace exclusive_scan_detail
 
-THRUST_INLINE_CONSTANT exclusive_scan_detail::exclusive_scan_fn exclusive_scan{};
+_CCCL_GLOBAL_CONSTANT exclusive_scan_detail::exclusive_scan_fn exclusive_scan{};
 
 } // namespace async
 

thrust/thrust/async/sort.h

@@ -146,7 +146,7 @@ struct stable_sort_fn final
 
 } // namespace stable_sort_detail
 
-THRUST_INLINE_CONSTANT stable_sort_detail::stable_sort_fn stable_sort{};
+_CCCL_GLOBAL_CONSTANT stable_sort_detail::stable_sort_fn stable_sort{};
 
 namespace fallback
 {
@@ -259,7 +259,7 @@ struct sort_fn final
 
 } // namespace sort_detail
 
-THRUST_INLINE_CONSTANT sort_detail::sort_fn sort{};
+_CCCL_GLOBAL_CONSTANT sort_detail::sort_fn sort{};
 
 /*! \endcond
  */

thrust/thrust/async/transform.h

@@ -125,7 +125,7 @@ struct transform_fn final
 
 } // namespace transform_detail
 
-THRUST_INLINE_CONSTANT transform_detail::transform_fn transform{};
+_CCCL_GLOBAL_CONSTANT transform_detail::transform_fn transform{};
 
 /*! \endcond
  */

thrust/thrust/detail/config/cpp_compatibility.h

@@ -16,7 +16,7 @@
 
 #pragma once
 
-// Internal config header that is only included through thrust/detail/config/config.h
+#include <cuda/__cccl_config>
 
 #if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
 #  pragma GCC system_header
@@ -26,29 +26,9 @@
 #  pragma system_header
 #endif // no system header
 
-#include <thrust/detail/config/cpp_dialect.h> // IWYU pragma: export
-
-#include <cuda/std/cstddef>
-
+// deprecated [Since 2.8.0]
 #define THRUST_NODISCARD _CCCL_NODISCARD
-
-// FIXME: Combine THRUST_INLINE_CONSTANT and
-// THRUST_INLINE_INTEGRAL_MEMBER_CONSTANT into one macro when NVCC properly
-// supports `constexpr` globals in host and device code.
-#if defined(__CUDA_ARCH__) || defined(_NVHPC_CUDA)
-// FIXME: Add this when NVCC supports inline variables.
-// #  if   _CCCL_STD_VER >= 2017
-// #    define THRUST_INLINE_CONSTANT                 inline constexpr
-// #    define THRUST_INLINE_INTEGRAL_MEMBER_CONSTANT inline constexpr
-#  define THRUST_INLINE_CONSTANT                 static const _CCCL_DEVICE
-#  define THRUST_INLINE_INTEGRAL_MEMBER_CONSTANT static constexpr
-
-#else
-// FIXME: Add this when NVCC supports inline variables.
-// #  if   _CCCL_STD_VER >= 2017
-// #    define THRUST_INLINE_CONSTANT                 inline constexpr
-// #    define THRUST_INLINE_INTEGRAL_MEMBER_CONSTANT inline constexpr
-#  define THRUST_INLINE_CONSTANT                 static constexpr
-#  define THRUST_INLINE_INTEGRAL_MEMBER_CONSTANT static constexpr
-
-#endif
+// deprecated [Since 2.8.0]
+#define THRUST_INLINE_CONSTANT _CCCL_GLOBAL_CONSTANT
+// deprecated [Since 2.8.0]
+#define THRUST_INLINE_INTEGRAL_MEMBER_CONSTANT static constexpr

thrust/thrust/detail/select_system.h

@@ -64,7 +64,7 @@ struct select_system_fn final
 
 } // namespace select_system_detail
 
-THRUST_INLINE_CONSTANT select_system_detail::select_system_fn select_system{};
+_CCCL_GLOBAL_CONSTANT select_system_detail::select_system_fn select_system{};
 
 } // namespace detail
 

thrust/thrust/detail/seq.h

@@ -49,6 +49,6 @@ struct seq_t
 
 } // namespace detail
 
-THRUST_INLINE_CONSTANT detail::seq_t seq;
+_CCCL_GLOBAL_CONSTANT detail::seq_t seq;
 
 THRUST_NAMESPACE_END

thrust/thrust/detail/static_assert.h

@@ -45,7 +45,7 @@ namespace detail
 template <typename, bool x>
 struct depend_on_instantiation
 {
-  THRUST_INLINE_INTEGRAL_MEMBER_CONSTANT bool value = x;
+  static constexpr bool value = x;
 };
 
 #if _CCCL_STD_VER >= 2017

thrust/thrust/execution_policy.h

@@ -336,7 +336,7 @@ static const detail::host_t host;
  *  \see host_execution_policy
  *  \see thrust::device
  */
-THRUST_INLINE_CONSTANT detail::device_t device;
+_CCCL_GLOBAL_CONSTANT detail::device_t device;
 
 // define seq for the purpose of Doxygenating it
 // it is actually defined elsewhere

thrust/thrust/functional.h

@@ -1046,43 +1046,43 @@ namespace placeholders
 
 /*! \p thrust::placeholders::_1 is the placeholder for the first function parameter.
  */
-THRUST_INLINE_CONSTANT thrust::detail::functional::placeholder<0>::type _1;
+_CCCL_GLOBAL_CONSTANT thrust::detail::functional::placeholder<0>::type _1;
 
 /*! \p thrust::placeholders::_2 is the placeholder for the second function parameter.
  */
-THRUST_INLINE_CONSTANT thrust::detail::functional::placeholder<1>::type _2;
+_CCCL_GLOBAL_CONSTANT thrust::detail::functional::placeholder<1>::type _2;
 
 /*! \p thrust::placeholders::_3 is the placeholder for the third function parameter.
  */
-THRUST_INLINE_CONSTANT thrust::detail::functional::placeholder<2>::type _3;
+_CCCL_GLOBAL_CONSTANT thrust::detail::functional::placeholder<2>::type _3;
 
 /*! \p thrust::placeholders::_4 is the placeholder for the fourth function parameter.
  */
-THRUST_INLINE_CONSTANT thrust::detail::functional::placeholder<3>::type _4;
+_CCCL_GLOBAL_CONSTANT thrust::detail::functional::placeholder<3>::type _4;
 
 /*! \p thrust::placeholders::_5 is the placeholder for the fifth function parameter.
  */
-THRUST_INLINE_CONSTANT thrust::detail::functional::placeholder<4>::type _5;
+_CCCL_GLOBAL_CONSTANT thrust::detail::functional::placeholder<4>::type _5;
 
 /*! \p thrust::placeholders::_6 is the placeholder for the sixth function parameter.
  */
-THRUST_INLINE_CONSTANT thrust::detail::functional::placeholder<5>::type _6;
+_CCCL_GLOBAL_CONSTANT thrust::detail::functional::placeholder<5>::type _6;
 
 /*! \p thrust::placeholders::_7 is the placeholder for the seventh function parameter.
  */
-THRUST_INLINE_CONSTANT thrust::detail::functional::placeholder<6>::type _7;
+_CCCL_GLOBAL_CONSTANT thrust::detail::functional::placeholder<6>::type _7;
 
 /*! \p thrust::placeholders::_8 is the placeholder for the eighth function parameter.
  */
-THRUST_INLINE_CONSTANT thrust::detail::functional::placeholder<7>::type _8;
+_CCCL_GLOBAL_CONSTANT thrust::detail::functional::placeholder<7>::type _8;
 
 /*! \p thrust::placeholders::_9 is the placeholder for the ninth function parameter.
  */
-THRUST_INLINE_CONSTANT thrust::detail::functional::placeholder<8>::type _9;
+_CCCL_GLOBAL_CONSTANT thrust::detail::functional::placeholder<8>::type _9;
 
 /*! \p thrust::placeholders::_10 is the placeholder for the tenth function parameter.
  */
-THRUST_INLINE_CONSTANT thrust::detail::functional::placeholder<9>::type _10;
+_CCCL_GLOBAL_CONSTANT thrust::detail::functional::placeholder<9>::type _10;
 
 } // namespace placeholders
 

thrust/thrust/future.h

@@ -164,7 +164,7 @@ using device_future = device_unique_eager_future<T>;
 struct new_stream_t final
 {};
 
-THRUST_INLINE_CONSTANT new_stream_t new_stream{};
+_CCCL_GLOBAL_CONSTANT new_stream_t new_stream{};
 
 ///////////////////////////////////////////////////////////////////////////////
 

thrust/thrust/system/cpp/detail/par.h

@@ -47,7 +47,7 @@ struct par_t
 
 } // namespace detail
 
-THRUST_INLINE_CONSTANT detail::par_t par;
+_CCCL_GLOBAL_CONSTANT detail::par_t par;
 
 } // namespace cpp
 } // namespace system

thrust/thrust/system/cuda/detail/future.inl

@@ -58,7 +58,7 @@ namespace detail
 struct nonowning_t final
 {};
 
-THRUST_INLINE_CONSTANT nonowning_t nonowning{};
+_CCCL_GLOBAL_CONSTANT nonowning_t nonowning{};
 
 ///////////////////////////////////////////////////////////////////////////////
 

thrust/thrust/system/cuda/detail/par.h

@@ -166,7 +166,7 @@ struct par_nosync_t
   }
 };
 
-THRUST_INLINE_CONSTANT par_t par;
+_CCCL_GLOBAL_CONSTANT par_t par;
 
 /*! \p thrust::cuda::par_nosync is a parallel execution policy targeting Thrust's CUDA device backend.
  *  Similar to \p thrust::cuda::par it allows execution of Thrust algorithms in a specific CUDA stream.
@@ -215,7 +215,7 @@ THRUST_INLINE_CONSTANT par_t par;
  *  \endcode
  *
  */
-THRUST_INLINE_CONSTANT par_nosync_t par_nosync;
+_CCCL_GLOBAL_CONSTANT par_nosync_t par_nosync;
 } // namespace cuda_cub
 
 namespace system

thrust/thrust/system/detail/sequential/execution_policy.h

@@ -70,7 +70,7 @@ struct execution_policy : thrust::execution_policy<Derived>
   }
 };
 
-THRUST_INLINE_CONSTANT tag seq;
+_CCCL_GLOBAL_CONSTANT tag seq;
 
 } // namespace sequential
 } // namespace detail