1 // Only test codegen on target side, as private clause does not require any action on the host side
2 // Test target codegen - host bc file has to be created first.
3 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
4 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64
5 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
6 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64
7 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
8 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32
9 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
10 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32
11
12 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
13 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck --check-prefix SIMD-ONLY0 %s
14 // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
15 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
16 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
17 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck --check-prefix SIMD-ONLY0 %s
18 // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
19 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
20 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
21
22 // expected-no-diagnostics
23 #ifndef HEADER
24 #define HEADER
25
26 template<typename tx, typename ty>
27 struct TT{
28 tx X;
29 ty Y;
30 };
31
32 // TCHECK: [[TT:%.+]] = type { i64, i8 }
33 // TCHECK: [[S1:%.+]] = type { double }
34
foo(int n)35 int foo(int n) {
36 int a = 0;
37 short aa = 0;
38 float b[10];
39 float bn[n];
40 double c[5][10];
41 double cn[5][n];
42 TT<long long, char> d;
43
44 #pragma omp target private(a)
45 {
46 }
47
48 // TCHECK: define weak void @__omp_offloading_{{.+}}()
49 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
50 // TCHECK-NOT: store {{.+}}, {{.+}} [[A]],
51 // TCHECK: ret void
52
53 #pragma omp target private(a)
54 {
55 a = 1;
56 }
57
58 // TCHECK: define weak void @__omp_offloading_{{.+}}()
59 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
60 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]],
61 // TCHECK: ret void
62
63 #pragma omp target private(a, aa)
64 {
65 a = 1;
66 aa = 1;
67 }
68
69 // TCHECK: define weak void @__omp_offloading_{{.+}}()
70 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
71 // TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}},
72 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]],
73 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2]],
74 // TCHECK: ret void
75
76 #pragma omp target private(a, b, bn, c, cn, d)
77 {
78 a = 1;
79 b[2] = 1.0;
80 bn[3] = 1.0;
81 c[1][2] = 1.0;
82 cn[1][3] = 1.0;
83 d.X = 1;
84 d.Y = 1;
85 }
86 // make sure that private variables are generated in all cases and that we use those instances for operations inside the
87 // target region
88 // TCHECK: define weak void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[VLA:%.+]], i{{[0-9]+}} [[VLA1:%.+]], i{{[0-9]+}} [[VLA3:%.+]])
89 // TCHECK: [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}},
90 // TCHECK: [[VLA_ADDR2:%.+]] = alloca i{{[0-9]+}},
91 // TCHECK: [[VLA_ADDR4:%.+]] = alloca i{{[0-9]+}},
92 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
93 // TCHECK: [[B:%.+]] = alloca [10 x float],
94 // TCHECK: [[SSTACK:%.+]] = alloca i8*,
95 // TCHECK: [[C:%.+]] = alloca [5 x [10 x double]],
96 // TCHECK: [[D:%.+]] = alloca [[TT]],
97 // TCHECK: store i{{[0-9]+}} [[VLA]], i{{[0-9]+}}* [[VLA_ADDR]],
98 // TCHECK: store i{{[0-9]+}} [[VLA1]], i{{[0-9]+}}* [[VLA_ADDR2]],
99 // TCHECK: store i{{[0-9]+}} [[VLA3]], i{{[0-9]+}}* [[VLA_ADDR4]],
100 // TCHECK: [[VLA_ADDR_REF:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR]],
101 // TCHECK: [[VLA_ADDR_REF2:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR2]],
102 // TCHECK: [[VLA_ADDR_REF4:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR4]],
103 // TCHECK: [[RET_STACK:%.+]] = call i8* @llvm.stacksave()
104 // TCHECK: store i8* [[RET_STACK]], i8** [[SSTACK]],
105 // TCHECK: [[VLA5:%.+]] = alloca float, i{{[0-9]+}} [[VLA_ADDR_REF]],
106 // TCHECK: [[VLA6_SIZE:%.+]] = mul{{.+}} i{{[0-9]+}} [[VLA_ADDR_REF2]], [[VLA_ADDR_REF4]]
107 // TCHECK: [[VLA6:%.+]] = alloca double, i{{[0-9]+}} [[VLA6_SIZE]],
108
109 // a = 1
110 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]],
111
112 // b[2] = 1.0
113 // TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x float], [10 x float]* [[B]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
114 // TCHECK: store float 1.0{{.*}}, float* [[B_GEP]],
115
116 // bn[3] = 1.0
117 // TCHECK: [[BN_GEP:%.+]] = getelementptr inbounds float, float* [[VLA5]], i{{[0-9]+}} 3
118 // TCHECK: store float 1.0{{.*}}, float* [[BN_GEP]],
119
120 // c[1][2] = 1.0
121 // TCHECK: [[C_GEP1:%.+]] = getelementptr inbounds [5 x [10 x double]], [5 x [10 x double]]* [[C]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
122 // TCHECK: [[C_GEP2:%.+]] = getelementptr inbounds [10 x double], [10 x double]* [[C_GEP1]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
123 // TCHECK: store double 1.0{{.*}}, double* [[C_GEP2]],
124
125 // cn[1][3] = 1.0
126 // TCHECK: [[CN_IND:%.+]] = mul{{.+}} i{{[0-9]+}} 1, [[VLA_ADDR_REF4]]
127 // TCHECK: [[CN_GEP_IND:%.+]] = getelementptr inbounds double, double* [[VLA6]], i{{[0-9]+}} [[CN_IND]]
128 // TCHECK: [[CN_GEP_3:%.+]] = getelementptr inbounds double, double* [[CN_GEP_IND]], i{{[0-9]+}} 3
129 // TCHECK: store double 1.0{{.*}}, double* [[CN_GEP_3]],
130
131 // d.X = 1
132 // [[X_FIELD:%.+]] = getelementptr inbounds [[TT]] [[TT]]* [[D]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
133 // store i{{[0-9]+}} 1, i{{[0-9]+}}* [[X_FIELD]],
134
135 // d.Y = 1
136 // [[Y_FIELD:%.+]] = getelementptr inbounds [[TT]] [[TT]]* [[D]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
137 // store i{{[0-9]+}} 1, i{{[0-9]+}}* [[Y_FIELD]],
138
139 // finish
140 // [[RELOAD_SSTACK:%.+]] = load i8*, i8** [[SSTACK]],
141 // call ovid @llvm.stackrestore(i8* [[RELOAD_SSTACK]])
142 // ret void
143
144 return a;
145 }
146
147
148 template<typename tx>
ftemplate(int n)149 tx ftemplate(int n) {
150 tx a = 0;
151 short aa = 0;
152 tx b[10];
153
154 #pragma omp target private(a,aa,b)
155 {
156 a = 1;
157 aa = 1;
158 b[2] = 1;
159 }
160
161 return a;
162 }
163
164 static
fstatic(int n)165 int fstatic(int n) {
166 int a = 0;
167 short aa = 0;
168 char aaa = 0;
169 int b[10];
170
171 #pragma omp target private(a,aa,aaa,b)
172 {
173 a = 1;
174 aa = 1;
175 aaa = 1;
176 b[2] = 1;
177 }
178
179 return a;
180 }
181
182 // TCHECK: define weak void @__omp_offloading_{{.+}}()
183 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
184 // TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}},
185 // TCHECK: [[A3:%.+]] = alloca i{{[0-9]+}},
186 // TCHECK: [[B:%.+]] = alloca [10 x i{{[0-9]+}}],
187 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]],
188 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2]],
189 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A3]],
190 // TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[B]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
191 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]],
192 // TCHECK: ret void
193
194 struct S1 {
195 double a;
196
r1S1197 int r1(int n){
198 int b = n+1;
199 short int c[2][n];
200
201 #pragma omp target private(b,c)
202 {
203 this->a = (double)b + 1.5;
204 c[1][1] = ++a;
205 }
206
207 return c[1][1] + (int)b;
208 }
209
210 // TCHECK: define weak void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i{{[0-9]+}} [[VLA:%.+]], i{{[0-9]+}} [[VLA1:%.+]])
211 // TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*,
212 // TCHECK: [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}},
213 // TCHECK: [[VLA_ADDR2:%.+]] = alloca i{{[0-9]+}},
214 // TCHECK: [[B:%.+]] = alloca i{{[0-9]+}},
215 // TCHECK: [[SSTACK:%.+]] = alloca i8*,
216 // TCHECK: store [[S1]]* [[TH]], [[S1]]** [[TH_ADDR]],
217 // TCHECK: store i{{[0-9]+}} [[VLA]], i{{[0-9]+}}* [[VLA_ADDR]],
218 // TCHECK: store i{{[0-9]+}} [[VLA1]], i{{[0-9]+}}* [[VLA_ADDR2]],
219 // TCHECK: [[TH_ADDR_REF:%.+]] = load [[S1]]*, [[S1]]** [[TH_ADDR]],
220 // TCHECK: [[VLA_ADDR_REF:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR]],
221 // TCHECK: [[VLA_ADDR_REF2:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR2]],
222 // TCHECK: [[RET_STACK:%.+]] = call i8* @llvm.stacksave()
223 // TCHECK: store i8* [[RET_STACK:%.+]], i8** [[SSTACK]],
224
225 // this->a = (double)b + 1.5;
226 // TCHECK: [[VLA_IND:%.+]] = mul{{.+}} i{{[0-9]+}} [[VLA_ADDR_REF]], [[VLA_ADDR_REF2]]
227 // TCHECK: [[VLA3:%.+]] = alloca i{{[0-9]+}}, i{{[0-9]+}} [[VLA_IND]],
228 // TCHECK: [[B_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[B]],
229 // TCHECK: [[B_CONV:%.+]] = sitofp i{{[0-9]+}} [[B_VAL]] to double
230 // TCHECK: [[NEW_A_VAL:%.+]] = fadd double [[B_CONV]], 1.5{{.+}}+00
231 // TCHECK: [[A_FIELD:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
232 // TCHECK: store double [[NEW_A_VAL]], double* [[A_FIELD]],
233
234 // c[1][1] = ++a;
235 // TCHECK: [[A_FIELD4:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
236 // TCHECK: [[A_FIELD4_VAL:%.+]] = load double, double* [[A_FIELD4]],
237 // TCHECK: [[A_FIELD_INC:%.+]] = fadd double [[A_FIELD4_VAL]], 1.0{{.+}}+00
238 // TCHECK: store double [[A_FIELD_INC]], double* [[A_FIELD4]],
239 // TCHECK: [[A_FIELD_INC_CONV:%.+]] = fptosi double [[A_FIELD_INC]] to i{{[0-9]+}}
240 // TCHECK: [[C_IND:%.+]] = mul{{.+}} i{{[0-9]+}} 1, [[VLA_ADDR_REF2]]
241 // TCHECK: [[C_1_REF:%.+]] = getelementptr inbounds i{{[0-9]+}}, i{{[0-9]+}}* [[VLA3]], i{{[0-9]+}} [[C_IND]]
242 // TCHECK: [[C_1_1_REF:%.+]] = getelementptr inbounds i{{[0-9]+}}, i{{[0-9]+}}* [[C_1_REF]], i{{[0-9]+}} 1
243 // TCHECK: store i{{[0-9]+}} [[A_FIELD_INC_CONV]], i{{[0-9]+}}* [[C_1_1_REF]],
244
245 // finish
246 // TCHECK: [[RELOAD_SSTACK:%.+]] = load i8*, i8** [[SSTACK]],
247 // TCHECK: call void @llvm.stackrestore(i8* [[RELOAD_SSTACK]])
248 // TCHECK: ret void
249 };
250
251
bar(int n)252 int bar(int n){
253 int a = 0;
254 a += foo(n);
255 S1 S;
256 a += S.r1(n);
257 a += fstatic(n);
258 a += ftemplate<int>(n);
259
260 return a;
261 }
262
263 // template
264 // TCHECK: define weak void @__omp_offloading_{{.+}}()
265 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
266 // TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}},
267 // TCHECK: [[B:%.+]] = alloca [10 x i{{[0-9]+}}],
268 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]],
269 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2]],
270 // TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[B]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
271 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]],
272 // TCHECK: ret void
273
274 #endif
275