Move to a CCCL wide assert definition

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);
   }
 

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;
   }

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;
   }

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;
   }

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>

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);
 }

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));
 }

libcudacxx/include/cuda/std/__algorithm/clamp.h

@@ -21,15 +21,15 @@
 #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
 
 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;
 }
 

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

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;
 

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;

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

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;                                   \
   }