diff --git a/cudax/include/cuda/experimental/__hierarchy/hierarchy_levels.cuh b/cudax/include/cuda/experimental/__hierarchy/hierarchy_levels.cuh
index 5764f95b52..9965a3bd11 100644
--- a/cudax/include/cuda/experimental/__hierarchy/hierarchy_levels.cuh
+++ b/cudax/include/cuda/experimental/__hierarchy/hierarchy_levels.cuh
@@ -140,7 +140,7 @@ struct grid_level
   using allowed_above = allowed_levels<>;
   using allowed_below = allowed_levels<block_level, cluster_level>;
 };
-_LIBCUDACXX_CPO_ACCESSIBILITY grid_level grid;
+_CCCL_GLOBAL_CONSTANT grid_level grid;
 
 /**
  * @brief Type representing the cluster level in CUDA thread hierarchy
@@ -158,7 +158,7 @@ struct cluster_level
   using allowed_above = allowed_levels<grid_level>;
   using allowed_below = allowed_levels<block_level>;
 };
-_LIBCUDACXX_CPO_ACCESSIBILITY cluster_level cluster;
+_CCCL_GLOBAL_CONSTANT cluster_level cluster;
 
 /**
  * @brief Type representing the block level in CUDA thread hierarchy
@@ -176,7 +176,7 @@ struct block_level
   using allowed_above = allowed_levels<grid_level, cluster_level>;
   using allowed_below = allowed_levels<thread_level>;
 };
-_LIBCUDACXX_CPO_ACCESSIBILITY block_level block;
+_CCCL_GLOBAL_CONSTANT block_level block;
 
 /**
  * @brief Type representing the thread level in CUDA thread hierarchy
@@ -194,7 +194,7 @@ struct thread_level
   using allowed_above = allowed_levels<block_level>;
   using allowed_below = allowed_levels<>;
 };
-_LIBCUDACXX_CPO_ACCESSIBILITY thread_level thread;
+_CCCL_GLOBAL_CONSTANT thread_level thread;
 
 template <typename Level>
 constexpr bool is_core_cuda_hierarchy_level =
diff --git a/libcudacxx/include/cuda/__memory_resource/get_property.h b/libcudacxx/include/cuda/__memory_resource/get_property.h
index c845cc68e8..4208d07232 100644
--- a/libcudacxx/include/cuda/__memory_resource/get_property.h
+++ b/libcudacxx/include/cuda/__memory_resource/get_property.h
@@ -176,7 +176,7 @@ _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto get_property = __get_property::__fn{};
+_CCCL_GLOBAL_CONSTANT auto get_property = __get_property::__fn{};
 } // namespace __cpo
 
 _LIBCUDACXX_END_NAMESPACE_CUDA
diff --git a/libcudacxx/include/cuda/std/__cccl/dialect.h b/libcudacxx/include/cuda/std/__cccl/dialect.h
index 7b15b6ef81..0c8989fbc6 100644
--- a/libcudacxx/include/cuda/std/__cccl/dialect.h
+++ b/libcudacxx/include/cuda/std/__cccl/dialect.h
@@ -90,4 +90,25 @@
 #  define _CCCL_TRAIT(__TRAIT, ...) __TRAIT<__VA_ARGS__>::value
 #endif // _CCCL_STD_VER <= 2014
 
+// In nvcc prior to 11.3 global variables could not be marked constexpr
+#if defined(_CCCL_CUDACC_BELOW_11_3)
+#  define _CCCL_CONSTEXPR_GLOBAL const
+#else // ^^^ _CCCL_CUDACC_BELOW_11_3 ^^^ / vvv !_CCCL_CUDACC_BELOW_11_3 vvv
+#  define _CCCL_CONSTEXPR_GLOBAL constexpr
+#endif // !_CCCL_CUDACC_BELOW_11_3
+
+// Inline variables are only available from C++17 onwards
+#if _CCCL_STD_VER >= 2017 && defined(__cpp_inline_variables) && (__cpp_inline_variables >= 201606L)
+#  define _CCCL_INLINE_VAR inline
+#else // ^^^ C++14 ^^^ / vvv C++17 vvv
+#  define _CCCL_INLINE_VAR
+#endif // _CCCL_STD_VER <= 2014
+
+// We need to treat host and device separately
+#if defined(__CUDA_ARCH__)
+#  define _CCCL_GLOBAL_CONSTANT _CCCL_DEVICE _CCCL_CONSTEXPR_GLOBAL
+#else // ^^^ __CUDA_ARCH__ ^^^ / vvv !__CUDA_ARCH__ vvv
+#  define _CCCL_GLOBAL_CONSTANT _CCCL_INLINE_VAR constexpr
+#endif // __CUDA_ARCH__
+
 #endif // __CCCL_DIALECT_H
diff --git a/libcudacxx/include/cuda/std/__concepts/swappable.h b/libcudacxx/include/cuda/std/__concepts/swappable.h
index adf54861dc..f148b54dc6 100644
--- a/libcudacxx/include/cuda/std/__concepts/swappable.h
+++ b/libcudacxx/include/cuda/std/__concepts/swappable.h
@@ -158,7 +158,7 @@ _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto swap = __swap::__fn{};
+_CCCL_GLOBAL_CONSTANT auto swap = __swap::__fn{};
 } // namespace __cpo
 _LIBCUDACXX_END_NAMESPACE_RANGES
 
diff --git a/libcudacxx/include/cuda/std/__expected/unexpect.h b/libcudacxx/include/cuda/std/__expected/unexpect.h
index fd8f103447..fcca0c777b 100644
--- a/libcudacxx/include/cuda/std/__expected/unexpect.h
+++ b/libcudacxx/include/cuda/std/__expected/unexpect.h
@@ -28,7 +28,7 @@ struct unexpect_t
   explicit unexpect_t() = default;
 };
 
-_LIBCUDACXX_CPO_ACCESSIBILITY unexpect_t unexpect{};
+_CCCL_GLOBAL_CONSTANT unexpect_t unexpect{};
 
 _LIBCUDACXX_END_NAMESPACE_STD
 
diff --git a/libcudacxx/include/cuda/std/__iterator/access.h b/libcudacxx/include/cuda/std/__iterator/access.h
index f720e7730d..6ff5bf0710 100644
--- a/libcudacxx/include/cuda/std/__iterator/access.h
+++ b/libcudacxx/include/cuda/std/__iterator/access.h
@@ -53,7 +53,7 @@ struct __fn
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto begin = __begin::__fn{};
+_CCCL_GLOBAL_CONSTANT auto begin = __begin::__fn{};
 } // namespace __cpo
 
 namespace __end
@@ -84,7 +84,7 @@ struct __fn
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto end = __end::__fn{};
+_CCCL_GLOBAL_CONSTANT auto end = __end::__fn{};
 } // namespace __cpo
 
 #if _CCCL_STD_VER >= 2014
@@ -104,7 +104,7 @@ struct __fn
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto cbegin = __cbegin::__fn{};
+_CCCL_GLOBAL_CONSTANT auto cbegin = __cbegin::__fn{};
 } // namespace __cpo
 
 namespace __cend
@@ -122,7 +122,7 @@ struct __fn
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto cend = __cend::__fn{};
+_CCCL_GLOBAL_CONSTANT auto cend = __cend::__fn{};
 } // namespace __cpo
 
 #endif // _CCCL_STD_VER >= 2014
diff --git a/libcudacxx/include/cuda/std/__iterator/advance.h b/libcudacxx/include/cuda/std/__iterator/advance.h
index 7e5368156b..ed6d711c0e 100644
--- a/libcudacxx/include/cuda/std/__iterator/advance.h
+++ b/libcudacxx/include/cuda/std/__iterator/advance.h
@@ -238,7 +238,7 @@ _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto advance = __advance::__fn{};
+_CCCL_GLOBAL_CONSTANT auto advance = __advance::__fn{};
 } // namespace __cpo
 
 _LIBCUDACXX_END_NAMESPACE_RANGES
diff --git a/libcudacxx/include/cuda/std/__iterator/distance.h b/libcudacxx/include/cuda/std/__iterator/distance.h
index 9b31aaa44a..2ce304f95a 100644
--- a/libcudacxx/include/cuda/std/__iterator/distance.h
+++ b/libcudacxx/include/cuda/std/__iterator/distance.h
@@ -116,7 +116,7 @@ _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto distance = __distance::__fn{};
+_CCCL_GLOBAL_CONSTANT auto distance = __distance::__fn{};
 } // namespace __cpo
 _LIBCUDACXX_END_NAMESPACE_RANGES
 
diff --git a/libcudacxx/include/cuda/std/__iterator/iter_move.h b/libcudacxx/include/cuda/std/__iterator/iter_move.h
index d447e29fa0..070da03f11 100644
--- a/libcudacxx/include/cuda/std/__iterator/iter_move.h
+++ b/libcudacxx/include/cuda/std/__iterator/iter_move.h
@@ -119,7 +119,7 @@ struct __fn
 _LIBCUDACXX_END_NAMESPACE_CPO
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto iter_move = __iter_move::__fn{};
+_CCCL_GLOBAL_CONSTANT auto iter_move = __iter_move::__fn{};
 } // namespace __cpo
 _LIBCUDACXX_END_NAMESPACE_RANGES
 
diff --git a/libcudacxx/include/cuda/std/__iterator/iter_swap.h b/libcudacxx/include/cuda/std/__iterator/iter_swap.h
index 76d02c6989..f989719586 100644
--- a/libcudacxx/include/cuda/std/__iterator/iter_swap.h
+++ b/libcudacxx/include/cuda/std/__iterator/iter_swap.h
@@ -119,7 +119,7 @@ _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto iter_swap = __iter_swap::__fn{};
+_CCCL_GLOBAL_CONSTANT auto iter_swap = __iter_swap::__fn{};
 } // namespace __cpo
 _LIBCUDACXX_END_NAMESPACE_RANGES
 
diff --git a/libcudacxx/include/cuda/std/__iterator/next.h b/libcudacxx/include/cuda/std/__iterator/next.h
index 5adfee7270..c4e70d2c4a 100644
--- a/libcudacxx/include/cuda/std/__iterator/next.h
+++ b/libcudacxx/include/cuda/std/__iterator/next.h
@@ -90,7 +90,7 @@ _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto next = __next::__fn{};
+_CCCL_GLOBAL_CONSTANT auto next = __next::__fn{};
 } // namespace __cpo
 _LIBCUDACXX_END_NAMESPACE_RANGES
 
diff --git a/libcudacxx/include/cuda/std/__iterator/prev.h b/libcudacxx/include/cuda/std/__iterator/prev.h
index 9d8255ec62..5a84b44c59 100644
--- a/libcudacxx/include/cuda/std/__iterator/prev.h
+++ b/libcudacxx/include/cuda/std/__iterator/prev.h
@@ -81,7 +81,7 @@ _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto prev = __prev::__fn{};
+_CCCL_GLOBAL_CONSTANT auto prev = __prev::__fn{};
 } // namespace __cpo
 _LIBCUDACXX_END_NAMESPACE_RANGES
 
diff --git a/libcudacxx/include/cuda/std/__iterator/reverse_access.h b/libcudacxx/include/cuda/std/__iterator/reverse_access.h
index 5ef118b11c..927b3a3d0e 100644
--- a/libcudacxx/include/cuda/std/__iterator/reverse_access.h
+++ b/libcudacxx/include/cuda/std/__iterator/reverse_access.h
@@ -65,7 +65,7 @@ struct __fn
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto rbegin = __rbegin::__fn{};
+_CCCL_GLOBAL_CONSTANT auto rbegin = __rbegin::__fn{};
 } // namespace __cpo
 
 namespace __rend
@@ -104,7 +104,7 @@ struct __fn
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto rend = __rend::__fn{};
+_CCCL_GLOBAL_CONSTANT auto rend = __rend::__fn{};
 } // namespace __cpo
 
 namespace __crbegin
@@ -122,7 +122,7 @@ struct __fn
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto crbegin = __crbegin::__fn{};
+_CCCL_GLOBAL_CONSTANT auto crbegin = __crbegin::__fn{};
 } // namespace __cpo
 
 namespace __crend
@@ -140,7 +140,7 @@ struct __fn
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto crend = __crend::__fn{};
+_CCCL_GLOBAL_CONSTANT auto crend = __crend::__fn{};
 } // namespace __cpo
 
 #endif // _CCCL_STD_VER >= 2014
diff --git a/libcudacxx/include/cuda/std/__iterator/unreachable_sentinel.h b/libcudacxx/include/cuda/std/__iterator/unreachable_sentinel.h
index 0543e33ad1..2304d5a15b 100644
--- a/libcudacxx/include/cuda/std/__iterator/unreachable_sentinel.h
+++ b/libcudacxx/include/cuda/std/__iterator/unreachable_sentinel.h
@@ -82,7 +82,7 @@ struct unreachable_sentinel_t : __unreachable_sentinel_detail::__unreachable_bas
 
 _LIBCUDACXX_END_NAMESPACE_RANGES_ABI
 
-_LIBCUDACXX_CPO_ACCESSIBILITY unreachable_sentinel_t unreachable_sentinel{};
+_CCCL_GLOBAL_CONSTANT unreachable_sentinel_t unreachable_sentinel{};
 _LIBCUDACXX_END_NAMESPACE_STD
 
 #endif // _CCCL_STD_VER > 2014
diff --git a/libcudacxx/include/cuda/std/__mdspan/standard_layout_static_array.h b/libcudacxx/include/cuda/std/__mdspan/standard_layout_static_array.h
index fc4c355d7d..64440bfacc 100644
--- a/libcudacxx/include/cuda/std/__mdspan/standard_layout_static_array.h
+++ b/libcudacxx/include/cuda/std/__mdspan/standard_layout_static_array.h
@@ -77,12 +77,11 @@ namespace __detail
 
 struct __construct_psa_from_dynamic_exts_values_tag_t
 {};
-_LIBCUDACXX_CPO_ACCESSIBILITY __construct_psa_from_dynamic_exts_values_tag_t
-  __construct_psa_from_dynamic_exts_values_tag;
+_CCCL_GLOBAL_CONSTANT __construct_psa_from_dynamic_exts_values_tag_t __construct_psa_from_dynamic_exts_values_tag;
 
 struct __construct_psa_from_all_exts_values_tag_t
 {};
-_LIBCUDACXX_CPO_ACCESSIBILITY __construct_psa_from_all_exts_values_tag_t __construct_psa_from_all_exts_values_tag;
+_CCCL_GLOBAL_CONSTANT __construct_psa_from_all_exts_values_tag_t __construct_psa_from_all_exts_values_tag;
 
 struct __construct_psa_from_all_exts_array_tag_t
 {};
diff --git a/libcudacxx/include/cuda/std/__ranges/access.h b/libcudacxx/include/cuda/std/__ranges/access.h
index c2b065075a..a2bc2b4133 100644
--- a/libcudacxx/include/cuda/std/__ranges/access.h
+++ b/libcudacxx/include/cuda/std/__ranges/access.h
@@ -125,7 +125,7 @@ _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto begin = __begin::__fn{};
+_CCCL_GLOBAL_CONSTANT auto begin = __begin::__fn{};
 } // namespace __cpo
 
 // [range.range]
@@ -213,7 +213,7 @@ _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto end = __end::__fn{};
+_CCCL_GLOBAL_CONSTANT auto end = __end::__fn{};
 } // namespace __cpo
 
 // [range.access.cbegin]
@@ -243,7 +243,7 @@ _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto cbegin = __cbegin::__fn{};
+_CCCL_GLOBAL_CONSTANT auto cbegin = __cbegin::__fn{};
 } // namespace __cpo
 
 // [range.access.cend]
@@ -273,7 +273,7 @@ _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto cend = __cend::__fn{};
+_CCCL_GLOBAL_CONSTANT auto cend = __cend::__fn{};
 } // namespace __cpo
 #endif // _CCCL_STD_VER > 2014 && !_CCCL_COMPILER_MSVC_2017
 
diff --git a/libcudacxx/include/cuda/std/__ranges/data.h b/libcudacxx/include/cuda/std/__ranges/data.h
index 63a89d6c07..14d3473f3c 100644
--- a/libcudacxx/include/cuda/std/__ranges/data.h
+++ b/libcudacxx/include/cuda/std/__ranges/data.h
@@ -97,7 +97,7 @@ _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto data = __data::__fn{};
+_CCCL_GLOBAL_CONSTANT auto data = __data::__fn{};
 } // namespace __cpo
 
 // [range.prim.cdata]
@@ -127,7 +127,7 @@ _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto cdata = __cdata::__fn{};
+_CCCL_GLOBAL_CONSTANT auto cdata = __cdata::__fn{};
 } // namespace __cpo
 
 #endif // _CCCL_STD_VER >= 2017 && !_CCCL_COMPILER_MSVC_2017
diff --git a/libcudacxx/include/cuda/std/__ranges/empty.h b/libcudacxx/include/cuda/std/__ranges/empty.h
index 84a9cc2e27..581faa8bbd 100644
--- a/libcudacxx/include/cuda/std/__ranges/empty.h
+++ b/libcudacxx/include/cuda/std/__ranges/empty.h
@@ -102,7 +102,7 @@ _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto empty = __empty::__fn{};
+_CCCL_GLOBAL_CONSTANT auto empty = __empty::__fn{};
 } // namespace __cpo
 
 #endif // _CCCL_STD_VER >= 2017 && !_CCCL_COMPILER_MSVC_2017
diff --git a/libcudacxx/include/cuda/std/__ranges/rbegin.h b/libcudacxx/include/cuda/std/__ranges/rbegin.h
index 9399228a4d..8fe23be3b2 100644
--- a/libcudacxx/include/cuda/std/__ranges/rbegin.h
+++ b/libcudacxx/include/cuda/std/__ranges/rbegin.h
@@ -130,7 +130,7 @@ _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto rbegin = __rbegin::__fn{};
+_CCCL_GLOBAL_CONSTANT auto rbegin = __rbegin::__fn{};
 } // namespace __cpo
 
 // [range.access.crbegin]
@@ -160,7 +160,7 @@ _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto crbegin = __crbegin::__fn{};
+_CCCL_GLOBAL_CONSTANT auto crbegin = __crbegin::__fn{};
 } // namespace __cpo
 
 #endif // _CCCL_STD_VER >= 2017 && && !_CCCL_COMPILER_MSVC_2017
diff --git a/libcudacxx/include/cuda/std/__ranges/rend.h b/libcudacxx/include/cuda/std/__ranges/rend.h
index 4a0e490e73..458480b14d 100644
--- a/libcudacxx/include/cuda/std/__ranges/rend.h
+++ b/libcudacxx/include/cuda/std/__ranges/rend.h
@@ -137,7 +137,7 @@ _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto rend = __rend::__fn{};
+_CCCL_GLOBAL_CONSTANT auto rend = __rend::__fn{};
 } // namespace __cpo
 
 // [range.access.crend]
@@ -167,7 +167,7 @@ _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto crend = __crend::__fn{};
+_CCCL_GLOBAL_CONSTANT auto crend = __crend::__fn{};
 } // namespace __cpo
 
 #endif // _CCCL_STD_VER >= 2017 && !_CCCL_COMPILER_MSVC_2017
diff --git a/libcudacxx/include/cuda/std/__ranges/size.h b/libcudacxx/include/cuda/std/__ranges/size.h
index feaa2b5148..849e55440c 100644
--- a/libcudacxx/include/cuda/std/__ranges/size.h
+++ b/libcudacxx/include/cuda/std/__ranges/size.h
@@ -157,7 +157,7 @@ _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto size = __size::__fn{};
+_CCCL_GLOBAL_CONSTANT auto size = __size::__fn{};
 } // namespace __cpo
 
 // [range.prim.ssize]
@@ -198,7 +198,7 @@ _LIBCUDACXX_END_NAMESPACE_CPO
 
 inline namespace __cpo
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY auto ssize = __ssize::__fn{};
+_CCCL_GLOBAL_CONSTANT auto ssize = __ssize::__fn{};
 } // namespace __cpo
 
 #endif // _CCCL_STD_VER >= 2017 && !_CCCL_COMPILER_MSVC_2017
diff --git a/libcudacxx/include/cuda/std/__utility/in_place.h b/libcudacxx/include/cuda/std/__utility/in_place.h
index 9bbd07a42b..cbd0aeb90e 100644
--- a/libcudacxx/include/cuda/std/__utility/in_place.h
+++ b/libcudacxx/include/cuda/std/__utility/in_place.h
@@ -32,7 +32,7 @@ struct _LIBCUDACXX_TYPE_VIS in_place_t
 {
   explicit in_place_t() = default;
 };
-_LIBCUDACXX_CPO_ACCESSIBILITY in_place_t in_place{};
+_CCCL_GLOBAL_CONSTANT in_place_t in_place{};
 
 template <class _Tp>
 struct _LIBCUDACXX_TEMPLATE_VIS in_place_type_t
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__config b/libcudacxx/include/cuda/std/detail/libcxx/include/__config
index d69f5216f9..fc823c46fd 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/__config
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__config
@@ -1743,19 +1743,6 @@ __sanitizer_annotate_contiguous_container(const void*, const void*, const void*,
 #    define _LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(_ClassName) static_assert(true, "")
 #  endif
 
-#  if (defined(__CUDACC_VER_MAJOR__) && __CUDACC_VER_MAJOR__ <= 11) \
-    && (defined(__CUDACC_VER_MINOR__) && __CUDACC_VER_MINOR__ <= 2)
-#    define _LIBCUDACXX_CONSTEXPR_GLOBAL const
-#  else
-#    define _LIBCUDACXX_CONSTEXPR_GLOBAL constexpr
-#  endif
-
-#  if defined(__CUDA_ARCH__)
-#    define _LIBCUDACXX_CPO_ACCESSIBILITY _CCCL_DEVICE _LIBCUDACXX_CONSTEXPR_GLOBAL
-#  else
-#    define _LIBCUDACXX_CPO_ACCESSIBILITY _LIBCUDACXX_INLINE_VAR constexpr
-#  endif
-
 // Older nvcc do not handle the constraint of `construct_at` in earlier std modes
 // So to preserve our performance optimization we default to the unconstrained
 // `__construct_at` and only in C++20 use `construct_at`
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/optional b/libcudacxx/include/cuda/std/detail/libcxx/include/optional
index c77ed212da..75190c3999 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/optional
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/optional
@@ -265,7 +265,7 @@ struct nullopt_t
   _LIBCUDACXX_INLINE_VISIBILITY constexpr explicit nullopt_t(__secret_tag, __secret_tag) noexcept {}
 };
 
-_LIBCUDACXX_CPO_ACCESSIBILITY nullopt_t nullopt{nullopt_t::__secret_tag{}, nullopt_t::__secret_tag{}};
+_CCCL_GLOBAL_CONSTANT nullopt_t nullopt{nullopt_t::__secret_tag{}, nullopt_t::__secret_tag{}};
 
 struct __optional_construct_from_invoke_tag
 {};
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/tuple b/libcudacxx/include/cuda/std/detail/libcxx/include/tuple
index 62802506dd..46b4f4db74 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/tuple
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/tuple
@@ -1037,7 +1037,7 @@ struct __ignore_t
 
 namespace
 {
-_LIBCUDACXX_CPO_ACCESSIBILITY __ignore_t<unsigned char> ignore{};
+_CCCL_GLOBAL_CONSTANT __ignore_t<unsigned char> ignore{};
 } // namespace
 
 template <class... _Tp>
diff --git a/libcudacxx/test/libcudacxx/std/ranges/range.access/begin.pass.cpp b/libcudacxx/test/libcudacxx/std/ranges/range.access/begin.pass.cpp
index bd1eed8799..2cfd0d1ab2 100644
--- a/libcudacxx/test/libcudacxx/std/ranges/range.access/begin.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/ranges/range.access/begin.pass.cpp
@@ -371,14 +371,14 @@ ASSERT_NOEXCEPT(cuda::std::ranges::begin(cuda::std::declval<int (&)[10]>()));
 ASSERT_NOEXCEPT(cuda::std::ranges::cbegin(cuda::std::declval<int (&)[10]>()));
 
 #if !defined(TEST_COMPILER_MSVC_2019) // broken noexcept
-_LIBCUDACXX_CPO_ACCESSIBILITY struct NoThrowMemberBegin
+_CCCL_GLOBAL_CONSTANT struct NoThrowMemberBegin
 {
   __host__ __device__ ThrowingIterator<int> begin() const noexcept; // auto(t.begin()) doesn't throw
 } ntmb;
 static_assert(noexcept(cuda::std::ranges::begin(ntmb)));
 static_assert(noexcept(cuda::std::ranges::cbegin(ntmb)));
 
-_LIBCUDACXX_CPO_ACCESSIBILITY struct NoThrowADLBegin
+_CCCL_GLOBAL_CONSTANT struct NoThrowADLBegin
 {
   __host__ __device__ friend ThrowingIterator<int> begin(NoThrowADLBegin&) noexcept; // auto(begin(t)) doesn't throw
   __host__ __device__ friend ThrowingIterator<int> begin(const NoThrowADLBegin&) noexcept;
@@ -388,7 +388,7 @@ static_assert(noexcept(cuda::std::ranges::cbegin(ntab)));
 #endif // !TEST_COMPILER_MSVC_2019
 
 #if !defined(TEST_COMPILER_ICC)
-_LIBCUDACXX_CPO_ACCESSIBILITY struct NoThrowMemberBeginReturnsRef
+_CCCL_GLOBAL_CONSTANT struct NoThrowMemberBeginReturnsRef
 {
   __host__ __device__ ThrowingIterator<int>& begin() const noexcept; // auto(t.begin()) may throw
 } ntmbrr;
@@ -396,7 +396,7 @@ static_assert(!noexcept(cuda::std::ranges::begin(ntmbrr)));
 static_assert(!noexcept(cuda::std::ranges::cbegin(ntmbrr)));
 #endif // !TEST_COMPILER_ICC
 
-_LIBCUDACXX_CPO_ACCESSIBILITY struct BeginReturnsArrayRef
+_CCCL_GLOBAL_CONSTANT struct BeginReturnsArrayRef
 {
   __host__ __device__ auto begin() const noexcept -> int (&)[10];
 } brar;
diff --git a/libcudacxx/test/libcudacxx/std/ranges/range.access/end.pass.cpp b/libcudacxx/test/libcudacxx/std/ranges/range.access/end.pass.cpp
index ef15449a4d..69f79cb6ba 100644
--- a/libcudacxx/test/libcudacxx/std/ranges/range.access/end.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/ranges/range.access/end.pass.cpp
@@ -420,7 +420,7 @@ ASSERT_NOEXCEPT(cuda::std::ranges::end(cuda::std::declval<int (&)[10]>()));
 ASSERT_NOEXCEPT(cuda::std::ranges::cend(cuda::std::declval<int (&)[10]>()));
 
 #if !defined(TEST_COMPILER_MSVC_2019) // broken noexcept
-_LIBCUDACXX_CPO_ACCESSIBILITY struct NoThrowMemberEnd
+_CCCL_GLOBAL_CONSTANT struct NoThrowMemberEnd
 {
   __host__ __device__ ThrowingIterator<int> begin() const;
   __host__ __device__ ThrowingIterator<int> end() const noexcept; // auto(t.end()) doesn't throw
@@ -428,7 +428,7 @@ _LIBCUDACXX_CPO_ACCESSIBILITY struct NoThrowMemberEnd
 static_assert(noexcept(cuda::std::ranges::end(ntme)));
 static_assert(noexcept(cuda::std::ranges::cend(ntme)));
 
-_LIBCUDACXX_CPO_ACCESSIBILITY struct NoThrowADLEnd
+_CCCL_GLOBAL_CONSTANT struct NoThrowADLEnd
 {
   __host__ __device__ ThrowingIterator<int> begin() const;
   __host__ __device__ friend ThrowingIterator<int> end(NoThrowADLEnd&) noexcept; // auto(end(t)) doesn't throw
@@ -439,7 +439,7 @@ static_assert(noexcept(cuda::std::ranges::cend(ntae)));
 #endif // !TEST_COMPILER_MSVC_2019
 
 #if !defined(TEST_COMPILER_ICC)
-_LIBCUDACXX_CPO_ACCESSIBILITY struct NoThrowMemberEndReturnsRef
+_CCCL_GLOBAL_CONSTANT struct NoThrowMemberEndReturnsRef
 {
   __host__ __device__ ThrowingIterator<int> begin() const;
   __host__ __device__ ThrowingIterator<int>& end() const noexcept; // auto(t.end()) may throw
@@ -448,7 +448,7 @@ static_assert(!noexcept(cuda::std::ranges::end(ntmerr)));
 static_assert(!noexcept(cuda::std::ranges::cend(ntmerr)));
 #endif // !TEST_COMPILER_ICC
 
-_LIBCUDACXX_CPO_ACCESSIBILITY struct EndReturnsArrayRef
+_CCCL_GLOBAL_CONSTANT struct EndReturnsArrayRef
 {
   __host__ __device__ auto begin() const noexcept -> int (&)[10];
   __host__ __device__ auto end() const noexcept -> int (&)[10];
diff --git a/libcudacxx/test/libcudacxx/std/ranges/range.access/rbegin.pass.cpp b/libcudacxx/test/libcudacxx/std/ranges/range.access/rbegin.pass.cpp
index d4e3245b74..373ac14a95 100644
--- a/libcudacxx/test/libcudacxx/std/ranges/range.access/rbegin.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/ranges/range.access/rbegin.pass.cpp
@@ -571,14 +571,14 @@ ASSERT_NOEXCEPT(cuda::std::ranges::rbegin(cuda::std::declval<int (&)[10]>()));
 ASSERT_NOEXCEPT(cuda::std::ranges::crbegin(cuda::std::declval<int (&)[10]>()));
 
 #if !defined(TEST_COMPILER_MSVC_2019) // broken noexcept
-_LIBCUDACXX_CPO_ACCESSIBILITY struct NoThrowMemberRBegin
+_CCCL_GLOBAL_CONSTANT struct NoThrowMemberRBegin
 {
   __host__ __device__ ThrowingIterator<int> rbegin() const noexcept; // auto(t.rbegin()) doesn't throw
 } ntmb;
 static_assert(noexcept(cuda::std::ranges::rbegin(ntmb)));
 static_assert(noexcept(cuda::std::ranges::crbegin(ntmb)));
 
-_LIBCUDACXX_CPO_ACCESSIBILITY struct NoThrowADLRBegin
+_CCCL_GLOBAL_CONSTANT struct NoThrowADLRBegin
 {
   __host__ __device__ friend ThrowingIterator<int> rbegin(NoThrowADLRBegin&) noexcept; // auto(rbegin(t)) doesn't throw
   __host__ __device__ friend ThrowingIterator<int> rbegin(const NoThrowADLRBegin&) noexcept;
@@ -588,7 +588,7 @@ static_assert(noexcept(cuda::std::ranges::crbegin(ntab)));
 #endif // !TEST_COMPILER_MSVC_2019
 
 #if !defined(TEST_COMPILER_ICC)
-_LIBCUDACXX_CPO_ACCESSIBILITY struct NoThrowMemberRBeginReturnsRef
+_CCCL_GLOBAL_CONSTANT struct NoThrowMemberRBeginReturnsRef
 {
   __host__ __device__ ThrowingIterator<int>& rbegin() const noexcept; // auto(t.rbegin()) may throw
 } ntmbrr;
@@ -596,7 +596,7 @@ static_assert(!noexcept(cuda::std::ranges::rbegin(ntmbrr)));
 static_assert(!noexcept(cuda::std::ranges::crbegin(ntmbrr)));
 #endif // !TEST_COMPILER_ICC
 
-_LIBCUDACXX_CPO_ACCESSIBILITY struct RBeginReturnsArrayRef
+_CCCL_GLOBAL_CONSTANT struct RBeginReturnsArrayRef
 {
   __host__ __device__ auto rbegin() const noexcept -> int (&)[10];
 } brar;
@@ -604,7 +604,7 @@ static_assert(noexcept(cuda::std::ranges::rbegin(brar)));
 static_assert(noexcept(cuda::std::ranges::crbegin(brar)));
 
 #if !defined(TEST_COMPILER_ICC)
-_LIBCUDACXX_CPO_ACCESSIBILITY struct NoThrowBeginThrowingEnd
+_CCCL_GLOBAL_CONSTANT struct NoThrowBeginThrowingEnd
 {
   __host__ __device__ int* begin() const noexcept;
   __host__ __device__ int* end() const;
@@ -613,7 +613,7 @@ static_assert(!noexcept(cuda::std::ranges::rbegin(ntbte)));
 static_assert(!noexcept(cuda::std::ranges::crbegin(ntbte)));
 #endif // !TEST_COMPILER_ICC
 
-_LIBCUDACXX_CPO_ACCESSIBILITY struct NoThrowEndThrowingBegin
+_CCCL_GLOBAL_CONSTANT struct NoThrowEndThrowingBegin
 {
   __host__ __device__ int* begin() const;
   __host__ __device__ int* end() const noexcept;
diff --git a/libcudacxx/test/libcudacxx/std/ranges/range.access/rend.pass.cpp b/libcudacxx/test/libcudacxx/std/ranges/range.access/rend.pass.cpp
index 44d281ac84..239b9d5387 100644
--- a/libcudacxx/test/libcudacxx/std/ranges/range.access/rend.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/ranges/range.access/rend.pass.cpp
@@ -629,7 +629,7 @@ ASSERT_NOEXCEPT(cuda::std::ranges::rend(cuda::std::declval<int (&)[10]>()));
 ASSERT_NOEXCEPT(cuda::std::ranges::crend(cuda::std::declval<int (&)[10]>()));
 
 #if !defined(TEST_COMPILER_MSVC_2019)
-_LIBCUDACXX_CPO_ACCESSIBILITY struct NoThrowMemberREnd
+_CCCL_GLOBAL_CONSTANT struct NoThrowMemberREnd
 {
   __host__ __device__ ThrowingIterator<int> rbegin() const;
   __host__ __device__ ThrowingIterator<int> rend() const noexcept; // auto(t.rend()) doesn't throw
@@ -637,7 +637,7 @@ _LIBCUDACXX_CPO_ACCESSIBILITY struct NoThrowMemberREnd
 static_assert(noexcept(cuda::std::ranges::rend(ntmre)));
 static_assert(noexcept(cuda::std::ranges::crend(ntmre)));
 
-_LIBCUDACXX_CPO_ACCESSIBILITY struct NoThrowADLREnd
+_CCCL_GLOBAL_CONSTANT struct NoThrowADLREnd
 {
   __host__ __device__ ThrowingIterator<int> rbegin() const;
   __host__ __device__ friend ThrowingIterator<int> rend(NoThrowADLREnd&) noexcept; // auto(rend(t)) doesn't throw
@@ -648,7 +648,7 @@ static_assert(noexcept(cuda::std::ranges::crend(ntare)));
 #endif // !TEST_COMPILER_MSVC_2019
 
 #if !defined(TEST_COMPILER_ICC)
-_LIBCUDACXX_CPO_ACCESSIBILITY struct NoThrowMemberREndReturnsRef
+_CCCL_GLOBAL_CONSTANT struct NoThrowMemberREndReturnsRef
 {
   __host__ __device__ ThrowingIterator<int> rbegin() const;
   __host__ __device__ ThrowingIterator<int>& rend() const noexcept; // auto(t.rend()) may throw
@@ -657,7 +657,7 @@ static_assert(!noexcept(cuda::std::ranges::rend(ntmrerr)));
 static_assert(!noexcept(cuda::std::ranges::crend(ntmrerr)));
 #endif // !TEST_COMPILER_ICC
 
-_LIBCUDACXX_CPO_ACCESSIBILITY struct REndReturnsArrayRef
+_CCCL_GLOBAL_CONSTANT struct REndReturnsArrayRef
 {
   __host__ __device__ auto rbegin() const noexcept -> int (&)[10];
   __host__ __device__ auto rend() const noexcept -> int (&)[10];
@@ -665,7 +665,7 @@ _LIBCUDACXX_CPO_ACCESSIBILITY struct REndReturnsArrayRef
 static_assert(noexcept(cuda::std::ranges::rend(rerar)));
 static_assert(noexcept(cuda::std::ranges::crend(rerar)));
 
-_LIBCUDACXX_CPO_ACCESSIBILITY struct NoThrowBeginThrowingEnd
+_CCCL_GLOBAL_CONSTANT struct NoThrowBeginThrowingEnd
 {
   __host__ __device__ int* begin() const noexcept;
   __host__ __device__ int* end() const;
@@ -674,7 +674,7 @@ static_assert(noexcept(cuda::std::ranges::rend(ntbte)));
 static_assert(noexcept(cuda::std::ranges::crend(ntbte)));
 
 #if !defined(TEST_COMPILER_ICC)
-_LIBCUDACXX_CPO_ACCESSIBILITY struct NoThrowEndThrowingBegin
+_CCCL_GLOBAL_CONSTANT struct NoThrowEndThrowingBegin
 {
   __host__ __device__ int* begin() const;
   __host__ __device__ int* end() const noexcept;
diff --git a/libcudacxx/test/support/test_macros.h b/libcudacxx/test/support/test_macros.h
index e6c8e42bc1..a066348d05 100644
--- a/libcudacxx/test/support/test_macros.h
+++ b/libcudacxx/test/support/test_macros.h
@@ -461,7 +461,7 @@ __host__ __device__ constexpr bool unused(T&&...)
 #  endif // not MSVC
 #endif
 
-#define TEST_CONSTEXPR_GLOBAL _LIBCUDACXX_CONSTEXPR_GLOBAL
+#define TEST_CONSTEXPR_GLOBAL _CCCL_CONSTEXPR_GLOBAL
 
 // Some convenience macros for checking nvcc versions
 #if defined(__CUDACC__) && _CCCL_CUDACC_VER < 1103000
diff --git a/thrust/thrust/optional.h b/thrust/thrust/optional.h
index a0add00e61..f761f24f85 100644
--- a/thrust/thrust/optional.h
+++ b/thrust/thrust/optional.h
@@ -736,7 +736,7 @@ struct nullopt_t
 /// foo(thrust::nullopt); //pass an empty optional
 /// ```
 #ifdef __CUDA_ARCH__
-__device__ static _LIBCUDACXX_CONSTEXPR_GLOBAL
+__device__ static _CCCL_CONSTEXPR_GLOBAL
 #else
 static constexpr
 #endif // __CUDA_ARCH__