diff --git a/libcudacxx/include/cuda/__functional/maximum.h b/libcudacxx/include/cuda/__functional/maximum.h
index 307512b2292..997bc058a33 100644
--- a/libcudacxx/include/cuda/__functional/maximum.h
+++ b/libcudacxx/include/cuda/__functional/maximum.h
@@ -26,6 +26,7 @@
 #include <cuda/std/__type_traits/common_type.h>
 #include <cuda/std/__type_traits/is_extended_floating_point.h>
 #include <cuda/std/__type_traits/is_floating_point.h>
+#include <cuda/std/__utility/ctad_support.h>
 
 #include <cuda/std/__cccl/prologue.h>
 
@@ -48,7 +49,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT maximum
     }
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(maximum);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(maximum);
 
 template <>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT maximum<void>
diff --git a/libcudacxx/include/cuda/__functional/minimum.h b/libcudacxx/include/cuda/__functional/minimum.h
index d53dfccd549..2c902a3e8ad 100644
--- a/libcudacxx/include/cuda/__functional/minimum.h
+++ b/libcudacxx/include/cuda/__functional/minimum.h
@@ -26,6 +26,7 @@
 #include <cuda/std/__type_traits/common_type.h>
 #include <cuda/std/__type_traits/is_extended_floating_point.h>
 #include <cuda/std/__type_traits/is_floating_point.h>
+#include <cuda/std/__utility/ctad_support.h>
 
 #include <cuda/std/__cccl/prologue.h>
 
@@ -48,7 +49,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT minimum
     }
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(minimum);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(minimum);
 
 template <>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT minimum<void>
diff --git a/libcudacxx/include/cuda/__mdspan/host_device_mdspan.h b/libcudacxx/include/cuda/__mdspan/host_device_mdspan.h
index 7c10352a88f..0464551e279 100644
--- a/libcudacxx/include/cuda/__mdspan/host_device_mdspan.h
+++ b/libcudacxx/include/cuda/__mdspan/host_device_mdspan.h
@@ -32,6 +32,7 @@
 #include <cuda/std/__type_traits/remove_all_extents.h>
 #include <cuda/std/__type_traits/remove_pointer.h>
 #include <cuda/std/__type_traits/remove_reference.h>
+#include <cuda/std/__utility/delegate_constructors.h>
 #include <cuda/std/mdspan>
 
 #include <cuda/std/__cccl/prologue.h>
@@ -45,7 +46,7 @@ template <typename _ElementType,
 class host_mdspan : public ::cuda::std::mdspan<_ElementType, _Extents, _LayoutPolicy, host_accessor<_AccessorPolicy>>
 {
 public:
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(
+  _CCCL_DELEGATE_CONSTRUCTORS(
     host_mdspan, ::cuda::std::mdspan, _ElementType, _Extents, _LayoutPolicy, host_accessor<_AccessorPolicy>);
 
   _CCCL_API friend constexpr void swap(host_mdspan& __x, host_mdspan& __y) noexcept
@@ -105,7 +106,7 @@ class device_mdspan
     : public ::cuda::std::mdspan<_ElementType, _Extents, _LayoutPolicy, device_accessor<_AccessorPolicy>>
 {
 public:
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(
+  _CCCL_DELEGATE_CONSTRUCTORS(
     device_mdspan, ::cuda::std::mdspan, _ElementType, _Extents, _LayoutPolicy, device_accessor<_AccessorPolicy>);
 
   _CCCL_API friend constexpr void swap(device_mdspan& __x, device_mdspan& __y) noexcept
@@ -167,7 +168,7 @@ class managed_mdspan
     : public ::cuda::std::mdspan<_ElementType, _Extents, _LayoutPolicy, managed_accessor<_AccessorPolicy>>
 {
 public:
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(
+  _CCCL_DELEGATE_CONSTRUCTORS(
     managed_mdspan, ::cuda::std::mdspan, _ElementType, _Extents, _LayoutPolicy, managed_accessor<_AccessorPolicy>);
 
   _CCCL_API friend constexpr void swap(managed_mdspan& __x, managed_mdspan& __y) noexcept
diff --git a/libcudacxx/include/cuda/__mdspan/restrict_mdspan.h b/libcudacxx/include/cuda/__mdspan/restrict_mdspan.h
index 2d064173ce6..c47fc9428e9 100644
--- a/libcudacxx/include/cuda/__mdspan/restrict_mdspan.h
+++ b/libcudacxx/include/cuda/__mdspan/restrict_mdspan.h
@@ -32,6 +32,7 @@
 #include <cuda/std/__type_traits/remove_all_extents.h>
 #include <cuda/std/__type_traits/remove_pointer.h>
 #include <cuda/std/__type_traits/remove_reference.h>
+#include <cuda/std/__utility/delegate_constructors.h>
 #include <cuda/std/mdspan>
 
 #include <cuda/std/__cccl/prologue.h>
@@ -46,7 +47,7 @@ class restrict_mdspan
     : public ::cuda::std::mdspan<_ElementType, _Extents, _LayoutPolicy, restrict_accessor<_AccessorPolicy>>
 {
 public:
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(
+  _CCCL_DELEGATE_CONSTRUCTORS(
     restrict_mdspan, ::cuda::std::mdspan, _ElementType, _Extents, _LayoutPolicy, restrict_accessor<_AccessorPolicy>);
 
   _CCCL_API friend constexpr void swap(restrict_mdspan& __x, restrict_mdspan& __y) noexcept
diff --git a/libcudacxx/include/cuda/__memory_resource/any_resource.h b/libcudacxx/include/cuda/__memory_resource/any_resource.h
index f4804d7df73..a705cf4b5c9 100644
--- a/libcudacxx/include/cuda/__memory_resource/any_resource.h
+++ b/libcudacxx/include/cuda/__memory_resource/any_resource.h
@@ -26,6 +26,7 @@
 #include <cuda/__memory_resource/resource.h>
 #include <cuda/__utility/basic_any.h>
 #include <cuda/std/__concepts/concept_macros.h>
+#include <cuda/std/__utility/delegate_constructors.h>
 #include <cuda/std/__utility/forward.h>
 #include <cuda/std/optional>
 
@@ -208,7 +209,7 @@ struct _CCCL_DECLSPEC_EMPTY_BASES any_synchronous_resource
     , __with_try_get_property<any_synchronous_resource<_Properties...>>
 {
   // Inherit constructors from __basic_any
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(any_synchronous_resource, ::cuda::__basic_any, __iresource<_Properties...>);
+  _CCCL_DELEGATE_CONSTRUCTORS(any_synchronous_resource, ::cuda::__basic_any, __iresource<_Properties...>);
 
   // any_resource is convertible to any_synchronous_resource
   _CCCL_TEMPLATE(class... _OtherProperties)
@@ -238,7 +239,7 @@ struct _CCCL_DECLSPEC_EMPTY_BASES any_resource
     , __with_try_get_property<any_resource<_Properties...>>
 {
   // Inherit constructors from __basic_any
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(any_resource, ::cuda::__basic_any, __iasync_resource<_Properties...>);
+  _CCCL_DELEGATE_CONSTRUCTORS(any_resource, ::cuda::__basic_any, __iasync_resource<_Properties...>);
 
   using default_queries = ::cuda::mr::properties_list<_Properties...>;
 
@@ -266,7 +267,7 @@ struct _CCCL_DECLSPEC_EMPTY_BASES synchronous_resource_ref
     , __with_try_get_property<synchronous_resource_ref<_Properties...>>
 {
   // Inherit constructors from __basic_any
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(synchronous_resource_ref, ::cuda::__basic_any, __iresource<_Properties...>&);
+  _CCCL_DELEGATE_CONSTRUCTORS(synchronous_resource_ref, ::cuda::__basic_any, __iresource<_Properties...>&);
 
   synchronous_resource_ref(const synchronous_resource_ref& __other) noexcept = default;
 
@@ -327,7 +328,7 @@ struct _CCCL_DECLSPEC_EMPTY_BASES resource_ref
     , __with_try_get_property<resource_ref<_Properties...>>
 {
   // Inherit other constructors from __basic_any
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(resource_ref, ::cuda::__basic_any, __iasync_resource<_Properties...>&);
+  _CCCL_DELEGATE_CONSTRUCTORS(resource_ref, ::cuda::__basic_any, __iasync_resource<_Properties...>&);
 
   resource_ref(const resource_ref& __other) noexcept = default;
 
diff --git a/libcudacxx/include/cuda/std/__atomic/api/owned.h b/libcudacxx/include/cuda/std/__atomic/api/owned.h
index cd2f01dd757..3f6951dea7e 100644
--- a/libcudacxx/include/cuda/std/__atomic/api/owned.h
+++ b/libcudacxx/include/cuda/std/__atomic/api/owned.h
@@ -44,9 +44,9 @@ struct __atomic_common
 
   __atomic_storage_t<_Tp> __a;
 
-#if defined(_LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE)
-  static constexpr bool is_always_lock_free = _LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), 0);
-#endif // defined(_LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE)
+#if defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE)
+  static constexpr bool is_always_lock_free = _CCCL_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), 0);
+#endif // defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE)
 
   _LIBCUDACXX_ATOMIC_COMMON_IMPL(, )
   _LIBCUDACXX_ATOMIC_COMMON_IMPL(, volatile)
@@ -63,9 +63,9 @@ struct __atomic_arithmetic
 
   __atomic_storage_t<_Tp> __a;
 
-#if defined(_LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE)
-  static constexpr bool is_always_lock_free = _LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), 0);
-#endif // defined(_LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE)
+#if defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE)
+  static constexpr bool is_always_lock_free = _CCCL_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), 0);
+#endif // defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE)
 
   _LIBCUDACXX_ATOMIC_COMMON_IMPL(, )
   _LIBCUDACXX_ATOMIC_COMMON_IMPL(, volatile)
@@ -85,9 +85,9 @@ struct __atomic_bitwise
 
   __atomic_storage_t<_Tp> __a;
 
-#if defined(_LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE)
-  static constexpr bool is_always_lock_free = _LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), 0);
-#endif // defined(_LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE)
+#if defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE)
+  static constexpr bool is_always_lock_free = _CCCL_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), 0);
+#endif // defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE)
 
   _LIBCUDACXX_ATOMIC_COMMON_IMPL(, )
   _LIBCUDACXX_ATOMIC_COMMON_IMPL(, volatile)
@@ -110,9 +110,9 @@ struct __atomic_pointer
 
   __atomic_storage_t<_Tp> __a;
 
-#if defined(_LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE)
-  static constexpr bool is_always_lock_free = _LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), 0);
-#endif // defined(_LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE)
+#if defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE)
+  static constexpr bool is_always_lock_free = _CCCL_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), 0);
+#endif // defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE)
 
   _LIBCUDACXX_ATOMIC_COMMON_IMPL(, )
   _LIBCUDACXX_ATOMIC_COMMON_IMPL(, volatile)
diff --git a/libcudacxx/include/cuda/std/__atomic/api/reference.h b/libcudacxx/include/cuda/std/__atomic/api/reference.h
index a19d260d516..f6747f4a4b9 100644
--- a/libcudacxx/include/cuda/std/__atomic/api/reference.h
+++ b/libcudacxx/include/cuda/std/__atomic/api/reference.h
@@ -42,9 +42,9 @@ struct __atomic_ref_common
 
   __atomic_ref_storage<_Tp> __a;
 
-#if defined(_LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE)
-  static constexpr bool is_always_lock_free = _LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), 0);
-#endif // defined(_LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE)
+#if defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE)
+  static constexpr bool is_always_lock_free = _CCCL_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), 0);
+#endif // defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE)
 
   _LIBCUDACXX_ATOMIC_COMMON_IMPL(const, )
 };
@@ -58,9 +58,9 @@ struct __atomic_ref_arithmetic
 
   __atomic_ref_storage<_Tp> __a;
 
-#if defined(_LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE)
-  static constexpr bool is_always_lock_free = _LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), 0);
-#endif // defined(_LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE)
+#if defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE)
+  static constexpr bool is_always_lock_free = _CCCL_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), 0);
+#endif // defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE)
 
   _LIBCUDACXX_ATOMIC_COMMON_IMPL(const, )
   _LIBCUDACXX_ATOMIC_ARITHMETIC_IMPL(const, )
@@ -75,9 +75,9 @@ struct __atomic_ref_bitwise
 
   __atomic_ref_storage<_Tp> __a;
 
-#if defined(_LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE)
-  static constexpr bool is_always_lock_free = _LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), 0);
-#endif // defined(_LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE)
+#if defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE)
+  static constexpr bool is_always_lock_free = _CCCL_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), 0);
+#endif // defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE)
 
   _LIBCUDACXX_ATOMIC_COMMON_IMPL(const, )
   _LIBCUDACXX_ATOMIC_ARITHMETIC_IMPL(const, )
@@ -93,9 +93,9 @@ struct __atomic_ref_pointer
 
   __atomic_ref_storage<_Tp> __a;
 
-#if defined(_LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE)
-  static constexpr bool is_always_lock_free = _LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), 0);
-#endif // defined(_LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE)
+#if defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE)
+  static constexpr bool is_always_lock_free = _CCCL_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), 0);
+#endif // defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE)
 
   _LIBCUDACXX_ATOMIC_COMMON_IMPL(const, )
   _LIBCUDACXX_ATOMIC_POINTER_IMPL(const, )
diff --git a/libcudacxx/include/cuda/std/__atomic/functions/cuda_local.h b/libcudacxx/include/cuda/std/__atomic/functions/cuda_local.h
index c88367d0d15..a63fb7f660f 100644
--- a/libcudacxx/include/cuda/std/__atomic/functions/cuda_local.h
+++ b/libcudacxx/include/cuda/std/__atomic/functions/cuda_local.h
@@ -41,13 +41,13 @@ _CCCL_BEGIN_NAMESPACE_CUDA_STD
 
 _CCCL_DEVICE inline bool __cuda_is_local(const volatile void* __ptr)
 {
-#  if defined(_LIBCUDACXX_ATOMIC_UNSAFE_AUTOMATIC_STORAGE) && !defined(_LIBCUDACXX_FORCE_PTX_AUTOMATIC_STORAGE_PATH)
+#  if defined(_CCCL_ATOMIC_UNSAFE_AUTOMATIC_STORAGE) && !defined(_LIBCUDACXX_FORCE_PTX_AUTOMATIC_STORAGE_PATH)
   return false;
-#  else // ^^^ _LIBCUDACXX_ATOMIC_UNSAFE_AUTOMATIC_STORAGE && !defined(_LIBCUDACXX_FORCE_PTX_AUTOMATIC_STORAGE_PATH) ^^^
-        // / vvv !_LIBCUDACXX_ATOMIC_UNSAFE_AUTOMATIC_STORAGE || defined(_LIBCUDACXX_FORCE_PTX_AUTOMATIC_STORAGE_PATH)
+#  else // ^^^ _CCCL_ATOMIC_UNSAFE_AUTOMATIC_STORAGE && !defined(_LIBCUDACXX_FORCE_PTX_AUTOMATIC_STORAGE_PATH) ^^^
+        // / vvv !_CCCL_ATOMIC_UNSAFE_AUTOMATIC_STORAGE || defined(_LIBCUDACXX_FORCE_PTX_AUTOMATIC_STORAGE_PATH)
         // vvv
   return ::cuda::device::is_address_from(__ptr, ::cuda::device::address_space::local);
-#  endif // ^^^ !_LIBCUDACXX_ATOMIC_UNSAFE_AUTOMATIC_STORAGE || defined(_LIBCUDACXX_FORCE_PTX_AUTOMATIC_STORAGE_PATH)
+#  endif // ^^^ !_CCCL_ATOMIC_UNSAFE_AUTOMATIC_STORAGE || defined(_LIBCUDACXX_FORCE_PTX_AUTOMATIC_STORAGE_PATH)
          // ^^^
 }
 
diff --git a/libcudacxx/include/cuda/std/__atomic/order.h b/libcudacxx/include/cuda/std/__atomic/order.h
index babaa892bd7..d1935c81e56 100644
--- a/libcudacxx/include/cuda/std/__atomic/order.h
+++ b/libcudacxx/include/cuda/std/__atomic/order.h
@@ -28,18 +28,17 @@
 
 _CCCL_BEGIN_NAMESPACE_CUDA_STD
 
-#define _LIBCUDACXX_CHECK_STORE_MEMORY_ORDER(__m)                                              \
-  _LIBCUDACXX_DIAGNOSE_WARNING(                                                                \
-    __m == memory_order_consume || __m == memory_order_acquire || __m == memory_order_acq_rel, \
-    "memory order argument to atomic operation is invalid")
-
-#define _LIBCUDACXX_CHECK_LOAD_MEMORY_ORDER(__m)                                           \
-  _LIBCUDACXX_DIAGNOSE_WARNING(__m == memory_order_release || __m == memory_order_acq_rel, \
-                               "memory order argument to atomic operation is invalid")
-
-#define _LIBCUDACXX_CHECK_EXCHANGE_MEMORY_ORDER(__m, __f)                                  \
-  _LIBCUDACXX_DIAGNOSE_WARNING(__f == memory_order_release || __f == memory_order_acq_rel, \
-                               "memory order argument to atomic operation is invalid")
+#define _LIBCUDACXX_CHECK_STORE_MEMORY_ORDER(__m)                                                                   \
+  _CCCL_DIAGNOSE_WARNING(__m == memory_order_consume || __m == memory_order_acquire || __m == memory_order_acq_rel, \
+                         "memory order argument to atomic operation is invalid")
+
+#define _LIBCUDACXX_CHECK_LOAD_MEMORY_ORDER(__m)                                     \
+  _CCCL_DIAGNOSE_WARNING(__m == memory_order_release || __m == memory_order_acq_rel, \
+                         "memory order argument to atomic operation is invalid")
+
+#define _LIBCUDACXX_CHECK_EXCHANGE_MEMORY_ORDER(__m, __f)                            \
+  _CCCL_DIAGNOSE_WARNING(__f == memory_order_release || __f == memory_order_acq_rel, \
+                         "memory order argument to atomic operation is invalid")
 
 #ifndef __ATOMIC_RELAXED
 #  define __ATOMIC_RELAXED 0
diff --git a/libcudacxx/include/cuda/std/__atomic/platform.h b/libcudacxx/include/cuda/std/__atomic/platform.h
index 1f62b394866..0464e19b0cb 100644
--- a/libcudacxx/include/cuda/std/__atomic/platform.h
+++ b/libcudacxx/include/cuda/std/__atomic/platform.h
@@ -66,13 +66,13 @@
 
 _CCCL_BEGIN_NAMESPACE_CUDA_STD
 
-#if defined(_LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE)
+#if defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE)
 template <typename _Tp>
 struct __atomic_is_always_lock_free
 {
   enum
   {
-    __value = _LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), 0)
+    __value = _CCCL_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), 0)
   };
 };
 #else
@@ -84,7 +84,7 @@ struct __atomic_is_always_lock_free
     __value = sizeof(_Tp) <= 8
   };
 };
-#endif // defined(_LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE)
+#endif // defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE)
 
 _CCCL_END_NAMESPACE_CUDA_STD
 
