211 lines · cpp
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 42// REQUIRES: amdgpu-registered-target3// REQUIRES: spirv-registered-target4// RUN: %clang_cc1 %s -emit-llvm -O0 -o - \5// RUN: -triple=amdgcn-amd-amdhsa | FileCheck --check-prefix=GCN %s6// RUN: %clang_cc1 %s -emit-llvm -O0 -o - \7// RUN: -triple=spirv64-amd-amdhsa | FileCheck --check-prefix=AMDGCNSPIRV %s8 9// CHECK-LABEL: define dso_local void @_Z25test_memory_fence_successv(10// CHECK-SAME: ) #[[ATTR0:[0-9]+]] {11// CHECK-NEXT: entry:12// CHECK-NEXT: fence syncscope("workgroup") seq_cst13// CHECK-NEXT: fence syncscope("agent") acquire14// CHECK-NEXT: fence seq_cst15// CHECK-NEXT: fence syncscope("agent") acq_rel16// CHECK-NEXT: fence syncscope("workgroup") release17// CHECK-NEXT: ret void18// GCN-LABEL: define dso_local void @_Z25test_memory_fence_successv(19// GCN-SAME: ) #[[ATTR0:[0-9]+]] {20// GCN-NEXT: entry:21// GCN-NEXT: fence syncscope("workgroup") seq_cst22// GCN-NEXT: fence syncscope("agent") acquire23// GCN-NEXT: fence seq_cst24// GCN-NEXT: fence syncscope("agent") acq_rel25// GCN-NEXT: fence syncscope("workgroup") release26// GCN-NEXT: ret void27//28// AMDGCNSPIRV-LABEL: define spir_func void @_Z25test_memory_fence_successv(29// AMDGCNSPIRV-SAME: ) addrspace(4) #[[ATTR0:[0-9]+]] {30// AMDGCNSPIRV-NEXT: entry:31// AMDGCNSPIRV-NEXT: fence syncscope("workgroup") seq_cst32// AMDGCNSPIRV-NEXT: fence syncscope("device") acquire33// AMDGCNSPIRV-NEXT: fence seq_cst34// AMDGCNSPIRV-NEXT: fence syncscope("device") acq_rel35// AMDGCNSPIRV-NEXT: fence syncscope("workgroup") release36// AMDGCNSPIRV-NEXT: ret void37//38void test_memory_fence_success() {39 40 __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup");41 42 __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent");43 44 __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "");45 46 __builtin_amdgcn_fence(4, "agent");47 48 __builtin_amdgcn_fence(3, "workgroup");49}50 51// CHECK-LABEL: define dso_local void @_Z10test_localv(52// CHECK-SAME: ) #[[ATTR0]] {53// CHECK-NEXT: entry:54// CHECK-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META3:![0-9]+]]55// CHECK-NEXT: fence syncscope("agent") acquire, !mmra [[META3]]56// CHECK-NEXT: fence seq_cst, !mmra [[META3]]57// CHECK-NEXT: fence syncscope("agent") acq_rel, !mmra [[META3]]58// CHECK-NEXT: fence syncscope("workgroup") release, !mmra [[META3]]59// CHECK-NEXT: ret void60// GCN-LABEL: define dso_local void @_Z10test_localv(61// GCN-SAME: ) #[[ATTR0]] {62// GCN-NEXT: entry:63// GCN-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META3:![0-9]+]]64// GCN-NEXT: fence syncscope("agent") acquire, !mmra [[META3]]65// GCN-NEXT: fence seq_cst, !mmra [[META3]]66// GCN-NEXT: fence syncscope("agent") acq_rel, !mmra [[META3]]67// GCN-NEXT: fence syncscope("workgroup") release, !mmra [[META3]]68// GCN-NEXT: ret void69//70// AMDGCNSPIRV-LABEL: define spir_func void @_Z10test_localv(71// AMDGCNSPIRV-SAME: ) addrspace(4) #[[ATTR0]] {72// AMDGCNSPIRV-NEXT: entry:73// AMDGCNSPIRV-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META3:![0-9]+]]74// AMDGCNSPIRV-NEXT: fence syncscope("device") acquire, !mmra [[META3]]75// AMDGCNSPIRV-NEXT: fence seq_cst, !mmra [[META3]]76// AMDGCNSPIRV-NEXT: fence syncscope("device") acq_rel, !mmra [[META3]]77// AMDGCNSPIRV-NEXT: fence syncscope("workgroup") release, !mmra [[META3]]78// AMDGCNSPIRV-NEXT: ret void79//80void test_local() {81 __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "local");82 83 __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent", "local");84 85 __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "", "local");86 87 __builtin_amdgcn_fence(4, "agent", "local");88 89 __builtin_amdgcn_fence(3, "workgroup", "local");90}91 92 93// CHECK-LABEL: define dso_local void @_Z11test_globalv(94// CHECK-SAME: ) #[[ATTR0]] {95// CHECK-NEXT: entry:96// CHECK-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META4:![0-9]+]]97// CHECK-NEXT: fence syncscope("agent") acquire, !mmra [[META4]]98// CHECK-NEXT: fence seq_cst, !mmra [[META4]]99// CHECK-NEXT: fence syncscope("agent") acq_rel, !mmra [[META4]]100// CHECK-NEXT: fence syncscope("workgroup") release, !mmra [[META4]]101// CHECK-NEXT: ret void102// GCN-LABEL: define dso_local void @_Z11test_globalv(103// GCN-SAME: ) #[[ATTR0]] {104// GCN-NEXT: entry:105// GCN-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META4:![0-9]+]]106// GCN-NEXT: fence syncscope("agent") acquire, !mmra [[META4]]107// GCN-NEXT: fence seq_cst, !mmra [[META4]]108// GCN-NEXT: fence syncscope("agent") acq_rel, !mmra [[META4]]109// GCN-NEXT: fence syncscope("workgroup") release, !mmra [[META4]]110// GCN-NEXT: ret void111//112// AMDGCNSPIRV-LABEL: define spir_func void @_Z11test_globalv(113// AMDGCNSPIRV-SAME: ) addrspace(4) #[[ATTR0]] {114// AMDGCNSPIRV-NEXT: entry:115// AMDGCNSPIRV-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META4:![0-9]+]]116// AMDGCNSPIRV-NEXT: fence syncscope("device") acquire, !mmra [[META4]]117// AMDGCNSPIRV-NEXT: fence seq_cst, !mmra [[META4]]118// AMDGCNSPIRV-NEXT: fence syncscope("device") acq_rel, !mmra [[META4]]119// AMDGCNSPIRV-NEXT: fence syncscope("workgroup") release, !mmra [[META4]]120// AMDGCNSPIRV-NEXT: ret void121//122void test_global() {123 __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "global");124 125 __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent", "global");126 127 __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "", "global");128 129 __builtin_amdgcn_fence(4, "agent", "global");130 131 __builtin_amdgcn_fence(3, "workgroup", "global");132}133 134// CHECK-LABEL: define dso_local void @_Z10test_imagev(135// CHECK-SAME: ) #[[ATTR0]] {136// CHECK-NEXT: entry:137// CHECK-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META3]]138// CHECK-NEXT: fence syncscope("agent") acquire, !mmra [[META3]]139// CHECK-NEXT: fence seq_cst, !mmra [[META3]]140// CHECK-NEXT: fence syncscope("agent") acq_rel, !mmra [[META3]]141// CHECK-NEXT: fence syncscope("workgroup") release, !mmra [[META3]]142// CHECK-NEXT: ret void143// GCN-LABEL: define dso_local void @_Z10test_imagev(144// GCN-SAME: ) #[[ATTR0]] {145// GCN-NEXT: entry:146// GCN-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META3]]147// GCN-NEXT: fence syncscope("agent") acquire, !mmra [[META3]]148// GCN-NEXT: fence seq_cst, !mmra [[META3]]149// GCN-NEXT: fence syncscope("agent") acq_rel, !mmra [[META3]]150// GCN-NEXT: fence syncscope("workgroup") release, !mmra [[META3]]151// GCN-NEXT: ret void152//153// AMDGCNSPIRV-LABEL: define spir_func void @_Z10test_imagev(154// AMDGCNSPIRV-SAME: ) addrspace(4) #[[ATTR0]] {155// AMDGCNSPIRV-NEXT: entry:156// AMDGCNSPIRV-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META3]]157// AMDGCNSPIRV-NEXT: fence syncscope("device") acquire, !mmra [[META3]]158// AMDGCNSPIRV-NEXT: fence seq_cst, !mmra [[META3]]159// AMDGCNSPIRV-NEXT: fence syncscope("device") acq_rel, !mmra [[META3]]160// AMDGCNSPIRV-NEXT: fence syncscope("workgroup") release, !mmra [[META3]]161// AMDGCNSPIRV-NEXT: ret void162//163void test_image() {164 __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "local");165 166 __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent", "local");167 168 __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "", "local");169 170 __builtin_amdgcn_fence(4, "agent", "local");171 172 __builtin_amdgcn_fence(3, "workgroup", "local");173}174 175// CHECK-LABEL: define dso_local void @_Z10test_mixedv(176// CHECK-SAME: ) #[[ATTR0]] {177// CHECK-NEXT: entry:178// CHECK-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META5:![0-9]+]]179// CHECK-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META5]]180// CHECK-NEXT: ret void181// GCN-LABEL: define dso_local void @_Z10test_mixedv(182// GCN-SAME: ) #[[ATTR0]] {183// GCN-NEXT: entry:184// GCN-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META5:![0-9]+]]185// GCN-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META5]]186// GCN-NEXT: ret void187//188// AMDGCNSPIRV-LABEL: define spir_func void @_Z10test_mixedv(189// AMDGCNSPIRV-SAME: ) addrspace(4) #[[ATTR0]] {190// AMDGCNSPIRV-NEXT: entry:191// AMDGCNSPIRV-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META5:![0-9]+]]192// AMDGCNSPIRV-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META5]]193// AMDGCNSPIRV-NEXT: ret void194//195void test_mixed() {196 __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "local", "global");197 __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "local", "local", "global", "local", "local");198}199// CHECK: [[META3]] = !{!"amdgpu-synchronize-as", !"local"}200// CHECK: [[META4]] = !{!"amdgpu-synchronize-as", !"global"}201// CHECK: [[META5]] = !{[[META4]], [[META3]]}202//.203// GCN: [[META3]] = !{!"amdgpu-synchronize-as", !"local"}204// GCN: [[META4]] = !{!"amdgpu-synchronize-as", !"global"}205// GCN: [[META5]] = !{[[META4]], [[META3]]}206//.207// AMDGCNSPIRV: [[META3]] = !{!"amdgpu-synchronize-as", !"local"}208// AMDGCNSPIRV: [[META4]] = !{!"amdgpu-synchronize-as", !"global"}209// AMDGCNSPIRV: [[META5]] = !{[[META4]], [[META3]]}210//.211