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 
24 // CK0-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}.region_id = weak constant i8 0
25 // CK0-64: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
26 // CK0-32: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
27 // CK0: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 35]
28 // CK0-64: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
29 // CK0-32: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
30 // CK0: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33]
31 // CK0-64: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
32 // CK0-32: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
33 // CK0: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34]
34 
35 class C {
36 public:
37   int a;
38   double *b;
39 };
40 
41 #pragma omp declare mapper(id: C s) map(s.a, s.b[0:2])
42 
43 // CK0-LABEL: define {{.*}}void @.omp_mapper.{{.*}}C.id{{.*}}(i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
44 // CK0: store i8* %{{[^,]+}}, i8** [[HANDLEADDR:%[^,]+]]
45 // CK0: store i8* %{{[^,]+}}, i8** [[BPTRADDR:%[^,]+]]
46 // CK0: store i8* %{{[^,]+}}, i8** [[VPTRADDR:%[^,]+]]
47 // CK0: store i64 %{{[^,]+}}, i{{64|32}}* [[SIZEADDR:%[^,]+]]
48 // CK0: store i64 %{{[^,]+}}, i64* [[TYPEADDR:%[^,]+]]
49 // CK0-DAG: [[SIZE:%.+]] = load i64, i64* [[SIZEADDR]]
50 // CK0-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]]
51 // CK0-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]]
52 // CK0-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C**
53 // CK0-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i64 [[SIZE]]
54 // CK0-DAG: [[BPTR:%.+]] = load i8*, i8** [[BPTRADDR]]
55 // CK0-DAG: [[BEGIN:%.+]] = load i8*, i8** [[VPTRADDR]]
56 // CK0: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
57 // CK0: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
58 
59 // CK0: [[INITEVALDEL]]
60 // CK0: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
61 // CK0: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
62 // CK0: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]]
63 // CK0: [[INIT]]
64 // CK0-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
65 // CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
66 // CK0-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
67 // CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]])
68 // CK0: br label %[[LHEAD:[^,]+]]
69 
70 // CK0: [[LHEAD]]
71 // CK0: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]]
72 // CK0: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
73 // CK0: [[LBODY]]
74 // CK0: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
75 // CK0: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]]
76 // CK0-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0
77 // CK0-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1
78 // CK0-DAG: [[BBEGIN2:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1
79 // CK0-DAG: [[BARRBEGIN:%.+]] = load double*, double** [[BBEGIN2]]
80 // CK0-DAG: [[BARRBEGINGEP:%.+]] = getelementptr inbounds double, double* [[BARRBEGIN]], i[[sz:64|32]] 0
81 // CK0-DAG: [[BEND:%.+]] = getelementptr double*, double** [[BBEGIN]], i32 1
82 // CK0-DAG: [[ABEGINV:%.+]] = bitcast i32* [[ABEGIN]] to i8*
83 // CK0-DAG: [[BENDV:%.+]] = bitcast double** [[BEND]] to i8*
84 // CK0-DAG: [[ABEGINI:%.+]] = ptrtoint i8* [[ABEGINV]] to i64
85 // CK0-DAG: [[BENDI:%.+]] = ptrtoint i8* [[BENDV]] to i64
86 // CK0-DAG: [[CSIZE:%.+]] = sub i64 [[BENDI]], [[ABEGINI]]
87 // CK0-DAG: [[CUSIZE:%.+]] = sdiv exact i64 [[CSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
88 // CK0-DAG: [[BPTRADDR0BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
89 // CK0-DAG: [[PTRADDR0BC:%.+]] = bitcast i32* [[ABEGIN]] to i8*
90 // CK0-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(i8* [[HANDLE]])
91 // CK0-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48
92 // CK0-DAG: br label %[[MEMBER:[^,]+]]
93 // CK0-DAG: [[MEMBER]]
94 // CK0-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
95 // CK0-DAG: [[MEMBERCOM]]
96 // CK0-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]]
97 // CK0-DAG: br label %[[LTYPE]]
98 // CK0-DAG: [[LTYPE]]
99 // CK0-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
100 // CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
101 // CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
102 // CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
103 // CK0-DAG: [[ALLOC]]
104 // CK0-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
105 // CK0-DAG: br label %[[TYEND:[^,]+]]
106 // CK0-DAG: [[ALLOCELSE]]
107 // CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
108 // CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
109 // CK0-DAG: [[TO]]
110 // CK0-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
111 // CK0-DAG: br label %[[TYEND]]
112 // CK0-DAG: [[TOELSE]]
113 // CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
114 // CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
115 // CK0-DAG: [[FROM]]
116 // CK0-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
117 // CK0-DAG: br label %[[TYEND]]
118 // CK0-DAG: [[TYEND]]
119 // CK0-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
120 // CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR0BC]], i8* [[PTRADDR0BC]], i64 [[CUSIZE]], i64 [[PHITYPE0]])
121 // CK0-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
122 // CK0-DAG: [[PTRADDR1BC:%.+]] = bitcast i32* [[ABEGIN]] to i8*
123 // CK0-DAG: br label %[[MEMBER:[^,]+]]
124 // CK0-DAG: [[MEMBER]]
125 // CK0-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
126 // CK0-DAG: [[MEMBERCOM]]
127 // 281474976710659 == 0x1,000,000,003
128 // CK0-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]]
129 // CK0-DAG: br label %[[LTYPE]]
130 // CK0-DAG: [[LTYPE]]
131 // CK0-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710659, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
132 // CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
133 // CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
134 // CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
135 // CK0-DAG: [[ALLOC]]
136 // CK0-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
137 // CK0-DAG: br label %[[TYEND:[^,]+]]
138 // CK0-DAG: [[ALLOCELSE]]
139 // CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
140 // CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
141 // CK0-DAG: [[TO]]
142 // CK0-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
143 // CK0-DAG: br label %[[TYEND]]
144 // CK0-DAG: [[TOELSE]]
145 // CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
146 // CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
147 // CK0-DAG: [[FROM]]
148 // CK0-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
149 // CK0-DAG: br label %[[TYEND]]
150 // CK0-DAG: [[TYEND]]
151 // CK0-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
152 // CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 4, i64 [[TYPE1]])
153 // CK0-DAG: [[BPTRADDR2BC:%.+]] = bitcast double** [[BBEGIN]] to i8*
154 // CK0-DAG: [[PTRADDR2BC:%.+]] = bitcast double* [[BARRBEGINGEP]] to i8*
155 // CK0-DAG: br label %[[MEMBER:[^,]+]]
156 // CK0-DAG: [[MEMBER]]
157 // CK0-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
158 // CK0-DAG: [[MEMBERCOM]]
159 // 281474976710675 == 0x1,000,000,013
160 // CK0-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]]
161 // CK0-DAG: br label %[[LTYPE]]
162 // CK0-DAG: [[LTYPE]]
163 // CK0-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710675, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
164 // CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
165 // CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
166 // CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
167 // CK0-DAG: [[ALLOC]]
168 // CK0-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
169 // CK0-DAG: br label %[[TYEND:[^,]+]]
170 // CK0-DAG: [[ALLOCELSE]]
171 // CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
172 // CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
173 // CK0-DAG: [[TO]]
174 // CK0-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
175 // CK0-DAG: br label %[[TYEND]]
176 // CK0-DAG: [[TOELSE]]
177 // CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
178 // CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
179 // CK0-DAG: [[FROM]]
180 // CK0-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
181 // CK0-DAG: br label %[[TYEND]]
182 // CK0-DAG: [[TYEND]]
183 // CK0-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
184 // CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR2BC]], i8* [[PTRADDR2BC]], i64 16, i64 [[TYPE2]])
185 // CK0: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1
186 // CK0: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]]
187 // CK0: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
188 
189 // CK0: [[LEXIT]]
190 // CK0: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
191 // CK0: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
192 // CK0: [[EVALDEL]]
193 // CK0: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
194 // CK0: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
195 // CK0: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]]
196 // CK0: [[DEL]]
197 // CK0-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
198 // CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
199 // CK0-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
200 // CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]])
201 // CK0: br label %[[DONE]]
202 // CK0: [[DONE]]
203 // CK0: ret void
204 
205 
206 // CK0-LABEL: define {{.*}}void @{{.*}}foo{{.*}}
207 void foo(int a){
208   int i = a;
209   C c;
210   c.a = a;
211 
212   // CK0-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
213   // CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
214   // CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
215   // CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
216   // CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
217   // CK0-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C**
218   // CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
219   // CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
220   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
221   // CK0: call void [[KERNEL:@.+]](%class.C* [[VAL]])
222   #pragma omp target map(mapper(id),tofrom: c)
223   {
224    ++c.a;
225   }
226 
227   // CK0-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[TGEPBP:%.+]], i8** [[TGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[TSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[TTYPES]]{{.+}})
228   // CK0-DAG: [[TGEPBP]] = getelementptr inbounds {{.+}}[[TBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
229   // CK0-DAG: [[TGEPP]] = getelementptr inbounds {{.+}}[[TP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
230   // CK0-DAG: [[TBP0:%.+]] = getelementptr inbounds {{.+}}[[TBP]], i{{.+}} 0, i{{.+}} 0
231   // CK0-DAG: [[TP0:%.+]] = getelementptr inbounds {{.+}}[[TP]], i{{.+}} 0, i{{.+}} 0
232   // CK0-DAG: [[TCBP0:%.+]] = bitcast i8** [[TBP0]] to %class.C**
233   // CK0-DAG: [[TCP0:%.+]] = bitcast i8** [[TP0]] to %class.C**
234   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[TCBP0]]
235   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[TCP0]]
236   #pragma omp target update to(mapper(id): c)
237 
238   // CK0-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[FSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[FTYPES]]{{.+}})
239   // CK0-DAG: [[FGEPBP]] = getelementptr inbounds {{.+}}[[FBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
240   // CK0-DAG: [[FGEPP]] = getelementptr inbounds {{.+}}[[FP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
241   // CK0-DAG: [[FBP0:%.+]] = getelementptr inbounds {{.+}}[[FBP]], i{{.+}} 0, i{{.+}} 0
242   // CK0-DAG: [[FP0:%.+]] = getelementptr inbounds {{.+}}[[FP]], i{{.+}} 0, i{{.+}} 0
243   // CK0-DAG: [[FCBP0:%.+]] = bitcast i8** [[FBP0]] to %class.C**
244   // CK0-DAG: [[FCP0:%.+]] = bitcast i8** [[FP0]] to %class.C**
245   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[FCBP0]]
246   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[FCP0]]
247   #pragma omp target update from(mapper(id): c)
248 }
249 
250 
251 // CK0: define internal void [[KERNEL]](%class.C* {{.+}}[[ARG:%.+]])
252 // CK0: [[ADDR:%.+]] = alloca %class.C*,
253 // CK0: store %class.C* [[ARG]], %class.C** [[ADDR]]
254 // CK0: [[CADDR:%.+]] = load %class.C*, %class.C** [[ADDR]]
255 // CK0: [[CAADDR:%.+]] = getelementptr inbounds %class.C, %class.C* [[CADDR]], i32 0, i32 0
256 // CK0: [[VAL:%[^,]+]] = load i32, i32* [[CAADDR]]
257 // CK0: {{.+}} = add nsw i32 [[VAL]], 1
258 // CK0: }
259 
260 #endif
261 
262 
263 ///==========================================================================///
264 // 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
265 // 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
266 // 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
267 // 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
268 // 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
269 // 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
270 
271 // 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
272 // 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
273 // 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
274 // 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
275 // 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
276 // 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
277 
278 #ifdef CK1
279 
280 template <class T>
281 class C {
282 public:
283   T a;
284 };
285 
286 #pragma omp declare mapper(id: C<int> s) map(s.a)
287 
288 // CK1-LABEL: define {{.*}}void @.omp_mapper.{{.*}}C{{.*}}.id{{.*}}(i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
289 // CK1: store i8* %{{[^,]+}}, i8** [[HANDLEADDR:%[^,]+]]
290 // CK1: store i8* %{{[^,]+}}, i8** [[BPTRADDR:%[^,]+]]
291 // CK1: store i8* %{{[^,]+}}, i8** [[VPTRADDR:%[^,]+]]
292 // CK1: store i64 %{{[^,]+}}, i{{64|32}}* [[SIZEADDR:%[^,]+]]
293 // CK1: store i64 %{{[^,]+}}, i64* [[TYPEADDR:%[^,]+]]
294 // CK1-DAG: [[SIZE:%.+]] = load i64, i64* [[SIZEADDR]]
295 // CK1-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]]
296 // CK1-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]]
297 // CK1-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C**
298 // CK1-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i64 [[SIZE]]
299 // CK1-DAG: [[BPTR:%.+]] = load i8*, i8** [[BPTRADDR]]
300 // CK1-DAG: [[BEGIN:%.+]] = load i8*, i8** [[VPTRADDR]]
301 // CK1: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
302 // CK1: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
303 
304 // CK1: [[INITEVALDEL]]
305 // CK1: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
306 // CK1: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
307 // CK1: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]]
308 // CK1: [[INIT]]
309 // CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4
310 // CK1-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
311 // CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]])
312 // CK1: br label %[[LHEAD:[^,]+]]
313 
314 // CK1: [[LHEAD]]
315 // CK1: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]]
316 // CK1: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
317 // CK1: [[LBODY]]
318 // CK1: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
319 // CK1: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]]
320 // CK1-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0
321 // CK1-DAG: [[AEND:%.+]] = getelementptr i32, i32* [[ABEGIN]], i32 1
322 // CK1-DAG: [[ABEGINV:%.+]] = bitcast i32* [[ABEGIN]] to i8*
323 // CK1-DAG: [[AENDV:%.+]] = bitcast i32* [[AEND]] to i8*
324 // CK1-DAG: [[ABEGINI:%.+]] = ptrtoint i8* [[ABEGINV]] to i64
325 // CK1-DAG: [[AENDI:%.+]] = ptrtoint i8* [[AENDV]] to i64
326 // CK1-DAG: [[CSIZE:%.+]] = sub i64 [[AENDI]], [[ABEGINI]]
327 // CK1-DAG: [[CUSIZE:%.+]] = sdiv exact i64 [[CSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
328 // CK1-DAG: [[BPTRADDR0BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
329 // CK1-DAG: [[PTRADDR0BC:%.+]] = bitcast i32* [[ABEGIN]] to i8*
330 // CK1-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(i8* [[HANDLE]])
331 // CK1-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48
332 // CK1-DAG: br label %[[MEMBER:[^,]+]]
333 // CK1-DAG: [[MEMBER]]
334 // CK1-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
335 // CK1-DAG: [[MEMBERCOM]]
336 // CK1-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]]
337 // CK1-DAG: br label %[[LTYPE]]
338 // CK1-DAG: [[LTYPE]]
339 // CK1-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
340 // CK1-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
341 // CK1-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
342 // CK1-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
343 // CK1-DAG: [[ALLOC]]
344 // CK1-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
345 // CK1-DAG: br label %[[TYEND:[^,]+]]
346 // CK1-DAG: [[ALLOCELSE]]
347 // CK1-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
348 // CK1-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
349 // CK1-DAG: [[TO]]
350 // CK1-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
351 // CK1-DAG: br label %[[TYEND]]
352 // CK1-DAG: [[TOELSE]]
353 // CK1-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
354 // CK1-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
355 // CK1-DAG: [[FROM]]
356 // CK1-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
357 // CK1-DAG: br label %[[TYEND]]
358 // CK1-DAG: [[TYEND]]
359 // CK1-DAG: [[TYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
360 // CK1-64: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR0BC]], i8* [[PTRADDR0BC]], i64 [[CUSIZE]], i64 [[TYPE0]])
361 // CK1-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
362 // CK1-DAG: [[PTRADDR1BC:%.+]] = bitcast i32* [[ABEGIN]] to i8*
363 // CK1-DAG: br label %[[MEMBER:[^,]+]]
364 // CK1-DAG: [[MEMBER]]
365 // CK1-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
366 // CK1-DAG: [[MEMBERCOM]]
367 // 281474976710659 == 0x1,000,000,003
368 // CK1-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]]
369 // CK1-DAG: br label %[[LTYPE]]
370 // CK1-DAG: [[LTYPE]]
371 // CK1-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710659, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
372 // CK1-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
373 // CK1-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
374 // CK1-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
375 // CK1-DAG: [[ALLOC]]
376 // CK1-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
377 // CK1-DAG: br label %[[TYEND:[^,]+]]
378 // CK1-DAG: [[ALLOCELSE]]
379 // CK1-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
380 // CK1-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
381 // CK1-DAG: [[TO]]
382 // CK1-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
383 // CK1-DAG: br label %[[TYEND]]
384 // CK1-DAG: [[TOELSE]]
385 // CK1-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
386 // CK1-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
387 // CK1-DAG: [[FROM]]
388 // CK1-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
389 // CK1-DAG: br label %[[TYEND]]
390 // CK1-DAG: [[TYEND]]
391 // CK1-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
392 // CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 4, i64 [[TYPE1]])
393 // CK1: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1
394 // CK1: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]]
395 // CK1: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
396 
397 // CK1: [[LEXIT]]
398 // CK1: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
399 // CK1: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
400 // CK1: [[EVALDEL]]
401 // CK1: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
402 // CK1: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
403 // CK1: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]]
404 // CK1: [[DEL]]
405 // CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4
406 // CK1-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
407 // CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]])
408 // CK1: br label %[[DONE]]
409 // CK1: [[DONE]]
410 // CK1: ret void
411 
412 #endif
413 
414 #endif
415