1 // REQUIRES: amdgpu-registered-target
2 // RUN: %clang_cc1 %s -x hip -fcuda-is-device -emit-llvm -O0 -o - \
3 // RUN: -triple=amdgcn-amd-amdhsa | opt -S | FileCheck %s
4
test_non_volatile_parameter32(__UINT32_TYPE__ * ptr)5 __attribute__((device)) void test_non_volatile_parameter32(__UINT32_TYPE__ *ptr) {
6 // CHECK-LABEL: test_non_volatile_parameter32
7 __UINT32_TYPE__ res;
8 // CHECK: %ptr.addr = alloca i32*, align 8, addrspace(5)
9 // CHECK-NEXT: %ptr.addr.ascast = addrspacecast i32* addrspace(5)* %ptr.addr to i32**
10 // CHECK-NEXT: %res = alloca i32, align 4, addrspace(5)
11 // CHECK-NEXT: %res.ascast = addrspacecast i32 addrspace(5)* %res to i32*
12 // CHECK-NEXT: store i32* %ptr, i32** %ptr.addr.ascast, align 8
13 // CHECK-NEXT: %0 = load i32*, i32** %ptr.addr.ascast, align 8
14 // CHECK-NEXT: %1 = load i32*, i32** %ptr.addr.ascast, align 8
15 // CHECK-NEXT: %2 = load i32, i32* %1, align 4
16 // CHECK-NEXT: %3 = call i32 @llvm.amdgcn.atomic.inc.i32.p0i32(i32* %0, i32 %2, i32 7, i32 2, i1 false)
17 // CHECK-NEXT: store i32 %3, i32* %res.ascast, align 4
18 res = __builtin_amdgcn_atomic_inc32(ptr, *ptr, __ATOMIC_SEQ_CST, "workgroup");
19
20 // CHECK: %4 = load i32*, i32** %ptr.addr.ascast, align 8
21 // CHECK-NEXT: %5 = load i32*, i32** %ptr.addr.ascast, align 8
22 // CHECK-NEXT: %6 = load i32, i32* %5, align 4
23 // CHECK-NEXT: %7 = call i32 @llvm.amdgcn.atomic.dec.i32.p0i32(i32* %4, i32 %6, i32 7, i32 2, i1 false)
24 // CHECK-NEXT: store i32 %7, i32* %res.ascast, align 4
25 res = __builtin_amdgcn_atomic_dec32(ptr, *ptr, __ATOMIC_SEQ_CST, "workgroup");
26 }
27
test_non_volatile_parameter64(__UINT64_TYPE__ * ptr)28 __attribute__((device)) void test_non_volatile_parameter64(__UINT64_TYPE__ *ptr) {
29 // CHECK-LABEL: test_non_volatile_parameter64
30 __UINT64_TYPE__ res;
31 // CHECK: %ptr.addr = alloca i64*, align 8, addrspace(5)
32 // CHECK-NEXT: %ptr.addr.ascast = addrspacecast i64* addrspace(5)* %ptr.addr to i64**
33 // CHECK-NEXT: %res = alloca i64, align 8, addrspace(5)
34 // CHECK-NEXT: %res.ascast = addrspacecast i64 addrspace(5)* %res to i64*
35 // CHECK-NEXT: store i64* %ptr, i64** %ptr.addr.ascast, align 8
36 // CHECK-NEXT: %0 = load i64*, i64** %ptr.addr.ascast, align 8
37 // CHECK-NEXT: %1 = load i64*, i64** %ptr.addr.ascast, align 8
38 // CHECK-NEXT: %2 = load i64, i64* %1, align 8
39 // CHECK-NEXT: %3 = call i64 @llvm.amdgcn.atomic.inc.i64.p0i64(i64* %0, i64 %2, i32 7, i32 2, i1 false)
40 // CHECK-NEXT: store i64 %3, i64* %res.ascast, align 8
41 res = __builtin_amdgcn_atomic_inc64(ptr, *ptr, __ATOMIC_SEQ_CST, "workgroup");
42
43 // CHECK: %4 = load i64*, i64** %ptr.addr.ascast, align 8
44 // CHECK-NEXT: %5 = load i64*, i64** %ptr.addr.ascast, align 8
45 // CHECK-NEXT: %6 = load i64, i64* %5, align 8
46 // CHECK-NEXT: %7 = call i64 @llvm.amdgcn.atomic.dec.i64.p0i64(i64* %4, i64 %6, i32 7, i32 2, i1 false)
47 // CHECK-NEXT: store i64 %7, i64* %res.ascast, align 8
48 res = __builtin_amdgcn_atomic_dec64(ptr, *ptr, __ATOMIC_SEQ_CST, "workgroup");
49 }
50
test_volatile_parameter32(volatile __UINT32_TYPE__ * ptr)51 __attribute__((device)) void test_volatile_parameter32(volatile __UINT32_TYPE__ *ptr) {
52 // CHECK-LABEL: test_volatile_parameter32
53 __UINT32_TYPE__ res;
54 // CHECK: %ptr.addr = alloca i32*, align 8, addrspace(5)
55 // CHECK-NEXT: %ptr.addr.ascast = addrspacecast i32* addrspace(5)* %ptr.addr to i32**
56 // CHECK-NEXT: %res = alloca i32, align 4, addrspace(5)
57 // CHECK-NEXT: %res.ascast = addrspacecast i32 addrspace(5)* %res to i32*
58 // CHECK-NEXT: store i32* %ptr, i32** %ptr.addr.ascast, align 8
59 // CHECK-NEXT: %0 = load i32*, i32** %ptr.addr.ascast, align 8
60 // CHECK-NEXT: %1 = load i32*, i32** %ptr.addr.ascast, align 8
61 // CHECK-NEXT: %2 = load volatile i32, i32* %1, align 4
62 // CHECK-NEXT: %3 = call i32 @llvm.amdgcn.atomic.inc.i32.p0i32(i32* %0, i32 %2, i32 7, i32 2, i1 true)
63 // CHECK-NEXT: store i32 %3, i32* %res.ascast, align 4
64 res = __builtin_amdgcn_atomic_inc32(ptr, *ptr, __ATOMIC_SEQ_CST, "workgroup");
65
66 // CHECK: %4 = load i32*, i32** %ptr.addr.ascast, align 8
67 // CHECK-NEXT: %5 = load i32*, i32** %ptr.addr.ascast, align 8
68 // CHECK-NEXT: %6 = load volatile i32, i32* %5, align 4
69 // CHECK-NEXT: %7 = call i32 @llvm.amdgcn.atomic.dec.i32.p0i32(i32* %4, i32 %6, i32 7, i32 2, i1 true)
70 // CHECK-NEXT: store i32 %7, i32* %res.ascast, align 4
71 res = __builtin_amdgcn_atomic_dec32(ptr, *ptr, __ATOMIC_SEQ_CST, "workgroup");
72 }
73
test_volatile_parameter64(volatile __UINT64_TYPE__ * ptr)74 __attribute__((device)) void test_volatile_parameter64(volatile __UINT64_TYPE__ *ptr) {
75 // CHECK-LABEL: test_volatile_parameter64
76 __UINT64_TYPE__ res;
77 // CHECK: %ptr.addr = alloca i64*, align 8, addrspace(5)
78 // CHECK-NEXT: %ptr.addr.ascast = addrspacecast i64* addrspace(5)* %ptr.addr to i64**
79 // CHECK-NEXT: %res = alloca i64, align 8, addrspace(5)
80 // CHECK-NEXT: %res.ascast = addrspacecast i64 addrspace(5)* %res to i64*
81 // CHECK-NEXT: store i64* %ptr, i64** %ptr.addr.ascast, align 8
82 // CHECK-NEXT: %0 = load i64*, i64** %ptr.addr.ascast, align 8
83 // CHECK-NEXT: %1 = load i64*, i64** %ptr.addr.ascast, align 8
84 // CHECK-NEXT: %2 = load volatile i64, i64* %1, align 8
85 // CHECK-NEXT: %3 = call i64 @llvm.amdgcn.atomic.inc.i64.p0i64(i64* %0, i64 %2, i32 7, i32 2, i1 true)
86 // CHECK-NEXT: store i64 %3, i64* %res.ascast, align 8
87 res = __builtin_amdgcn_atomic_inc64(ptr, *ptr, __ATOMIC_SEQ_CST, "workgroup");
88
89 // CHECK: %4 = load i64*, i64** %ptr.addr.ascast, align 8
90 // CHECK-NEXT: %5 = load i64*, i64** %ptr.addr.ascast, align 8
91 // CHECK-NEXT: %6 = load volatile i64, i64* %5, align 8
92 // CHECK-NEXT: %7 = call i64 @llvm.amdgcn.atomic.dec.i64.p0i64(i64* %4, i64 %6, i32 7, i32 2, i1 true)
93 // CHECK-NEXT: store i64 %7, i64* %res.ascast, align 8
94 res = __builtin_amdgcn_atomic_dec64(ptr, *ptr, __ATOMIC_SEQ_CST, "workgroup");
95 }
96
test_shared32()97 __attribute__((device)) void test_shared32() {
98 // CHECK-LABEL: test_shared32
99 __attribute__((shared)) __UINT32_TYPE__ val;
100
101 // CHECK: %0 = load i32, i32* addrspacecast (i32 addrspace(3)* @_ZZ13test_shared32vE3val to i32*), align 4
102 // CHECK-NEXT: %1 = call i32 @llvm.amdgcn.atomic.inc.i32.p0i32(i32* addrspacecast (i32 addrspace(3)* @_ZZ13test_shared32vE3val to i32*), i32 %0, i32 7, i32 2, i1 false)
103 // CHECK-NEXT: store i32 %1, i32* addrspacecast (i32 addrspace(3)* @_ZZ13test_shared32vE3val to i32*), align 4
104 val = __builtin_amdgcn_atomic_inc32(&val, val, __ATOMIC_SEQ_CST, "workgroup");
105
106 // CHECK: %2 = load i32, i32* addrspacecast (i32 addrspace(3)* @_ZZ13test_shared32vE3val to i32*), align 4
107 // CHECK-NEXT: %3 = call i32 @llvm.amdgcn.atomic.dec.i32.p0i32(i32* addrspacecast (i32 addrspace(3)* @_ZZ13test_shared32vE3val to i32*), i32 %2, i32 7, i32 2, i1 false)
108 // CHECK-NEXT: store i32 %3, i32* addrspacecast (i32 addrspace(3)* @_ZZ13test_shared32vE3val to i32*), align 4
109 val = __builtin_amdgcn_atomic_dec32(&val, val, __ATOMIC_SEQ_CST, "workgroup");
110 }
111
test_shared64()112 __attribute__((device)) void test_shared64() {
113 // CHECK-LABEL: test_shared64
114 __attribute__((shared)) __UINT64_TYPE__ val;
115
116 // CHECK: %0 = load i64, i64* addrspacecast (i64 addrspace(3)* @_ZZ13test_shared64vE3val to i64*), align 8
117 // CHECK-NEXT: %1 = call i64 @llvm.amdgcn.atomic.inc.i64.p0i64(i64* addrspacecast (i64 addrspace(3)* @_ZZ13test_shared64vE3val to i64*), i64 %0, i32 7, i32 2, i1 false)
118 // CHECK-NEXT: store i64 %1, i64* addrspacecast (i64 addrspace(3)* @_ZZ13test_shared64vE3val to i64*), align 8
119 val = __builtin_amdgcn_atomic_inc64(&val, val, __ATOMIC_SEQ_CST, "workgroup");
120
121 // CHECK: %2 = load i64, i64* addrspacecast (i64 addrspace(3)* @_ZZ13test_shared64vE3val to i64*), align 8
122 // CHECK-NEXT: %3 = call i64 @llvm.amdgcn.atomic.dec.i64.p0i64(i64* addrspacecast (i64 addrspace(3)* @_ZZ13test_shared64vE3val to i64*), i64 %2, i32 7, i32 2, i1 false)
123 // CHECK-NEXT: store i64 %3, i64* addrspacecast (i64 addrspace(3)* @_ZZ13test_shared64vE3val to i64*), align 8
124 val = __builtin_amdgcn_atomic_dec64(&val, val, __ATOMIC_SEQ_CST, "workgroup");
125 }
126
127 __attribute__((device)) __UINT32_TYPE__ global_val32;
test_global32()128 __attribute__((device)) void test_global32() {
129 // CHECK-LABEL: test_global32
130 // CHECK: %0 = load i32, i32* addrspacecast (i32 addrspace(1)* @global_val32 to i32*), align 4
131 // CHECK-NEXT: %1 = call i32 @llvm.amdgcn.atomic.inc.i32.p0i32(i32* addrspacecast (i32 addrspace(1)* @global_val32 to i32*), i32 %0, i32 7, i32 2, i1 false)
132 // CHECK-NEXT: store i32 %1, i32* addrspacecast (i32 addrspace(1)* @global_val32 to i32*), align 4
133 global_val32 = __builtin_amdgcn_atomic_inc32(&global_val32, global_val32, __ATOMIC_SEQ_CST, "workgroup");
134
135 // CHECK: %2 = load i32, i32* addrspacecast (i32 addrspace(1)* @global_val32 to i32*), align 4
136 // CHECK-NEXT: %3 = call i32 @llvm.amdgcn.atomic.dec.i32.p0i32(i32* addrspacecast (i32 addrspace(1)* @global_val32 to i32*), i32 %2, i32 7, i32 2, i1 false)
137 // CHECK-NEXT: store i32 %3, i32* addrspacecast (i32 addrspace(1)* @global_val32 to i32*), align 4
138 global_val32 = __builtin_amdgcn_atomic_dec32(&global_val32, global_val32, __ATOMIC_SEQ_CST, "workgroup");
139 }
140
141 __attribute__((device)) __UINT64_TYPE__ global_val64;
test_global64()142 __attribute__((device)) void test_global64() {
143 // CHECK-LABEL: test_global64
144 // CHECK: %0 = load i64, i64* addrspacecast (i64 addrspace(1)* @global_val64 to i64*), align 8
145 // CHECK-NEXT: %1 = call i64 @llvm.amdgcn.atomic.inc.i64.p0i64(i64* addrspacecast (i64 addrspace(1)* @global_val64 to i64*), i64 %0, i32 7, i32 2, i1 false)
146 // CHECK-NEXT: store i64 %1, i64* addrspacecast (i64 addrspace(1)* @global_val64 to i64*), align 8
147 global_val64 = __builtin_amdgcn_atomic_inc64(&global_val64, global_val64, __ATOMIC_SEQ_CST, "workgroup");
148
149 // CHECK: %2 = load i64, i64* addrspacecast (i64 addrspace(1)* @global_val64 to i64*), align 8
150 // CHECK-NEXT: %3 = call i64 @llvm.amdgcn.atomic.dec.i64.p0i64(i64* addrspacecast (i64 addrspace(1)* @global_val64 to i64*), i64 %2, i32 7, i32 2, i1 false)
151 // CHECK-NEXT: store i64 %3, i64* addrspacecast (i64 addrspace(1)* @global_val64 to i64*), align 8
152 global_val64 = __builtin_amdgcn_atomic_dec64(&global_val64, global_val64, __ATOMIC_SEQ_CST, "workgroup");
153 }
154
155 __attribute__((constant)) __UINT32_TYPE__ cval32;
test_constant32()156 __attribute__((device)) void test_constant32() {
157 // CHECK-LABEL: test_constant32
158 __UINT32_TYPE__ local_val;
159
160 // CHECK: %0 = load i32, i32* addrspacecast (i32 addrspace(4)* @cval32 to i32*), align 4
161 // CHECK-NEXT: %1 = call i32 @llvm.amdgcn.atomic.inc.i32.p0i32(i32* addrspacecast (i32 addrspace(4)* @cval32 to i32*), i32 %0, i32 7, i32 2, i1 false)
162 // CHECK-NEXT: store i32 %1, i32* %local_val.ascast, align 4
163 local_val = __builtin_amdgcn_atomic_inc32(&cval32, cval32, __ATOMIC_SEQ_CST, "workgroup");
164
165 // CHECK: %2 = load i32, i32* addrspacecast (i32 addrspace(4)* @cval32 to i32*), align 4
166 // CHECK-NEXT: %3 = call i32 @llvm.amdgcn.atomic.dec.i32.p0i32(i32* addrspacecast (i32 addrspace(4)* @cval32 to i32*), i32 %2, i32 7, i32 2, i1 false)
167 // CHECK-NEXT: store i32 %3, i32* %local_val.ascast, align 4
168 local_val = __builtin_amdgcn_atomic_dec32(&cval32, cval32, __ATOMIC_SEQ_CST, "workgroup");
169 }
170
171 __attribute__((constant)) __UINT64_TYPE__ cval64;
test_constant64()172 __attribute__((device)) void test_constant64() {
173 // CHECK-LABEL: test_constant64
174 __UINT64_TYPE__ local_val;
175
176 // CHECK: %0 = load i64, i64* addrspacecast (i64 addrspace(4)* @cval64 to i64*), align 8
177 // CHECK-NEXT: %1 = call i64 @llvm.amdgcn.atomic.inc.i64.p0i64(i64* addrspacecast (i64 addrspace(4)* @cval64 to i64*), i64 %0, i32 7, i32 2, i1 false)
178 // CHECK-NEXT: store i64 %1, i64* %local_val.ascast, align 8
179 local_val = __builtin_amdgcn_atomic_inc64(&cval64, cval64, __ATOMIC_SEQ_CST, "workgroup");
180
181 // CHECK: %2 = load i64, i64* addrspacecast (i64 addrspace(4)* @cval64 to i64*), align 8
182 // CHECK-NEXT: %3 = call i64 @llvm.amdgcn.atomic.dec.i64.p0i64(i64* addrspacecast (i64 addrspace(4)* @cval64 to i64*), i64 %2, i32 7, i32 2, i1 false)
183 // CHECK-NEXT: store i64 %3, i64* %local_val.ascast, align 8
184 local_val = __builtin_amdgcn_atomic_dec64(&cval64, cval64, __ATOMIC_SEQ_CST, "workgroup");
185 }
186
test_order32()187 __attribute__((device)) void test_order32() {
188 // CHECK-LABEL: test_order32
189 __attribute__((shared)) __UINT32_TYPE__ val;
190
191 // CHECK: call i32 @llvm.amdgcn.atomic.inc.i32.p0i32(i32* addrspacecast (i32 addrspace(3)* @_ZZ12test_order32vE3val to i32*), i32 {{.*}}, i32 2, i32 2, i1 false)
192 val = __builtin_amdgcn_atomic_inc32(&val, val, __ATOMIC_RELAXED, "workgroup");
193
194 // CHECK: call i32 @llvm.amdgcn.atomic.inc.i32.p0i32(i32* addrspacecast (i32 addrspace(3)* @_ZZ12test_order32vE3val to i32*), i32 {{.*}}, i32 4, i32 2, i1 false)
195 val = __builtin_amdgcn_atomic_inc32(&val, val, __ATOMIC_CONSUME, "workgroup");
196
197 // CHECK: call i32 @llvm.amdgcn.atomic.inc.i32.p0i32(i32* addrspacecast (i32 addrspace(3)* @_ZZ12test_order32vE3val to i32*), i32 {{.*}}, i32 4, i32 2, i1 false)
198 val = __builtin_amdgcn_atomic_inc32(&val, val, __ATOMIC_ACQUIRE, "workgroup");
199
200 // CHECK: call i32 @llvm.amdgcn.atomic.dec.i32.p0i32(i32* addrspacecast (i32 addrspace(3)* @_ZZ12test_order32vE3val to i32*), i32 {{.*}}, i32 5, i32 2, i1 false)
201 val = __builtin_amdgcn_atomic_dec32(&val, val, __ATOMIC_RELEASE, "workgroup");
202
203 // CHECK: call i32 @llvm.amdgcn.atomic.dec.i32.p0i32(i32* addrspacecast (i32 addrspace(3)* @_ZZ12test_order32vE3val to i32*), i32 {{.*}}, i32 6, i32 2, i1 false)
204 val = __builtin_amdgcn_atomic_dec32(&val, val, __ATOMIC_ACQ_REL, "workgroup");
205
206 // CHECK: call i32 @llvm.amdgcn.atomic.dec.i32.p0i32(i32* addrspacecast (i32 addrspace(3)* @_ZZ12test_order32vE3val to i32*), i32 {{.*}}, i32 7, i32 2, i1 false)
207 val = __builtin_amdgcn_atomic_dec32(&val, val, __ATOMIC_SEQ_CST, "workgroup");
208 }
209
test_order64()210 __attribute__((device)) void test_order64() {
211 // CHECK-LABEL: test_order64
212 __attribute__((shared)) __UINT64_TYPE__ val;
213
214 // CHECK: call i64 @llvm.amdgcn.atomic.inc.i64.p0i64(i64* addrspacecast (i64 addrspace(3)* @_ZZ12test_order64vE3val to i64*), i64 {{.*}}, i32 2, i32 2, i1 false)
215 val = __builtin_amdgcn_atomic_inc64(&val, val, __ATOMIC_RELAXED, "workgroup");
216
217 // CHECK: call i64 @llvm.amdgcn.atomic.dec.i64.p0i64(i64* addrspacecast (i64 addrspace(3)* @_ZZ12test_order64vE3val to i64*), i64 {{.*}}, i32 4, i32 2, i1 false)
218 val = __builtin_amdgcn_atomic_dec64(&val, val, __ATOMIC_CONSUME, "workgroup");
219
220 // CHECK: call i64 @llvm.amdgcn.atomic.inc.i64.p0i64(i64* addrspacecast (i64 addrspace(3)* @_ZZ12test_order64vE3val to i64*), i64 {{.*}}, i32 4, i32 2, i1 false)
221 val = __builtin_amdgcn_atomic_inc64(&val, val, __ATOMIC_ACQUIRE, "workgroup");
222
223 // CHECK: call i64 @llvm.amdgcn.atomic.dec.i64.p0i64(i64* addrspacecast (i64 addrspace(3)* @_ZZ12test_order64vE3val to i64*), i64 {{.*}}, i32 5, i32 2, i1 false)
224 val = __builtin_amdgcn_atomic_dec64(&val, val, __ATOMIC_RELEASE, "workgroup");
225
226 // CHECK: call i64 @llvm.amdgcn.atomic.dec.i64.p0i64(i64* addrspacecast (i64 addrspace(3)* @_ZZ12test_order64vE3val to i64*), i64 {{.*}}, i32 6, i32 2, i1 false)
227 val = __builtin_amdgcn_atomic_dec64(&val, val, __ATOMIC_ACQ_REL, "workgroup");
228
229 // CHECK: call i64 @llvm.amdgcn.atomic.dec.i64.p0i64(i64* addrspacecast (i64 addrspace(3)* @_ZZ12test_order64vE3val to i64*), i64 {{.*}}, i32 7, i32 2, i1 false)
230 val = __builtin_amdgcn_atomic_dec64(&val, val, __ATOMIC_SEQ_CST, "workgroup");
231 }
232
test_scope32()233 __attribute__((device)) void test_scope32() {
234 // CHECK-LABEL: test_scope32
235 __attribute__((shared)) __UINT32_TYPE__ val;
236
237 // CHECK: %1 = call i32 @llvm.amdgcn.atomic.inc.i32.p0i32(i32* addrspacecast (i32 addrspace(3)* @_ZZ12test_scope32vE3val to i32*), i32 %0, i32 7, i32 1, i1 false)
238 val = __builtin_amdgcn_atomic_inc32(&val, val, __ATOMIC_SEQ_CST, "");
239
240 // CHECK: %3 = call i32 @llvm.amdgcn.atomic.dec.i32.p0i32(i32* addrspacecast (i32 addrspace(3)* @_ZZ12test_scope32vE3val to i32*), i32 %2, i32 7, i32 2, i1 false)
241 val = __builtin_amdgcn_atomic_dec32(&val, val, __ATOMIC_SEQ_CST, "workgroup");
242
243 // CHECK: %5 = call i32 @llvm.amdgcn.atomic.dec.i32.p0i32(i32* addrspacecast (i32 addrspace(3)* @_ZZ12test_scope32vE3val to i32*), i32 %4, i32 7, i32 3, i1 false)
244 val = __builtin_amdgcn_atomic_dec32(&val, val, __ATOMIC_SEQ_CST, "agent");
245
246 // CHECK: %7 = call i32 @llvm.amdgcn.atomic.dec.i32.p0i32(i32* addrspacecast (i32 addrspace(3)* @_ZZ12test_scope32vE3val to i32*), i32 %6, i32 7, i32 4, i1 false)
247 val = __builtin_amdgcn_atomic_dec32(&val, val, __ATOMIC_SEQ_CST, "wavefront");
248 }
249
test_scope64()250 __attribute__((device)) void test_scope64() {
251 // CHECK-LABEL: test_scope64
252 __attribute__((shared)) __UINT64_TYPE__ val;
253
254 // CHECK: %1 = call i64 @llvm.amdgcn.atomic.inc.i64.p0i64(i64* addrspacecast (i64 addrspace(3)* @_ZZ12test_scope64vE3val to i64*), i64 %0, i32 7, i32 1, i1 false)
255 val = __builtin_amdgcn_atomic_inc64(&val, val, __ATOMIC_SEQ_CST, "");
256
257 // CHECK: %3 = call i64 @llvm.amdgcn.atomic.dec.i64.p0i64(i64* addrspacecast (i64 addrspace(3)* @_ZZ12test_scope64vE3val to i64*), i64 %2, i32 7, i32 2, i1 false)
258 val = __builtin_amdgcn_atomic_dec64(&val, val, __ATOMIC_SEQ_CST, "workgroup");
259
260 // CHECK: %5 = call i64 @llvm.amdgcn.atomic.dec.i64.p0i64(i64* addrspacecast (i64 addrspace(3)* @_ZZ12test_scope64vE3val to i64*), i64 %4, i32 7, i32 3, i1 false)
261 val = __builtin_amdgcn_atomic_dec64(&val, val, __ATOMIC_SEQ_CST, "agent");
262
263 // CHECK: %7 = call i64 @llvm.amdgcn.atomic.dec.i64.p0i64(i64* addrspacecast (i64 addrspace(3)* @_ZZ12test_scope64vE3val to i64*), i64 %6, i32 7, i32 4, i1 false)
264 val = __builtin_amdgcn_atomic_dec64(&val, val, __ATOMIC_SEQ_CST, "wavefront");
265 }
266