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 | |
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> |
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 |
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 | |
197 | 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 | |
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 | |