diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/CMakeLists.txt b/libcudacxx/include/cuda/std/detail/libcxx/include/CMakeLists.txt
index 3ca73fe9bc6..364e6b7494f 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/CMakeLists.txt
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/CMakeLists.txt
@@ -243,6 +243,7 @@ set(files
   __memory/allocator_arg_t.h
   __memory/allocator_destructor.h
   __memory/allocator_traits.h
+  __memory/allocator.h
   __memory/construct_at.h
   __memory/pointer_traits.h
   __memory/uses_allocator.h
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__config b/libcudacxx/include/cuda/std/detail/libcxx/include/__config
index 8309ff800b0..5a92cfbc1ae 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/__config
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__config
@@ -1626,7 +1626,8 @@ typedef unsigned int   char32_t;
 
 #if !defined(_LIBCUDACXX_HAS_NO_ALIGNED_ALLOCATION) && \
     (defined(_LIBCUDACXX_HAS_NO_LIBRARY_ALIGNED_ALLOCATION) || \
-    (!defined(__cpp_aligned_new) || __cpp_aligned_new < 201606))
+    (!defined(__cpp_aligned_new) || __cpp_aligned_new < 201606)) \
+  || defined(__cuda_std__) // FIXME: Properly handle aligned allocations
 #  define _LIBCUDACXX_HAS_NO_ALIGNED_ALLOCATION
 #endif
 
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 00000000000..dd4ba3876a3
--- /dev/null
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__memory/allocator.h
@@ -0,0 +1,330 @@
+// -*- 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 _CCCL_STD_VER >= 2020 && !defined(_CCCL_COMPILER_NVRTC)
+#  include <memory>
+#endif // _CCCL_STD_VER >= 2020
+
+_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 _Tp*
+  allocate(size_t __n)
+  {
+    if (__n > allocator_traits<allocator>::max_size(*this))
+    {
+      __throw_bad_array_new_length();
+    }
+#if _CCCL_STD_VER >= 2020 && !defined(_CCCL_COMPILER_NVRTC)
+    if (_CUDA_VSTD::is_constant_evaluated())
+    {
+      return ::std::allocator<_Tp>{}.allocate(__n);
+    }
+    else
+#endif // _CCCL_STD_VER >= 2020 && !_CCCL_COMPILER_NVRTC
+    {
+      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_STD_VER >= 2023
+
+  _CCCL_EXEC_CHECK_DISABLE
+  _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 void deallocate(_Tp* __p, size_t __n) noexcept
+  {
+#if _CCCL_STD_VER >= 2020 && !defined(_CCCL_COMPILER_NVRTC)
+    if (_CUDA_VSTD::is_constant_evaluated())
+    {
+      return ::std::allocator<_Tp>{}.deallocate(__p, __n);
+    }
+    else
+#endif // _CCCL_STD_VER >= 2020 && !_CCCL_COMPILER_NVRTC
+    {
+      _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 b/libcudacxx/include/cuda/std/detail/libcxx/include/memory
index f903016996e..c4b145d4614 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/memory
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/memory
@@ -667,6 +667,7 @@ void* align(size_t alignment, size_t size, void*& ptr, size_t& space);
 #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>
@@ -746,247 +747,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
-
-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,
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
index 11b8d3cc322..3c17bef4be4 100644
--- 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
@@ -17,11 +17,8 @@
 //     ...
 // };
 
-// Currently no suppport for std::allocator
-// XFAIL: true
-
 #include <cuda/std/__memory>
-#include <cuda/std/new>
+#include <cuda/std/__new>
 #include <cuda/std/type_traits>
 #include <cuda/std/cassert>
 #include <cuda/std/utility>
@@ -174,9 +171,7 @@ int main(int, char**)
 {
     test();
 
-#if TEST_STD_VER >= 2020  \
- && !defined(TEST_COMPILER_NVCC) \
- && !defined(TEST_COMPILER_NVRTC)
+#if TEST_STD_VER >= 2020
     static_assert(test());
 #endif // TEST_STD_VER >= 2020
 
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
index fa91a7890c6..e48d4027719 100644
--- 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
@@ -18,7 +18,6 @@
 // };
 
 // Currently no suppport for std::allocator
-// XFAIL: true
 
 #include <cuda/std/__memory>
 #include <cuda/std/cassert>
@@ -131,9 +130,7 @@ __host__ __device__ TEST_CONSTEXPR_CXX20 bool test()
 int main(int, char**)
 {
     test();
-#if TEST_STD_VER >= 2020  \
- && !defined(TEST_COMPILER_NVCC) \
- && !defined(TEST_COMPILER_NVRTC)
+#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/PR50299.compile.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/PR50299.compile.pass.cpp
new file mode 100644
index 00000000000..1738c7d4209
--- /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> a;
+
+int main(int, char**)
+{
+    unused(a);
+    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 00000000000..cd7baf53795
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.ctor.pass.cpp
@@ -0,0 +1,51 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 00000000000..e2fb258cfcc
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.dtor.pass.cpp
@@ -0,0 +1,44 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 00000000000..e4269f86687
--- /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 00000000000..82718f39d90
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.members/allocate.constexpr.size.verify.cpp
@@ -0,0 +1,46 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/cstddef>
+#include <cuda/std/__memory>
+
+#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 00000000000..0e070e6715d
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.members/allocate.pass.cpp
@@ -0,0 +1,121 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 "test_macros.h"
+#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_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 TEST_STD_VER >= 2020 \
+ && !defined(TEST_COMPILER_NVCC) \
+ && !defined(TEST_COMPILER_NVRTC)
+    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 // TEST_STD_VER >= 2020
+
+  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 00000000000..1def4d734f5
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.members/allocate.size.pass.cpp
@@ -0,0 +1,58 @@
+//===----------------------------------------------------------------------===//
+//
+// 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"
+
+template <typename T>
+__host__ __device__ void test_max(cuda::std::size_t count)
+{
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    cuda::std::allocator<T> a;
+    try {
+        TEST_IGNORE_NODISCARD a.allocate(count);
+        assert(false);
+    } catch (const cuda::std::bad_array_new_length &) {
+    }
+#else
+    unused(count);
+#endif // TEST_HAS_NO_EXCEPTIONS
+}
+
+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
+}
+
+int main(int, char**)
+{
+    test<double>();
+    test<const double>();
+
+    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 00000000000..e619c2d1f13
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.members/allocate.verify.cpp
@@ -0,0 +1,29 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 00000000000..58f9bfa22c9
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator.members/allocate_at_least.pass.cpp
@@ -0,0 +1,111 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/cassert>
+#include <cuda/std/concepts>
+#include <cuda/std/__memory>
+
+#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 00000000000..35b48dabc8a
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator_pointers.pass.cpp
@@ -0,0 +1,126 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 00000000000..286b74c614a
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator_types.deprecated_in_cxx17.verify.cpp
@@ -0,0 +1,58 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 00000000000..74250563830
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator_types.pass.cpp
@@ -0,0 +1,68 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/type_traits>
+#include <cuda/std/cstddef>
+
+#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 <= 17
+    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
+}
+
+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 00000000000..bdbd4b58585
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/memory/default.allocator/allocator_types.removed_in_cxx20.verify.cpp
@@ -0,0 +1,51 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 00000000000..881ce9e7c4e
--- /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/support/count_new.h b/libcudacxx/test/support/count_new.h
index e8052544f95..875c1e424ba 100644
--- a/libcudacxx/test/support/count_new.h
+++ b/libcudacxx/test/support/count_new.h
@@ -9,79 +9,89 @@
 #ifndef COUNT_NEW_H
 #define COUNT_NEW_H
 
-# include <cstdlib>
-# include <cassert>
+# include <cuda/std/cstdlib>
+# include <cuda/std/cassert>
+
+#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(TEST_COMPILER_NVRTC)
 #define DISABLE_NEW_COUNT
 #endif
 
 namespace detail
 {
    TEST_NORETURN
-   inline void throw_bad_alloc_helper() {
+   __host__ __device__ inline void throw_bad_alloc_helper() {
 #ifndef TEST_HAS_NO_EXCEPTIONS
-       throw std::bad_alloc();
+       throw cuda::std::bad_alloc();
 #else
-       std::abort();
+       _LIBCUDACXX_UNREACHABLE();
 #endif
    }
 }
 
+// All checks return true when disable_checking is enabled.
+#ifdef DISABLE_NEW_COUNT
+STATIC_TEST_GLOBAL_VAR const bool MemCounter_disable_checking = true;
+#else
+STATIC_TEST_GLOBAL_VAR const bool MemCounter_disable_checking = false;
+#endif
+
+// 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;
+            throw_after = MemCounter_never_throw_value;
             detail::throw_bad_alloc_helper();
-        } else if (throw_after != never_throw_value) {
+        } else if (throw_after != MemCounter_never_throw_value) {
             --throw_after;
         }
         ++new_called;
@@ -89,31 +99,31 @@ 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;
+            throw_after = MemCounter_never_throw_value;
             detail::throw_bad_alloc_helper();
         } else {
             // don't decrement throw_after here. newCalled will end up doing that.
@@ -123,39 +133,39 @@ 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;
@@ -175,212 +185,223 @@ 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() {
-  static MemCounter counter((MemCounter::MemCounterCtorArg_()));
+__host__ __device__ inline constexpr MemCounter* getGlobalMemCounter() {
   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)
+__host__ __device__ 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();
     return ret;
 }
 
-void  operator delete(void* p) TEST_NOEXCEPT
+__host__ __device__ void  operator delete(void* p) TEST_NOEXCEPT
+{
+    getGlobalMemCounter()->deleteCalled(p);
+    free(p);
+}
+
+#ifdef TEST_COMPILER_GCC
+__host__ __device__ void  operator delete(void* p, cuda::std::size_t) TEST_NOEXCEPT
 {
     getGlobalMemCounter()->deleteCalled(p);
-    std::free(p);
+    free(p);
 }
+#endif // TEST_COMPILER_GCC
 
-void* operator new[](std::size_t s) TEST_THROW_SPEC(std::bad_alloc)
+__host__ __device__ void* operator new[](cuda::std::size_t s)
 {
     getGlobalMemCounter()->newArrayCalled(s);
     return operator new(s);
 }
 
-void operator delete[](void* p) TEST_NOEXCEPT
+__host__ __device__ void operator delete[](void* p) TEST_NOEXCEPT
+{
+    getGlobalMemCounter()->deleteArrayCalled(p);
+    operator delete(p);
+}
+
+#ifdef TEST_COMPILER_GCC
+__host__ __device__ 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) || \
@@ -388,8 +409,8 @@ void operator delete[](void* p) TEST_NOEXCEPT
 #define USE_ALIGNED_ALLOC
 #endif
 
-void* operator new(std::size_t s, std::align_val_t av) TEST_THROW_SPEC(std::bad_alloc) {
-  const std::size_t a = static_cast<std::size_t>(av);
+__host__ __device__ void* operator new(cuda::std::size_t s, cuda::std::align_val_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
@@ -402,8 +423,8 @@ void* operator new(std::size_t s, std::align_val_t av) TEST_THROW_SPEC(std::bad_
   return ret;
 }
 
-void operator delete(void *p, std::align_val_t av) TEST_NOEXCEPT {
-  const std::size_t a = static_cast<std::size_t>(av);
+__host__ __device__ void operator delete(void *p, cuda::std::align_val_t av) TEST_NOEXCEPT {
+  const cuda::std::size_t a = static_cast<cuda::std::size_t>(av);
   getGlobalMemCounter()->alignedDeleteCalled(p, a);
   if (p) {
 #ifdef USE_ALIGNED_ALLOC
@@ -414,14 +435,14 @@ 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) {
-  const std::size_t a = static_cast<std::size_t>(av);
+__host__ __device__ void* operator new[](cuda::std::size_t s, cuda::std::align_val_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 {
-  const std::size_t a = static_cast<std::size_t>(av);
+__host__ __device__ void operator delete[](void *p, cuda::std::align_val_t av) TEST_NOEXCEPT {
+  const cuda::std::size_t a = static_cast<cuda::std::size_t>(av);
   getGlobalMemCounter()->alignedDeleteArrayCalled(p, a);
   return operator delete(p, av);
 }
@@ -431,31 +452,31 @@ void operator delete[](void *p, std::align_val_t av) TEST_NOEXCEPT {
 #endif // DISABLE_NEW_COUNT
 
 struct DisableAllocationGuard {
-    explicit DisableAllocationGuard(bool disable = true) : m_disabled(disable)
+    __host__ __device__ explicit DisableAllocationGuard(bool disable = true) : m_disabled(disable)
     {
         // Don't re-disable if already disabled.
         if (globalMemCounter.disable_allocations == true) m_disabled = false;
         if (m_disabled) globalMemCounter.disableAllocations();
     }
 
-    void release() {
+    __host__ __device__ void release() {
         if (m_disabled) globalMemCounter.enableAllocations();
         m_disabled = false;
     }
 
-    ~DisableAllocationGuard() {
+    __host__ __device__ ~DisableAllocationGuard() {
         release();
     }
 
 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),
@@ -463,23 +484,23 @@ struct RequireAllocationGuard {
     {
     }
 
-    void requireAtLeast(std::size_t N) { m_req_alloc = N; m_exactly = false; }
-    void requireExactly(std::size_t N) { m_req_alloc = N; m_exactly = true; }
+    __host__ __device__ void requireAtLeast(cuda::std::size_t N) { m_req_alloc = N; m_exactly = false; }
+    __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 */