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 simd 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 simd 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: store i32 2, i32* % |
115 | // LAMBDA: [[OMP_IS_LAST_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[OMP_IS_LAST]], |
116 | // LAMBDA: [[IS_LAST_IT:%.+]] = icmp ne i{{[0-9]+}} [[OMP_IS_LAST_VAL]], 0 |
117 | // LAMBDA: br i1 [[IS_LAST_IT]], label %[[OMP_LASTPRIV_BLOCK:.+]], label %[[OMP_LASTPRIV_DONE:.+]] |
118 | |
119 | // LAMBDA: [[OMP_LASTPRIV_BLOCK]]: |
120 | // LAMBDA-DAG: store {{.+}}, {{.+}} [[G_TGT]], |
121 | // LAMBDA-DAG: store {{.+}}, {{.+}} [[G1_TGT]], |
122 | // LAMBDA-64-DAG: store {{.+}}, {{.+}} [[SVAR_TGT]], |
123 | // LAMBDA-32-DAG: store {{.+}}, {{.+}} [[SVAR_ADDR]], |
124 | // LAMBDA-DAG: store {{.+}}, {{.+}} [[SFVAR_TGT]], |
125 | // LAMBDA: br label %[[OMP_LASTPRIV_DONE]] |
126 | // LAMBDA: [[OMP_LASTPRIV_DONE]]: |
127 | // LAMBDA: ret |
128 | |
129 | // LAMBDA: define{{.*}} internal{{.*}} void @[[LPAR_OUTL]](i32* noalias %{{.+}}, i32* noalias %{{.+}}, {{.+}}, {{.+}}, {{.+}} [[G1_IN:%.+]], {{.+}} [[SVAR_IN:%.+]], {{.+}} [[SFVAR_IN:%.+]], {{.+}} [[G_IN:%.+]]) |
130 | // skip tid and prev variables |
131 | // LAMBDA: alloca |
132 | // LAMBDA: alloca |
133 | // LAMBDA: alloca |
134 | // LAMBDA: alloca |
135 | // LAMBDA: [[G1_ADDR:%.+]] = alloca {{.+}}, |
136 | // LAMBDA: [[SVAR_ADDR:%.+]] = alloca {{.+}}, |
137 | // LAMBDA: [[SFVAR_ADDR:%.+]] = alloca {{.+}}, |
138 | // LAMBDA: [[G_ADDR:%.+]] = alloca {{.+}}, |
139 | // LAMBDA-64: [[G1_REF:%.+]] = alloca double*, |
140 | // loop variables |
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: {{.+}} = alloca i{{[0-9]+}}, |
146 | // LAMBDA: [[OMP_IS_LAST:%.+]] = alloca i{{[0-9]+}}, |
147 | |
148 | // LAMBDA-DAG: store {{.+}} [[G_IN]], {{.+}} [[G_ADDR]], |
149 | // LAMBDA-DAG: store {{.+}} [[G1_IN]], {{.+}} [[G1_ADDR]], |
150 | // LAMBDA-DAG: store {{.+}} [[SVAR_IN]], {{.+}} [[SVAR_ADDR]], |
151 | // LAMBDA-DAG: store {{.+}} [[SFVAR_IN]], {{.+}} [[SFVAR_ADDR]], |
152 | |
153 | // LAMBDA-64-DAG: [[G_TGT:%.+]] = bitcast {{.+}} [[G_ADDR]] to |
154 | // LAMBDA-32-DAG: [[G_TGT:%.+]] = load {{.+}}, {{.+}} [[G_ADDR]], |
155 | // LAMBDA-64-DAG: [[G1_TGT:%.+]] = load {{.+}}, {{.+}} [[G1_REF]], |
156 | // LAMBDA-32-DAG: [[G1_TGT:%.+]] = load {{.+}}, {{.+}} [[G1_ADDR]], |
157 | // LAMBDA-64-DAG: [[SVAR_TGT:%.+]] = bitcast {{.+}} [[SVAR_ADDR]] to |
158 | // LAMBDA-DAG: [[SFVAR_TGT:%.+]] = bitcast {{.+}} [[SFVAR_ADDR]] to |
159 | |
160 | |
161 | // LAMBDA: call {{.*}}void @__kmpc_for_static_init_4( |
162 | // LAMBDA: call{{.*}} void [[INNER_LAMBDA:@.+]]( |
163 | // LAMBDA: call {{.*}}void @__kmpc_for_static_fini( |
164 | |
165 | // LAMBDA: store i32 2, i32* % |
166 | // LAMBDA: [[OMP_IS_LAST_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[OMP_IS_LAST]], |
167 | // LAMBDA: [[IS_LAST_IT:%.+]] = icmp ne i{{[0-9]+}} [[OMP_IS_LAST_VAL]], 0 |
168 | // LAMBDA: br i1 [[IS_LAST_IT]], label %[[OMP_LASTPRIV_BLOCK:.+]], label %[[OMP_LASTPRIV_DONE:.+]] |
169 | |
170 | // LAMBDA: [[OMP_LASTPRIV_BLOCK]]: |
171 | // LAMBDA-DAG: store {{.+}}, {{.+}} [[G_TGT]], |
172 | // LAMBDA-DAG: store {{.+}}, {{.+}} [[G1_TGT]], |
173 | // LAMBDA-64-DAG: store {{.+}}, {{.+}} [[SVAR_TGT]], |
174 | // LAMBDA-32-DAG: store {{.+}}, {{.+}} [[SVAR_ADDR]], |
175 | // LAMBDA-DAG: store {{.+}}, {{.+}} [[SFVAR_TGT]], |
176 | // LAMBDA: br label %[[OMP_LASTPRIV_DONE]] |
177 | // LAMBDA: [[OMP_LASTPRIV_DONE]]: |
178 | // LAMBDA: ret |
179 | |
180 | [&]() { |
181 | // LAMBDA: define {{.+}} void [[INNER_LAMBDA]](%{{.+}}* [[ARG_PTR:%.+]]) |
182 | // LAMBDA: store %{{.+}}* [[ARG_PTR]], %{{.+}}** [[ARG_PTR_REF:%.+]], |
183 | g = 2; |
184 | g1 = 2; |
185 | svar = 4; |
186 | sfvar = 8.0; |
187 | // LAMBDA: [[ARG_PTR:%.+]] = load %{{.+}}*, %{{.+}}** [[ARG_PTR_REF]] |
188 | // LAMBDA: [[G_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 |
189 | // LAMBDA: [[G_REF:%.+]] = load double*, double** [[G_PTR_REF]] |
190 | // LAMBDA: store double 2.0{{.+}}, double* [[G_REF]] |
191 | |
192 | // LAMBDA: [[TMP_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 |
193 | // LAMBDA: [[G1_REF:%.+]] = load double*, double** [[TMP_PTR_REF]] |
194 | // LAMBDA: store double 2.0{{.+}}, double* [[G1_REF]], |
195 | // LAMBDA: [[SVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 |
196 | // LAMBDA: [[SVAR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[SVAR_PTR_REF]] |
197 | // LAMBDA: store i{{[0-9]+}} 4, i{{[0-9]+}}* [[SVAR_REF]] |
198 | // LAMBDA: [[SFVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 3 |
199 | // LAMBDA: [[SFVAR_REF:%.+]] = load float*, float** [[SFVAR_PTR_REF]] |
200 | // LAMBDA: store float 8.0{{.+}}, float* [[SFVAR_REF]] |
201 | }(); |
202 | } |
203 | }(); |
204 | return 0; |
205 | #else |
206 | S<float> test; |
207 | int t_var = 0; |
208 | int vec[] = {1, 2}; |
209 | S<float> s_arr[] = {1, 2}; |
210 | S<float> &var = test; |
211 | |
212 | #pragma omp target teams distribute parallel for simd lastprivate(t_var, vec, s_arr, s_arr, var, var, svar) |
213 | for (int i = 0; i < 2; ++i) { |
214 | vec[i] = t_var; |
215 | s_arr[i] = var; |
216 | } |
217 | int i; |
218 | |
219 | return tmain<int>(); |
220 | #endif |
221 | } |
222 | |
223 | // CHECK: define{{.*}} i{{[0-9]+}} @main() |
224 | // CHECK: [[TEST:%.+]] = alloca [[S_FLOAT_TY]], |
225 | // CHECK: call {{.*}} [[S_FLOAT_TY_DEF_CONSTR:@.+]]([[S_FLOAT_TY]]* [[TEST]]) |
226 | // CHECK: call i{{[0-9]+}} @__tgt_target_teams( |
227 | // CHECK: call void [[OFFLOAD_FUN:@.+]]( |
228 | // CHECK: ret |
229 | |
230 | // CHECK: define{{.+}} [[OFFLOAD_FUN]]( |
231 | // CHECK: call void (%{{.+}}*, i{{[0-9]+}}, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)*, ...) @__kmpc_fork_teams( |
232 | // CHECK: ret |
233 | // |
234 | // 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:%.+]]) |
235 | // CHECK: {{.+}} = alloca i{{[0-9]+}}*, |
236 | // CHECK: {{.+}} = alloca i{{[0-9]+}}*, |
237 | // CHECK: [[VEC_ADDR:%.+]] = alloca [2 x i{{[0-9]+}}]*, |
238 | // CHECK: [[T_VAR_ADDR:%.+]] = alloca i{{[0-9]+}}, |
239 | // CHECK: [[S_ARR_ADDR:%.+]] = alloca [2 x [[S_FLOAT_TY]]]*, |
240 | // CHECK: [[VAR_ADDR:%.+]] = alloca [[S_FLOAT_TY]]*, |
241 | // CHECK: [[SVAR_ADDR:%.+]] = alloca i{{[0-9]+}}, |
242 | // skip loop variables |
243 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
244 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
245 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
246 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
247 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
248 | // CHECK: [[OMP_IS_LAST:%.+]] = alloca i{{[0-9]+}}, |
249 | |
250 | // copy from parameters to local address variables |
251 | // CHECK: store {{.+}} [[VEC_IN]], {{.+}} [[VEC_ADDR]], |
252 | // CHECK: store {{.+}} [[T_VAR_IN]], {{.+}} [[T_VAR_ADDR]], |
253 | // CHECK: store {{.+}} [[S_ARR_IN]], {{.+}} [[S_ARR_ADDR]], |
254 | // CHECK: store {{.+}} [[VAR_IN]], {{.+}} [[VAR_ADDR]], |
255 | // CHECK: store {{.+}} [[S_VAR_IN]], {{.+}} [[SVAR_ADDR]], |
256 | |
257 | // prepare lastprivate targets |
258 | // CHECK-64-DAG: [[TVAR_TGT:%.+]] = bitcast {{.+}} [[T_VAR_ADDR]] to |
259 | // CHECK-DAG: [[VEC_TGT:%.+]] = load {{.+}}, {{.+}} [[VEC_ADDR]], |
260 | // CHECK-DAG: [[S_ARR_TGT:%.+]] = load {{.+}}, {{.+}} [[S_ARR_ADDR]], |
261 | // CHECK-DAG: [[VAR_TGT:%.+]] = load {{.+}}, {{.+}} [[VAR_ADDR]], |
262 | // CHECK-64-DAG: [[SVAR_TGT:%.+]] = bitcast {{.+}} [[SVAR_ADDR]] to |
263 | |
264 | // the distribute loop |
265 | // CHECK: call void @__kmpc_for_static_init_4( |
266 | // CHECK: call void {{.*}} @__kmpc_fork_call({{.+}}, {{.+}}, {{.+}} @[[PAR_OUTL:.+]] to |
267 | // CHECK: call void @__kmpc_for_static_fini( |
268 | |
269 | // lastprivates |
270 | // CHECK: [[OMP_IS_LAST_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[OMP_IS_LAST]], |
271 | // CHECK: [[IS_LAST_IT:%.+]] = icmp ne i{{[0-9]+}} [[OMP_IS_LAST_VAL]], 0 |
272 | // CHECK: br i1 [[IS_LAST_IT]], label %[[OMP_LASTPRIV_BLOCK:.+]], label %[[OMP_LASTPRIV_DONE:.+]] |
273 | |
274 | // CHECK: [[OMP_LASTPRIV_BLOCK]]: |
275 | // CHECK-64-DAG: store {{.+}}, {{.+}} [[TVAR_TGT]], |
276 | // CHECK-32-DAG: store {{.+}}, {{.+}} [[T_VAR_ADDR]], |
277 | // CHECK-DAG: [[VEC_TGT_REF:%.+]] = bitcast {{.+}} [[VEC_TGT]] to |
278 | // CHECK-DAG: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[VEC_TGT_REF]], |
279 | // CHECK-DAG: [[S_ARR_BEGIN:%.+]] = getelementptr {{.+}} [[S_ARR_TGT]], |
280 | // CHECK-DAG: call void @llvm.memcpy.{{.+}}( |
281 | // CHECK-DAG: [[VAR_TGT_BCAST:%.+]] = bitcast {{.+}} [[VAR_TGT]] to |
282 | // CHECK-DAG: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[VAR_TGT_BCAST]], |
283 | // CHECK-64-DAG: store {{.+}}, {{.+}} [[SVAR_TGT]], |
284 | // CHECK-32-DAG: store {{.+}}, {{.+}} [[SVAR_ADDR]], |
285 | // CHECK: ret void |
286 | |
287 | // 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:%.+]]) |
288 | |
289 | // gbl and bound tid vars, prev lb and ub vars |
290 | // CHECK: {{.+}} = alloca i{{[0-9]+}}*, |
291 | // CHECK: {{.+}} = alloca i{{[0-9]+}}*, |
292 | // CHECK: alloca i{{[0-9]+}}, |
293 | // CHECK: alloca i{{[0-9]+}}, |
294 | |
295 | // CHECK: [[VEC_ADDR:%.+]] = alloca [2 x i{{[0-9]+}}]*, |
296 | // CHECK: [[T_VAR_ADDR:%.+]] = alloca i{{[0-9]+}}, |
297 | // CHECK: [[S_ARR_ADDR:%.+]] = alloca [2 x [[S_FLOAT_TY]]]*, |
298 | // CHECK: [[VAR_ADDR:%.+]] = alloca [[S_FLOAT_TY]]*, |
299 | // CHECK: [[SVAR_ADDR:%.+]] = alloca i{{[0-9]+}}, |
300 | // skip loop variables |
301 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
302 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
303 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
304 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
305 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
306 | // CHECK: [[OMP_IS_LAST:%.+]] = alloca i{{[0-9]+}}, |
307 | |
308 | // copy from parameters to local address variables |
309 | // CHECK: store {{.+}} [[VEC_IN]], {{.+}} [[VEC_ADDR]], |
310 | // CHECK: store {{.+}} [[T_VAR_IN]], {{.+}} [[T_VAR_ADDR]], |
311 | // CHECK: store {{.+}} [[S_ARR_IN]], {{.+}} [[S_ARR_ADDR]], |
312 | // CHECK: store {{.+}} [[VAR_IN]], {{.+}} [[VAR_ADDR]], |
313 | // CHECK: store {{.+}} [[S_VAR_IN]], {{.+}} [[SVAR_ADDR]], |
314 | |
315 | // prepare lastprivate targets |
316 | // CHECK-64-DAG: [[TVAR_TGT:%.+]] = bitcast {{.+}} [[T_VAR_ADDR]] to |
317 | // CHECK-DAG: [[VEC_TGT:%.+]] = load {{.+}}, {{.+}} [[VEC_ADDR]], |
318 | // CHECK-DAG: [[S_ARR_TGT:%.+]] = load {{.+}}, {{.+}} [[S_ARR_ADDR]], |
319 | // CHECK-DAG: [[VAR_TGT:%.+]] = load {{.+}}, {{.+}} [[VAR_ADDR]], |
320 | // CHECK-64-DAG: [[SVAR_TGT:%.+]] = bitcast {{.+}} [[SVAR_ADDR]] to |
321 | |
322 | // the distribute loop |
323 | // CHECK: call void @__kmpc_for_static_init_4( |
324 | // skip body: code generation routine is same as distribute parallel for lastprivate |
325 | // CHECK: call void @__kmpc_for_static_fini( |
326 | |
327 | // lastprivates |
328 | // CHECK: [[OMP_IS_LAST_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[OMP_IS_LAST]], |
329 | // CHECK: [[IS_LAST_IT:%.+]] = icmp ne i{{[0-9]+}} [[OMP_IS_LAST_VAL]], 0 |
330 | // CHECK: br i1 [[IS_LAST_IT]], label %[[OMP_LASTPRIV_BLOCK:.+]], label %[[OMP_LASTPRIV_DONE:.+]] |
331 | |
332 | // CHECK: [[OMP_LASTPRIV_BLOCK]]: |
333 | // CHECK-64-DAG: store {{.+}}, {{.+}} [[TVAR_TGT]], |
334 | // CHECK-32-DAG: store {{.+}}, {{.+}} [[T_VAR_ADDR]], |
335 | // CHECK-DAG: [[VEC_TGT_REF:%.+]] = bitcast {{.+}} [[VEC_TGT]] to |
336 | // CHECK-DAG: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[VEC_TGT_REF]], |
337 | // CHECK-DAG: [[S_ARR_BEGIN:%.+]] = getelementptr {{.+}} [[S_ARR_TGT]], |
338 | // CHECK-DAG: call void @llvm.memcpy.{{.+}}( |
339 | // CHECK-DAG: [[VAR_TGT_BCAST:%.+]] = bitcast {{.+}} [[VAR_TGT]] to |
340 | // CHECK-DAG: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[VAR_TGT_BCAST]], |
341 | // CHECK-64-DAG: store {{.+}}, {{.+}} [[SVAR_TGT]], |
342 | // CHECK-32-DAG: store {{.+}}, {{.+}} [[SVAR_ADDR]], |
343 | // CHECK: ret void |
344 | |
345 | // template tmain |
346 | // CHECK: define{{.*}} i{{[0-9]+}} [[TMAIN_INT:@.+]]() |
347 | // CHECK: [[TEST:%.+]] = alloca [[S_INT_TY]], |
348 | // CHECK: call {{.*}} [[S_INT_TY_DEF_CONSTR:@.+]]([[S_INT_TY]]* [[TEST]]) |
349 | // CHECK: call i{{[0-9]+}} @__tgt_target_teams( |
350 | // CHECK: call void [[OFFLOAD_FUN_1:@.+]]( |
351 | // CHECK: ret |
352 | |
353 | // CHECK: define internal void [[OFFLOAD_FUN_1]]( |
354 | // CHECK: call void (%{{.+}}*, i{{[0-9]+}}, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)*, ...) @__kmpc_fork_teams(%{{.+}}* @{{.+}}, i{{[0-9]+}} 4, |
355 | // CHECK: ret |
356 | |
357 | // 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:%.+]]) |
358 | // skip alloca of global_tid and bound_tid |
359 | // CHECK: {{.+}} = alloca i{{[0-9]+}}*, |
360 | // CHECK: {{.+}} = alloca i{{[0-9]+}}*, |
361 | // CHECK: [[VEC_ADDR1:%.+]] = alloca [2 x i{{[0-9]+}}]*, |
362 | // CHECK: [[T_VAR_ADDR1:%.+]] = alloca i{{[0-9]+}}, |
363 | // CHECK: [[S_ARR_ADDR1:%.+]] = alloca [2 x [[S_INT_TY]]]*, |
364 | // CHECK: [[VAR_ADDR1:%.+]] = alloca [[S_INT_TY]]*, |
365 | // skip loop variables |
366 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
367 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
368 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
369 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
370 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
371 | // CHECK: [[OMP_IS_LAST1:%.+]] = alloca i{{[0-9]+}}, |
372 | |
373 | // copy from parameters to local address variables |
374 | // CHECK: store {{.+}} [[VEC_IN1]], {{.+}} [[VEC_ADDR1]], |
375 | // CHECK: store {{.+}} [[T_VAR_IN1]], {{.+}} [[T_VAR_ADDR1]], |
376 | // CHECK: store {{.+}} [[S_ARR_IN1]], {{.+}} [[S_ARR_ADDR1]], |
377 | // CHECK: store {{.+}} [[VAR_IN1]], {{.+}} [[VAR_ADDR1]], |
378 | |
379 | // prepare lastprivate targets |
380 | // CHECK-64-DAG: [[T_VAR_TGT:%.+]] = bitcast {{.+}} [[T_VAR_ADDR1]] to |
381 | // CHECK-DAG: [[VEC_TGT:%.+]] = load {{.+}}, {{.+}} [[VEC_ADDR1]], |
382 | // CHECK-DAG: [[S_ARR_TGT:%.+]] = load {{.+}}, {{.+}} [[S_ARR_ADDR1]], |
383 | // CHECK-DAG: [[VAR_TGT:%.+]] = load {{.+}}, {{.+}} [[VAR_ADDR1]], |
384 | |
385 | // CHECK: call void @__kmpc_for_static_init_4( |
386 | // CHECK: call void {{.*}} @__kmpc_fork_call({{.+}}, {{.+}}, {{.+}} @[[TPAR_OUTL:.+]] to |
387 | // CHECK: call void @__kmpc_for_static_fini( |
388 | |
389 | // lastprivates |
390 | // CHECK: [[OMP_IS_LAST_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[OMP_IS_LAST1]], |
391 | // CHECK: [[IS_LAST_IT:%.+]] = icmp ne i{{[0-9]+}} [[OMP_IS_LAST_VAL]], 0 |
392 | // CHECK: br i1 [[IS_LAST_IT]], label %[[OMP_LASTPRIV_BLOCK:.+]], label %[[OMP_LASTPRIV_DONE:.+]] |
393 | |
394 | // CHECK: [[OMP_LASTPRIV_BLOCK]]: |
395 | // CHECK-64-DAG: store {{.+}}, {{.+}} [[T_VAR_TGT]], |
396 | // CHECK-32-DAG: store {{.+}}, {{.+}} [[T_VAR_ADDR1]], |
397 | // CHECK-DAG: [[VEC_TGT_BCAST:%.+]] = bitcast {{.+}} [[VEC_TGT]] to |
398 | // CHECK-DAG: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[VEC_TGT_BCAST]], |
399 | // CHECK-DAG: {{.+}} = getelementptr {{.+}} [[S_ARR_TGT]], |
400 | // CHECK: call void @llvm.memcpy.{{.+}}( |
401 | // CHECK-DAG: [[VAR_ADDR_BCAST:%.+]] = bitcast {{.+}} [[VAR_TGT]] to |
402 | // CHECK-DAG: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[VAR_ADDR_BCAST]], |
403 | // CHECK: ret void |
404 | |
405 | // 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:%.+]]) |
406 | // skip alloca of global_tid and bound_tid, and prev lb and ub vars |
407 | // CHECK: {{.+}} = alloca i{{[0-9]+}}*, |
408 | // CHECK: {{.+}} = alloca i{{[0-9]+}}*, |
409 | // CHECK: alloca i{{[0-9]+}}, |
410 | // CHECK: alloca i{{[0-9]+}}, |
411 | |
412 | // CHECK: [[VEC_ADDR1:%.+]] = alloca [2 x i{{[0-9]+}}]*, |
413 | // CHECK: [[T_VAR_ADDR1:%.+]] = alloca i{{[0-9]+}}, |
414 | // CHECK: [[S_ARR_ADDR1:%.+]] = alloca [2 x [[S_INT_TY]]]*, |
415 | // CHECK: [[VAR_ADDR1:%.+]] = alloca [[S_INT_TY]]*, |
416 | // skip loop variables |
417 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
418 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
419 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
420 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
421 | // CHECK: {{.+}} = alloca i{{[0-9]+}}, |
422 | // CHECK: [[OMP_IS_LAST1:%.+]] = alloca i{{[0-9]+}}, |
423 | |
424 | // copy from parameters to local address variables |
425 | // CHECK: store {{.+}} [[VEC_IN1]], {{.+}} [[VEC_ADDR1]], |
426 | // CHECK: store {{.+}} [[T_VAR_IN1]], {{.+}} [[T_VAR_ADDR1]], |
427 | // CHECK: store {{.+}} [[S_ARR_IN1]], {{.+}} [[S_ARR_ADDR1]], |
428 | // CHECK: store {{.+}} [[VAR_IN1]], {{.+}} [[VAR_ADDR1]], |
429 | |
430 | // prepare lastprivate targets |
431 | // CHECK-64-DAG: [[T_VAR_TGT:%.+]] = bitcast {{.+}} [[T_VAR_ADDR1]] to |
432 | // CHECK-DAG: [[VEC_TGT:%.+]] = load {{.+}}, {{.+}} [[VEC_ADDR1]], |
433 | // CHECK-DAG: [[S_ARR_TGT:%.+]] = load {{.+}}, {{.+}} [[S_ARR_ADDR1]], |
434 | // CHECK-DAG: [[VAR_TGT:%.+]] = load {{.+}}, {{.+}} [[VAR_ADDR1]], |
435 | |
436 | // CHECK: call void @__kmpc_for_static_init_4( |
437 | // skip body: code generation routine is same as distribute parallel for lastprivate |
438 | // CHECK: call void @__kmpc_for_static_fini( |
439 | |
440 | // lastprivates |
441 | // CHECK: [[OMP_IS_LAST_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[OMP_IS_LAST1]], |
442 | // CHECK: [[IS_LAST_IT:%.+]] = icmp ne i{{[0-9]+}} [[OMP_IS_LAST_VAL]], 0 |
443 | // CHECK: br i1 [[IS_LAST_IT]], label %[[OMP_LASTPRIV_BLOCK:.+]], label %[[OMP_LASTPRIV_DONE:.+]] |
444 | |
445 | // CHECK: [[OMP_LASTPRIV_BLOCK]]: |
446 | // CHECK-64-DAG: store {{.+}}, {{.+}} [[T_VAR_TGT]], |
447 | // CHECK-32-DAG: store {{.+}}, {{.+}} [[T_VAR_ADDR1]], |
448 | // CHECK-DAG: [[VEC_TGT_BCAST:%.+]] = bitcast {{.+}} [[VEC_TGT]] to |
449 | // CHECK-DAG: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[VEC_TGT_BCAST]], |
450 | // CHECK-DAG: {{.+}} = getelementptr {{.+}} [[S_ARR_TGT]], |
451 | // CHECK: call void @llvm.memcpy.{{.+}}( |
452 | // CHECK-DAG: [[VAR_ADDR_BCAST:%.+]] = bitcast {{.+}} [[VAR_TGT]] to |
453 | // CHECK-DAG: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[VAR_ADDR_BCAST]], |
454 | // CHECK: ret void |
455 | |
456 | #endif |
457 | |