[libclc][cuda] CTS fix: CUDA backend uses "success" atomic order for cas.

[JackAKirk]

CTS fix: CUDA backend uses "success" atomic order for cas.

There was a bug in the cas impl for nvptx in libclc that lead to CTS test failures for the CUDA backend.
This fixes the bug in a simple way by simply replacing the cases where the failure order differs from the success order (when failure order is either release or acquire), so that the failure order matches the success order (acq_rel). This is safe even if the cas performs the failure operation, because acq_rel can be used for both acquire (load) and release (store) atomic ops in ptx. I think that this is the only valid way to implement cas for nvptx, because the cas operation only takes one order argument.
Now the sycl cts passes for acq_rel atomics for the cuda backend.

[ldrumm]

Looks good. Do you have any references to NVIDIA docs that might back up your assertion of safety?

[JackAKirk]

Looks good. Do you have any references to NVIDIA docs that might back up your assertion of safety?

From https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#id70

• "acquire operation: A memory operation with .acquire or .acq_rel qualifier."
• "release operation: A memory operation with .release or .acq_rel qualifier."

For the case that failure is stronger than success. Justifying still using success order:

• From SYCL spec on definition of memory consistency model behavior:

"The SYCL memory consistency model is based upon the memory consistency model of the C++ core language. Where SYCL offers extensions to classes and functions that may affect memory consistency, the default behavior when these extensions are not used always matches the behavior of standard C++."

• and then a page later:

"
sycl::memory_order::acquire;
sycl::memory_order::release;
sycl::memory_order::acq_rel;
sycl::memory_order::seq_cst.
The meanings of these values are identical to those defined in the C++ core language.
"

Then looking at C++ definition:
https://en.cppreference.com/w/cpp/atomic/atomic/compare_exchange

"If failure is stronger than success or(until C++17) is one of std::memory_order_release and std::memory_order_acq_rel, the behavior is undefined."

Therefore the mapping that I described in the commit message, from the C++/SYCL interface (assuming default scope from atomic_ref constructor to satisfy the clause "the default behavior when these extensions are not used always matches the behavior of standard C++") to ptx interfaces:
compare_exchange_strong( T& expected, T desired, std::memory_order success, std::memory_order failure )

to atom{.sem}{.scope}{.space}.cas

Is I think the most sensible mapping that still satisfies the sycl function definition.