clang/AMDGPU: Set noalias.addrspace metadata on atomicrmw #102462

arsenm · 2024-08-08T13:19:12Z

No description provided.

arsenm · 2024-08-08T13:19:24Z

clang/AMDGPU: Set noalias.addrspace metadata on atomicrmw #102462 👈
[RFC] IR: Define noalias.addrspace metadata #102461 : 2 other dependent PRs (#102599 , #103938 )
main

This stack of pull requests is managed by Graphite. Learn more about stacking.

Join @arsenm and the rest of your teammates on Graphite

llvmbot · 2024-08-08T13:23:45Z

@llvm/pr-subscribers-backend-amdgpu

@llvm/pr-subscribers-clang-codegen

Author: Matt Arsenault (arsenm)

Changes

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

4 Files Affected:

(modified) clang/include/clang/Basic/LangOptions.h (+8)
(modified) clang/lib/CodeGen/Targets/AMDGPU.cpp (+11)
(modified) clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu (+59-56)
(modified) clang/test/CodeGenOpenCL/atomic-ops.cl (+10-8)

diff --git a/clang/include/clang/Basic/LangOptions.h b/clang/include/clang/Basic/LangOptions.h
index 91f1c2f2e6239e..9b7f1bddc0d1b3 100644
--- a/clang/include/clang/Basic/LangOptions.h
+++ b/clang/include/clang/Basic/LangOptions.h
@@ -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;
+  }
+
   /// Return the OpenCL C or C++ version as a VersionTuple.
   VersionTuple getOpenCLVersionTuple() const;
 
diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp
index 37e6af3d4196a8..c5797a405f7b30 100644
--- a/clang/lib/CodeGen/Targets/AMDGPU.cpp
+++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp
@@ -9,6 +9,7 @@
 #include "ABIInfoImpl.h"
 #include "TargetInfo.h"
 #include "clang/Basic/TargetOptions.h"
+#include "llvm/Support/AMDGPUAddrSpace.h"
 
 using namespace clang;
 using namespace clang::CodeGen;
