1

From my understanding, CUDA's atomicCAS has the following definition (this is one of the four)

int atomicCAS(int* address, int compare, int val);

and it compares atomically the values located at address (named in the doc old) in the global shared memory with compare and in case of equality assigns the value to val, otherwise does nothing. In both cases returns old.

Looking at SYCL API, I can only find compare_exchange_strong which, unfortunately, does not do what I'm looking for as, using the same naming as above, it compares old with compare and if not successful, alters compare (which is passed by reference).

aland
  • 4,829
  • 2
  • 24
  • 42
Elle
  • 305
  • 2
  • 10
  • 1
    Yes, C++ [`std::atomic`'s `compare_exchange_strong`](https://en.cppreference.com/w/cpp/atomic/atomic/compare_exchange) updates `expected` by reference, but it's intended to be a local variable. The update of that output arg is *not* atomic. It's just a different way of exposing the same underlying primitive operation, an atomic CAS. Changing code to use it is purely a matter of declaring local variables. If the compare was true, then `compare` is already equal to the `old` value, so not update of it is necessary. IDK if that's why you thought it was different or a problem for your code. – Peter Cordes May 29 '22 at 18:59

1 Answers1

2

As Peter Cordes noted in a comment, sycl::compare_exchange_strong is the right choice. From the SYCL 2020 rev. 4 description of compare_exchange_strong:

Atomically compares the value of the object referenced by this atomic_ref against the value of expected. If the values are equal, replaces the value of the referenced object with the value of desired; otherwise assigns the original value of the referenced object to expected.

So,

int old = compare;
ref.compare_exchange_strong(old, val);

is, in terms of updating ref, equivalent to

old = atomicCAS(address, compare, val);

If you are interested, you can see for yourself how hipSYCL implements sycl::compare_exchange_strong.

aland
  • 4,829
  • 2
  • 24
  • 42