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

[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?
+    return OpenCL || CUDA;
----------------
gonzalobg wrote:

> @gonzalobg -- Does NVIDIA define what happens if atomics are used on local address space?

I agree with @arsenm that this is a language property. In CUDA C++, just like in C++, the behavior of atomics to automatic variables is well-defined, e.g., this is ok:

```
__device__ void foo() {
   cuda::atomic<...> x(0);
   x.fetch_add(1); // OK
}
```

When compiling to PTX, however, `atom` requires global or shared statespaces (or generic to those).
That is, for local memory, LLVM must generate code that does not use `atom`.
But that's a problem for the LLVM NVPTX backend to solve.

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