1 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
2 
3 // expected-no-diagnostics
4 #ifndef HEADER
5 #define HEADER
6 
7 ///==========================================================================///
8 // RUN: %clang_cc1 -DCK0 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK0 --check-prefix CK0-64 %s
9 // RUN: %clang_cc1 -DCK0 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
10 // RUN: %clang_cc1 -DCK0 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK0 --check-prefix CK0-64 %s
11 // RUN: %clang_cc1 -DCK0 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK0 --check-prefix CK0-32 %s
12 // RUN: %clang_cc1 -DCK0 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
13 // RUN: %clang_cc1 -DCK0 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK0 --check-prefix CK0-32 %s
14 
15 // RUN: %clang_cc1 -DCK0 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
16 // RUN: %clang_cc1 -DCK0 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
17 // RUN: %clang_cc1 -DCK0 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
18 // RUN: %clang_cc1 -DCK0 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
19 // RUN: %clang_cc1 -DCK0 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
20 // RUN: %clang_cc1 -DCK0 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
21 
22 #ifdef CK0
23 // Mapper function code generation and runtime interface.
24 
25 // CK0: [[IDENT_T:%.+]] = type { i32, i32, i32, i32, i8* }
26 // CK0: [[ENTRY:%.+]] = type { i8*, i8*, i[[SZ:32|64]], i32, i32 }
27 // CK0: [[ANON_T:%.+]] = type { %class.C* }
28 // CK0: [[ANON_T_0:%.+]] = type { %class.C* }
29 // CK0: [[KMP_TASK_T_WITH_PRIVATES:%.+]] = type { [[KMP_TASK_T:%[^,]+]], [[KMP_PRIVATES_T:%.+]] }
30 // CK0: [[KMP_TASK_T]] = type { i8*, i32 (i32, i8*)*, i32, %{{[^,]+}}, %{{[^,]+}} }
31 // CK0-32: [[KMP_PRIVATES_T]] = type { [1 x i64], [1 x i8*], [1 x i8*], [1 x i8*] }
32 // CK0-64: [[KMP_PRIVATES_T]] = type { [1 x i8*], [1 x i8*], [1 x i64], [1 x i8*] }
33 // CK0: [[KMP_TASK_T_WITH_PRIVATES_1:%.+]] = type { [[KMP_TASK_T]], [[KMP_PRIVATES_T_2:%.+]] }
34 // CK0-32: [[KMP_PRIVATES_T_2]] = type { [1 x i64], [1 x i8*], [1 x i8*], [1 x i8*] }
35 // CK0-64: [[KMP_PRIVATES_T_2]] = type { [1 x i8*], [1 x i8*], [1 x i64], [1 x i8*] }
36 
37 // CK0-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}.region_id = weak constant i8 0
38 // CK0-64: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
39 // CK0-32: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
40 // CK0: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 35]
41 // CK0-64: [[NWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
42 // CK0-32: [[NWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
43 // CK0: [[NWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 35]
44 // CK0-64: [[TEAMSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
45 // CK0-32: [[TEAMSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
46 // CK0: [[TEAMTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33]
47 // CK0-64: [[TEAMNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
48 // CK0-32: [[TEAMNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
49 // CK0: [[TEAMNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33]
50 // CK0-64: [[EDSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
51 // CK0-32: [[EDSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
52 // CK0: [[EDTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33]
53 // CK0-64: [[EDNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
54 // CK0-32: [[EDNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
55 // CK0: [[EDNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33]
56 // CK0-64: [[EXDSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
57 // CK0-32: [[EXDSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
58 // CK0: [[EXDTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34]
59 // CK0-64: [[EXDNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
60 // CK0-32: [[EXDNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
61 // CK0: [[EXDNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34]
62 // CK0-64: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
63 // CK0-32: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
64 // CK0: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33]
65 // CK0-64: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
66 // CK0-32: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
67 // CK0: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34]
68 // CK0-64: [[FNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
69 // CK0-32: [[FNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
70 // CK0: [[FNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34]
71 
72 class C {
73 public:
74   int a;
75   double *b;
76 };
77 
78 #pragma omp declare mapper(id: C s) map(s.a, s.b[0:2])
79 
80 // CK0: define {{.*}}void [[MPRFUNC:@[.]omp_mapper[.].*C[.]id]](i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
81 // CK0: store i8* %{{[^,]+}}, i8** [[HANDLEADDR:%[^,]+]]
82 // CK0: store i8* %{{[^,]+}}, i8** [[BPTRADDR:%[^,]+]]
83 // CK0: store i8* %{{[^,]+}}, i8** [[VPTRADDR:%[^,]+]]
84 // CK0: store i64 %{{[^,]+}}, i{{64|32}}* [[SIZEADDR:%[^,]+]]
85 // CK0: store i64 %{{[^,]+}}, i64* [[TYPEADDR:%[^,]+]]
86 // CK0-DAG: [[BYTESIZE:%.+]] = load i64, i64* [[SIZEADDR]]
87 // CK0-64-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 16
88 // CK0-32-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 8
89 // CK0-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]]
90 // CK0-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]]
91 // CK0-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C**
92 // CK0-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i64 [[SIZE]]
93 // CK0-DAG: [[BPTR:%.+]] = load i8*, i8** [[BPTRADDR]]
94 // CK0-DAG: [[BEGIN:%.+]] = load i8*, i8** [[VPTRADDR]]
95 // CK0: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
96 // CK0: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
97 
98 // CK0: [[INITEVALDEL]]
99 // CK0: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
100 // CK0: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
101 // CK0: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]]
102 // CK0: [[INIT]]
103 // CK0-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
104 // CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
105 // CK0-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
106 // CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]])
107 // CK0: br label %[[LHEAD:[^,]+]]
108 
109 // CK0: [[LHEAD]]
110 // CK0: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]]
111 // CK0: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
112 // CK0: [[LBODY]]
113 // CK0: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
114 // CK0: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]]
115 // CK0-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0
116 // CK0-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1
117 // CK0-DAG: [[BBEGIN2:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1
118 // CK0-DAG: [[BARRBEGIN:%.+]] = load double*, double** [[BBEGIN2]]
119 // CK0-DAG: [[BARRBEGINGEP:%.+]] = getelementptr inbounds double, double* [[BARRBEGIN]], i[[sz:64|32]] 0
120 // CK0-DAG: [[BEND:%.+]] = getelementptr double*, double** [[BBEGIN]], i32 1
121 // CK0-DAG: [[ABEGINV:%.+]] = bitcast i32* [[ABEGIN]] to i8*
122 // CK0-DAG: [[BENDV:%.+]] = bitcast double** [[BEND]] to i8*
123 // CK0-DAG: [[ABEGINI:%.+]] = ptrtoint i8* [[ABEGINV]] to i64
124 // CK0-DAG: [[BENDI:%.+]] = ptrtoint i8* [[BENDV]] to i64
125 // CK0-DAG: [[CSIZE:%.+]] = sub i64 [[BENDI]], [[ABEGINI]]
126 // CK0-DAG: [[CUSIZE:%.+]] = sdiv exact i64 [[CSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
127 // CK0-DAG: [[BPTRADDR0BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
128 // CK0-DAG: [[PTRADDR0BC:%.+]] = bitcast i32* [[ABEGIN]] to i8*
129 // CK0-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(i8* [[HANDLE]])
130 // CK0-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48
131 // CK0-DAG: br label %[[MEMBER:[^,]+]]
132 // CK0-DAG: [[MEMBER]]
133 // CK0-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
134 // CK0-DAG: [[MEMBERCOM]]
135 // CK0-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]]
136 // CK0-DAG: br label %[[LTYPE]]
137 // CK0-DAG: [[LTYPE]]
138 // CK0-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
139 // CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
140 // CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
141 // CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
142 // CK0-DAG: [[ALLOC]]
143 // CK0-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
144 // CK0-DAG: br label %[[TYEND:[^,]+]]
145 // CK0-DAG: [[ALLOCELSE]]
146 // CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
147 // CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
148 // CK0-DAG: [[TO]]
149 // CK0-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
150 // CK0-DAG: br label %[[TYEND]]
151 // CK0-DAG: [[TOELSE]]
152 // CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
153 // CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
154 // CK0-DAG: [[FROM]]
155 // CK0-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
156 // CK0-DAG: br label %[[TYEND]]
157 // CK0-DAG: [[TYEND]]
158 // CK0-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
159 // CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR0BC]], i8* [[PTRADDR0BC]], i64 [[CUSIZE]], i64 [[PHITYPE0]])
160 // CK0-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
161 // CK0-DAG: [[PTRADDR1BC:%.+]] = bitcast i32* [[ABEGIN]] to i8*
162 // CK0-DAG: br label %[[MEMBER:[^,]+]]
163 // CK0-DAG: [[MEMBER]]
164 // CK0-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
165 // CK0-DAG: [[MEMBERCOM]]
166 // 281474976710659 == 0x1,000,000,003
167 // CK0-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]]
168 // CK0-DAG: br label %[[LTYPE]]
169 // CK0-DAG: [[LTYPE]]
170 // CK0-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710659, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
171 // CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
172 // CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
173 // CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
174 // CK0-DAG: [[ALLOC]]
175 // CK0-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
176 // CK0-DAG: br label %[[TYEND:[^,]+]]
177 // CK0-DAG: [[ALLOCELSE]]
178 // CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
179 // CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
180 // CK0-DAG: [[TO]]
181 // CK0-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
182 // CK0-DAG: br label %[[TYEND]]
183 // CK0-DAG: [[TOELSE]]
184 // CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
185 // CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
186 // CK0-DAG: [[FROM]]
187 // CK0-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
188 // CK0-DAG: br label %[[TYEND]]
189 // CK0-DAG: [[TYEND]]
190 // CK0-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
191 // CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 4, i64 [[TYPE1]])
192 // CK0-DAG: [[BPTRADDR2BC:%.+]] = bitcast double** [[BBEGIN]] to i8*
193 // CK0-DAG: [[PTRADDR2BC:%.+]] = bitcast double* [[BARRBEGINGEP]] to i8*
194 // CK0-DAG: br label %[[MEMBER:[^,]+]]
195 // CK0-DAG: [[MEMBER]]
196 // CK0-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
197 // CK0-DAG: [[MEMBERCOM]]
198 // 281474976710675 == 0x1,000,000,013
199 // CK0-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]]
200 // CK0-DAG: br label %[[LTYPE]]
201 // CK0-DAG: [[LTYPE]]
202 // CK0-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710675, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
203 // CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
204 // CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
205 // CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
206 // CK0-DAG: [[ALLOC]]
207 // CK0-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
208 // CK0-DAG: br label %[[TYEND:[^,]+]]
209 // CK0-DAG: [[ALLOCELSE]]
210 // CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
211 // CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
212 // CK0-DAG: [[TO]]
213 // CK0-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
214 // CK0-DAG: br label %[[TYEND]]
215 // CK0-DAG: [[TOELSE]]
216 // CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
217 // CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
218 // CK0-DAG: [[FROM]]
219 // CK0-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
220 // CK0-DAG: br label %[[TYEND]]
221 // CK0-DAG: [[TYEND]]
222 // CK0-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
223 // CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR2BC]], i8* [[PTRADDR2BC]], i64 16, i64 [[TYPE2]])
224 // CK0: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1
225 // CK0: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]]
226 // CK0: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
227 
228 // CK0: [[LEXIT]]
229 // CK0: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
230 // CK0: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
231 // CK0: [[EVALDEL]]
232 // CK0: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
233 // CK0: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
234 // CK0: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]]
235 // CK0: [[DEL]]
236 // CK0-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
237 // CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
238 // CK0-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
239 // CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]])
240 // CK0: br label %[[DONE]]
241 // CK0: [[DONE]]
242 // CK0: ret void
243 
244 
245 // CK0-LABEL: define {{.*}}void @{{.*}}foo{{.*}}
246 void foo(int a){
247   int i = a;
248   C c;
249   c.a = a;
250 
251   // CK0-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
252   // CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
253   // CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
254   // CK0-DAG: [[MPRGEP]] = bitcast [1 x i8*]* [[MPR:%[^,]+]] to i8**
255   // CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
256   // CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
257   // CK0-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i[[sz]] 0, i[[sz]] 0
258   // CK0-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C**
259   // CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
260   // CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
261   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
262   // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
263   // CK0: call void [[KERNEL_1:@.+]](%class.C* [[VAL]])
264   #pragma omp target map(mapper(id),tofrom: c)
265   {
266     ++c.a;
267   }
268 
269   // CK0: [[BP2GEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_BP2:%[^,]+]], i32 0, i32 0
270   // CK0: [[BP2CAST:%.+]] = bitcast i8** [[BP2GEP]] to %class.C**
271   // CK0: store %class.C* [[CADDR:%[^,]+]], %class.C** [[BP2CAST]], align
272   // CK0: [[P2GEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_P2:%[^,]+]], i32 0, i32 0
273   // CK0: [[P2CAST:%.+]] = bitcast i8** [[P2GEP]] to %class.C**
274   // CK0: store %class.C* [[CADDR]], %class.C** [[P2CAST]], align
275   // CK0: [[MAPPER2GEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_MAPPER2:%[^,]+]], i[[SZ]] 0, i[[SZ]] 0
276   // CK0: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MAPPER2GEP]], align
277   // CK0: [[BP2:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_BP2]], i32 0, i32 0
278   // CK0: [[P2:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_P2]], i32 0, i32 0
279   // CK0: [[MAPPER2:%.+]] = bitcast [1 x i8*]* [[OFFLOAD_MAPPER2]] to i8**
280   // CK0-32: [[TASK:%.+]] = call i8* @__kmpc_omp_target_task_alloc([[IDENT_T]]* {{@.+}}, i32 {{%.+}}, i32 1, i32 40, i32 4, i32 (i32, i8*)* bitcast (i32 (i32, [[KMP_TASK_T_WITH_PRIVATES]]*)* [[TASK_ENTRY:@.+]] to i32 (i32, i8*)*), i64 -1)
281   // CK0-64: [[TASK:%.+]] = call i8* @__kmpc_omp_target_task_alloc([[IDENT_T]]* {{@.+}}, i32 {{%.+}}, i32 1, i64 72, i64 8, i32 (i32, i8*)* bitcast (i32 (i32, [[KMP_TASK_T_WITH_PRIVATES]]*)* [[TASK_ENTRY:@.+]] to i32 (i32, i8*)*), i64 -1)
282   // CK0: [[TASK_CAST:%.+]] = bitcast i8* [[TASK]] to [[KMP_TASK_T_WITH_PRIVATES]]*
283   // CK0: [[TASK_WITH_PRIVATES:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES]], [[KMP_TASK_T_WITH_PRIVATES]]* [[TASK_CAST]], i32 0, i32 1
284   // CK0: {{.+}} = call i32 @__kmpc_omp_task([[IDENT_T]]* @1, i32 {{.+}}, i8* [[TASK]])
285   #pragma omp target map(mapper(id),tofrom: c) nowait
286   {
287     ++c.a;
288   }
289 
290   // CK0-DAG: call i32 @__tgt_target_teams_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[TEAMSIZES]]{{.+}}, {{.+}}[[TEAMTYPES]]{{.+}}, i8** [[MPRGEP:%.+]], i32 0, i32 0)
291   // CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
292   // CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
293   // CK0-DAG: [[MPRGEP]] = bitcast [1 x i8*]* [[MPR:%[^,]+]] to i8**
294   // CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
295   // CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
296   // CK0-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i[[sz]] 0, i[[sz]] 0
297   // CK0-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C**
298   // CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
299   // CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
300   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
301   // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
302   // CK0: call void [[KERNEL_3:@.+]](%class.C* [[VAL]])
303   #pragma omp target teams map(mapper(id),to: c)
304   {
305     ++c.a;
306   }
307 
308   // CK0-32: [[TASK_1:%.+]] = call i8* @__kmpc_omp_target_task_alloc([[IDENT_T]]* {{@.+}}, i32 {{%.+}}, i32 1, i32 40, i32 4, i32 (i32, i8*)* bitcast (i32 (i32, [[KMP_TASK_T_WITH_PRIVATES_1]]*)* [[TASK_ENTRY_1:@.+]] to i32 (i32, i8*)*), i64 -1)
309   // CK0-64: [[TASK_1:%.+]] = call i8* @__kmpc_omp_target_task_alloc([[IDENT_T]]* {{@.+}}, i32 {{%.+}}, i32 1, i64 72, i64 8, i32 (i32, i8*)* bitcast (i32 (i32, [[KMP_TASK_T_WITH_PRIVATES_1]]*)* [[TASK_ENTRY_1:@.+]] to i32 (i32, i8*)*), i64 -1)
310   // CK0: [[TASK_CAST_1:%.+]] = bitcast i8* [[TASK_1]] to [[KMP_TASK_T_WITH_PRIVATES_1]]*
311   // CK0: [[TASK_CAST_GET_1:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES_1]], [[KMP_TASK_T_WITH_PRIVATES_1]]* [[TASK_CAST_1]], i32 0, i32 0
312   // CK0: {{.+}} = getelementptr inbounds [[KMP_TASK_T]], [[KMP_TASK_T]]* [[TASK_CAST_GET_1]], i32 0, i32 0
313   // CK0: {{.+}} = call i32 @__kmpc_omp_task([[IDENT_T]]* @1, i32 {{.+}}, i8* [[TASK_1]])
314   #pragma omp target teams map(mapper(id),to: c) nowait
315   {
316     ++c.a;
317   }
318 
319   // CK0-DAG: call void @__tgt_target_data_begin_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EDSIZES]]{{.+}}, {{.+}}[[EDTYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
320   // CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
321   // CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
322   // CK0-DAG: [[MPRGEP]] = bitcast [1 x i8*]* [[MPR:%[^,]+]] to i8**
323   // CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
324   // CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
325   // CK0-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i[[sz]] 0, i[[sz]] 0
326   // CK0-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C**
327   // CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
328   // CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
329   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
330   // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
331   #pragma omp target enter data map(mapper(id),to: c)
332 
333   // CK0-DAG: call void @__tgt_target_data_begin_nowait_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EDNWSIZES]]{{.+}}, {{.+}}[[EDNWTYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
334   // CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
335   // CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
336   // CK0-DAG: [[MPRGEP]] = bitcast [1 x i8*]* [[MPR:%[^,]+]] to i8**
337   // CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
338   // CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
339   // CK0-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i[[sz]] 0, i[[sz]] 0
340   // CK0-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C**
341   // CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
342   // CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
343   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
344   // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
345   #pragma omp target enter data map(mapper(id),to: c) nowait
346 
347   // CK0-DAG: call void @__tgt_target_data_end_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EXDSIZES]]{{.+}}, {{.+}}[[EXDTYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
348   // CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
349   // CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
350   // CK0-DAG: [[MPRGEP]] = bitcast [1 x i8*]* [[MPR:%[^,]+]] to i8**
351   // CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
352   // CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
353   // CK0-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i[[sz]] 0, i[[sz]] 0
354   // CK0-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C**
355   // CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
356   // CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
357   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
358   // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
359   #pragma omp target exit data map(mapper(id),from: c)
360 
361   // CK0-DAG: call void @__tgt_target_data_end_nowait_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EXDNWSIZES]]{{.+}}, {{.+}}[[EXDNWTYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
362   // CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
363   // CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
364   // CK0-DAG: [[MPRGEP]] = bitcast [1 x i8*]* [[MPR:%[^,]+]] to i8**
365   // CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
366   // CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
367   // CK0-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i[[sz]] 0, i[[sz]] 0
368   // CK0-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C**
369   // CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
370   // CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
371   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
372   // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
373   #pragma omp target exit data map(mapper(id),from: c) nowait
374 
375   // CK0-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[TGEPBP:%.+]], i8** [[TGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[TSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[TTYPES]]{{.+}}, i8** [[TMPRGEP:%.+]])
376   // CK0-DAG: [[TGEPBP]] = getelementptr inbounds {{.+}}[[TBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
377   // CK0-DAG: [[TGEPP]] = getelementptr inbounds {{.+}}[[TP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
378   // CK0-DAG: [[TMPRGEP]] = bitcast [1 x i8*]* [[TMPR:%[^,]+]] to i8**
379   // CK0-DAG: [[TBP0:%.+]] = getelementptr inbounds {{.+}}[[TBP]], i{{.+}} 0, i{{.+}} 0
380   // CK0-DAG: [[TP0:%.+]] = getelementptr inbounds {{.+}}[[TP]], i{{.+}} 0, i{{.+}} 0
381   // CK0-DAG: [[TMPR1:%.+]] = getelementptr inbounds {{.+}}[[TMPR]], i[[sz]] 0, i[[sz]] 0
382   // CK0-DAG: [[TCBP0:%.+]] = bitcast i8** [[TBP0]] to %class.C**
383   // CK0-DAG: [[TCP0:%.+]] = bitcast i8** [[TP0]] to %class.C**
384   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[TCBP0]]
385   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[TCP0]]
386   // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[TMPR1]]
387   #pragma omp target update to(mapper(id): c)
388 
389   // CK0-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[FSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[FTYPES]]{{.+}}, i8** [[FMPRGEP:%.+]])
390   // CK0-DAG: [[FGEPBP]] = getelementptr inbounds {{.+}}[[FBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
391   // CK0-DAG: [[FGEPP]] = getelementptr inbounds {{.+}}[[FP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
392   // CK0-DAG: [[FMPRGEP]] = bitcast [1 x i8*]* [[FMPR:%[^,]+]] to i8**
393   // CK0-DAG: [[FBP0:%.+]] = getelementptr inbounds {{.+}}[[FBP]], i{{.+}} 0, i{{.+}} 0
394   // CK0-DAG: [[FP0:%.+]] = getelementptr inbounds {{.+}}[[FP]], i{{.+}} 0, i{{.+}} 0
395   // CK0-DAG: [[FMPR1:%.+]] = getelementptr inbounds {{.+}}[[FMPR]], i[[sz]] 0, i[[sz]] 0
396   // CK0-DAG: [[FCBP0:%.+]] = bitcast i8** [[FBP0]] to %class.C**
397   // CK0-DAG: [[FCP0:%.+]] = bitcast i8** [[FP0]] to %class.C**
398   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[FCBP0]]
399   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[FCP0]]
400   // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[FMPR1]]
401   #pragma omp target update from(mapper(id): c)
402 
403   // CK0-DAG: call void @__tgt_target_data_update_nowait_mapper(i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[FNWSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[FNWTYPES]]{{.+}}, i8** [[FMPRGEP:%.+]])
404   // CK0-DAG: [[FGEPBP]] = getelementptr inbounds {{.+}}[[FBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
405   // CK0-DAG: [[FGEPP]] = getelementptr inbounds {{.+}}[[FP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
406   // CK0-DAG: [[FMPRGEP]] = bitcast [1 x i8*]* [[FMPR:%[^,]+]] to i8**
407   // CK0-DAG: [[FBP0:%.+]] = getelementptr inbounds {{.+}}[[FBP]], i{{.+}} 0, i{{.+}} 0
408   // CK0-DAG: [[FP0:%.+]] = getelementptr inbounds {{.+}}[[FP]], i{{.+}} 0, i{{.+}} 0
409   // CK0-DAG: [[FMPR1:%.+]] = getelementptr inbounds {{.+}}[[FMPR]], i[[sz]] 0, i[[sz]] 0
410   // CK0-DAG: [[FCBP0:%.+]] = bitcast i8** [[FBP0]] to %class.C**
411   // CK0-DAG: [[FCP0:%.+]] = bitcast i8** [[FP0]] to %class.C**
412   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[FCBP0]]
413   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[FCP0]]
414   // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[FMPR1]]
415   #pragma omp target update from(mapper(id): c) nowait
416 }
417 
418 
419 // CK0: define internal void [[KERNEL_1]](%class.C* {{.+}}[[ARG:%.+]])
420 // CK0: [[ADDR:%.+]] = alloca %class.C*,
421 // CK0: store %class.C* [[ARG]], %class.C** [[ADDR]]
422 // CK0: [[CADDR:%.+]] = load %class.C*, %class.C** [[ADDR]]
423 // CK0: [[CAADDR:%.+]] = getelementptr inbounds %class.C, %class.C* [[CADDR]], i32 0, i32 0
424 // CK0: [[VAL:%[^,]+]] = load i32, i32* [[CAADDR]]
425 // CK0: {{.+}} = add nsw i32 [[VAL]], 1
426 // CK0: }
427 
428 // CK0: define internal void [[KERNEL_2:@.+]](%class.C* {{.+}}[[ARG:%.+]])
429 // CK0: [[ADDR:%.+]] = alloca %class.C*,
430 // CK0: store %class.C* [[ARG]], %class.C** [[ADDR]]
431 // CK0: [[CADDR:%.+]] = load %class.C*, %class.C** [[ADDR]]
432 // CK0: [[CAADDR:%.+]] = getelementptr inbounds %class.C, %class.C* [[CADDR]], i32 0, i32 0
433 // CK0: [[VAL:%[^,]+]] = load i32, i32* [[CAADDR]]
434 // CK0: {{.+}} = add nsw i32 [[VAL]], 1
435 // CK0: }
436 
437 // CK0: define internal void [[OUTLINED:@.+]](i32 {{.*}}{{[^,]+}}, [[ANON_T]]* noalias [[CTXARG:%.+]])
438 // CK0-DAG: call i32 @__tgt_target_nowait_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZEGEP:%[0-9]+]], {{.+}}[[NWTYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
439 // CK0-DAG: [[BPGEP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BPFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0
440 // CK0-DAG: [[PGEP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[PFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0
441 // CK0-DAG: [[SIZEGEP]] = getelementptr inbounds [1 x i64], [1 x i64]* [[SIZEFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0
442 // CK0-DAG: [[MPRGEP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[MPRFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0
443 // CK0-DAG: [[BPFPADDR]] = load [1 x i8*]*, [1 x i8*]** [[FPPTRADDR_BP:%.+]], align
444 // CK0-DAG: [[PFPADDR]] = load [1 x i8*]*, [1 x i8*]** [[FPPTRADDR_P:%.+]], align
445 // CK0-DAG: [[SIZEFPADDR]] = load [1 x i64]*, [1 x i64]** [[FPPTRADDR_SIZE:%.+]], align
446 // CK0-DAG: [[MPRFPADDR]] = load [1 x i8*]*, [1 x i8*]** [[FPPTRADDR_MPR:%.+]], align
447 // CK0-DAG: call void (i8*, ...) %1(i8* %2, {{.+}}[[FPPTRADDR_BP]], {{.+}}[[FPPTRADDR_P]], {{.+}}[[FPPTRADDR_SIZE]], {{.+}}[[FPPTRADDR_MPR]])
448 // CK0-DAG: call void [[KERNEL_2:@.+]](%class.C* [[KERNELARG:%.+]])
449 // CK0-DAG: [[KERNELARG]] = load %class.C*, %class.C** [[KERNELARGGEP:%.+]], align
450 // CK0-DAG: [[KERNELARGGEP]] = getelementptr inbounds [[ANON_T]], [[ANON_T]]* [[CTX:%.+]], i32 0, i32 0
451 // CK0-DAG: [[CTX]] = load [[ANON_T]]*, [[ANON_T]]** [[CTXADDR:%.+]], align
452 // CK0-DAG: store [[ANON_T]]* [[CTXARG]], [[ANON_T]]** [[CTXADDR]], align
453 // CK0: }
454 
455 // CK0: define internal {{.*}}i32 [[TASK_ENTRY]](i32 {{.*}}%0, [[KMP_TASK_T_WITH_PRIVATES]]* noalias %1)
456 // CK0: store [[KMP_TASK_T_WITH_PRIVATES]]* %1, [[KMP_TASK_T_WITH_PRIVATES]]** [[ADDR:%.+]], align
457 // CK0: [[TASK_T_WITH_PRIVATES:%.+]] = load [[KMP_TASK_T_WITH_PRIVATES]]*, [[KMP_TASK_T_WITH_PRIVATES]]** [[ADDR]], align
458 // CK0: [[TASKGEP:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES]], [[KMP_TASK_T_WITH_PRIVATES]]* [[TASK_T_WITH_PRIVATES]], i32 0, i32 0
459 // CK0: [[SHAREDSGEP:%.+]] = getelementptr inbounds [[KMP_TASK_T]], [[KMP_TASK_T]]* [[TASKGEP]], i32 0, i32 0
460 // CK0: [[SHAREDS:%.+]] = load i8*, i8** [[SHAREDSGEP]], align
461 // CK0: [[ANON:%.+]] = bitcast i8* [[SHAREDS]] to [[ANON_T]]*
462 // CK0: [[PRIVATESGEP:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES]], [[KMP_TASK_T_WITH_PRIVATES]]* [[TASK_T_WITH_PRIVATES]], i32 0, i32 1
463 // CK0: [[PRIVATES:%.+]] = bitcast [[KMP_PRIVATES_T]]* [[PRIVATESGEP]] to i8*
464 // CK0: [[TASK_WITH_PRIVATES:%.+]] = bitcast [[KMP_TASK_T_WITH_PRIVATES]]* [[TASK_T_WITH_PRIVATES]] to i8*
465 // CK0: call void [[OUTLINED]](i32 {{%.+}}, i32* {{%.+}}, i8* [[PRIVATES]], {{.+}}, i8* [[TASK_WITH_PRIVATES]], [[ANON_T]]* [[ANON]])
466 // CK0: }
467 
468 // CK0: define internal void [[OUTLINE_1:@.+]](i32 {{.*}}%.global_tid.{{.+}}, [[ANON_T_0]]* noalias [[CTXARG:%.+]])
469 // CK0-DAG: call i32 @__tgt_target_teams_nowait_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], i64* [[SIZEGEP:%[0-9]+]], {{.+}}[[TEAMNWTYPES]]{{.+}}, i8** [[MPRGEP:%.+]], i32 0, i32 0)
470 // CK0-DAG: [[BPGEP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BPFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0
471 // CK0-DAG: [[PGEP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[PFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0
472 // CK0-DAG: [[SIZEGEP]] = getelementptr inbounds [1 x i64], [1 x i64]* [[SIZEFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0
473 // CK0-DAG: [[MPRGEP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[MPRFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0
474 // CK0-DAG: [[BPFPADDR]] = load [1 x i8*]*, [1 x i8*]** [[FPPTRADDR_BP:%.+]], align
475 // CK0-DAG: [[PFPADDR]] = load [1 x i8*]*, [1 x i8*]** [[FPPTRADDR_P:%.+]], align
476 // CK0-DAG: [[SIZEFPADDR]] = load [1 x i64]*, [1 x i64]** [[FPPTRADDR_SIZE:%.+]], align
477 // CK0-DAG: [[MPRFPADDR]] = load [1 x i8*]*, [1 x i8*]** [[FPPTRADDR_MPR:%.+]], align
478 // CK0-DAG: call void (i8*, ...) %1(i8* %2, {{.+}}[[FPPTRADDR_BP]], {{.+}}[[FPPTRADDR_P]], {{.+}}[[FPPTRADDR_SIZE]], {{.+}}[[FPPTRADDR_MPR]])
479 // CK0-DAG: call void [[KERNEL_2:@.+]](%class.C* [[KERNELARG:%.+]])
480 // CK0-DAG: [[KERNELARG]] = load %class.C*, %class.C** [[KERNELARGGEP:%.+]], align
481 // CK0-DAG: [[KERNELARGGEP]] = getelementptr inbounds [[ANON_T_0]], [[ANON_T_0]]* [[CTX:%.+]], i32 0, i32 0
482 // CK0-DAG: [[CTX]] = load [[ANON_T_0]]*, [[ANON_T_0]]** [[CTXADDR:%.+]], align
483 // CK0-DAG: store [[ANON_T_0]]* [[CTXARG]], [[ANON_T_0]]** [[CTXADDR]], align
484 // CK0: }
485 
486 // CK0: define internal {{.*}}i32 [[TASK_ENTRY_1]](i32 {{.*}}%0, [[KMP_TASK_T_WITH_PRIVATES_1]]* noalias %1)
487 // CK0: store [[KMP_TASK_T_WITH_PRIVATES_1]]* %1, [[KMP_TASK_T_WITH_PRIVATES_1]]** [[ADDR:%.+]], align
488 // CK0: [[TASK_T_WITH_PRIVATES:%.+]] = load [[KMP_TASK_T_WITH_PRIVATES_1]]*, [[KMP_TASK_T_WITH_PRIVATES_1]]** [[ADDR]], align
489 // CK0: [[TASKGEP:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES_1]], [[KMP_TASK_T_WITH_PRIVATES_1]]* [[TASK_T_WITH_PRIVATES]], i32 0, i32 0
490 // CK0: [[SHAREDSGEP:%.+]] = getelementptr inbounds [[KMP_TASK_T]], [[KMP_TASK_T]]* [[TASKGEP]], i32 0, i32 0
491 // CK0: [[SHAREDS:%.+]] = load i8*, i8** [[SHAREDSGEP]], align
492 // CK0: [[ANON:%.+]] = bitcast i8* [[SHAREDS]] to [[ANON_T_0]]*
493 // CK0: [[PRIVATESGEP:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES_1]], [[KMP_TASK_T_WITH_PRIVATES_1]]* [[TASK_T_WITH_PRIVATES]], i32 0, i32 1
494 // CK0: [[PRIVATES:%.+]] = bitcast [[KMP_PRIVATES_T_2]]* [[PRIVATESGEP]] to i8*
495 // CK0: [[TASK_WITH_PRIVATES:%.+]] = bitcast [[KMP_TASK_T_WITH_PRIVATES_1]]* [[TASK_T_WITH_PRIVATES]] to i8*
496 // CK0: call void [[OUTLINE_1]](i32 {{%.+}}, i32* {{%.+}}, i8* [[PRIVATES]], {{.+}}, i8* [[TASK_WITH_PRIVATES]], [[ANON_T_0]]* [[ANON]])
497 // CK0: }
498 
499 #endif // CK0
500 
501 
502 ///==========================================================================///
503 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK1 --check-prefix CK1-64 %s
504 // RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
505 // RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK1 --check-prefix CK1-64 %s
506 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK1 --check-prefix CK1-32 %s
507 // RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
508 // RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK1 --check-prefix CK1-32 %s
509 
510 // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
511 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
512 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
513 // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
514 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
515 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
516 
517 #ifdef CK1
518 // C++ template
519 
520 template <class T>
521 class C {
522 public:
523   T a;
524 };
525 
526 #pragma omp declare mapper(id: C<int> s) map(s.a)
527 
528 // CK1-LABEL: define {{.*}}void @.omp_mapper.{{.*}}C{{.*}}.id{{.*}}(i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
529 // CK1: store i8* %{{[^,]+}}, i8** [[HANDLEADDR:%[^,]+]]
530 // CK1: store i8* %{{[^,]+}}, i8** [[BPTRADDR:%[^,]+]]
531 // CK1: store i8* %{{[^,]+}}, i8** [[VPTRADDR:%[^,]+]]
532 // CK1: store i64 %{{[^,]+}}, i{{64|32}}* [[SIZEADDR:%[^,]+]]
533 // CK1: store i64 %{{[^,]+}}, i64* [[TYPEADDR:%[^,]+]]
534 // CK1-DAG: [[BYTESIZE:%.+]] = load i64, i64* [[SIZEADDR]]
535 // CK1-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 4
536 // CK1-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]]
537 // CK1-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]]
538 // CK1-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C**
539 // CK1-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i64 [[SIZE]]
540 // CK1-DAG: [[BPTR:%.+]] = load i8*, i8** [[BPTRADDR]]
541 // CK1-DAG: [[BEGIN:%.+]] = load i8*, i8** [[VPTRADDR]]
542 // CK1: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
543 // CK1: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
544 
545 // CK1: [[INITEVALDEL]]
546 // CK1: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
547 // CK1: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
548 // CK1: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]]
549 // CK1: [[INIT]]
550 // CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4
551 // CK1-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
552 // CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]])
553 // CK1: br label %[[LHEAD:[^,]+]]
554 
555 // CK1: [[LHEAD]]
556 // CK1: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]]
557 // CK1: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
558 // CK1: [[LBODY]]
559 // CK1: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
560 // CK1: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]]
561 // CK1-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0
562 // CK1-DAG: [[AEND:%.+]] = getelementptr i32, i32* [[ABEGIN]], i32 1
563 // CK1-DAG: [[ABEGINV:%.+]] = bitcast i32* [[ABEGIN]] to i8*
564 // CK1-DAG: [[AENDV:%.+]] = bitcast i32* [[AEND]] to i8*
565 // CK1-DAG: [[ABEGINI:%.+]] = ptrtoint i8* [[ABEGINV]] to i64
566 // CK1-DAG: [[AENDI:%.+]] = ptrtoint i8* [[AENDV]] to i64
567 // CK1-DAG: [[CSIZE:%.+]] = sub i64 [[AENDI]], [[ABEGINI]]
568 // CK1-DAG: [[CUSIZE:%.+]] = sdiv exact i64 [[CSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
569 // CK1-DAG: [[BPTRADDR0BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
570 // CK1-DAG: [[PTRADDR0BC:%.+]] = bitcast i32* [[ABEGIN]] to i8*
571 // CK1-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(i8* [[HANDLE]])
572 // CK1-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48
573 // CK1-DAG: br label %[[MEMBER:[^,]+]]
574 // CK1-DAG: [[MEMBER]]
575 // CK1-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
576 // CK1-DAG: [[MEMBERCOM]]
577 // CK1-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]]
578 // CK1-DAG: br label %[[LTYPE]]
579 // CK1-DAG: [[LTYPE]]
580 // CK1-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
581 // CK1-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
582 // CK1-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
583 // CK1-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
584 // CK1-DAG: [[ALLOC]]
585 // CK1-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
586 // CK1-DAG: br label %[[TYEND:[^,]+]]
587 // CK1-DAG: [[ALLOCELSE]]
588 // CK1-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
589 // CK1-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
590 // CK1-DAG: [[TO]]
591 // CK1-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
592 // CK1-DAG: br label %[[TYEND]]
593 // CK1-DAG: [[TOELSE]]
594 // CK1-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
595 // CK1-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
596 // CK1-DAG: [[FROM]]
597 // CK1-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
598 // CK1-DAG: br label %[[TYEND]]
599 // CK1-DAG: [[TYEND]]
600 // CK1-DAG: [[TYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
601 // CK1-64: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR0BC]], i8* [[PTRADDR0BC]], i64 [[CUSIZE]], i64 [[TYPE0]])
602 // CK1-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
603 // CK1-DAG: [[PTRADDR1BC:%.+]] = bitcast i32* [[ABEGIN]] to i8*
604 // CK1-DAG: br label %[[MEMBER:[^,]+]]
605 // CK1-DAG: [[MEMBER]]
606 // CK1-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
607 // CK1-DAG: [[MEMBERCOM]]
608 // 281474976710659 == 0x1,000,000,003
609 // CK1-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]]
610 // CK1-DAG: br label %[[LTYPE]]
611 // CK1-DAG: [[LTYPE]]
612 // CK1-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710659, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
613 // CK1-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
614 // CK1-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
615 // CK1-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
616 // CK1-DAG: [[ALLOC]]
617 // CK1-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
618 // CK1-DAG: br label %[[TYEND:[^,]+]]
619 // CK1-DAG: [[ALLOCELSE]]
620 // CK1-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
621 // CK1-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
622 // CK1-DAG: [[TO]]
623 // CK1-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
624 // CK1-DAG: br label %[[TYEND]]
625 // CK1-DAG: [[TOELSE]]
626 // CK1-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
627 // CK1-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
628 // CK1-DAG: [[FROM]]
629 // CK1-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
630 // CK1-DAG: br label %[[TYEND]]
631 // CK1-DAG: [[TYEND]]
632 // CK1-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
633 // CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 4, i64 [[TYPE1]])
634 // CK1: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1
635 // CK1: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]]
636 // CK1: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
637 
638 // CK1: [[LEXIT]]
639 // CK1: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
640 // CK1: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
641 // CK1: [[EVALDEL]]
642 // CK1: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
643 // CK1: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
644 // CK1: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]]
645 // CK1: [[DEL]]
646 // CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4
647 // CK1-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
648 // CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]])
649 // CK1: br label %[[DONE]]
650 // CK1: [[DONE]]
651 // CK1: ret void
652 
653 #endif // CK1
654 
655 
656 ///==========================================================================///
657 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK2 --check-prefix CK2-64 %s
658 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
659 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK2 --check-prefix CK2-64 %s
660 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK2 --check-prefix CK2-32 %s
661 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
662 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK2 --check-prefix CK2-32 %s
663 
664 // RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
665 // RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
666 // RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
667 // RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
668 // RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
669 // RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
670 
671 #ifdef CK2
672 // Nested mappers.
673 
674 class B {
675 public:
676   double a;
677 };
678 
679 class C {
680 public:
681   double a;
682   B b;
683 };
684 
685 #pragma omp declare mapper(B s) map(s.a)
686 
687 #pragma omp declare mapper(id: C s) map(s.b)
688 
689 // CK2: define {{.*}}void [[BMPRFUNC:@[.]omp_mapper[.].*B[.]default]](i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
690 
691 // CK2-LABEL: define {{.*}}void @.omp_mapper.{{.*}}C{{.*}}.id(i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
692 // CK2: store i8* %{{[^,]+}}, i8** [[HANDLEADDR:%[^,]+]]
693 // CK2: store i8* %{{[^,]+}}, i8** [[BPTRADDR:%[^,]+]]
694 // CK2: store i8* %{{[^,]+}}, i8** [[VPTRADDR:%[^,]+]]
695 // CK2: store i64 %{{[^,]+}}, i{{64|32}}* [[SIZEADDR:%[^,]+]]
696 // CK2: store i64 %{{[^,]+}}, i64* [[TYPEADDR:%[^,]+]]
697 // CK2-DAG: [[BYTESIZE:%.+]] = load i64, i64* [[SIZEADDR]]
698 // CK2-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 16
699 // CK2-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]]
700 // CK2-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]]
701 // CK2-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C**
702 // CK2-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i64 [[SIZE]]
703 // CK2-DAG: [[BPTR:%.+]] = load i8*, i8** [[BPTRADDR]]
704 // CK2-DAG: [[BEGIN:%.+]] = load i8*, i8** [[VPTRADDR]]
705 // CK2: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
706 // CK2: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
707 
708 // CK2: [[INITEVALDEL]]
709 // CK2: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
710 // CK2: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
711 // CK2: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]]
712 // CK2: [[INIT]]
713 // CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
714 // CK2-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
715 // CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]])
716 // CK2: br label %[[LHEAD:[^,]+]]
717 
718 // CK2: [[LHEAD]]
719 // CK2: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]]
720 // CK2: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
721 // CK2: [[LBODY]]
722 // CK2: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
723 // CK2: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]]
724 // CK2-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1
725 // CK2-DAG: [[BEND:%.+]] = getelementptr %class.B, %class.B* [[BBEGIN]], i32 1
726 // CK2-DAG: [[BBEGINV:%.+]] = bitcast %class.B* [[BBEGIN]] to i8*
727 // CK2-DAG: [[BENDV:%.+]] = bitcast %class.B* [[BEND]] to i8*
728 // CK2-DAG: [[BBEGINI:%.+]] = ptrtoint i8* [[BBEGINV]] to i64
729 // CK2-DAG: [[BENDI:%.+]] = ptrtoint i8* [[BENDV]] to i64
730 // CK2-DAG: [[BSIZE:%.+]] = sub i64 [[BENDI]], [[BBEGINI]]
731 // CK2-DAG: [[BUSIZE:%.+]] = sdiv exact i64 [[BSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
732 // CK2-DAG: [[BPTRADDR0BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
733 // CK2-DAG: [[PTRADDR0BC:%.+]] = bitcast %class.B* [[BBEGIN]] to i8*
734 // CK2-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(i8* [[HANDLE]])
735 // CK2-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48
736 // CK2-DAG: br label %[[MEMBER:[^,]+]]
737 // CK2-DAG: [[MEMBER]]
738 // CK2-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
739 // CK2-DAG: [[MEMBERCOM]]
740 // CK2-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]]
741 // CK2-DAG: br label %[[LTYPE]]
742 // CK2-DAG: [[LTYPE]]
743 // CK2-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
744 // CK2-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
745 // CK2-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
746 // CK2-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
747 // CK2-DAG: [[ALLOC]]
748 // CK2-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
749 // CK2-DAG: br label %[[TYEND:[^,]+]]
750 // CK2-DAG: [[ALLOCELSE]]
751 // CK2-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
752 // CK2-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
753 // CK2-DAG: [[TO]]
754 // CK2-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
755 // CK2-DAG: br label %[[TYEND]]
756 // CK2-DAG: [[TOELSE]]
757 // CK2-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
758 // CK2-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
759 // CK2-DAG: [[FROM]]
760 // CK2-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
761 // CK2-DAG: br label %[[TYEND]]
762 // CK2-DAG: [[TYEND]]
763 // CK2-DAG: [[TYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
764 // CK2-64: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR0BC]], i8* [[PTRADDR0BC]], i64 [[BUSIZE]], i64 [[TYPE0]])
765 // CK2-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
766 // CK2-DAG: [[PTRADDR1BC:%.+]] = bitcast %class.B* [[BBEGIN]] to i8*
767 // CK2-DAG: br label %[[MEMBER:[^,]+]]
768 // CK2-DAG: [[MEMBER]]
769 // CK2-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
770 // CK2-DAG: [[MEMBERCOM]]
771 // 281474976710659 == 0x1,000,000,003
772 // CK2-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]]
773 // CK2-DAG: br label %[[LTYPE]]
774 // CK2-DAG: [[LTYPE]]
775 // CK2-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710659, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
776 // CK2-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
777 // CK2-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
778 // CK2-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
779 // CK2-DAG: [[ALLOC]]
780 // CK2-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
781 // CK2-DAG: br label %[[TYEND:[^,]+]]
782 // CK2-DAG: [[ALLOCELSE]]
783 // CK2-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
784 // CK2-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
785 // CK2-DAG: [[TO]]
786 // CK2-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
787 // CK2-DAG: br label %[[TYEND]]
788 // CK2-DAG: [[TOELSE]]
789 // CK2-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
790 // CK2-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
791 // CK2-DAG: [[FROM]]
792 // CK2-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
793 // CK2-DAG: br label %[[TYEND]]
794 // CK2-DAG: [[TYEND]]
795 // CK2-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
796 // CK2: call void [[BMPRFUNC]](i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 8, i64 [[TYPE1]])
797 // CK2: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1
798 // CK2: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]]
799 // CK2: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
800 
801 // CK2: [[LEXIT]]
802 // CK2: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
803 // CK2: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
804 // CK2: [[EVALDEL]]
805 // CK2: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
806 // CK2: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
807 // CK2: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]]
808 // CK2: [[DEL]]
809 // CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
810 // CK2-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
811 // CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]])
812 // CK2: br label %[[DONE]]
813 // CK2: [[DONE]]
814 // CK2: ret void
815 
816 #endif // CK2
817 
818 
819 ///==========================================================================///
820 // RUN: %clang_cc1 -DCK3 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK3 %s
821 // RUN: %clang_cc1 -DCK3 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
822 // RUN: %clang_cc1 -DCK3 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK3 %s
823 // RUN: %clang_cc1 -DCK3 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK3 %s
824 // RUN: %clang_cc1 -DCK3 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
825 // RUN: %clang_cc1 -DCK3 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK3 %s
826 
827 // RUN: %clang_cc1 -DCK3 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
828 // RUN: %clang_cc1 -DCK3 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
829 // RUN: %clang_cc1 -DCK3 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
830 // RUN: %clang_cc1 -DCK3 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
831 // RUN: %clang_cc1 -DCK3 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
832 // RUN: %clang_cc1 -DCK3 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
833 
834 #ifdef CK3
835 // map of array sections and nested components.
836 
837 // CK3-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}.region_id = weak constant i8 0
838 // CK3: [[TYPES:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710659, i64 35]
839 
840 class C {
841 public:
842   int a;
843   double *b;
844 };
845 
846 class B {
847 public:
848   C c;
849 };
850 
851 #pragma omp declare mapper(id: C s) map(s.a, s.b[0:2])
852 
853 // CK3: define {{.*}}void [[MPRFUNC:@[.]omp_mapper[.].*C[.]id]](i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
854 
855 // CK3-LABEL: define {{.*}}void @{{.*}}foo{{.*}}
856 void foo(int a){
857   // CK3-DAG: [[CVAL:%.+]] = alloca [10 x %class.C]
858   // CK3-DAG: [[BVAL:%.+]] = alloca %class.B
859   C c[10];
860   B b;
861 
862   // CK3-DAG: [[BC:%.+]] = getelementptr inbounds %class.B, %class.B* [[BVAL]], i32 0, i32 0
863   // CK3-DAG: [[BCEND:%.+]] = getelementptr %class.C, %class.C* [[BC]], i32 1
864   // CK3-DAG: [[BCC:%.+]] = bitcast %class.C* [[BC]] to i8*
865   // CK3-DAG: [[BCENDC:%.+]] = bitcast %class.C* [[BCEND]] to i8*
866   // CK3-DAG: [[BCI:%.+]] = ptrtoint i8* [[BCC]] to i64
867   // CK3-DAG: [[BCENDI:%.+]] = ptrtoint i8* [[BCENDC]] to i64
868   // CK3-DAG: [[BSIZE:%.+]] = sub i64 [[BCENDI]], [[BCI]]
869   // CK3-DAG: [[BSIZED:%.+]] = sdiv exact i64 [[BSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
870 
871   // CK3-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 3, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], i64* [[SGEP:%[^,]+]], {{.+}}[[TYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
872   // CK3-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
873   // CK3-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
874   // CK3-DAG: [[SGEP]] = getelementptr inbounds {{.+}}[[SIZES:%[^,]+]], i32 0, i32 0
875   // CK3-DAG: [[MPRGEP]] = bitcast [3 x i8*]* [[MPR:%[^,]+]] to i8**
876   // CK3-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
877   // CK3-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
878   // CK3-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[SIZES]], i32 0, i32 0
879   // CK3-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i{{64|32}} 0, i{{64|32}} 0
880   // CK3-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.B**
881   // CK3-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
882   // CK3-DAG: store %class.B* [[BVAL]], %class.B** [[CBP1]]
883   // CK3-DAG: store %class.C* [[BC]], %class.C** [[CP1]]
884   // CK3-DAG: store i64 [[BSIZED]], i64* [[S1]]
885   // CK3-DAG: store i8* null, i8** [[MPR1]]
886   // CK3-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1
887   // CK3-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1
888   // CK3-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[SIZES]], i32 0, i32 1
889   // CK3-DAG: [[MPR2:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i{{64|32}} 0, i{{64|32}} 1
890   // CK3-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to %class.B**
891   // CK3-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to %class.C**
892   // CK3-DAG: store %class.B* [[BVAL]], %class.B** [[CBP2]]
893   // CK3-DAG: store %class.C* [[BC]], %class.C** [[CP2]]
894   // CK3-64-DAG: store i64 16, i64* [[S2]]
895   // CK3-32-DAG: store i64 8, i64* [[S2]]
896   // CK3-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR2]]
897   // CK3-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 2
898   // CK3-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 2
899   // CK3-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[SIZES]], i32 0, i32 2
900   // CK3-DAG: [[MPR3:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i{{64|32}} 0, i{{64|32}} 2
901   // CK3-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to [10 x %class.C]**
902   // CK3-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to %class.C**
903   // CK3-DAG: store [10 x %class.C]* [[CVAL]], [10 x %class.C]** [[CBP3]]
904   // CK3-DAG: [[CVALGEP:%.+]] = getelementptr inbounds {{.+}}[[CVAL]], i{{64|32}} 0, i{{64|32}} 0
905   // CK3-DAG: store %class.C* [[CVALGEP]], %class.C** [[CP3]]
906   // CK3-64-DAG: store i64 160, i64* [[S3]]
907   // CK3-32-DAG: store i64 80, i64* [[S3]]
908   // CK3-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR3]]
909   // CK3: call void [[KERNEL:@.+]](%class.B* [[BVAL]], [10 x %class.C]* [[CVAL]])
910   #pragma omp target map(mapper(id),tofrom: c[0:10], b.c)
911   for (int i = 0; i < 10; i++) {
912     b.c.a += ++c[i].a;
913   }
914 }
915 
916 
917 // CK3: define internal void [[KERNEL]](%class.B* {{[^,]+}}, [10 x %class.C]* {{[^,]+}})
918 
919 #endif // CK3
920 
921 ///==========================================================================///
922 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK4 --check-prefix CK4-64 %s
923 // RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
924 // RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK4 --check-prefix CK4-64 %s
925 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK4 --check-prefix CK4-32 %s
926 // RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
927 // RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK4 --check-prefix CK4-32 %s
928 
929 // RUN: %clang_cc1 -DCK4 -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
930 // RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
931 // RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
932 // RUN: %clang_cc1 -DCK4 -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
933 // RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
934 // RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
935 
936 #ifdef CK4
937 // Mapper function code generation and runtime interface.
938 
939 // CK4-64: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
940 // CK4-32: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
941 // PRESENT=0x1000 | TARGET_PARAM=0x20 | TO=0x1 = 0x1021
942 // CK4: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1021]]]
943 
944 // CK4-64: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
945 // CK4-32: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
946 // PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 = 0x1022
947 // CK4: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1022]]]
948 
949 class C {
950 public:
951   int a;
952   double *b;
953 };
954 
955 #pragma omp declare mapper(id: C s) map(s.a, s.b[0:2])
956 
957 // CK4: define {{.*}}void [[MPRFUNC:@[.]omp_mapper[.].*C[.]id]](i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
958 // CK4: store i8* %{{[^,]+}}, i8** [[HANDLEADDR:%[^,]+]]
959 // CK4: store i8* %{{[^,]+}}, i8** [[BPTRADDR:%[^,]+]]
960 // CK4: store i8* %{{[^,]+}}, i8** [[VPTRADDR:%[^,]+]]
961 // CK4: store i64 %{{[^,]+}}, i{{64|32}}* [[SIZEADDR:%[^,]+]]
962 // CK4: store i64 %{{[^,]+}}, i64* [[TYPEADDR:%[^,]+]]
963 // CK4-DAG: [[BYTESIZE:%.+]] = load i64, i64* [[SIZEADDR]]
964 // CK4-64-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 16
965 // CK4-32-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 8
966 // CK4-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]]
967 // CK4-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]]
968 // CK4-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C**
969 // CK4-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i64 [[SIZE]]
970 // CK4-DAG: [[BPTR:%.+]] = load i8*, i8** [[BPTRADDR]]
971 // CK4-DAG: [[BEGIN:%.+]] = load i8*, i8** [[VPTRADDR]]
972 // CK4: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
973 // CK4: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
974 
975 // CK4: [[INITEVALDEL]]
976 // CK4: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
977 // CK4: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
978 // CK4: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]]
979 // CK4: [[INIT]]
980 // CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
981 // CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
982 // CK4-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
983 // CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]])
984 // CK4: br label %[[LHEAD:[^,]+]]
985 
986 // CK4: [[LHEAD]]
987 // CK4: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]]
988 // CK4: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
989 // CK4: [[LBODY]]
990 // CK4: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
991 // CK4: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]]
992 // CK4-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0
993 // CK4-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1
994 // CK4-DAG: [[BBEGIN2:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1
995 // CK4-DAG: [[BARRBEGIN:%.+]] = load double*, double** [[BBEGIN2]]
996 // CK4-DAG: [[BARRBEGINGEP:%.+]] = getelementptr inbounds double, double* [[BARRBEGIN]], i[[sz:64|32]] 0
997 // CK4-DAG: [[BEND:%.+]] = getelementptr double*, double** [[BBEGIN]], i32 1
998 // CK4-DAG: [[ABEGINV:%.+]] = bitcast i32* [[ABEGIN]] to i8*
999 // CK4-DAG: [[BENDV:%.+]] = bitcast double** [[BEND]] to i8*
1000 // CK4-DAG: [[ABEGINI:%.+]] = ptrtoint i8* [[ABEGINV]] to i64
1001 // CK4-DAG: [[BENDI:%.+]] = ptrtoint i8* [[BENDV]] to i64
1002 // CK4-DAG: [[CSIZE:%.+]] = sub i64 [[BENDI]], [[ABEGINI]]
1003 // CK4-DAG: [[CUSIZE:%.+]] = sdiv exact i64 [[CSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
1004 // CK4-DAG: [[BPTRADDR0BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
1005 // CK4-DAG: [[PTRADDR0BC:%.+]] = bitcast i32* [[ABEGIN]] to i8*
1006 // CK4-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(i8* [[HANDLE]])
1007 // CK4-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48
1008 // CK4-DAG: br label %[[MEMBER:[^,]+]]
1009 // CK4-DAG: [[MEMBER]]
1010 // CK4-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
1011 // CK4-DAG: [[MEMBERCOM]]
1012 // CK4-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]]
1013 // CK4-DAG: br label %[[LTYPE]]
1014 // CK4-DAG: [[LTYPE]]
1015 // CK4-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
1016 // CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
1017 // CK4-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
1018 // CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
1019 // CK4-DAG: [[ALLOC]]
1020 // CK4-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
1021 // CK4-DAG: br label %[[TYEND:[^,]+]]
1022 // CK4-DAG: [[ALLOCELSE]]
1023 // CK4-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
1024 // CK4-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
1025 // CK4-DAG: [[TO]]
1026 // CK4-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
1027 // CK4-DAG: br label %[[TYEND]]
1028 // CK4-DAG: [[TOELSE]]
1029 // CK4-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
1030 // CK4-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
1031 // CK4-DAG: [[FROM]]
1032 // CK4-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
1033 // CK4-DAG: br label %[[TYEND]]
1034 // CK4-DAG: [[TYEND]]
1035 // CK4-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
1036 // CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR0BC]], i8* [[PTRADDR0BC]], i64 [[CUSIZE]], i64 [[PHITYPE0]])
1037 // CK4-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
1038 // CK4-DAG: [[PTRADDR1BC:%.+]] = bitcast i32* [[ABEGIN]] to i8*
1039 // CK4-DAG: br label %[[MEMBER:[^,]+]]
1040 // CK4-DAG: [[MEMBER]]
1041 // CK4-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
1042 // CK4-DAG: [[MEMBERCOM]]
1043 // 281474976710659 == 0x1,000,000,003
1044 // CK4-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]]
1045 // CK4-DAG: br label %[[LTYPE]]
1046 // CK4-DAG: [[LTYPE]]
1047 // CK4-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710659, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
1048 // CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
1049 // CK4-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
1050 // CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
1051 // CK4-DAG: [[ALLOC]]
1052 // CK4-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
1053 // CK4-DAG: br label %[[TYEND:[^,]+]]
1054 // CK4-DAG: [[ALLOCELSE]]
1055 // CK4-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
1056 // CK4-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
1057 // CK4-DAG: [[TO]]
1058 // CK4-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
1059 // CK4-DAG: br label %[[TYEND]]
1060 // CK4-DAG: [[TOELSE]]
1061 // CK4-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
1062 // CK4-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
1063 // CK4-DAG: [[FROM]]
1064 // CK4-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
1065 // CK4-DAG: br label %[[TYEND]]
1066 // CK4-DAG: [[TYEND]]
1067 // CK4-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
1068 // CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 4, i64 [[TYPE1]])
1069 // CK4-DAG: [[BPTRADDR2BC:%.+]] = bitcast double** [[BBEGIN]] to i8*
1070 // CK4-DAG: [[PTRADDR2BC:%.+]] = bitcast double* [[BARRBEGINGEP]] to i8*
1071 // CK4-DAG: br label %[[MEMBER:[^,]+]]
1072 // CK4-DAG: [[MEMBER]]
1073 // CK4-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
1074 // CK4-DAG: [[MEMBERCOM]]
1075 // 281474976710675 == 0x1,000,000,013
1076 // CK4-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]]
1077 // CK4-DAG: br label %[[LTYPE]]
1078 // CK4-DAG: [[LTYPE]]
1079 // CK4-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710675, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
1080 // CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
1081 // CK4-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
1082 // CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
1083 // CK4-DAG: [[ALLOC]]
1084 // CK4-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
1085 // CK4-DAG: br label %[[TYEND:[^,]+]]
1086 // CK4-DAG: [[ALLOCELSE]]
1087 // CK4-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
1088 // CK4-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
1089 // CK4-DAG: [[TO]]
1090 // CK4-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
1091 // CK4-DAG: br label %[[TYEND]]
1092 // CK4-DAG: [[TOELSE]]
1093 // CK4-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
1094 // CK4-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
1095 // CK4-DAG: [[FROM]]
1096 // CK4-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
1097 // CK4-DAG: br label %[[TYEND]]
1098 // CK4-DAG: [[TYEND]]
1099 // CK4-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
1100 // CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR2BC]], i8* [[PTRADDR2BC]], i64 16, i64 [[TYPE2]])
1101 // CK4: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1
1102 // CK4: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]]
1103 // CK4: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
1104 
1105 // CK4: [[LEXIT]]
1106 // CK4: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
1107 // CK4: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
1108 // CK4: [[EVALDEL]]
1109 // CK4: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
1110 // CK4: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
1111 // CK4: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]]
1112 // CK4: [[DEL]]
1113 // CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
1114 // CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
1115 // CK4-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
1116 // CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]])
1117 // CK4: br label %[[DONE]]
1118 // CK4: [[DONE]]
1119 // CK4: ret void
1120 
1121 
1122 // CK4-LABEL: define {{.*}}void @{{.*}}foo{{.*}}
1123 void foo(int a){
1124   int i = a;
1125   C c;
1126   c.a = a;
1127 
1128   // CK4-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[TGEPBP:%.+]], i8** [[TGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[TSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[TTYPES]]{{.+}}, i8** [[TMPRGEP:%.+]])
1129   // CK4-DAG: [[TGEPBP]] = getelementptr inbounds {{.+}}[[TBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
1130   // CK4-DAG: [[TGEPP]] = getelementptr inbounds {{.+}}[[TP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
1131   // CK4-DAG: [[TMPRGEP]] = bitcast [1 x i8*]* [[TMPR:%[^,]+]] to i8**
1132   // CK4-DAG: [[TBP0:%.+]] = getelementptr inbounds {{.+}}[[TBP]], i{{.+}} 0, i{{.+}} 0
1133   // CK4-DAG: [[TP0:%.+]] = getelementptr inbounds {{.+}}[[TP]], i{{.+}} 0, i{{.+}} 0
1134   // CK4-DAG: [[TMPR1:%.+]] = getelementptr inbounds {{.+}}[[TMPR]], i[[sz]] 0, i[[sz]] 0
1135   // CK4-DAG: [[TCBP0:%.+]] = bitcast i8** [[TBP0]] to %class.C**
1136   // CK4-DAG: [[TCP0:%.+]] = bitcast i8** [[TP0]] to %class.C**
1137   // CK4-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[TCBP0]]
1138   // CK4-DAG: store %class.C* [[VAL]], %class.C** [[TCP0]]
1139   // CK4-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[TMPR1]]
1140   #pragma omp target update to(present, mapper(id): c)
1141 
1142   // CK4-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[FSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[FTYPES]]{{.+}}, i8** [[FMPRGEP:%.+]])
1143   // CK4-DAG: [[FGEPBP]] = getelementptr inbounds {{.+}}[[FBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
1144   // CK4-DAG: [[FGEPP]] = getelementptr inbounds {{.+}}[[FP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
1145   // CK4-DAG: [[FMPRGEP]] = bitcast [1 x i8*]* [[FMPR:%[^,]+]] to i8**
1146   // CK4-DAG: [[FBP0:%.+]] = getelementptr inbounds {{.+}}[[FBP]], i{{.+}} 0, i{{.+}} 0
1147   // CK4-DAG: [[FP0:%.+]] = getelementptr inbounds {{.+}}[[FP]], i{{.+}} 0, i{{.+}} 0
1148   // CK4-DAG: [[FMPR1:%.+]] = getelementptr inbounds {{.+}}[[FMPR]], i[[sz]] 0, i[[sz]] 0
1149   // CK4-DAG: [[FCBP0:%.+]] = bitcast i8** [[FBP0]] to %class.C**
1150   // CK4-DAG: [[FCP0:%.+]] = bitcast i8** [[FP0]] to %class.C**
1151   // CK4-DAG: store %class.C* [[VAL]], %class.C** [[FCBP0]]
1152   // CK4-DAG: store %class.C* [[VAL]], %class.C** [[FCP0]]
1153   // CK4-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[FMPR1]]
1154   #pragma omp target update from(mapper(id), present: c)
1155 }
1156 
1157 #endif // CK4
1158 
1159 #endif // HEADER
1160