diff --git a/libcudacxx/include/cuda/std/__atomic/types/locked.h b/libcudacxx/include/cuda/std/__atomic/types/locked.h
index 90e6c2582cf..764456b4539 100644
--- a/libcudacxx/include/cuda/std/__atomic/types/locked.h
+++ b/libcudacxx/include/cuda/std/__atomic/types/locked.h
@@ -39,7 +39,7 @@ struct __atomic_locked_storage
   static constexpr __atomic_tag __tag = __atomic_tag::__atomic_locked_tag;
 
   _Tp __a_value;
-  mutable __atomic_storage<_LIBCUDACXX_ATOMIC_FLAG_TYPE> __a_lock;
+  mutable __atomic_storage<_CCCL_ATOMIC_FLAG_TYPE> __a_lock;
 
   _CCCL_HIDE_FROM_ABI explicit constexpr __atomic_locked_storage() noexcept = default;
 
@@ -51,24 +51,24 @@ struct __atomic_locked_storage
   template <typename _Sco>
   _CCCL_HOST_DEVICE inline void __lock(_Sco) const volatile noexcept
   {
-    while (1 == __atomic_exchange_dispatch(&__a_lock, _LIBCUDACXX_ATOMIC_FLAG_TYPE(true), memory_order_acquire, _Sco{}))
+    while (1 == __atomic_exchange_dispatch(&__a_lock, _CCCL_ATOMIC_FLAG_TYPE(true), memory_order_acquire, _Sco{}))
       /*spin*/;
   }
   template <typename _Sco>
   _CCCL_HOST_DEVICE inline void __lock(_Sco) const noexcept
   {
-    while (1 == __atomic_exchange_dispatch(&__a_lock, _LIBCUDACXX_ATOMIC_FLAG_TYPE(true), memory_order_acquire, _Sco{}))
+    while (1 == __atomic_exchange_dispatch(&__a_lock, _CCCL_ATOMIC_FLAG_TYPE(true), memory_order_acquire, _Sco{}))
       /*spin*/;
   }
   template <typename _Sco>
   _CCCL_HOST_DEVICE inline void __unlock(_Sco) const volatile noexcept
   {
-    __atomic_store_dispatch(&__a_lock, _LIBCUDACXX_ATOMIC_FLAG_TYPE(false), memory_order_release, _Sco{});
+    __atomic_store_dispatch(&__a_lock, _CCCL_ATOMIC_FLAG_TYPE(false), memory_order_release, _Sco{});
   }
   template <typename _Sco>
   _CCCL_HOST_DEVICE inline void __unlock(_Sco) const noexcept
   {
-    __atomic_store_dispatch(&__a_lock, _LIBCUDACXX_ATOMIC_FLAG_TYPE(false), memory_order_release, _Sco{});
+    __atomic_store_dispatch(&__a_lock, _CCCL_ATOMIC_FLAG_TYPE(false), memory_order_release, _Sco{});
   }
 };
 
diff --git a/libcudacxx/include/cuda/std/__expected/expected_base.h b/libcudacxx/include/cuda/std/__expected/expected_base.h
index 6ab02b1ecc0..32525b05a47 100644
--- a/libcudacxx/include/cuda/std/__expected/expected_base.h
+++ b/libcudacxx/include/cuda/std/__expected/expected_base.h
@@ -45,6 +45,7 @@
 #include <cuda/std/__type_traits/is_trivially_move_assignable.h>
 #include <cuda/std/__type_traits/is_trivially_move_constructible.h>
 #include <cuda/std/__type_traits/is_void.h>
+#include <cuda/std/__utility/delegate_constructors.h>
 #include <cuda/std/__utility/exception_guard.h>
 #include <cuda/std/__utility/forward.h>
 #include <cuda/std/__utility/in_place.h>
@@ -437,7 +438,7 @@ _CCCL_DIAG_POP
 template <class _Tp, class _Err>
 struct __expected_storage : __expected_destruct<_Tp, _Err>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__expected_storage, __expected_destruct, _Tp, _Err);
+  _CCCL_DELEGATE_CONSTRUCTORS(__expected_storage, __expected_destruct, _Tp, _Err);
 
   _CCCL_EXEC_CHECK_DISABLE
   _CCCL_TEMPLATE(class _T1, class _T2, class... _Args)
@@ -532,13 +533,13 @@ inline constexpr __smf_availability __expected_can_copy_construct =
 template <class _Tp, class _Err, __smf_availability = __expected_can_copy_construct<_Tp, _Err>>
 struct __expected_copy : __expected_storage<_Tp, _Err>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__expected_copy, __expected_storage, _Tp, _Err);
+  _CCCL_DELEGATE_CONSTRUCTORS(__expected_copy, __expected_storage, _Tp, _Err);
 };
 
 template <class _Tp, class _Err>
 struct __expected_copy<_Tp, _Err, __smf_availability::__available> : __expected_storage<_Tp, _Err>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__expected_copy, __expected_storage, _Tp, _Err);
+  _CCCL_DELEGATE_CONSTRUCTORS(__expected_copy, __expected_storage, _Tp, _Err);
 
   _CCCL_API inline _CCCL_CONSTEXPR_CXX20 __expected_copy(const __expected_copy& __other) noexcept(
     is_nothrow_copy_constructible_v<_Tp> && is_nothrow_copy_constructible_v<_Err>)
@@ -562,7 +563,7 @@ struct __expected_copy<_Tp, _Err, __smf_availability::__available> : __expected_
 template <class _Tp, class _Err>
 struct __expected_copy<_Tp, _Err, __smf_availability::__deleted> : __expected_storage<_Tp, _Err>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__expected_copy, __expected_storage, _Tp, _Err);
+  _CCCL_DELEGATE_CONSTRUCTORS(__expected_copy, __expected_storage, _Tp, _Err);
 
   __expected_copy(const __expected_copy&)                                = delete;
   _CCCL_HIDE_FROM_ABI __expected_copy(__expected_copy&&)                 = default;
@@ -581,13 +582,13 @@ inline constexpr __smf_availability __expected_can_move_construct =
 template <class _Tp, class _Err, __smf_availability = __expected_can_move_construct<_Tp, _Err>>
 struct __expected_move : __expected_copy<_Tp, _Err>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__expected_move, __expected_copy, _Tp, _Err);
+  _CCCL_DELEGATE_CONSTRUCTORS(__expected_move, __expected_copy, _Tp, _Err);
 };
 
 template <class _Tp, class _Err>
 struct __expected_move<_Tp, _Err, __smf_availability::__available> : __expected_copy<_Tp, _Err>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__expected_move, __expected_copy, _Tp, _Err);
+  _CCCL_DELEGATE_CONSTRUCTORS(__expected_move, __expected_copy, _Tp, _Err);
 
   _CCCL_HIDE_FROM_ABI __expected_move(const __expected_move&) = default;
 
@@ -614,7 +615,7 @@ struct __expected_move<_Tp, _Err, __smf_availability::__available> : __expected_
 template <class _Tp, class _Err>
 struct __expected_move<_Tp, _Err, __smf_availability::__deleted> : __expected_copy<_Tp, _Err>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__expected_move, __expected_copy, _Tp, _Err);
+  _CCCL_DELEGATE_CONSTRUCTORS(__expected_move, __expected_copy, _Tp, _Err);
 
   _CCCL_HIDE_FROM_ABI __expected_move(const __expected_move&)            = default;
   __expected_move(__expected_move&&)                                     = delete;
@@ -639,13 +640,13 @@ inline constexpr __smf_availability __expected_can_copy_assign =
 template <class _Tp, class _Err, __smf_availability = __expected_can_copy_assign<_Tp, _Err>>
 struct __expected_copy_assign : __expected_move<_Tp, _Err>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__expected_copy_assign, __expected_move, _Tp, _Err);
+  _CCCL_DELEGATE_CONSTRUCTORS(__expected_copy_assign, __expected_move, _Tp, _Err);
 };
 
 template <class _Tp, class _Err>
 struct __expected_copy_assign<_Tp, _Err, __smf_availability::__available> : __expected_move<_Tp, _Err>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__expected_copy_assign, __expected_move, _Tp, _Err);
+  _CCCL_DELEGATE_CONSTRUCTORS(__expected_copy_assign, __expected_move, _Tp, _Err);
 
   _CCCL_HIDE_FROM_ABI __expected_copy_assign(const __expected_copy_assign&) = default;
   _CCCL_HIDE_FROM_ABI __expected_copy_assign(__expected_copy_assign&&)      = default;
@@ -683,7 +684,7 @@ struct __expected_copy_assign<_Tp, _Err, __smf_availability::__available> : __ex
 template <class _Tp, class _Err>
 struct __expected_copy_assign<_Tp, _Err, __smf_availability::__deleted> : __expected_move<_Tp, _Err>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__expected_copy_assign, __expected_move, _Tp, _Err);
+  _CCCL_DELEGATE_CONSTRUCTORS(__expected_copy_assign, __expected_move, _Tp, _Err);
 
   _CCCL_HIDE_FROM_ABI __expected_copy_assign(const __expected_copy_assign&)       = default;
   _CCCL_HIDE_FROM_ABI __expected_copy_assign(__expected_copy_assign&&)            = default;
@@ -706,13 +707,13 @@ inline constexpr __smf_availability __expected_can_move_assign =
 template <class _Tp, class _Err, __smf_availability = __expected_can_move_assign<_Tp, _Err>>
 struct __expected_move_assign : __expected_copy_assign<_Tp, _Err>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__expected_move_assign, __expected_copy_assign, _Tp, _Err);
+  _CCCL_DELEGATE_CONSTRUCTORS(__expected_move_assign, __expected_copy_assign, _Tp, _Err);
 };
 
 template <class _Tp, class _Err>
 struct __expected_move_assign<_Tp, _Err, __smf_availability::__available> : __expected_copy_assign<_Tp, _Err>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__expected_move_assign, __expected_copy_assign, _Tp, _Err);
+  _CCCL_DELEGATE_CONSTRUCTORS(__expected_move_assign, __expected_copy_assign, _Tp, _Err);
 
   _CCCL_HIDE_FROM_ABI __expected_move_assign(const __expected_move_assign&)            = default;
   _CCCL_HIDE_FROM_ABI __expected_move_assign(__expected_move_assign&&)                 = default;
@@ -749,7 +750,7 @@ struct __expected_move_assign<_Tp, _Err, __smf_availability::__available> : __ex
 template <class _Tp, class _Err>
 struct __expected_move_assign<_Tp, _Err, __smf_availability::__deleted> : __expected_copy_assign<_Tp, _Err>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__expected_move_assign, __expected_copy_assign, _Tp, _Err);
+  _CCCL_DELEGATE_CONSTRUCTORS(__expected_move_assign, __expected_copy_assign, _Tp, _Err);
 
   _CCCL_HIDE_FROM_ABI __expected_move_assign(const __expected_move_assign&)            = default;
   _CCCL_HIDE_FROM_ABI __expected_move_assign(__expected_move_assign&&)                 = default;
@@ -914,7 +915,7 @@ _CCCL_DIAG_POP
 template <class _Err>
 struct __expected_storage<void, _Err> : __expected_destruct<void, _Err>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__expected_storage, __expected_destruct, void, _Err);
+  _CCCL_DELEGATE_CONSTRUCTORS(__expected_storage, __expected_destruct, void, _Err);
 
   _CCCL_EXEC_CHECK_DISABLE
   static _CCCL_API inline _CCCL_CONSTEXPR_CXX20 void __swap_val_unex_impl(
@@ -931,7 +932,7 @@ struct __expected_storage<void, _Err> : __expected_destruct<void, _Err>
 template <class _Err>
 struct __expected_copy<void, _Err, __smf_availability::__available> : __expected_storage<void, _Err>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__expected_copy, __expected_storage, void, _Err);
+  _CCCL_DELEGATE_CONSTRUCTORS(__expected_copy, __expected_storage, void, _Err);
 
   _CCCL_API inline _CCCL_CONSTEXPR_CXX20
   __expected_copy(const __expected_copy& __other) noexcept(is_nothrow_copy_constructible_v<_Err>)
@@ -951,7 +952,7 @@ struct __expected_copy<void, _Err, __smf_availability::__available> : __expected
 template <class _Err>
 struct __expected_move<void, _Err, __smf_availability::__available> : __expected_copy<void, _Err>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__expected_move, __expected_copy, void, _Err);
+  _CCCL_DELEGATE_CONSTRUCTORS(__expected_move, __expected_copy, void, _Err);
 
   _CCCL_HIDE_FROM_ABI __expected_move(const __expected_move&) = default;
 
@@ -973,7 +974,7 @@ struct __expected_move<void, _Err, __smf_availability::__available> : __expected
 template <class _Err>
 struct __expected_copy_assign<void, _Err, __smf_availability::__available> : __expected_move<void, _Err>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__expected_copy_assign, __expected_move, void, _Err);
+  _CCCL_DELEGATE_CONSTRUCTORS(__expected_copy_assign, __expected_move, void, _Err);
 
   _CCCL_HIDE_FROM_ABI __expected_copy_assign(const __expected_copy_assign&) = default;
   _CCCL_HIDE_FROM_ABI __expected_copy_assign(__expected_copy_assign&&)      = default;
@@ -1010,7 +1011,7 @@ struct __expected_copy_assign<void, _Err, __smf_availability::__available> : __e
 template <class _Err>
 struct __expected_move_assign<void, _Err, __smf_availability::__available> : __expected_copy_assign<void, _Err>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__expected_move_assign, __expected_copy_assign, void, _Err);
+  _CCCL_DELEGATE_CONSTRUCTORS(__expected_move_assign, __expected_copy_assign, void, _Err);
 
   _CCCL_HIDE_FROM_ABI __expected_move_assign(const __expected_move_assign&)            = default;
   _CCCL_HIDE_FROM_ABI __expected_move_assign(__expected_move_assign&&)                 = default;
diff --git a/libcudacxx/include/cuda/std/__format/format_context.h b/libcudacxx/include/cuda/std/__format/format_context.h
index fc2a78d9638..5bd3babb990 100644
--- a/libcudacxx/include/cuda/std/__format/format_context.h
+++ b/libcudacxx/include/cuda/std/__format/format_context.h
@@ -26,6 +26,7 @@
 #include <cuda/std/__fwd/format.h>
 #include <cuda/std/__iterator/back_insert_iterator.h>
 #include <cuda/std/__iterator/concepts.h>
+#include <cuda/std/__utility/ctad_support.h>
 #include <cuda/std/__utility/move.h>
 
 #include <cuda/std/__cccl/prologue.h>
@@ -76,7 +77,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT basic_format_context
   basic_format_args<basic_format_context> __args_;
 };
 
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(basic_format_context);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(basic_format_context);
 
 template <class _OutIt, class _CharT>
 [[nodiscard]] _CCCL_API basic_format_context<_OutIt, _CharT>
diff --git a/libcudacxx/include/cuda/std/__format/format_parse_context.h b/libcudacxx/include/cuda/std/__format/format_parse_context.h
index c6f0109d120..d317163b77e 100644
--- a/libcudacxx/include/cuda/std/__format/format_parse_context.h
+++ b/libcudacxx/include/cuda/std/__format/format_parse_context.h
@@ -24,6 +24,7 @@
 #include <cuda/std/__format/format_error.h>
 #include <cuda/std/__fwd/format.h>
 #include <cuda/std/__type_traits/is_constant_evaluated.h>
+#include <cuda/std/__utility/ctad_support.h>
 #include <cuda/std/string_view>
 
 #include <cuda/std/__cccl/prologue.h>
@@ -115,7 +116,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT basic_format_parse_context
   size_t __num_args_;
 };
 
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(basic_format_parse_context);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(basic_format_parse_context);
 
 _CCCL_END_NAMESPACE_CUDA_STD
 
diff --git a/libcudacxx/include/cuda/std/__functional/bind_back.h b/libcudacxx/include/cuda/std/__functional/bind_back.h
index ba608589bb4..52238096b9e 100644
--- a/libcudacxx/include/cuda/std/__functional/bind_back.h
+++ b/libcudacxx/include/cuda/std/__functional/bind_back.h
@@ -29,6 +29,7 @@
 #include <cuda/std/__type_traits/enable_if.h>
 #include <cuda/std/__type_traits/is_constructible.h>
 #include <cuda/std/__type_traits/is_move_constructible.h>
+#include <cuda/std/__utility/delegate_constructors.h>
 #include <cuda/std/__utility/forward.h>
 #include <cuda/std/__utility/integer_sequence.h>
 #include <cuda/std/tuple>
@@ -56,7 +57,7 @@ struct __bind_back_op<_NBound, index_sequence<_Ip...>>
 template <class _Fn, class _BoundArgs>
 struct __bind_back_t : __perfect_forward<__bind_back_op<tuple_size_v<_BoundArgs>>, _Fn, _BoundArgs>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(
+  _CCCL_DELEGATE_CONSTRUCTORS(
     __bind_back_t, __perfect_forward, __bind_back_op<tuple_size_v<_BoundArgs>>, _Fn, _BoundArgs);
 };
 
diff --git a/libcudacxx/include/cuda/std/__functional/bind_front.h b/libcudacxx/include/cuda/std/__functional/bind_front.h
index cea72f5b48a..09dc97c1fbc 100644
--- a/libcudacxx/include/cuda/std/__functional/bind_front.h
+++ b/libcudacxx/include/cuda/std/__functional/bind_front.h
@@ -29,6 +29,7 @@
 #include <cuda/std/__type_traits/is_constructible.h>
 #include <cuda/std/__type_traits/is_move_constructible.h>
 #include <cuda/std/__type_traits/is_nothrow_constructible.h>
+#include <cuda/std/__utility/delegate_constructors.h>
 #include <cuda/std/__utility/forward.h>
 
 #include <cuda/std/__cccl/prologue.h>
@@ -49,7 +50,7 @@ struct __bind_front_op
 template <class _Fn, class... _BoundArgs>
 struct __bind_front_t : __perfect_forward<__bind_front_op, _Fn, _BoundArgs...>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__bind_front_t, __perfect_forward, __bind_front_op, _Fn, _BoundArgs...);
+  _CCCL_DELEGATE_CONSTRUCTORS(__bind_front_t, __perfect_forward, __bind_front_op, _Fn, _BoundArgs...);
 };
 
 template <class _Fn, class... _Args>
diff --git a/libcudacxx/include/cuda/std/__functional/compose.h b/libcudacxx/include/cuda/std/__functional/compose.h
index df17e47355b..46d2ac0b10f 100644
--- a/libcudacxx/include/cuda/std/__functional/compose.h
+++ b/libcudacxx/include/cuda/std/__functional/compose.h
@@ -24,6 +24,7 @@
 #include <cuda/std/__functional/invoke.h>
 #include <cuda/std/__functional/perfect_forward.h>
 #include <cuda/std/__type_traits/decay.h>
+#include <cuda/std/__utility/delegate_constructors.h>
 #include <cuda/std/__utility/forward.h>
 
 #include <cuda/std/__cccl/prologue.h>
@@ -49,7 +50,7 @@ struct __compose_op
 template <class _Fn1, class _Fn2>
 struct __compose_t : __perfect_forward<__compose_op, _Fn1, _Fn2>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__compose_t, __perfect_forward, __compose_op, _Fn1, _Fn2);
