================ @@ -1,22 +1,113 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 4 // REQUIRES: amdgpu-registered-target // RUN: %clang_cc1 %s -emit-llvm -O0 -o - \ -// RUN: -triple=amdgcn-amd-amdhsa | opt -S | FileCheck %s +// RUN: -triple=amdgcn-amd-amdhsa | FileCheck %s +// CHECK-LABEL: define dso_local void @_Z25test_memory_fence_successv( +// CHECK-SAME: ) #[[ATTR0:[0-9]+]] { +// CHECK-NEXT: entry: +// CHECK-NEXT: fence syncscope("workgroup") seq_cst +// CHECK-NEXT: fence syncscope("agent") acquire +// CHECK-NEXT: fence seq_cst +// CHECK-NEXT: fence syncscope("agent") acq_rel +// CHECK-NEXT: fence syncscope("workgroup") release +// CHECK-NEXT: ret void +// void test_memory_fence_success() { - // CHECK-LABEL: test_memory_fence_success - // CHECK: fence syncscope("workgroup") seq_cst __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup"); - // CHECK: fence syncscope("agent") acquire __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent"); - // CHECK: fence seq_cst __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, ""); - // CHECK: fence syncscope("agent") acq_rel __builtin_amdgcn_fence(4, "agent"); - // CHECK: fence syncscope("workgroup") release __builtin_amdgcn_fence(3, "workgroup"); } + +// CHECK-LABEL: define dso_local void @_Z10test_localv( +// CHECK-SAME: ) #[[ATTR0]] { +// CHECK-NEXT: entry: +// CHECK-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META3:![0-9]+]] +// CHECK-NEXT: fence syncscope("agent") acquire, !mmra [[META3]] +// CHECK-NEXT: fence seq_cst, !mmra [[META3]] +// CHECK-NEXT: fence syncscope("agent") acq_rel, !mmra [[META3]] +// CHECK-NEXT: fence syncscope("workgroup") release, !mmra [[META3]] +// CHECK-NEXT: ret void +// +void test_local() { + __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "local"); + + __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent", "local"); + + __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "", "local"); + + __builtin_amdgcn_fence(4, "agent", "local"); + + __builtin_amdgcn_fence(3, "workgroup", "local"); +} + + +// CHECK-LABEL: define dso_local void @_Z11test_globalv( +// CHECK-SAME: ) #[[ATTR0]] { +// CHECK-NEXT: entry: +// CHECK-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META4:![0-9]+]] +// CHECK-NEXT: fence syncscope("agent") acquire, !mmra [[META4]] +// CHECK-NEXT: fence seq_cst, !mmra [[META4]] +// CHECK-NEXT: fence syncscope("agent") acq_rel, !mmra [[META4]] +// CHECK-NEXT: fence syncscope("workgroup") release, !mmra [[META4]] +// CHECK-NEXT: ret void +// +void test_global() { + __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "global"); + + __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent", "global"); + + __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "", "global"); + + __builtin_amdgcn_fence(4, "agent", "global"); + + __builtin_amdgcn_fence(3, "workgroup", "global"); +} + +// CHECK-LABEL: define dso_local void @_Z10test_imagev( +// CHECK-SAME: ) #[[ATTR0]] { +// CHECK-NEXT: entry: +// CHECK-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META5:![0-9]+]] +// CHECK-NEXT: fence syncscope("agent") acquire, !mmra [[META5]] +// CHECK-NEXT: fence seq_cst, !mmra [[META5]] +// CHECK-NEXT: fence syncscope("agent") acq_rel, !mmra [[META5]] +// CHECK-NEXT: fence syncscope("workgroup") release, !mmra [[META5]] +// CHECK-NEXT: ret void +// +void test_image() { + __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "image"); + + __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent", "image"); + + __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "", "image"); + + __builtin_amdgcn_fence(4, "agent", "image"); + + __builtin_amdgcn_fence(3, "workgroup", "image"); +} + +// CHECK-LABEL: define dso_local void @_Z10test_mixedv( +// CHECK-SAME: ) #[[ATTR0]] { +// CHECK-NEXT: entry: +// CHECK-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META6:![0-9]+]] +// CHECK-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META7:![0-9]+]] +// CHECK-NEXT: ret void +// +void test_mixed() { + __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "image", "global"); + __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "image", "local", "global"); +} ---------------- arsenm wrote:
Maybe test repeated AS name https://github.com/llvm/llvm-project/pull/78572 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits