1 // expected-no-diagnostics
2 #ifndef HEADER
3 #define HEADER
4 // Test host codegen.
5 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64
6 // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
7 // RUN: %clang_cc1 -DCK1 -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 CK1 --check-prefix CK1-64
8 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32
9 // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
10 // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -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 CK1 --check-prefix CK1-32
11 #ifdef CK1
12
13 int Gbla;
14 long long Gblb;
15 int &Gblc = Gbla;
16
17 // CK1-LABEL: teams_argument_global_local
teams_argument_global_local(int a)18 int teams_argument_global_local(int a){
19 int comp = 1;
20
21 int la = 23;
22 float lc = 25.0;
23
24 // CK1: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0)
25 // CK1: call void @{{.+}}(i{{64|32}} %{{.+}})
26 #pragma omp target
27 #pragma omp teams
28 {
29 ++comp;
30 }
31
32 // CK1: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0)
33 // CK1: call void @{{.+}}(i{{64|32}} %{{.+}})
34 #pragma omp target
35 {{{
36 #pragma omp teams
37 {
38 ++comp;
39 }
40 }}}
41
42 // CK1-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 [[NT:%[^,]+]], i32 0)
43 // CK1-DAG: [[NT]] = load i32, i32* [[NTA:%[^,]+]],
44
45 // CK1: call void @{{.+}}(i{{64|32}} %{{.+}})
46 #pragma omp target
47 #pragma omp teams num_teams(la)
48 {
49 ++comp;
50 }
51
52 // CK1-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 [[NT:%[^,]+]])
53 // CK1-DAG: [[NT]] = load i32, i32* [[NTA:%[^,]+]],
54
55 // CK1: call void @{{.+}}(i{{64|32}} %{{.+}})
56 #pragma omp target
57 #pragma omp teams thread_limit(la)
58 {
59 ++comp;
60 }
61
62 // CK1-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 5, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 [[NT:%[^,]+]], i32 [[TL:%[^,]+]])
63
64 // CK1-DAG: [[NT]] = add nsw i32 [[NTA:%[^,]+]], [[NTB:%[^,]+]]
65 // CK1-DAG: [[NTA]] = load i32, i32* @Gbla,
66 // CK1-DAG: [[NTB]] = load i32, i32* %{{.+}},
67
68 // CK1-DAG: [[TL]] = trunc i64 [[TLA:%[^,]+]] to i32
69 // CK1-DAG: [[TLA]] = add nsw i64 [[TLB:%[^,]+]], [[TLC:%[^,]+]]
70 // CK1-DAG: [[TLC]] = fptosi float [[TLD:%[^,]+]] to i64
71 // CK1-DAG: [[TLD]] = load float, float* %{{.+}},
72 // CK1-DAG: [[TLB]] = load i64, i64* @Gblb,
73
74 // CK1: call void @{{.+}}(i{{.+}} {{.+}}, i{{.+}} {{.+}}, i{{.+}} {{.+}}, i{{.+}} {{.+}}, i{{.+}} {{.+}})
75 #pragma omp target
76 #pragma omp teams num_teams(Gbla+a) thread_limit(Gblb+(long long)lc)
77 {
78 ++comp;
79 }
80
81 // CK1-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 {{.+}}, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 [[NT:%[^,]+]], i32 [[TL:%[^,]+]])
82
83 // CK1-DAG: [[NT]] = add nsw i32 [[NTA:%[^,]+]], 1
84 // CK1-DAG: [[NTA]] = load i32, i32* @Gbla,
85
86 // CK1-DAG: [[TL]] = add nsw i32 [[TLA:%[^,]+]], 2
87 // CK1-DAG: [[TLA]] = load i32, i32* @Gbla,
88
89 // CK1: call void @{{.+}}(i{{.+}} {{.+}}
90 #pragma omp target
91 #pragma omp teams num_teams(Gblc+1) thread_limit(Gblc+2)
92 {
93 comp += Gblc;
94 }
95
96 return comp;
97 }
98
99 #endif // CK1
100
101 // Test host codegen.
102 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64
103 // RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
104 // RUN: %clang_cc1 -DCK2 -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 CK2 --check-prefix CK2-64
105 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32
106 // RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
107 // RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -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 CK2 --check-prefix CK2-32
108 #ifdef CK2
109
110 // CK2-DAG: [[SSI:%.+]] = type { i32, float }
111 // CK2-DAG: [[SSL:%.+]] = type { i64, float }
112 template <typename T>
113 struct SS{
114 T a;
115 float b;
116 };
117
118 SS<int> Gbla;
119 SS<long long> Gblb;
120
121 // CK2-LABEL: teams_template_arg
teams_template_arg(void)122 int teams_template_arg(void) {
123 int comp = 1;
124
125 SS<int> la;
126 SS<long long> lb;
127
128 // CK2-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 [[NT:%[^,]+]], i32 [[TL:%[^,]+]])
129
130 // CK2-DAG: [[NT]] = load i32, i32* getelementptr inbounds ([[SSI]], [[SSI]]* @Gbla, i32 0, i32 0)
131
132 // CK2-DAG: [[TL]] = trunc i64 [[TLA:%[^,]+]] to i32
133 // CK2-DAG: [[TLA]] = fptosi float [[TLB:%[^,]+]] to i64
134 // CK2-DAG: [[TLB]] = load float, float* [[TLC:%[^,]+]],
135 // CK2-DAG: [[TLC]] = getelementptr inbounds [[SSI]], [[SSI]]* %{{.+}}, i32 0, i32 1
136
137 // CK2: call void @{{.+}}({{.+}} {{.+}}, {{.+}} {{.+}}, {{.+}} {{.+}})
138 #pragma omp target
139 #pragma omp teams num_teams(Gbla.a) thread_limit((long long)la.b)
140 {
141 ++comp;
142 }
143
144 // CK2-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 [[NT:%[^,]+]], i32 [[TL:%[^,]+]])
145
146 // CK2-DAG: [[TL]] = trunc i64 [[TLD:%[^,]+]] to i32
147 // CK2-DAG: [[TLD]] = load i64, i64* getelementptr inbounds ([[SSL]], [[SSL]]* @Gblb, i32 0, i32 0),
148
149 // CK2-DAG: [[NT]] = trunc i64 [[NTA:%[^,]+]] to i32
150 // CK2-DAG: [[NTA]] = fptosi float [[NTB:%[^,]+]] to i64
151 // CK2-DAG: [[NTB]] = load float, float* [[NTC:%[^,]+]],
152 // CK2-DAG: [[NTC]] = getelementptr inbounds [[SSL]], [[SSL]]* %{{.+}}, i32 0, i32 1
153
154 // CK2: call void @{{.+}}({{.+}} {{.+}}, {{.+}} {{.+}}, {{.+}} {{.+}})
155 #pragma omp target
156 #pragma omp teams num_teams((long long)lb.b) thread_limit(Gblb.a)
157 {
158 ++comp;
159 }
160 return comp;
161 }
162 #endif // CK2
163
164 // Test host codegen.
165 // RUN: %clang_cc1 -DCK3 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-64
166 // RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
167 // RUN: %clang_cc1 -DCK3 -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 CK3 --check-prefix CK3-64
168 // RUN: %clang_cc1 -DCK3 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-32
169 // RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
170 // RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -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 CK3 --check-prefix CK3-32
171 #ifdef CK3
172
173 // CK3: [[SSI:%.+]] = type { i32, float }
174 // CK3-LABEL: teams_template_struct
175
176 template <typename T, int X, long long Y>
177 struct SS{
178 T a;
179 float b;
180
fooSS181 int foo(void) {
182 int comp = 1;
183
184 // CK3-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 [[NT:%[^,]+]], i32 123)
185
186 // CK3-DAG: [[NT]] = load i32, i32* [[NTA:%[^,]+]],
187 // CK3-DAG: [[NTA]] = getelementptr inbounds [[SSI]], [[SSI]]* [[NTB:%[^,]+]], i32 0, i32 0
188 // CK3-DAG: [[NTB]] = load [[SSI]]*, [[SSI]]** %{{.+}},
189
190 // CK3: call void @{{.+}}({{.+}} {{.+}}, {{.+}} {{.+}})
191 #pragma omp target
192 #pragma omp teams num_teams(a) thread_limit(X)
193 {
194 ++comp;
195 }
196
197 // CK3-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 456, i32 [[TL:%[^,]+]])
198
199 // CK3-DAG: [[TL]] = add nsw i32 [[TLA:%[^,]+]], 123
200 // CK3-DAG: [[TLA]] = fptosi float [[TLB:%[^,]+]] to i32
201 // CK3-DAG: [[TLB]] = load float, float* [[TLC:%[^,]+]],
202 // CK3-DAG: [[TLC]] = getelementptr inbounds [[SSI]], [[SSI]]* [[THIS:%[^,]+]], i32 0, i32 1
203
204 // CK3: call void @{{.+}}({{.+}} {{.+}}, {{.+}} {{.+}})
205 #pragma omp target
206 #pragma omp teams num_teams(Y) thread_limit((int)b+X)
207 {
208 ++comp;
209 }
210 return comp;
211 }
212 };
213
teams_template_struct(void)214 int teams_template_struct(void) {
215 SS<int, 123, 456> V;
216 return V.foo();
217
218 }
219 #endif // CK3
220
221 // Test target codegen - host bc file has to be created first.
222 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
223 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-64
224 // RUN: %clang_cc1 -DCK4 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
225 // RUN: %clang_cc1 -DCK4 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-64
226 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
227 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-32
228 // RUN: %clang_cc1 -DCK4 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
229 // RUN: %clang_cc1 -DCK4 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-32
230
231 #ifdef CK4
232
233 // CK4-DAG: %ident_t = type { i32, i32, i32, i32, i8* }
234 // CK4-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
235 // CK4-DAG: [[DEF_LOC_0:@.+]] = private unnamed_addr constant %ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
236 // CK4-DEBUG-DAG: [[LOC1:@.+]] = private unnamed_addr constant [{{.+}} x i8] c";{{.*}}teams_codegen.cpp;main;[[@LINE+14]];9;;\00"
237 // CK4-DEBUG-DAG: [[LOC2:@.+]] = private unnamed_addr constant [{{.+}} x i8] c";{{.*}}teams_codegen.cpp;tmain;[[@LINE+7]];9;;\00"
238
239 template <typename T>
tmain(T argc)240 int tmain(T argc) {
241 #pragma omp target
242 #pragma omp teams
243 argc = 0;
244 return 0;
245 }
246
main(int argc,char ** argv)247 int main (int argc, char **argv) {
248 #pragma omp target
249 #pragma omp teams
250 argc = 0;
251 return tmain(argv);
252 }
253
254 // CK4: define {{.*}}void @{{[^,]+}}(i{{.+}} %[[ARGC:.+]])
255 // CK4: [[ARGCADDR:%.+]] = alloca i{{.+}}
256 // CK4: store i{{.+}} %[[ARGC]], i{{.+}}* [[ARGCADDR]]
257 // CK4-64: [[CONV:%.+]] = bitcast i64* [[ARGCADDR]] to i32*
258 // CK4-64: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* {{.+}} to void (i32*, i32*, ...)*), i32* [[CONV]])
259 // CK4-32: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* {{.+}} to void (i32*, i32*, ...)*), i32* [[ARGCADDR]])
260 // CK4: ret void
261 // CK4-NEXT: }
262
263 // CK4: define {{.*}}void @{{[^,]+}}(i8** [[ARGC1:%.+]])
264 // CK4: [[ARGCADDR1:%.+]] = alloca i8**
265 // CK4: store i8** [[ARGC1]], i8*** [[ARGCADDR1]]
266 // CK4: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i8***)* {{.+}} to void (i32*, i32*, ...)*), i8*** [[ARGCADDR1]])
267
268
269 #endif // CK4
270
271 // Test target codegen - host bc file has to be created first.
272 // RUN: %clang_cc1 -DCK5 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
273 // RUN: %clang_cc1 -DCK5 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CK5 --check-prefix CK5-64
274 // RUN: %clang_cc1 -DCK5 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
275 // RUN: %clang_cc1 -DCK5 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK5 --check-prefix CK5-64
276 // RUN: %clang_cc1 -DCK5 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
277 // RUN: %clang_cc1 -DCK5 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK5 --check-prefix CK5-32
278 // RUN: %clang_cc1 -DCK5 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
279 // RUN: %clang_cc1 -DCK5 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK5 --check-prefix CK5-32
280
281 // expected-no-diagnostics
282 #ifdef CK5
283
284 // CK5-DAG: %ident_t = type { i32, i32, i32, i32, i8* }
285 // CK5-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
286 // CK5-DAG: [[DEF_LOC_0:@.+]] = private unnamed_addr constant %ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
287 // CK5-DEBUG-DAG: [[LOC1:@.+]] = private unnamed_addr constant [{{.+}} x i8] c";{{.*}}teams_codegen.cpp;main;[[@LINE+14]];9;;\00"
288 // CK5-DEBUG-DAG: [[LOC2:@.+]] = private unnamed_addr constant [{{.+}} x i8] c";{{.*}}teams_codegen.cpp;tmain;[[@LINE+7]];9;;\00"
289
290 template <typename T>
tmain(T argc)291 int tmain(T argc) {
292 int a = 10;
293 int b = 5;
294 #pragma omp target
295 #pragma omp teams num_teams(a) thread_limit(b)
296 {
297 argc = 0;
298 }
299 return 0;
300 }
301
main(int argc,char ** argv)302 int main (int argc, char **argv) {
303 int a = 20;
304 int b = 5;
305 #pragma omp target
306 #pragma omp teams num_teams(a) thread_limit(b)
307 {
308 argc = 0;
309 }
310 return tmain(argv);
311 }
312
313 // CK5: define {{.*}}void @{{[^,]+}}(i{{.+}} [[AP:%.+]], i{{.+}} [[BP:%.+]], i{{.+}} [[ARGC:.+]])
314 // CK5: [[AADDR:%.+]] = alloca i{{.+}}
315 // CK5: [[BADDR:%.+]] = alloca i{{.+}}
316 // CK5: [[ARGCADDR:%.+]] = alloca i{{.+}}
317 // CK5: [[GBL_TH_NUM:%.+]] = call i32 @__kmpc_global_thread_num(%ident_t* [[DEF_LOC_0]])
318 // CK5: store i{{.+}} [[AP]], i{{.+}}* [[AADDR]]
319 // CK5: store i{{.+}} [[BP]], i{{.+}}* [[BADDR]]
320 // CK5: store i{{.+}} [[ARGC]], i{{.+}}* [[ARGCADDR]]
321 // CK5-64: [[ACONV:%.+]] = bitcast i64* [[AADDR]] to i32*
322 // CK5-64: [[BCONV:%.+]] = bitcast i64* [[BADDR]] to i32*
323 // CK5-64: [[CONV:%.+]] = bitcast i64* [[ARGCADDR]] to i32*
324 // CK5-64: [[ACONVVAL:%.+]] = load i32, i32* [[ACONV]]
325 // CK5-64: [[BCONVVAL:%.+]] = load i32, i32* [[BCONV]]
326 // CK5-32: [[ACONVVAL:%.+]] = load i32, i32* [[AADDR]]
327 // CK5-32: [[BCONVVAL:%.+]] = load i32, i32* [[BADDR]]
328 // CK5: {{.+}} = call i32 @__kmpc_push_num_teams(%ident_t* [[DEF_LOC_0]], i32 [[GBL_TH_NUM]], i32 [[ACONVVAL]], i32 [[BCONVVAL]])
329 // CK5-64: call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[CONV]])
330 // CK5-32: call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[ARGCADDR]])
331
332 // CK5: define {{.*}}void @{{[^,]+}}(i{{.+}} [[AP:%.+]], i{{.+}} [[BP:%.+]], i{{.+}}** [[ARGC:%.+]])
333 // CK5: [[AADDR:%.+]] = alloca i{{.+}}
334 // CK5: [[BADDR:%.+]] = alloca i{{.+}}
335 // CK5: [[ARGCADDR:%.+]] = alloca i{{.+}}**
336 // CK5: [[GBL_TH_NUM:%.+]] = call i32 @__kmpc_global_thread_num(%ident_t* [[DEF_LOC_0]])
337 // CK5: store i{{.+}} [[AP]], i{{.+}}* [[AADDR]]
338 // CK5: store i{{.+}} [[BP]], i{{.+}}* [[BADDR]]
339 // CK5: store i{{.+}}** [[ARGC]], i{{.+}}*** [[ARGCADDR]]
340 // CK5-64: [[ACONV:%.+]] = bitcast i64* [[AADDR]] to i32*
341 // CK5-64: [[BCONV:%.+]] = bitcast i64* [[BADDR]] to i32*
342 // CK5-64: [[ACONVVAL:%.+]] = load i32, i32* [[ACONV]]
343 // CK5-64: [[BCONVVAL:%.+]] = load i32, i32* [[BCONV]]
344 // CK5-64: {{.+}} = call i32 @__kmpc_push_num_teams(%ident_t* [[DEF_LOC_0]], i32 [[GBL_TH_NUM]], i32 [[ACONVVAL]], i32 [[BCONVVAL]])
345 // CK5-32: [[A_VAL:%.+]] = load i32, i32* [[AADDR]]
346 // CK5-32: [[B_VAL:%.+]] = load i32, i32* [[BADDR]]
347 // CK5-32: {{.+}} = call i32 @__kmpc_push_num_teams(%ident_t* [[DEF_LOC_0]], i32 [[GBL_TH_NUM]], i32 [[A_VAL]], i32 [[B_VAL]])
348 // CK5: call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i{{.+}})* @.omp_outlined.{{.+}} to void (i32*, i32*, ...)*), i{{.+}}*** [[ARGCADDR]])
349 // CK5: ret void
350 // CK5-NEXT: }
351
352 #endif // CK5
353 #endif
354