================ @@ -117,13 +117,44 @@ void test_update_dpp(global int* out, int arg1, int arg2) } // CHECK-LABEL: @test_ds_fadd -// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.ds.fadd.f32(ptr addrspace(3) %out, float %src, i32 0, i32 0, i1 false) +// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src monotonic, align 4{{$}} +// CHECK: atomicrmw volatile fadd ptr addrspace(3) %out, float %src monotonic, align 4{{$}} + +// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src acquire, align 4{{$}} +// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src acquire, align 4{{$}} +// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src release, align 4{{$}} +// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src acq_rel, align 4{{$}} +// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src seq_cst, align 4{{$}} +// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src seq_cst, align 4{{$}} + +// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src syncscope("agent") monotonic, align 4{{$}} +// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src syncscope("workgroup") monotonic, align 4{{$}} +// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src syncscope("wavefront") monotonic, align 4{{$}} +// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src syncscope("singlethread") monotonic, align 4{{$}} +// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src monotonic, align 4{{$}} #if !defined(__SPIRV__) void test_ds_faddf(local float *out, float src) { #else -void test_ds_faddf(__attribute__((address_space(3))) float *out, float src) { + void test_ds_faddf(__attribute__((address_space(3))) float *out, float src) { #endif + *out = __builtin_amdgcn_ds_faddf(out, src, 0, 0, false); + *out = __builtin_amdgcn_ds_faddf(out, src, 0, 0, true); + + // Test all orders. + *out = __builtin_amdgcn_ds_faddf(out, src, 1, 0, false); ---------------- yxsamliu wrote:
better use predefined macros ``` // Define macros for the C11 / C++11 memory orderings Builder.defineMacro("__ATOMIC_RELAXED", "0"); Builder.defineMacro("__ATOMIC_CONSUME", "1"); Builder.defineMacro("__ATOMIC_ACQUIRE", "2"); Builder.defineMacro("__ATOMIC_RELEASE", "3"); Builder.defineMacro("__ATOMIC_ACQ_REL", "4"); Builder.defineMacro("__ATOMIC_SEQ_CST", "5"); // Define macros for the clang atomic scopes. Builder.defineMacro("__MEMORY_SCOPE_SYSTEM", "0"); Builder.defineMacro("__MEMORY_SCOPE_DEVICE", "1"); Builder.defineMacro("__MEMORY_SCOPE_WRKGRP", "2"); Builder.defineMacro("__MEMORY_SCOPE_WVFRNT", "3"); Builder.defineMacro("__MEMORY_SCOPE_SINGLE", "4"); ``` https://github.com/llvm/llvm-project/pull/95395 _______________________________________________ llvm-branch-commits mailing list llvm-branch-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits