diff --git a/cub/cub/thread/thread_sort.cuh b/cub/cub/thread/thread_sort.cuh
index 7d9e8622f82..e21d9dd3ecd 100644
--- a/cub/cub/thread/thread_sort.cuh
+++ b/cub/cub/thread/thread_sort.cuh
@@ -45,6 +45,7 @@
 CUB_NAMESPACE_BEGIN
 
 template <typename T>
+CCCL_DEPRECATED_BECAUSE("Use cuda::std::swap")
 _CCCL_DEVICE _CCCL_FORCEINLINE void Swap(T& lhs, T& rhs)
 {
   T temp = lhs;
@@ -95,10 +96,11 @@ StableOddEvenSort(KeyT (&keys)[ITEMS_PER_THREAD], ValueT (&items)[ITEMS_PER_THRE
     {
       if (compare_op(keys[j + 1], keys[j]))
       {
-        Swap(keys[j], keys[j + 1]);
+        using ::cuda::std::swap;
+        swap(keys[j], keys[j + 1]);
         if (!KEYS_ONLY)
         {
-          Swap(items[j], items[j + 1]);
+          swap(items[j], items[j + 1]);
         }
       }
     } // inner loop
diff --git a/docs/repo.toml b/docs/repo.toml
index 87655b3c65d..f5be6925fb5 100644
--- a/docs/repo.toml
+++ b/docs/repo.toml
@@ -273,6 +273,7 @@ doxygen_predefined = [
   "_CCCL_DIAG_SUPPRESS_ICC(x)=",
   "_CCCL_DIAG_SUPPRESS_MSVC(x)=",
   "_CCCL_DIAG_SUPPRESS_NVHPC(x)=",
+  "_CCCL_GLOBAL_CONSTANT=constexpr",
   "CUDASTF_HOST=",
   "CUDASTF_DEVICE=",
   "CUDASTF_HOST_DEVICE=",
diff --git a/thrust/thrust/async/copy.h b/thrust/thrust/async/copy.h
index d52526444a0..f981e1ac17d 100644
--- a/thrust/thrust/async/copy.h
+++ b/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
  */
diff --git a/thrust/thrust/async/for_each.h b/thrust/thrust/async/for_each.h
index 17376343fed..656d43d4f15 100644
--- a/thrust/thrust/async/for_each.h
+++ b/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
  */
diff --git a/thrust/thrust/async/reduce.h b/thrust/thrust/async/reduce.h
index d4cd7cd7c3f..0a56a2ba51c 100644
--- a/thrust/thrust/async/reduce.h
+++ b/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
  */
diff --git a/thrust/thrust/async/scan.h b/thrust/thrust/async/scan.h
index fcbb41ad43c..5f287b64660 100644
--- a/thrust/thrust/async/scan.h
+++ b/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
 
diff --git a/thrust/thrust/async/sort.h b/thrust/thrust/async/sort.h
index bd294048311..7b76c0f9c00 100644
--- a/thrust/thrust/async/sort.h
+++ b/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
  */
diff --git a/thrust/thrust/async/transform.h b/thrust/thrust/async/transform.h
index 901518c12e8..0e97e0bce51 100644
--- a/thrust/thrust/async/transform.h
+++ b/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
  */
diff --git a/thrust/thrust/detail/config/cpp_compatibility.h b/thrust/thrust/detail/config/cpp_compatibility.h
index 6094319dae0..36f74614eec 100644
--- a/thrust/thrust/detail/config/cpp_compatibility.h
+++ b/thrust/thrust/detail/config/cpp_compatibility.h
@@ -26,32 +26,12 @@
 #  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
 
 // These definitions were intended for internal use only and are now obsolete.
 // If you relied on them, consider porting your code to use the functionality
diff --git a/thrust/thrust/detail/select_system.h b/thrust/thrust/detail/select_system.h
index 65450582967..65f4593cc7a 100644
--- a/thrust/thrust/detail/select_system.h
+++ b/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
 
diff --git a/thrust/thrust/detail/seq.h b/thrust/thrust/detail/seq.h
index 992a9b09100..7adf956b376 100644
--- a/thrust/thrust/detail/seq.h
+++ b/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
diff --git a/thrust/thrust/detail/static_assert.h b/thrust/thrust/detail/static_assert.h
index a49d5d96eb1..d21dcbdfbf2 100644
--- a/thrust/thrust/detail/static_assert.h
+++ b/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
diff --git a/thrust/thrust/execution_policy.h b/thrust/thrust/execution_policy.h
index 56199aa5e91..8f733215e9a 100644
--- a/thrust/thrust/execution_policy.h
+++ b/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
diff --git a/thrust/thrust/functional.h b/thrust/thrust/functional.h
index a654a8b79b7..816c8744b89 100644
--- a/thrust/thrust/functional.h
+++ b/thrust/thrust/functional.h
@@ -1307,43 +1307,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
 
diff --git a/thrust/thrust/future.h b/thrust/thrust/future.h
index 39f25a0b0c0..eb408da58b8 100644
--- a/thrust/thrust/future.h
+++ b/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{};
 
 ///////////////////////////////////////////////////////////////////////////////
 
diff --git a/thrust/thrust/system/cpp/detail/par.h b/thrust/thrust/system/cpp/detail/par.h
index 5e32d6d0149..162f51215aa 100644
--- a/thrust/thrust/system/cpp/detail/par.h
+++ b/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
diff --git a/thrust/thrust/system/cuda/detail/future.inl b/thrust/thrust/system/cuda/detail/future.inl
index 343a18156ec..081fa0f2a02 100644
--- a/thrust/thrust/system/cuda/detail/future.inl
+++ b/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{};
 
 ///////////////////////////////////////////////////////////////////////////////
 
diff --git a/thrust/thrust/system/cuda/detail/par.h b/thrust/thrust/system/cuda/detail/par.h
index 5541ca60074..f52bae2ed81 100644
--- a/thrust/thrust/system/cuda/detail/par.h
+++ b/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
diff --git a/thrust/thrust/system/detail/sequential/execution_policy.h b/thrust/thrust/system/detail/sequential/execution_policy.h
index d9d52aad85f..d4834ebc89b 100644
--- a/thrust/thrust/system/detail/sequential/execution_policy.h
+++ b/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