PI CUDA ERROR when using sycl::atomic_ref

[guoci]

Describe the bug
A clear and concise description of what the bug is.
I get an error when running the program below.

PI CUDA ERROR:
	Value:           719
	Name:            CUDA_ERROR_LAUNCH_FAILED
	Description:     unspecified launch failure
	Function:        wait
	Source Location: /root/intel-llvm-mirror/sycl/plugins/cuda/pi_cuda.cpp:653

terminate called after throwing an instance of 'sycl::_V1::runtime_error'
  what():  Native API failed. Native API returns: -999 (Unknown PI error) -999 (Unknown PI error)

To Reproduce
Please describe the steps to reproduce the behavior:

1. Include code snippet as short as possible

#include <sycl/sycl.hpp>
int main() {
    int i = 0;
    sycl::queue q{sycl::gpu_selector_v};
    sycl::buffer<int> buf_mod{&i, 1};
    q.submit([&](sycl::handler &h) {
        sycl::accessor mod{buf_mod, h, sycl::write_only};
        h.single_task([=] {
            sycl::atomic_ref<int, sycl::memory_order::seq_cst,
                    sycl::memory_scope::device, sycl::access::address_space::global_space> a{mod[0]};
            ++a;
        });
    }).wait_and_throw();
    std::cout << i << std::endl;
}

2. Specify the command which should be used to compile the program
   clang++ -O3 -DNDEBUG main.cpp -fsycl -fsycl-targets=nvptx64-nvidia-cuda
3. Specify the comment which should be used to launch the program
   ./a.out
4. Indicate what is wrong and what was expected
   program should run without errors.
   Environment (please complete the following information):

• OS: [e.g Windows/Linux]
  Linux

• Target device and vendor: [e.g. Intel GPU]
  Nvidia GPU

• DPC++ version: [e.g. commit hash or output of clang++ --version]
  Intel(R) oneAPI DPC++/C++ Compiler 2023.2.0 (2023.2.0.20230622)
  Target: x86_64-unknown-linux-gnu
  Thread model: posix
  InstalledDir: /opt/intel/oneapi/compiler/2023.2.0/linux/bin-llvm

• Dependencies version: [e.g. low-level runtime versions (like NEO 20.04)]

Additional context
Add any other context about the problem here.
If the sycl::atomic_ref is using sycl::memory_order::acq_rel or sycl::memory_order::relaxed, then it runs without errors.

[JackAKirk]

seq_cst isn't implemented in the cuda backend atm. However there should at least be a proper error saying this.

Is this feature important to you?

[guoci]

I am fine without it.
On a side note, sycl::memory_order::acq_rel works although it seems that CUDA only has relaxed atomics.

[JackAKirk]

I am fine without it. On a side note, sycl::memory_order::acq_rel works although it seems that CUDA only has relaxed atomics.

cuda supports non relaxed atomics for sm70 and above: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#memory-consistency-model

The only difference is seq_cst, which doesn't have the same level of native support, but can apparently be implemented via fences.

Here is a summary I wrote:

The sycl::memory_order parameter corresponds with the ptx .sem qualifier that can be used on all atomic operations:

"The .sem qualifier requires sm_70 or higher. It specifies a memory synchronizing effect as described in the Memory Consistency Model. If the .sem qualifier is absent, .relaxed is assumed by default."

There is a memory-order correspondence between all the possible values of the .sem qualifier and those of sycl::memory_order except that sycl::memory_order::seq_cst is not supported by ptx:

sycl::memory_order::relaxed  → .sem = .relaxed
sycl::memory_order::acquire  →  .sem = .acquire
sycl::memory_order::release  →  .sem = .release
sycl::memory_order::acq_rel → .sem = .acq_rel

[JackAKirk]

I've mapped sycl seq_cst to cuda backend. Details here: #12516 (comment)