1 // Test target codegen - host bc file has to be created first.
2 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
3 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -disable-llvm-optzns | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix SEQ
4 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -disable-llvm-optzns -fopenmp-cuda-parallel-target-regions | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix PAR
5 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
6 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - -disable-llvm-optzns | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix SEQ
7 // RUN: %clang_cc1 -verify -fopenmp -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - -disable-llvm-optzns | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix SEQ
8 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - -disable-llvm-optzns -fopenmp-cuda-parallel-target-regions | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix PAR
9 // RUN: %clang_cc1 -verify -fopenmp -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - -disable-llvm-optzns -fopenmp-cuda-parallel-target-regions | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix PAR
10 // expected-no-diagnostics
11 #ifndef HEADER
12 #define HEADER
13
14 template<typename tx>
ftemplate(int n)15 tx ftemplate(int n) {
16 tx a = 0;
17 short aa = 0;
18 tx b[10];
19
20 #pragma omp target if(0)
21 {
22 #pragma omp parallel
23 {
24 int a = 41;
25 }
26 a += 1;
27 }
28
29 #pragma omp target
30 {
31 #pragma omp parallel
32 {
33 int a = 42;
34 }
35 #pragma omp parallel if(0)
36 {
37 int a = 43;
38 }
39 #pragma omp parallel if(1)
40 {
41 int a = 44;
42 }
43 a += 1;
44 }
45
46 #pragma omp target if(n>40)
47 {
48 #pragma omp parallel if(n>1000)
49 {
50 int a = 45;
51 #pragma omp barrier
52 }
53 a += 1;
54 aa += 1;
55 b[2] += 1;
56 }
57
58 #pragma omp target
59 {
60 #pragma omp parallel
61 {
62 #pragma omp critical
63 ++a;
64 }
65 ++a;
66 }
67 return a;
68 }
69
bar(int n)70 int bar(int n){
71 int a = 0;
72
73 a += ftemplate<int>(n);
74
75 return a;
76 }
77
78 // SEQ: [[MEM_TY:%.+]] = type { [128 x i8] }
79 // SEQ-DAG: [[SHARED_GLOBAL_RD:@.+]] = weak addrspace(3) global [[MEM_TY]] undef
80 // SEQ-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* undef
81 // SEQ-DAG: [[KERNEL_SIZE:@.+]] = internal unnamed_addr constant i{{64|32}} 4
82 // SEQ-DAG: [[KERNEL_SHARED:@.+]] = internal unnamed_addr constant i16 1
83
84 // CHECK-NOT: define {{.*}}void {{@__omp_offloading_.+template.+l20}}_worker()
85
86 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l29}}_worker()
87 // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
88 // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
89 // CHECK: store i8* null, i8** [[OMP_WORK_FN]],
90 // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]],
91 // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
92 //
93 // CHECK: [[AWAIT_WORK]]
94 // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
95 // CHECK: [[KPR:%.+]] = call i1 @__kmpc_kernel_parallel(i8** [[OMP_WORK_FN]])
96 // CHECK: [[KPRB:%.+]] = zext i1 [[KPR]] to i8
97 // store i8 [[KPRB]], i8* [[OMP_EXEC_STATUS]], align 1
98 // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]],
99 // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null
100 // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]]
101 //
102 // CHECK: [[SEL_WORKERS]]
103 // CHECK: [[ST:%.+]] = load i8, i8* [[OMP_EXEC_STATUS]]
104 // CHECK: [[IS_ACTIVE:%.+]] = icmp ne i8 [[ST]], 0
105 // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]]
106 //
107 // CHECK: [[EXEC_PARALLEL]]
108 // CHECK: [[WF1:%.+]] = load i8*, i8** [[OMP_WORK_FN]],
109 // CHECK: [[WM1:%.+]] = icmp eq i8* [[WF1]], bitcast (void (i16, i32)* [[PARALLEL_FN1:@.+]]_wrapper to i8*)
110 // CHECK: br i1 [[WM1]], label {{%?}}[[EXEC_PFN1:.+]], label {{%?}}[[CHECK_NEXT1:.+]]
111 //
112 // CHECK: [[EXEC_PFN1]]
113 // CHECK: call void [[PARALLEL_FN1]]_wrapper(
114 // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
115 //
116 // CHECK: [[CHECK_NEXT1]]
117 // CHECK: [[WF2:%.+]] = load i8*, i8** [[OMP_WORK_FN]],
118 // CHECK: [[WM2:%.+]] = icmp eq i8* [[WF2]], bitcast (void (i16, i32)* [[PARALLEL_FN2:@.+]]_wrapper to i8*)
119 // CHECK: br i1 [[WM2]], label {{%?}}[[EXEC_PFN2:.+]], label {{%?}}[[CHECK_NEXT2:.+]]
120 //
121 // CHECK: [[EXEC_PFN2]]
122 // CHECK: call void [[PARALLEL_FN2]]_wrapper(
123 // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
124 //
125 // CHECK: [[CHECK_NEXT2]]
126 // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
127 //
128 // CHECK: [[TERM_PARALLEL]]
129 // CHECK: call void @__kmpc_kernel_end_parallel()
130 // CHECK: br label {{%?}}[[BAR_PARALLEL]]
131 //
132 // CHECK: [[BAR_PARALLEL]]
133 // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
134 // CHECK: br label {{%?}}[[AWAIT_WORK]]
135 //
136 // CHECK: [[EXIT]]
137 // CHECK: ret void
138
139 // CHECK: define {{.*}}void [[T6:@__omp_offloading_.+template.+l29]](i[[SZ:32|64]]
140 // Create local storage for each capture.
141 // CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]],
142 // CHECK-DAG: store i[[SZ]] [[ARG_A:%.+]], i[[SZ]]* [[LOCAL_A]]
143 // Store captures in the context.
144 // CHECK-64-DAG:[[REF_A:%.+]] = bitcast i[[SZ]]* [[LOCAL_A]] to i32*
145 //
146 // CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
147 // CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
148 // CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
149 // CHECK-DAG: [[TH_LIMIT:%.+]] = sub nuw i32 [[NTH]], [[WS]]
150 // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[TH_LIMIT]]
151 // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[CHECK_MASTER:.+]]
152 //
153 // CHECK: [[WORKER]]
154 // CHECK: {{call|invoke}} void [[T6]]_worker()
155 // CHECK: br label {{%?}}[[EXIT:.+]]
156 //
157 // CHECK: [[CHECK_MASTER]]
158 // CHECK-DAG: [[CMTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
159 // CHECK-DAG: [[CMNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
160 // CHECK-DAG: [[CMWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
161 // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]],
162 // CHECK: br i1 [[IS_MASTER]], label {{%?}}[[MASTER:.+]], label {{%?}}[[EXIT]]
163 //
164 // CHECK: [[MASTER]]
165 // CHECK-DAG: [[MNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
166 // CHECK-DAG: [[MWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
167 // CHECK: [[MTMP1:%.+]] = sub nuw i32 [[MNTH]], [[MWS]]
168 // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]]
169 // CHECK: call void @__kmpc_kernel_prepare_parallel(i8* bitcast (void (i16, i32)* [[PARALLEL_FN1]]_wrapper to i8*))
170 // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
171 // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
172 // CHECK: call void @__kmpc_serialized_parallel(
173 // CHECK: {{call|invoke}} void [[PARALLEL_FN3:@.+]](
174 // CHECK: call void @__kmpc_end_serialized_parallel(
175 // CHECK: call void @__kmpc_kernel_prepare_parallel(i8* bitcast (void (i16, i32)* [[PARALLEL_FN2]]_wrapper to i8*))
176 // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
177 // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
178 // CHECK-64-DAG: load i32, i32* [[REF_A]]
179 // CHECK-32-DAG: load i32, i32* [[LOCAL_A]]
180 // CHECK: br label {{%?}}[[TERMINATE:.+]]
181 //
182 // CHECK: [[TERMINATE]]
183 // CHECK: call void @__kmpc_kernel_deinit(
184 // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
185 // CHECK: br label {{%?}}[[EXIT]]
186 //
187 // CHECK: [[EXIT]]
188 // CHECK: ret void
189
190 // CHECK-DAG: define internal void [[PARALLEL_FN1]](
191 // CHECK: [[A:%.+]] = alloca i[[SZ:32|64]],
192 // CHECK: store i[[SZ]] 42, i[[SZ]]* %a,
193 // CHECK: ret void
194
195 // CHECK-DAG: define internal void [[PARALLEL_FN3]](
196 // CHECK: [[A:%.+]] = alloca i[[SZ:32|64]],
197 // CHECK: store i[[SZ]] 43, i[[SZ]]* %a,
198 // CHECK: ret void
199
200 // CHECK-DAG: define internal void [[PARALLEL_FN2]](
201 // CHECK: [[A:%.+]] = alloca i[[SZ:32|64]],
202 // CHECK: store i[[SZ]] 44, i[[SZ]]* %a,
203 // CHECK: ret void
204
205 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l46}}_worker()
206 // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
207 // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
208 // CHECK: store i8* null, i8** [[OMP_WORK_FN]],
209 // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]],
210 // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
211 //
212 // CHECK: [[AWAIT_WORK]]
213 // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
214 // CHECK: [[KPR:%.+]] = call i1 @__kmpc_kernel_parallel(i8** [[OMP_WORK_FN]])
215 // CHECK: [[KPRB:%.+]] = zext i1 [[KPR]] to i8
216 // store i8 [[KPRB]], i8* [[OMP_EXEC_STATUS]], align 1
217 // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]],
218 // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null
219 // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]]
220 //
221 // CHECK: [[SEL_WORKERS]]
222 // CHECK: [[ST:%.+]] = load i8, i8* [[OMP_EXEC_STATUS]]
223 // CHECK: [[IS_ACTIVE:%.+]] = icmp ne i8 [[ST]], 0
224 // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]]
225 //
226 // CHECK: [[EXEC_PARALLEL]]
227 // CHECK: [[WF:%.+]] = load i8*, i8** [[OMP_WORK_FN]],
228 // CHECK: [[WM:%.+]] = icmp eq i8* [[WF]], bitcast (void (i16, i32)* [[PARALLEL_FN4:@.+]]_wrapper to i8*)
229 // CHECK: br i1 [[WM]], label {{%?}}[[EXEC_PFN:.+]], label {{%?}}[[CHECK_NEXT:.+]]
230 //
231 // CHECK: [[EXEC_PFN]]
232 // CHECK: call void [[PARALLEL_FN4]]_wrapper(
233 // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
234 //
235 // CHECK: [[CHECK_NEXT]]
236 // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
237 //
238 // CHECK: [[TERM_PARALLEL]]
239 // CHECK: call void @__kmpc_kernel_end_parallel()
240 // CHECK: br label {{%?}}[[BAR_PARALLEL]]
241 //
242 // CHECK: [[BAR_PARALLEL]]
243 // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
244 // CHECK: br label {{%?}}[[AWAIT_WORK]]
245 //
246 // CHECK: [[EXIT]]
247 // CHECK: ret void
248
249 // CHECK: define {{.*}}void [[T6:@__omp_offloading_.+template.+l46]](i[[SZ:32|64]]
250 // Create local storage for each capture.
251 // CHECK: [[LOCAL_N:%.+]] = alloca i[[SZ]],
252 // CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]],
253 // CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]],
254 // CHECK: [[LOCAL_B:%.+]] = alloca [10 x i32]*
255 // CHECK-DAG: store i[[SZ]] [[ARG_N:%.+]], i[[SZ]]* [[LOCAL_N]]
256 // CHECK-DAG: store i[[SZ]] [[ARG_A:%.+]], i[[SZ]]* [[LOCAL_A]]
257 // CHECK-DAG: store i[[SZ]] [[ARG_AA:%.+]], i[[SZ]]* [[LOCAL_AA]]
258 // CHECK-DAG: store [10 x i32]* [[ARG_B:%.+]], [10 x i32]** [[LOCAL_B]]
259 // Store captures in the context.
260 // CHECK-64-DAG:[[REF_N:%.+]] = bitcast i[[SZ]]* [[LOCAL_N]] to i32*
261 // CHECK-64-DAG:[[REF_A:%.+]] = bitcast i[[SZ]]* [[LOCAL_A]] to i32*
262 // CHECK-DAG: [[REF_AA:%.+]] = bitcast i[[SZ]]* [[LOCAL_AA]] to i16*
263 // CHECK-DAG: [[REF_B:%.+]] = load [10 x i32]*, [10 x i32]** [[LOCAL_B]],
264 //
265 // CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
266 // CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
267 // CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
268 // CHECK-DAG: [[TH_LIMIT:%.+]] = sub nuw i32 [[NTH]], [[WS]]
269 // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[TH_LIMIT]]
270 // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[CHECK_MASTER:.+]]
271 //
272 // CHECK: [[WORKER]]
273 // CHECK: {{call|invoke}} void [[T6]]_worker()
274 // CHECK: br label {{%?}}[[EXIT:.+]]
275 //
276 // CHECK: [[CHECK_MASTER]]
277 // CHECK-DAG: [[CMTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
278 // CHECK-DAG: [[CMNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
279 // CHECK-DAG: [[CMWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
280 // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]],
281 // CHECK: br i1 [[IS_MASTER]], label {{%?}}[[MASTER:.+]], label {{%?}}[[EXIT]]
282 //
283 // CHECK: [[MASTER]]
284 // CHECK-DAG: [[MNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
285 // CHECK-DAG: [[MWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
286 // CHECK: [[MTMP1:%.+]] = sub nuw i32 [[MNTH]], [[MWS]]
287 // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]]
288 // CHECK-64: [[N:%.+]] = load i32, i32* [[REF_N]],
289 // CHECK-32: [[N:%.+]] = load i32, i32* [[LOCAL_N]],
290 // CHECK: [[CMP:%.+]] = icmp sgt i32 [[N]], 1000
291 // CHECK: br i1 [[CMP]], label {{%?}}[[IF_THEN:.+]], label {{%?}}[[IF_ELSE:.+]]
292 //
293 // CHECK: [[IF_THEN]]
294 // CHECK: call void @__kmpc_kernel_prepare_parallel(i8* bitcast (void (i16, i32)* [[PARALLEL_FN4]]_wrapper to i8*))
295 // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
296 // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
297 // CHECK: br label {{%?}}[[IF_END:.+]]
298 //
299 // CHECK: [[IF_ELSE]]
300 // CHECK: call void @__kmpc_serialized_parallel(
301 // CHECK: {{call|invoke}} void [[PARALLEL_FN4]](
302 // CHECK: call void @__kmpc_end_serialized_parallel(
303 // br label [[IF_END]]
304 //
305 // CHECK: [[IF_END]]
306 // CHECK-64-DAG: load i32, i32* [[REF_A]]
307 // CHECK-32-DAG: load i32, i32* [[LOCAL_A]]
308 // CHECK-DAG: load i16, i16* [[REF_AA]]
309 // CHECK-DAG: getelementptr inbounds [10 x i32], [10 x i32]* [[REF_B]], i[[SZ]] 0, i[[SZ]] 2
310 //
311 // CHECK: br label {{%?}}[[TERMINATE:.+]]
312 //
313 // CHECK: [[TERMINATE]]
314 // CHECK: call void @__kmpc_kernel_deinit(
315 // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
316 // CHECK: br label {{%?}}[[EXIT]]
317 //
318 // CHECK: [[EXIT]]
319 // CHECK: ret void
320
321 // CHECK: define internal void [[PARALLEL_FN4]](
322 // CHECK: [[A:%.+]] = alloca i[[SZ:32|64]],
323 // CHECK: store i[[SZ]] 45, i[[SZ]]* %a,
324 // CHECK: call void @__kmpc_barrier(%struct.ident_t* @{{.+}}, i32 %{{.+}})
325 // CHECK: ret void
326
327 // CHECK: declare void @__kmpc_barrier(%struct.ident_t*, i32) #[[#CONVERGENT:]]
328
329 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l58}}_worker()
330 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l58}}(
331 // CHECK-32: [[A_ADDR:%.+]] = alloca i32,
332 // CHECK-64: [[A_ADDR:%.+]] = alloca i64,
333 // CHECK-64: [[CONV:%.+]] = bitcast i64* [[A_ADDR]] to i32*
334 // SEQ: [[IS_SHARED:%.+]] = load i16, i16* [[KERNEL_SHARED]],
335 // SEQ: [[SIZE:%.+]] = load i{{64|32}}, i{{64|32}}* [[KERNEL_SIZE]],
336 // SEQ: call void @__kmpc_get_team_static_memory(i16 0, i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([[MEM_TY]], [[MEM_TY]] addrspace(3)* [[SHARED_GLOBAL_RD]], i32 0, i32 0, i32 0) to i8*), i{{64|32}} [[SIZE]], i16 [[IS_SHARED]], i8** addrspacecast (i8* addrspace(3)* [[KERNEL_PTR]] to i8**))
337 // SEQ: [[KERNEL_RD:%.+]] = load i8*, i8* addrspace(3)* [[KERNEL_PTR]],
338 // SEQ: [[STACK:%.+]] = getelementptr inbounds i8, i8* [[KERNEL_RD]], i{{64|32}} 0
339 // PAR: [[STACK:%.+]] = call i8* @__kmpc_data_sharing_push_stack(i{{32|64}} 4, i16 1)
340 // CHECK: [[BC:%.+]] = bitcast i8* [[STACK]] to %struct._globalized_locals_ty*
341 // CHECK-32: [[A:%.+]] = load i32, i32* [[A_ADDR]],
342 // CHECK-64: [[A:%.+]] = load i32, i32* [[CONV]],
343 // CHECK: [[GLOBAL_A_ADDR:%.+]] = getelementptr inbounds %struct._globalized_locals_ty, %struct._globalized_locals_ty* [[BC]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
344 // CHECK: store i32 [[A]], i32* [[GLOBAL_A_ADDR]],
345 // SEQ: [[IS_SHARED:%.+]] = load i16, i16* [[KERNEL_SHARED]],
346 // SEQ: call void @__kmpc_restore_team_static_memory(i16 0, i16 [[IS_SHARED]])
347 // PAR: call void @__kmpc_data_sharing_pop_stack(i8* [[STACK]])
348
349 // CHECK-LABEL: define internal void @{{.+}}(i32* noalias %{{.+}}, i32* noalias %{{.+}}, i32* nonnull align {{[0-9]+}} dereferenceable{{.*}})
350 // CHECK: [[CC:%.+]] = alloca i32,
351 // CHECK: [[MASK:%.+]] = call i32 @__kmpc_warp_active_thread_mask(){{$}}
352 // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
353 // CHECK: [[NUM_THREADS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
354 // CHECK: store i32 0, i32* [[CC]],
355 // CHECK: br label
356
357 // CHECK: [[CC_VAL:%.+]] = load i32, i32* [[CC]],
358 // CHECK: [[RES:%.+]] = icmp slt i32 [[CC_VAL]], [[NUM_THREADS]]
359 // CHECK: br i1 [[RES]], label
360
361 // CHECK: [[CC_VAL:%.+]] = load i32, i32* [[CC]],
362 // CHECK: [[RES:%.+]] = icmp eq i32 [[TID]], [[CC_VAL]]
363 // CHECK: br i1 [[RES]], label
364
365 // CHECK: call void @__kmpc_critical(
366 // CHECK: load i32, i32*
367 // CHECK: add nsw i32
368 // CHECK: store i32
369 // CHECK: call void @__kmpc_end_critical(
370
371 // CHECK: call void @__kmpc_syncwarp(i32 [[MASK]]){{$}}
372 // CHECK: [[NEW_CC_VAL:%.+]] = add nsw i32 [[CC_VAL]], 1
373 // CHECK: store i32 [[NEW_CC_VAL]], i32* [[CC]],
374 // CHECK: br label
375
376
377 // CHECK: declare i32 @__kmpc_warp_active_thread_mask() #[[#CONVERGENT:]]
378 // CHECK: declare void @__kmpc_syncwarp(i32) #[[#CONVERGENT:]]
379
380 // CHECK: attributes #[[#CONVERGENT:]] = {{.*}} convergent {{.*}}
381
382 #endif
383