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 // 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 - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
7 // expected-no-diagnostics
8 #ifndef HEADER
9 #define HEADER
10 
11 // Check that the execution mode of all 6 target regions is set to Generic Mode.
12 // CHECK-DAG: {{@__omp_offloading_.+l102}}_exec_mode = weak constant i8 1
13 // CHECK-DAG: {{@__omp_offloading_.+l179}}_exec_mode = weak constant i8 1
14 // CHECK-DAG: {{@__omp_offloading_.+l289}}_exec_mode = weak constant i8 1
15 // CHECK-DAG: {{@__omp_offloading_.+l326}}_exec_mode = weak constant i8 1
16 // CHECK-DAG: {{@__omp_offloading_.+l344}}_exec_mode = weak constant i8 1
17 // CHECK-DAG: {{@__omp_offloading_.+l309}}_exec_mode = weak constant i8 1
18 
19 __thread int id;
20 
21 int baz(int f);
22 
23 template<typename tx, typename ty>
24 struct TT{
25   tx X;
26   ty Y;
27 };
28 
29 int foo(int n) {
30   int a = 0;
31   short aa = 0;
32   float b[10];
33   float bn[n];
34   double c[5][10];
35   double cn[5][n];
36   TT<long long, char> d;
37 
38   // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l102}}_worker()
39   // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
40   // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
41   // CHECK: store i8* null, i8** [[OMP_WORK_FN]],
42   // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]],
43   // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
44   //
45   // CHECK: [[AWAIT_WORK]]
46   // CHECK: call void @llvm.nvvm.barrier0()
47   // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]],
48   // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null
49   // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]]
50   //
51   // CHECK: [[SEL_WORKERS]]
52   // CHECK: [[ST:%.+]] = load i8, i8* [[OMP_EXEC_STATUS]],
53   // CHECK: [[IS_ACTIVE:%.+]] = icmp ne i8 [[ST]], 0
54   // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]]
55   //
56   // CHECK: [[EXEC_PARALLEL]]
57   // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
58   //
59   // CHECK: [[TERM_PARALLEL]]
60   // CHECK: br label {{%?}}[[BAR_PARALLEL]]
61   //
62   // CHECK: [[BAR_PARALLEL]]
63   // CHECK: call void @llvm.nvvm.barrier0()
64   // CHECK: br label {{%?}}[[AWAIT_WORK]]
65   //
66   // CHECK: [[EXIT]]
67   // CHECK: ret void
68 
69   // CHECK: define {{.*}}void [[T1:@__omp_offloading_.+foo.+l102]]()
70   // CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
71   // CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
72   // CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
73   // CHECK-DAG: [[TH_LIMIT:%.+]] = sub nuw i32 [[NTH]], [[WS]]
74   // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[TH_LIMIT]]
75   // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[CHECK_MASTER:.+]]
76   //
77   // CHECK: [[WORKER]]
78   // CHECK: {{call|invoke}} void [[T1]]_worker()
79   // CHECK: br label {{%?}}[[EXIT:.+]]
80   //
81   // CHECK: [[CHECK_MASTER]]
82   // CHECK-DAG: [[CMTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
83   // CHECK-DAG: [[CMNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
84   // CHECK-DAG: [[CMWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
85   // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]],
86   // CHECK: br i1 [[IS_MASTER]], label {{%?}}[[MASTER:.+]], label {{%?}}[[EXIT]]
87   //
88   // CHECK: [[MASTER]]
89   // CHECK-DAG: [[MNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
90   // CHECK-DAG: [[MWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
91   // CHECK: [[MTMP1:%.+]] = sub nuw i32 [[MNTH]], [[MWS]]
92   // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]]
93   // CHECK: br label {{%?}}[[TERMINATE:.+]]
94   //
95   // CHECK: [[TERMINATE]]
96   // CHECK: call void @__kmpc_kernel_deinit(
97   // CHECK: call void @llvm.nvvm.barrier0()
98   // CHECK: br label {{%?}}[[EXIT]]
99   //
100   // CHECK: [[EXIT]]
101   // CHECK: ret void
102   #pragma omp target
103   {
104   }
105 
106   // CHECK-NOT: define {{.*}}void [[T2:@__omp_offloading_.+foo.+]]_worker()
107   #pragma omp target if(0)
108   {
109   }
110 
111   // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l179}}_worker()
112   // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
113   // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
114   // CHECK: store i8* null, i8** [[OMP_WORK_FN]],
115   // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]],
116   // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
117   //
118   // CHECK: [[AWAIT_WORK]]
119   // CHECK: call void @llvm.nvvm.barrier0()
120   // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]],
121   // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null
122   // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]]
123   //
124   // CHECK: [[SEL_WORKERS]]
125   // CHECK: [[ST:%.+]] = load i8, i8* [[OMP_EXEC_STATUS]],
126   // CHECK: [[IS_ACTIVE:%.+]] = icmp ne i8 [[ST]], 0
127   // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]]
128   //
129   // CHECK: [[EXEC_PARALLEL]]
130   // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
131   //
132   // CHECK: [[TERM_PARALLEL]]
133   // CHECK: br label {{%?}}[[BAR_PARALLEL]]
134   //
135   // CHECK: [[BAR_PARALLEL]]
136   // CHECK: call void @llvm.nvvm.barrier0()
137   // CHECK: br label {{%?}}[[AWAIT_WORK]]
138   //
139   // CHECK: [[EXIT]]
140   // CHECK: ret void
141 
142   // CHECK: define {{.*}}void [[T2:@__omp_offloading_.+foo.+l179]](i[[SZ:32|64]] [[ARG1:%[a-zA-Z_]+]], i[[SZ:32|64]] [[ID:%[a-zA-Z_]+]])
143   // CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]],
144   // CHECK: store i[[SZ]] [[ARG1]], i[[SZ]]* [[AA_ADDR]],
145   // CHECK: [[AA_CADDR:%.+]] = bitcast i[[SZ]]* [[AA_ADDR]] to i16*
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 [[T2]]_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: load i16, i16* [[AA_CADDR]],
170   // CHECK: br label {{%?}}[[TERMINATE:.+]]
171   //
172   // CHECK: [[TERMINATE]]
173   // CHECK: call void @__kmpc_kernel_deinit(
174   // CHECK: call void @llvm.nvvm.barrier0()
175   // CHECK: br label {{%?}}[[EXIT]]
176   //
177   // CHECK: [[EXIT]]
178   // CHECK: ret void
179   #pragma omp target if(1)
180   {
181     aa += 1;
182     id = aa;
183   }
184 
185   // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l289}}_worker()
186   // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
187   // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
188   // CHECK: store i8* null, i8** [[OMP_WORK_FN]],
189   // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]],
190   // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
191   //
192   // CHECK: [[AWAIT_WORK]]
193   // CHECK: call void @llvm.nvvm.barrier0()
194   // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]],
195   // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null
196   // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]]
197   //
198   // CHECK: [[SEL_WORKERS]]
199   // CHECK: [[ST:%.+]] = load i8, i8* [[OMP_EXEC_STATUS]],
200   // CHECK: [[IS_ACTIVE:%.+]] = icmp ne i8 [[ST]], 0
201   // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]]
202   //
203   // CHECK: [[EXEC_PARALLEL]]
204   // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
205   //
206   // CHECK: [[TERM_PARALLEL]]
207   // CHECK: br label {{%?}}[[BAR_PARALLEL]]
208   //
209   // CHECK: [[BAR_PARALLEL]]
210   // CHECK: call void @llvm.nvvm.barrier0()
211   // CHECK: br label {{%?}}[[AWAIT_WORK]]
212   //
213   // CHECK: [[EXIT]]
214   // CHECK: ret void
215 
216   // CHECK: define {{.*}}void [[T3:@__omp_offloading_.+foo.+l289]](i[[SZ]]
217   // Create local storage for each capture.
218   // CHECK:    [[LOCAL_A:%.+]] = alloca i[[SZ]]
219   // CHECK:    [[LOCAL_B:%.+]] = alloca [10 x float]*
220   // CHECK:    [[LOCAL_VLA1:%.+]] = alloca i[[SZ]]
221   // CHECK:    [[LOCAL_BN:%.+]] = alloca float*
222   // CHECK:    [[LOCAL_C:%.+]] = alloca [5 x [10 x double]]*
223   // CHECK:    [[LOCAL_VLA2:%.+]] = alloca i[[SZ]]
224   // CHECK:    [[LOCAL_VLA3:%.+]] = alloca i[[SZ]]
225   // CHECK:    [[LOCAL_CN:%.+]] = alloca double*
226   // CHECK:    [[LOCAL_D:%.+]] = alloca [[TT:%.+]]*
227   // CHECK-DAG: store i[[SZ]] [[ARG_A:%.+]], i[[SZ]]* [[LOCAL_A]]
228   // CHECK-DAG: store [10 x float]* [[ARG_B:%.+]], [10 x float]** [[LOCAL_B]]
229   // CHECK-DAG: store i[[SZ]] [[ARG_VLA1:%.+]], i[[SZ]]* [[LOCAL_VLA1]]
230   // CHECK-DAG: store float* [[ARG_BN:%.+]], float** [[LOCAL_BN]]
231   // CHECK-DAG: store [5 x [10 x double]]* [[ARG_C:%.+]], [5 x [10 x double]]** [[LOCAL_C]]
232   // CHECK-DAG: store i[[SZ]] [[ARG_VLA2:%.+]], i[[SZ]]* [[LOCAL_VLA2]]
233   // CHECK-DAG: store i[[SZ]] [[ARG_VLA3:%.+]], i[[SZ]]* [[LOCAL_VLA3]]
234   // CHECK-DAG: store double* [[ARG_CN:%.+]], double** [[LOCAL_CN]]
235   // CHECK-DAG: store [[TT]]* [[ARG_D:%.+]], [[TT]]** [[LOCAL_D]]
236   //
237   // CHECK-64-DAG: [[REF_A:%.+]] = bitcast i64* [[LOCAL_A]] to i32*
238   // CHECK-DAG:    [[REF_B:%.+]] = load [10 x float]*, [10 x float]** [[LOCAL_B]],
239   // CHECK-DAG:    [[VAL_VLA1:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA1]],
240   // CHECK-DAG:    [[REF_BN:%.+]] = load float*, float** [[LOCAL_BN]],
241   // CHECK-DAG:    [[REF_C:%.+]] = load [5 x [10 x double]]*, [5 x [10 x double]]** [[LOCAL_C]],
242   // CHECK-DAG:    [[VAL_VLA2:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA2]],
243   // CHECK-DAG:    [[VAL_VLA3:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA3]],
244   // CHECK-DAG:    [[REF_CN:%.+]] = load double*, double** [[LOCAL_CN]],
245   // CHECK-DAG:    [[REF_D:%.+]] = load [[TT]]*, [[TT]]** [[LOCAL_D]],
246   //
247   // CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
248   // CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
249   // CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
250   // CHECK-DAG: [[TH_LIMIT:%.+]] = sub nuw i32 [[NTH]], [[WS]]
251   // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[TH_LIMIT]]
252   // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[CHECK_MASTER:.+]]
253   //
254   // CHECK: [[WORKER]]
255   // CHECK: {{call|invoke}} void [[T3]]_worker()
256   // CHECK: br label {{%?}}[[EXIT:.+]]
257   //
258   // CHECK: [[CHECK_MASTER]]
259   // CHECK-DAG: [[CMTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
260   // CHECK-DAG: [[CMNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
261   // CHECK-DAG: [[CMWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
262   // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]],
263   // CHECK: br i1 [[IS_MASTER]], label {{%?}}[[MASTER:.+]], label {{%?}}[[EXIT]]
264   //
265   // CHECK: [[MASTER]]
266   // CHECK-DAG: [[MNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
267   // CHECK-DAG: [[MWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
268   // CHECK: [[MTMP1:%.+]] = sub nuw i32 [[MNTH]], [[MWS]]
269   // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]]
270   //
271   // Use captures.
272   // CHECK-64-DAG:  load i32, i32* [[REF_A]]
273   // CHECK-32-DAG:  load i32, i32* [[LOCAL_A]]
274   // CHECK-DAG:  getelementptr inbounds [10 x float], [10 x float]* [[REF_B]], i[[SZ]] 0, i[[SZ]] 2
275   // CHECK-DAG:  getelementptr inbounds float, float* [[REF_BN]], i[[SZ]] 3
276   // CHECK-DAG:  getelementptr inbounds [5 x [10 x double]], [5 x [10 x double]]* [[REF_C]], i[[SZ]] 0, i[[SZ]] 1
277   // CHECK-DAG:  getelementptr inbounds double, double* [[REF_CN]], i[[SZ]] %{{.+}}
278   // CHECK-DAG:     getelementptr inbounds [[TT]], [[TT]]* [[REF_D]], i32 0, i32 0
279   //
280   // CHECK: br label {{%?}}[[TERMINATE:.+]]
281   //
282   // CHECK: [[TERMINATE]]
283   // CHECK: call void @__kmpc_kernel_deinit(
284   // CHECK: call void @llvm.nvvm.barrier0()
285   // CHECK: br label {{%?}}[[EXIT]]
286   //
287   // CHECK: [[EXIT]]
288   // CHECK: ret void
289   #pragma omp target if(n>20)
290   {
291     a += 1;
292     b[2] += 1.0;
293     bn[3] += 1.0;
294     c[1][2] += 1.0;
295     cn[1][3] += 1.0;
296     d.X += 1;
297     d.Y += 1;
298   }
299 
300   return a;
301 }
302 
303 template<typename tx>
304 tx ftemplate(int n) {
305   tx a = 0;
306   short aa = 0;
307   tx b[10];
308 
309   #pragma omp target if(n>40)
310   {
311     a += 1;
312     aa += 1;
313     b[2] += 1;
314   }
315 
316   return a;
317 }
318 
319 static
320 int fstatic(int n) {
321   int a = 0;
322   short aa = 0;
323   char aaa = 0;
324   int b[10];
325 
326   #pragma omp target if(n>50)
327   {
328     a += 1;
329     aa += 1;
330     aaa += 1;
331     b[2] += 1;
332   }
333 
334   return a;
335 }
336 
337 struct S1 {
338   double a;
339 
340   int r1(int n){
341     int b = n+1;
342     short int c[2][n];
343 
344     #pragma omp target if(n>60)
345     {
346       this->a = (double)b + 1.5;
347       c[1][1] = ++a;
348       baz(a);
349     }
350 
351     return c[1][1] + (int)b;
352   }
353 };
354 
355 int bar(int n){
356   int a = 0;
357 
358   a += foo(n);
359 
360   S1 S;
361   a += S.r1(n);
362 
363   a += fstatic(n);
364 
365   a += ftemplate<int>(n);
366 
367   return a;
368 }
369 
370 int baz(int f) {
371 #pragma omp parallel
372   f = 2;
373   return f;
374 }
375 
376   // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+static.+326}}_worker()
377   // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
378   // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
379   // CHECK: store i8* null, i8** [[OMP_WORK_FN]],
380   // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]],
381   // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
382   //
383   // CHECK: [[AWAIT_WORK]]
384   // CHECK: call void @llvm.nvvm.barrier0()
385   // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]],
386   // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null
387   // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]]
388   //
389   // CHECK: [[SEL_WORKERS]]
390   // CHECK: [[ST:%.+]] = load i8, i8* [[OMP_EXEC_STATUS]],
391   // CHECK: [[IS_ACTIVE:%.+]] = icmp ne i8 [[ST]], 0
392   // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]]
393   //
394   // CHECK: [[EXEC_PARALLEL]]
395   // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
396   //
397   // CHECK: [[TERM_PARALLEL]]
398   // CHECK: br label {{%?}}[[BAR_PARALLEL]]
399   //
400   // CHECK: [[BAR_PARALLEL]]
401   // CHECK: call void @llvm.nvvm.barrier0()
402   // CHECK: br label {{%?}}[[AWAIT_WORK]]
403   //
404   // CHECK: [[EXIT]]
405   // CHECK: ret void
406 
407   // CHECK: define {{.*}}void [[T4:@__omp_offloading_.+static.+l326]](i[[SZ]]
408   // Create local storage for each capture.
409   // CHECK:  [[LOCAL_A:%.+]] = alloca i[[SZ]]
410   // CHECK:  [[LOCAL_AA:%.+]] = alloca i[[SZ]]
411   // CHECK:  [[LOCAL_AAA:%.+]] = alloca i[[SZ]]
412   // CHECK:  [[LOCAL_B:%.+]] = alloca [10 x i32]*
413   // CHECK-DAG:  store i[[SZ]] [[ARG_A:%.+]], i[[SZ]]* [[LOCAL_A]]
414   // CHECK-DAG:  store i[[SZ]] [[ARG_AA:%.+]], i[[SZ]]* [[LOCAL_AA]]
415   // CHECK-DAG:  store i[[SZ]] [[ARG_AAA:%.+]], i[[SZ]]* [[LOCAL_AAA]]
416   // CHECK-DAG:  store [10 x i32]* [[ARG_B:%.+]], [10 x i32]** [[LOCAL_B]]
417   // Store captures in the context.
418   // CHECK-64-DAG:   [[REF_A:%.+]] = bitcast i[[SZ]]* [[LOCAL_A]] to i32*
419   // CHECK-DAG:      [[REF_AA:%.+]] = bitcast i[[SZ]]* [[LOCAL_AA]] to i16*
420   // CHECK-DAG:      [[REF_AAA:%.+]] = bitcast i[[SZ]]* [[LOCAL_AAA]] to i8*
421   // CHECK-DAG:      [[REF_B:%.+]] = load [10 x i32]*, [10 x i32]** [[LOCAL_B]],
422   //
423   // CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
424   // CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
425   // CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
426   // CHECK-DAG: [[TH_LIMIT:%.+]] = sub nuw i32 [[NTH]], [[WS]]
427   // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[TH_LIMIT]]
428   // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[CHECK_MASTER:.+]]
429   //
430   // CHECK: [[WORKER]]
431   // CHECK: {{call|invoke}} void [[T4]]_worker()
432   // CHECK: br label {{%?}}[[EXIT:.+]]
433   //
434   // CHECK: [[CHECK_MASTER]]
435   // CHECK-DAG: [[CMTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
436   // CHECK-DAG: [[CMNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
437   // CHECK-DAG: [[CMWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
438   // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]],
439   // CHECK: br i1 [[IS_MASTER]], label {{%?}}[[MASTER:.+]], label {{%?}}[[EXIT]]
440   //
441   // CHECK: [[MASTER]]
442   // CHECK-DAG: [[MNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
443   // CHECK-DAG: [[MWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
444   // CHECK: [[MTMP1:%.+]] = sub nuw i32 [[MNTH]], [[MWS]]
445   // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]]
446   // CHECK-64-DAG: load i32, i32* [[REF_A]]
447   // CHECK-32-DAG: load i32, i32* [[LOCAL_A]]
448   // CHECK-DAG:    load i16, i16* [[REF_AA]]
449   // CHECK-DAG:    getelementptr inbounds [10 x i32], [10 x i32]* [[REF_B]], i[[SZ]] 0, i[[SZ]] 2
450   // CHECK: br label {{%?}}[[TERMINATE:.+]]
451   //
452   // CHECK: [[TERMINATE]]
453   // CHECK: call void @__kmpc_kernel_deinit(
454   // CHECK: call void @llvm.nvvm.barrier0()
455   // CHECK: br label {{%?}}[[EXIT]]
456   //
457   // CHECK: [[EXIT]]
458   // CHECK: ret void
459 
460 
461 
462   // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+S1.+l344}}_worker()
463   // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
464   // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
465   // CHECK: [[GTID:%.+]] = call i32 @__kmpc_global_thread_num(%struct.ident_t*
466   // CHECK: store i8* null, i8** [[OMP_WORK_FN]],
467   // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]],
468   // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
469   //
470   // CHECK: [[AWAIT_WORK]]
471   // CHECK: call void @llvm.nvvm.barrier0()
472   // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]],
473   // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null
474   // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]]
475   //
476   // CHECK: [[SEL_WORKERS]]
477   // CHECK: [[ST:%.+]] = load i8, i8* [[OMP_EXEC_STATUS]],
478   // CHECK: [[IS_ACTIVE:%.+]] = icmp ne i8 [[ST]], 0
479   // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]]
480   //
481   // CHECK: [[EXEC_PARALLEL]]
482   // CHECK: [[WORK_FN:%.+]] = bitcast i8* [[WORK]] to void (i16, i32)*
483   // CHECK: call void [[WORK_FN]](i16 0, i32 [[GTID]])
484   // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
485   //
486   // CHECK: [[TERM_PARALLEL]]
487   // CHECK: br label {{%?}}[[BAR_PARALLEL]]
488   //
489   // CHECK: [[BAR_PARALLEL]]
490   // CHECK: call void @llvm.nvvm.barrier0()
491   // CHECK: br label {{%?}}[[AWAIT_WORK]]
492   //
493   // CHECK: [[EXIT]]
494   // CHECK: ret void
495 
496   // CHECK: define {{.*}}void [[T5:@__omp_offloading_.+S1.+l344]](
497   // Create local storage for each capture.
498   // CHECK:       [[LOCAL_THIS:%.+]] = alloca [[S1:%struct.*]]*
499   // CHECK:       [[LOCAL_B:%.+]] = alloca i[[SZ]]
500   // CHECK:       [[LOCAL_VLA1:%.+]] = alloca i[[SZ]]
501   // CHECK:       [[LOCAL_VLA2:%.+]] = alloca i[[SZ]]
502   // CHECK:       [[LOCAL_C:%.+]] = alloca i16*
503   // CHECK-DAG:   store [[S1]]* [[ARG_THIS:%.+]], [[S1]]** [[LOCAL_THIS]]
504   // CHECK-DAG:   store i[[SZ]] [[ARG_B:%.+]], i[[SZ]]* [[LOCAL_B]]
505   // CHECK-DAG:   store i[[SZ]] [[ARG_VLA1:%.+]], i[[SZ]]* [[LOCAL_VLA1]]
506   // CHECK-DAG:   store i[[SZ]] [[ARG_VLA2:%.+]], i[[SZ]]* [[LOCAL_VLA2]]
507   // CHECK-DAG:   store i16* [[ARG_C:%.+]], i16** [[LOCAL_C]]
508   // Store captures in the context.
509   // CHECK-DAG:   [[REF_THIS:%.+]] = load [[S1]]*, [[S1]]** [[LOCAL_THIS]],
510   // CHECK-64-DAG:[[REF_B:%.+]] = bitcast i[[SZ]]* [[LOCAL_B]] to i32*
511   // CHECK-DAG:   [[VAL_VLA1:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA1]],
512   // CHECK-DAG:   [[VAL_VLA2:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA2]],
513   // CHECK-DAG:   [[REF_C:%.+]] = load i16*, i16** [[LOCAL_C]],
514   //
515   // CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
516   // CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
517   // CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
518   // CHECK-DAG: [[TH_LIMIT:%.+]] = sub nuw i32 [[NTH]], [[WS]]
519   // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[TH_LIMIT]]
520   // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[CHECK_MASTER:.+]]
521   //
522   // CHECK: [[WORKER]]
523   // CHECK: {{call|invoke}} void [[T5]]_worker()
524   // CHECK: br label {{%?}}[[EXIT:.+]]
525   //
526   // CHECK: [[CHECK_MASTER]]
527   // CHECK-DAG: [[CMTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
528   // CHECK-DAG: [[CMNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
529   // CHECK-DAG: [[CMWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
530   // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]],
531   // CHECK: br i1 [[IS_MASTER]], label {{%?}}[[MASTER:.+]], label {{%?}}[[EXIT]]
532   //
533   // CHECK: [[MASTER]]
534   // CHECK-DAG: [[MNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
535   // CHECK-DAG: [[MWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
536   // CHECK: [[MTMP1:%.+]] = sub nuw i32 [[MNTH]], [[MWS]]
537   // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]]
538   // Use captures.
539   // CHECK-DAG:   getelementptr inbounds [[S1]], [[S1]]* [[REF_THIS]], i32 0, i32 0
540   // CHECK-64-DAG:load i32, i32* [[REF_B]]
541   // CHECK-32-DAG:load i32, i32* [[LOCAL_B]]
542   // CHECK-DAG:   getelementptr inbounds i16, i16* [[REF_C]], i[[SZ]] %{{.+}}
543   // CHECK: call i32 [[BAZ:@.*baz.*]](i32 %
544   // CHECK: br label {{%?}}[[TERMINATE:.+]]
545   //
546   // CHECK: [[TERMINATE]]
547   // CHECK: call void @__kmpc_kernel_deinit(
548   // CHECK: call void @llvm.nvvm.barrier0()
549   // CHECK: br label {{%?}}[[EXIT]]
550   //
551   // CHECK: [[EXIT]]
552   // CHECK: ret void
553 
554   // CHECK: define i32 [[BAZ]](i32 [[F:%.*]])
555   // CHECK: [[ZERO_ADDR:%.+]] = alloca i32,
556   // CHECK: [[GTID:%.+]] = call i32 @__kmpc_global_thread_num(%struct.ident_t*
557   // CHECK: [[GTID_ADDR:%.+]] = alloca i32,
558   // CHECK: store i32 0, i32* [[ZERO_ADDR]]
559   // CHECK: [[PTR:%.+]] = call i8* @__kmpc_data_sharing_push_stack(i{{64|32}} 4, i16 0)
560   // CHECK: [[REC_ADDR:%.+]] = bitcast i8* [[PTR]] to %struct._globalized_locals_ty*
561   // CHECK: [[F_PTR:%.+]] = getelementptr inbounds %struct._globalized_locals_ty, %struct._globalized_locals_ty* [[REC_ADDR]], i32 0, i32 0
562   // CHECK: store i32 [[F]], i32* [[F_PTR]],
563   // CHECK: store i32 [[GTID]], i32* [[GTID_ADDR]],
564   // CHECK: icmp eq i32
565   // CHECK: br i1
566 
567   // CHECK: call void @__kmpc_kernel_prepare_parallel(i8* bitcast (void (i16, i32)* @{{.+}} to i8*), i16 1)
568   // CHECK: call void @__kmpc_begin_sharing_variables(i8*** [[SHARED_PTR:%.+]], i{{64|32}} 1)
569   // CHECK: [[SHARED:%.+]] = load i8**, i8*** [[SHARED_PTR]],
570   // CHECK: [[REF:%.+]] = getelementptr inbounds i8*, i8** [[SHARED]], i{{64|32}} 0
571   // CHECK: [[F_REF:%.+]] = bitcast i32* [[F_PTR]] to i8*
572   // CHECK: store i8* [[F_REF]], i8** [[REF]],
573   // CHECK: call void @llvm.nvvm.barrier0()
574   // CHECK: call void @llvm.nvvm.barrier0()
575   // CHECK: call void @__kmpc_end_sharing_variables()
576   // CHECK: br label
577 
578   // CHECK: [[RES:%.+]] = call i16 @__kmpc_parallel_level(%struct.ident_t* @{{.+}}, i32 [[GTID]])
579   // CHECK: icmp ne i16 [[RES]], 0
580   // CHECK: br i1
581 
582   // CHECK: call void @__kmpc_serialized_parallel(%struct.ident_t* @{{.+}}, i32 [[GTID]])
583   // CHECK: call void @__omp_outlined__(i32* [[GTID_ADDR]], i32* [[ZERO_ADDR]], i32* [[F_PTR]])
584   // CHECK: call void @__kmpc_end_serialized_parallel(%struct.ident_t* @{{.+}}, i32 [[GTID]])
585   // CHECK: br label
586 
587   // CHECK: call void @__omp_outlined__(i32* [[GTID_ADDR]], i32* [[ZERO_ADDR]], i32* [[F_PTR]])
588   // CHECK: br label
589 
590   // CHECK: [[RES:%.+]] = load i32, i32* [[F_PTR]],
591   // CHECK: call void @__kmpc_data_sharing_pop_stack(i8* [[PTR]])
592   // CHECK: ret i32 [[RES]]
593 
594 
595   // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l309}}_worker()
596   // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
597   // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
598   // CHECK: store i8* null, i8** [[OMP_WORK_FN]],
599   // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]],
600   // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
601   //
602   // CHECK: [[AWAIT_WORK]]
603   // CHECK: call void @llvm.nvvm.barrier0()
604   // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]],
605   // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null
606   // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]]
607   //
608   // CHECK: [[SEL_WORKERS]]
609   // CHECK: [[ST:%.+]] = load i8, i8* [[OMP_EXEC_STATUS]],
610   // CHECK: [[IS_ACTIVE:%.+]] = icmp ne i8 [[ST]], 0
611   // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]]
612   //
613   // CHECK: [[EXEC_PARALLEL]]
614   // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
615   //
616   // CHECK: [[TERM_PARALLEL]]
617   // CHECK: br label {{%?}}[[BAR_PARALLEL]]
618   //
619   // CHECK: [[BAR_PARALLEL]]
620   // CHECK: call void @llvm.nvvm.barrier0()
621   // CHECK: br label {{%?}}[[AWAIT_WORK]]
622   //
623   // CHECK: [[EXIT]]
624   // CHECK: ret void
625 
626   // CHECK: define {{.*}}void [[T6:@__omp_offloading_.+template.+l309]](i[[SZ]]
627   // Create local storage for each capture.
628   // CHECK:  [[LOCAL_A:%.+]] = alloca i[[SZ]]
629   // CHECK:  [[LOCAL_AA:%.+]] = alloca i[[SZ]]
630   // CHECK:  [[LOCAL_B:%.+]] = alloca [10 x i32]*
631   // CHECK-DAG:  store i[[SZ]] [[ARG_A:%.+]], i[[SZ]]* [[LOCAL_A]]
632   // CHECK-DAG:  store i[[SZ]] [[ARG_AA:%.+]], i[[SZ]]* [[LOCAL_AA]]
633   // CHECK-DAG:   store [10 x i32]* [[ARG_B:%.+]], [10 x i32]** [[LOCAL_B]]
634   // Store captures in the context.
635   // CHECK-64-DAG:[[REF_A:%.+]] = bitcast i[[SZ]]* [[LOCAL_A]] to i32*
636   // CHECK-DAG:   [[REF_AA:%.+]] = bitcast i[[SZ]]* [[LOCAL_AA]] to i16*
637   // CHECK-DAG:   [[REF_B:%.+]] = load [10 x i32]*, [10 x i32]** [[LOCAL_B]],
638   //
639   // CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
640   // CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
641   // CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
642   // CHECK-DAG: [[TH_LIMIT:%.+]] = sub nuw i32 [[NTH]], [[WS]]
643   // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[TH_LIMIT]]
644   // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[CHECK_MASTER:.+]]
645   //
646   // CHECK: [[WORKER]]
647   // CHECK: {{call|invoke}} void [[T6]]_worker()
648   // CHECK: br label {{%?}}[[EXIT:.+]]
649   //
650   // CHECK: [[CHECK_MASTER]]
651   // CHECK-DAG: [[CMTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
652   // CHECK-DAG: [[CMNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
653   // CHECK-DAG: [[CMWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
654   // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]],
655   // CHECK: br i1 [[IS_MASTER]], label {{%?}}[[MASTER:.+]], label {{%?}}[[EXIT]]
656   //
657   // CHECK: [[MASTER]]
658   // CHECK-DAG: [[MNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
659   // CHECK-DAG: [[MWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
660   // CHECK: [[MTMP1:%.+]] = sub nuw i32 [[MNTH]], [[MWS]]
661   // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]]
662   //
663   // CHECK-64-DAG: load i32, i32* [[REF_A]]
664   // CHECK-32-DAG: load i32, i32* [[LOCAL_A]]
665   // CHECK-DAG:    load i16, i16* [[REF_AA]]
666   // CHECK-DAG:    getelementptr inbounds [10 x i32], [10 x i32]* [[REF_B]], i[[SZ]] 0, i[[SZ]] 2
667   //
668   // CHECK: br label {{%?}}[[TERMINATE:.+]]
669   //
670   // CHECK: [[TERMINATE]]
671   // CHECK: call void @__kmpc_kernel_deinit(
672   // CHECK: call void @llvm.nvvm.barrier0()
673   // CHECK: br label {{%?}}[[EXIT]]
674   //
675   // CHECK: [[EXIT]]
676   // CHECK: ret void
677 #endif
678