+  _CCCL_DELEGATE_CONSTRUCTORS(__compose_t, __perfect_forward, __compose_op, _Fn1, _Fn2);
 };
 
 template <class _Fn1, class _Fn2>
diff --git a/libcudacxx/include/cuda/std/__functional/default_searcher.h b/libcudacxx/include/cuda/std/__functional/default_searcher.h
index a498c1a7c50..9928540e343 100644
--- a/libcudacxx/include/cuda/std/__functional/default_searcher.h
+++ b/libcudacxx/include/cuda/std/__functional/default_searcher.h
@@ -25,6 +25,7 @@
 #include <cuda/std/__functional/identity.h>
 #include <cuda/std/__functional/operations.h>
 #include <cuda/std/__iterator/iterator_traits.h>
+#include <cuda/std/__utility/ctad_support.h>
 #include <cuda/std/__utility/pair.h>
 
 #include <cuda/std/__cccl/prologue.h>
@@ -64,7 +65,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT default_searcher
   _ForwardIterator __last_;
   _BinaryPredicate __pred_;
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(default_searcher);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(default_searcher);
 
 #endif // __cuda_std__
 
diff --git a/libcudacxx/include/cuda/std/__functional/operations.h b/libcudacxx/include/cuda/std/__functional/operations.h
index a2cfb3fed75..6fba0512d98 100644
--- a/libcudacxx/include/cuda/std/__functional/operations.h
+++ b/libcudacxx/include/cuda/std/__functional/operations.h
@@ -23,6 +23,7 @@
 
 #include <cuda/std/__functional/binary_function.h>
 #include <cuda/std/__functional/unary_function.h>
+#include <cuda/std/__utility/ctad_support.h>
 #include <cuda/std/__utility/forward.h>
 
 #include <cuda/std/__cccl/prologue.h>
@@ -41,7 +42,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT plus : __binary_function<_Tp, _Tp, _Tp>
     return __x + __y;
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(plus);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(plus);
 
 template <>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT plus<void>
@@ -67,7 +68,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT minus : __binary_function<_Tp, _Tp, _Tp>
     return __x - __y;
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(minus);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(minus);
 
 template <>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT minus<void>
@@ -93,7 +94,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT multiplies : __binary_function<_Tp, _Tp, _T
     return __x * __y;
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(multiplies);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(multiplies);
 
 template <>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT multiplies<void>
@@ -119,7 +120,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT divides : __binary_function<_Tp, _Tp, _Tp>
     return __x / __y;
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(divides);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(divides);
 
 template <>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT divides<void>
@@ -145,7 +146,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT modulus : __binary_function<_Tp, _Tp, _Tp>
     return __x % __y;
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(modulus);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(modulus);
 
 template <>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT modulus<void>
@@ -171,7 +172,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT negate : __unary_function<_Tp, _Tp>
     return -__x;
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(negate);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(negate);
 
 template <>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT negate<void>
@@ -198,7 +199,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT bit_and : __binary_function<_Tp, _Tp, _Tp>
     return __x & __y;
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(bit_and);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(bit_and);
 
 template <>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT bit_and<void>
@@ -223,7 +224,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT bit_not : __unary_function<_Tp, _Tp>
     return ~__x;
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(bit_not);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(bit_not);
 
 template <>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT bit_not<void>
@@ -248,7 +249,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT bit_or : __binary_function<_Tp, _Tp, _Tp>
     return __x | __y;
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(bit_or);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(bit_or);
 
 template <>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT bit_or<void>
@@ -274,7 +275,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT bit_xor : __binary_function<_Tp, _Tp, _Tp>
     return __x ^ __y;
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(bit_xor);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(bit_xor);
 
 template <>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT bit_xor<void>
@@ -302,7 +303,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT equal_to : __binary_function<_Tp, _Tp, bool
     return __x == __y;
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(equal_to);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(equal_to);
 
 template <>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT equal_to<void>
@@ -328,7 +329,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT not_equal_to : __binary_function<_Tp, _Tp,
     return __x != __y;
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(not_equal_to);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(not_equal_to);
 
 template <>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT not_equal_to<void>
@@ -354,7 +355,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT less : __binary_function<_Tp, _Tp, bool>
     return __x < __y;
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(less);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(less);
 
 template <>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT less<void>
@@ -380,7 +381,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT less_equal : __binary_function<_Tp, _Tp, bo
     return __x <= __y;
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(less_equal);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(less_equal);
 
 template <>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT less_equal<void>
@@ -406,7 +407,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT greater_equal : __binary_function<_Tp, _Tp,
     return __x >= __y;
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(greater_equal);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(greater_equal);
 
 template <>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT greater_equal<void>
@@ -432,7 +433,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT greater : __binary_function<_Tp, _Tp, bool>
     return __x > __y;
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(greater);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(greater);
 
 template <>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT greater<void>
@@ -460,7 +461,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT logical_and : __binary_function<_Tp, _Tp, b
     return __x && __y;
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(logical_and);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(logical_and);
 
 template <>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT logical_and<void>
@@ -486,7 +487,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT logical_not : __unary_function<_Tp, bool>
     return !__x;
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(logical_not);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(logical_not);
 
 template <>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT logical_not<void>
@@ -511,7 +512,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT logical_or : __binary_function<_Tp, _Tp, bo
     return __x || __y;
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(logical_or);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(logical_or);
 
 template <>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT logical_or<void>
diff --git a/libcudacxx/include/cuda/std/__internal/atomic.h b/libcudacxx/include/cuda/std/__internal/atomic.h
new file mode 100644
index 00000000000..dd412117875
--- /dev/null
+++ b/libcudacxx/include/cuda/std/__internal/atomic.h
@@ -0,0 +1,55 @@
+//===---------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===---------------------------------------------------------------------===//
+
+#ifndef _CUDA_STD___INTERNAL_ATOMIC_H
+#define _CUDA_STD___INTERNAL_ATOMIC_H
+
+#include <cuda/__cccl_config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/std/__internal/features.h>
+
+#if _CCCL_HAS_CUDA_COMPILER()
+#  define _CCCL_ATOMIC_ALWAYS_LOCK_FREE(size, ptr) (size <= 8)
+#elif _CCCL_COMPILER(CLANG) || _CCCL_COMPILER(GCC)
+#  define _CCCL_ATOMIC_ALWAYS_LOCK_FREE(...) __atomic_always_lock_free(__VA_ARGS__)
+#endif // _CCCL_CUDA_COMPILER
+
+// Enable bypassing automatic storage checks in atomics when using CTK 12.2 and below and if NDEBUG is defined.
+// A compiler bug prevents the safe use of `__is_local` and PTX spacep until after 13.0.
+#ifndef _CCCL_ATOMIC_UNSAFE_AUTOMATIC_STORAGE
+#  if _CCCL_CUDACC_BELOW(13, 1) && !defined(NDEBUG)
+#    define _CCCL_ATOMIC_UNSAFE_AUTOMATIC_STORAGE
+#  endif // _CCCL_CUDACC_BELOW(13, 1)
+#endif // _CCCL_ATOMIC_UNSAFE_AUTOMATIC_STORAGE
+
+#define _CCCL_ATOMIC_FLAG_TYPE int
+
+// Clang provides 128b atomics as a builtin
+#if defined(CCCL_ENABLE_EXPERIMENTAL_HOST_ATOMICS_128B)
+#  define _CCCL_HOST_128_ATOMICS_ENABLED() 1
+#  define _CCCL_HOST_128_ATOMICS_MAYBE()   0
+// GCC does not provide 128b atomics, but they may be available as a library, this requires opt-in usage.
+// See: https://gcc.gnu.org/onlinedocs/gcc/x86-Options.html "-mcx16" for more
+#elif _CCCL_COMPILER(CLANG) || _CCCL_COMPILER(GCC)
+#  define _CCCL_HOST_128_ATOMICS_ENABLED() 0
+#  define _CCCL_HOST_128_ATOMICS_MAYBE()   1
+#else
+#  define _CCCL_HOST_128_ATOMICS_ENABLED() 0
+#  define _CCCL_HOST_128_ATOMICS_MAYBE()   0
+#endif
+
+#endif // _CUDA_STD___INTERNAL_ATOMIC_H
diff --git a/libcudacxx/include/cuda/std/__internal/features.h b/libcudacxx/include/cuda/std/__internal/features.h
index a36c62db5d9..f8d7b218c9c 100644
--- a/libcudacxx/include/cuda/std/__internal/features.h
+++ b/libcudacxx/include/cuda/std/__internal/features.h
@@ -21,7 +21,6 @@
 #endif // no system header
 
 #define _LIBCUDACXX_HAS_CXX20_CHRONO_LITERALS() (!_CCCL_COMPILER(CLANG) || _CCCL_STD_VER >= 2020)
-#define _LIBCUDACXX_HAS_EXTERNAL_ATOMIC_IMP()   1
 #define _LIBCUDACXX_HAS_MONOTONIC_CLOCK()       0
 #define _LIBCUDACXX_HAS_SPACESHIP_OPERATOR()    0
 
@@ -69,18 +68,38 @@
 #  define _LIBCUDACXX_HAS_NVBF16() 0
 #endif // _CCCL_HAS_NVBF16() && _CCCL_CTK_AT_LEAST(12, 2)
 
-// Clang provides 128b atomics as a builtin
-#if defined(CCCL_ENABLE_EXPERIMENTAL_HOST_ATOMICS_128B)
-#  define _CCCL_HOST_128_ATOMICS_ENABLED() 1
-#  define _CCCL_HOST_128_ATOMICS_MAYBE()   0
-// GCC does not provide 128b atomics, but they may be available as a library, this requires opt-in usage.
-// See: https://gcc.gnu.org/onlinedocs/gcc/x86-Options.html "-mcx16" for more
-#elif _CCCL_COMPILER(CLANG) || _CCCL_COMPILER(GCC)
-#  define _CCCL_HOST_128_ATOMICS_ENABLED() 0
-#  define _CCCL_HOST_128_ATOMICS_MAYBE()   1
+#if _CCCL_COMPILER(MSVC)
+#  define _CCCL_ALIGNAS_TYPE(x) alignas(x)
+#  define _CCCL_ALIGNAS(x)      __declspec(align(x))
+#elif _CCCL_HAS_FEATURE(cxx_alignas)
+#  define _CCCL_ALIGNAS_TYPE(x) alignas(x)
+#  define _CCCL_ALIGNAS(x)      alignas(x)
 #else
-#  define _CCCL_HOST_128_ATOMICS_ENABLED() 0
-#  define _CCCL_HOST_128_ATOMICS_MAYBE()   0
+#  define _CCCL_ALIGNAS_TYPE(x) __attribute__((__aligned__(alignof(x))))
+#  define _CCCL_ALIGNAS(x)      __attribute__((__aligned__(x)))
+#endif // !_CCCL_COMPILER(MSVC) && !_CCCL_HAS_FEATURE(cxx_alignas)
+
+// We can only expose constexpr allocations if the compiler supports it
+// For now disable constexpr allocation support until we can actually use
+#if 0 && defined(__cpp_constexpr_dynamic_alloc) && defined(__cpp_lib_constexpr_dynamic_alloc) && _CCCL_STD_VER >= 2020 \
+  && !_CCCL_COMPILER(NVRTC)
+#  define _CCCL_HAS_CONSTEXPR_ALLOCATION
+#  define _CCCL_CONSTEXPR_CXX20_ALLOCATION constexpr
+#else // ^^^ __cpp_constexpr_dynamic_alloc ^^^ / vvv !__cpp_constexpr_dynamic_alloc vvv
+#  define _CCCL_CONSTEXPR_CXX20_ALLOCATION
+#endif
+
+// Enable removed C++17 features
+#if defined(_LIBCUDACXX_ENABLE_CXX17_REMOVED_FEATURES)
+#  define _LIBCUDACXX_ENABLE_CXX17_REMOVED_BINDERS
+#endif // _LIBCUDACXX_ENABLE_CXX17_REMOVED_FEATURES
+
+#ifndef _CCCL_DISABLE_ADDITIONAL_DIAGNOSTICS
+#  define _CCCL_DIAGNOSE_WARNING(_COND, _MSG) _CCCL_DIAGNOSE_IF(_COND, _MSG, "warning")
+#  define _CCCL_DIAGNOSE_ERROR(_COND, _MSG)   _CCCL_DIAGNOSE_IF(_COND, _MSG, "error")
+#else
+#  define _CCCL_DIAGNOSE_WARNING(_COND, _MSG)
+#  define _CCCL_DIAGNOSE_ERROR(_COND, _MSG)
 #endif
 
 #endif // _CUDA_STD___INTERNAL_FEATURES_H
diff --git a/libcudacxx/include/cuda/std/__internal/namespaces.h b/libcudacxx/include/cuda/std/__internal/namespaces.h
index cfc8bd982ba..0ad4c8285d1 100644
--- a/libcudacxx/include/cuda/std/__internal/namespaces.h
+++ b/libcudacxx/include/cuda/std/__internal/namespaces.h
@@ -20,6 +20,8 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__internal/version.h>
+
 // During the header testing, we want to check if the code is wrapped by the prologue/epilogue
 #if defined(_CCCL_HEADER_TEST)
 #  define _LIBCUDACXX_PROLOGUE_INCLUDE_CHECK() \
diff --git a/libcudacxx/include/cuda/std/__internal/thread_api.h b/libcudacxx/include/cuda/std/__internal/thread_api.h
new file mode 100644
index 00000000000..ac8b6a0b7a9
--- /dev/null
+++ b/libcudacxx/include/cuda/std/__internal/thread_api.h
@@ -0,0 +1,58 @@
+//===---------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===---------------------------------------------------------------------===//
+
+#ifndef _CUDA_STD___INTERNAL_THREAD_API_H
+#define _CUDA_STD___INTERNAL_THREAD_API_H
+
+#include <cuda/__cccl_config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+// Thread API
+#ifndef _CCCL_HAS_THREAD_API_EXTERNAL
+#  if _CCCL_COMPILER(NVRTC) || defined(__EMSCRIPTEN__)
+#    define _CCCL_HAS_THREAD_API_EXTERNAL
+#  endif
+#endif // _CCCL_HAS_THREAD_API_EXTERNAL
+
+#ifndef _CCCL_HAS_THREAD_API_CUDA
+#  if ((_CCCL_DEVICE_COMPILATION() && !_CCCL_CUDA_COMPILER(NVHPC)) || defined(__EMSCRIPTEN__))
+#    define _CCCL_HAS_THREAD_API_CUDA
+#  endif // ((_CCCL_DEVICE_COMPILATION() && !_CCCL_CUDA_COMPILER(NVHPC)) || defined(__EMSCRIPTEN__))
+#endif // _CCCL_HAS_THREAD_API_CUDA
+
+#ifndef _CCCL_HAS_THREAD_API_WIN32
+#  if _CCCL_COMPILER(MSVC) && !defined(_CCCL_HAS_THREAD_API_CUDA)
+#    define _CCCL_HAS_THREAD_API_WIN32
+#  endif // _CCCL_COMPILER(MSVC) && !defined(_CCCL_HAS_THREAD_API_CUDA)
+#endif // _CCCL_HAS_THREAD_API_WIN32
+
+#if !defined(_CCCL_HAS_THREAD_API_PTHREAD) && !defined(_CCCL_HAS_THREAD_API_WIN32) \
+  && !defined(_CCCL_HAS_THREAD_API_EXTERNAL)
+#  if defined(__GNU__) || _CCCL_OS(LINUX) || _CCCL_OS(APPLE) || _CCCL_OS(QNX) \
+    || (defined(__MINGW32__) && _CCCL_HAS_INCLUDE(<pthread.h>))
+#    define _CCCL_HAS_THREAD_API_PTHREAD
+#  elif defined(_WIN32)
+#    define _CCCL_HAS_THREAD_API_WIN32
+#  else
+#    define _CCCL_UNSUPPORTED_THREAD_API
+#  endif // _CCCL_HAS_THREAD_API
+#endif
+
+#ifndef __STDCPP_THREADS__
+#  define __STDCPP_THREADS__ 1
+#endif // __STDCPP_THREADS__
+
+#endif // _CUDA_STD___INTERNAL_THREAD_API_H
diff --git a/libcudacxx/include/cuda/std/__internal/version.h b/libcudacxx/include/cuda/std/__internal/version.h
new file mode 100644
index 00000000000..d7be937c0da
--- /dev/null
+++ b/libcudacxx/include/cuda/std/__internal/version.h
@@ -0,0 +1,52 @@
+//===---------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===---------------------------------------------------------------------===//
+
+#ifndef _CUDA_STD___INTERNAL_VERSION_H
+#define _CUDA_STD___INTERNAL_VERSION_H
+
+#include <cuda/__cccl_config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/std/__cccl/version.h> // IWYU pragma: export
+
+#define _LIBCUDACXX_CUDA_API_VERSION       CCCL_VERSION
+#define _LIBCUDACXX_CUDA_API_VERSION_MAJOR CCCL_MAJOR_VERSION
+#define _LIBCUDACXX_CUDA_API_VERSION_MINOR CCCL_MINOR_VERSION
+#define _LIBCUDACXX_CUDA_API_VERSION_PATCH CCCL_PATCH_VERSION
+
+#ifndef _LIBCUDACXX_CUDA_ABI_VERSION_LATEST
+#  define _LIBCUDACXX_CUDA_ABI_VERSION_LATEST 4
+#endif
+
+#ifdef _LIBCUDACXX_CUDA_ABI_VERSION
+#  if _LIBCUDACXX_CUDA_ABI_VERSION != 4
+#    error Unsupported libcu++ ABI version requested. Only version 4 is allowed.
+#  endif
+#else
+#  define _LIBCUDACXX_CUDA_ABI_VERSION _LIBCUDACXX_CUDA_ABI_VERSION_LATEST
+#endif
+
+#if (_LIBCUDACXX_CUDA_ABI_VERSION < 4) && !defined(LIBCUDACXX_IGNORE_DEPRECATED_ABI)
+#  error "libcu++ ABIs older than version 4 are deprecated, define LIBCUDACXX_IGNORE_DEPRECATED_ABI to ignore"
+#endif
+
+#ifdef _LIBCUDACXX_PIPELINE_ASSUMED_ABI_VERSION
+#  if _LIBCUDACXX_PIPELINE_ASSUMED_ABI_VERSION != _LIBCUDACXX_CUDA_ABI_VERSION
+#    error cuda_pipeline.h has assumed a different libcu++ ABI version than provided by this library. To fix this, please include a libcu++ header before including cuda_pipeline.h, or upgrade to a version of the toolkit this version of libcu++ shipped in.
+#  endif
+#endif
+
+#endif // _CUDA_STD___INTERNAL_VERSION_H
diff --git a/libcudacxx/include/cuda/std/__iterator/back_insert_iterator.h b/libcudacxx/include/cuda/std/__iterator/back_insert_iterator.h
index 7aae18a6fe1..87474af66a1 100644
--- a/libcudacxx/include/cuda/std/__iterator/back_insert_iterator.h
+++ b/libcudacxx/include/cuda/std/__iterator/back_insert_iterator.h
@@ -25,6 +25,7 @@
 #include <cuda/std/__iterator/iterator.h>
 #include <cuda/std/__iterator/iterator_traits.h>
 #include <cuda/std/__memory/addressof.h>
