diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__config b/libcudacxx/include/cuda/std/detail/libcxx/include/__config
index 54d4dbe628..13124e9653 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/__config
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__config
@@ -1623,23 +1623,24 @@ typedef unsigned int   char32_t;
 #  endif
 #endif // defined(__APPLE__)
 
-#if !defined(_LIBCUDACXX_HAS_NO_ALIGNED_ALLOCATION) && \
-    (defined(_LIBCUDACXX_HAS_NO_LIBRARY_ALIGNED_ALLOCATION) || \
-    (!defined(__cpp_aligned_new) || __cpp_aligned_new < 201606))
-#  define _LIBCUDACXX_HAS_NO_ALIGNED_ALLOCATION
-#endif
+#  if !defined(_LIBCUDACXX_HAS_NO_ALIGNED_ALLOCATION)                     \
+      && (defined(_LIBCUDACXX_HAS_NO_LIBRARY_ALIGNED_ALLOCATION)          \
+          || (!defined(__cpp_aligned_new) || __cpp_aligned_new < 201606)) \
+    || defined(__cuda_std__) // FIXME: Properly handle aligned allocations
+#    define _LIBCUDACXX_HAS_NO_ALIGNED_ALLOCATION
+#  endif
 
-#if defined(__APPLE__) || defined(__FreeBSD__)
-#define _LIBCUDACXX_HAS_DEFAULTRUNELOCALE
-#endif
+#  if defined(__APPLE__) || defined(__FreeBSD__)
+#    define _LIBCUDACXX_HAS_DEFAULTRUNELOCALE
+#  endif
 
-#if defined(__APPLE__) || defined(__FreeBSD__) || defined(__sun__)
-#define _LIBCUDACXX_WCTYPE_IS_MASK
-#endif
+#  if defined(__APPLE__) || defined(__FreeBSD__) || defined(__sun__)
+#    define _LIBCUDACXX_WCTYPE_IS_MASK
+#  endif
 
-#if _CCCL_STD_VER <= 2017 || !defined(__cpp_char8_t)
-#define _LIBCUDACXX_NO_HAS_CHAR8_T
-#endif
+#  if _CCCL_STD_VER <= 2017 || !defined(__cpp_char8_t)
+#    define _LIBCUDACXX_NO_HAS_CHAR8_T
+#  endif
 
 // Deprecation macros.
 //
