1 // expected-no-diagnostics
2 #ifndef HEADER
3 #define HEADER
4 
5 ///
6 /// Implicit maps.
7 ///
8 
9 ///==========================================================================///
10 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64
11 // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
12 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK1 --check-prefix CK1-64
13 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK1 --check-prefix CK1-32
14 // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
15 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK1 --check-prefix CK1-32
16 #ifdef CK1
17 
18 // CK1-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
19 // Map types: OMP_MAP_BYCOPY = 128
20 // CK1-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
21 
22 // CK1-LABEL: implicit_maps_integer
implicit_maps_integer(int a)23 void implicit_maps_integer (int a){
24   int i = a;
25 
26   // CK1-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
27   // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
28   // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
29   // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
30   // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
31   // CK1-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
32   // CK1-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
33   // CK1-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
34   // CK1-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
35   // CK1-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]],
36   // CK1-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32*
37   // CK1-64-DAG: store i32 {{.+}}, i32* [[CADDR]],
38 
39   // CK1: call void [[KERNEL:@.+]](i[[sz]] [[VAL]])
40   #pragma omp target
41   {
42    ++i;
43   }
44 }
45 
46 // CK1: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]])
47 // CK1: [[ADDR:%.+]] = alloca i[[sz]],
48 // CK1: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]],
49 // CK1-64: [[CADDR:%.+]] = bitcast i64* [[ADDR]] to i32*
50 // CK1-64: {{.+}} = load i32, i32* [[CADDR]],
51 // CK1-32: {{.+}} = load i32, i32* [[ADDR]],
52 
53 #endif
54 ///==========================================================================///
55 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64
56 // RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
57 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK2 --check-prefix CK2-64
58 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK2 --check-prefix CK2-32
59 // RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
60 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK2 --check-prefix CK2-32
61 #ifdef CK2
62 
63 // CK2-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
64 // Map types: OMP_MAP_BYCOPY = 128
65 // CK2-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
66 
67 // CK2-LABEL: implicit_maps_integer_reference
implicit_maps_integer_reference(int a)68 void implicit_maps_integer_reference (int a){
69   int &i = a;
70   // CK2-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
71   // CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
72   // CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
73   // CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
74   // CK2-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
75   // CK2-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
76   // CK2-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
77   // CK2-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
78   // CK2-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
79   // CK2-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]],
80   // CK2-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32*
81   // CK2-64-DAG: store i32 {{.+}}, i32* [[CADDR]],
82 
83   // CK2: call void [[KERNEL:@.+]](i[[sz]] [[VAL]])
84   #pragma omp target
85   {
86    ++i;
87   }
88 }
89 
90 // CK2: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]])
91 // CK2: [[ADDR:%.+]] = alloca i[[sz]],
92 // CK2: [[REF:%.+]] = alloca i32*,
93 // CK2: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]],
94 // CK2-64: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32*
95 // CK2-64: store i32* [[CADDR]], i32** [[REF]],
96 // CK2-64: [[RVAL:%.+]] = load i32*, i32** [[REF]],
97 // CK2-64: {{.+}} = load i32, i32* [[RVAL]],
98 // CK2-32: store i32* [[ADDR]], i32** [[REF]],
99 // CK2-32: [[RVAL:%.+]] = load i32*, i32** [[REF]],
100 // CK2-32: {{.+}} = load i32, i32* [[RVAL]],
101 
102 #endif
103 ///==========================================================================///
104 // RUN: %clang_cc1 -DCK3 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-64
105 // RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
106 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK3 --check-prefix CK3-64
107 // RUN: %clang_cc1 -DCK3 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK3 --check-prefix CK3-32
108 // RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
109 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK3 --check-prefix CK3-32
110 #ifdef CK3
111 
112 // CK3-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
113 // Map types: OMP_MAP_BYCOPY = 128
114 // CK3-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
115 
116 // CK3-LABEL: implicit_maps_parameter
implicit_maps_parameter(int a)117 void implicit_maps_parameter (int a){
118 
119   // CK3-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
120   // CK3-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
121   // CK3-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
122   // CK3-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
123   // CK3-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
124   // CK3-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
125   // CK3-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
126   // CK3-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
127   // CK3-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
128   // CK3-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]],
129   // CK3-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32*
130   // CK3-64-DAG: store i32 {{.+}}, i32* [[CADDR]],
131 
132   // CK3: call void [[KERNEL:@.+]](i[[sz]] [[VAL]])
133   #pragma omp target
134   {
135    ++a;
136   }
137 }
138 
139 // CK3: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]])
140 // CK3: [[ADDR:%.+]] = alloca i[[sz]],
141 // CK3: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]],
142 // CK3-64: [[CADDR:%.+]] = bitcast i64* [[ADDR]] to i32*
143 // CK3-64: {{.+}} = load i32, i32* [[CADDR]],
144 // CK3-32: {{.+}} = load i32, i32* [[ADDR]],
145 
146 #endif
147 ///==========================================================================///
148 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-64
149 // RUN: %clang_cc1 -DCK4 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
150 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK4 --check-prefix CK4-64
151 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK4 --check-prefix CK4-32
152 // RUN: %clang_cc1 -DCK4 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
153 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK4 --check-prefix CK4-32
154 #ifdef CK4
155 
156 // CK4-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
157 // Map types: OMP_MAP_BYCOPY = 128
158 // CK4-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
159 
160 // CK4-LABEL: implicit_maps_nested_integer
implicit_maps_nested_integer(int a)161 void implicit_maps_nested_integer (int a){
162   int i = a;
163 
164   // The captures in parallel are by reference. Only the capture in target is by
165   // copy.
166 
167   // CK4: call void {{.+}}@__kmpc_fork_call({{.+}} [[KERNELP1:@.+]] to void (i32*, i32*, ...)*), i32* {{.+}})
168   // CK4: define internal void [[KERNELP1]](i32* {{[^,]+}}, i32* {{[^,]+}}, i32* {{[^,]+}})
169   #pragma omp parallel
170   {
171     // CK4-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
172     // CK4-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
173     // CK4-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
174     // CK4-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
175     // CK4-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
176     // CK4-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
177     // CK4-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
178     // CK4-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
179     // CK4-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
180     // CK4-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]],
181     // CK4-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32*
182     // CK4-64-DAG: store i32 {{.+}}, i32* [[CADDR]],
183 
184     // CK4: call void [[KERNEL:@.+]](i[[sz]] [[VAL]])
185     #pragma omp target
186     {
187       #pragma omp parallel
188       {
189         ++i;
190       }
191     }
192   }
193 }
194 
195 // CK4: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]])
196 // CK4: [[ADDR:%.+]] = alloca i[[sz]],
197 // CK4: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]],
198 // CK4-64: [[CADDR:%.+]] = bitcast i64* [[ADDR]] to i32*
199 // CK4-64: call void {{.+}}@__kmpc_fork_call({{.+}} [[KERNELP2:@.+]] to void (i32*, i32*, ...)*), i32* [[CADDR]])
200 // CK4-32: call void {{.+}}@__kmpc_fork_call({{.+}} [[KERNELP2:@.+]] to void (i32*, i32*, ...)*), i32* [[ADDR]])
201 // CK4: define internal void [[KERNELP2]](i32* {{[^,]+}}, i32* {{[^,]+}}, i32* {{[^,]+}})
202 #endif
203 ///==========================================================================///
204 // RUN: %clang_cc1 -DCK5 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK5 --check-prefix CK5-64
205 // RUN: %clang_cc1 -DCK5 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
206 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK5 --check-prefix CK5-64
207 // RUN: %clang_cc1 -DCK5 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK5 --check-prefix CK5-32
208 // RUN: %clang_cc1 -DCK5 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
209 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK5 --check-prefix CK5-32
210 #ifdef CK5
211 
212 // CK5-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
213 // Map types: OMP_MAP_BYCOPY = 128
214 // CK5-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
215 
216 // CK5-LABEL: implicit_maps_nested_integer_and_enum
implicit_maps_nested_integer_and_enum(int a)217 void implicit_maps_nested_integer_and_enum (int a){
218   enum Bla {
219     SomeEnum = 0x09
220   };
221 
222   // Using an enum should not change the mapping information.
223   int  i = a;
224 
225   // CK5-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
226   // CK5-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
227   // CK5-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
228   // CK5-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
229   // CK5-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
230   // CK5-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
231   // CK5-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
232   // CK5-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
233   // CK5-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
234   // CK5-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]],
235   // CK5-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32*
236   // CK5-64-DAG: store i32 {{.+}}, i32* [[CADDR]],
237 
238   // CK5: call void [[KERNEL:@.+]](i[[sz]] [[VAL]])
239   #pragma omp target
240   {
241     ++i;
242     i += SomeEnum;
243   }
244 }
245 
246 // CK5: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]])
247 // CK5: [[ADDR:%.+]] = alloca i[[sz]],
248 // CK5: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]],
249 // CK5-64: [[CADDR:%.+]] = bitcast i64* [[ADDR]] to i32*
250 // CK5-64: {{.+}} = load i32, i32* [[CADDR]],
251 // CK5-32: {{.+}} = load i32, i32* [[ADDR]],
252 
253 #endif
254 ///==========================================================================///
255 // RUN: %clang_cc1 -DCK6 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK6 --check-prefix CK6-64
256 // RUN: %clang_cc1 -DCK6 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
257 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK6 --check-prefix CK6-64
258 // RUN: %clang_cc1 -DCK6 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK6 --check-prefix CK6-32
259 // RUN: %clang_cc1 -DCK6 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
260 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK6 --check-prefix CK6-32
261 #ifdef CK6
262 // CK6-DAG: [[GBL:@Gi]] = global i32 0
263 // CK6-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
264 // Map types: OMP_MAP_BYCOPY = 128
265 // CK6-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
266 
267 // CK6-LABEL: implicit_maps_host_global
268 int Gi;
implicit_maps_host_global(int a)269 void implicit_maps_host_global (int a){
270   // CK6-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
271   // CK6-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
272   // CK6-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
273   // CK6-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
274   // CK6-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
275   // CK6-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
276   // CK6-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
277   // CK6-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
278   // CK6-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
279   // CK6-64-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]],
280   // CK6-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32*
281   // CK6-64-DAG: store i32 [[GBLVAL:%.+]], i32* [[CADDR]],
282   // CK6-64-DAG: [[GBLVAL]] = load i32, i32* [[GBL]],
283   // CK6-32-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[GBLVAL:%.+]],
284 
285   // CK6: call void [[KERNEL:@.+]](i[[sz]] [[VAL]])
286   #pragma omp target
287   {
288     ++Gi;
289   }
290 }
291 
292 // CK6: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]])
293 // CK6: [[ADDR:%.+]] = alloca i[[sz]],
294 // CK6: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]],
295 // CK6-64: [[CADDR:%.+]] = bitcast i64* [[ADDR]] to i32*
296 // CK6-64: {{.+}} = load i32, i32* [[CADDR]],
297 // CK6-32: {{.+}} = load i32, i32* [[ADDR]],
298 
299 #endif
300 ///==========================================================================///
301 // RUN: %clang_cc1 -DCK7 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK7 --check-prefix CK7-64
302 // RUN: %clang_cc1 -DCK7 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
303 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK7  --check-prefix CK7-64
304 // RUN: %clang_cc1 -DCK7 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK7  --check-prefix CK7-32
305 // RUN: %clang_cc1 -DCK7 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
306 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK7  --check-prefix CK7-32
307 #ifdef CK7
308 
309 // For a 32-bit targets, the value doesn't fit the size of the pointer,
310 // therefore it is passed by reference with a map 'to' specification.
311 
312 // CK7-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 8]
313 // Map types: OMP_MAP_BYCOPY = 128
314 // CK7-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
315 // Map types: OMP_MAP_TO = 1
316 // CK7-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 1]
317 
318 // CK7-LABEL: implicit_maps_double
implicit_maps_double(int a)319 void implicit_maps_double (int a){
320   double d = (double)a;
321 
322   // CK7-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
323   // CK7-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
324   // CK7-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
325   // CK7-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
326   // CK7-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
327 
328   // CK7-64-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
329   // CK7-64-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
330   // CK7-64-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
331   // CK7-64-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
332   // CK7-64-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]],
333   // CK7-64-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to double*
334   // CK7-64-64-DAG: store double {{.+}}, double* [[CADDR]],
335 
336   // CK7-32-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
337   // CK7-32-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
338   // CK7-32-DAG: [[VALBP]] = bitcast double* [[DECL:%.+]] to i8*
339   // CK7-32-DAG: [[VALP]] = bitcast double* [[DECL]] to i8*
340 
341   // CK7-64: call void [[KERNEL:@.+]](i[[sz]] [[VAL]])
342   // CK7-32: call void [[KERNEL:@.+]](double* [[DECL]])
343   #pragma omp target
344   {
345     d += 1.0;
346   }
347 }
348 
349 // CK7-64: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]])
350 // CK7-64: [[ADDR:%.+]] = alloca i[[sz]],
351 // CK7-64: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]],
352 // CK7-64: [[CADDR:%.+]] = bitcast i64* [[ADDR]] to double*
353 // CK7-64: {{.+}} = load double, double* [[CADDR]],
354 
355 // CK7-32: define internal void [[KERNEL]](double* {{.+}}[[ARG:%.+]])
356 // CK7-32: [[ADDR:%.+]] = alloca double*,
357 // CK7-32: store double* [[ARG]], double** [[ADDR]],
358 // CK7-32: [[REF:%.+]] = load double*, double** [[ADDR]],
359 // CK7-32: {{.+}} = load double, double* [[REF]],
360 
361 #endif
362 ///==========================================================================///
363 // RUN: %clang_cc1 -DCK8 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK8
364 // RUN: %clang_cc1 -DCK8 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
365 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK8
366 // RUN: %clang_cc1 -DCK8 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK8
367 // RUN: %clang_cc1 -DCK8 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
368 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK8
369 #ifdef CK8
370 
371 // CK8-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
372 // Map types: OMP_MAP_BYCOPY = 128
373 // CK8-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
374 
375 // CK8-LABEL: implicit_maps_float
implicit_maps_float(int a)376 void implicit_maps_float (int a){
377   float f = (float)a;
378 
379   // CK8-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
380   // CK8-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
381   // CK8-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
382   // CK8-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
383   // CK8-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
384   // CK8-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
385   // CK8-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
386   // CK8-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
387   // CK8-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
388   // CK8-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]],
389   // CK8-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to float*
390   // CK8-DAG: store float {{.+}}, float* [[CADDR]],
391 
392   // CK8: call void [[KERNEL:@.+]](i[[sz]] [[VAL]])
393   #pragma omp target
394   {
395     f += 1.0;
396   }
397 }
398 
399 // CK8: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]])
400 // CK8: [[ADDR:%.+]] = alloca i[[sz]],
401 // CK8: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]],
402 // CK8: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to float*
403 // CK8: {{.+}} = load float, float* [[CADDR]],
404 
405 #endif
406 ///==========================================================================///
407 // RUN: %clang_cc1 -DCK9 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK9
408 // RUN: %clang_cc1 -DCK9 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
409 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK9
410 // RUN: %clang_cc1 -DCK9 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK9
411 // RUN: %clang_cc1 -DCK9 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
412 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK9
413 #ifdef CK9
414 
415 // CK9-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 16]
416 // Map types: OMP_MAP_TO + OMP_MAP_FROM = 2 + 1
417 // CK9-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 3]
418 
419 // CK9-LABEL: implicit_maps_array
implicit_maps_array(int a)420 void implicit_maps_array (int a){
421   double darr[2] = {(double)a, (double)a};
422 
423   // CK9-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
424   // CK9-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
425   // CK9-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
426   // CK9-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
427   // CK9-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
428   // CK9-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
429   // CK9-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
430   // CK9-DAG: [[VALBP]] = bitcast [2 x double]* [[DECL:%.+]] to i8*
431   // CK9-DAG: [[VALP]] = bitcast [2 x double]* [[DECL]] to i8*
432 
433   // CK9: call void [[KERNEL:@.+]]([2 x double]* [[DECL]])
434   #pragma omp target
435   {
436     darr[0] += 1.0;
437     darr[1] += 1.0;
438   }
439 }
440 
441 // CK9: define internal void [[KERNEL]]([2 x double]* {{.+}}[[ARG:%.+]])
442 // CK9: [[ADDR:%.+]] = alloca [2 x double]*,
443 // CK9: store [2 x double]* [[ARG]], [2 x double]** [[ADDR]],
444 // CK9: [[REF:%.+]] = load [2 x double]*, [2 x double]** [[ADDR]],
445 // CK9: {{.+}} = getelementptr inbounds [2 x double], [2 x double]* [[REF]], i[[sz]] 0, i[[sz]] 0
446 #endif
447 ///==========================================================================///
448 // RUN: %clang_cc1 -DCK10 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK10
449 // RUN: %clang_cc1 -DCK10 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
450 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK10
451 // RUN: %clang_cc1 -DCK10 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK10
452 // RUN: %clang_cc1 -DCK10 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
453 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK10
454 #ifdef CK10
455 
456 // CK10-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}]
457 // Map types: OMP_MAP_BYCOPY | OMP_MAP_PTR = 128 + 32
458 // CK10-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 160]
459 
460 // CK10-LABEL: implicit_maps_pointer
implicit_maps_pointer()461 void implicit_maps_pointer (){
462   double *ddyn;
463 
464   // CK10-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
465   // CK10-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
466   // CK10-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
467   // CK10-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
468   // CK10-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
469   // CK10-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
470   // CK10-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
471   // CK10-DAG: [[VALBP]] = bitcast double* [[PTR:%.+]] to i8*
472   // CK10-DAG: [[VALP]] = bitcast double* [[PTR]] to i8*
473 
474   // CK10: call void [[KERNEL:@.+]](double* [[PTR]])
475   #pragma omp target
476   {
477     ddyn[0] += 1.0;
478     ddyn[1] += 1.0;
479   }
480 }
481 
482 // CK10: define internal void [[KERNEL]](double* {{.*}}[[ARG:%.+]])
483 // CK10: [[ADDR:%.+]] = alloca double*,
484 // CK10: store double* [[ARG]], double** [[ADDR]],
485 // CK10: [[REF:%.+]] = load double*, double** [[ADDR]],
486 // CK10: {{.+}} = getelementptr inbounds double, double* [[REF]], i[[sz]] 0
487 
488 #endif
489 ///==========================================================================///
490 // RUN: %clang_cc1 -DCK11 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK11
491 // RUN: %clang_cc1 -DCK11 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
492 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK11
493 // RUN: %clang_cc1 -DCK11 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK11
494 // RUN: %clang_cc1 -DCK11 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
495 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK11
496 #ifdef CK11
497 
498 // CK11-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 16]
499 // Map types: OMP_MAP_TO = 1
500 // CK11-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 1]
501 
502 // CK11-LABEL: implicit_maps_double_complex
implicit_maps_double_complex(int a)503 void implicit_maps_double_complex (int a){
504   double _Complex dc = (double)a;
505 
506   // CK11-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
507   // CK11-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
508   // CK11-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
509   // CK11-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
510   // CK11-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
511   // CK11-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
512   // CK11-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
513   // CK11-DAG: [[VALBP]] = bitcast { double, double }* [[PTR:%.+]] to i8*
514   // CK11-DAG: [[VALP]] = bitcast { double, double }* [[PTR]] to i8*
515 
516   // CK11: call void [[KERNEL:@.+]]({ double, double }* [[PTR]])
517   #pragma omp target
518   {
519    dc *= dc;
520   }
521 }
522 
523 // CK11: define internal void [[KERNEL]]({ double, double }* {{.*}}[[ARG:%.+]])
524 // CK11: [[ADDR:%.+]] = alloca { double, double }*,
525 // CK11: store { double, double }* [[ARG]], { double, double }** [[ADDR]],
526 // CK11: [[REF:%.+]] = load { double, double }*, { double, double }** [[ADDR]],
527 // CK11: {{.+}} = getelementptr inbounds { double, double }, { double, double }* [[REF]], i32 0, i32 0
528 #endif
529 ///==========================================================================///
530 // RUN: %clang_cc1 -DCK12 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK12 --check-prefix CK12-64
531 // RUN: %clang_cc1 -DCK12 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
532 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK12 --check-prefix CK12-64
533 // RUN: %clang_cc1 -DCK12 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK12 --check-prefix CK12-32
534 // RUN: %clang_cc1 -DCK12 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
535 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK12 --check-prefix CK12-32
536 #ifdef CK12
537 
538 // For a 32-bit targets, the value doesn't fit the size of the pointer,
539 // therefore it is passed by reference with a map 'to' specification.
540 
541 // CK12-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 8]
542 // Map types: OMP_MAP_BYCOPY = 128
543 // CK12-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
544 // Map types: OMP_MAP_TO = 1
545 // CK12-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 1]
546 
547 // CK12-LABEL: implicit_maps_float_complex
implicit_maps_float_complex(int a)548 void implicit_maps_float_complex (int a){
549   float _Complex fc = (float)a;
550 
551   // CK12-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
552   // CK12-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
553   // CK12-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
554   // CK12-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
555   // CK12-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
556 
557   // CK12-64-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
558   // CK12-64-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
559   // CK12-64-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
560   // CK12-64-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
561   // CK12-64-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]],
562   // CK12-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to { float, float }*
563   // CK12-64-DAG: store { float, float } {{.+}}, { float, float }* [[CADDR]],
564 
565   // CK12-32-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
566   // CK12-32-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
567   // CK12-32-DAG: [[VALBP]] = bitcast { float, float }* [[DECL:%.+]] to i8*
568   // CK12-32-DAG: [[VALP]] = bitcast { float, float }* [[DECL]] to i8*
569 
570   // CK12-64: call void [[KERNEL:@.+]](i[[sz]] [[VAL]])
571   // CK12-32: call void [[KERNEL:@.+]]({ float, float }* [[DECL]])
572   #pragma omp target
573   {
574     fc *= fc;
575   }
576 }
577 
578 // CK12-64: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]])
579 // CK12-64: [[ADDR:%.+]] = alloca i[[sz]],
580 // CK12-64: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]],
581 // CK12-64: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to { float, float }*
582 // CK12-64: {{.+}} = getelementptr inbounds { float, float }, { float, float }* [[CADDR]], i32 0, i32 0
583 
584 // CK12-32: define internal void [[KERNEL]]({ float, float }* {{.+}}[[ARG:%.+]])
585 // CK12-32: [[ADDR:%.+]] = alloca { float, float }*,
586 // CK12-32: store { float, float }* [[ARG]], { float, float }** [[ADDR]],
587 // CK12-32: [[REF:%.+]] = load { float, float }*, { float, float }** [[ADDR]],
588 // CK12-32: {{.+}} = getelementptr inbounds { float, float }, { float, float }* [[REF]], i32 0, i32 0
589 #endif
590 ///==========================================================================///
591 // RUN: %clang_cc1 -DCK13 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK13
592 // RUN: %clang_cc1 -DCK13 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
593 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK13
594 // RUN: %clang_cc1 -DCK13 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK13
595 // RUN: %clang_cc1 -DCK13 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
596 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK13
597 #ifdef CK13
598 
599 // We don't have a constant map size for VLAs.
600 // Map types:
601 //  - OMP_MAP_BYCOPY = 128 (vla size)
602 //  - OMP_MAP_BYCOPY = 128 (vla size)
603 //  - OMP_MAP_TO + OMP_MAP_FROM = 2 + 1
604 // CK13-DAG: [[TYPES:@.+]] = {{.+}}constant [3 x i32] [i32 128, i32 128, i32 3]
605 
606 // CK13-LABEL: implicit_maps_variable_length_array
implicit_maps_variable_length_array(int a)607 void implicit_maps_variable_length_array (int a){
608   double vla[2][a];
609 
610   // CK13-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 3, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], i[[sz:64|32]]* [[SGEP:%[^,]+]], {{.+}}[[TYPES]]{{.+}})
611   // CK13-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
612   // CK13-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
613   // CK13-DAG: [[SGEP]] = getelementptr inbounds {{.+}}[[SS:%[^,]+]], i32 0, i32 0
614 
615   // CK13-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
616   // CK13-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
617   // CK13-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[SS]], i32 0, i32 0
618   // CK13-DAG: store i8* inttoptr (i[[sz]] 2 to i8*), i8** [[BP0]],
619   // CK13-DAG: store i8* inttoptr (i[[sz]] 2 to i8*), i8** [[P0]],
620   // CK13-DAG: store i[[sz]] {{8|4}}, i[[sz]]* [[S0]],
621 
622   // CK13-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1
623   // CK13-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1
624   // CK13-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[SS]], i32 0, i32 1
625   // CK13-DAG: store i8* [[VALBP1:%.+]], i8** [[BP1]],
626   // CK13-DAG: store i8* [[VALP1:%.+]], i8** [[P1]],
627   // CK13-DAG: [[VALBP1]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
628   // CK13-DAG: [[VALP1]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
629   // CK13-DAG: store i[[sz]] {{8|4}}, i[[sz]]* [[S1]],
630 
631   // CK13-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 2
632   // CK13-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 2
633   // CK13-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[SS]], i32 0, i32 2
634   // CK13-DAG: store i8* [[VALBP2:%.+]], i8** [[BP2]],
635   // CK13-DAG: store i8* [[VALP2:%.+]], i8** [[P2]],
636   // CK13-DAG: store i[[sz]] [[VALS2:%.+]], i[[sz]]* [[S2]],
637   // CK13-DAG: [[VALBP2]] = bitcast double* [[DECL:%.+]] to i8*
638   // CK13-DAG: [[VALP2]] = bitcast double* [[DECL]] to i8*
639   // CK13-DAG: [[VALS2]] = mul nuw i[[sz]] %{{.+}}, 8
640 
641   // CK13: call void [[KERNEL:@.+]](i[[sz]] {{.+}}, i[[sz]] {{.+}}, double* [[DECL]])
642   #pragma omp target
643   {
644     vla[1][3] += 1.0;
645   }
646 }
647 
648 // CK13: define internal void [[KERNEL]](i[[sz]] [[VLA0:%.+]], i[[sz]] [[VLA1:%.+]], double* {{.+}}[[ARG:%.+]])
649 // CK13: [[ADDR0:%.+]] = alloca i[[sz]],
650 // CK13: [[ADDR1:%.+]] = alloca i[[sz]],
651 // CK13: [[ADDR2:%.+]] = alloca double*,
652 // CK13: store i[[sz]] [[VLA0]], i[[sz]]* [[ADDR0]],
653 // CK13: store i[[sz]] [[VLA1]], i[[sz]]* [[ADDR1]],
654 // CK13: store double* [[ARG]], double** [[ADDR2]],
655 // CK13: {{.+}} = load i[[sz]],  i[[sz]]* [[ADDR0]],
656 // CK13: {{.+}} = load i[[sz]],  i[[sz]]* [[ADDR1]],
657 // CK13: [[REF:%.+]] = load double*, double** [[ADDR2]],
658 // CK13: {{.+}} = getelementptr inbounds double, double* [[REF]], i[[sz]] %{{.+}}
659 #endif
660 ///==========================================================================///
661 // RUN: %clang_cc1 -DCK14 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK14 --check-prefix CK14-64
662 // RUN: %clang_cc1 -DCK14 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
663 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK14 --check-prefix CK14-64
664 // RUN: %clang_cc1 -DCK14 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK14 --check-prefix CK14-32
665 // RUN: %clang_cc1 -DCK14 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
666 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK14 --check-prefix CK14-32
667 #ifdef CK14
668 
669 // CK14-DAG: [[ST:%.+]] = type { i32, double }
670 // CK14-DAG: [[SIZES:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{16|12}}, i{{64|32}} 4]
671 // Map types:
672 // - OMP_MAP_TO | OMP_MAP_FROM = 1 + 2
673 // - OMP_MAP_BYCOPY = 128
674 // CK14-DAG: [[TYPES:@.+]] = {{.+}}constant [2 x i32] [i32 3, i32 128]
675 
676 class SSS {
677 public:
678   int a;
679   double b;
680 
foo(int c)681   void foo(int c) {
682     #pragma omp target
683     {
684       a += c;
685       b += (double)c;
686     }
687   }
688 
SSS(int a,double b)689   SSS(int a, double b) : a(a), b(b) {}
690 };
691 
692 // CK14-LABEL: implicit_maps_class
implicit_maps_class(int a)693 void implicit_maps_class (int a){
694   SSS sss(a, (double)a);
695 
696   // CK14: define {{.*}}void @{{.+}}foo{{.+}}([[ST]]* {{[^,]+}}, i32 {{[^,]+}})
697   // CK14-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
698   // CK14-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
699   // CK14-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
700 
701   // CK14-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
702   // CK14-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
703   // CK14-DAG: store i8* [[VALBP0:%.+]], i8** [[BP0]],
704   // CK14-DAG: store i8* [[VALP0:%.+]], i8** [[P0]],
705   // CK14-DAG: [[VALBP0]] = bitcast [[ST]]* [[DECL:%.+]] to i8*
706   // CK14-DAG: [[VALP0]] = bitcast [[ST]]* [[DECL]] to i8*
707 
708   // CK14-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1
709   // CK14-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1
710   // CK14-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
711   // CK14-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
712   // CK14-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
713   // CK14-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
714   // CK14-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]],
715   // CK14-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32*
716   // CK14-64-DAG: store i32 {{.+}}, i32* [[CADDR]],
717 
718   // CK14: call void [[KERNEL:@.+]]([[ST]]* [[DECL]], i[[sz]] {{.+}})
719   sss.foo(123);
720 }
721 
722 // CK14: define internal void [[KERNEL]]([[ST]]* [[THIS:%.+]], i[[sz]] [[ARG:%.+]])
723 // CK14: [[ADDR0:%.+]] = alloca [[ST]]*,
724 // CK14: [[ADDR1:%.+]] = alloca i[[sz]],
725 // CK14: store [[ST]]* [[THIS]], [[ST]]** [[ADDR0]],
726 // CK14: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR1]],
727 // CK14: [[REF0:%.+]] = load [[ST]]*, [[ST]]** [[ADDR0]],
728 // CK14-64: [[CADDR1:%.+]] = bitcast i[[sz]]* [[ADDR1]] to i32*
729 // CK14-64: {{.+}} = load i32,  i32* [[CADDR1]],
730 // CK14-32: {{.+}} = load i32, i32* [[ADDR1]],
731 // CK14: {{.+}} = getelementptr inbounds [[ST]], [[ST]]* [[REF0]], i32 0, i32 0
732 
733 #endif
734 ///==========================================================================///
735 // RUN: %clang_cc1 -DCK15 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK15 --check-prefix CK15-64
736 // RUN: %clang_cc1 -DCK15 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
737 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK15 --check-prefix CK15-64
738 // RUN: %clang_cc1 -DCK15 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK15 --check-prefix CK15-32
739 // RUN: %clang_cc1 -DCK15 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
740 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK15 --check-prefix CK15-32
741 #ifdef CK15
742 
743 // CK15: [[ST:%.+]] = type { i32, double, i32* }
744 // CK15: [[SIZES:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{24|16}}, i{{64|32}} 4]
745 // Map types:
746 // - OMP_MAP_TO | OMP_MAP_FROM = 1 + 2
747 // - OMP_MAP_BYCOPY = 128
748 // CK15: [[TYPES:@.+]] = {{.+}}constant [2 x i32] [i32 3, i32 128]
749 
750 // CK15: [[SIZES2:@.+]] = {{.+}}constant [2 x i[[sz]]] [i{{64|32}} {{24|16}}, i{{64|32}} 4]
751 // Map types:
752 // - OMP_MAP_TO | OMP_MAP_FROM = 1 + 2
753 // - OMP_MAP_BYCOPY = 128
754 // CK15: [[TYPES2:@.+]] = {{.+}}constant [2 x i32] [i32 3, i32 128]
755 
756 template<int x>
757 class SSST {
758 public:
759   int a;
760   double b;
761   int &r;
762 
foo(int c)763   void foo(int c) {
764     #pragma omp target
765     {
766       a += c + x;
767       b += (double)(c + x);
768       r += x;
769     }
770   }
771   template<int y>
bar(int c)772   void bar(int c) {
773     #pragma omp target
774     {
775       a += c + x + y;
776       b += (double)(c + x + y);
777       r += x + y;
778     }
779   }
780 
SSST(int a,double b,int & r)781   SSST(int a, double b, int &r) : a(a), b(b), r(r) {}
782 };
783 
784 // CK15-LABEL: implicit_maps_templated_class
implicit_maps_templated_class(int a)785 void implicit_maps_templated_class (int a){
786   SSST<123> ssst(a, (double)a, a);
787 
788   // CK15: define {{.*}}void @{{.+}}foo{{.+}}([[ST]]* {{[^,]+}}, i32 {{[^,]+}})
789   // CK15-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
790   // CK15-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
791   // CK15-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
792 
793   // CK15-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
794   // CK15-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
795   // CK15-DAG: store i8* [[VALBP0:%.+]], i8** [[BP0]],
796   // CK15-DAG: store i8* [[VALP0:%.+]], i8** [[P0]],
797   // CK15-DAG: [[VALBP0]] = bitcast [[ST]]* [[DECL:%.+]] to i8*
798   // CK15-DAG: [[VALP0]] = bitcast [[ST]]* [[DECL]] to i8*
799 
800   // CK15-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1
801   // CK15-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1
802   // CK15-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
803   // CK15-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
804   // CK15-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
805   // CK15-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
806   // CK15-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]],
807   // CK15-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32*
808   // CK15-64-DAG: store i32 {{.+}}, i32* [[CADDR]],
809 
810   // CK15: call void [[KERNEL:@.+]]([[ST]]* [[DECL]], i[[sz]] {{.+}})
811   ssst.foo(456);
812 
813   // CK15: define {{.*}}void @{{.+}}bar{{.+}}([[ST]]* {{[^,]+}}, i32 {{[^,]+}})
814   // CK15-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES2]]{{.+}}, {{.+}}[[TYPES2]]{{.+}})
815   // CK15-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
816   // CK15-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
817 
818   // CK15-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
819   // CK15-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
820   // CK15-DAG: store i8* [[VALBP0:%.+]], i8** [[BP0]],
821   // CK15-DAG: store i8* [[VALP0:%.+]], i8** [[P0]],
822   // CK15-DAG: [[VALBP0]] = bitcast [[ST]]* [[DECL:%.+]] to i8*
823   // CK15-DAG: [[VALP0]] = bitcast [[ST]]* [[DECL]] to i8*
824 
825   // CK15-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1
826   // CK15-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1
827   // CK15-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
828   // CK15-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
829   // CK15-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
830   // CK15-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
831   // CK15-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]],
832   // CK15-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32*
833   // CK15-64-DAG: store i32 {{.+}}, i32* [[CADDR]],
834 
835   // CK15: call void [[KERNEL2:@.+]]([[ST]]* [[DECL]], i[[sz]] {{.+}})
836   ssst.bar<210>(789);
837 }
838 
839 // CK15: define internal void [[KERNEL]]([[ST]]* [[THIS:%.+]], i[[sz]] [[ARG:%.+]])
840 // CK15: [[ADDR0:%.+]] = alloca [[ST]]*,
841 // CK15: [[ADDR1:%.+]] = alloca i[[sz]],
842 // CK15: store [[ST]]* [[THIS]], [[ST]]** [[ADDR0]],
843 // CK15: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR1]],
844 // CK15: [[REF0:%.+]] = load [[ST]]*, [[ST]]** [[ADDR0]],
845 // CK15-64: [[CADDR1:%.+]] = bitcast i[[sz]]* [[ADDR1]] to i32*
846 // CK15-64: {{.+}} = load i32,  i32* [[CADDR1]],
847 // CK15-32: {{.+}} = load i32, i32* [[ADDR1]],
848 // CK15: {{.+}} = getelementptr inbounds [[ST]], [[ST]]* [[REF0]], i32 0, i32 0
849 
850 // CK15: define internal void [[KERNEL2]]([[ST]]* [[THIS:%.+]], i[[sz]] [[ARG:%.+]])
851 // CK15: [[ADDR0:%.+]] = alloca [[ST]]*,
852 // CK15: [[ADDR1:%.+]] = alloca i[[sz]],
853 // CK15: store [[ST]]* [[THIS]], [[ST]]** [[ADDR0]],
854 // CK15: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR1]],
855 // CK15: [[REF0:%.+]] = load [[ST]]*, [[ST]]** [[ADDR0]],
856 // CK15-64: [[CADDR1:%.+]] = bitcast i[[sz]]* [[ADDR1]] to i32*
857 // CK15-64: {{.+}} = load i32,  i32* [[CADDR1]],
858 // CK15-32: {{.+}} = load i32, i32* [[ADDR1]],
859 // CK15: {{.+}} = getelementptr inbounds [[ST]], [[ST]]* [[REF0]], i32 0, i32 0
860 
861 #endif
862 ///==========================================================================///
863 // RUN: %clang_cc1 -DCK16 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK16 --check-prefix CK16-64
864 // RUN: %clang_cc1 -DCK16 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
865 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK16 --check-prefix CK16-64
866 // RUN: %clang_cc1 -DCK16 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK16 --check-prefix CK16-32
867 // RUN: %clang_cc1 -DCK16 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
868 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK16 --check-prefix CK16-32
869 #ifdef CK16
870 
871 // CK16-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
872 // Map types:
873 // - OMP_MAP_BYCOPY = 128
874 // CK16-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
875 
876 template<int y>
foo(int d)877 int foo(int d) {
878   int res = d;
879   #pragma omp target
880   {
881     res += y;
882   }
883   return res;
884 }
885 // CK16-LABEL: implicit_maps_templated_function
implicit_maps_templated_function(int a)886 void implicit_maps_templated_function (int a){
887   int i = a;
888 
889   // CK16: define {{.*}}i32 @{{.+}}foo{{.+}}(i32 {{[^,]+}})
890   // CK16-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
891   // CK16-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
892   // CK16-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
893 
894   // CK16-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
895   // CK16-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
896   // CK16-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
897   // CK16-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
898   // CK16-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
899   // CK16-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
900   // CK16-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]],
901   // CK16-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32*
902   // CK16-64-DAG: store i32 {{.+}}, i32* [[CADDR]],
903 
904   // CK16: call void [[KERNEL:@.+]](i[[sz]] [[VAL]])
905   i = foo<543>(i);
906 }
907 // CK16: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]])
908 // CK16: [[ADDR:%.+]] = alloca i[[sz]],
909 // CK16: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]],
910 // CK16-64: [[CADDR:%.+]] = bitcast i64* [[ADDR]] to i32*
911 // CK16-64: {{.+}} = load i32, i32* [[CADDR]],
912 // CK16-32: {{.+}} = load i32, i32* [[ADDR]],
913 
914 #endif
915 ///==========================================================================///
916 // RUN: %clang_cc1 -DCK17 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK17
917 // RUN: %clang_cc1 -DCK17 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
918 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK17
919 // RUN: %clang_cc1 -DCK17 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK17
920 // RUN: %clang_cc1 -DCK17 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
921 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK17
922 #ifdef CK17
923 
924 // CK17-DAG: [[ST:%.+]] = type { i32, double }
925 // CK17-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{16|12}}]
926 // Map types: OMP_MAP_TO + OMP_MAP_FROM = 2 + 1
927 // CK17-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 3]
928 
929 class SSS {
930 public:
931   int a;
932   double b;
933 };
934 
935 // CK17-LABEL: implicit_maps_struct
implicit_maps_struct(int a)936 void implicit_maps_struct (int a){
937   SSS s = {a, (double)a};
938 
939   // CK17-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
940   // CK17-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
941   // CK17-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
942   // CK17-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
943   // CK17-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
944   // CK17-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
945   // CK17-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
946   // CK17-DAG: [[VALBP]] = bitcast [[ST]]* [[DECL:%.+]] to i8*
947   // CK17-DAG: [[VALP]] = bitcast [[ST]]* [[DECL]] to i8*
948 
949   // CK17: call void [[KERNEL:@.+]]([[ST]]* [[DECL]])
950   #pragma omp target
951   {
952     s.a += 1;
953     s.b += 1.0;
954   }
955 }
956 
957 // CK17: define internal void [[KERNEL]]([[ST]]* {{.+}}[[ARG:%.+]])
958 // CK17: [[ADDR:%.+]] = alloca [[ST]]*,
959 // CK17: store [[ST]]* [[ARG]], [[ST]]** [[ADDR]],
960 // CK17: [[REF:%.+]] = load [[ST]]*, [[ST]]** [[ADDR]],
961 // CK17: {{.+}} = getelementptr inbounds [[ST]], [[ST]]* [[REF]], i32 0, i32 0
962 #endif
963 ///==========================================================================///
964 // RUN: %clang_cc1 -DCK18 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK18 --check-prefix CK18-64
965 // RUN: %clang_cc1 -DCK18 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
966 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK18 --check-prefix CK18-64
967 // RUN: %clang_cc1 -DCK18 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK18 --check-prefix CK18-32
968 // RUN: %clang_cc1 -DCK18 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
969 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK18 --check-prefix CK18-32
970 #ifdef CK18
971 
972 // CK18-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
973 // Map types:
974 // - OMP_MAP_BYCOPY = 128
975 // CK18-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
976 
977 template<typename T>
foo(T d)978 int foo(T d) {
979   #pragma omp target
980   {
981     d += (T)1;
982   }
983   return d;
984 }
985 // CK18-LABEL: implicit_maps_template_type_capture
implicit_maps_template_type_capture(int a)986 void implicit_maps_template_type_capture (int a){
987   int i = a;
988 
989   // CK18: define {{.*}}i32 @{{.+}}foo{{.+}}(i32 {{[^,]+}})
990   // CK18-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
991   // CK18-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
992   // CK18-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
993 
994   // CK18-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
995   // CK18-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
996   // CK18-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
997   // CK18-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
998   // CK18-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
999   // CK18-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
1000   // CK18-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]],
1001   // CK18-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32*
1002   // CK18-64-DAG: store i32 {{.+}}, i32* [[CADDR]],
1003 
1004   // CK18: call void [[KERNEL:@.+]](i[[sz]] [[VAL]])
1005   i = foo(i);
1006 }
1007 // CK18: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]])
1008 // CK18: [[ADDR:%.+]] = alloca i[[sz]],
1009 // CK18: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]],
1010 // CK18-64: [[CADDR:%.+]] = bitcast i64* [[ADDR]] to i32*
1011 // CK18-64: {{.+}} = load i32, i32* [[CADDR]],
1012 // CK18-32: {{.+}} = load i32, i32* [[ADDR]],
1013 
1014 #endif
1015 #endif
1016