diff --git a/cudax/include/cuda/experimental/__memory_resource/any_resource.cuh b/cudax/include/cuda/experimental/__memory_resource/any_resource.cuh
index eafd626739..fcb1ed47ce 100644
--- a/cudax/include/cuda/experimental/__memory_resource/any_resource.cuh
+++ b/cudax/include/cuda/experimental/__memory_resource/any_resource.cuh
@@ -118,7 +118,7 @@ public:
           nullptr, _CUDA_VSTD::exchange(__other.__static_vtable, nullptr))
       , __vtable(__other)
   {
-    _LIBCUDACXX_ASSERT(this->__static_vtable != nullptr, "copying from a moved-from object");
+    _CCCL_ASSERT(this->__static_vtable != nullptr, "copying from a moved-from object");
     this->__static_vtable->__move_fn(&this->__object, &__other.__object);
   }
 
@@ -129,7 +129,7 @@ public:
           nullptr, _CUDA_VSTD::exchange(__other.__static_vtable, nullptr))
       , __vtable(__other)
   {
-    _LIBCUDACXX_ASSERT(this->__static_vtable != nullptr, "copying from a moved-from object");
+    _CCCL_ASSERT(this->__static_vtable != nullptr, "copying from a moved-from object");
     this->__static_vtable->__move_fn(&this->__object, &__other.__object);
   }
 
@@ -158,7 +158,7 @@ public:
       : _CUDA_VMR::_Resource_base<_Alloc_type, _CUDA_VMR::_WrapperType::_Owning>(nullptr, __other.__static_vtable)
       , __vtable(__other)
   {
-    _LIBCUDACXX_ASSERT(this->__static_vtable != nullptr, "copying from a moved-from object");
+    _CCCL_ASSERT(this->__static_vtable != nullptr, "copying from a moved-from object");
     this->__static_vtable->__copy_fn(&this->__object, &__other.__object);
   }
 
diff --git a/libcudacxx/include/cuda/__memory_resource/device_memory_resource.h b/libcudacxx/include/cuda/__memory_resource/device_memory_resource.h
index 485cc080dc..617aab83c0 100644
--- a/libcudacxx/include/cuda/__memory_resource/device_memory_resource.h
+++ b/libcudacxx/include/cuda/__memory_resource/device_memory_resource.h
@@ -85,8 +85,7 @@ class device_memory_resource
   void deallocate(void* __ptr, const size_t, const size_t __alignment = default_cuda_malloc_alignment) const noexcept
   {
     // We need to ensure that the provided alignment matches the minimal provided alignment
-    _LIBCUDACXX_ASSERT(__is_valid_alignment(__alignment),
-                       "Invalid alignment passed to device_memory_resource::deallocate.");
+    _CCCL_ASSERT(__is_valid_alignment(__alignment), "Invalid alignment passed to device_memory_resource::deallocate.");
     _CCCL_ASSERT_CUDA_API(::cudaFree, "device_memory_resource::deallocate failed", __ptr);
     (void) __alignment;
   }
diff --git a/libcudacxx/include/cuda/__memory_resource/managed_memory_resource.h b/libcudacxx/include/cuda/__memory_resource/managed_memory_resource.h
index 977012669a..6ce6e2f002 100644
--- a/libcudacxx/include/cuda/__memory_resource/managed_memory_resource.h
+++ b/libcudacxx/include/cuda/__memory_resource/managed_memory_resource.h
@@ -51,7 +51,7 @@ class managed_memory_resource
   constexpr managed_memory_resource(const unsigned int __flags = cudaMemAttachGlobal) noexcept
       : __flags_(__flags & __available_flags)
   {
-    _LIBCUDACXX_ASSERT(__flags_ == __flags, "Unexpected flags passed to managed_memory_resource");
+    _CCCL_ASSERT(__flags_ == __flags, "Unexpected flags passed to managed_memory_resource");
   }
 
   //! @brief Allocate CUDA unified memory of size at least \p __bytes.
@@ -80,8 +80,7 @@ class managed_memory_resource
   void deallocate(void* __ptr, const size_t, const size_t __alignment = default_cuda_malloc_alignment) const noexcept
   {
     // We need to ensure that the provided alignment matches the minimal provided alignment
-    _LIBCUDACXX_ASSERT(__is_valid_alignment(__alignment),
-                       "Invalid alignment passed to managed_memory_resource::deallocate.");
+    _CCCL_ASSERT(__is_valid_alignment(__alignment), "Invalid alignment passed to managed_memory_resource::deallocate.");
     _CCCL_ASSERT_CUDA_API(::cudaFree, "managed_memory_resource::deallocate failed", __ptr);
     (void) __alignment;
   }
diff --git a/libcudacxx/include/cuda/__memory_resource/pinned_memory_resource.h b/libcudacxx/include/cuda/__memory_resource/pinned_memory_resource.h
index 557acc9a67..5cb567120a 100644
--- a/libcudacxx/include/cuda/__memory_resource/pinned_memory_resource.h
+++ b/libcudacxx/include/cuda/__memory_resource/pinned_memory_resource.h
@@ -53,7 +53,7 @@ class pinned_memory_resource
   constexpr pinned_memory_resource(const unsigned int __flags = cudaHostAllocDefault) noexcept
       : __flags_(__flags & __available_flags)
   {
-    _LIBCUDACXX_ASSERT(__flags_ == __flags, "Unexpected flags passed to pinned_memory_resource");
+    _CCCL_ASSERT(__flags_ == __flags, "Unexpected flags passed to pinned_memory_resource");
   }
 
   //! @brief Allocate host memory of size at least \p __bytes.
@@ -83,8 +83,7 @@ class pinned_memory_resource
   deallocate(void* __ptr, const size_t, const size_t __alignment = default_cuda_malloc_host_alignment) const noexcept
   {
     // We need to ensure that the provided alignment matches the minimal provided alignment
-    _LIBCUDACXX_ASSERT(__is_valid_alignment(__alignment),
-                       "Invalid alignment passed to pinned_memory_resource::deallocate.");
+    _CCCL_ASSERT(__is_valid_alignment(__alignment), "Invalid alignment passed to pinned_memory_resource::deallocate.");
     _CCCL_ASSERT_CUDA_API(::cudaFreeHost, "pinned_memory_resource::deallocate failed", __ptr);
     (void) __alignment;
   }
diff --git a/libcudacxx/include/cuda/annotated_ptr b/libcudacxx/include/cuda/annotated_ptr
index 35675b6b01..8e06473f94 100644
--- a/libcudacxx/include/cuda/annotated_ptr
+++ b/libcudacxx/include/cuda/annotated_ptr
@@ -144,9 +144,9 @@
 
 #include <cuda/barrier>
 #include <cuda/discard_memory>
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/cstdint>
 #include <cuda/std/detail/__access_property>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 
 _LIBCUDACXX_BEGIN_NAMESPACE_CUDA
 
@@ -362,8 +362,8 @@ public:
   {
     NV_IF_TARGET(
       NV_IS_DEVICE,
-      (_LIBCUDACXX_ASSERT(
-         (std::is_same<_Property, access_property::shared>::value && __isShared(__p)) || __isGlobal(__p), "");))
+      (_CCCL_ASSERT((std::is_same<_Property, access_property::shared>::value && __isShared(__p)) || __isGlobal(__p),
+                    "");))
   }
 
   template <typename _RuntimeProperty>
@@ -380,7 +380,7 @@ public:
         || std::is_same<_RuntimeProperty, access_property::persisting>::value
         || std::is_same<_RuntimeProperty, access_property>::value,
       "This method requires RuntimeProperty=global|normal|streaming|persisting|access_property");
-    NV_IF_TARGET(NV_IS_DEVICE, (_LIBCUDACXX_ASSERT((__isGlobal(__p) == true), "");))
+    NV_IF_TARGET(NV_IS_DEVICE, (_CCCL_ASSERT((__isGlobal(__p) == true), "");))
   }
 
   template <class _TTp, class _Prop>
