Skip to content

Commit

Permalink
Revert to _LIBCUDACXX_DEBUG_ASSERT
Browse files Browse the repository at this point in the history
  • Loading branch information
@ahendriksen
ahendriksen committed Aug 31, 2023
1 parent ba0b89e commit ec483b2
Show file tree
Hide file tree
Showing 2 changed files with 18 additions and 18 deletions.
12 changes: 6 additions & 6 deletions libcudacxx/include/cuda/std/detail/libcxx/include/__cuda/barrier.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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);
}
};
Expand Down Expand Up @@ -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, (
Expand Down Expand Up @@ -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, (
Expand Down
24 changes: 12 additions & 12 deletions libcudacxx/include/cuda/std/detail/libcxx/include/barrier
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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();
Expand All @@ -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
Expand Down Expand Up @@ -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);
Expand Down Expand Up @@ -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;
Expand All @@ -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
Expand Down

0 comments on commit ec483b2

Please sign in to comment.