+#include <cuda/std/__utility/ctad_support.h>
 #include <cuda/std/__utility/move.h>
 #include <cuda/std/cstddef>
 
@@ -148,7 +149,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT back_insert_iterator
   }
 };
 _CCCL_SUPPRESS_DEPRECATED_POP
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(back_insert_iterator);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(back_insert_iterator);
 
 template <class _Container>
 [[nodiscard]] _CCCL_API constexpr back_insert_iterator<_Container> back_inserter(_Container& __x) noexcept
diff --git a/libcudacxx/include/cuda/std/__iterator/counted_iterator.h b/libcudacxx/include/cuda/std/__iterator/counted_iterator.h
index 649b27ee94e..38ac7b7137e 100644
--- a/libcudacxx/include/cuda/std/__iterator/counted_iterator.h
+++ b/libcudacxx/include/cuda/std/__iterator/counted_iterator.h
@@ -41,6 +41,7 @@
 #include <cuda/std/__type_traits/is_nothrow_default_constructible.h>
 #include <cuda/std/__type_traits/is_nothrow_move_constructible.h>
 #include <cuda/std/__type_traits/void_t.h>
+#include <cuda/std/__utility/ctad_support.h>
 #include <cuda/std/__utility/exception_guard.h>
 #include <cuda/std/__utility/move.h>
 
@@ -319,8 +320,8 @@ class counted_iterator
   }
 
 #if _CCCL_STD_VER <= 2017
-  [[nodiscard]] _CCCL_API friend constexpr bool
-  operator==(const counted_iterator& __lhs, const counted_iterator& __rhs) noexcept
+  [[nodiscard]]
+  _CCCL_API friend constexpr bool operator==(const counted_iterator& __lhs, const counted_iterator& __rhs) noexcept
   {
     return __lhs.__count_ == __rhs.__count_;
   }
@@ -413,7 +414,7 @@ class counted_iterator
     return ::cuda::std::ranges::iter_swap(__x.__current_, __y.__current_);
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(counted_iterator);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(counted_iterator);
 _LIBCUDACXX_END_HIDDEN_FRIEND_NAMESPACE(counted_iterator)
 
 // Not a hidden friend because of MSVC
diff --git a/libcudacxx/include/cuda/std/__iterator/front_insert_iterator.h b/libcudacxx/include/cuda/std/__iterator/front_insert_iterator.h
index e64011b43c2..d05437eab80 100644
--- a/libcudacxx/include/cuda/std/__iterator/front_insert_iterator.h
+++ b/libcudacxx/include/cuda/std/__iterator/front_insert_iterator.h
@@ -24,6 +24,7 @@
 #include <cuda/std/__iterator/iterator.h>
 #include <cuda/std/__iterator/iterator_traits.h>
 #include <cuda/std/__memory/addressof.h>
+#include <cuda/std/__utility/ctad_support.h>
 #include <cuda/std/__utility/move.h>
 #include <cuda/std/cstddef>
 
@@ -84,7 +85,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT front_insert_iterator
   }
 };
 _CCCL_SUPPRESS_DEPRECATED_POP
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(front_insert_iterator);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(front_insert_iterator);
 
 template <class _Container>
 [[nodiscard]] _CCCL_API constexpr front_insert_iterator<_Container> front_inserter(_Container& __x) noexcept
diff --git a/libcudacxx/include/cuda/std/__iterator/move_iterator.h b/libcudacxx/include/cuda/std/__iterator/move_iterator.h
index 6e2e476cdf2..4d91b7d2a4b 100644
--- a/libcudacxx/include/cuda/std/__iterator/move_iterator.h
+++ b/libcudacxx/include/cuda/std/__iterator/move_iterator.h
@@ -45,6 +45,7 @@
 #include <cuda/std/__type_traits/is_nothrow_move_constructible.h>
 #include <cuda/std/__type_traits/is_reference.h>
 #include <cuda/std/__type_traits/remove_reference.h>
+#include <cuda/std/__utility/ctad_support.h>
 #include <cuda/std/__utility/declval.h>
 #include <cuda/std/__utility/move.h>
 
@@ -408,7 +409,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT move_iterator : public __move_iter_category_
     return ::cuda::std::ranges::iter_swap(__x.__current_, __y.__current_);
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(move_iterator);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(move_iterator);
 _LIBCUDACXX_END_HIDDEN_FRIEND_NAMESPACE(move_iterator)
 
 // Some compilers have issues determining __is_fancy_pointer
diff --git a/libcudacxx/include/cuda/std/__memory/allocate_at_least.h b/libcudacxx/include/cuda/std/__memory/allocate_at_least.h
index 28acde2764a..5be5e3a974a 100644
--- a/libcudacxx/include/cuda/std/__memory/allocate_at_least.h
+++ b/libcudacxx/include/cuda/std/__memory/allocate_at_least.h
@@ -23,6 +23,7 @@
 #endif // no system header
 
 #include <cuda/std/__memory/allocator_traits.h>
+#include <cuda/std/__utility/ctad_support.h>
 #include <cuda/std/cstddef>
 
 #include <cuda/std/__cccl/prologue.h>
@@ -36,7 +37,7 @@ struct allocation_result
   _Pointer ptr;
   size_t count;
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(allocation_result);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(allocation_result);
 
 template <class _Alloc>
 [[nodiscard]] _CCCL_API constexpr allocation_result<typename allocator_traits<_Alloc>::pointer>
diff --git a/libcudacxx/include/cuda/std/__optional/optional_base.h b/libcudacxx/include/cuda/std/__optional/optional_base.h
index 6c338a8963c..edbe6f1be61 100644
--- a/libcudacxx/include/cuda/std/__optional/optional_base.h
+++ b/libcudacxx/include/cuda/std/__optional/optional_base.h
@@ -42,6 +42,7 @@
 #include <cuda/std/__type_traits/is_trivially_move_assignable.h>
 #include <cuda/std/__type_traits/is_trivially_move_constructible.h>
 #include <cuda/std/__type_traits/remove_cv.h>
+#include <cuda/std/__utility/delegate_constructors.h>
 #include <cuda/std/__utility/forward.h>
 #include <cuda/std/__utility/in_place.h>
 #include <cuda/std/__utility/move.h>
@@ -188,7 +189,7 @@ struct __optional_destruct_base<_Tp, true>
 template <class _Tp>
 struct __optional_storage_base : __optional_destruct_base<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__optional_storage_base, __optional_destruct_base, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__optional_storage_base, __optional_destruct_base, _Tp);
 
   using value_type = _Tp;
 
@@ -267,13 +268,13 @@ inline constexpr __smf_availability __optional_can_copy_construct =
 template <class _Tp, __smf_availability = __optional_can_copy_construct<_Tp>>
 struct __optional_copy_base : __optional_storage_base<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__optional_copy_base, __optional_storage_base, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__optional_copy_base, __optional_storage_base, _Tp);
 };
 
 template <class _Tp>
 struct __optional_copy_base<_Tp, __smf_availability::__available> : __optional_storage_base<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__optional_copy_base, __optional_storage_base, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__optional_copy_base, __optional_storage_base, _Tp);
 
   // This ctor shouldn't need to initialize the base explicitly, but g++ 9 considers it to be uninitialized
   // during constexpr evaluation if it isn't initialized explicitly. This can be replaced with the pattern
@@ -293,7 +294,7 @@ struct __optional_copy_base<_Tp, __smf_availability::__available> : __optional_s
 template <class _Tp>
 struct __optional_copy_base<_Tp, __smf_availability::__deleted> : __optional_storage_base<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__optional_copy_base, __optional_storage_base, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__optional_copy_base, __optional_storage_base, _Tp);
   _CCCL_HIDE_FROM_ABI __optional_copy_base(const __optional_copy_base&)            = delete;
   _CCCL_HIDE_FROM_ABI __optional_copy_base(__optional_copy_base&&)                 = default;
   _CCCL_HIDE_FROM_ABI __optional_copy_base& operator=(const __optional_copy_base&) = default;
@@ -310,13 +311,13 @@ inline constexpr __smf_availability __optional_can_move_construct =
 template <class _Tp, __smf_availability = __optional_can_move_construct<_Tp>>
 struct __optional_move_base : __optional_copy_base<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__optional_move_base, __optional_copy_base, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__optional_move_base, __optional_copy_base, _Tp);
 };
 
 template <class _Tp>
 struct __optional_move_base<_Tp, __smf_availability::__available> : __optional_copy_base<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__optional_move_base, __optional_copy_base, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__optional_move_base, __optional_copy_base, _Tp);
 
   _CCCL_HIDE_FROM_ABI __optional_move_base(const __optional_move_base&) = default;
 
@@ -332,7 +333,7 @@ struct __optional_move_base<_Tp, __smf_availability::__available> : __optional_c
 template <class _Tp>
 struct __optional_move_base<_Tp, __smf_availability::__deleted> : __optional_copy_base<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__optional_move_base, __optional_copy_base, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__optional_move_base, __optional_copy_base, _Tp);
 
   _CCCL_HIDE_FROM_ABI __optional_move_base(const __optional_move_base&)            = default;
   _CCCL_HIDE_FROM_ABI __optional_move_base(__optional_move_base&&)                 = delete;
@@ -351,13 +352,13 @@ inline constexpr __smf_availability __optional_can_copy_assign =
 template <class _Tp, __smf_availability = __optional_can_copy_assign<_Tp>>
 struct __optional_copy_assign_base : __optional_move_base<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__optional_copy_assign_base, __optional_move_base, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__optional_copy_assign_base, __optional_move_base, _Tp);
 };
 
 template <class _Tp>
 struct __optional_copy_assign_base<_Tp, __smf_availability::__available> : __optional_move_base<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__optional_copy_assign_base, __optional_move_base, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__optional_copy_assign_base, __optional_move_base, _Tp);
 
   _CCCL_HIDE_FROM_ABI __optional_copy_assign_base(const __optional_copy_assign_base&) = default;
   _CCCL_HIDE_FROM_ABI __optional_copy_assign_base(__optional_copy_assign_base&&)      = default;
@@ -374,7 +375,7 @@ struct __optional_copy_assign_base<_Tp, __smf_availability::__available> : __opt
 template <class _Tp>
 struct __optional_copy_assign_base<_Tp, __smf_availability::__deleted> : __optional_move_base<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__optional_copy_assign_base, __optional_move_base, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__optional_copy_assign_base, __optional_move_base, _Tp);
 
   _CCCL_HIDE_FROM_ABI __optional_copy_assign_base(const __optional_copy_assign_base&)            = default;
   _CCCL_HIDE_FROM_ABI __optional_copy_assign_base(__optional_copy_assign_base&&)                 = default;
@@ -393,13 +394,13 @@ inline constexpr __smf_availability __optional_can_move_assign =
 template <class _Tp, __smf_availability = __optional_can_move_assign<_Tp>>
 struct __optional_move_assign_base : __optional_copy_assign_base<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__optional_move_assign_base, __optional_copy_assign_base, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__optional_move_assign_base, __optional_copy_assign_base, _Tp);
 };
 
 template <class _Tp>
 struct __optional_move_assign_base<_Tp, __smf_availability::__available> : __optional_copy_assign_base<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__optional_move_assign_base, __optional_copy_assign_base, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__optional_move_assign_base, __optional_copy_assign_base, _Tp);
 
   _CCCL_HIDE_FROM_ABI __optional_move_assign_base(const __optional_move_assign_base& __opt)      = default;
   _CCCL_HIDE_FROM_ABI __optional_move_assign_base(__optional_move_assign_base&&)                 = default;
@@ -416,7 +417,7 @@ struct __optional_move_assign_base<_Tp, __smf_availability::__available> : __opt
 template <class _Tp>
 struct __optional_move_assign_base<_Tp, __smf_availability::__deleted> : __optional_copy_assign_base<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__optional_move_assign_base, __optional_copy_assign_base, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__optional_move_assign_base, __optional_copy_assign_base, _Tp);
 
   _CCCL_HIDE_FROM_ABI __optional_move_assign_base(const __optional_move_assign_base& __opt)      = default;
   _CCCL_HIDE_FROM_ABI __optional_move_assign_base(__optional_move_assign_base&&)                 = default;
diff --git a/libcudacxx/include/cuda/std/__ranges/compressed_movable_box.h b/libcudacxx/include/cuda/std/__ranges/compressed_movable_box.h
index cf40659908e..4f32ac6ed66 100644
--- a/libcudacxx/include/cuda/std/__ranges/compressed_movable_box.h
+++ b/libcudacxx/include/cuda/std/__ranges/compressed_movable_box.h
@@ -46,6 +46,7 @@
 #include <cuda/std/__type_traits/is_trivially_move_assignable.h>
 #include <cuda/std/__type_traits/is_trivially_move_constructible.h>
 #include <cuda/std/__type_traits/remove_cv.h>
+#include <cuda/std/__utility/delegate_constructors.h>
 #include <cuda/std/__utility/forward.h>
 #include <cuda/std/__utility/in_place.h>
 #include <cuda/std/__utility/move.h>
@@ -346,7 +347,7 @@ inline constexpr __smf_availability __compressed_box_copy_construct_available =
 template <size_t _Index, class _Tp, __smf_availability = __compressed_box_copy_construct_available<_Tp>>
 struct _CCCL_DECLSPEC_EMPTY_BASES __compressed_box_copy_base : __compressed_box_base<_Index, _Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__compressed_box_copy_base, __compressed_box_base, _Index, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__compressed_box_copy_base, __compressed_box_base, _Index, _Tp);
 };
 
 //! @brief We must ensure we only copy when engaged
@@ -354,7 +355,7 @@ template <size_t _Index, class _Tp>
 struct _CCCL_DECLSPEC_EMPTY_BASES
 __compressed_box_copy_base<_Index, _Tp, __smf_availability::__available> : __compressed_box_base<_Index, _Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__compressed_box_copy_base, __compressed_box_base, _Index, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__compressed_box_copy_base, __compressed_box_base, _Index, _Tp);
 
   _CCCL_API _CCCL_CONSTEXPR_CXX20
   __compressed_box_copy_base(const __compressed_box_copy_base& __other) noexcept(is_nothrow_copy_constructible_v<_Tp>)
@@ -376,7 +377,7 @@ template <size_t _Index, class _Tp>
 struct _CCCL_DECLSPEC_EMPTY_BASES
 __compressed_box_copy_base<_Index, _Tp, __smf_availability::__deleted> : __compressed_box_base<_Index, _Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__compressed_box_copy_base, __compressed_box_base, _Index, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__compressed_box_copy_base, __compressed_box_base, _Index, _Tp);
   _CCCL_HIDE_FROM_ABI __compressed_box_copy_base(const __compressed_box_copy_base&)            = delete;
   _CCCL_HIDE_FROM_ABI __compressed_box_copy_base(__compressed_box_copy_base&&)                 = default;
   _CCCL_HIDE_FROM_ABI __compressed_box_copy_base& operator=(const __compressed_box_copy_base&) = default;
@@ -396,7 +397,7 @@ inline constexpr __smf_availability __compressed_box_move_construct_available =
 template <size_t _Index, class _Tp, __smf_availability = __compressed_box_move_construct_available<_Tp>>
 struct _CCCL_DECLSPEC_EMPTY_BASES __compressed_box_move_base : __compressed_box_copy_base<_Index, _Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__compressed_box_move_base, __compressed_box_copy_base, _Index, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__compressed_box_move_base, __compressed_box_copy_base, _Index, _Tp);
 };
 
 //! @brief We must ensure we only move when engaged
@@ -404,7 +405,7 @@ template <size_t _Index, class _Tp>
 struct _CCCL_DECLSPEC_EMPTY_BASES
 __compressed_box_move_base<_Index, _Tp, __smf_availability::__available> : __compressed_box_copy_base<_Index, _Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__compressed_box_move_base, __compressed_box_copy_base, _Index, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__compressed_box_move_base, __compressed_box_copy_base, _Index, _Tp);
 
   _CCCL_HIDE_FROM_ABI __compressed_box_move_base(const __compressed_box_move_base&) = default;
   _CCCL_API _CCCL_CONSTEXPR_CXX20
@@ -426,7 +427,7 @@ template <size_t _Index, class _Tp>
 struct _CCCL_DECLSPEC_EMPTY_BASES
 __compressed_box_move_base<_Index, _Tp, __smf_availability::__deleted> : __compressed_box_copy_base<_Index, _Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__compressed_box_move_base, __compressed_box_copy_base, _Index, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__compressed_box_move_base, __compressed_box_copy_base, _Index, _Tp);
   _CCCL_HIDE_FROM_ABI __compressed_box_move_base(const __compressed_box_move_base&)            = default;
   _CCCL_HIDE_FROM_ABI __compressed_box_move_base(__compressed_box_move_base&&)                 = delete;
   _CCCL_HIDE_FROM_ABI __compressed_box_move_base& operator=(const __compressed_box_move_base&) = default;
@@ -444,7 +445,7 @@ inline constexpr __smf_availability __compressed_box_copy_assign_available =
 template <size_t _Index, class _Tp, __smf_availability = __compressed_box_copy_assign_available<_Tp>>
 struct _CCCL_DECLSPEC_EMPTY_BASES __compressed_box_copy_assign_base : __compressed_box_move_base<_Index, _Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__compressed_box_copy_assign_base, __compressed_box_move_base, _Index, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__compressed_box_copy_assign_base, __compressed_box_move_base, _Index, _Tp);
 };
 
 //! @brief If we can either assign or copy construct we do so
@@ -452,7 +453,7 @@ template <size_t _Index, class _Tp>
 struct _CCCL_DECLSPEC_EMPTY_BASES __compressed_box_copy_assign_base<_Index, _Tp, __smf_availability::__available>
     : __compressed_box_move_base<_Index, _Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__compressed_box_copy_assign_base, __compressed_box_move_base, _Index, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__compressed_box_copy_assign_base, __compressed_box_move_base, _Index, _Tp);
 
   _CCCL_HIDE_FROM_ABI __compressed_box_copy_assign_base(const __compressed_box_copy_assign_base&) = default;
   _CCCL_HIDE_FROM_ABI __compressed_box_copy_assign_base(__compressed_box_copy_assign_base&&)      = default;
@@ -521,7 +522,7 @@ template <size_t _Index, class _Tp>
 struct _CCCL_DECLSPEC_EMPTY_BASES
 __compressed_box_copy_assign_base<_Index, _Tp, __smf_availability::__deleted> : __compressed_box_move_base<_Index, _Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__compressed_box_copy_assign_base, __compressed_box_move_base, _Index, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__compressed_box_copy_assign_base, __compressed_box_move_base, _Index, _Tp);
   _CCCL_HIDE_FROM_ABI __compressed_box_copy_assign_base(const __compressed_box_copy_assign_base&)            = default;
   _CCCL_HIDE_FROM_ABI __compressed_box_copy_assign_base(__compressed_box_copy_assign_base&&)                 = default;
   _CCCL_HIDE_FROM_ABI __compressed_box_copy_assign_base& operator=(const __compressed_box_copy_assign_base&) = delete;
