================
@@ -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
_______________________________________________
llvm-branch-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
