xref: /llvm-project/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp (revision a5ea6760674762cb597cf328dc467f1296633da0) 
1 // expected-no-diagnostics
2 #ifndef HEADER
3 #define HEADER
4 
5 ///==========================================================================///
6 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64
7 // RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
8 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK1 --check-prefix CK1-64
9 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK1 --check-prefix CK1-32
10 // RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
11 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK1 --check-prefix CK1-32
12 
13 // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
14 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
15 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
16 // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
17 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
18 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
19 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
20 #ifdef CK1
21 
22 double *g;
23 
24 // CK1: @g ={{.*}} global ptr
25 // CK1: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 19, i64 64]
26 // CK1: [[MTYPE01:@.+]] = {{.*}}constant [1 x i64] [i64 67]
27 // CK1: [[MTYPE03:@.+]] = {{.*}}constant [1 x i64] [i64 67]
28 // CK1: [[MTYPE04:@.+]] = {{.*}}constant [1 x i64] [i64 67]
29 // CK1: [[MTYPE05:@.+]] = {{.*}}constant [1 x i64] [i64 67]
30 // CK1: [[MTYPE06:@.+]] = {{.*}}constant [1 x i64] [i64 67]
31 // CK1: [[MTYPE07:@.+]] = {{.*}}constant [1 x i64] [i64 67]
32 // CK1: [[MTYPE08:@.+]] = {{.*}}constant [2 x i64] [i64 67, i64 3]
33 // CK1: [[MTYPE09:@.+]] = {{.*}}constant [2 x i64] [i64 67, i64 67]
34 // CK1: [[MTYPE10:@.+]] = {{.*}}constant [2 x i64] [i64 67, i64 67]
35 // CK1: [[MTYPE11:@.+]] = {{.*}}constant [2 x i64] [i64 3, i64 64]
36 // CK1: [[MTYPE12:@.+]] = {{.*}}constant [2 x i64] [i64 3, i64 64]
37 
38 // CK1-LABEL: @_Z3foo
39 template<typename T>
40 void foo(float *&lr, T *&tr) {
41   float *l;
42   T *t;
43 
44   // CK1:     [[T:%.+]] = load ptr, ptr [[DECL:@g]],
45   // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x ptr], ptr %{{.+}}, i32 0, i32 1
46   // CK1:     store ptr [[T]], ptr [[BP]],
47   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE00]]
48   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
49   // CK1-NOT: store ptr [[VAL]], ptr [[DECL]],
50   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
51   // CK1:     [[TT:%.+]] = load ptr, ptr [[PVT]],
52   // CK1:     getelementptr inbounds double, ptr [[TT]], i32 1
53   #pragma omp target data map(g[:10]) use_device_ptr(g)
54   {
55     ++g;
56   }
57   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE00]]
58   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
59   // CK1:     getelementptr inbounds double, ptr [[TTT]], i32 1
60   ++g;
61 
62   // CK1:     [[T1:%.+]] = load ptr, ptr [[DECL:%.+]],
63   // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x ptr], ptr %{{.+}}, i32 0, i32 0
64   // CK1:     store ptr [[T1]], ptr [[BP]],
65   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE01]]
66   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
67   // CK1-NOT: store ptr [[VAL]], ptr [[DECL]],
68   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
69   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
70   // CK1:     getelementptr inbounds float, ptr [[TT1]], i32 1
71   #pragma omp target data map(l[:10]) use_device_ptr(l)
72   {
73     ++l;
74   }
75   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE01]]
76   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
77   // CK1:     getelementptr inbounds float, ptr [[TTT]], i32 1
78   ++l;
79 
80   // CK1-NOT: call void @__tgt_target
81   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
82   // CK1:     getelementptr inbounds float, ptr [[TTT]], i32 1
83   #pragma omp target data map(l[:10]) use_device_ptr(l) if(0)
84   {
85     ++l;
86   }
87   // CK1-NOT: call void @__tgt_target
88   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
89   // CK1:     getelementptr inbounds float, ptr [[TTT]], i32 1
90   ++l;
91 
92   // CK1:     [[T1:%.+]] = load ptr, ptr [[DECL:%.+]],
93   // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x ptr], ptr %{{.+}}, i32 0, i32 0
94   // CK1:     store ptr [[T1]], ptr [[BP]],
95   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE03]]
96   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
97   // CK1-NOT: store ptr [[VAL]], ptr [[DECL]],
98   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
99   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
100   // CK1:     getelementptr inbounds float, ptr [[TT1]], i32 1
101   #pragma omp target data map(l[:10]) use_device_ptr(l) if(1)
102   {
103     ++l;
104   }
105   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE03]]
106   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
107   // CK1:     getelementptr inbounds float, ptr [[TTT]], i32 1
108   ++l;
109 
110   // CK1:     [[CMP:%.+]] = icmp ne ptr %{{.+}}, null
111   // CK1:     br i1 [[CMP]], label %[[BTHEN:.+]], label %[[BELSE:.+]]
112 
113   // CK1:     [[BTHEN]]:
114   // CK1:     [[T1:%.+]] = load ptr, ptr [[DECL:%.+]],
115   // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x ptr], ptr %{{.+}}, i32 0, i32 0
116   // CK1:     store ptr [[T1]], ptr [[BP]],
117   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE04]]
118   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
119   // CK1-NOT: store ptr [[VAL]], ptr [[DECL]],
120   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
121   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
122   // CK1:     getelementptr inbounds float, ptr [[TT1]], i32 1
123   // CK1:     br label %[[BEND:.+]]
124 
125   // CK1:     [[BELSE]]:
126   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
127   // CK1:     getelementptr inbounds float, ptr [[TTT]], i32 1
128   // CK1:     br label %[[BEND]]
129   #pragma omp target data map(l[:10]) use_device_ptr(l) if(lr != 0)
130   {
131     ++l;
132   }
133   // CK1:     [[BEND]]:
134   // CK1:     [[CMP:%.+]] = icmp ne ptr %{{.+}}, null
135   // CK1:     br i1 [[CMP]], label %[[BTHEN:.+]], label %[[BELSE:.+]]
136 
137   // CK1:     [[BTHEN]]:
138   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE04]]
139   // CK1:     br label %[[BEND:.+]]
140 
141   // CK1:     [[BELSE]]:
142   // CK1:     br label %[[BEND]]
143 
144   // CK1:     [[BEND]]:
145   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
146   // CK1:     getelementptr inbounds float, ptr [[TTT]], i32 1
147   ++l;
148 
149   // CK1:     [[T2:%.+]] = load ptr, ptr [[DECL:%.+]],
150   // CK1:     [[T1:%.+]] = load ptr, ptr [[T2]],
151   // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x ptr], ptr %{{.+}}, i32 0, i32 0
152   // CK1:     store ptr [[T1]], ptr [[BP]],
153   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE05]]
154   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
155   // CK1:     store ptr [[VAL]], ptr [[PVTV:%.+]],
156   // CK1-NOT: store ptr [[PVTV]], ptr [[DECL]],
157   // CK1:     store ptr [[PVTV]], ptr [[PVT:%.+]],
158   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
159   // CK1:     [[TT2:%.+]] = load ptr, ptr [[TT1]],
160   // CK1:     getelementptr inbounds float, ptr [[TT2]], i32 1
161   #pragma omp target data map(lr[:10]) use_device_ptr(lr)
162   {
163     ++lr;
164   }
165   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE05]]
166   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
167   // CK1:     [[TTTT:%.+]] = load ptr, ptr [[TTT]],
168   // CK1:     getelementptr inbounds float, ptr [[TTTT]], i32 1
169   ++lr;
170 
171   // CK1:     [[T1:%.+]] = load ptr, ptr [[DECL:%.+]],
172   // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x ptr], ptr %{{.+}}, i32 0, i32 0
173   // CK1:     store ptr [[T1]], ptr [[BP]],
174   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE06]]
175   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
176   // CK1-NOT: store ptr [[VAL]], ptr [[DECL]],
177   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
178   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
179   // CK1:     getelementptr inbounds i32, ptr [[TT1]], i32 1
180   #pragma omp target data map(t[:10]) use_device_ptr(t)
181   {
182     ++t;
183   }
184   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE06]]
185   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
186   // CK1:     getelementptr inbounds i32, ptr [[TTT]], i32 1
187   ++t;
188 
189   // CK1:     [[T2:%.+]] = load ptr, ptr [[DECL:%.+]],
190   // CK1:     [[T1:%.+]] = load ptr, ptr [[T2]],
191   // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x ptr], ptr %{{.+}}, i32 0, i32 0
192   // CK1:     store ptr [[T1]], ptr [[BP]],
193   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE07]]
194   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
195   // CK1:     store ptr [[VAL]], ptr [[PVTV:%.+]],
196   // CK1-NOT: store ptr [[PVTV]], ptr [[DECL]],
197   // CK1:     store ptr [[PVTV]], ptr [[PVT:%.+]],
198   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
199   // CK1:     [[TT2:%.+]] = load ptr, ptr [[TT1]],
200   // CK1:     getelementptr inbounds i32, ptr [[TT2]], i32 1
201   #pragma omp target data map(tr[:10]) use_device_ptr(tr)
202   {
203     ++tr;
204   }
205   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE07]]
206   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
207   // CK1:     [[TTTT:%.+]] = load ptr, ptr [[TTT]],
208   // CK1:     getelementptr inbounds i32, ptr [[TTTT]], i32 1
209   ++tr;
210 
211   // CK1:     [[T1:%.+]] = load ptr, ptr [[DECL:%.+]],
212   // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x ptr], ptr %{{.+}}, i32 0, i32 0
213   // CK1:     store ptr [[T1]], ptr [[BP]],
214   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE08]]
215   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
216   // CK1-NOT: store ptr [[VAL]], ptr [[DECL]],
217   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
218   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
219   // CK1:     getelementptr inbounds float, ptr [[TT1]], i32 1
220   #pragma omp target data map(l[:10], t[:10]) use_device_ptr(l)
221   {
222     ++l; ++t;
223   }
224   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE08]]
225   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
226   // CK1:     getelementptr inbounds float, ptr [[TTT]], i32 1
227   ++l; ++t;
228 
229 
230   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE09]]
231   // CK1:     [[_VAL:%.+]] = load ptr, ptr {{%.+}},
232   // CK1:     store ptr [[_VAL]], ptr [[_PVT:%.+]],
233   // CK1:     [[VAL:%.+]] = load ptr, ptr {{%.+}},
234   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
235   // CK1:     [[_TT1:%.+]] = load ptr, ptr [[_PVT]],
236   // CK1:     getelementptr inbounds float, ptr [[_TT1]], i32 1
237   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
238   // CK1:     getelementptr inbounds i32, ptr [[TT1]], i32 1
239   #pragma omp target data map(l[:10], t[:10]) use_device_ptr(l) use_device_ptr(t)
240   {
241     ++l; ++t;
242   }
243   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE09]]
244   // CK1:     [[_TTT:%.+]] = load ptr, ptr {{%.+}},
245   // CK1:     getelementptr inbounds float, ptr [[_TTT]], i32 1
246   // CK1:     [[TTT:%.+]] = load ptr, ptr {{%.+}},
247   // CK1:     getelementptr inbounds i32, ptr [[TTT]], i32 1
248   ++l; ++t;
249 
250   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE10]]
251   // CK1:     [[_VAL:%.+]] = load ptr, ptr {{%.+}},
252   // CK1:     store ptr [[_VAL]], ptr [[_PVT:%.+]],
253   // CK1:     [[VAL:%.+]] = load ptr, ptr {{%.+}},
254   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
255   // CK1:     [[_TT1:%.+]] = load ptr, ptr [[_PVT]],
256   // CK1:     getelementptr inbounds float, ptr [[_TT1]], i32 1
257   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
258   // CK1:     getelementptr inbounds i32, ptr [[TT1]], i32 1
259   #pragma omp target data map(l[:10], t[:10]) use_device_ptr(l,t)
260   {
261     ++l; ++t;
262   }
263   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE10]]
264   // CK1:     [[_TTT:%.+]] = load ptr, ptr {{%.+}},
265   // CK1:     getelementptr inbounds float, ptr [[_TTT]], i32 1
266   // CK1:     [[TTT:%.+]] = load ptr, ptr {{%.+}},
267   // CK1:     getelementptr inbounds i32, ptr [[TTT]], i32 1
268   ++l; ++t;
269 
270   // CK1:     [[T1:%.+]] = load ptr, ptr [[DECL:%.+]],
271   // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x ptr], ptr %{{.+}}, i32 0, i32 1
272   // CK1:     store ptr [[T1]], ptr [[BP]],
273   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE11]]
274   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
275   // CK1-NOT: store ptr [[VAL]], ptr [[DECL]],
276   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
277   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
278   // CK1:     getelementptr inbounds i32, ptr [[TT1]], i32 1
279   #pragma omp target data map(l[:10]) use_device_ptr(t)
280   {
281     ++l; ++t;
282   }
283   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE11]]
284   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
285   // CK1:     getelementptr inbounds i32, ptr [[TTT]], i32 1
286   ++l; ++t;
287 
288   // CK1:     [[T2:%.+]] = load ptr, ptr [[DECL:%.+]],
289   // CK1:     [[T1:%.+]] = load ptr, ptr [[T2]],
290   // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x ptr], ptr %{{.+}}, i32 0, i32 1
291   // CK1:     store ptr [[T1]], ptr [[BP]],
292   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE12]]
293   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
294   // CK1:     store ptr [[VAL]], ptr [[PVTV:%.+]],
295   // CK1-NOT: store ptr [[PVTV]], ptr [[DECL]],
296   // CK1:     store ptr [[PVTV]], ptr [[PVT:%.+]],
297   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
298   // CK1:     [[TT2:%.+]] = load ptr, ptr [[TT1]],
299   // CK1:     getelementptr inbounds i32, ptr [[TT2]], i32 1
300   #pragma omp target data map(l[:10]) use_device_ptr(tr)
301   {
302     ++l; ++tr;
303   }
304   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE12]]
305   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
306   // CK1:     [[TTTT:%.+]] = load ptr, ptr [[TTT]],
307   // CK1:     getelementptr inbounds i32, ptr [[TTTT]], i32 1
308   ++l; ++tr;
309 
310 }
311 
312 void bar(float *&a, int *&b) {
313   foo<int>(a,b);
314 }
315 
316 #endif
317 ///==========================================================================///
318 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64
319 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
320 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK2 --check-prefix CK2-64
321 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK2 --check-prefix CK2-32
322 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
323 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK2 --check-prefix CK2-32
324 
325 // RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY1 %s
326 // RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
327 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s
328 // RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY1 %s
329 // RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
330 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s
331 // SIMD-ONLY1-NOT: {{__kmpc|__tgt}}
332 #ifdef CK2
333 
334 // CK2: [[ST:%.+]] = type { ptr, ptr }
335 // CK2: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 0, i64 281474976710739]
336 // CK2: [[MTYPE01:@.+]] = {{.*}}constant [2 x i64] [i64 0, i64 281474976710739]
337 // CK2: [[MTYPE02:@.+]] = {{.*}}constant [3 x i64] [i64 3, i64 0, i64 562949953421392]
338 // CK2: [[MTYPE03:@.+]] = {{.*}}constant [3 x i64] [i64 0, i64 281474976710739, i64 281474976710736]
339 
340 template <typename T>
341 struct ST {
342   T *a;
343   double *&b;
344   ST(double *&b) : a(0), b(b) {}
345 
346   // CK2-LABEL: @{{.*}}foo{{.*}}
347   void foo(double *&arg) {
348     int *la = 0;
349 
350     // CK2:     [[BP:%.+]] = getelementptr inbounds [2 x ptr], ptr %{{.+}}, i32 0, i32 1
351     // CK2:     store ptr [[RVAL:%.+]], ptr [[BP]],
352     // CK2:     call void @__tgt_target_data_begin{{.+}}[[MTYPE00]]
353     // CK2:     [[VAL:%.+]] = load ptr, ptr [[BP]],
354     // CK2:     store ptr [[VAL]], ptr [[PVT:%.+]],
355     // CK2:     store ptr [[PVT]], ptr [[PVT2:%.+]],
356     // CK2:     [[TT1:%.+]] = load ptr, ptr [[PVT2]],
357     // CK2:     [[TT2:%.+]] = load ptr, ptr [[TT1]],
358     // CK2:     getelementptr inbounds double, ptr [[TT2]], i32 1
359     #pragma omp target data map(a[:10]) use_device_ptr(a)
360     {
361       a++;
362     }
363     // CK2:     call void @__tgt_target_data_end{{.+}}[[MTYPE00]]
364     // CK2:     [[DECL:%.+]] = getelementptr inbounds [[ST]], ptr %this1, i32 0, i32 0
365     // CK2:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
366     // CK2:     getelementptr inbounds double, ptr [[TTT]], i32 1
367     a++;
368 
369     // CK2:     [[BP:%.+]] = getelementptr inbounds [2 x ptr], ptr %{{.+}}, i32 0, i32 1
370     // CK2:     store ptr [[RVAL:%.+]], ptr [[BP]],
371     // CK2:     call void @__tgt_target_data_begin{{.+}}[[MTYPE01]]
372     // CK2:     [[VAL:%.+]] = load ptr, ptr [[BP]],
373     // CK2:     store ptr [[VAL]], ptr [[PVT:%.+]],
374     // CK2:     store ptr [[PVT]], ptr [[PVT2:%.+]],
375     // CK2:     [[TT1:%.+]] = load ptr, ptr [[PVT2]],
376     // CK2:     [[TT2:%.+]] = load ptr, ptr [[TT1]],
377     // CK2:     getelementptr inbounds double, ptr [[TT2]], i32 1
378     #pragma omp target data map(b[:10]) use_device_ptr(b)
379     {
380       b++;
381     }
382     // CK2:     call void @__tgt_target_data_end{{.+}}[[MTYPE01]]
383     // CK2:     [[DECL:%.+]] = getelementptr inbounds [[ST]], ptr %{{.+}}, i32 0, i32 1
384     // CK2:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
385     // CK2:     [[TTTT:%.+]] = load ptr, ptr [[TTT]],
386     // CK2:     getelementptr inbounds double, ptr [[TTTT]], i32 1
387     b++;
388 
389     // CK2:     [[BP:%.+]] = getelementptr inbounds [3 x ptr], ptr %{{.+}}, i32 0, i32 2
390     // CK2:     store ptr [[RVAL:%.+]], ptr [[BP]],
391     // CK2:     call void @__tgt_target_data_begin{{.+}}[[MTYPE02]]
392     // CK2:     [[VAL:%.+]] = load ptr, ptr [[BP]],
393     // CK2:     store ptr [[VAL]], ptr [[PVT:%.+]],
394     // CK2:     store ptr [[PVT]], ptr [[PVT2:%.+]],
395     // CK2:     [[TT1:%.+]] = load ptr, ptr [[PVT2]],
396     // CK2:     [[TT2:%.+]] = load ptr, ptr [[TT1]],
397     // CK2:     getelementptr inbounds double, ptr [[TT2]], i32 1
398     #pragma omp target data map(la[:10]) use_device_ptr(a)
399     {
400       a++;
401       la++;
402     }
403     // CK2:     call void @__tgt_target_data_end{{.+}}[[MTYPE02]]
404     // CK2:     [[DECL:%.+]] = getelementptr inbounds [[ST]], ptr %this1, i32 0, i32 0
405     // CK2:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
406     // CK2:     getelementptr inbounds double, ptr [[TTT]], i32 1
407     a++;
408     la++;
409 
410     // CK2:     [[BP1:%.+]] = getelementptr inbounds [3 x ptr], ptr %{{.+}}, i32 0, i32 1
411     // CK2:     store ptr [[RVAL1:%.+]], ptr [[BP1]],
412     // CK2:     [[BP2:%.+]] = getelementptr inbounds [3 x ptr], ptr %{{.+}}, i32 0, i32 2
413     // CK2:     store ptr [[RVAL2:%.+]], ptr [[BP2]],
414     // CK2:     call void @__tgt_target_data_begin{{.+}}[[MTYPE03]]
415     // CK2:     [[VAL2:%.+]] = load ptr, ptr [[BP2]],
416     // CK2:     store ptr [[VAL2]], ptr [[PVT2:%.+]],
417     // CK2:     store ptr [[PVT2]], ptr [[_PVT2:%.+]],
418     // CK2:     [[VAL1:%.+]] = load ptr, ptr [[BP1]],
419     // CK2:     store ptr [[VAL1]], ptr [[PVT1:%.+]],
420     // CK2:     store ptr [[PVT1]], ptr [[_PVT1:%.+]],
421     // CK2:     [[TT2:%.+]] = load ptr, ptr [[_PVT2]],
422     // CK2:     [[_TT2:%.+]] = load ptr, ptr [[TT2]],
423     // CK2:     getelementptr inbounds double, ptr [[_TT2]], i32 1
424     // CK2:     [[TT1:%.+]] = load ptr, ptr [[_PVT1]],
425     // CK2:     [[_TT1:%.+]] = load ptr, ptr [[TT1]],
426     // CK2:     getelementptr inbounds double, ptr [[_TT1]], i32 1
427     #pragma omp target data map(b[:10]) use_device_ptr(a, b)
428     {
429       a++;
430       b++;
431     }
432     // CK2:     call void @__tgt_target_data_end{{.+}}[[MTYPE03]]
433     // CK2:     [[DECL:%.+]] = getelementptr inbounds [[ST]], ptr %this1, i32 0, i32 0
434     // CK2:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
435     // CK2:     getelementptr inbounds double, ptr [[TTT]], i32 1
436     // CK2:     [[_DECL:%.+]] = getelementptr inbounds [[ST]], ptr %this1, i32 0, i32 1
437     // CK2:     [[_TTT:%.+]] = load ptr, ptr [[_DECL]],
438     // CK2:     [[_TTTT:%.+]] = load ptr, ptr [[_TTT]],
439     // CK2:     getelementptr inbounds double, ptr [[_TTTT]], i32 1
440     a++;
441     b++;
442   }
443 };
444 
445 void bar(double *arg){
446   ST<double> A(arg);
447   A.foo(arg);
448   ++arg;
449 }
450 #endif
451 #endif
452