@@ -539,7 +540,7 @@ inline constexpr __smf_availability __compressed_box_move_assign_available =
 template <size_t _Index, class _Tp, __smf_availability = __compressed_box_move_assign_available<_Tp>>
 struct _CCCL_DECLSPEC_EMPTY_BASES __compressed_box_move_assign_base : __compressed_box_copy_assign_base<_Index, _Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__compressed_box_move_assign_base, __compressed_box_copy_assign_base, _Index, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__compressed_box_move_assign_base, __compressed_box_copy_assign_base, _Index, _Tp);
 };
 
 //! @brief If we can either assign or copy construct we do so
@@ -547,7 +548,7 @@ template <size_t _Index, class _Tp>
 struct _CCCL_DECLSPEC_EMPTY_BASES __compressed_box_move_assign_base<_Index, _Tp, __smf_availability::__available>
     : __compressed_box_copy_assign_base<_Index, _Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__compressed_box_move_assign_base, __compressed_box_copy_assign_base, _Index, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__compressed_box_move_assign_base, __compressed_box_copy_assign_base, _Index, _Tp);
 
   _CCCL_HIDE_FROM_ABI __compressed_box_move_assign_base(const __compressed_box_move_assign_base&)            = default;
   _CCCL_HIDE_FROM_ABI __compressed_box_move_assign_base(__compressed_box_move_assign_base&&)                 = default;
@@ -616,7 +617,7 @@ template <size_t _Index, class _Tp>
 struct _CCCL_DECLSPEC_EMPTY_BASES __compressed_box_move_assign_base<_Index, _Tp, __smf_availability::__deleted>
     : __compressed_box_copy_assign_base<_Index, _Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__compressed_box_move_assign_base, __compressed_box_copy_assign_base, _Index, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__compressed_box_move_assign_base, __compressed_box_copy_assign_base, _Index, _Tp);
   _CCCL_HIDE_FROM_ABI __compressed_box_move_assign_base(const __compressed_box_move_assign_base&)            = default;
   _CCCL_HIDE_FROM_ABI __compressed_box_move_assign_base(__compressed_box_move_assign_base&&)                 = default;
   _CCCL_HIDE_FROM_ABI __compressed_box_move_assign_base& operator=(const __compressed_box_move_assign_base&) = default;
@@ -626,7 +627,7 @@ struct _CCCL_DECLSPEC_EMPTY_BASES __compressed_box_move_assign_base<_Index, _Tp,
 template <size_t _Index, class _Tp>
 struct _CCCL_DECLSPEC_EMPTY_BASES __compressed_box : __compressed_box_move_assign_base<_Index, _Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__compressed_box, __compressed_box_move_assign_base, _Index, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__compressed_box, __compressed_box_move_assign_base, _Index, _Tp);
 };
 
 template <class...>
diff --git a/libcudacxx/include/cuda/std/__ranges/movable_box.h b/libcudacxx/include/cuda/std/__ranges/movable_box.h
index 5864ba0788b..3878f951e7f 100644
--- a/libcudacxx/include/cuda/std/__ranges/movable_box.h
+++ b/libcudacxx/include/cuda/std/__ranges/movable_box.h
@@ -32,6 +32,7 @@
 #include <cuda/std/__type_traits/is_nothrow_copy_constructible.h>
 #include <cuda/std/__type_traits/is_nothrow_default_constructible.h>
 #include <cuda/std/__type_traits/is_nothrow_move_constructible.h>
+#include <cuda/std/__utility/delegate_constructors.h>
 #include <cuda/std/__utility/move.h>
 #include <cuda/std/optional>
 
@@ -139,7 +140,7 @@ struct __mb_optional_destruct_base<_Tp, true>
 template <class _Tp, bool = copy_constructible<_Tp>>
 struct __mb_optional_copy_assign : __mb_optional_destruct_base<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__mb_optional_copy_assign, __mb_optional_destruct_base, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__mb_optional_copy_assign, __mb_optional_destruct_base, _Tp);
 
   _CCCL_HIDE_FROM_ABI constexpr __mb_optional_copy_assign(const __mb_optional_copy_assign&) = delete;
   _CCCL_HIDE_FROM_ABI constexpr __mb_optional_copy_assign(__mb_optional_copy_assign&&)      = default;
@@ -151,7 +152,7 @@ struct __mb_optional_copy_assign : __mb_optional_destruct_base<_Tp>
 template <class _Tp>
 struct __mb_optional_copy_assign<_Tp, true> : __mb_optional_destruct_base<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__mb_optional_copy_assign, __mb_optional_destruct_base, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__mb_optional_copy_assign, __mb_optional_destruct_base, _Tp);
 
   _CCCL_HIDE_FROM_ABI constexpr __mb_optional_copy_assign(const __mb_optional_copy_assign&) = default;
   _CCCL_HIDE_FROM_ABI constexpr __mb_optional_copy_assign(__mb_optional_copy_assign&&)      = default;
@@ -178,13 +179,13 @@ struct __mb_optional_copy_assign<_Tp, true> : __mb_optional_destruct_base<_Tp>
 template <class _Tp, bool = movable<_Tp>>
 struct __mb_optional_move_assign : __mb_optional_copy_assign<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__mb_optional_move_assign, __mb_optional_copy_assign, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__mb_optional_move_assign, __mb_optional_copy_assign, _Tp);
 };
 
 template <class _Tp>
 struct __mb_optional_move_assign<_Tp, false> : __mb_optional_copy_assign<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__mb_optional_move_assign, __mb_optional_copy_assign, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__mb_optional_move_assign, __mb_optional_copy_assign, _Tp);
 
   _CCCL_HIDE_FROM_ABI constexpr __mb_optional_move_assign(const __mb_optional_move_assign&)            = default;
   _CCCL_HIDE_FROM_ABI constexpr __mb_optional_move_assign(__mb_optional_move_assign&&)                 = default;
@@ -211,7 +212,7 @@ struct __mb_optional_move_assign<_Tp, false> : __mb_optional_copy_assign<_Tp>
 template <class _Tp>
 struct __mb_optional_base : __mb_optional_move_assign<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__mb_optional_base, __mb_optional_move_assign, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__mb_optional_base, __mb_optional_move_assign, _Tp);
 
   [[nodiscard]] _CCCL_API constexpr _Tp const& operator*() const noexcept
   {
@@ -295,13 +296,13 @@ struct __mb_holder_base<_Tp, true>
 template <class _Tp, bool = copyable<_Tp>>
 struct __mb_copy_assign : __mb_holder_base<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__mb_copy_assign, __mb_holder_base, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__mb_copy_assign, __mb_holder_base, _Tp);
 };
 
 template <class _Tp>
 struct __mb_copy_assign<_Tp, false> : __mb_holder_base<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__mb_copy_assign, __mb_holder_base, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__mb_copy_assign, __mb_holder_base, _Tp);
 
   _CCCL_HIDE_FROM_ABI constexpr __mb_copy_assign(const __mb_copy_assign&) = default;
   _CCCL_HIDE_FROM_ABI constexpr __mb_copy_assign(__mb_copy_assign&&)      = default;
@@ -323,13 +324,13 @@ struct __mb_copy_assign<_Tp, false> : __mb_holder_base<_Tp>
 template <class _Tp, bool = movable<_Tp>>
 struct __mb_move_assign : __mb_copy_assign<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__mb_move_assign, __mb_copy_assign, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__mb_move_assign, __mb_copy_assign, _Tp);
 };
 
 template <class _Tp>
 struct __mb_move_assign<_Tp, false> : __mb_copy_assign<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__mb_move_assign, __mb_copy_assign, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__mb_move_assign, __mb_copy_assign, _Tp);
 
   _CCCL_HIDE_FROM_ABI constexpr __mb_move_assign(const __mb_move_assign&)            = default;
   _CCCL_HIDE_FROM_ABI constexpr __mb_move_assign(__mb_move_assign&&)                 = default;
@@ -353,7 +354,7 @@ struct __mb_move_assign<_Tp, false> : __mb_copy_assign<_Tp>
 template <class _Tp>
 struct __mb_base : __mb_move_assign<_Tp>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__mb_base, __mb_move_assign, _Tp);
+  _CCCL_DELEGATE_CONSTRUCTORS(__mb_base, __mb_move_assign, _Tp);
 
   [[nodiscard]] _CCCL_API constexpr _Tp const& operator*() const noexcept
   {
diff --git a/libcudacxx/include/cuda/std/__ranges/owning_view.h b/libcudacxx/include/cuda/std/__ranges/owning_view.h
index b8f12040443..2be3f4d164e 100644
--- a/libcudacxx/include/cuda/std/__ranges/owning_view.h
+++ b/libcudacxx/include/cuda/std/__ranges/owning_view.h
@@ -32,6 +32,7 @@
 #include <cuda/std/__type_traits/enable_if.h>
 #include <cuda/std/__type_traits/is_nothrow_default_constructible.h>
 #include <cuda/std/__type_traits/remove_cvref.h>
+#include <cuda/std/__utility/ctad_support.h>
 #include <cuda/std/__utility/move.h>
 #include <cuda/std/initializer_list>
 
@@ -150,7 +151,7 @@ class owning_view : public view_interface<owning_view<_Rp>>
     return ::cuda::std::ranges::data(__r_);
   }
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(owning_view);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(owning_view);
 
 template <class _Rp>
 inline constexpr bool enable_borrowed_range<owning_view<_Rp>> = enable_borrowed_range<_Rp>;
diff --git a/libcudacxx/include/cuda/std/__ranges/range_adaptor.h b/libcudacxx/include/cuda/std/__ranges/range_adaptor.h
index 69b8ff553e2..c363a88e624 100644
--- a/libcudacxx/include/cuda/std/__ranges/range_adaptor.h
+++ b/libcudacxx/include/cuda/std/__ranges/range_adaptor.h
@@ -32,6 +32,7 @@
 #include <cuda/std/__type_traits/enable_if.h>
 #include <cuda/std/__type_traits/is_nothrow_constructible.h>
 #include <cuda/std/__type_traits/remove_cvref.h>
+#include <cuda/std/__utility/ctad_support.h>
 #include <cuda/std/__utility/forward.h>
 #include <cuda/std/__utility/move.h>
 
@@ -59,7 +60,7 @@ struct __pipeable
       : _Fn(::cuda::std::move(__f))
   {}
 };
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(__pipeable);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(__pipeable);
 
 template <class _Tp>
 _CCCL_HOST_DEVICE _Tp __derived_from_range_adaptor_closure(__range_adaptor_closure<_Tp>*);
diff --git a/libcudacxx/include/cuda/std/__thread/threading_support.h b/libcudacxx/include/cuda/std/__thread/threading_support.h
index ed75d87318a..cab8003eb5c 100644
--- a/libcudacxx/include/cuda/std/__thread/threading_support.h
+++ b/libcudacxx/include/cuda/std/__thread/threading_support.h
@@ -23,17 +23,17 @@
 #include <cuda/std/__chrono/duration.h>
 #include <cuda/std/__chrono/high_resolution_clock.h>
 
-#if defined(_LIBCUDACXX_HAS_THREAD_API_EXTERNAL)
+#if defined(_CCCL_HAS_THREAD_API_EXTERNAL)
 #  include <cuda/std/__thread/threading_support_external.h>
-#endif // _LIBCUDACXX_HAS_THREAD_API_EXTERNAL
+#endif // _CCCL_HAS_THREAD_API_EXTERNAL
 
-#if defined(_LIBCUDACXX_HAS_THREAD_API_CUDA)
+#if defined(_CCCL_HAS_THREAD_API_CUDA)
 #  include <cuda/std/__thread/threading_support_cuda.h>
-#elif defined(_LIBCUDACXX_HAS_THREAD_API_PTHREAD)
+#elif defined(_CCCL_HAS_THREAD_API_PTHREAD)
 #  include <cuda/std/__thread/threading_support_pthread.h>
-#elif defined(_LIBCUDACXX_HAS_THREAD_API_WIN32)
+#elif defined(_CCCL_HAS_THREAD_API_WIN32)
 #  include <cuda/std/__thread/threading_support_win32.h>
-#else // ^^^ _LIBCUDACXX_HAS_THREAD_API_WIN32 ^^^ / vvv Unknown Thread API vvv
+#else // ^^^ _CCCL_HAS_THREAD_API_WIN32 ^^^ / vvv Unknown Thread API vvv
 #  error "Unknown Thread API"
 #endif // Unknown Thread API
 
diff --git a/libcudacxx/include/cuda/std/__thread/threading_support_cuda.h b/libcudacxx/include/cuda/std/__thread/threading_support_cuda.h
index 65fb4823560..285c93c9fa5 100644
--- a/libcudacxx/include/cuda/std/__thread/threading_support_cuda.h
+++ b/libcudacxx/include/cuda/std/__thread/threading_support_cuda.h
@@ -20,7 +20,7 @@
 #  pragma system_header
 #endif // no system header
 
-#if defined(_LIBCUDACXX_HAS_THREAD_API_CUDA)
+#if defined(_CCCL_HAS_THREAD_API_CUDA)
 
 #  include <cuda/std/__chrono/duration.h>
 #  include <cuda/std/climits>
@@ -42,6 +42,6 @@ _CCCL_END_NAMESPACE_CUDA_STD
 
 #  include <cuda/std/__cccl/epilogue.h>
 
-#endif // _LIBCUDACXX_HAS_THREAD_API_CUDA
+#endif // _CCCL_HAS_THREAD_API_CUDA
 
 #endif // _CUDA_STD___THREAD_THREADING_SUPPORT_CUDA_H
diff --git a/libcudacxx/include/cuda/std/__thread/threading_support_external.h b/libcudacxx/include/cuda/std/__thread/threading_support_external.h
index 4392f3a6c35..76059a9c51e 100644
--- a/libcudacxx/include/cuda/std/__thread/threading_support_external.h
+++ b/libcudacxx/include/cuda/std/__thread/threading_support_external.h
@@ -20,7 +20,7 @@
 #  pragma system_header
 #endif // no system header
 
-#if defined(_LIBCUDACXX_HAS_THREAD_API_EXTERNAL)
+#if defined(_CCCL_HAS_THREAD_API_EXTERNAL)
 
 #  include <cuda/std/__chrono/duration.h>
 
@@ -36,6 +36,6 @@ _CCCL_END_NAMESPACE_CUDA_STD
 
 #  include <cuda/std/__cccl/epilogue.h>
 
-#endif // _LIBCUDACXX_HAS_THREAD_API_EXTERNAL
+#endif // _CCCL_HAS_THREAD_API_EXTERNAL
 
 #endif // _CUDA_STD___THREAD_THREADING_SUPPORT_EXTERNAL_H
diff --git a/libcudacxx/include/cuda/std/__thread/threading_support_pthread.h b/libcudacxx/include/cuda/std/__thread/threading_support_pthread.h
index 703927bb52d..3bc6b1540c5 100644
--- a/libcudacxx/include/cuda/std/__thread/threading_support_pthread.h
+++ b/libcudacxx/include/cuda/std/__thread/threading_support_pthread.h
@@ -20,7 +20,7 @@
 #  pragma system_header
 #endif // no system header
 
-#if defined(_LIBCUDACXX_HAS_THREAD_API_PTHREAD)
+#if defined(_CCCL_HAS_THREAD_API_PTHREAD)
 
 #  include <cuda/std/__chrono/duration.h>
 #  include <cuda/std/__utility/cmp.h>
@@ -138,6 +138,6 @@ _CCCL_END_NAMESPACE_CUDA_STD
 
 #  include <cuda/std/__cccl/epilogue.h>
 
-#endif // !_LIBCUDACXX_HAS_THREAD_API_PTHREAD
+#endif // !_CCCL_HAS_THREAD_API_PTHREAD
 
 #endif // _CUDA_STD___THREAD_THREADING_SUPPORT_PTHREAD_H
diff --git a/libcudacxx/include/cuda/std/__thread/threading_support_win32.h b/libcudacxx/include/cuda/std/__thread/threading_support_win32.h
index 4283f12ccc3..e0b6a4c6469 100644
--- a/libcudacxx/include/cuda/std/__thread/threading_support_win32.h
+++ b/libcudacxx/include/cuda/std/__thread/threading_support_win32.h
@@ -20,7 +20,7 @@
 #  pragma system_header
 #endif // no system header
 
-#if defined(_LIBCUDACXX_HAS_THREAD_API_WIN32)
+#if defined(_CCCL_HAS_THREAD_API_WIN32)
 
 #  include <cuda/std/__chrono/duration.h>
 
@@ -82,6 +82,6 @@ _CCCL_END_NAMESPACE_CUDA_STD
 
 #  include <cuda/std/__cccl/epilogue.h>
 
-#endif // _LIBCUDACXX_HAS_THREAD_API_WIN32
+#endif // _CCCL_HAS_THREAD_API_WIN32
 
 #endif // _CUDA_STD___THREAD_THREADING_SUPPORT_H
diff --git a/libcudacxx/include/cuda/std/__type_traits/aligned_storage.h b/libcudacxx/include/cuda/std/__type_traits/aligned_storage.h
index 29dc03b46a1..58148a52f10 100644
--- a/libcudacxx/include/cuda/std/__type_traits/aligned_storage.h
+++ b/libcudacxx/include/cuda/std/__type_traits/aligned_storage.h
@@ -27,12 +27,18 @@
 
 #include <cuda/std/__cccl/prologue.h>
 
+#if !_CCCL_COMPILER(NVRTC)
+#  define _CCCL_PREFERRED_ALIGNOF(_Tp) __alignof(_Tp)
+#else // ^^^ !_CCCL_COMPILER(NVRTC) ^^^ / vvv _CCCL_COMPILER(NVRTC) vvv
+#  define _CCCL_PREFERRED_ALIGNOF(_Tp) alignof(_Tp)
+#endif // ^^^^ _CCCL_COMPILER(NVRTC) ^^^
+
 _CCCL_BEGIN_NAMESPACE_CUDA_STD
 
 template <class _Tp>
 struct __align_type
 {
-  static const size_t value = _LIBCUDACXX_PREFERRED_ALIGNOF(_Tp);
+  static const size_t value = _CCCL_PREFERRED_ALIGNOF(_Tp);
   using type                = _Tp;
 };
 
diff --git a/libcudacxx/include/cuda/std/__type_traits/aligned_union.h b/libcudacxx/include/cuda/std/__type_traits/aligned_union.h
index c12b1d05756..31a137b5f1c 100644
--- a/libcudacxx/include/cuda/std/__type_traits/aligned_union.h
+++ b/libcudacxx/include/cuda/std/__type_traits/aligned_union.h
@@ -47,7 +47,7 @@ template <size_t _Len, class _Type0, class... _Types>
 struct aligned_union
 {
   static const size_t alignment_value =
-    __static_max<_LIBCUDACXX_PREFERRED_ALIGNOF(_Type0), _LIBCUDACXX_PREFERRED_ALIGNOF(_Types)...>::value;
+    __static_max<_CCCL_PREFERRED_ALIGNOF(_Type0), _CCCL_PREFERRED_ALIGNOF(_Types)...>::value;
   static const size_t __len = __static_max<_Len, sizeof(_Type0), sizeof(_Types)...>::value;
   using type                = typename aligned_storage<__len, alignment_value>::type;
 };
