[clang] [RFC] Add clang atomic control options and pragmas (PR #102569)

[Matt Arsenault via cfe-commits]

================
@@ -1,50 +1,48 @@
 // RUN: %clang_cc1 -x hip %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa \
 // RUN:   -fcuda-is-device -target-cpu gfx906 -fnative-half-type \
-// RUN:   -fnative-half-arguments-and-returns | FileCheck -check-prefixes=CHECK,SAFEIR %s
+// RUN:   -fnative-half-arguments-and-returns | FileCheck -check-prefixes=FUN,CHECK,SAFEIR %s
 
 // RUN: %clang_cc1 -x hip %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa \
 // RUN:   -fcuda-is-device -target-cpu gfx906 -fnative-half-type \
-// RUN:   -fnative-half-arguments-and-returns -munsafe-fp-atomics | FileCheck -check-prefixes=CHECK,UNSAFEIR %s
+// RUN:   -fnative-half-arguments-and-returns -munsafe-fp-atomics | FileCheck -check-prefixes=FUN,CHECK,UNSAFEIR %s
 
 // RUN: %clang_cc1 -x hip %s -O3 -S -o - -triple=amdgcn-amd-amdhsa \
 // RUN:   -fcuda-is-device -target-cpu gfx1100 -fnative-half-type \
-// RUN:   -fnative-half-arguments-and-returns | FileCheck -check-prefix=SAFE %s
+// RUN:   -fnative-half-arguments-and-returns | FileCheck -check-prefixes=FUN,SAFE %s
 
 // RUN: %clang_cc1 -x hip %s -O3 -S -o - -triple=amdgcn-amd-amdhsa \
 // RUN:   -fcuda-is-device -target-cpu gfx940 -fnative-half-type \
 // RUN:   -fnative-half-arguments-and-returns -munsafe-fp-atomics \
-// RUN:   | FileCheck -check-prefix=UNSAFE %s
+// RUN:   | FileCheck -check-prefixes=FUN,UNSAFE %s
 
 // REQUIRES: amdgpu-registered-target
 
 #include "Inputs/cuda.h"
 #include <stdatomic.h>
 
 __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]+$}}
-
-  // SAFE: _Z4ffp1Pf
-  // SAFE: global_atomic_cmpswap
-  // SAFE: global_atomic_cmpswap
-  // SAFE: global_atomic_cmpswap
-  // SAFE: global_atomic_cmpswap
-  // SAFE: global_atomic_cmpswap
+  // FUN-LABEL: @_Z4ffp1Pf
+  // SAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 4, [[DEFMD:!amdgpu.no.fine.grained.memory ![0-9]+, !amdgpu.no.remote.memory ![0-9]+$]]
+  // SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 4, [[DEFMD]]
----------------
arsenm wrote:

This is losing the checks for end of line 

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