@@ -2156,6 +2157,15 @@ extern "C" _LIBCUDACXX_FUNC_VIS void __sanitizer_annotate_contiguous_container(
   _CUDA_VSTD::__construct_at(_CUDA_VSTD::addressof(_LOCATION), __VA_ARGS__)
 #endif
 
+// We can only expose constexpr allocations if the compiler supports it
+#  if defined(__cpp_constexpr_dynamic_alloc) && defined(__cpp_lib_constexpr_dynamic_alloc) && _CCCL_STD_VER >= 2020 \
+    && !defined(_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 defined(_CCCL_COMPILER_NVRTC) || (defined(_LIBCUDACXX_CUDACC_BELOW_11_3) && defined(_CCCL_COMPILER_CLANG))
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__functional_base b/libcudacxx/include/cuda/std/detail/libcxx/include/__functional_base
index 76e8e64217..2de26d2ec0 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/__functional_base
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__functional_base
@@ -25,6 +25,8 @@
 #include <cuda/std/detail/libcxx/include/__functional/reference_wrapper.h>
 #include <cuda/std/detail/libcxx/include/__functional/unary_function.h>
 #include <cuda/std/detail/libcxx/include/__functional/weak_result_type.h>
+#include <cuda/std/detail/libcxx/include/__memory/allocator_arg_t.h>
+#include <cuda/std/detail/libcxx/include/__memory/uses_allocator.h>
 #include <cuda/std/detail/libcxx/include/__type_traits/integral_constant.h>
 #include <cuda/std/detail/libcxx/include/__type_traits/is_constructible.h>
 #include <cuda/std/detail/libcxx/include/__type_traits/is_convertible.h>
@@ -37,93 +39,6 @@
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
-// allocator_arg_t
-
-struct _LIBCUDACXX_TEMPLATE_VIS allocator_arg_t { explicit allocator_arg_t() = default; };
-
-#if defined(_LIBCUDACXX_BUILDING_LIBRARY)
-extern _LIBCUDACXX_EXPORTED_FROM_ABI const allocator_arg_t allocator_arg;
-#else
-/* _LIBCUDACXX_INLINE_VAR */ constexpr allocator_arg_t allocator_arg = allocator_arg_t();
-#endif
-
-// uses_allocator
-
-template <class _Tp>
-struct __has_allocator_type
-{
-private:
-    struct __two {char __lx; char __lxx;};
-    template <class _Up> _LIBCUDACXX_INLINE_VISIBILITY static __two __test(...);
-    template <class _Up> _LIBCUDACXX_INLINE_VISIBILITY static char __test(typename _Up::allocator_type* = 0);
-public:
-    static const bool value = sizeof(__test<_Tp>(0)) == 1;
-};
-
-template <class _Tp, class _Alloc, bool = __has_allocator_type<_Tp>::value>
-struct __uses_allocator
-    : public integral_constant<bool,
-        is_convertible<_Alloc, typename _Tp::allocator_type>::value>
-{
-};
-
-template <class _Tp, class _Alloc>
-struct __uses_allocator<_Tp, _Alloc, false>
-    : public false_type
-{
-};
-
-template <class _Tp, class _Alloc>
-struct _LIBCUDACXX_TEMPLATE_VIS uses_allocator
-    : public __uses_allocator<_Tp, _Alloc>
-{
-};
-
-#if _CCCL_STD_VER > 2014
-template <class _Tp, class _Alloc>
-_LIBCUDACXX_INLINE_VAR constexpr size_t uses_allocator_v = uses_allocator<_Tp, _Alloc>::value;
-#endif
-
-// allocator construction
-
-template <class _Tp, class _Alloc, class ..._Args>
-struct __uses_alloc_ctor_imp
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE __remove_cvref_t<_Alloc> _RawAlloc;
-    static const bool __ua = uses_allocator<_Tp, _RawAlloc>::value;
-    static const bool __ic =
-        is_constructible<_Tp, allocator_arg_t, _Alloc, _Args...>::value;
-    static const int value = __ua ? 2 - __ic : 0;
-};
-
-template <class _Tp, class _Alloc, class ..._Args>
-struct __uses_alloc_ctor
-    : integral_constant<int, __uses_alloc_ctor_imp<_Tp, _Alloc, _Args...>::value>
-    {};
-
-template <class _Tp, class _Allocator, class... _Args>
-inline _LIBCUDACXX_INLINE_VISIBILITY
-void __user_alloc_construct_impl (integral_constant<int, 0>, _Tp *__storage, const _Allocator &, _Args &&... __args )
-{
-    new (__storage) _Tp (_CUDA_VSTD::forward<_Args>(__args)...);
-}
-
-// FIXME: This should have a version which takes a non-const alloc.
-template <class _Tp, class _Allocator, class... _Args>
-inline _LIBCUDACXX_INLINE_VISIBILITY
-void __user_alloc_construct_impl (integral_constant<int, 1>, _Tp *__storage, const _Allocator &__a, _Args &&... __args )
-{
-    new (__storage) _Tp (allocator_arg, __a, _CUDA_VSTD::forward<_Args>(__args)...);
-}
-
-// FIXME: This should have a version which takes a non-const alloc.
-template <class _Tp, class _Allocator, class... _Args>
-inline _LIBCUDACXX_INLINE_VISIBILITY
-void __user_alloc_construct_impl (integral_constant<int, 2>, _Tp *__storage, const _Allocator &__a, _Args &&... __args )
-{
-    new (__storage) _Tp (_CUDA_VSTD::forward<_Args>(__args)..., __a);
-}
-
 _LIBCUDACXX_END_NAMESPACE_STD
 
 #include <cuda/std/detail/libcxx/include/__pragma_pop>
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/align.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/align.h
new file mode 100644
index 0000000000..5c669494a0
--- /dev/null
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/align.h
@@ -0,0 +1,55 @@
+// -*- C++ -*-
+//===----------------------------------------------------------------------===//
+//
+// 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___MEMORY_ALIGN_H
+#define _LIBCUDACXX___MEMORY_ALIGN_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/cstddef>
+#include <cuda/std/cstdint>
+
+_CCCL_DIAG_PUSH
+_CCCL_DIAG_SUPPRESS_MSVC(4146) // unary minus operator applied to unsigned type, result still unsigned
+
+_LIBCUDACXX_BEGIN_NAMESPACE_STD
+
+inline _LIBCUDACXX_INLINE_VISIBILITY void* align(size_t __alignment, size_t __size, void*& __ptr, size_t& __space)
+{
+  if (__space < __size)
+  {
+    return nullptr;
+  }
+
+  const auto __intptr  = reinterpret_cast<uintptr_t>(__ptr);
+  const auto __aligned = (__intptr - 1u + __alignment) & -__alignment;
+  const auto __diff    = __aligned - __intptr;
+  if (__diff > (__space - __size))
+  {
+    return nullptr;
+  }
+  __space -= __diff;
+  return __ptr = reinterpret_cast<void*>(__aligned);
+}
+
+_LIBCUDACXX_END_NAMESPACE_STD
+
+_CCCL_DIAG_POP
+
+#endif // _LIBCUDACXX___MEMORY_ALIGN_H
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/allocate_at_least.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/allocate_at_least.h
new file mode 100644
index 0000000000..469978163a
--- /dev/null
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/allocate_at_least.h
@@ -0,0 +1,80 @@
+// -*- C++ -*-
+//===----------------------------------------------------------------------===//
+//
+// 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___MEMORY_ALLOCATE_AT_LEAST_H
+#define _LIBCUDACXX___MEMORY_ALLOCATE_AT_LEAST_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/cstddef>
+#include <cuda/std/detail/libcxx/include/__memory/allocator_traits.h>
+
+_LIBCUDACXX_BEGIN_NAMESPACE_STD
+
+#if _CCCL_STD_VER >= 2023
+template <class _Pointer>
+struct allocation_result
+{
+  _Pointer ptr;
+  size_t count;
+};
+_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(allocation_result);
+
+template <class _Alloc>
+_LIBCUDACXX_NODISCARD_ATTRIBUTE _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr allocation_result<
+  typename allocator_traits<_Alloc>::pointer>
+allocate_at_least(_Alloc& __alloc, size_t __n)
+{
+  if constexpr (requires { __alloc.allocate_at_least(__n); })
+  {
+    return __alloc.allocate_at_least(__n);
+  }
+  else
+  {
+    return {__alloc.allocate(__n), __n};
+  }
+}
+
+template <class _Alloc>
+_LIBCUDACXX_NODISCARD_ATTRIBUTE _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr auto
+__allocate_at_least(_Alloc& __alloc, size_t __n)
+{
+  return _CUDA_VSTD::allocate_at_least(__alloc, __n);
+}
+#else // ^^^ _CCCL_STD_VER >= 2023 ^^^ / vvv _CCCL_STD_VER < 2023 vvv
+template <class _Pointer>
+struct __allocation_result
+{
+  _Pointer ptr;
+  size_t count;
+};
+
+template <class _Alloc>
+_LIBCUDACXX_NODISCARD_ATTRIBUTE _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr __allocation_result<
+  typename allocator_traits<_Alloc>::pointer>
+__allocate_at_least(_Alloc& __alloc, size_t __n)
+{
+  return {__alloc.allocate(__n), __n};
+}
+
+#endif // _CCCL_STD_VER >= 2023
+
+_LIBCUDACXX_END_NAMESPACE_STD
+
+#endif // _LIBCUDACXX___MEMORY_ALLOCATE_AT_LEAST_H
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/allocation_guard.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/allocation_guard.h
new file mode 100644
index 0000000000..3d65c2ce8b
--- /dev/null
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/allocation_guard.h
@@ -0,0 +1,97 @@
+// -*- C++ -*-
+//===----------------------------------------------------------------------===//
+//
+// 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) 2023 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBCUDACXX___MEMORY_ALLOCATION_GUARD_H
+#define _LIBCUDACXX___MEMORY_ALLOCATION_GUARD_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/cstddef>
+#include <cuda/std/detail/libcxx/include/__memory/allocator_traits.h>
+#include <cuda/std/detail/libcxx/include/__utility/move.h>
+
+_LIBCUDACXX_BEGIN_NAMESPACE_STD
+
+// Helper class to allocate memory using an Allocator in an exception safe
+// manner.
+//
+// The intended usage of this class is as follows:
+//
+// 0
+// 1     __allocation_guard<SomeAllocator> guard(alloc, 10);
+// 2     do_some_initialization_that_may_throw(guard.__get());
+// 3     save_allocated_pointer_in_a_noexcept_operation(guard.__release_ptr());
+// 4
+//
+// If line (2) throws an exception during initialization of the memory, the
+// guard's destructor will be called, and the memory will be released using
+// Allocator deallocation. Otherwise, we release the memory from the guard on
+// line (3) in an operation that can't throw -- after that, the guard is not
+// responsible for the memory anymore.
+//
+// This is similar to a unique_ptr, except it's easier to use with a
+// custom allocator.
+
+_CCCL_NV_DIAG_SUPPRESS(2659) // constexpr non-static member function will not be implicitly 'const' in C++14
+
+template <class _Alloc>
+struct __allocation_guard
+{
+  using _Pointer = typename allocator_traits<_Alloc>::pointer;
+  using _Size    = typename allocator_traits<_Alloc>::size_type;
+
+  template <class _AllocT> // we perform the allocator conversion inside the constructor
+  _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 explicit __allocation_guard(
+    _AllocT __alloc, _Size __n)
+      : __alloc_(_CUDA_VSTD::move(__alloc))
+      , __n_(__n)
+      , __ptr_(allocator_traits<_Alloc>::allocate(__alloc_, __n_)) // initialization order is important
+  {}
+
+  _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 ~__allocation_guard() noexcept
+  {
+    if (__ptr_ != nullptr)
+    {
+      allocator_traits<_Alloc>::deallocate(__alloc_, __ptr_, __n_);
+    }
+  }
+
+  _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 _Pointer __release_ptr() noexcept
+  { // not called __release() because it's a keyword in objective-c++
+    _Pointer __tmp = __ptr_;
+    __ptr_         = nullptr;
+    return __tmp;
+  }
+
+  _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr _Pointer __get() const noexcept
+  {
+    return __ptr_;
+  }
+
+private:
+  _Alloc __alloc_;
+  _Size __n_;
+  _Pointer __ptr_;
+};
+
+_CCCL_NV_DIAG_DEFAULT(2659)
+
+_LIBCUDACXX_END_NAMESPACE_STD
+
+#endif // _LIBCUDACXX___MEMORY_ALLOCATION_GUARD_H
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/allocator.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/allocator.h
new file mode 100644
index 0000000000..de0ba61a8f
--- /dev/null
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/allocator.h
@@ -0,0 +1,328 @@
+// -*- C++ -*-
+//===----------------------------------------------------------------------===//
+//
+// 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___MEMORY_ALLOCATOR_H
+#define _LIBCUDACXX___MEMORY_ALLOCATOR_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/cstddef>
+#include <cuda/std/detail/libcxx/include/__memory/addressof.h>
+#include <cuda/std/detail/libcxx/include/__memory/allocate_at_least.h>
+#include <cuda/std/detail/libcxx/include/__memory/allocator_traits.h>
+#include <cuda/std/detail/libcxx/include/__type_traits/is_constant_evaluated.h>
+#include <cuda/std/detail/libcxx/include/__type_traits/is_same.h>
+#include <cuda/std/detail/libcxx/include/__type_traits/is_void.h>
+#include <cuda/std/detail/libcxx/include/__type_traits/is_volatile.h>
+#include <cuda/std/detail/libcxx/include/__utility/forward.h>
+#include <cuda/std/detail/libcxx/include/cstdlib>
+#include <cuda/std/detail/libcxx/include/new>
+
+#if defined(_CCCL_HAS_CONSTEXPR_ALLOCATION) && !defined(_CCCL_COMPILER_NVRTC)
+#  include <memory>
+#endif // _CCCL_HAS_CONSTEXPR_ALLOCATION && !_CCCL_COMPILER_NVRTC
+
+_LIBCUDACXX_BEGIN_NAMESPACE_STD
+
+template <class _Tp>
+class allocator;
+
+#if _CCCL_STD_VER <= 2017
+// These specializations shouldn't be marked _LIBCUDACXX_DEPRECATED_IN_CXX17.
+// Specializing allocator<void> is deprecated, but not using it.
+template <>
+class _LIBCUDACXX_TEMPLATE_VIS allocator<void>
+{
+public:
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 typedef void* pointer;
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 typedef const void* const_pointer;
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 typedef void value_type;
+
+  template <class _Up>
+  struct _LIBCUDACXX_DEPRECATED_IN_CXX17 rebind
+  {
+    typedef allocator<_Up> other;
+  };
+};
+
+template <>
+class _LIBCUDACXX_TEMPLATE_VIS allocator<const void>
+{
+public:
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 typedef const void* pointer;
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 typedef const void* const_pointer;
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 typedef const void value_type;
+
+  template <class _Up>
+  struct _LIBCUDACXX_DEPRECATED_IN_CXX17 rebind
+  {
+    typedef allocator<_Up> other;
+  };
+};
+#endif // _CCCL_STD_VER <= 2017
+
+// This class provides a non-trivial default constructor to the class that derives from it
+// if the condition is satisfied.
+//
+// The second template parameter exists to allow giving a unique type to __non_trivial_if,
+// which makes it possible to avoid breaking the ABI when making this a base class of an
+// existing class. Without that, imagine we have classes D1 and D2, both of which used to
+// have no base classes, but which now derive from __non_trivial_if. The layout of a class
+// that inherits from both D1 and D2 will change because the two __non_trivial_if base
+// classes are not allowed to share the same address.
+//
+// By making those __non_trivial_if base classes unique, we work around this problem and
+// it is safe to start deriving from __non_trivial_if in existing classes.
+template <bool _Cond, class _Unique>
+struct __non_trivial_if
+{};
+
+template <class _Unique>
+struct __non_trivial_if<true, _Unique>
+{
+  _LIBCUDACXX_INLINE_VISIBILITY constexpr __non_trivial_if() noexcept {}
+};
+
+// allocator
+//
+// Note: For ABI compatibility between C++20 and previous standards, we make
+//       allocator<void> trivial in C++20.
+
+template <class _Tp>
+class _LIBCUDACXX_TEMPLATE_VIS allocator : private __non_trivial_if<!_LIBCUDACXX_TRAIT(is_void, _Tp), allocator<_Tp>>
+{
+  static_assert(!_LIBCUDACXX_TRAIT(is_volatile, _Tp), "std::allocator does not support volatile types");
+
+public:
+  typedef size_t size_type;
+  typedef ptrdiff_t difference_type;
+  typedef _Tp value_type;
+  typedef true_type propagate_on_container_move_assignment;
+  typedef true_type is_always_equal;
+
+  _CCCL_CONSTEXPR_CXX20 allocator() noexcept = default;
+
+  template <class _Up>
+  _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 allocator(const allocator<_Up>&) noexcept
+  {}
+
+  _CCCL_EXEC_CHECK_DISABLE
+  _LIBCUDACXX_NODISCARD_AFTER_CXX17 _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20_ALLOCATION _Tp*
+  allocate(size_t __n)
+  {
+    if (__n > allocator_traits<allocator>::max_size(*this))
+    {
+      __throw_bad_array_new_length();
+    }
+#if defined(_CCCL_HAS_CONSTEXPR_ALLOCATION)
+    if (__libcpp_is_constant_evaluated())
+    {
+      return ::std::allocator<_Tp>{}.allocate(__n);
+    }
+#endif // _CCCL_HAS_CONSTEXPR_ALLOCATION
+    {
+      return static_cast<_Tp*>(_CUDA_VSTD::__libcpp_allocate(__n * sizeof(_Tp), _LIBCUDACXX_ALIGNOF(_Tp)));
+    }
+  }
+
+#if _CCCL_STD_VER >= 2023
+  _LIBCUDACXX_NODISCARD_ATTRIBUTE _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr allocation_result<
+    _Tp*>
+  allocate_at_least(size_t __n)
+  {
+    return {allocate(__n), __n};
+  }
+#endif // _CCCL_HAS_CONSTEXPR_ALLOCATION
+
+  _CCCL_EXEC_CHECK_DISABLE
+  _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20_ALLOCATION void deallocate(_Tp* __p, size_t __n) noexcept
+  {
+#if defined(_CCCL_HAS_CONSTEXPR_ALLOCATION)
+    if (__libcpp_is_constant_evaluated())
+    {
+      return ::std::allocator<_Tp>{}.deallocate(__p, __n);
+    }
+    else
+#endif // _CCCL_STD_VER >= 2020
+    {
+      _CUDA_VSTD::__libcpp_deallocate((void*) __p, __n * sizeof(_Tp), _LIBCUDACXX_ALIGNOF(_Tp));
+    }
+  }
+
+  // C++20 Removed members
+#if _CCCL_STD_VER <= 2017
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 typedef _Tp* pointer;
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 typedef const _Tp* const_pointer;
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 typedef _Tp& reference;
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 typedef const _Tp& const_reference;
+
+  template <class _Up>
+  struct _LIBCUDACXX_DEPRECATED_IN_CXX17 rebind
+  {
+    typedef allocator<_Up> other;
+  };
+
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 _LIBCUDACXX_INLINE_VISIBILITY pointer address(reference __x) const noexcept
+  {
+    return _CUDA_VSTD::addressof(__x);
+  }
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 _LIBCUDACXX_INLINE_VISIBILITY const_pointer address(const_reference __x) const noexcept
+  {
+    return _CUDA_VSTD::addressof(__x);
+  }
+
+  _LIBCUDACXX_NODISCARD_AFTER_CXX17 _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_DEPRECATED_IN_CXX17 _Tp*
+  allocate(size_t __n, const void*)
+  {
+    return allocate(__n);
+  }
+
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 _LIBCUDACXX_INLINE_VISIBILITY size_type max_size() const noexcept
+  {
+    return size_type(~0) / sizeof(_Tp);
+  }
+
+  template <class _Up, class... _Args>
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 _LIBCUDACXX_INLINE_VISIBILITY void construct(_Up* __p, _Args&&... __args)
+  {
+    ::new ((void*) __p) _Up(_CUDA_VSTD::forward<_Args>(__args)...);
+  }
+
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 _LIBCUDACXX_INLINE_VISIBILITY void destroy(pointer __p)
+  {
+    __p->~_Tp();
+  }
+#endif // _CCCL_STD_VER <= 2017
+};
+
+template <class _Tp>
+class _LIBCUDACXX_TEMPLATE_VIS allocator<const _Tp>
+    : private __non_trivial_if<!_LIBCUDACXX_TRAIT(is_void, _Tp), allocator<const _Tp>>
+{
+  static_assert(!_LIBCUDACXX_TRAIT(is_volatile, _Tp), "std::allocator does not support volatile types");
+
+public:
+  typedef size_t size_type;
+  typedef ptrdiff_t difference_type;
+  typedef const _Tp value_type;
+  typedef true_type propagate_on_container_move_assignment;
+  typedef true_type is_always_equal;
+
+  _CCCL_CONSTEXPR_CXX20 allocator() noexcept = default;
+
+  template <class _Up>
+  _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 allocator(const allocator<_Up>&) noexcept
+  {}
+
+  _LIBCUDACXX_NODISCARD_AFTER_CXX17 _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 const _Tp* allocate(size_t __n)
+  {
+    if (__n > allocator_traits<allocator>::max_size(*this))
+    {
+      __throw_bad_array_new_length();
+    }
+    if (__libcpp_is_constant_evaluated())
+    {
+      return static_cast<const _Tp*>(::operator new(__n * sizeof(_Tp)));
+    }
+    else
+    {
+      return static_cast<const _Tp*>(_CUDA_VSTD::__libcpp_allocate(__n * sizeof(_Tp), _LIBCUDACXX_ALIGNOF(_Tp)));
+    }
+  }
+
+#if _CCCL_STD_VER >= 2023
+  _LIBCUDACXX_NODISCARD_ATTRIBUTE _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr allocation_result<
+    const _Tp*>
+  allocate_at_least(size_t __n)
+  {
+    return {allocate(__n), __n};
+  }
+#endif // _CCCL_STD_VER >= 2023
+
+  _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 void deallocate(const _Tp* __p, size_t __n)
+  {
+    if (__libcpp_is_constant_evaluated())
+    {
+      ::operator delete(const_cast<_Tp*>(__p));
+    }
+    else
+    {
+      _CUDA_VSTD::__libcpp_deallocate((void*) const_cast<_Tp*>(__p), __n * sizeof(_Tp), _LIBCUDACXX_ALIGNOF(_Tp));
+    }
+  }
+
+  // C++20 Removed members
+#if _CCCL_STD_VER <= 2017
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 typedef const _Tp* pointer;
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 typedef const _Tp* const_pointer;
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 typedef const _Tp& reference;
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 typedef const _Tp& const_reference;
+
+  template <class _Up>
+  struct _LIBCUDACXX_DEPRECATED_IN_CXX17 rebind
+  {
+    typedef allocator<_Up> other;
+  };
+
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 _LIBCUDACXX_INLINE_VISIBILITY const_pointer address(const_reference __x) const noexcept
+  {
+    return _CUDA_VSTD::addressof(__x);
+  }
+
+  _LIBCUDACXX_NODISCARD_AFTER_CXX17 _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_DEPRECATED_IN_CXX17 const _Tp*
+  allocate(size_t __n, const void*)
+  {
+    return allocate(__n);
+  }
+
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 _LIBCUDACXX_INLINE_VISIBILITY size_type max_size() const noexcept
+  {
+    return size_type(~0) / sizeof(_Tp);
+  }
+
+  template <class _Up, class... _Args>
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 _LIBCUDACXX_INLINE_VISIBILITY void construct(_Up* __p, _Args&&... __args)
+  {
+    ::new ((void*) __p) _Up(_CUDA_VSTD::forward<_Args>(__args)...);
+  }
+
+  _LIBCUDACXX_DEPRECATED_IN_CXX17 _LIBCUDACXX_INLINE_VISIBILITY void destroy(pointer __p)
+  {
+    __p->~_Tp();
+  }
+#endif // _CCCL_STD_VER <= 2017
+};
+
+template <class _Tp, class _Up>
+inline _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 bool
+operator==(const allocator<_Tp>&, const allocator<_Up>&) noexcept
+{
+  return true;
+}
+
+template <class _Tp, class _Up>
+inline _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 bool
+operator!=(const allocator<_Tp>&, const allocator<_Up>&) noexcept
+{
+  return false;
+}
+
+_LIBCUDACXX_END_NAMESPACE_STD
+
+#endif // _LIBCUDACXX___MEMORY_ALLOCATOR_H
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/allocator_arg_t.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/allocator_arg_t.h
new file mode 100644
index 0000000000..277ab32ea9
--- /dev/null
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/allocator_arg_t.h
@@ -0,0 +1,80 @@
+// -*- C++ -*-
+//===----------------------------------------------------------------------===//
+//
+// 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___FUNCTIONAL_ALLOCATOR_ARG_T_H
+#define _LIBCUDACXX___FUNCTIONAL_ALLOCATOR_ARG_T_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/detail/libcxx/include/__memory/uses_allocator.h>
+#include <cuda/std/detail/libcxx/include/__type_traits/integral_constant.h>
+#include <cuda/std/detail/libcxx/include/__type_traits/is_constructible.h>
+#include <cuda/std/detail/libcxx/include/__type_traits/remove_cvref.h>
+#include <cuda/std/detail/libcxx/include/__utility/forward.h>
+
+_LIBCUDACXX_BEGIN_NAMESPACE_STD
+
+struct _LIBCUDACXX_TEMPLATE_VIS allocator_arg_t
+{
+  explicit allocator_arg_t() = default;
+};
+
+_LIBCUDACXX_INLINE_VAR constexpr allocator_arg_t allocator_arg = allocator_arg_t();
+
+// allocator construction
+
+template <class _Tp, class _Alloc, class... _Args>
+struct __uses_alloc_ctor_imp
+{
+  typedef _LIBCUDACXX_NODEBUG __remove_cvref_t<_Alloc> _RawAlloc;
+  static const bool __ua = uses_allocator<_Tp, _RawAlloc>::value;
+  static const bool __ic = is_constructible<_Tp, allocator_arg_t, _Alloc, _Args...>::value;
+  static const int value = __ua ? 2 - __ic : 0;
+};
+
+template <class _Tp, class _Alloc, class... _Args>
+struct __uses_alloc_ctor : integral_constant<int, __uses_alloc_ctor_imp<_Tp, _Alloc, _Args...>::value>
+{};
+
+template <class _Tp, class _Allocator, class... _Args>
+inline _LIBCUDACXX_INLINE_VISIBILITY void
+__user_alloc_construct_impl(integral_constant<int, 0>, _Tp* __storage, const _Allocator&, _Args&&... __args)
+{
+  new (__storage) _Tp(_CUDA_VSTD::forward<_Args>(__args)...);
+}
+
+// FIXME: This should have a version which takes a non-const alloc.
+template <class _Tp, class _Allocator, class... _Args>
+inline _LIBCUDACXX_INLINE_VISIBILITY void
+__user_alloc_construct_impl(integral_constant<int, 1>, _Tp* __storage, const _Allocator& __a, _Args&&... __args)
+{
+  new (__storage) _Tp(allocator_arg, __a, _CUDA_VSTD::forward<_Args>(__args)...);
+}
+
+// FIXME: This should have a version which takes a non-const alloc.
+template <class _Tp, class _Allocator, class... _Args>
+inline _LIBCUDACXX_INLINE_VISIBILITY void
+__user_alloc_construct_impl(integral_constant<int, 2>, _Tp* __storage, const _Allocator& __a, _Args&&... __args)
+{
+  new (__storage) _Tp(_CUDA_VSTD::forward<_Args>(__args)..., __a);
+}
+
+_LIBCUDACXX_END_NAMESPACE_STD
+
+#endif // _LIBCUDACXX___FUNCTIONAL_ALLOCATOR_ARG_T_H
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/allocator_destructor.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/allocator_destructor.h
new file mode 100644
index 0000000000..dfd3bc8743
--- /dev/null
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/allocator_destructor.h
@@ -0,0 +1,55 @@
+// -*- C++ -*-
+//===----------------------------------------------------------------------===//
+//
+// 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___MEMORY_ALLOCATOR_DESTRUCTOR_H
+#define _LIBCUDACXX___MEMORY_ALLOCATOR_DESTRUCTOR_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/detail/libcxx/include/__memory/allocator_traits.h>
+
+_LIBCUDACXX_BEGIN_NAMESPACE_STD
+
+template <class _Alloc>
+class __allocator_destructor
+{
+  typedef _LIBCUDACXX_NODEBUG_TYPE allocator_traits<_Alloc> __alloc_traits;
+
+public:
+  typedef _LIBCUDACXX_NODEBUG_TYPE typename __alloc_traits::pointer pointer;
+  typedef _LIBCUDACXX_NODEBUG_TYPE typename __alloc_traits::size_type size_type;
+
+private:
+  _Alloc& __alloc_;
+  size_type __s_;
+
+public:
+  _LIBCUDACXX_INLINE_VISIBILITY __allocator_destructor(_Alloc& __a, size_type __s) noexcept
+      : __alloc_(__a)
+      , __s_(__s)
+  {}
+  _LIBCUDACXX_INLINE_VISIBILITY void operator()(pointer __p) noexcept
+  {
+    __alloc_traits::deallocate(__alloc_, __p, __s_);
+  }
+};
+
+_LIBCUDACXX_END_NAMESPACE_STD
+
+#endif // _LIBCUDACXX___MEMORY_ALLOCATOR_DESTRUCTOR_H
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/allocator_traits.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/allocator_traits.h
new file mode 100644
index 0000000000..df56f4d8ca
--- /dev/null
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/allocator_traits.h
@@ -0,0 +1,604 @@
+// -*- C++ -*-
+//===----------------------------------------------------------------------===//
+//
+// 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) 2023 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBCUDACXX___MEMORY_ALLOCATOR_TRAITS_H
+#define _LIBCUDACXX___MEMORY_ALLOCATOR_TRAITS_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/detail/__pragma_push>
+#include <cuda/std/detail/libcxx/include/__memory/construct_at.h>
+#include <cuda/std/detail/libcxx/include/__memory/pointer_traits.h>
+#include <cuda/std/detail/libcxx/include/__type_traits/enable_if.h>
+#include <cuda/std/detail/libcxx/include/__type_traits/is_copy_constructible.h>
+#include <cuda/std/detail/libcxx/include/__type_traits/is_empty.h>
+#include <cuda/std/detail/libcxx/include/__type_traits/is_move_constructible.h>
+#include <cuda/std/detail/libcxx/include/__type_traits/make_unsigned.h>
+#include <cuda/std/detail/libcxx/include/__type_traits/remove_reference.h>
+#include <cuda/std/detail/libcxx/include/__type_traits/void_t.h>
+#include <cuda/std/detail/libcxx/include/__utility/declval.h>
+#include <cuda/std/detail/libcxx/include/__utility/forward.h>
+#include <cuda/std/detail/libcxx/include/cstring>
+#include <cuda/std/limits>
+
+_CCCL_NV_DIAG_SUPPRESS(1215)
+
+_LIBCUDACXX_BEGIN_NAMESPACE_STD
+
+#if _CCCL_STD_VER <= 2014
+#  define _LIBCUDACXX_ALLOCATOR_TRAITS_HAS_XXX(NAME, PROPERTY)     \
+    template <class _Tp, class = void>                             \
+    struct NAME : false_type                                       \
+    {};                                                            \
+    template <class _Tp>                                           \
+    struct NAME<_Tp, __void_t<typename _Tp::PROPERTY>> : true_type \
+    {}
+
+#else // ^^^ _CCCL_STD_VER <= 2014 ^^^ / vvv _CCCL_STD_VER >= 2017 vvv
+#  define _LIBCUDACXX_ALLOCATOR_TRAITS_HAS_XXX(NAME, PROPERTY) \
+    template <class _Tp, class = void>                         \
+    inline constexpr bool NAME##_v = false;                    \
+    template <class _Tp>                                       \
+    inline constexpr bool NAME##_v<_Tp, __void_t<typename _Tp::PROPERTY>> = true;
+#endif // _CCCL_STD_VER >= 2017
+
+// __pointer
+_LIBCUDACXX_ALLOCATOR_TRAITS_HAS_XXX(__has_pointer, pointer);
+template <class _Tp,
+          class _Alloc,
+          class _RawAlloc = __libcpp_remove_reference_t<_Alloc>,
+          bool            = _LIBCUDACXX_TRAIT(__has_pointer, _RawAlloc)>
+struct __pointer
+{
+  using type _LIBCUDACXX_NODEBUG_TYPE = typename _RawAlloc::pointer;
+};
+template <class _Tp, class _Alloc, class _RawAlloc>
+struct __pointer<_Tp, _Alloc, _RawAlloc, false>
+{
+  using type _LIBCUDACXX_NODEBUG_TYPE = _Tp*;
+};
+
+// __const_pointer
+_LIBCUDACXX_ALLOCATOR_TRAITS_HAS_XXX(__has_const_pointer, const_pointer);
+template <class _Tp, class _Ptr, class _Alloc, bool = _LIBCUDACXX_TRAIT(__has_const_pointer, _Alloc)>
+struct __const_pointer
+{
+  using type _LIBCUDACXX_NODEBUG_TYPE = typename _Alloc::const_pointer;
+};
+template <class _Tp, class _Ptr, class _Alloc>
+struct __const_pointer<_Tp, _Ptr, _Alloc, false>
+{
+  using type _LIBCUDACXX_NODEBUG_TYPE = typename pointer_traits<_Ptr>::template rebind<const _Tp>;
+};
+
+// __void_pointer
+_LIBCUDACXX_ALLOCATOR_TRAITS_HAS_XXX(__has_void_pointer, void_pointer);
+template <class _Ptr, class _Alloc, bool = _LIBCUDACXX_TRAIT(__has_void_pointer, _Alloc)>
+struct __void_pointer
+{
+  using type _LIBCUDACXX_NODEBUG_TYPE = typename _Alloc::void_pointer;
+};
+template <class _Ptr, class _Alloc>
+struct __void_pointer<_Ptr, _Alloc, false>
+{
+  using type _LIBCUDACXX_NODEBUG_TYPE = typename pointer_traits<_Ptr>::template rebind<void>;
+};
+
+// __const_void_pointer
+_LIBCUDACXX_ALLOCATOR_TRAITS_HAS_XXX(__has_const_void_pointer, const_void_pointer);
+template <class _Ptr, class _Alloc, bool = _LIBCUDACXX_TRAIT(__has_const_void_pointer, _Alloc)>
+struct __const_void_pointer
+{
+  using type _LIBCUDACXX_NODEBUG_TYPE = typename _Alloc::const_void_pointer;
+};
+template <class _Ptr, class _Alloc>
+struct __const_void_pointer<_Ptr, _Alloc, false>
+{
+  using type _LIBCUDACXX_NODEBUG_TYPE = typename pointer_traits<_Ptr>::template rebind<const void>;
+};
+
+// __size_type
+_LIBCUDACXX_ALLOCATOR_TRAITS_HAS_XXX(__has_size_type, size_type);
+template <class _Alloc, class _DiffType, bool = _LIBCUDACXX_TRAIT(__has_size_type, _Alloc)>
+struct __size_type : make_unsigned<_DiffType>
+{};
+template <class _Alloc, class _DiffType>
+struct __size_type<_Alloc, _DiffType, true>
+{
+  using type _LIBCUDACXX_NODEBUG_TYPE = typename _Alloc::size_type;
+};
+
+// __alloc_traits_difference_type
+_LIBCUDACXX_ALLOCATOR_TRAITS_HAS_XXX(__has_alloc_traits_difference_type, difference_type);
+template <class _Alloc, class _Ptr, bool = _LIBCUDACXX_TRAIT(__has_alloc_traits_difference_type, _Alloc)>
+struct __alloc_traits_difference_type
+{
+  using type _LIBCUDACXX_NODEBUG_TYPE = typename pointer_traits<_Ptr>::difference_type;
+};
+template <class _Alloc, class _Ptr>
+struct __alloc_traits_difference_type<_Alloc, _Ptr, true>
+{
+  using type _LIBCUDACXX_NODEBUG_TYPE = typename _Alloc::difference_type;
+};
+
+// __propagate_on_container_copy_assignment
+_LIBCUDACXX_ALLOCATOR_TRAITS_HAS_XXX(__has_propagate_on_container_copy_assignment,
+                                     propagate_on_container_copy_assignment);
+template <class _Alloc, bool = _LIBCUDACXX_TRAIT(__has_propagate_on_container_copy_assignment, _Alloc)>
+struct __propagate_on_container_copy_assignment : false_type
+{};
+template <class _Alloc>
+struct __propagate_on_container_copy_assignment<_Alloc, true>
+{
+  using type _LIBCUDACXX_NODEBUG_TYPE = typename _Alloc::propagate_on_container_copy_assignment;
+};
+
+// __propagate_on_container_move_assignment
+_LIBCUDACXX_ALLOCATOR_TRAITS_HAS_XXX(__has_propagate_on_container_move_assignment,
+                                     propagate_on_container_move_assignment);
+template <class _Alloc, bool = _LIBCUDACXX_TRAIT(__has_propagate_on_container_move_assignment, _Alloc)>
+struct __propagate_on_container_move_assignment : false_type
+{};
+template <class _Alloc>
+struct __propagate_on_container_move_assignment<_Alloc, true>
+{
+  using type _LIBCUDACXX_NODEBUG_TYPE = typename _Alloc::propagate_on_container_move_assignment;
+};
+
+// __propagate_on_container_swap
+_LIBCUDACXX_ALLOCATOR_TRAITS_HAS_XXX(__has_propagate_on_container_swap, propagate_on_container_swap);
+template <class _Alloc, bool = _LIBCUDACXX_TRAIT(__has_propagate_on_container_swap, _Alloc)>
+struct __propagate_on_container_swap : false_type
+{};
+template <class _Alloc>
+struct __propagate_on_container_swap<_Alloc, true>
+{
+  using type _LIBCUDACXX_NODEBUG_TYPE = typename _Alloc::propagate_on_container_swap;
+};
+
+// __is_always_equal
+_LIBCUDACXX_ALLOCATOR_TRAITS_HAS_XXX(__has_is_always_equal, is_always_equal);
+template <class _Alloc, bool = _LIBCUDACXX_TRAIT(__has_is_always_equal, _Alloc)>
+struct __is_always_equal : is_empty<_Alloc>
+{};
+template <class _Alloc>
+struct __is_always_equal<_Alloc, true>
+{
+  using type _LIBCUDACXX_NODEBUG_TYPE = typename _Alloc::is_always_equal;
+};
+
+// __allocator_traits_rebind
+_CCCL_SUPPRESS_DEPRECATED_PUSH
+template <class _Tp, class _Up, class = void>
+struct __has_rebind_other : false_type
+{};
+template <class _Tp, class _Up>
+struct __has_rebind_other<_Tp, _Up, __void_t<typename _Tp::template rebind<_Up>::other>> : true_type
+{};
+
+template <class _Tp, class _Up, bool = __has_rebind_other<_Tp, _Up>::value>
+struct __allocator_traits_rebind
+{
+  static_assert(__has_rebind_other<_Tp, _Up>::value, "This allocator has to implement rebind");
+  using type _LIBCUDACXX_NODEBUG_TYPE = typename _Tp::template rebind<_Up>::other;
+};
+template <template <class, class...> class _Alloc, class _Tp, class... _Args, class _Up>
+struct __allocator_traits_rebind<_Alloc<_Tp, _Args...>, _Up, true>
+{
+  using type _LIBCUDACXX_NODEBUG_TYPE = typename _Alloc<_Tp, _Args...>::template rebind<_Up>::other;
+};
+template <template <class, class...> class _Alloc, class _Tp, class... _Args, class _Up>
+struct __allocator_traits_rebind<_Alloc<_Tp, _Args...>, _Up, false>
+{
+  using type _LIBCUDACXX_NODEBUG_TYPE = _Alloc<_Up, _Args...>;
+};
+_CCCL_SUPPRESS_DEPRECATED_POP
+
+template <class _Alloc, class _Tp>
+using __allocator_traits_rebind_t = typename __allocator_traits_rebind<_Alloc, _Tp>::type;
+
+_CCCL_SUPPRESS_DEPRECATED_PUSH
+
+// __has_allocate_hint
+template <class _Alloc, class _SizeType, class _ConstVoidPtr, class = void>
+struct __has_allocate_hint : false_type
+{};
+
+template <class _Alloc, class _SizeType, class _ConstVoidPtr>
+struct __has_allocate_hint<_Alloc,
+                           _SizeType,
+                           _ConstVoidPtr,
+                           decltype((void) _CUDA_VSTD::declval<_Alloc>().allocate(
+                             _CUDA_VSTD::declval<_SizeType>(), _CUDA_VSTD::declval<_ConstVoidPtr>()))> : true_type
+{};
+
+// __has_construct
+template <class, class _Alloc, class... _Args>
+struct __has_construct_impl : false_type
+{};
+
+template <class _Alloc, class... _Args>
+struct __has_construct_impl<decltype((void) _CUDA_VSTD::declval<_Alloc>().construct(_CUDA_VSTD::declval<_Args>()...)),
+                            _Alloc,
+                            _Args...> : true_type
+{};
+
+template <class _Alloc, class... _Args>
+struct __has_construct : __has_construct_impl<void, _Alloc, _Args...>
+{};
+
+// __has_destroy
+template <class _Alloc, class _Pointer, class = void>
+struct __has_destroy : false_type
+{};
+
+template <class _Alloc, class _Pointer>
+struct __has_destroy<_Alloc,
+                     _Pointer,
+                     decltype((void) _CUDA_VSTD::declval<_Alloc>().destroy(_CUDA_VSTD::declval<_Pointer>()))>
+    : true_type
+{};
+
+// __has_max_size
+template <class _Alloc, class = void>
+struct __has_max_size : false_type
+{};
+
+template <class _Alloc>
+struct __has_max_size<_Alloc, decltype((void) _CUDA_VSTD::declval<_Alloc&>().max_size())> : true_type
+{};
+
+// __has_select_on_container_copy_construction
+template <class _Alloc, class = void>
+struct __has_select_on_container_copy_construction : false_type
+{};
+
+template <class _Alloc>
+struct __has_select_on_container_copy_construction<
+  _Alloc,
+  decltype((void) _CUDA_VSTD::declval<_Alloc>().select_on_container_copy_construction())> : true_type
+{};
+
+template <class _Tp>
+inline _LIBCUDACXX_INLINE_VISIBILITY constexpr _Tp* __to_raw_pointer(_Tp* __p) noexcept
+{
+  return __p;
+}
+
+#if _CCCL_STD_VER <= 2017
+template <class _Pointer>
+inline _LIBCUDACXX_INLINE_VISIBILITY typename pointer_traits<_Pointer>::element_type*
+__to_raw_pointer(_Pointer __p) noexcept
+{
+  return _CUDA_VSTD::__to_raw_pointer(__p.operator->());
+}
+#else // ^^^ C++17 ^^^ / vvv C++20 vvv
+template <class _Pointer>
+inline _LIBCUDACXX_INLINE_VISIBILITY auto __to_raw_pointer(const _Pointer& __p) noexcept
+  -> decltype(pointer_traits<_Pointer>::to_address(__p))
+{
+  return pointer_traits<_Pointer>::to_address(__p);
+}
+
+template <class _Pointer, class... _None>
+inline _LIBCUDACXX_INLINE_VISIBILITY auto __to_raw_pointer(const _Pointer& __p, _None...) noexcept
+{
+  return _CUDA_VSTD::__to_raw_pointer(__p.operator->());
+}
+#endif // _CCCL_STD_VER >= 2020
+
+_CCCL_SUPPRESS_DEPRECATED_POP
+
+// __is_default_allocator
+template <class _Tp>
+struct __is_default_allocator : false_type
+{};
+
+template <class>
+class allocator;
+
+template <class _Tp>
+struct __is_default_allocator<allocator<_Tp>> : true_type
+{};
+// __is_cpp17_move_insertable
+template <class _Alloc, class = void>
+struct __is_cpp17_move_insertable : is_move_constructible<typename _Alloc::value_type>
+{};
+
+template <class _Alloc>
+struct __is_cpp17_move_insertable<
+  _Alloc,
+  __enable_if_t<!__is_default_allocator<_Alloc>::value
+                && __has_construct<_Alloc, typename _Alloc::value_type*, typename _Alloc::value_type&&>::value>>
+    : true_type
+{};
+
+// __is_cpp17_copy_insertable
+template <class _Alloc, class = void>
+struct __is_cpp17_copy_insertable
+    : integral_constant<bool,
+                        is_copy_constructible<typename _Alloc::value_type>::value
+                          && __is_cpp17_move_insertable<_Alloc>::value>
+{};
+
+template <class _Alloc>
+struct __is_cpp17_copy_insertable<
+  _Alloc,
+  __enable_if_t<!__is_default_allocator<_Alloc>::value
+                && __has_construct<_Alloc, typename _Alloc::value_type*, const typename _Alloc::value_type&>::value>>
+    : __is_cpp17_move_insertable<_Alloc>
+{};
+
+template <class _Alloc>
+struct _LIBCUDACXX_TEMPLATE_VIS allocator_traits
+{
+  using allocator_type     = _Alloc;
+  using value_type         = typename allocator_type::value_type;
+  using pointer            = typename __pointer<value_type, allocator_type>::type;
+  using const_pointer      = typename __const_pointer<value_type, pointer, allocator_type>::type;
+  using void_pointer       = typename __void_pointer<pointer, allocator_type>::type;
+  using const_void_pointer = typename __const_void_pointer<pointer, allocator_type>::type;
+  using difference_type    = typename __alloc_traits_difference_type<allocator_type, pointer>::type;
+  using size_type          = typename __size_type<allocator_type, difference_type>::type;
+  using propagate_on_container_copy_assignment =
+    typename __propagate_on_container_copy_assignment<allocator_type>::type;
+  using propagate_on_container_move_assignment =
+    typename __propagate_on_container_move_assignment<allocator_type>::type;
+  using propagate_on_container_swap = typename __propagate_on_container_swap<allocator_type>::type;
+  using is_always_equal             = typename __is_always_equal<allocator_type>::type;
+
+  template <class _Tp>
+  using rebind_alloc = __allocator_traits_rebind_t<allocator_type, _Tp>;
+  template <class _Tp>
+  using rebind_traits = allocator_traits<rebind_alloc<_Tp>>;
+
+  _LIBCUDACXX_NODISCARD_AFTER_CXX17 _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 static pointer
+  allocate(allocator_type& __a, size_type __n)
+  {
+    return __a.allocate(__n);
+  }
+
+  template <class _Ap = _Alloc, __enable_if_t<__has_allocate_hint<_Ap, size_type, const_void_pointer>::value, int> = 0>
+  _LIBCUDACXX_NODISCARD_AFTER_CXX17 _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 static pointer
+  allocate(allocator_type& __a, size_type __n, const_void_pointer __hint)
+  {
+    _CCCL_SUPPRESS_DEPRECATED_PUSH
+    return __a.allocate(__n, __hint);
+    _CCCL_SUPPRESS_DEPRECATED_POP
+  }
+  template <class _Ap                                                                           = _Alloc,
+            class                                                                               = void,
+            __enable_if_t<!__has_allocate_hint<_Ap, size_type, const_void_pointer>::value, int> = 0>
+  _LIBCUDACXX_NODISCARD_AFTER_CXX17 _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 static pointer
+  allocate(allocator_type& __a, size_type __n, const_void_pointer)
+  {
+    return __a.allocate(__n);
+  }
+
+  _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 static void
+  deallocate(allocator_type& __a, pointer __p, size_type __n) noexcept
+  {
+    __a.deallocate(__p, __n);
+  }
+
+  template <class _Tp, class... _Args, __enable_if_t<__has_construct<allocator_type, _Tp*, _Args...>::value, int> = 0>
+  _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 static void
+  construct(allocator_type& __a, _Tp* __p, _Args&&... __args)
+  {
+    _CCCL_SUPPRESS_DEPRECATED_PUSH
+    __a.construct(__p, _CUDA_VSTD::forward<_Args>(__args)...);
+    _CCCL_SUPPRESS_DEPRECATED_POP
+  }
+  template <class _Tp,
+            class... _Args,
+            class                                                                       = void,
+            __enable_if_t<!__has_construct<allocator_type, _Tp*, _Args...>::value, int> = 0>
+  _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 static void construct(allocator_type&, _Tp* __p, _Args&&... __args)
+  {
+#if _CCCL_STD_VER >= 2020
+    _CUDA_VSTD::construct_at(__p, _CUDA_VSTD::forward<_Args>(__args)...);
+#else
+    ::new ((void*) __p) _Tp(_CUDA_VSTD::forward<_Args>(__args)...);
+#endif
+  }
+
+  template <class _Tp, __enable_if_t<__has_destroy<allocator_type, _Tp*>::value, int> = 0>
+  _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 static void destroy(allocator_type& __a, _Tp* __p)
+  {
+    _CCCL_SUPPRESS_DEPRECATED_PUSH
+    __a.destroy(__p);
+    _CCCL_SUPPRESS_DEPRECATED_POP
+  }
+  template <class _Tp, class = void, __enable_if_t<!__has_destroy<allocator_type, _Tp*>::value, int> = 0>
+  _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 static void destroy(allocator_type&, _Tp* __p)
+  {
+#if _CCCL_STD_VER >= 2020
+    _CUDA_VSTD::destroy_at(__p);
+#else
+    __p->~_Tp();
+#endif
+  }
+
+  _CCCL_SUPPRESS_DEPRECATED_PUSH
+  template <class _Ap = _Alloc, __enable_if_t<__has_max_size<const _Ap>::value, int> = 0>
+  _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 static size_type max_size(const allocator_type& __a) noexcept
+  {
+    return __a.max_size();
+  }
+  _CCCL_SUPPRESS_DEPRECATED_POP
+  template <class _Ap = _Alloc, class = void, __enable_if_t<!__has_max_size<const _Ap>::value, int> = 0>
+  _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 static size_type max_size(const allocator_type&) noexcept
+  {
+    return numeric_limits<size_type>::max() / sizeof(value_type);
+  }
+
+  template <class _Ap = _Alloc, __enable_if_t<__has_select_on_container_copy_construction<const _Ap>::value, int> = 0>
+  _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 static allocator_type
+  select_on_container_copy_construction(const allocator_type& __a)
+  {
+    return __a.select_on_container_copy_construction();
+  }
+  template <class _Ap                                                                          = _Alloc,
+            class                                                                              = void,
+            __enable_if_t<!__has_select_on_container_copy_construction<const _Ap>::value, int> = 0>
+  _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 static allocator_type
+  select_on_container_copy_construction(const allocator_type& __a)
+  {
+    return __a;
+  }
+
+  template <class _Ptr>
+  _LIBCUDACXX_INLINE_VISIBILITY static void
+  __construct_forward_with_exception_guarantees(allocator_type& __a, _Ptr __begin1, _Ptr __end1, _Ptr& __begin2)
+  {
+    static_assert(__is_cpp17_move_insertable<allocator_type>::value,
+                  "The specified type does not meet the requirements of Cpp17MoveInsertible");
+    for (; __begin1 != __end1; ++__begin1, (void) ++__begin2)
+    {
+      construct(__a,
+                _CUDA_VSTD::__to_raw_pointer(__begin2),
+#ifdef _LIBCUDACXX_NO_EXCEPTIONS
+                _CUDA_VSTD::move(*__begin1)
+#else
+                _CUDA_VSTD::move_if_noexcept(*__begin1)
+#endif
+      );
+    }
+  }
+
+  template <class _Tp>
+  _LIBCUDACXX_INLINE_VISIBILITY static __enable_if_t<
+    (__is_default_allocator<allocator_type>::value || !__has_construct<allocator_type, _Tp*, _Tp>::value)
+      && is_trivially_move_constructible<_Tp>::value,
+    void>
+  __construct_forward_with_exception_guarantees(allocator_type&, _Tp* __begin1, _Tp* __end1, _Tp*& __begin2)
+  {
+    ptrdiff_t _Np = __end1 - __begin1;
+    if (_Np > 0)
+    {
+      _CUDA_VSTD::memcpy(__begin2, __begin1, _Np * sizeof(_Tp));
+      __begin2 += _Np;
+    }
+  }
+
+  template <class _Iter, class _Ptr>
+  _LIBCUDACXX_INLINE_VISIBILITY static void
+  __construct_range_forward(allocator_type& __a, _Iter __begin1, _Iter __end1, _Ptr& __begin2)
+  {
+    for (; __begin1 != __end1; ++__begin1, (void) ++__begin2)
+    {
+      construct(__a, _CUDA_VSTD::__to_raw_pointer(__begin2), *__begin1);
+    }
+  }
+
+  template <class _SourceTp,
+            class _DestTp,
+            class _RawSourceTp = __remove_const_t<_SourceTp>,
+            class _RawDestTp   = __remove_const_t<_DestTp>>
+  _LIBCUDACXX_INLINE_VISIBILITY static __enable_if_t<
+    is_trivially_move_constructible<_DestTp>::value && is_same<_RawSourceTp, _RawDestTp>::value
+      && (__is_default_allocator<allocator_type>::value || !__has_construct<allocator_type, _DestTp*, _SourceTp&>::value),
+    void>
+  __construct_range_forward(allocator_type&, _SourceTp* __begin1, _SourceTp* __end1, _DestTp*& __begin2)
+  {
+    ptrdiff_t _Np = __end1 - __begin1;
+    if (_Np > 0)
+    {
+      _CUDA_VSTD::memcpy(const_cast<_RawDestTp*>(__begin2), __begin1, _Np * sizeof(_DestTp));
+      __begin2 += _Np;
+    }
+  }
+
+  template <class _Ptr>
+  _LIBCUDACXX_INLINE_VISIBILITY static void
+  __construct_backward_with_exception_guarantees(allocator_type& __a, _Ptr __begin1, _Ptr __end1, _Ptr& __end2)
+  {
+    static_assert(__is_cpp17_move_insertable<allocator_type>::value,
+                  "The specified type does not meet the requirements of Cpp17MoveInsertable");
+    while (__end1 != __begin1)
+    {
+      construct(__a,
+                _CUDA_VSTD::__to_raw_pointer(__end2 - 1),
+#ifdef _LIBCUDACXX_NO_EXCEPTIONS
+                _CUDA_VSTD::move(*--__end1)
+#else
+                _CUDA_VSTD::move_if_noexcept(*--__end1)
+#endif
+      );
+      --__end2;
+    }
+  }
+
+  template <class _Tp>
+  _LIBCUDACXX_INLINE_VISIBILITY static __enable_if_t<
+    (__is_default_allocator<allocator_type>::value || !__has_construct<allocator_type, _Tp*, _Tp>::value)
+      && is_trivially_move_constructible<_Tp>::value,
+    void>
+  __construct_backward_with_exception_guarantees(allocator_type&, _Tp* __begin1, _Tp* __end1, _Tp*& __end2)
+  {
+    ptrdiff_t _Np = __end1 - __begin1;
+    __end2 -= _Np;
+    if (_Np > 0)
+    {
+      _CUDA_VSTD::memcpy(__end2, __begin1, _Np * sizeof(_Tp));
+    }
+  }
+};
+
+template <class _Traits, class _Tp>
+using __rebind_alloc _LIBCUDACXX_NODEBUG = typename _Traits::template rebind_alloc<_Tp>;
+
+template <class _Traits, class _Tp>
+struct __rebind_alloc_helper
+{
+  typedef _LIBCUDACXX_NODEBUG_TYPE typename _Traits::template rebind_alloc<_Tp> type;
+};
+
+// ASan choices
+#ifndef _LIBCUDACXX_HAS_NO_ASAN
+#  define _LIBCUDACXX_HAS_ASAN_CONTAINER_ANNOTATIONS_FOR_ALL_ALLOCATORS 1
+#endif
+
+#ifdef _LIBCUDACXX_HAS_ASAN_CONTAINER_ANNOTATIONS_FOR_ALL_ALLOCATORS
+template <class _Alloc>
+struct __asan_annotate_container_with_allocator
+#  if defined(_LIBCUDACXX_CLANG_VER) && _LIBCUDACXX_CLANG_VER >= 1600
+    : true_type
+{};
+#  else
+    // TODO LLVM18: Remove the special-casing
+    : false_type
+{
+};
+#  endif
+
+template <class _Tp>
+struct __asan_annotate_container_with_allocator<allocator<_Tp>> : true_type
+{};
+#endif
+
+#undef _LIBCUDACXX_ALLOCATOR_TRAITS_HAS_XXX
+
+_LIBCUDACXX_END_NAMESPACE_STD
+
+_CCCL_NV_DIAG_DEFAULT(1215)
+
+#include <cuda/std/detail/__pragma_pop>
+
+#endif // _LIBCUDACXX___MEMORY_ALLOCATOR_TRAITS_H
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/uses_allocator.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/uses_allocator.h
new file mode 100644
index 0000000000..1e0444f86d
--- /dev/null
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/uses_allocator.h
@@ -0,0 +1,70 @@
+// -*- C++ -*-
+//===----------------------------------------------------------------------===//
+//
+// 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___MEMORY_USES_ALLOCATOR_H
+#define _LIBCUDACXX___MEMORY_USES_ALLOCATOR_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/cstddef>
+#include <cuda/std/detail/libcxx/include/__type_traits/is_convertible.h>
+#include <cuda/std/detail/libcxx/include/__type_traits/void_t.h>
+
+_LIBCUDACXX_BEGIN_NAMESPACE_STD
+
+#if _CCCL_STD_VER <= 2014
+template <class _Tp, class = void>
+struct __has_allocator_type : false_type
+{};
+template <class _Tp>
+struct __has_allocator_type<_Tp, __void_t<typename _Tp::allocator_type>> : true_type
+{};
+
+template <class _Tp, class _Alloc, bool = _LIBCUDACXX_TRAIT(__has_allocator_type, _Tp)>
+struct __uses_allocator : false_type
+{};
+template <class _Tp, class _Alloc>
+struct __uses_allocator<_Tp, _Alloc, true> : is_convertible<_Alloc, typename _Tp::allocator_type>
+{};
+#else // ^^^ _CCCL_STD_VER <= 2014 ^^^ / vvv _CCCL_STD_VER >= 2017 vvv
+template <class _Tp, class = void>
+_LIBCUDACXX_INLINE_VAR constexpr bool __has_allocator_type_v = false;
+template <class _Tp>
+_LIBCUDACXX_INLINE_VAR constexpr bool __has_allocator_type_v<_Tp, __void_t<typename _Tp::allocator_type>> = true;
+
+template <class _Tp, class _Alloc, bool = _LIBCUDACXX_TRAIT(__has_allocator_type, _Tp)>
+_LIBCUDACXX_INLINE_VAR constexpr bool __uses_allocator_v = false;
+template <class _Tp, class _Alloc>
+_LIBCUDACXX_INLINE_VAR constexpr bool __uses_allocator_v<_Tp, _Alloc, true> =
+  is_convertible_v<_Alloc, typename _Tp::allocator_type>;
+#endif // _CCCL_STD_VER >= 2017
+
+template <class _Tp, class _Alloc>
+struct _LIBCUDACXX_TEMPLATE_VIS uses_allocator
+    : public integral_constant<bool, _LIBCUDACXX_TRAIT(__uses_allocator, _Tp, _Alloc)>
+{};
+
+#if _CCCL_STD_VER >= 2014
+template <class _Tp, class _Alloc>
+_LIBCUDACXX_INLINE_VAR constexpr bool uses_allocator_v = _LIBCUDACXX_TRAIT(__uses_allocator, _Tp, _Alloc);
+#endif // _CCCL_STD_VER >= 2014
+
+_LIBCUDACXX_END_NAMESPACE_STD
+
+#endif // _LIBCUDACXX___MEMORY_USES_ALLOCATOR_H
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__new/allocate.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__new/allocate.h
index da553efed8..b746c86506 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/__new/allocate.h
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__new/allocate.h
@@ -59,7 +59,8 @@ inline _LIBCUDACXX_INLINE_VISIBILITY constexpr bool __is_overaligned_for_new(siz
 template <class... _Args>
 _LIBCUDACXX_INLINE_VISIBILITY void* __libcpp_operator_new(_Args... __args)
 {
-#if __has_builtin(__builtin_operator_new) && __has_builtin(__builtin_operator_delete)
+  // Those builtins are not usable on device and the tests crash when using them
+#if 0 && __has_builtin(__builtin_operator_new) && __has_builtin(__builtin_operator_delete)
   return __builtin_operator_new(__args...);
 #else // ^^^ use builtin ^^^ / vvv no builtin
   return ::operator new(__args...);
@@ -69,7 +70,8 @@ _LIBCUDACXX_INLINE_VISIBILITY void* __libcpp_operator_new(_Args... __args)
 template <class... _Args>
 _LIBCUDACXX_INLINE_VISIBILITY void __libcpp_operator_delete(_Args... __args)
 {
-#if __has_builtin(__builtin_operator_new) && __has_builtin(__builtin_operator_delete)
+  // Those builtins are not usable on device and the tests crash when using them
+#if 0 && __has_builtin(__builtin_operator_new) && __has_builtin(__builtin_operator_delete)
   __builtin_operator_delete(__args...);
 #else // ^^^ use builtin ^^^ / vvv no builtin
   ::operator delete(__args...);
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/memory b/libcudacxx/include/cuda/std/detail/libcxx/include/memory
index 81aa132e42..1803f61fea 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/memory
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/memory
@@ -647,24 +647,30 @@ void* align(size_t alignment, size_t size, void*& ptr, size_t& space);
 }  // std
 
 */
-
-#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
+#ifndef __cuda_std__
+#include <__config>
+#include <typeinfo>
+#include <cstring>
+#if !defined(_LIBCUDACXX_HAS_NO_ATOMIC_HEADER)
+#  include <atomic>
+#endif
+#endif //__cuda_std__
 
 #include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/detail/libcxx/include/__functional_base>
 #include <cuda/std/detail/libcxx/include/__iterator/iterator_traits.h>
 #include <cuda/std/detail/libcxx/include/__iterator/iterator.h>
 #include <cuda/std/detail/libcxx/include/__memory/addressof.h>
+#include <cuda/std/detail/libcxx/include/__memory/align.h>
+#include <cuda/std/detail/libcxx/include/__memory/allocate_at_least.h>
+#include <cuda/std/detail/libcxx/include/__memory/allocation_guard.h>
+#include <cuda/std/detail/libcxx/include/__memory/allocator_arg_t.h>
+#include <cuda/std/detail/libcxx/include/__memory/allocator_destructor.h>
+#include <cuda/std/detail/libcxx/include/__memory/allocator_traits.h>
+#include <cuda/std/detail/libcxx/include/__memory/allocator.h>
 #include <cuda/std/detail/libcxx/include/__memory/construct_at.h>
 #include <cuda/std/detail/libcxx/include/__memory/pointer_traits.h>
+#include <cuda/std/detail/libcxx/include/__memory/uses_allocator.h>
 #include <cuda/std/detail/libcxx/include/__memory/voidify.h>
 #include <cuda/std/detail/libcxx/include/__tuple_dir/tuple_indices.h>
 #include <cuda/std/detail/libcxx/include/__type_traits/decay.h>
@@ -731,978 +737,6 @@ _ValueType __libcpp_acquire_load(_ValueType const* __value) {
 #endif
 }
 
-// addressof moved to <type_traits>
-
-template <class _Tp> class allocator;
-
-template <>
-class _LIBCUDACXX_TEMPLATE_VIS allocator<void>
-{
-public:
-    typedef void*             pointer;
-    typedef const void*       const_pointer;
-    typedef void              value_type;
-
-    template <class _Up> struct rebind {typedef allocator<_Up> other;};
-};
-
-template <>
-class _LIBCUDACXX_TEMPLATE_VIS allocator<const void>
-{
-public:
-    typedef const void*       pointer;
-    typedef const void*       const_pointer;
-    typedef const void        value_type;
-
-    template <class _Up> struct rebind {typedef allocator<_Up> other;};
-};
-
-// allocator_traits
-
-template <class _Tp, class = void>
-struct __has_pointer_type : false_type {};
-
-template <class _Tp>
-struct __has_pointer_type<_Tp,
-          __void_t<typename _Tp::pointer>> : true_type {};
-
-namespace __pointer_type_imp
-{
-
-template <class _Tp, class _Dp, bool = __has_pointer_type<_Dp>::value>
-struct __pointer_type
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename _Dp::pointer type;
-};
-
-template <class _Tp, class _Dp>
-struct __pointer_type<_Tp, _Dp, false>
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE _Tp* type;
-};
-
-}  // __pointer_type_imp
-
-template <class _Tp, class _Dp>
-struct __pointer_type
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename __pointer_type_imp::__pointer_type<_Tp, __libcpp_remove_reference_t<_Dp>>::type type;
-};
-
-template <class _Tp, class = void>
-struct __has_const_pointer : false_type {};
-
-template <class _Tp>
-struct __has_const_pointer<_Tp,
-            __void_t<typename _Tp::const_pointer>> : true_type {};
-
-template <class _Tp, class _Ptr, class _Alloc, bool = __has_const_pointer<_Alloc>::value>
-struct __const_pointer
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename _Alloc::const_pointer type;
-};
-
-template <class _Tp, class _Ptr, class _Alloc>
-struct __const_pointer<_Tp, _Ptr, _Alloc, false>
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename pointer_traits<_Ptr>::template rebind<const _Tp> type;
-};
-
-template <class _Tp, class = void>
-struct __has_void_pointer : false_type {};
-
-template <class _Tp>
-struct __has_void_pointer<_Tp,
-               __void_t<typename _Tp::void_pointer>> : true_type {};
-
-template <class _Ptr, class _Alloc, bool = __has_void_pointer<_Alloc>::value>
-struct __void_pointer
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename _Alloc::void_pointer type;
-};
-
-template <class _Ptr, class _Alloc>
-struct __void_pointer<_Ptr, _Alloc, false>
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename pointer_traits<_Ptr>::template rebind<void> type;
-};
-
-template <class _Tp, class = void>
-struct __has_const_void_pointer : false_type {};
-
-template <class _Tp>
-struct __has_const_void_pointer<_Tp,
-            __void_t<typename _Tp::const_void_pointer>> : true_type {};
-
-template <class _Ptr, class _Alloc, bool = __has_const_void_pointer<_Alloc>::value>
-struct __const_void_pointer
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename _Alloc::const_void_pointer type;
-};
-
-template <class _Ptr, class _Alloc>
-struct __const_void_pointer<_Ptr, _Alloc, false>
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename pointer_traits<_Ptr>::template rebind<const void> type;
-};
-
-template <class _Tp>
-inline _LIBCUDACXX_INLINE_VISIBILITY constexpr
-_Tp*
-__to_raw_pointer(_Tp* __p) noexcept
-{
-    return __p;
-}
-
-#if _CCCL_STD_VER <= 2017
-template <class _Pointer>
-inline _LIBCUDACXX_INLINE_VISIBILITY
-typename pointer_traits<_Pointer>::element_type*
-__to_raw_pointer(_Pointer __p) noexcept
-{
-    return _CUDA_VSTD::__to_raw_pointer(__p.operator->());
-}
-#else
-template <class _Pointer>
-inline _LIBCUDACXX_INLINE_VISIBILITY
-auto
-__to_raw_pointer(const _Pointer& __p) noexcept
--> decltype(pointer_traits<_Pointer>::to_address(__p))
-{
-    return pointer_traits<_Pointer>::to_address(__p);
-}
-
-template <class _Pointer, class... _None>
-inline _LIBCUDACXX_INLINE_VISIBILITY
-auto
-__to_raw_pointer(const _Pointer& __p, _None...) noexcept
-{
-    return _CUDA_VSTD::__to_raw_pointer(__p.operator->());
-}
-#endif
-
-template <class _Tp, class = void>
-struct __has_size_type : false_type {};
-
-template <class _Tp>
-struct __has_size_type<_Tp,
-               __void_t<typename _Tp::size_type>> : true_type {};
-
-template <class _Alloc, class _DiffType, bool = __has_size_type<_Alloc>::value>
-struct __size_type
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename make_unsigned<_DiffType>::type type;
-};
-
-template <class _Alloc, class _DiffType>
-struct __size_type<_Alloc, _DiffType, true>
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename _Alloc::size_type type;
-};
-
-template <class _Tp, class = void>
-struct __has_propagate_on_container_copy_assignment : false_type {};
-
-template <class _Tp>
-struct __has_propagate_on_container_copy_assignment<_Tp,
-    __void_t<typename _Tp::propagate_on_container_copy_assignment>>
-        : true_type {};
-
-template <class _Alloc, bool = __has_propagate_on_container_copy_assignment<_Alloc>::value>
-struct __propagate_on_container_copy_assignment
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE false_type type;
-};
-
-template <class _Alloc>
-struct __propagate_on_container_copy_assignment<_Alloc, true>
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename _Alloc::propagate_on_container_copy_assignment type;
-};
-
-template <class _Tp, class = void>
-struct __has_propagate_on_container_move_assignment : false_type {};
-
-template <class _Tp>
-struct __has_propagate_on_container_move_assignment<_Tp,
-           __void_t<typename _Tp::propagate_on_container_move_assignment>>
-               : true_type {};
-
-template <class _Alloc, bool = __has_propagate_on_container_move_assignment<_Alloc>::value>
-struct __propagate_on_container_move_assignment
-{
-    typedef false_type type;
-};
-
-template <class _Alloc>
-struct __propagate_on_container_move_assignment<_Alloc, true>
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename _Alloc::propagate_on_container_move_assignment type;
-};
-
-template <class _Tp, class = void>
-struct __has_propagate_on_container_swap : false_type {};
-
-template <class _Tp>
-struct __has_propagate_on_container_swap<_Tp,
-           __void_t<typename _Tp::propagate_on_container_swap>>
-               : true_type {};
-
-template <class _Alloc, bool = __has_propagate_on_container_swap<_Alloc>::value>
-struct __propagate_on_container_swap
-{
-    typedef false_type type;
-};
-
-template <class _Alloc>
-struct __propagate_on_container_swap<_Alloc, true>
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename _Alloc::propagate_on_container_swap type;
-};
-
-template <class _Tp, class = void>
-struct __has_is_always_equal : false_type {};
-
-template <class _Tp>
-struct __has_is_always_equal<_Tp,
-           __void_t<typename _Tp::is_always_equal>>
-               : true_type {};
-
-template <class _Alloc, bool = __has_is_always_equal<_Alloc>::value>
-struct __is_always_equal
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename _CUDA_VSTD::is_empty<_Alloc>::type type;
-};
-
-template <class _Alloc>
-struct __is_always_equal<_Alloc, true>
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename _Alloc::is_always_equal type;
-};
-
-template <class _Tp, class _Up, bool = __has_rebind<_Tp, _Up>::value>
-struct __has_rebind_other
-{
-private:
-    struct __two {char __lx; char __lxx;};
-    template <class _Xp> static __two __test(...);
-    template <class _Xp> static char __test(typename _Xp::template rebind<_Up>::other* = 0);
-public:
-    static const bool value = sizeof(__test<_Tp>(0)) == 1;
-};
-
-template <class _Tp, class _Up>
-struct __has_rebind_other<_Tp, _Up, false>
-{
-    static const bool value = false;
-};
-
-template <class _Tp, class _Up, bool = __has_rebind_other<_Tp, _Up>::value>
-struct __allocator_traits_rebind
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename _Tp::template rebind<_Up>::other type;
-};
-
-#ifndef _LIBCUDACXX_HAS_NO_VARIADICS
-
-template <template <class, class...> class _Alloc, class _Tp, class ..._Args, class _Up>
-struct __allocator_traits_rebind<_Alloc<_Tp, _Args...>, _Up, true>
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename _Alloc<_Tp, _Args...>::template rebind<_Up>::other type;
-};
-
-template <template <class, class...> class _Alloc, class _Tp, class ..._Args, class _Up>
-struct __allocator_traits_rebind<_Alloc<_Tp, _Args...>, _Up, false>
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE _Alloc<_Up, _Args...> type;
-};
-
-#else  // _LIBCUDACXX_HAS_NO_VARIADICS
-
-template <template <class> class _Alloc, class _Tp, class _Up>
-struct __allocator_traits_rebind<_Alloc<_Tp>, _Up, true>
-{
-    typedef typename _Alloc<_Tp>::template rebind<_Up>::other type;
-};
-
-template <template <class> class _Alloc, class _Tp, class _Up>
-struct __allocator_traits_rebind<_Alloc<_Tp>, _Up, false>
-{
-    typedef _Alloc<_Up> type;
-};
-
-template <template <class, class> class _Alloc, class _Tp, class _A0, class _Up>
-struct __allocator_traits_rebind<_Alloc<_Tp, _A0>, _Up, true>
-{
-    typedef typename _Alloc<_Tp, _A0>::template rebind<_Up>::other type;
-};
-
-template <template <class, class> class _Alloc, class _Tp, class _A0, class _Up>
-struct __allocator_traits_rebind<_Alloc<_Tp, _A0>, _Up, false>
-{
-    typedef _Alloc<_Up, _A0> type;
-};
-
-template <template <class, class, class> class _Alloc, class _Tp, class _A0,
-                                         class _A1, class _Up>
-struct __allocator_traits_rebind<_Alloc<_Tp, _A0, _A1>, _Up, true>
-{
-    typedef typename _Alloc<_Tp, _A0, _A1>::template rebind<_Up>::other type;
-};
-
-template <template <class, class, class> class _Alloc, class _Tp, class _A0,
-                                         class _A1, class _Up>
-struct __allocator_traits_rebind<_Alloc<_Tp, _A0, _A1>, _Up, false>
-{
-    typedef _Alloc<_Up, _A0, _A1> type;
-};
-
-template <template <class, class, class, class> class _Alloc, class _Tp, class _A0,
-                                                class _A1, class _A2, class _Up>
-struct __allocator_traits_rebind<_Alloc<_Tp, _A0, _A1, _A2>, _Up, true>
-{
-    typedef typename _Alloc<_Tp, _A0, _A1, _A2>::template rebind<_Up>::other type;
-};
-
-template <template <class, class, class, class> class _Alloc, class _Tp, class _A0,
-                                                class _A1, class _A2, class _Up>
-struct __allocator_traits_rebind<_Alloc<_Tp, _A0, _A1, _A2>, _Up, false>
-{
-    typedef _Alloc<_Up, _A0, _A1, _A2> type;
-};
-
-#endif  // _LIBCUDACXX_HAS_NO_VARIADICS
-
-template <class _Alloc, class _SizeType, class _ConstVoidPtr>
-auto
-__has_allocate_hint_test(_Alloc&& __a, _SizeType&& __sz, _ConstVoidPtr&& __p)
-    -> decltype((void)__a.allocate(__sz, __p), true_type());
-
-template <class _Alloc, class _SizeType, class _ConstVoidPtr>
-auto
-__has_allocate_hint_test(const _Alloc& __a, _SizeType&& __sz, _ConstVoidPtr&& __p)
-    -> false_type;
-
-template <class _Alloc, class _SizeType, class _ConstVoidPtr>
-struct __has_allocate_hint
-    : integral_constant<bool,
-        is_same<
-            decltype(_CUDA_VSTD::__has_allocate_hint_test(declval<_Alloc>(),
-                                          declval<_SizeType>(),
-                                          declval<_ConstVoidPtr>())),
-            true_type>::value>
-{
-};
-
-template <class _Alloc, class _Tp, class ..._Args>
-decltype(_CUDA_VSTD::declval<_Alloc>().construct(_CUDA_VSTD::declval<_Tp*>(),
-                                           _CUDA_VSTD::declval<_Args>()...),
-                                           true_type())
-__has_construct_test(_Alloc&& __a, _Tp* __p, _Args&& ...__args);
-
-template <class _Alloc, class _Pointer, class ..._Args>
-false_type
-__has_construct_test(const _Alloc& __a, _Pointer&& __p, _Args&& ...__args);
-
-template <class _Alloc, class _Pointer, class ..._Args>
-struct __has_construct
-    : integral_constant<bool,
-        is_same<
-            decltype(_CUDA_VSTD::__has_construct_test(declval<_Alloc>(),
-                                          declval<_Pointer>(),
-                                          declval<_Args>()...)),
-            true_type>::value>
-{
-};
-
-template <class _Alloc, class _Pointer>
-auto
-__has_destroy_test(_Alloc&& __a, _Pointer&& __p)
-    -> decltype(__a.destroy(__p), true_type());
-
-template <class _Alloc, class _Pointer>
-auto
-__has_destroy_test(const _Alloc& __a, _Pointer&& __p)
-    -> false_type;
-
-template <class _Alloc, class _Pointer>
-struct __has_destroy
-    : integral_constant<bool,
-        is_same<
-            decltype(_CUDA_VSTD::__has_destroy_test(declval<_Alloc>(),
-                                        declval<_Pointer>())),
-            true_type>::value>
-{
-};
-
-template <class _Alloc>
-auto
-__has_max_size_test(_Alloc&& __a)
-    -> decltype(__a.max_size(), true_type());
-
-template <class _Alloc>
-auto
-__has_max_size_test(const volatile _Alloc& __a)
-    -> false_type;
-
-template <class _Alloc>
-struct __has_max_size
-    : integral_constant<bool,
-        is_same<
-            decltype(_CUDA_VSTD::__has_max_size_test(declval<_Alloc&>())),
-            true_type>::value>
-{
-};
-
-template <class _Alloc>
-auto
-__has_select_on_container_copy_construction_test(_Alloc&& __a)
-    -> decltype(__a.select_on_container_copy_construction(), true_type());
-
-template <class _Alloc>
-auto
-__has_select_on_container_copy_construction_test(const volatile _Alloc& __a)
-    -> false_type;
-
-template <class _Alloc>
-struct __has_select_on_container_copy_construction
-    : integral_constant<bool,
-        is_same<
-            decltype(_CUDA_VSTD::__has_select_on_container_copy_construction_test(declval<_Alloc&>())),
-            true_type>::value>
-{
-};
-
-template <class _Alloc, class _Ptr, bool = __has_difference_type<_Alloc>::value>
-struct __alloc_traits_difference_type
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename pointer_traits<_Ptr>::difference_type type;
-};
-
-template <class _Alloc, class _Ptr>
-struct __alloc_traits_difference_type<_Alloc, _Ptr, true>
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename _Alloc::difference_type type;
-};
-
-template <class _Tp>
-struct __is_default_allocator : false_type {};
-
-template <class _Tp>
-struct __is_default_allocator<_CUDA_VSTD::allocator<_Tp> > : true_type {};
-
-
-
-template <class _Alloc,
-    bool = __has_construct<_Alloc, typename _Alloc::value_type*,  typename _Alloc::value_type&&>::value && !__is_default_allocator<_Alloc>::value
-    >
-struct __is_cpp17_move_insertable;
-template <class _Alloc>
-struct __is_cpp17_move_insertable<_Alloc, true> : std::true_type {};
-template <class _Alloc>
-struct __is_cpp17_move_insertable<_Alloc, false> : std::is_move_constructible<typename _Alloc::value_type> {};
-
-template <class _Alloc,
-    bool = __has_construct<_Alloc, typename _Alloc::value_type*, const typename _Alloc::value_type&>::value && !__is_default_allocator<_Alloc>::value
-    >
-struct __is_cpp17_copy_insertable;
-template <class _Alloc>
-struct __is_cpp17_copy_insertable<_Alloc, true> : __is_cpp17_move_insertable<_Alloc> {};
-template <class _Alloc>
-struct __is_cpp17_copy_insertable<_Alloc, false> : integral_constant<bool,
-    std::is_copy_constructible<typename _Alloc::value_type>::value &&
-    __is_cpp17_move_insertable<_Alloc>::value>
-  {};
-
-
-
-template <class _Alloc>
-struct _LIBCUDACXX_TEMPLATE_VIS allocator_traits
-{
-    typedef _Alloc                              allocator_type;
-    typedef typename allocator_type::value_type value_type;
-
-    typedef typename __pointer_type<value_type, allocator_type>::type pointer;
-    typedef typename __const_pointer<value_type, pointer, allocator_type>::type const_pointer;
-    typedef typename __void_pointer<pointer, allocator_type>::type void_pointer;
-    typedef typename __const_void_pointer<pointer, allocator_type>::type const_void_pointer;
-
-    typedef typename __alloc_traits_difference_type<allocator_type, pointer>::type difference_type;
-    typedef typename __size_type<allocator_type, difference_type>::type size_type;
-
-    typedef typename __propagate_on_container_copy_assignment<allocator_type>::type
-                     propagate_on_container_copy_assignment;
-    typedef typename __propagate_on_container_move_assignment<allocator_type>::type
-                     propagate_on_container_move_assignment;
-    typedef typename __propagate_on_container_swap<allocator_type>::type
-                     propagate_on_container_swap;
-    typedef typename __is_always_equal<allocator_type>::type
-                     is_always_equal;
-
-    template <class _Tp> using rebind_alloc =
-                  typename __allocator_traits_rebind<allocator_type, _Tp>::type;
-    template <class _Tp> using rebind_traits = allocator_traits<rebind_alloc<_Tp> >;
-
-    _LIBCUDACXX_NODISCARD_AFTER_CXX17 _LIBCUDACXX_INLINE_VISIBILITY
-    static pointer allocate(allocator_type& __a, size_type __n)
-        {return __a.allocate(__n);}
-    _LIBCUDACXX_NODISCARD_AFTER_CXX17 _LIBCUDACXX_INLINE_VISIBILITY
-    static pointer allocate(allocator_type& __a, size_type __n, const_void_pointer __hint)
-        {return __allocate(__a, __n, __hint,
-            __has_allocate_hint<allocator_type, size_type, const_void_pointer>());}
-
-    _LIBCUDACXX_INLINE_VISIBILITY
-    static void deallocate(allocator_type& __a, pointer __p, size_type __n) noexcept
-        {__a.deallocate(__p, __n);}
-
-#ifndef _LIBCUDACXX_HAS_NO_VARIADICS
-    template <class _Tp, class... _Args>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        static void construct(allocator_type& __a, _Tp* __p, _Args&&... __args)
-            {__construct(__has_construct<allocator_type, _Tp*, _Args...>(),
-                         __a, __p, _CUDA_VSTD::forward<_Args>(__args)...);}
-#else  // _LIBCUDACXX_HAS_NO_VARIADICS
-    template <class _Tp>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        static void construct(allocator_type&, _Tp* __p)
-            {
-                ::new ((void*)__p) _Tp();
-            }
-    template <class _Tp, class _A0>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        static void construct(allocator_type& __a, _Tp* __p, const _A0& __a0)
-            {
-                __construct(__has_construct<allocator_type, _Tp*, const _A0&>(),
-                            __a, __p, __a0);
-            }
-    template <class _Tp, class _A0, class _A1>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        static void construct(allocator_type&, _Tp* __p, const _A0& __a0,
-                              const _A1& __a1)
-            {
-                ::new ((void*)__p) _Tp(__a0, __a1);
-            }
-    template <class _Tp, class _A0, class _A1, class _A2>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        static void construct(allocator_type&, _Tp* __p, const _A0& __a0,
-                              const _A1& __a1, const _A2& __a2)
-            {
-                ::new ((void*)__p) _Tp(__a0, __a1, __a2);
-            }
-#endif  // _LIBCUDACXX_HAS_NO_VARIADICS
-
-    template <class _Tp>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        static void destroy(allocator_type& __a, _Tp* __p)
-            {__destroy(__has_destroy<allocator_type, _Tp*>(), __a, __p);}
-
-    _LIBCUDACXX_INLINE_VISIBILITY
-    static size_type max_size(const allocator_type& __a) noexcept
-        {return __max_size(__has_max_size<const allocator_type>(), __a);}
-
-    _LIBCUDACXX_INLINE_VISIBILITY
-    static allocator_type
-        select_on_container_copy_construction(const allocator_type& __a)
-            {return __select_on_container_copy_construction(
-                __has_select_on_container_copy_construction<const allocator_type>(),
-                __a);}
-
-    template <class _Ptr>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        static
-        void
-        __construct_forward_with_exception_guarantees(allocator_type& __a, _Ptr __begin1, _Ptr __end1, _Ptr& __begin2)
-        {
-            static_assert(__is_cpp17_move_insertable<allocator_type>::value,
-              "The specified type does not meet the requirements of Cpp17MoveInsertible");
-            for (; __begin1 != __end1; ++__begin1, (void) ++__begin2)
-              construct(__a, _CUDA_VSTD::__to_raw_pointer(__begin2),
-#ifdef _LIBCUDACXX_NO_EXCEPTIONS
-                        _CUDA_VSTD::move(*__begin1)
-#else
-                        _CUDA_VSTD::move_if_noexcept(*__begin1)
-#endif
-                        );
-        }
-
-    template <class _Tp>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        static
-        typename enable_if
-        <
-            (__is_default_allocator<allocator_type>::value
-                || !__has_construct<allocator_type, _Tp*, _Tp>::value) &&
-             is_trivially_move_constructible<_Tp>::value,
-            void
-        >::type
-        __construct_forward_with_exception_guarantees(allocator_type&, _Tp* __begin1, _Tp* __end1, _Tp*& __begin2)
-        {
-            ptrdiff_t _Np = __end1 - __begin1;
-            if (_Np > 0)
-            {
-                _CUDA_VSTD::memcpy(__begin2, __begin1, _Np * sizeof(_Tp));
-                __begin2 += _Np;
-            }
-        }
-
-    template <class _Iter, class _Ptr>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        static
-        void
-        __construct_range_forward(allocator_type& __a, _Iter __begin1, _Iter __end1, _Ptr& __begin2)
-        {
-            for (; __begin1 != __end1; ++__begin1, (void) ++__begin2)
-                construct(__a, _CUDA_VSTD::__to_raw_pointer(__begin2), *__begin1);
-        }
-
-    template <class _SourceTp, class _DestTp,
-              class _RawSourceTp = __remove_const_t<_SourceTp>,
-              class _RawDestTp = __remove_const_t<_DestTp>>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        static
-        typename enable_if
-        <
-            is_trivially_move_constructible<_DestTp>::value &&
-            is_same<_RawSourceTp, _RawDestTp>::value &&
-            (__is_default_allocator<allocator_type>::value ||
-             !__has_construct<allocator_type, _DestTp*, _SourceTp&>::value),
-            void
-        >::type
-        __construct_range_forward(allocator_type&, _SourceTp* __begin1, _SourceTp* __end1, _DestTp*& __begin2)
-        {
-            ptrdiff_t _Np = __end1 - __begin1;
-            if (_Np > 0)
-            {
-                _CUDA_VSTD::memcpy(const_cast<_RawDestTp*>(__begin2), __begin1, _Np * sizeof(_DestTp));
-                __begin2 += _Np;
-            }
-        }
-
-    template <class _Ptr>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        static
-        void
-        __construct_backward_with_exception_guarantees(allocator_type& __a, _Ptr __begin1, _Ptr __end1, _Ptr& __end2)
-        {
-            static_assert(__is_cpp17_move_insertable<allocator_type>::value,
-              "The specified type does not meet the requirements of Cpp17MoveInsertable");
-            while (__end1 != __begin1)
-            {
-              construct(__a, _CUDA_VSTD::__to_raw_pointer(__end2 - 1),
-#ifdef _LIBCUDACXX_NO_EXCEPTIONS
-                        _CUDA_VSTD::move(*--__end1)
-#else
-                        _CUDA_VSTD::move_if_noexcept(*--__end1)
-#endif
-                        );
-              --__end2;
-            }
-        }
-
-    template <class _Tp>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        static
-        typename enable_if
-        <
-            (__is_default_allocator<allocator_type>::value
-                || !__has_construct<allocator_type, _Tp*, _Tp>::value) &&
-             is_trivially_move_constructible<_Tp>::value,
-            void
-        >::type
-        __construct_backward_with_exception_guarantees(allocator_type&, _Tp* __begin1, _Tp* __end1, _Tp*& __end2)
-        {
-            ptrdiff_t _Np = __end1 - __begin1;
-            __end2 -= _Np;
-            if (_Np > 0)
-                _CUDA_VSTD::memcpy(__end2, __begin1, _Np * sizeof(_Tp));
-        }
-
-private:
-
-    _LIBCUDACXX_INLINE_VISIBILITY
-    static pointer __allocate(allocator_type& __a, size_type __n,
-        const_void_pointer __hint, true_type)
-        {return __a.allocate(__n, __hint);}
-    _LIBCUDACXX_INLINE_VISIBILITY
-    static pointer __allocate(allocator_type& __a, size_type __n,
-        const_void_pointer, false_type)
-        {return __a.allocate(__n);}
-
-#ifndef _LIBCUDACXX_HAS_NO_VARIADICS
-    template <class _Tp, class... _Args>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        static void __construct(true_type, allocator_type& __a, _Tp* __p, _Args&&... __args)
-            {__a.construct(__p, _CUDA_VSTD::forward<_Args>(__args)...);}
-    template <class _Tp, class... _Args>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        static void __construct(false_type, allocator_type&, _Tp* __p, _Args&&... __args)
-            {
-                ::new ((void*)__p) _Tp(_CUDA_VSTD::forward<_Args>(__args)...);
-            }
-#else  // _LIBCUDACXX_HAS_NO_VARIADICS
-    template <class _Tp, class _A0>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        static void __construct(true_type, allocator_type& __a, _Tp* __p,
-                                const _A0& __a0)
-            {__a.construct(__p, __a0);}
-    template <class _Tp, class _A0>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        static void __construct(false_type, allocator_type&, _Tp* __p,
-                                const _A0& __a0)
-            {
-                ::new ((void*)__p) _Tp(__a0);
-            }
-#endif  // _LIBCUDACXX_HAS_NO_VARIADICS
-
-    template <class _Tp>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        static void __destroy(true_type, allocator_type& __a, _Tp* __p)
-            {__a.destroy(__p);}
-    template <class _Tp>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        static void __destroy(false_type, allocator_type&, _Tp* __p)
-            {
-                __p->~_Tp();
-            }
-
-    _LIBCUDACXX_INLINE_VISIBILITY
-    static size_type __max_size(true_type, const allocator_type& __a) noexcept
-            {return __a.max_size();}
-    _LIBCUDACXX_INLINE_VISIBILITY
-    static size_type __max_size(false_type, const allocator_type&) noexcept
-            {return numeric_limits<size_type>::max() / sizeof(value_type);}
-
-    _LIBCUDACXX_INLINE_VISIBILITY
-    static allocator_type
-        __select_on_container_copy_construction(true_type, const allocator_type& __a)
-            {return __a.select_on_container_copy_construction();}
-    _LIBCUDACXX_INLINE_VISIBILITY
-    static allocator_type
-        __select_on_container_copy_construction(false_type, const allocator_type& __a)
-            {return __a;}
-};
-
-template <class _Traits, class _Tp>
-struct __rebind_alloc_helper
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename _Traits::template rebind_alloc<_Tp>        type;
-};
-
-// allocator
-
-template <class _Tp>
-class _LIBCUDACXX_TEMPLATE_VIS allocator
-{
-public:
-    typedef size_t            size_type;
-    typedef ptrdiff_t         difference_type;
-    typedef _Tp*              pointer;
-    typedef const _Tp*        const_pointer;
-    typedef _Tp&              reference;
-    typedef const _Tp&        const_reference;
-    typedef _Tp               value_type;
-
-    typedef true_type propagate_on_container_move_assignment;
-    typedef true_type is_always_equal;
-
-    template <class _Up> struct rebind {typedef allocator<_Up> other;};
-
-    _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20
-    allocator() noexcept {}
-
-    template <class _Up>
-    _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20
-    allocator(const allocator<_Up>&) noexcept {}
-
-    _LIBCUDACXX_INLINE_VISIBILITY pointer address(reference __x) const noexcept
-        {return _CUDA_VSTD::addressof(__x);}
-    _LIBCUDACXX_INLINE_VISIBILITY const_pointer address(const_reference __x) const noexcept
-        {return _CUDA_VSTD::addressof(__x);}
-    _LIBCUDACXX_NODISCARD_AFTER_CXX17 _LIBCUDACXX_INLINE_VISIBILITY
-    pointer allocate(size_type __n, allocator<void>::const_pointer = 0)
-        {
-        if (__n > max_size())
-            __throw_length_error("allocator<T>::allocate(size_t n)"
-                                 " 'n' exceeds maximum supported size");
-        return static_cast<pointer>(_CUDA_VSTD::__libcpp_allocate(__n * sizeof(_Tp), _LIBCUDACXX_ALIGNOF(_Tp)));
-        }
-    _LIBCUDACXX_INLINE_VISIBILITY void deallocate(pointer __p, size_type __n) noexcept
-        {_CUDA_VSTD::__libcpp_deallocate((void*)__p, __n * sizeof(_Tp), _LIBCUDACXX_ALIGNOF(_Tp));}
-    _LIBCUDACXX_INLINE_VISIBILITY size_type max_size() const noexcept
-        {return size_type(~0) / sizeof(_Tp);}
-#if !defined(_LIBCUDACXX_HAS_NO_RVALUE_REFERENCES) && !defined(_LIBCUDACXX_HAS_NO_VARIADICS)
-    template <class _Up, class... _Args>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        void
-        construct(_Up* __p, _Args&&... __args)
-        {
-            ::new((void*)__p) _Up(_CUDA_VSTD::forward<_Args>(__args)...);
-        }
-#else  // !defined(_LIBCUDACXX_HAS_NO_RVALUE_REFERENCES) && !defined(_LIBCUDACXX_HAS_NO_VARIADICS)
-        _LIBCUDACXX_INLINE_VISIBILITY
-        void
-        construct(pointer __p)
-        {
-            ::new((void*)__p) _Tp();
-        }
-# if defined(_LIBCUDACXX_HAS_NO_RVALUE_REFERENCES)
-
-    template <class _A0>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        void
-        construct(pointer __p, _A0& __a0)
-        {
-            ::new((void*)__p) _Tp(__a0);
-        }
-    template <class _A0>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        void
-        construct(pointer __p, const _A0& __a0)
-        {
-            ::new((void*)__p) _Tp(__a0);
-        }
-# endif  // defined(_LIBCUDACXX_HAS_NO_RVALUE_REFERENCES)
-    template <class _A0, class _A1>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        void
-        construct(pointer __p, _A0& __a0, _A1& __a1)
-        {
-            ::new((void*)__p) _Tp(__a0, __a1);
-        }
-    template <class _A0, class _A1>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        void
-        construct(pointer __p, const _A0& __a0, _A1& __a1)
-        {
-            ::new((void*)__p) _Tp(__a0, __a1);
-        }
-    template <class _A0, class _A1>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        void
-        construct(pointer __p, _A0& __a0, const _A1& __a1)
-        {
-            ::new((void*)__p) _Tp(__a0, __a1);
-        }
-    template <class _A0, class _A1>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        void
-        construct(pointer __p, const _A0& __a0, const _A1& __a1)
-        {
-            ::new((void*)__p) _Tp(__a0, __a1);
-        }
-#endif  // !defined(_LIBCUDACXX_HAS_NO_RVALUE_REFERENCES) && !defined(_LIBCUDACXX_HAS_NO_VARIADICS)
-    _LIBCUDACXX_INLINE_VISIBILITY void destroy(pointer __p) {__p->~_Tp();}
-};
-
-template <class _Tp>
-class _LIBCUDACXX_TEMPLATE_VIS allocator<const _Tp>
-{
-public:
-    typedef size_t            size_type;
-    typedef ptrdiff_t         difference_type;
-    typedef const _Tp*        pointer;
-    typedef const _Tp*        const_pointer;
-    typedef const _Tp&        reference;
-    typedef const _Tp&        const_reference;
-    typedef const _Tp         value_type;
-
-    typedef true_type propagate_on_container_move_assignment;
-    typedef true_type is_always_equal;
-
-    template <class _Up> struct rebind {typedef allocator<_Up> other;};
-
-    _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20
-    allocator() noexcept {}
-
-    template <class _Up>
-    _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20
-    allocator(const allocator<_Up>&) noexcept {}
-
-    _LIBCUDACXX_INLINE_VISIBILITY const_pointer address(const_reference __x) const noexcept
-        {return _CUDA_VSTD::addressof(__x);}
-    _LIBCUDACXX_INLINE_VISIBILITY pointer allocate(size_type __n, allocator<void>::const_pointer = 0)
-    {
-        if (__n > max_size())
-            __throw_length_error("allocator<const T>::allocate(size_t n)"
-                                 " 'n' exceeds maximum supported size");
-        return static_cast<pointer>(_CUDA_VSTD::__libcpp_allocate(__n * sizeof(_Tp), _LIBCUDACXX_ALIGNOF(_Tp)));
-    }
-    _LIBCUDACXX_INLINE_VISIBILITY void deallocate(pointer __p, size_type __n) noexcept
-        {_CUDA_VSTD::__libcpp_deallocate((void*) const_cast<_Tp *>(__p), __n * sizeof(_Tp), _LIBCUDACXX_ALIGNOF(_Tp));}
-    _LIBCUDACXX_INLINE_VISIBILITY size_type max_size() const noexcept
-        {return size_type(~0) / sizeof(_Tp);}
-#if !defined(_LIBCUDACXX_HAS_NO_RVALUE_REFERENCES) && !defined(_LIBCUDACXX_HAS_NO_VARIADICS)
-    template <class _Up, class... _Args>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        void
-        construct(_Up* __p, _Args&&... __args)
-        {
-            ::new((void*)__p) _Up(_CUDA_VSTD::forward<_Args>(__args)...);
-        }
-#else  // !defined(_LIBCUDACXX_HAS_NO_RVALUE_REFERENCES) && !defined(_LIBCUDACXX_HAS_NO_VARIADICS)
-        _LIBCUDACXX_INLINE_VISIBILITY
-        void
-        construct(pointer __p)
-        {
-            ::new((void*) const_cast<_Tp *>(__p)) _Tp();
-        }
-# if defined(_LIBCUDACXX_HAS_NO_RVALUE_REFERENCES)
-
-    template <class _A0>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        void
-        construct(pointer __p, _A0& __a0)
-        {
-            ::new((void*) const_cast<_Tp *>(__p)) _Tp(__a0);
-        }
-    template <class _A0>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        void
-        construct(pointer __p, const _A0& __a0)
-        {
-            ::new((void*) const_cast<_Tp *>(__p)) _Tp(__a0);
-        }
-# endif  // defined(_LIBCUDACXX_HAS_NO_RVALUE_REFERENCES)
-    template <class _A0, class _A1>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        void
-        construct(pointer __p, _A0& __a0, _A1& __a1)
-        {
-            ::new((void*) const_cast<_Tp *>(__p)) _Tp(__a0, __a1);
-        }
-    template <class _A0, class _A1>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        void
-        construct(pointer __p, const _A0& __a0, _A1& __a1)
-        {
-            ::new((void*) const_cast<_Tp *>(__p)) _Tp(__a0, __a1);
-        }
-    template <class _A0, class _A1>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        void
-        construct(pointer __p, _A0& __a0, const _A1& __a1)
-        {
-            ::new((void*) const_cast<_Tp *>(__p)) _Tp(__a0, __a1);
-        }
-    template <class _A0, class _A1>
-        _LIBCUDACXX_INLINE_VISIBILITY
-        void
-        construct(pointer __p, const _A0& __a0, const _A1& __a1)
-        {
-            ::new((void*) const_cast<_Tp *>(__p)) _Tp(__a0, __a1);
-        }
-#endif  // !defined(_LIBCUDACXX_HAS_NO_RVALUE_REFERENCES) && !defined(_LIBCUDACXX_HAS_NO_VARIADICS)
-    _LIBCUDACXX_INLINE_VISIBILITY void destroy(pointer __p) {__p->~_Tp();}
-};
-
-template <class _Tp, class _Up>
-inline _LIBCUDACXX_INLINE_VISIBILITY
-bool operator==(const allocator<_Tp>&, const allocator<_Up>&) noexcept {return true;}
-
-template <class _Tp, class _Up>
-inline _LIBCUDACXX_INLINE_VISIBILITY
-bool operator!=(const allocator<_Tp>&, const allocator<_Up>&) noexcept {return false;}
-
 template <class _OutputIterator, class _Tp>
 class _LIBCUDACXX_TEMPLATE_VIS raw_storage_iterator
     : public iterator<output_iterator_tag,