diff --git a/libcudacxx/include/cuda/std/__utility/ctad_support.h b/libcudacxx/include/cuda/std/__utility/ctad_support.h
new file mode 100644
index 00000000000..cb072c2c788
--- /dev/null
+++ b/libcudacxx/include/cuda/std/__utility/ctad_support.h
@@ -0,0 +1,27 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _CUDA_STD___UTILITY_CTAD_SUPPORT_H
+#define _CUDA_STD___UTILITY_CTAD_SUPPORT_H
+
+#include <cuda/std/detail/__config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#define _CCCL_CTAD_SUPPORTED_FOR_TYPE(_ClassName) \
+  template <class... _Tag>                        \
+  _ClassName(typename _Tag::__allow_ctad...)->_ClassName<_Tag...>
+
+#endif // _CUDA_STD___UTILITY_CTAD_SUPPORT_H
diff --git a/libcudacxx/include/cuda/std/__utility/delegate_constructors.h b/libcudacxx/include/cuda/std/__utility/delegate_constructors.h
new file mode 100644
index 00000000000..e9c4982c611
--- /dev/null
+++ b/libcudacxx/include/cuda/std/__utility/delegate_constructors.h
@@ -0,0 +1,51 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _CUDA_STD___UTILITY_DELEGATE_CONSTRUCTORS_H
+#define _CUDA_STD___UTILITY_DELEGATE_CONSTRUCTORS_H
+
+#include <cuda/std/detail/__config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/std/__type_traits/is_constructible.h>
+#include <cuda/std/__type_traits/is_default_constructible.h>
+#include <cuda/std/__type_traits/is_nothrow_constructible.h>
+#include <cuda/std/__type_traits/is_nothrow_default_constructible.h>
+#include <cuda/std/__utility/forward.h>
+
+#include <cuda/std/__cccl/prologue.h>
+
+// NVRTC has a bug that prevented the use of delegated constructors, as it did not accept execution space annotations.
+// This creates a whole lot of boilerplate that we can avoid through a macro (see nvbug3961621)
+#if _CCCL_COMPILER(NVRTC, <, 12, 6)
+#  define _CCCL_DELEGATE_CONSTRUCTORS(__class, __baseclass, ...)                                                       \
+    using __base = __baseclass<__VA_ARGS__>;                                                                           \
+    _CCCL_TEMPLATE(class... _Args)                                                                                     \
+    _CCCL_REQUIRES(::cuda::std::is_constructible_v<__base, _Args...>)                                                  \
+    _CCCL_API constexpr __class(_Args&&... __args) noexcept(::cuda::std::is_nothrow_constructible_v<__base, _Args...>) \
+        : __base(::cuda::std::forward<_Args>(__args)...)                                                               \
+    {}                                                                                                                 \
+    _CCCL_HIDE_FROM_ABI constexpr __class() noexcept(::cuda::std::is_nothrow_default_constructible_v<__base>) = default;
+#else // ^^^ workaround ^^^ / vvv no workaround vvv
+#  define _CCCL_DELEGATE_CONSTRUCTORS(__class, __baseclass, ...) \
+    using __base = __baseclass<__VA_ARGS__>;                     \
+    using __base::__base;                                        \
+    _CCCL_HIDE_FROM_ABI constexpr __class() noexcept(::cuda::std::is_nothrow_default_constructible_v<__base>) = default;
+#endif // ^^^ no workaround ^^^
+
+#include <cuda/std/__cccl/epilogue.h>
+
+#endif // _CUDA_STD___UTILITY_DELEGATE_CONSTRUCTORS_H
diff --git a/libcudacxx/include/cuda/std/__utility/exception_guard.h b/libcudacxx/include/cuda/std/__utility/exception_guard.h
index 8ba321a8667..4fe888f2e94 100644
--- a/libcudacxx/include/cuda/std/__utility/exception_guard.h
+++ b/libcudacxx/include/cuda/std/__utility/exception_guard.h
@@ -21,6 +21,7 @@
 #endif // no system header
 
 #include <cuda/std/__type_traits/is_nothrow_move_constructible.h>
+#include <cuda/std/__utility/ctad_support.h>
 #include <cuda/std/__utility/exchange.h>
 #include <cuda/std/__utility/move.h>
 
@@ -105,7 +106,7 @@ struct __exception_guard_exceptions
   bool __completed_;
 };
 
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(__exception_guard_exceptions);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(__exception_guard_exceptions);
 
 template <class _Rollback>
 struct __exception_guard_noexceptions
@@ -138,7 +139,7 @@ struct __exception_guard_noexceptions
   bool __completed_ = false;
 };
 
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(__exception_guard_noexceptions);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(__exception_guard_noexceptions);
 
 #if !_CCCL_HAS_EXCEPTIONS()
 template <class _Rollback>
diff --git a/libcudacxx/include/cuda/std/atomic b/libcudacxx/include/cuda/std/atomic
index ad8cee1d67b..b6619996d43 100644
--- a/libcudacxx/include/cuda/std/atomic
+++ b/libcudacxx/include/cuda/std/atomic
@@ -17,7 +17,7 @@
 #ifdef _LIBCUDACXX_HAS_NO_ATOMIC_HEADER
 # error <atomic> is not implemented
 #endif
-#ifdef _LIBCUDACXX_UNSUPPORTED_THREAD_API
+#ifdef _CCCL_UNSUPPORTED_THREAD_API
 # error "<atomic> is not supported on this system"
 #endif
 #ifdef kill_dependency
@@ -562,40 +562,40 @@ atomic_fetch_xor_explicit(atomic<_Tp>* __o, _Tp __op, memory_order __m) noexcept
 
 struct atomic_flag
 {
-  __atomic_storage_t<_LIBCUDACXX_ATOMIC_FLAG_TYPE> __a;
+  __atomic_storage_t<_CCCL_ATOMIC_FLAG_TYPE> __a;
 
   _CCCL_API inline bool test(memory_order __m = memory_order_seq_cst) const volatile noexcept
   {
-    return _LIBCUDACXX_ATOMIC_FLAG_TYPE(true) == __atomic_load_dispatch(&__a, __m, __thread_scope_system_tag{});
+    return _CCCL_ATOMIC_FLAG_TYPE(true) == __atomic_load_dispatch(&__a, __m, __thread_scope_system_tag{});
   }
   _CCCL_API inline bool test(memory_order __m = memory_order_seq_cst) const noexcept
   {
-    return _LIBCUDACXX_ATOMIC_FLAG_TYPE(true) == __atomic_load_dispatch(&__a, __m, __thread_scope_system_tag{});
+    return _CCCL_ATOMIC_FLAG_TYPE(true) == __atomic_load_dispatch(&__a, __m, __thread_scope_system_tag{});
   }
 
   _CCCL_API inline bool test_and_set(memory_order __m = memory_order_seq_cst) volatile noexcept
   {
-    return __atomic_exchange_dispatch(&__a, _LIBCUDACXX_ATOMIC_FLAG_TYPE(true), __m, __thread_scope_system_tag{});
+    return __atomic_exchange_dispatch(&__a, _CCCL_ATOMIC_FLAG_TYPE(true), __m, __thread_scope_system_tag{});
   }
   _CCCL_API inline bool test_and_set(memory_order __m = memory_order_seq_cst) noexcept
   {
-    return __atomic_exchange_dispatch(&__a, _LIBCUDACXX_ATOMIC_FLAG_TYPE(true), __m, __thread_scope_system_tag{});
+    return __atomic_exchange_dispatch(&__a, _CCCL_ATOMIC_FLAG_TYPE(true), __m, __thread_scope_system_tag{});
   }
   _CCCL_API inline void clear(memory_order __m = memory_order_seq_cst) volatile noexcept
   {
-    __atomic_store_dispatch(&__a, _LIBCUDACXX_ATOMIC_FLAG_TYPE(false), __m, __thread_scope_system_tag{});
+    __atomic_store_dispatch(&__a, _CCCL_ATOMIC_FLAG_TYPE(false), __m, __thread_scope_system_tag{});
   }
   _CCCL_API inline void clear(memory_order __m = memory_order_seq_cst) noexcept
   {
-    __atomic_store_dispatch(&__a, _LIBCUDACXX_ATOMIC_FLAG_TYPE(false), __m, __thread_scope_system_tag{});
+    __atomic_store_dispatch(&__a, _CCCL_ATOMIC_FLAG_TYPE(false), __m, __thread_scope_system_tag{});
   }
 
-  _CCCL_API inline void wait(_LIBCUDACXX_ATOMIC_FLAG_TYPE __v, memory_order __m = memory_order_seq_cst) const
+  _CCCL_API inline void wait(_CCCL_ATOMIC_FLAG_TYPE __v, memory_order __m = memory_order_seq_cst) const
     volatile noexcept
   {
     __atomic_wait(&__a, __v, __m, __thread_scope_system_tag{});
   }
-  _CCCL_API inline void wait(_LIBCUDACXX_ATOMIC_FLAG_TYPE __v, memory_order __m = memory_order_seq_cst) const noexcept
+  _CCCL_API inline void wait(_CCCL_ATOMIC_FLAG_TYPE __v, memory_order __m = memory_order_seq_cst) const noexcept
   {
     __atomic_wait(&__a, __v, __m, __thread_scope_system_tag{});
   }
diff --git a/libcudacxx/include/cuda/std/detail/__config b/libcudacxx/include/cuda/std/detail/__config
index d8ac073f48d..5cf896c89c2 100644
--- a/libcudacxx/include/cuda/std/detail/__config
+++ b/libcudacxx/include/cuda/std/detail/__config
@@ -11,35 +11,11 @@
 #ifndef __cuda_std__
 #define __cuda_std__
 
-#include <cuda/std/__cccl/version.h> // IWYU pragma: export
-
-#define _LIBCUDACXX_CUDA_API_VERSION       CCCL_VERSION
-#define _LIBCUDACXX_CUDA_API_VERSION_MAJOR CCCL_MAJOR_VERSION
-#define _LIBCUDACXX_CUDA_API_VERSION_MINOR CCCL_MINOR_VERSION
-#define _LIBCUDACXX_CUDA_API_VERSION_PATCH CCCL_PATCH_VERSION
-
-#ifndef _LIBCUDACXX_CUDA_ABI_VERSION_LATEST
-#  define _LIBCUDACXX_CUDA_ABI_VERSION_LATEST 4
-#endif
-
-#ifdef _LIBCUDACXX_CUDA_ABI_VERSION
-#  if _LIBCUDACXX_CUDA_ABI_VERSION != 4
-#    error Unsupported libcu++ ABI version requested. Only version 4 is allowed.
-#  endif
-#else
-#  define _LIBCUDACXX_CUDA_ABI_VERSION _LIBCUDACXX_CUDA_ABI_VERSION_LATEST
-#endif
-
-#if (_LIBCUDACXX_CUDA_ABI_VERSION < 4) && !defined(LIBCUDACXX_IGNORE_DEPRECATED_ABI)
-#  error "libcu++ ABIs older than version 4 are deprecated, define LIBCUDACXX_IGNORE_DEPRECATED_ABI to ignore"
-#endif
-
-#ifdef _LIBCUDACXX_PIPELINE_ASSUMED_ABI_VERSION
-#  if _LIBCUDACXX_PIPELINE_ASSUMED_ABI_VERSION != _LIBCUDACXX_CUDA_ABI_VERSION
-#    error cuda_pipeline.h has assumed a different libcu++ ABI version than provided by this library. To fix this, please include a libcu++ header before including cuda_pipeline.h, or upgrade to a version of the toolkit this version of libcu++ shipped in.
-#  endif
-#endif
-
-#include <cuda/std/detail/libcxx/include/__config> // IWYU pragma: export
+#include <cuda/std/__internal/atomic.h>
+#include <cuda/std/__internal/cpp_dialect.h>
+#include <cuda/std/__internal/features.h>
+#include <cuda/std/__internal/namespaces.h>
+#include <cuda/std/__internal/thread_api.h>
+#include <cuda/std/__internal/version.h>
 
 #endif //__cuda_std__
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__config b/libcudacxx/include/cuda/std/detail/libcxx/include/__config
deleted file mode 100644
index 64e509f780a..00000000000
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/__config
+++ /dev/null
@@ -1,207 +0,0 @@
-//===----------------------------------------------------------------------===//
-//
-// Part of libcu++, the C++ Standard Library for your entire system,
-// under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef _LIBCUDACXX_CONFIG
-#define _LIBCUDACXX_CONFIG
-
-#include <cuda/__cccl_config> // IWYU pragma: export
-
-#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
-#  pragma GCC system_header
-#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
-#  pragma clang system_header
-#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
-#  pragma system_header
-#endif // no system header
-
-#include <cuda/std/__internal/cpp_dialect.h>
-#include <cuda/std/__internal/features.h>
-#include <cuda/std/__internal/namespaces.h>
-
-#ifdef __cplusplus
-
-// __config may be included in `extern "C"` contexts, switch back to include <nv/target>
-extern "C++" {
-#  include <nv/target>
-}
-
-#  define _LIBCUDACXX_VERSION 10000
-
-#  ifndef __has_extension
-#    define __has_extension(__x) 0
-#  endif
-
-#  if !_CCCL_COMPILER(NVRTC)
-#    define _LIBCUDACXX_PREFERRED_ALIGNOF(_Tp) __alignof(_Tp)
-#  else // ^^^ !_CCCL_COMPILER(NVRTC) ^^^ / vvv _CCCL_COMPILER(NVRTC) vvv
-#    define _LIBCUDACXX_PREFERRED_ALIGNOF(_Tp) alignof(_Tp)
-#  endif // ^^^^ _CCCL_COMPILER(NVRTC) ^^^
-
-#  if _CCCL_COMPILER(MSVC)
-#    define _CCCL_ALIGNAS_TYPE(x) alignas(x)
-#    define _CCCL_ALIGNAS(x)      __declspec(align(x))
-#  elif _CCCL_HAS_FEATURE(cxx_alignas)
-#    define _CCCL_ALIGNAS_TYPE(x) alignas(x)
-#    define _CCCL_ALIGNAS(x)      alignas(x)
-#  else
-#    define _CCCL_ALIGNAS_TYPE(x) __attribute__((__aligned__(alignof(x))))
-#    define _CCCL_ALIGNAS(x)      __attribute__((__aligned__(x)))
-#  endif // !_CCCL_COMPILER(MSVC) && !_CCCL_HAS_FEATURE(cxx_alignas)
-
-#  if _CCCL_HAS_CUDA_COMPILER()
-#    define _LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE(size, ptr) (size <= 8)
-#  elif _CCCL_COMPILER(CLANG) || _CCCL_COMPILER(GCC)
-#    define _LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE(...) __atomic_always_lock_free(__VA_ARGS__)
-#  endif // _CCCL_CUDA_COMPILER
-
-#  ifdef _LIBCUDACXX_DEBUG
-#    if _LIBCUDACXX_DEBUG == 0
-#      define _LIBCUDACXX_DEBUG_LEVEL 1
-#    elif _LIBCUDACXX_DEBUG == 1
-#      define _LIBCUDACXX_DEBUG_LEVEL 2
-#    else
-#      error Supported values for _LIBCUDACXX_DEBUG are 0 and 1
-#    endif
-#  endif
-
-// Thread API
-#  ifndef _LIBCUDACXX_HAS_THREAD_API_EXTERNAL
-#    if _CCCL_COMPILER(NVRTC) || defined(__EMSCRIPTEN__)
-#      define _LIBCUDACXX_HAS_THREAD_API_EXTERNAL
-#    endif
-#  endif // _LIBCUDACXX_HAS_THREAD_API_EXTERNAL
-
-#  ifndef _LIBCUDACXX_HAS_THREAD_API_CUDA
-#    if ((_CCCL_DEVICE_COMPILATION() && !_CCCL_CUDA_COMPILER(NVHPC)) || defined(__EMSCRIPTEN__))
-#      define _LIBCUDACXX_HAS_THREAD_API_CUDA
-#    endif // __cuda_std__
-#  endif // _LIBCUDACXX_HAS_THREAD_API_CUDA
-
-#  ifndef _LIBCUDACXX_HAS_THREAD_API_WIN32
-#    if _CCCL_COMPILER(MSVC) && !defined(_LIBCUDACXX_HAS_THREAD_API_CUDA)
-#      define _LIBCUDACXX_HAS_THREAD_API_WIN32
-#    endif
-#  endif // _LIBCUDACXX_HAS_THREAD_API_WIN32
-
-#  if !defined(_LIBCUDACXX_HAS_THREAD_API_PTHREAD) && !defined(_LIBCUDACXX_HAS_THREAD_API_WIN32) \
-    && !defined(_LIBCUDACXX_HAS_THREAD_API_EXTERNAL)
-#    if defined(__GNU__) || _CCCL_OS(LINUX) || _CCCL_OS(APPLE) || _CCCL_OS(QNX) \
-      || (defined(__MINGW32__) && _CCCL_HAS_INCLUDE(<pthread.h>))
-#      define _LIBCUDACXX_HAS_THREAD_API_PTHREAD
-#    elif defined(_WIN32)
-#      define _LIBCUDACXX_HAS_THREAD_API_WIN32
-#    else
-#      define _LIBCUDACXX_UNSUPPORTED_THREAD_API
-#    endif // _LIBCUDACXX_HAS_THREAD_API
-#  endif
-
-#  if !defined(__STDCPP_THREADS__)
-#    define __STDCPP_THREADS__ 1
-#  endif
-
-// TODO: Support C11 Atomics?
-// #if _CCCL_HAS_FEATURE(cxx_atomic) || __has_extension(c_atomic) || _CCCL_HAS_KEYWORD(_Atomic)
-// #  define _LIBCUDACXX_HAS_C_ATOMIC_IMP
-#  if _CCCL_COMPILER(CLANG)
-#    define _LIBCUDACXX_HAS_GCC_ATOMIC_IMP
-#  elif _CCCL_COMPILER(GCC)
-#    define _LIBCUDACXX_HAS_GCC_ATOMIC_IMP
-#  elif _CCCL_COMPILER(NVHPC)
-#    define _LIBCUDACXX_HAS_GCC_ATOMIC_IMP
-#  elif _CCCL_COMPILER(MSVC)
-#    define _LIBCUDACXX_HAS_MSVC_ATOMIC_IMPL
-#  endif
-
-// Enable bypassing automatic storage checks in atomics when using CTK 12.2 and below and if NDEBUG is defined.
-// A compiler bug prevents the safe use of `__is_local` and PTX spacep until after 13.0.
-#  if _CCCL_CUDACC_BELOW(13, 1) && !defined(NDEBUG)
-#    define _LIBCUDACXX_ATOMIC_UNSAFE_AUTOMATIC_STORAGE
-#  endif // _CCCL_CUDACC_BELOW(13, 1)
-
-// CUDA Atomics supersede host atomics in order to insert the host/device dispatch layer
-#  if _CCCL_HAS_CUDA_COMPILER() || _CCCL_COMPILER(NVHPC)
-#    define _LIBCUDACXX_HAS_CUDA_ATOMIC_IMPL
-#  endif // _CCCL_HAS_CUDA_COMPILER() || _CCCL_COMPILER(NVHPC)
-
-#  if (!defined(_LIBCUDACXX_HAS_C_ATOMIC_IMP) && !defined(_LIBCUDACXX_HAS_GCC_ATOMIC_IMP) \
-       && !_LIBCUDACXX_HAS_EXTERNAL_ATOMIC_IMP())
-#    define _LIBCUDACXX_HAS_NO_ATOMIC_HEADER
-#  else
-#    ifdef __cuda_std__
-#      undef _LIBCUDACXX_ATOMIC_FLAG_TYPE
-#      define _LIBCUDACXX_ATOMIC_FLAG_TYPE int
-#    endif
-#    ifndef _LIBCUDACXX_ATOMIC_FLAG_TYPE
-#      define _LIBCUDACXX_ATOMIC_FLAG_TYPE bool
-#    endif
-#  endif
-
-#  if !defined(_LIBCUDACXX_DISABLE_ADDITIONAL_DIAGNOSTICS)
-#    define _LIBCUDACXX_DIAGNOSE_WARNING(_COND, _MSG) _CCCL_DIAGNOSE_IF(_COND, _MSG, "warning")
-#    define _LIBCUDACXX_DIAGNOSE_ERROR(_COND, _MSG)   _CCCL_DIAGNOSE_IF(_COND, _MSG, "error")
-#  else
-#    define _LIBCUDACXX_DIAGNOSE_WARNING(_COND, _MSG)
-#    define _LIBCUDACXX_DIAGNOSE_ERROR(_COND, _MSG)
-#  endif
-
-#  if defined(_LIBCUDACXX_ENABLE_CXX17_REMOVED_FEATURES)
-#    define _LIBCUDACXX_ENABLE_CXX17_REMOVED_AUTO_PTR
-#    define _LIBCUDACXX_ENABLE_CXX17_REMOVED_UNEXPECTED_FUNCTIONS
-#    define _LIBCUDACXX_ENABLE_CXX17_REMOVED_RANDOM_SHUFFLE
-#    define _LIBCUDACXX_ENABLE_CXX17_REMOVED_BINDERS
-#  endif // _LIBCUDACXX_ENABLE_CXX17_REMOVED_FEATURES
-
-// There are a handful of public standard library types that are intended to
-// support CTAD but don't need any explicit deduction guides to do so. This
-// macro is used to mark them as such, which suppresses the
-// '-Wctad-maybe-unsupported' compiler warning when CTAD is used in user code
-// with these classes.
-#  if (!_CCCL_COMPILER(GCC) || _CCCL_COMPILER(GCC, >, 6))
-#    define _LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(_ClassName) \
-      template <class... _Tag>                              \
-      _ClassName(typename _Tag::__allow_ctad...)->_ClassName<_Tag...>
-#  else
-#    define _LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(_ClassName) static_assert(true, "")
-#  endif
-
-// We can only expose constexpr allocations if the compiler supports it
-// For now disable constexpr allocation support until we can actually use
-#  if 0 && defined(__cpp_constexpr_dynamic_alloc) && defined(__cpp_lib_constexpr_dynamic_alloc) \
-    && _CCCL_STD_VER >= 2020 && !_CCCL_COMPILER(NVRTC)
-#    define _CCCL_HAS_CONSTEXPR_ALLOCATION
-#    define _CCCL_CONSTEXPR_CXX20_ALLOCATION constexpr
-#  else // ^^^ __cpp_constexpr_dynamic_alloc ^^^ / vvv !__cpp_constexpr_dynamic_alloc vvv
-#    define _CCCL_CONSTEXPR_CXX20_ALLOCATION
-#  endif
-
-// NVRTC has a bug that prevented the use of delegated constructors, as it did not accept execution space annotations.
-// This creates a whole lot of boilerplate that we can avoid through a macro (see nvbug3961621)
-#  if _CCCL_COMPILER(NVRTC, <, 12, 6)
-#    define _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__class, __baseclass, ...)                                          \
-      using __base = __baseclass<__VA_ARGS__>;                                                                    \
-      _CCCL_TEMPLATE(class... _Args)                                                                              \
-      _CCCL_REQUIRES(::cuda::std::is_constructible_v<__base, _Args...>)                                           \
-      _CCCL_API constexpr __class(_Args&&... __args) noexcept(                                                    \
-        ::cuda::std::is_nothrow_constructible_v<__base, _Args...>)                                                \
-          : __base(::cuda::std::forward<_Args>(__args)...)                                                        \
-      {}                                                                                                          \
-      _CCCL_HIDE_FROM_ABI constexpr __class() noexcept(::cuda::std::is_nothrow_default_constructible_v<__base>) = \
-        default;
-#  else // ^^^ workaround ^^^ / vvv no workaround vvv
-#    define _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__class, __baseclass, ...)                                          \
-      using __base = __baseclass<__VA_ARGS__>;                                                                    \
-      using __base::__base;                                                                                       \
-      _CCCL_HIDE_FROM_ABI constexpr __class() noexcept(::cuda::std::is_nothrow_default_constructible_v<__base>) = \
-        default;
-#  endif // ^^^ no workaround ^^^
-
-#endif // __cplusplus
-
-#endif // _LIBCUDACXX_CONFIG
diff --git a/libcudacxx/include/cuda/std/inplace_vector b/libcudacxx/include/cuda/std/inplace_vector
index a270ceab45c..39df45ff6e2 100644
--- a/libcudacxx/include/cuda/std/inplace_vector
+++ b/libcudacxx/include/cuda/std/inplace_vector
@@ -63,6 +63,7 @@
 #include <cuda/std/__type_traits/is_trivially_destructible.h>
 #include <cuda/std/__type_traits/is_trivially_move_assignable.h>
 #include <cuda/std/__type_traits/is_trivially_move_constructible.h>