diff --git a/libcudacxx/include/cuda/barrier b/libcudacxx/include/cuda/barrier
index 7bee17fb26..d10befdad2 100644
--- a/libcudacxx/include/cuda/barrier
+++ b/libcudacxx/include/cuda/barrier
@@ -57,9 +57,9 @@ _LIBCUDACXX_BEGIN_NAMESPACE_CUDA_DEVICE_EXPERIMENTAL
 inline _CCCL_DEVICE void cp_async_bulk_global_to_shared(
   void* __dest, const void* __src, _CUDA_VSTD::uint32_t __size, ::cuda::barrier<::cuda::thread_scope_block>& __bar)
 {
-  _LIBCUDACXX_ASSERT(__size % 16 == 0, "Size must be multiple of 16.");
-  _LIBCUDACXX_ASSERT(__isShared(__dest), "Destination must be shared memory address.");
-  _LIBCUDACXX_ASSERT(__isGlobal(__src), "Source must be global memory address.");
+  _CCCL_ASSERT(__size % 16 == 0, "Size must be multiple of 16.");
+  _CCCL_ASSERT(__isShared(__dest), "Destination must be shared memory address.");
+  _CCCL_ASSERT(__isGlobal(__src), "Source must be global memory address.");
 
   _CUDA_VPTX::cp_async_bulk(
     _CUDA_VPTX::space_cluster,
@@ -73,9 +73,9 @@ inline _CCCL_DEVICE void cp_async_bulk_global_to_shared(
 // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk
 inline _CCCL_DEVICE void cp_async_bulk_shared_to_global(void* __dest, const void* __src, _CUDA_VSTD::uint32_t __size)
 {
-  _LIBCUDACXX_ASSERT(__size % 16 == 0, "Size must be multiple of 16.");
-  _LIBCUDACXX_ASSERT(__isGlobal(__dest), "Destination must be global memory address.");
-  _LIBCUDACXX_ASSERT(__isShared(__src), "Source must be shared memory address.");
+  _CCCL_ASSERT(__size % 16 == 0, "Size must be multiple of 16.");
+  _CCCL_ASSERT(__isGlobal(__dest), "Destination must be global memory address.");
+  _CCCL_ASSERT(__isShared(__src), "Source must be shared memory address.");
 
   _CUDA_VPTX::cp_async_bulk(_CUDA_VPTX::space_global, _CUDA_VPTX::space_shared, __dest, __src, __size);
 }
diff --git a/libcudacxx/include/cuda/cmath b/libcudacxx/include/cuda/cmath
index 61529d8ce2..1321c24b96 100644
--- a/libcudacxx/include/cuda/cmath
+++ b/libcudacxx/include/cuda/cmath
@@ -33,9 +33,9 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 _Tp ceil_div(con
 {
   _CCCL_IF_CONSTEXPR (_CCCL_TRAIT(_CUDA_VSTD::is_signed, _Tp))
   {
-    _LIBCUDACXX_ASSERT(__a >= _Tp(0), "cuda::ceil_div: a must be non negative");
+    _CCCL_ASSERT(__a >= _Tp(0), "cuda::ceil_div: a must be non negative");
   }
-  _LIBCUDACXX_ASSERT(__b > _Tp(0), "cuda::ceil_div: b must be positive");
+  _CCCL_ASSERT(__b > _Tp(0), "cuda::ceil_div: b must be positive");
   const _Tp __res = static_cast<_Tp>(__a / __b);
   return static_cast<_Tp>(__res + (__res * __b != __a));
 }
diff --git a/libcudacxx/include/cuda/std/__algorithm/clamp.h b/libcudacxx/include/cuda/std/__algorithm/clamp.h
index 357fde9c22..6e3bb4656c 100644
--- a/libcudacxx/include/cuda/std/__algorithm/clamp.h
+++ b/libcudacxx/include/cuda/std/__algorithm/clamp.h
@@ -21,7 +21,7 @@
 #endif // no system header
 
 #include <cuda/std/__algorithm/comp.h>
-#include <cuda/std/detail/libcxx/include/__assert>
+#include <cuda/std/__cccl/assert.h>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
@@ -29,7 +29,7 @@ template <class _Tp, class _Compare>
 _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 const _Tp&
 clamp(const _Tp& __v, const _Tp& __lo, const _Tp& __hi, _Compare __comp)
 {
-  _LIBCUDACXX_ASSERT(!__comp(__hi, __lo), "Bad bounds passed to std::clamp");
+  _CCCL_ASSERT(!__comp(__hi, __lo), "Bad bounds passed to std::clamp");
   return __comp(__v, __lo) ? __lo : __comp(__hi, __v) ? __hi : __v;
 }
 
diff --git a/libcudacxx/include/cuda/std/__algorithm/comp_ref_type.h b/libcudacxx/include/cuda/std/__algorithm/comp_ref_type.h
index 075c741a4d..1b16ffa6fb 100644
--- a/libcudacxx/include/cuda/std/__algorithm/comp_ref_type.h
+++ b/libcudacxx/include/cuda/std/__algorithm/comp_ref_type.h
@@ -20,8 +20,8 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__utility/declval.h>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
@@ -60,7 +60,7 @@ struct __debug_less
   _CCCL_CONSTEXPR_CXX14 decltype((void) declval<_Compare&>()(declval<_LHS&>(), declval<_RHS&>()))
   __do_compare_assert(int, _LHS& __l, _RHS& __r)
   {
-    _LIBCUDACXX_ASSERT(!__comp_(__l, __r), "Comparator does not induce a strict weak ordering");
+    _CCCL_ASSERT(!__comp_(__l, __r), "Comparator does not induce a strict weak ordering");
     (void) __l;
     (void) __r;
   }
@@ -72,7 +72,7 @@ struct __debug_less
 
 // Pass the comparator by lvalue reference. Or in debug mode, using a
 // debugging wrapper that stores a reference.
-#ifdef _LIBCUDACXX_ENABLE_DEBUG_MODE
+#ifdef _CCCL_ENABLE_DEBUG_MODE
 template <class _Comp>
 using __comp_ref_type = __debug_less<_Comp>;
 #else
diff --git a/libcudacxx/include/cuda/std/__algorithm/pop_heap.h b/libcudacxx/include/cuda/std/__algorithm/pop_heap.h
index 6b15a27655..be998d8085 100644
--- a/libcudacxx/include/cuda/std/__algorithm/pop_heap.h
+++ b/libcudacxx/include/cuda/std/__algorithm/pop_heap.h
@@ -25,11 +25,11 @@
 #include <cuda/std/__algorithm/iterator_operations.h>
 #include <cuda/std/__algorithm/push_heap.h>
 #include <cuda/std/__algorithm/sift_down.h>
+#include <cuda/std/__cccl/assert.h>
 #include <cuda/std/__iterator/iterator_traits.h>
 #include <cuda/std/__type_traits/is_copy_assignable.h>
 #include <cuda/std/__type_traits/is_copy_constructible.h>
 #include <cuda/std/__utility/move.h>
-#include <cuda/std/detail/libcxx/include/__assert>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
@@ -41,7 +41,7 @@ _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 void __pop_heap(
   typename iterator_traits<_RandomAccessIterator>::difference_type __len)
 {
   // Calling `pop_heap` on an empty range is undefined behavior, but in practice it will be a no-op.
-  _LIBCUDACXX_ASSERT(__len > 0, "The heap given to pop_heap must be non-empty");
+  _CCCL_ASSERT(__len > 0, "The heap given to pop_heap must be non-empty");
 
   __comp_ref_type<_Compare> __comp_ref = __comp;
 
diff --git a/libcudacxx/include/cuda/std/__algorithm/sift_down.h b/libcudacxx/include/cuda/std/__algorithm/sift_down.h
index d988081dab..f31b80881e 100644
--- a/libcudacxx/include/cuda/std/__algorithm/sift_down.h
+++ b/libcudacxx/include/cuda/std/__algorithm/sift_down.h
@@ -21,9 +21,9 @@
 #endif // no system header
 
 #include <cuda/std/__algorithm/iterator_operations.h>
+#include <cuda/std/__cccl/assert.h>
 #include <cuda/std/__iterator/iterator_traits.h>
 #include <cuda/std/__utility/move.h>
-#include <cuda/std/detail/libcxx/include/__assert>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
@@ -99,7 +99,7 @@ _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 _RandomAccessIterator __floyd_si
   typename iterator_traits<_RandomAccessIterator>::difference_type __len)
 {
   using difference_type = typename iterator_traits<_RandomAccessIterator>::difference_type;
-  _LIBCUDACXX_ASSERT(__len >= 2, "shouldn't be called unless __len >= 2");
+  _CCCL_ASSERT(__len >= 2, "shouldn't be called unless __len >= 2");
 
   _RandomAccessIterator __hole    = __first;
   _RandomAccessIterator __child_i = __first;
diff --git a/libcudacxx/include/cuda/std/__cccl/assert.h b/libcudacxx/include/cuda/std/__cccl/assert.h
new file mode 100644
index 0000000000..d8dd2a12f3
--- /dev/null
+++ b/libcudacxx/include/cuda/std/__cccl/assert.h
@@ -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.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CCCL_ASSERT_H
+#define __CCCL_ASSERT_H
+
+#include <cuda/__cccl_config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/std/cassert>
+#include <cuda/std/cstdlib>
+
+#if __has_attribute(__format__)
+// The attribute uses 1-based indices for ordinary and static member functions.
+// The attribute uses 2-based indices for non-static member functions.
+#  define _CCCL_ATTRIBUTE_FORMAT(archetype, format_string_index, first_format_arg_index) \
+    __attribute__((__format__(archetype, format_string_index, first_format_arg_index)))
+#else
+#  define _CCCL_ATTRIBUTE_FORMAT(archetype, format_string_index, first_format_arg_index) /* nothing */
+#endif
+
+_CCCL_NORETURN _CCCL_ATTRIBUTE_FORMAT(__printf__, 1, 2)
+  _LIBCUDACXX_HIDE_FROM_ABI void __cccl_verbose_abort(const char*, ...)
+{
+  NV_IF_ELSE_TARGET(NV_IS_HOST, (::abort();), (__trap();))
+  _LIBCUDACXX_UNREACHABLE();
+}
+
+#if defined(_DEBUG) || defined(DEBUG)
+#  ifndef _CCCL_ENABLE_DEBUG_MODE
+#    define _CCCL_ENABLE_DEBUG_MODE
+#  endif // !_CCCL_ENABLE_DEBUG_MODE
+#endif // _DEBUG || DEBUG
+
+// Compatibility with older cccl versions
+#if defined(_LIBCUDACXX_ENABLE_DEBUG_MODE) || defined(_LIBCUDACXX_DEBUG) || defined(_LIBCUDACXX_ENABLE_ASSERTIONS)
+#  ifndef _CCCL_ENABLE_DEBUG_MODE
+#    define _CCCL_ENABLE_DEBUG_MODE
+#  endif // !_CCCL_ENABLE_DEBUG_MODE
+#endif // _LIBCUDACXX_ENABLE_DEBUG_MODE || _LIBCUDACXX_DEBUG || _LIBCUDACXX_ENABLE_ASSERTIONS
+
+// Automatically enable assertions when the debug mode is enabled manually.
+#ifdef _CCCL_ENABLE_DEBUG_MODE
+#  ifndef _CCCL_ENABLE_ASSERTIONS
+#    define _CCCL_ENABLE_ASSERTIONS
+#  endif // !_CCCL_ENABLE_ASSERTIONS
+#endif // _CCCL_ENABLE_DEBUG_MODE
+
+// Assert that is enabled unconditionally
+#define _CCCL_VERIFY(expression, message)            \
+  __builtin_expect(static_cast<bool>(expression), 1) \
+    ? (void) 0                                       \
+    : __cccl_verbose_abort("%s:%d: assertion %s failed: %s", __FILE__, __LINE__, #expression, message)
+
+#ifdef _CCCL_ENABLE_ASSERTIONS
+#  define _CCCL_ASSERT(expression, message) _CCCL_VERIFY(expression, message)
+#else // ^^^ _CCCL_ENABLE_ASSERTIONS ^^^ / vvv !_CCCL_ENABLE_ASSERTIONS vvv
+#  define _CCCL_ASSERT(expression, message) ((void) 0)
+#endif // !_CCCL_ENABLE_ASSERTIONS
+
+#endif // __CCCL_ASSERT_H
diff --git a/libcudacxx/include/cuda/std/__cuda/api_wrapper.h b/libcudacxx/include/cuda/std/__cuda/api_wrapper.h
index b5cbddc0d0..da7c0857d0 100644
--- a/libcudacxx/include/cuda/std/__cuda/api_wrapper.h
+++ b/libcudacxx/include/cuda/std/__cuda/api_wrapper.h
@@ -43,7 +43,7 @@
 #define _CCCL_ASSERT_CUDA_API(_NAME, _MSG, ...)        \
   {                                                    \
     const ::cudaError_t __status = _NAME(__VA_ARGS__); \
-    _LIBCUDACXX_ASSERT(__status == cudaSuccess, _MSG); \
+    _CCCL_ASSERT(__status == cudaSuccess, _MSG);       \
     (void) __status;                                   \
   }
 
diff --git a/libcudacxx/include/cuda/std/__cuda/barrier.h b/libcudacxx/include/cuda/std/__cuda/barrier.h
index 9b6dd972b8..f2c16001c4 100644
--- a/libcudacxx/include/cuda/std/__cuda/barrier.h
+++ b/libcudacxx/include/cuda/std/__cuda/barrier.h
@@ -26,9 +26,9 @@
 #endif // no system header
 
 #include <cuda/std/__atomic/api/owned.h>
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__type_traits/void_t.h> // _CUDA_VSTD::void_t
 #include <cuda/std/cstdlib> // _LIBCUDACXX_UNREACHABLE
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 
 #if defined(_CCCL_CUDA_COMPILER)
 #  include <cuda/ptx> // cuda::ptx::*
@@ -98,13 +98,13 @@ class barrier : public _CUDA_VSTD::__barrier_base<_CompletionF, _Sco>
 
   _LIBCUDACXX_HIDE_FROM_ABI friend void init(barrier* __b, _CUDA_VSTD::ptrdiff_t __expected)
   {
-    _LIBCUDACXX_ASSERT(__expected >= 0, "Cannot initialize barrier with negative arrival count");
+    _CCCL_ASSERT(__expected >= 0, "Cannot initialize barrier with negative arrival count");
     new (__b) barrier(__expected);
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI friend void init(barrier* __b, _CUDA_VSTD::ptrdiff_t __expected, _CompletionF __completion)
   {
-    _LIBCUDACXX_ASSERT(__expected >= 0, "Cannot initialize barrier with negative arrival count");
+    _CCCL_ASSERT(__expected >= 0, "Cannot initialize barrier with negative arrival count");
     new (__b) barrier(__expected, __completion);
   }
 };
@@ -197,7 +197,7 @@ class barrier<thread_scope_block, _CUDA_VSTD::__empty_completion> : public __blo
 
   _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI arrival_token arrive(_CUDA_VSTD::ptrdiff_t __update = 1)
   {
-    _LIBCUDACXX_ASSERT(__update >= 0, "Arrival count update must be non-negative.");
+    _CCCL_ASSERT(__update >= 0, "Arrival count update must be non-negative.");
     arrival_token __token = {};
     NV_DISPATCH_TARGET(
       NV_PROVIDES_SM_90,
@@ -550,12 +550,12 @@ _CCCL_NODISCARD _CCCL_DEVICE inline barrier<thread_scope_block>::arrival_token b
   _CUDA_VSTD::ptrdiff_t __arrive_count_update,
   _CUDA_VSTD::ptrdiff_t __transaction_count_update)
 {
-  _LIBCUDACXX_ASSERT(__isShared(barrier_native_handle(__b)), "Barrier must be located in local shared memory.");
-  _LIBCUDACXX_ASSERT(1 <= __arrive_count_update, "Arrival count update must be at least one.");
-  _LIBCUDACXX_ASSERT(__arrive_count_update <= (1 << 20) - 1, "Arrival count update cannot exceed 2^20 - 1.");
-  _LIBCUDACXX_ASSERT(__transaction_count_update >= 0, "Transaction count update must be non-negative.");
+  _CCCL_ASSERT(__isShared(barrier_native_handle(__b)), "Barrier must be located in local shared memory.");
+  _CCCL_ASSERT(1 <= __arrive_count_update, "Arrival count update must be at least one.");
+  _CCCL_ASSERT(__arrive_count_update <= (1 << 20) - 1, "Arrival count update cannot exceed 2^20 - 1.");
+  _CCCL_ASSERT(__transaction_count_update >= 0, "Transaction count update must be non-negative.");
   // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#contents-of-the-mbarrier-object
-  _LIBCUDACXX_ASSERT(__transaction_count_update <= (1 << 20) - 1, "Transaction count update cannot exceed 2^20 - 1.");
+  _CCCL_ASSERT(__transaction_count_update <= (1 << 20) - 1, "Transaction count update cannot exceed 2^20 - 1.");
 
   barrier<thread_scope_block>::arrival_token __token = {};
   // On architectures pre-sm90, arrive_tx is not supported.
@@ -599,10 +599,10 @@ extern "C" _CCCL_DEVICE void __cuda_ptx_barrier_expect_tx_is_not_supported_befor
 _CCCL_DEVICE inline void
 barrier_expect_tx(barrier<thread_scope_block>& __b, _CUDA_VSTD::ptrdiff_t __transaction_count_update)
 {
-  _LIBCUDACXX_ASSERT(__isShared(barrier_native_handle(__b)), "Barrier must be located in local shared memory.");
-  _LIBCUDACXX_ASSERT(__transaction_count_update >= 0, "Transaction count update must be non-negative.");
+  _CCCL_ASSERT(__isShared(barrier_native_handle(__b)), "Barrier must be located in local shared memory.");
+  _CCCL_ASSERT(__transaction_count_update >= 0, "Transaction count update must be non-negative.");
   // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#contents-of-the-mbarrier-object
-  _LIBCUDACXX_ASSERT(__transaction_count_update <= (1 << 20) - 1, "Transaction count update cannot exceed 2^20 - 1.");
+  _CCCL_ASSERT(__transaction_count_update <= (1 << 20) - 1, "Transaction count update cannot exceed 2^20 - 1.");
 
   // We do not check for the statespace of the barrier here. This is
   // on purpose. This allows debugging tools like memcheck/racecheck
@@ -640,9 +640,9 @@ _CCCL_DEVICE inline async_contract_fulfillment memcpy_async_tx(
 #    endif
   static_assert(16 <= _Alignment, "mempcy_async_tx expects arguments to be at least 16 byte aligned.");
 
-  _LIBCUDACXX_ASSERT(__isShared(barrier_native_handle(__b)), "Barrier must be located in local shared memory.");
-  _LIBCUDACXX_ASSERT(__isShared(__dest), "dest must point to shared memory.");
-  _LIBCUDACXX_ASSERT(__isGlobal(__src), "src must point to global memory.");
+  _CCCL_ASSERT(__isShared(barrier_native_handle(__b)), "Barrier must be located in local shared memory.");
+  _CCCL_ASSERT(__isShared(__dest), "dest must point to shared memory.");
+  _CCCL_ASSERT(__isGlobal(__src), "src must point to global memory.");
 
   NV_IF_ELSE_TARGET(
     NV_PROVIDES_SM_90,
@@ -1104,8 +1104,8 @@ _CCCL_NODISCARD _CCCL_DEVICE inline __completion_mechanism __dispatch_memcpy_asy
     NV_PROVIDES_SM_90,
     (const bool __can_use_complete_tx = __allowed_completions & uint32_t(__completion_mechanism::__mbarrier_complete_tx);
      _LIBCUDACXX_UNUSED_VAR(__can_use_complete_tx);
-     _LIBCUDACXX_ASSERT(__can_use_complete_tx == (nullptr != __bar_handle),
-                        "Pass non-null bar_handle if and only if can_use_complete_tx.");
+     _CCCL_ASSERT(__can_use_complete_tx == (nullptr != __bar_handle),
+                  "Pass non-null bar_handle if and only if can_use_complete_tx.");
      _CCCL_IF_CONSTEXPR (_Align >= 16) {
        if (__can_use_complete_tx && __isShared(__bar_handle))
        {
@@ -1178,8 +1178,8 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI __completion_mechanism __dispatch_memc
   _CUDA_VSTD::size_t __size,
   _CUDA_VSTD::uint32_t __allowed_completions)
 {
-  _LIBCUDACXX_ASSERT(!(__allowed_completions & uint32_t(__completion_mechanism::__mbarrier_complete_tx)),
-                     "Cannot allow mbarrier_complete_tx completion mechanism when not passing a barrier. ");
+  _CCCL_ASSERT(!(__allowed_completions & uint32_t(__completion_mechanism::__mbarrier_complete_tx)),
+               "Cannot allow mbarrier_complete_tx completion mechanism when not passing a barrier. ");
   return __dispatch_memcpy_async<_Align>(__group, __dest_char, __src_char, __size, __allowed_completions, nullptr);
 }
 
diff --git a/libcudacxx/include/cuda/std/__expected/expected.h b/libcudacxx/include/cuda/std/__expected/expected.h
index 42f0e5b223..ed52e8aed5 100644
--- a/libcudacxx/include/cuda/std/__expected/expected.h
+++ b/libcudacxx/include/cuda/std/__expected/expected.h
@@ -19,6 +19,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h>
 #include <cuda/std/__concepts/invocable.h>
 #include <cuda/std/__expected/bad_expected_access.h>
 #include <cuda/std/__expected/expected_base.h>
@@ -61,7 +62,6 @@
 #include <cuda/std/__utility/move.h>
 #include <cuda/std/__utility/swap.h>
 #include <cuda/std/cstdlib>
-#include <cuda/std/detail/libcxx/include/__assert>
 #include <cuda/std/initializer_list>
 
 #if _CCCL_STD_VER > 2011
@@ -488,37 +488,37 @@ class expected : private __expected_move_assign<_Tp, _Err>
   // [expected.object.obs], observers
   _LIBCUDACXX_HIDE_FROM_ABI constexpr const _Tp* operator->() const noexcept
   {
-    _LIBCUDACXX_ASSERT(this->__has_val_, "expected::operator-> requires the expected to contain a value");
+    _CCCL_ASSERT(this->__has_val_, "expected::operator-> requires the expected to contain a value");
     return _CUDA_VSTD::addressof(this->__union_.__val_);
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp* operator->() noexcept
   {
-    _LIBCUDACXX_ASSERT(this->__has_val_, "expected::operator-> requires the expected to contain a value");
+    _CCCL_ASSERT(this->__has_val_, "expected::operator-> requires the expected to contain a value");
     return _CUDA_VSTD::addressof(this->__union_.__val_);
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr const _Tp& operator*() const& noexcept
   {
-    _LIBCUDACXX_ASSERT(this->__has_val_, "expected::operator* requires the expected to contain a value");
+    _CCCL_ASSERT(this->__has_val_, "expected::operator* requires the expected to contain a value");
     return this->__union_.__val_;
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp& operator*() & noexcept
   {
-    _LIBCUDACXX_ASSERT(this->__has_val_, "expected::operator* requires the expected to contain a value");
+    _CCCL_ASSERT(this->__has_val_, "expected::operator* requires the expected to contain a value");
     return this->__union_.__val_;
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr const _Tp&& operator*() const&& noexcept
   {
-    _LIBCUDACXX_ASSERT(this->__has_val_, "expected::operator* requires the expected to contain a value");
+    _CCCL_ASSERT(this->__has_val_, "expected::operator* requires the expected to contain a value");
     return _CUDA_VSTD::move(this->__union_.__val_);
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp&& operator*() && noexcept
   {
-    _LIBCUDACXX_ASSERT(this->__has_val_, "expected::operator* requires the expected to contain a value");
+    _CCCL_ASSERT(this->__has_val_, "expected::operator* requires the expected to contain a value");
     return _CUDA_VSTD::move(this->__union_.__val_);
   }
 
@@ -580,25 +580,25 @@ class expected : private __expected_move_assign<_Tp, _Err>
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr const _Err& error() const& noexcept
   {
-    _LIBCUDACXX_ASSERT(!this->__has_val_, "expected::error requires the expected to contain an error");
+    _CCCL_ASSERT(!this->__has_val_, "expected::error requires the expected to contain an error");
     return this->__union_.__unex_;
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr _Err& error() & noexcept
   {
-    _LIBCUDACXX_ASSERT(!this->__has_val_, "expected::error requires the expected to contain an error");
+    _CCCL_ASSERT(!this->__has_val_, "expected::error requires the expected to contain an error");
     return this->__union_.__unex_;
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr const _Err&& error() const&& noexcept
   {
-    _LIBCUDACXX_ASSERT(!this->__has_val_, "expected::error requires the expected to contain an error");
+    _CCCL_ASSERT(!this->__has_val_, "expected::error requires the expected to contain an error");
     return _CUDA_VSTD::move(this->__union_.__unex_);
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr _Err&& error() && noexcept
   {
-    _LIBCUDACXX_ASSERT(!this->__has_val_, "expected::error requires the expected to contain an error");
+    _CCCL_ASSERT(!this->__has_val_, "expected::error requires the expected to contain an error");
     return _CUDA_VSTD::move(this->__union_.__unex_);
   }
 
@@ -1436,7 +1436,7 @@ class expected<void, _Err> : private __expected_move_assign<void, _Err>
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr void operator*() const noexcept
   {
-    _LIBCUDACXX_ASSERT(this->__has_val_, "expected::operator* requires the expected to contain a value");
+    _CCCL_ASSERT(this->__has_val_, "expected::operator* requires the expected to contain a value");
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr void value() const&
@@ -1461,25 +1461,25 @@ class expected<void, _Err> : private __expected_move_assign<void, _Err>
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr const _Err& error() const& noexcept
   {
-    _LIBCUDACXX_ASSERT(!this->__has_val_, "expected::error requires the expected to contain an error");
+    _CCCL_ASSERT(!this->__has_val_, "expected::error requires the expected to contain an error");
     return this->__union_.__unex_;
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr _Err& error() & noexcept
   {
-    _LIBCUDACXX_ASSERT(!this->__has_val_, "expected::error requires the expected to contain an error");
+    _CCCL_ASSERT(!this->__has_val_, "expected::error requires the expected to contain an error");
     return this->__union_.__unex_;
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr const _Err&& error() const&& noexcept
   {
-    _LIBCUDACXX_ASSERT(!this->__has_val_, "expected::error requires the expected to contain an error");
+    _CCCL_ASSERT(!this->__has_val_, "expected::error requires the expected to contain an error");
     return _CUDA_VSTD::move(this->__union_.__unex_);
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr _Err&& error() && noexcept
   {
-    _LIBCUDACXX_ASSERT(!this->__has_val_, "expected::error requires the expected to contain an error");
+    _CCCL_ASSERT(!this->__has_val_, "expected::error requires the expected to contain an error");
     return _CUDA_VSTD::move(this->__union_.__unex_);
   }
 
diff --git a/libcudacxx/include/cuda/std/__expected/expected_base.h b/libcudacxx/include/cuda/std/__expected/expected_base.h
index be5fd87afb..f321238f3c 100644
--- a/libcudacxx/include/cuda/std/__expected/expected_base.h
+++ b/libcudacxx/include/cuda/std/__expected/expected_base.h
@@ -19,6 +19,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h>
 #include <cuda/std/__concepts/__concept_macros.h>
 #include <cuda/std/__concepts/invocable.h>
 #include <cuda/std/__expected/unexpect.h>
@@ -49,7 +50,6 @@
 #include <cuda/std/__utility/forward.h>
 #include <cuda/std/__utility/in_place.h>
 #include <cuda/std/__utility/move.h>
-#include <cuda/std/detail/libcxx/include/__assert>
 
 #if _CCCL_STD_VER > 2011
 
diff --git a/libcudacxx/include/cuda/std/__functional/function.h b/libcudacxx/include/cuda/std/__functional/function.h
index 2d7c6931a1..e86e20ff97 100644
--- a/libcudacxx/include/cuda/std/__functional/function.h
+++ b/libcudacxx/include/cuda/std/__functional/function.h
@@ -21,6 +21,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h>
 #include <cuda/std/__exception/terminate.h>
 #include <cuda/std/__functional/binary_function.h>
 #include <cuda/std/__functional/invoke.h>
@@ -47,7 +48,6 @@
 #include <cuda/std/__utility/move.h>
 #include <cuda/std/__utility/piecewise_construct.h>
 #include <cuda/std/__utility/swap.h>
-#include <cuda/std/detail/libcxx/include/__assert>
 #include <cuda/std/tuple>
 
 #ifndef __cuda_std__
@@ -938,10 +938,10 @@ class __func<_Rp1 (^)(_ArgTypes1...), _Alloc, _Rp(_ArgTypes...)> : public __base
 
   virtual __base<_Rp(_ArgTypes...)>* __clone() const
   {
-    _LIBCUDACXX_ASSERT(false,
-                       "Block pointers are just pointers, so they should always fit into "
-                       "std::function's small buffer optimization. This function should "
-                       "never be invoked.");
+    _CCCL_ASSERT(false,
+                 "Block pointers are just pointers, so they should always fit into "
+                 "std::function's small buffer optimization. This function should "
+                 "never be invoked.");
     return nullptr;
   }
 
@@ -963,10 +963,10 @@ class __func<_Rp1 (^)(_ArgTypes1...), _Alloc, _Rp(_ArgTypes...)> : public __base
 
   virtual void destroy_deallocate() noexcept
   {
-    _LIBCUDACXX_ASSERT(false,
-                       "Block pointers are just pointers, so they should always fit into "
-                       "std::function's small buffer optimization. This function should "
-                       "never be invoked.");
+    _CCCL_ASSERT(false,
+                 "Block pointers are just pointers, so they should always fit into "
+                 "std::function's small buffer optimization. This function should "
+                 "never be invoked.");
   }
 
   virtual _Rp operator()(_ArgTypes&&... __arg)
diff --git a/libcudacxx/include/cuda/std/__iterator/advance.h b/libcudacxx/include/cuda/std/__iterator/advance.h
index 0a84feb53b..97e3ae200f 100644
--- a/libcudacxx/include/cuda/std/__iterator/advance.h
+++ b/libcudacxx/include/cuda/std/__iterator/advance.h
@@ -21,6 +21,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h>
 #include <cuda/std/__concepts/assignable.h>
 #include <cuda/std/__concepts/same_as.h>
 #include <cuda/std/__iterator/concepts.h>
@@ -29,7 +30,6 @@
 #include <cuda/std/__utility/convert_to_integral.h>
 #include <cuda/std/__utility/move.h>
 #include <cuda/std/cstdlib>
-#include <cuda/std/detail/libcxx/include/__assert>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
@@ -78,8 +78,8 @@ _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 void advance(_InputIter& __i, _D
 {
   typedef typename iterator_traits<_InputIter>::difference_type _Difference;
   _Difference __n = static_cast<_Difference>(_CUDA_VSTD::__convert_to_integral(__orig_n));
-  _LIBCUDACXX_ASSERT(__n >= 0 || __is_cpp17_bidirectional_iterator<_InputIter>::value,
-                     "Attempt to advance(it, n) with negative n on a non-bidirectional iterator");
+  _CCCL_ASSERT(__n >= 0 || __is_cpp17_bidirectional_iterator<_InputIter>::value,
+               "Attempt to advance(it, n) with negative n on a non-bidirectional iterator");
   _CUDA_VSTD::__advance(__i, __n, typename iterator_traits<_InputIter>::iterator_category());
 }
 
@@ -129,8 +129,7 @@ struct __fn
   _LIBCUDACXX_REQUIRES(input_or_output_iterator<_Ip>)
   _LIBCUDACXX_HIDE_FROM_ABI constexpr void operator()(_Ip& __i, iter_difference_t<_Ip> __n) const
   {
-    _LIBCUDACXX_ASSERT(__n >= 0 || bidirectional_iterator<_Ip>,
-                       "If `n < 0`, then `bidirectional_iterator<I>` must be true.");
+    _CCCL_ASSERT(__n >= 0 || bidirectional_iterator<_Ip>, "If `n < 0`, then `bidirectional_iterator<I>` must be true.");
 
     // If `I` models `random_access_iterator`, equivalent to `i += n`.
     if constexpr (random_access_iterator<_Ip>)
@@ -190,8 +189,8 @@ struct __fn
   _LIBCUDACXX_HIDE_FROM_ABI constexpr iter_difference_t<_Ip>
   operator()(_Ip& __i, iter_difference_t<_Ip> __n, _Sp __bound_sentinel) const
   {
-    _LIBCUDACXX_ASSERT((__n >= 0) || (bidirectional_iterator<_Ip> && same_as<_Ip, _Sp>),
-                       "If `n < 0`, then `bidirectional_iterator<I> && same_as<I, S>` must be true.");
+    _CCCL_ASSERT((__n >= 0) || (bidirectional_iterator<_Ip> && same_as<_Ip, _Sp>),
+                 "If `n < 0`, then `bidirectional_iterator<I> && same_as<I, S>` must be true.");
     // If `S` and `I` model `sized_sentinel_for<S, I>`:
     if constexpr (sized_sentinel_for<_Sp, _Ip>)
     {
diff --git a/libcudacxx/include/cuda/std/__iterator/bounded_iter.h b/libcudacxx/include/cuda/std/__iterator/bounded_iter.h
index 2c2fc3f67b..c21073b7d8 100644
--- a/libcudacxx/include/cuda/std/__iterator/bounded_iter.h
+++ b/libcudacxx/include/cuda/std/__iterator/bounded_iter.h
@@ -21,12 +21,12 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h>
 #include <cuda/std/__iterator/iterator_traits.h>
 #include <cuda/std/__memory/pointer_traits.h>
 #include <cuda/std/__type_traits/enable_if.h>
 #include <cuda/std/__type_traits/is_convertible.h>
 #include <cuda/std/__utility/move.h>
-#include <cuda/std/detail/libcxx/include/__assert>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
@@ -88,7 +88,7 @@ struct __bounded_iter
       , __begin_(__begin)
       , __end_(__end)
   {
-    _LIBCUDACXX_ASSERT(__begin <= __end, "__bounded_iter(current, begin, end): [begin, end) is not a valid range");
+    _CCCL_ASSERT(__begin <= __end, "__bounded_iter(current, begin, end): [begin, end) is not a valid range");
   }
 
   template <class _It>
@@ -100,22 +100,21 @@ struct __bounded_iter
   // These operations check that the iterator is dereferenceable, that is within [begin, end).
   _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 reference operator*() const noexcept
   {
-    _LIBCUDACXX_ASSERT(__in_bounds(__current_),
-                       "__bounded_iter::operator*: Attempt to dereference an out-of-range iterator");
+    _CCCL_ASSERT(__in_bounds(__current_), "__bounded_iter::operator*: Attempt to dereference an out-of-range iterator");
     return *__current_;
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 pointer operator->() const noexcept
   {
-    _LIBCUDACXX_ASSERT(__in_bounds(__current_),
-                       "__bounded_iter::operator->: Attempt to dereference an out-of-range iterator");
+    _CCCL_ASSERT(__in_bounds(__current_),
+                 "__bounded_iter::operator->: Attempt to dereference an out-of-range iterator");
     return _CUDA_VSTD::__to_address(__current_);
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 reference operator[](difference_type __n) const noexcept
   {
-    _LIBCUDACXX_ASSERT(__in_bounds(__current_ + __n),
-                       "__bounded_iter::operator[]: Attempt to index an iterator out-of-range");
+    _CCCL_ASSERT(__in_bounds(__current_ + __n),
+                 "__bounded_iter::operator[]: Attempt to index an iterator out-of-range");
     return __current_[__n];
   }
 
diff --git a/libcudacxx/include/cuda/std/__iterator/next.h b/libcudacxx/include/cuda/std/__iterator/next.h
index d31050637b..fe32b1467b 100644
--- a/libcudacxx/include/cuda/std/__iterator/next.h
+++ b/libcudacxx/include/cuda/std/__iterator/next.h
@@ -21,12 +21,12 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h>
 #include <cuda/std/__iterator/advance.h>
 #include <cuda/std/__iterator/concepts.h>
 #include <cuda/std/__iterator/incrementable_traits.h>
 #include <cuda/std/__iterator/iterator_traits.h>
 #include <cuda/std/__type_traits/enable_if.h>
-#include <cuda/std/detail/libcxx/include/__assert>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
@@ -34,8 +34,8 @@ template <class _InputIter>
 _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 __enable_if_t<__is_cpp17_input_iterator<_InputIter>::value, _InputIter>
 next(_InputIter __x, typename iterator_traits<_InputIter>::difference_type __n = 1)
 {
-  _LIBCUDACXX_ASSERT(__n >= 0 || __is_cpp17_bidirectional_iterator<_InputIter>::value,
-                     "Attempt to next(it, n) with negative n on a non-bidirectional iterator");
+  _CCCL_ASSERT(__n >= 0 || __is_cpp17_bidirectional_iterator<_InputIter>::value,
+               "Attempt to next(it, n) with negative n on a non-bidirectional iterator");
 
   _CUDA_VSTD::advance(__x, __n);
   return __x;
diff --git a/libcudacxx/include/cuda/std/__iterator/prev.h b/libcudacxx/include/cuda/std/__iterator/prev.h
index a69f4f456c..a16804bd06 100644
--- a/libcudacxx/include/cuda/std/__iterator/prev.h
+++ b/libcudacxx/include/cuda/std/__iterator/prev.h
@@ -21,12 +21,12 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h>
 #include <cuda/std/__iterator/advance.h>
 #include <cuda/std/__iterator/concepts.h>
 #include <cuda/std/__iterator/incrementable_traits.h>
 #include <cuda/std/__iterator/iterator_traits.h>
 #include <cuda/std/__type_traits/enable_if.h>
-#include <cuda/std/detail/libcxx/include/__assert>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
@@ -34,8 +34,8 @@ template <class _InputIter>
 _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 __enable_if_t<__is_cpp17_input_iterator<_InputIter>::value, _InputIter>
 prev(_InputIter __x, typename iterator_traits<_InputIter>::difference_type __n = 1)
 {
-  _LIBCUDACXX_ASSERT(__n <= 0 || __is_cpp17_bidirectional_iterator<_InputIter>::value,
-                     "Attempt to prev(it, +n) on a non-bidi iterator");
+  _CCCL_ASSERT(__n <= 0 || __is_cpp17_bidirectional_iterator<_InputIter>::value,
+               "Attempt to prev(it, +n) on a non-bidi iterator");
   _CUDA_VSTD::advance(__x, -__n);
   return __x;
 }
diff --git a/libcudacxx/include/cuda/std/__mdspan/layout_left.h b/libcudacxx/include/cuda/std/__mdspan/layout_left.h
index be598ebbf7..27a6f1cc64 100644
--- a/libcudacxx/include/cuda/std/__mdspan/layout_left.h
+++ b/libcudacxx/include/cuda/std/__mdspan/layout_left.h
@@ -54,6 +54,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h>
 #include <cuda/std/__mdspan/extents.h>
 #include <cuda/std/__mdspan/layout_stride.h>
 #include <cuda/std/__mdspan/macros.h>
@@ -62,7 +63,6 @@
 #include <cuda/std/__type_traits/is_nothrow_constructible.h>
 #include <cuda/std/__utility/integer_sequence.h>
 #include <cuda/std/cstddef>
-#include <cuda/std/detail/libcxx/include/__assert>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
diff --git a/libcudacxx/include/cuda/std/__mdspan/layout_right.h b/libcudacxx/include/cuda/std/__mdspan/layout_right.h
index 5c2c51d7c4..475cb99e8c 100644
--- a/libcudacxx/include/cuda/std/__mdspan/layout_right.h
+++ b/libcudacxx/include/cuda/std/__mdspan/layout_right.h
@@ -54,6 +54,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h>
 #include <cuda/std/__mdspan/extents.h>
 #include <cuda/std/__mdspan/layout_stride.h>
 #include <cuda/std/__mdspan/macros.h>
@@ -62,7 +63,6 @@
 #include <cuda/std/__type_traits/is_nothrow_constructible.h>
 #include <cuda/std/__utility/integer_sequence.h>
 #include <cuda/std/cstddef>
-#include <cuda/std/detail/libcxx/include/__assert>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
diff --git a/libcudacxx/include/cuda/std/__mdspan/macros.h b/libcudacxx/include/cuda/std/__mdspan/macros.h
index 831bfa45a3..f98c837565 100644
--- a/libcudacxx/include/cuda/std/__mdspan/macros.h
+++ b/libcudacxx/include/cuda/std/__mdspan/macros.h
@@ -660,7 +660,7 @@ _LIBCUDACXX_END_NAMESPACE_STD
     if (!(_COND))                                          \
     __throw_runtime_error(_MESSAGE)
 #else // ^^^ !_CCCL_NO_EXCEPTIONS ^^^ / vvv _CCCL_NO_EXCEPTIONS vvv
-#  define _LIBCUDACXX_THROW_RUNTIME_ERROR(_COND, _MESSAGE) _LIBCUDACXX_ASSERT(_COND, _MESSAGE)
+#  define _LIBCUDACXX_THROW_RUNTIME_ERROR(_COND, _MESSAGE) _CCCL_ASSERT(_COND, _MESSAGE)
 #endif // _CCCL_NO_EXCEPTIONS
 
 #endif // _LIBCUDACXX___MDSPAN_MACROS_HPP
diff --git a/libcudacxx/include/cuda/std/__memory/construct_at.h b/libcudacxx/include/cuda/std/__memory/construct_at.h
index 3d01e03919..2f9be970d5 100644
--- a/libcudacxx/include/cuda/std/__memory/construct_at.h
+++ b/libcudacxx/include/cuda/std/__memory/construct_at.h
@@ -21,6 +21,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h>
 #include <cuda/std/__concepts/__concept_macros.h>
 #include <cuda/std/__iterator/access.h>
 #include <cuda/std/__memory/addressof.h>
@@ -37,7 +38,6 @@
 #include <cuda/std/__utility/declval.h>
 #include <cuda/std/__utility/forward.h>
 #include <cuda/std/__utility/move.h>
-#include <cuda/std/detail/libcxx/include/__assert>
 
 #ifdef _CCCL_CUDA_COMPILER_CLANG
 #  include <new>
@@ -113,7 +113,7 @@ _LIBCUDACXX_HIDE_FROM_ABI
 _CCCL_CONSTEXPR_CXX20 __enable_if_t<!__detail::__can_optimize_construct_at<_Tp, _Args...>::value, _Tp*>
 construct_at(_Tp* __location, _Args&&... __args)
 {
-  _LIBCUDACXX_ASSERT(__location != nullptr, "null pointer given to construct_at");
+  _CCCL_ASSERT(__location != nullptr, "null pointer given to construct_at");
   // Need to go through `std::construct_at` as that is the explicitly blessed function
   if (__libcpp_is_constant_evaluated())
   {
@@ -130,7 +130,7 @@ _LIBCUDACXX_HIDE_FROM_ABI
 _CCCL_CONSTEXPR_CXX20 __enable_if_t<__detail::__can_optimize_construct_at<_Tp, _Args...>::value, _Tp*>
 construct_at(_Tp* __location, _Args&&... __args)
 {
-  _LIBCUDACXX_ASSERT(__location != nullptr, "null pointer given to construct_at");
+  _CCCL_ASSERT(__location != nullptr, "null pointer given to construct_at");
   // Need to go through `std::construct_at` as that is the explicitly blessed function
   if (__libcpp_is_constant_evaluated())
   {
@@ -148,7 +148,7 @@ _LIBCUDACXX_HIDE_FROM_ABI
 _CCCL_CONSTEXPR_CXX20 __enable_if_t<!__detail::__can_optimize_construct_at<_Tp, _Args...>::value, _Tp*>
 __construct_at(_Tp* __location, _Args&&... __args)
 {
-  _LIBCUDACXX_ASSERT(__location != nullptr, "null pointer given to construct_at");
+  _CCCL_ASSERT(__location != nullptr, "null pointer given to construct_at");
 #if _CCCL_STD_VER >= 2020
   // Need to go through `std::construct_at` as that is the explicitly blessed function
   if (__libcpp_is_constant_evaluated())
@@ -165,7 +165,7 @@ _LIBCUDACXX_HIDE_FROM_ABI
 _CCCL_CONSTEXPR_CXX20 __enable_if_t<__detail::__can_optimize_construct_at<_Tp, _Args...>::value, _Tp*>
 __construct_at(_Tp* __location, _Args&&... __args)
 {
-  _LIBCUDACXX_ASSERT(__location != nullptr, "null pointer given to construct_at");
+  _CCCL_ASSERT(__location != nullptr, "null pointer given to construct_at");
 #if _CCCL_STD_VER >= 2020
   // Need to go through `std::construct_at` as that is the explicitly blessed function
   if (__libcpp_is_constant_evaluated())
@@ -190,7 +190,7 @@ template <class _Tp,
           __enable_if_t<!is_trivially_destructible<_Tp>::value, int> = 0>
 _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 void __destroy_at(_Tp* __loc)
 {
-  _LIBCUDACXX_ASSERT(__loc != nullptr, "null pointer given to destroy_at");
+  _CCCL_ASSERT(__loc != nullptr, "null pointer given to destroy_at");
   __loc->~_Tp();
 }
 
@@ -200,7 +200,7 @@ template <class _Tp,
           __enable_if_t<is_trivially_destructible<_Tp>::value, int> = 0>
 _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 void __destroy_at(_Tp* __loc)
 {
-  _LIBCUDACXX_ASSERT(__loc != nullptr, "null pointer given to destroy_at");
+  _CCCL_ASSERT(__loc != nullptr, "null pointer given to destroy_at");
   (void) __loc;
 }
 
@@ -208,7 +208,7 @@ _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 void __destroy_at(_Tp* __loc)
 template <class _Tp, __enable_if_t<is_array<_Tp>::value, int> = 0>
 _LIBCUDACXX_HIDE_FROM_ABI constexpr void __destroy_at(_Tp* __loc)
 {
-  _LIBCUDACXX_ASSERT(__loc != nullptr, "null pointer given to destroy_at");
+  _CCCL_ASSERT(__loc != nullptr, "null pointer given to destroy_at");
   _CUDA_VSTD::__destroy(_CUDA_VSTD::begin(*__loc), _CUDA_VSTD::end(*__loc));
 }
 #endif // _CCCL_STD_VER >= 2020
@@ -241,7 +241,7 @@ __reverse_destroy(_BidirectionalIterator __first, _BidirectionalIterator __last)
 template <class _Tp, enable_if_t<!is_array_v<_Tp>, int> = 0>
 _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX20 void destroy_at(_Tp* __loc) noexcept
 {
-  _LIBCUDACXX_ASSERT(__loc != nullptr, "null pointer given to destroy_at");
+  _CCCL_ASSERT(__loc != nullptr, "null pointer given to destroy_at");
   __loc->~_Tp();
 }
 
diff --git a/libcudacxx/include/cuda/std/__memory_ b/libcudacxx/include/cuda/std/__memory_
index 0cf9e7f213..6e456a1441 100644
--- a/libcudacxx/include/cuda/std/__memory_
+++ b/libcudacxx/include/cuda/std/__memory_
@@ -21,6 +21,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__memory/addressof.h>
 #include <cuda/std/__memory/align.h>
 #include <cuda/std/__memory/allocate_at_least.h>
@@ -33,7 +34,6 @@
 #include <cuda/std/__memory/uninitialized_algorithms.h>
 #include <cuda/std/__memory/unique_ptr.h>
 #include <cuda/std/__memory/uses_allocator.h>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 
 // standard-mandated includes
 #include <cuda/std/version>
diff --git a/libcudacxx/include/cuda/std/__new_ b/libcudacxx/include/cuda/std/__new_
index 9d7c1fa853..d932b94c0e 100644
--- a/libcudacxx/include/cuda/std/__new_
+++ b/libcudacxx/include/cuda/std/__new_
@@ -21,10 +21,10 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__new/allocate.h>
 #include <cuda/std/__new/bad_alloc.h>
 #include <cuda/std/__new/launder.h>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/version>
 
 #endif // _CUDA_STD_NEW
diff --git a/libcudacxx/include/cuda/std/__numeric/gcd_lcm.h b/libcudacxx/include/cuda/std/__numeric/gcd_lcm.h
index 9212d1d069..e123160b15 100644
--- a/libcudacxx/include/cuda/std/__numeric/gcd_lcm.h
+++ b/libcudacxx/include/cuda/std/__numeric/gcd_lcm.h
@@ -100,7 +100,7 @@ _CCCL_CONSTEXPR_CXX14 _LIBCUDACXX_HIDE_FROM_ABI __common_type_t<_Tp, _Up> lcm(_T
   using _Rp  = __common_type_t<_Tp, _Up>;
   _Rp __val1 = __ct_abs<_Rp, _Tp>()(__m) / _CUDA_VSTD::gcd(__m, __n);
   _Rp __val2 = __ct_abs<_Rp, _Up>()(__n);
-  _LIBCUDACXX_ASSERT((numeric_limits<_Rp>::max() / __val1 > __val2), "Overflow in lcm");
+  _CCCL_ASSERT((numeric_limits<_Rp>::max() / __val1 > __val2), "Overflow in lcm");
   return __val1 * __val2;
 }
 
diff --git a/libcudacxx/include/cuda/std/__ranges/subrange.h b/libcudacxx/include/cuda/std/__ranges/subrange.h
index 50abc34f20..0a166688f5 100644
--- a/libcudacxx/include/cuda/std/__ranges/subrange.h
+++ b/libcudacxx/include/cuda/std/__ranges/subrange.h
@@ -20,6 +20,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h>
 #include <cuda/std/__concepts/constructible.h>
 #include <cuda/std/__concepts/convertible_to.h>
 #include <cuda/std/__concepts/copyable.h>
@@ -51,7 +52,6 @@
 #include <cuda/std/__type_traits/remove_pointer.h>
 #include <cuda/std/__utility/move.h>
 #include <cuda/std/cstdlib>
-#include <cuda/std/detail/libcxx/include/__assert>
 
 #if _CCCL_STD_VER >= 2017 && !defined(_CCCL_COMPILER_MSVC_2017)
 
@@ -276,8 +276,8 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT subrange : public view_interface<subrange<_I
   {
     if constexpr (sized_sentinel_for<_Sent, _Iter>)
     {
-      _LIBCUDACXX_ASSERT((__end_ - __begin_) == static_cast<iter_difference_t<_Iter>>(__n),
-                         "_CUDA_VSTD::_CUDA_VRANGES::subrange was passed an invalid size hint");
+      _CCCL_ASSERT((__end_ - __begin_) == static_cast<iter_difference_t<_Iter>>(__n),
+                   "_CUDA_VSTD::_CUDA_VRANGES::subrange was passed an invalid size hint");
     }
   }
 
diff --git a/libcudacxx/include/cuda/std/__ranges/view_interface.h b/libcudacxx/include/cuda/std/__ranges/view_interface.h
index afcec90ca4..3c64f9810c 100644
--- a/libcudacxx/include/cuda/std/__ranges/view_interface.h
+++ b/libcudacxx/include/cuda/std/__ranges/view_interface.h
@@ -20,6 +20,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h>
 #include <cuda/std/__concepts/derived_from.h>
 #include <cuda/std/__concepts/same_as.h>
 #include <cuda/std/__iterator/concepts.h>
@@ -34,7 +35,6 @@
 #include <cuda/std/__type_traits/is_reference.h>
 #include <cuda/std/__type_traits/remove_cv.h>
 #include <cuda/std/__utility/declval.h>
-#include <cuda/std/detail/libcxx/include/__assert>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_RANGES
 
@@ -135,7 +135,7 @@ class view_interface
   _LIBCUDACXX_REQUIRES(forward_range<_D2>)
   _LIBCUDACXX_HIDE_FROM_ABI constexpr decltype(auto) front()
   {
-    _LIBCUDACXX_ASSERT(!empty(), "Precondition `!empty()` not satisfied. `.front()` called on an empty view.");
+    _CCCL_ASSERT(!empty(), "Precondition `!empty()` not satisfied. `.front()` called on an empty view.");
     return *_CUDA_VRANGES::begin(__derived());
   }
 
@@ -143,7 +143,7 @@ class view_interface
   _LIBCUDACXX_REQUIRES(forward_range<const _D2>)
   _LIBCUDACXX_HIDE_FROM_ABI constexpr decltype(auto) front() const
   {
-    _LIBCUDACXX_ASSERT(!empty(), "Precondition `!empty()` not satisfied. `.front()` called on an empty view.");
+    _CCCL_ASSERT(!empty(), "Precondition `!empty()` not satisfied. `.front()` called on an empty view.");
     return *_CUDA_VRANGES::begin(__derived());
   }
 
@@ -151,7 +151,7 @@ class view_interface
   _LIBCUDACXX_REQUIRES(bidirectional_range<_D2> _LIBCUDACXX_AND common_range<_D2>)
   _LIBCUDACXX_HIDE_FROM_ABI constexpr decltype(auto) back()
   {
-    _LIBCUDACXX_ASSERT(!empty(), "Precondition `!empty()` not satisfied. `.back()` called on an empty view.");
+    _CCCL_ASSERT(!empty(), "Precondition `!empty()` not satisfied. `.back()` called on an empty view.");
     return *_CUDA_VRANGES::prev(_CUDA_VRANGES::end(__derived()));
   }
 
@@ -159,7 +159,7 @@ class view_interface
   _LIBCUDACXX_REQUIRES(bidirectional_range<const _D2> _LIBCUDACXX_AND common_range<const _D2>)
   _LIBCUDACXX_HIDE_FROM_ABI constexpr decltype(auto) back() const
   {
-    _LIBCUDACXX_ASSERT(!empty(), "Precondition `!empty()` not satisfied. `.back()` called on an empty view.");
+    _CCCL_ASSERT(!empty(), "Precondition `!empty()` not satisfied. `.back()` called on an empty view.");
     return *_CUDA_VRANGES::prev(_CUDA_VRANGES::end(__derived()));
   }
 
diff --git a/libcudacxx/include/cuda/std/__utility/exception_guard.h b/libcudacxx/include/cuda/std/__utility/exception_guard.h
index aa5457badc..1e4d23aca9 100644
--- a/libcudacxx/include/cuda/std/__utility/exception_guard.h
+++ b/libcudacxx/include/cuda/std/__utility/exception_guard.h
@@ -20,10 +20,10 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h>
 #include <cuda/std/__type_traits/is_nothrow_move_constructible.h>
 #include <cuda/std/__utility/exchange.h>
 #include <cuda/std/__utility/move.h>
-#include <cuda/std/detail/libcxx/include/__assert>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
@@ -132,7 +132,7 @@ struct __exception_guard_noexceptions
 
   _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX20 _LIBCUDACXX_NODEBUG_TYPE ~__exception_guard_noexceptions()
   {
-    _LIBCUDACXX_ASSERT(__completed_, "__exception_guard not completed with exceptions disabled");
+    _CCCL_ASSERT(__completed_, "__exception_guard not completed with exceptions disabled");
   }
 
 private:
diff --git a/libcudacxx/include/cuda/std/bit b/libcudacxx/include/cuda/std/bit
index b9fb6c358c..ab4519f2a8 100644
--- a/libcudacxx/include/cuda/std/bit
+++ b/libcudacxx/include/cuda/std/bit
@@ -29,7 +29,7 @@
 #include <cuda/std/__bit/integral.h>
 #include <cuda/std/__bit/popcount.h>
 #include <cuda/std/__bit/rotate.h>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/version>
 
 #endif // _CUDA_STD_BIT
diff --git a/libcudacxx/include/cuda/std/bitset b/libcudacxx/include/cuda/std/bitset
index 025ddcdfe6..7482b9a1bc 100644
--- a/libcudacxx/include/cuda/std/bitset
+++ b/libcudacxx/include/cuda/std/bitset
@@ -26,12 +26,12 @@
 #include <cuda/std/__algorithm/fill.h>
 #include <cuda/std/__algorithm/find.h>
 #include <cuda/std/__bit/reference.h>
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__functional/hash.h>
 #include <cuda/std/__functional/unary_function.h>
 #include <cuda/std/__type_traits/is_char_like_type.h>
 #include <cuda/std/climits>
 #include <cuda/std/cstddef>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/detail/libcxx/include/__string>
 #include <cuda/std/detail/libcxx/include/stdexcept>
 #if defined(_LIBCUDACXX_HAS_STRING_VIEW)
diff --git a/libcudacxx/include/cuda/std/concepts b/libcudacxx/include/cuda/std/concepts
index 744e722826..a75ddc4599 100644
--- a/libcudacxx/include/cuda/std/concepts
+++ b/libcudacxx/include/cuda/std/concepts
@@ -20,6 +20,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__concepts/__concept_macros.h>
 #include <cuda/std/__concepts/_One_of.h>
 #include <cuda/std/__concepts/all_of.h>
@@ -45,7 +46,6 @@
 #include <cuda/std/__concepts/semiregular.h>
 #include <cuda/std/__concepts/swappable.h>
 #include <cuda/std/__concepts/totally_ordered.h>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/version>
 
 #endif // _CUDA_STD_CONCEPTS
diff --git a/libcudacxx/include/cuda/std/detail/__annotated_ptr b/libcudacxx/include/cuda/std/detail/__annotated_ptr
index 2b29905af0..1b08c8c996 100644
--- a/libcudacxx/include/cuda/std/detail/__annotated_ptr
+++ b/libcudacxx/include/cuda/std/detail/__annotated_ptr
@@ -140,7 +140,7 @@ _CCCL_DEVICE void* __associate_address_space(void* __ptr, _Property __prop)
   if (std::is_same<_Property, access_property::shared>::value == true)
   {
     bool __b = __isShared(__ptr);
-    _LIBCUDACXX_ASSERT(__b, "");
+    _CCCL_ASSERT(__b, "");
 #if !defined(_CCCL_CUDACC_BELOW_11_2)
     __builtin_assume(__b);
 #else // ^^^ !_CCCL_CUDACC_BELOW_11_2 ^^^ / vvv _CCCL_CUDACC_BELOW_11_2 vvv
@@ -154,7 +154,7 @@ _CCCL_DEVICE void* __associate_address_space(void* __ptr, _Property __prop)
            || std::is_same<_Property, access_property>::value)
   {
     bool __b = __isGlobal(__ptr);
-    _LIBCUDACXX_ASSERT(__b, "");
+    _CCCL_ASSERT(__b, "");
 #if !defined(_CCCL_CUDACC_BELOW_11_2)
     __builtin_assume(__b);
 #else // ^^^ !_CCCL_CUDACC_BELOW_11_2 ^^^ / vvv _CCCL_CUDACC_BELOW_11_2 vvv
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__assert b/libcudacxx/include/cuda/std/detail/libcxx/include/__assert
deleted file mode 100644
index 5a99a98cfc..0000000000
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/__assert
+++ /dev/null
@@ -1,64 +0,0 @@
-// -*- C++ -*-
-//===----------------------------------------------------------------------===//
-//
-// 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) 2023 NVIDIA CORPORATION & AFFILIATES.
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef _LIBCUDACXX___ASSERT
-#define _LIBCUDACXX___ASSERT
-
-#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/__verbose_abort>
-
-// This is for backwards compatibility with code that might have been enabling
-// assertions through the Debug mode previously.
-// TODO: In LLVM 16, make it an error to define _LIBCUDACXX_DEBUG
-#if defined(_LIBCUDACXX_DEBUG)
-#  ifndef _LIBCUDACXX_ENABLE_ASSERTIONS
-#    define _LIBCUDACXX_ENABLE_ASSERTIONS 1
-#  endif
-#endif
-
-// Automatically enable assertions when the debug mode is enabled.
-#if defined(_LIBCUDACXX_ENABLE_DEBUG_MODE)
-#  ifndef _LIBCUDACXX_ENABLE_ASSERTIONS
-#    define _LIBCUDACXX_ENABLE_ASSERTIONS 1
-#  endif
-#endif
-
-#ifndef _LIBCUDACXX_ENABLE_ASSERTIONS
-#  define _LIBCUDACXX_ENABLE_ASSERTIONS _LIBCUDACXX_ENABLE_ASSERTIONS_DEFAULT
-#endif
-
-#if _LIBCUDACXX_ENABLE_ASSERTIONS != 0 && _LIBCUDACXX_ENABLE_ASSERTIONS != 1
-#  error "_LIBCUDACXX_ENABLE_ASSERTIONS must be set to 0 or 1"
-#endif
-
-#if _LIBCUDACXX_ENABLE_ASSERTIONS
-#  define _LIBCUDACXX_ASSERT(expression, message)                                                             \
-    (_CCCL_DIAG_PUSH _CCCL_DIAG_SUPPRESS_CLANG("-Wassume") __builtin_expect(static_cast<bool>(expression), 1) \
-       ? (void) 0                                                                                             \
-       : _CUDA_VSTD::__libcpp_verbose_abort(                                                                  \
-           "%s:%d: assertion %s failed: %s", __FILE__, __LINE__, #expression, message) _CCCL_DIAG_POP)
-#elif 0 // !defined(_LIBCUDACXX_ASSERTIONS_DISABLE_ASSUME) && __has_builtin(__builtin_assume)
-#  define _LIBCUDACXX_ASSERT(expression, message)                                                          \
-    (_CCCL_DIAG_PUSH _CCCL_DIAG_SUPPRESS_CLANG("-Wassume") __builtin_assume(static_cast<bool>(expression)) \
-       _CCCL_DIAG_POP)
-#else
-#  define _LIBCUDACXX_ASSERT(expression, message) ((void) 0)
-#endif
-
-#endif // _LIBCUDACXX___ASSERT
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__config b/libcudacxx/include/cuda/std/detail/libcxx/include/__config
index 5a86aef2f7..841df865f6 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/__config
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__config
@@ -1329,15 +1329,6 @@ __sanitizer_annotate_contiguous_container(const void*, const void*, const void*,
 
 #  define _LIBCUDACXX_UNUSED_VAR(x) ((void) (x))
 
-#  if __has_attribute(__format__)
-// The attribute uses 1-based indices for ordinary and static member functions.
-// The attribute uses 2-based indices for non-static member functions.
-#    define _LIBCUDACXX_ATTRIBUTE_FORMAT(archetype, format_string_index, first_format_arg_index) \
-      __attribute__((__format__(archetype, format_string_index, first_format_arg_index)))
-#  else
-#    define _LIBCUDACXX_ATTRIBUTE_FORMAT(archetype, format_string_index, first_format_arg_index) /* nothing */
-#  endif
-
 #  ifndef _LIBCUDACXX_SYS_CLOCK_DURATION
 #    define _LIBCUDACXX_SYS_CLOCK_DURATION nanoseconds
 #  endif // _LIBCUDACXX_SYS_CLOCK_DURATION
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__string b/libcudacxx/include/cuda/std/detail/libcxx/include/__string
index 1d87272e76..5e6023ec3a 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/__string
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__string
@@ -68,8 +68,8 @@ template <> struct char_traits<char8_t>;  // c++20
 #include <cuda/std/__algorithm/iterator_operations.h>
 #include <cuda/std/__algorithm/min.h>
 #include <cuda/std/__algorithm/search.h>
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__fwd/string.h>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/detail/libcxx/include/iosfwd>
 
 _CCCL_PUSH_MACROS
@@ -206,7 +206,7 @@ _LIBCUDACXX_HIDE_FROM_ABI _CharT* char_traits<_CharT>::move(char_type* __s1, con
 template <class _CharT>
 _LIBCUDACXX_HIDE_FROM_ABI _CharT* char_traits<_CharT>::copy(char_type* __s1, const char_type* __s2, size_t __n)
 {
-  _LIBCUDACXX_ASSERT(__s2 < __s1 || __s2 >= __s1 + __n, "char_traits::copy overlapped range");
+  _CCCL_ASSERT(__s2 < __s1 || __s2 >= __s1 + __n, "char_traits::copy overlapped range");
   char_type* __r = __s1;
   for (; __n; --__n, ++__s1, ++__s2)
   {
@@ -289,7 +289,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT char_traits<char>
   }
   _LIBCUDACXX_HIDE_FROM_ABI static char_type* copy(char_type* __s1, const char_type* __s2, size_t __n) noexcept
   {
-    _LIBCUDACXX_ASSERT(__s2 < __s1 || __s2 >= __s1 + __n, "char_traits::copy overlapped range");
+    _CCCL_ASSERT(__s2 < __s1 || __s2 >= __s1 + __n, "char_traits::copy overlapped range");
     return __n == 0 ? __s1 : (char_type*) memcpy(__s1, __s2, __n);
   }
   _LIBCUDACXX_HIDE_FROM_ABI static char_type* assign(char_type* __s, size_t __n, char_type __a) noexcept
@@ -404,7 +404,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT char_traits<wchar_t>
   }
   _LIBCUDACXX_HIDE_FROM_ABI static char_type* copy(char_type* __s1, const char_type* __s2, size_t __n) noexcept
   {
-    _LIBCUDACXX_ASSERT(__s2 < __s1 || __s2 >= __s1 + __n, "char_traits::copy overlapped range");
+    _CCCL_ASSERT(__s2 < __s1 || __s2 >= __s1 + __n, "char_traits::copy overlapped range");
     return __n == 0 ? __s1 : (char_type*) wmemcpy(__s1, __s2, __n);
   }
   _LIBCUDACXX_HIDE_FROM_ABI static char_type* assign(char_type* __s, size_t __n, char_type __a) noexcept
@@ -561,7 +561,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT char_traits<char8_t>
 
   _LIBCUDACXX_HIDE_FROM_ABI static char_type* copy(char_type* __s1, const char_type* __s2, size_t __n) noexcept
   {
-    _LIBCUDACXX_ASSERT(__s2 < __s1 || __s2 >= __s1 + __n, "char_traits::copy overlapped range");
+    _CCCL_ASSERT(__s2 < __s1 || __s2 >= __s1 + __n, "char_traits::copy overlapped range");
     return __n == 0 ? __s1 : (char_type*) memcpy(__s1, __s2, __n);
   }
 
@@ -768,7 +768,7 @@ char_traits<char16_t>::move(char_type* __s1, const char_type* __s2, size_t __n)
 _LIBCUDACXX_HIDE_FROM_ABI char16_t*
 char_traits<char16_t>::copy(char_type* __s1, const char_type* __s2, size_t __n) noexcept
 {
-  _LIBCUDACXX_ASSERT(__s2 < __s1 || __s2 >= __s1 + __n, "char_traits::copy overlapped range");
+  _CCCL_ASSERT(__s2 < __s1 || __s2 >= __s1 + __n, "char_traits::copy overlapped range");
   char_type* __r = __s1;
   for (; __n; --__n, ++__s1, ++__s2)
   {
@@ -909,7 +909,7 @@ char_traits<char32_t>::move(char_type* __s1, const char_type* __s2, size_t __n)
 _LIBCUDACXX_HIDE_FROM_ABI char32_t*
 char_traits<char32_t>::copy(char_type* __s1, const char_type* __s2, size_t __n) noexcept
 {
-  _LIBCUDACXX_ASSERT(__s2 < __s1 || __s2 >= __s1 + __n, "char_traits::copy overlapped range");
+  _CCCL_ASSERT(__s2 < __s1 || __s2 >= __s1 + __n, "char_traits::copy overlapped range");
   char_type* __r = __s1;
   for (; __n; --__n, ++__s1, ++__s2)
   {
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__threading_support b/libcudacxx/include/cuda/std/detail/libcxx/include/__threading_support
index bcfffa5358..fee2ff5ac8 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/__threading_support
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__threading_support
@@ -20,10 +20,10 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__functional/hash.h>
 #include <cuda/std/chrono>
 #include <cuda/std/climits>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/detail/libcxx/include/iosfwd>
 
 _CCCL_PUSH_MACROS
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__verbose_abort b/libcudacxx/include/cuda/std/detail/libcxx/include/__verbose_abort
deleted file mode 100644
index 3e9d3ca80c..0000000000
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/__verbose_abort
+++ /dev/null
@@ -1,37 +0,0 @@
-// -*- C++ -*-
-//===----------------------------------------------------------------------===//
-//
-// 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) 2023 NVIDIA CORPORATION & AFFILIATES.
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef _LIBCUDACXX___VERBOSE_ABORT
-#define _LIBCUDACXX___VERBOSE_ABORT
-
-#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/cstdlib>
-
-_LIBCUDACXX_BEGIN_NAMESPACE_STD
-
-_CCCL_NORETURN _LIBCUDACXX_ATTRIBUTE_FORMAT(__printf__, 1, 2)
-  _LIBCUDACXX_HIDE_FROM_ABI void __libcpp_verbose_abort(const char*, ...)
-{
-  NV_IF_ELSE_TARGET(NV_IS_HOST, (::abort();), (__trap();))
-  _LIBCUDACXX_UNREACHABLE();
-}
-
-_LIBCUDACXX_END_NAMESPACE_STD
-
-#endif // _LIBCUDACXX___VERBOSE_ABORT
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/algorithm b/libcudacxx/include/cuda/std/detail/libcxx/include/algorithm
index 6fbfbc7c61..620599e973 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/algorithm
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/algorithm
@@ -733,6 +733,7 @@ template <class BidirectionalIterator, class Compare>
 #include <cuda/std/__algorithm/unique.h>
 #include <cuda/std/__algorithm/unique_copy.h>
 #include <cuda/std/__algorithm/upper_bound.h>
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__iterator/distance.h>
 #include <cuda/std/__iterator/iterator_traits.h>
 #include <cuda/std/__iterator/move_iterator.h>
@@ -751,7 +752,6 @@ template <class BidirectionalIterator, class Compare>
 #include <cuda/std/__type_traits/remove_const.h>
 #include <cuda/std/bit>
 #include <cuda/std/cstddef>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/detail/libcxx/include/cstring>
 #include <cuda/std/functional>
 #include <cuda/std/initializer_list>
@@ -1147,7 +1147,7 @@ _LIBCUDACXX_HIDE_FROM_ABI _SampleIterator __sample(
                   || __is_cpp17_random_access_iterator<_SampleIterator>::value,
                 "SampleIterator must meet the requirements of RandomAccessIterator");
   typedef typename common_type<_Distance, _Difference>::type _CommonType;
-  _LIBCUDACXX_ASSERT(__n >= 0, "N must be a positive number.");
+  _CCCL_ASSERT(__n >= 0, "N must be a positive number.");
   return _CUDA_VSTD::__sample(__first, __last, __output_iter, _CommonType(__n), __g, _PopCategory());
 }
 
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/array b/libcudacxx/include/cuda/std/detail/libcxx/include/array
index 61afff0000..f32fd2e9a6 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/array
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/array
@@ -123,6 +123,7 @@ template <size_t I, class T, size_t N> const T&& get(const array<T, N>&&) noexce
 #include <cuda/std/__algorithm/fill_n.h>
 #include <cuda/std/__algorithm/lexicographical_compare.h>
 #include <cuda/std/__algorithm/swap_ranges.h>
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__iterator/reverse_iterator.h>
 #include <cuda/std/__tuple_dir/sfinae_helpers.h>
 #include <cuda/std/__type_traits/conditional.h>
@@ -140,7 +141,6 @@ template <size_t I, class T, size_t N> const T&& get(const array<T, N>&&) noexce
 #include <cuda/std/__utility/unreachable.h>
 #include <cuda/std/cstdint>
 #include <cuda/std/cstdlib>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/detail/libcxx/include/stdexcept>
 #include <cuda/std/limits>
 
@@ -265,12 +265,12 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT array
   // element access:
   _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 reference operator[](size_type __n) noexcept
   {
-    _LIBCUDACXX_ASSERT(__n < _Size, "out-of-bounds access in std::array<T, N>");
+    _CCCL_ASSERT(__n < _Size, "out-of-bounds access in std::array<T, N>");
     return __elems_[__n];
   }
   _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 const_reference operator[](size_type __n) const noexcept
   {
-    _LIBCUDACXX_ASSERT(__n < _Size, "out-of-bounds access in std::array<T, N>");
+    _CCCL_ASSERT(__n < _Size, "out-of-bounds access in std::array<T, N>");
     return __elems_[__n];
   }
 
@@ -435,14 +435,14 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT array<_Tp, 0>
   // element access:
   _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 reference operator[](size_type) noexcept
   {
-    _LIBCUDACXX_ASSERT(false, "cannot call array<T, 0>::operator[] on a zero-sized array");
+    _CCCL_ASSERT(false, "cannot call array<T, 0>::operator[] on a zero-sized array");
     _LIBCUDACXX_UNREACHABLE();
     return *data();
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 const_reference operator[](size_type) const noexcept
   {
-    _LIBCUDACXX_ASSERT(false, "cannot call array<T, 0>::operator[] on a zero-sized array");
+    _CCCL_ASSERT(false, "cannot call array<T, 0>::operator[] on a zero-sized array");
     _LIBCUDACXX_UNREACHABLE();
     return *data();
   }
@@ -461,28 +461,28 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT array<_Tp, 0>
 
   _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 reference front() noexcept
   {
-    _LIBCUDACXX_ASSERT(false, "cannot call array<T, 0>::front() on a zero-sized array");
+    _CCCL_ASSERT(false, "cannot call array<T, 0>::front() on a zero-sized array");
     _LIBCUDACXX_UNREACHABLE();
     return *data();
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 const_reference front() const noexcept
   {
-    _LIBCUDACXX_ASSERT(false, "cannot call array<T, 0>::front() on a zero-sized array");
+    _CCCL_ASSERT(false, "cannot call array<T, 0>::front() on a zero-sized array");
     _LIBCUDACXX_UNREACHABLE();
     return *data();
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 reference back() noexcept
   {
-    _LIBCUDACXX_ASSERT(false, "cannot call array<T, 0>::back() on a zero-sized array");
+    _CCCL_ASSERT(false, "cannot call array<T, 0>::back() on a zero-sized array");
     _LIBCUDACXX_UNREACHABLE();
     return *data();
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 const_reference back() const noexcept
   {
-    _LIBCUDACXX_ASSERT(false, "cannot call array<T, 0>::back() on a zero-sized array");
+    _CCCL_ASSERT(false, "cannot call array<T, 0>::back() on a zero-sized array");
     _LIBCUDACXX_UNREACHABLE();
     return *data();
   }
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/barrier b/libcudacxx/include/cuda/std/detail/libcxx/include/barrier
index 54e203ac99..f8767a4c4d 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/barrier
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/barrier
@@ -53,11 +53,11 @@ namespace std
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__new_>
 #include <cuda/std/atomic>
 #include <cuda/std/chrono>
 #include <cuda/std/cstddef>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 
 _CCCL_PUSH_MACROS
 
@@ -104,7 +104,7 @@ class alignas(64) __barrier_base
               __current_expected = __expected, __last_node = (__current_expected >> 1);
     for (size_t __round = 0;; ++__round)
     {
-      _LIBCUDACXX_ASSERT(__round <= 63, "");
+      _CCCL_ASSERT(__round <= 63, "");
       if (__current_expected == 1)
       {
         return true;
@@ -121,7 +121,7 @@ class alignas(64) __barrier_base
           __libcpp_thread_favorite_barrier_index = __current;
         }
 #  endif
-        _LIBCUDACXX_ASSERT(__current <= __last_node, "");
+        _CCCL_ASSERT(__current <= __last_node, "");
         __phase_t expect = __old_phase;
         if (__current == __last_node && (__current_expected & 1))
         {
@@ -130,7 +130,7 @@ class alignas(64) __barrier_base
           {
             break; // I'm 1 in 1, go to next __round
           }
-          _LIBCUDACXX_ASSERT(expect == __full_step, "");
+          _CCCL_ASSERT(expect == __full_step, "");
         }
         else if (__state[__current].__tickets[__round].__phase.compare_exchange_strong(
                    expect, __half_step, memory_order_acq_rel))
@@ -144,9 +144,9 @@ class alignas(64) __barrier_base
           {
             break; // I'm 2 in 2, go to next __round
           }
-          _LIBCUDACXX_ASSERT(expect == __full_step, "");
+          _CCCL_ASSERT(expect == __full_step, "");
         }
-        _LIBCUDACXX_ASSERT(__round == 0 && expect == __full_step, "");
+        _CCCL_ASSERT(__round == 0 && expect == __full_step, "");
       }
       __current_expected = (__current_expected >> 1) + (__current_expected & 1);
       __current &= ~(1 << __round);
@@ -164,7 +164,7 @@ public:
       , __phase(0)
       , __state((__expected + 1) >> 1)
   {
-    _LIBCUDACXX_ASSERT(__expected >= 0, "");
+    _CCCL_ASSERT(__expected >= 0, "");
   }
 
   _CCCL_HIDE_FROM_ABI ~__barrier_base() = default;
@@ -174,7 +174,7 @@ public:
 
   _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI arrival_token arrive(ptrdiff_t update = 1)
   {
-    _LIBCUDACXX_ASSERT(update > 0, "");
+    _CCCL_ASSERT(update > 0, "");
     auto __old_phase = __phase.load(memory_order_relaxed);
     for (; update; --update)
     {
@@ -312,7 +312,7 @@ public:
     auto const __result       = __arrived.fetch_sub(__update, memory_order_acq_rel) - __update;
     auto const __new_expected = __expected.load(memory_order_relaxed);
 
-    _LIBCUDACXX_ASSERT(__result >= 0, "");
+    _CCCL_ASSERT(__result >= 0, "");
 
     if (0 == __result)
     {
@@ -372,7 +372,7 @@ private:
 #  if _CCCL_STD_VER > 2011
     // This debug assert is not supported in C++11 due to resulting in a
     // multi-statement constexpr function.
-    _LIBCUDACXX_ASSERT(__count >= 0, "Count must be non-negative.");
+    _CCCL_ASSERT(__count >= 0, "Count must be non-negative.");
 #  endif // _CCCL_STD_VER > 2011
     return (((1u << 31) - __count) << 32) | ((1u << 31) - __count);
   }
@@ -397,7 +397,7 @@ public:
   __barrier_base(ptrdiff_t __count, __empty_completion = __empty_completion())
       : __phase_arrived_expected(__init(__count))
   {
-    _LIBCUDACXX_ASSERT(__count >= 0, "");
+    _CCCL_ASSERT(__count >= 0, "");
   }
 
   _CCCL_HIDE_FROM_ABI ~__barrier_base() = default;
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/chrono b/libcudacxx/include/cuda/std/detail/libcxx/include/chrono
index 5274c69dde..2add525051 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/chrono
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/chrono
@@ -839,13 +839,13 @@ constexpr chrono::year                                  operator ""y(unsigned lo
 #  endif // _CCCL_COMPILER_NVRTC
 #endif // __cuda_std__
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__type_traits/common_type.h>
 #include <cuda/std/__type_traits/enable_if.h>
 #include <cuda/std/__type_traits/integral_constant.h>
 #include <cuda/std/__type_traits/is_convertible.h>
 #include <cuda/std/__type_traits/is_floating_point.h>
 #include <cuda/std/ctime>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/limits>
 #include <cuda/std/ratio>
 
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/complex b/libcudacxx/include/cuda/std/detail/libcxx/include/complex
index 5f9693a662..86a5c01dc9 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/complex
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/complex
@@ -242,6 +242,7 @@ template<class T> complex<T> tanh (const complex<T>&);
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__complex/vector_support.h>
 #include <cuda/std/__type_traits/enable_if.h>
 #include <cuda/std/__type_traits/is_constructible.h>
@@ -252,7 +253,6 @@ template<class T> complex<T> tanh (const complex<T>&);
 #include <cuda/std/climits>
 #include <cuda/std/cmath>
 #include <cuda/std/cstdint>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/type_traits>
 #include <cuda/std/version>
 
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/cstddef b/libcudacxx/include/cuda/std/detail/libcxx/include/cstddef
index ee109ef536..07e16d8cc7 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/cstddef
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/cstddef
@@ -43,10 +43,10 @@ Types:
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__cuda/cstddef_prelude.h>
 #include <cuda/std/__type_traits/enable_if.h>
 #include <cuda/std/__type_traits/is_integral.h>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/version>
 
 _CCCL_PUSH_MACROS
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/iosfwd b/libcudacxx/include/cuda/std/detail/libcxx/include/iosfwd
index 7f9f4a2665..ce7d2653be 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/iosfwd
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/iosfwd
@@ -101,8 +101,8 @@ typedef fpos<char_traits<wchar_t>::state_type> wstreampos;
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__fwd/string.h>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/version>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/latch b/libcudacxx/include/cuda/std/detail/libcxx/include/latch
index 3ec4935321..8041fce50d 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/latch
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/latch
@@ -49,9 +49,8 @@ namespace std
 #endif // no system header
 
 #include <cuda/std/__atomic/api/owned.h>
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/atomic>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
-#include <cuda/std/detail/libcxx/include/
 #include <cuda/std/limits>
 
 _CCCL_PUSH_MACROS
@@ -84,9 +83,9 @@ public:
 
   _LIBCUDACXX_HIDE_FROM_ABI void count_down(ptrdiff_t __update = 1)
   {
-    _LIBCUDACXX_ASSERT(__update > 0, "");
+    _CCCL_ASSERT(__update > 0, "");
     auto const __old = __counter.fetch_sub(__update, memory_order_release);
-    _LIBCUDACXX_ASSERT(__old >= __update, "");
+    _CCCL_ASSERT(__old >= __update, "");
     if (__old == __update)
     {
       __counter.notify_all();
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/limits b/libcudacxx/include/cuda/std/detail/libcxx/include/limits
index 1d6a409aee..42e1c88c2d 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/limits
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/limits
@@ -112,8 +112,8 @@ template<> class numeric_limits<cv long double>;
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/climits>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/type_traits>
 #include <cuda/std/version>
 
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/optional b/libcudacxx/include/cuda/std/detail/libcxx/include/optional
index 29e2d0f91b..b1868a3d9a 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/optional
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/optional
@@ -169,6 +169,7 @@ template<class T>
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__concepts/__concept_macros.h>
 #include <cuda/std/__concepts/invocable.h>
 #include <cuda/std/__exception/terminate.h>
@@ -199,9 +200,7 @@ template<class T>
 #include <cuda/std/__utility/move.h>
 #include <cuda/std/__utility/swap.h>
 #include <cuda/std/cstdlib>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/detail/libcxx/include/__availability>
-#include <cuda/std/detail/libcxx/include/__verbose_abort>
 #include <cuda/std/initializer_list>
 
 // standard-mandated includes
@@ -397,7 +396,7 @@ struct __optional_storage_base : __optional_destruct_base<_Tp>
   template <class... _Args>
   _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX20 void __construct(_Args&&... __args)
   {
-    _LIBCUDACXX_ASSERT(!has_value(), "__construct called for engaged __optional_storage");
+    _CCCL_ASSERT(!has_value(), "__construct called for engaged __optional_storage");
 #  if _CCCL_STD_VER > 2017
     _CUDA_VSTD::construct_at(_CUDA_VSTD::addressof(this->__val_), _CUDA_VSTD::forward<_Args>(__args)...);
 #  else
@@ -851,37 +850,37 @@ public:
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr add_pointer_t<value_type const> operator->() const
   {
-    _LIBCUDACXX_ASSERT(this->has_value(), "optional operator-> called on a disengaged value");
+    _CCCL_ASSERT(this->has_value(), "optional operator-> called on a disengaged value");
     return _CUDA_VSTD::addressof(this->__get());
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr add_pointer_t<value_type> operator->()
   {
-    _LIBCUDACXX_ASSERT(this->has_value(), "optional operator-> called on a disengaged value");
+    _CCCL_ASSERT(this->has_value(), "optional operator-> called on a disengaged value");
     return _CUDA_VSTD::addressof(this->__get());
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr const value_type& operator*() const& noexcept
   {
-    _LIBCUDACXX_ASSERT(this->has_value(), "optional operator* called on a disengaged value");
+    _CCCL_ASSERT(this->has_value(), "optional operator* called on a disengaged value");
     return this->__get();
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr value_type& operator*() & noexcept
   {
-    _LIBCUDACXX_ASSERT(this->has_value(), "optional operator* called on a disengaged value");
+    _CCCL_ASSERT(this->has_value(), "optional operator* called on a disengaged value");
     return this->__get();
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr value_type&& operator*() && noexcept
   {
-    _LIBCUDACXX_ASSERT(this->has_value(), "optional operator* called on a disengaged value");
+    _CCCL_ASSERT(this->has_value(), "optional operator* called on a disengaged value");
     return _CUDA_VSTD::move(this->__get());
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr const value_type&& operator*() const&& noexcept
   {
-    _LIBCUDACXX_ASSERT(this->has_value(), "optional operator* called on a disengaged value");
+    _CCCL_ASSERT(this->has_value(), "optional operator* called on a disengaged value");
     return _CUDA_VSTD::move(this->__get());
   }
 
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/ratio b/libcudacxx/include/cuda/std/detail/libcxx/include/ratio
index e59dbd5826..f3c69f00fa 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/ratio
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/ratio
@@ -87,10 +87,10 @@ typedef ratio<1000000000000000000000000, 1> yotta;  // not supported
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__type_traits/integral_constant.h>
 #include <cuda/std/climits>
 #include <cuda/std/cstdint>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/type_traits>
 
 _CCCL_PUSH_MACROS
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/semaphore b/libcudacxx/include/cuda/std/detail/libcxx/include/semaphore
index 581fcd9b0a..e0ca53ad88 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/semaphore
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/semaphore
@@ -55,9 +55,9 @@ using binary_semaphore = counting_semaphore<1>;
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__type_traits/conditional.h>
 #include <cuda/std/atomic>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/type_traits>
 
 _CCCL_PUSH_MACROS
@@ -221,7 +221,7 @@ public:
 
   _LIBCUDACXX_HIDE_FROM_ABI void release(ptrdiff_t __update = 1)
   {
-    _LIBCUDACXX_ASSERT(__update == 1, "");
+    _CCCL_ASSERT(__update == 1, "");
     __available.store(1, memory_order_release);
     __available.notify_one();
     (void) __update;
@@ -283,7 +283,7 @@ class __sem_semaphore_base
     bool const __post_two = __back_amount > 1;
     auto const __success =
       (!__post_one || __libcpp_semaphore_post(&__semaphore)) && (!__post_two || __libcpp_semaphore_post(&__semaphore));
-    _LIBCUDACXX_ASSERT(__success, "");
+    _CCCL_ASSERT(__success, "");
     if (!__post_one || !__post_two)
     {
       __backbuffer.fetch_add(!__post_one ? 2 : 1, memory_order_relaxed);
@@ -354,12 +354,12 @@ _LIBCUDACXX_HIDE_FROM_ABI void __release_slow(ptrdiff_t __post_amount)
   }
   auto const __success =
     (!__post_one || __libcpp_semaphore_post(&__semaphore)) && (!__post_two || __libcpp_semaphore_post(&__semaphore));
-  _LIBCUDACXX_ASSERT(__success, "");
+  _CCCL_ASSERT(__success, "");
 #  else
     for (; __post_amount; --__post_amount)
     {
       auto const __success = __libcpp_semaphore_post(&__semaphore);
-      _LIBCUDACXX_ASSERT(__success, "");
+      _CCCL_ASSERT(__success, "");
     }
 #  endif
 }
@@ -387,23 +387,23 @@ _LIBCUDACXX_HIDE_FROM_ABI __sem_semaphore_base(ptrdiff_t __count = 0)
     , __backbuffer(0)
 #  endif
 {
-  _LIBCUDACXX_ASSERT(__count <= max(), "");
+  _CCCL_ASSERT(__count <= max(), "");
   auto const __success =
 #  ifndef _LIBCUDACXX_HAS_NO_SEMAPHORE_FRONT_BUFFER
     __libcpp_semaphore_init(&__semaphore, 0);
 #  else
       __libcpp_semaphore_init(&__semaphore, __count);
 #  endif
-  _LIBCUDACXX_ASSERT(__success, "");
+  _CCCL_ASSERT(__success, "");
 }
 
 _LIBCUDACXX_HIDE_FROM_ABI ~__sem_semaphore_base()
 {
 #  ifndef _LIBCUDACXX_HAS_NO_SEMAPHORE_FRONT_BUFFER
-  _LIBCUDACXX_ASSERT(0 == (__frontbuffer.load(memory_order_relaxed) & ~0u), "");
+  _CCCL_ASSERT(0 == (__frontbuffer.load(memory_order_relaxed) & ~0u), "");
 #  endif
   auto const __success = __libcpp_semaphore_destroy(&__semaphore);
-  _LIBCUDACXX_ASSERT(__success, "");
+  _CCCL_ASSERT(__success, "");
 }
 
 __sem_semaphore_base(const __sem_semaphore_base&)            = delete;
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/span b/libcudacxx/include/cuda/std/detail/libcxx/include/span
index e2c0a9dd47..4320120195 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/span
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/span
@@ -138,6 +138,7 @@ template<class R>
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__concepts/convertible_to.h>
 #include <cuda/std/__concepts/equality_comparable.h>
 #include <cuda/std/__fwd/array.h>
@@ -171,7 +172,6 @@ template<class R>
 #include <cuda/std/__utility/declval.h>
 #include <cuda/std/array>
 #include <cuda/std/cstddef> // for ptrdiff_t
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 
 // standard-mandated includes
 #include <cuda/std/version>
@@ -335,7 +335,7 @@ public:
       : __data_{_CUDA_VSTD::to_address(__first)}
   {
     (void) __count;
-    _LIBCUDACXX_ASSERT(_Extent == __count, "size mismatch in span's constructor (iterator, len)");
+    _CCCL_ASSERT(_Extent == __count, "size mismatch in span's constructor (iterator, len)");
   }
 
   _LIBCUDACXX_TEMPLATE(class _It, class _End)
@@ -345,22 +345,22 @@ public:
       : __data_{_CUDA_VSTD::to_address(__first)}
   {
     (void) __last;
-    _LIBCUDACXX_ASSERT((__last - __first >= 0), "invalid range in span's constructor (iterator, sentinel)");
-    _LIBCUDACXX_ASSERT(__last - __first == _Extent,
-                       "invalid range in span's constructor (iterator, sentinel): last - first != extent");
+    _CCCL_ASSERT((__last - __first >= 0), "invalid range in span's constructor (iterator, sentinel)");
+    _CCCL_ASSERT(__last - __first == _Extent,
+                 "invalid range in span's constructor (iterator, sentinel): last - first != extent");
   }
 #  else // ^^^ C++17 ^^^ / vvv C++14 vvv
   _LIBCUDACXX_HIDE_FROM_ABI constexpr span(pointer __ptr, size_type __count)
       : __data_{__ptr}
   {
     (void) __count;
-    _LIBCUDACXX_ASSERT(_Extent == __count, "size mismatch in span's constructor (ptr, len)");
+    _CCCL_ASSERT(_Extent == __count, "size mismatch in span's constructor (ptr, len)");
   }
   _LIBCUDACXX_HIDE_FROM_ABI constexpr span(pointer __f, pointer __l)
       : __data_{__f}
   {
     (void) __l;
-    _LIBCUDACXX_ASSERT(_Extent == distance(__f, __l), "size mismatch in span's constructor (ptr, ptr)");
+    _CCCL_ASSERT(_Extent == distance(__f, __l), "size mismatch in span's constructor (ptr, ptr)");
   }
 #  endif // _CCCL_STD_VER <= 2014 || _CCCL_COMPILER_MSVC_2017
 
@@ -393,7 +393,7 @@ public:
   _LIBCUDACXX_HIDE_FROM_ABI constexpr explicit span(_Range&& __r)
       : __data_{_CUDA_VRANGES::data(__r)}
   {
-    _LIBCUDACXX_ASSERT(_CUDA_VRANGES::size(__r) == _Extent, "size mismatch in span's constructor (range)");
+    _CCCL_ASSERT(_CUDA_VRANGES::size(__r) == _Extent, "size mismatch in span's constructor (range)");
   }
 #  else // ^^^ C++17 ^^^ / vvv C++14 vvv
   _LIBCUDACXX_TEMPLATE(class _Container)
@@ -401,7 +401,7 @@ public:
   _LIBCUDACXX_HIDE_FROM_ABI constexpr span(_Container& __c) noexcept(noexcept(_CUDA_VSTD::data(__c)))
       : __data_{_CUDA_VSTD::data(__c)}
   {
-    _LIBCUDACXX_ASSERT(_Extent == _CUDA_VSTD::size(__c), "size mismatch in span's constructor (other span)");
+    _CCCL_ASSERT(_Extent == _CUDA_VSTD::size(__c), "size mismatch in span's constructor (other span)");
   }
 
   _LIBCUDACXX_TEMPLATE(class _Container)
@@ -409,7 +409,7 @@ public:
   _LIBCUDACXX_HIDE_FROM_ABI constexpr span(const _Container& __c) noexcept(noexcept(_CUDA_VSTD::data(__c)))
       : __data_{_CUDA_VSTD::data(__c)}
   {
-    _LIBCUDACXX_ASSERT(_Extent == _CUDA_VSTD::size(__c), "size mismatch in span's constructor (other span)");
+    _CCCL_ASSERT(_Extent == _CUDA_VSTD::size(__c), "size mismatch in span's constructor (other span)");
   }
 #  endif // _CCCL_STD_VER <= 2014 || _CCCL_COMPILER_MSVC_2017
 
@@ -425,7 +425,7 @@ public:
   _LIBCUDACXX_HIDE_FROM_ABI constexpr explicit span(const span<_OtherElementType, dynamic_extent>& __other) noexcept
       : __data_{__other.data()}
   {
-    _LIBCUDACXX_ASSERT(_Extent == __other.size(), "size mismatch in span's constructor (other span)");
+    _CCCL_ASSERT(_Extent == __other.size(), "size mismatch in span's constructor (other span)");
   }
 
   //  ~span() noexcept = default;
@@ -446,13 +446,13 @@ public:
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr span<element_type, dynamic_extent> first(size_type __count) const noexcept
   {
-    _LIBCUDACXX_ASSERT(__count <= size(), "span<T, N>::first(count): count out of range");
+    _CCCL_ASSERT(__count <= size(), "span<T, N>::first(count): count out of range");
     return {data(), __count};
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr span<element_type, dynamic_extent> last(size_type __count) const noexcept
   {
-    _LIBCUDACXX_ASSERT(__count <= size(), "span<T, N>::last(count): count out of range");
+    _CCCL_ASSERT(__count <= size(), "span<T, N>::last(count): count out of range");
     return {data() + size() - __count, __count};
   }
 
@@ -471,14 +471,14 @@ public:
   _LIBCUDACXX_HIDE_FROM_ABI constexpr span<element_type, dynamic_extent>
   subspan(size_type __offset, size_type __count = dynamic_extent) const noexcept
   {
-    _LIBCUDACXX_ASSERT(__offset <= size(), "span<T, N>::subspan(offset, count): offset out of range");
-    _LIBCUDACXX_ASSERT(__count <= size() || __count == dynamic_extent,
-                       "span<T, N>::subspan(offset, count): count out of range");
+    _CCCL_ASSERT(__offset <= size(), "span<T, N>::subspan(offset, count): offset out of range");
+    _CCCL_ASSERT(__count <= size() || __count == dynamic_extent,
+                 "span<T, N>::subspan(offset, count): count out of range");
     if (__count == dynamic_extent)
     {
       return {data() + __offset, size() - __offset};
     }
-    _LIBCUDACXX_ASSERT(__count <= size() - __offset, "span<T, N>::subspan(offset, count): offset + count out of range");
+    _CCCL_ASSERT(__count <= size() - __offset, "span<T, N>::subspan(offset, count): offset + count out of range");
     return {data() + __offset, __count};
   }
 
@@ -497,19 +497,19 @@ public:
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr reference operator[](size_type __idx) const noexcept
   {
-    _LIBCUDACXX_ASSERT(__idx < size(), "span<T, N>::operator[](index): index out of range");
+    _CCCL_ASSERT(__idx < size(), "span<T, N>::operator[](index): index out of range");
     return __data_[__idx];
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr reference front() const noexcept
   {
-    _LIBCUDACXX_ASSERT(!empty(), "span<T, N>::front() on empty span");
+    _CCCL_ASSERT(!empty(), "span<T, N>::front() on empty span");
     return __data_[0];
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr reference back() const noexcept
   {
-    _LIBCUDACXX_ASSERT(!empty(), "span<T, N>::back() on empty span");
+    _CCCL_ASSERT(!empty(), "span<T, N>::back() on empty span");
     return __data_[size() - 1];
   }
 
@@ -591,7 +591,7 @@ public:
       : __data_(_CUDA_VSTD::to_address(__first))
       , __size_(__last - __first)
   {
-    _LIBCUDACXX_ASSERT(__last - __first >= 0, "invalid range in span's constructor (iterator, sentinel)");
+    _CCCL_ASSERT(__last - __first >= 0, "invalid range in span's constructor (iterator, sentinel)");
   }
 
 #  else // ^^^ C++17 ^^^ / vvv C++14 vvv
@@ -660,26 +660,26 @@ public:
   template <size_t _Count>
   _LIBCUDACXX_HIDE_FROM_ABI constexpr span<element_type, _Count> first() const noexcept
   {
-    _LIBCUDACXX_ASSERT(_Count == 0 || _Count <= size(), "span<T>::first<Count>(): Count out of range");
+    _CCCL_ASSERT(_Count == 0 || _Count <= size(), "span<T>::first<Count>(): Count out of range");
     return span<element_type, _Count>{data(), _Count};
   }
 
   template <size_t _Count>
   _LIBCUDACXX_HIDE_FROM_ABI constexpr span<element_type, _Count> last() const noexcept
   {
-    _LIBCUDACXX_ASSERT(_Count == 0 || _Count <= size(), "span<T>::last<Count>(): Count out of range");
+    _CCCL_ASSERT(_Count == 0 || _Count <= size(), "span<T>::last<Count>(): Count out of range");
     return span<element_type, _Count>{data() + size() - _Count, _Count};
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr span<element_type, dynamic_extent> first(size_type __count) const noexcept
   {
-    _LIBCUDACXX_ASSERT(__count <= size(), "span<T>::first(count): count out of range");
+    _CCCL_ASSERT(__count <= size(), "span<T>::first(count): count out of range");
     return {data(), __count};
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr span<element_type, dynamic_extent> last(size_type __count) const noexcept
   {
-    _LIBCUDACXX_ASSERT(__count <= size(), "span<T>::last(count): count out of range");
+    _CCCL_ASSERT(__count <= size(), "span<T>::last(count): count out of range");
     return {data() + size() - __count, __count};
   }
 
@@ -689,23 +689,22 @@ public:
   template <size_t _Offset, size_t _Count = dynamic_extent>
   _LIBCUDACXX_HIDE_FROM_ABI constexpr __subspan_t<_Offset, _Count> subspan() const noexcept
   {
-    _LIBCUDACXX_ASSERT(_Offset == 0 || _Offset <= size(), "span<T>::subspan<Offset, Count>(): Offset out of range");
-    _LIBCUDACXX_ASSERT(_Count == dynamic_extent || _Count == 0 || _Count <= size() - _Offset,
-                       "span<T>::subspan<Offset, Count>(): Offset + Count out of range");
+    _CCCL_ASSERT(_Offset == 0 || _Offset <= size(), "span<T>::subspan<Offset, Count>(): Offset out of range");
+    _CCCL_ASSERT(_Count == dynamic_extent || _Count == 0 || _Count <= size() - _Offset,
+                 "span<T>::subspan<Offset, Count>(): Offset + Count out of range");
     return __subspan_t<_Offset, _Count>{data() + _Offset, _Count == dynamic_extent ? size() - _Offset : _Count};
   }
 
   constexpr span<element_type, dynamic_extent> _LIBCUDACXX_HIDE_FROM_ABI
   subspan(size_type __offset, size_type __count = dynamic_extent) const noexcept
   {
-    _LIBCUDACXX_ASSERT(__offset <= size(), "span<T>::subspan(offset, count): offset out of range");
-    _LIBCUDACXX_ASSERT(__count <= size() || __count == dynamic_extent,
-                       "span<T>::subspan(offset, count): count out of range");
+    _CCCL_ASSERT(__offset <= size(), "span<T>::subspan(offset, count): offset out of range");
+    _CCCL_ASSERT(__count <= size() || __count == dynamic_extent, "span<T>::subspan(offset, count): count out of range");
     if (__count == dynamic_extent)
     {
       return {data() + __offset, size() - __offset};
     }
-    _LIBCUDACXX_ASSERT(__count <= size() - __offset, "span<T>::subspan(offset, count): offset + count out of range");
+    _CCCL_ASSERT(__count <= size() - __offset, "span<T>::subspan(offset, count): offset + count out of range");
     return {data() + __offset, __count};
   }
 
@@ -724,19 +723,19 @@ public:
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr reference operator[](size_type __idx) const noexcept
   {
-    _LIBCUDACXX_ASSERT(__idx < size(), "span<T>::operator[](index): index out of range");
+    _CCCL_ASSERT(__idx < size(), "span<T>::operator[](index): index out of range");
     return __data_[__idx];
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr reference front() const noexcept
   {
-    _LIBCUDACXX_ASSERT(!empty(), "span<T>::front() on empty span");
+    _CCCL_ASSERT(!empty(), "span<T>::front() on empty span");
     return __data_[0];
   }
 
   _LIBCUDACXX_HIDE_FROM_ABI constexpr reference back() const noexcept
   {
-    _LIBCUDACXX_ASSERT(!empty(), "span<T>::back() on empty span");
+    _CCCL_ASSERT(!empty(), "span<T>::back() on empty span");
     return __data_[size() - 1];
   }
 
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/stdexcept b/libcudacxx/include/cuda/std/detail/libcxx/include/stdexcept
index 0c5ee77e8f..e91a489df3 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/stdexcept
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/stdexcept
@@ -22,8 +22,8 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__exception/terminate.h>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 
 #ifndef _CCCL_NO_EXCEPTIONS
 #  include <stdexcept>
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/tuple b/libcudacxx/include/cuda/std/detail/libcxx/include/tuple
index f8392629a4..4967530968 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/tuple
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/tuple
@@ -155,6 +155,7 @@ template <class... Types>
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__functional/unwrap_ref.h>
 #include <cuda/std/__fwd/array.h>
 #include <cuda/std/__memory/allocator_arg_t.h>
@@ -178,7 +179,6 @@ template <class... Types>
 #include <cuda/std/climits>
 #include <cuda/std/cstddef>
 #include <cuda/std/cstdint>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/type_traits>
 #include <cuda/std/utility>
 
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/variant b/libcudacxx/include/cuda/std/detail/libcxx/include/variant
index fccb314034..984cd2af4c 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/variant
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/variant
@@ -215,7 +215,7 @@ C++20
 #  pragma system_header
 #endif // no system header
 
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/detail/libcxx/include/__availability>
 #ifndef _LIBCUDACXX_HAS_NO_SPACESHIP_OPERATOR
 #  include <cuda/std/__compare/common_comparison_category.h>
diff --git a/libcudacxx/include/cuda/std/expected b/libcudacxx/include/cuda/std/expected
index 213b2497e6..ee024d8dcd 100644
--- a/libcudacxx/include/cuda/std/expected
+++ b/libcudacxx/include/cuda/std/expected
@@ -21,11 +21,11 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__expected/bad_expected_access.h>
 #include <cuda/std/__expected/expected.h>
 #include <cuda/std/__expected/unexpect.h>
 #include <cuda/std/__expected/unexpected.h>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/version>
 
 #endif //_CUDA_STD_EXPECTED
diff --git a/libcudacxx/include/cuda/std/functional b/libcudacxx/include/cuda/std/functional
index 47f5806f64..8100a70dad 100644
--- a/libcudacxx/include/cuda/std/functional
+++ b/libcudacxx/include/cuda/std/functional
@@ -23,6 +23,7 @@
 
 #include <cuda/std/__algorithm/max.h>
 #include <cuda/std/__algorithm/min.h>
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__functional/binary_function.h>
 #include <cuda/std/__functional/binary_negate.h>
 #include <cuda/std/__functional/bind.h>
@@ -50,7 +51,6 @@
 #include <cuda/std/__functional/unary_negate.h>
 #include <cuda/std/__functional/unwrap_ref.h>
 #include <cuda/std/__functional/weak_result_type.h>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/detail/libcxx/include/iosfwd> // for forward declarations of vector and string
 #include <cuda/std/version>
 
diff --git a/libcudacxx/include/cuda/std/inplace_vector b/libcudacxx/include/cuda/std/inplace_vector
index 6c2963c8a1..4faa9cf577 100644
--- a/libcudacxx/include/cuda/std/inplace_vector
+++ b/libcudacxx/include/cuda/std/inplace_vector
@@ -32,6 +32,7 @@
 #  include <cuda/std/__algorithm/remove_if.h>
 #  include <cuda/std/__algorithm/rotate.h>
 #  include <cuda/std/__algorithm/swap_ranges.h>
+#  include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #  include <cuda/std/__exception/terminate.h>
 #  include <cuda/std/__iterator/advance.h>
 #  include <cuda/std/__iterator/concepts.h>
@@ -66,7 +67,6 @@
 #  include <cuda/std/__utility/forward.h>
 #  include <cuda/std/__utility/move.h>
 #  include <cuda/std/cstdint>
-#  include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #  include <cuda/std/detail/libcxx/include/stdexcept>
 #  include <cuda/std/initializer_list>
 #  include <cuda/std/limits>
diff --git a/libcudacxx/include/cuda/std/iterator b/libcudacxx/include/cuda/std/iterator
index 67a6b5faeb..3d0a516a3f 100644
--- a/libcudacxx/include/cuda/std/iterator
+++ b/libcudacxx/include/cuda/std/iterator
@@ -21,6 +21,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__iterator/access.h>
 #include <cuda/std/__iterator/advance.h>
 #include <cuda/std/__iterator/back_insert_iterator.h>
@@ -57,7 +58,6 @@
 #include <cuda/std/__iterator/sortable.h>
 #include <cuda/std/__iterator/unreachable_sentinel.h>
 #include <cuda/std/__iterator/wrap_iter.h>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/detail/libcxx/include/iosfwd> // for forward declarations of vector and string.
 #include <cuda/std/version>
 
diff --git a/libcudacxx/include/cuda/std/mdspan b/libcudacxx/include/cuda/std/mdspan
index 798e1a69d3..52a5489c57 100644
--- a/libcudacxx/include/cuda/std/mdspan
+++ b/libcudacxx/include/cuda/std/mdspan
@@ -21,6 +21,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__mdspan/default_accessor.h>
 #include <cuda/std/__mdspan/dynamic_extent.h>
 #include <cuda/std/__mdspan/extents.h>
@@ -32,7 +33,6 @@
 #include <cuda/std/__mdspan/mdspan.h>
 #include <cuda/std/__mdspan/static_array.h>
 #include <cuda/std/__mdspan/submdspan.h>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/version>
 
 #endif // _CUDA_STD_MDSPAN
diff --git a/libcudacxx/include/cuda/std/numeric b/libcudacxx/include/cuda/std/numeric
index 4057620006..f0f13995b0 100644
--- a/libcudacxx/include/cuda/std/numeric
+++ b/libcudacxx/include/cuda/std/numeric
@@ -21,6 +21,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__numeric/accumulate.h>
 #include <cuda/std/__numeric/adjacent_difference.h>
 #include <cuda/std/__numeric/exclusive_scan.h>
@@ -34,7 +35,6 @@
 #include <cuda/std/__numeric/transform_exclusive_scan.h>
 #include <cuda/std/__numeric/transform_inclusive_scan.h>
 #include <cuda/std/__numeric/transform_reduce.h>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 
 // standard mandated include
 #include <cuda/std/version>
diff --git a/libcudacxx/include/cuda/std/ranges b/libcudacxx/include/cuda/std/ranges
index bd68edc1cf..78839bbe02 100644
--- a/libcudacxx/include/cuda/std/ranges
+++ b/libcudacxx/include/cuda/std/ranges
@@ -25,6 +25,7 @@
 _CCCL_DIAG_PUSH
 _CCCL_DIAG_SUPPRESS_MSVC(4848)
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__ranges/access.h>
 #include <cuda/std/__ranges/concepts.h>
 #include <cuda/std/__ranges/dangling.h>
@@ -38,7 +39,6 @@ _CCCL_DIAG_SUPPRESS_MSVC(4848)
 #include <cuda/std/__ranges/subrange.h>
 #include <cuda/std/__ranges/view_interface.h>
 #include <cuda/std/__ranges/views.h>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 
 // standard-mandated includes
 #include <cuda/std/version>
diff --git a/libcudacxx/include/cuda/std/type_traits b/libcudacxx/include/cuda/std/type_traits
index 09729c7180..316b5f27e6 100644
--- a/libcudacxx/include/cuda/std/type_traits
+++ b/libcudacxx/include/cuda/std/type_traits
@@ -22,6 +22,7 @@
 #endif // no system header
 
 #include <cuda/std/__algorithm/iter_swap.h>
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__functional/identity.h>
 #include <cuda/std/__functional/invoke.h>
 #include <cuda/std/__fwd/hash.h>
@@ -160,7 +161,6 @@
 #include <cuda/std/__utility/swap.h>
 #include <cuda/std/cstddef>
 #include <cuda/std/cstdint>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/version>
 
 #endif // _CUDA_STD_TYPE_TRAITS
diff --git a/libcudacxx/include/cuda/std/utility b/libcudacxx/include/cuda/std/utility
index 9576806c5a..5e491ce5b4 100644
--- a/libcudacxx/include/cuda/std/utility
+++ b/libcudacxx/include/cuda/std/utility
@@ -31,6 +31,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__cccl/assert.h> // all public C++ headers provide the assertion handler
 #include <cuda/std/__functional/binary_function.h>
 #include <cuda/std/__functional/hash.h>
 #include <cuda/std/__functional/reference_wrapper.h>
@@ -57,7 +58,6 @@
 #include <cuda/std/__utility/swap.h>
 #include <cuda/std/__utility/to_underlying.h>
 #include <cuda/std/__utility/unreachable.h>
-#include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
 #include <cuda/std/limits>
 
 // standard-mandated includes
diff --git a/libcudacxx/test/internal_headers/CMakeLists.txt b/libcudacxx/test/internal_headers/CMakeLists.txt
index 5595f96ed1..c72300eef4 100644
--- a/libcudacxx/test/internal_headers/CMakeLists.txt
+++ b/libcudacxx/test/internal_headers/CMakeLists.txt
@@ -41,7 +41,7 @@ function(libcudacxx_create_internal_header_test header_name, headertest_src, fal
     $<$<COMPILE_LANGUAGE:CUDA>:${headertest_warning_levels_device}>
     $<$<COMPILE_LANGUAGE:CXX>:${headertest_warning_levels_host}>
     -DLIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE)
-  target_compile_definitions(headertest_${header_name} PRIVATE "-D_LIBCUDACXX_ENABLE_DEBUG_MODE")
+  target_compile_definitions(headertest_${header_name} PRIVATE "-D_CCCL_ENABLE_DEBUG_MODE")
   target_link_libraries(headertest_${header_name} CUDA::cudart)
 
   if(fallback)
diff --git a/libcudacxx/test/libcudacxx/CMakeLists.txt b/libcudacxx/test/libcudacxx/CMakeLists.txt
index 94b5afbc10..e628eb2266 100644
--- a/libcudacxx/test/libcudacxx/CMakeLists.txt
+++ b/libcudacxx/test/libcudacxx/CMakeLists.txt
@@ -42,7 +42,7 @@ endif()
 # enable exceptions and assertions in tests
 string(APPEND LIBCUDACXX_TEST_COMPILER_FLAGS
        " -DLIBCUDACXX_ENABLE_EXCEPTIONS"
-       " -D_LIBCUDACXX_ENABLE_DEBUG_MODE")
+       " -D_CCCL_ENABLE_DEBUG_MODE")
 
 if (NOT MSVC AND NOT ${CMAKE_CUDA_COMPILER_ID} STREQUAL "Clang")
   set(LIBCUDACXX_WARNING_LEVEL "--compiler-options=-Wall --compiler-options=-Wextra")
diff --git a/libcudacxx/test/libcudacxx/libcxx/iterators/bounded_iter/dereference.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/iterators/bounded_iter/dereference.pass.cpp
index f241a9a9de..aa65ee614e 100644
--- a/libcudacxx/test/libcudacxx/libcxx/iterators/bounded_iter/dereference.pass.cpp
+++ b/libcudacxx/test/libcudacxx/libcxx/iterators/bounded_iter/dereference.pass.cpp
@@ -14,7 +14,7 @@
 
 // REQUIRES: has-unix-headers
 // UNSUPPORTED: c++03
-// ADDITIONAL_COMPILE_FLAGS: -D_LIBCUDACXX_ENABLE_ASSERTIONS=1
+// ADDITIONAL_COMPILE_FLAGS: -D_CCCL_ENABLE_ASSERTIONS
 
 #include <cuda/std/iterator>
 
@@ -64,15 +64,15 @@ __host__ __device__ void test_death()
   cuda::std::__bounded_iter<Iter> const oob  = cuda::std::__make_bounded_iter(Iter(e), Iter(b), Iter(e));
 
   // operator*
-  TEST_LIBCUDACXX_ASSERT_FAILURE(*oob, "__bounded_iter::operator*: Attempt to dereference an out-of-range iterator");
+  TEST_CCCL_ASSERT_FAILURE(*oob, "__bounded_iter::operator*: Attempt to dereference an out-of-range iterator");
   // operator->
-  TEST_LIBCUDACXX_ASSERT_FAILURE(oob->x, "__bounded_iter::operator->: Attempt to dereference an out-of-range iterator");
+  TEST_CCCL_ASSERT_FAILURE(oob->x, "__bounded_iter::operator->: Attempt to dereference an out-of-range iterator");
   // operator[]
-  TEST_LIBCUDACXX_ASSERT_FAILURE(iter[-1], "__bounded_iter::operator[]: Attempt to index an iterator out-of-range");
-  TEST_LIBCUDACXX_ASSERT_FAILURE(iter[5], "__bounded_iter::operator[]: Attempt to index an iterator out-of-range");
-  TEST_LIBCUDACXX_ASSERT_FAILURE(oob[0], "__bounded_iter::operator[]: Attempt to index an iterator out-of-range");
-  TEST_LIBCUDACXX_ASSERT_FAILURE(oob[1], "__bounded_iter::operator[]: Attempt to index an iterator out-of-range");
-  TEST_LIBCUDACXX_ASSERT_FAILURE(oob[-6], "__bounded_iter::operator[]: Attempt to index an iterator out-of-range");
+  TEST_CCCL_ASSERT_FAILURE(iter[-1], "__bounded_iter::operator[]: Attempt to index an iterator out-of-range");
+  TEST_CCCL_ASSERT_FAILURE(iter[5], "__bounded_iter::operator[]: Attempt to index an iterator out-of-range");
+  TEST_CCCL_ASSERT_FAILURE(oob[0], "__bounded_iter::operator[]: Attempt to index an iterator out-of-range");
+  TEST_CCCL_ASSERT_FAILURE(oob[1], "__bounded_iter::operator[]: Attempt to index an iterator out-of-range");
+  TEST_CCCL_ASSERT_FAILURE(oob[-6], "__bounded_iter::operator[]: Attempt to index an iterator out-of-range");
 }
 
 int main(int, char**)
diff --git a/libcudacxx/test/public_headers/CMakeLists.txt b/libcudacxx/test/public_headers/CMakeLists.txt
index 9588d174bb..b6358553ef 100644
--- a/libcudacxx/test/public_headers/CMakeLists.txt
+++ b/libcudacxx/test/public_headers/CMakeLists.txt
@@ -46,7 +46,7 @@ function(libcudacxx_add_public_header_test header)
                          PRIVATE
                          ${headertest_warning_levels_device}
                          -DLIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE)
-  target_compile_definitions(headertest_${header_name} PRIVATE "-D_LIBCUDACXX_ENABLE_DEBUG_MODE")
+  target_compile_definitions(headertest_${header_name} PRIVATE "-D_CCCL_ENABLE_DEBUG_MODE")
 
   # Ensure that if this is an atomic header, we only include the right architectures
   string(REGEX MATCH "atomic|barrier|latch|semaphore|annotated_ptr|pipeline" match "${header}")
diff --git a/libcudacxx/test/public_headers_host_only/CMakeLists.txt b/libcudacxx/test/public_headers_host_only/CMakeLists.txt
index d970e34a4d..a9eceb63ad 100644
--- a/libcudacxx/test/public_headers_host_only/CMakeLists.txt
+++ b/libcudacxx/test/public_headers_host_only/CMakeLists.txt
@@ -28,7 +28,7 @@ function(libcudacxx_add_std_header_test header)
   add_library(headertest_std_${header_name} SHARED "${headertest_src}.cpp")
   target_include_directories(headertest_std_${header_name} PRIVATE "${libcudacxx_SOURCE_DIR}/include")
   target_compile_options(headertest_std_${header_name} PRIVATE ${headertest_warning_levels_host})
-  target_compile_definitions(headertest_std_${header_name} PRIVATE "-D_LIBCUDACXX_ENABLE_DEBUG_MODE")
+  target_compile_definitions(headertest_std_${header_name} PRIVATE "-D_CCCL_ENABLE_DEBUG_MODE")
 
   add_dependencies(libcudacxx.test.public_headers_host_only headertest_std_${header_name})
 endfunction()
diff --git a/libcudacxx/test/support/check_assertion.h b/libcudacxx/test/support/check_assertion.h
index 96c0286888..6def8f701e 100644
--- a/libcudacxx/test/support/check_assertion.h
+++ b/libcudacxx/test/support/check_assertion.h
@@ -391,12 +391,12 @@ inline bool ExpectDeath(const char* stmt, Func&& func)
     },                                     \
     Matcher)))
 
-#define TEST_LIBCUDACXX_ASSERT_FAILURE(expr, message) \
-  assert((ExpectDeath(                                \
-    #expr,                                            \
-    [&]() {                                           \
-      (void) (expr);                                  \
-    },                                                \
+#define TEST_CCCL_ASSERT_FAILURE(expr, message) \
+  assert((ExpectDeath(                          \
+    #expr,                                      \
+    [&]() {                                     \
+      (void) (expr);                            \
+    },                                          \
     AssertionInfoMatcher(message))))
 
 #endif // TEST_SUPPORT_CHECK_ASSERTION_H
diff --git a/libcudacxx/test/support/container_debug_tests.h b/libcudacxx/test/support/container_debug_tests.h
index 6b847190ee..d7d278563b 100644
--- a/libcudacxx/test/support/container_debug_tests.h
+++ b/libcudacxx/test/support/container_debug_tests.h
@@ -15,7 +15,7 @@
 #  error This header may only be used for libc++ tests
 #endif
 
-#ifndef _LIBCUDACXX_ENABLE_DEBUG_MODE
+#ifndef _CCCL_ENABLE_DEBUG_MODE
 #  error The library must be built with the debug mode enabled in order to use this header
 #endif
 
diff --git a/libcudacxx/test/utils/libcudacxx/test/format.py b/libcudacxx/test/utils/libcudacxx/test/format.py
index 3a58447989..efd55f6751 100644
--- a/libcudacxx/test/utils/libcudacxx/test/format.py
+++ b/libcudacxx/test/utils/libcudacxx/test/format.py
@@ -141,13 +141,13 @@ def _execute(self, test, lit_config):
             test_cxx.compile_flags += [('-D%s' % mdef.strip()) for
                                        mdef in extra_modules_defines]
             test_cxx.addWarningFlagIfSupported('-Wno-macro-redefined')
-            # FIXME: libc++ debug tests #define _LIBCUDACXX_ASSERT to override it
+            # FIXME: libc++ debug tests #define _CCCL_ASSERT to override it
             # If we see this we need to build the test against uniquely built
             # modules.
             if is_libcxx_test:
                 with open(test.getSourcePath(), 'rb') as f:
                     contents = f.read()
-                if b'#define _LIBCUDACXX_ASSERT' in contents:
+                if b'#define _CCCL_ASSERT' in contents:
                     test_cxx.useModules(False)
 
         if is_objcxx_test: