Skip to content

Commit

Permalink
Allow infinite waiting (#200)
Browse files Browse the repository at this point in the history
chhwang authored Oct 23, 2023

Verified

This commit was created on GitHub.com and signed with GitHub’s verified signature. The key has expired.
1 parent 6f43282 commit 7686e15
Showing 2 changed files with 4 additions and 4 deletions.
6 changes: 3 additions & 3 deletions include/mscclpp/poll.hpp
Original file line number Diff line number Diff line change
@@ -17,7 +17,7 @@ extern "C" __device__ void __assert_fail(const char *__assertion, const char *__
int64_t __spin_cnt = 0; \
__status = 0; \
while (__cond) { \
if (__spin_cnt++ == __max_spin_cnt) { \
if (__max_spin_cnt >= 0 && __spin_cnt++ == __max_spin_cnt) { \
__status = 1; \
break; \
} \
@@ -29,7 +29,7 @@ extern "C" __device__ void __assert_fail(const char *__assertion, const char *__
do { \
int64_t __spin_cnt = 0; \
while (__cond) { \
if (__spin_cnt++ == __max_spin_cnt) { \
if (__max_spin_cnt >= 0 && __spin_cnt++ == __max_spin_cnt) { \
__assert_fail(#__cond, __FILE__, __LINE__, __PRETTY_FUNCTION__); \
} \
} \
@@ -46,7 +46,7 @@ extern "C" __device__ void __assert_fail(const char *__assertion, const char *__
} else if (!(__cond2)) { \
break; \
} \
if (__spin_cnt++ == __max_spin_cnt) { \
if (__max_spin_cnt >= 0 && __spin_cnt++ == __max_spin_cnt) { \
__assert_fail(#__cond1 #__cond2, __FILE__, __LINE__, __PRETTY_FUNCTION__); \
} \
} \
2 changes: 1 addition & 1 deletion src/semaphore.cc
Original file line number Diff line number Diff line change
@@ -78,7 +78,7 @@ MSCCLPP_API_CPP void Host2HostSemaphore::wait(int64_t maxSpinCount) {
int64_t spinCount = 0;
while (cuda::atomic_ref<uint64_t, cuda::thread_scope_system>{*(uint64_t*)localInboundSemaphore_.get()}.load(
cuda::memory_order_acquire) < (*expectedInboundSemaphore_)) {
if (spinCount++ == maxSpinCount) {
if (maxSpinCount >= 0 && spinCount++ == maxSpinCount) {
throw Error("Host2HostSemaphore::wait timed out", ErrorCode::Timeout);
}
}

0 comments on commit 7686e15

Please sign in to comment.