@@ -2074,7 +1108,7 @@ class _LIBCUDACXX_TEMPLATE_VIS unique_ptr {
 public:
   typedef _Tp element_type;
   typedef _Dp deleter_type;
-  typedef _LIBCUDACXX_NODEBUG_TYPE typename __pointer_type<_Tp, deleter_type>::type pointer;
+  typedef _LIBCUDACXX_NODEBUG_TYPE typename __pointer<_Tp, deleter_type>::type pointer;
 
   static_assert(!is_rvalue_reference<deleter_type>::value,
                 "the specified deleter type cannot be an rvalue reference");
@@ -2276,7 +1310,7 @@ class _LIBCUDACXX_TEMPLATE_VIS unique_ptr<_Tp[], _Dp> {
 public:
   typedef _Tp element_type;
   typedef _Dp deleter_type;
-  typedef typename __pointer_type<_Tp, deleter_type>::type pointer;
+  typedef typename __pointer<_Tp, deleter_type>::type pointer;
 
 private:
   __compressed_pair<pointer, deleter_type> __ptr_;
@@ -2737,25 +1771,6 @@ public:
         {__process(__p, integral_constant<bool, is_trivially_destructible<_Tp>::value>());}
 };
 
-template <class _Alloc>
-class __allocator_destructor
-{
-    typedef _LIBCUDACXX_NODEBUG_TYPE allocator_traits<_Alloc> __alloc_traits;
-public:
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename __alloc_traits::pointer pointer;
-    typedef _LIBCUDACXX_NODEBUG_TYPE typename __alloc_traits::size_type size_type;
-private:
-    _Alloc& __alloc_;
-    size_type __s_;
-public:
-    _LIBCUDACXX_INLINE_VISIBILITY __allocator_destructor(_Alloc& __a, size_type __s)
-             noexcept
-        : __alloc_(__a), __s_(__s) {}
-    _LIBCUDACXX_INLINE_VISIBILITY
-    void operator()(pointer __p) noexcept
-        {__alloc_traits::deallocate(__alloc_, __p, __s_);}
-};
-
 template <class _InputIterator, class _ForwardIterator>
 _ForwardIterator
 uninitialized_copy(_InputIterator __f, _InputIterator __l, _ForwardIterator __r)
@@ -4904,8 +3919,6 @@ undeclare_reachable(_Tp* __p)
     return static_cast<_Tp*>(__undeclare_reachable(__p));
 }
 
