From c838d50991fc96da38e6558a56781bb15b5d0f68 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Mon, 16 Sep 2024 13:27:52 +0200 Subject: [PATCH 01/20] Work towards enabling assertions --- .../device/dispatch/dispatch_transform.cuh | 12 +- .../cuda/experimental/__async/stop_token.cuh | 4 +- .../cuda/experimental/__async/variant.cuh | 8 +- .../__memory_resource/any_resource.cuh | 6 +- .../async_memory_resource.cuh | 6 +- libcudacxx/include/cuda/__cmath/ceil_div.h | 7 +- .../device_memory_resource.h | 3 +- .../managed_memory_resource.h | 5 +- .../pinned_memory_resource.h | 5 +- libcudacxx/include/cuda/annotated_ptr | 10 +- libcudacxx/include/cuda/barrier | 12 +- .../include/cuda/std/__algorithm/clamp.h | 4 +- .../cuda/std/__algorithm/comp_ref_type.h | 6 +- .../include/cuda/std/__algorithm/pop_heap.h | 4 +- .../include/cuda/std/__algorithm/sift_down.h | 4 +- .../cuda/std/__algorithm/unwrap_iter.h | 4 - libcudacxx/include/cuda/std/__cccl/assert.h | 124 ++++++++ .../include/cuda/std/__cuda/api_wrapper.h | 2 +- libcudacxx/include/cuda/std/__cuda/barrier.h | 39 ++- .../include/cuda/std/__expected/expected.h | 32 +- .../cuda/std/__expected/expected_base.h | 2 +- .../include/cuda/std/__functional/function.h | 19 +- .../include/cuda/std/__iterator/advance.h | 13 +- .../cuda/std/__iterator/bounded_iter.h | 15 +- libcudacxx/include/cuda/std/__iterator/next.h | 6 +- libcudacxx/include/cuda/std/__iterator/prev.h | 6 +- .../include/cuda/std/__iterator/wrap_iter.h | 88 +----- .../include/cuda/std/__mdspan/layout_left.h | 2 +- .../include/cuda/std/__mdspan/layout_right.h | 2 +- libcudacxx/include/cuda/std/__mdspan/macros.h | 2 +- .../include/cuda/std/__memory/construct_at.h | 18 +- libcudacxx/include/cuda/std/__memory_ | 2 +- libcudacxx/include/cuda/std/__new_ | 2 +- .../include/cuda/std/__numeric/gcd_lcm.h | 2 +- .../include/cuda/std/__ranges/subrange.h | 6 +- .../cuda/std/__ranges/view_interface.h | 10 +- .../cuda/std/__utility/exception_guard.h | 4 +- libcudacxx/include/cuda/std/bit | 2 +- libcudacxx/include/cuda/std/bitset | 2 +- libcudacxx/include/cuda/std/concepts | 2 +- .../include/cuda/std/detail/__annotated_ptr | 4 +- .../cuda/std/detail/libcxx/include/__assert | 66 ---- .../cuda/std/detail/libcxx/include/__config | 9 - .../cuda/std/detail/libcxx/include/__debug | 288 ------------------ .../cuda/std/detail/libcxx/include/__string | 15 +- .../detail/libcxx/include/__threading_support | 2 +- .../std/detail/libcxx/include/__verbose_abort | 59 ---- .../cuda/std/detail/libcxx/include/algorithm | 5 +- .../cuda/std/detail/libcxx/include/array | 18 +- .../cuda/std/detail/libcxx/include/barrier | 23 +- .../cuda/std/detail/libcxx/include/chrono | 2 +- .../cuda/std/detail/libcxx/include/complex | 2 +- .../cuda/std/detail/libcxx/include/cstddef | 2 +- .../cuda/std/detail/libcxx/include/iosfwd | 2 +- .../cuda/std/detail/libcxx/include/latch | 7 +- .../cuda/std/detail/libcxx/include/limits | 2 +- .../cuda/std/detail/libcxx/include/optional | 17 +- .../cuda/std/detail/libcxx/include/ratio | 2 +- .../cuda/std/detail/libcxx/include/semaphore | 18 +- .../cuda/std/detail/libcxx/include/span | 69 +++-- .../cuda/std/detail/libcxx/include/stdexcept | 2 +- .../cuda/std/detail/libcxx/include/tuple | 2 +- .../cuda/std/detail/libcxx/include/variant | 2 +- libcudacxx/include/cuda/std/expected | 2 +- libcudacxx/include/cuda/std/functional | 2 +- libcudacxx/include/cuda/std/inplace_vector | 2 +- libcudacxx/include/cuda/std/iterator | 2 +- libcudacxx/include/cuda/std/mdspan | 2 +- libcudacxx/include/cuda/std/numeric | 2 +- libcudacxx/include/cuda/std/ranges | 2 +- libcudacxx/include/cuda/std/type_traits | 2 +- libcudacxx/include/cuda/std/utility | 3 +- .../test/internal_headers/CMakeLists.txt | 1 + libcudacxx/test/libcudacxx/CMakeLists.txt | 3 + .../bounded_iter/dereference.pass.cpp | 16 +- .../libcxx/iterators/unwrap_iter.pass.cpp | 2 +- libcudacxx/test/public_headers/CMakeLists.txt | 1 + .../public_headers_host_only/CMakeLists.txt | 1 + libcudacxx/test/support/check_assertion.h | 12 +- .../test/support/container_debug_tests.h | 2 +- .../test/utils/libcudacxx/test/format.py | 4 +- 81 files changed, 395 insertions(+), 786 deletions(-) create mode 100644 libcudacxx/include/cuda/std/__cccl/assert.h delete mode 100644 libcudacxx/include/cuda/std/detail/libcxx/include/__assert delete mode 100644 libcudacxx/include/cuda/std/detail/libcxx/include/__debug delete mode 100644 libcudacxx/include/cuda/std/detail/libcxx/include/__verbose_abort diff --git a/cub/cub/device/dispatch/dispatch_transform.cuh b/cub/cub/device/dispatch/dispatch_transform.cuh index 8fb596da07..0e6231d74b 100644 --- a/cub/cub/device/dispatch/dispatch_transform.cuh +++ b/cub/cub/device/dispatch/dispatch_transform.cuh @@ -168,7 +168,7 @@ template _CCCL_HOST_DEVICE _CCCL_FORCEINLINE constexpr auto round_up_to_po2_multiple(Integral x, Integral mult) -> Integral { #if _CCCL_STD_VER > 2011 - _LIBCUDACXX_ASSERT(::cuda::std::has_single_bit(static_cast<::cuda::std::__make_unsigned_t>(mult)), ""); + _CCCL_ASSERT(::cuda::std::has_single_bit(static_cast<::cuda::std::__make_unsigned_t>(mult)), ""); #endif // _CCCL_STD_VER > 2011 return (x + mult - 1) & ~(mult - 1); } @@ -177,7 +177,7 @@ template _CCCL_HOST_DEVICE _CCCL_FORCEINLINE const char* round_down_ptr(const T* ptr, unsigned alignment) { #if _CCCL_STD_VER > 2011 - _LIBCUDACXX_ASSERT(::cuda::std::has_single_bit(alignment), ""); + _CCCL_ASSERT(::cuda::std::has_single_bit(alignment), ""); #endif // _CCCL_STD_VER > 2011 return reinterpret_cast( reinterpret_cast<::cuda::std::uintptr_t>(ptr) & ~::cuda::std::uintptr_t{alignment - 1}); @@ -278,8 +278,8 @@ _CCCL_DEVICE void bulk_copy_tile( const char* src = aligned_ptr.ptr + global_offset * sizeof(T); char* dst = smem + smem_offset; - _LIBCUDACXX_ASSERT(reinterpret_cast(src) % bulk_copy_alignment == 0, ""); - _LIBCUDACXX_ASSERT(reinterpret_cast(dst) % bulk_copy_alignment == 0, ""); + _CCCL_ASSERT(reinterpret_cast(src) % bulk_copy_alignment == 0, ""); + _CCCL_ASSERT(reinterpret_cast(dst) % bulk_copy_alignment == 0, ""); // TODO(bgruber): we could precompute bytes_to_copy on the host const int bytes_to_copy = round_up_to_po2_multiple( @@ -303,8 +303,8 @@ _CCCL_DEVICE void bulk_copy_tile_fallback( { const T* src = aligned_ptr.ptr_to_elements() + global_offset; T* dst = reinterpret_cast(smem + smem_offset + aligned_ptr.head_padding); - _LIBCUDACXX_ASSERT(reinterpret_cast(src) % alignof(T) == 0, ""); - _LIBCUDACXX_ASSERT(reinterpret_cast(dst) % alignof(T) == 0, ""); + _CCCL_ASSERT(reinterpret_cast(src) % alignof(T) == 0, ""); + _CCCL_ASSERT(reinterpret_cast(dst) % alignof(T) == 0, ""); const int bytes_to_copy = static_cast(sizeof(T)) * tile_size; cooperative_groups::memcpy_async(cooperative_groups::this_thread_block(), dst, src, bytes_to_copy); diff --git a/cudax/include/cuda/experimental/__async/stop_token.cuh b/cudax/include/cuda/experimental/__async/stop_token.cuh index 2a3e93a8d7..237585af21 100644 --- a/cudax/include/cuda/experimental/__async/stop_token.cuh +++ b/cudax/include/cuda/experimental/__async/stop_token.cuh @@ -303,8 +303,8 @@ _CCCL_HOST_DEVICE inline void __inplace_stop_callback_base::__register_callback( _CCCL_HOST_DEVICE inline inplace_stop_source::~inplace_stop_source() { - _LIBCUDACXX_ASSERT((__state_.load(_CUDA_VSTD::memory_order_relaxed) & __locked_flag) == 0, ""); - _LIBCUDACXX_ASSERT(__callbacks_ == nullptr, ""); + _CCCL_ASSERT((__state_.load(_CUDA_VSTD::memory_order_relaxed) & __locked_flag) == 0, ""); + _CCCL_ASSERT(__callbacks_ == nullptr, ""); } _CCCL_HOST_DEVICE inline auto inplace_stop_source::request_stop() noexcept -> bool diff --git a/cudax/include/cuda/experimental/__async/variant.cuh b/cudax/include/cuda/experimental/__async/variant.cuh index 2762566c41..fbca1df166 100644 --- a/cudax/include/cuda/experimental/__async/variant.cuh +++ b/cudax/include/cuda/experimental/__async/variant.cuh @@ -139,7 +139,7 @@ public: { // make this local in case destroying the sub-object destroys *this const auto index = __self.__index_; - _LIBCUDACXX_ASSERT(index != __npos, ""); + _CCCL_ASSERT(index != __npos, ""); ((_Idx == index ? static_cast<_Fn&&>(__fn)(static_cast<_As&&>(__as)..., static_cast<_Self&&>(__self).template __get<_Idx>()) : void()), @@ -149,21 +149,21 @@ public: template _CCCL_HOST_DEVICE __at<_Ny>&& __get() && noexcept { - _LIBCUDACXX_ASSERT(_Ny == __index_, ""); + _CCCL_ASSERT(_Ny == __index_, ""); return static_cast<__at<_Ny>&&>(*static_cast<__at<_Ny>*>(__ptr())); } template _CCCL_HOST_DEVICE __at<_Ny>& __get() & noexcept { - _LIBCUDACXX_ASSERT(_Ny == __index_, ""); + _CCCL_ASSERT(_Ny == __index_, ""); return *static_cast<__at<_Ny>*>(__ptr()); } template _CCCL_HOST_DEVICE const __at<_Ny>& __get() const& noexcept { - _LIBCUDACXX_ASSERT(_Ny == __index_, ""); + _CCCL_ASSERT(_Ny == __index_, ""); return *static_cast*>(__ptr()); } }; diff --git a/cudax/include/cuda/experimental/__memory_resource/any_resource.cuh b/cudax/include/cuda/experimental/__memory_resource/any_resource.cuh index e25ebac35f..b300e39d94 100644 --- a/cudax/include/cuda/experimental/__memory_resource/any_resource.cuh +++ b/cudax/include/cuda/experimental/__memory_resource/any_resource.cuh @@ -151,7 +151,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); } @@ -162,7 +162,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); } @@ -191,7 +191,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/cudax/include/cuda/experimental/__memory_resource/async_memory_resource.cuh b/cudax/include/cuda/experimental/__memory_resource/async_memory_resource.cuh index b9fd038dd9..ab9451bb6d 100644 --- a/cudax/include/cuda/experimental/__memory_resource/async_memory_resource.cuh +++ b/cudax/include/cuda/experimental/__memory_resource/async_memory_resource.cuh @@ -164,8 +164,7 @@ public: //! properly synchronize all relevant streams before calling `deallocate`. void deallocate(void* __ptr, const size_t, const size_t __alignment = _CUDA_VMR::default_cuda_malloc_alignment) { - _LIBCUDACXX_ASSERT(__is_valid_alignment(__alignment), - "Invalid alignment passed to async_memory_resource::deallocate."); + _CCCL_ASSERT(__is_valid_alignment(__alignment), "Invalid alignment passed to async_memory_resource::deallocate."); _CCCL_ASSERT_CUDA_API( ::cudaFreeAsync, "async_memory_resource::deallocate failed", __ptr, __async_memory_resource_sync_stream().get()); __async_memory_resource_sync_stream().wait(); @@ -221,8 +220,7 @@ public: void deallocate_async(void* __ptr, const size_t __bytes, const size_t __alignment, const ::cuda::stream_ref __stream) { // We need to ensure that the provided alignment matches the minimal provided alignment - _LIBCUDACXX_ASSERT(__is_valid_alignment(__alignment), - "Invalid alignment passed to async_memory_resource::deallocate."); + _CCCL_ASSERT(__is_valid_alignment(__alignment), "Invalid alignment passed to async_memory_resource::deallocate."); deallocate_async(__ptr, __bytes, __stream); (void) __alignment; } diff --git a/libcudacxx/include/cuda/__cmath/ceil_div.h b/libcudacxx/include/cuda/__cmath/ceil_div.h index d6ee4f74f3..5687b3370a 100644 --- a/libcudacxx/include/cuda/__cmath/ceil_div.h +++ b/libcudacxx/include/cuda/__cmath/ceil_div.h @@ -29,7 +29,6 @@ #include #include #include -#include _LIBCUDACXX_BEGIN_NAMESPACE_CUDA @@ -44,7 +43,7 @@ template = 0> _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 _Tp ceil_div(const _Tp __a, const _Up __b) noexcept { - _LIBCUDACXX_DEBUG_ASSERT(__b > _Tp(0), "cuda::ceil_div: b must be positive"); + _CCCL_ASSERT(__b > _Tp(0), "cuda::ceil_div: b must be positive"); using _UCommon = _CUDA_VSTD::__make_unsigned_t<_CUDA_VSTD::__common_type_t<_Tp, _Up>>; const auto __res = static_cast<_UCommon>(__a) / static_cast<_UCommon>(__b); return static_cast<_Tp>(__res + (__res * static_cast<_UCommon>(__b) != static_cast<_UCommon>(__a))); @@ -61,8 +60,8 @@ template = 0> _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 _Tp ceil_div(const _Tp __a, const _Up __b) noexcept { - _LIBCUDACXX_DEBUG_ASSERT(__a >= _Tp(0), "cuda::ceil_div: a must be non negative"); - _LIBCUDACXX_DEBUG_ASSERT(__b > _Tp(0), "cuda::ceil_div: b must be positive"); + _CCCL_ASSERT(__a >= _Tp(0), "cuda::ceil_div: a must be non negative"); + _CCCL_ASSERT(__b > _Tp(0), "cuda::ceil_div: b must be positive"); using _UCommon = _CUDA_VSTD::__make_unsigned_t<_CUDA_VSTD::__common_type_t<_Tp, _Up>>; // Due to the precondition `__a >= 0` we can safely cast to unsigned without danger of overflowing return static_cast<_Tp>((static_cast<_UCommon>(__a) + static_cast<_UCommon>(__b) - 1) / static_cast<_UCommon>(__b)); 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 a04be90c1f..b501164ea9 100644 --- a/libcudacxx/include/cuda/annotated_ptr +++ b/libcudacxx/include/cuda/annotated_ptr @@ -144,6 +144,7 @@ #include #include +#include // all public C++ headers provide the assertion handler #include #include @@ -359,9 +360,10 @@ public: _CCCL_HOST_DEVICE explicit annotated_ptr(pointer __p) : __repr(__p) { - NV_IF_TARGET( - NV_IS_DEVICE, - (_LIBCUDACXX_DEBUG_ASSERT((std::is_same<_Property, shared>::value && __isShared(__p) || __isGlobal(__p)), "");)) + NV_IF_TARGET(NV_IS_DEVICE, + (_CCCL_ASSERT((std::is_same<_Property, access_property::shared>::value && __isShared((void*) __p)) + || __isGlobal((void*) __p), + "");)) } template @@ -378,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_DEBUG_ASSERT((__isGlobal(__p) == true), "");)) + NV_IF_TARGET(NV_IS_DEVICE, (_CCCL_ASSERT((__isGlobal((void*) __p) == true), "");)) } template diff --git a/libcudacxx/include/cuda/barrier b/libcudacxx/include/cuda/barrier index 99117dde90..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_DEBUG_ASSERT(__size % 16 == 0, "Size must be multiple of 16."); - _LIBCUDACXX_DEBUG_ASSERT(__isShared(__dest), "Destination must be shared memory address."); - _LIBCUDACXX_DEBUG_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_DEBUG_ASSERT(__size % 16 == 0, "Size must be multiple of 16."); - _LIBCUDACXX_DEBUG_ASSERT(__isGlobal(__dest), "Destination must be global memory address."); - _LIBCUDACXX_DEBUG_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/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 -#include +#include _LIBCUDACXX_BEGIN_NAMESPACE_STD @@ -29,7 +29,7 @@ template _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 75575972a3..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 // all public C++ headers provide the assertion handler #include -#include _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_DEBUG_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 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 #include #include +#include #include #include #include #include -#include _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 +#include #include #include -#include _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/__algorithm/unwrap_iter.h b/libcudacxx/include/cuda/std/__algorithm/unwrap_iter.h index a97b443345..490241019d 100644 --- a/libcudacxx/include/cuda/std/__algorithm/unwrap_iter.h +++ b/libcudacxx/include/cuda/std/__algorithm/unwrap_iter.h @@ -52,8 +52,6 @@ struct __unwrap_iter_impl } }; -#ifndef _LIBCUDACXX_ENABLE_DEBUG_MODE - // It's a contiguous iterator, so we can use a raw pointer instead template struct __unwrap_iter_impl<_Iter, true> @@ -71,8 +69,6 @@ struct __unwrap_iter_impl<_Iter, true> } }; -#endif // !_LIBCUDACXX_ENABLE_DEBUG_MODE - template , __enable_if_t::value, int> = 0> diff --git a/libcudacxx/include/cuda/std/__cccl/assert.h b/libcudacxx/include/cuda/std/__cccl/assert.h new file mode 100644 index 0000000000..b774225a61 --- /dev/null +++ b/libcudacxx/include/cuda/std/__cccl/assert.h @@ -0,0 +1,124 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef __CCCL_ASSERT_H +#define __CCCL_ASSERT_H + +#include +#include + +#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 + +#if defined(_DEBUG) || defined(DEBUG) +# ifndef _CCCL_ENABLE_DEBUG_MODE +# define _CCCL_ENABLE_DEBUG_MODE +# endif // !_CCCL_ENABLE_DEBUG_MODE +#endif // _DEBUG || DEBUG + +// Automatically enable assertions when debug mode is enabled +#ifdef _CCCL_ENABLE_DEBUG_MODE +# ifndef CCCL_ENABLE_ASSERTIONS +# define CCCL_ENABLE_ASSERTIONS +# endif // !CCCL_ENABLE_ASSERTIONS +#endif // _CCCL_ENABLE_DEBUG_MODE + +//! Ensure that we switch on host assertions when all assertions are enabled +#ifndef CCCL_ENABLE_HOST_ASSERTIONS +# ifdef CCCL_ENABLE_ASSERTIONS +# define CCCL_ENABLE_HOST_ASSERTIONS +# endif // CCCL_ENABLE_ASSERTIONS +#endif // !CCCL_ENABLE_HOST_ASSERTIONS + +//! Ensure that we switch on device assertions when all assertions are enabled +#ifndef CCCL_ENABLE_DEVICE_ASSERTIONS +# ifdef CCCL_ENABLE_ASSERTIONS +# define CCCL_ENABLE_DEVICE_ASSERTIONS +# endif // CCCL_ENABLE_ASSERTIONS +#endif // !CCCL_ENABLE_DEVICE_ASSERTIONS + +//! Use internal nvcc implementation on device or the host library for clang-cuda +#ifdef __CUDA_ARCH__ +# ifdef CCCL_ENABLE_DEVICE_ASSERTIONS +# if defined(_CCCL_CUDA_COMPILER_NVCC) //! Use __assert_fail to implement device side asserts +# include +# define _CCCL_ASSERT_IMPL(expression, message) \ + __builtin_expect(static_cast(expression), 1) \ + ? (void) 0 \ + : __assert_fail(message, __FILE__, __LINE__, __func__) +# else // ^^^ _CCCL_CUDA_COMPILER_NVCC ^^^ / vvv !_CCCL_CUDA_COMPILER_NVCC vvv +# if __has_include(<__assert>) // libc++ uses _LIBCPP_ASSERT from <__assert> +# include <__assert> +# define _CCCL_ASSERT_IMPL(expression, message) _LIBCPP_ASSERT(expression, message) +# else // libstdc++ uses __glibcxx_assert from +# if defined(_GLIBCXX_RELEASE) && _GLIBCXX_RELEASE >= 12 +// libstdc++ does not fully qualify its use of `__is_constant_evaluated` +// It was introduced in the assert handling in 5e8a30d +// libstdc++ : Redefine __glibcxx_assert to work in C++ 23 constexpr +_LIBCUDACXX_BEGIN_NAMESPACE_STD +using ::std::__is_constant_evaluated; +_LIBCUDACXX_END_NAMESPACE_STD +# endif // _GLIBCXX_RELEASE >= 12 +# define _CCCL_ASSERT_IMPL(expression, message) __glibcxx_assert(expression) +# endif // libstdc++ +# endif // !_CCCL_CUDA_COMPILER_NVCC +# else // ^^^ CCCL_ENABLE_DEVICE_ASSERTIONS ^^^ / vvv !CCCL_ENABLE_DEVICE_ASSERTIONS vvv +# define _CCCL_ASSERT_IMPL +# endif // !CCCL_ENABLE_DEVICE_ASSERTIONS +#endif // __CUDA_ARCH__ + +//! Use the different standard library implementations to implement host side asserts +#ifndef __CUDA_ARCH__ +# ifdef CCCL_ENABLE_HOST_ASSERTIONS +# if __has_include() // MSVC uses _STL_VERIFY from +# include +# define _CCCL_ASSERT_IMPL(expression, message) _STL_VERIFY(expression, message) +# elif __has_include(<__assert>) // libc++ uses _LIBCPP_ASSERT from <__assert> +# include <__assert> +# define _CCCL_ASSERT_IMPL(expression, message) _LIBCPP_ASSERT(expression, message) +# else // libstdc++ uses __glibcxx_assert from +# if defined(_GLIBCXX_RELEASE) && _GLIBCXX_RELEASE >= 12 +// libstdc++ does not fully qualify its use of `__is_constant_evaluated` +// It was introduced in the assert handling in 5e8a30d +// libstdc++ : Redefine __glibcxx_assert to work in C++ 23 constexpr +_LIBCUDACXX_BEGIN_NAMESPACE_STD +using ::std::__is_constant_evaluated; +_LIBCUDACXX_END_NAMESPACE_STD +# endif // _GLIBCXX_RELEASE >= 12 +# define _CCCL_ASSERT_IMPL(expression, message) __glibcxx_assert(expression) +# endif // libstdc++ +# endif // !CCCL_ENABLE_HOST_ASSERTIONS +#endif // !__CUDA_ARCH__ + +//! _CCCL_VERIFY is an unconditionally enabled assertion that is reserved for the most important checks +//! Some compilers warn about `is_constant_evaluated()` in a non constexpr function, so silence that right away +#if defined(_CCCL_COMPILER_ICC) +# define _CCCL_VERIFY(expression, message) \ + _CCCL_DIAG_PUSH _CCCL_DIAG_SUPPRESS_ICC(4190 3060) _CCCL_ASSERT_IMPL(expression, message) _CCCL_DIAG_POP +#else // ^^^ _CCCL_COMPILER_ICC ^^^ / vvv !_CCCL_COMPILER_ICC vvv +# define _CCCL_VERIFY(expression, message) _CCCL_ASSERT_IMPL(expression, message) +#endif // !_CCCL_COMPILER_ICC + +#if defined(CCCL_ENABLE_HOST_ASSERTIONS) || defined(CCCL_ENABLE_DEVICE_ASSERTIONS) +# define _CCCL_ASSERT(expression, message) _CCCL_VERIFY(expression, message) +#else // ^^^ CCCL_ENABLE_HOST_ASSERTIONS || CCCL_ENABLE_DEVICE_ASSERTIONS ^^^ / vvv No assertions vvv +# define _CCCL_ASSERT(expression, message) ((void) 0) +#endif // No 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 97ad47bfb0..1074cd8acc 100644 --- a/libcudacxx/include/cuda/std/__cuda/barrier.h +++ b/libcudacxx/include/cuda/std/__cuda/barrier.h @@ -26,6 +26,7 @@ #endif // no system header #include +#include // all public C++ headers provide the assertion handler #include // _CUDA_VSTD::void_t #if defined(_CCCL_CUDA_COMPILER) @@ -96,13 +97,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_DEBUG_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_DEBUG_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); } }; @@ -195,7 +196,7 @@ class barrier : public __blo _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI arrival_token arrive(_CUDA_VSTD::ptrdiff_t __update = 1) { - _LIBCUDACXX_DEBUG_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, @@ -548,13 +549,12 @@ _CCCL_NODISCARD _CCCL_DEVICE inline barrier::arrival_token b _CUDA_VSTD::ptrdiff_t __arrive_count_update, _CUDA_VSTD::ptrdiff_t __transaction_count_update) { - _LIBCUDACXX_DEBUG_ASSERT(__isShared(barrier_native_handle(__b)), "Barrier must be located in local shared memory."); - _LIBCUDACXX_DEBUG_ASSERT(1 <= __arrive_count_update, "Arrival count update must be at least one."); - _LIBCUDACXX_DEBUG_ASSERT(__arrive_count_update <= (1 << 20) - 1, "Arrival count update cannot exceed 2^20 - 1."); - _LIBCUDACXX_DEBUG_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_DEBUG_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::arrival_token __token = {}; // On architectures pre-sm90, arrive_tx is not supported. @@ -598,11 +598,10 @@ extern "C" _CCCL_DEVICE void __cuda_ptx_barrier_expect_tx_is_not_supported_befor _CCCL_DEVICE inline void barrier_expect_tx(barrier& __b, _CUDA_VSTD::ptrdiff_t __transaction_count_update) { - _LIBCUDACXX_DEBUG_ASSERT(__isShared(barrier_native_handle(__b)), "Barrier must be located in local shared memory."); - _LIBCUDACXX_DEBUG_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_DEBUG_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 +639,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_DEBUG_ASSERT(__isShared(barrier_native_handle(__b)), "Barrier must be located in local shared memory."); - _LIBCUDACXX_DEBUG_ASSERT(__isShared(__dest), "dest must point to shared memory."); - _LIBCUDACXX_DEBUG_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 +1103,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_DEBUG_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 +1177,8 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI __completion_mechanism __dispatch_memc _CUDA_VSTD::size_t __size, _CUDA_VSTD::uint32_t __allowed_completions) { - _LIBCUDACXX_DEBUG_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 2e229bda4c..587e2c0d08 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 #include #include #include @@ -60,7 +61,6 @@ #include #include #include -#include #include #if _CCCL_STD_VER > 2011 @@ -487,37 +487,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_); } @@ -579,25 +579,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_); } @@ -1435,7 +1435,7 @@ class expected : private __expected_move_assign _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& @@ -1460,25 +1460,25 @@ class expected : private __expected_move_assign _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 #include #include #include @@ -49,7 +50,6 @@ #include #include #include -#include #if _CCCL_STD_VER > 2011 diff --git a/libcudacxx/include/cuda/std/__functional/function.h b/libcudacxx/include/cuda/std/__functional/function.h index 719bdcbd23..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 #include #include #include @@ -47,8 +48,6 @@ #include #include #include -#include -#include #include #ifndef __cuda_std__ @@ -939,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; } @@ -964,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 c8061d422a..bd07995e2d 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 #include #include #include @@ -28,7 +29,6 @@ #include #include #include -#include _LIBCUDACXX_BEGIN_NAMESPACE_STD @@ -77,8 +77,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()); } @@ -128,8 +128,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` must be true."); + _CCCL_ASSERT(__n >= 0 || bidirectional_iterator<_Ip>, "If `n < 0`, then `bidirectional_iterator` must be true."); // If `I` models `random_access_iterator`, equivalent to `i += n`. if constexpr (random_access_iterator<_Ip>) @@ -189,8 +188,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 && same_as` must be true."); + _CCCL_ASSERT((__n >= 0) || (bidirectional_iterator<_Ip> && same_as<_Ip, _Sp>), + "If `n < 0`, then `bidirectional_iterator && same_as` must be true."); // If `S` and `I` model `sized_sentinel_for`: 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 #include #include #include #include #include -#include _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 @@ -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 #include #include #include #include #include -#include _LIBCUDACXX_BEGIN_NAMESPACE_STD @@ -34,8 +34,8 @@ template _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 #include #include #include #include #include -#include _LIBCUDACXX_BEGIN_NAMESPACE_STD @@ -34,8 +34,8 @@ template _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/__iterator/wrap_iter.h b/libcudacxx/include/cuda/std/__iterator/wrap_iter.h index fa9c6358d7..0760192de8 100644 --- a/libcudacxx/include/cuda/std/__iterator/wrap_iter.h +++ b/libcudacxx/include/cuda/std/__iterator/wrap_iter.h @@ -27,15 +27,6 @@ #include #include #include -#include - -#ifndef _LIBCUDACXX_CONSTEXPR_IF_NODEBUG -# if defined(_LIBCUDACXX_DEBUG) -# define _LIBCUDACXX_CONSTEXPR_IF_NODEBUG -# else -# define _LIBCUDACXX_CONSTEXPR_IF_NODEBUG constexpr -# endif -#endif _LIBCUDACXX_BEGIN_NAMESPACE_STD @@ -59,67 +50,23 @@ class __wrap_iter public: _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 __wrap_iter() noexcept : __i_() - { - _CUDA_VSTD::__debug_db_insert_i(this); - } + {} template _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 __wrap_iter(const __wrap_iter<_Up>& __u, typename enable_if::value>::type* = nullptr) noexcept : __i_(__u.base()) - { -#ifdef _LIBCUDACXX_ENABLE_DEBUG_MODE - if (!__libcpp_is_constant_evaluated()) - { - __get_db()->__iterator_copy(this, _CUDA_VSTD::addressof(__u)); - } -#endif - } -#ifdef _LIBCUDACXX_ENABLE_DEBUG_MODE - _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 __wrap_iter(const __wrap_iter& __x) - : __i_(__x.base()) - { - if (!__libcpp_is_constant_evaluated()) - { - __get_db()->__iterator_copy(this, _CUDA_VSTD::addressof(__x)); - } - } - _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 __wrap_iter& operator=(const __wrap_iter& __x) - { - if (this != _CUDA_VSTD::addressof(__x)) - { - if (!__libcpp_is_constant_evaluated()) - { - __get_db()->__iterator_copy(this, _CUDA_VSTD::addressof(__x)); - } - __i_ = __x.__i_; - } - return *this; - } - _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX20 ~__wrap_iter() - { - if (!__libcpp_is_constant_evaluated()) - { - __get_db()->__erase_i(this); - } - } -#endif + {} _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 reference operator*() const noexcept { - _LIBCUDACXX_DEBUG_ASSERT(__get_const_db()->__dereferenceable(this), - "Attempted to dereference a non-dereferenceable iterator"); return *__i_; } _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 pointer operator->() const noexcept { - _LIBCUDACXX_DEBUG_ASSERT(__get_const_db()->__dereferenceable(this), - "Attempted to dereference a non-dereferenceable iterator"); return _CUDA_VSTD::__to_address(__i_); } _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 __wrap_iter& operator++() noexcept { - _LIBCUDACXX_DEBUG_ASSERT(__get_const_db()->__dereferenceable(this), - "Attempted to increment a non-incrementable iterator"); ++__i_; return *this; } @@ -132,8 +79,6 @@ class __wrap_iter _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 __wrap_iter& operator--() noexcept { - _LIBCUDACXX_DEBUG_ASSERT(__get_const_db()->__decrementable(this), - "Attempted to decrement a non-decrementable iterator"); --__i_; return *this; } @@ -151,8 +96,6 @@ class __wrap_iter } _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 __wrap_iter& operator+=(difference_type __n) noexcept { - _LIBCUDACXX_DEBUG_ASSERT(__get_const_db()->__addable(this, __n), - "Attempted to add/subtract an iterator outside its valid range"); __i_ += __n; return *this; } @@ -167,8 +110,6 @@ class __wrap_iter } _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 reference operator[](difference_type __n) const noexcept { - _LIBCUDACXX_DEBUG_ASSERT(__get_const_db()->__subscriptable(this, __n), - "Attempted to subscript an iterator outside its valid range"); return __i_[__n]; } @@ -177,21 +118,10 @@ class __wrap_iter return __i_; } -// private: -#if _LIBCUDACXX_DEBUG_LEVEL >= 2 - _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_CONSTEXPR_IF_NODEBUG __wrap_iter(const void* __p, iterator_type __x) - : __i_(__x) - { - if (!__libcpp_is_constant_evaluated()) - { - __get_db()->__insert_ic(this, __p); - } - } -#else - _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_CONSTEXPR_IF_NODEBUG __wrap_iter(iterator_type __x) noexcept +private: + _LIBCUDACXX_HIDE_FROM_ABI constexpr __wrap_iter(iterator_type __x) noexcept : __i_(__x) {} -#endif template friend class __wrap_iter; @@ -221,9 +151,6 @@ template _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 bool operator<(const __wrap_iter<_Iter1>& __x, const __wrap_iter<_Iter1>& __y) noexcept { - _LIBCUDACXX_DEBUG_ASSERT( - __get_const_db()->__less_than_comparable(_CUDA_VSTD::addressof(__x), _CUDA_VSTD::addressof(__y)), - "Attempted to compare incomparable iterators"); return __x.base() < __y.base(); } @@ -231,8 +158,6 @@ template _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 bool operator<(const __wrap_iter<_Iter1>& __x, const __wrap_iter<_Iter2>& __y) noexcept { - _LIBCUDACXX_DEBUG_ASSERT(__get_const_db()->__less_than_comparable(&__x, &__y), - "Attempted to compare incomparable iterators"); return __x.base() < __y.base(); } @@ -296,9 +221,6 @@ template _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 auto operator-(const __wrap_iter<_Iter1>& __x, const __wrap_iter<_Iter2>& __y) noexcept -> decltype(__x.base() - __y.base()) { - _LIBCUDACXX_DEBUG_ASSERT( - __get_const_db()->__less_than_comparable(_CUDA_VSTD::addressof(__x), _CUDA_VSTD::addressof(__y)), - "Attempted to subtract incompatible iterators"); return __x.base() - __y.base(); } @@ -331,6 +253,4 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT pointer_traits<__wrap_iter<_It>> _LIBCUDACXX_END_NAMESPACE_STD -#undef _LIBCUDACXX_CONSTEXPR_IF_NODEBUG - #endif // _LIBCUDACXX___ITERATOR_WRAP_ITER_H 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 #include #include #include @@ -62,7 +63,6 @@ #include #include #include -#include _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 #include #include #include @@ -62,7 +63,6 @@ #include #include #include -#include _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 #include #include #include @@ -37,7 +38,6 @@ #include #include #include -#include #ifdef _CCCL_CUDA_COMPILER_CLANG # include @@ -113,7 +113,7 @@ _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX20 __enable_if_t::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::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 ::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 ::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 ::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 , 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 // all public C++ headers provide the assertion handler #include #include #include @@ -33,7 +34,6 @@ #include #include #include -#include // all public C++ headers provide the assertion handler // standard-mandated includes #include 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 // all public C++ headers provide the assertion handler #include #include #include -#include // all public C++ headers provide the assertion handler #include #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 b75a4cc0df..dfd4b430b1 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 #include #include #include @@ -50,7 +51,6 @@ #include #include #include -#include #if _CCCL_STD_VER >= 2017 && !defined(_CCCL_COMPILER_MSVC_2017) @@ -275,8 +275,8 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT subrange : public view_interface) { - _LIBCUDACXX_ASSERT((__end_ - __begin_) == static_cast>(__n), - "_CUDA_VSTD::_CUDA_VRANGES::subrange was passed an invalid size hint"); + _CCCL_ASSERT((__end_ - __begin_) == static_cast>(__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 #include #include #include @@ -34,7 +35,6 @@ #include #include #include -#include _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) _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 _LIBCUDACXX_AND common_range) _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 #include #include #include -#include _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 #include #include -#include // all public C++ headers provide the assertion handler +#include // all public C++ headers provide the assertion handler #include #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 #include #include +#include // all public C++ headers provide the assertion handler #include #include #include #include #include -#include // all public C++ headers provide the assertion handler #include #include #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 // all public C++ headers provide the assertion handler #include #include #include @@ -45,7 +46,6 @@ #include #include #include -#include // all public C++ headers provide the assertion handler #include #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 3568b3b746..0000000000 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/__assert +++ /dev/null @@ -1,66 +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 - -#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 - -// 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(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(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 82fa6cf575..6fdffc57e6 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/__config +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__config @@ -1338,15 +1338,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/__debug b/libcudacxx/include/cuda/std/detail/libcxx/include/__debug deleted file mode 100644 index ecac8a9167..0000000000 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/__debug +++ /dev/null @@ -1,288 +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___DEBUG -#define _LIBCUDACXX___DEBUG - -#include - -#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 -#include -#include - -#if defined(_LIBCUDACXX_ENABLE_DEBUG_MODE) && !defined(_LIBCUDACXX_DEBUG_RANDOMIZE_UNSPECIFIED_STABILITY) -# define _LIBCUDACXX_DEBUG_RANDOMIZE_UNSPECIFIED_STABILITY -#endif - -#if defined(_LIBCUDACXX_ENABLE_DEBUG_MODE) && !defined(_LIBCUDACXX_DEBUG_ITERATOR_BOUNDS_CHECKING) -# define _LIBCUDACXX_DEBUG_ITERATOR_BOUNDS_CHECKING -#endif - -#ifdef _LIBCUDACXX_ENABLE_DEBUG_MODE -# define _LIBCUDACXX_DEBUG_ASSERT(x, m) _LIBCUDACXX_ASSERT(::std::__libcpp_is_constant_evaluated() || (x), m) -#else -# define _LIBCUDACXX_DEBUG_ASSERT(x, m) ((void) 0) -#endif - -#if defined(_LIBCUDACXX_ENABLE_DEBUG_MODE) || defined(_LIBCUDACXX_BUILDING_LIBRARY) - -_LIBCUDACXX_BEGIN_NAMESPACE_STD - -struct _CCCL_TYPE_VISIBILITY_DEFAULT __c_node; - -struct _CCCL_TYPE_VISIBILITY_DEFAULT __i_node -{ - void* __i_; - __i_node* __next_; - __c_node* __c_; - - __i_node(const __i_node&) = delete; - __i_node& operator=(const __i_node&) = delete; - - _LIBCUDACXX_HIDE_FROM_ABI __i_node(void* __i, __i_node* __next, __c_node* __c) - : __i_(__i) - , __next_(__next) - , __c_(__c) - {} - ~__i_node(); -}; - -struct _CCCL_TYPE_VISIBILITY_DEFAULT __c_node -{ - void* __c_; - __c_node* __next_; - __i_node** beg_; - __i_node** end_; - __i_node** cap_; - - __c_node(const __c_node&) = delete; - __c_node& operator=(const __c_node&) = delete; - - _LIBCUDACXX_HIDE_FROM_ABI explicit __c_node(void* __c, __c_node* __next) - : __c_(__c) - , __next_(__next) - , beg_(nullptr) - , end_(nullptr) - , cap_(nullptr) - {} - virtual ~__c_node(); - - virtual bool __dereferenceable(const void*) const = 0; - virtual bool __decrementable(const void*) const = 0; - virtual bool __addable(const void*, ptrdiff_t) const = 0; - virtual bool __subscriptable(const void*, ptrdiff_t) const = 0; - - void __add(__i_node* __i); - _CCCL_VISIBILITY_HIDDEN void __remove(__i_node* __i); -}; - -template -struct _C_node : public __c_node -{ - explicit _C_node(void* __c, __c_node* __n) - : __c_node(__c, __n) - {} - - bool __dereferenceable(const void*) const override; - bool __decrementable(const void*) const override; - bool __addable(const void*, ptrdiff_t) const override; - bool __subscriptable(const void*, ptrdiff_t) const override; -}; - -template -inline bool _C_node<_Cont>::__dereferenceable(const void* __i) const -{ - typedef typename _Cont::const_iterator iterator; - const iterator* __j = static_cast(__i); - _Cont* _Cp = static_cast<_Cont*>(__c_); - return _Cp->__dereferenceable(__j); -} - -template -inline bool _C_node<_Cont>::__decrementable(const void* __i) const -{ - typedef typename _Cont::const_iterator iterator; - const iterator* __j = static_cast(__i); - _Cont* _Cp = static_cast<_Cont*>(__c_); - return _Cp->__decrementable(__j); -} - -template -inline bool _C_node<_Cont>::__addable(const void* __i, ptrdiff_t __n) const -{ - typedef typename _Cont::const_iterator iterator; - const iterator* __j = static_cast(__i); - _Cont* _Cp = static_cast<_Cont*>(__c_); - return _Cp->__addable(__j, __n); -} - -template -inline bool _C_node<_Cont>::__subscriptable(const void* __i, ptrdiff_t __n) const -{ - typedef typename _Cont::const_iterator iterator; - const iterator* __j = static_cast(__i); - _Cont* _Cp = static_cast<_Cont*>(__c_); - return _Cp->__subscriptable(__j, __n); -} - -class _CCCL_TYPE_VISIBILITY_DEFAULT __libcpp_db -{ - __c_node** __cbeg_; - __c_node** __cend_; - size_t __csz_; - __i_node** __ibeg_; - __i_node** __iend_; - size_t __isz_; - - explicit __libcpp_db(); - -public: - __libcpp_db(const __libcpp_db&) = delete; - __libcpp_db& operator=(const __libcpp_db&) = delete; - - ~__libcpp_db(); - - class __db_c_iterator; - class __db_c_const_iterator; - class __db_i_iterator; - class __db_i_const_iterator; - - __db_c_const_iterator __c_end() const; - __db_i_const_iterator __i_end() const; - - typedef __c_node*(_InsertConstruct) (void*, void*, __c_node*); - - template - _LIBCUDACXX_HIDE_FROM_ABI static __c_node* __create_C_node(void* __mem, void* __c, __c_node* __next) - { - return ::new (__mem) _C_node<_Cont>(__c, __next); - } - - template - _LIBCUDACXX_HIDE_FROM_ABI void __insert_c(_Cont* __c) - { - __insert_c(static_cast(__c), &__create_C_node<_Cont>); - } - - void __insert_i(void* __i); - void __insert_c(void* __c, _InsertConstruct* __fn); - void __erase_c(void* __c); - - void __insert_ic(void* __i, const void* __c); - void __iterator_copy(void* __i, const void* __i0); - void __erase_i(void* __i); - - void* __find_c_from_i(void* __i) const; - void __invalidate_all(void* __c); - __c_node* __find_c_and_lock(void* __c) const; - __c_node* __find_c(void* __c) const; - void unlock() const; - - void swap(void* __c1, void* __c2); - - bool __dereferenceable(const void* __i) const; - bool __decrementable(const void* __i) const; - bool __addable(const void* __i, ptrdiff_t __n) const; - bool __subscriptable(const void* __i, ptrdiff_t __n) const; - bool __less_than_comparable(const void* __i, const void* __j) const; - -private: - _CCCL_VISIBILITY_HIDDEN __i_node* __insert_iterator(void* __i); - _CCCL_VISIBILITY_HIDDEN __i_node* __find_iterator(const void* __i) const; - - friend _LIBCUDACXX_HIDE_FROM_ABI __libcpp_db* __get_db(); -}; - -_LIBCUDACXX_HIDE_FROM_ABI __libcpp_db* __get_db(); -_LIBCUDACXX_HIDE_FROM_ABI const __libcpp_db* __get_const_db(); - -_LIBCUDACXX_END_NAMESPACE_STD - -#endif // defined(_LIBCUDACXX_ENABLE_DEBUG_MODE) || defined(_LIBCUDACXX_BUILDING_LIBRARY) - -_LIBCUDACXX_BEGIN_NAMESPACE_STD - -template -_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 void __debug_db_insert_c(_Tp* __c) -{ -#ifdef _LIBCUDACXX_ENABLE_DEBUG_MODE - if (!__libcpp_is_constant_evaluated()) - { - __get_db()->__insert_c(__c); - } -#else - (void) (__c); -#endif -} - -template -_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 void __debug_db_insert_i(_Tp* __i) -{ -#ifdef _LIBCUDACXX_ENABLE_DEBUG_MODE - if (!__libcpp_is_constant_evaluated()) - { - __get_db()->__insert_i(__i); - } -#else - (void) (__i); -#endif -} - -template -_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 void __debug_db_erase_c(_Tp* __c) -{ -#ifdef _LIBCUDACXX_ENABLE_DEBUG_MODE - if (!__libcpp_is_constant_evaluated()) - { - __get_db()->__erase_c(__c); - } -#else - (void) (__c); -#endif -} - -template -_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 void __debug_db_swap(_Tp* __lhs, _Tp* __rhs) -{ -#ifdef _LIBCUDACXX_ENABLE_DEBUG_MODE - if (!__libcpp_is_constant_evaluated()) - { - __get_db()->swap(__lhs, __rhs); - } -#else - (void) (__lhs); - (void) (__rhs); -#endif -} - -template -_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 void __debug_db_invalidate_all(_Tp* __c) -{ -#ifdef _LIBCUDACXX_ENABLE_DEBUG_MODE - if (!__libcpp_is_constant_evaluated()) - { - __get_db()->__invalidate_all(__c); - } -#else - (void) (__c); -#endif -} - -_LIBCUDACXX_END_NAMESPACE_STD - -#endif // _LIBCUDACXX___DEBUG diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__string b/libcudacxx/include/cuda/std/detail/libcxx/include/__string index 8e6f13469e..5e6023ec3a 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/__string +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__string @@ -68,9 +68,8 @@ template <> struct char_traits; // c++20 #include #include #include +#include // all public C++ headers provide the assertion handler #include -#include // all public C++ headers provide the assertion handler -#include #include _CCCL_PUSH_MACROS @@ -207,7 +206,7 @@ _LIBCUDACXX_HIDE_FROM_ABI _CharT* char_traits<_CharT>::move(char_type* __s1, con template _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) { @@ -290,7 +289,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT char_traits } _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 @@ -405,7 +404,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT char_traits } _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 @@ -562,7 +561,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT char_traits _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); } @@ -769,7 +768,7 @@ char_traits::move(char_type* __s1, const char_type* __s2, size_t __n) _LIBCUDACXX_HIDE_FROM_ABI char16_t* char_traits::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) { @@ -910,7 +909,7 @@ char_traits::move(char_type* __s1, const char_type* __s2, size_t __n) _LIBCUDACXX_HIDE_FROM_ABI char32_t* char_traits::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 // all public C++ headers provide the assertion handler #include #include #include -#include // all public C++ headers provide the assertion handler #include _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 b44c0506ef..0000000000 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/__verbose_abort +++ /dev/null @@ -1,59 +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 - -#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 - -// Provide a default implementation of __libcpp_verbose_abort if we know that neither the built -// library nor the user is providing one. Otherwise, just declare it and use the one from the -// built library or the one provided by the user. -// -// We can't provide a great implementation because it needs to be pretty much -// dependency-free (this is included everywhere else in the library). -#if defined(_LIBCUDACXX_HAS_NO_VERBOSE_ABORT_IN_LIBRARY) \ - && !defined(_LIBCUDACXX_AVAILABILITY_CUSTOM_VERBOSE_ABORT_PROVIDED) - -extern "C" void abort(); - -_LIBCUDACXX_BEGIN_NAMESPACE_STD - -_CCCL_NORETURN _LIBCUDACXX_ATTRIBUTE_FORMAT(__printf__, 1, 2) - _LIBCUDACXX_HIDE_FROM_ABI void __libcpp_verbose_abort(const char*, ...) -{ - ::abort(); - __builtin_unreachable(); // never reached, but needed to tell the compiler that the function never returns -} - -_LIBCUDACXX_END_NAMESPACE_STD - -#else - -_LIBCUDACXX_BEGIN_NAMESPACE_STD - -_CCCL_NORETURN -_CCCL_HOST_DEVICE _LIBCUDACXX_ATTRIBUTE_FORMAT(__printf__, 1, 2) void __libcpp_verbose_abort(const char* __format, ...); - -_LIBCUDACXX_END_NAMESPACE_STD - -#endif - -#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 3641e69bd5..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 #include #include #include +#include // all public C++ headers provide the assertion handler #include #include #include @@ -751,8 +752,6 @@ template #include #include #include -#include // all public C++ headers provide the assertion handler -#include #include #include #include @@ -1148,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 c8091fe41c..ff75040a7a 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/array +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/array @@ -123,6 +123,7 @@ template const T&& get(const array&&) noexce #include #include #include +#include // all public C++ headers provide the assertion handler #include #include #include @@ -139,7 +140,6 @@ template const T&& get(const array&&) noexce #include #include #include -#include // all public C++ headers provide the assertion handler #include #include @@ -264,12 +264,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"); + _CCCL_ASSERT(__n < _Size, "out-of-bounds access in std::array"); 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"); + _CCCL_ASSERT(__n < _Size, "out-of-bounds access in std::array"); return __elems_[__n]; } @@ -434,14 +434,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::operator[] on a zero-sized array"); + _CCCL_ASSERT(false, "cannot call array::operator[] on a zero-sized array"); _CCCL_UNREACHABLE(); return *data(); } _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 const_reference operator[](size_type) const noexcept { - _LIBCUDACXX_ASSERT(false, "cannot call array::operator[] on a zero-sized array"); + _CCCL_ASSERT(false, "cannot call array::operator[] on a zero-sized array"); _CCCL_UNREACHABLE(); return *data(); } @@ -460,28 +460,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::front() on a zero-sized array"); + _CCCL_ASSERT(false, "cannot call array::front() on a zero-sized array"); _CCCL_UNREACHABLE(); return *data(); } _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 const_reference front() const noexcept { - _LIBCUDACXX_ASSERT(false, "cannot call array::front() on a zero-sized array"); + _CCCL_ASSERT(false, "cannot call array::front() on a zero-sized array"); _CCCL_UNREACHABLE(); return *data(); } _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 reference back() noexcept { - _LIBCUDACXX_ASSERT(false, "cannot call array::back() on a zero-sized array"); + _CCCL_ASSERT(false, "cannot call array::back() on a zero-sized array"); _CCCL_UNREACHABLE(); return *data(); } _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 const_reference back() const noexcept { - _LIBCUDACXX_ASSERT(false, "cannot call array::back() on a zero-sized array"); + _CCCL_ASSERT(false, "cannot call array::back() on a zero-sized array"); _CCCL_UNREACHABLE(); return *data(); } diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/barrier b/libcudacxx/include/cuda/std/detail/libcxx/include/barrier index a3102d6087..f8767a4c4d 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/barrier +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/barrier @@ -53,12 +53,11 @@ namespace std # pragma system_header #endif // no system header +#include // all public C++ headers provide the assertion handler #include #include #include #include -#include // all public C++ headers provide the assertion handler -#include _CCCL_PUSH_MACROS @@ -105,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; @@ -122,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)) { @@ -131,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)) @@ -145,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); @@ -165,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; @@ -175,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) { @@ -313,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_DEBUG_ASSERT(__result >= 0, ""); + _CCCL_ASSERT(__result >= 0, ""); if (0 == __result) { @@ -373,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_DEBUG_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); } @@ -398,7 +397,7 @@ public: __barrier_base(ptrdiff_t __count, __empty_completion = __empty_completion()) : __phase_arrived_expected(__init(__count)) { - _LIBCUDACXX_DEBUG_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 // all public C++ headers provide the assertion handler #include #include #include #include #include #include -#include // all public C++ headers provide the assertion handler #include #include 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 complex tanh (const complex&); # pragma system_header #endif // no system header +#include // all public C++ headers provide the assertion handler #include #include #include @@ -252,7 +253,6 @@ template complex tanh (const complex&); #include #include #include -#include // all public C++ headers provide the assertion handler #include #include 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 // all public C++ headers provide the assertion handler #include #include #include -#include // all public C++ headers provide the assertion handler #include _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::state_type> wstreampos; # pragma system_header #endif // no system header +#include // all public C++ headers provide the assertion handler #include -#include // all public C++ headers provide the assertion handler #include _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 e74c77cc1a..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 +#include // all public C++ headers provide the assertion handler #include -#include // all public C++ headers provide the assertion handler -#include #include _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; # pragma system_header #endif // no system header +#include // all public C++ headers provide the assertion handler #include -#include // all public C++ headers provide the assertion handler #include #include diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/optional b/libcudacxx/include/cuda/std/detail/libcxx/include/optional index f8663cf319..ba05d15fb8 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/optional +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/optional @@ -169,6 +169,7 @@ template # pragma system_header #endif // no system header +#include // all public C++ headers provide the assertion handler #include #include #include @@ -198,9 +199,7 @@ template #include #include #include -#include // all public C++ headers provide the assertion handler #include -#include #include // standard-mandated includes @@ -396,7 +395,7 @@ struct __optional_storage_base : __optional_destruct_base<_Tp> template _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 @@ -850,37 +849,37 @@ public: _LIBCUDACXX_HIDE_FROM_ABI constexpr add_pointer_t 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 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 // all public C++ headers provide the assertion handler #include #include #include -#include // all public C++ headers provide the assertion handler #include _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 // all public C++ headers provide the assertion handler #include #include -#include // all public C++ headers provide the assertion handler #include _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 12479eb37b..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 # pragma system_header #endif // no system header +#include // all public C++ headers provide the assertion handler #include #include #include @@ -171,7 +172,6 @@ template #include #include #include // for ptrdiff_t -#include // all public C++ headers provide the assertion handler // standard-mandated includes #include @@ -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 first(size_type __count) const noexcept { - _LIBCUDACXX_ASSERT(__count <= size(), "span::first(count): count out of range"); + _CCCL_ASSERT(__count <= size(), "span::first(count): count out of range"); return {data(), __count}; } _LIBCUDACXX_HIDE_FROM_ABI constexpr span last(size_type __count) const noexcept { - _LIBCUDACXX_ASSERT(__count <= size(), "span::last(count): count out of range"); + _CCCL_ASSERT(__count <= size(), "span::last(count): count out of range"); return {data() + size() - __count, __count}; } @@ -471,14 +471,14 @@ public: _LIBCUDACXX_HIDE_FROM_ABI constexpr span subspan(size_type __offset, size_type __count = dynamic_extent) const noexcept { - _LIBCUDACXX_ASSERT(__offset <= size(), "span::subspan(offset, count): offset out of range"); - _LIBCUDACXX_ASSERT(__count <= size() || __count == dynamic_extent, - "span::subspan(offset, count): count out of range"); + _CCCL_ASSERT(__offset <= size(), "span::subspan(offset, count): offset out of range"); + _CCCL_ASSERT(__count <= size() || __count == dynamic_extent, + "span::subspan(offset, count): count out of range"); if (__count == dynamic_extent) { return {data() + __offset, size() - __offset}; } - _LIBCUDACXX_ASSERT(__count <= size() - __offset, "span::subspan(offset, count): offset + count out of range"); + _CCCL_ASSERT(__count <= size() - __offset, "span::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::operator[](index): index out of range"); + _CCCL_ASSERT(__idx < size(), "span::operator[](index): index out of range"); return __data_[__idx]; } _LIBCUDACXX_HIDE_FROM_ABI constexpr reference front() const noexcept { - _LIBCUDACXX_ASSERT(!empty(), "span::front() on empty span"); + _CCCL_ASSERT(!empty(), "span::front() on empty span"); return __data_[0]; } _LIBCUDACXX_HIDE_FROM_ABI constexpr reference back() const noexcept { - _LIBCUDACXX_ASSERT(!empty(), "span::back() on empty span"); + _CCCL_ASSERT(!empty(), "span::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 _LIBCUDACXX_HIDE_FROM_ABI constexpr span first() const noexcept { - _LIBCUDACXX_ASSERT(_Count <= size(), "span::first(): Count out of range"); + _CCCL_ASSERT(_Count == 0 || _Count <= size(), "span::first(): Count out of range"); return span{data(), _Count}; } template _LIBCUDACXX_HIDE_FROM_ABI constexpr span last() const noexcept { - _LIBCUDACXX_ASSERT(_Count <= size(), "span::last(): Count out of range"); + _CCCL_ASSERT(_Count == 0 || _Count <= size(), "span::last(): Count out of range"); return span{data() + size() - _Count, _Count}; } _LIBCUDACXX_HIDE_FROM_ABI constexpr span first(size_type __count) const noexcept { - _LIBCUDACXX_ASSERT(__count <= size(), "span::first(count): count out of range"); + _CCCL_ASSERT(__count <= size(), "span::first(count): count out of range"); return {data(), __count}; } _LIBCUDACXX_HIDE_FROM_ABI constexpr span last(size_type __count) const noexcept { - _LIBCUDACXX_ASSERT(__count <= size(), "span::last(count): count out of range"); + _CCCL_ASSERT(__count <= size(), "span::last(count): count out of range"); return {data() + size() - __count, __count}; } @@ -689,23 +689,22 @@ public: template _LIBCUDACXX_HIDE_FROM_ABI constexpr __subspan_t<_Offset, _Count> subspan() const noexcept { - _LIBCUDACXX_ASSERT(_Offset <= size(), "span::subspan(): Offset out of range"); - _LIBCUDACXX_ASSERT(_Count == dynamic_extent || _Count <= size() - _Offset, - "span::subspan(): Offset + Count out of range"); + _CCCL_ASSERT(_Offset == 0 || _Offset <= size(), "span::subspan(): Offset out of range"); + _CCCL_ASSERT(_Count == dynamic_extent || _Count == 0 || _Count <= size() - _Offset, + "span::subspan(): Offset + Count out of range"); return __subspan_t<_Offset, _Count>{data() + _Offset, _Count == dynamic_extent ? size() - _Offset : _Count}; } constexpr span _LIBCUDACXX_HIDE_FROM_ABI subspan(size_type __offset, size_type __count = dynamic_extent) const noexcept { - _LIBCUDACXX_ASSERT(__offset <= size(), "span::subspan(offset, count): offset out of range"); - _LIBCUDACXX_ASSERT(__count <= size() || __count == dynamic_extent, - "span::subspan(offset, count): count out of range"); + _CCCL_ASSERT(__offset <= size(), "span::subspan(offset, count): offset out of range"); + _CCCL_ASSERT(__count <= size() || __count == dynamic_extent, "span::subspan(offset, count): count out of range"); if (__count == dynamic_extent) { return {data() + __offset, size() - __offset}; } - _LIBCUDACXX_ASSERT(__count <= size() - __offset, "span::subspan(offset, count): offset + count out of range"); + _CCCL_ASSERT(__count <= size() - __offset, "span::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::operator[](index): index out of range"); + _CCCL_ASSERT(__idx < size(), "span::operator[](index): index out of range"); return __data_[__idx]; } _LIBCUDACXX_HIDE_FROM_ABI constexpr reference front() const noexcept { - _LIBCUDACXX_ASSERT(!empty(), "span::front() on empty span"); + _CCCL_ASSERT(!empty(), "span::front() on empty span"); return __data_[0]; } _LIBCUDACXX_HIDE_FROM_ABI constexpr reference back() const noexcept { - _LIBCUDACXX_ASSERT(!empty(), "span::back() on empty span"); + _CCCL_ASSERT(!empty(), "span::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 // all public C++ headers provide the assertion handler #include -#include // all public C++ headers provide the assertion handler #ifndef _CCCL_NO_EXCEPTIONS # include diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/tuple b/libcudacxx/include/cuda/std/detail/libcxx/include/tuple index b02ede9c81..67176216b0 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/tuple +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/tuple @@ -155,6 +155,7 @@ template # pragma system_header #endif // no system header +#include // all public C++ headers provide the assertion handler #include #include #include @@ -178,7 +179,6 @@ template #include #include #include -#include // all public C++ headers provide the assertion handler #include #include diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/variant b/libcudacxx/include/cuda/std/detail/libcxx/include/variant index df74afebe9..c5be059b72 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 // all public C++ headers provide the assertion handler +#include // all public C++ headers provide the assertion handler #include #ifndef _LIBCUDACXX_HAS_NO_SPACESHIP_OPERATOR # include 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 // all public C++ headers provide the assertion handler #include #include #include #include -#include // all public C++ headers provide the assertion handler #include #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 #include +#include // all public C++ headers provide the assertion handler #include #include #include @@ -50,7 +51,6 @@ #include #include #include -#include // all public C++ headers provide the assertion handler #include // for forward declarations of vector and string #include diff --git a/libcudacxx/include/cuda/std/inplace_vector b/libcudacxx/include/cuda/std/inplace_vector index 771beef955..a34d927b5a 100644 --- a/libcudacxx/include/cuda/std/inplace_vector +++ b/libcudacxx/include/cuda/std/inplace_vector @@ -32,6 +32,7 @@ # include # include # include +# include // all public C++ headers provide the assertion handler # include # include # include @@ -66,7 +67,6 @@ # include # include # include -# include // all public C++ headers provide the assertion handler # include # include # include 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 // all public C++ headers provide the assertion handler #include #include #include @@ -57,7 +58,6 @@ #include #include #include -#include // all public C++ headers provide the assertion handler #include // for forward declarations of vector and string. #include 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 // all public C++ headers provide the assertion handler #include #include #include @@ -32,7 +33,6 @@ #include #include #include -#include // all public C++ headers provide the assertion handler #include #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 // all public C++ headers provide the assertion handler #include #include #include @@ -34,7 +35,6 @@ #include #include #include -#include // all public C++ headers provide the assertion handler // standard mandated include #include 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 // all public C++ headers provide the assertion handler #include #include #include @@ -38,7 +39,6 @@ _CCCL_DIAG_SUPPRESS_MSVC(4848) #include #include #include -#include // all public C++ headers provide the assertion handler // standard-mandated includes #include diff --git a/libcudacxx/include/cuda/std/type_traits b/libcudacxx/include/cuda/std/type_traits index 97e80a2e4d..9ac236f42e 100644 --- a/libcudacxx/include/cuda/std/type_traits +++ b/libcudacxx/include/cuda/std/type_traits @@ -22,6 +22,7 @@ #endif // no system header #include +#include // all public C++ headers provide the assertion handler #include #include #include @@ -161,7 +162,6 @@ #include #include #include -#include // all public C++ headers provide the assertion handler #include #endif // _CUDA_STD_TYPE_TRAITS diff --git a/libcudacxx/include/cuda/std/utility b/libcudacxx/include/cuda/std/utility index 766e3e6a33..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 // all public C++ headers provide the assertion handler #include #include #include @@ -57,8 +58,6 @@ #include #include #include -#include // all public C++ headers provide the assertion handler -#include #include // standard-mandated includes diff --git a/libcudacxx/test/internal_headers/CMakeLists.txt b/libcudacxx/test/internal_headers/CMakeLists.txt index c9c060c2dd..4120e3d22e 100644 --- a/libcudacxx/test/internal_headers/CMakeLists.txt +++ b/libcudacxx/test/internal_headers/CMakeLists.txt @@ -41,6 +41,7 @@ function(libcudacxx_create_internal_header_test header_name, headertest_src, fal $<$:${headertest_warning_levels_device}> $<$:${headertest_warning_levels_host}> -DLIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE) + target_compile_definitions(headertest_${header_name} PRIVATE CCCL_ENABLE_ASSERTIONS) 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 605c98e888..b98c11d1d6 100644 --- a/libcudacxx/test/libcudacxx/CMakeLists.txt +++ b/libcudacxx/test/libcudacxx/CMakeLists.txt @@ -39,6 +39,9 @@ else() # NOT LIBCUDACXX_TEST_WITH_NVRTC set(LIBCUDACXX_TEST_COMPILER_FLAGS "-DLIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE") endif() +# enable exceptions and assertions in tests +string(APPEND LIBCUDACXX_TEST_COMPILER_FLAGS " -DCCCL_ENABLE_ASSERTIONS") + if (NOT MSVC AND NOT ${CMAKE_CUDA_COMPILER_ID} STREQUAL "Clang") set(LIBCUDACXX_WARNING_LEVEL "--compiler-options=-Wall --compiler-options=-Wextra") endif() 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 @@ -64,15 +64,15 @@ __host__ __device__ void test_death() cuda::std::__bounded_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/libcudacxx/libcxx/iterators/unwrap_iter.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/iterators/unwrap_iter.pass.cpp index 3f151ad7e8..7d70944cad 100644 --- a/libcudacxx/test/libcudacxx/libcxx/iterators/unwrap_iter.pass.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/iterators/unwrap_iter.pass.cpp @@ -65,7 +65,7 @@ TEST_CONSTEXPR_CXX20 bool test() return true; } -#endif +#endif // _LIBCUDACXX_HAS_STRING int main(int, char**) { diff --git a/libcudacxx/test/public_headers/CMakeLists.txt b/libcudacxx/test/public_headers/CMakeLists.txt index df84e5c95c..a9b3471cba 100644 --- a/libcudacxx/test/public_headers/CMakeLists.txt +++ b/libcudacxx/test/public_headers/CMakeLists.txt @@ -46,6 +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 CCCL_ENABLE_ASSERTIONS) # 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 8b41cd681c..224e49b712 100644 --- a/libcudacxx/test/public_headers_host_only/CMakeLists.txt +++ b/libcudacxx/test/public_headers_host_only/CMakeLists.txt @@ -28,6 +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 CCCL_ENABLE_ASSERTIONS) 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: From d7f10435e18b1c6665ce42c4947470983445d221 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Mon, 16 Sep 2024 13:53:39 +0200 Subject: [PATCH 02/20] Try to work better with nvhpc --- libcudacxx/include/cuda/std/__cccl/assert.h | 87 ++++++++++----------- 1 file changed, 40 insertions(+), 47 deletions(-) diff --git a/libcudacxx/include/cuda/std/__cccl/assert.h b/libcudacxx/include/cuda/std/__cccl/assert.h index b774225a61..8a5c483a2f 100644 --- a/libcudacxx/include/cuda/std/__cccl/assert.h +++ b/libcudacxx/include/cuda/std/__cccl/assert.h @@ -24,6 +24,8 @@ #include +#include + #if defined(_DEBUG) || defined(DEBUG) # ifndef _CCCL_ENABLE_DEBUG_MODE # define _CCCL_ENABLE_DEBUG_MODE @@ -51,60 +53,51 @@ # endif // CCCL_ENABLE_ASSERTIONS #endif // !CCCL_ENABLE_DEVICE_ASSERTIONS -//! Use internal nvcc implementation on device or the host library for clang-cuda -#ifdef __CUDA_ARCH__ -# ifdef CCCL_ENABLE_DEVICE_ASSERTIONS -# if defined(_CCCL_CUDA_COMPILER_NVCC) //! Use __assert_fail to implement device side asserts -# include -# define _CCCL_ASSERT_IMPL(expression, message) \ - __builtin_expect(static_cast(expression), 1) \ - ? (void) 0 \ - : __assert_fail(message, __FILE__, __LINE__, __func__) -# else // ^^^ _CCCL_CUDA_COMPILER_NVCC ^^^ / vvv !_CCCL_CUDA_COMPILER_NVCC vvv -# if __has_include(<__assert>) // libc++ uses _LIBCPP_ASSERT from <__assert> -# include <__assert> -# define _CCCL_ASSERT_IMPL(expression, message) _LIBCPP_ASSERT(expression, message) -# else // libstdc++ uses __glibcxx_assert from -# if defined(_GLIBCXX_RELEASE) && _GLIBCXX_RELEASE >= 12 -// libstdc++ does not fully qualify its use of `__is_constant_evaluated` -// It was introduced in the assert handling in 5e8a30d -// libstdc++ : Redefine __glibcxx_assert to work in C++ 23 constexpr -_LIBCUDACXX_BEGIN_NAMESPACE_STD -using ::std::__is_constant_evaluated; -_LIBCUDACXX_END_NAMESPACE_STD -# endif // _GLIBCXX_RELEASE >= 12 -# define _CCCL_ASSERT_IMPL(expression, message) __glibcxx_assert(expression) -# endif // libstdc++ -# endif // !_CCCL_CUDA_COMPILER_NVCC -# else // ^^^ CCCL_ENABLE_DEVICE_ASSERTIONS ^^^ / vvv !CCCL_ENABLE_DEVICE_ASSERTIONS vvv -# define _CCCL_ASSERT_IMPL -# endif // !CCCL_ENABLE_DEVICE_ASSERTIONS -#endif // __CUDA_ARCH__ - //! Use the different standard library implementations to implement host side asserts -#ifndef __CUDA_ARCH__ -# ifdef CCCL_ENABLE_HOST_ASSERTIONS -# if __has_include() // MSVC uses _STL_VERIFY from -# include -# define _CCCL_ASSERT_IMPL(expression, message) _STL_VERIFY(expression, message) -# elif __has_include(<__assert>) // libc++ uses _LIBCPP_ASSERT from <__assert> -# include <__assert> -# define _CCCL_ASSERT_IMPL(expression, message) _LIBCPP_ASSERT(expression, message) -# else // libstdc++ uses __glibcxx_assert from -# if defined(_GLIBCXX_RELEASE) && _GLIBCXX_RELEASE >= 12 +#ifdef CCCL_ENABLE_HOST_ASSERTIONS +# if __has_include() // MSVC uses _STL_VERIFY from +# include +# define _CCCL_ASSERT_HOST(expression, message) _STL_VERIFY(expression, message) +# elif __has_include(<__assert>) // libc++ uses _LIBCPP_ASSERT from <__assert> +# include <__assert> +# define _CCCL_ASSERT_HOST(expression, message) _LIBCPP_ASSERT(expression, message) +# elif __has_include() // libstdc++ uses __glibcxx_assert from +# if defined(_GLIBCXX_RELEASE) && _GLIBCXX_RELEASE >= 12 // libstdc++ does not fully qualify its use of `__is_constant_evaluated` // It was introduced in the assert handling in 5e8a30d // libstdc++ : Redefine __glibcxx_assert to work in C++ 23 constexpr _LIBCUDACXX_BEGIN_NAMESPACE_STD using ::std::__is_constant_evaluated; _LIBCUDACXX_END_NAMESPACE_STD -# endif // _GLIBCXX_RELEASE >= 12 -# define _CCCL_ASSERT_IMPL(expression, message) __glibcxx_assert(expression) -# endif // libstdc++ -# endif // !CCCL_ENABLE_HOST_ASSERTIONS -#endif // !__CUDA_ARCH__ +# endif // _GLIBCXX_RELEASE >= 12 +# define _CCCL_ASSERT_HOST(expression, message) __glibcxx_assert(expression) +# else // ^^^ libstdc++ ^^^ / vvv Unknown standard library vvv +# error "Unknown host standard library used." +# endif // Unknown standard library +#else // ^^^ CCCL_ENABLE_HOST_ASSERTIONS ^^^ / vvv !CCCL_ENABLE_HOST_ASSERTIONS vvv +# define _CCCL_ASSERT_HOST(expression, message) ((void) 0) +#endif // !CCCL_ENABLE_HOST_ASSERTIONS + +//! Use internal nvcc implementation on device or the host library for other cuda compilers +#ifdef CCCL_ENABLE_DEVICE_ASSERTIONS +# if defined(_CCCL_CUDA_COMPILER_NVCC) //! Use __assert_fail to implement device side asserts +# include +# define _CCCL_ASSERT_DEVICE(expression, message) \ + __builtin_expect(static_cast(expression), 1) \ + ? (void) 0 \ + : __assert_fail(message, __FILE__, __LINE__, __func__) +# elif defined(_CCCL_CUDA_COMPILER) +# define _CCCL_ASSERT_DEVICE(expression, message) _CCCL_ASSERT_HOST(expression, message) +# endif // _CCCL_CUDA_COMPILER +#else // ^^^ CCCL_ENABLE_DEVICE_ASSERTIONS ^^^ / vvv !CCCL_ENABLE_DEVICE_ASSERTIONS vvv +# define _CCCL_ASSERT_DEVICE(expression, message) ((void) 0) +#endif // !CCCL_ENABLE_DEVICE_ASSERTIONS + +//! Use the right assert in the right situation +#define _CCCL_ASSERT_IMPL(expression, message) \ + NV_IF_ELSE_TARGET( \ + NV_IS_DEVICE, (_CCCL_ASSERT_DEVICE(expression, message);), (_CCCL_ASSERT_HOST(expression, message);)) //! _CCCL_VERIFY is an unconditionally enabled assertion that is reserved for the most important checks //! Some compilers warn about `is_constant_evaluated()` in a non constexpr function, so silence that right away From dc79cb22270329804f886f5801bae37a18dc4554 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Mon, 16 Sep 2024 14:48:57 +0200 Subject: [PATCH 03/20] More specializations --- libcudacxx/include/cuda/std/__cccl/assert.h | 16 ++++++++++++---- 1 file changed, 12 insertions(+), 4 deletions(-) diff --git a/libcudacxx/include/cuda/std/__cccl/assert.h b/libcudacxx/include/cuda/std/__cccl/assert.h index 8a5c483a2f..7c1dff555c 100644 --- a/libcudacxx/include/cuda/std/__cccl/assert.h +++ b/libcudacxx/include/cuda/std/__cccl/assert.h @@ -94,10 +94,18 @@ _LIBCUDACXX_END_NAMESPACE_STD # define _CCCL_ASSERT_DEVICE(expression, message) ((void) 0) #endif // !CCCL_ENABLE_DEVICE_ASSERTIONS -//! Use the right assert in the right situation -#define _CCCL_ASSERT_IMPL(expression, message) \ - NV_IF_ELSE_TARGET( \ - NV_IS_DEVICE, (_CCCL_ASSERT_DEVICE(expression, message);), (_CCCL_ASSERT_HOST(expression, message);)) +//! Try to work with NVHPC +#if defined(_CCCL_CUDA_COMPILER_NVHPC) +# define _CCCL_ASSERT_IMPL(expression, message) \ + NV_IF_ELSE_TARGET( \ + NV_IS_DEVICE, (_CCCL_ASSERT_DEVICE(expression, message);), (_CCCL_ASSERT_HOST(expression, message);)) +#else // ^^^ _CCCL_CUDA_COMPILER_NVHPC ^^^ / vvv !_CCCL_CUDA_COMPILER_NVHPC vvv +# ifdef __CUDA_ARCH__ +# define _CCCL_ASSERT_IMPL(expression, message) _CCCL_ASSERT_DEVICE(expression, message) +# else // ^^^ __CUDA_ARCH__ ^^^ / vvv !__CUDA_ARCH__ vvv +# define _CCCL_ASSERT_IMPL(expression, message) _CCCL_ASSERT_HOST(expression, message) +# endif // !__CUDA_ARCH__ +#endif // !_CCCL_CUDA_COMPILER_NVHPC //! _CCCL_VERIFY is an unconditionally enabled assertion that is reserved for the most important checks //! Some compilers warn about `is_constant_evaluated()` in a non constexpr function, so silence that right away From 9f95036810a6bff4e63bf0d0ed39ef9dc1d3b3e0 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Mon, 16 Sep 2024 13:42:31 +0000 Subject: [PATCH 04/20] Suppress warning about is_constant_evaluated --- libcudacxx/include/cuda/std/__cccl/assert.h | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/libcudacxx/include/cuda/std/__cccl/assert.h b/libcudacxx/include/cuda/std/__cccl/assert.h index 7c1dff555c..92983c1ddc 100644 --- a/libcudacxx/include/cuda/std/__cccl/assert.h +++ b/libcudacxx/include/cuda/std/__cccl/assert.h @@ -101,9 +101,11 @@ _LIBCUDACXX_END_NAMESPACE_STD NV_IS_DEVICE, (_CCCL_ASSERT_DEVICE(expression, message);), (_CCCL_ASSERT_HOST(expression, message);)) #else // ^^^ _CCCL_CUDA_COMPILER_NVHPC ^^^ / vvv !_CCCL_CUDA_COMPILER_NVHPC vvv # ifdef __CUDA_ARCH__ -# define _CCCL_ASSERT_IMPL(expression, message) _CCCL_ASSERT_DEVICE(expression, message) +# define _CCCL_ASSERT_IMPL(expression, message) \ + _CCCL_NV_DIAG_SUPPRESS(3060) _CCCL_ASSERT_DEVICE(expression, message) _CCCL_NV_DIAG_DEFAULT(3060) # else // ^^^ __CUDA_ARCH__ ^^^ / vvv !__CUDA_ARCH__ vvv -# define _CCCL_ASSERT_IMPL(expression, message) _CCCL_ASSERT_HOST(expression, message) +# define _CCCL_ASSERT_IMPL(expression, message) \ + _CCCL_NV_DIAG_SUPPRESS(3060) _CCCL_ASSERT_HOST(expression, message) _CCCL_NV_DIAG_DEFAULT(3060) # endif // !__CUDA_ARCH__ #endif // !_CCCL_CUDA_COMPILER_NVHPC From 3c679d3d1fc60d4870652d2b23cc73f80ecf1798 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Mon, 16 Sep 2024 15:11:03 +0000 Subject: [PATCH 05/20] Avoid tests failing due to excessive constant folding --- .../partial_sort_copy.pass.cpp | 25 +++++++++++-------- .../partial_sort_copy_comp.pass.cpp | 25 +++++++++++-------- 2 files changed, 30 insertions(+), 20 deletions(-) diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.sort/partial.sort.copy/partial_sort_copy.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.sort/partial.sort.copy/partial_sort_copy.pass.cpp index f67f570090..f2302f5f43 100644 --- a/libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.sort/partial.sort.copy/partial_sort_copy.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.sort/partial.sort.copy/partial_sort_copy.pass.cpp @@ -20,6 +20,7 @@ #include #include +#include #include "MoveOnly.h" #include "test_iterators.h" @@ -77,14 +78,20 @@ __host__ __device__ TEST_CONSTEXPR_CXX14 bool test() assert(j == 75); test, random_access_iterator>(); - test, int*>(); - test>(); - test(); + if (!cuda::std::__libcpp_is_constant_evaluated()) // This breaks some compilers due to excessive constant folding + { + test, int*>(); + test>(); + test(); + } test, random_access_iterator>(); - test, MoveOnly*>(); - test>(); - test(); + if (!cuda::std::__libcpp_is_constant_evaluated()) // This breaks some compilers due to excessive constant folding + { + test, MoveOnly*>(); + test>(); + test(); + } return true; } @@ -92,11 +99,9 @@ __host__ __device__ TEST_CONSTEXPR_CXX14 bool test() int main(int, char**) { test(); -#if !defined(TEST_COMPILER_CLANG) && !defined(TEST_COMPILER_MSVC) // Over constexpr evaluation limit -# if TEST_STD_VER >= 2014 +#if TEST_STD_VER >= 2014 static_assert(test(), ""); -# endif // TEST_STD_VER >= 2014 -#endif // !TEST_COMPILER_CLANG && !TEST_COMPILER_MSVC +#endif // TEST_STD_VER >= 2014 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.sort/partial.sort.copy/partial_sort_copy_comp.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.sort/partial.sort.copy/partial_sort_copy_comp.pass.cpp index 5cda73bb48..a8a9b010fa 100644 --- a/libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.sort/partial.sort.copy/partial_sort_copy_comp.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.sort/partial.sort.copy/partial_sort_copy_comp.pass.cpp @@ -23,6 +23,7 @@ #include #include #include +#include #include "MoveOnly.h" #include "test_iterators.h" @@ -83,14 +84,20 @@ __host__ __device__ TEST_CONSTEXPR_CXX14 bool test() assert(j == 75); test, random_access_iterator>(); - test, int*>(); - test>(); - test(); + if (!cuda::std::__libcpp_is_constant_evaluated()) // This breaks some compilers due to excessive constant folding + { + test, int*>(); + test>(); + test(); + } test, random_access_iterator>(); - test, MoveOnly*>(); - test>(); - test(); + if (!cuda::std::__libcpp_is_constant_evaluated()) // This breaks some compilers due to excessive constant folding + { + test, MoveOnly*>(); + test>(); + test(); + } return true; } @@ -98,11 +105,9 @@ __host__ __device__ TEST_CONSTEXPR_CXX14 bool test() int main(int, char**) { test(); -#if !defined(TEST_COMPILER_CLANG) && !defined(TEST_COMPILER_MSVC) // Over constexpr evaluation limit -# if TEST_STD_VER >= 2014 +#if TEST_STD_VER >= 2014 static_assert(test(), ""); -# endif // TEST_STD_VER >= 2014 -#endif // !TEST_COMPILER_CLANG && !TEST_COMPILER_MSVC +#endif // TEST_STD_VER >= 2014 return 0; } From 822256089e38e087c7e5d255ccba163d793b4f5f Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Tue, 17 Sep 2024 09:44:31 +0200 Subject: [PATCH 06/20] Fix assert handling --- libcudacxx/include/cuda/__cmath/ceil_div.h | 1 + 1 file changed, 1 insertion(+) diff --git a/libcudacxx/include/cuda/__cmath/ceil_div.h b/libcudacxx/include/cuda/__cmath/ceil_div.h index 5687b3370a..6943fd906b 100644 --- a/libcudacxx/include/cuda/__cmath/ceil_div.h +++ b/libcudacxx/include/cuda/__cmath/ceil_div.h @@ -21,6 +21,7 @@ # pragma system_header #endif // no system header +#include #include #include #include From c76e60bb1e141dff80db12ef7b2304506a5cb468 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Tue, 17 Sep 2024 09:24:13 +0000 Subject: [PATCH 07/20] Avoid pointless comparison warnings --- .../include/cuda/std/detail/libcxx/include/span | 11 +++++++---- 1 file changed, 7 insertions(+), 4 deletions(-) diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/span b/libcudacxx/include/cuda/std/detail/libcxx/include/span index 4320120195..7c58623a58 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/span +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/span @@ -660,14 +660,16 @@ public: template _LIBCUDACXX_HIDE_FROM_ABI constexpr span first() const noexcept { - _CCCL_ASSERT(_Count == 0 || _Count <= size(), "span::first(): Count out of range"); + // ternary avoids "pointless comparison of unsigned integer with zero" warning + _CCCL_ASSERT(_Count == 0 ? true : _Count <= size(), "span::first(): Count out of range"); return span{data(), _Count}; } template _LIBCUDACXX_HIDE_FROM_ABI constexpr span last() const noexcept { - _CCCL_ASSERT(_Count == 0 || _Count <= size(), "span::last(): Count out of range"); + // ternary avoids "pointless comparison of unsigned integer with zero" warning + _CCCL_ASSERT(_Count == 0 ? true : _Count <= size(), "span::last(): Count out of range"); return span{data() + size() - _Count, _Count}; } @@ -689,8 +691,9 @@ public: template _LIBCUDACXX_HIDE_FROM_ABI constexpr __subspan_t<_Offset, _Count> subspan() const noexcept { - _CCCL_ASSERT(_Offset == 0 || _Offset <= size(), "span::subspan(): Offset out of range"); - _CCCL_ASSERT(_Count == dynamic_extent || _Count == 0 || _Count <= size() - _Offset, + // ternary avoids "pointless comparison of unsigned integer with zero" warning + _CCCL_ASSERT(_Offset == 0 ? true : _Offset <= size(), "span::subspan(): Offset out of range"); + _CCCL_ASSERT(_Count == dynamic_extent || _Count == 0 ? true : _Count <= size() - _Offset, "span::subspan(): Offset + Count out of range"); return __subspan_t<_Offset, _Count>{data() + _Offset, _Count == dynamic_extent ? size() - _Offset : _Count}; } From b376a14b062ede3b2754f0ff8dff41a8f6d72008 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Tue, 17 Sep 2024 11:50:57 +0000 Subject: [PATCH 08/20] Avoid issues with ICC and warning suppression --- libcudacxx/include/cuda/std/__cccl/assert.h | 21 ++++++--------------- 1 file changed, 6 insertions(+), 15 deletions(-) diff --git a/libcudacxx/include/cuda/std/__cccl/assert.h b/libcudacxx/include/cuda/std/__cccl/assert.h index 92983c1ddc..d00ca6530e 100644 --- a/libcudacxx/include/cuda/std/__cccl/assert.h +++ b/libcudacxx/include/cuda/std/__cccl/assert.h @@ -94,30 +94,21 @@ _LIBCUDACXX_END_NAMESPACE_STD # define _CCCL_ASSERT_DEVICE(expression, message) ((void) 0) #endif // !CCCL_ENABLE_DEVICE_ASSERTIONS -//! Try to work with NVHPC -#if defined(_CCCL_CUDA_COMPILER_NVHPC) -# define _CCCL_ASSERT_IMPL(expression, message) \ - NV_IF_ELSE_TARGET( \ +//! _CCCL_VERIFY is an unconditionally enabled assertion that is reserved for the most important checks +#if defined(_CCCL_CUDA_COMPILER_NVHPC) // NVHPC needs to use NV_IF_TARGET +# define _CCCL_VERIFY(expression, message) \ + NV_IF_ELSE_TARGET( \ NV_IS_DEVICE, (_CCCL_ASSERT_DEVICE(expression, message);), (_CCCL_ASSERT_HOST(expression, message);)) #else // ^^^ _CCCL_CUDA_COMPILER_NVHPC ^^^ / vvv !_CCCL_CUDA_COMPILER_NVHPC vvv # ifdef __CUDA_ARCH__ -# define _CCCL_ASSERT_IMPL(expression, message) \ +# define _CCCL_VERIFY(expression, message) \ _CCCL_NV_DIAG_SUPPRESS(3060) _CCCL_ASSERT_DEVICE(expression, message) _CCCL_NV_DIAG_DEFAULT(3060) # else // ^^^ __CUDA_ARCH__ ^^^ / vvv !__CUDA_ARCH__ vvv -# define _CCCL_ASSERT_IMPL(expression, message) \ +# define _CCCL_VERIFY(expression, message) \ _CCCL_NV_DIAG_SUPPRESS(3060) _CCCL_ASSERT_HOST(expression, message) _CCCL_NV_DIAG_DEFAULT(3060) # endif // !__CUDA_ARCH__ #endif // !_CCCL_CUDA_COMPILER_NVHPC -//! _CCCL_VERIFY is an unconditionally enabled assertion that is reserved for the most important checks -//! Some compilers warn about `is_constant_evaluated()` in a non constexpr function, so silence that right away -#if defined(_CCCL_COMPILER_ICC) -# define _CCCL_VERIFY(expression, message) \ - _CCCL_DIAG_PUSH _CCCL_DIAG_SUPPRESS_ICC(4190 3060) _CCCL_ASSERT_IMPL(expression, message) _CCCL_DIAG_POP -#else // ^^^ _CCCL_COMPILER_ICC ^^^ / vvv !_CCCL_COMPILER_ICC vvv -# define _CCCL_VERIFY(expression, message) _CCCL_ASSERT_IMPL(expression, message) -#endif // !_CCCL_COMPILER_ICC - #if defined(CCCL_ENABLE_HOST_ASSERTIONS) || defined(CCCL_ENABLE_DEVICE_ASSERTIONS) # define _CCCL_ASSERT(expression, message) _CCCL_VERIFY(expression, message) #else // ^^^ CCCL_ENABLE_HOST_ASSERTIONS || CCCL_ENABLE_DEVICE_ASSERTIONS ^^^ / vvv No assertions vvv From 7118890608abb359b6a6f633f6b43f3282ca8a3b Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Tue, 17 Sep 2024 14:10:37 +0200 Subject: [PATCH 09/20] Move the diagnostic suppression to the libstdc++ definition --- libcudacxx/include/cuda/std/__cccl/assert.h | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/libcudacxx/include/cuda/std/__cccl/assert.h b/libcudacxx/include/cuda/std/__cccl/assert.h index d00ca6530e..130862dab9 100644 --- a/libcudacxx/include/cuda/std/__cccl/assert.h +++ b/libcudacxx/include/cuda/std/__cccl/assert.h @@ -71,7 +71,9 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD using ::std::__is_constant_evaluated; _LIBCUDACXX_END_NAMESPACE_STD # endif // _GLIBCXX_RELEASE >= 12 -# define _CCCL_ASSERT_HOST(expression, message) __glibcxx_assert(expression) +# define _CCCL_ASSERT_HOST(expression, message) \ + _CCCL_NV_DIAG_SUPPRESS(3060) __glibcxx_assert(expression); \ + _CCCL_NV_DIAG_DEFAULT(3060) # else // ^^^ libstdc++ ^^^ / vvv Unknown standard library vvv # error "Unknown host standard library used." # endif // Unknown standard library @@ -101,11 +103,9 @@ _LIBCUDACXX_END_NAMESPACE_STD NV_IS_DEVICE, (_CCCL_ASSERT_DEVICE(expression, message);), (_CCCL_ASSERT_HOST(expression, message);)) #else // ^^^ _CCCL_CUDA_COMPILER_NVHPC ^^^ / vvv !_CCCL_CUDA_COMPILER_NVHPC vvv # ifdef __CUDA_ARCH__ -# define _CCCL_VERIFY(expression, message) \ - _CCCL_NV_DIAG_SUPPRESS(3060) _CCCL_ASSERT_DEVICE(expression, message) _CCCL_NV_DIAG_DEFAULT(3060) +# define _CCCL_VERIFY(expression, message) _CCCL_ASSERT_DEVICE(expression, message) # else // ^^^ __CUDA_ARCH__ ^^^ / vvv !__CUDA_ARCH__ vvv -# define _CCCL_VERIFY(expression, message) \ - _CCCL_NV_DIAG_SUPPRESS(3060) _CCCL_ASSERT_HOST(expression, message) _CCCL_NV_DIAG_DEFAULT(3060) +# define _CCCL_VERIFY(expression, message) _CCCL_ASSERT_HOST(expression, message) # endif // !__CUDA_ARCH__ #endif // !_CCCL_CUDA_COMPILER_NVHPC From 62c514e9971acf691075750509eb7bf1b10baf25 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Tue, 17 Sep 2024 14:26:34 +0200 Subject: [PATCH 10/20] Try and make everything work on different compilers --- libcudacxx/include/cuda/std/__cccl/assert.h | 90 ++++++++++++--------- 1 file changed, 51 insertions(+), 39 deletions(-) diff --git a/libcudacxx/include/cuda/std/__cccl/assert.h b/libcudacxx/include/cuda/std/__cccl/assert.h index 130862dab9..ed64e3f751 100644 --- a/libcudacxx/include/cuda/std/__cccl/assert.h +++ b/libcudacxx/include/cuda/std/__cccl/assert.h @@ -54,65 +54,77 @@ #endif // !CCCL_ENABLE_DEVICE_ASSERTIONS //! Use the different standard library implementations to implement host side asserts -#ifdef CCCL_ENABLE_HOST_ASSERTIONS -# if __has_include() // MSVC uses _STL_VERIFY from -# include -# define _CCCL_ASSERT_HOST(expression, message) _STL_VERIFY(expression, message) -# elif __has_include(<__assert>) // libc++ uses _LIBCPP_ASSERT from <__assert> -# include <__assert> -# define _CCCL_ASSERT_HOST(expression, message) _LIBCPP_ASSERT(expression, message) -# elif __has_include() // libstdc++ uses __glibcxx_assert from -# if defined(_GLIBCXX_RELEASE) && _GLIBCXX_RELEASE >= 12 +//! _CCCL_ASSERT_IMPL_HOST should never be used directly +#if __has_include() // MSVC uses _STL_VERIFY from +# include +# define _CCCL_ASSERT_IMPL_HOST(expression, message) _STL_VERIFY(expression, message) +#elif __has_include(<__assert>) // libc++ uses _LIBCPP_ASSERT from <__assert> +# include <__assert> +# define _CCCL_ASSERT_IMPL_HOST(expression, message) _LIBCPP_ASSERT(expression, message) +#elif __has_include() // libstdc++ uses __glibcxx_assert from +# include +# if defined(_GLIBCXX_RELEASE) && _GLIBCXX_RELEASE >= 12 // libstdc++ does not fully qualify its use of `__is_constant_evaluated` // It was introduced in the assert handling in 5e8a30d // libstdc++ : Redefine __glibcxx_assert to work in C++ 23 constexpr _LIBCUDACXX_BEGIN_NAMESPACE_STD using ::std::__is_constant_evaluated; _LIBCUDACXX_END_NAMESPACE_STD -# endif // _GLIBCXX_RELEASE >= 12 -# define _CCCL_ASSERT_HOST(expression, message) \ - _CCCL_NV_DIAG_SUPPRESS(3060) __glibcxx_assert(expression); \ - _CCCL_NV_DIAG_DEFAULT(3060) -# else // ^^^ libstdc++ ^^^ / vvv Unknown standard library vvv -# error "Unknown host standard library used." -# endif // Unknown standard library +# endif // _GLIBCXX_RELEASE >= 12 +//! libstdc++ uses `is_constant_evaluated` in its assert definition which triggers a warning in non constexpr functions +# define _CCCL_ASSERT_IMPL_HOST(expression, message) \ + _CCCL_NV_DIAG_SUPPRESS(3060) __glibcxx_assert(expression); \ + _CCCL_NV_DIAG_DEFAULT(3060) +#else // ^^^ libstdc++ ^^^ / vvv Unknown standard library vvv +# error "Unknown host standard library used." +#endif // Unknown standard library + +//! Use internal nvcc implementation on device or the host library for other cuda compilers +//! _CCCL_ASSERT_IMPL_DEVICE should never be used directly +#if defined(_CCCL_CUDA_COMPILER_NVCC) //! Use __assert_fail to implement device side asserts +# include +# define _CCCL_ASSERT_IMPL_DEVICE(expression, message) \ + __builtin_expect(static_cast(expression), 1) ? (void) 0 : __assert_fail(message, __FILE__, __LINE__, __func__) +#elif defined(_CCCL_CUDA_COMPILER) +# define _CCCL_ASSERT_IMPL_DEVICE(expression, message) _CCCL_ASSERT_IMPL_HOST(expression, message) +#else // ^^^ _CCCL_CUDA_COMPILER ^^^ / vvv !_CCCL_CUDA_COMPILER vvv +# define _CCCL_ASSERT_IMPL_DEVICE(expression, message) ((void) 0) +#endif // !_CCCL_CUDA_COMPILER + +//! _CCCL_ASSERT_HOST is enabled conditionally depending on CCCL_ENABLE_HOST_ASSERTIONS +#ifdef CCCL_ENABLE_HOST_ASSERTIONS +# define _CCCL_ASSERT_HOST(expression, message) _CCCL_ASSERT_IMPL_HOST(expression, message) #else // ^^^ CCCL_ENABLE_HOST_ASSERTIONS ^^^ / vvv !CCCL_ENABLE_HOST_ASSERTIONS vvv # define _CCCL_ASSERT_HOST(expression, message) ((void) 0) #endif // !CCCL_ENABLE_HOST_ASSERTIONS -//! Use internal nvcc implementation on device or the host library for other cuda compilers +//! _CCCL_ASSERT_DEVICE is enabled conditionally depending on CCCL_ENABLE_DEVICE_ASSERTIONS #ifdef CCCL_ENABLE_DEVICE_ASSERTIONS -# if defined(_CCCL_CUDA_COMPILER_NVCC) //! Use __assert_fail to implement device side asserts -# include -# define _CCCL_ASSERT_DEVICE(expression, message) \ - __builtin_expect(static_cast(expression), 1) \ - ? (void) 0 \ - : __assert_fail(message, __FILE__, __LINE__, __func__) -# elif defined(_CCCL_CUDA_COMPILER) -# define _CCCL_ASSERT_DEVICE(expression, message) _CCCL_ASSERT_HOST(expression, message) -# endif // _CCCL_CUDA_COMPILER +# define _CCCL_ASSERT_DEVICE(expression, message) _CCCL_ASSERT_IMPL_DEVICE(expression, message) #else // ^^^ CCCL_ENABLE_DEVICE_ASSERTIONS ^^^ / vvv !CCCL_ENABLE_DEVICE_ASSERTIONS vvv # define _CCCL_ASSERT_DEVICE(expression, message) ((void) 0) #endif // !CCCL_ENABLE_DEVICE_ASSERTIONS -//! _CCCL_VERIFY is an unconditionally enabled assertion that is reserved for the most important checks -#if defined(_CCCL_CUDA_COMPILER_NVHPC) // NVHPC needs to use NV_IF_TARGET +//! _CCCL_VERIFY is enabled unconditionally and reserved for critical checks that are required to always be on +//! _CCCL_ASSERT is enabled conditionally depending on CCCL_ENABLE_HOST_ASSERTIONS and CCCL_ENABLE_DEVICE_ASSERTIONS +#if defined(_CCCL_CUDA_COMPILER_NVHPC) // NVHPC needs to use NV_IF_TARGET instead of __CUDA_ARCH__ # define _CCCL_VERIFY(expression, message) \ + NV_IF_ELSE_TARGET( \ + NV_IS_DEVICE, (_CCCL_ASSERT_IMPL_DEVICE(expression, message);), (_CCCL_ASSERT_IMPL_HOST(expression, message);)) +# define _CCCL_ASSERT(expression, message) \ NV_IF_ELSE_TARGET( \ NV_IS_DEVICE, (_CCCL_ASSERT_DEVICE(expression, message);), (_CCCL_ASSERT_HOST(expression, message);)) -#else // ^^^ _CCCL_CUDA_COMPILER_NVHPC ^^^ / vvv !_CCCL_CUDA_COMPILER_NVHPC vvv +#elif defined(_CCCL_CUDA_COMPILER) # ifdef __CUDA_ARCH__ -# define _CCCL_VERIFY(expression, message) _CCCL_ASSERT_DEVICE(expression, message) +# define _CCCL_VERIFY(expression, message) _CCCL_ASSERT_IMPL_DEVICE(expression, message) +# define _CCCL_ASSERT(expression, message) _CCCL_ASSERT_DEVICE(expression, message) # else // ^^^ __CUDA_ARCH__ ^^^ / vvv !__CUDA_ARCH__ vvv -# define _CCCL_VERIFY(expression, message) _CCCL_ASSERT_HOST(expression, message) +# define _CCCL_VERIFY(expression, message) _CCCL_ASSERT_IMPL_HOST(expression, message) +# define _CCCL_ASSERT(expression, message) _CCCL_ASSERT_HOST(expression, message) # endif // !__CUDA_ARCH__ -#endif // !_CCCL_CUDA_COMPILER_NVHPC - -#if defined(CCCL_ENABLE_HOST_ASSERTIONS) || defined(CCCL_ENABLE_DEVICE_ASSERTIONS) -# define _CCCL_ASSERT(expression, message) _CCCL_VERIFY(expression, message) -#else // ^^^ CCCL_ENABLE_HOST_ASSERTIONS || CCCL_ENABLE_DEVICE_ASSERTIONS ^^^ / vvv No assertions vvv -# define _CCCL_ASSERT(expression, message) ((void) 0) -#endif // No assertions +#else // ^^^ _CCCL_CUDA_COMPILER ^^^ / vvv !_CCCL_CUDA_COMPILER vvv +# define _CCCL_VERIFY(expression, message) _CCCL_ASSERT_IMPL_HOST(expression, message) +# define _CCCL_ASSERT(expression, message) _CCCL_ASSERT_HOST(expression, message) +#endif // !_CCCL_CUDA_COMPILER #endif // __CCCL_ASSERT_H From 24d584327acb42d14e8437003646cb166ea4c183 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Tue, 17 Sep 2024 15:15:19 +0200 Subject: [PATCH 11/20] Improve comment --- libcudacxx/include/cuda/std/__cccl/assert.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libcudacxx/include/cuda/std/__cccl/assert.h b/libcudacxx/include/cuda/std/__cccl/assert.h index ed64e3f751..c54cd235ac 100644 --- a/libcudacxx/include/cuda/std/__cccl/assert.h +++ b/libcudacxx/include/cuda/std/__cccl/assert.h @@ -64,7 +64,7 @@ #elif __has_include() // libstdc++ uses __glibcxx_assert from # include # if defined(_GLIBCXX_RELEASE) && _GLIBCXX_RELEASE >= 12 -// libstdc++ does not fully qualify its use of `__is_constant_evaluated` +// libstdc++ does not fully qualify its use of `__is_constant_evaluated`, so we need to pull it into cuda::std // It was introduced in the assert handling in 5e8a30d // libstdc++ : Redefine __glibcxx_assert to work in C++ 23 constexpr _LIBCUDACXX_BEGIN_NAMESPACE_STD From a9b3234ff2ae2e9f31a12664c39bfc632871cc94 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Wed, 18 Sep 2024 11:33:38 +0200 Subject: [PATCH 12/20] Fix MSVC issues --- libcudacxx/include/cuda/__cccl_config | 1 + libcudacxx/include/cuda/std/__cccl/assert.h | 13 ++++++++++--- 2 files changed, 11 insertions(+), 3 deletions(-) diff --git a/libcudacxx/include/cuda/__cccl_config b/libcudacxx/include/cuda/__cccl_config index 241f4f8ae2..15f616e0e2 100644 --- a/libcudacxx/include/cuda/__cccl_config +++ b/libcudacxx/include/cuda/__cccl_config @@ -11,6 +11,7 @@ #ifndef _CUDA__CCCL_CONFIG #define _CUDA__CCCL_CONFIG +#include // IWYU pragma: export #include // IWYU pragma: export #include // IWYU pragma: export #include // IWYU pragma: export diff --git a/libcudacxx/include/cuda/std/__cccl/assert.h b/libcudacxx/include/cuda/std/__cccl/assert.h index c54cd235ac..e7db0939d8 100644 --- a/libcudacxx/include/cuda/std/__cccl/assert.h +++ b/libcudacxx/include/cuda/std/__cccl/assert.h @@ -79,12 +79,19 @@ _LIBCUDACXX_END_NAMESPACE_STD # error "Unknown host standard library used." #endif // Unknown standard library -//! Use internal nvcc implementation on device or the host library for other cuda compilers +//! Use custom implementations with nvcc on device and the host ones with clang-cuda and nvhpc //! _CCCL_ASSERT_IMPL_DEVICE should never be used directly #if defined(_CCCL_CUDA_COMPILER_NVCC) //! Use __assert_fail to implement device side asserts # include -# define _CCCL_ASSERT_IMPL_DEVICE(expression, message) \ - __builtin_expect(static_cast(expression), 1) ? (void) 0 : __assert_fail(message, __FILE__, __LINE__, __func__) +# if defined(_CCCL_COMPILER_MSVC) +# define _CCCL_ASSERT_IMPL_DEVICE(expression, message) \ + __builtin_expect(static_cast(expression), 1) ? (void) 0 : _wassert(_CRT_WIDE(#message), __FILEW__, __LINE__) +# else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv +# define _CCCL_ASSERT_IMPL_DEVICE(expression, message) \ + __builtin_expect(static_cast(expression), 1) \ + ? (void) 0 \ + : __assert_fail(message, __FILE__, __LINE__, __func__) +# endif // !_CCCL_COMPILER_MSVC #elif defined(_CCCL_CUDA_COMPILER) # define _CCCL_ASSERT_IMPL_DEVICE(expression, message) _CCCL_ASSERT_IMPL_HOST(expression, message) #else // ^^^ _CCCL_CUDA_COMPILER ^^^ / vvv !_CCCL_CUDA_COMPILER vvv From 2791d144700cd85b4e980028afdd66705f7889ec Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Wed, 18 Sep 2024 11:56:50 +0200 Subject: [PATCH 13/20] Add tests --- .../libcxx/asserts/assert_device.runfail.cpp | 26 ++++++++++++++++++ .../asserts/assert_device_disabled.pass.cpp | 25 +++++++++++++++++ .../assert_device_enable_device.runfail.cpp | 26 ++++++++++++++++++ .../assert_device_enable_host.pass.cpp | 26 ++++++++++++++++++ .../libcxx/asserts/assert_host.runfail.cpp | 27 +++++++++++++++++++ .../asserts/assert_host_disable.pass.cpp | 25 +++++++++++++++++ .../assert_host_enable_device.pass.cpp | 26 ++++++++++++++++++ .../assert_host_enable_host.runfail.cpp | 26 ++++++++++++++++++ .../verify_device_disabled.runfail.cpp | 25 +++++++++++++++++ .../asserts/verify_host_disable.runfail.cpp | 25 +++++++++++++++++ 10 files changed, 257 insertions(+) create mode 100644 libcudacxx/test/libcudacxx/libcxx/asserts/assert_device.runfail.cpp create mode 100644 libcudacxx/test/libcudacxx/libcxx/asserts/assert_device_disabled.pass.cpp create mode 100644 libcudacxx/test/libcudacxx/libcxx/asserts/assert_device_enable_device.runfail.cpp create mode 100644 libcudacxx/test/libcudacxx/libcxx/asserts/assert_device_enable_host.pass.cpp create mode 100644 libcudacxx/test/libcudacxx/libcxx/asserts/assert_host.runfail.cpp create mode 100644 libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_disable.pass.cpp create mode 100644 libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_enable_device.pass.cpp create mode 100644 libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_enable_host.runfail.cpp create mode 100644 libcudacxx/test/libcudacxx/libcxx/asserts/verify_device_disabled.runfail.cpp create mode 100644 libcudacxx/test/libcudacxx/libcxx/asserts/verify_host_disable.runfail.cpp diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device.runfail.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device.runfail.cpp new file mode 100644 index 0000000000..c2db432788 --- /dev/null +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device.runfail.cpp @@ -0,0 +1,26 @@ +//===----------------------------------------------------------------------===// +// +// 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. +// +//===----------------------------------------------------------------------===// + +// We compile with CCCL_ENABLE_ASSERTIONS +#ifndef CCCL_ENABLE_ASSERTIONS +# error "Should be compiled with CCCL_ENABLE_ASSERTIONS +#endif // !CCCL_ENABLE_ASSERTIONS +#include + +__host__ __device__ bool failed_on_device() +{ + NV_IF_ELSE_TARGET(NV_IS_DEVICE, return false;, return true;) +} + +int main(int, char**) +{ + _CCCL_ASSERT(failed_on_device(), "Should fail on device"); + return 0; +} diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device_disabled.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device_disabled.pass.cpp new file mode 100644 index 0000000000..d10bde6f0a --- /dev/null +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device_disabled.pass.cpp @@ -0,0 +1,25 @@ +//===----------------------------------------------------------------------===// +// +// 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. +// +//===----------------------------------------------------------------------===// + +// We compile with CCCL_ENABLE_ASSERTIONS, but want to disable assertions +#undef CCCL_ENABLE_ASSERTIONS + +#include + +__host__ __device__ bool failed_on_device() +{ + NV_IF_ELSE_TARGET(NV_IS_DEVICE, return false;, return true;) +} + +int main(int, char**) +{ + _CCCL_ASSERT(failed_on_device(), "Should fail on device"); + return 0; +} diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device_enable_device.runfail.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device_enable_device.runfail.cpp new file mode 100644 index 0000000000..3cf7d6f7e0 --- /dev/null +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device_enable_device.runfail.cpp @@ -0,0 +1,26 @@ +//===----------------------------------------------------------------------===// +// +// 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. +// +//===----------------------------------------------------------------------===// + +// We compile with CCCL_ENABLE_ASSERTIONS, but want to enable only device assertions +#undef CCCL_ENABLE_ASSERTIONS +#define CCCL_ENABLE_DEVICE_ASSERTIONS + +#include + +__host__ __device__ bool failed_on_device() +{ + NV_IF_ELSE_TARGET(NV_IS_DEVICE, return false;, return true;) +} + +int main(int, char**) +{ + _CCCL_ASSERT(failed_on_device(), "Should fail on device"); + return 0; +} diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device_enable_host.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device_enable_host.pass.cpp new file mode 100644 index 0000000000..612e42bcbf --- /dev/null +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device_enable_host.pass.cpp @@ -0,0 +1,26 @@ +//===----------------------------------------------------------------------===// +// +// 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. +// +//===----------------------------------------------------------------------===// + +// We compile with CCCL_ENABLE_ASSERTIONS, but want to enable only host assertions +#undef CCCL_ENABLE_ASSERTIONS +#define CCCL_ENABLE_HOST_ASSERTIONS + +#include + +__host__ __device__ bool failed_on_device() +{ + NV_IF_ELSE_TARGET(NV_IS_DEVICE, return false;, return true;) +} + +int main(int, char**) +{ + _CCCL_ASSERT(failed_on_device(), "Should fail on device"); + return 0; +} diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host.runfail.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host.runfail.cpp new file mode 100644 index 0000000000..e0d9e37618 --- /dev/null +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host.runfail.cpp @@ -0,0 +1,27 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// We compile with CCCL_ENABLE_ASSERTIONS +#ifndef CCCL_ENABLE_ASSERTIONS +# error "Should be compiled with CCCL_ENABLE_ASSERTIONS +#endif // !CCCL_ENABLE_ASSERTIONS + +#include + +__host__ __device__ bool failed_on_host() +{ + NV_IF_ELSE_TARGET(NV_IS_DEVICE, return true;, return false;) +} + +int main(int, char**) +{ + _CCCL_ASSERT(failed_on_host(), "Should fail on host"); + return 0; +} diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_disable.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_disable.pass.cpp new file mode 100644 index 0000000000..fcc257619a --- /dev/null +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_disable.pass.cpp @@ -0,0 +1,25 @@ +//===----------------------------------------------------------------------===// +// +// 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. +// +//===----------------------------------------------------------------------===// + +// We compile with CCCL_ENABLE_ASSERTIONS, but want to disable assertions +#undef CCCL_ENABLE_ASSERTIONS + +#include + +__host__ __device__ bool failed_on_host() +{ + NV_IF_ELSE_TARGET(NV_IS_DEVICE, return true;, return false;) +} + +int main(int, char**) +{ + _CCCL_ASSERT(failed_on_host(), "Should fail on host"); + return 0; +} diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_enable_device.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_enable_device.pass.cpp new file mode 100644 index 0000000000..293d194b8d --- /dev/null +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_enable_device.pass.cpp @@ -0,0 +1,26 @@ +//===----------------------------------------------------------------------===// +// +// 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. +// +//===----------------------------------------------------------------------===// + +// We compile with CCCL_ENABLE_ASSERTIONS, but want to enable only device assertions +#undef CCCL_ENABLE_ASSERTIONS +#define CCCL_ENABLE_DEVICE_ASSERTIONS + +#include + +__host__ __device__ bool failed_on_host() +{ + NV_IF_ELSE_TARGET(NV_IS_DEVICE, return true;, return false;) +} + +int main(int, char**) +{ + _CCCL_ASSERT(failed_on_host(), "Should fail on host"); + return 0; +} diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_enable_host.runfail.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_enable_host.runfail.cpp new file mode 100644 index 0000000000..3b6b146cbe --- /dev/null +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_enable_host.runfail.cpp @@ -0,0 +1,26 @@ +//===----------------------------------------------------------------------===// +// +// 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. +// +//===----------------------------------------------------------------------===// + +// We compile with CCCL_ENABLE_ASSERTIONS, but want to enable only host assertions +#undef CCCL_ENABLE_ASSERTIONS +#define CCCL_ENABLE_HOST_ASSERTIONS + +#include + +__host__ __device__ bool failed_on_host() +{ + NV_IF_ELSE_TARGET(NV_IS_DEVICE, return true;, return false;) +} + +int main(int, char**) +{ + _CCCL_ASSERT(failed_on_host(), "Should fail on host"); + return 0; +} diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/verify_device_disabled.runfail.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/verify_device_disabled.runfail.cpp new file mode 100644 index 0000000000..2003ff6310 --- /dev/null +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/verify_device_disabled.runfail.cpp @@ -0,0 +1,25 @@ +//===----------------------------------------------------------------------===// +// +// 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. +// +//===----------------------------------------------------------------------===// + +// _CCCL_VERIFY is always on +#undef CCCL_ENABLE_ASSERTIONS + +#include + +__host__ __device__ bool failed_on_device() +{ + NV_IF_ELSE_TARGET(NV_IS_DEVICE, return false;, return true;) +} + +int main(int, char**) +{ + _CCCL_VERIFY(failed_on_device(), "Should fail on device"); + return 0; +} diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/verify_host_disable.runfail.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/verify_host_disable.runfail.cpp new file mode 100644 index 0000000000..60bd8b1488 --- /dev/null +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/verify_host_disable.runfail.cpp @@ -0,0 +1,25 @@ +//===----------------------------------------------------------------------===// +// +// 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. +// +//===----------------------------------------------------------------------===// + +// _CCCL_VERIFY is always on +#undef CCCL_ENABLE_ASSERTIONS + +#include + +__host__ __device__ bool failed_on_host() +{ + NV_IF_ELSE_TARGET(NV_IS_DEVICE, return true;, return false;) +} + +int main(int, char**) +{ + _CCCL_VERIFY(failed_on_host(), "Should fail on host"); + return 0; +} From 2711e9eeed7ca5bc2d619e427a22e254c0d507e8 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Wed, 18 Sep 2024 13:08:40 +0200 Subject: [PATCH 14/20] Do not include it in __cccl_config yet We need to pull in some more infrastructure for that --- libcudacxx/include/cuda/__cccl_config | 1 - libcudacxx/include/cuda/std/cassert | 2 ++ 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/libcudacxx/include/cuda/__cccl_config b/libcudacxx/include/cuda/__cccl_config index 15f616e0e2..241f4f8ae2 100644 --- a/libcudacxx/include/cuda/__cccl_config +++ b/libcudacxx/include/cuda/__cccl_config @@ -11,7 +11,6 @@ #ifndef _CUDA__CCCL_CONFIG #define _CUDA__CCCL_CONFIG -#include // IWYU pragma: export #include // IWYU pragma: export #include // IWYU pragma: export #include // IWYU pragma: export diff --git a/libcudacxx/include/cuda/std/cassert b/libcudacxx/include/cuda/std/cassert index fea5e74b8c..b6364f2947 100644 --- a/libcudacxx/include/cuda/std/cassert +++ b/libcudacxx/include/cuda/std/cassert @@ -27,4 +27,6 @@ # include #endif //_CCCL_COMPILER_NVRTC +#include + #endif // _CUDA_STD_CASSERT From ed5a7dc5df59d2f1e3aff082e54318e42d7f2c64 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Thu, 19 Sep 2024 12:48:50 +0200 Subject: [PATCH 15/20] runfail tests need to be executed --- .../test/libcudacxx/libcxx/asserts/assert_device.runfail.cpp | 2 ++ .../libcxx/asserts/assert_device_enable_device.runfail.cpp | 2 ++ .../test/libcudacxx/libcxx/asserts/assert_host.runfail.cpp | 2 ++ .../libcxx/asserts/assert_host_enable_host.runfail.cpp | 2 ++ .../libcxx/asserts/verify_device_disabled.runfail.cpp | 2 ++ .../libcudacxx/libcxx/asserts/verify_host_disable.runfail.cpp | 2 ++ 6 files changed, 12 insertions(+) diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device.runfail.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device.runfail.cpp index c2db432788..6f5f3cea56 100644 --- a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device.runfail.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device.runfail.cpp @@ -8,6 +8,8 @@ // //===----------------------------------------------------------------------===// +// UNSUPPORTED: no_execute + // We compile with CCCL_ENABLE_ASSERTIONS #ifndef CCCL_ENABLE_ASSERTIONS # error "Should be compiled with CCCL_ENABLE_ASSERTIONS diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device_enable_device.runfail.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device_enable_device.runfail.cpp index 3cf7d6f7e0..3162822941 100644 --- a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device_enable_device.runfail.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device_enable_device.runfail.cpp @@ -8,6 +8,8 @@ // //===----------------------------------------------------------------------===// +// UNSUPPORTED: no_execute + // We compile with CCCL_ENABLE_ASSERTIONS, but want to enable only device assertions #undef CCCL_ENABLE_ASSERTIONS #define CCCL_ENABLE_DEVICE_ASSERTIONS diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host.runfail.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host.runfail.cpp index e0d9e37618..ef2c72eaa7 100644 --- a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host.runfail.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host.runfail.cpp @@ -8,6 +8,8 @@ // //===----------------------------------------------------------------------===// +// UNSUPPORTED: no_execute + // We compile with CCCL_ENABLE_ASSERTIONS #ifndef CCCL_ENABLE_ASSERTIONS # error "Should be compiled with CCCL_ENABLE_ASSERTIONS diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_enable_host.runfail.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_enable_host.runfail.cpp index 3b6b146cbe..52522b779f 100644 --- a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_enable_host.runfail.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_enable_host.runfail.cpp @@ -8,6 +8,8 @@ // //===----------------------------------------------------------------------===// +// UNSUPPORTED: no_execute + // We compile with CCCL_ENABLE_ASSERTIONS, but want to enable only host assertions #undef CCCL_ENABLE_ASSERTIONS #define CCCL_ENABLE_HOST_ASSERTIONS diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/verify_device_disabled.runfail.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/verify_device_disabled.runfail.cpp index 2003ff6310..6039bec59a 100644 --- a/libcudacxx/test/libcudacxx/libcxx/asserts/verify_device_disabled.runfail.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/verify_device_disabled.runfail.cpp @@ -8,6 +8,8 @@ // //===----------------------------------------------------------------------===// +// UNSUPPORTED: no_execute + // _CCCL_VERIFY is always on #undef CCCL_ENABLE_ASSERTIONS diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/verify_host_disable.runfail.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/verify_host_disable.runfail.cpp index 60bd8b1488..6f85d65a6b 100644 --- a/libcudacxx/test/libcudacxx/libcxx/asserts/verify_host_disable.runfail.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/verify_host_disable.runfail.cpp @@ -8,6 +8,8 @@ // //===----------------------------------------------------------------------===// +// UNSUPPORTED: no_execute + // _CCCL_VERIFY is always on #undef CCCL_ENABLE_ASSERTIONS From 8d413a819ab96735dcfdea1b0ead34d93630b4bd Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Thu, 19 Sep 2024 14:10:44 +0200 Subject: [PATCH 16/20] No host with nvrtc --- libcudacxx/include/cuda/std/__cccl/assert.h | 2 ++ 1 file changed, 2 insertions(+) diff --git a/libcudacxx/include/cuda/std/__cccl/assert.h b/libcudacxx/include/cuda/std/__cccl/assert.h index e7db0939d8..bf81d4056e 100644 --- a/libcudacxx/include/cuda/std/__cccl/assert.h +++ b/libcudacxx/include/cuda/std/__cccl/assert.h @@ -75,6 +75,8 @@ _LIBCUDACXX_END_NAMESPACE_STD # define _CCCL_ASSERT_IMPL_HOST(expression, message) \ _CCCL_NV_DIAG_SUPPRESS(3060) __glibcxx_assert(expression); \ _CCCL_NV_DIAG_DEFAULT(3060) +#elif defined(_CCCL_COMPILER_NVRTC) // There is no host standard library in nvrtc +# define _CCCL_ASSERT_IMPL_HOST(expression, message) ((void) 0) #else // ^^^ libstdc++ ^^^ / vvv Unknown standard library vvv # error "Unknown host standard library used." #endif // Unknown standard library From f492b166eccc5093e40bc97461ee7717d4a543d0 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Thu, 19 Sep 2024 14:35:26 +0200 Subject: [PATCH 17/20] Fix the msvc header include --- libcudacxx/include/cuda/std/__cccl/assert.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/libcudacxx/include/cuda/std/__cccl/assert.h b/libcudacxx/include/cuda/std/__cccl/assert.h index bf81d4056e..026e61f391 100644 --- a/libcudacxx/include/cuda/std/__cccl/assert.h +++ b/libcudacxx/include/cuda/std/__cccl/assert.h @@ -55,8 +55,8 @@ //! Use the different standard library implementations to implement host side asserts //! _CCCL_ASSERT_IMPL_HOST should never be used directly -#if __has_include() // MSVC uses _STL_VERIFY from -# include +#if __has_include() // MSVC uses _STL_VERIFY from +# include # define _CCCL_ASSERT_IMPL_HOST(expression, message) _STL_VERIFY(expression, message) #elif __has_include(<__assert>) // libc++ uses _LIBCPP_ASSERT from <__assert> # include <__assert> From 03d9f5e4f93a626adce028d51ffa51cd968a5934 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Thu, 19 Sep 2024 20:34:57 +0200 Subject: [PATCH 18/20] Fix typo --- .../test/libcudacxx/libcxx/asserts/assert_device.runfail.cpp | 2 +- .../test/libcudacxx/libcxx/asserts/assert_host.runfail.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device.runfail.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device.runfail.cpp index 6f5f3cea56..5ab7b52693 100644 --- a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device.runfail.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_device.runfail.cpp @@ -12,7 +12,7 @@ // We compile with CCCL_ENABLE_ASSERTIONS #ifndef CCCL_ENABLE_ASSERTIONS -# error "Should be compiled with CCCL_ENABLE_ASSERTIONS +# error "Should be compiled with CCCL_ENABLE_ASSERTIONS" #endif // !CCCL_ENABLE_ASSERTIONS #include diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host.runfail.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host.runfail.cpp index ef2c72eaa7..381fce3d43 100644 --- a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host.runfail.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host.runfail.cpp @@ -12,7 +12,7 @@ // We compile with CCCL_ENABLE_ASSERTIONS #ifndef CCCL_ENABLE_ASSERTIONS -# error "Should be compiled with CCCL_ENABLE_ASSERTIONS +# error "Should be compiled with CCCL_ENABLE_ASSERTIONS" #endif // !CCCL_ENABLE_ASSERTIONS #include From 388aff2f3dc12bde7e9bf86ec38230560c5dbd70 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Thu, 19 Sep 2024 20:38:06 +0200 Subject: [PATCH 19/20] Host asserts are not meant for nvrtc --- .../test/libcudacxx/libcxx/asserts/assert_host.runfail.cpp | 1 + .../test/libcudacxx/libcxx/asserts/assert_host_disable.pass.cpp | 2 ++ .../libcxx/asserts/assert_host_enable_device.pass.cpp | 2 ++ .../libcxx/asserts/assert_host_enable_host.runfail.cpp | 1 + .../libcudacxx/libcxx/asserts/verify_host_disable.runfail.cpp | 1 + 5 files changed, 7 insertions(+) diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host.runfail.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host.runfail.cpp index 381fce3d43..8551a23038 100644 --- a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host.runfail.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host.runfail.cpp @@ -9,6 +9,7 @@ //===----------------------------------------------------------------------===// // UNSUPPORTED: no_execute +// UNSUPPORTED: nvrtc // We compile with CCCL_ENABLE_ASSERTIONS #ifndef CCCL_ENABLE_ASSERTIONS diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_disable.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_disable.pass.cpp index fcc257619a..ea630f7bd0 100644 --- a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_disable.pass.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_disable.pass.cpp @@ -8,6 +8,8 @@ // //===----------------------------------------------------------------------===// +// UNSUPPORTED: nvrtc + // We compile with CCCL_ENABLE_ASSERTIONS, but want to disable assertions #undef CCCL_ENABLE_ASSERTIONS diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_enable_device.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_enable_device.pass.cpp index 293d194b8d..e246056e9b 100644 --- a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_enable_device.pass.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_enable_device.pass.cpp @@ -8,6 +8,8 @@ // //===----------------------------------------------------------------------===// +// UNSUPPORTED: nvrtc + // We compile with CCCL_ENABLE_ASSERTIONS, but want to enable only device assertions #undef CCCL_ENABLE_ASSERTIONS #define CCCL_ENABLE_DEVICE_ASSERTIONS diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_enable_host.runfail.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_enable_host.runfail.cpp index 52522b779f..27f04913e9 100644 --- a/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_enable_host.runfail.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/assert_host_enable_host.runfail.cpp @@ -9,6 +9,7 @@ //===----------------------------------------------------------------------===// // UNSUPPORTED: no_execute +// UNSUPPORTED: nvrtc // We compile with CCCL_ENABLE_ASSERTIONS, but want to enable only host assertions #undef CCCL_ENABLE_ASSERTIONS diff --git a/libcudacxx/test/libcudacxx/libcxx/asserts/verify_host_disable.runfail.cpp b/libcudacxx/test/libcudacxx/libcxx/asserts/verify_host_disable.runfail.cpp index 6f85d65a6b..d08ce0a454 100644 --- a/libcudacxx/test/libcudacxx/libcxx/asserts/verify_host_disable.runfail.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/asserts/verify_host_disable.runfail.cpp @@ -9,6 +9,7 @@ //===----------------------------------------------------------------------===// // UNSUPPORTED: no_execute +// UNSUPPORTED: nvrtc // _CCCL_VERIFY is always on #undef CCCL_ENABLE_ASSERTIONS From fa5ebe70ca0f102ee567745561d0f958b639c320 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Thu, 19 Sep 2024 20:40:22 +0200 Subject: [PATCH 20/20] Disable constexpr tests for MSVC2017 due to constant folding --- .../alg.sort/partial.sort.copy/partial_sort_copy.pass.cpp | 4 ++-- .../partial.sort.copy/partial_sort_copy_comp.pass.cpp | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.sort/partial.sort.copy/partial_sort_copy.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.sort/partial.sort.copy/partial_sort_copy.pass.cpp index f2302f5f43..d3217b85ba 100644 --- a/libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.sort/partial.sort.copy/partial_sort_copy.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.sort/partial.sort.copy/partial_sort_copy.pass.cpp @@ -99,9 +99,9 @@ __host__ __device__ TEST_CONSTEXPR_CXX14 bool test() int main(int, char**) { test(); -#if TEST_STD_VER >= 2014 +#if TEST_STD_VER >= 2014 && !defined(TEST_COMPILER_MSVC_2017) static_assert(test(), ""); -#endif // TEST_STD_VER >= 2014 +#endif // TEST_STD_VER >= 2014 && ! TEST_COMPILER_MSVC_2017 return 0; } diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.sort/partial.sort.copy/partial_sort_copy_comp.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.sort/partial.sort.copy/partial_sort_copy_comp.pass.cpp index a8a9b010fa..45bced305b 100644 --- a/libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.sort/partial.sort.copy/partial_sort_copy_comp.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.sort/partial.sort.copy/partial_sort_copy_comp.pass.cpp @@ -105,9 +105,9 @@ __host__ __device__ TEST_CONSTEXPR_CXX14 bool test() int main(int, char**) { test(); -#if TEST_STD_VER >= 2014 +#if TEST_STD_VER >= 2014 && !defined(TEST_COMPILER_MSVC_2017) static_assert(test(), ""); -#endif // TEST_STD_VER >= 2014 +#endif // TEST_STD_VER >= 2014 && ! TEST_COMPILER_MSVC_2017 return 0; }