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