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