1 // Test host codegen.
2 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
3 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
4 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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
5 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
6 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
7 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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 CHECK --check-prefix CHECK-32
8 
9 // Test target codegen - host bc file has to be created first.
10 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
11 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -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 TCHECK --check-prefix TCHECK-64
12 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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
13 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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 TCHECK --check-prefix TCHECK-64
14 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
15 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -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 TCHECK --check-prefix TCHECK-32
16 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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
17 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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 TCHECK --check-prefix TCHECK-32
18 
19 // expected-no-diagnostics
20 #ifndef HEADER
21 #define HEADER
22 
23 // CHECK-DAG: %ident_t = type { i32, i32, i32, i32, i8* }
24 // CHECK-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
25 // CHECK-DAG: [[DEF_LOC:@.+]] = 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) }
26 
27 // CHECK-DAG: [[S1:%.+]] = type { double }
28 // CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]], i32, i32 }
29 // CHECK-DAG: [[DEVTY:%.+]] = type { i8*, i8*, [[ENTTY]]*, [[ENTTY]]* }
30 // CHECK-DAG: [[DSCTY:%.+]] = type { i32, [[DEVTY]]*, [[ENTTY]]*, [[ENTTY]]* }
31 
32 // TCHECK: [[ENTTY:%.+]] = type { i8*, i8*, i{{32|64}}, i32, i32 }
33 
34 // CHECK-DAG: $[[REGFN:\.omp_offloading\..+]] = comdat
35 
36 // We have 6 target regions
37 
38 // CHECK-DAG: @{{.*}} = private constant i8 0
39 // CHECK-DAG: @{{.*}} = private constant i8 0
40 // CHECK-DAG: @{{.*}} = private constant i8 0
41 // CHECK-DAG: @{{.*}} = private constant i8 0
42 // CHECK-DAG: @{{.*}} = private constant i8 0
43 // CHECK-DAG: @{{.*}} = private constant i8 0
44 
45 // TCHECK: @{{.+}} = constant [[ENTTY]]
46 // TCHECK: @{{.+}} = constant [[ENTTY]]
47 // TCHECK: @{{.+}} = constant [[ENTTY]]
48 // TCHECK: @{{.+}} = constant [[ENTTY]]
49 // TCHECK: @{{.+}} = constant [[ENTTY]]
50 // TCHECK: @{{.+}} = constant [[ENTTY]]
51 
52 // Check if offloading descriptor is created.
53 // CHECK: [[ENTBEGIN:@.+]] = external constant [[ENTTY]]
54 // CHECK: [[ENTEND:@.+]] = external constant [[ENTTY]]
55 // CHECK: [[DEVBEGIN:@.+]] = external constant i8
56 // CHECK: [[DEVEND:@.+]] = external constant i8
57 // CHECK: [[IMAGES:@.+]] = internal unnamed_addr constant [1 x [[DEVTY]]] [{{.+}} { i8* [[DEVBEGIN]], i8* [[DEVEND]], [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }], comdat($[[REGFN]])
58 // CHECK: [[DESC:@.+]] = internal constant [[DSCTY]] { i32 1, [[DEVTY]]* getelementptr inbounds ([1 x [[DEVTY]]], [1 x [[DEVTY]]]* [[IMAGES]], i32 0, i32 0), [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }, comdat($[[REGFN]])
59 
60 // Check target registration is registered as a Ctor.
61 // CHECK: appending global [1 x { i32, void ()*, i8* }] [{ i32, void ()*, i8* } { i32 0, void ()* bitcast (void (i8*)* @[[REGFN]] to void ()*), i8* bitcast (void (i8*)* @[[REGFN]] to i8*) }]
62 
63 
64 template<typename tx>
65 tx ftemplate(int n) {
66   tx a = 0;
67 
68   #pragma omp target teams thread_limit(tx(20))
69   {
70   }
71 
72   short b = 1;
73   #pragma omp target teams num_teams(b) thread_limit(1024)
74   {
75     a += b;
76   }
77 
78   return a;
79 }
80 
81 static
82 int fstatic(int n) {
83 
84   #pragma omp target teams num_teams(n) thread_limit(n*32)
85   {
86   }
87 
88   #pragma omp target teams thread_limit(32+n)
89   {
90   }
91 
92   return n+1;
93 }
94 
95 struct S1 {
96   double a;
97 
98   int r1(int n){
99     int b = 1;
100 
101     #pragma omp target teams thread_limit(n-b)
102     {
103       this->a = (double)b + 1.5;
104     }
105 
106     #pragma omp target teams thread_limit(1024)
107     {
108       this->a = 2.5;
109     }
110 
111     return (int)a;
112   }
113 };
114 
115 // CHECK: define {{.*}}@{{.*}}bar{{.*}}
116 int bar(int n){
117   int a = 0;
118 
119   S1 S;
120   // CHECK: call {{.*}}i32 [[FS1:@.+]]([[S1]]* {{.*}}, i32 {{.*}})
121   a += S.r1(n);
122 
123   // CHECK: call {{.*}}i32 [[FSTATIC:@.+]](i32 {{.*}})
124   a += fstatic(n);
125 
126   // CHECK: call {{.*}}i32 [[FTEMPLATE:@.+]](i32 {{.*}})
127   a += ftemplate<int>(n);
128 
129   return a;
130 }
131 
132 
133 
134 //
135 // CHECK: define {{.*}}[[FS1]]([[S1]]* {{%.+}}, i32 {{[^%]*}}[[PARM:%.+]])
136 //
137 // CHECK-DAG:   store i32 [[PARM]], i32* [[N_ADDR:%.+]], align
138 // CHECK:       store i32 1, i32* [[B:%.+]], align
139 // CHECK:       [[NV:%.+]] = load i32, i32* [[N_ADDR]], align
140 // CHECK:       [[BV:%.+]] = load i32, i32* [[B]], align
141 // CHECK:       [[SUB:%.+]] = sub nsw i32 [[NV]], [[BV]]
142 // CHECK:       store i32 [[SUB]], i32* [[CAPE_ADDR:%.+]], align
143 // CHECK:       [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR]], align
144 // CHECK-64:    [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i32*
145 // CHECK-64:    store i32 [[CEV]], i32* [[CONV]], align
146 // CHECK-32:    store i32 [[CEV]], i32* [[CAPEC_ADDR:%.+]], align
147 // CHECK:       [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align
148 // CHECK:       [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align
149 //
150 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i32 0, i32 [[TL]])
151 // CHECK:       store i32 [[RET]], i32* [[RHV:%.+]], align
152 // CHECK:       [[RET2:%.+]] = load i32, i32* [[RHV]], align
153 // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0
154 // CHECK:       br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
155 //
156 // CHECK:       [[FAIL]]
157 // CHECK:       call void [[HVT1:@.+]]([[S1]]* {{%.+}}, i[[SZ]] {{%.+}}, i[[SZ]] [[ARG]])
158 // CHECK:       br label {{%?}}[[END]]
159 // CHECK:       [[END]]
160 //
161 //
162 //
163 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, {{.+}}, i32 0, i32 1024)
164 // CHECK:       store i32 [[RET]], i32* [[RHV:%.+]], align
165 // CHECK:       [[RET2:%.+]] = load i32, i32* [[RHV]], align
166 // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0
167 // CHECK:       br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
168 //
169 // CHECK:       [[FAIL]]
170 // CHECK:       call void [[HVT2:@.+]]([[S1]]* {{[^,]+}})
171 // CHECK:       br label {{%?}}[[END]]
172 // CHECK:       [[END]]
173 //
174 
175 
176 
177 
178 
179 
180 //
181 // CHECK: define {{.*}}[[FSTATIC]](i32 {{[^%]*}}[[PARM:%.+]])
182 //
183 // CHECK-DAG:   store i32 [[PARM]], i32* [[N_ADDR:%.+]], align
184 // CHECK:       [[NV:%.+]] = load i32, i32* [[N_ADDR]], align
185 // CHECK:       store i32 [[NV]], i32* [[CAPE_ADDR1:%.+]], align
186 // CHECK:       [[NV:%.+]] = load i32, i32* [[N_ADDR]], align
187 // CHECK:       [[MUL:%.+]] = mul nsw i32 [[NV]], 32
188 // CHECK:       store i32 [[MUL]], i32* [[CAPE_ADDR2:%.+]], align
189 // CHECK:       [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR1]], align
190 // CHECK-64:    [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR1:%.+]] to i32*
191 // CHECK-64:    store i32 [[CEV]], i32* [[CONV]], align
192 // CHECK-32:    store i32 [[CEV]], i32* [[CAPEC_ADDR1:%.+]], align
193 // CHECK:       [[ARG1:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR1]], align
194 // CHECK:       [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR2]], align
195 // CHECK-64:    [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR2:%.+]] to i32*
196 // CHECK-64:    store i32 [[CEV]], i32* [[CONV]], align
197 // CHECK-32:    store i32 [[CEV]], i32* [[CAPEC_ADDR2:%.+]], align
198 // CHECK:       [[ARG2:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR2]], align
199 // CHECK:       [[TEAMS:%.+]] = load i32, i32* [[CAPE_ADDR1]], align
200 // CHECK:       [[TL:%.+]] = load i32, i32* [[CAPE_ADDR2]], align
201 //
202 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 2, {{.*}}, i32 [[TEAMS]], i32 [[TL]])
203 // CHECK:       store i32 [[RET]], i32* [[RHV:%.+]], align
204 // CHECK:       [[RET2:%.+]] = load i32, i32* [[RHV]], align
205 // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0
206 // CHECK:       br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
207 //
208 // CHECK:       [[FAIL]]
209 // CHECK:       call void [[HVT3:@.+]](i[[SZ]] [[ARG1]], i[[SZ]] [[ARG2]])
210 // CHECK:       br label {{%?}}[[END]]
211 // CHECK:       [[END]]
212 //
213 //
214 //
215 // CHECK:       [[NV:%.+]] = load i32, i32* [[N_ADDR]], align
216 // CHECK:       [[ADD:%.+]] = add nsw i32 32, [[NV]]
217 // CHECK:       store i32 [[ADD]], i32* [[CAPE_ADDR:%.+]], align
218 // CHECK:       [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR]], align
219 // CHECK-64:    [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i32*
220 // CHECK-64:    store i32 [[CEV]], i32* [[CONV]], align
221 // CHECK-32:    store i32 [[CEV]], i32* [[CAPEC_ADDR:%.+]], align
222 // CHECK:       [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align
223 // CHECK:       [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align
224 //
225 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i32 0, i32 [[TL]])
226 // CHECK:       store i32 [[RET]], i32* [[RHV:%.+]], align
227 // CHECK:       [[RET2:%.+]] = load i32, i32* [[RHV]], align
228 // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0
229 // CHECK:       br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
230 //
231 // CHECK:       [[FAIL]]
232 // CHECK:       call void [[HVT4:@.+]](i[[SZ]] [[ARG]])
233 // CHECK:       br label {{%?}}[[END]]
234 // CHECK:       [[END]]
235 //
236 
237 
238 
239 
240 
241 
242 //
243 // CHECK: define {{.*}}[[FTEMPLATE]]
244 //
245 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 0, {{.*}}, i32 0, i32 20)
246 // CHECK-NEXT:  store i32 [[RET]], i32* [[RHV:%.+]], align
247 // CHECK-NEXT:  [[RET2:%.+]] = load i32, i32* [[RHV]], align
248 // CHECK-NEXT:  [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0
249 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
250 //
251 // CHECK:       [[FAIL]]
252 // CHECK:       call void [[HVT5:@.+]]()
253 // CHECK:       br label {{%?}}[[END]]
254 //
255 // CHECK:       [[END]]
256 //
257 //
258 //
259 // CHECK:       store i16 1, i16* [[B:%.+]], align
260 // CHECK:       [[BV:%.+]] = load i16, i16* [[B]], align
261 // CHECK:       store i16 [[BV]], i16* [[CAPE_ADDR:%.+]], align
262 // CHECK:       [[CEV:%.+]] = load i16, i16* [[CAPE_ADDR]], align
263 // CHECK:       [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i16*
264 // CHECK:       store i16 [[CEV]], i16* [[CONV]], align
265 // CHECK:       [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align
266 // CHECK:       [[T:%.+]] = load i16, i16* [[CAPE_ADDR]], align
267 // CHECK:       [[TEAMS:%.+]] = sext i16 [[T]] to i32
268 //
269 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i32 [[TEAMS]], i32 1024)
270 // CHECK:       store i32 [[RET]], i32* [[RHV:%.+]], align
271 // CHECK:       [[RET2:%.+]] = load i32, i32* [[RHV]], align
272 // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0
273 // CHECK:       br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
274 //
275 // CHECK:       [[FAIL]]
276 // CHECK:       call void [[HVT6:@.+]](i[[SZ]] {{%.+}}, i[[SZ]] {{%.+}}, i[[SZ]] [[ARG]])
277 // CHECK:       br label {{%?}}[[END]]
278 // CHECK:       [[END]]
279 //
280 
281 
282 
283 
284 
285 
286 // Check that the offloading functions are emitted and that the parallel function
287 // is appropriately guarded.
288 
289 // CHECK:       define internal void [[HVT1]]([[S1]]* {{%.+}}, i[[SZ]] [[PARM1:%.+]], i[[SZ]] [[PARM2:%.+]])
290 // CHECK-DAG:   store i[[SZ]] [[PARM2]], i[[SZ]]* [[CAPE_ADDR:%.+]], align
291 // CHECK-64:    [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i32*
292 // CHECK-64:    [[TL:%.+]] = load i32, i32* [[CONV]], align
293 // CHECK-32:    [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align
294 // CHECK:       call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 0, i32 [[TL]])
295 // CHECK:       call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 2,
296 //
297 //
298 
299 
300 // CHECK:       define internal void [[HVT2]]([[S1]]* {{%.+}})
301 // CHECK:       call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 0, i32 1024)
302 // CHECK:       call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 1,
303 //
304 //
305 
306 
307 
308 
309 
310 
311 
312 
313 // CHECK:       define internal void [[HVT3]](i[[SZ]] [[PARM1:%.+]], i[[SZ]] [[PARM2:%.+]])
314 // CHECK-DAG:   store i[[SZ]] [[PARM1]], i[[SZ]]* [[CAPE_ADDR1:%.+]], align
315 // CHECK-DAG:   store i[[SZ]] [[PARM2]], i[[SZ]]* [[CAPE_ADDR2:%.+]], align
316 // CHECK-64:    [[CONV1:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR1]] to i32*
317 // CHECK-64:    [[CONV2:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR2]] to i32*
318 // CHECK-64:    [[NT:%.+]] = load i32, i32* [[CONV1]], align
319 // CHECK-64:    [[TL:%.+]] = load i32, i32* [[CONV2]], align
320 // CHECK-32:    [[NT:%.+]] = load i32, i32* [[CAPE_ADDR1]], align
321 // CHECK-32:    [[TL:%.+]] = load i32, i32* [[CAPE_ADDR2]], align
322 // CHECK:       call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 [[NT]], i32 [[TL]])
323 // CHECK:       call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 0,
324 //
325 //
326 // CHECK:       define internal void [[HVT4]](i[[SZ]] [[PARM:%.+]])
327 // CHECK-DAG:   store i[[SZ]] [[PARM]], i[[SZ]]* [[CAPE_ADDR:%.+]], align
328 // CHECK-64:    [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i32*
329 // CHECK-64:    [[TL:%.+]] = load i32, i32* [[CONV]], align
330 // CHECK-32:    [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align
331 // CHECK:       call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 0, i32 [[TL]])
332 // CHECK:       call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 0,
333 //
334 //
335 
336 
337 
338 
339 
340 // CHECK:       define internal void [[HVT5]](
341 // CHECK:       call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 0, i32 20)
342 // CHECK:       call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 0,
343 //
344 //
345 
346 
347 // CHECK:       define internal void [[HVT6]](i[[SZ]] [[PARM1:%.+]], i[[SZ]] [[PARM2:%.+]], i[[SZ]] [[PARM3:%.+]])
348 // CHECK-DAG:   store i[[SZ]] [[PARM3]], i[[SZ]]* [[CAPE_ADDR:%.+]], align
349 // CHECK:       [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i16*
350 // CHECK:       [[T:%.+]] = load i16, i16* [[CONV]], align
351 // CHECK:       [[NT:%.+]] = sext i16 [[T]] to i32
352 // CHECK:       call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 [[NT]], i32 1024)
353 // CHECK:       call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 2,
354 //
355 //
356 
357 
358 
359 #endif
360