1 | // RUN: %clang_cc1 -DLAMBDA -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix LAMBDA --check-prefix LAMBDA-64 |
2 | // RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s |
3 | // RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix LAMBDA --check-prefix LAMBDA-64 |
4 | // RUN: %clang_cc1 -DLAMBDA -verify -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix LAMBDA --check-prefix LAMBDA-32 |
5 | // RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s |
6 | // RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix LAMBDA --check-prefix LAMBDA-32 |
7 | |
8 | // RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 |
9 | // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s |
10 | // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 |
11 | // RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 |
12 | // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s |
13 | // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 |
14 | |
15 | // RUN: %clang_cc1 -DLAMBDA -verify -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix SIMD-ONLY |
16 | // RUN: %clang_cc1 -DLAMBDA -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s |
17 | // RUN: %clang_cc1 -DLAMBDA -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY |
18 | // RUN: %clang_cc1 -DLAMBDA -verify -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix SIMD-ONLY |
19 | // RUN: %clang_cc1 -DLAMBDA -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s |
20 | // RUN: %clang_cc1 -DLAMBDA -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY |
21 | |
22 | // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix SIMD-ONLY |
23 | // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s |
24 | // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY |
25 | // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix SIMD-ONLY |
26 | // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s |
27 | // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY |
28 | // SIMD-ONLY-NOT: {{__kmpc|__tgt}} |
29 | |
30 | // expected-no-diagnostics |
31 | #ifndef HEADER |
32 | #define HEADER |
33 | |
34 | template <class T> |
35 | struct S { |
36 | T f; |
37 | S(T a) : f(a) {} |
38 | S() : f() {} |
39 | operator T() { return T(); } |
40 | ~S() {} |
41 | }; |
42 | |
43 | // CHECK: [[S_FLOAT_TY:%.+]] = type { float } |
44 | // CHECK: [[S_INT_TY:%.+]] = type { i{{[0-9]+}} } |
45 | template <typename T> |
46 | T tmain() { |
47 | S<T> test; |
48 | T t_var = T(); |
49 | T vec[] = {1, 2}; |
50 | S<T> s_arr[] = {1, 2}; |
51 | S<T> &var = test; |
52 | #pragma omp target teams distribute parallel for lastprivate(t_var, vec, s_arr, s_arr, var, var) |
53 | for (int i = 0; i < 2; ++i) { |
54 | vec[i] = t_var; |
55 | s_arr[i] = var; |
56 | } |
57 | return T(); |
58 | } |
59 | |
60 | int main() { |
61 | static int svar; |
62 | volatile double g; |
63 | volatile double &g1 = g; |
64 | |
65 | #ifdef LAMBDA |
66 | // LAMBDA-LABEL: @main |
67 | // LAMBDA: call{{.*}} void [[OUTER_LAMBDA:@.+]]( |
68 | [&]() { |
69 | static float sfvar; |
70 | // LAMBDA: define{{.*}} internal{{.*}} void [[OUTER_LAMBDA]]( |
71 | // LAMBDA: call i{{[0-9]+}} @__tgt_target_teams( |
72 | // LAMBDA: call void [[OFFLOADING_FUN:@.+]]( |
73 | |
74 | // LAMBDA: define{{.+}} void [[OFFLOADING_FUN]]( |
75 | // LAMBDA: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, i32 4, {{.+}}* [[OMP_OUTLINED:@.+]] to {{.+}}) |
76 | #pragma omp target teams distribute parallel for lastprivate(g, g1, svar, sfvar) |
77 | for (int i = 0; i < 2; ++i) { |
78 | // LAMBDA: define{{.*}} internal{{.*}} void [[OMP_OUTLINED]](i32* {{.+}}, i32* {{.+}}, {{.+}} [[G1_IN:%.+]], {{.+}} [[SVAR_IN:%.+]], {{.+}} [[SFVAR_IN:%.+]], {{.+}} [[G_IN:%.+]]) |
79 | // skip gbl and bound tid |
80 | // LAMBDA: alloca |
81 | // LAMBDA: alloca |
82 | // LAMBDA: [[G1_ADDR:%.+]] = alloca {{.+}}, |
83 | // LAMBDA: [[SVAR_ADDR:%.+]] = alloca {{.+}}, |
84 | // LAMBDA: [[SFVAR_ADDR:%.+]] = alloca {{.+}}, |
85 | // LAMBDA: [[G_ADDR:%.+]] = alloca {{.+}}, |
86 | // LAMBDA-64: [[G1_REF:%.+]] = alloca double*, |
87 | // loop variables |
88 | // LAMBDA: {{.+}} = alloca i{{[0-9]+}}, |
89 | // LAMBDA: {{.+}} = alloca i{{[0-9]+}}, |
90 | // LAMBDA: {{.+}} = alloca i{{[0-9]+}}, |
91 | // LAMBDA: {{.+}} = alloca i{{[0-9]+}}, |
92 | // LAMBDA: {{.+}} = alloca i{{[0-9]+}}, |
93 | // LAMBDA: [[OMP_IS_LAST:%.+]] = alloca i{{[0-9]+}}, |
94 | |
95 | // LAMBDA-DAG: store {{.+}} [[G_IN]], {{.+}} [[G_ADDR]], |
96 | // LAMBDA-DAG: store {{.+}} [[G1_IN]], {{.+}} [[G1_ADDR]], |
97 | // LAMBDA-DAG: store {{.+}} [[SVAR_IN]], {{.+}} [[SVAR_ADDR]], |
98 | // LAMBDA-DAG: store {{.+}} [[SFVAR_IN]], {{.+}} [[SFVAR_ADDR]], |
99 | |
100 | // LAMBDA-64-DAG: [[G_TGT:%.+]] = bitcast {{.+}} [[G_ADDR]] to |
101 | // LAMBDA-32-DAG: [[G_TGT:%.+]] = load {{.+}}, {{.+}} [[G_ADDR]], |
102 | // LAMBDA-64-DAG: [[G1_TGT:%.+]] = load {{.+}}, {{.+}} [[G1_REF]], |
103 | // LAMBDA-32-DAG: [[G1_TGT:%.+]] = load {{.+}}, {{.+}} [[G1_ADDR]], |
104 | // LAMBDA-64-DAG: [[SVAR_TGT:%.+]] = bitcast {{.+}} [[SVAR_ADDR]] to |
105 | // LAMBDA-DAG: [[SFVAR_TGT:%.+]] = bitcast {{.+}} [[SFVAR_ADDR]] to |
106 | |
107 | g1 = 1; |
108 | svar = 3; |
109 | sfvar = 4.0; |
110 | // LAMBDA: call {{.*}}void @__kmpc_for_static_init_4( |
111 | // LAMBDA: call void {{.*}} @__kmpc_fork_call({{.+}}, {{.+}}, {{.+}} @[[LPAR_OUTL:.+]] to |
112 | // LAMBDA: call {{.*}}void @__kmpc_for_static_fini( |
113 | |
114 | // LAMBDA: [[OMP_IS_LAST_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[OMP_IS_LAST]], |
115 | // LAMBDA: [[IS_LAST_IT:%.+]] = icmp ne i{{[0-9]+}} [[OMP_IS_LAST_VAL]], 0 |
116 | // LAMBDA: br i1 [[IS_LAST_IT]], label %[[OMP_LASTPRIV_BLOCK:.+]], label %[[OMP_LASTPRIV_DONE:.+]] |
117 | |
118 | // LAMBDA: [[OMP_LASTPRIV_BLOCK]]: |
119 | // LAMBDA-DAG: store {{.+}}, {{.+}} [[G_TGT]], |
120 | // LAMBDA-DAG: store {{.+}}, {{.+}} [[G1_TGT]], |
121 | // LAMBDA-64-DAG: store {{.+}}, {{.+}} [[SVAR_TGT]], |
122 | // LAMBDA-32-DAG: store {{.+}}, {{.+}} [[SVAR_ADDR]], |
123 | // LAMBDA-DAG: store {{.+}}, {{.+}} [[SFVAR_TGT]], |
124 | // LAMBDA: br label %[[OMP_LASTPRIV_DONE]] |
125 | // LAMBDA: [[OMP_LASTPRIV_DONE]]: |
126 | // LAMBDA: ret |
127 | |
128 | // LAMBDA: define{{.*}} internal{{.*}} void @[[LPAR_OUTL]](i32* noalias %{{.+}}, i32* noalias %{{.+}}, {{.+}}, {{.+}}, {{.+}} [[G1_IN:%.+]], {{.+}} [[SVAR_IN:%.+]], {{.+}} [[SFVAR_IN:%.+]], {{.+}} [[G_IN:%.+]]) |
129 | // skip tid and prev variables |
130 | // LAMBDA: alloca |
131 | // LAMBDA: alloca |
132 | // LAMBDA: alloca |
133 | // LAMBDA: alloca |
134 | // LAMBDA: [[G1_ADDR:%.+]] = alloca {{.+}}, |
135 | // LAMBDA: [[SVAR_ADDR:%.+]] = alloca {{.+}}, |
136 | // LAMBDA: [[SFVAR_ADDR:%.+]] = alloca {{.+}}, |
137 | // LAMBDA: [[G_ADDR:%.+]] = alloca {{.+}}, |
138 | // LAMBDA-64: [[G1_REF:%.+]] = alloca double*, |
139 | // loop variables |
140 | // LAMBDA: {{.+}} = alloca i{{[0-9]+}}, |
141 | // LAMBDA: {{.+}} = alloca i{{[0-9]+}}, |
142 | // LAMBDA: {{.+}} = alloca i{{[0-9]+}}, |
143 | // LAMBDA: {{.+}} = alloca i{{[0-9]+}}, |
144 | // LAMBDA: {{.+}} = alloca i{{[0-9]+}}, |
145 | // LAMBDA: [[OMP_IS_LAST:%.+]] = alloca i{{[0-9]+}}, |
146 | |
147 | // LAMBDA-DAG: store {{.+}} [[G_IN]], {{.+}} [[G_ADDR]], |
148 | // LAMBDA-DAG: store {{.+}} [[G1_IN]], {{.+}} [[G1_ADDR]], |
149 | // LAMBDA-DAG: store {{.+}} [[SVAR_IN]], {{.+}} [[SVAR_ADDR]], |
150 | // LAMBDA-DAG: store {{.+}} [[SFVAR_IN]], {{.+}} [[SFVAR_ADDR]], |
151 | |
152 | // LAMBDA-64-DAG: [[G_TGT:%.+]] = bitcast {{.+}} [[G_ADDR]] to |
153 | // LAMBDA-32-DAG: [[G_TGT:%.+]] = load {{.+}}, {{.+}} [[G_ADDR]], |
154 | // LAMBDA-64-DAG: [[G1_TGT:%.+]] = load {{.+}}, {{.+}} [[G1_REF]], |
155 | // LAMBDA-32-DAG: [[G1_TGT:%.+]] = load {{.+}}, {{.+}} [[G1_ADDR]], |
156 | // LAMBDA-64-DAG: [[SVAR_TGT:%.+]] = bitcast {{.+}} [[SVAR_ADDR]] to |
157 | // LAMBDA-DAG: [[SFVAR_TGT:%.+]] = bitcast {{.+}} [[SFVAR_ADDR]] to |
158 | |
159 | |
160 | // LAMBDA: call {{.*}}void @__kmpc_for_static_init_4( |
161 | // LAMBDA: call{{.*}} void [[INNER_LAMBDA:@.+]]( |
162 | // LAMBDA: call {{.*}}void @__kmpc_for_static_fini( |
163 | |
164 | // LAMBDA: [[OMP_IS_LAST_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[OMP_IS_LAST]], |
165 | // LAMBDA: [[IS_LAST_IT:%.+]] = icmp ne i{{[0-9]+}} [[OMP_IS_LAST_VAL]], 0 |
166 | // LAMBDA: br i1 [[IS_LAST_IT]], label %[[OMP_LASTPRIV_BLOCK:.+]], label %[[OMP_LASTPRIV_DONE:.+]] |
167 | |
168 | // LAMBDA: [[OMP_LASTPRIV_BLOCK]]: |
169 | // LAMBDA-DAG: store {{.+}}, {{.+}} [[G_TGT]], |
170 | // LAMBDA-DAG: store {{.+}}, {{.+}} [[G1_TGT]], |
171 | // LAMBDA-64-DAG: store {{.+}}, {{.+}} [[SVAR_TGT]], |
172 | // LAMBDA-32-DAG: store {{.+}}, {{.+}} [[SVAR_ADDR]], |
173 | // LAMBDA-DAG: store {{.+}}, {{.+}} [[SFVAR_TGT]], |
174 | // LAMBDA: br label %[[OMP_LASTPRIV_DONE]] |
175 | // LAMBDA: [[OMP_LASTPRIV_DONE]]: |
176 | // LAMBDA: ret |
177 | |
178 | [&]() { |
179 | // LAMBDA: define {{.+}} void [[INNER_LAMBDA]](%{{.+}}* [[ARG_PTR:%.+]]) |
180 | // LAMBDA: store %{{.+}}* [[ARG_PTR]], %{{.+}}** [[ARG_PTR_REF:%.+]], |
181 | g = 2; |
182 | g1 = 2; |
183 | svar = 4; |
184 | sfvar = 8.0; |
185 | // LAMBDA: [[ARG_PTR:%.+]] = load %{{.+}}*, %{{.+}}** [[ARG_PTR_REF]] |
186 | // LAMBDA: [[G_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 |
187 | // LAMBDA: [[G_REF:%.+]] = load double*, double** [[G_PTR_REF]] |
188 | // LAMBDA: store double 2.0{{.+}}, double* [[G_REF]] |
189 | |
190 | // LAMBDA: [[TMP_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 |
191 | // LAMBDA: [[G1_REF:%.+]] = load double*, double** [[TMP_PTR_REF]] |
192 | // LAMBDA: store double 2.0{{.+}}, double* [[G1_REF]], |
193 | // LAMBDA: [[SVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 |
194 | // LAMBDA: [[SVAR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[SVAR_PTR_REF]] |
195 | // LAMBDA: store i{{[0-9]+}} 4, i{{[0-9]+}}* [[SVAR_REF]] |
196 | // LAMBDA: [[SFVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 3 |
197 | // LAMBDA: [[SFVAR_REF:%.+]] = load float*, float** [[SFVAR_PTR_REF]] |
198 | // LAMBDA: store float 8.0{{.+}}, float* [[SFVAR_REF]] |
199 | }(); |
200 | } |
201 | }(); |
202 | return 0; |
203 | #else |
204 | S<float> test; |
205 | int t_var = 0; |
206 | int vec[] = {1, 2}; |
207 | S<float> s_arr[] = {1, 2}; |
208 | S<float> &var = test; |
209 | |
210 | #pragma omp target teams distribute parallel for lastprivate(t_var, vec, s_arr, s_arr, var, var, svar) |
211 | for (int i = 0; i < 2; ++i) { |
212 | vec[i] = t_var; |
213 | s_arr[i] = var; |
214 | } |
215 | int i; |
216 | |
217 | return tmain<int>(); |
218 | #endif |
219 | } |
220 | |
221 | // CHECK: define{{.*}} i{{[0-9]+}} @main() |
222 | // CHECK: [[TEST:%.+]] = alloca [[S_FLOAT_TY]], |
223 | // CHECK: call {{.*}} [[S_FLOAT_TY_DEF_CONSTR:@.+]]([[S_FLOAT_TY]]* [[TEST]]) |
224 | // CHECK: call i{{[0-9]+}} @__tgt_target_teams( |
225 | // CHECK: call void [[OFFLOAD_FUN:@.+]]( |
226 | // CHECK: ret |
227 | |
228 | // CHECK: define{{.+}} [[OFFLOAD_FUN]]( |
229 | // CHECK: call void (%{{.+}}*, i{{[0-9]+}}, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)*, ...) @__kmpc_fork_teams( |
230 | // CHECK: ret |
231 | // |
232 | // CHECK: define internal void [[OMP_OUTLINED:@.+]](i{{[0-9]+}}* noalias [[GTID_ADDR:%.+]], i{{[0-9]+}}* noalias %{{.+}}, [2 x i{{[0-9]+}}]*{{.+}} [[VEC_IN:%.+]], i{{[0-9]+}}{{.+}} [[T_VAR_IN:%.+]], [2 x [[S_FLOAT_TY]]]*{{.+}} [[S_ARR_IN:%.+]], [[S_FLOAT_TY]]*{{.+}} [[VAR_IN:%.+]], i{{[0-9]+}}{{.*}} [[S_VAR_IN:%.+]]) |
233 | // CHECK: {{.+}} = alloca i{{[0-9]+}}*, |
234 | // CHECK: {{.+}} = alloca i{{[0-9]+}}*, |
235 | // CHECK: [[VEC_ADDR:%.+]] = alloca [2 x i{{[0-9]+}}]*, |
236 | // CHECK: [[T_VAR_ADDR:%.+]] = alloca i{{[0-9]+}}, |
237 | // CHECK: [[S_ARR_ADDR:%.+]] = alloca [2 x [[S_FLOAT_TY]]]*, |
238 | // CHECK: [[VAR_ADDR:%.+]] = alloca [[S_FLOAT_TY]]*, |
239 | // CHECK: [[SVAR_ADDR:%.+]] = alloca i{{[0-9]+}}, |
240 | // skip loop variables |
241 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
242 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
243 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
244 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
245 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
246 | // CHECK: [[OMP_IS_LAST:%.+]] = alloca i{{[0-9]+}}, |
247 | |
248 | // copy from parameters to local address variables |
249 | // CHECK: store {{.+}} [[VEC_IN]], {{.+}} [[VEC_ADDR]], |
250 | // CHECK: store {{.+}} [[T_VAR_IN]], {{.+}} [[T_VAR_ADDR]], |
251 | // CHECK: store {{.+}} [[S_ARR_IN]], {{.+}} [[S_ARR_ADDR]], |
252 | // CHECK: store {{.+}} [[VAR_IN]], {{.+}} [[VAR_ADDR]], |
253 | // CHECK: store {{.+}} [[S_VAR_IN]], {{.+}} [[SVAR_ADDR]], |
254 | |
255 | // prepare lastprivate targets |
256 | // CHECK-64-DAG: [[TVAR_TGT:%.+]] = bitcast {{.+}} [[T_VAR_ADDR]] to |
257 | // CHECK-DAG: [[VEC_TGT:%.+]] = load {{.+}}, {{.+}} [[VEC_ADDR]], |
258 | // CHECK-DAG: [[S_ARR_TGT:%.+]] = load {{.+}}, {{.+}} [[S_ARR_ADDR]], |
259 | // CHECK-DAG: [[VAR_TGT:%.+]] = load {{.+}}, {{.+}} [[VAR_ADDR]], |
260 | // CHECK-64-DAG: [[SVAR_TGT:%.+]] = bitcast {{.+}} [[SVAR_ADDR]] to |
261 | |
262 | // the distribute loop |
263 | // CHECK: call void @__kmpc_for_static_init_4( |
264 | // CHECK: call void {{.*}} @__kmpc_fork_call({{.+}}, {{.+}}, {{.+}} @[[PAR_OUTL:.+]] to |
265 | // CHECK: call void @__kmpc_for_static_fini( |
266 | |
267 | // lastprivates |
268 | // CHECK: [[OMP_IS_LAST_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[OMP_IS_LAST]], |
269 | // CHECK: [[IS_LAST_IT:%.+]] = icmp ne i{{[0-9]+}} [[OMP_IS_LAST_VAL]], 0 |
270 | // CHECK: br i1 [[IS_LAST_IT]], label %[[OMP_LASTPRIV_BLOCK:.+]], label %[[OMP_LASTPRIV_DONE:.+]] |
271 | |
272 | // CHECK: [[OMP_LASTPRIV_BLOCK]]: |
273 | // CHECK-64-DAG: store {{.+}}, {{.+}} [[TVAR_TGT]], |
274 | // CHECK-32-DAG: store {{.+}}, {{.+}} [[T_VAR_ADDR]], |
275 | // CHECK-DAG: [[VEC_TGT_REF:%.+]] = bitcast {{.+}} [[VEC_TGT]] to |
276 | // CHECK-DAG: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[VEC_TGT_REF]], |
277 | // CHECK-DAG: [[S_ARR_BEGIN:%.+]] = getelementptr {{.+}} [[S_ARR_TGT]], |
278 | // CHECK-DAG: call void @llvm.memcpy.{{.+}}( |
279 | // CHECK-DAG: [[VAR_TGT_BCAST:%.+]] = bitcast {{.+}} [[VAR_TGT]] to |
280 | // CHECK-DAG: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[VAR_TGT_BCAST]], |
281 | // CHECK-64-DAG: store {{.+}}, {{.+}} [[SVAR_TGT]], |
282 | // CHECK-32-DAG: store {{.+}}, {{.+}} [[SVAR_ADDR]], |
283 | // CHECK: ret void |
284 | |
285 | // CHECK: define internal void [[OMP_OUTLINED:@.+]](i{{[0-9]+}}* noalias [[GTID_ADDR:%.+]], i{{[0-9]+}}* noalias %{{.+}}, {{.+}}, {{.+}}, [2 x i{{[0-9]+}}]*{{.+}} [[VEC_IN:%.+]], i{{[0-9]+}}{{.+}} [[T_VAR_IN:%.+]], [2 x [[S_FLOAT_TY]]]*{{.+}} [[S_ARR_IN:%.+]], [[S_FLOAT_TY]]*{{.+}} [[VAR_IN:%.+]], i{{[0-9]+}}{{.*}} [[S_VAR_IN:%.+]]) |
286 | |
287 | // gbl and bound tid vars, prev lb and ub vars |
288 | // CHECK: {{.+}} = alloca i{{[0-9]+}}*, |
289 | // CHECK: {{.+}} = alloca i{{[0-9]+}}*, |
290 | // CHECK: alloca i{{[0-9]+}}, |
291 | // CHECK: alloca i{{[0-9]+}}, |
292 | |
293 | // CHECK: [[VEC_ADDR:%.+]] = alloca [2 x i{{[0-9]+}}]*, |
294 | // CHECK: [[T_VAR_ADDR:%.+]] = alloca i{{[0-9]+}}, |
295 | // CHECK: [[S_ARR_ADDR:%.+]] = alloca [2 x [[S_FLOAT_TY]]]*, |
296 | // CHECK: [[VAR_ADDR:%.+]] = alloca [[S_FLOAT_TY]]*, |
297 | // CHECK: [[SVAR_ADDR:%.+]] = alloca i{{[0-9]+}}, |
298 | // skip loop variables |
299 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
300 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
301 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
302 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
303 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
304 | // CHECK: [[OMP_IS_LAST:%.+]] = alloca i{{[0-9]+}}, |
305 | |
306 | // copy from parameters to local address variables |
307 | // CHECK: store {{.+}} [[VEC_IN]], {{.+}} [[VEC_ADDR]], |
308 | // CHECK: store {{.+}} [[T_VAR_IN]], {{.+}} [[T_VAR_ADDR]], |
309 | // CHECK: store {{.+}} [[S_ARR_IN]], {{.+}} [[S_ARR_ADDR]], |
310 | // CHECK: store {{.+}} [[VAR_IN]], {{.+}} [[VAR_ADDR]], |
311 | // CHECK: store {{.+}} [[S_VAR_IN]], {{.+}} [[SVAR_ADDR]], |
312 | |
313 | // prepare lastprivate targets |
314 | // CHECK-64-DAG: [[TVAR_TGT:%.+]] = bitcast {{.+}} [[T_VAR_ADDR]] to |
315 | // CHECK-DAG: [[VEC_TGT:%.+]] = load {{.+}}, {{.+}} [[VEC_ADDR]], |
316 | // CHECK-DAG: [[S_ARR_TGT:%.+]] = load {{.+}}, {{.+}} [[S_ARR_ADDR]], |
317 | // CHECK-DAG: [[VAR_TGT:%.+]] = load {{.+}}, {{.+}} [[VAR_ADDR]], |
318 | // CHECK-64-DAG: [[SVAR_TGT:%.+]] = bitcast {{.+}} [[SVAR_ADDR]] to |
319 | |
320 | // the distribute loop |
321 | // CHECK: call void @__kmpc_for_static_init_4( |
322 | // skip body: code generation routine is same as distribute parallel for lastprivate |
323 | // CHECK: call void @__kmpc_for_static_fini( |
324 | |
325 | // lastprivates |
326 | // CHECK: [[OMP_IS_LAST_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[OMP_IS_LAST]], |
327 | // CHECK: [[IS_LAST_IT:%.+]] = icmp ne i{{[0-9]+}} [[OMP_IS_LAST_VAL]], 0 |
328 | // CHECK: br i1 [[IS_LAST_IT]], label %[[OMP_LASTPRIV_BLOCK:.+]], label %[[OMP_LASTPRIV_DONE:.+]] |
329 | |
330 | // CHECK: [[OMP_LASTPRIV_BLOCK]]: |
331 | // CHECK-64-DAG: store {{.+}}, {{.+}} [[TVAR_TGT]], |
332 | // CHECK-32-DAG: store {{.+}}, {{.+}} [[T_VAR_ADDR]], |
333 | // CHECK-DAG: [[VEC_TGT_REF:%.+]] = bitcast {{.+}} [[VEC_TGT]] to |
334 | // CHECK-DAG: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[VEC_TGT_REF]], |
335 | // CHECK-DAG: [[S_ARR_BEGIN:%.+]] = getelementptr {{.+}} [[S_ARR_TGT]], |
336 | // CHECK-DAG: call void @llvm.memcpy.{{.+}}( |
337 | // CHECK-DAG: [[VAR_TGT_BCAST:%.+]] = bitcast {{.+}} [[VAR_TGT]] to |
338 | // CHECK-DAG: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[VAR_TGT_BCAST]], |
339 | // CHECK-64-DAG: store {{.+}}, {{.+}} [[SVAR_TGT]], |
340 | // CHECK-32-DAG: store {{.+}}, {{.+}} [[SVAR_ADDR]], |
341 | // CHECK: ret void |
342 | |
343 | // template tmain |
344 | // CHECK: define{{.*}} i{{[0-9]+}} [[TMAIN_INT:@.+]]() |
345 | // CHECK: [[TEST:%.+]] = alloca [[S_INT_TY]], |
346 | // CHECK: call {{.*}} [[S_INT_TY_DEF_CONSTR:@.+]]([[S_INT_TY]]* [[TEST]]) |
347 | // CHECK: call i{{[0-9]+}} @__tgt_target_teams( |
348 | // CHECK: call void [[OFFLOAD_FUN_1:@.+]]( |
349 | // CHECK: ret |
350 | |
351 | // CHECK: define internal void [[OFFLOAD_FUN_1]]( |
352 | // CHECK: call void (%{{.+}}*, i{{[0-9]+}}, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)*, ...) @__kmpc_fork_teams(%{{.+}}* @{{.+}}, i{{[0-9]+}} 4, |
353 | // CHECK: ret |
354 | |
355 | // CHECK: define internal void [[OMP_OUTLINED_1:@.+]](i{{[0-9]+}}* noalias [[GTID_ADDR1:%.+]], i{{[0-9]+}}* noalias %{{.+}}, [2 x i{{[0-9]+}}]*{{.+}} [[VEC_IN1:%.+]], i{{[0-9]+}}{{.+}} [[T_VAR_IN1:%.+]], [2 x [[S_INT_TY]]]*{{.+}} [[S_ARR_IN1:%.+]], [[S_INT_TY]]*{{.+}} [[VAR_IN1:%.+]]) |
356 | // skip alloca of global_tid and bound_tid |
357 | // CHECK: {{.+}} = alloca i{{[0-9]+}}*, |
358 | // CHECK: {{.+}} = alloca i{{[0-9]+}}*, |
359 | // CHECK: [[VEC_ADDR1:%.+]] = alloca [2 x i{{[0-9]+}}]*, |
360 | // CHECK: [[T_VAR_ADDR1:%.+]] = alloca i{{[0-9]+}}, |
361 | // CHECK: [[S_ARR_ADDR1:%.+]] = alloca [2 x [[S_INT_TY]]]*, |
362 | // CHECK: [[VAR_ADDR1:%.+]] = alloca [[S_INT_TY]]*, |
363 | // skip loop variables |
364 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
365 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
366 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
367 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
368 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
369 | // CHECK: [[OMP_IS_LAST1:%.+]] = alloca i{{[0-9]+}}, |
370 | |
371 | // copy from parameters to local address variables |
372 | // CHECK: store {{.+}} [[VEC_IN1]], {{.+}} [[VEC_ADDR1]], |
373 | // CHECK: store {{.+}} [[T_VAR_IN1]], {{.+}} [[T_VAR_ADDR1]], |
374 | // CHECK: store {{.+}} [[S_ARR_IN1]], {{.+}} [[S_ARR_ADDR1]], |
375 | // CHECK: store {{.+}} [[VAR_IN1]], {{.+}} [[VAR_ADDR1]], |
376 | |
377 | // prepare lastprivate targets |
378 | // CHECK-64-DAG: [[T_VAR_TGT:%.+]] = bitcast {{.+}} [[T_VAR_ADDR1]] to |
379 | // CHECK-DAG: [[VEC_TGT:%.+]] = load {{.+}}, {{.+}} [[VEC_ADDR1]], |
380 | // CHECK-DAG: [[S_ARR_TGT:%.+]] = load {{.+}}, {{.+}} [[S_ARR_ADDR1]], |
381 | // CHECK-DAG: [[VAR_TGT:%.+]] = load {{.+}}, {{.+}} [[VAR_ADDR1]], |
382 | |
383 | // CHECK: call void @__kmpc_for_static_init_4( |
384 | // CHECK: call void {{.*}} @__kmpc_fork_call({{.+}}, {{.+}}, {{.+}} @[[TPAR_OUTL:.+]] to |
385 | // CHECK: call void @__kmpc_for_static_fini( |
386 | |
387 | // lastprivates |
388 | // CHECK: [[OMP_IS_LAST_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[OMP_IS_LAST1]], |
389 | // CHECK: [[IS_LAST_IT:%.+]] = icmp ne i{{[0-9]+}} [[OMP_IS_LAST_VAL]], 0 |
390 | // CHECK: br i1 [[IS_LAST_IT]], label %[[OMP_LASTPRIV_BLOCK:.+]], label %[[OMP_LASTPRIV_DONE:.+]] |
391 | |
392 | // CHECK: [[OMP_LASTPRIV_BLOCK]]: |
393 | // CHECK-64-DAG: store {{.+}}, {{.+}} [[T_VAR_TGT]], |
394 | // CHECK-32-DAG: store {{.+}}, {{.+}} [[T_VAR_ADDR1]], |
395 | // CHECK-DAG: [[VEC_TGT_BCAST:%.+]] = bitcast {{.+}} [[VEC_TGT]] to |
396 | // CHECK-DAG: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[VEC_TGT_BCAST]], |
397 | // CHECK-DAG: {{.+}} = getelementptr {{.+}} [[S_ARR_TGT]], |
398 | // CHECK: call void @llvm.memcpy.{{.+}}( |
399 | // CHECK-DAG: [[VAR_ADDR_BCAST:%.+]] = bitcast {{.+}} [[VAR_TGT]] to |
400 | // CHECK-DAG: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[VAR_ADDR_BCAST]], |
401 | // CHECK: ret void |
402 | |
403 | // CHECK: define internal void [[TPAR_OUTL:@.+]](i{{[0-9]+}}* noalias [[GTID_ADDR1:%.+]], i{{[0-9]+}}* noalias %{{.+}}, {{.+}}, {{.+}}, [2 x i{{[0-9]+}}]*{{.+}} [[VEC_IN1:%.+]], i{{[0-9]+}}{{.+}} [[T_VAR_IN1:%.+]], [2 x [[S_INT_TY]]]*{{.+}} [[S_ARR_IN1:%.+]], [[S_INT_TY]]*{{.+}} [[VAR_IN1:%.+]]) |
404 | // skip alloca of global_tid and bound_tid, and prev lb and ub vars |
405 | // CHECK: {{.+}} = alloca i{{[0-9]+}}*, |
406 | // CHECK: {{.+}} = alloca i{{[0-9]+}}*, |
407 | // CHECK: alloca i{{[0-9]+}}, |
408 | // CHECK: alloca i{{[0-9]+}}, |
409 | |
410 | // CHECK: [[VEC_ADDR1:%.+]] = alloca [2 x i{{[0-9]+}}]*, |
411 | // CHECK: [[T_VAR_ADDR1:%.+]] = alloca i{{[0-9]+}}, |
412 | // CHECK: [[S_ARR_ADDR1:%.+]] = alloca [2 x [[S_INT_TY]]]*, |
413 | // CHECK: [[VAR_ADDR1:%.+]] = alloca [[S_INT_TY]]*, |
414 | // skip loop variables |
415 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
416 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
417 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
418 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
419 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
420 | // CHECK: [[OMP_IS_LAST1:%.+]] = alloca i{{[0-9]+}}, |
421 | |
422 | // copy from parameters to local address variables |
423 | // CHECK: store {{.+}} [[VEC_IN1]], {{.+}} [[VEC_ADDR1]], |
424 | // CHECK: store {{.+}} [[T_VAR_IN1]], {{.+}} [[T_VAR_ADDR1]], |
425 | // CHECK: store {{.+}} [[S_ARR_IN1]], {{.+}} [[S_ARR_ADDR1]], |
426 | // CHECK: store {{.+}} [[VAR_IN1]], {{.+}} [[VAR_ADDR1]], |
427 | |
428 | // prepare lastprivate targets |
429 | // CHECK-64-DAG: [[T_VAR_TGT:%.+]] = bitcast {{.+}} [[T_VAR_ADDR1]] to |
430 | // CHECK-DAG: [[VEC_TGT:%.+]] = load {{.+}}, {{.+}} [[VEC_ADDR1]], |
431 | // CHECK-DAG: [[S_ARR_TGT:%.+]] = load {{.+}}, {{.+}} [[S_ARR_ADDR1]], |
432 | // CHECK-DAG: [[VAR_TGT:%.+]] = load {{.+}}, {{.+}} [[VAR_ADDR1]], |
433 | |
434 | // CHECK: call void @__kmpc_for_static_init_4( |
435 | // skip body: code generation routine is same as distribute parallel for lastprivate |
436 | // CHECK: call void @__kmpc_for_static_fini( |
437 | |
438 | // lastprivates |
439 | // CHECK: [[OMP_IS_LAST_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[OMP_IS_LAST1]], |
440 | // CHECK: [[IS_LAST_IT:%.+]] = icmp ne i{{[0-9]+}} [[OMP_IS_LAST_VAL]], 0 |
441 | // CHECK: br i1 [[IS_LAST_IT]], label %[[OMP_LASTPRIV_BLOCK:.+]], label %[[OMP_LASTPRIV_DONE:.+]] |
442 | |
443 | // CHECK: [[OMP_LASTPRIV_BLOCK]]: |
444 | // CHECK-64-DAG: store {{.+}}, {{.+}} [[T_VAR_TGT]], |
445 | // CHECK-32-DAG: store {{.+}}, {{.+}} [[T_VAR_ADDR1]], |
446 | // CHECK-DAG: [[VEC_TGT_BCAST:%.+]] = bitcast {{.+}} [[VEC_TGT]] to |
447 | // CHECK-DAG: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[VEC_TGT_BCAST]], |
448 | // CHECK-DAG: {{.+}} = getelementptr {{.+}} [[S_ARR_TGT]], |
449 | // CHECK: call void @llvm.memcpy.{{.+}}( |
450 | // CHECK-DAG: [[VAR_ADDR_BCAST:%.+]] = bitcast {{.+}} [[VAR_TGT]] to |
451 | // CHECK-DAG: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[VAR_ADDR_BCAST]], |
452 | // CHECK: ret void |
453 | |
454 | #endif |
455 | |