[llvm-branch-commits] [clang] clang/AMDGPU: Set noalias.addrspace metadata on atomicrmw (PR #102462)

[Harald van Dijk via llvm-branch-commits]

================
@@ -647,6 +647,14 @@ class LangOptions : public LangOptionsBase {
     return ConvergentFunctions;
   }
 
+  /// Return true if atomicrmw operations targeting allocations in private
+  /// memory are undefined.
+  bool threadPrivateMemoryAtomicsAreUndefined() const {
+    // Should be false for OpenMP.
+    // TODO: Should this be true for SYCL?
----------------
hvdijk wrote:

I think this is meant to be true for SYCL in the future, but not just yet? [SYCL 2020 4.15.3](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:atomic-references) specifies

> The `sycl::atomic_ref` class also has a template parameter `AddressSpace`, which allows the application to make an assertion about the address space of the object of type `T` that it references. The default value for this parameter is `access::address_space::generic_space`, which indicates that the object could be in either the global or local address spaces. If the application knows the address space, it can set this template parameter to either `access::address_space::global_space` or `access::address_space::local_space` as an assertion to the implementation. Specifying the address space via this template parameter may allow the implementation to perform certain optimizations. Specifying an address space that does not match the object’s actual address space results in undefined behavior.

It does not specifically call out the private address space as being undefined, but it says an address space that does not match what is specified is undefined, and provides no way to specify the private address space, so I think the end result is the same.

However, at the moment, the [legacy atomic types](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:atom-types-depr) are also still available and the same logic cannot be applied to those, so barring any explicit statement that the private address space is undefined, I think it will be necessary to assume for now that this is well-defined.

https://github.com/llvm/llvm-project/pull/102462