+#include <cuda/std/__utility/delegate_constructors.h>
 #include <cuda/std/__utility/forward.h>
 #include <cuda/std/__utility/move.h>
 #include <cuda/std/cstdint>
@@ -165,7 +166,7 @@ struct __inplace_vector_storage : public __inplace_vector_destruct_base<_Tp, _Ca
   using iterator       = _Tp*;
   using const_iterator = const _Tp*;
 
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__inplace_vector_storage, __inplace_vector_destruct_base, _Tp, _Capacity);
+  _CCCL_DELEGATE_CONSTRUCTORS(__inplace_vector_storage, __inplace_vector_destruct_base, _Tp, _Capacity);
 
   // [containers.sequences.inplace.vector.members] size/capacity
   [[nodiscard]] _CCCL_API constexpr size_type size() const noexcept
@@ -336,7 +337,7 @@ protected:
 template <class _Tp, size_t _Capacity, bool = is_trivially_copy_constructible_v<_Tp>>
 struct __inplace_vector_copy : __inplace_vector_storage<_Tp, _Capacity>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__inplace_vector_copy, __inplace_vector_storage, _Tp, _Capacity);
+  _CCCL_DELEGATE_CONSTRUCTORS(__inplace_vector_copy, __inplace_vector_storage, _Tp, _Capacity);
 
   _CCCL_API inline __inplace_vector_copy(const __inplace_vector_copy& __other) noexcept(
     is_nothrow_copy_constructible_v<_Tp>)
@@ -353,7 +354,7 @@ struct __inplace_vector_copy : __inplace_vector_storage<_Tp, _Capacity>
 template <class _Tp, size_t _Capacity>
 struct __inplace_vector_copy<_Tp, _Capacity, true> : __inplace_vector_storage<_Tp, _Capacity>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__inplace_vector_copy, __inplace_vector_storage, _Tp, _Capacity);
+  _CCCL_DELEGATE_CONSTRUCTORS(__inplace_vector_copy, __inplace_vector_storage, _Tp, _Capacity);
 
   _CCCL_HIDE_FROM_ABI __inplace_vector_copy(const __inplace_vector_copy&)            = default;
   _CCCL_HIDE_FROM_ABI __inplace_vector_copy(__inplace_vector_copy&&)                 = default;
@@ -365,7 +366,7 @@ struct __inplace_vector_copy<_Tp, _Capacity, true> : __inplace_vector_storage<_T
 template <class _Tp, size_t _Capacity, bool = is_trivially_move_constructible_v<_Tp>>
 struct __inplace_vector_move : __inplace_vector_copy<_Tp, _Capacity>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__inplace_vector_move, __inplace_vector_copy, _Tp, _Capacity);
+  _CCCL_DELEGATE_CONSTRUCTORS(__inplace_vector_move, __inplace_vector_copy, _Tp, _Capacity);
 
   _CCCL_HIDE_FROM_ABI __inplace_vector_move(const __inplace_vector_move&) = default;
 
@@ -382,7 +383,7 @@ struct __inplace_vector_move : __inplace_vector_copy<_Tp, _Capacity>
 template <class _Tp, size_t _Capacity>
 struct __inplace_vector_move<_Tp, _Capacity, true> : __inplace_vector_copy<_Tp, _Capacity>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__inplace_vector_move, __inplace_vector_copy, _Tp, _Capacity);
+  _CCCL_DELEGATE_CONSTRUCTORS(__inplace_vector_move, __inplace_vector_copy, _Tp, _Capacity);
 
   _CCCL_HIDE_FROM_ABI __inplace_vector_move(const __inplace_vector_move&)            = default;
   _CCCL_HIDE_FROM_ABI __inplace_vector_move(__inplace_vector_move&&)                 = default;
@@ -397,7 +398,7 @@ template <class _Tp,
           bool = is_trivially_copy_constructible_v<_Tp> && is_trivially_copy_assignable_v<_Tp>>
 struct __inplace_vector_copy_assign : __inplace_vector_move<_Tp, _Capacity>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__inplace_vector_copy_assign, __inplace_vector_move, _Tp, _Capacity);
+  _CCCL_DELEGATE_CONSTRUCTORS(__inplace_vector_copy_assign, __inplace_vector_move, _Tp, _Capacity);
 
   _CCCL_HIDE_FROM_ABI __inplace_vector_copy_assign(const __inplace_vector_copy_assign&) = default;
   _CCCL_HIDE_FROM_ABI __inplace_vector_copy_assign(__inplace_vector_copy_assign&&)      = default;
@@ -423,7 +424,7 @@ struct __inplace_vector_copy_assign : __inplace_vector_move<_Tp, _Capacity>
 template <class _Tp, size_t _Capacity>
 struct __inplace_vector_copy_assign<_Tp, _Capacity, true> : __inplace_vector_move<_Tp, _Capacity>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__inplace_vector_copy_assign, __inplace_vector_move, _Tp, _Capacity);
+  _CCCL_DELEGATE_CONSTRUCTORS(__inplace_vector_copy_assign, __inplace_vector_move, _Tp, _Capacity);
 
   _CCCL_HIDE_FROM_ABI __inplace_vector_copy_assign(const __inplace_vector_copy_assign&)            = default;
   _CCCL_HIDE_FROM_ABI __inplace_vector_copy_assign(__inplace_vector_copy_assign&&)                 = default;
@@ -438,7 +439,7 @@ template <class _Tp,
           bool = is_trivially_move_constructible_v<_Tp> && is_trivially_move_assignable_v<_Tp>>
 struct __inplace_vector_move_assign : __inplace_vector_copy_assign<_Tp, _Capacity>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__inplace_vector_move_assign, __inplace_vector_copy_assign, _Tp, _Capacity);
+  _CCCL_DELEGATE_CONSTRUCTORS(__inplace_vector_move_assign, __inplace_vector_copy_assign, _Tp, _Capacity);
 
   _CCCL_HIDE_FROM_ABI __inplace_vector_move_assign(const __inplace_vector_move_assign&)            = default;
   _CCCL_HIDE_FROM_ABI __inplace_vector_move_assign(__inplace_vector_move_assign&&)                 = default;
@@ -464,7 +465,7 @@ struct __inplace_vector_move_assign : __inplace_vector_copy_assign<_Tp, _Capacit
 template <class _Tp, size_t _Capacity>
 struct __inplace_vector_move_assign<_Tp, _Capacity, true> : __inplace_vector_copy_assign<_Tp, _Capacity>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__inplace_vector_move_assign, __inplace_vector_copy_assign, _Tp, _Capacity);
+  _CCCL_DELEGATE_CONSTRUCTORS(__inplace_vector_move_assign, __inplace_vector_copy_assign, _Tp, _Capacity);
 
   _CCCL_HIDE_FROM_ABI __inplace_vector_move_assign(const __inplace_vector_move_assign&)            = default;
   _CCCL_HIDE_FROM_ABI __inplace_vector_move_assign(__inplace_vector_move_assign&&)                 = default;
@@ -478,7 +479,7 @@ template <class _Tp,
           __inplace_vector_specialization _Spec = __select_inplace_vector_specialization<_Tp, _Capacity>()>
 struct __inplace_vector_base : __inplace_vector_move_assign<_Tp, _Capacity>
 {
-  _LIBCUDACXX_DELEGATE_CONSTRUCTORS(__inplace_vector_base, __inplace_vector_move_assign, _Tp, _Capacity);
+  _CCCL_DELEGATE_CONSTRUCTORS(__inplace_vector_base, __inplace_vector_move_assign, _Tp, _Capacity);
 };
 
 template <class _Tp, size_t _Capacity>
diff --git a/libcudacxx/include/cuda/std/string_view b/libcudacxx/include/cuda/std/string_view
index 8acf7a651e9..a8fdd4415d2 100644
--- a/libcudacxx/include/cuda/std/string_view
+++ b/libcudacxx/include/cuda/std/string_view
@@ -54,6 +54,7 @@
 #include <cuda/std/__type_traits/is_trivially_copyable.h>
 #include <cuda/std/__type_traits/remove_cvref.h>
 #include <cuda/std/__type_traits/type_identity.h>
+#include <cuda/std/__utility/ctad_support.h>
 #include <cuda/std/__utility/declval.h>
 #include <cuda/std/limits>
 #include <cuda/std/version>
@@ -808,7 +809,7 @@ private:
   size_type __size_;
 };
 
-_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(basic_string_view);
+_CCCL_CTAD_SUPPORTED_FOR_TYPE(basic_string_view);
 
 _CCCL_TEMPLATE(class _It, class _End)
 _CCCL_REQUIRES(contiguous_iterator<_It> _CCCL_AND sized_sentinel_for<_End, _It>)
diff --git a/libcudacxx/test/atomic_codegen/CMakeLists.txt b/libcudacxx/test/atomic_codegen/CMakeLists.txt
index 7fcd058e2c2..4da9273a438 100644
--- a/libcudacxx/test/atomic_codegen/CMakeLists.txt
+++ b/libcudacxx/test/atomic_codegen/CMakeLists.txt
@@ -32,7 +32,7 @@ foreach(test_path IN LISTS libcudacxx_atomic_codegen_tests)
         atomic_codegen_${test_name}
         PROPERTIES
           CUDA_ARCHITECTURES "${atomic_codegen_cuda_arch}"
-          COMPILE_DEFINITIONS "_LIBCUDACXX_ATOMIC_UNSAFE_AUTOMATIC_STORAGE=1"
+          COMPILE_DEFINITIONS "_CCCL_ATOMIC_UNSAFE_AUTOMATIC_STORAGE=1"
     )
 
     # Clang stopped emitting PTX in clang20. Add flags to re-enable it.
