xref: /aosp_15_r20/external/clang/test/OpenMP/teams_private_codegen.cpp (revision 67e74705e28f6214e480b399dd47ea732279e315)
1*67e74705SXin Li // 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*67e74705SXin Li // 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*67e74705SXin Li // 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*67e74705SXin Li // 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*67e74705SXin Li // 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*67e74705SXin Li // 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*67e74705SXin Li 
8*67e74705SXin Li // 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*67e74705SXin Li // 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*67e74705SXin Li // 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*67e74705SXin Li // 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*67e74705SXin Li // 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*67e74705SXin Li // 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*67e74705SXin Li // expected-no-diagnostics
15*67e74705SXin Li #ifndef HEADER
16*67e74705SXin Li #define HEADER
17*67e74705SXin Li template <class T>
18*67e74705SXin Li struct S {
19*67e74705SXin Li   T f;
SS20*67e74705SXin Li   S(T a) : f(a) {}
SS21*67e74705SXin Li   S() : f() {}
operator TS22*67e74705SXin Li   operator T() { return T(); }
~SS23*67e74705SXin Li   ~S() {}
24*67e74705SXin Li };
25*67e74705SXin Li 
26*67e74705SXin Li volatile int g __attribute__((aligned(128))) = 1212;
27*67e74705SXin Li 
28*67e74705SXin Li struct SS {
29*67e74705SXin Li   int a;
30*67e74705SXin Li   int b : 4;
31*67e74705SXin Li   int &c;
SSSS32*67e74705SXin Li   SS(int &d) : a(0), b(0), c(d) {
33*67e74705SXin Li #pragma omp target
34*67e74705SXin Li #pragma omp teams private(a, b, c)
35*67e74705SXin Li #ifdef LAMBDA
36*67e74705SXin Li     [&]() {
37*67e74705SXin Li       ++this->a, --b, (this)->c /= 1;
38*67e74705SXin Li     }();
39*67e74705SXin Li #else
40*67e74705SXin Li     ++this->a, --b, c /= 1;
41*67e74705SXin Li #endif
42*67e74705SXin Li   }
43*67e74705SXin Li };
44*67e74705SXin Li 
45*67e74705SXin Li template<typename T>
46*67e74705SXin Li struct SST {
47*67e74705SXin Li   T a;
SSTSST48*67e74705SXin Li   SST() : a(T()) {
49*67e74705SXin Li #pragma omp target
50*67e74705SXin Li #pragma omp teams private(a)
51*67e74705SXin Li #ifdef LAMBDA
52*67e74705SXin Li     [&]() {
53*67e74705SXin Li       [&]() {
54*67e74705SXin Li         ++this->a;
55*67e74705SXin Li       }();
56*67e74705SXin Li     }();
57*67e74705SXin Li #else
58*67e74705SXin Li     ++(this)->a;
59*67e74705SXin Li #endif
60*67e74705SXin Li   }
61*67e74705SXin Li };
62*67e74705SXin Li 
63*67e74705SXin Li // CHECK: [[SS_TY:%.+]] = type { i{{[0-9]+}}, i8
64*67e74705SXin Li // LAMBDA: [[SS_TY:%.+]] = type { i{{[0-9]+}}, i8
65*67e74705SXin Li // LAMBDA: [[CAP_0_TY:%.+]] = type { [[SS_TY]]*, i{{[0-9]+}}*,
66*67e74705SXin Li // LAMBDA: [[CAP_1_TY:%.+]] = type { i{{[0-9]+}}*, i{{[0-9]+}}* }
67*67e74705SXin Li // CHECK: [[S_FLOAT_TY:%.+]] = type { float }
68*67e74705SXin Li // CHECK: [[S_INT_TY:%.+]] = type { i{{[0-9]+}} }
69*67e74705SXin Li // CHECK: [[SST_TY:%.+]] = type { i{{[0-9]+}} }
70*67e74705SXin Li template <typename T>
tmain()71*67e74705SXin Li T tmain() {
72*67e74705SXin Li   S<T> test;
73*67e74705SXin Li   SST<T> sst;
74*67e74705SXin Li   T t_var __attribute__((aligned(128))) = T();
75*67e74705SXin Li   T vec[] __attribute__((aligned(128))) = {1, 2};
76*67e74705SXin Li   S<T> s_arr[] __attribute__((aligned(128))) = {1, 2};
77*67e74705SXin Li   S<T> var __attribute__((aligned(128))) (3);
78*67e74705SXin Li #pragma omp target
79*67e74705SXin Li #pragma omp teams private(t_var, vec, s_arr, var)
80*67e74705SXin Li   {
81*67e74705SXin Li     vec[0] = t_var;
82*67e74705SXin Li     s_arr[0] = var;
83*67e74705SXin Li   }
84*67e74705SXin Li   return T();
85*67e74705SXin Li }
86*67e74705SXin Li 
main()87*67e74705SXin Li int main() {
88*67e74705SXin Li   static int sivar;
89*67e74705SXin Li   SS ss(sivar);
90*67e74705SXin Li #ifdef LAMBDA
91*67e74705SXin Li   // LAMBDA: [[G:@.+]] = global i{{[0-9]+}} 1212,
92*67e74705SXin Li   // LAMBDA: define {{.+}} @main()
93*67e74705SXin Li   // LAMBDA: alloca [[SS_TY]],
94*67e74705SXin Li   // LAMBDA: alloca [[CAP_TY:%.+]],
95*67e74705SXin Li 
96*67e74705SXin Li   // LAMBDA: call{{.*}} [[ST_CONSTR_INIT:@.+]]([[SS_TY]]*
97*67e74705SXin Li   // LAMBDA: call{{.*}} void [[OUTER_LAMBDA:@[^(]+]]([[CAP_TY]]*
98*67e74705SXin Li 
99*67e74705SXin Li   // lambda and target region in main
100*67e74705SXin Li   // LAMBDA: define {{.+}} [[OUTER_LAMBDA]]([[CAP_TY]]* {{.+}})
101*67e74705SXin Li   // LAMBDA: call void @[[OMP_OFFLOADING:.+]](i{{[0-9]+}}* {{.+}}, i{{[0-9]+}} {{.+}}
102*67e74705SXin Li 
103*67e74705SXin Li   // target region in struct constructor
104*67e74705SXin Li   // LAMBDA: define{{.*}} void [[ST_CONSTR:@.+]]([[SS_TY]]* %this,
105*67e74705SXin Li   // LAMBDA: call void [[OMP_OFFLOADING_1:@.+]]([[SS_TY]]
106*67e74705SXin Li 
107*67e74705SXin Li   // offloading function in struct constructor
108*67e74705SXin Li   // LAMBDA: define{{.*}} void [[OMP_OFFLOADING_1]]([[SS_TY]]
109*67e74705SXin Li   // LAMBDA: call void (%{{.+}}*, i{{[0-9]+}}, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)*, ...) @__kmpc_fork_teams(%{{.+}}* @{{.+}}, i{{[0-9]+}} 1, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)* bitcast (void (i{{[0-9]+}}*, i{{[0-9]+}}*, [[SS_TY]]*)* [[OMP_OUTLINED:@.+]] to void
110*67e74705SXin Li 
111*67e74705SXin Li   // outlined teams region in struct constructor
112*67e74705SXin Li   // LAMBDA: define{{.*}} void [[OMP_OUTLINED]](i{{[0-9]+}}* {{.+}}, i{{[0-9]+}}* {{.+}}, [[SS_TY]]*
113*67e74705SXin Li   // LAMBDA: [[THIS_ADDR:%.+]] = alloca [[SS_TY]]*,
114*67e74705SXin Li   // LAMBDA: [[A_PRIV:%.+]] = alloca i{{[0-9]+}},
115*67e74705SXin Li   // LAMBDA: [[B_PRIV:%.+]] = alloca i{{[0-9]+}},
116*67e74705SXin Li   // LAMBDA: [[C_PRIV:%.+]] = alloca i{{[0-9]+}},
117*67e74705SXin Li   // LAMBDA: [[THIS_REF:%.+]] = load [[SS_TY]]*, [[SS_TY]]** [[THIS_ADDR]],
118*67e74705SXin Li   // LAMBDA: store i{{[0-9]+}}* [[A_PRIV]], i{{[0-9]+}}** [[A_TMP_REF:%.+]],
119*67e74705SXin Li   // LAMBDA: store i{{[0-9]+}}* [[C_PRIV]], i{{[0-9]+}}** [[C_TMP_REF:%.+]],
120*67e74705SXin Li   // LAMBDA: [[CAP_THIS_REF:%.+]] = getelementptr {{.+}} [[CAP_0_TY]], [[CAP_0_TY]]* {{.+}}, i{{[0-9]+}} 0, i{{[0-9]+}} 0
121*67e74705SXin Li   // LAMBDA: store [[SS_TY]]* [[THIS_REF]], [[SS_TY]]** [[CAP_THIS_REF]],
122*67e74705SXin Li   // LAMBDA: [[CAP_A_REF:%.+]] = getelementptr {{.+}} [[CAP_0_TY]], [[CAP_0_TY]]* {{.+}}, i{{[0-9]+}} 0, i{{[0-9]+}} 1
123*67e74705SXin Li   // LAMBDA: [[A_TMP_VAL:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[A_TMP_REF]],
124*67e74705SXin Li   // LAMBDA: store i{{[0-9]+}}* [[A_TMP_VAL]], i{{[0-9]+}}** [[CAP_A_REF]],
125*67e74705SXin Li   // LAMBDA: [[CAP_B_REF:%.+]] = getelementptr {{.+}} [[CAP_0_TY]], [[CAP_0_TY]]* {{.+}}, i{{[0-9]+}} 0, i{{[0-9]+}} 2
126*67e74705SXin Li   // LAMBDA: store i{{[0-9]+}}* [[B_PRIV]], i{{[0-9]+}}** [[CAP_B_REF]],
127*67e74705SXin Li   // LAMBDA: [[CAP_C_REF:%.+]] = getelementptr {{.+}} [[CAP_0_TY]], [[CAP_0_TY]]* {{.+}}, i{{[0-9]+}} 0, i{{[0-9]+}} 3
128*67e74705SXin Li   // LAMBDA: [[C_TMP_VAL:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[C_TMP_REF]],
129*67e74705SXin Li   // LAMBDA: store i{{[0-9]+}}* [[C_TMP_VAL]], i{{[0-9]+}}** [[CAP_C_REF]],
130*67e74705SXin Li   // call void [[INNER_LAMBDA_CONSTR:@.+]]([[CAP_0_TY]]*
131*67e74705SXin Li 
132*67e74705SXin Li   // inner lambda in struct constructor
133*67e74705SXin Li   // define{{.*}} void [[INNER_LAMBDA_CONSTR]]([[CAP_0_TY]]*
134*67e74705SXin Li   // LAMBDA: [[CAP_A_REF_1:%.+]] = getelementptr {{.+}} [[CAP_0_TY]], [[CAP_0_TY]]* {{.+}}, i{{[0-9]+}} 0, i{{[0-9]+}} 1
135*67e74705SXin Li   // LAMBDA: [[A_REF_FROM_CAP:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[CAP_A_REF_1]],
136*67e74705SXin Li   // LAMBDA: [[A_VAL_FROM_CAP:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_REF_FROM_CAP]],
137*67e74705SXin Li   // LAMBDA: [[A_INC_VAL:%.+]] = add {{.+}} i{{[0-9]+}} [[A_VAL_FROM_CAP]], 1
138*67e74705SXin Li   // LAMBDA: store i{{[0-9]+}} [[A_INC_VAL]], i{{[0-9]+}}* [[A_REF_FROM_CAP]],
139*67e74705SXin Li 
140*67e74705SXin Li   // LAMBDA: [[CAP_B_REF_1:%.+]] = getelementptr {{.+}} [[CAP_0_TY]], [[CAP_0_TY]]* {{.+}}, i{{[0-9]+}} 0, i{{[0-9]+}} 2
141*67e74705SXin Li   // LAMBDA: [[B_REF_FROM_CAP:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[CAP_B_REF_1]],
142*67e74705SXin Li   // LAMBDA: [[B_VAL_FROM_CAP:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[B_REF_FROM_CAP]],
143*67e74705SXin Li   // LAMBDA: [[B_DEC_VAL:%.+]] = add {{.+}} i{{[0-9]+}} [[B_VAL_FROM_CAP]], -1
144*67e74705SXin Li   // LAMBDA: store i{{[0-9]+}} [[B_DEC_VAL]], i{{[0-9]+}}* [[B_REF_FROM_CAP]],
145*67e74705SXin Li 
146*67e74705SXin Li   // LAMBDA: [[CAP_C_REF_1:%.+]] = getelementptr {{.+}} [[CAP_0_TY]], [[CAP_0_TY]]* {{.+}}, i{{[0-9]+}} 0, i{{[0-9]+}} 3
147*67e74705SXin Li   // LAMBDA: [[C_REF_FROM_CAP:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[CAP_C_REF_1]],
148*67e74705SXin Li   // LAMBDA: [[C_VAL_FROM_CAP:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[C_REF_FROM_CAP]],
149*67e74705SXin Li   // LAMBDA: [[C_DEC_VAL:%.+]] = sdiv{{.*}} i{{[0-9]+}} [[C_VAL_FROM_CAP]], 1
150*67e74705SXin Li   // LAMBDA: store i{{[0-9]+}} [[C_DEC_VAL]], i{{[0-9]+}}* [[C_REF_FROM_CAP]],
151*67e74705SXin Li   // ret
152*67e74705SXin Li 
153*67e74705SXin Li   [&]() {
154*67e74705SXin Li #pragma omp target
155*67e74705SXin Li #pragma omp teams private(g, sivar)
156*67e74705SXin Li   {
157*67e74705SXin Li     // LAMBDA: define{{.+}} @[[OMP_OFFLOADING]](i{{[0-9]+}}* {{.+}} [[G_IN:%.+]], i{{[0-9]+}} [[SIVAR_IN:%.+]]
158*67e74705SXin Li     // LAMBDA: call void (%{{.+}}*, i{{[0-9]+}}, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)*, ...) @__kmpc_fork_teams(%{{.+}}* @{{.+}}, i{{[0-9]+}} 0, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)* bitcast (void (i{{[0-9]+}}*, i{{[0-9]+}}*)* [[OMP_OUTLINED_1:@.+]] to void
159*67e74705SXin Li 
160*67e74705SXin Li     // LAMBDA: define {{.+}} [[OMP_OUTLINED_1]](i{{[0-9]+}}* {{.+}}, i{{[0-9]+}}* {{.+}}
161*67e74705SXin Li     // LAMBDA: [[G_LOC_OUTER:%.+]] = alloca i{{[0-9]+}},
162*67e74705SXin Li     // LAMBDA: [[SIVAR_LOC_OUTER:%.+]] = alloca i{{[0-9]+}},
163*67e74705SXin Li     // LAMBDA: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[G_LOC_OUTER]]
164*67e74705SXin Li     // LAMBDA: store i{{[0-9]+}} 2, i{{[0-9]+}}* [[SIVAR_LOC_OUTER]]
165*67e74705SXin Li     // LAMBDA: call{{.*}} void [[INNER_LAMBDA:@[^(]+]]([[CAP_1_TY]]*
166*67e74705SXin Li     // LAMBDA: ret
167*67e74705SXin Li     g = 1;
168*67e74705SXin Li     sivar = 2;
169*67e74705SXin Li     [&]() {
170*67e74705SXin Li       // LAMBDA: define {{.+}} [[INNER_LAMBDA]]([[CAP_1_TY]]* {{.+}})
171*67e74705SXin Li       g = 2;
172*67e74705SXin Li       sivar = 4;
173*67e74705SXin Li       // LAMBDA: store i{{[0-9]+}} 2, i{{[0-9]+}}*
174*67e74705SXin Li       // LAMBDA: store i{{[0-9]+}} 4, i{{[0-9]+}}*
175*67e74705SXin Li     }();
176*67e74705SXin Li   }
177*67e74705SXin Li   }();
178*67e74705SXin Li   return 0;
179*67e74705SXin Li #else
180*67e74705SXin Li   S<float> test;
181*67e74705SXin Li   int t_var = 0;
182*67e74705SXin Li   int vec[] = {1, 2};
183*67e74705SXin Li   S<float> s_arr[] = {1, 2};
184*67e74705SXin Li   S<float> var(3);
185*67e74705SXin Li #pragma omp target
186*67e74705SXin Li #pragma omp teams private(t_var, vec, s_arr, var, sivar)
187*67e74705SXin Li   {
188*67e74705SXin Li     vec[0] = t_var;
189*67e74705SXin Li     s_arr[0] = var;
190*67e74705SXin Li     sivar = 3;
191*67e74705SXin Li   }
192*67e74705SXin Li   return tmain<int>();
193*67e74705SXin Li #endif
194*67e74705SXin Li }
195*67e74705SXin Li 
196*67e74705SXin Li // CHECK: define{{.*}} i{{[0-9]+}} @main()
197*67e74705SXin Li // CHECK: [[TEST:%.+]] = alloca [[S_FLOAT_TY]],
198*67e74705SXin Li // CHECK: call {{.*}} [[S_FLOAT_TY_DEF_CONSTR:@.+]]([[S_FLOAT_TY]]* [[TEST]])
199*67e74705SXin Li // CHECK: [[OFF_IN:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* {{%.+}},
200*67e74705SXin Li // CHECK: call void @[[OMP_OFFLOADING:.+]](i{{[0-9]+}} [[OFF_IN]]
201*67e74705SXin Li // CHECK: = call{{.*}} i{{.+}} [[TMAIN_INT:@.+]]()
202*67e74705SXin Li // CHECK: call void [[S_FLOAT_TY_DESTR:@.+]]([[S_FLOAT_TY]]*
203*67e74705SXin Li // CHECK: ret
204*67e74705SXin Li 
205*67e74705SXin Li // target region in main function
206*67e74705SXin Li // CHECK: define{{.+}} @[[OMP_OFFLOADING]](i{{[0-9]+}}
207*67e74705SXin Li // CHECK: call void (%{{.+}}*, i{{[0-9]+}}, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)*, ...) @__kmpc_fork_teams(%{{.+}}* @{{.+}}, i{{[0-9]+}} 0, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)* bitcast (void (i{{[0-9]+}}*, i{{[0-9]+}}*)* [[OMP_OUTLINED:@.+]] to void
208*67e74705SXin Li 
209*67e74705SXin Li // CHECK: define internal void [[OMP_OUTLINED]](i{{[0-9]+}}* noalias [[GTID_ADDR:%.+]], i{{[0-9]+}}* noalias %{{.+}})
210*67e74705SXin Li // CHECK: [[T_VAR_PRIV:%.+]] = alloca i{{[0-9]+}},
211*67e74705SXin Li // CHECK: [[VEC_PRIV:%.+]] = alloca [2 x i{{[0-9]+}}],
212*67e74705SXin Li // CHECK: [[S_ARR_PRIV:%.+]] = alloca [2 x [[S_FLOAT_TY]]],
213*67e74705SXin Li // CHECK: [[VAR_PRIV:%.+]] = alloca [[S_FLOAT_TY]],
214*67e74705SXin Li // CHECK: [[SIVAR_PRIV:%.+]] = alloca i{{[0-9]+}},
215*67e74705SXin Li // CHECK: store i{{[0-9]+}}* [[GTID_ADDR]], i{{[0-9]+}}** [[GTID_ADDR_REF:%.+]]
216*67e74705SXin Li // CHECK-NOT: [[T_VAR_PRIV]]
217*67e74705SXin Li // CHECK-NOT: [[VEC_PRIV]]
218*67e74705SXin Li // CHECK: {{.+}}:
219*67e74705SXin Li // CHECK: [[S_ARR_PRIV_ITEM:%.+]] = phi [[S_FLOAT_TY]]*
220*67e74705SXin Li // CHECK: call {{.*}} [[S_FLOAT_TY_DEF_CONSTR]]([[S_FLOAT_TY]]* [[S_ARR_PRIV_ITEM]])
221*67e74705SXin Li // CHECK-NOT: [[T_VAR_PRIV]]
222*67e74705SXin Li // CHECK-NOT: [[VEC_PRIV]]
223*67e74705SXin Li // CHECK: call {{.*}} [[S_FLOAT_TY_DEF_CONSTR]]([[S_FLOAT_TY]]* [[VAR_PRIV]])
224*67e74705SXin Li // CHECK-DAG: call void [[S_FLOAT_TY_DESTR]]([[S_FLOAT_TY]]* [[VAR_PRIV]])
225*67e74705SXin Li // CHECK-DAG: call void [[S_FLOAT_TY_DESTR]]([[S_FLOAT_TY]]*
226*67e74705SXin Li // CHECK: ret void
227*67e74705SXin Li 
228*67e74705SXin Li // template tmain
229*67e74705SXin Li // CHECK: define{{.*}} i{{[0-9]+}} [[TMAIN_INT]]()
230*67e74705SXin Li // CHECK: [[TEST:%.+]] = alloca [[S_INT_TY]],
231*67e74705SXin Li // CHECK: call {{.*}} [[S_INT_TY_DEF_CONSTR:@.+]]([[S_INT_TY]]* [[TEST]])
232*67e74705SXin Li // CHECK: call void [[S_INT_TY_CONSTR:@.+]]([[S_INT_TY]]* {{.+}}, i{{[0-9]+}}{{.*}} 3)
233*67e74705SXin Li // CHECK: call void [[OMP_OFFLOADING_TMAIN:@.+]]()
234*67e74705SXin Li 
235*67e74705SXin Li // target in SS constructor
236*67e74705SXin Li // CHECK: define{{.+}} [[OMP_OFFLOADING_SS:@.+]]([[SS_TY]]*
237*67e74705SXin Li // CHECK: call void (%{{.+}}*, i{{[0-9]+}}, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)*, ...) @__kmpc_fork_teams(%{{.+}}* @{{.+}}, i{{[0-9]+}} 1, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)* bitcast (void (i{{[0-9]+}}*, i{{[0-9]+}}*, [[SS_TY]]*)* [[OMP_OUTLINED_SS:@.+]] to void
238*67e74705SXin Li 
239*67e74705SXin Li // CHECK: define{{.*}} void [[OMP_OUTLINED_SS]](i{{[0-9]+}}* {{.+}}, i{{[0-9]+}}* {{.+}}, [[SS_TY]]*
240*67e74705SXin Li // CHECK: [[A_PRIV:%.+]] = alloca i{{[0-9]+}},
241*67e74705SXin Li // CHECK: [[B_PRIV:%.+]] = alloca i{{[0-9]+}},
242*67e74705SXin Li // CHECK: [[C_PRIV:%.+]] = alloca i{{[0-9]+}},
243*67e74705SXin Li // CHECK: store i{{[0-9]+}}* [[A_PRIV]], i{{[0-9]+}}** [[A_REF:%.+]],
244*67e74705SXin Li // CHECK: store i{{[0-9]+}}* [[C_PRIV]], i{{[0-9]+}}** [[C_REF:%.+]],
245*67e74705SXin Li // CHECK: [[A_REF_VAL:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[A_REF]]
246*67e74705SXin Li // CHECK: [[A_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_REF_VAL]]
247*67e74705SXin Li // CHECK: [[A_INC:%.+]] = add{{.*}} i{{[0-9]+}} [[A_VAL]], 1
248*67e74705SXin Li // CHECK: store i{{[0-9]+}} [[A_INC]], i{{[0-9]+}}* [[A_REF_VAL]],
249*67e74705SXin Li // CHECK: [[B_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[B_PRIV]]
250*67e74705SXin Li // CHECK: [[B_DEC:%.+]] = add{{.*}} i{{[0-9]+}} [[B_VAL]], -1
251*67e74705SXin Li // CHECK: store i{{[0-9]+}} [[B_DEC]], i{{[0-9]+}}* [[B_PRIV]],
252*67e74705SXin Li // CHECK: [[C_REF_VAL:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[C_REF]]
253*67e74705SXin Li // CHECK: [[C_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[C_REF_VAL]]
254*67e74705SXin Li // CHECK: [[C_DIV:%.+]] = sdiv i{{[0-9]+}} [[C_VAL]], 1
255*67e74705SXin Li // CHECK: store i{{[0-9]+}} [[C_DIV]], i{{[0-9]+}}* [[C_REF_VAL]],
256*67e74705SXin Li // CHECK: ret
257*67e74705SXin Li 
258*67e74705SXin Li // target in tmain template
259*67e74705SXin Li // CHECK: define{{.+}} [[OMP_OFFLOADING_TMAIN]]()
260*67e74705SXin Li // CHECK: call void (%{{.+}}*, i{{[0-9]+}}, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)*, ...) @__kmpc_fork_teams(%{{.+}}* @{{.+}}, i{{[0-9]+}} 0, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)* bitcast (void (i{{[0-9]+}}*, i{{[0-9]+}}*)* [[OMP_OUTLINED_TMAIN:@.+]] to void
261*67e74705SXin Li 
262*67e74705SXin Li // CHECK: define{{.*}} void [[OMP_OUTLINED_TMAIN]](i{{[0-9]+}}* noalias [[GTID_ADDR:%.+]], i{{[0-9]+}}* noalias %{{.+}})
263*67e74705SXin Li // CHECK: [[T_VAR_PRIV:%.+]] = alloca i{{[0-9]+}}, align 128
264*67e74705SXin Li // CHECK: [[VEC_PRIV:%.+]] = alloca [2 x i{{[0-9]+}}], align 128
265*67e74705SXin Li // CHECK: [[S_ARR_PRIV:%.+]] = alloca [2 x [[S_INT_TY]]], align 128
266*67e74705SXin Li // CHECK: [[VAR_PRIV:%.+]] = alloca [[S_INT_TY]], align 128
267*67e74705SXin Li // CHECK: store i{{[0-9]+}}* [[GTID_ADDR]], i{{[0-9]+}}** [[GTID_ADDR_REF:%.+]]
268*67e74705SXin Li // CHECK-NOT: [[T_VAR_PRIV]]
269*67e74705SXin Li // CHECK-NOT: [[VEC_PRIV]]
270*67e74705SXin Li // CHECK: {{.+}}:
271*67e74705SXin Li // CHECK: [[S_ARR_PRIV_ITEM:%.+]] = phi [[S_INT_TY]]*
272*67e74705SXin Li // CHECK: call {{.*}} [[S_INT_TY_DEF_CONSTR]]([[S_INT_TY]]* [[S_ARR_PRIV_ITEM]])
273*67e74705SXin Li // CHECK-NOT: [[T_VAR_PRIV]]
274*67e74705SXin Li // CHECK-NOT: [[VEC_PRIV]]
275*67e74705SXin Li // CHECK: call {{.*}} [[S_INT_TY_DEF_CONSTR]]([[S_INT_TY]]* [[VAR_PRIV]])
276*67e74705SXin Li // CHECK-DAG: call void [[S_INT_TY_DESTR:@.+]]([[S_INT_TY]]* [[VAR_PRIV]])
277*67e74705SXin Li // CHECK-DAG: call void [[S_INT_TY_DESTR]]([[S_INT_TY]]*
278*67e74705SXin Li // CHECK: ret
279*67e74705SXin Li 
280*67e74705SXin Li // SST constructor
281*67e74705SXin Li // CHECK: define{{.+}} [[SST_CONST:@.+]]([[SST_TY]]* {{.+}})
282*67e74705SXin Li // CHECK: call void [[OMP_OFFLOADING_SST:@.+]]([[SST_TY]]* {{.+}})
283*67e74705SXin Li 
284*67e74705SXin Li // target in SST constructor
285*67e74705SXin Li // CHECK: define{{.+}} [[OMP_OFFLOADING_SST]]([[SST_TY]]* {{.+}})
286*67e74705SXin Li // CHECK: call void (%{{.+}}*, i{{[0-9]+}}, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)*, ...) @__kmpc_fork_teams(%{{.+}}* @{{.+}}, i{{[0-9]+}} 1, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)* bitcast (void (i{{[0-9]+}}*, i{{[0-9]+}}*, [[SST_TY]]*)* [[OMP_OUTLINED_SST:@.+]] to void
287*67e74705SXin Li 
288*67e74705SXin Li // CHECK: define{{.+}} [[OMP_OUTLINED_SST]](i{{[0-9]+}}* {{.+}}, i{{[0-9]+}}* noalias %{{.+}}, [[SST_TY]]* {{.+}})
289*67e74705SXin Li // CHECK: [[A_PRIV_1:%.+]] = alloca i{{[0-9]+}},
290*67e74705SXin Li // CHECK: store i{{[0-9]+}}* [[A_PRIV_1]], i{{[0-9]+}}** [[A_REF_1:%.+]],
291*67e74705SXin Li // CHECK: [[A_REF_VAL_1:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[A_REF_1]]
292*67e74705SXin Li // CHECK: [[A_VAL_1:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_REF_VAL_1]]
293*67e74705SXin Li // CHECK: [[A_INC_1:%.+]] = add{{.*}} i{{[0-9]+}} [[A_VAL_1]], 1
294*67e74705SXin Li // CHECK: store i{{[0-9]+}} [[A_INC_1]], i{{[0-9]+}}* [[A_REF_VAL_1]],
295*67e74705SXin Li // CHECK: ret
296*67e74705SXin Li 
297*67e74705SXin Li #endif
298*67e74705SXin Li 
299