From ec483b274fea3c4d1afe5693c11143cabc1eb553 Mon Sep 17 00:00:00 2001 From: Allard Hendriksen Date: Thu, 31 Aug 2023 15:06:22 +0200 Subject: [PATCH] Revert to _LIBCUDACXX_DEBUG_ASSERT --- .../detail/libcxx/include/__cuda/barrier.h | 12 +++++----- .../cuda/std/detail/libcxx/include/barrier | 24 +++++++++---------- 2 files changed, 18 insertions(+), 18 deletions(-) diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__cuda/barrier.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__cuda/barrier.h index f4284bd3632..21920dc5116 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/__cuda/barrier.h +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__cuda/barrier.h @@ -67,13 +67,13 @@ class barrier : public _CUDA_VSTD::__barrier_base<_CompletionF, _Sco> { _LIBCUDACXX_INLINE_VISIBILITY friend void init(barrier * __b, _CUDA_VSTD::ptrdiff_t __expected) { - _LIBCUDACXX_ASSERT(__expected >= 0, "Cannot initialize barrier with negative arrival count"); + _LIBCUDACXX_DEBUG_ASSERT(__expected >= 0, "Cannot initialize barrier with negative arrival count"); new (__b) barrier(__expected); } _LIBCUDACXX_INLINE_VISIBILITY friend void init(barrier * __b, _CUDA_VSTD::ptrdiff_t __expected, _CompletionF __completion) { - _LIBCUDACXX_ASSERT(__expected >= 0, "Cannot initialize barrier with negative arrival count"); + _LIBCUDACXX_DEBUG_ASSERT(__expected >= 0, "Cannot initialize barrier with negative arrival count"); new (__b) barrier(__expected, __completion); } }; @@ -178,10 +178,10 @@ friend class _CUDA_VSTD::__barrier_poll_tester_parity; _LIBCUDACXX_NODISCARD_ATTRIBUTE _LIBCUDACXX_INLINE_VISIBILITY arrival_token arrive_tx(_CUDA_VSTD::ptrdiff_t __arrive_count_update, _CUDA_VSTD::ptrdiff_t __transaction_count_update) { - _LIBCUDACXX_ASSERT(__arrive_count_update >= 0, "Arrival count update must be non-negative."); - _LIBCUDACXX_ASSERT(__transaction_count_update >= 0, "Transaction count update must be non-negative."); + _LIBCUDACXX_DEBUG_ASSERT(__arrive_count_update >= 0, "Arrival count update must be non-negative."); + _LIBCUDACXX_DEBUG_ASSERT(__transaction_count_update >= 0, "Transaction count update must be non-negative."); // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#contents-of-the-mbarrier-object - _LIBCUDACXX_ASSERT(__transaction_count_update <= (1 << 20) - 1, "Transaction count update cannot exceed 2^20 - 1."); + _LIBCUDACXX_DEBUG_ASSERT(__transaction_count_update <= (1 << 20) - 1, "Transaction count update cannot exceed 2^20 - 1."); arrival_token __token = {}; NV_DISPATCH_TARGET( NV_PROVIDES_SM_90, ( @@ -217,7 +217,7 @@ friend class _CUDA_VSTD::__barrier_poll_tester_parity; _LIBCUDACXX_NODISCARD_ATTRIBUTE _LIBCUDACXX_INLINE_VISIBILITY arrival_token arrive(_CUDA_VSTD::ptrdiff_t __update = 1) { - _LIBCUDACXX_ASSERT(__update >= 0, "Arrival count update must be non-negative."); + _LIBCUDACXX_DEBUG_ASSERT(__update >= 0, "Arrival count update must be non-negative."); arrival_token __token = {}; NV_DISPATCH_TARGET( NV_PROVIDES_SM_90, ( diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/barrier b/libcudacxx/include/cuda/std/detail/libcxx/include/barrier index f158fbbb020..fc3424d4efb 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/barrier +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/barrier @@ -190,10 +190,10 @@ public: _LIBCUDACXX_NODISCARD_ATTRIBUTE _LIBCUDACXX_INLINE_VISIBILITY arrival_token arrive_tx(_CUDA_VSTD::ptrdiff_t __arrive_count_update, _CUDA_VSTD::ptrdiff_t __transaction_count_update) { - _LIBCUDACXX_ASSERT(__arrive_count_update >= 0, "Arrival count update must be non-negative."); - _LIBCUDACXX_ASSERT(__transaction_count_update >= 0, "Transaction count update must be non-negative."); + _LIBCUDACXX_DEBUG_ASSERT(__arrive_count_update >= 0, "Arrival count update must be non-negative."); + _LIBCUDACXX_DEBUG_ASSERT(__transaction_count_update >= 0, "Transaction count update must be non-negative."); // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#contents-of-the-mbarrier-object - _LIBCUDACXX_ASSERT( + _LIBCUDACXX_DEBUG_ASSERT( __transaction_count_update <= (1 << 20) - 1, "Transaction count update cannot exceed 2^20 - 1."); // Only the thread-block scope specialization can keep track of the @@ -339,7 +339,7 @@ public: auto const __result = __arrived.fetch_sub(__update, memory_order_acq_rel) - __update; auto const __new_expected = __expected.load(memory_order_relaxed); - _LIBCUDACXX_ASSERT(__result >= 0, ""); + _LIBCUDACXX_DEBUG_ASSERT(__result >= 0, ""); if(0 == __result) { __completion(); @@ -352,10 +352,10 @@ public: _LIBCUDACXX_NODISCARD_ATTRIBUTE _LIBCUDACXX_INLINE_VISIBILITY arrival_token arrive_tx(_CUDA_VSTD::ptrdiff_t __arrive_count_update, _CUDA_VSTD::ptrdiff_t __transaction_count_update) { - _LIBCUDACXX_ASSERT(__arrive_count_update >= 0, "Arrival count update must be non-negative."); - _LIBCUDACXX_ASSERT(__transaction_count_update >= 0, "Transaction count update must be non-negative."); + _LIBCUDACXX_DEBUG_ASSERT(__arrive_count_update >= 0, "Arrival count update must be non-negative."); + _LIBCUDACXX_DEBUG_ASSERT(__transaction_count_update >= 0, "Transaction count update must be non-negative."); // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#contents-of-the-mbarrier-object - _LIBCUDACXX_ASSERT( + _LIBCUDACXX_DEBUG_ASSERT( __transaction_count_update <= (1 << 20) - 1, "Transaction count update cannot exceed 2^20 - 1."); // Only the thread-block scope specialization can keep track of the @@ -418,7 +418,7 @@ private: static _LIBCUDACXX_INLINE_VISIBILITY constexpr uint64_t __init(ptrdiff_t __count) noexcept { - _LIBCUDACXX_ASSERT(__count >= 0, ""); + _LIBCUDACXX_DEBUG_ASSERT(__count >= 0, ""); return (((1u << 31) - __count) << 32) | ((1u << 31) - __count); @@ -446,7 +446,7 @@ public: _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 __barrier_base(ptrdiff_t __count, __empty_completion = __empty_completion()) : __phase_arrived_expected(__init(__count)) { - _LIBCUDACXX_ASSERT(__count >= 0, ""); + _LIBCUDACXX_DEBUG_ASSERT(__count >= 0, ""); } ~__barrier_base() = default; @@ -468,10 +468,10 @@ public: _LIBCUDACXX_NODISCARD_ATTRIBUTE _LIBCUDACXX_INLINE_VISIBILITY arrival_token arrive_tx(_CUDA_VSTD::ptrdiff_t __arrive_count_update, _CUDA_VSTD::ptrdiff_t __transaction_count_update) { - _LIBCUDACXX_ASSERT(__arrive_count_update >= 0, "Arrival count update must be non-negative."); - _LIBCUDACXX_ASSERT(__transaction_count_update >= 0, "Transaction count update must be non-negative."); + _LIBCUDACXX_DEBUG_ASSERT(__arrive_count_update >= 0, "Arrival count update must be non-negative."); + _LIBCUDACXX_DEBUG_ASSERT(__transaction_count_update >= 0, "Transaction count update must be non-negative."); // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#contents-of-the-mbarrier-object - _LIBCUDACXX_ASSERT( + _LIBCUDACXX_DEBUG_ASSERT( __transaction_count_update <= (1 << 20) - 1, "Transaction count update cannot exceed 2^20 - 1."); // Only the thread-block scope specialization can keep track of the