Fix SYCL atomics for local memory

[masterleinad]

Fixes #4582. @brian-kelley Can you please check if this fixes the local atomics issue for you?

[masterleinad]

Also, see intel/llvm#4669 (comment). We would likely replace the checks here with address_space_cast as soon as it's available in the compiler or use generic_space as soon as it's available.
IMHO, passing the address space from the outside is likely neither an option simplicity (sycl::local_ptr, sycl::global_ptr) nor explicitly.

[brian-kelley]

@masterleinad Wow, that was fast. Yes, it did fix local atomics :)

[masterleinad]

Now, we just need to find a way to make that work for SYCL+CUDA...

[masterleinad]

I decided, for now, to only fix the behavior for Intel GPUs. I'm happy to create a corresponding pull request on the desul side if we are happy with this.

[dalg24]

Why is it not

auto g = __SYCL_GenericCastToPtrExplicit_ToGlobal<TYPE>(dest);
if (g) { ...

[masterleinad]

Do you want both checks?
For one check, I choose the space to test pretty arbitrarily. Do you have any arguments for rather checking for the global space?

[dalg24]

No I was just wondering whether we should check that the pointer is indeed global.

[masterleinad]

I don't feel overly strong about this. As of know, there is https://github.com/intel/llvm/blob/a3e9aab6a25a9fc69e805f1a93f409375ce1f7d1/sycl/include/CL/sycl/atomic_ref.hpp#L42-L49 but also https://github.com/intel/llvm/blob/a3e9aab6a25a9fc69e805f1a93f409375ce1f7d1/sycl/include/CL/sycl/atomic_ref.hpp#L116-L120. All that is to say that really only global_space and local_space are supported (and we would get rid of all of this as soon as generic_space is available).
At this point, I slightly prefer not to have another check for every call SYCL atomics. If you (or others) prefer we can add the check and adopt https://github.com/intel/llvm/blob/a3e9aab6a25a9fc69e805f1a93f409375ce1f7d1/sycl/include/sycl/ext/oneapi/sub_group.hpp#L442-L443.

[masterleinad]

Retest this please.

[masterleinad]

Retest this please.

[dalg24]

I decided, for now, to only fix the behavior for Intel GPUs. I'm happy to create a corresponding pull request on the desul side if we are happy with this.

Please do so

[masterleinad]

Please do so

See desul/desul#48.