[clang] clang/AMDGPU: Use atomicrmw for ds fmin/fmax builtins (PR #96738)

[Yaxun Liu via cfe-commits]

================
@@ -158,23 +158,85 @@ void test_ds_faddf(local float *out, float src) {
 }
 
 // CHECK-LABEL: @test_ds_fmin
-// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.ds.fmin.f32(ptr addrspace(3) %out, float %src, i32 0, i32 0, i1 false)
+// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src monotonic, align 4{{$}}
+// CHECK: atomicrmw volatile fmin ptr addrspace(3) %out, float %src monotonic, align 4{{$}}
+
+// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src acquire, align 4{{$}}
+// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src acquire, align 4{{$}}
+// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src release, align 4{{$}}
+// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src acq_rel, align 4{{$}}
+// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src seq_cst, align 4{{$}}
+// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src seq_cst, align 4{{$}}
+
+// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src syncscope("agent") monotonic, align 4{{$}}
+// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src syncscope("workgroup") monotonic, align 4{{$}}
+// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src syncscope("wavefront") monotonic, align 4{{$}}
+// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src syncscope("singlethread") monotonic, align 4{{$}}
+// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src monotonic, align 4{{$}}
+
 #if !defined(__SPIRV__)
 void test_ds_fminf(local float *out, float src) {
 #else
 void test_ds_fminf(__attribute__((address_space(3))) float *out, float src) {
 #endif
   *out = __builtin_amdgcn_ds_fminf(out, src, 0, 0, false);
+  *out = __builtin_amdgcn_ds_fminf(out, src, 0, 0, true);
+
+  // Test all orders.
+  *out = __builtin_amdgcn_ds_fminf(out, src, 1, 0, false);
----------------
yxsamliu wrote:

pls use clang predefined macros for atomic memory order and scope. same as below

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