Is there a DPC++/SYCL equivalent of CUDA's atomicCAS?

Question

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).

Answer 1

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.