Fixes CUDA generic atomic operation by fixing atomic load - #2054
Conversation
- Original way seemed to cause some race condition
with the messages feature and the CUDA execution policy
|
Is there a reason for changing just the implementation for atomicGeneric and not using the refactored atomic_ref implementation in cuda_atomicCAS_loop? |
cuda_atomicCAS_loop has been around for much longer than atomicGeneric and is used by other operations such as atomicMin and atomicMax without any noticeable problems. This new implementation has not been thoroughly tested or benchmarked, so I'm hesitant to make changes to cuda_atomicCAS_loop directly. |
|
Regarding |
artv3
left a comment
There was a problem hiding this comment.
LGTM, just to double check -- we have tests for this right?
Yes, there are tests. |
I don't understand the cause yet. Though the old implementation was using atomicOr(acc, 0) to do the load, so maybe there was something weird happening with that? |
Maybe good to capture in an issue if we don't have one already |
@bechols97 did some more testing and it appears that using cuda::atomic_ref::load in place of the atomicOr for the load fixed the issues. I'm guessing atomic_ref will be faster than the atomicOr, but even if not, it seems like for correctness we should be using atomic_ref for the load. |
| std::enable_if_t<cuda_useBuiltinCommon<T>::value, bool> = true> | ||
| RAJA_INLINE __device__ T cuda_atomicLoad(T* acc) | ||
| { | ||
| return cuda_atomicOr(acc, static_cast<T>(0)); |
There was a problem hiding this comment.
Using cuda::atomic_ref instead of this fixed the bug @bechols97 was seeing.
There was a problem hiding this comment.
Interesting, I'm pretty sure I added this and it was passing our tests before.
There was a problem hiding this comment.
Yeah, I'm pretty sure we have testing for it, though not tests that use the result in other atomic operations.
Summary