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