1 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
2 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
3 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
4 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
5 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
6 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
7 // expected-no-diagnostics
8 #ifndef HEADER
9 #define HEADER
10 
11 
12 // CHECK-DAG: [[GA:@.+]] = global double 1.000000e+00
13 // CHECK-DAG: [[GB:@.+]] = global double 2.000000e+00
14 // CHECK-DAG: [[GC:@.+]] = global double 3.000000e+00
15 // CHECK-DAG: [[GD:@.+]] = global double 4.000000e+00
16 // CHECK-DAG: [[FA:@.+]] = internal global float 5.000000e+00
17 // CHECK-DAG: [[FB:@.+]] = internal global float 6.000000e+00
18 // CHECK-DAG: [[FC:@.+]] = internal global float 7.000000e+00
19 // CHECK-DAG: [[FD:@.+]] = internal global float 8.000000e+00
20 // CHECK-DAG: [[BA:@.+]] = internal global float 9.000000e+00
21 // CHECK-DAG: [[BB:@.+]] = internal global float 1.000000e+01
22 // CHECK-DAG: [[BC:@.+]] = internal global float 1.100000e+01
23 // CHECK-DAG: [[BD:@.+]] = internal global float 1.200000e+01
24 // CHECK-DAG: [[TBA:@.+]] = {{.*}}global float 1.700000e+01
25 // CHECK-DAG: [[TBB:@.+]] = {{.*}}global float 1.800000e+01
26 // CHECK-DAG: [[TBC:@.+]] = {{.*}}global float 1.900000e+01
27 // CHECK-DAG: [[TBD:@.+]] = {{.*}}global float 2.000000e+01
28 
29 double Ga = 1.0;
30 double Gb = 2.0;
31 double Gc = 3.0;
32 double Gd = 4.0;
33 
34 // CHECK: define {{.*}} @{{.*}}foo{{.*}}(
35 // CHECK-SAME: i16 {{[^,]*}}[[A:%[^,]+]],
36 // CHECK-SAME: i16 {{[^,]*}}[[B:%[^,]+]],
37 // CHECK-SAME: i16 {{[^,]*}}[[C:%[^,]+]],
38 // CHECK-SAME: i16 {{[^,]*}}[[D:%[^,]+]])
39 // CHECK: [[LA:%.+]] = alloca i16,
40 // CHECK: [[LB:%.+]] = alloca i16,
41 // CHECK: [[LC:%.+]] = alloca i16,
42 // CHECK: [[LD:%.+]] = alloca i16,
43 int foo(short a, short b, short c, short d){
44   static float Sa = 5.0;
45   static float Sb = 6.0;
46   static float Sc = 7.0;
47   static float Sd = 8.0;
48 
49   // CHECK-DAG:    [[VALLB:%.+]] = load i16, i16* [[LB]],
50   // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* [[GB]],
51   // CHECK-DAG:    [[VALFB:%.+]] = load float, float* [[FB]],
52   // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* [[GC]],
53   // CHECK-DAG:    [[VALLC:%.+]] = load i16, i16* [[LC]],
54   // CHECK-DAG:    [[VALFC:%.+]] = load float, float* [[FC]],
55   // CHECK-DAG:    [[VALLD:%.+]] = load i16, i16* [[LD]],
56   // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* [[GD]],
57   // CHECK-DAG:    [[VALFD:%.+]] = load float, float* [[FD]],
58 
59   // 3 local vars being captured.
60 
61   // CHECK-DAG: store i16 [[VALLB]], i16* [[CONVLB:%.+]],
62   // CHECK-DAG: [[CONVLB]] = bitcast i[[sz:64|32]]* [[CADDRLB:%.+]] to i16*
63   // CHECK-DAG: [[CVALLB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLB]],
64   // CHECK-DAG: store i[[sz]] [[CVALLB]], i[[sz]]* [[CBP:%.+]],
65   // CHECK-DAG: [[CBP]] = bitcast i8** [[GEPLB:%.+]] to i[[sz]]*
66   // CHECK-DAG: [[GEPLB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
67 
68   // CHECK-DAG: store i16 [[VALLC]], i16* [[CONVLC:%.+]],
69   // CHECK-DAG: [[CONVLC]] = bitcast i[[sz]]* [[CADDRLC:%.+]] to i16*
70   // CHECK-DAG: [[CVALLC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLC]],
71   // CHECK-DAG: store i[[sz]] [[CVALLC]], i[[sz]]* [[CBP:%.+]],
72   // CHECK-DAG: [[CBP]] = bitcast i8** [[GEPLC:%.+]] to i[[sz]]*
73   // CHECK-DAG: [[GEPLC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
74 
75   // CHECK-DAG: store i16 [[VALLD]], i16* [[CONVLD:%.+]],
76   // CHECK-DAG: [[CONVLD]] = bitcast i[[sz]]* [[CADDRLD:%.+]] to i16*
77   // CHECK-DAG: [[CVALLD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLD]],
78   // CHECK-DAG: store i[[sz]] [[CVALLD]], i[[sz]]* [[CBP:%.+]],
79   // CHECK-DAG: [[CBP]] = bitcast i8** [[GEPLD:%.+]] to i[[sz]]*
80   // CHECK-DAG: [[GEPLD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
81 
82   // 3 static vars being captured.
83 
84   // CHECK-DAG: store float [[VALFB]], float* [[CONVFB:%.+]],
85   // CHECK-DAG: [[CONVFB]] = bitcast i[[sz]]* [[CADDRFB:%.+]] to float*
86   // CHECK-DAG: [[CVALFB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFB]],
87   // CHECK-DAG: store i[[sz]] [[CVALFB]], i[[sz]]* [[CBP:%.+]],
88   // CHECK-DAG: [[CBP]] = bitcast i8** [[GEPFB:%.+]] to i[[sz]]*
89   // CHECK-DAG: [[GEPFB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
90 
91   // CHECK-DAG: store float [[VALFC]], float* [[CONVFC:%.+]],
92   // CHECK-DAG: [[CONVFC]] = bitcast i[[sz]]* [[CADDRFC:%.+]] to float*
93   // CHECK-DAG: [[CVALFC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFC]],
94   // CHECK-DAG: store i[[sz]] [[CVALFC]], i[[sz]]* [[CBP:%.+]],
95   // CHECK-DAG: [[CBP]] = bitcast i8** [[GEPFC:%.+]] to i[[sz]]*
96   // CHECK-DAG: [[GEPFC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
97 
98   // CHECK-DAG: store float [[VALFD]], float* [[CONVFD:%.+]],
99   // CHECK-DAG: [[CONVFD]] = bitcast i[[sz]]* [[CADDRFD:%.+]] to float*
100   // CHECK-DAG: [[CVALFD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFD]],
101   // CHECK-DAG: store i[[sz]] [[CVALFD]], i[[sz]]* [[CBP:%.+]],
102   // CHECK-DAG: [[CBP]] = bitcast i8** [[GEPFD:%.+]] to i[[sz]]*
103   // CHECK-DAG: [[GEPFD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
104 
105   // 3 static global vars being captured.
106 
107   // CHECK-64-DAG: store double [[VALGB]], double* [[CONVGB:%.+]],
108   // CHECK-64-DAG: [[CONVGB]] = bitcast i[[sz]]* [[CADDRGB:%.+]] to double*
109   // CHECK-64-DAG: [[CVALGB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGB]],
110   // CHECK-64-DAG: store i[[sz]] [[CVALGB]], i[[sz]]* [[CBP:%.+]],
111   // CHECK-64-DAG: [[CBP]] = bitcast i8** [[GEPGB:%.+]] to i[[sz]]*
112   // CHECK-32-DAG: store double* @Gb, double** [[CBP:%.+]],
113   // CHECK-32-DAG: [[CBP]] = bitcast i8** [[GEPGB:%.+]] to double**
114   // CHECK-DAG: [[GEPGB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
115 
116   // CHECK-64-DAG: store double [[VALGC]], double* [[CONVGC:%.+]],
117   // CHECK-64-DAG: [[CONVGC]] = bitcast i[[sz]]* [[CADDRGC:%.+]] to double*
118   // CHECK-64-DAG: [[CVALGC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGC]],
119   // CHECK-64-DAG: store i[[sz]] [[CVALGC]], i[[sz]]* [[CBP:%.+]],
120   // CHECK-64-DAG: [[CBP]] = bitcast i8** [[GEPGC:%.+]] to i[[sz]]*
121   // CHECK-32-DAG: store double* @Gc, double** [[CBP:%.+]],
122   // CHECK-32-DAG: [[CBP]] = bitcast i8** [[GEPGC:%.+]] to double**
123   // CHECK-DAG: [[GEPGC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
124 
125   // CHECK-64-DAG: store double [[VALGD]], double* [[CONVGD:%.+]],
126   // CHECK-64-DAG: [[CONVGD]] = bitcast i[[sz]]* [[CADDRGD:%.+]] to double*
127   // CHECK-64-DAG: [[CVALGD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGD]],
128   // CHECK-64-DAG: store i[[sz]] [[CVALGD]], i[[sz]]* [[CBP:%.+]],
129   // CHECK-64-DAG: [[CBP]] = bitcast i8** [[GEPGD:%.+]] to i[[sz]]*
130   // CHECK-32-DAG: store double* @Gd, double** [[CBP:%.+]],
131   // CHECK-32-DAG: [[CBP]] = bitcast i8** [[GEPGD:%.+]] to double**
132   // CHECK-DAG: [[GEPGD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
133 
134   // CHECK: call i32 @__tgt_target
135   // CHECK: call void [[OFFLOADF:@.+]](
136   // Capture b, Gb, Sb, Gc, c, Sc, d, Gd, Sd
137   #pragma omp target if(Ga>0.0 && a>0 && Sa>0.0)
138   {
139     b += 1;
140     Gb += 1.0;
141     Sb += 1.0;
142 
143     // CHECK: define internal void [[OFFLOADF]]({{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}})
144     // The parallel region only uses 3 captures.
145     // CHECK:     call {{.*}}@__kmpc_fork_call(%ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}}), {{.+}}* %{{.+}}, {{.+}}* %{{.+}}, {{.+}}* %{{.+}})
146     // CHECK:     call void @.omp_outlined.(i32* %{{.+}}, i32* %{{.+}}, {{.+}}* %{{.+}}, {{.+}}* %{{.+}}, {{.+}}* %{{.+}})
147     // Capture d, Gd, Sd,
148 
149     // CHECK: define internal void [[PARF]](i32* noalias %{{.*}}, i32* noalias %{{.*}},
150     #pragma omp parallel if(Gc>0.0 && c>0 && Sc>0.0)
151     {
152       d += 1;
153       Gd += 1.0;
154       Sd += 1.0;
155     }
156   }
157   return a + b + c + d + (int)Sa + (int)Sb + (int)Sc + (int)Sd;
158 }
159 
160 // CHECK: define {{.*}} @{{.*}}bar{{.*}}(
161 // CHECK-SAME: i16 {{[^,]*}}[[A:%[^,]+]],
162 // CHECK-SAME: i16 {{[^,]*}}[[B:%[^,]+]],
163 // CHECK-SAME: i16 {{[^,]*}}[[C:%[^,]+]],
164 // CHECK-SAME: i16 {{[^,]*}}[[D:%[^,]+]])
165 // CHECK: [[LA:%.+]] = alloca i16
166 // CHECK: [[LB:%.+]] = alloca i16
167 // CHECK: [[LC:%.+]] = alloca i16
168 // CHECK: [[LD:%.+]] = alloca i16
169 int bar(short a, short b, short c, short d){
170   static float Sa = 9.0;
171   static float Sb = 10.0;
172   static float Sc = 11.0;
173   static float Sd = 12.0;
174 
175   // CHECK: call void {{.*}}@__kmpc_fork_call(%ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}}), i16* %{{.+}}, i16* %{{.+}}, i16* %{{.+}}, i16* %{{.+}})
176   // CHECK: define internal void [[PARF]](i32* noalias %{{.*}}, i32* noalias %{{.*}}, i16* dereferenceable(2) [[A:%.+]], i16* dereferenceable(2) [[B:%.+]], i16* dereferenceable(2) [[C:%.+]], i16* dereferenceable(2) [[D:%.+]])
177   // Capture a, b, c, d
178   // CHECK: [[ALLOCLA:%.+]] = alloca i16
179   // CHECK: [[ALLOCLB:%.+]] = alloca i16
180   // CHECK: [[ALLOCLC:%.+]] = alloca i16
181   // CHECK: [[ALLOCLD:%.+]] = alloca i16
182   // CHECK: [[LLA:%.+]] = load i16*, i16** [[ALLOCLA]],
183   // CHECK: [[LLB:%.+]] = load i16*, i16** [[ALLOCLB]],
184   // CHECK: [[LLC:%.+]] = load i16*, i16** [[ALLOCLC]],
185   // CHECK: [[LLD:%.+]] = load i16*, i16** [[ALLOCLD]],
186   #pragma omp parallel
187   {
188     // CHECK-DAG:    [[VALLB:%.+]] = load i16, i16* [[LLB]],
189     // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* [[GB]],
190     // CHECK-DAG:    [[VALFB:%.+]] = load float, float* [[BB]],
191     // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* [[GC]],
192     // CHECK-DAG:    [[VALLC:%.+]] = load i16, i16* [[LLC]],
193     // CHECK-DAG:    [[VALFC:%.+]] = load float, float* [[BC]],
194     // CHECK-DAG:    [[VALLD:%.+]] = load i16, i16* [[LLD]],
195     // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* [[GD]],
196     // CHECK-DAG:    [[VALFD:%.+]] = load float, float* [[BD]],
197 
198     // 3 local vars being captured.
199 
200     // CHECK-DAG: store i16 [[VALLB]], i16* [[CONVLB:%.+]],
201     // CHECK-DAG: [[CONVLB]] = bitcast i[[sz:64|32]]* [[CADDRLB:%.+]] to i16*
202     // CHECK-DAG: [[CVALLB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLB]],
203     // CHECK-DAG: store i[[sz]] [[CVALLB]], i[[sz]]* [[CBP:%.+]],
204     // CHECK-DAG: [[CBP]] = bitcast i8** [[GEPLB:%.+]] to i[[sz]]*
205     // CHECK-DAG: [[GEPLB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
206 
207     // CHECK-DAG: store i16 [[VALLC]], i16* [[CONVLC:%.+]],
208     // CHECK-DAG: [[CONVLC]] = bitcast i[[sz]]* [[CADDRLC:%.+]] to i16*
209     // CHECK-DAG: [[CVALLC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLC]],
210     // CHECK-DAG: store i[[sz]] [[CVALLC]], i[[sz]]* [[CBP:%.+]],
211     // CHECK-DAG: [[CBP]] = bitcast i8** [[GEPLC:%.+]] to i[[sz]]*
212     // CHECK-DAG: [[GEPLC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
213 
214     // CHECK-DAG: store i16 [[VALLD]], i16* [[CONVLD:%.+]],
215     // CHECK-DAG: [[CONVLD]] = bitcast i[[sz]]* [[CADDRLD:%.+]] to i16*
216     // CHECK-DAG: [[CVALLD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLD]],
217     // CHECK-DAG: store i[[sz]] [[CVALLD]], i[[sz]]* [[CBP:%.+]],
218     // CHECK-DAG: [[CBP]] = bitcast i8** [[GEPLD:%.+]] to i[[sz]]*
219     // CHECK-DAG: [[GEPLD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
220 
221     // 3 static vars being captured.
222 
223     // CHECK-DAG: store float [[VALFB]], float* [[CONVFB:%.+]],
224     // CHECK-DAG: [[CONVFB]] = bitcast i[[sz]]* [[CADDRFB:%.+]] to float*
225     // CHECK-DAG: [[CVALFB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFB]],
226     // CHECK-DAG: store i[[sz]] [[CVALFB]], i[[sz]]* [[CBP:%.+]],
227     // CHECK-DAG: [[CBP]] = bitcast i8** [[GEPFB:%.+]] to i[[sz]]*
228     // CHECK-DAG: [[GEPFB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
229 
230     // CHECK-DAG: store float [[VALFC]], float* [[CONVFC:%.+]],
231     // CHECK-DAG: [[CONVFC]] = bitcast i[[sz]]* [[CADDRFC:%.+]] to float*
232     // CHECK-DAG: [[CVALFC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFC]],
233     // CHECK-DAG: store i[[sz]] [[CVALFC]], i[[sz]]* [[CBP:%.+]],
234     // CHECK-DAG: [[CBP]] = bitcast i8** [[GEPFC:%.+]] to i[[sz]]*
235     // CHECK-DAG: [[GEPFC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
236 
237     // CHECK-DAG: store float [[VALFD]], float* [[CONVFD:%.+]],
238     // CHECK-DAG: [[CONVFD]] = bitcast i[[sz]]* [[CADDRFD:%.+]] to float*
239     // CHECK-DAG: [[CVALFD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFD]],
240     // CHECK-DAG: store i[[sz]] [[CVALFD]], i[[sz]]* [[CBP:%.+]],
241     // CHECK-DAG: [[CBP]] = bitcast i8** [[GEPFD:%.+]] to i[[sz]]*
242     // CHECK-DAG: [[GEPFD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
243 
244     // 3 static global vars being captured.
245 
246     // CHECK-64-DAG: store double [[VALGB]], double* [[CONVGB:%.+]],
247     // CHECK-64-DAG: [[CONVGB]] = bitcast i[[sz]]* [[CADDRGB:%.+]] to double*
248     // CHECK-64-DAG: [[CVALGB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGB]],
249     // CHECK-64-DAG: store i[[sz]] [[CVALGB]], i[[sz]]* [[CBP:%.+]],
250     // CHECK-64-DAG: [[CBP]] = bitcast i8** [[GEPGB:%.+]] to i[[sz]]*
251     // CHECK-32-DAG: store double* @Gb, double** [[CBP:%.+]],
252     // CHECK-32-DAG: [[CBP]] = bitcast i8** [[GEPGB:%.+]] to double**
253     // CHECK-DAG: [[GEPGB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
254 
255     // CHECK-64-DAG: store double [[VALGC]], double* [[CONVGC:%.+]],
256     // CHECK-64-DAG: [[CONVGC]] = bitcast i[[sz]]* [[CADDRGC:%.+]] to double*
257     // CHECK-64-DAG: [[CVALGC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGC]],
258     // CHECK-64-DAG: store i[[sz]] [[CVALGC]], i[[sz]]* [[CBP:%.+]],
259     // CHECK-64-DAG: [[CBP]] = bitcast i8** [[GEPGC:%.+]] to i[[sz]]*
260     // CHECK-32-DAG: store double* @Gc, double** [[CBP:%.+]],
261     // CHECK-32-DAG: [[CBP]] = bitcast i8** [[GEPGC:%.+]] to double**
262     // CHECK-DAG: [[GEPGC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
263 
264     // CHECK-64-DAG: store double [[VALGD]], double* [[CONVGD:%.+]],
265     // CHECK-64-DAG: [[CONVGD]] = bitcast i[[sz]]* [[CADDRGD:%.+]] to double*
266     // CHECK-64-DAG: [[CVALGD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGD]],
267     // CHECK-64-DAG: store i[[sz]] [[CVALGD]], i[[sz]]* [[CBP:%.+]],
268     // CHECK-64-DAG: [[CBP]] = bitcast i8** [[GEPGD:%.+]] to i[[sz]]*
269     // CHECK-32-DAG: store double* @Gd, double** [[CBP:%.+]],
270     // CHECK-32-DAG: [[CBP]] = bitcast i8** [[GEPGD:%.+]] to double**
271     // CHECK-DAG: [[GEPGD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
272 
273     // CHECK: call i32 @__tgt_target
274     // CHECK: call void [[OFFLOADF:@.+]](
275     // Capture b, Gb, Sb, Gc, c, Sc, d, Gd, Sd
276     #pragma omp target if(Ga>0.0 && a>0 && Sa>0.0)
277     {
278       b += 1;
279       Gb += 1.0;
280       Sb += 1.0;
281 
282       // CHECK: define internal void [[OFFLOADF]]({{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}})
283       // CHECK: call void {{.*}}@__kmpc_fork_call(%ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}})
284 
285       // CHECK: define internal void [[PARF]](i32* noalias %{{.*}}, i32* noalias %{{.*}}, {{.+}}* dereferenceable({{.+}}) %{{.+}}, {{.+}}* dereferenceable({{.+}}) %{{.+}}, {{.+}}* dereferenceable({{.+}}) %{{.+}})
286       // Capture d, Gd, Sd
287       #pragma omp parallel if(Gc>0.0 && c>0 && Sc>0.0)
288       {
289         d += 1;
290         Gd += 1.0;
291         Sd += 1.0;
292       }
293     }
294   }
295   return a + b + c + d + (int)Sa + (int)Sb + (int)Sc + (int)Sd;
296 }
297 
298 ///
299 /// Tests with template functions.
300 ///
301 
302 // CHECK: define {{.*}} @{{.*}}tbar2{{.*}}(
303 
304 // CHECK: define {{.*}} @{{.*}}tbar{{.*}}(
305 // CHECK-SAME: i16 {{[^,]*}}[[A:%[^,]+]],
306 // CHECK-SAME: i16 {{[^,]*}}[[B:%[^,]+]],
307 // CHECK-SAME: i16 {{[^,]*}}[[C:%[^,]+]],
308 // CHECK-SAME: i16 {{[^,]*}}[[D:%[^,]+]])
309 // CHECK: [[LA:%.+]] = alloca i16
310 // CHECK: [[LB:%.+]] = alloca i16
311 // CHECK: [[LC:%.+]] = alloca i16
312 // CHECK: [[LD:%.+]] = alloca i16
313 template<typename T>
314 int tbar(T a, T b, T c, T d){
315   static float Sa = 17.0;
316   static float Sb = 18.0;
317   static float Sc = 19.0;
318   static float Sd = 20.0;
319 
320   // CHECK: call void {{.*}}@__kmpc_fork_call(%ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}}), i16* %{{.+}}, i16* %{{.+}}, i16* %{{.+}}, i16* %{{.+}})
321   // CHECK: define internal void [[PARF]](i32* noalias %{{.*}}, i32* noalias %{{.*}}, i16* dereferenceable(2) [[A:%.+]], i16* dereferenceable(2) [[B:%.+]], i16* dereferenceable(2) [[C:%.+]], i16* dereferenceable(2) [[D:%.+]])
322   // Capture a, b, c, d
323   // CHECK: [[ALLOCLA:%.+]] = alloca i16
324   // CHECK: [[ALLOCLB:%.+]] = alloca i16
325   // CHECK: [[ALLOCLC:%.+]] = alloca i16
326   // CHECK: [[ALLOCLD:%.+]] = alloca i16
327   // CHECK: [[LLA:%.+]] = load i16*, i16** [[ALLOCLA]],
328   // CHECK: [[LLB:%.+]] = load i16*, i16** [[ALLOCLB]],
329   // CHECK: [[LLC:%.+]] = load i16*, i16** [[ALLOCLC]],
330   // CHECK: [[LLD:%.+]] = load i16*, i16** [[ALLOCLD]],
331   #pragma omp parallel
332   {
333     // CHECK-DAG:    [[VALLB:%.+]] = load i16, i16* [[LLB]],
334     // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* [[GB]],
335     // CHECK-DAG:    [[VALFB:%.+]] = load float, float* [[TBB]],
336     // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* [[GC]],
337     // CHECK-DAG:    [[VALLC:%.+]] = load i16, i16* [[LLC]],
338     // CHECK-DAG:    [[VALFC:%.+]] = load float, float* [[TBC]],
339     // CHECK-DAG:    [[VALLD:%.+]] = load i16, i16* [[LLD]],
340     // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* [[GD]],
341     // CHECK-DAG:    [[VALFD:%.+]] = load float, float* [[TBD]],
342 
343     // 3 local vars being captured.
344 
345     // CHECK-DAG: store i16 [[VALLB]], i16* [[CONVLB:%.+]],
346     // CHECK-DAG: [[CONVLB]] = bitcast i[[sz:64|32]]* [[CADDRLB:%.+]] to i16*
347     // CHECK-DAG: [[CVALLB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLB]],
348     // CHECK-DAG: store i[[sz]] [[CVALLB]], i[[sz]]* [[CBP:%.+]],
349     // CHECK-DAG: [[CBP]] = bitcast i8** [[GEPLB:%.+]] to i[[sz]]*
350     // CHECK-DAG: [[GEPLB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
351 
352     // CHECK-DAG: store i16 [[VALLC]], i16* [[CONVLC:%.+]],
353     // CHECK-DAG: [[CONVLC]] = bitcast i[[sz]]* [[CADDRLC:%.+]] to i16*
354     // CHECK-DAG: [[CVALLC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLC]],
355     // CHECK-DAG: store i[[sz]] [[CVALLC]], i[[sz]]* [[CBP:%.+]],
356     // CHECK-DAG: [[CBP]] = bitcast i8** [[GEPLC:%.+]] to i[[sz]]*
357     // CHECK-DAG: [[GEPLC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
358 
359     // CHECK-DAG: store i16 [[VALLD]], i16* [[CONVLD:%.+]],
360     // CHECK-DAG: [[CONVLD]] = bitcast i[[sz]]* [[CADDRLD:%.+]] to i16*
361     // CHECK-DAG: [[CVALLD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLD]],
362     // CHECK-DAG: store i[[sz]] [[CVALLD]], i[[sz]]* [[CBP:%.+]],
363     // CHECK-DAG: [[CBP]] = bitcast i8** [[GEPLD:%.+]] to i[[sz]]*
364     // CHECK-DAG: [[GEPLD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
365 
366     // 3 static vars being captured.
367 
368     // CHECK-DAG: store float [[VALFB]], float* [[CONVFB:%.+]],
369     // CHECK-DAG: [[CONVFB]] = bitcast i[[sz]]* [[CADDRFB:%.+]] to float*
370     // CHECK-DAG: [[CVALFB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFB]],
371     // CHECK-DAG: store i[[sz]] [[CVALFB]], i[[sz]]* [[CBP:%.+]],
372     // CHECK-DAG: [[CBP]] = bitcast i8** [[GEPFB:%.+]] to i[[sz]]*
373     // CHECK-DAG: [[GEPFB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
374 
375     // CHECK-DAG: store float [[VALFC]], float* [[CONVFC:%.+]],
376     // CHECK-DAG: [[CONVFC]] = bitcast i[[sz]]* [[CADDRFC:%.+]] to float*
377     // CHECK-DAG: [[CVALFC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFC]],
378     // CHECK-DAG: store i[[sz]] [[CVALFC]], i[[sz]]* [[CBP:%.+]],
379     // CHECK-DAG: [[CBP]] = bitcast i8** [[GEPFC:%.+]] to i[[sz]]*
380     // CHECK-DAG: [[GEPFC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
381 
382     // CHECK-DAG: store float [[VALFD]], float* [[CONVFD:%.+]],
383     // CHECK-DAG: [[CONVFD]] = bitcast i[[sz]]* [[CADDRFD:%.+]] to float*
384     // CHECK-DAG: [[CVALFD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFD]],
385     // CHECK-DAG: store i[[sz]] [[CVALFD]], i[[sz]]* [[CBP:%.+]],
386     // CHECK-DAG: [[CBP]] = bitcast i8** [[GEPFD:%.+]] to i[[sz]]*
387     // CHECK-DAG: [[GEPFD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
388 
389     // 3 static global vars being captured.
390 
391     // CHECK-64-DAG: store double [[VALGB]], double* [[CONVGB:%.+]],
392     // CHECK-64-DAG: [[CONVGB]] = bitcast i[[sz]]* [[CADDRGB:%.+]] to double*
393     // CHECK-64-DAG: [[CVALGB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGB]],
394     // CHECK-64-DAG: store i[[sz]] [[CVALGB]], i[[sz]]* [[CBP:%.+]],
395     // CHECK-64-DAG: [[CBP]] = bitcast i8** [[GEPGB:%.+]] to i[[sz]]*
396     // CHECK-32-DAG: store double* @Gb, double** [[CBP:%.+]],
397     // CHECK-32-DAG: [[CBP]] = bitcast i8** [[GEPGB:%.+]] to double**
398     // CHECK-DAG: [[GEPGB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
399 
400     // CHECK-64-DAG: store double [[VALGC]], double* [[CONVGC:%.+]],
401     // CHECK-64-DAG: [[CONVGC]] = bitcast i[[sz]]* [[CADDRGC:%.+]] to double*
402     // CHECK-64-DAG: [[CVALGC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGC]],
403     // CHECK-64-DAG: store i[[sz]] [[CVALGC]], i[[sz]]* [[CBP:%.+]],
404     // CHECK-64-DAG: [[CBP]] = bitcast i8** [[GEPGC:%.+]] to i[[sz]]*
405     // CHECK-32-DAG: store double* @Gc, double** [[CBP:%.+]],
406     // CHECK-32-DAG: [[CBP]] = bitcast i8** [[GEPGC:%.+]] to double**
407     // CHECK-DAG: [[GEPGC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
408 
409     // CHECK-64-DAG: store double [[VALGD]], double* [[CONVGD:%.+]],
410     // CHECK-64-DAG: [[CONVGD]] = bitcast i[[sz]]* [[CADDRGD:%.+]] to double*
411     // CHECK-64-DAG: [[CVALGD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGD]],
412     // CHECK-64-DAG: store i[[sz]] [[CVALGD]], i[[sz]]* [[CBP:%.+]],
413     // CHECK-64-DAG: [[CBP]] = bitcast i8** [[GEPGD:%.+]] to i[[sz]]*
414     // CHECK-32-DAG: store double* @Gd, double** [[CBP:%.+]],
415     // CHECK-32-DAG: [[CBP]] = bitcast i8** [[GEPGD:%.+]] to double**
416     // CHECK-DAG: [[GEPGD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
417 
418     // CHECK: call i32 @__tgt_target
419     // CHECK: call void [[OFFLOADF:@.+]](
420     // Capture b, Gb, Sb, Gc, c, Sc, d, Gd, Sd
421     #pragma omp target if(Ga>0.0 && a>0 && Sa>0.0)
422     {
423       b += 1;
424       Gb += 1.0;
425       Sb += 1.0;
426 
427       // CHECK: define internal void [[OFFLOADF]]({{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}})
428       // CHECK: call void {{.*}}@__kmpc_fork_call(%ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}})
429 
430       // CHECK: define internal void [[PARF]](i32* noalias %{{.*}}, i32* noalias %{{.*}}, {{.+}}* dereferenceable({{.+}}) %{{.+}}, {{.+}}* dereferenceable({{.+}}) %{{.+}}, {{.+}}* dereferenceable({{.+}}) %{{.+}})
431       // Capture d, Gd, Sd
432       #pragma omp parallel if(Gc>0.0 && c>0 && Sc>0.0)
433       {
434         d += 1;
435         Gd += 1.0;
436         Sd += 1.0;
437       }
438     }
439   }
440   return a + b + c + d + (int)Sa + (int)Sb + (int)Sc + (int)Sd;
441 }
442 
443 int tbar2(short a, short b, short c, short d){
444   return tbar(a, b, c, d);
445 }
446 
447 #endif
448