Skip to content

Fixes CUDA generic atomic operation by fixing atomic load - #2054

Merged
adayton1 merged 6 commits into
developfrom
bugfix/bechols97/cuda_generic_atomic
Jul 11, 2026
Merged

Fixes CUDA generic atomic operation by fixing atomic load#2054
adayton1 merged 6 commits into
developfrom
bugfix/bechols97/cuda_generic_atomic

Conversation

@adayton1

@adayton1 adayton1 commented Jul 9, 2026

Copy link
Copy Markdown
Member

Summary

  • This PR is a bugfix from @bechols97
  • It does the following:
    • Uses cuda::atomic_ref::load for atomicLoad instead of atomicOr with a value of 0
    • Uses cuda::atomic_ref::store for atomicStore instead of atomicExchange
    • The old implementation of atomicLoad in CUDA was giving incorrect results in Feature: Device messages #1832
    • Performance should be equivalent or better using atomic_ref, but even if not, correctness matters most
    • The RAJA_ENABLE_CUDA_ATOMIC_REF was never getting defined as intended, but now that we require CUDA 12, cuda::atomic_ref should always be available. The cuda/atomic header gives a very clear error message if an architecture older than sm_60 is used.

bechols97 added 2 commits July 9, 2026 14:15
  - Original way seemed to cause some race condition
    with the messages feature and the CUDA execution policy
@smithsg84

Copy link
Copy Markdown
Member

Is there a reason for changing just the implementation for atomicGeneric and not using the refactored atomic_ref implementation in cuda_atomicCAS_loop?

@adayton1

adayton1 commented Jul 9, 2026

Copy link
Copy Markdown
Member Author

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.

@artv3

artv3 commented Jul 10, 2026

Copy link
Copy Markdown
Member

Regarding old implementation of atomicGeneric in CUDA is giving incorrect results is that a CUDA related bug? or why is that happening?

@artv3 artv3 left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM, just to double check -- we have tests for this right?

Comment thread include/RAJA/policy/cuda/atomic.hpp Outdated
@adayton1

Copy link
Copy Markdown
Member Author

LGTM, just to double check -- we have tests for this right?

Yes, there are tests.

@adayton1

Copy link
Copy Markdown
Member Author

Regarding old implementation of atomicGeneric in CUDA is giving incorrect results is that a CUDA related bug? or why is that happening?

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?

@artv3

artv3 commented Jul 10, 2026

Copy link
Copy Markdown
Member

Regarding old implementation of atomicGeneric in CUDA is giving incorrect results is that a CUDA related bug? or why is that happening?

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

@adayton1

Copy link
Copy Markdown
Member Author

Regarding old implementation of atomicGeneric in CUDA is giving incorrect results is that a CUDA related bug? or why is that happening?

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?

@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));

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Using cuda::atomic_ref instead of this fixed the bug @bechols97 was seeing.

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Interesting, I'm pretty sure I added this and it was passing our tests before.

@adayton1 adayton1 Jul 10, 2026

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yeah, I'm pretty sure we have testing for it, though not tests that use the result in other atomic operations.

Comment thread include/RAJA/policy/cuda/atomic.hpp Outdated
Comment thread include/RAJA/policy/cuda/atomic.hpp Outdated
Comment thread include/RAJA/policy/cuda/atomic.hpp Outdated
@adayton1 adayton1 changed the title Fixes CUDA generic atomic operation Fixes CUDA generic atomic operation by fixing atomic load Jul 10, 2026
@bechols97 bechols97 mentioned this pull request Jul 10, 2026
6 tasks
@adayton1
adayton1 enabled auto-merge (squash) July 10, 2026 23:26
@adayton1
adayton1 merged commit 2be7c5a into develop Jul 11, 2026
21 checks passed
@adayton1
adayton1 deleted the bugfix/bechols97/cuda_generic_atomic branch July 11, 2026 02:21
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

6 participants