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 - -Wno-openmp-target | 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 - -Wno-openmp-target | 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 - -Wno-openmp-target | 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 - -Wno-openmp-target | FileCheck %s --check-prefix LAMBDA --check-prefix LAMBDA-32 |
7 | |
8 | // 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 - -Wno-openmp-target | FileCheck --check-prefix SIMD-ONLY0 %s |
9 | // 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 |
10 | // 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 - -Wno-openmp-target | FileCheck --check-prefix SIMD-ONLY0 %s |
11 | // 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 - -Wno-openmp-target | FileCheck --check-prefix SIMD-ONLY0 %s |
12 | // 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 |
13 | // 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 - -Wno-openmp-target | FileCheck --check-prefix SIMD-ONLY0 %s |
14 | // SIMD-ONLY0-NOT: {{__kmpc|__tgt}} |
15 | |
16 | // RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - -Wno-openmp-target | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 |
17 | // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s |
18 | // 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 - -Wno-openmp-target | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 |
19 | // RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - -Wno-openmp-target | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 |
20 | // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s |
21 | // 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 - -Wno-openmp-target | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 |
22 | |
23 | // 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 - -Wno-openmp-target | FileCheck --check-prefix SIMD-ONLY1 %s |
24 | // 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 |
25 | // 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 - -Wno-openmp-target | FileCheck --check-prefix SIMD-ONLY1 %s |
26 | // 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 - -Wno-openmp-target | FileCheck --check-prefix SIMD-ONLY1 %s |
27 | // 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 |
28 | // 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 - -Wno-openmp-target | FileCheck --check-prefix SIMD-ONLY1 %s |
29 | // SIMD-ONLY1-NOT: {{__kmpc|__tgt}} |
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 |
53 | #pragma omp teams |
54 | #pragma omp distribute firstprivate(t_var, vec, s_arr, s_arr, var, var) |
55 | for (int i = 0; i < 2; ++i) { |
56 | vec[i] = t_var; |
57 | s_arr[i] = var; |
58 | } |
59 | return T(); |
60 | } |
61 | |
62 | int main() { |
63 | static int svar; |
64 | volatile double g; |
65 | volatile double &g1 = g; |
66 | |
67 | #ifdef LAMBDA |
68 | // LAMBDA-LABEL: @main |
69 | // LAMBDA: call{{.*}} void [[OUTER_LAMBDA:@.+]]( |
70 | [&]() { |
71 | static float sfvar; |
72 | // LAMBDA: define{{.*}} internal{{.*}} void [[OUTER_LAMBDA]]( |
73 | // LAMBDA: call i{{[0-9]+}} @__tgt_target_teams( |
74 | // LAMBDA: call void [[OFFLOADING_FUN:@.+]]( |
75 | |
76 | // LAMBDA: define{{.+}} void [[OFFLOADING_FUN]]( |
77 | // LAMBDA: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED:@.+]] to {{.+}}) |
78 | #pragma omp target |
79 | #pragma omp teams |
80 | #pragma omp distribute firstprivate(g, g1, svar, sfvar) |
81 | for (int i = 0; i < 2; ++i) { |
82 | // LAMBDA: define internal{{.*}} void [[OMP_OUTLINED]](i32* noalias %{{.+}}, i32* noalias %{{.+}}, double*{{.*}} [[G_IN:%.+]], double*{{.*}} [[G1_IN:%.+]], i32*{{.*}} [[SVAR_IN:%.+]], float*{{.*}} [[SFVAR_IN:%.+]]) |
83 | // Private alloca's for conversion |
84 | // LAMBDA: [[G_ADDR:%.+]] = alloca double*, |
85 | // LAMBDA: [[G1_ADDR:%.+]] = alloca double*, |
86 | // LAMBDA: [[SVAR_ADDR:%.+]] = alloca i{{[0-9]+}}*, |
87 | // LAMBDA: [[SFVAR_ADDR:%.+]] = alloca float*, |
88 | // LAMBDA: [[G1_REF:%.+]] = alloca double*, |
89 | |
90 | // Actual private variables to be used in the body (tmp is used for the reference type) |
91 | // LAMBDA: [[G_PRIVATE:%.+]] = alloca double, |
92 | // LAMBDA: [[G1_PRIVATE:%.+]] = alloca double, |
93 | // LAMBDA: [[TMP_PRIVATE:%.+]] = alloca double*, |
94 | // LAMBDA: [[SVAR_PRIVATE:%.+]] = alloca i{{[0-9]+}}, |
95 | // LAMBDA: [[SFVAR_PRIVATE:%.+]] = alloca float, |
96 | |
97 | // Store input parameter addresses into private alloca's for conversion |
98 | // LAMBDA: store double* [[G_IN]], double** [[G_ADDR]], |
99 | // LAMBDA: store double* [[G1_IN]], double** [[G1_ADDR]], |
100 | // LAMBDA: store i{{[0-9]+}}* [[SVAR_IN]], i{{[0-9]+}}** [[SVAR_ADDR]], |
101 | // LAMBDA: store float* [[SFVAR_IN]], float** [[SFVAR_ADDR]], |
102 | |
103 | // LAMBDA-DAG: [[G_ADDR_VAL:%.+]] = load double*, double** [[G_ADDR]], |
104 | // LAMBDA-DAG: [[SVAR_ADDR_VAL:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[SVAR_ADDR]], |
105 | // LAMBDA-DAG: [[SFVAR_ADDR_VAL:%.+]] = load float*, float** [[SFVAR_ADDR]], |
106 | // LAMBDA-DAG: [[G1_ADDR_VAL:%.+]] = load double*, double** [[G1_ADDR]], |
107 | // LAMBDA-DAG: store double* [[G1_ADDR_VAL]], double** [[G1_REF]], |
108 | // LAMBDA-DAG: [[G_CONV_VAL:%.+]] = load{{.*}} double, double* [[G_ADDR_VAL]], |
109 | // LAMBDA-DAG: store double [[G_CONV_VAL]], double* [[G_PRIVATE]], |
110 | // LAMBDA-DAG: [[TMP_VAL:%.+]] = load double*, double** [[G1_REF]], |
111 | // LAMBDA-DAG: [[TMP_VAL_VAL:%.+]] = load{{.*}} double, double* [[TMP_VAL]], |
112 | // LAMBDA-DAG: store double [[TMP_VAL_VAL]], double* [[G1_PRIVATE]], |
113 | // LAMBDA-DAG: store double* [[G1_PRIVATE]], double** [[TMP_PRIVATE]], |
114 | // LAMBDA-DAG: [[SVAR_CONV_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[SVAR_ADDR_VAL]], |
115 | // LAMBDA-DAG: store i{{[0-9]+}} [[SVAR_CONV_VAL]], i{{[0-9]+}}* [[SVAR_PRIVATE]], |
116 | // LAMBDA-DAG: [[SFVAR_CONV_VAL:%.+]] = load float, float* [[SFVAR_ADDR_VAL]], |
117 | // LAMBDA-DAG: store float [[SFVAR_CONV_VAL]], float* [[SFVAR_PRIVATE]], |
118 | // LAMBDA: call {{.*}}void @__kmpc_for_static_init_4( |
119 | g += 1; |
120 | g1 += 1; |
121 | svar += 3; |
122 | sfvar += 4.0; |
123 | // LAMBDA-DAG: [[G_VAL:%.+]] = load double, double* [[G_PRIVATE]], |
124 | // LAMBDA-DAG: [[G_NEXT:%.+]] = fadd double [[G_VAL]], 1.{{.+}} |
125 | // LAMBDA-DAG: store double [[G_NEXT]], double* [[G_PRIVATE]], |
126 | // LAMBDA-DAG: [[TMP_VAL1:%.+]] = load double*, double** [[TMP_PRIVATE]], |
127 | // LAMBDA-DAG: [[TMP_VAL_VAL1:%.+]] = load{{.*}} double, double* [[TMP_VAL1]], |
128 | // LAMBDA-DAG: [[TMP_ADD:%.+]] = fadd double [[TMP_VAL_VAL1]], 1.{{.+}} |
129 | // LAMBDA-DAG: store{{.*}} double [[TMP_ADD]], double* [[TMP_VAL1]], |
130 | // LAMBDA-DAG: [[SVAR_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[SVAR_PRIVATE]], |
131 | // LAMBDA-DAG: [[SVAR_ADD:%.+]] = add{{.*}} i{{[0-9]+}} [[SVAR_VAL]], 3 |
132 | // LAMBDA-DAG: store i{{[0-9]+}} [[SVAR_ADD]], i{{[0-9]+}}* [[SVAR_PRIVATE]], |
133 | // LAMBDA-DAG: [[SFVAR_VAL:%.+]] = load float, float* [[SFVAR_PRIVATE]], |
134 | // LAMBDA-DAG: [[SFVAR_CONV_VAL1:%.+]] = fpext float [[SFVAR_VAL]] to double |
135 | // LAMBDA-DAG: [[SFVAR_ADD:%.+]] = fadd double [[SFVAR_CONV_VAL1]], 4.{{.+}} |
136 | // LAMBDA-DAG: [[SFVAR_CONV_VAL2:%.+]] = fptrunc double [[SFVAR_ADD]] to float |
137 | // LAMBDA-DAG: store float [[SFVAR_CONV_VAL2:%.+]], float* [[SFVAR_PRIVATE]], |
138 | |
139 | // call inner lambda (use refs to private alloca's) |
140 | // LAMBDA: [[GEP_0:%.+]] = getelementptr{{.+}}, i{{[0-9]+}} 0, i{{[0-9]+}} 0 |
141 | // LAMBDA: store double* [[G_PRIVATE]], double** [[GEP_0]], |
142 | // LAMBDA: [[GEP_1:%.+]] = getelementptr{{.+}}, i{{[0-9]+}} 0, i{{[0-9]+}} 1 |
143 | // LAMBDA: [[TMP_PAR:%.+]] = load double*, double** [[TMP_PRIVATE]], |
144 | // LAMBDA: store double* [[TMP_PAR]], double** [[GEP_1]], |
145 | // LAMBDA: [[GEP_2:%.+]] = getelementptr{{.+}}, i{{[0-9]+}} 0, i{{[0-9]+}} 2 |
146 | // LAMBDA: store i{{[0-9]+}}* [[SVAR_PRIVATE]], i{{[0-9]+}}** [[GEP_2]], |
147 | // LAMBDA: [[GEP_3:%.+]] = getelementptr{{.+}}, i{{[0-9]+}} 0, i{{[0-9]+}} 3 |
148 | // LAMBDA: store float* [[SFVAR_PRIVATE]], float** [[GEP_3]], |
149 | // LAMBDA: call{{.*}} void [[INNER_LAMBDA:@.+]](%{{.+}}* {{.+}}) |
150 | // LAMBDA: call {{.*}}void @__kmpc_for_static_fini( |
151 | [&]() { |
152 | // LAMBDA: define {{.+}} void [[INNER_LAMBDA]](%{{.+}}* [[ARG_PTR:%.+]]) |
153 | // LAMBDA: store %{{.+}}* [[ARG_PTR]], %{{.+}}** [[ARG_PTR_REF:%.+]], |
154 | g += 2; |
155 | g1 += 2; |
156 | svar += 4; |
157 | sfvar += 8.0; |
158 | // LAMBDA-DAG: [[ARG_PTR:%.+]] = load %{{.+}}*, %{{.+}}** [[ARG_PTR_REF]] |
159 | // LAMBDA-DAG: [[G_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 |
160 | // LAMBDA-DAG: [[G_REF:%.+]] = load double*, double** [[G_PTR_REF]], |
161 | // LAMBDA-DAG: [[G_REF_VAL:%.+]] = load double, double* [[G_REF]], |
162 | // LAMBDA-DAG: [[G_REF_ADD:%.+]] = fadd double [[G_REF_VAL]], 2.{{.+}} |
163 | // LAMBDA-DAG: store double [[G_REF_ADD]], double* [[G_REF]] |
164 | |
165 | // LAMBDA-DAG: [[TMP_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 |
166 | // LAMBDA-DAG: [[G1_REF:%.+]] = load double*, double** [[TMP_PTR_REF]] |
167 | // LAMBDA-DAG: [[G1_REF_VAL:%.+]] = load double, double* [[G1_REF]], |
168 | // LAMBDA-DAG: [[G1_ADD:%.+]] = fadd double [[G1_REF_VAL]], 2.{{.+}} |
169 | // LAMBDA-DAG: store double [[G1_ADD]], double* [[G1_REF]], |
170 | |
171 | // LAMBDA-DAG: [[SVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 |
172 | // LAMBDA-DAG: [[SVAR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[SVAR_PTR_REF]] |
173 | // LAMBDA-DAG: [[SVAR_REF_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[SVAR_REF]] |
174 | // LAMBDA-DAG: [[SVAR_ADD:%.+]] = add{{.*}} i{{[0-9]+}} [[SVAR_REF_VAL]], 4 |
175 | // LAMBDA-DAG: store i{{[0-9]+}} [[SVAR_ADD]], i{{[0-9]+}}* [[SVAR_REF]] |
176 | |
177 | // LAMBDA-DAG: [[SFVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 3 |
178 | // LAMBDA-DAG: [[SFVAR_REF:%.+]] = load float*, float** [[SFVAR_PTR_REF]] |
179 | // LAMBDA-DAG: [[SFVAR_REF_VAL:%.+]] = load float, float* [[SFVAR_REF]] |
180 | // LAMBDA-DAG: [[SFVAR_REF_CONV:%.+]] = fpext float [[SFVAR_REF_VAL]] to double |
181 | // LAMBDA-DAG: [[SFVAR_ADD:%.+]] = fadd double [[SFVAR_REF_CONV]], 8.{{.+}} |
182 | // LAMBDA-DAG: [[SFVAR_ADD_CONV:%.+]] = fptrunc double [[SFVAR_ADD]] to float |
183 | // LAMBDA-DAG: store float [[SFVAR_ADD_CONV]], float* [[SFVAR_REF]], |
184 | }(); |
185 | } |
186 | }(); |
187 | return 0; |
188 | #else |
189 | S<float> test; |
190 | int t_var = 0; |
191 | int vec[] = {1, 2}; |
192 | S<float> s_arr[] = {1, 2}; |
193 | S<float> &var = test; |
194 | |
195 | #pragma omp target |
196 | #pragma omp teams |
197 | #pragma omp distribute firstprivate(t_var, vec, s_arr, s_arr, var, var, svar) |
198 | for (int i = 0; i < 2; ++i) { |
199 | vec[i] = t_var; |
200 | s_arr[i] = var; |
201 | } |
202 | return tmain<int>(); |
203 | #endif |
204 | } |
205 | |
206 | // CHECK: define{{.*}} i{{[0-9]+}} @main() |
207 | // CHECK: [[TEST:%.+]] = alloca [[S_FLOAT_TY]], |
208 | // CHECK: call {{.*}} [[S_FLOAT_TY_DEF_CONSTR:@.+]]([[S_FLOAT_TY]]* [[TEST]]) |
209 | // CHECK: call i{{[0-9]+}} @__tgt_target_teams( |
210 | // CHECK: call void [[OFFLOAD_FUN:@.+]]( |
211 | // CHECK: ret |
212 | |
213 | // CHECK: define{{.+}} [[OFFLOAD_FUN]]( |
214 | // CHECK: call void (%{{.+}}*, i{{[0-9]+}}, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)*, ...) @__kmpc_fork_teams(%{{.+}}* @{{.+}}, i{{[0-9]+}} 5, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)* bitcast (void (i{{[0-9]+}}*, i{{[0-9]+}}*, i{{[0-9]+}}*, [2 x i{{[0-9]+}}]*, [2 x [[S_FLOAT_TY]]]*, [[S_FLOAT_TY]]*, i{{[0-9]+}}*)* [[OMP_OUTLINED:@.+]] to void |
215 | // CHECK: ret |
216 | // |
217 | // CHECK: define internal void [[OMP_OUTLINED]](i{{[0-9]+}}*{{.+}}, i{{[0-9]+}}*{{.+}}, i{{[0-9]+}}*{{.+}} [[T_VAR_IN:%.+]], [2 x i{{[0-9]+}}]*{{.*}} [[VEC_IN:%.+]], [2 x [[S_FLOAT_TY]]]*{{.*}} [[S_ARR_IN:%.+]], [[S_FLOAT_TY]]*{{.*}} [[VAR_IN:%.+]], i{{[0-9]+}}*{{.+}} [[SVAR_IN:%.+]]) |
218 | |
219 | // CHECK: alloca i{{[0-9]+}}*, |
220 | // CHECK: alloca i{{[0-9]+}}*, |
221 | // CHECK: [[T_VAR_ADDR:%.+]] = alloca i{{[0-9]+}}*, |
222 | // CHECK: [[VEC_ADDR:%.+]] = alloca [2 x i{{[0-9]+}}]*, |
223 | // CHECK: [[S_ARR_ADDR:%.+]] = alloca [2 x [[S_FLOAT_TY]]]*, |
224 | // CHECK: [[VAR_ADDR:%.+]] = alloca [[S_FLOAT_TY]]*, |
225 | // CHECK: [[SVAR_ADDR:%.+]] = alloca i{{[0-9]+}}*, |
226 | // CHECK: [[TMP:%.+]] = alloca [[S_FLOAT_TY]]*, |
227 | |
228 | // discard omp loop variables |
229 | // CHECK: {{.*}} = alloca i{{[0-9]+}}, |
230 | // CHECK: {{.*}} = alloca i{{[0-9]+}}, |
231 | // CHECK: {{.*}} = alloca i{{[0-9]+}}, |
232 | // CHECK: {{.*}} = alloca i{{[0-9]+}}, |
233 | // CHECK: {{.*}} = alloca i{{[0-9]+}}, |
234 | // CHECK: {{.*}} = alloca i{{[0-9]+}}, |
235 | |
236 | // CHECK-DAG: [[T_VAR_PRIV:%.+]] = alloca i{{[0-9]+}}, |
237 | // CHECK-DAG: [[VEC_PRIV:%.+]] = alloca [2 x i{{[0-9]+}}], |
238 | // CHECK-DAG: [[S_ARR_PRIV:%.+]] = alloca [2 x [[S_FLOAT_TY]]], |
239 | // CHECK-DAG: [[VAR_PRIV:%.+]] = alloca [[S_FLOAT_TY]], |
240 | // CHECK-DAG: [[TMP_PRIV:%.+]] = alloca [[S_FLOAT_TY]]*, |
241 | // CHECK: [[SVAR_PRIV:%.+]] = alloca i{{[0-9]+}}, |
242 | |
243 | // CHECK: store i{{[0-9]+}}* [[T_VAR_IN]], i{{[0-9]+}}** [[T_VAR_ADDR]], |
244 | // CHECK: store [2 x i{{[0-9]+}}]* [[VEC_IN]], [2 x i{{[0-9]+}}]** [[VEC_ADDR]], |
245 | // CHECK: store [2 x [[S_FLOAT_TY]]]* [[S_ARR_IN]], [2 x [[S_FLOAT_TY]]]** [[S_ARR_ADDR]], |
246 | // CHECK: store [[S_FLOAT_TY]]* [[VAR_IN]], [[S_FLOAT_TY]]** [[VAR_ADDR]], |
247 | // CHECK: store i{{[0-9]+}}* [[SVAR_IN]], i{{[0-9]+}}** [[SVAR_ADDR]], |
248 | |
249 | // init t_var |
250 | // CHECK-DAG: [[T_VAR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[T_VAR_ADDR]], |
251 | // CHECK-DAG: [[T_VAR_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[T_VAR_REF]], |
252 | // CHECK-DAG: store i{{[0-9]+}} [[T_VAR_VAL]], i{{[0-9]+}}* [[T_VAR_PRIV]], |
253 | |
254 | // init vec |
255 | // CHECK-DAG: [[VEC_ADDR_VAL:%.+]] = load [2 x i{{[0-9]+}}]*, [2 x i{{[0-9]+}}]** [[VEC_ADDR]], |
256 | // CHECK-DAG: [[VEC_ADDR_VAL_BCAST:%.+]] = bitcast [2 x i{{[0-9]+}}]* [[VEC_ADDR_VAL]] to i{{[0-9]+}}* |
257 | // CHECK-DAG: [[VEC_PRIV_BCAST:%.+]] = bitcast [2 x i{{[0-9]+}}]* [[VEC_PRIV]] to i{{[0-9]+}}* |
258 | // CHECK-DAG: call void @llvm.memcpy.{{.*}}(i{{[0-9]+}}* align {{[0-9]+}} [[VEC_PRIV_BCAST]], i{{[0-9]+}}* align {{[0-9]+}} [[VEC_ADDR_VAL_BCAST]],{{.+}}) |
259 | |
260 | // init s_arr |
261 | // CHECK-DAG: [[S_ARR_ADDR_VAL:%.+]] = load [2 x [[S_FLOAT_TY]]]*, [2 x [[S_FLOAT_TY]]]** [[S_ARR_ADDR]], |
262 | // CHECK-DAG: [[S_ARR_ADDR_BCAST:%.+]] = bitcast [2 x [[S_FLOAT_TY]]]* [[S_ARR_ADDR_VAL]] to [[S_FLOAT_TY]]* |
263 | // CHECK-DAG: [[S_ARR_PRIV_BGN:%.+]] = getelementptr{{.+}} [2 x [[S_FLOAT_TY]]], [2 x [[S_FLOAT_TY]]]* [[S_ARR_PRIV]]{{.+}} |
264 | // CHECK-DAG: [[S_ARR_PRIV_NEXT:%.+]] = getelementptr [[S_FLOAT_TY]], [[S_FLOAT_TY]]* [[S_ARR_PRIV_BGN]]{{.+}} |
265 | // CHECK-DAG: [[S_ARR_IS_EMPTY:%.+]] = icmp eq [[S_FLOAT_TY]]* [[S_ARR_PRIV_BGN]], [[S_ARR_PRIV_NEXT]] |
266 | // CHECK-DAG: br i1 [[S_ARR_IS_EMPTY]], label %[[S_ARR_CPY_DONE:.+]], label %[[S_ARR_CPY_BODY:.+]] |
267 | |
268 | // CHECK-DAG: [[S_ARR_CPY_BODY]]: |
269 | // CHECK-DAG: [[S_ARR_SRC_PAST:%.+]] = phi{{.+}} [ [[S_ARR_ADDR_BCAST]],{{.+}} ], [ [[S_ARR_SRC:%.+]],{{.+}} ] |
270 | // CHECK-DAG: [[S_ARR_DST_PAST:%.+]] = phi{{.+}} [ [[S_ARR_PRIV_BGN]],{{.+}} ], [ [[S_ARR_DST:%.+]],{{.+}} ] |
271 | // CHECK-DAG: [[S_ARR_SRC_BCAST:%.+]] = bitcast{{.+}} [[S_ARR_SRC_PAST]] to{{.+}} |
272 | // CHECK-DAG: [[S_ARR_DST_BCAST:%.+]] = bitcast{{.+}} [[S_ARR_DST_PAST]] to{{.+}} |
273 | // CHECK-DAG: call{{.+}} @llvm.memcpy.{{.+}}({{.+}}* align {{[0-9]+}} [[S_ARR_DST_BCAST]], {{.+}}* align {{[0-9]+}} [[S_ARR_SRC_BCAST]]{{.+}}) |
274 | // CHECK-DAG: [[S_ARR_SRC]] = getelementptr{{.+}} |
275 | // CHECK-DAG: [[S_ARR_DST]] = getelementptr{{.+}} |
276 | // CHECK-DAG: [[S_ARR_CPY_FIN:%.+]] = icmp{{.+}} [[S_ARR_DST]], [[S_ARR_PRIV_NEXT]] |
277 | // CHECK-DAG: br i1 [[S_ARR_CPY_FIN]], label %[[S_ARR_CPY_DONE]], label %[[S_ARR_CPY_BODY]] |
278 | // CHECK-DAG: [[S_ARR_CPY_DONE]]: |
279 | |
280 | // init var |
281 | // CHECK-DAG: [[VAR_ADDR_VAL:%.+]] = load [[S_FLOAT_TY]]*, [[S_FLOAT_TY]]** [[VAR_ADDR]], |
282 | // CHECK-DAG: store{{.+}} [[VAR_ADDR_VAL]],{{.+}} [[TMP]], |
283 | // CHECK-DAG: [[TMP_VAL:%.+]] = load [[S_FLOAT_TY]]*, [[S_FLOAT_TY]]** [[TMP]], |
284 | // CHECK-DAG: [[VAR_PRIV_BCAST:%.+]] = bitcast [[S_FLOAT_TY]]* [[VAR_PRIV]] to{{.+}} |
285 | // CHECK-DAG: [[TMP_BCAST:%.+]] = bitcast [[S_FLOAT_TY]]* [[TMP_VAL]] to{{.+}} |
286 | // CHECK-DAG: call{{.+}} @llvm.memcpy.{{.+}}({{.+}}* align {{[0-9]+}} [[VAR_PRIV_BCAST]], {{.+}}* align {{[0-9]+}} [[TMP_BCAST]],{{.+}}) |
287 | // CHECK-DAG: store [[S_FLOAT_TY]]* [[VAR_PRIV]], [[S_FLOAT_TY]]** [[TMP_PRIV]], |
288 | |
289 | // init svar |
290 | // CHECK-DAG: [[SVAR_REF:%.+]] = load{{.+}},{{.+}} [[SVAR_ADDR]], |
291 | // CHECK-DAG: [[SVAR_VAL:%.+]] = load{{.+}},{{.+}} [[SVAR_REF]], |
292 | // CHECK-DAG: store{{.+}} [[SVAR_VAL]],{{.+}} [[SVAR_PRIV]], |
293 | |
294 | // CHECK-DAG: store i{{[0-9]+}} 0, i{{[0-9]+}}* %.omp{{.+}}, |
295 | // CHECK-DAG: store i{{[0-9]+}} 1, i{{[0-9]+}}* %.omp{{.+}}, |
296 | // CHECK-DAG: store i{{[0-9]+}} 1, i{{[0-9]+}}* %.omp{{.+}}, |
297 | // CHECK-DAG: store i{{[0-9]+}} 0, i{{[0-9]+}}* %.omp{{.+}}, |
298 | |
299 | // CHECK: call void @__kmpc_for_static_init_4( |
300 | // CHECK: call void @__kmpc_for_static_fini( |
301 | // CHECK: ret void |
302 | |
303 | // Template |
304 | // CHECK: define{{.*}} i{{[0-9]+}} [[TMAIN_INT:@.+]]() |
305 | // CHECK: [[TEST:%.+]] = alloca [[S_INT_TY]], |
306 | // CHECK: call {{.*}} [[S_INT_TY_DEF_CONSTR:@.+]]([[S_INT_TY]]* [[TEST]]) |
307 | // CHECK: call i{{[0-9]+}} @__tgt_target_teams( |
308 | // CHECK: call void [[OFFLOAD_FUN_1:@.+]]( |
309 | // CHECK: ret |
310 | |
311 | // CHECK: define{{.+}} [[OFFLOAD_FUN_1]]( |
312 | // CHECK: call void (%{{.+}}*, i{{[0-9]+}}, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)*, ...) @__kmpc_fork_teams(%{{.+}}* @{{.+}}, i{{[0-9]+}} 4, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)* bitcast (void (i{{[0-9]+}}*, i{{[0-9]+}}*, i{{[0-9]+}}*, [2 x i{{[0-9]+}}]*, [2 x [[S_INT_TY]]]*, [[S_INT_TY]]*)* [[OMP_OUTLINED_1:@.+]] to void |
313 | // CHECK: ret |
314 | // |
315 | // CHECK: define internal void [[OMP_OUTLINED_1]](i{{[0-9]+}}*{{.+}}, i{{[0-9]+}}*{{.+}}, i{{[0-9]+}}*{{.+}} [[T_VAR_IN:%.+]], [2 x i{{[0-9]+}}]*{{.*}} [[VEC_IN:%.+]], [2 x [[S_INT_TY]]]*{{.*}} [[S_ARR_IN:%.+]], [[S_INT_TY]]*{{.*}} [[VAR_IN:%.+]]) |
316 | |
317 | // CHECK: alloca i{{[0-9]+}}*, |
318 | // CHECK: alloca i{{[0-9]+}}*, |
319 | // CHECK: [[T_VAR_ADDR:%.+]] = alloca i{{[0-9]+}}*, |
320 | // CHECK: [[VEC_ADDR:%.+]] = alloca [2 x i{{[0-9]+}}]*, |
321 | // CHECK: [[S_ARR_ADDR:%.+]] = alloca [2 x [[S_INT_TY]]]*, |
322 | // CHECK: [[VAR_ADDR:%.+]] = alloca [[S_INT_TY]]*, |
323 | // CHECK: [[TMP:%.+]] = alloca [[S_INT_TY]]*, |
324 | |
325 | // discard omp loop variables |
326 | // CHECK: {{.*}} = alloca i{{[0-9]+}}, |
327 | // CHECK: {{.*}} = alloca i{{[0-9]+}}, |
328 | // CHECK: {{.*}} = alloca i{{[0-9]+}}, |
329 | // CHECK: {{.*}} = alloca i{{[0-9]+}}, |
330 | // CHECK: {{.*}} = alloca i{{[0-9]+}}, |
331 | // CHECK: {{.*}} = alloca i{{[0-9]+}}, |
332 | |
333 | // CHECK-DAG: [[T_VAR_PRIV:%.+]] = alloca i{{[0-9]+}}, |
334 | // CHECK-DAG: [[VEC_PRIV:%.+]] = alloca [2 x i{{[0-9]+}}], |
335 | // CHECK-DAG: [[S_ARR_PRIV:%.+]] = alloca [2 x [[S_INT_TY]]], |
336 | // CHECK-DAG: [[VAR_PRIV:%.+]] = alloca [[S_INT_TY]], |
337 | // CHECK-DAG: [[TMP_PRIV:%.+]] = alloca [[S_INT_TY]]*, |
338 | |
339 | // CHECK: store i{{[0-9]+}}* [[T_VAR_IN]], i{{[0-9]+}}** [[T_VAR_ADDR]], |
340 | // CHECK: store [2 x i{{[0-9]+}}]* [[VEC_IN]], [2 x i{{[0-9]+}}]** [[VEC_ADDR]], |
341 | // CHECK: store [2 x [[S_INT_TY]]]* [[S_ARR_IN]], [2 x [[S_INT_TY]]]** [[S_ARR_ADDR]], |
342 | // CHECK: store [[S_INT_TY]]* [[VAR_IN]], [[S_INT_TY]]** [[VAR_ADDR]], |
343 | |
344 | // init t_var |
345 | // CHECK-DAG: [[T_VAR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[T_VAR_ADDR]], |
346 | // CHECK-DAG: [[T_VAR:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[T_VAR_REF]], |
347 | // CHECK-DAG: store i{{[0-9]+}} [[T_VAR]], i{{[0-9]+}}* [[T_VAR_PRIV]], |
348 | |
349 | // init vec |
350 | // CHECK-DAG: [[VEC_ADDR_VAL:%.+]] = load [2 x i{{[0-9]+}}]*, [2 x i{{[0-9]+}}]** [[VEC_ADDR]], |
351 | // CHECK-DAG: [[VEC_ADDR_VAL_BCAST:%.+]] = bitcast [2 x i{{[0-9]+}}]* [[VEC_ADDR_VAL]] to i{{[0-9]+}}* |
352 | // CHECK-DAG: [[VEC_PRIV_BCAST:%.+]] = bitcast [2 x i{{[0-9]+}}]* [[VEC_PRIV]] to i{{[0-9]+}}* |
353 | // CHECK-DAG: call void @llvm.memcpy.{{.*}}(i{{[0-9]+}}* align {{[0-9]+}} [[VEC_PRIV_BCAST]], i{{[0-9]+}}* align {{[0-9]+}} [[VEC_ADDR_VAL_BCAST]],{{.+}}) |
354 | |
355 | // init s_arr |
356 | // CHECK-DAG: [[S_ARR_ADDR_VAL:%.+]] = load [2 x [[S_INT_TY]]]*, [2 x [[S_INT_TY]]]** [[S_ARR_ADDR]], |
357 | // CHECK-DAG: [[S_ARR_ADDR_BCAST:%.+]] = bitcast [2 x [[S_INT_TY]]]* [[S_ARR_ADDR_VAL]] to [[S_INT_TY]]* |
358 | // CHECK-DAG: [[S_ARR_PRIV_BGN:%.+]] = getelementptr{{.+}} [2 x [[S_INT_TY]]], [2 x [[S_INT_TY]]]* [[S_ARR_PRIV]]{{.+}} |
359 | // CHECK-DAG: [[S_ARR_PRIV_NEXT:%.+]] = getelementptr [[S_INT_TY]], [[S_INT_TY]]* [[S_ARR_PRIV_BGN]]{{.+}} |
360 | // CHECK-DAG: [[S_ARR_IS_EMPTY:%.+]] = icmp eq [[S_INT_TY]]* [[S_ARR_PRIV_BGN]], [[S_ARR_PRIV_NEXT]] |
361 | // CHECK-DAG: br i1 [[S_ARR_IS_EMPTY]], label %[[S_ARR_CPY_DONE:.+]], label %[[S_ARR_CPY_BODY:.+]] |
362 | |
363 | // CHECK-DAG: [[S_ARR_CPY_BODY]]: |
364 | // CHECK-DAG: [[S_ARR_SRC_PAST:%.+]] = phi{{.+}} [ [[S_ARR_ADDR_BCAST]],{{.+}} ], [ [[S_ARR_SRC:%.+]],{{.+}} ] |
365 | // CHECK-DAG: [[S_ARR_DST_PAST:%.+]] = phi{{.+}} [ [[S_ARR_PRIV_BGN]],{{.+}} ], [ [[S_ARR_DST:%.+]],{{.+}} ] |
366 | // CHECK-DAG: [[S_ARR_SRC_BCAST:%.+]] = bitcast{{.+}} [[S_ARR_SRC_PAST]] to{{.+}} |
367 | // CHECK-DAG: [[S_ARR_DST_BCAST:%.+]] = bitcast{{.+}} [[S_ARR_DST_PAST]] to{{.+}} |
368 | // CHECK-DAG: call{{.+}} @llvm.memcpy.{{.+}}({{.+}}* align {{[0-9]+}} [[S_ARR_DST_BCAST]], {{.+}}* align {{[0-9]+}} [[S_ARR_SRC_BCAST]]{{.+}}) |
369 | // CHECK-DAG: [[S_ARR_SRC]] = getelementptr{{.+}} |
370 | // CHECK-DAG: [[S_ARR_DST]] = getelementptr{{.+}} |
371 | // CHECK-DAG: [[S_ARR_CPY_FIN:%.+]] = icmp{{.+}} [[S_ARR_DST]], [[S_ARR_PRIV_NEXT]] |
372 | // CHECK-DAG: br i1 [[S_ARR_CPY_FIN]], label %[[S_ARR_CPY_DONE]], label %[[S_ARR_CPY_BODY]] |
373 | // CHECK-DAG: [[S_ARR_CPY_DONE]]: |
374 | |
375 | // init var |
376 | // CHECK-DAG: [[VAR_ADDR_VAL:%.+]] = load [[S_INT_TY]]*, [[S_INT_TY]]** [[VAR_ADDR]], |
377 | // CHECK-DAG: store{{.+}} [[VAR_ADDR_VAL]],{{.+}} [[TMP]], |
378 | // CHECK-DAG: [[TMP_VAL:%.+]] = load [[S_INT_TY]]*, [[S_INT_TY]]** [[TMP]], |
379 | // CHECK-DAG: [[VAR_PRIV_BCAST:%.+]] = bitcast [[S_INT_TY]]* [[VAR_PRIV]] to{{.+}} |
380 | // CHECK-DAG: [[TMP_BCAST:%.+]] = bitcast [[S_INT_TY]]* [[TMP_VAL]] to{{.+}} |
381 | // CHECK-DAG: call{{.+}} @llvm.memcpy.{{.+}}({{.+}}* align {{[0-9]+}} [[VAR_PRIV_BCAST]], {{.+}}* align {{[0-9]+}} [[TMP_BCAST]],{{.+}}) |
382 | // CHECK-DAG: store [[S_INT_TY]]* [[VAR_PRIV]], [[S_INT_TY]]** [[TMP_PRIV]], |
383 | |
384 | // CHECK-DAG: store i{{[0-9]+}} 0, i{{[0-9]+}}* %.omp{{.+}}, |
385 | // CHECK-DAG: store i{{[0-9]+}} 1, i{{[0-9]+}}* %.omp{{.+}}, |
386 | // CHECK-DAG: store i{{[0-9]+}} 1, i{{[0-9]+}}* %.omp{{.+}}, |
387 | // CHECK-DAG: store i{{[0-9]+}} 0, i{{[0-9]+}}* %.omp{{.+}}, |
388 | |
389 | // CHECK: call void @__kmpc_for_static_init_4( |
390 | // CHECK: call void @__kmpc_for_static_fini( |
391 | // CHECK: ret void |
392 | #endif |
393 | |