-_LIBCUDACXX_FUNC_VIS void* align(size_t __align, size_t __sz, void*& __ptr, size_t& __space);
-
 // --- Helper for container swap --
 template <typename _Alloc>
 inline _LIBCUDACXX_INLINE_VISIBILITY
diff --git a/libcudacxx/include/cuda/stream_ref b/libcudacxx/include/cuda/stream_ref
index 6f30611202..2e0bd81faa 100644
--- a/libcudacxx/include/cuda/stream_ref
+++ b/libcudacxx/include/cuda/stream_ref
@@ -53,7 +53,6 @@ private:
 #include <cuda/std/cstddef>
 #include <cuda/std/__cuda/api_wrapper.h>
 #include <cuda/std/detail/libcxx/include/__exception/cuda_error.h>
-#include <cuda/std/detail/libcxx/include/__exception/terminate.h>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_CUDA
 
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.tag/allocator_arg.fail.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.tag/allocator_arg.fail.cpp
new file mode 100644
index 0000000000..5a019b6de6
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.tag/allocator_arg.fail.cpp
@@ -0,0 +1,34 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: c++98, c++03
+
+// Before GCC 6, aggregate initialization kicks in.
+// See https://stackoverflow.com/q/41799015/627587.
+// UNSUPPORTED: gcc-5
+
+// <memory>
+
+// struct allocator_arg_t { explicit allocator_arg_t() = default; };
+// const allocator_arg_t allocator_arg = allocator_arg_t();
+
+// This test checks for LWG 2510.
+
+#include <cuda/std/__memory>
+
+cuda::std::allocator_arg_t f()
+{
+  return {};
+} // expected-error 1 {{chosen constructor is explicit in copy-initialization}}
+
+int main(int, char**)
+{
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.tag/allocator_arg.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.tag/allocator_arg.pass.cpp
new file mode 100644
index 0000000000..caa22431e1
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.tag/allocator_arg.pass.cpp
@@ -0,0 +1,27 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+
+// <memory>
+
+// struct allocator_arg_t { explicit allocator_arg_t() = default; };
+// const allocator_arg_t allocator_arg = allocator_arg_t();
+
+#include <cuda/std/__memory>
+
+#include "test_macros.h"
+
+__host__ __device__ void test(cuda::std::allocator_arg_t) {}
+
+int main(int, char**)
+{
+  test(cuda::std::allocator_arg);
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocate_at_least.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocate_at_least.pass.cpp
new file mode 100644
index 0000000000..305f61dc21
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocate_at_least.pass.cpp
@@ -0,0 +1,82 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: c++03, c++11, c++14, c++17, c++20
+
+// <memory>
+
+// template<class Allocator>
+// [[nodiscard]] constexpr allocation_result<typename allocator_traits<Allocator>::pointer>
+//   allocate_at_least(Allocator& a, size_t n);
+
+#include <cuda/std/__memory>
+#include <cuda/std/cassert>
+#include <cuda/std/concepts>
+#include <cuda/std/cstddef>
+
+// check that cuda::std::allocation_result exists and isn't restricted to pointers
+using AllocResult = cuda::std::allocation_result<int>;
+
+template <class T>
+struct no_allocate_at_least
+{
+  using value_type = T;
+  T t;
+
+  constexpr T* allocate(cuda::std::size_t)
+  {
+    return &t;
+  }
+  constexpr void deallocate(T*, cuda::std::size_t) {}
+};
+
+template <class T>
+struct has_allocate_at_least
+{
+  using value_type = T;
+  T t1;
+  T t2;
+
+  constexpr T* allocate(cuda::std::size_t)
+  {
+    return &t1;
+  }
+  constexpr void deallocate(T*, cuda::std::size_t) {}
+  constexpr cuda::std::allocation_result<T*> allocate_at_least(cuda::std::size_t)
+  {
+    return {&t2, 2};
+  }
+};
+
+constexpr bool test()
+{
+  { // check that cuda::std::allocate_at_least forwards to allocator::allocate if no allocate_at_least exists
+    no_allocate_at_least<int> alloc;
+    cuda::std::same_as<cuda::std::allocation_result<int*>> decltype(auto) ret = cuda::std::allocate_at_least(alloc, 1);
+    assert(ret.count == 1);
+    assert(ret.ptr == &alloc.t);
+  }
+
+  { // check that cuda::std::allocate_at_least forwards to allocator::allocate_at_least if allocate_at_least exists
+    has_allocate_at_least<int> alloc;
+    cuda::std::same_as<cuda::std::allocation_result<int*>> decltype(auto) ret = cuda::std::allocate_at_least(alloc, 1);
+    assert(ret.count == 2);
+    assert(ret.ptr == &alloc.t2);
+  }
+
+  return true;
+}
+
+int main(int, char**)
+{
+  test();
+  static_assert(test());
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/allocate.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/allocate.pass.cpp
new file mode 100644
index 0000000000..4813be37ca
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/allocate.pass.cpp
@@ -0,0 +1,64 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     static constexpr pointer allocate(allocator_type& a, size_type n);
+//     ...
+// };
+
+#include <cuda/std/__memory>
+#include <cuda/std/cassert>
+#include <cuda/std/cstdint>
+
+#include "incomplete_type_helper.h"
+#include "test_macros.h"
+
+template <class T>
+struct A
+{
+  typedef T value_type;
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 A() {}
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 value_type* allocate(cuda::std::size_t n)
+  {
+    assert(n == 10);
+    return &storage;
+  }
+
+  value_type storage;
+};
+
+__host__ __device__ TEST_CONSTEXPR_CXX20 bool test()
+{
+  {
+    A<int> a;
+    assert(cuda::std::allocator_traits<A<int>>::allocate(a, 10) == &a.storage);
+  }
+  {
+    typedef A<IncompleteHolder*> Alloc;
+    Alloc a;
+    assert(cuda::std::allocator_traits<Alloc>::allocate(a, 10) == &a.storage);
+  }
+
+  return true;
+}
+
+int main(int, char**)
+{
+  test();
+#if TEST_STD_VER >= 2020
+  static_assert(test());
+#endif // TEST_STD_VER >= 2020
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/allocate.verify.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/allocate.verify.cpp
new file mode 100644
index 0000000000..1192f10474
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/allocate.verify.cpp
@@ -0,0 +1,39 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     static constexpr pointer allocate(allocator_type& a, size_type n);
+//     ...
+// };
+
+// UNSUPPORTED: c++03, c++11, c++14, c++17
+
+#include <cuda/std/__memory>
+#include <cuda/std/cstddef>
+
+template <class T>
+struct A
+{
+  typedef T value_type;
+  value_type* allocate(cuda::std::size_t n);
+  value_type* allocate(cuda::std::size_t n, const void* p);
+};
+
+void f()
+{
+  A<int> a;
+  cuda::std::allocator_traits<A<int>>::allocate(a, 10); // expected-warning {{ignoring return value of function declared
+                                                        // with 'nodiscard' attribute}}
+  cuda::std::allocator_traits<A<int>>::allocate(a, 10, nullptr); // expected-warning {{ignoring return value of function
+                                                                 // declared with 'nodiscard' attribute}}
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/allocate_hint.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/allocate_hint.pass.cpp
new file mode 100644
index 0000000000..00b5f6baca
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/allocate_hint.pass.cpp
@@ -0,0 +1,88 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     static constexpr pointer allocate(allocator_type& a, size_type n, const_void_pointer hint);
+//     ...
+// };
+
+#include <cuda/std/__memory>
+#include <cuda/std/cassert>
+#include <cuda/std/cstdint>
+
+#include "incomplete_type_helper.h"
+#include "test_macros.h"
+
+template <class T>
+struct A
+{
+  typedef T value_type;
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 A() {}
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 value_type* allocate(cuda::std::size_t n)
+  {
+    assert(n == 10);
+    return &storage;
+  }
+
+  value_type storage;
+};
+
+template <class T>
+struct B
+{
+  typedef T value_type;
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 value_type* allocate(cuda::std::size_t n, const void* p)
+  {
+    assert(n == 11);
+    assert(p == nullptr);
+    return &storage;
+  }
+
+  value_type storage;
+};
+
+__host__ __device__ TEST_CONSTEXPR_CXX20 bool test()
+{
+  {
+    A<int> a;
+    assert(cuda::std::allocator_traits<A<int>>::allocate(a, 10, nullptr) == &a.storage);
+  }
+  {
+    typedef A<IncompleteHolder*> Alloc;
+    Alloc a;
+    assert(cuda::std::allocator_traits<Alloc>::allocate(a, 10, nullptr) == &a.storage);
+  }
+  {
+    B<int> b;
+    assert(cuda::std::allocator_traits<B<int>>::allocate(b, 11, nullptr) == &b.storage);
+  }
+  {
+    typedef B<IncompleteHolder*> Alloc;
+    Alloc b;
+    assert(cuda::std::allocator_traits<Alloc>::allocate(b, 11, nullptr) == &b.storage);
+  }
+
+  return true;
+}
+
+int main(int, char**)
+{
+  test();
+#if TEST_STD_VER >= 2020
+  static_assert(test());
+#endif // TEST_STD_VER >= 2020
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/construct.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/construct.pass.cpp
new file mode 100644
index 0000000000..3a352b5eec
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/construct.pass.cpp
@@ -0,0 +1,200 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     template <class Ptr, class... Args>
+//     static constexpr void construct(allocator_type& a, Ptr p, Args&&... args);
+//     ...
+// };
+
+#include <cuda/std/__memory>
+#include <cuda/std/__new>
+#include <cuda/std/cassert>
+#include <cuda/std/type_traits>
+#include <cuda/std/utility>
+
+#include "incomplete_type_helper.h"
+#include "test_macros.h"
+
+template <class T>
+struct A
+{
+  typedef T value_type;
+};
+
+template <class T>
+struct Alloc
+{
+  typedef T value_type;
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 Alloc() {}
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 value_type* allocate(cuda::std::size_t n)
+  {
+    assert(n == 1);
+    return &storage;
+  }
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 void deallocate(value_type*, cuda::std::size_t) {}
+
+  value_type storage;
+};
+
+template <class T>
+struct B
+{
+  typedef T value_type;
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 B(int& count)
+      : count_(count)
+  {}
+
+  template <class U, class... Args>
+  __host__ __device__ TEST_CONSTEXPR_CXX20 void construct(U* p, Args&&... args)
+  {
+    ++count_;
+#if TEST_STD_VER >= 2020
+    cuda::std::construct_at(p, cuda::std::forward<Args>(args)...);
+#else // ^^^ C++20 ^^^ / vvv C++17 vvv
+    ::new ((void*) p) U(cuda::std::forward<Args>(args)...);
+#endif // TEST_STD_VER >= 2020
+  }
+
+  int& count_;
+};
+
+struct A0
+{
+  A0() = default;
+  __host__ __device__ TEST_CONSTEXPR_CXX20 A0(int* count)
+  {
+    ++*count;
+  }
+};
+
+struct A1
+{
+  A1() = default;
+  __host__ __device__ TEST_CONSTEXPR_CXX20 A1(int* count, char c)
+  {
+    assert(c == 'c');
+    ++*count;
+  }
+};
+
+struct A2
+{
+  A2() = default;
+  __host__ __device__ TEST_CONSTEXPR_CXX20 A2(int* count, char c, int i)
+  {
+    assert(c == 'd');
+    assert(i == 5);
+    ++*count;
+  }
+};
+
+__host__ __device__ TEST_CONSTEXPR_CXX20 bool test()
+{
+  {
+    int A0_count = 0;
+    A<A0> a;
+    Alloc<A0> alloc;
+    A0* a0 = alloc.allocate(1);
+    assert(A0_count == 0);
+    cuda::std::allocator_traits<A<A0>>::construct(a, a0, &A0_count);
+    assert(A0_count == 1);
+    alloc.deallocate(a0, 1);
+  }
+  {
+    int A1_count = 0;
+    A<A1> a;
+    Alloc<A1> alloc;
+    A1* a1 = alloc.allocate(1);
+    assert(A1_count == 0);
+    cuda::std::allocator_traits<A<A1>>::construct(a, a1, &A1_count, 'c');
+    assert(A1_count == 1);
+    alloc.deallocate(a1, 1);
+  }
+  {
+    int A2_count = 0;
+    A<A2> a;
+    Alloc<A2> alloc;
+    A2* a2 = alloc.allocate(1);
+    assert(A2_count == 0);
+    cuda::std::allocator_traits<A<A2>>::construct(a, a2, &A2_count, 'd', 5);
+    assert(A2_count == 1);
+    alloc.deallocate(a2, 1);
+  }
+  {
+    typedef IncompleteHolder* VT;
+    typedef A<VT> Alloc2;
+    Alloc2 a;
+    Alloc<VT> alloc;
+    VT* vt = alloc.allocate(1);
+    cuda::std::allocator_traits<Alloc2>::construct(a, vt, nullptr);
+    alloc.deallocate(vt, 1);
+  }
+
+  {
+    int A0_count    = 0;
+    int b_construct = 0;
+    B<A0> b(b_construct);
+    Alloc<A0> alloc;
+    A0* a0 = alloc.allocate(1);
+    assert(A0_count == 0);
+    assert(b_construct == 0);
+    cuda::std::allocator_traits<B<A0>>::construct(b, a0, &A0_count);
+    assert(A0_count == 1);
+    assert(b_construct == 1);
+    alloc.deallocate(a0, 1);
+  }
+  {
+    int A1_count    = 0;
+    int b_construct = 0;
+    B<A1> b(b_construct);
+    Alloc<A1> alloc;
+    A1* a1 = alloc.allocate(1);
+    assert(A1_count == 0);
+    assert(b_construct == 0);
+    cuda::std::allocator_traits<B<A1>>::construct(b, a1, &A1_count, 'c');
+    assert(A1_count == 1);
+    assert(b_construct == 1);
+    alloc.deallocate(a1, 1);
+  }
+  {
+    int A2_count    = 0;
+    int b_construct = 0;
+    B<A2> b(b_construct);
+    Alloc<A2> alloc;
+    A2* a2 = alloc.allocate(1);
+    assert(A2_count == 0);
+    assert(b_construct == 0);
+    cuda::std::allocator_traits<B<A2>>::construct(b, a2, &A2_count, 'd', 5);
+    assert(A2_count == 1);
+    assert(b_construct == 1);
+    alloc.deallocate(a2, 1);
+  }
+
+  return true;
+}
+
+int main(int, char**)
+{
+  test();
+
+#if TEST_STD_VER >= 2020
+  static_assert(test());
+#endif // TEST_STD_VER >= 2020
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/deallocate.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/deallocate.pass.cpp
new file mode 100644
index 0000000000..89b4a7078d
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/deallocate.pass.cpp
@@ -0,0 +1,73 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     static constexpr void deallocate(allocator_type& a, pointer p, size_type n);
+//     ...
+// };
+
+#include <cuda/std/__memory>
+#include <cuda/std/cassert>
+#include <cuda/std/cstddef>
+
+#include "incomplete_type_helper.h"
+#include "test_macros.h"
+
+template <class T>
+struct A
+{
+  typedef T value_type;
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 A(int& called)
+      : called_(called)
+  {}
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 void deallocate(value_type* p, cuda::std::size_t n)
+  {
+    assert(p == &storage);
+    assert(n == 10);
+    ++called_;
+  }
+
+  int& called_;
+
+  value_type storage;
+};
+
+__host__ __device__ TEST_CONSTEXPR_CXX20 bool test()
+{
+  {
+    int called = 0;
+    A<int> a(called);
+    cuda::std::allocator_traits<A<int>>::deallocate(a, &a.storage, 10);
+    assert(called == 1);
+  }
+  {
+    int called = 0;
+    typedef A<IncompleteHolder*> Alloc;
+    Alloc a(called);
+    cuda::std::allocator_traits<Alloc>::deallocate(a, &a.storage, 10);
+    assert(called == 1);
+  }
+
+  return true;
+}
+
+int main(int, char**)
+{
+  test();
+#if TEST_STD_VER >= 2020
+  static_assert(test());
+#endif // TEST_STD_VER >= 2020
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/destroy.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/destroy.pass.cpp
new file mode 100644
index 0000000000..26fd3cf1a3
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/destroy.pass.cpp
@@ -0,0 +1,156 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     template <class Ptr>
+//     static constexpr void destroy(allocator_type& a, Ptr p);
+//     ...
+// };
+
+// Currently no suppport for std::allocator
+
+#include <cuda/std/__memory>
+#include <cuda/std/cassert>
+#include <cuda/std/cstddef>
+
+#include "incomplete_type_helper.h"
+#include "test_macros.h"
+
+template <class T>
+struct NoDestroy
+{
+  typedef T value_type;
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 T* allocate(cuda::std::size_t n)
+  {
+    return cuda::std::allocator<T>().allocate(n);
+  }
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 void deallocate(T* p, cuda::std::size_t n)
+  {
+    return cuda::std::allocator<T>().deallocate(p, n);
+  }
+};
+
+template <class T>
+struct CountDestroy
+{
+  __host__ __device__ TEST_CONSTEXPR explicit CountDestroy(int* counter)
+      : counter_(counter)
+  {}
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 ~CountDestroy() {}
+
+  typedef T value_type;
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 T* allocate(cuda::std::size_t n)
+  {
+    return &storage;
+  }
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 void deallocate(T* p, cuda::std::size_t n) {}
+
+  template <class U, class... Args>
+  __host__ __device__ TEST_CONSTEXPR_CXX20 void construct(U* p, Args&&... args)
+  {
+    assert(p == nullptr);
+    cuda::std::__construct_at(&storage, cuda::std::forward<Args>(args)...);
+  }
+
+  template <class U>
+  __host__ __device__ TEST_CONSTEXPR_CXX20 void destroy(U* p)
+  {
+    assert(p == nullptr);
+    ++*counter_;
+    storage.~U();
+  }
+
+  int* counter_;
+  union
+  {
+    char dummy_{};
+    value_type storage;
+  };
+};
+
+struct CountDestructor
+{
+  __host__ __device__ TEST_CONSTEXPR explicit CountDestructor(int* counter)
+      : counter_(counter)
+  {}
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 ~CountDestructor()
+  {
+    ++*counter_;
+  }
+
+  int* counter_;
+};
+
+__host__ __device__ TEST_CONSTEXPR_CXX20 bool test()
+{
+  if (!cuda::std::__libcpp_is_constant_evaluated())
+  {
+    using Alloc     = NoDestroy<CountDestructor>;
+    int destructors = 0;
+    Alloc alloc;
+    CountDestructor* pool = cuda::std::allocator_traits<Alloc>::allocate(alloc, 1);
+
+    cuda::std::allocator_traits<Alloc>::construct(alloc, pool, &destructors);
+    assert(destructors == 0);
+
+    cuda::std::allocator_traits<Alloc>::destroy(alloc, pool);
+    assert(destructors == 1);
+
+    cuda::std::allocator_traits<Alloc>::deallocate(alloc, pool, 1);
+  }
+#if !defined(TEST_COMPILER_MSVC) && TEST_STD_VER >= 2020 // incomplete type not allowed
+  if (!cuda::std::__libcpp_is_constant_evaluated())
+  {
+    typedef IncompleteHolder* T;
+    typedef NoDestroy<T> Alloc;
+    Alloc alloc;
+    T* pool = cuda::std::allocator_traits<Alloc>::allocate(alloc, 1);
+    cuda::std::allocator_traits<Alloc>::construct(alloc, pool, nullptr);
+    cuda::std::allocator_traits<Alloc>::destroy(alloc, pool);
+    cuda::std::allocator_traits<Alloc>::deallocate(alloc, pool, 1);
+  }
+#endif // !defined(TEST_COMPILER_MSVC) && TEST_STD_VER >= 2020
+  {
+    using Alloc            = CountDestroy<CountDestructor>;
+    int destroys_called    = 0;
+    int destructors_called = 0;
+    Alloc alloc(&destroys_called);
+    CountDestructor* fake_ptr = nullptr;
+
+    cuda::std::allocator_traits<Alloc>::construct(alloc, fake_ptr, &destructors_called);
+    assert(destroys_called == 0);
+    assert(destructors_called == 0);
+
+    cuda::std::allocator_traits<Alloc>::destroy(alloc, fake_ptr);
+    assert(destroys_called == 1);
+    assert(destructors_called == 1);
+
+    cuda::std::allocator_traits<Alloc>::deallocate(alloc, fake_ptr, 1);
+  }
+  return true;
+}
+
+int main(int, char**)
+{
+  test();
+#if TEST_STD_VER >= 2020
+  static_assert(test());
+#endif // TEST_STD_VER >= 2020
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/incomplete_type_helper.h b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/incomplete_type_helper.h
new file mode 100644
index 0000000000..f0d2049b06
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/incomplete_type_helper.h
@@ -0,0 +1,19 @@
+#ifndef TEST_INCOMPLETE_TYPE_HELPER_H
+#define TEST_INCOMPLETE_TYPE_HELPER_H
+
+#include "min_allocator.h"
+
+namespace NS
+{
+struct Incomplete;
+}
+
+template <class T>
+struct Holder
+{
+  T value;
+};
+
+typedef Holder<NS::Incomplete> IncompleteHolder;
+
+#endif
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/max_size.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/max_size.pass.cpp
new file mode 100644
index 0000000000..5bfa21f77c
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/max_size.pass.cpp
@@ -0,0 +1,88 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     static constexpr size_type max_size(const allocator_type& a) noexcept;
+//     ...
+// };
+
+#include <cuda/std/__memory>
+#include <cuda/std/cassert>
+#include <cuda/std/limits>
+#include <cuda/std/type_traits>
+
+#include "incomplete_type_helper.h"
+#include "test_macros.h"
+
+template <class T>
+struct A
+{
+  typedef T value_type;
+};
+
+template <class T>
+struct B
+{
+  typedef T value_type;
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 cuda::std::size_t max_size() const
+  {
+    return 100;
+  }
+};
+
+__host__ __device__ TEST_CONSTEXPR_CXX20 bool test()
+{
+  {
+    B<int> b;
+    assert(cuda::std::allocator_traits<B<int>>::max_size(b) == 100);
+  }
+  {
+    const B<int> b = {};
+    assert(cuda::std::allocator_traits<B<int>>::max_size(b) == 100);
+  }
+  {
+    typedef IncompleteHolder* VT;
+    typedef B<VT> Alloc;
+    Alloc a;
+    assert(cuda::std::allocator_traits<Alloc>::max_size(a) == 100);
+  }
+  {
+    A<int> a;
+    assert(cuda::std::allocator_traits<A<int>>::max_size(a)
+           == cuda::std::numeric_limits<cuda::std::size_t>::max() / sizeof(int));
+  }
+  {
+    const A<int> a = {};
+    assert(cuda::std::allocator_traits<A<int>>::max_size(a)
+           == cuda::std::numeric_limits<cuda::std::size_t>::max() / sizeof(int));
+  }
+  {
+    cuda::std::allocator<int> a;
+    static_assert(noexcept(cuda::std::allocator_traits<cuda::std::allocator<int>>::max_size(a)) == true, "");
+    unused(a);
+  }
+
+  return true;
+}
+
+int main(int, char**)
+{
+  test();
+
+#if TEST_STD_VER >= 2020
+  static_assert(test());
+#endif // TEST_STD_VER >= 2020
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/select_on_container_copy_construction.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/select_on_container_copy_construction.pass.cpp
new file mode 100644
index 0000000000..3c89366a74
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.members/select_on_container_copy_construction.pass.cpp
@@ -0,0 +1,88 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     static constexpr allocator_type
+//         select_on_container_copy_construction(const allocator_type& a);
+//     ...
+// };
+
+#include <cuda/std/__memory>
+#include <cuda/std/cassert>
+#include <cuda/std/type_traits>
+
+#include "incomplete_type_helper.h"
+#include "test_macros.h"
+
+template <class T>
+struct A
+{
+  typedef T value_type;
+  int id;
+  __host__ __device__ TEST_CONSTEXPR_CXX20 explicit A(int i = 0)
+      : id(i)
+  {}
+};
+
+template <class T>
+struct B
+{
+  typedef T value_type;
+
+  int id;
+  __host__ __device__ TEST_CONSTEXPR_CXX20 explicit B(int i = 0)
+      : id(i)
+  {}
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 B select_on_container_copy_construction() const
+  {
+    return B(100);
+  }
+};
+
+__host__ __device__ TEST_CONSTEXPR_CXX20 bool test()
+{
+  {
+    A<int> a;
+    assert(cuda::std::allocator_traits<A<int>>::select_on_container_copy_construction(a).id == 0);
+  }
+  {
+    const A<int> a(0);
+    assert(cuda::std::allocator_traits<A<int>>::select_on_container_copy_construction(a).id == 0);
+  }
+  {
+    typedef IncompleteHolder* VT;
+    typedef A<VT> Alloc;
+    Alloc a;
+    assert(cuda::std::allocator_traits<Alloc>::select_on_container_copy_construction(a).id == 0);
+  }
+  {
+    B<int> b;
+    assert(cuda::std::allocator_traits<B<int>>::select_on_container_copy_construction(b).id == 100);
+  }
+  {
+    const B<int> b(0);
+    assert(cuda::std::allocator_traits<B<int>>::select_on_container_copy_construction(b).id == 100);
+  }
+
+  return true;
+}
+
+int main(int, char**)
+{
+  test();
+#if TEST_STD_VER >= 2020
+  static_assert(test());
+#endif // TEST_STD_VER >= 2020
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/const_pointer.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/const_pointer.pass.cpp
new file mode 100644
index 0000000000..27a63e6805
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/const_pointer.pass.cpp
@@ -0,0 +1,75 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     typedef Alloc::const_pointer
+//           | pointer_traits<pointer>::rebind<const value_type>
+//     ...
+// };
+
+#include <cuda/std/__memory>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+template <class T>
+struct Ptr
+{};
+
+template <class T>
+struct A
+{
+  typedef T value_type;
+  typedef Ptr<T> pointer;
+};
+
+template <class T>
+struct B
+{
+  typedef T value_type;
+};
+
+template <class T>
+struct CPtr
+{};
+
+template <class T>
+struct C
+{
+  typedef T value_type;
+  typedef CPtr<T> pointer;
+  typedef CPtr<const T> const_pointer;
+};
+
+#if !defined(TEST_COMPILER_MSVC_2017) // const_pointer is inaccessible
+template <class T>
+struct D
+{
+  typedef T value_type;
+
+private:
+  typedef void const_pointer;
+};
+#endif // !TEST_COMPILER_MSVC_2017
+
+int main(int, char**)
+{
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<A<char>>::const_pointer, Ptr<const char>>::value), "");
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<B<char>>::const_pointer, const char*>::value), "");
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<C<char>>::const_pointer, CPtr<const char>>::value), "");
+#if !defined(TEST_COMPILER_MSVC_2017) // const_pointer is inaccessible
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<D<char>>::const_pointer, const char*>::value), "");
+#endif // !TEST_COMPILER_MSVC_2017
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/const_void_pointer.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/const_void_pointer.pass.cpp
new file mode 100644
index 0000000000..9a668f7744
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/const_void_pointer.pass.cpp
@@ -0,0 +1,76 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     typedef Alloc::const_void_pointer
+//           | pointer_traits<pointer>::rebind<const void>
+//                                          const_void_pointer;
+//     ...
+// };
+
+#include <cuda/std/__memory>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+template <class T>
+struct Ptr
+{};
+
+template <class T>
+struct A
+{
+  typedef T value_type;
+  typedef Ptr<T> pointer;
+};
+
+template <class T>
+struct B
+{
+  typedef T value_type;
+};
+
+template <class T>
+struct CPtr
+{};
+
+template <class T>
+struct C
+{
+  typedef T value_type;
+  typedef CPtr<const void> const_void_pointer;
+};
+
+#if !defined(TEST_COMPILER_MSVC_2017) // const_void_pointer is inaccessible
+template <class T>
+struct D
+{
+  typedef T value_type;
+
+private:
+  typedef int const_void_pointer;
+};
+#endif // !TEST_COMPILER_MSVC_2017
+
+int main(int, char**)
+{
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<A<char>>::const_void_pointer, Ptr<const void>>::value),
+                "");
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<B<char>>::const_void_pointer, const void*>::value), "");
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<C<char>>::const_void_pointer, CPtr<const void>>::value),
+                "");
+#if !defined(TEST_COMPILER_MSVC_2017) // const_void_pointer is inaccessible
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<D<char>>::const_void_pointer, const void*>::value), "");
+#endif // !TEST_COMPILER_MSVC_2017
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/difference_type.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/difference_type.pass.cpp
new file mode 100644
index 0000000000..ed278b610e
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/difference_type.pass.cpp
@@ -0,0 +1,91 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     typedef Alloc::difference_type
+//           | pointer_traits<pointer>::difference_type         difference_type;
+//     ...
+// };
+
+#include <cuda/std/__memory>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+template <class T>
+struct A
+{
+  typedef T value_type;
+  typedef short difference_type;
+};
+
+template <class T>
+struct B
+{
+  typedef T value_type;
+};
+
+template <class T>
+struct C
+{
+  typedef T value_type;
+  struct pointer
+  {};
+  struct const_pointer
+  {};
+  struct void_pointer
+  {};
+  struct const_void_pointer
+  {};
+};
+
+#if !defined(TEST_COMPILER_MSVC_2017) // difference_type is inaccessible
+template <class T>
+struct D
+{
+  typedef T value_type;
+
+private:
+  typedef void difference_type;
+};
+#endif // !TEST_COMPILER_MSVC_2017
+
+namespace cuda
+{
+namespace std
+{
+
+template <>
+struct pointer_traits<C<char>::pointer>
+{
+  typedef C<char>::pointer pointer;
+  typedef char element_type;
+  typedef signed char difference_type;
+};
+
+} // namespace std
+} // namespace cuda
+
+int main(int, char**)
+{
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<A<char>>::difference_type, short>::value), "");
+  static_assert(
+    (cuda::std::is_same<cuda::std::allocator_traits<B<char>>::difference_type, cuda::std::ptrdiff_t>::value), "");
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<C<char>>::difference_type, signed char>::value), "");
+#if !defined(TEST_COMPILER_MSVC_2017) // difference_type is inaccessible
+  static_assert(
+    (cuda::std::is_same<cuda::std::allocator_traits<D<char>>::difference_type, cuda::std::ptrdiff_t>::value), "");
+#endif // !TEST_COMPILER_MSVC_2017
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/is_always_equal.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/is_always_equal.pass.cpp
new file mode 100644
index 0000000000..e614ab17f6
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/is_always_equal.pass.cpp
@@ -0,0 +1,63 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//   typedef Alloc::is_always_equal
+//         | is_empty                     is_always_equal;
+//     ...
+// };
+
+#include <cuda/std/__memory>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+template <class T>
+struct A
+{
+  typedef T value_type;
+  typedef cuda::std::true_type is_always_equal;
+};
+
+template <class T>
+struct B
+{
+  typedef T value_type;
+};
+
+template <class T>
+struct C
+{
+  typedef T value_type;
+  int not_empty_; // some random member variable
+};
+
+int main(int, char**)
+{
+  static_assert(
+    (cuda::std::is_same<cuda::std::allocator_traits<A<char>>::is_always_equal, cuda::std::true_type>::value), "");
+  static_assert(
+    (cuda::std::is_same<cuda::std::allocator_traits<B<char>>::is_always_equal, cuda::std::true_type>::value), "");
+  static_assert(
+    (cuda::std::is_same<cuda::std::allocator_traits<C<char>>::is_always_equal, cuda::std::false_type>::value), "");
+
+  static_assert(
+    (cuda::std::is_same<cuda::std::allocator_traits<A<const char>>::is_always_equal, cuda::std::true_type>::value), "");
+  static_assert(
+    (cuda::std::is_same<cuda::std::allocator_traits<B<const char>>::is_always_equal, cuda::std::true_type>::value), "");
+  static_assert(
+    (cuda::std::is_same<cuda::std::allocator_traits<C<const char>>::is_always_equal, cuda::std::false_type>::value),
+    "");
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/pointer.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/pointer.pass.cpp
new file mode 100644
index 0000000000..d41f5574bd
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/pointer.pass.cpp
@@ -0,0 +1,61 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     typedef Alloc::pointer | value_type* pointer;
+//     ...
+// };
+
+#include <cuda/std/__memory>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+template <class T>
+struct Ptr
+{};
+
+template <class T>
+struct A
+{
+  typedef T value_type;
+  typedef Ptr<T> pointer;
+};
+
+template <class T>
+struct B
+{
+  typedef T value_type;
+};
+
+#if !defined(TEST_COMPILER_MSVC_2017) // pointer is inaccessible
+template <class T>
+struct C
+{
+  typedef T value_type;
+
+private:
+  typedef void pointer;
+};
+#endif // !TEST_COMPILER_MSVC_2017
+
+int main(int, char**)
+{
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<A<char>>::pointer, Ptr<char>>::value), "");
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<B<char>>::pointer, char*>::value), "");
+#if !defined(TEST_COMPILER_MSVC_2017) // pointer is inaccessible
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<C<char>>::pointer, char*>::value), "");
+#endif // !TEST_COMPILER_MSVC_2017
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/propagate_on_container_copy_assignment.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/propagate_on_container_copy_assignment.pass.cpp
new file mode 100644
index 0000000000..b9361a8feb
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/propagate_on_container_copy_assignment.pass.cpp
@@ -0,0 +1,64 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     typedef Alloc::propagate_on_container_copy_assignment
+//           | false_type                   propagate_on_container_copy_assignment;
+//     ...
+// };
+
+#include <cuda/std/__memory>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+template <class T>
+struct A
+{
+  typedef T value_type;
+  typedef cuda::std::true_type propagate_on_container_copy_assignment;
+};
+
+template <class T>
+struct B
+{
+  typedef T value_type;
+};
+
+#if !defined(TEST_COMPILER_MSVC_2017) // propagate_on_container_copy_assignment is inaccessible
+template <class T>
+struct C
+{
+  typedef T value_type;
+
+private:
+  typedef cuda::std::true_type propagate_on_container_copy_assignment;
+};
+#endif // !TEST_COMPILER_MSVC_2017
+
+int main(int, char**)
+{
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<A<char>>::propagate_on_container_copy_assignment,
+                                    cuda::std::true_type>::value),
+                "");
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<B<char>>::propagate_on_container_copy_assignment,
+                                    cuda::std::false_type>::value),
+                "");
+#if !defined(TEST_COMPILER_MSVC_2017) // propagate_on_container_copy_assignment is inaccessible
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<C<char>>::propagate_on_container_copy_assignment,
+                                    cuda::std::false_type>::value),
+                "");
+#endif // !TEST_COMPILER_MSVC_2017
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/propagate_on_container_move_assignment.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/propagate_on_container_move_assignment.pass.cpp
new file mode 100644
index 0000000000..c91b2f6c6b
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/propagate_on_container_move_assignment.pass.cpp
@@ -0,0 +1,64 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     typedef Alloc::propagate_on_container_move_assignment
+//           | false_type                   propagate_on_container_move_assignment;
+//     ...
+// };
+
+#include <cuda/std/__memory>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+template <class T>
+struct A
+{
+  typedef T value_type;
+  typedef cuda::std::true_type propagate_on_container_move_assignment;
+};
+
+template <class T>
+struct B
+{
+  typedef T value_type;
+};
+
+#if !defined(TEST_COMPILER_MSVC_2017) // propagate_on_container_move_assignment is inaccessible
+template <class T>
+struct C
+{
+  typedef T value_type;
+
+private:
+  typedef cuda::std::true_type propagate_on_container_move_assignment;
+};
+#endif // !TEST_COMPILER_MSVC_2017
+
+int main(int, char**)
+{
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<A<char>>::propagate_on_container_move_assignment,
+                                    cuda::std::true_type>::value),
+                "");
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<B<char>>::propagate_on_container_move_assignment,
+                                    cuda::std::false_type>::value),
+                "");
+#if !defined(TEST_COMPILER_MSVC_2017) // propagate_on_container_move_assignment is inaccessible
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<C<char>>::propagate_on_container_move_assignment,
+                                    cuda::std::false_type>::value),
+                "");
+#endif // !TEST_COMPILER_MSVC_2017
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/propagate_on_container_swap.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/propagate_on_container_swap.pass.cpp
new file mode 100644
index 0000000000..872662ddb1
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/propagate_on_container_swap.pass.cpp
@@ -0,0 +1,64 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     typedef Alloc::propagate_on_container_swap
+//           | false_type                   propagate_on_container_swap;
+//     ...
+// };
+
+#include <cuda/std/__memory>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+template <class T>
+struct A
+{
+  typedef T value_type;
+  typedef cuda::std::true_type propagate_on_container_swap;
+};
+
+template <class T>
+struct B
+{
+  typedef T value_type;
+};
+
+#if !defined(TEST_COMPILER_MSVC_2017) // propagate_on_container_swap is inaccessible
+template <class T>
+struct C
+{
+  typedef T value_type;
+
+private:
+  typedef cuda::std::true_type propagate_on_container_swap;
+};
+#endif // !TEST_COMPILER_MSVC_2017
+
+int main(int, char**)
+{
+  static_assert(
+    (cuda::std::is_same<cuda::std::allocator_traits<A<char>>::propagate_on_container_swap, cuda::std::true_type>::value),
+    "");
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<B<char>>::propagate_on_container_swap,
+                                    cuda::std::false_type>::value),
+                "");
+#if !defined(TEST_COMPILER_MSVC_2017) // propagate_on_container_swap is inaccessible
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<C<char>>::propagate_on_container_swap,
+                                    cuda::std::false_type>::value),
+                "");
+#endif // !TEST_COMPILER_MSVC_2017
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/rebind_alloc.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/rebind_alloc.pass.cpp
new file mode 100644
index 0000000000..47f4c17a1d
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/rebind_alloc.pass.cpp
@@ -0,0 +1,124 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     template <class T> using rebind_alloc  = Alloc::rebind<U>::other | Alloc<T, Args...>;
+//     ...
+// };
+
+#include <cuda/std/__memory>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+template <class T>
+struct ReboundA
+{};
+
+template <class T>
+struct A
+{
+  typedef T value_type;
+
+  template <class U>
+  struct rebind
+  {
+    typedef ReboundA<U> other;
+  };
+};
+
+template <class T, class U>
+struct ReboundB
+{};
+
+template <class T, class U>
+struct B
+{
+  typedef T value_type;
+
+  template <class V>
+  struct rebind
+  {
+    typedef ReboundB<V, U> other;
+  };
+};
+
+template <class T>
+struct C
+{
+  typedef T value_type;
+};
+
+template <class T, class U>
+struct D
+{
+  typedef T value_type;
+};
+
+template <class T>
+struct E
+{
+  typedef T value_type;
+
+  template <class U>
+  struct rebind
+  {
+    typedef ReboundA<U> otter;
+  };
+};
+
+#if !defined(TEST_COMPILER_MSVC_2017) // rebind is inaccessible
+template <class T>
+struct F
+{
+  typedef T value_type;
+
+private:
+  template <class>
+  struct rebind
+  {
+    typedef void other;
+  };
+};
+
+template <class T>
+struct G
+{
+  typedef T value_type;
+  template <class>
+  struct rebind
+  {
+  private:
+    typedef void other;
+  };
+};
+#endif // !TEST_COMPILER_MSVC_2017
+
+int main(int, char**)
+{
+  static_assert(
+    (cuda::std::is_same<cuda::std::allocator_traits<A<char>>::rebind_alloc<double>, ReboundA<double>>::value), "");
+  static_assert(
+    (cuda::std::is_same<cuda::std::allocator_traits<B<int, char>>::rebind_alloc<double>, ReboundB<double, char>>::value),
+    "");
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<C<char>>::rebind_alloc<double>, C<double>>::value), "");
+  static_assert(
+    (cuda::std::is_same<cuda::std::allocator_traits<D<int, char>>::rebind_alloc<double>, D<double, char>>::value), "");
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<E<char>>::rebind_alloc<double>, E<double>>::value), "");
+#if !defined(TEST_COMPILER_MSVC_2017) // rebind is inaccessible
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<F<char>>::rebind_alloc<double>, F<double>>::value), "");
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<G<char>>::rebind_alloc<double>, G<double>>::value), "");
+#endif // !TEST_COMPILER_MSVC_2017
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/size_type.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/size_type.pass.cpp
new file mode 100644
index 0000000000..8828e4269b
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/size_type.pass.cpp
@@ -0,0 +1,89 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     typedef Alloc::size_type | size_t    size_type;
+//     ...
+// };
+
+#include <cuda/std/__memory>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+template <class T>
+struct A
+{
+  typedef T value_type;
+  typedef unsigned short size_type;
+};
+
+template <class T>
+struct B
+{
+  typedef T value_type;
+};
+
+template <class T>
+struct C
+{
+  typedef T value_type;
+  struct pointer
+  {};
+  struct const_pointer
+  {};
+  struct void_pointer
+  {};
+  struct const_void_pointer
+  {};
+};
+
+#if !defined(TEST_COMPILER_MSVC_2017) // rebind is inaccessible
+template <class T>
+struct D
+{
+  typedef T value_type;
+  typedef short difference_type;
+
+private:
+  typedef void size_type;
+};
+#endif // !TEST_COMPILER_MSVC_2017
+
+namespace cuda
+{
+namespace std
+{
+
+template <>
+struct pointer_traits<C<char>::pointer>
+{
+  typedef signed char difference_type;
+};
+
+} // namespace std
+} // namespace cuda
+
+int main(int, char**)
+{
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<A<char>>::size_type, unsigned short>::value), "");
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<B<char>>::size_type,
+                                    cuda::std::make_unsigned<cuda::std::ptrdiff_t>::type>::value),
+                "");
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<C<char>>::size_type, unsigned char>::value), "");
+#if !defined(TEST_COMPILER_MSVC_2017) // rebind is inaccessible
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<D<char>>::size_type, unsigned short>::value), "");
+#endif // !TEST_COMPILER_MSVC_2017
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/void_pointer.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/void_pointer.pass.cpp
new file mode 100644
index 0000000000..ed12e38d53
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator.traits.types/void_pointer.pass.cpp
@@ -0,0 +1,75 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     typedef Alloc::void_pointer
+//           | pointer_traits<pointer>::rebind<void>
+//                                          void_pointer;
+//     ...
+// };
+
+#include <cuda/std/__memory>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+template <class T>
+struct Ptr
+{};
+
+template <class T>
+struct A
+{
+  typedef T value_type;
+  typedef Ptr<T> pointer;
+};
+
+template <class T>
+struct B
+{
+  typedef T value_type;
+};
+
+template <class T>
+struct CPtr
+{};
+
+template <class T>
+struct C
+{
+  typedef T value_type;
+  typedef CPtr<void> void_pointer;
+};
+
+#if !defined(TEST_COMPILER_MSVC_2017) // void_pointer is inaccessible
+template <class T>
+struct D
+{
+  typedef T value_type;
+
+private:
+  typedef void void_pointer;
+};
+#endif // !TEST_COMPILER_MSVC_2017
+
+int main(int, char**)
+{
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<A<char>>::void_pointer, Ptr<void>>::value), "");
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<B<char>>::void_pointer, void*>::value), "");
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<C<char>>::void_pointer, CPtr<void>>::value), "");
+#if !defined(TEST_COMPILER_MSVC_2017) // void_pointer is inaccessible
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<D<char>>::void_pointer, void*>::value), "");
+#endif // !TEST_COMPILER_MSVC_2017
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator_type.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator_type.pass.cpp
new file mode 100644
index 0000000000..b6b8ed322a
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/allocator_type.pass.cpp
@@ -0,0 +1,35 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     typedef Alloc allocator_type;
+//     ...
+// };
+
+#include <cuda/std/__memory>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+template <class T>
+struct A
+{
+  typedef T value_type;
+};
+
+int main(int, char**)
+{
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<A<char>>::allocator_type, A<char>>::value), "");
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/rebind_traits.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/rebind_traits.pass.cpp
new file mode 100644
index 0000000000..73e5edb8dc
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/rebind_traits.pass.cpp
@@ -0,0 +1,99 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     template <class T> using rebind_traits = allocator_traits<rebind_alloc<T>>;
+//     ...
+// };
+
+#include <cuda/std/__memory>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+template <class T>
+struct ReboundA
+{};
+
+template <class T>
+struct A
+{
+  typedef T value_type;
+
+  template <class U>
+  struct rebind
+  {
+    typedef ReboundA<U> other;
+  };
+};
+
+template <class T, class U>
+struct ReboundB
+{};
+
+template <class T, class U>
+struct B
+{
+  typedef T value_type;
+
+  template <class V>
+  struct rebind
+  {
+    typedef ReboundB<V, U> other;
+  };
+};
+
+template <class T>
+struct C
+{
+  typedef T value_type;
+};
+
+template <class T, class U>
+struct D
+{
+  typedef T value_type;
+};
+
+template <class T>
+struct E
+{
+  typedef T value_type;
+
+  template <class U>
+  struct rebind
+  {
+    typedef ReboundA<U> otter;
+  };
+};
+
+int main(int, char**)
+{
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<A<char>>::rebind_traits<double>,
+                                    cuda::std::allocator_traits<ReboundA<double>>>::value),
+                "");
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<B<int, char>>::rebind_traits<double>,
+                                    cuda::std::allocator_traits<ReboundB<double, char>>>::value),
+                "");
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<C<char>>::rebind_traits<double>,
+                                    cuda::std::allocator_traits<C<double>>>::value),
+                "");
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<D<int, char>>::rebind_traits<double>,
+                                    cuda::std::allocator_traits<D<double, char>>>::value),
+                "");
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<E<char>>::rebind_traits<double>,
+                                    cuda::std::allocator_traits<E<double>>>::value),
+                "");
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/value_type.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/value_type.pass.cpp
new file mode 100644
index 0000000000..9bcc9a3851
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.traits/value_type.pass.cpp
@@ -0,0 +1,35 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// <memory>
+
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     typedef typename Alloc::value_type value_type;
+//     ...
+// };
+
+#include <cuda/std/__memory>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+template <class T>
+struct A
+{
+  typedef T value_type;
+};
+
+int main(int, char**)
+{
+  static_assert((cuda::std::is_same<cuda::std::allocator_traits<A<char>>::value_type, char>::value), "");
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.uses/allocator.uses.trait/uses_allocator.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.uses/allocator.uses.trait/uses_allocator.pass.cpp
new file mode 100644
index 0000000000..13962ec2b1
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/allocator.uses/allocator.uses.trait/uses_allocator.pass.cpp
@@ -0,0 +1,96 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+// <memory>
+
+// template <class T, class Alloc> struct uses_allocator;
+
+#include <cuda/std/__memory>
+#if defined(_LIBCUDACXX_HAS_VECTOR)
+#  include <cuda/std/vector>
+#endif // _LIBCUDACXX_HAS_VECTOR
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+struct A
+{};
+
+struct B
+{
+  typedef int allocator_type;
+};
+
+#if !defined(TEST_COMPILER_NVRTC)
+struct C
+{
+  static int allocator_type;
+};
+#endif // !TEST_COMPILER_NVRTC
+
+struct D
+{
+  __host__ __device__ static int allocator_type()
+  {
+    return 0;
+  }
+};
+
+struct E
+{
+private:
+  typedef int allocator_type;
+};
+
+template <bool Expected, class T, class A>
+__host__ __device__ void test()
+{
+  static_assert((cuda::std::uses_allocator<T, A>::value == Expected), "");
+  static_assert(
+    cuda::std::is_base_of<cuda::std::integral_constant<bool, Expected>, cuda::std::uses_allocator<T, A>>::value, "");
+#if TEST_STD_VER >= 2017
+  ASSERT_SAME_TYPE(decltype(cuda::std::uses_allocator_v<T, A>), const bool);
+  static_assert((cuda::std::uses_allocator_v<T, A> == Expected), "");
+#endif // TEST_STD_VER >= 2017
+}
+
+int main(int, char**)
+{
+  test<false, int, cuda::std::allocator<int>>();
+#if defined(_LIBCUDACXX_HAS_VECTOR)
+  test<true, cuda::std::vector<int>, cuda::std::allocator<int>>();
+#endif //_LIBCUDACXX_HAS_VECTOR
+  test<false, A, cuda::std::allocator<int>>();
+  test<false, B, cuda::std::allocator<int>>();
+  test<true, B, double>();
+#if !defined(TEST_COMPILER_NVRTC)
+  test<false, C, decltype(C::allocator_type)>();
+#endif // !TEST_COMPILER_NVRTC
+  test<false, D, decltype(D::allocator_type)>();
+#if !defined(TEST_COMPILER_GCC) && !defined(TEST_COMPILER_MSVC_2017) // E::allocator_type is private
+  test<false, E, int>();
+#endif // !TEST_COMPILER_GCC && !TEST_COMPILER_MSVC_2017
+
+  static_assert((!cuda::std::uses_allocator<int, cuda::std::allocator<int>>::value), "");
+#if defined(_LIBCUDACXX_HAS_VECTOR)
+  static_assert((cuda::std::uses_allocator<cuda::std::vector<int>, cuda::std::allocator<int>>::value), "");
+#endif // _LIBCUDACXX_HAS_VECTOR
+  static_assert((!cuda::std::uses_allocator<A, cuda::std::allocator<int>>::value), "");
+  static_assert((!cuda::std::uses_allocator<B, cuda::std::allocator<int>>::value), "");
+  static_assert((cuda::std::uses_allocator<B, double>::value), "");
+#if !defined(TEST_COMPILER_NVRTC)
+  static_assert((!cuda::std::uses_allocator<C, decltype(C::allocator_type)>::value), "");
+  static_assert((!cuda::std::uses_allocator<D, decltype(D::allocator_type)>::value), "");
+#endif // !TEST_COMPILER_NVRTC
+#if !defined(TEST_COMPILER_GCC) && !defined(TEST_COMPILER_MSVC_2017) // E::allocator_type is private
+  static_assert((!cuda::std::uses_allocator<E, int>::value), "");
+#endif // !TEST_COMPILER_GCC && !TEST_COMPILER_MSVC_2017
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/PR50299.compile.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/PR50299.compile.pass.cpp
new file mode 100644
index 0000000000..a28206a873
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/PR50299.compile.pass.cpp
@@ -0,0 +1,30 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+
+// <memory>
+
+// Make sure we can use cuda::std::allocator<void> in all Standard modes. While the
+// explicit specialization for cuda::std::allocator<void> was deprecated, using that
+// specialization was neither deprecated nor removed (in C++20 it should simply
+// start using the primary template).
+//
+// See https://llvm.org/PR50299.
+
+#include <cuda/std/__memory>
+
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::std::allocator<void> alloc;
+
+int main(int, char**)
+{
+  unused(alloc);
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.ctor.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.ctor.pass.cpp
new file mode 100644
index 0000000000..142775780d
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.ctor.pass.cpp
@@ -0,0 +1,55 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+
+// <memory>
+//
+// template <class T>
+// class allocator
+// {
+// public: // All of these are constexpr after C++17
+//  allocator() noexcept;
+//  allocator(const allocator&) noexcept;
+//  template<class U> allocator(const allocator<U>&) noexcept;
+// ...
+// };
+
+#include <cuda/std/__memory>
+#include <cuda/std/cstddef>
+
+#include "test_macros.h"
+
+template <class T>
+__host__ __device__ TEST_CONSTEXPR_CXX20 bool test()
+{
+  typedef cuda::std::allocator<T> A1;
+  typedef cuda::std::allocator<long> A2;
+
+  A1 a1;
+  A1 a1_copy = a1;
+  unused(a1_copy);
+  A2 a2 = a1;
+  unused(a2);
+
+  return true;
+}
+
+int main(int, char**)
+{
+  test<char>();
+  test<char const>();
+  test<void>();
+
+#if TEST_STD_VER >= 2020
+  static_assert(test<char>());
+  static_assert(test<char const>());
+  static_assert(test<void>());
+#endif // TEST_STD_VER >= 2020
+  return 0;
+}
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
new file mode 100644
index 0000000000..6619f32872
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.dtor.pass.cpp
@@ -0,0 +1,45 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: c++03, c++11, c++14, c++17
+
+// template <class T>
+// constexpr allocator<T>::~allocator();
+
+#include <cuda/std/__memory>
+
+#include "test_macros.h"
+
+template <typename T>
+__host__ __device__ constexpr bool test()
+{
+  cuda::std::allocator<T> alloc;
+  unused(alloc);
+
+  // destructor called here
+  return true;
+}
+
+int main(int, char**)
+{
+  test<int>();
+  test<void>();
+#ifdef _LIBCUDACXX_VERSION // extension
+  test<int const>();
+#endif // _LIBCUDACXX_VERSION
+
+  static_assert(test<int>());
+  static_assert(test<void>());
+#ifdef _LIBCUDACXX_VERSION // extension
+  static_assert(test<int const>());
+#endif // _LIBCUDACXX_VERSION
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.globals/eq.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.globals/eq.pass.cpp
new file mode 100644
index 0000000000..87a220398d
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.globals/eq.pass.cpp
@@ -0,0 +1,47 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+
+// <memory>
+
+// allocator:
+
+// template <class T1, class T2>
+//   constexpr bool
+//   operator==(const allocator<T1>&, const allocator<T2>&) throw();
+//
+// template <class T1, class T2>
+//   constexpr bool
+//   operator!=(const allocator<T1>&, const allocator<T2>&) throw();
+
+#include <cuda/std/__memory>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+__host__ __device__ TEST_CONSTEXPR_CXX20 bool test()
+{
+  cuda::std::allocator<int> a1;
+  cuda::std::allocator<int> a2;
+  assert(a1 == a2);
+  assert(!(a1 != a2));
+
+  return true;
+}
+
+int main(int, char**)
+{
+  test();
+
+#if TEST_STD_VER >= 2020
+  static_assert(test());
+#endif // TEST_STD_VER >= 2020
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.members/allocate.constexpr.size.verify.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.members/allocate.constexpr.size.verify.cpp
new file mode 100644
index 0000000000..d63bf79787
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.members/allocate.constexpr.size.verify.cpp
@@ -0,0 +1,48 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+
+// <memory>
+
+// allocator:
+// constexpr T* allocate(size_type n);
+
+// UNSUPPORTED: c++03, c++11, c++14, c++17
+
+#include <cuda/std/__memory>
+#include <cuda/std/cstddef>
+
+#include "test_macros.h"
+
+template <typename T>
+__host__ __device__ constexpr bool test()
+{
+  typedef cuda::std::allocator<T> A;
+  typedef cuda::std::allocator_traits<A> AT;
+  A a;
+  TEST_IGNORE_NODISCARD a.allocate(AT::max_size(a) + 1); // just barely too large
+  TEST_IGNORE_NODISCARD a.allocate(AT::max_size(a) * 2); // significantly too large
+  TEST_IGNORE_NODISCARD a.allocate(((cuda::std::size_t) -1) / sizeof(T) + 1); // multiply will overflow
+  TEST_IGNORE_NODISCARD a.allocate((cuda::std::size_t) -1); // way too large
+
+  return true;
+}
+
+__host__ __device__ void f()
+{
+  static_assert(test<double>()); // expected-error-re {{{{(static_assert|static assertion)}} expression is not an
+                                 // integral constant expression}}
+  LIBCPP_STATIC_ASSERT(test<const double>()); // expected-error-re {{{{(static_assert|static assertion)}} expression is
+                                              // not an integral constant expression}}
+}
+
+int main(int, char**)
+{
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.members/allocate.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.members/allocate.pass.cpp
new file mode 100644
index 0000000000..651edf46b3
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.members/allocate.pass.cpp
@@ -0,0 +1,136 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+
+// <memory>
+
+// allocator:
+// constexpr T* allocate(size_t n);
+
+#define _LIBCUDACXX_DISABLE_DEPRECATION_WARNINGS
+
+#include <cuda/std/__memory>
+#include <cuda/std/cassert>
+#include <cuda/std/cstddef> // for cuda::std::max_align_t
+
+#include "count_new.h"
+#include "test_macros.h"
+
+#if defined(TEST_COMPILER_MSVC)
+#  pragma warning(disable : 4324)
+#endif // TEST_COMPILER_MSVC
+
+#ifdef TEST_HAS_NO_ALIGNED_ALLOCATION
+static const bool UsingAlignedNew = false;
+#else
+static const bool UsingAlignedNew = true;
+#endif
+
+#ifdef __STDCPP_DEFAULT_NEW_ALIGNMENT__
+STATIC_TEST_GLOBAL_VAR const cuda::std::size_t MaxAligned = __STDCPP_DEFAULT_NEW_ALIGNMENT__;
+#else
+STATIC_TEST_GLOBAL_VAR const cuda::std::size_t MaxAligned = cuda::std::alignment_of<cuda::std::max_align_t>::value;
+#endif
+
+STATIC_TEST_GLOBAL_VAR const cuda::std::size_t OverAligned = MaxAligned * 2;
+
+STATIC_TEST_GLOBAL_VAR int AlignedType_constructed = 0;
+
+template <cuda::std::size_t Align>
+struct TEST_ALIGNAS(Align) AlignedType
+{
+  char data;
+  __host__ __device__ AlignedType()
+  {
+    ++AlignedType_constructed;
+  }
+  __host__ __device__ AlignedType(AlignedType const&)
+  {
+    ++AlignedType_constructed;
+  }
+  __host__ __device__ ~AlignedType()
+  {
+    --AlignedType_constructed;
+  }
+};
+
+template <cuda::std::size_t Align>
+__host__ __device__ void test_aligned()
+{
+  typedef AlignedType<Align> T;
+  AlignedType_constructed = 0;
+  globalMemCounter.reset();
+  cuda::std::allocator<T> a;
+  const bool IsOverAlignedType = Align > MaxAligned;
+  const bool ExpectAligned     = IsOverAlignedType && UsingAlignedNew;
+  {
+    assert(globalMemCounter.checkOutstandingNewEq(0));
+    assert(AlignedType_constructed == 0);
+    globalMemCounter.last_new_size  = 0;
+    globalMemCounter.last_new_align = 0;
+    T* ap                           = a.allocate(3);
+#if !defined(TEST_COMPILER_NVCC) && !defined(TEST_COMPILER_NVRTC)
+    DoNotOptimize(ap);
+#else
+    const auto meow = reinterpret_cast<uintptr_t>(ap) + 2;
+    (void) meow;
+#endif // !TEST_COMPILER_NVCC && !TEST_COMPILER_NVRTC
+    // assert(globalMemCounter.checkOutstandingNewEq(1));
+    assert(globalMemCounter.checkNewCalledEq(1));
+    assert(globalMemCounter.checkAlignedNewCalledEq(ExpectAligned));
+    assert(globalMemCounter.checkLastNewSizeEq(3 * sizeof(T)));
+    assert(globalMemCounter.checkLastNewAlignEq(ExpectAligned ? Align : 0));
+    assert(AlignedType_constructed == 0);
+    globalMemCounter.last_delete_align = 0;
+    a.deallocate(ap, 3);
+    // assert(globalMemCounter.checkOutstandingNewEq(0));
+    assert(globalMemCounter.checkDeleteCalledEq(1));
+    assert(globalMemCounter.checkAlignedDeleteCalledEq(ExpectAligned));
+    assert(globalMemCounter.checkLastDeleteAlignEq(ExpectAligned ? Align : 0));
+    assert(AlignedType_constructed == 0);
+  }
+}
+
+#if TEST_STD_VER >= 2020
+template <cuda::std::size_t Align>
+__host__ __device__ constexpr bool test_aligned_constexpr()
+{
+  typedef AlignedType<Align> T;
+  cuda::std::allocator<T> a;
+  T* ap = a.allocate(3);
+  a.deallocate(ap, 3);
+
+  return true;
+}
+#endif // TEST_STD_VER >= 2020
+
+int main(int, char**)
+{
+  test_aligned<1>();
+  test_aligned<2>();
+  test_aligned<4>();
+  test_aligned<8>();
+  test_aligned<16>();
+  test_aligned<MaxAligned>();
+  test_aligned<OverAligned>();
+  test_aligned<OverAligned * 2>();
+
+#if defined(_CCCL_HAS_CONSTEXPR_ALLOCATION)
+  static_assert(test_aligned_constexpr<1>());
+  static_assert(test_aligned_constexpr<2>());
+  static_assert(test_aligned_constexpr<4>());
+  static_assert(test_aligned_constexpr<8>());
+  static_assert(test_aligned_constexpr<16>());
+  static_assert(test_aligned_constexpr<MaxAligned>());
+  static_assert(test_aligned_constexpr<OverAligned>());
+  static_assert(test_aligned_constexpr<OverAligned * 2>());
+#endif // _CCCL_HAS_CONSTEXPR_ALLOCATION
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.members/allocate.size.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.members/allocate.size.pass.cpp
new file mode 100644
index 0000000000..94ee545a7a
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.members/allocate.size.pass.cpp
@@ -0,0 +1,60 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-exceptions
+// <memory>
+
+// allocator:
+// constexpr T* allocate(size_t n);
+
+#define _LIBCUDACXX_DISABLE_DEPRECATION_WARNINGS
+
+#include <cuda/std/__memory>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+#ifndef TEST_HAS_NO_EXCEPTIONS
+template <typename T>
+__host__ __device__ void test_max(cuda::std::size_t count)
+{
+  cuda::std::allocator<T> a;
+  try
+  {
+    TEST_IGNORE_NODISCARD a.allocate(count);
+    assert(false);
+  }
+  catch (const cuda::std::bad_array_new_length&)
+  {}
+}
+
+template <typename T>
+__host__ __device__ void test()
+{
+  // Bug 26812 -- allocating too large
+  typedef cuda::std::allocator<T> A;
+  typedef cuda::std::allocator_traits<A> AT;
+  A a;
+  test_max<T>(AT::max_size(a) + 1); // just barely too large
+  test_max<T>(AT::max_size(a) * 2); // significantly too large
+  test_max<T>(((cuda::std::size_t) -1) / sizeof(T) + 1); // multiply will overflow
+  test_max<T>((cuda::std::size_t) -1); // way too large
+}
+#endif // !TEST_HAS_NO_EXCEPTIONS
+
+int main(int, char**)
+{
+#ifndef TEST_HAS_NO_EXCEPTIONS
+  NV_IF_TARGET(NV_IS_HOST, (test<double>();));
+  NV_IF_TARGET(NV_IS_HOST, (test<const double>();));
+#endif // !TEST_HAS_NO_EXCEPTIONS
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.members/allocate.verify.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.members/allocate.verify.cpp
new file mode 100644
index 0000000000..766fe63347
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.members/allocate.verify.cpp
@@ -0,0 +1,30 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: c++03, c++11, c++14, c++17
+// UNSUPPORTED: nvrtc
+
+// <memory>
+
+// allocator:
+// T* allocate(size_t n);
+
+#include <cuda/std/__memory>
+
+__host__ __device__ void f()
+{
+  cuda::std::allocator<int> a;
+  a.allocate(3); // expected-warning {{ignoring return value of function declared with 'nodiscard' attribute}}
+}
+
+int main(int, char**)
+{
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.members/allocate_at_least.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.members/allocate_at_least.pass.cpp
new file mode 100644
index 0000000000..68fb6712ac
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.members/allocate_at_least.pass.cpp
@@ -0,0 +1,123 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: c++03, c++11, c++14, c++17, c++20
+
+// <memory>
+
+// allocation_result<T*> allocate_at_least(size_t n)
+
+#include <cuda/std/__memory>
+#include <cuda/std/cassert>
+#include <cuda/std/concepts>
+
+#include "count_new.h"
+
+#ifdef TEST_HAS_NO_ALIGNED_ALLOCATION
+static const bool UsingAlignedNew = false;
+#else
+static const bool UsingAlignedNew = true;
+#endif
+
+#ifdef __STDCPP_DEFAULT_NEW_ALIGNMENT__
+static const cuda::std::size_t MaxAligned = __STDCPP_DEFAULT_NEW_ALIGNMENT__;
+#else
+static const cuda::std::size_t MaxAligned = cuda::std::alignment_of<cuda::std::max_align_t>::value;
+#endif
+
+static const cuda::std::size_t OverAligned = MaxAligned * 2;
+
+template <cuda::std::size_t Align>
+struct alignas(Align) AlignedType
+{
+  char data;
+  static int constructed;
+  __host__ __device__ AlignedType()
+  {
+    ++constructed;
+  }
+  __host__ __device__ AlignedType(AlignedType const&)
+  {
+    ++constructed;
+  }
+  __host__ __device__ ~AlignedType()
+  {
+    --constructed;
+  }
+};
+template <cuda::std::size_t Align>
+int AlignedType<Align>::constructed = 0;
+
+template <cuda::std::size_t Align>
+__host__ __device__ void test_aligned()
+{
+  typedef AlignedType<Align> T;
+  T::constructed = 0;
+  globalMemCounter.reset();
+  cuda::std::allocator<T> a;
+  const bool IsOverAlignedType = Align > MaxAligned;
+  const bool ExpectAligned     = IsOverAlignedType && UsingAlignedNew;
+  {
+    assert(globalMemCounter.checkOutstandingNewEq(0));
+    assert(T::constructed == 0);
+    globalMemCounter.last_new_size                                         = 0;
+    globalMemCounter.last_new_align                                        = 0;
+    cuda::std::same_as<cuda::std::allocation_result<T*>> decltype(auto) ap = a.allocate_at_least(3);
+    assert(ap.count >= 3);
+    DoNotOptimize(ap);
+    assert(globalMemCounter.checkOutstandingNewEq(1));
+    assert(globalMemCounter.checkNewCalledEq(1));
+    assert(globalMemCounter.checkAlignedNewCalledEq(ExpectAligned));
+    assert(globalMemCounter.checkLastNewSizeEq(3 * sizeof(T)));
+    assert(globalMemCounter.checkLastNewAlignEq(ExpectAligned ? Align : 0));
+    assert(T::constructed == 0);
+    globalMemCounter.last_delete_align = 0;
+    a.deallocate(ap.ptr, 3);
+    assert(globalMemCounter.checkOutstandingNewEq(0));
+    assert(globalMemCounter.checkDeleteCalledEq(1));
+    assert(globalMemCounter.checkAlignedDeleteCalledEq(ExpectAligned));
+    assert(globalMemCounter.checkLastDeleteAlignEq(ExpectAligned ? Align : 0));
+    assert(T::constructed == 0);
+  }
+}
+
+template <cuda::std::size_t Align>
+__host__ __device__ constexpr bool test_aligned_constexpr()
+{
+  typedef AlignedType<Align> T;
+  cuda::std::allocator<T> a;
+  cuda::std::same_as<cuda::std::allocation_result<T*>> decltype(auto) ap = a.allocate_at_least(3);
+  assert(ap.count >= 3);
+  a.deallocate(ap.ptr, 3);
+
+  return true;
+}
+
+int main(int, char**)
+{
+  test_aligned<1>();
+  test_aligned<2>();
+  test_aligned<4>();
+  test_aligned<8>();
+  test_aligned<16>();
+  test_aligned<MaxAligned>();
+  test_aligned<OverAligned>();
+  test_aligned<OverAligned * 2>();
+
+  static_assert(test_aligned_constexpr<1>());
+  static_assert(test_aligned_constexpr<2>());
+  static_assert(test_aligned_constexpr<4>());
+  static_assert(test_aligned_constexpr<8>());
+  static_assert(test_aligned_constexpr<16>());
+  static_assert(test_aligned_constexpr<MaxAligned>());
+  static_assert(test_aligned_constexpr<OverAligned>());
+  static_assert(test_aligned_constexpr<OverAligned * 2>());
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator_pointers.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator_pointers.pass.cpp
new file mode 100644
index 0000000000..04cb34af8b
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator_pointers.pass.cpp
@@ -0,0 +1,129 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: c++03
+
+#include <cuda/std/__memory>
+#include <cuda/std/cassert>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+// <memory>
+//
+// template <class Alloc>
+// struct allocator_traits
+// {
+//     typedef Alloc                        allocator_type;
+//     typedef typename allocator_type::value_type
+//                                          value_type;
+//
+//     typedef Alloc::pointer | value_type* pointer;
+//     typedef Alloc::const_pointer
+//           | pointer_traits<pointer>::rebind<const value_type>
+//                                          const_pointer;
+//     typedef Alloc::void_pointer
+//           | pointer_traits<pointer>::rebind<void>
+//                                          void_pointer;
+//     typedef Alloc::const_void_pointer
+//           | pointer_traits<pointer>::rebind<const void>
+//                                          const_void_pointer;
+
+template <typename Alloc>
+__host__ __device__ void test_pointer()
+{
+  typename cuda::std::allocator_traits<Alloc>::pointer vp;
+  typename cuda::std::allocator_traits<Alloc>::const_pointer cvp;
+
+  unused(vp); // Prevent unused warning
+  unused(cvp); // Prevent unused warning
+
+  static_assert(cuda::std::is_same<bool, decltype(vp == vp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(vp != vp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(vp > vp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(vp >= vp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(vp < vp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(vp <= vp)>::value, "");
+
+  static_assert(cuda::std::is_same<bool, decltype(vp == cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp == vp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(vp != cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp != vp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(vp > cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp > vp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(vp >= cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp >= vp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(vp < cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp < vp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(vp <= cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp <= vp)>::value, "");
+
+  static_assert(cuda::std::is_same<bool, decltype(cvp == cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp != cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp > cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp >= cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp < cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp <= cvp)>::value, "");
+}
+
+template <typename Alloc>
+__host__ __device__ void test_void_pointer()
+{
+  typename cuda::std::allocator_traits<Alloc>::void_pointer vp;
+  typename cuda::std::allocator_traits<Alloc>::const_void_pointer cvp;
+
+  unused(vp); // Prevent unused warning
+  unused(cvp); // Prevent unused warning
+
+  static_assert(cuda::std::is_same<bool, decltype(vp == vp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(vp != vp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(vp > vp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(vp >= vp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(vp < vp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(vp <= vp)>::value, "");
+
+  static_assert(cuda::std::is_same<bool, decltype(vp == cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp == vp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(vp != cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp != vp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(vp > cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp > vp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(vp >= cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp >= vp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(vp < cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp < vp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(vp <= cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp <= vp)>::value, "");
+
+  static_assert(cuda::std::is_same<bool, decltype(cvp == cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp != cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp > cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp >= cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp < cvp)>::value, "");
+  static_assert(cuda::std::is_same<bool, decltype(cvp <= cvp)>::value, "");
+}
+
+struct Foo
+{
+  int x;
+};
+
+int main(int, char**)
+{
+  test_pointer<cuda::std::allocator<char>>();
+  test_pointer<cuda::std::allocator<int>>();
+  test_pointer<cuda::std::allocator<Foo>>();
+
+  test_void_pointer<cuda::std::allocator<char>>();
+  test_void_pointer<cuda::std::allocator<int>>();
+  test_void_pointer<cuda::std::allocator<Foo>>();
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator_types.deprecated_in_cxx17.verify.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator_types.deprecated_in_cxx17.verify.cpp
new file mode 100644
index 0000000000..575043fc24
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator_types.deprecated_in_cxx17.verify.cpp
@@ -0,0 +1,65 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+
+// <memory>
+
+// Check that the following nested types are deprecated in C++17:
+
+// template <class T>
+// class allocator
+// {
+// public:
+//     typedef T*                                           pointer;
+//     typedef const T*                                     const_pointer;
+//     typedef typename add_lvalue_reference<T>::type       reference;
+//     typedef typename add_lvalue_reference<const T>::type const_reference;
+//
+//     template <class U> struct rebind {typedef allocator<U> other;};
+// ...
+// };
+
+// REQUIRES: c++17
+
+#include <cuda/std/__memory>
+
+__host__ __device__ void f()
+{
+  {
+    typedef cuda::std::allocator<char>::pointer Pointer; // expected-warning {{'pointer' is deprecated}}
+    typedef cuda::std::allocator<char>::const_pointer ConstPointer; // expected-warning {{'const_pointer' is
+                                                                    // deprecated}}
+    typedef cuda::std::allocator<char>::reference Reference; // expected-warning {{'reference' is deprecated}}
+    typedef cuda::std::allocator<char>::const_reference ConstReference; // expected-warning {{'const_reference' is
+                                                                        // deprecated}}
+    typedef cuda::std::allocator<char>::rebind<int>::other Rebind; // expected-warning {{'rebind<int>' is deprecated}}
+  }
+  {
+    typedef cuda::std::allocator<char const>::pointer Pointer; // expected-warning {{'pointer' is deprecated}}
+    typedef cuda::std::allocator<char const>::const_pointer ConstPointer; // expected-warning {{'const_pointer' is
+                                                                          // deprecated}}
+    typedef cuda::std::allocator<char const>::reference Reference; // expected-warning {{'reference' is deprecated}}
+    typedef cuda::std::allocator<char const>::const_reference ConstReference; // expected-warning {{'const_reference' is
+                                                                              // deprecated}}
+    typedef cuda::std::allocator<char const>::rebind<int>::other Rebind; // expected-warning {{'rebind<int>' is
+                                                                         // deprecated}}
+  }
+  {
+    typedef cuda::std::allocator<void>::pointer Pointer; // expected-warning {{'pointer' is deprecated}}
+    typedef cuda::std::allocator<void>::const_pointer ConstPointer; // expected-warning {{'const_pointer' is
+                                                                    // deprecated}}
+    // reference and const_reference are not provided by cuda::std::allocator<void>
+    typedef cuda::std::allocator<void>::rebind<int>::other Rebind; // expected-warning {{'rebind<int>' is deprecated}}
+  }
+}
+
+int main(int, char**)
+{
+  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
new file mode 100644
index 0000000000..67ac0ccd98
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator_types.pass.cpp
@@ -0,0 +1,71 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+
+// <memory>
+
+// Check that the nested types of cuda::std::allocator are provided:
+
+// template <class T>
+// class allocator
+// {
+// public:
+//     typedef size_t    size_type;
+//     typedef ptrdiff_t difference_type;
+//     typedef T         value_type;
+//
+//     typedef T*        pointer;           // deprecated in C++17, removed in C++20
+//     typedef T const*  const_pointer;     // deprecated in C++17, removed in C++20
+//     typedef T&        reference;         // deprecated in C++17, removed in C++20
+//     typedef T const&  const_reference;   // deprecated in C++17, removed in C++20
+//     template< class U > struct rebind { typedef allocator<U> other; }; // deprecated in C++17, removed in C++20
+//
+//     typedef true_type propagate_on_container_move_assignment;
+//     typedef true_type is_always_equal;
+// ...
+// };
+
+// ADDITIONAL_COMPILE_FLAGS: -D_LIBCUDACXX_DISABLE_DEPRECATION_WARNINGS
+
+#include <cuda/std/__memory>
+#include <cuda/std/cstddef>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+struct U;
+
+template <typename T>
+__host__ __device__ void test()
+{
+  typedef cuda::std::allocator<T> Alloc;
+  static_assert((cuda::std::is_same<typename Alloc::size_type, cuda::std::size_t>::value), "");
+  static_assert((cuda::std::is_same<typename Alloc::difference_type, cuda::std::ptrdiff_t>::value), "");
+  static_assert((cuda::std::is_same<typename Alloc::value_type, T>::value), "");
+  static_assert(
+    (cuda::std::is_same<typename Alloc::propagate_on_container_move_assignment, cuda::std::true_type>::value), "");
+  static_assert((cuda::std::is_same<typename Alloc::is_always_equal, cuda::std::true_type>::value), "");
+
+#if TEST_STD_VER <= 2017
+  static_assert((cuda::std::is_same<typename Alloc::pointer, T*>::value), "");
+  static_assert((cuda::std::is_same<typename Alloc::const_pointer, T const*>::value), "");
+  static_assert((cuda::std::is_same<typename Alloc::reference, T&>::value), "");
+  static_assert((cuda::std::is_same<typename Alloc::const_reference, T const&>::value), "");
+  static_assert((cuda::std::is_same<typename Alloc::template rebind<U>::other, cuda::std::allocator<U>>::value), "");
+#endif // TEST_STD_VER <= 2017
+}
+
+int main(int, char**)
+{
+  test<char>();
+#ifdef _LIBCUDACXX_VERSION
+  test<char const>(); // extension
+#endif
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator_types.removed_in_cxx20.verify.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator_types.removed_in_cxx20.verify.cpp
new file mode 100644
index 0000000000..69124c5392
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator_types.removed_in_cxx20.verify.cpp
@@ -0,0 +1,53 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+
+// <memory>
+
+// Check that the following nested types are removed in C++20:
+
+// template <class T>
+// class allocator
+// {
+// public:
+//     typedef T*                                           pointer;
+//     typedef const T*                                     const_pointer;
+//     typedef typename add_lvalue_reference<T>::type       reference;
+//     typedef typename add_lvalue_reference<const T>::type const_reference;
+//
+//     template <class U> struct rebind {typedef allocator<U> other;};
+// ...
+// };
+
+// UNSUPPORTED: c++03, c++11, c++14, c++17
+
+#include <cuda/std/__memory>
+
+template <typename T>
+__host__ __device__ void check()
+{
+  typedef typename cuda::std::allocator<T>::pointer AP; // expected-error 3 {{no type named 'pointer'}}
+  typedef typename cuda::std::allocator<T>::const_pointer ACP; // expected-error 3 {{no type named 'const_pointer'}}
+  typedef typename cuda::std::allocator<T>::reference AR; // expected-error 3 {{no type named 'reference'}}
+  typedef typename cuda::std::allocator<T>::const_reference ACR; // expected-error 3 {{no type named 'const_reference'}}
+  typedef typename cuda::std::allocator<T>::template rebind<int>::other ARO; // expected-error 3 {{no member named
+                                                                             // 'rebind'}}
+}
+
+__host__ __device__ void f()
+{
+  check<char>();
+  check<char const>();
+  check<void>();
+}
+
+int main(int, char**)
+{
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator_types.void.compile.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator_types.void.compile.pass.cpp
new file mode 100644
index 0000000000..9c04594cb1
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator_types.void.compile.pass.cpp
@@ -0,0 +1,42 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+
+// Check that the nested types of cuda::std::allocator<void> are provided.
+// After C++17, those are not provided in the primary template and the
+// explicit specialization doesn't exist anymore, so this test is moot.
+
+// REQUIRES: c++03 || c++11 || c++14 || c++17
+
+// template <>
+// class allocator<void>
+// {
+// public:
+//     typedef void*                                 pointer;
+//     typedef const void*                           const_pointer;
+//     typedef void                                  value_type;
+//
+//     template <class _Up> struct rebind {typedef allocator<_Up> other;};
+// };
+
+// ADDITIONAL_COMPILE_FLAGS: -D_LIBCUDACXX_DISABLE_DEPRECATION_WARNINGS
+
+#include <cuda/std/__memory>
+#include <cuda/std/type_traits>
+
+static_assert((cuda::std::is_same<cuda::std::allocator<void>::pointer, void*>::value), "");
+static_assert((cuda::std::is_same<cuda::std::allocator<void>::const_pointer, const void*>::value), "");
+static_assert((cuda::std::is_same<cuda::std::allocator<void>::value_type, void>::value), "");
+static_assert((cuda::std::is_same<cuda::std::allocator<void>::rebind<int>::other, cuda::std::allocator<int>>::value),
+              "");
+
+int main(int, char**)
+{
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/utilities/memory/ptr.align/align.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/ptr.align/align.pass.cpp
new file mode 100644
index 0000000000..da0c7fe520
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/ptr.align/align.pass.cpp
@@ -0,0 +1,89 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+// #include <memory>
+
+// void* align(size_t alignment, size_t size, void*& ptr, size_t& space);
+
+#include <cuda/std/__memory>
+#include <cuda/std/cassert>
+#include <cuda/std/cstddef>
+
+#include "test_macros.h"
+
+int main(int, char**)
+{
+  const unsigned N = 20;
+  char buf[N];
+  void* r;
+  void* p             = &buf[0];
+  cuda::std::size_t s = N;
+  r                   = cuda::std::align(4, 10, p, s);
+  assert(p == &buf[0]);
+  assert(r == p);
+  assert(s == N);
+
+  p = &buf[1];
+  s = N;
+  r = cuda::std::align(4, 10, p, s);
+  assert(p == &buf[4]);
+  assert(r == p);
+  assert(s == N - 3);
+
+  p = &buf[2];
+  s = N;
+  r = cuda::std::align(4, 10, p, s);
+  assert(p == &buf[4]);
+  assert(r == p);
+  assert(s == N - 2);
+
+  p = &buf[3];
+  s = N;
+  r = cuda::std::align(4, 10, p, s);
+  assert(p == &buf[4]);
+  assert(r == p);
+  assert(s == N - 1);
+
+  p = &buf[4];
+  s = N;
+  r = cuda::std::align(4, 10, p, s);
+  assert(p == &buf[4]);
+  assert(r == p);
+  assert(s == N);
+
+  p = &buf[0];
+  s = N;
+  r = cuda::std::align(4, N, p, s);
+  assert(p == &buf[0]);
+  assert(r == p);
+  assert(s == N);
+
+  p = &buf[1];
+  s = N - 1;
+  r = cuda::std::align(4, N - 4, p, s);
+  assert(p == &buf[4]);
+  assert(r == p);
+  assert(s == N - 4);
+
+  p = &buf[1];
+  s = N - 1;
+  r = cuda::std::align(4, N - 3, p, s);
+  assert(p == &buf[1]);
+  assert(r == nullptr);
+  assert(s == N - 1);
+
+  p = &buf[0];
+  s = N;
+  r = cuda::std::align(1, N + 1, p, s);
+  assert(p == &buf[0]);
+  assert(r == nullptr);
+  assert(s == N);
+
+  return 0;
+}
diff --git a/libcudacxx/test/support/count_new.h b/libcudacxx/test/support/count_new.h
index 3cdc7599aa..29c20f6a79 100644
--- a/libcudacxx/test/support/count_new.h
+++ b/libcudacxx/test/support/count_new.h
@@ -9,86 +9,80 @@
 #ifndef COUNT_NEW_H
 #define COUNT_NEW_H
 
-#include <cassert>
-#include <cstdlib>
-#include <new>
+#include <cuda/std/cassert>
+#include <cuda/std/cstdlib>
+
+#if !defined(_LIBCUDACXX_COMPILER_NVRTC)
+#  include <new>
+#endif // !_LIBCUDACXX_COMPILER_NVRTC
 
 #include "test_macros.h"
 
-#if defined(TEST_HAS_SANITIZERS)
+#if defined(TEST_HAS_SANITIZERS) || defined(__CUDA_ARCH__) // we cannot overload operator {new, delete} on device
 #  define DISABLE_NEW_COUNT
 #endif
 
-namespace detail
-{
-TEST_NORETURN
-inline void throw_bad_alloc_helper()
-{
-#ifndef TEST_HAS_NO_EXCEPTIONS
-  throw std::bad_alloc();
+// All checks return true when disable_checking is enabled.
+#ifdef DISABLE_NEW_COUNT
+STATIC_TEST_GLOBAL_VAR const bool MemCounter_disable_checking = true;
 #else
-  std::abort();
+STATIC_TEST_GLOBAL_VAR const bool MemCounter_disable_checking = false;
 #endif
-}
-} // namespace detail
+
+// number of allocations to throw after. Default (unsigned)-1. If
+// throw_after has the default value it will never be decremented.
+STATIC_TEST_GLOBAL_VAR const unsigned MemCounter_never_throw_value = static_cast<unsigned>(-1);
 
 class MemCounter
 {
 public:
+  MemCounter() = default;
+
   // Make MemCounter super hard to accidentally construct or copy.
   class MemCounterCtorArg_
   {};
-  explicit MemCounter(MemCounterCtorArg_)
-  {
-    reset();
-  }
+  __host__ __device__ explicit MemCounter(MemCounterCtorArg_) {}
 
 private:
-  MemCounter(MemCounter const&);
-  MemCounter& operator=(MemCounter const&);
+  __host__ __device__ MemCounter(MemCounter const&)            = delete;
+  __host__ __device__ MemCounter& operator=(MemCounter const&) = delete;
 
 public:
-  // All checks return true when disable_checking is enabled.
-  static const bool disable_checking;
-
   // Disallow any allocations from occurring. Useful for testing that
   // code doesn't perform any allocations.
-  bool disable_allocations;
-
-  // number of allocations to throw after. Default (unsigned)-1. If
-  // throw_after has the default value it will never be decremented.
-  static const unsigned never_throw_value = static_cast<unsigned>(-1);
-  unsigned throw_after;
-
-  int outstanding_new;
-  int new_called;
-  int delete_called;
-  int aligned_new_called;
-  int aligned_delete_called;
-  std::size_t last_new_size;
-  std::size_t last_new_align;
-  std::size_t last_delete_align;
-
-  int outstanding_array_new;
-  int new_array_called;
-  int delete_array_called;
-  int aligned_new_array_called;
-  int aligned_delete_array_called;
-  std::size_t last_new_array_size;
-  std::size_t last_new_array_align;
-  std::size_t last_delete_array_align;
+  bool disable_allocations = false;
+
+  unsigned throw_after = MemCounter_never_throw_value;
+
+  int outstanding_new                 = 0;
+  int new_called                      = 0;
+  int delete_called                   = 0;
+  int aligned_new_called              = 0;
+  int aligned_delete_called           = 0;
+  cuda::std::size_t last_new_size     = 0;
+  cuda::std::size_t last_new_align    = 0;
+  cuda::std::size_t last_delete_align = 0;
+
+  int outstanding_array_new                 = 0;
+  int new_array_called                      = 0;
+  int delete_array_called                   = 0;
+  int aligned_new_array_called              = 0;
+  int aligned_delete_array_called           = 0;
+  cuda::std::size_t last_new_array_size     = 0;
+  cuda::std::size_t last_new_array_align    = 0;
+  cuda::std::size_t last_delete_array_align = 0;
 
 public:
-  void newCalled(std::size_t s)
+  __host__ __device__ void newCalled(cuda::std::size_t s)
   {
     assert(disable_allocations == false);
     assert(s);
     if (throw_after == 0)
     {
-      throw_after = never_throw_value;
-      detail::throw_bad_alloc_helper();
+      throw_after = MemCounter_never_throw_value;
+      cuda::std::__throw_bad_alloc();
     }
-    else if (throw_after != never_throw_value)
+    else if (throw_after != MemCounter_never_throw_value)
     {
       --throw_after;
     }
@@ -97,35 +91,35 @@ class MemCounter
     last_new_size = s;
   }
 
-  void alignedNewCalled(std::size_t s, std::size_t a)
+  __host__ __device__ void alignedNewCalled(cuda::std::size_t s, cuda::std::size_t a)
   {
     newCalled(s);
     ++aligned_new_called;
     last_new_align = a;
   }
 
-  void deleteCalled(void* p)
+  __host__ __device__ void deleteCalled(void* p)
   {
     assert(p);
     --outstanding_new;
     ++delete_called;
   }
 
-  void alignedDeleteCalled(void* p, std::size_t a)
+  __host__ __device__ void alignedDeleteCalled(void* p, cuda::std::size_t a)
   {
     deleteCalled(p);
     ++aligned_delete_called;
     last_delete_align = a;
   }
 
-  void newArrayCalled(std::size_t s)
+  __host__ __device__ void newArrayCalled(cuda::std::size_t s)
   {
     assert(disable_allocations == false);
     assert(s);
     if (throw_after == 0)
     {
-      throw_after = never_throw_value;
-      detail::throw_bad_alloc_helper();
+      throw_after = MemCounter_never_throw_value;
+      cuda::std::__throw_bad_alloc();
     }
     else
     {
@@ -136,41 +130,41 @@ class MemCounter
     last_new_array_size = s;
   }
 
-  void alignedNewArrayCalled(std::size_t s, std::size_t a)
+  __host__ __device__ void alignedNewArrayCalled(cuda::std::size_t s, cuda::std::size_t a)
   {
     newArrayCalled(s);
     ++aligned_new_array_called;
     last_new_array_align = a;
   }
 
-  void deleteArrayCalled(void* p)
+  __host__ __device__ void deleteArrayCalled(void* p)
   {
     assert(p);
     --outstanding_array_new;
     ++delete_array_called;
   }
 
-  void alignedDeleteArrayCalled(void* p, std::size_t a)
+  __host__ __device__ void alignedDeleteArrayCalled(void* p, cuda::std::size_t a)
   {
     deleteArrayCalled(p);
     ++aligned_delete_array_called;
     last_delete_array_align = a;
   }
 
-  void disableAllocations()
+  __host__ __device__ void disableAllocations()
   {
     disable_allocations = true;
   }
 
-  void enableAllocations()
+  __host__ __device__ void enableAllocations()
   {
     disable_allocations = false;
   }
 
-  void reset()
+  __host__ __device__ void reset()
   {
     disable_allocations = false;
-    throw_after         = never_throw_value;
+    throw_after         = MemCounter_never_throw_value;
 
     outstanding_new       = 0;
     new_called            = 0;
@@ -190,194 +184,189 @@ class MemCounter
   }
 
 public:
-  bool checkOutstandingNewEq(int n) const
+  __host__ __device__ bool checkOutstandingNewEq(int n) const
   {
-    return disable_checking || n == outstanding_new;
+    return MemCounter_disable_checking || n == outstanding_new;
   }
 
-  bool checkOutstandingNewNotEq(int n) const
+  __host__ __device__ bool checkOutstandingNewNotEq(int n) const
   {
-    return disable_checking || n != outstanding_new;
+    return MemCounter_disable_checking || n != outstanding_new;
   }
 
-  bool checkNewCalledEq(int n) const
+  __host__ __device__ bool checkNewCalledEq(int n) const
   {
-    return disable_checking || n == new_called;
+    return MemCounter_disable_checking || n == new_called;
   }
 
-  bool checkNewCalledNotEq(int n) const
+  __host__ __device__ bool checkNewCalledNotEq(int n) const
   {
-    return disable_checking || n != new_called;
+    return MemCounter_disable_checking || n != new_called;
   }
 
-  bool checkNewCalledGreaterThan(int n) const
+  __host__ __device__ bool checkNewCalledGreaterThan(int n) const
   {
-    return disable_checking || new_called > n;
+    return MemCounter_disable_checking || new_called > n;
   }
 
-  bool checkDeleteCalledEq(int n) const
+  __host__ __device__ bool checkDeleteCalledEq(int n) const
   {
-    return disable_checking || n == delete_called;
+    return MemCounter_disable_checking || n == delete_called;
   }
 
-  bool checkDeleteCalledNotEq(int n) const
+  __host__ __device__ bool checkDeleteCalledNotEq(int n) const
   {
-    return disable_checking || n != delete_called;
+    return MemCounter_disable_checking || n != delete_called;
   }
 
-  bool checkAlignedNewCalledEq(int n) const
+  __host__ __device__ bool checkAlignedNewCalledEq(int n) const
   {
-    return disable_checking || n == aligned_new_called;
+    return MemCounter_disable_checking || n == aligned_new_called;
   }
 
-  bool checkAlignedNewCalledNotEq(int n) const
+  __host__ __device__ bool checkAlignedNewCalledNotEq(int n) const
   {
-    return disable_checking || n != aligned_new_called;
+    return MemCounter_disable_checking || n != aligned_new_called;
   }
 
-  bool checkAlignedNewCalledGreaterThan(int n) const
+  __host__ __device__ bool checkAlignedNewCalledGreaterThan(int n) const
   {
-    return disable_checking || aligned_new_called > n;
+    return MemCounter_disable_checking || aligned_new_called > n;
   }
 
-  bool checkAlignedDeleteCalledEq(int n) const
+  __host__ __device__ bool checkAlignedDeleteCalledEq(int n) const
   {
-    return disable_checking || n == aligned_delete_called;
+    return MemCounter_disable_checking || n == aligned_delete_called;
   }
 
-  bool checkAlignedDeleteCalledNotEq(int n) const
+  __host__ __device__ bool checkAlignedDeleteCalledNotEq(int n) const
   {
-    return disable_checking || n != aligned_delete_called;
+    return MemCounter_disable_checking || n != aligned_delete_called;
   }
 
-  bool checkLastNewSizeEq(std::size_t n) const
+  __host__ __device__ bool checkLastNewSizeEq(cuda::std::size_t n) const
   {
-    return disable_checking || n == last_new_size;
+    return MemCounter_disable_checking || n == last_new_size;
   }
 
-  bool checkLastNewSizeNotEq(std::size_t n) const
+  __host__ __device__ bool checkLastNewSizeNotEq(cuda::std::size_t n) const
   {
-    return disable_checking || n != last_new_size;
+    return MemCounter_disable_checking || n != last_new_size;
   }
 
-  bool checkLastNewAlignEq(std::size_t n) const
+  __host__ __device__ bool checkLastNewAlignEq(cuda::std::size_t n) const
   {
-    return disable_checking || n == last_new_align;
+    return MemCounter_disable_checking || n == last_new_align;
   }
 
-  bool checkLastNewAlignNotEq(std::size_t n) const
+  __host__ __device__ bool checkLastNewAlignNotEq(cuda::std::size_t n) const
   {
-    return disable_checking || n != last_new_align;
+    return MemCounter_disable_checking || n != last_new_align;
   }
 
-  bool checkLastDeleteAlignEq(std::size_t n) const
+  __host__ __device__ bool checkLastDeleteAlignEq(cuda::std::size_t n) const
   {
-    return disable_checking || n == last_delete_align;
+    return MemCounter_disable_checking || n == last_delete_align;
   }
 
-  bool checkLastDeleteAlignNotEq(std::size_t n) const
+  __host__ __device__ bool checkLastDeleteAlignNotEq(cuda::std::size_t n) const
   {
-    return disable_checking || n != last_delete_align;
+    return MemCounter_disable_checking || n != last_delete_align;
   }
 
-  bool checkOutstandingArrayNewEq(int n) const
+  __host__ __device__ bool checkOutstandingArrayNewEq(int n) const
   {
-    return disable_checking || n == outstanding_array_new;
+    return MemCounter_disable_checking || n == outstanding_array_new;
   }
 
-  bool checkOutstandingArrayNewNotEq(int n) const
+  __host__ __device__ bool checkOutstandingArrayNewNotEq(int n) const
   {
-    return disable_checking || n != outstanding_array_new;
+    return MemCounter_disable_checking || n != outstanding_array_new;
   }
 
-  bool checkNewArrayCalledEq(int n) const
+  __host__ __device__ bool checkNewArrayCalledEq(int n) const
   {
-    return disable_checking || n == new_array_called;
+    return MemCounter_disable_checking || n == new_array_called;
   }
 
-  bool checkNewArrayCalledNotEq(int n) const
+  __host__ __device__ bool checkNewArrayCalledNotEq(int n) const
   {
-    return disable_checking || n != new_array_called;
+    return MemCounter_disable_checking || n != new_array_called;
   }
 
-  bool checkDeleteArrayCalledEq(int n) const
+  __host__ __device__ bool checkDeleteArrayCalledEq(int n) const
   {
-    return disable_checking || n == delete_array_called;
+    return MemCounter_disable_checking || n == delete_array_called;
   }
 
-  bool checkDeleteArrayCalledNotEq(int n) const
+  __host__ __device__ bool checkDeleteArrayCalledNotEq(int n) const
   {
-    return disable_checking || n != delete_array_called;
+    return MemCounter_disable_checking || n != delete_array_called;
   }
 
-  bool checkAlignedNewArrayCalledEq(int n) const
+  __host__ __device__ bool checkAlignedNewArrayCalledEq(int n) const
   {
-    return disable_checking || n == aligned_new_array_called;
+    return MemCounter_disable_checking || n == aligned_new_array_called;
   }
 
-  bool checkAlignedNewArrayCalledNotEq(int n) const
+  __host__ __device__ bool checkAlignedNewArrayCalledNotEq(int n) const
   {
-    return disable_checking || n != aligned_new_array_called;
+    return MemCounter_disable_checking || n != aligned_new_array_called;
   }
 
-  bool checkAlignedNewArrayCalledGreaterThan(int n) const
+  __host__ __device__ bool checkAlignedNewArrayCalledGreaterThan(int n) const
   {
-    return disable_checking || aligned_new_array_called > n;
+    return MemCounter_disable_checking || aligned_new_array_called > n;
   }
 
-  bool checkAlignedDeleteArrayCalledEq(int n) const
+  __host__ __device__ bool checkAlignedDeleteArrayCalledEq(int n) const
   {
-    return disable_checking || n == aligned_delete_array_called;
+    return MemCounter_disable_checking || n == aligned_delete_array_called;
   }
 
-  bool checkAlignedDeleteArrayCalledNotEq(int n) const
+  __host__ __device__ bool checkAlignedDeleteArrayCalledNotEq(int n) const
   {
-    return disable_checking || n != aligned_delete_array_called;
+    return MemCounter_disable_checking || n != aligned_delete_array_called;
   }
 
-  bool checkLastNewArraySizeEq(std::size_t n) const
+  __host__ __device__ bool checkLastNewArraySizeEq(cuda::std::size_t n) const
   {
-    return disable_checking || n == last_new_array_size;
+    return MemCounter_disable_checking || n == last_new_array_size;
   }
 
-  bool checkLastNewArraySizeNotEq(std::size_t n) const
+  __host__ __device__ bool checkLastNewArraySizeNotEq(cuda::std::size_t n) const
   {
-    return disable_checking || n != last_new_array_size;
+    return MemCounter_disable_checking || n != last_new_array_size;
   }
 
-  bool checkLastNewArrayAlignEq(std::size_t n) const
+  __host__ __device__ bool checkLastNewArrayAlignEq(cuda::std::size_t n) const
   {
-    return disable_checking || n == last_new_array_align;
+    return MemCounter_disable_checking || n == last_new_array_align;
   }
 
-  bool checkLastNewArrayAlignNotEq(std::size_t n) const
+  __host__ __device__ bool checkLastNewArrayAlignNotEq(cuda::std::size_t n) const
   {
-    return disable_checking || n != last_new_array_align;
+    return MemCounter_disable_checking || n != last_new_array_align;
   }
 };
 
-#ifdef DISABLE_NEW_COUNT
-const bool MemCounter::disable_checking = true;
-#else
-const bool MemCounter::disable_checking = false;
-#endif
+STATIC_TEST_GLOBAL_VAR MemCounter counter{};
 
-inline MemCounter* getGlobalMemCounter()
+__host__ __device__ inline constexpr MemCounter* getGlobalMemCounter()
 {
-  static MemCounter counter((MemCounter::MemCounterCtorArg_()));
   return &counter;
 }
 
-MemCounter& globalMemCounter = *getGlobalMemCounter();
+STATIC_TEST_GLOBAL_VAR MemCounter& globalMemCounter = *getGlobalMemCounter();
 
 #ifndef DISABLE_NEW_COUNT
-void* operator new(std::size_t s) TEST_THROW_SPEC(std::bad_alloc)
+void* operator new(cuda::std::size_t s)
 {
   getGlobalMemCounter()->newCalled(s);
-  void* ret = std::malloc(s);
+  void* ret = malloc(s);
   if (ret == nullptr)
   {
-    detail::throw_bad_alloc_helper();
+    cuda::std::__throw_bad_alloc();
   }
   return ret;
 }
@@ -385,10 +374,18 @@ void* operator new(std::size_t s) TEST_THROW_SPEC(std::bad_alloc)
 void operator delete(void* p) TEST_NOEXCEPT
 {
   getGlobalMemCounter()->deleteCalled(p);
-  std::free(p);
+  free(p);
+}
+
+#  ifdef TEST_COMPILER_GCC
+void operator delete(void* p, cuda::std::size_t) TEST_NOEXCEPT
+{
+  getGlobalMemCounter()->deleteCalled(p);
+  free(p);
 }
+#  endif // TEST_COMPILER_GCC
 
-void* operator new[](std::size_t s) TEST_THROW_SPEC(std::bad_alloc)
+void* operator new[](cuda::std::size_t s)
 {
   getGlobalMemCounter()->newArrayCalled(s);
   return operator new(s);
@@ -400,14 +397,22 @@ void operator delete[](void* p) TEST_NOEXCEPT
   operator delete(p);
 }
 
+#  ifdef TEST_COMPILER_GCC
+void operator delete[](void* p, cuda::std::size_t) TEST_NOEXCEPT
+{
+  getGlobalMemCounter()->deleteArrayCalled(p);
+  operator delete(p);
+}
+#  endif // TEST_COMPILER_GCC
+
 #  ifndef TEST_HAS_NO_ALIGNED_ALLOCATION
 #    if defined(_LIBCUDACXX_MSVCRT_LIKE) || (!defined(_LIBCUDACXX_VERSION) && defined(_WIN32))
 #      define USE_ALIGNED_ALLOC
 #    endif
 
-void* operator new(std::size_t s, std::align_val_t av) TEST_THROW_SPEC(std::bad_alloc)
+void* operator new(cuda::std::size_t s, cuda::std::align_val_t av)
 {
-  const std::size_t a = static_cast<std::size_t>(av);
+  const cuda::std::size_t a = static_cast<cuda::std::size_t>(av);
   getGlobalMemCounter()->alignedNewCalled(s, a);
   void* ret;
 #    ifdef USE_ALIGNED_ALLOC
@@ -417,14 +422,14 @@ void* operator new(std::size_t s, std::align_val_t av) TEST_THROW_SPEC(std::bad_
 #    endif
   if (ret == nullptr)
   {
-    detail::throw_bad_alloc_helper();
+    cuda::std::__throw_bad_alloc();
   }
   return ret;
 }
 
-void operator delete(void* p, std::align_val_t av) TEST_NOEXCEPT
+void operator delete(void* p, cuda::std::align_val_t av) TEST_NOEXCEPT
 {
-  const std::size_t a = static_cast<std::size_t>(av);
+  const cuda::std::size_t a = static_cast<cuda::std::size_t>(av);
   getGlobalMemCounter()->alignedDeleteCalled(p, a);
   if (p)
   {
@@ -436,16 +441,16 @@ void operator delete(void* p, std::align_val_t av) TEST_NOEXCEPT
   }
 }
 
-void* operator new[](std::size_t s, std::align_val_t av) TEST_THROW_SPEC(std::bad_alloc)
+void* operator new[](cuda::std::size_t s, cuda::std::align_val_t av)
 {
-  const std::size_t a = static_cast<std::size_t>(av);
+  const cuda::std::size_t a = static_cast<cuda::std::size_t>(av);
   getGlobalMemCounter()->alignedNewArrayCalled(s, a);
   return operator new(s, av);
 }
 
-void operator delete[](void* p, std::align_val_t av) TEST_NOEXCEPT
+void operator delete[](void* p, cuda::std::align_val_t av) TEST_NOEXCEPT
 {
-  const std::size_t a = static_cast<std::size_t>(av);
+  const cuda::std::size_t a = static_cast<cuda::std::size_t>(av);
   getGlobalMemCounter()->alignedDeleteArrayCalled(p, a);
   return operator delete(p, av);
 }
@@ -456,7 +461,7 @@ void operator delete[](void* p, std::align_val_t av) TEST_NOEXCEPT
 
 struct DisableAllocationGuard
 {
-  explicit DisableAllocationGuard(bool disable = true)
+  __host__ __device__ explicit DisableAllocationGuard(bool disable = true)
       : m_disabled(disable)
   {
     // Don't re-disable if already disabled.
@@ -470,7 +475,7 @@ struct DisableAllocationGuard
     }
   }
 
-  void release()
+  __host__ __device__ void release()
   {
     if (m_disabled)
     {
@@ -479,7 +484,7 @@ struct DisableAllocationGuard
     m_disabled = false;
   }
 
-  ~DisableAllocationGuard()
+  __host__ __device__ ~DisableAllocationGuard()
   {
     release();
   }
@@ -487,45 +492,45 @@ struct DisableAllocationGuard
 private:
   bool m_disabled;
 
-  DisableAllocationGuard(DisableAllocationGuard const&);
-  DisableAllocationGuard& operator=(DisableAllocationGuard const&);
+  __host__ __device__ DisableAllocationGuard(DisableAllocationGuard const&)            = delete;
+  __host__ __device__ DisableAllocationGuard& operator=(DisableAllocationGuard const&) = delete;
 };
 
 struct RequireAllocationGuard
 {
-  explicit RequireAllocationGuard(std::size_t RequireAtLeast = 1)
+  __host__ __device__ explicit RequireAllocationGuard(cuda::std::size_t RequireAtLeast = 1)
       : m_req_alloc(RequireAtLeast)
       , m_new_count_on_init(globalMemCounter.new_called)
       , m_outstanding_new_on_init(globalMemCounter.outstanding_new)
       , m_exactly(false)
   {}
 
-  void requireAtLeast(std::size_t N)
+  __host__ __device__ void requireAtLeast(cuda::std::size_t N)
   {
     m_req_alloc = N;
     m_exactly   = false;
   }
-  void requireExactly(std::size_t N)
+  __host__ __device__ void requireExactly(cuda::std::size_t N)
   {
     m_req_alloc = N;
     m_exactly   = true;
   }
 
-  ~RequireAllocationGuard()
+  __host__ __device__ ~RequireAllocationGuard()
   {
     assert(globalMemCounter.checkOutstandingNewEq(static_cast<int>(m_outstanding_new_on_init)));
-    std::size_t Expect = m_new_count_on_init + m_req_alloc;
+    cuda::std::size_t Expect = m_new_count_on_init + m_req_alloc;
     assert(globalMemCounter.checkNewCalledEq(static_cast<int>(Expect))
            || (!m_exactly && globalMemCounter.checkNewCalledGreaterThan(static_cast<int>(Expect))));
   }
 
 private:
-  std::size_t m_req_alloc;
-  const std::size_t m_new_count_on_init;
-  const std::size_t m_outstanding_new_on_init;
+  cuda::std::size_t m_req_alloc;
+  const cuda::std::size_t m_new_count_on_init;
+  const cuda::std::size_t m_outstanding_new_on_init;
   bool m_exactly;
-  RequireAllocationGuard(RequireAllocationGuard const&);
-  RequireAllocationGuard& operator=(RequireAllocationGuard const&);
+  __host__ __device__ RequireAllocationGuard(RequireAllocationGuard const&)            = delete;
+  __host__ __device__ RequireAllocationGuard& operator=(RequireAllocationGuard const&) = delete;
 };
 
 #endif /* COUNT_NEW_H */
diff --git a/libcudacxx/test/support/min_allocator.h b/libcudacxx/test/support/min_allocator.h
index 8291ffecad..b233417aee 100644
--- a/libcudacxx/test/support/min_allocator.h
+++ b/libcudacxx/test/support/min_allocator.h
@@ -9,12 +9,11 @@
 #ifndef MIN_ALLOCATOR_H
 #define MIN_ALLOCATOR_H
 
+#include <cuda/std/__memory>
 #include <cuda/std/cassert>
-
-#include <climits>
-#include <cstddef>
-#include <cstdlib>
-#include <memory>
+#include <cuda/std/climits>
+#include <cuda/std/cstddef>
+#include <cuda/std/cstdlib>
 
 #include "test_macros.h"
 
@@ -24,29 +23,29 @@ class bare_allocator
 public:
   typedef T value_type;
 
-  bare_allocator() TEST_NOEXCEPT {}
+  __host__ __device__ bare_allocator() TEST_NOEXCEPT {}
 
   template <class U>
-  bare_allocator(bare_allocator<U>) TEST_NOEXCEPT
+  __host__ __device__ bare_allocator(bare_allocator<U>) TEST_NOEXCEPT
   {}
 
-  T* allocate(std::size_t n)
+  __host__ __device__ T* allocate(cuda::std::size_t n)
   {
     return static_cast<T*>(::operator new(n * sizeof(T)));
   }
 
-  void deallocate(T* p, std::size_t)
+  __host__ __device__ void deallocate(T* p, cuda::std::size_t)
   {
     return ::operator delete(static_cast<void*>(p));
   }
 
-  friend bool operator==(bare_allocator, bare_allocator)
+  __host__ __device__ friend bool operator==(bare_allocator, bare_allocator)
   {
     return true;
   }
-  friend bool operator!=(bare_allocator x, bare_allocator y)
+  __host__ __device__ friend bool operator!=(bare_allocator, bare_allocator)
   {
-    return !(x == y);
+    return false;
   }
 };
 
@@ -56,10 +55,10 @@ class no_default_allocator
   no_default_allocator() = delete;
   struct construct_tag
   {};
-  explicit no_default_allocator(construct_tag) {}
+  __host__ __device__ explicit no_default_allocator(construct_tag) {}
 
 public:
-  static no_default_allocator create()
+  __host__ __device__ static no_default_allocator create()
   {
     construct_tag tag;
     return no_default_allocator(tag);
@@ -69,154 +68,143 @@ class no_default_allocator
   typedef T value_type;
 
   template <class U>
-  no_default_allocator(no_default_allocator<U>) TEST_NOEXCEPT
+  __host__ __device__ no_default_allocator(no_default_allocator<U>) TEST_NOEXCEPT
   {}
 
-  T* allocate(std::size_t n)
+  __host__ __device__ T* allocate(cuda::std::size_t n)
   {
     return static_cast<T*>(::operator new(n * sizeof(T)));
   }
 
-  void deallocate(T* p, std::size_t)
+  __host__ __device__ void deallocate(T* p, cuda::std::size_t)
   {
     return ::operator delete(static_cast<void*>(p));
   }
 
-  friend bool operator==(no_default_allocator, no_default_allocator)
+  __host__ __device__ friend bool operator==(no_default_allocator, no_default_allocator)
   {
     return true;
   }
-  friend bool operator!=(no_default_allocator x, no_default_allocator y)
+  __host__ __device__ friend bool operator!=(no_default_allocator, no_default_allocator)
   {
-    return !(x == y);
+    return false;
   }
 };
 
+STATIC_TEST_GLOBAL_VAR size_t malloc_allocator_base_outstanding_bytes         = 0;
+STATIC_TEST_GLOBAL_VAR size_t malloc_allocator_base_alloc_count               = 0;
+STATIC_TEST_GLOBAL_VAR size_t malloc_allocator_base_dealloc_count             = 0;
+STATIC_TEST_GLOBAL_VAR bool malloc_allocator_base_disable_default_constructor = false;
+
 struct malloc_allocator_base
 {
-  static size_t outstanding_bytes;
-  static size_t alloc_count;
-  static size_t dealloc_count;
-  static bool disable_default_constructor;
-
-  static size_t outstanding_alloc()
+  __host__ __device__ static size_t outstanding_alloc()
   {
-    assert(alloc_count >= dealloc_count);
-    return (alloc_count - dealloc_count);
+    assert(malloc_allocator_base_alloc_count >= malloc_allocator_base_dealloc_count);
+    return (malloc_allocator_base_alloc_count - malloc_allocator_base_dealloc_count);
   }
 
-  static void reset()
+  __host__ __device__ static void reset()
   {
     assert(outstanding_alloc() == 0);
-    disable_default_constructor = false;
-    outstanding_bytes           = 0;
-    alloc_count                 = 0;
-    dealloc_count               = 0;
+    malloc_allocator_base_disable_default_constructor = false;
+    malloc_allocator_base_outstanding_bytes           = 0;
+    malloc_allocator_base_alloc_count                 = 0;
+    malloc_allocator_base_dealloc_count               = 0;
   }
 };
 
-size_t malloc_allocator_base::outstanding_bytes         = 0;
-size_t malloc_allocator_base::alloc_count               = 0;
-size_t malloc_allocator_base::dealloc_count             = 0;
-bool malloc_allocator_base::disable_default_constructor = false;
-
 template <class T>
 class malloc_allocator : public malloc_allocator_base
 {
 public:
   typedef T value_type;
 
-  malloc_allocator() TEST_NOEXCEPT
+  __host__ __device__ malloc_allocator() TEST_NOEXCEPT
   {
-    assert(!disable_default_constructor);
+    assert(!malloc_allocator_base_disable_default_constructor);
   }
 
   template <class U>
-  malloc_allocator(malloc_allocator<U>) TEST_NOEXCEPT
+  __host__ __device__ malloc_allocator(malloc_allocator<U>) TEST_NOEXCEPT
   {}
 
-  T* allocate(std::size_t n)
+  __host__ __device__ T* allocate(cuda::std::size_t n)
   {
     const size_t nbytes = n * sizeof(T);
-    ++alloc_count;
-    outstanding_bytes += nbytes;
-    return static_cast<T*>(std::malloc(nbytes));
+    ++malloc_allocator_base_alloc_count;
+    malloc_allocator_base_outstanding_bytes += nbytes;
+    return static_cast<T*>(cuda::std::malloc(nbytes));
   }
 
-  void deallocate(T* p, std::size_t n)
+  __host__ __device__ void deallocate(T* p, cuda::std::size_t n)
   {
     const size_t nbytes = n * sizeof(T);
-    ++dealloc_count;
-    outstanding_bytes -= nbytes;
-    std::free(static_cast<void*>(p));
+    ++malloc_allocator_base_dealloc_count;
+    malloc_allocator_base_outstanding_bytes -= nbytes;
+    cuda::std::free(static_cast<void*>(p));
   }
 
-  friend bool operator==(malloc_allocator, malloc_allocator)
+  __host__ __device__ friend bool operator==(malloc_allocator, malloc_allocator)
   {
     return true;
   }
-  friend bool operator!=(malloc_allocator x, malloc_allocator y)
+  __host__ __device__ friend bool operator!=(malloc_allocator, malloc_allocator)
   {
-    return !(x == y);
+    return false;
   }
 };
 
+STATIC_TEST_GLOBAL_VAR bool cpp03_allocator_construct_called = false;
 template <class T>
 struct cpp03_allocator : bare_allocator<T>
 {
   typedef T value_type;
   typedef value_type* pointer;
 
-  static bool construct_called;
-
   // Returned value is not used but it's not prohibited.
-  pointer construct(pointer p, const value_type& val)
+  __host__ __device__ pointer construct(pointer p, const value_type& val)
   {
     ::new (p) value_type(val);
-    construct_called = true;
+    cpp03_allocator_construct_called = true;
     return p;
   }
 
-  std::size_t max_size() const
+  __host__ __device__ cuda::std::size_t max_size() const
   {
     return UINT_MAX / sizeof(T);
   }
 };
-template <class T>
-bool cpp03_allocator<T>::construct_called = false;
 
+STATIC_TEST_GLOBAL_VAR bool cpp03_overload_allocator_construct_called = false;
 template <class T>
 struct cpp03_overload_allocator : bare_allocator<T>
 {
   typedef T value_type;
   typedef value_type* pointer;
 
-  static bool construct_called;
-
-  void construct(pointer p, const value_type& val)
+  __host__ __device__ void construct(pointer p, const value_type& val)
   {
-    construct(p, val, std::is_class<T>());
+    construct(p, val, cuda::std::is_class<T>());
   }
-  void construct(pointer p, const value_type& val, std::true_type)
+  __host__ __device__ void construct(pointer p, const value_type& val, cuda::std::true_type)
   {
     ::new (p) value_type(val);
-    construct_called = true;
+    cpp03_overload_allocator_construct_called = true;
   }
-  void construct(pointer p, const value_type& val, std::false_type)
+  __host__ __device__ void construct(pointer p, const value_type& val, cuda::std::false_type)
   {
     ::new (p) value_type(val);
-    construct_called = true;
+    cpp03_overload_allocator_construct_called = true;
   }
 
-  std::size_t max_size() const
+  __host__ __device__ cuda::std::size_t max_size() const
   {
     return UINT_MAX / sizeof(T);
   }
 };
-template <class T>
-bool cpp03_overload_allocator<T>::construct_called = false;
 
-template <class T, class = std::integral_constant<size_t, 0>>
+template <class T, class = cuda::std::integral_constant<size_t, 0>>
 class min_pointer;
 template <class T, class ID>
 class min_pointer<const T, ID>;
@@ -234,21 +222,21 @@ class min_pointer<const void, ID>
 
 public:
   min_pointer() TEST_NOEXCEPT = default;
-  min_pointer(std::nullptr_t) TEST_NOEXCEPT : ptr_(nullptr) {}
+  __host__ __device__ min_pointer(cuda::std::nullptr_t) TEST_NOEXCEPT : ptr_(nullptr) {}
   template <class T>
-  min_pointer(min_pointer<T, ID> p) TEST_NOEXCEPT : ptr_(p.ptr_)
+  __host__ __device__ min_pointer(min_pointer<T, ID> p) TEST_NOEXCEPT : ptr_(p.ptr_)
   {}
 
-  explicit operator bool() const
+  __host__ __device__ explicit operator bool() const
   {
     return ptr_ != nullptr;
   }
 
-  friend bool operator==(min_pointer x, min_pointer y)
+  __host__ __device__ friend bool operator==(min_pointer x, min_pointer y)
   {
     return x.ptr_ == y.ptr_;
   }
-  friend bool operator!=(min_pointer x, min_pointer y)
+  __host__ __device__ friend bool operator!=(min_pointer x, min_pointer y)
   {
     return !(x == y);
   }
@@ -263,21 +251,21 @@ class min_pointer<void, ID>
 
 public:
   min_pointer() TEST_NOEXCEPT = default;
-  min_pointer(std::nullptr_t) TEST_NOEXCEPT : ptr_(nullptr) {}
-  template <class T, class = typename std::enable_if<!std::is_const<T>::value>::type>
-  min_pointer(min_pointer<T, ID> p) TEST_NOEXCEPT : ptr_(p.ptr_)
+  __host__ __device__ min_pointer(cuda::std::nullptr_t) TEST_NOEXCEPT : ptr_(nullptr) {}
+  template <class T, class = typename cuda::std::enable_if<!cuda::std::is_const<T>::value>::type>
+  __host__ __device__ min_pointer(min_pointer<T, ID> p) TEST_NOEXCEPT : ptr_(p.ptr_)
   {}
 
-  explicit operator bool() const
+  __host__ __device__ explicit operator bool() const
   {
     return ptr_ != nullptr;
   }
 
-  friend bool operator==(min_pointer x, min_pointer y)
+  __host__ __device__ friend bool operator==(min_pointer x, min_pointer y)
   {
     return x.ptr_ == y.ptr_;
   }
-  friend bool operator!=(min_pointer x, min_pointer y)
+  __host__ __device__ friend bool operator!=(min_pointer x, min_pointer y)
   {
     return !(x == y);
   }
@@ -290,124 +278,124 @@ class min_pointer
 {
   T* ptr_;
 
-  explicit min_pointer(T* p) TEST_NOEXCEPT : ptr_(p) {}
+  __host__ __device__ explicit min_pointer(T* p) TEST_NOEXCEPT : ptr_(p) {}
 
 public:
   min_pointer() TEST_NOEXCEPT = default;
-  min_pointer(std::nullptr_t) TEST_NOEXCEPT : ptr_(nullptr) {}
-  explicit min_pointer(min_pointer<void, ID> p) TEST_NOEXCEPT : ptr_(static_cast<T*>(p.ptr_)) {}
+  __host__ __device__ min_pointer(cuda::std::nullptr_t) TEST_NOEXCEPT : ptr_(nullptr) {}
+  __host__ __device__ explicit min_pointer(min_pointer<void, ID> p) TEST_NOEXCEPT : ptr_(static_cast<T*>(p.ptr_)) {}
 
-  explicit operator bool() const
+  __host__ __device__ explicit operator bool() const
   {
     return ptr_ != nullptr;
   }
 
-  typedef std::ptrdiff_t difference_type;
+  typedef cuda::std::ptrdiff_t difference_type;
   typedef T& reference;
   typedef T* pointer;
   typedef T value_type;
-  typedef std::random_access_iterator_tag iterator_category;
+  typedef cuda::std::random_access_iterator_tag iterator_category;
 
-  reference operator*() const
+  __host__ __device__ reference operator*() const
   {
     return *ptr_;
   }
-  pointer operator->() const
+  __host__ __device__ pointer operator->() const
   {
     return ptr_;
   }
 
-  min_pointer& operator++()
+  __host__ __device__ min_pointer& operator++()
   {
     ++ptr_;
     return *this;
   }
-  min_pointer operator++(int)
+  __host__ __device__ min_pointer operator++(int)
   {
     min_pointer tmp(*this);
     ++ptr_;
     return tmp;
   }
 
-  min_pointer& operator--()
+  __host__ __device__ min_pointer& operator--()
   {
     --ptr_;
     return *this;
   }
-  min_pointer operator--(int)
+  __host__ __device__ min_pointer operator--(int)
   {
     min_pointer tmp(*this);
     --ptr_;
     return tmp;
   }
 
-  min_pointer& operator+=(difference_type n)
+  __host__ __device__ min_pointer& operator+=(difference_type n)
   {
     ptr_ += n;
     return *this;
   }
-  min_pointer& operator-=(difference_type n)
+  __host__ __device__ min_pointer& operator-=(difference_type n)
   {
     ptr_ -= n;
     return *this;
   }
 
-  min_pointer operator+(difference_type n) const
+  __host__ __device__ min_pointer operator+(difference_type n) const
   {
     min_pointer tmp(*this);
     tmp += n;
     return tmp;
   }
 
-  friend min_pointer operator+(difference_type n, min_pointer x)
+  __host__ __device__ friend min_pointer operator+(difference_type n, min_pointer x)
   {
     return x + n;
   }
 
-  min_pointer operator-(difference_type n) const
+  __host__ __device__ min_pointer operator-(difference_type n) const
   {
     min_pointer tmp(*this);
     tmp -= n;
     return tmp;
   }
 
-  friend difference_type operator-(min_pointer x, min_pointer y)
+  __host__ __device__ friend difference_type operator-(min_pointer x, min_pointer y)
   {
     return x.ptr_ - y.ptr_;
   }
 
-  reference operator[](difference_type n) const
+  __host__ __device__ reference operator[](difference_type n) const
   {
     return ptr_[n];
   }
 
-  friend bool operator<(min_pointer x, min_pointer y)
+  __host__ __device__ friend bool operator<(min_pointer x, min_pointer y)
   {
     return x.ptr_ < y.ptr_;
   }
-  friend bool operator>(min_pointer x, min_pointer y)
+  __host__ __device__ friend bool operator>(min_pointer x, min_pointer y)
   {
     return y < x;
   }
-  friend bool operator<=(min_pointer x, min_pointer y)
+  __host__ __device__ friend bool operator<=(min_pointer x, min_pointer y)
   {
     return !(y < x);
   }
-  friend bool operator>=(min_pointer x, min_pointer y)
+  __host__ __device__ friend bool operator>=(min_pointer x, min_pointer y)
   {
     return !(x < y);
   }
 
-  static min_pointer pointer_to(T& t)
+  __host__ __device__ static min_pointer pointer_to(T& t)
   {
-    return min_pointer(std::addressof(t));
+    return min_pointer(cuda::std::addressof(t));
   }
 
-  friend bool operator==(min_pointer x, min_pointer y)
+  __host__ __device__ friend bool operator==(min_pointer x, min_pointer y)
   {
     return x.ptr_ == y.ptr_;
   }
-  friend bool operator!=(min_pointer x, min_pointer y)
+  __host__ __device__ friend bool operator!=(min_pointer x, min_pointer y)
   {
     return !(x == y);
   }
@@ -422,133 +410,133 @@ class min_pointer<const T, ID>
 {
   const T* ptr_;
 
-  explicit min_pointer(const T* p)
+  __host__ __device__ explicit min_pointer(const T* p)
       : ptr_(p)
   {}
 
 public:
   min_pointer() TEST_NOEXCEPT = default;
-  min_pointer(std::nullptr_t)
+  __host__ __device__ min_pointer(cuda::std::nullptr_t)
       : ptr_(nullptr)
   {}
-  min_pointer(min_pointer<T, ID> p)
+  __host__ __device__ min_pointer(min_pointer<T, ID> p)
       : ptr_(p.ptr_)
   {}
-  explicit min_pointer(min_pointer<const void, ID> p)
+  __host__ __device__ explicit min_pointer(min_pointer<const void, ID> p)
       : ptr_(static_cast<const T*>(p.ptr_))
   {}
 
-  explicit operator bool() const
+  __host__ __device__ explicit operator bool() const
   {
     return ptr_ != nullptr;
   }
 
-  typedef std::ptrdiff_t difference_type;
+  typedef cuda::std::ptrdiff_t difference_type;
   typedef const T& reference;
   typedef const T* pointer;
   typedef const T value_type;
-  typedef std::random_access_iterator_tag iterator_category;
+  typedef cuda::std::random_access_iterator_tag iterator_category;
 
-  reference operator*() const
+  __host__ __device__ reference operator*() const
   {
     return *ptr_;
   }
-  pointer operator->() const
+  __host__ __device__ pointer operator->() const
   {
     return ptr_;
   }
 
-  min_pointer& operator++()
+  __host__ __device__ min_pointer& operator++()
   {
     ++ptr_;
     return *this;
   }
-  min_pointer operator++(int)
+  __host__ __device__ min_pointer operator++(int)
   {
     min_pointer tmp(*this);
     ++ptr_;
     return tmp;
   }
 
-  min_pointer& operator--()
+  __host__ __device__ min_pointer& operator--()
   {
     --ptr_;
     return *this;
   }
-  min_pointer operator--(int)
+  __host__ __device__ min_pointer operator--(int)
   {
     min_pointer tmp(*this);
     --ptr_;
     return tmp;
   }
 
-  min_pointer& operator+=(difference_type n)
+  __host__ __device__ min_pointer& operator+=(difference_type n)
   {
     ptr_ += n;
     return *this;
   }
-  min_pointer& operator-=(difference_type n)
+  __host__ __device__ min_pointer& operator-=(difference_type n)
   {
     ptr_ -= n;
     return *this;
   }
 
-  min_pointer operator+(difference_type n) const
+  __host__ __device__ min_pointer operator+(difference_type n) const
   {
     min_pointer tmp(*this);
     tmp += n;
     return tmp;
   }
 
-  friend min_pointer operator+(difference_type n, min_pointer x)
+  __host__ __device__ friend min_pointer operator+(difference_type n, min_pointer x)
   {
     return x + n;
   }
 
-  min_pointer operator-(difference_type n) const
+  __host__ __device__ min_pointer operator-(difference_type n) const
   {
     min_pointer tmp(*this);
     tmp -= n;
     return tmp;
   }
 
-  friend difference_type operator-(min_pointer x, min_pointer y)
+  __host__ __device__ friend difference_type operator-(min_pointer x, min_pointer y)
   {
     return x.ptr_ - y.ptr_;
   }
 
-  reference operator[](difference_type n) const
+  __host__ __device__ reference operator[](difference_type n) const
   {
     return ptr_[n];
   }
 
-  friend bool operator<(min_pointer x, min_pointer y)
+  __host__ __device__ friend bool operator<(min_pointer x, min_pointer y)
   {
     return x.ptr_ < y.ptr_;
   }
-  friend bool operator>(min_pointer x, min_pointer y)
+  __host__ __device__ friend bool operator>(min_pointer x, min_pointer y)
   {
     return y < x;
   }
-  friend bool operator<=(min_pointer x, min_pointer y)
+  __host__ __device__ friend bool operator<=(min_pointer x, min_pointer y)
   {
     return !(y < x);
   }
-  friend bool operator>=(min_pointer x, min_pointer y)
+  __host__ __device__ friend bool operator>=(min_pointer x, min_pointer y)
   {
     return !(x < y);
   }
 
-  static min_pointer pointer_to(const T& t)
+  __host__ __device__ static min_pointer pointer_to(const T& t)
   {
-    return min_pointer(std::addressof(t));
+    return min_pointer(cuda::std::addressof(t));
   }
 
-  friend bool operator==(min_pointer x, min_pointer y)
+  __host__ __device__ friend bool operator==(min_pointer x, min_pointer y)
   {
     return x.ptr_ == y.ptr_;
   }
-  friend bool operator!=(min_pointer x, min_pointer y)
+  __host__ __device__ friend bool operator!=(min_pointer x, min_pointer y)
   {
     return !(x == y);
   }
@@ -557,25 +545,25 @@ class min_pointer<const T, ID>
 };
 
 template <class T, class ID>
-inline bool operator==(min_pointer<T, ID> x, std::nullptr_t)
+__host__ __device__ inline bool operator==(min_pointer<T, ID> x, cuda::std::nullptr_t)
 {
   return !static_cast<bool>(x);
 }
 
 template <class T, class ID>
-inline bool operator==(std::nullptr_t, min_pointer<T, ID> x)
+__host__ __device__ inline bool operator==(cuda::std::nullptr_t, min_pointer<T, ID> x)
 {
   return !static_cast<bool>(x);
 }
 
 template <class T, class ID>
-inline bool operator!=(min_pointer<T, ID> x, std::nullptr_t)
+__host__ __device__ inline bool operator!=(min_pointer<T, ID> x, cuda::std::nullptr_t)
 {
   return static_cast<bool>(x);
 }
 
 template <class T, class ID>
-inline bool operator!=(std::nullptr_t, min_pointer<T, ID> x)
+__host__ __device__ inline bool operator!=(cuda::std::nullptr_t, min_pointer<T, ID> x)
 {
   return static_cast<bool>(x);
 }
@@ -589,26 +577,26 @@ class min_allocator
 
   min_allocator() = default;
   template <class U>
-  min_allocator(min_allocator<U>)
+  __host__ __device__ min_allocator(min_allocator<U>)
   {}
 
-  pointer allocate(std::ptrdiff_t n)
+  __host__ __device__ pointer allocate(cuda::std::ptrdiff_t n)
   {
     return pointer(static_cast<T*>(::operator new(n * sizeof(T))));
   }
 
-  void deallocate(pointer p, std::ptrdiff_t)
+  __host__ __device__ void deallocate(pointer p, cuda::std::ptrdiff_t)
   {
     return ::operator delete(p.ptr_);
   }
 
-  friend bool operator==(min_allocator, min_allocator)
+  __host__ __device__ friend bool operator==(min_allocator, min_allocator)
   {
     return true;
   }
-  friend bool operator!=(min_allocator x, min_allocator y)
+  __host__ __device__ friend bool operator!=(min_allocator, min_allocator)
   {
-    return !(x == y);
+    return false;
   }
 };
 
@@ -618,29 +606,29 @@ class explicit_allocator
 public:
   typedef T value_type;
 
-  explicit_allocator() TEST_NOEXCEPT {}
+  __host__ __device__ explicit_allocator() TEST_NOEXCEPT {}
 
   template <class U>
-  explicit explicit_allocator(explicit_allocator<U>) TEST_NOEXCEPT
+  __host__ __device__ explicit explicit_allocator(explicit_allocator<U>) TEST_NOEXCEPT
   {}
 
-  T* allocate(std::size_t n)
+  __host__ __device__ T* allocate(cuda::std::size_t n)
   {
     return static_cast<T*>(::operator new(n * sizeof(T)));
   }
 
-  void deallocate(T* p, std::size_t)
+  __host__ __device__ void deallocate(T* p, cuda::std::size_t)
   {
     return ::operator delete(static_cast<void*>(p));
   }
 
-  friend bool operator==(explicit_allocator, explicit_allocator)
+  __host__ __device__ friend bool operator==(explicit_allocator, explicit_allocator)
   {
     return true;
   }
-  friend bool operator!=(explicit_allocator x, explicit_allocator y)
+  __host__ __device__ friend bool operator!=(explicit_allocator, explicit_allocator)
   {
-    return !(x == y);
+    return false;
   }
 };
 
diff --git a/libcudacxx/test/support/test_macros.h b/libcudacxx/test/support/test_macros.h
index ca566ac177..1a946f7741 100644
--- a/libcudacxx/test/support/test_macros.h
+++ b/libcudacxx/test/support/test_macros.h
@@ -171,6 +171,12 @@
 #else
 #  define TEST_CONSTEXPR_CXX20
 #endif
+#if defined(__cpp_constexpr_dynamic_alloc) && defined(__cpp_lib_constexpr_dynamic_alloc) && TEST_STD_VER >= 2020 \
+  && !defined(TEST_COMPILER_NVRTC)
+#  define TEST_CONSTEXPR_CXX20_ALLOC constexpr
+#else
+#  define TEST_CONSTEXPR_CXX20_ALLOC
+#endif
 #if TEST_STD_VER > 2014
 #  define TEST_THROW_SPEC(...)
 #else