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