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