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 #ifdef CK1
13 
14 double *g;
15 
16 // CK1: @g = global double*
17 // CK1: [[SIZES00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}]
18 // CK1: [[TYPES00:@.+]] = {{.+}}constant [1 x i32] [i32 288]
19 
20 // CK1: [[SIZES01:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]
21 // CK1: [[TYPES01:@.+]] = {{.+}}constant [1 x i32] [i32 288]
22 
23 // CK1: [[SIZES02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]
24 // CK1: [[TYPES02:@.+]] = {{.+}}constant [1 x i32] [i32 288]
25 
26 // CK1: [[SIZES03:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]
27 // CK1: [[TYPES03:@.+]] = {{.+}}constant [1 x i32] [i32 288]
28 
29 // CK1: [[SIZES04:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]
30 // CK1: [[TYPES04:@.+]] = {{.+}}constant [1 x i32] [i32 288]
31 
32 // CK1: [[SIZES05:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]
33 // CK1: [[TYPES05:@.+]] = {{.+}}constant [1 x i32] [i32 288]
34 
35 // CK1: [[SIZES06:@.+]] = {{.+}}constant [2 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] {{8|4}}]
36 // CK1: [[TYPES06:@.+]] = {{.+}}constant [2 x i32] [i32 288, i32 288]
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-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES00]]{{.+}}, {{.+}}[[TYPES00]]{{.+}})
45   // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
46   // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
47   // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
48   // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
49   // CK1-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
50   // CK1-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
51   // CK1-DAG: [[VALBP]] = bitcast double* [[VAL:%.+]] to i8*
52   // CK1-DAG: [[VALP]] = bitcast double* [[VAL]] to i8*
53   // CK1-DAG: [[VAL]] = load double*, double** [[ADDR:@g]],
54 
55   // CK1: call void [[KERNEL:@.+]](double* [[VAL]])
56   #pragma omp target is_device_ptr(g)
57   {
58     ++g;
59   }
60 
61   // CK1-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES01]]{{.+}}, {{.+}}[[TYPES01]]{{.+}})
62   // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
63   // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
64   // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
65   // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
66   // CK1-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
67   // CK1-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
68   // CK1-DAG: [[VALBP]] = bitcast float* [[VAL:%.+]] to i8*
69   // CK1-DAG: [[VALP]] = bitcast float* [[VAL]] to i8*
70   // CK1-DAG: [[VAL]] = load float*, float** [[ADDR:%.+]],
71 
72   // CK1: call void [[KERNEL:@.+]](float* [[VAL]])
73   #pragma omp target is_device_ptr(l)
74   {
75     ++l;
76   }
77 
78   // CK1-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES02]]{{.+}}, {{.+}}[[TYPES02]]{{.+}})
79   // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
80   // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
81   // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
82   // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
83   // CK1-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
84   // CK1-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
85   // CK1-DAG: [[VALBP]] = bitcast i32* [[VAL:%.+]] to i8*
86   // CK1-DAG: [[VALP]] = bitcast i32* [[VAL]] to i8*
87   // CK1-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]],
88 
89   // CK1: call void [[KERNEL:@.+]](i32* [[VAL]])
90   #pragma omp target is_device_ptr(t)
91   {
92     ++t;
93   }
94 
95   // CK1-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES03]]{{.+}}, {{.+}}[[TYPES03]]{{.+}})
96   // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
97   // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
98   // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
99   // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
100   // CK1-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
101   // CK1-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
102   // CK1-DAG: [[VALBP]] = bitcast float* [[VAL:%.+]] to i8*
103   // CK1-DAG: [[VALP]] = bitcast float* [[VAL]] to i8*
104   // CK1-DAG: [[VAL]] = load float*, float** [[ADDR:%.+]],
105   // CK1-DAG: [[ADDR]] = load float**, float*** [[ADDR2:%.+]],
106 
107   // CK1: call void [[KERNEL:@.+]](float* [[VAL]])
108   #pragma omp target is_device_ptr(lr)
109   {
110     ++lr;
111   }
112 
113   // CK1-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES04]]{{.+}}, {{.+}}[[TYPES04]]{{.+}})
114   // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
115   // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
116   // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
117   // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
118   // CK1-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
119   // CK1-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
120   // CK1-DAG: [[VALBP]] = bitcast i32* [[VAL:%.+]] to i8*
121   // CK1-DAG: [[VALP]] = bitcast i32* [[VAL]] to i8*
122   // CK1-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]],
123   // CK1-DAG: [[ADDR]] = load i32**, i32*** [[ADDR2:%.+]],
124 
125   // CK1: call void [[KERNEL:@.+]](i32* [[VAL]])
126   #pragma omp target is_device_ptr(tr)
127   {
128     ++tr;
129   }
130 
131   // CK1-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES05]]{{.+}}, {{.+}}[[TYPES05]]{{.+}})
132   // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
133   // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
134   // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
135   // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
136   // CK1-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
137   // CK1-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
138   // CK1-DAG: [[VALBP]] = bitcast i32* [[VAL:%.+]] to i8*
139   // CK1-DAG: [[VALP]] = bitcast i32* [[VAL]] to i8*
140   // CK1-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]],
141   // CK1-DAG: [[ADDR]] = load i32**, i32*** [[ADDR2:%.+]],
142 
143   // CK1: call void [[KERNEL:@.+]](i32* [[VAL]])
144   #pragma omp target is_device_ptr(tr,lr)
145   {
146     ++tr;
147   }
148 
149   // CK1-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES06]]{{.+}}, {{.+}}[[TYPES06]]{{.+}})
150   // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
151   // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
152   // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
153   // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
154   // CK1-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
155   // CK1-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
156   // CK1-DAG: [[VALBP]] = bitcast i32* [[VAL:%.+]] to i8*
157   // CK1-DAG: [[VALP]] = bitcast i32* [[VAL]] to i8*
158   // CK1-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]],
159   // CK1-DAG: [[ADDR]] = load i32**, i32*** [[ADDR2:%.+]],
160 
161   // CK1-DAG: [[_BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1
162   // CK1-DAG: [[_P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1
163   // CK1-DAG: store i8* [[_VALBP:%.+]], i8** [[_BP1]],
164   // CK1-DAG: store i8* [[_VALP:%.+]], i8** [[_P1]],
165   // CK1-DAG: [[_VALBP]] = bitcast float* [[_VAL:%.+]] to i8*
166   // CK1-DAG: [[_VALP]] = bitcast float* [[_VAL]] to i8*
167   // CK1-DAG: [[_VAL]] = load float*, float** [[_ADDR:%.+]],
168   // CK1-DAG: [[_ADDR]] = load float**, float*** [[_ADDR2:%.+]],
169 
170   // CK1: call void [[KERNEL:@.+]](i32* [[VAL]], float* [[_VAL]])
171   #pragma omp target is_device_ptr(tr,lr)
172   {
173     ++tr,++lr;
174   }
175 }
176 
177 void bar(float *&a, int *&b) {
178   foo<int>(a,b);
179 }
180 
181 #endif
182 ///==========================================================================///
183 // 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
184 // 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
185 // 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
186 // 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
187 // 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
188 // 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
189 #ifdef CK2
190 
191 // CK2: [[ST:%.+]] = type { double*, double** }
192 
193 // CK2: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}]
194 // CK2: [[MTYPE00:@.+]] = {{.+}}constant [1 x i32] [i32 33]
195 
196 // CK2: [[SIZE01:@.+]] = {{.+}}constant [2 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] {{8|4}}]
197 // CK2: [[MTYPE01:@.+]] = {{.+}}constant [2 x i32] [i32 32, i32 17]
198 
199 // CK2: [[SIZE02:@.+]] = {{.+}}constant [3 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] {{8|4}}, i[[sz]] {{8|4}}]
200 // CK2: [[MTYPE02:@.+]] = {{.+}}constant [3 x i32] [i32 33, i32 0, i32 17]
201 
202 template <typename T>
203 struct ST {
204   T *a;
205   double *&b;
206   ST(double *&b) : a(0), b(b) {}
207 
208   // CK2-LABEL: @{{.*}}foo{{.*}}
209   void foo(double *&arg) {
210     int *la = 0;
211 
212     // CK2-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
213     // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
214     // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
215 
216     // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
217     // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
218     // CK2-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
219     // CK2-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
220     // CK2-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[VAR0:%.+]] to i8*
221     // CK2-DAG: [[CPVAL0]] = bitcast double** [[SEC0:%.+]] to i8*
222     // CK2-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[VAR0]], i{{.+}} 0, i{{.+}} 0
223     #pragma omp target is_device_ptr(a)
224     {
225       a++;
226     }
227 
228     // CK2-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE01]]{{.+}})
229     // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
230     // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
231 
232     // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
233     // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
234     // CK2-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
235     // CK2-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
236     // CK2-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[VAR0:%.+]] to i8*
237     // CK2-DAG: [[CPVAL0]] = bitcast double*** [[SEC0:%.+]] to i8*
238     // CK2-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
239 
240     // CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
241     // CK2-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
242     // CK2-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
243     // CK2-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
244     // CK2-DAG: [[CBPVAL1]] = bitcast double*** [[SEC0]] to i8*
245     // CK2-DAG: [[CPVAL1]] = bitcast double** [[SEC1:%.+]] to i8*
246     // CK2-DAG: [[SEC1]] = load double**, double*** [[SEC0]]
247     #pragma omp target is_device_ptr(b)
248     {
249       b++;
250     }
251 
252     // CK2-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE02]]{{.+}})
253     // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
254     // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
255 
256     // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
257     // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
258     // CK2-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
259     // CK2-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
260     // CK2-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[VAR0:%.+]] to i8*
261     // CK2-DAG: [[CPVAL0]] = bitcast double*** [[SEC0:%.+]] to i8*
262     // CK2-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
263 
264     // CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
265     // CK2-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
266     // CK2-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
267     // CK2-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
268     // CK2-DAG: [[CBPVAL1]] = bitcast double*** [[SEC0]] to i8*
269     // CK2-DAG: [[CPVAL1]] = bitcast double** [[SEC1:%.+]] to i8*
270     // CK2-DAG: [[SEC1]] = load double**, double*** [[SEC0]]
271 
272     // CK2-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
273     // CK2-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
274     // CK2-DAG: store i8* [[CBPVAL2:%[^,]+]], i8** [[BP2]]
275     // CK2-DAG: store i8* [[CPVAL2:%[^,]+]], i8** [[P2]]
276     // CK2-DAG: [[CBPVAL2]] = bitcast [[ST]]* [[VAR2:%.+]] to i8*
277     // CK2-DAG: [[CPVAL2]] = bitcast double** [[SEC2:%.+]] to i8*
278     // CK2-DAG: [[SEC2]] = getelementptr {{.*}}[[ST]]* [[VAR2]], i{{.+}} 0, i{{.+}} 0
279     #pragma omp target is_device_ptr(a, b)
280     {
281       a++;
282       b++;
283     }
284   }
285 };
286 
287 void bar(double *arg){
288   ST<double> A(arg);
289   A.foo(arg);
290   ++arg;
291 }
292 #endif
293 #endif
294