diff --git a/libcudacxx/test/libcudacxx/cuda/atomics/atomic.local.pass.cpp b/libcudacxx/test/libcudacxx/cuda/atomics/atomic.local.pass.cpp
index 3138e396d80..ee8b0e254eb 100644
--- a/libcudacxx/test/libcudacxx/cuda/atomics/atomic.local.pass.cpp
+++ b/libcudacxx/test/libcudacxx/cuda/atomics/atomic.local.pass.cpp
@@ -120,7 +120,7 @@ __device__ inline void tests()
 
 int main(int arg, char** argv)
 {
-#if !defined(_LIBCUDACXX_ATOMIC_UNSAFE_AUTOMATIC_STORAGE)
+#if !defined(_CCCL_ATOMIC_UNSAFE_AUTOMATIC_STORAGE)
   NV_IF_ELSE_TARGET(
     NV_IS_HOST,
     (cuda_thread_count = 64;),
diff --git a/libcudacxx/test/libcudacxx/libcxx/atomics/version.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/atomics/version.pass.cpp
index 8aed2af7d24..1b063db2dc4 100644
--- a/libcudacxx/test/libcudacxx/libcxx/atomics/version.pass.cpp
+++ b/libcudacxx/test/libcudacxx/libcxx/atomics/version.pass.cpp
@@ -15,8 +15,8 @@
 
 #include "test_macros.h"
 
-#ifndef _LIBCUDACXX_VERSION
-#  error _LIBCUDACXX_VERSION not defined
+#ifndef _CUDA_STD_VERSION
+#  error _CUDA_STD_VERSION not defined
 #endif
 
 int main(int, char**)
diff --git a/libcudacxx/test/libcudacxx/libcxx/utilities/meta/version.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/utilities/meta/version.pass.cpp
index 9a1418f207f..4e133bf1e0c 100644
--- a/libcudacxx/test/libcudacxx/libcxx/utilities/meta/version.pass.cpp
+++ b/libcudacxx/test/libcudacxx/libcxx/utilities/meta/version.pass.cpp
@@ -12,8 +12,8 @@
 
 #include "test_macros.h"
 
-#ifndef _LIBCUDACXX_VERSION
-#  error _LIBCUDACXX_VERSION not defined
+#ifndef _CUDA_STD_VERSION
+#  error _CUDA_STD_VERSION not defined
 #endif
 
 int main(int, char**)
diff --git a/libcudacxx/test/libcudacxx/std/atomics/atomics.lockfree/isalwayslockfree.pass.cpp b/libcudacxx/test/libcudacxx/std/atomics/atomics.lockfree/isalwayslockfree.pass.cpp
index 1eaa902ca34..efd34a67a95 100644
--- a/libcudacxx/test/libcudacxx/std/atomics/atomics.lockfree/isalwayslockfree.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/atomics/atomics.lockfree/isalwayslockfree.pass.cpp
@@ -106,7 +106,7 @@ __host__ __device__ void run()
   CHECK_ALWAYS_LOCK_FREE(float);
   CHECK_ALWAYS_LOCK_FREE(double);
   // CHECK_ALWAYS_LOCK_FREE(long double); // long double is not supported
-#if _CCCL_HAS_ATTRIBUTE(vector_size) && defined(_LIBCUDACXX_VERSION) && !_CCCL_CUDA_COMPILATION()
+#if _CCCL_HAS_ATTRIBUTE(vector_size) && defined(_CUDA_STD_VERSION) && !_CCCL_CUDA_COMPILATION()
   // NOTE: NVCC doesn't support the vector_size attribute in device code.
   CHECK_ALWAYS_LOCK_FREE(int __attribute__((vector_size(1 * sizeof(int)))));
   CHECK_ALWAYS_LOCK_FREE(int __attribute__((vector_size(2 * sizeof(int)))));
@@ -123,7 +123,7 @@ __host__ __device__ void run()
   CHECK_ALWAYS_LOCK_FREE(double __attribute__((vector_size(4 * sizeof(double)))));
   CHECK_ALWAYS_LOCK_FREE(double __attribute__((vector_size(16 * sizeof(double)))));
   CHECK_ALWAYS_LOCK_FREE(double __attribute__((vector_size(32 * sizeof(double)))));
-#endif // _CCCL_HAS_ATTRIBUTE(vector_size) && defined(_LIBCUDACXX_VERSION)
+#endif // _CCCL_HAS_ATTRIBUTE(vector_size) && defined(_CUDA_STD_VERSION)
   CHECK_ALWAYS_LOCK_FREE(struct Empty{});
   CHECK_ALWAYS_LOCK_FREE(struct OneInt { int i; });
   CHECK_ALWAYS_LOCK_FREE(struct IntArr2 { int i[2]; });
diff --git a/libcudacxx/test/libcudacxx/std/thread/thread.barrier/version.pass.cpp b/libcudacxx/test/libcudacxx/std/thread/thread.barrier/version.pass.cpp
index ae7c002403c..010f00f98d8 100644
--- a/libcudacxx/test/libcudacxx/std/thread/thread.barrier/version.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/thread/thread.barrier/version.pass.cpp
@@ -15,8 +15,8 @@
 
 #include "test_macros.h"
 
-#ifndef _LIBCUDACXX_VERSION
-#  error _LIBCUDACXX_VERSION not defined
+#ifndef _CUDA_STD_VERSION
+#  error _CUDA_STD_VERSION not defined
 #endif
 
 int main(int, char**)
diff --git a/libcudacxx/test/libcudacxx/std/thread/thread.latch/version.pass.cpp b/libcudacxx/test/libcudacxx/std/thread/thread.latch/version.pass.cpp
index 34c01b69353..cd156d6a391 100644
--- a/libcudacxx/test/libcudacxx/std/thread/thread.latch/version.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/thread/thread.latch/version.pass.cpp
@@ -15,8 +15,8 @@
 
 #include "test_macros.h"
 
-#ifndef _LIBCUDACXX_VERSION
-#  error _LIBCUDACXX_VERSION not defined
+#ifndef _CUDA_STD_VERSION
+#  error _CUDA_STD_VERSION not defined
 #endif
 
 int main(int, char**)
diff --git a/libcudacxx/test/libcudacxx/std/thread/thread.semaphore/version.pass.cpp b/libcudacxx/test/libcudacxx/std/thread/thread.semaphore/version.pass.cpp
index ccde773de0d..8c86f0c112c 100644
--- a/libcudacxx/test/libcudacxx/std/thread/thread.semaphore/version.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/thread/thread.semaphore/version.pass.cpp
@@ -15,8 +15,8 @@
 
 #include "test_macros.h"
 
-#ifndef _LIBCUDACXX_VERSION
-#  error _LIBCUDACXX_VERSION not defined
+#ifndef _CUDA_STD_VERSION
+#  error _CUDA_STD_VERSION not defined
 #endif
 
 int main(int, char**)
diff --git a/libcudacxx/test/libcudacxx/std/utilities/function.objects/func.not_fn/not_fn.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/function.objects/func.not_fn/not_fn.pass.cpp
index a20c50f36c9..8c951146e15 100644
--- a/libcudacxx/test/libcudacxx/std/utilities/function.objects/func.not_fn/not_fn.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/utilities/function.objects/func.not_fn/not_fn.pass.cpp
@@ -388,11 +388,11 @@ __host__ __device__ void constructor_tests()
     assert(ret() == false);
     auto ret2 = cuda::std::not_fn(value2);
     assert(ret2() == true);
-#if defined(_LIBCUDACXX_VERSION)
+#if defined(_CUDA_STD_VERSION)
     ret = ret2;
     assert(ret() == true);
     assert(ret2() == true);
-#endif // _LIBCUDACXX_VERSION
+#endif // _CUDA_STD_VERSION
   }
   {
     using T = MoveAssignableWrapper;
@@ -407,10 +407,10 @@ __host__ __device__ void constructor_tests()
     assert(ret() == false);
     auto ret2 = cuda::std::not_fn(cuda::std::move(value2));
     assert(ret2() == true);
-#if defined(_LIBCUDACXX_VERSION)
+#if defined(_CUDA_STD_VERSION)
     ret = cuda::std::move(ret2);
     assert(ret() == true);
-#endif // _LIBCUDACXX_VERSION
+#endif // _CUDA_STD_VERSION
   }
 }
 
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.dtor.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.dtor.pass.cpp
index e599063f968..d991340b5ae 100644
--- a/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.dtor.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.dtor.pass.cpp
@@ -31,15 +31,15 @@ int main(int, char**)
 {
   test<int>();
   test<void>();
-#ifdef _LIBCUDACXX_VERSION // extension
+#ifdef _CUDA_STD_VERSION // extension
   test<int const>();
-#endif // _LIBCUDACXX_VERSION
+#endif // _CUDA_STD_VERSION
 
   static_assert(test<int>());
   static_assert(test<void>());
-#ifdef _LIBCUDACXX_VERSION // extension
+#ifdef _CUDA_STD_VERSION // extension
   static_assert(test<int const>());
-#endif // _LIBCUDACXX_VERSION
+#endif // _CUDA_STD_VERSION
 
   return 0;
 }
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator_types.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator_types.pass.cpp
index 70986bda06c..cfd9ab3e226 100644
--- a/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator_types.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator_types.pass.cpp
@@ -64,7 +64,7 @@ __host__ __device__ void test()
 int main(int, char**)
 {
   test<char>();
-#ifdef _LIBCUDACXX_VERSION
+#ifdef _CUDA_STD_VERSION
   test<char const>(); // extension
 #endif
   return 0;
diff --git a/libcudacxx/test/libcudacxx/std/utilities/template.bitset/bitset.members/index_const.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/template.bitset/bitset.members/index_const.pass.cpp
index 6f835dcdaa7..31715a9da78 100644
--- a/libcudacxx/test/libcudacxx/std/utilities/template.bitset/bitset.members/index_const.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/utilities/template.bitset/bitset.members/index_const.pass.cpp
@@ -26,7 +26,7 @@ __host__ __device__ constexpr void test_index_const()
     {
       assert(v[N / 2] == v.test(N / 2));
     }
-#if !defined(_LIBCUDACXX_VERSION) || defined(_LIBCUDACXX_ABI_BITSET_span_BOOL_CONST_SUBSCRIPT_RETURN_BOOL)
+#if !defined(_CUDA_STD_VERSION) || defined(_LIBCUDACXX_ABI_BITSET_span_BOOL_CONST_SUBSCRIPT_RETURN_BOOL)
     static_assert(cuda::std::is_same_v<decltype(v[0]), bool>);
 #else
     static_assert(cuda::std::is_same_v<decltype(v[0]), typename cuda::std::bitset<N>::const_reference>);
@@ -50,7 +50,7 @@ __host__ __device__ constexpr bool test()
   const auto& set = set_;
   auto b          = set[0];
   set_[0]         = true;
-#if !defined(_LIBCUDACXX_VERSION) || defined(_LIBCUDACXX_ABI_BITSET_span_BOOL_CONST_SUBSCRIPT_RETURN_BOOL)
+#if !defined(_CUDA_STD_VERSION) || defined(_LIBCUDACXX_ABI_BITSET_span_BOOL_CONST_SUBSCRIPT_RETURN_BOOL)
   assert(!b);
 #else
   assert(b);
diff --git a/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/PR31384.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/PR31384.pass.cpp
index 60a0a1e97f9..559ed1791e4 100644
--- a/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/PR31384.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/PR31384.pass.cpp
@@ -82,7 +82,7 @@ int main(int, char**)
   count = 0;
   {
     // FIXME: Libc++ incorrectly rejects this code.
-#ifndef _LIBCUDACXX_VERSION
+#ifndef _CUDA_STD_VERSION
     cuda::std::tuple<Implicit> foo = ExplicitDerived<int>{42};
     static_assert(cuda::std::is_convertible<ExplicitDerived<int>, cuda::std::tuple<Implicit>>::value,
                   "correct STLs accept this");
diff --git a/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/UTypes.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/UTypes.pass.cpp
index 1187a339a3f..3366dd964e5 100644
--- a/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/UTypes.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/UTypes.pass.cpp
@@ -68,7 +68,7 @@ __host__ __device__ void test_default_constructible_extension_sfinae()
     static_assert(cuda::std::is_constructible<NestedTuple, MoveOnly, Tuple, MoveOnly, MoveOnly>::value, "");
   }
   // testing extensions
-#ifdef _LIBCUDACXX_VERSION
+#ifdef _CUDA_STD_VERSION
   {
     using Tuple       = cuda::std::tuple<MoveOnly, int>;
     using NestedTuple = cuda::std::tuple<MoveOnly, Tuple, MoveOnly, MoveOnly>;
@@ -98,7 +98,7 @@ int main(int, char**)
     assert(cuda::std::get<2>(t) == 2);
   }
   // extensions, MSVC issues
-#if defined(_LIBCUDACXX_VERSION) && !TEST_COMPILER(MSVC)
+#if defined(_CUDA_STD_VERSION) && !TEST_COMPILER(MSVC)
   {
     using E   = MoveOnly;
     using Tup = cuda::std::tuple<E, E, E>;
diff --git a/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/alloc_copy.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/alloc_copy.pass.cpp
index b3563065153..10de3e5d616 100644
--- a/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/alloc_copy.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/alloc_copy.pass.cpp
@@ -51,7 +51,7 @@ int main(int, char**)
     assert(cuda::std::get<0>(t) == 2);
   }
 // testing extensions
-#ifdef _LIBCUDACXX_VERSION
+#ifdef _CUDA_STD_VERSION
   {
     using T = cuda::std::tuple<alloc_first, alloc_last>;
     T t0(2, 3);
diff --git a/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/alloc_move.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/alloc_move.pass.cpp
index 7cb5a7a9cc5..4fde18d7675 100644
--- a/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/alloc_move.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/alloc_move.pass.cpp
@@ -52,7 +52,7 @@ int main(int, char**)
     assert(cuda::std::get<0>(t) == 1);
   }
 // testing extensions
-#ifdef _LIBCUDACXX_VERSION
+#ifdef _CUDA_STD_VERSION
   {
     using T = cuda::std::tuple<MoveOnly, alloc_first>;
     T t0(0, 1);
diff --git a/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/const_Types.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/const_Types.pass.cpp
index db6de56fded..afc9d6ccc2f 100644
--- a/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/const_Types.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/const_Types.pass.cpp
@@ -152,7 +152,7 @@ int main(int, char**)
   // extensions
   // cuda::std::string not supported
   /*
-#ifdef _LIBCUDACXX_VERSION
+#ifdef _CUDA_STD_VERSION
   {
       cuda::std::tuple<int, char*, cuda::std::string> t(2);
       assert(cuda::std::get<0>(t) == 2);
diff --git a/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/test_lazy_sfinae.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/test_lazy_sfinae.pass.cpp
index 1c83b896152..73e45a18106 100644
--- a/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/test_lazy_sfinae.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.cnstr/test_lazy_sfinae.pass.cpp
@@ -65,7 +65,7 @@ struct AllowAssertT
 // will cause a static assertion.
 __host__ __device__ void test_tuple_like_lazy_sfinae()
 {
-#if defined(_LIBCUDACXX_VERSION)
+#if defined(_CUDA_STD_VERSION)
   // This test requires libc++'s reduced arity initialization.
   using T1 = ConstructibleFromT<cuda::std::pair<int, int>>;
   using T2 = CtorAssertsT<int>;
diff --git a/libcudacxx/test/libcudacxx/std/utilities/variant/variant.variant/variant.swap/swap.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/variant/variant.variant/variant.swap/swap.pass.cpp
index da37dcfceca..b15103a937e 100644
--- a/libcudacxx/test/libcudacxx/std/utilities/variant/variant.variant/variant.swap/swap.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/utilities/variant/variant.variant/variant.swap/swap.pass.cpp
@@ -434,7 +434,7 @@ void test_exceptions_different_alternatives()
     assert(cuda::std::get<1>(v2).value == 100);
   }
 // FIXME: The tests below are just very libc++ specific
-#  ifdef _LIBCUDACXX_VERSION
+#  ifdef _CUDA_STD_VERSION
   {
     using T1 = ThrowsOnSecondMove;
     using T2 = NonThrowingNonNoexceptType;
@@ -466,7 +466,7 @@ void test_exceptions_different_alternatives()
     assert(v1.valueless_by_exception());
     assert(cuda::std::get<0>(v2).value == 42);
   }
-#  endif // _LIBCUDACXX_VERSION
+#  endif // _CUDA_STD_VERSION
 }
 #endif // TEST_HAS_EXCEPTIONS()
 
@@ -585,10 +585,10 @@ __host__ __device__ void test_swap_noexcept()
   }
 }
 
-#ifdef _LIBCUDACXX_VERSION
+#ifdef _CUDA_STD_VERSION
 // This is why variant should SFINAE member swap. :-)
 template class cuda::std::variant<int, NotSwappable>;
-#endif // _LIBCUDACXX_VERSION
+#endif // _CUDA_STD_VERSION
 
 int main(int, char**)
 {
diff --git a/libcudacxx/test/support/check_assertion.h b/libcudacxx/test/support/check_assertion.h
index ca713c1b583..3d933c94693 100644
--- a/libcudacxx/test/support/check_assertion.h
+++ b/libcudacxx/test/support/check_assertion.h
@@ -27,7 +27,7 @@
 #include "test_macros.h"
 #include <sys/wait.h>
 
-#ifndef _LIBCUDACXX_VERSION
+#ifndef _CUDA_STD_VERSION
 #  error "This header may only be used for libc++ tests"
 #endif
 
diff --git a/libcudacxx/test/support/container_debug_tests.h b/libcudacxx/test/support/container_debug_tests.h
index 7eff11f9879..911ffa806d2 100644
--- a/libcudacxx/test/support/container_debug_tests.h
+++ b/libcudacxx/test/support/container_debug_tests.h
@@ -11,7 +11,7 @@
 #define TEST_SUPPORT_CONTAINER_DEBUG_TESTS_H
 
 #include <ciso646>
-#ifndef _LIBCUDACXX_VERSION
+#ifndef _CUDA_STD_VERSION
 #  error This header may only be used for libc++ tests
 #endif
 
diff --git a/libcudacxx/test/support/deduction_guides_sfinae_checks.h b/libcudacxx/test/support/deduction_guides_sfinae_checks.h
index f3c3b5b0d99..ee1d011844c 100644
--- a/libcudacxx/test/support/deduction_guides_sfinae_checks.h
+++ b/libcudacxx/test/support/deduction_guides_sfinae_checks.h
@@ -54,9 +54,9 @@ __host__ __device__ constexpr void SequenceContainerDeductionGuidesSfinaeAway()
   // containers because they have constructors of the form `(size_type count,
   // const value_type& value)`. These constructors would be used when passing
   // two integral types and would deduce `value_type` to be an integral type.
-#ifdef _LIBCUDACXX_VERSION
+#ifdef _CUDA_STD_VERSION
   using OutputIter = cuda::std::insert_iterator<InstantiatedContainer>;
-#endif // _LIBCUDACXX_VERSION
+#endif // _CUDA_STD_VERSION
 
   // (iter, iter)
   //
@@ -97,9 +97,9 @@ __host__ __device__ constexpr void AssociativeContainerDeductionGuidesSfinaeAway
   // The only requirement in the Standard is that integral types cannot be
   // considered input iterators, beyond that it is unspecified.
   using BadIter = int;
-#ifdef _LIBCUDACXX_VERSION
+#ifdef _CUDA_STD_VERSION
   using OutputIter = cuda::std::insert_iterator<InstantiatedContainer>;
-#endif // _LIBCUDACXX_VERSION
+#endif // _CUDA_STD_VERSION
   using AllocAsComp = Alloc;
 
   // (iter, iter)
@@ -171,9 +171,9 @@ __host__ __device__ constexpr void UnorderedContainerDeductionGuidesSfinaeAway()
   // The only requirement in the Standard is that integral types cannot be
   // considered input iterators, beyond that it is unspecified.
   using BadIter = int;
-#ifdef _LIBCUDACXX_VERSION
+#ifdef _CUDA_STD_VERSION
   using OutputIter = cuda::std::insert_iterator<InstantiatedContainer>;
-#endif // _LIBCUDACXX_VERSION
+#endif // _CUDA_STD_VERSION
   using AllocAsHash = Alloc;
   using AllocAsPred = Alloc;
 
diff --git a/libcudacxx/test/support/experimental_any_helpers.h b/libcudacxx/test/support/experimental_any_helpers.h
index e3b0296f49c..f9c77506b3a 100644
--- a/libcudacxx/test/support/experimental_any_helpers.h
+++ b/libcudacxx/test/support/experimental_any_helpers.h
@@ -34,7 +34,7 @@ struct IsSmallObject
 template <class Type>
 bool isSmallType()
 {
-#if defined(_LIBCUDACXX_VERSION)
+#if defined(_CUDA_STD_VERSION)
   return std::experimental::__any_imp::_IsSmallObject<Type>::value;
 #else
   return IsSmallObject<Type>::value;
diff --git a/libcudacxx/test/support/filesystem_include.h b/libcudacxx/test/support/filesystem_include.h
index 0ddd1536932..abb13e08935 100644
--- a/libcudacxx/test/support/filesystem_include.h
+++ b/libcudacxx/test/support/filesystem_include.h
@@ -5,7 +5,7 @@
 
 #include "test_macros.h"
 
-#if defined(_LIBCUDACXX_VERSION)
+#if defined(_CUDA_STD_VERSION)
 namespace fs = std::__fs::filesystem;
 #else
 namespace fs = std::filesystem;
diff --git a/libcudacxx/test/support/msvc_stdlib_force_include.h b/libcudacxx/test/support/msvc_stdlib_force_include.h
index 77f8283df7d..20588549f21 100644
--- a/libcudacxx/test/support/msvc_stdlib_force_include.h
+++ b/libcudacxx/test/support/msvc_stdlib_force_include.h
@@ -24,7 +24,7 @@
 #include <stdlib.h>
 #include <vcruntime.h>
 
-#if defined(_LIBCUDACXX_VERSION)
+#if defined(_CUDA_STD_VERSION)
 #  error This header may not be used when targeting libc++
 #endif
 
diff --git a/libcudacxx/test/utils/libcudacxx/test/config.py b/libcudacxx/test/utils/libcudacxx/test/config.py
index d84ab0f4470..eba7b924b12 100644
--- a/libcudacxx/test/utils/libcudacxx/test/config.py
+++ b/libcudacxx/test/utils/libcudacxx/test/config.py
@@ -1372,7 +1372,6 @@ def configure_debug_mode(self):
             return
         if debug_level not in ["0", "1"]:
             self.lit_config.fatal('Invalid value for debug_level "%s".' % debug_level)
-        self.cxx.compile_flags += ["-D_LIBCUDACXX_DEBUG=%s" % debug_level]
 
     def configure_warnings(self):
         default_enable_warnings = (