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