@@ -550,6 +551,16 @@ AMDGPUTargetCodeGenInfo::getLLVMSyncScopeID(const LangOptions &LangOpts,
 
 void AMDGPUTargetCodeGenInfo::setTargetAtomicMetadata(
     CodeGenFunction &CGF, llvm::AtomicRMWInst &RMW) const {
+
+  if (RMW.getPointerAddressSpace() == llvm::AMDGPUAS::FLAT_ADDRESS &&
+      CGF.CGM.getLangOpts().threadPrivateMemoryAtomicsAreUndefined()) {
+    llvm::MDBuilder MDHelper(CGF.getLLVMContext());
+    llvm::MDNode *ASRange = MDHelper.createRange(
+        llvm::APInt(32, llvm::AMDGPUAS::PRIVATE_ADDRESS),
+        llvm::APInt(32, llvm::AMDGPUAS::PRIVATE_ADDRESS + 1));
+    RMW.setMetadata(llvm::LLVMContext::MD_noalias_addrspace, ASRange);
+  }
+
   if (!CGF.getTarget().allowAMDGPUUnsafeFPAtomics())
     return;
 
diff --git a/clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu b/clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu
index 8bf8241e343e70..a5e9bd6df07143 100644
--- a/clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu
+++ b/clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu
@@ -22,19 +22,19 @@
 
 __global__ void ffp1(float *p) {
   // CHECK-LABEL: @_Z4ffp1Pf
-  // SAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 4{{$}}
-  // SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 4{{$}}
-  // SAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 4{{$}}
-  // SAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 4{{$}}
-  // SAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 4{{$}}
-  // SAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 4{{$}}
-
-  // UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+}}, !amdgpu.ignore.denormal.mode !{{[0-9]+$}}
-  // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
-  // UNSAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
-  // UNSAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
-  // UNSAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
-  // UNSAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // SAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 4, !noalias.addrspace ![[NO_PRIVATE:[0-9]+]]{{$}}
+  // SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 4, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+  // SAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 4, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+  // SAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 4, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+  // SAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 4, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+  // SAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 4, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+
+  // UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 4, !noalias.addrspace ![[NO_PRIVATE:[0-9]+]], !amdgpu.no.fine.grained.memory !{{[0-9]+}}, !amdgpu.ignore.denormal.mode !{{[0-9]+$}}
+  // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 4, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // UNSAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 4, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // UNSAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 4, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // UNSAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 4, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // UNSAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 4, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
 
   // SAFE: _Z4ffp1Pf
   // SAFE: global_atomic_cmpswap
@@ -62,19 +62,19 @@ __global__ void ffp1(float *p) {
 
 __global__ void ffp2(double *p) {
   // CHECK-LABEL: @_Z4ffp2Pd
-  // SAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 8{{$}}
-  // SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8{{$}}
-  // SAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 8{{$}}
-  // SAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 8{{$}}
-  // SAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 8{{$}}
-  // SAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 8{{$}}
-
-  // UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
-  // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
-  // UNSAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
-  // UNSAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
-  // UNSAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
-  // UNSAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // SAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+  // SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+  // SAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+  // SAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+  // SAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+  // SAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+
+  // UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // UNSAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // UNSAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // UNSAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // UNSAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
 
   // SAFE-LABEL: @_Z4ffp2Pd
   // SAFE: global_atomic_cmpswap_b64
@@ -102,19 +102,19 @@ __global__ void ffp2(double *p) {
 // long double is the same as double for amdgcn.
 __global__ void ffp3(long double *p) {
   // CHECK-LABEL: @_Z4ffp3Pe
-  // SAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 8{{$}}
-  // SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8{{$}}
-  // SAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 8{{$}}
-  // SAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 8{{$}}
-  // SAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 8{{$}}
-  // SAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 8{{$}}
-
-  // UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
-  // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
-  // UNSAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
-  // UNSAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
-  // UNSAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
-  // UNSAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // SAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+  // SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+  // SAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+  // SAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+  // SAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+  // SAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+
+  // UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // UNSAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // UNSAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // UNSAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // UNSAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
 
   // SAFE-LABEL: @_Z4ffp3Pe
   // SAFE: global_atomic_cmpswap_b64
@@ -139,34 +139,34 @@ __global__ void ffp3(long double *p) {
 __device__ double ffp4(double *p, float f) {
   // CHECK-LABEL: @_Z4ffp4Pdf
   // CHECK: fpext float {{.*}} to double
-  // SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8{{$}}
-  // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+  // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
   return __atomic_fetch_sub(p, f, memory_order_relaxed);
 }
 
 __device__ double ffp5(double *p, int i) {
   // CHECK-LABEL: @_Z4ffp5Pdi
   // CHECK: sitofp i32 {{.*}} to double
-  // SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8{{$}}
-  // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+  // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
   return __atomic_fetch_sub(p, i, memory_order_relaxed);
 }
 
 __global__ void ffp6(_Float16 *p) {
   // CHECK-LABEL: @_Z4ffp6PDF16
-  // SAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 2{{$}}
-  // SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 2{{$}}
-  // SAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 2{{$}}
-  // SAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 2{{$}}
-  // SAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 2{{$}}
-  // SAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 2{{$}}
-
-  // UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 2, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
-  // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 2, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
-  // UNSAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 2, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
-  // UNSAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 2, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
-  // UNSAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 2, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
-  // UNSAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 2, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // SAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 2, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+  // SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 2, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+  // SAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 2, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+  // SAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 2, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+  // SAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 2, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+  // SAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 2, !noalias.addrspace ![[NO_PRIVATE]]{{$}}
+
+  // UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 2, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 2, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // UNSAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 2, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // UNSAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 2, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // UNSAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 2, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+  // UNSAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 2, !noalias.addrspace ![[NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
 
   // SAFE: _Z4ffp6PDF16
   // SAFE: global_atomic_cmpswap
@@ -190,3 +190,6 @@ __global__ void ffp6(_Float16 *p) {
   __hip_atomic_fetch_max(p, 1.0f, memory_order_relaxed, __HIP_MEMORY_SCOPE_AGENT);
   __hip_atomic_fetch_min(p, 1.0f, memory_order_relaxed, __HIP_MEMORY_SCOPE_WORKGROUP);
 }
+
+// SAFEIR: ![[NO_PRIVATE]] = !{i32 5, i32 6}
+// UNSAFEIR: ![[NO_PRIVATE]] = !{i32 5, i32 6}
diff --git a/clang/test/CodeGenOpenCL/atomic-ops.cl b/clang/test/CodeGenOpenCL/atomic-ops.cl
index 5e2de38ac3d3e3..137d78e32b1da8 100644
--- a/clang/test/CodeGenOpenCL/atomic-ops.cl
+++ b/clang/test/CodeGenOpenCL/atomic-ops.cl
@@ -70,19 +70,19 @@ void test_addr(global atomic_int *ig, private atomic_int *ip, local atomic_int *
 
 void fi3(atomic_int *i, atomic_uint *ui) {
   // CHECK-LABEL: @fi3
-  // CHECK: atomicrmw and ptr %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
+  // CHECK: atomicrmw and ptr %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4, !noalias.addrspace [[NOPRIVATE:![0-9]+]]{{$}}
   int x = __opencl_atomic_fetch_and(i, 1, memory_order_seq_cst, memory_scope_work_group);
 
-  // CHECK: atomicrmw min ptr %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
+  // CHECK: atomicrmw min ptr %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4, !noalias.addrspace [[NOPRIVATE:![0-9]+]]{{$}}
   x = __opencl_atomic_fetch_min(i, 1, memory_order_seq_cst, memory_scope_work_group);
 
-  // CHECK: atomicrmw max ptr %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
+  // CHECK: atomicrmw max ptr %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4, !noalias.addrspace [[NOPRIVATE:![0-9]+]]{{$}}
   x = __opencl_atomic_fetch_max(i, 1, memory_order_seq_cst, memory_scope_work_group);
 
-  // CHECK: atomicrmw umin ptr %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
+  // CHECK: atomicrmw umin ptr %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4, !noalias.addrspace [[NOPRIVATE:![0-9]+]]{{$}}
   x = __opencl_atomic_fetch_min(ui, 1, memory_order_seq_cst, memory_scope_work_group);
 
-  // CHECK: atomicrmw umax ptr %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
+  // CHECK: atomicrmw umax ptr %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4, !noalias.addrspace [[NOPRIVATE:![0-9]+]]{{$}}
   x = __opencl_atomic_fetch_max(ui, 1, memory_order_seq_cst, memory_scope_work_group);
 }
 
@@ -186,19 +186,19 @@ void ff2(atomic_float *d) {
 
 float ff3(atomic_float *d) {
   // CHECK-LABEL: @ff3
-  // CHECK: atomicrmw xchg ptr {{.*}} syncscope("workgroup") seq_cst, align 4
+  // CHECK: atomicrmw xchg ptr {{.*}} syncscope("workgroup") seq_cst, align 4, !noalias.addrspace [[NOPRIVATE:![0-9]+]]{{$}}
   return __opencl_atomic_exchange(d, 2, memory_order_seq_cst, memory_scope_work_group);
 }
 
 float ff4(global atomic_float *d, float a) {
   // CHECK-LABEL: @ff4
-  // CHECK: atomicrmw fadd ptr addrspace(1) {{.*}} syncscope("workgroup-one-as") monotonic
+  // CHECK: atomicrmw fadd ptr addrspace(1) {{.*}} syncscope("workgroup-one-as") monotonic, align 4{{$}}
   return __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_work_group);
 }
 
 float ff5(global atomic_double *d, double a) {
   // CHECK-LABEL: @ff5
-  // CHECK: atomicrmw fadd ptr addrspace(1) {{.*}} syncscope("workgroup-one-as") monotonic
+  // CHECK: atomicrmw fadd ptr addrspace(1) {{.*}} syncscope("workgroup-one-as") monotonic, align 8{{$}}
   return __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_work_group);
 }
 
@@ -342,3 +342,5 @@ int test_volatile(volatile atomic_int *i) {
 }
 
 #endif
+
+// CHECK: [[NOPRIVATE]] = !{i32 5, i32 6}

shiltian · 2024-08-08T14:50:25Z

clang/include/clang/Basic/LangOptions.h

@@ -647,6 +647,14 @@ class LangOptions : public LangOptionsBase {
    return ConvergentFunctions;
  }

+  /// Return true if atomicrmw operations targeting allocations in private


Do we want to have a check in target machine to tell if atomic operation on specific address space is legal? I'm thinking of adding atomic support in AAAddressSpace, and could drop the address space if an atomic operation in an inferred address space is not legal.

This has nothing to do with the target, this is language semantics. What the target natively supports is an incomprehensibly complicated question, and is a legalization issue. We could annotate some operations with this in infer address spaces, but there's only a few cases where it's probably useful there

clang/lib/CodeGen/Targets/AMDGPU.cpp

clang/include/clang/Basic/LangOptions.h

arsenm · 2024-09-06T06:38:17Z

ping

Legacy atomic* apis are allowed to not work for private, but not std::atomic style. Base this based on the language builtins, rather than the language directly. Also start handling cmpxchg.

arsenm · 2024-10-15T05:46:48Z

ping

llvm-ci · 2024-10-17T13:23:20Z

LLVM Buildbot has detected a new failure on builder sanitizer-aarch64-linux-fuzzer running on sanitizer-buildbot12 while building clang at step 2 "annotate".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/159/builds/8365

Here is the relevant piece of the build log for the reference

Step 2 (annotate) failure: 'python ../sanitizer_buildbot/sanitizers/zorg/buildbot/builders/sanitizers/buildbot_selector.py' (failure)
...
  specific short-term circumstances.  Projects should be ported to the NEW
  behavior and not rely on setting a policy to OLD.
Call Stack (most recent call first):
  CMakeLists.txt:6 (include)


-- Building with -fPIC
-- LLVM host triple: aarch64-unknown-linux-gnu
-- LLVM default target triple: aarch64-unknown-linux-gnu
-- Using libunwind testing configuration: /home/b/sanitizer-aarch64-linux-fuzzer/build/llvm-project/libunwind/test/configs/llvm-libunwind-shared.cfg.in
-- Failed to locate sphinx-build executable (missing: SPHINX_EXECUTABLE) 
-- Using libc++abi testing configuration: /home/b/sanitizer-aarch64-linux-fuzzer/build/llvm-project/libcxxabi/test/configs/llvm-libc++abi-shared.cfg.in
-- Using libc++ testing configuration: /home/b/sanitizer-aarch64-linux-fuzzer/build/llvm-project/libcxx/test/configs/llvm-libc++-shared.cfg.in
-- ABI list file not generated for configuration aarch64-unknown-linux-gnu.libcxxabi.v1.stable.exceptions.nonew, `check-cxx-abilist` will not be available.
CMake Deprecation Warning at /home/b/sanitizer-aarch64-linux-fuzzer/build/llvm-project/cmake/Modules/CMakePolicy.cmake:6 (cmake_policy):
  The OLD behavior for policy CMP0116 will be removed from a future version
  of CMake.

  The cmake-policies(7) manual explains that the OLD behaviors of all
  policies are deprecated and that a policy should be set to OLD only under
  specific short-term circumstances.  Projects should be ported to the NEW
  behavior and not rely on setting a policy to OLD.
Call Stack (most recent call first):
  /home/b/sanitizer-aarch64-linux-fuzzer/build/llvm-project/compiler-rt/CMakeLists.txt:12 (include)


-- Compiler-RT supported architectures: aarch64
-- Generated Sanitizer SUPPORTED_TOOLS list on "Linux" is "asan;lsan;hwasan;msan;tsan;ubsan"
-- sanitizer_common tests on "Linux" will run against "asan;lsan;hwasan;msan;tsan;ubsan"
-- Configuring done (2.4s)
-- Generating done (0.9s)
-- Build files have been written to: /home/b/sanitizer-aarch64-linux-fuzzer/build/llvm_build0/runtimes/runtimes-bins
[1054/1063] Building CXX object tools/clang/tools/libclang/CMakeFiles/libclang.dir/CIndex.cpp.o
[1055/1063] Performing build step for 'runtimes'
[0/7] Performing build step for 'libcxx_fuzzer_aarch64'
ninja: no work to do.
[3/7] Building CXX object compiler-rt/lib/hwasan/CMakeFiles/RTHwasan_dynamic_version_script_dummy.aarch64.dir/dummy.cpp.o
[5/7] Linking CXX shared library /home/b/sanitizer-aarch64-linux-fuzzer/build/llvm_build0/lib/clang/20/lib/aarch64-unknown-linux-gnu/libclang_rt.ubsan_standalone.so
[6/7] Linking CXX shared library /home/b/sanitizer-aarch64-linux-fuzzer/build/llvm_build0/lib/clang/20/lib/aarch64-unknown-linux-gnu/libclang_rt.hwasan.so
[7/7] Linking CXX shared library /home/b/sanitizer-aarch64-linux-fuzzer/build/llvm_build0/lib/clang/20/lib/aarch64-unknown-linux-gnu/libclang_rt.asan.so
[1056/1063] No install step for 'runtimes'
[1058/1063] Completed 'runtimes'
[1059/1063] Building CXX object tools/clang/tools/libclang/CMakeFiles/libclang.dir/CXExtractAPI.cpp.o
[1060/1063] Linking CXX shared library lib/libclang.so.20.0.0git
[1061/1063] Creating library symlink lib/libclang.so.20.0git lib/libclang.so
[1062/1063] Linking CXX executable bin/c-arcmt-test
[1063/1063] Linking CXX executable bin/c-index-test
dc1578d895703bd928ff4acfef85056a  llvm_build0/bin/clang
@@@BUILD_STEP get fuzzer-test-suite @@@
fatal: unable to access 'https://github.com/google/fuzzer-test-suite.git/': gnutls_handshake() failed: Illegal parameter
Step 7 (stage1 build all) failure: stage1 build all (failure)
...
CMake Deprecation Warning at /home/b/sanitizer-aarch64-linux-fuzzer/build/llvm-project/cmake/Modules/CMakePolicy.cmake:6 (cmake_policy):
  The OLD behavior for policy CMP0116 will be removed from a future version
  of CMake.

  The cmake-policies(7) manual explains that the OLD behaviors of all
  policies are deprecated and that a policy should be set to OLD only under
  specific short-term circumstances.  Projects should be ported to the NEW
  behavior and not rely on setting a policy to OLD.
Call Stack (most recent call first):
  CMakeLists.txt:6 (include)
-- Building with -fPIC
-- LLVM host triple: aarch64-unknown-linux-gnu
-- LLVM default target triple: aarch64-unknown-linux-gnu
-- Using libunwind testing configuration: /home/b/sanitizer-aarch64-linux-fuzzer/build/llvm-project/libunwind/test/configs/llvm-libunwind-shared.cfg.in
-- Failed to locate sphinx-build executable (missing: SPHINX_EXECUTABLE) 
-- Using libc++abi testing configuration: /home/b/sanitizer-aarch64-linux-fuzzer/build/llvm-project/libcxxabi/test/configs/llvm-libc++abi-shared.cfg.in
-- Using libc++ testing configuration: /home/b/sanitizer-aarch64-linux-fuzzer/build/llvm-project/libcxx/test/configs/llvm-libc++-shared.cfg.in
-- ABI list file not generated for configuration aarch64-unknown-linux-gnu.libcxxabi.v1.stable.exceptions.nonew, `check-cxx-abilist` will not be available.
CMake Deprecation Warning at /home/b/sanitizer-aarch64-linux-fuzzer/build/llvm-project/cmake/Modules/CMakePolicy.cmake:6 (cmake_policy):
  The OLD behavior for policy CMP0116 will be removed from a future version
  of CMake.

  The cmake-policies(7) manual explains that the OLD behaviors of all
  policies are deprecated and that a policy should be set to OLD only under
  specific short-term circumstances.  Projects should be ported to the NEW
  behavior and not rely on setting a policy to OLD.
Call Stack (most recent call first):
  /home/b/sanitizer-aarch64-linux-fuzzer/build/llvm-project/compiler-rt/CMakeLists.txt:12 (include)
-- Compiler-RT supported architectures: aarch64
-- Generated Sanitizer SUPPORTED_TOOLS list on "Linux" is "asan;lsan;hwasan;msan;tsan;ubsan"
-- sanitizer_common tests on "Linux" will run against "asan;lsan;hwasan;msan;tsan;ubsan"
-- Configuring done (2.4s)
-- Generating done (0.9s)
-- Build files have been written to: /home/b/sanitizer-aarch64-linux-fuzzer/build/llvm_build0/runtimes/runtimes-bins
[1054/1063] Building CXX object tools/clang/tools/libclang/CMakeFiles/libclang.dir/CIndex.cpp.o
[1055/1063] Performing build step for 'runtimes'
[0/7] Performing build step for 'libcxx_fuzzer_aarch64'
ninja: no work to do.
[3/7] Building CXX object compiler-rt/lib/hwasan/CMakeFiles/RTHwasan_dynamic_version_script_dummy.aarch64.dir/dummy.cpp.o
[5/7] Linking CXX shared library /home/b/sanitizer-aarch64-linux-fuzzer/build/llvm_build0/lib/clang/20/lib/aarch64-unknown-linux-gnu/libclang_rt.ubsan_standalone.so
[6/7] Linking CXX shared library /home/b/sanitizer-aarch64-linux-fuzzer/build/llvm_build0/lib/clang/20/lib/aarch64-unknown-linux-gnu/libclang_rt.hwasan.so
[7/7] Linking CXX shared library /home/b/sanitizer-aarch64-linux-fuzzer/build/llvm_build0/lib/clang/20/lib/aarch64-unknown-linux-gnu/libclang_rt.asan.so
[1056/1063] No install step for 'runtimes'
[1058/1063] Completed 'runtimes'
[1059/1063] Building CXX object tools/clang/tools/libclang/CMakeFiles/libclang.dir/CXExtractAPI.cpp.o
[1060/1063] Linking CXX shared library lib/libclang.so.20.0.0git
[1061/1063] Creating library symlink lib/libclang.so.20.0git lib/libclang.so
[1062/1063] Linking CXX executable bin/c-arcmt-test
[1063/1063] Linking CXX executable bin/c-index-test
dc1578d895703bd928ff4acfef85056a  llvm_build0/bin/clang

arsenm mentioned this pull request Aug 8, 2024

[RFC] IR: Define noalias.addrspace metadata #102461

Merged

arsenm requested review from AlexVlx, Artem-B, bogner, Endilll, frasercrmck, hvdijk, jdoerfert, jh7370, jhuber6, jyknight, lei137, shiltian and yxsamliu August 8, 2024 13:23

arsenm marked this pull request as ready for review August 8, 2024 13:23

llvmbot added clang Clang issues not falling into any other category backend:AMDGPU clang:frontend Language frontend issues, e.g. anything involving "Sema" clang:codegen IR generation bugs: mangling, exceptions, etc. labels Aug 8, 2024

shiltian reviewed Aug 8, 2024

View reviewed changes

Artem-B reviewed Aug 8, 2024

View reviewed changes

clang/lib/CodeGen/Targets/AMDGPU.cpp Outdated Show resolved Hide resolved

clang/include/clang/Basic/LangOptions.h Show resolved Hide resolved

hvdijk reviewed Aug 9, 2024

View reviewed changes

clang/include/clang/Basic/LangOptions.h Show resolved Hide resolved

arsenm force-pushed the users/arsenm/noalias-addrspace-metadata branch from e403d08 to 1e084c6 Compare August 9, 2024 11:06

arsenm mentioned this pull request Aug 9, 2024

AMDGPU: Add noalias.addrspace metadata when autoupgrading atomic intrinsics #102599

Merged

arsenm force-pushed the users/arsenm/noalias-addrspace-metadata branch from bff0427 to 32ee4c9 Compare August 14, 2024 13:21

This was referenced Aug 14, 2024

Local: Handle noalias_addrspace in combineMetadata #103938

Merged

Local: Handle noalias.addrspace in copyMetadataForLoad #103939

Merged

arsenm force-pushed the users/arsenm/noalias-addrspace-metadata branch from 32ee4c9 to 3e97987 Compare September 4, 2024 05:23

arsenm force-pushed the users/arsenm/clang-amdgpu-set-noalias-addrspace-atomicrmw branch from a8852f0 to 39392c6 Compare September 4, 2024 05:24

arsenm force-pushed the users/arsenm/noalias-addrspace-metadata branch 2 times, most recently from 16dd092 to c3dbcec Compare September 6, 2024 17:40

arsenm force-pushed the users/arsenm/noalias-addrspace-metadata branch 2 times, most recently from 79e1e8b to f84e99b Compare September 20, 2024 11:54

arsenm force-pushed the users/arsenm/clang-amdgpu-set-noalias-addrspace-atomicrmw branch from 39392c6 to a9840f7 Compare September 20, 2024 12:00

arsenm force-pushed the users/arsenm/noalias-addrspace-metadata branch from f84e99b to 29b8b60 Compare September 24, 2024 04:35

arsenm force-pushed the users/arsenm/clang-amdgpu-set-noalias-addrspace-atomicrmw branch from a9840f7 to 32d917d Compare September 24, 2024 04:36

arsenm force-pushed the users/arsenm/noalias-addrspace-metadata branch 2 times, most recently from fb9d412 to d7b148d Compare September 30, 2024 07:09

arsenm force-pushed the users/arsenm/noalias-addrspace-metadata branch from d7b148d to 4a6df18 Compare October 7, 2024 09:15

Base automatically changed from users/arsenm/noalias-addrspace-metadata to main October 7, 2024 19:21

arsenm added 2 commits October 8, 2024 11:44

clang/AMDGPU: Set noalias.addrspace metadata on atomicrmw

5f30c2e

Change behavior per-builtin

cb6691a

Legacy atomic* apis are allowed to not work for private, but not std::atomic style. Base this based on the language builtins, rather than the language directly. Also start handling cmpxchg.

arsenm force-pushed the users/arsenm/clang-amdgpu-set-noalias-addrspace-atomicrmw branch from 32d917d to cb6691a Compare October 8, 2024 09:59

shiltian approved these changes Oct 15, 2024

View reviewed changes

arsenm merged commit 51b4ada into main Oct 17, 2024
8 checks passed

arsenm deleted the users/arsenm/clang-amdgpu-set-noalias-addrspace-atomicrmw branch October 17, 2024 13:10

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

clang/AMDGPU: Set noalias.addrspace metadata on atomicrmw #102462

clang/AMDGPU: Set noalias.addrspace metadata on atomicrmw #102462

Uh oh!

arsenm commented Aug 8, 2024

Uh oh!

arsenm commented Aug 8, 2024 •

edited

Loading

Uh oh!

llvmbot commented Aug 8, 2024 •

edited

Loading

Uh oh!

shiltian Aug 8, 2024

Uh oh!

arsenm Aug 8, 2024

Uh oh!

Uh oh!

Uh oh!

Uh oh!

arsenm commented Sep 6, 2024

Uh oh!

arsenm commented Oct 15, 2024

Uh oh!

Uh oh!

llvm-ci commented Oct 17, 2024

Uh oh!

Uh oh!

clang/AMDGPU: Set noalias.addrspace metadata on atomicrmw #102462

clang/AMDGPU: Set noalias.addrspace metadata on atomicrmw #102462

Uh oh!

Conversation

arsenm commented Aug 8, 2024

Uh oh!

arsenm commented Aug 8, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmbot commented Aug 8, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

shiltian Aug 8, 2024

Choose a reason for hiding this comment

Uh oh!

arsenm Aug 8, 2024

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

arsenm commented Sep 6, 2024

Uh oh!

arsenm commented Oct 15, 2024

Uh oh!

Uh oh!

llvm-ci commented Oct 17, 2024

Uh oh!

Uh oh!

arsenm commented Aug 8, 2024 •

edited

Loading

llvmbot commented Aug 8, 2024 •

edited

Loading