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 - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
4 // 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
5 // 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 - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
6 // expected-no-diagnostics
7 #ifndef HEADER
8 #define HEADER
9
10 // CHECK-DAG: [[OMP_NT:@.+]] = common addrspace(3) global i32 0
11 // CHECK-DAG: [[OMP_WID:@.+]] = common addrspace(3) global i64 0
12
13 template<typename tx, typename ty>
14 struct TT{
15 tx X;
16 ty Y;
17 };
18
foo(int n)19 int foo(int n) {
20 int a = 0;
21 short aa = 0;
22 float b[10];
23 float bn[n];
24 double c[5][10];
25 double cn[5][n];
26 TT<long long, char> d;
27
28 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l86}}_worker()
29 // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
30 //
31 // CHECK: [[AWAIT_WORK]]
32 // CHECK: call void @llvm.nvvm.barrier0()
33 // CHECK: [[WORK:%.+]] = load i64, i64 addrspace(3)* [[OMP_WID]],
34 // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i64 [[WORK]], 0
35 // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]]
36 //
37 // CHECK: [[SEL_WORKERS]]
38 // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
39 // CHECK: [[NT:%.+]] = load i32, i32 addrspace(3)* [[OMP_NT]]
40 // CHECK: [[IS_ACTIVE:%.+]] = icmp slt i32 [[TID]], [[NT]]
41 // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]]
42 //
43 // CHECK: [[EXEC_PARALLEL]]
44 // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
45 //
46 // CHECK: [[TERM_PARALLEL]]
47 // CHECK: br label {{%?}}[[BAR_PARALLEL]]
48 //
49 // CHECK: [[BAR_PARALLEL]]
50 // CHECK: call void @llvm.nvvm.barrier0()
51 // CHECK: br label {{%?}}[[AWAIT_WORK]]
52 //
53 // CHECK: [[EXIT]]
54 // CHECK: ret void
55
56 // CHECK: define {{.*}}void [[T1:@__omp_offloading_.+foo.+l86]]()
57 // CHECK: [[NTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
58 // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
59 // CHECK: [[A:%.+]] = sub i32 [[WS]], 1
60 // CHECK: [[B:%.+]] = sub i32 [[NTID]], 1
61 // CHECK: [[MID:%.+]] = and i32 [[B]],
62 // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
63 // CHECK: [[EXCESS:%.+]] = icmp ugt i32 [[TID]], [[MID]]
64 // CHECK: br i1 [[EXCESS]], label {{%?}}[[EXIT:.+]], label {{%?}}[[CHECK_WORKER:.+]]
65 //
66 // CHECK: [[CHECK_WORKER]]
67 // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[MID]]
68 // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[MASTER:.+]]
69 //
70 // CHECK: [[WORKER]]
71 // CHECK: call void [[T1]]_worker()
72 // CHECK: br label {{%?}}[[EXIT]]
73 //
74 // CHECK: [[MASTER]]
75 // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
76 // CHECK: call void @__kmpc_kernel_init(i32 0, i32 [[TID]])
77 // CHECK: br label {{%?}}[[TERM:.+]]
78 //
79 // CHECK: [[TERM]]
80 // CHECK: store i64 0, i64 addrspace(3)* [[OMP_WID]],
81 // CHECK: call void @llvm.nvvm.barrier0()
82 // CHECK: br label {{%?}}[[EXIT]]
83 //
84 // CHECK: [[EXIT]]
85 // CHECK: ret void
86 #pragma omp target
87 {
88 }
89
90 // CHECK-NOT: define {{.*}}void [[T2:@__omp_offloading_.+foo.+]]_worker()
91 #pragma omp target if(0)
92 {
93 }
94
95 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l157}}_worker()
96 // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
97 //
98 // CHECK: [[AWAIT_WORK]]
99 // CHECK: call void @llvm.nvvm.barrier0()
100 // CHECK: [[WORK:%.+]] = load i64, i64 addrspace(3)* [[OMP_WID]],
101 // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i64 [[WORK]], 0
102 // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]]
103 //
104 // CHECK: [[SEL_WORKERS]]
105 // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
106 // CHECK: [[NT:%.+]] = load i32, i32 addrspace(3)* [[OMP_NT]]
107 // CHECK: [[IS_ACTIVE:%.+]] = icmp slt i32 [[TID]], [[NT]]
108 // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]]
109 //
110 // CHECK: [[EXEC_PARALLEL]]
111 // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
112 //
113 // CHECK: [[TERM_PARALLEL]]
114 // CHECK: br label {{%?}}[[BAR_PARALLEL]]
115 //
116 // CHECK: [[BAR_PARALLEL]]
117 // CHECK: call void @llvm.nvvm.barrier0()
118 // CHECK: br label {{%?}}[[AWAIT_WORK]]
119 //
120 // CHECK: [[EXIT]]
121 // CHECK: ret void
122
123 // CHECK: define {{.*}}void [[T3:@__omp_offloading_.+foo.+l157]](i[[SZ:32|64]] [[ARG1:%.+]])
124 // CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]],
125 // CHECK: store i[[SZ]] [[ARG1]], i[[SZ]]* [[AA_ADDR]],
126 // CHECK: [[AA_CADDR:%.+]] = bitcast i[[SZ]]* [[AA_ADDR]] to i16*
127 // CHECK: [[NTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
128 // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
129 // CHECK: [[A:%.+]] = sub i32 [[WS]], 1
130 // CHECK: [[B:%.+]] = sub i32 [[NTID]], 1
131 // CHECK: [[MID:%.+]] = and i32 [[B]],
132 // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
133 // CHECK: [[EXCESS:%.+]] = icmp ugt i32 [[TID]], [[MID]]
134 // CHECK: br i1 [[EXCESS]], label {{%?}}[[EXIT:.+]], label {{%?}}[[CHECK_WORKER:.+]]
135 //
136 // CHECK: [[CHECK_WORKER]]
137 // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[MID]]
138 // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[MASTER:.+]]
139 //
140 // CHECK: [[WORKER]]
141 // CHECK: call void [[T3]]_worker()
142 // CHECK: br label {{%?}}[[EXIT]]
143 //
144 // CHECK: [[MASTER]]
145 // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
146 // CHECK: call void @__kmpc_kernel_init(i32 0, i32 [[TID]])
147 // CHECK: load i16, i16* [[AA_CADDR]],
148 // CHECK: br label {{%?}}[[TERM:.+]]
149 //
150 // CHECK: [[TERM]]
151 // CHECK: store i64 0, i64 addrspace(3)* [[OMP_WID]],
152 // CHECK: call void @llvm.nvvm.barrier0()
153 // CHECK: br label {{%?}}[[EXIT]]
154 //
155 // CHECK: [[EXIT]]
156 // CHECK: ret void
157 #pragma omp target if(1)
158 {
159 aa += 1;
160 }
161
162 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l260}}_worker()
163 // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
164 //
165 // CHECK: [[AWAIT_WORK]]
166 // CHECK: call void @llvm.nvvm.barrier0()
167 // CHECK: [[WORK:%.+]] = load i64, i64 addrspace(3)* [[OMP_WID]],
168 // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i64 [[WORK]], 0
169 // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]]
170 //
171 // CHECK: [[SEL_WORKERS]]
172 // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
173 // CHECK: [[NT:%.+]] = load i32, i32 addrspace(3)* [[OMP_NT]]
174 // CHECK: [[IS_ACTIVE:%.+]] = icmp slt i32 [[TID]], [[NT]]
175 // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]]
176 //
177 // CHECK: [[EXEC_PARALLEL]]
178 // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
179 //
180 // CHECK: [[TERM_PARALLEL]]
181 // CHECK: br label {{%?}}[[BAR_PARALLEL]]
182 //
183 // CHECK: [[BAR_PARALLEL]]
184 // CHECK: call void @llvm.nvvm.barrier0()
185 // CHECK: br label {{%?}}[[AWAIT_WORK]]
186 //
187 // CHECK: [[EXIT]]
188 // CHECK: ret void
189
190 // CHECK: define {{.*}}void [[T4:@__omp_offloading_.+foo.+l260]](i[[SZ]]
191 // Create local storage for each capture.
192 // CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]
193 // CHECK: [[LOCAL_B:%.+]] = alloca [10 x float]*
194 // CHECK: [[LOCAL_VLA1:%.+]] = alloca i[[SZ]]
195 // CHECK: [[LOCAL_BN:%.+]] = alloca float*
196 // CHECK: [[LOCAL_C:%.+]] = alloca [5 x [10 x double]]*
197 // CHECK: [[LOCAL_VLA2:%.+]] = alloca i[[SZ]]
198 // CHECK: [[LOCAL_VLA3:%.+]] = alloca i[[SZ]]
199 // CHECK: [[LOCAL_CN:%.+]] = alloca double*
200 // CHECK: [[LOCAL_D:%.+]] = alloca [[TT:%.+]]*
201 // CHECK-DAG: store i[[SZ]] [[ARG_A:%.+]], i[[SZ]]* [[LOCAL_A]]
202 // CHECK-DAG: store [10 x float]* [[ARG_B:%.+]], [10 x float]** [[LOCAL_B]]
203 // CHECK-DAG: store i[[SZ]] [[ARG_VLA1:%.+]], i[[SZ]]* [[LOCAL_VLA1]]
204 // CHECK-DAG: store float* [[ARG_BN:%.+]], float** [[LOCAL_BN]]
205 // CHECK-DAG: store [5 x [10 x double]]* [[ARG_C:%.+]], [5 x [10 x double]]** [[LOCAL_C]]
206 // CHECK-DAG: store i[[SZ]] [[ARG_VLA2:%.+]], i[[SZ]]* [[LOCAL_VLA2]]
207 // CHECK-DAG: store i[[SZ]] [[ARG_VLA3:%.+]], i[[SZ]]* [[LOCAL_VLA3]]
208 // CHECK-DAG: store double* [[ARG_CN:%.+]], double** [[LOCAL_CN]]
209 // CHECK-DAG: store [[TT]]* [[ARG_D:%.+]], [[TT]]** [[LOCAL_D]]
210 //
211 // CHECK-64-DAG: [[REF_A:%.+]] = bitcast i64* [[LOCAL_A]] to i32*
212 // CHECK-DAG: [[REF_B:%.+]] = load [10 x float]*, [10 x float]** [[LOCAL_B]],
213 // CHECK-DAG: [[VAL_VLA1:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA1]],
214 // CHECK-DAG: [[REF_BN:%.+]] = load float*, float** [[LOCAL_BN]],
215 // CHECK-DAG: [[REF_C:%.+]] = load [5 x [10 x double]]*, [5 x [10 x double]]** [[LOCAL_C]],
216 // CHECK-DAG: [[VAL_VLA2:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA2]],
217 // CHECK-DAG: [[VAL_VLA3:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA3]],
218 // CHECK-DAG: [[REF_CN:%.+]] = load double*, double** [[LOCAL_CN]],
219 // CHECK-DAG: [[REF_D:%.+]] = load [[TT]]*, [[TT]]** [[LOCAL_D]],
220 //
221 // CHECK: [[NTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
222 // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
223 // CHECK: [[A:%.+]] = sub i32 [[WS]], 1
224 // CHECK: [[B:%.+]] = sub i32 [[NTID]], 1
225 // CHECK: [[MID:%.+]] = and i32 [[B]],
226 // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
227 // CHECK: [[EXCESS:%.+]] = icmp ugt i32 [[TID]], [[MID]]
228 // CHECK: br i1 [[EXCESS]], label {{%?}}[[EXIT:.+]], label {{%?}}[[CHECK_WORKER:.+]]
229 //
230 // CHECK: [[CHECK_WORKER]]
231 // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[MID]]
232 // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[MASTER:.+]]
233 //
234 // CHECK: [[WORKER]]
235 // CHECK: call void [[T4]]_worker()
236 // CHECK: br label {{%?}}[[EXIT]]
237 //
238 // CHECK: [[MASTER]]
239 // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
240 // CHECK: call void @__kmpc_kernel_init(i32 0, i32 [[TID]])
241 //
242 // Use captures.
243 // CHECK-64-DAG: load i32, i32* [[REF_A]]
244 // CHECK-32-DAG: load i32, i32* [[LOCAL_A]]
245 // CHECK-DAG: getelementptr inbounds [10 x float], [10 x float]* [[REF_B]], i[[SZ]] 0, i[[SZ]] 2
246 // CHECK-DAG: getelementptr inbounds float, float* [[REF_BN]], i[[SZ]] 3
247 // CHECK-DAG: getelementptr inbounds [5 x [10 x double]], [5 x [10 x double]]* [[REF_C]], i[[SZ]] 0, i[[SZ]] 1
248 // CHECK-DAG: getelementptr inbounds double, double* [[REF_CN]], i[[SZ]] %{{.+}}
249 // CHECK-DAG: getelementptr inbounds [[TT]], [[TT]]* [[REF_D]], i32 0, i32 0
250 //
251 // CHECK: br label {{%?}}[[TERM:.+]]
252 //
253 // CHECK: [[TERM]]
254 // CHECK: store i64 0, i64 addrspace(3)* [[OMP_WID]],
255 // CHECK: call void @llvm.nvvm.barrier0()
256 // CHECK: br label {{%?}}[[EXIT]]
257 //
258 // CHECK: [[EXIT]]
259 // CHECK: ret void
260 #pragma omp target if(n>20)
261 {
262 a += 1;
263 b[2] += 1.0;
264 bn[3] += 1.0;
265 c[1][2] += 1.0;
266 cn[1][3] += 1.0;
267 d.X += 1;
268 d.Y += 1;
269 }
270
271 return a;
272 }
273
274 template<typename tx>
ftemplate(int n)275 tx ftemplate(int n) {
276 tx a = 0;
277 short aa = 0;
278 tx b[10];
279
280 #pragma omp target if(n>40)
281 {
282 a += 1;
283 aa += 1;
284 b[2] += 1;
285 }
286
287 return a;
288 }
289
290 static
fstatic(int n)291 int fstatic(int n) {
292 int a = 0;
293 short aa = 0;
294 char aaa = 0;
295 int b[10];
296
297 #pragma omp target if(n>50)
298 {
299 a += 1;
300 aa += 1;
301 aaa += 1;
302 b[2] += 1;
303 }
304
305 return a;
306 }
307
308 struct S1 {
309 double a;
310
r1S1311 int r1(int n){
312 int b = n+1;
313 short int c[2][n];
314
315 #pragma omp target if(n>60)
316 {
317 this->a = (double)b + 1.5;
318 c[1][1] = ++a;
319 }
320
321 return c[1][1] + (int)b;
322 }
323 };
324
bar(int n)325 int bar(int n){
326 int a = 0;
327
328 a += foo(n);
329
330 S1 S;
331 a += S.r1(n);
332
333 a += fstatic(n);
334
335 a += ftemplate<int>(n);
336
337 return a;
338 }
339
340 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+static.+l297}}_worker()
341 // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
342 //
343 // CHECK: [[AWAIT_WORK]]
344 // CHECK: call void @llvm.nvvm.barrier0()
345 // CHECK: [[WORK:%.+]] = load i64, i64 addrspace(3)* [[OMP_WID]],
346 // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i64 [[WORK]], 0
347 // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]]
348 //
349 // CHECK: [[SEL_WORKERS]]
350 // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
351 // CHECK: [[NT:%.+]] = load i32, i32 addrspace(3)* [[OMP_NT]]
352 // CHECK: [[IS_ACTIVE:%.+]] = icmp slt i32 [[TID]], [[NT]]
353 // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]]
354 //
355 // CHECK: [[EXEC_PARALLEL]]
356 // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
357 //
358 // CHECK: [[TERM_PARALLEL]]
359 // CHECK: br label {{%?}}[[BAR_PARALLEL]]
360 //
361 // CHECK: [[BAR_PARALLEL]]
362 // CHECK: call void @llvm.nvvm.barrier0()
363 // CHECK: br label {{%?}}[[AWAIT_WORK]]
364 //
365 // CHECK: [[EXIT]]
366 // CHECK: ret void
367
368 // CHECK: define {{.*}}void [[T5:@__omp_offloading_.+static.+l297]](i[[SZ]]
369 // Create local storage for each capture.
370 // CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]
371 // CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]]
372 // CHECK: [[LOCAL_AAA:%.+]] = alloca i[[SZ]]
373 // CHECK: [[LOCAL_B:%.+]] = alloca [10 x i32]*
374 // CHECK-DAG: store i[[SZ]] [[ARG_A:%.+]], i[[SZ]]* [[LOCAL_A]]
375 // CHECK-DAG: store i[[SZ]] [[ARG_AA:%.+]], i[[SZ]]* [[LOCAL_AA]]
376 // CHECK-DAG: store i[[SZ]] [[ARG_AAA:%.+]], i[[SZ]]* [[LOCAL_AAA]]
377 // CHECK-DAG: store [10 x i32]* [[ARG_B:%.+]], [10 x i32]** [[LOCAL_B]]
378 // Store captures in the context.
379 // CHECK-64-DAG: [[REF_A:%.+]] = bitcast i[[SZ]]* [[LOCAL_A]] to i32*
380 // CHECK-DAG: [[REF_AA:%.+]] = bitcast i[[SZ]]* [[LOCAL_AA]] to i16*
381 // CHECK-DAG: [[REF_AAA:%.+]] = bitcast i[[SZ]]* [[LOCAL_AAA]] to i8*
382 // CHECK-DAG: [[REF_B:%.+]] = load [10 x i32]*, [10 x i32]** [[LOCAL_B]],
383 //
384 // CHECK: [[NTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
385 // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
386 // CHECK: [[A:%.+]] = sub i32 [[WS]], 1
387 // CHECK: [[B:%.+]] = sub i32 [[NTID]], 1
388 // CHECK: [[MID:%.+]] = and i32 [[B]],
389 // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
390 // CHECK: [[EXCESS:%.+]] = icmp ugt i32 [[TID]], [[MID]]
391 // CHECK: br i1 [[EXCESS]], label {{%?}}[[EXIT:.+]], label {{%?}}[[CHECK_WORKER:.+]]
392 //
393 // CHECK: [[CHECK_WORKER]]
394 // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[MID]]
395 // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[MASTER:.+]]
396 //
397 // CHECK: [[WORKER]]
398 // CHECK: call void [[T5]]_worker()
399 // CHECK: br label {{%?}}[[EXIT]]
400 //
401 // CHECK: [[MASTER]]
402 // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
403 // CHECK: call void @__kmpc_kernel_init(i32 0, i32 [[TID]])
404 //
405 // CHECK-64-DAG: load i32, i32* [[REF_A]]
406 // CHECK-32-DAG: load i32, i32* [[LOCAL_A]]
407 // CHECK-DAG: load i16, i16* [[REF_AA]]
408 // CHECK-DAG: getelementptr inbounds [10 x i32], [10 x i32]* [[REF_B]], i[[SZ]] 0, i[[SZ]] 2
409 //
410 // CHECK: br label {{%?}}[[TERM:.+]]
411 //
412 // CHECK: [[TERM]]
413 // CHECK: store i64 0, i64 addrspace(3)* [[OMP_WID]],
414 // CHECK: call void @llvm.nvvm.barrier0()
415 // CHECK: br label {{%?}}[[EXIT]]
416 //
417 // CHECK: [[EXIT]]
418 // CHECK: ret void
419
420
421
422 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+S1.+l315}}_worker()
423 // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
424 //
425 // CHECK: [[AWAIT_WORK]]
426 // CHECK: call void @llvm.nvvm.barrier0()
427 // CHECK: [[WORK:%.+]] = load i64, i64 addrspace(3)* [[OMP_WID]],
428 // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i64 [[WORK]], 0
429 // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]]
430 //
431 // CHECK: [[SEL_WORKERS]]
432 // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
433 // CHECK: [[NT:%.+]] = load i32, i32 addrspace(3)* [[OMP_NT]]
434 // CHECK: [[IS_ACTIVE:%.+]] = icmp slt i32 [[TID]], [[NT]]
435 // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]]
436 //
437 // CHECK: [[EXEC_PARALLEL]]
438 // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
439 //
440 // CHECK: [[TERM_PARALLEL]]
441 // CHECK: br label {{%?}}[[BAR_PARALLEL]]
442 //
443 // CHECK: [[BAR_PARALLEL]]
444 // CHECK: call void @llvm.nvvm.barrier0()
445 // CHECK: br label {{%?}}[[AWAIT_WORK]]
446 //
447 // CHECK: [[EXIT]]
448 // CHECK: ret void
449
450 // CHECK: define {{.*}}void [[T6:@__omp_offloading_.+S1.+l315]](
451 // Create local storage for each capture.
452 // CHECK: [[LOCAL_THIS:%.+]] = alloca [[S1:%struct.*]]*
453 // CHECK: [[LOCAL_B:%.+]] = alloca i[[SZ]]
454 // CHECK: [[LOCAL_VLA1:%.+]] = alloca i[[SZ]]
455 // CHECK: [[LOCAL_VLA2:%.+]] = alloca i[[SZ]]
456 // CHECK: [[LOCAL_C:%.+]] = alloca i16*
457 // CHECK-DAG: store [[S1]]* [[ARG_THIS:%.+]], [[S1]]** [[LOCAL_THIS]]
458 // CHECK-DAG: store i[[SZ]] [[ARG_B:%.+]], i[[SZ]]* [[LOCAL_B]]
459 // CHECK-DAG: store i[[SZ]] [[ARG_VLA1:%.+]], i[[SZ]]* [[LOCAL_VLA1]]
460 // CHECK-DAG: store i[[SZ]] [[ARG_VLA2:%.+]], i[[SZ]]* [[LOCAL_VLA2]]
461 // CHECK-DAG: store i16* [[ARG_C:%.+]], i16** [[LOCAL_C]]
462 // Store captures in the context.
463 // CHECK-DAG: [[REF_THIS:%.+]] = load [[S1]]*, [[S1]]** [[LOCAL_THIS]],
464 // CHECK-64-DAG:[[REF_B:%.+]] = bitcast i[[SZ]]* [[LOCAL_B]] to i32*
465 // CHECK-DAG: [[VAL_VLA1:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA1]],
466 // CHECK-DAG: [[VAL_VLA2:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA2]],
467 // CHECK-DAG: [[REF_C:%.+]] = load i16*, i16** [[LOCAL_C]],
468 // CHECK: [[NTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
469 // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
470 // CHECK: [[A:%.+]] = sub i32 [[WS]], 1
471 // CHECK: [[B:%.+]] = sub i32 [[NTID]], 1
472 // CHECK: [[MID:%.+]] = and i32 [[B]],
473 // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
474 // CHECK: [[EXCESS:%.+]] = icmp ugt i32 [[TID]], [[MID]]
475 // CHECK: br i1 [[EXCESS]], label {{%?}}[[EXIT:.+]], label {{%?}}[[CHECK_WORKER:.+]]
476 //
477 // CHECK: [[CHECK_WORKER]]
478 // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[MID]]
479 // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[MASTER:.+]]
480 //
481 // CHECK: [[WORKER]]
482 // CHECK: call void [[T6]]_worker()
483 // CHECK: br label {{%?}}[[EXIT]]
484 //
485 // CHECK: [[MASTER]]
486 // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
487 // CHECK: call void @__kmpc_kernel_init(i32 0, i32 [[TID]])
488 // Use captures.
489 // CHECK-DAG: getelementptr inbounds [[S1]], [[S1]]* [[REF_THIS]], i32 0, i32 0
490 // CHECK-64-DAG:load i32, i32* [[REF_B]]
491 // CHECK-32-DAG:load i32, i32* [[LOCAL_B]]
492 // CHECK-DAG: getelementptr inbounds i16, i16* [[REF_C]], i[[SZ]] %{{.+}}
493 // CHECK: br label {{%?}}[[TERM:.+]]
494 //
495 // CHECK: [[TERM]]
496 // CHECK: store i64 0, i64 addrspace(3)* [[OMP_WID]],
497 // CHECK: call void @llvm.nvvm.barrier0()
498 // CHECK: br label {{%?}}[[EXIT]]
499 //
500 // CHECK: [[EXIT]]
501 // CHECK: ret void
502
503
504
505 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l280}}_worker()
506 // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
507 //
508 // CHECK: [[AWAIT_WORK]]
509 // CHECK: call void @llvm.nvvm.barrier0()
510 // CHECK: [[WORK:%.+]] = load i64, i64 addrspace(3)* [[OMP_WID]],
511 // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i64 [[WORK]], 0
512 // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]]
513 //
514 // CHECK: [[SEL_WORKERS]]
515 // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
516 // CHECK: [[NT:%.+]] = load i32, i32 addrspace(3)* [[OMP_NT]]
517 // CHECK: [[IS_ACTIVE:%.+]] = icmp slt i32 [[TID]], [[NT]]
518 // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]]
519 //
520 // CHECK: [[EXEC_PARALLEL]]
521 // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
522 //
523 // CHECK: [[TERM_PARALLEL]]
524 // CHECK: br label {{%?}}[[BAR_PARALLEL]]
525 //
526 // CHECK: [[BAR_PARALLEL]]
527 // CHECK: call void @llvm.nvvm.barrier0()
528 // CHECK: br label {{%?}}[[AWAIT_WORK]]
529 //
530 // CHECK: [[EXIT]]
531 // CHECK: ret void
532
533 // CHECK: define {{.*}}void [[T7:@__omp_offloading_.+template.+l280]](i[[SZ]]
534 // Create local storage for each capture.
535 // CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]
536 // CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]]
537 // CHECK: [[LOCAL_B:%.+]] = alloca [10 x i32]*
538 // CHECK-DAG: store i[[SZ]] [[ARG_A:%.+]], i[[SZ]]* [[LOCAL_A]]
539 // CHECK-DAG: store i[[SZ]] [[ARG_AA:%.+]], i[[SZ]]* [[LOCAL_AA]]
540 // CHECK-DAG: store [10 x i32]* [[ARG_B:%.+]], [10 x i32]** [[LOCAL_B]]
541 // Store captures in the context.
542 // CHECK-64-DAG:[[REF_A:%.+]] = bitcast i[[SZ]]* [[LOCAL_A]] to i32*
543 // CHECK-DAG: [[REF_AA:%.+]] = bitcast i[[SZ]]* [[LOCAL_AA]] to i16*
544 // CHECK-DAG: [[REF_B:%.+]] = load [10 x i32]*, [10 x i32]** [[LOCAL_B]],
545 //
546 // CHECK: [[NTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
547 // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
548 // CHECK: [[A:%.+]] = sub i32 [[WS]], 1
549 // CHECK: [[B:%.+]] = sub i32 [[NTID]], 1
550 // CHECK: [[MID:%.+]] = and i32 [[B]],
551 // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
552 // CHECK: [[EXCESS:%.+]] = icmp ugt i32 [[TID]], [[MID]]
553 // CHECK: br i1 [[EXCESS]], label {{%?}}[[EXIT:.+]], label {{%?}}[[CHECK_WORKER:.+]]
554 //
555 // CHECK: [[CHECK_WORKER]]
556 // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[MID]]
557 // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[MASTER:.+]]
558 //
559 // CHECK: [[WORKER]]
560 // CHECK: call void [[T7]]_worker()
561 // CHECK: br label {{%?}}[[EXIT]]
562 //
563 // CHECK: [[MASTER]]
564 // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
565 // CHECK: call void @__kmpc_kernel_init(i32 0, i32 [[TID]])
566 //
567 // CHECK-64-DAG: load i32, i32* [[REF_A]]
568 // CHECK-32-DAG: load i32, i32* [[LOCAL_A]]
569 // CHECK-DAG: load i16, i16* [[REF_AA]]
570 // CHECK-DAG: getelementptr inbounds [10 x i32], [10 x i32]* [[REF_B]], i[[SZ]] 0, i[[SZ]] 2
571 //
572 // CHECK: br label {{%?}}[[TERM:.+]]
573 //
574 // CHECK: [[TERM]]
575 // CHECK: store i64 0, i64 addrspace(3)* [[OMP_WID]],
576 // CHECK: call void @llvm.nvvm.barrier0()
577 // CHECK: br label {{%?}}[[EXIT]]
578 //
579 // CHECK: [[EXIT]]
580 // CHECK: ret void
581 #endif
582