1 // expected-no-diagnostics
2 #ifndef HEADER
3 #define HEADER
4 
5 ///==========================================================================///
6 // RUN: %clang_cc1 -no-opaque-pointers -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 -no-opaque-pointers -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 -no-opaque-pointers -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 -no-opaque-pointers -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 -no-opaque-pointers -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 -no-opaque-pointers -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 -no-opaque-pointers -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 -no-opaque-pointers -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 -no-opaque-pointers -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 -no-opaque-pointers -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 -no-opaque-pointers -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 -no-opaque-pointers -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: [[SIZES00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}]
26 // CK1: [[TYPES00:@.+]] = {{.+}}constant [1 x i64] [i64 288]
27 
28 // CK1: [[SIZES01:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]
29 // CK1: [[TYPES01:@.+]] = {{.+}}constant [1 x i64] [i64 288]
30 
31 // CK1: [[SIZES02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]
32 // CK1: [[TYPES02:@.+]] = {{.+}}constant [1 x i64] [i64 288]
33 
34 // CK1: [[SIZES03:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]
35 // CK1: [[TYPES03:@.+]] = {{.+}}constant [1 x i64] [i64 288]
36 
37 // CK1: [[SIZES04:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]
38 // CK1: [[TYPES04:@.+]] = {{.+}}constant [1 x i64] [i64 288]
39 
40 // CK1: [[SIZES05:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]
41 // CK1: [[TYPES05:@.+]] = {{.+}}constant [1 x i64] [i64 288]
42 
43 // CK1: [[SIZES06:@.+]] = {{.+}}constant [2 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] {{8|4}}]
44 // CK1: [[TYPES06:@.+]] = {{.+}}constant [2 x i64] [i64 288, i64 288]
45 
46 // CK1-LABEL: @_Z3foo{{.*}}(
47 template<typename T>
48 void foo(float *&lr, T *&tr) {
49   float *l;
50   T *t;
51 
52 // CK1-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]])
53 // CK1-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
54 // CK1-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]]
55 // CK1-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
56 // CK1-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]]
57 // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
58 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
59 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
60 // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
61 // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double**
62 // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double**
63 // CK1-DAG: store double* [[VAL:%.+]], double** [[CBP1]]
64 // CK1-DAG: store double* [[VAL]], double** [[CP1]]
65 // CK1-DAG: [[VAL]] = load double*, double** [[ADDR:@g]],
66 
67 // CK1: call void [[KERNEL:@.+]](double* [[VAL]])
68 #pragma omp target is_device_ptr(g)
69   {
70     ++g;
71   }
72 
73 // CK1-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]])
74 // CK1-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
75 // CK1-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]]
76 // CK1-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
77 // CK1-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]]
78 // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
79 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
80 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
81 // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
82 // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to float**
83 // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to float**
84 // CK1-DAG: store float* [[VAL:%.+]], float** [[CBP1]]
85 // CK1-DAG: store float* [[VAL]], float** [[CP1]]
86 // CK1-DAG: [[VAL]] = load float*, float** [[ADDR:%.+]],
87 
88 // CK1: call void [[KERNEL:@.+]](float* [[VAL]])
89 #pragma omp target is_device_ptr(l)
90   {
91     ++l;
92   }
93 
94 // CK1-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]])
95 // CK1-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
96 // CK1-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]]
97 // CK1-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
98 // CK1-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]]
99 // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
100 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
101 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
102 // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
103 // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32**
104 // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
105 // CK1-DAG: store i32* [[VAL:%.+]], i32** [[CBP1]]
106 // CK1-DAG: store i32* [[VAL]], i32** [[CP1]]
107 // CK1-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]],
108 
109 // CK1: call void [[KERNEL:@.+]](i32* [[VAL]])
110 #pragma omp target is_device_ptr(t)
111   {
112     ++t;
113   }
114 
115 // CK1-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]])
116 // CK1-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
117 // CK1-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]]
118 // CK1-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
119 // CK1-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]]
120 // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
121 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
122 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
123 // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
124 // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to float**
125 // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to float**
126 // CK1-DAG: store float* [[VAL:%.+]], float** [[CBP1]]
127 // CK1-DAG: store float* [[VAL]], float** [[CP1]]
128 // CK1-DAG: [[VAL]] = load float*, float** [[ADDR:%.+]],
129 // CK1-DAG: [[ADDR]] = load float**, float*** [[ADDR2:%.+]],
130 
131 // CK1: call void [[KERNEL:@.+]](float* [[VAL]])
132 #pragma omp target is_device_ptr(lr)
133   {
134     ++lr;
135   }
136 
137 // CK1-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]])
138 // CK1-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
139 // CK1-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]]
140 // CK1-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
141 // CK1-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]]
142 // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
143 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
144 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
145 // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
146 // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32**
147 // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
148 // CK1-DAG: store i32* [[VAL:%.+]], i32** [[CBP1]]
149 // CK1-DAG: store i32* [[VAL]], i32** [[CP1]]
150 // CK1-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]],
151 // CK1-DAG: [[ADDR]] = load i32**, i32*** [[ADDR2:%.+]],
152 
153 // CK1: call void [[KERNEL:@.+]](i32* [[VAL]])
154 #pragma omp target is_device_ptr(tr)
155   {
156     ++tr;
157   }
158 
159 // CK1-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]])
160 // CK1-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
161 // CK1-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]]
162 // CK1-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
163 // CK1-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]]
164 // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
165 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
166 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
167 // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
168 // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32**
169 // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
170 // CK1-DAG: store i32* [[VAL:%.+]], i32** [[CBP1]]
171 // CK1-DAG: store i32* [[VAL]], i32** [[CP1]]
172 // CK1-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]],
173 // CK1-DAG: [[ADDR]] = load i32**, i32*** [[ADDR2:%.+]],
174 
175 // CK1: call void [[KERNEL:@.+]](i32* [[VAL]])
176 #pragma omp target is_device_ptr(tr, lr)
177   {
178     ++tr;
179   }
180 
181 // CK1-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]])
182 // CK1-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
183 // CK1-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]]
184 // CK1-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
185 // CK1-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]]
186 // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
187 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
188 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
189 // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
190 // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32**
191 // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
192 // CK1-DAG: store i32* [[VAL:%.+]], i32** [[CBP1]]
193 // CK1-DAG: store i32* [[VAL]], i32** [[CP1]]
194 // CK1-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]],
195 // CK1-DAG: [[ADDR]] = load i32**, i32*** [[ADDR2:%.+]],
196 
197 // CK1-DAG: [[_BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1
198 // CK1-DAG: [[_P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1
199 // CK1-DAG: [[_CBP1:%.+]] = bitcast i8** [[_BP1]] to float**
200 // CK1-DAG: [[_CP1:%.+]] = bitcast i8** [[_P1]] to float**
201 // CK1-DAG: store float* [[_VAL:%.+]], float** [[_CBP1]]
202 // CK1-DAG: store float* [[_VAL]], float** [[_CP1]]
203 // CK1-DAG: [[_VAL]] = load float*, float** [[_ADDR:%.+]],
204 // CK1-DAG: [[_ADDR]] = load float**, float*** [[_ADDR2:%.+]],
205 
206 // CK1: call void [[KERNEL:@.+]](i32* [[VAL]], float* [[_VAL]])
207 #pragma omp target is_device_ptr(tr, lr)
208   {
209     ++tr,++lr;
210   }
211 }
212 
213 void bar(float *&a, int *&b) {
214   foo<int>(a,b);
215 }
216 
217 #endif
218 ///==========================================================================///
219 // RUN: %clang_cc1 -no-opaque-pointers -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
220 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
221 // RUN: %clang_cc1 -no-opaque-pointers -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
222 // RUN: %clang_cc1 -no-opaque-pointers -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
223 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
224 // RUN: %clang_cc1 -no-opaque-pointers -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
225 
226 // RUN: %clang_cc1 -no-opaque-pointers -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
227 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
228 // RUN: %clang_cc1 -no-opaque-pointers -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
229 // RUN: %clang_cc1 -no-opaque-pointers -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
230 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
231 // RUN: %clang_cc1 -no-opaque-pointers -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
232 // SIMD-ONLY1-NOT: {{__kmpc|__tgt}}
233 #ifdef CK2
234 
235 // CK2: [[ST:%.+]] = type { double*, double** }
236 
237 template <typename T>
238 struct ST {
239   T *a;
240   double *&b;
241   ST(double *&b) : a(0), b(b) {}
242 
243   // CK2-LABEL: @{{.*}}foo{{.*}}
244   void foo(double *&arg) {
245     int *la = 0;
246 
247 // CK2-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]])
248 // CK2-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
249 // CK2-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]]
250 // CK2-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
251 // CK2-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]]
252 // CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
253 // CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
254 
255 // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
256 // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
257 // CK2-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
258 // CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]**
259 // CK2-DAG: store [[ST]]* [[VAR0:%.+]], [[ST]]** [[CBP0]]
260 // CK2-DAG: store [[ST]]* [[VAR0]], [[ST]]** [[CP0]]
261 #pragma omp target is_device_ptr(a)
262     {
263       a++;
264     }
265 
266 // CK2-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]])
267 // CK2-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
268 // CK2-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]]
269 // CK2-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
270 // CK2-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]]
271 // CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
272 // CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
273 
274 // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
275 // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
276 // CK2-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
277 // CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]**
278 // CK2-DAG: store [[ST]]* [[VAR0:%.+]], [[ST]]** [[CBP0]]
279 // CK2-DAG: store [[ST]]* [[VAR0]], [[ST]]** [[CP0]]
280 #pragma omp target is_device_ptr(b)
281     {
282       b++;
283     }
284 
285 // CK2-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]])
286 // CK2-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
287 // CK2-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]]
288 // CK2-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
289 // CK2-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]]
290 // CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
291 // CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
292 
293 // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
294 // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
295 // CK2-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
296 // CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]**
297 // CK2-DAG: store [[ST]]* [[VAR0:%.+]], [[ST]]** [[CBP0]]
298 // CK2-DAG: store [[ST]]* [[VAR0]], [[ST]]** [[CP0]]
299 #pragma omp target is_device_ptr(a, b)
300     {
301       a++;
302       b++;
303     }
304   }
305 };
306 
307 void bar(double *arg){
308   ST<double> A(arg);
309   A.foo(arg);
310   ++arg;
311 }
312 #endif
313 ///==========================================================================///
314 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-64
315 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
316 // RUN: %clang_cc1 -no-opaque-pointers -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 CK3 --check-prefix CK3-64
317 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK3 --check-prefix CK3-32
318 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
319 // RUN: %clang_cc1 -no-opaque-pointers -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 CK3 --check-prefix CK3-32
320 
321 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -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
322 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
323 // RUN: %clang_cc1 -no-opaque-pointers -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
324 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -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
325 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
326 // RUN: %clang_cc1 -no-opaque-pointers -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
327 // SIMD-ONLY1-NOT: {{__kmpc|__tgt}}
328 #ifdef CK3
329 
330 // CK3-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[SZ:64|32]]] [i{{64|32}} {{8|4}}]
331 // OMP_MAP_TARGET_PARAM = 0x20 | OMP_MAP_TO = 0x1 = 0x21
332 // CK3-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x21]]]
333 void bar() {
334   __attribute__((aligned(64))) double *ptr;
335   // CK3-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]])
336   // CK3-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
337   // CK3-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]]
338   // CK3-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
339   // CK3-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]]
340   // CK3-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
341   // CK3-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
342   // CK3-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
343   // CK3-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
344   // CK3-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double***
345   // CK3-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double***
346   // CK3-DAG: store double** [[PTR:%.+]], double*** [[CBP1]]
347   // CK3-DAG: store double** [[PTR]], double*** [[CP1]]
348 
349   // CK3: call void [[KERNEL:@.+]](double** [[PTR]])
350 #pragma omp target is_device_ptr(ptr)
351   *ptr = 0;
352 }
353 #endif
354 #endif
355