xref: /llvm-project/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp (revision 7eca38ce76d5d1915f4ab7e665964062c0b37697) 
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 nuw 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 nuw 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 nuw 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 nuw 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 nuw 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 nuw 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 nuw 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 nuw 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 nuw 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 nuw 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:     br i1 [[CMP]], label %[[BTHEN:.+]], label %[[BELSE:.+]]
135 
136   // CK1:     [[BTHEN]]:
137   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE04]]
138   // CK1:     br label %[[BEND:.+]]
139 
140   // CK1:     [[BELSE]]:
141   // CK1:     br label %[[BEND]]
142 
143   // CK1:     [[BEND]]:
144   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
145   // CK1:     getelementptr inbounds nuw float, ptr [[TTT]], i32 1
146   ++l;
147 
148   // CK1:     [[T2:%.+]] = load ptr, ptr [[DECL:%.+]],
149   // CK1:     [[T1:%.+]] = load ptr, ptr [[T2]],
150   // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x ptr], ptr %{{.+}}, i32 0, i32 0
151   // CK1:     store ptr [[T1]], ptr [[BP]],
152   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE05]]
153   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
154   // CK1:     store ptr [[VAL]], ptr [[PVTV:%.+]],
155   // CK1-NOT: store ptr [[PVTV]], ptr [[DECL]],
156   // CK1:     store ptr [[PVTV]], ptr [[PVT:%.+]],
157   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
158   // CK1:     [[TT2:%.+]] = load ptr, ptr [[TT1]],
159   // CK1:     getelementptr inbounds nuw float, ptr [[TT2]], i32 1
160   #pragma omp target data map(lr[:10]) use_device_ptr(lr)
161   {
162     ++lr;
163   }
164   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE05]]
165   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
166   // CK1:     [[TTTT:%.+]] = load ptr, ptr [[TTT]],
167   // CK1:     getelementptr inbounds nuw float, ptr [[TTTT]], i32 1
168   ++lr;
169 
170   // CK1:     [[T1:%.+]] = load ptr, ptr [[DECL:%.+]],
171   // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x ptr], ptr %{{.+}}, i32 0, i32 0
172   // CK1:     store ptr [[T1]], ptr [[BP]],
173   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE06]]
174   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
175   // CK1-NOT: store ptr [[VAL]], ptr [[DECL]],
176   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
177   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
178   // CK1:     getelementptr inbounds nuw i32, ptr [[TT1]], i32 1
179   #pragma omp target data map(t[:10]) use_device_ptr(t)
180   {
181     ++t;
182   }
183   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE06]]
184   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
185   // CK1:     getelementptr inbounds nuw i32, ptr [[TTT]], i32 1
186   ++t;
187 
188   // CK1:     [[T2:%.+]] = load ptr, ptr [[DECL:%.+]],
189   // CK1:     [[T1:%.+]] = load ptr, ptr [[T2]],
190   // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x ptr], ptr %{{.+}}, i32 0, i32 0
191   // CK1:     store ptr [[T1]], ptr [[BP]],
192   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE07]]
193   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
194   // CK1:     store ptr [[VAL]], ptr [[PVTV:%.+]],
195   // CK1-NOT: store ptr [[PVTV]], ptr [[DECL]],
196   // CK1:     store ptr [[PVTV]], ptr [[PVT:%.+]],
197   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
198   // CK1:     [[TT2:%.+]] = load ptr, ptr [[TT1]],
199   // CK1:     getelementptr inbounds nuw i32, ptr [[TT2]], i32 1
200   #pragma omp target data map(tr[:10]) use_device_ptr(tr)
201   {
202     ++tr;
203   }
204   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE07]]
205   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
206   // CK1:     [[TTTT:%.+]] = load ptr, ptr [[TTT]],
207   // CK1:     getelementptr inbounds nuw i32, ptr [[TTTT]], i32 1
208   ++tr;
209 
210   // CK1:     [[T1:%.+]] = load ptr, ptr [[DECL:%.+]],
211   // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x ptr], ptr %{{.+}}, i32 0, i32 0
212   // CK1:     store ptr [[T1]], ptr [[BP]],
213   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE08]]
214   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
215   // CK1-NOT: store ptr [[VAL]], ptr [[DECL]],
216   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
217   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
218   // CK1:     getelementptr inbounds nuw float, ptr [[TT1]], i32 1
219   #pragma omp target data map(l[:10], t[:10]) use_device_ptr(l)
220   {
221     ++l; ++t;
222   }
223   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE08]]
224   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
225   // CK1:     getelementptr inbounds nuw float, ptr [[TTT]], i32 1
226   ++l; ++t;
227 
228 
229   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE09]]
230   // CK1:     [[_VAL:%.+]] = load ptr, ptr {{%.+}},
231   // CK1:     store ptr [[_VAL]], ptr [[_PVT:%.+]],
232   // CK1:     [[VAL:%.+]] = load ptr, ptr {{%.+}},
233   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
234   // CK1:     [[_TT1:%.+]] = load ptr, ptr [[_PVT]],
235   // CK1:     getelementptr inbounds nuw float, ptr [[_TT1]], i32 1
236   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
237   // CK1:     getelementptr inbounds nuw i32, ptr [[TT1]], i32 1
238   #pragma omp target data map(l[:10], t[:10]) use_device_ptr(l) use_device_ptr(t)
239   {
240     ++l; ++t;
241   }
242   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE09]]
243   // CK1:     [[_TTT:%.+]] = load ptr, ptr {{%.+}},
244   // CK1:     getelementptr inbounds nuw float, ptr [[_TTT]], i32 1
245   // CK1:     [[TTT:%.+]] = load ptr, ptr {{%.+}},
246   // CK1:     getelementptr inbounds nuw i32, ptr [[TTT]], i32 1
247   ++l; ++t;
248 
249   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE10]]
250   // CK1:     [[_VAL:%.+]] = load ptr, ptr {{%.+}},
251   // CK1:     store ptr [[_VAL]], ptr [[_PVT:%.+]],
252   // CK1:     [[VAL:%.+]] = load ptr, ptr {{%.+}},
253   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
254   // CK1:     [[_TT1:%.+]] = load ptr, ptr [[_PVT]],
255   // CK1:     getelementptr inbounds nuw float, ptr [[_TT1]], i32 1
256   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
257   // CK1:     getelementptr inbounds nuw i32, ptr [[TT1]], i32 1
258   #pragma omp target data map(l[:10], t[:10]) use_device_ptr(l,t)
259   {
260     ++l; ++t;
261   }
262   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE10]]
263   // CK1:     [[_TTT:%.+]] = load ptr, ptr {{%.+}},
264   // CK1:     getelementptr inbounds nuw float, ptr [[_TTT]], i32 1
265   // CK1:     [[TTT:%.+]] = load ptr, ptr {{%.+}},
266   // CK1:     getelementptr inbounds nuw i32, ptr [[TTT]], i32 1
267   ++l; ++t;
268 
269   // CK1:     [[T1:%.+]] = load ptr, ptr [[DECL:%.+]],
270   // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x ptr], ptr %{{.+}}, i32 0, i32 1
271   // CK1:     store ptr [[T1]], ptr [[BP]],
272   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE11]]
273   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
274   // CK1-NOT: store ptr [[VAL]], ptr [[DECL]],
275   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
276   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
277   // CK1:     getelementptr inbounds nuw i32, ptr [[TT1]], i32 1
278   #pragma omp target data map(l[:10]) use_device_ptr(t)
279   {
280     ++l; ++t;
281   }
282   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE11]]
283   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
284   // CK1:     getelementptr inbounds nuw i32, ptr [[TTT]], i32 1
285   ++l; ++t;
286 
287   // CK1:     [[T2:%.+]] = load ptr, ptr [[DECL:%.+]],
288   // CK1:     [[T1:%.+]] = load ptr, ptr [[T2]],
289   // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x ptr], ptr %{{.+}}, i32 0, i32 1
290   // CK1:     store ptr [[T1]], ptr [[BP]],
291   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE12]]
292   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
293   // CK1:     store ptr [[VAL]], ptr [[PVTV:%.+]],
294   // CK1-NOT: store ptr [[PVTV]], ptr [[DECL]],
295   // CK1:     store ptr [[PVTV]], ptr [[PVT:%.+]],
296   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
297   // CK1:     [[TT2:%.+]] = load ptr, ptr [[TT1]],
298   // CK1:     getelementptr inbounds nuw i32, ptr [[TT2]], i32 1
299   #pragma omp target data map(l[:10]) use_device_ptr(tr)
300   {
301     ++l; ++tr;
302   }
303   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE12]]
304   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
305   // CK1:     [[TTTT:%.+]] = load ptr, ptr [[TTT]],
306   // CK1:     getelementptr inbounds nuw i32, ptr [[TTTT]], i32 1
307   ++l; ++tr;
308 
309 }
310 
311 void bar(float *&a, int *&b) {
312   foo<int>(a,b);
313 }
314 
315 #endif
316 ///==========================================================================///
317 // 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
318 // 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
319 // 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
320 // 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
321 // 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
322 // 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
323 
324 // 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
325 // 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
326 // 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
327 // 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
328 // 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
329 // 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
330 // SIMD-ONLY1-NOT: {{__kmpc|__tgt}}
331 #ifdef CK2
332 
333 // CK2: [[ST:%.+]] = type { ptr, ptr }
334 // CK2: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 0, i64 281474976710739]
335 // CK2: [[MTYPE01:@.+]] = {{.*}}constant [2 x i64] [i64 0, i64 281474976710739]
336 // CK2: [[MTYPE02:@.+]] = {{.*}}constant [3 x i64] [i64 3, i64 0, i64 562949953421392]
337 // CK2: [[MTYPE03:@.+]] = {{.*}}constant [3 x i64] [i64 0, i64 281474976710739, i64 281474976710736]
338 
339 template <typename T>
340 struct ST {
341   T *a;
342   double *&b;
343   ST(double *&b) : a(0), b(b) {}
344 
345   // CK2-LABEL: @{{.*}}foo{{.*}}
346   void foo(double *&arg) {
347     int *la = 0;
348 
349     // CK2:     [[BP:%.+]] = getelementptr inbounds [2 x ptr], ptr %{{.+}}, i32 0, i32 1
350     // CK2:     store ptr [[RVAL:%.+]], ptr [[BP]],
351     // CK2:     call void @__tgt_target_data_begin{{.+}}[[MTYPE00]]
352     // CK2:     [[VAL:%.+]] = load ptr, ptr [[BP]],
353     // CK2:     store ptr [[VAL]], ptr [[PVT:%.+]],
354     // CK2:     store ptr [[PVT]], ptr [[PVT2:%.+]],
355     // CK2:     [[TT1:%.+]] = load ptr, ptr [[PVT2]],
356     // CK2:     [[TT2:%.+]] = load ptr, ptr [[TT1]],
357     // CK2:     getelementptr inbounds nuw double, ptr [[TT2]], i32 1
358     #pragma omp target data map(a[:10]) use_device_ptr(a)
359     {
360       a++;
361     }
362     // CK2:     call void @__tgt_target_data_end{{.+}}[[MTYPE00]]
363     // CK2:     [[DECL:%.+]] = getelementptr inbounds nuw [[ST]], ptr %this1, i32 0, i32 0
364     // CK2:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
365     // CK2:     getelementptr inbounds nuw double, ptr [[TTT]], i32 1
366     a++;
367 
368     // CK2:     [[BP:%.+]] = getelementptr inbounds [2 x ptr], ptr %{{.+}}, i32 0, i32 1
369     // CK2:     store ptr [[RVAL:%.+]], ptr [[BP]],
370     // CK2:     call void @__tgt_target_data_begin{{.+}}[[MTYPE01]]
371     // CK2:     [[VAL:%.+]] = load ptr, ptr [[BP]],
372     // CK2:     store ptr [[VAL]], ptr [[PVT:%.+]],
373     // CK2:     store ptr [[PVT]], ptr [[PVT2:%.+]],
374     // CK2:     [[TT1:%.+]] = load ptr, ptr [[PVT2]],
375     // CK2:     [[TT2:%.+]] = load ptr, ptr [[TT1]],
376     // CK2:     getelementptr inbounds nuw double, ptr [[TT2]], i32 1
377     #pragma omp target data map(b[:10]) use_device_ptr(b)
378     {
379       b++;
380     }
381     // CK2:     call void @__tgt_target_data_end{{.+}}[[MTYPE01]]
382     // CK2:     [[DECL:%.+]] = getelementptr inbounds nuw [[ST]], ptr %{{.+}}, i32 0, i32 1
383     // CK2:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
384     // CK2:     [[TTTT:%.+]] = load ptr, ptr [[TTT]],
385     // CK2:     getelementptr inbounds nuw double, ptr [[TTTT]], i32 1
386     b++;
387 
388     // CK2:     [[BP:%.+]] = getelementptr inbounds [3 x ptr], ptr %{{.+}}, i32 0, i32 2
389     // CK2:     store ptr [[RVAL:%.+]], ptr [[BP]],
390     // CK2:     call void @__tgt_target_data_begin{{.+}}[[MTYPE02]]
391     // CK2:     [[VAL:%.+]] = load ptr, ptr [[BP]],
392     // CK2:     store ptr [[VAL]], ptr [[PVT:%.+]],
393     // CK2:     store ptr [[PVT]], ptr [[PVT2:%.+]],
394     // CK2:     [[TT1:%.+]] = load ptr, ptr [[PVT2]],
395     // CK2:     [[TT2:%.+]] = load ptr, ptr [[TT1]],
396     // CK2:     getelementptr inbounds nuw double, ptr [[TT2]], i32 1
397     #pragma omp target data map(la[:10]) use_device_ptr(a)
398     {
399       a++;
400       la++;
401     }
402     // CK2:     call void @__tgt_target_data_end{{.+}}[[MTYPE02]]
403     // CK2:     [[DECL:%.+]] = getelementptr inbounds nuw [[ST]], ptr %this1, i32 0, i32 0
404     // CK2:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
405     // CK2:     getelementptr inbounds nuw double, ptr [[TTT]], i32 1
406     a++;
407     la++;
408 
409     // CK2:     [[BP1:%.+]] = getelementptr inbounds [3 x ptr], ptr %{{.+}}, i32 0, i32 1
410     // CK2:     store ptr [[RVAL1:%.+]], ptr [[BP1]],
411     // CK2:     [[BP2:%.+]] = getelementptr inbounds [3 x ptr], ptr %{{.+}}, i32 0, i32 2
412     // CK2:     store ptr [[RVAL2:%.+]], ptr [[BP2]],
413     // CK2:     call void @__tgt_target_data_begin{{.+}}[[MTYPE03]]
414     // CK2:     [[VAL1:%.+]] = load ptr, ptr [[BP1]],
415     // CK2:     store ptr [[VAL1]], ptr [[PVT1:%.+]],
416     // CK2:     [[VAL2:%.+]] = load ptr, ptr [[BP2]],
417     // CK2:     store ptr [[VAL2]], ptr [[PVT2:%.+]],
418     // CK2:     store ptr [[PVT2]], ptr [[_PVT2:%.+]],
419     // CK2:     store ptr [[PVT1]], ptr [[_PVT1:%.+]],
420     // CK2:     [[TT2:%.+]] = load ptr, ptr [[_PVT2]],
421     // CK2:     [[_TT2:%.+]] = load ptr, ptr [[TT2]],
422     // CK2:     getelementptr inbounds nuw double, ptr [[_TT2]], i32 1
423     // CK2:     [[TT1:%.+]] = load ptr, ptr [[_PVT1]],
424     // CK2:     [[_TT1:%.+]] = load ptr, ptr [[TT1]],
425     // CK2:     getelementptr inbounds nuw double, ptr [[_TT1]], i32 1
426     #pragma omp target data map(b[:10]) use_device_ptr(a, b)
427     {
428       a++;
429       b++;
430     }
431     // CK2:     call void @__tgt_target_data_end{{.+}}[[MTYPE03]]
432     // CK2:     [[DECL:%.+]] = getelementptr inbounds nuw [[ST]], ptr %this1, i32 0, i32 0
433     // CK2:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
434     // CK2:     getelementptr inbounds nuw double, ptr [[TTT]], i32 1
435     // CK2:     [[_DECL:%.+]] = getelementptr inbounds nuw [[ST]], ptr %this1, i32 0, i32 1
436     // CK2:     [[_TTT:%.+]] = load ptr, ptr [[_DECL]],
437     // CK2:     [[_TTTT:%.+]] = load ptr, ptr [[_TTT]],
438     // CK2:     getelementptr inbounds nuw double, ptr [[_TTTT]], i32 1
439     a++;
440     b++;
441   }
442 };
443 
444 void bar(double *arg){
445   ST<double> A(arg);
446   A.foo(arg);
447   ++arg;
448 }
449 #endif
450 #endif
451