1 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 4
2 // REQUIRES: amdgpu-registered-target
3 // RUN: %clang_cc1 %s -emit-llvm -O0 -o - \
4 // RUN: -triple=amdgcn-amd-amdhsa | FileCheck %s
5
6 // CHECK-LABEL: define dso_local void @_Z25test_memory_fence_successv(
7 // CHECK-SAME: ) #[[ATTR0:[0-9]+]] {
8 // CHECK-NEXT: entry:
9 // CHECK-NEXT: fence syncscope("workgroup") seq_cst
10 // CHECK-NEXT: fence syncscope("agent") acquire
11 // CHECK-NEXT: fence seq_cst
12 // CHECK-NEXT: fence syncscope("agent") acq_rel
13 // CHECK-NEXT: fence syncscope("workgroup") release
14 // CHECK-NEXT: ret void
15 //
test_memory_fence_success()16 void test_memory_fence_success() {
17
18 __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup");
19
20 __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent");
21
22 __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "");
23
24 __builtin_amdgcn_fence(4, "agent");
25
26 __builtin_amdgcn_fence(3, "workgroup");
27 }
28
29 // CHECK-LABEL: define dso_local void @_Z10test_localv(
30 // CHECK-SAME: ) #[[ATTR0]] {
31 // CHECK-NEXT: entry:
32 // CHECK-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META3:![0-9]+]]
33 // CHECK-NEXT: fence syncscope("agent") acquire, !mmra [[META3]]
34 // CHECK-NEXT: fence seq_cst, !mmra [[META3]]
35 // CHECK-NEXT: fence syncscope("agent") acq_rel, !mmra [[META3]]
36 // CHECK-NEXT: fence syncscope("workgroup") release, !mmra [[META3]]
37 // CHECK-NEXT: ret void
38 //
test_local()39 void test_local() {
40 __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "local");
41
42 __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent", "local");
43
44 __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "", "local");
45
46 __builtin_amdgcn_fence(4, "agent", "local");
47
48 __builtin_amdgcn_fence(3, "workgroup", "local");
49 }
50
51
52 // CHECK-LABEL: define dso_local void @_Z11test_globalv(
53 // CHECK-SAME: ) #[[ATTR0]] {
54 // CHECK-NEXT: entry:
55 // CHECK-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META4:![0-9]+]]
56 // CHECK-NEXT: fence syncscope("agent") acquire, !mmra [[META4]]
57 // CHECK-NEXT: fence seq_cst, !mmra [[META4]]
58 // CHECK-NEXT: fence syncscope("agent") acq_rel, !mmra [[META4]]
59 // CHECK-NEXT: fence syncscope("workgroup") release, !mmra [[META4]]
60 // CHECK-NEXT: ret void
61 //
test_global()62 void test_global() {
63 __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "global");
64
65 __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent", "global");
66
67 __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "", "global");
68
69 __builtin_amdgcn_fence(4, "agent", "global");
70
71 __builtin_amdgcn_fence(3, "workgroup", "global");
72 }
73
74 // CHECK-LABEL: define dso_local void @_Z10test_imagev(
75 // CHECK-SAME: ) #[[ATTR0]] {
76 // CHECK-NEXT: entry:
77 // CHECK-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META3]]
78 // CHECK-NEXT: fence syncscope("agent") acquire, !mmra [[META3]]
79 // CHECK-NEXT: fence seq_cst, !mmra [[META3]]
80 // CHECK-NEXT: fence syncscope("agent") acq_rel, !mmra [[META3]]
81 // CHECK-NEXT: fence syncscope("workgroup") release, !mmra [[META3]]
82 // CHECK-NEXT: ret void
83 //
test_image()84 void test_image() {
85 __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "local");
86
87 __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent", "local");
88
89 __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "", "local");
90
91 __builtin_amdgcn_fence(4, "agent", "local");
92
93 __builtin_amdgcn_fence(3, "workgroup", "local");
94 }
95
96 // CHECK-LABEL: define dso_local void @_Z10test_mixedv(
97 // CHECK-SAME: ) #[[ATTR0]] {
98 // CHECK-NEXT: entry:
99 // CHECK-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META5:![0-9]+]]
100 // CHECK-NEXT: fence syncscope("workgroup") seq_cst, !mmra [[META5]]
101 // CHECK-NEXT: ret void
102 //
test_mixed()103 void test_mixed() {
104 __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "local", "global");
105 __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "local", "local", "global", "local", "local");
106 }
107 //.
108 // CHECK: [[META3]] = !{!"amdgpu-as", !"local"}
109 // CHECK: [[META4]] = !{!"amdgpu-as", !"global"}
110 // CHECK: [[META5]] = !{[[META4]], [[META3]]}
111 //.
112