1 ; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --function-signature --check-attributes --include-generated-funcs
2 ; RUN: opt -S -passes=openmp-opt < %s | FileCheck %s
3 ; RUN: opt -openmp-opt-disable-state-machine-rewrite -S -passes=openmp-opt < %s | FileCheck %s --check-prefix=CHECK-DISABLED
8 ;; void unknown_pure(void) __attribute__((pure));
9 ;; void unknown_no_openmp(void) __attribute__((assume("omp_no_openmp")));
12 ;; void no_parallel_region_in_here(void) {
17 ;; void no_state_machine_needed() {
18 ;; #pragma omp target teams
20 ;; no_parallel_region_in_here();
21 ;; unknown_no_openmp();
25 ;; void simple_state_machine() {
26 ;; #pragma omp target teams
28 ;; unknown_no_openmp();
29 ;; #pragma omp parallel
31 ;; no_parallel_region_in_here();
32 ;; #pragma omp parallel
37 ;; void simple_state_machine_interprocedural_after(void);
38 ;; void simple_state_machine_interprocedural_before(void) {
39 ;; #pragma omp parallel
42 ;; void simple_state_machine_interprocedural() {
43 ;; #pragma omp target teams
45 ;; unknown_no_openmp();
46 ;; simple_state_machine_interprocedural_before();
47 ;; no_parallel_region_in_here();
48 ;; #pragma omp parallel
50 ;; simple_state_machine_interprocedural_after();
53 ;; void simple_state_machine_interprocedural_after(void) {
54 ;; #pragma omp parallel
58 ;; void simple_state_machine_with_fallback() {
59 ;; #pragma omp target teams
61 ;; #pragma omp parallel
64 ;; #pragma omp parallel
69 ;; void simple_state_machine_no_openmp_attr() {
70 ;; #pragma omp target teams
72 ;; #pragma omp parallel
74 ;; unknown_no_openmp();
75 ;; #pragma omp parallel
80 ;; void simple_state_machine_pure() {
81 ;; #pragma omp target teams
83 ;; unknown_no_openmp();
84 ;; #pragma omp parallel
87 ;; #pragma omp parallel
92 ;; int omp_get_thread_num();
93 ;; void simple_state_machine_interprocedural_nested_recursive_after(int);
94 ;; void simple_state_machine_interprocedural_nested_recursive_after_after(void);
95 ;; void simple_state_machine_interprocedural_nested_recursive() {
96 ;; #pragma omp target teams
98 ;; simple_state_machine_interprocedural_nested_recursive_after(
99 ;; omp_get_thread_num());
103 ;; void simple_state_machine_interprocedural_nested_recursive_after(int a) {
106 ;; simple_state_machine_interprocedural_nested_recursive_after(a - 1);
107 ;; simple_state_machine_interprocedural_nested_recursive_after_after();
109 ;; void simple_state_machine_interprocedural_nested_recursive_after_after(void) {
110 ;; #pragma omp parallel
114 ;; __attribute__((weak)) void weak_callee_empty(void) {}
115 ;; void no_state_machine_weak_callee() {
116 ;; #pragma omp target teams
117 ;; { weak_callee_empty(); }
120 target triple = "nvptx64"
122 %struct.ident_t = type { i32, i32, i32, i32, i8* }
124 @0 = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00", align 1
125 @1 = private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @0, i32 0, i32 0) }, align 8
126 @__omp_offloading_14_a36502b_no_state_machine_needed_l14_exec_mode = weak constant i8 1
127 @__omp_offloading_14_a36502b_simple_state_machine_l22_exec_mode = weak constant i8 1
128 @__omp_offloading_14_a36502b_simple_state_machine_interprocedural_l39_exec_mode = weak constant i8 1
129 @__omp_offloading_14_a36502b_simple_state_machine_with_fallback_l55_exec_mode = weak constant i8 1
130 @__omp_offloading_14_a36502b_simple_state_machine_no_openmp_attr_l66_exec_mode = weak constant i8 1
131 @__omp_offloading_14_a36502b_simple_state_machine_pure_l77_exec_mode = weak constant i8 1
132 @__omp_offloading_14_a36502b_simple_state_machine_interprocedural_nested_recursive_l92_exec_mode = weak constant i8 1
133 @__omp_offloading_14_a36502b_no_state_machine_weak_callee_l112_exec_mode = weak constant i8 1
134 @2 = private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 2, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @0, i32 0, i32 0) }, align 8
135 @G = external global i32, align 4
136 @3 = private unnamed_addr constant %struct.ident_t { i32 0, i32 322, i32 2, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @0, i32 0, i32 0) }, align 8
137 @llvm.compiler.used = appending global [8 x i8*] [i8* @__omp_offloading_14_a36502b_no_state_machine_needed_l14_exec_mode, i8* @__omp_offloading_14_a36502b_simple_state_machine_l22_exec_mode, i8* @__omp_offloading_14_a36502b_simple_state_machine_interprocedural_l39_exec_mode, i8* @__omp_offloading_14_a36502b_simple_state_machine_with_fallback_l55_exec_mode, i8* @__omp_offloading_14_a36502b_simple_state_machine_no_openmp_attr_l66_exec_mode, i8* @__omp_offloading_14_a36502b_simple_state_machine_pure_l77_exec_mode, i8* @__omp_offloading_14_a36502b_simple_state_machine_interprocedural_nested_recursive_l92_exec_mode, i8* @__omp_offloading_14_a36502b_no_state_machine_weak_callee_l112_exec_mode], section "llvm.metadata"
139 define weak void @__omp_offloading_14_a36502b_no_state_machine_needed_l14() #0 {
141 %.zero.addr = alloca i32, align 4
142 %.threadid_temp. = alloca i32, align 4
143 store i32 0, i32* %.zero.addr, align 4
144 %0 = call i32 @__kmpc_target_init(%struct.ident_t* @1, i1 false, i1 true, i1 true)
145 %exec_user_code = icmp eq i32 %0, -1
146 br i1 %exec_user_code, label %user_code.entry, label %worker.exit
148 user_code.entry: ; preds = %entry
149 %1 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @1)
150 store i32 %1, i32* %.threadid_temp., align 4
151 call void @__omp_outlined__(i32* %.threadid_temp., i32* %.zero.addr) #3
152 call void @__kmpc_target_deinit(%struct.ident_t* @1, i1 false, i1 true)
155 worker.exit: ; preds = %entry
159 declare i32 @__kmpc_target_init(%struct.ident_t*, i1, i1, i1)
161 define internal void @__omp_outlined__(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 {
163 %.global_tid..addr = alloca i32*, align 8
164 %.bound_tid..addr = alloca i32*, align 8
165 store i32* %.global_tid., i32** %.global_tid..addr, align 8
166 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
167 call void @no_parallel_region_in_here() #7
168 call void @unknown_no_openmp() #8
172 define hidden void @no_parallel_region_in_here() #1 {
174 %0 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @2)
175 %1 = call i32 @__kmpc_single(%struct.ident_t* @2, i32 %0)
176 %2 = icmp ne i32 %1, 0
177 br i1 %2, label %omp_if.then, label %omp_if.end
179 omp_if.then: ; preds = %entry
180 store i32 0, i32* @G, align 4
181 call void @__kmpc_end_single(%struct.ident_t* @2, i32 %0)
184 omp_if.end: ; preds = %omp_if.then, %entry
185 call void @__kmpc_barrier(%struct.ident_t* @3, i32 %0)
189 declare void @unknown_no_openmp() #2
191 declare i32 @__kmpc_global_thread_num(%struct.ident_t*) #3
193 declare void @__kmpc_target_deinit(%struct.ident_t*, i1, i1)
195 define weak void @__omp_offloading_14_a36502b_simple_state_machine_l22() #0 {
197 %.zero.addr = alloca i32, align 4
198 %.threadid_temp. = alloca i32, align 4
199 store i32 0, i32* %.zero.addr, align 4
200 %0 = call i32 @__kmpc_target_init(%struct.ident_t* @1, i1 false, i1 true, i1 true)
201 %exec_user_code = icmp eq i32 %0, -1
202 br i1 %exec_user_code, label %user_code.entry, label %worker.exit
204 user_code.entry: ; preds = %entry
205 %1 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @1)
206 store i32 %1, i32* %.threadid_temp., align 4
207 call void @__omp_outlined__1(i32* %.threadid_temp., i32* %.zero.addr) #3
208 call void @__kmpc_target_deinit(%struct.ident_t* @1, i1 false, i1 true)
211 worker.exit: ; preds = %entry
215 define internal void @__omp_outlined__1(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 {
217 %.global_tid..addr = alloca i32*, align 8
218 %.bound_tid..addr = alloca i32*, align 8
219 %captured_vars_addrs = alloca [0 x i8*], align 8
220 %captured_vars_addrs1 = alloca [0 x i8*], align 8
221 store i32* %.global_tid., i32** %.global_tid..addr, align 8
222 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
223 call void @unknown_no_openmp() #8
224 %0 = load i32*, i32** %.global_tid..addr, align 8
225 %1 = load i32, i32* %0, align 4
226 %2 = bitcast [0 x i8*]* %captured_vars_addrs to i8**
227 call void @__kmpc_parallel_51(%struct.ident_t* @1, i32 %1, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__2 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__2_wrapper to i8*), i8** %2, i64 0)
228 call void @no_parallel_region_in_here() #7
229 %3 = bitcast [0 x i8*]* %captured_vars_addrs1 to i8**
230 call void @__kmpc_parallel_51(%struct.ident_t* @1, i32 %1, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__3 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__3_wrapper to i8*), i8** %3, i64 0)
234 define internal void @__omp_outlined__2(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 {
236 %.global_tid..addr = alloca i32*, align 8
237 %.bound_tid..addr = alloca i32*, align 8
238 store i32* %.global_tid., i32** %.global_tid..addr, align 8
239 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
244 declare void @p0() #4
246 define internal void @__omp_outlined__2_wrapper(i16 zeroext %0, i32 %1) #0 {
248 %.addr = alloca i16, align 2
249 %.addr1 = alloca i32, align 4
250 %.zero.addr = alloca i32, align 4
251 %global_args = alloca i8**, align 8
252 store i32 0, i32* %.zero.addr, align 4
253 store i16 %0, i16* %.addr, align 2
254 store i32 %1, i32* %.addr1, align 4
255 call void @__kmpc_get_shared_variables(i8*** %global_args)
256 call void @__omp_outlined__2(i32* %.addr1, i32* %.zero.addr) #3
260 declare void @__kmpc_get_shared_variables(i8***)
262 declare void @__kmpc_parallel_51(%struct.ident_t*, i32, i32, i32, i32, i8*, i8*, i8**, i64)
264 define internal void @__omp_outlined__3(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 {
266 %.global_tid..addr = alloca i32*, align 8
267 %.bound_tid..addr = alloca i32*, align 8
268 store i32* %.global_tid., i32** %.global_tid..addr, align 8
269 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
274 declare void @p1() #4
276 define internal void @__omp_outlined__3_wrapper(i16 zeroext %0, i32 %1) #0 {
278 %.addr = alloca i16, align 2
279 %.addr1 = alloca i32, align 4
280 %.zero.addr = alloca i32, align 4
281 %global_args = alloca i8**, align 8
282 store i32 0, i32* %.zero.addr, align 4
283 store i16 %0, i16* %.addr, align 2
284 store i32 %1, i32* %.addr1, align 4
285 call void @__kmpc_get_shared_variables(i8*** %global_args)
286 call void @__omp_outlined__3(i32* %.addr1, i32* %.zero.addr) #3
290 define weak void @__omp_offloading_14_a36502b_simple_state_machine_interprocedural_l39() #0 {
292 %.zero.addr = alloca i32, align 4
293 %.threadid_temp. = alloca i32, align 4
294 store i32 0, i32* %.zero.addr, align 4
295 %0 = call i32 @__kmpc_target_init(%struct.ident_t* @1, i1 false, i1 true, i1 true)
296 %exec_user_code = icmp eq i32 %0, -1
297 br i1 %exec_user_code, label %user_code.entry, label %worker.exit
299 user_code.entry: ; preds = %entry
300 %1 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @1)
301 store i32 %1, i32* %.threadid_temp., align 4
302 call void @__omp_outlined__4(i32* %.threadid_temp., i32* %.zero.addr) #3
303 call void @__kmpc_target_deinit(%struct.ident_t* @1, i1 false, i1 true)
306 worker.exit: ; preds = %entry
310 define internal void @__omp_outlined__4(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 {
312 %.global_tid..addr = alloca i32*, align 8
313 %.bound_tid..addr = alloca i32*, align 8
314 %captured_vars_addrs = alloca [0 x i8*], align 8
315 store i32* %.global_tid., i32** %.global_tid..addr, align 8
316 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
317 call void @unknown_no_openmp() #8
318 call void @simple_state_machine_interprocedural_before() #7
319 call void @no_parallel_region_in_here() #7
320 %0 = load i32*, i32** %.global_tid..addr, align 8
321 %1 = load i32, i32* %0, align 4
322 %2 = bitcast [0 x i8*]* %captured_vars_addrs to i8**
323 call void @__kmpc_parallel_51(%struct.ident_t* @1, i32 %1, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__5 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__5_wrapper to i8*), i8** %2, i64 0)
324 call void @simple_state_machine_interprocedural_after() #7
328 define hidden void @simple_state_machine_interprocedural_before() #1 {
330 %captured_vars_addrs = alloca [0 x i8*], align 8
331 %0 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @2)
332 %1 = bitcast [0 x i8*]* %captured_vars_addrs to i8**
333 call void @__kmpc_parallel_51(%struct.ident_t* @2, i32 %0, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__17 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__17_wrapper to i8*), i8** %1, i64 0)
337 define internal void @__omp_outlined__5(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 {
339 %.global_tid..addr = alloca i32*, align 8
340 %.bound_tid..addr = alloca i32*, align 8
341 store i32* %.global_tid., i32** %.global_tid..addr, align 8
342 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
347 define internal void @__omp_outlined__5_wrapper(i16 zeroext %0, i32 %1) #0 {
349 %.addr = alloca i16, align 2
350 %.addr1 = alloca i32, align 4
351 %.zero.addr = alloca i32, align 4
352 %global_args = alloca i8**, align 8
353 store i32 0, i32* %.zero.addr, align 4
354 store i16 %0, i16* %.addr, align 2
355 store i32 %1, i32* %.addr1, align 4
356 call void @__kmpc_get_shared_variables(i8*** %global_args)
357 call void @__omp_outlined__5(i32* %.addr1, i32* %.zero.addr) #3
361 define hidden void @simple_state_machine_interprocedural_after() #1 {
363 %captured_vars_addrs = alloca [0 x i8*], align 8
364 %0 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @2)
365 %1 = bitcast [0 x i8*]* %captured_vars_addrs to i8**
366 call void @__kmpc_parallel_51(%struct.ident_t* @2, i32 %0, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__18 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__18_wrapper to i8*), i8** %1, i64 0)
370 define weak void @__omp_offloading_14_a36502b_simple_state_machine_with_fallback_l55() #0 {
372 %.zero.addr = alloca i32, align 4
373 %.threadid_temp. = alloca i32, align 4
374 store i32 0, i32* %.zero.addr, align 4
375 %0 = call i32 @__kmpc_target_init(%struct.ident_t* @1, i1 false, i1 true, i1 true)
376 %exec_user_code = icmp eq i32 %0, -1
377 br i1 %exec_user_code, label %user_code.entry, label %worker.exit
379 user_code.entry: ; preds = %entry
380 %1 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @1)
381 store i32 %1, i32* %.threadid_temp., align 4
382 call void @__omp_outlined__6(i32* %.threadid_temp., i32* %.zero.addr) #3
383 call void @__kmpc_target_deinit(%struct.ident_t* @1, i1 false, i1 true)
386 worker.exit: ; preds = %entry
390 define internal void @__omp_outlined__6(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 {
392 %.global_tid..addr = alloca i32*, align 8
393 %.bound_tid..addr = alloca i32*, align 8
394 %captured_vars_addrs = alloca [0 x i8*], align 8
395 %captured_vars_addrs1 = alloca [0 x i8*], align 8
396 store i32* %.global_tid., i32** %.global_tid..addr, align 8
397 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
398 %0 = load i32*, i32** %.global_tid..addr, align 8
399 %1 = load i32, i32* %0, align 4
400 %2 = bitcast [0 x i8*]* %captured_vars_addrs to i8**
401 call void @__kmpc_parallel_51(%struct.ident_t* @1, i32 %1, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__7 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__7_wrapper to i8*), i8** %2, i64 0)
402 %call = call i32 @unknown() #7
403 %3 = bitcast [0 x i8*]* %captured_vars_addrs1 to i8**
404 call void @__kmpc_parallel_51(%struct.ident_t* @1, i32 %1, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__8 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__8_wrapper to i8*), i8** %3, i64 0)
408 define internal void @__omp_outlined__7(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 {
410 %.global_tid..addr = alloca i32*, align 8
411 %.bound_tid..addr = alloca i32*, align 8
412 store i32* %.global_tid., i32** %.global_tid..addr, align 8
413 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
418 define internal void @__omp_outlined__7_wrapper(i16 zeroext %0, i32 %1) #0 {
420 %.addr = alloca i16, align 2
421 %.addr1 = alloca i32, align 4
422 %.zero.addr = alloca i32, align 4
423 %global_args = alloca i8**, align 8
424 store i32 0, i32* %.zero.addr, align 4
425 store i16 %0, i16* %.addr, align 2
426 store i32 %1, i32* %.addr1, align 4
427 call void @__kmpc_get_shared_variables(i8*** %global_args)
428 call void @__omp_outlined__7(i32* %.addr1, i32* %.zero.addr) #3
432 declare i32 @unknown() #4
434 define internal void @__omp_outlined__8(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 {
436 %.global_tid..addr = alloca i32*, align 8
437 %.bound_tid..addr = alloca i32*, align 8
438 store i32* %.global_tid., i32** %.global_tid..addr, align 8
439 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
444 define internal void @__omp_outlined__8_wrapper(i16 zeroext %0, i32 %1) #0 {
446 %.addr = alloca i16, align 2
447 %.addr1 = alloca i32, align 4
448 %.zero.addr = alloca i32, align 4
449 %global_args = alloca i8**, align 8
450 store i32 0, i32* %.zero.addr, align 4
451 store i16 %0, i16* %.addr, align 2
452 store i32 %1, i32* %.addr1, align 4
453 call void @__kmpc_get_shared_variables(i8*** %global_args)
454 call void @__omp_outlined__8(i32* %.addr1, i32* %.zero.addr) #3
458 define weak void @__omp_offloading_14_a36502b_simple_state_machine_no_openmp_attr_l66() #0 {
460 %.zero.addr = alloca i32, align 4
461 %.threadid_temp. = alloca i32, align 4
462 store i32 0, i32* %.zero.addr, align 4
463 %0 = call i32 @__kmpc_target_init(%struct.ident_t* @1, i1 false, i1 true, i1 true)
464 %exec_user_code = icmp eq i32 %0, -1
465 br i1 %exec_user_code, label %user_code.entry, label %worker.exit
467 user_code.entry: ; preds = %entry
468 %1 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @1)
469 store i32 %1, i32* %.threadid_temp., align 4
470 call void @__omp_outlined__9(i32* %.threadid_temp., i32* %.zero.addr) #3
471 call void @__kmpc_target_deinit(%struct.ident_t* @1, i1 false, i1 true)
474 worker.exit: ; preds = %entry
478 define internal void @__omp_outlined__9(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 {
480 %.global_tid..addr = alloca i32*, align 8
481 %.bound_tid..addr = alloca i32*, align 8
482 %captured_vars_addrs = alloca [0 x i8*], align 8
483 %captured_vars_addrs1 = alloca [0 x i8*], align 8
484 store i32* %.global_tid., i32** %.global_tid..addr, align 8
485 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
486 %0 = load i32*, i32** %.global_tid..addr, align 8
487 %1 = load i32, i32* %0, align 4
488 %2 = bitcast [0 x i8*]* %captured_vars_addrs to i8**
489 call void @__kmpc_parallel_51(%struct.ident_t* @1, i32 %1, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__10 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__10_wrapper to i8*), i8** %2, i64 0)
490 call void @unknown_no_openmp() #8
491 %3 = bitcast [0 x i8*]* %captured_vars_addrs1 to i8**
492 call void @__kmpc_parallel_51(%struct.ident_t* @1, i32 %1, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__11 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__11_wrapper to i8*), i8** %3, i64 0)
496 define internal void @__omp_outlined__10(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 {
498 %.global_tid..addr = alloca i32*, align 8
499 %.bound_tid..addr = alloca i32*, align 8
500 store i32* %.global_tid., i32** %.global_tid..addr, align 8
501 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
506 define internal void @__omp_outlined__10_wrapper(i16 zeroext %0, i32 %1) #0 {
508 %.addr = alloca i16, align 2
509 %.addr1 = alloca i32, align 4
510 %.zero.addr = alloca i32, align 4
511 %global_args = alloca i8**, align 8
512 store i32 0, i32* %.zero.addr, align 4
513 store i16 %0, i16* %.addr, align 2
514 store i32 %1, i32* %.addr1, align 4
515 call void @__kmpc_get_shared_variables(i8*** %global_args)
516 call void @__omp_outlined__10(i32* %.addr1, i32* %.zero.addr) #3
520 define internal void @__omp_outlined__11(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 {
522 %.global_tid..addr = alloca i32*, align 8
523 %.bound_tid..addr = alloca i32*, align 8
524 store i32* %.global_tid., i32** %.global_tid..addr, align 8
525 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
530 define internal void @__omp_outlined__11_wrapper(i16 zeroext %0, i32 %1) #0 {
532 %.addr = alloca i16, align 2
533 %.addr1 = alloca i32, align 4
534 %.zero.addr = alloca i32, align 4
535 %global_args = alloca i8**, align 8
536 store i32 0, i32* %.zero.addr, align 4
537 store i16 %0, i16* %.addr, align 2
538 store i32 %1, i32* %.addr1, align 4
539 call void @__kmpc_get_shared_variables(i8*** %global_args)
540 call void @__omp_outlined__11(i32* %.addr1, i32* %.zero.addr) #3
544 define weak void @__omp_offloading_14_a36502b_simple_state_machine_pure_l77() #0 {
546 %.zero.addr = alloca i32, align 4
547 %.threadid_temp. = alloca i32, align 4
548 store i32 0, i32* %.zero.addr, align 4
549 %0 = call i32 @__kmpc_target_init(%struct.ident_t* @1, i1 false, i1 true, i1 true)
550 %exec_user_code = icmp eq i32 %0, -1
551 br i1 %exec_user_code, label %user_code.entry, label %worker.exit
553 user_code.entry: ; preds = %entry
554 %1 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @1)
555 store i32 %1, i32* %.threadid_temp., align 4
556 call void @__omp_outlined__12(i32* %.threadid_temp., i32* %.zero.addr) #3
557 call void @__kmpc_target_deinit(%struct.ident_t* @1, i1 false, i1 true)
560 worker.exit: ; preds = %entry
564 define internal void @__omp_outlined__12(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 {
566 %.global_tid..addr = alloca i32*, align 8
567 %.bound_tid..addr = alloca i32*, align 8
568 %captured_vars_addrs = alloca [0 x i8*], align 8
569 %captured_vars_addrs1 = alloca [0 x i8*], align 8
570 store i32* %.global_tid., i32** %.global_tid..addr, align 8
571 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
572 call void @unknown_no_openmp() #8
573 %0 = load i32*, i32** %.global_tid..addr, align 8
574 %1 = load i32, i32* %0, align 4
575 %2 = bitcast [0 x i8*]* %captured_vars_addrs to i8**
576 call void @__kmpc_parallel_51(%struct.ident_t* @1, i32 %1, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__13 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__13_wrapper to i8*), i8** %2, i64 0)
577 call void @unknown_pure() #9
578 %3 = bitcast [0 x i8*]* %captured_vars_addrs1 to i8**
579 call void @__kmpc_parallel_51(%struct.ident_t* @1, i32 %1, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__14 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__14_wrapper to i8*), i8** %3, i64 0)
583 define internal void @__omp_outlined__13(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 {
585 %.global_tid..addr = alloca i32*, align 8
586 %.bound_tid..addr = alloca i32*, align 8
587 store i32* %.global_tid., i32** %.global_tid..addr, align 8
588 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
593 define internal void @__omp_outlined__13_wrapper(i16 zeroext %0, i32 %1) #0 {
595 %.addr = alloca i16, align 2
596 %.addr1 = alloca i32, align 4
597 %.zero.addr = alloca i32, align 4
598 %global_args = alloca i8**, align 8
599 store i32 0, i32* %.zero.addr, align 4
600 store i16 %0, i16* %.addr, align 2
601 store i32 %1, i32* %.addr1, align 4
602 call void @__kmpc_get_shared_variables(i8*** %global_args)
603 call void @__omp_outlined__13(i32* %.addr1, i32* %.zero.addr) #3
607 declare void @unknown_pure() #5
609 define internal void @__omp_outlined__14(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 {
611 %.global_tid..addr = alloca i32*, align 8
612 %.bound_tid..addr = alloca i32*, align 8
613 store i32* %.global_tid., i32** %.global_tid..addr, align 8
614 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
619 define internal void @__omp_outlined__14_wrapper(i16 zeroext %0, i32 %1) #0 {
621 %.addr = alloca i16, align 2
622 %.addr1 = alloca i32, align 4
623 %.zero.addr = alloca i32, align 4
624 %global_args = alloca i8**, align 8
625 store i32 0, i32* %.zero.addr, align 4
626 store i16 %0, i16* %.addr, align 2
627 store i32 %1, i32* %.addr1, align 4
628 call void @__kmpc_get_shared_variables(i8*** %global_args)
629 call void @__omp_outlined__14(i32* %.addr1, i32* %.zero.addr) #3
633 define weak void @__omp_offloading_14_a36502b_simple_state_machine_interprocedural_nested_recursive_l92() #0 {
635 %.zero.addr = alloca i32, align 4
636 %.threadid_temp. = alloca i32, align 4
637 store i32 0, i32* %.zero.addr, align 4
638 %0 = call i32 @__kmpc_target_init(%struct.ident_t* @1, i1 false, i1 true, i1 true)
639 %exec_user_code = icmp eq i32 %0, -1
640 br i1 %exec_user_code, label %user_code.entry, label %worker.exit
642 user_code.entry: ; preds = %entry
643 %1 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @1)
644 store i32 %1, i32* %.threadid_temp., align 4
645 call void @__omp_outlined__15(i32* %.threadid_temp., i32* %.zero.addr) #3
646 call void @__kmpc_target_deinit(%struct.ident_t* @1, i1 false, i1 true)
649 worker.exit: ; preds = %entry
653 define internal void @__omp_outlined__15(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 {
655 %.global_tid..addr = alloca i32*, align 8
656 %.bound_tid..addr = alloca i32*, align 8
657 store i32* %.global_tid., i32** %.global_tid..addr, align 8
658 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
659 %call = call i32 bitcast (i32 (...)* @omp_get_thread_num to i32 ()*)() #7
660 call void @simple_state_machine_interprocedural_nested_recursive_after(i32 %call) #7
664 define hidden void @simple_state_machine_interprocedural_nested_recursive_after(i32 %a) #1 {
666 %a.addr = alloca i32, align 4
667 store i32 %a, i32* %a.addr, align 4
668 %0 = load i32, i32* %a.addr, align 4
669 %cmp = icmp eq i32 %0, 0
670 br i1 %cmp, label %if.then, label %if.end
672 if.then: ; preds = %entry
675 if.end: ; preds = %entry
676 %1 = load i32, i32* %a.addr, align 4
677 %sub = sub nsw i32 %1, 1
678 call void @simple_state_machine_interprocedural_nested_recursive_after(i32 %sub) #7
679 call void @simple_state_machine_interprocedural_nested_recursive_after_after() #7
682 return: ; preds = %if.end, %if.then
686 declare i32 @omp_get_thread_num(...) #4
688 define weak void @__omp_offloading_14_a36502b_no_state_machine_weak_callee_l112() #0 {
690 %.zero.addr = alloca i32, align 4
691 %.threadid_temp. = alloca i32, align 4
692 store i32 0, i32* %.zero.addr, align 4
693 %0 = call i32 @__kmpc_target_init(%struct.ident_t* @1, i1 false, i1 true, i1 true)
694 %exec_user_code = icmp eq i32 %0, -1
695 br i1 %exec_user_code, label %user_code.entry, label %worker.exit
697 user_code.entry: ; preds = %entry
698 %1 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @1)
699 store i32 %1, i32* %.threadid_temp., align 4
700 call void @__omp_outlined__16(i32* %.threadid_temp., i32* %.zero.addr) #3
701 call void @__kmpc_target_deinit(%struct.ident_t* @1, i1 false, i1 true)
704 worker.exit: ; preds = %entry
708 define internal void @__omp_outlined__16(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 {
710 %.global_tid..addr = alloca i32*, align 8
711 %.bound_tid..addr = alloca i32*, align 8
712 store i32* %.global_tid., i32** %.global_tid..addr, align 8
713 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
714 call void @weak_callee_empty() #7
718 define weak hidden void @weak_callee_empty() #1 {
723 declare i32 @__kmpc_single(%struct.ident_t*, i32) #6
725 declare void @__kmpc_end_single(%struct.ident_t*, i32) #6
727 declare void @__kmpc_barrier(%struct.ident_t*, i32) #6
729 define internal void @__omp_outlined__17(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 {
731 %.global_tid..addr = alloca i32*, align 8
732 %.bound_tid..addr = alloca i32*, align 8
733 store i32* %.global_tid., i32** %.global_tid..addr, align 8
734 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
739 define internal void @__omp_outlined__17_wrapper(i16 zeroext %0, i32 %1) #0 {
741 %.addr = alloca i16, align 2
742 %.addr1 = alloca i32, align 4
743 %.zero.addr = alloca i32, align 4
744 %global_args = alloca i8**, align 8
745 store i32 0, i32* %.zero.addr, align 4
746 store i16 %0, i16* %.addr, align 2
747 store i32 %1, i32* %.addr1, align 4
748 call void @__kmpc_get_shared_variables(i8*** %global_args)
749 call void @__omp_outlined__17(i32* %.addr1, i32* %.zero.addr) #3
753 define internal void @__omp_outlined__18(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 {
755 %.global_tid..addr = alloca i32*, align 8
756 %.bound_tid..addr = alloca i32*, align 8
757 store i32* %.global_tid., i32** %.global_tid..addr, align 8
758 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
763 define internal void @__omp_outlined__18_wrapper(i16 zeroext %0, i32 %1) #0 {
765 %.addr = alloca i16, align 2
766 %.addr1 = alloca i32, align 4
767 %.zero.addr = alloca i32, align 4
768 %global_args = alloca i8**, align 8
769 store i32 0, i32* %.zero.addr, align 4
770 store i16 %0, i16* %.addr, align 2
771 store i32 %1, i32* %.addr1, align 4
772 call void @__kmpc_get_shared_variables(i8*** %global_args)
773 call void @__omp_outlined__18(i32* %.addr1, i32* %.zero.addr) #3
777 define hidden void @simple_state_machine_interprocedural_nested_recursive_after_after() #1 {
779 %captured_vars_addrs = alloca [0 x i8*], align 8
780 %0 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @2)
781 %1 = bitcast [0 x i8*]* %captured_vars_addrs to i8**
782 call void @__kmpc_parallel_51(%struct.ident_t* @2, i32 %0, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__19 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__19_wrapper to i8*), i8** %1, i64 0)
786 define internal void @__omp_outlined__19(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 {
788 %.global_tid..addr = alloca i32*, align 8
789 %.bound_tid..addr = alloca i32*, align 8
790 store i32* %.global_tid., i32** %.global_tid..addr, align 8
791 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
796 define internal void @__omp_outlined__19_wrapper(i16 zeroext %0, i32 %1) #0 {
798 %.addr = alloca i16, align 2
799 %.addr1 = alloca i32, align 4
800 %.zero.addr = alloca i32, align 4
801 %global_args = alloca i8**, align 8
802 store i32 0, i32* %.zero.addr, align 4
803 store i16 %0, i16* %.addr, align 2
804 store i32 %1, i32* %.addr1, align 4
805 call void @__kmpc_get_shared_variables(i8*** %global_args)
806 call void @__omp_outlined__19(i32* %.addr1, i32* %.zero.addr) #3
810 attributes #0 = { convergent noinline norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32,+sm_20" }
811 attributes #1 = { convergent noinline nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32,+sm_20" }
812 attributes #2 = { convergent "frame-pointer"="none" "llvm.assume"="omp_no_openmp" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32,+sm_20" }
813 attributes #3 = { nounwind }
814 attributes #4 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32,+sm_20" }
815 attributes #5 = { convergent nounwind readonly willreturn "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32,+sm_20" }
816 attributes #6 = { convergent nounwind }
817 attributes #7 = { convergent }
818 attributes #8 = { convergent "llvm.assume"="omp_no_openmp" }
819 attributes #9 = { convergent nounwind readonly willreturn }
821 !omp_offload.info = !{!0, !1, !2, !3, !4, !5, !6, !7}
822 !nvvm.annotations = !{!8, !9, !10, !11, !12, !13, !14, !15}
823 !llvm.module.flags = !{!16, !17, !18}
825 !0 = !{i32 0, i32 20, i32 171331627, !"simple_state_machine_interprocedural", i32 39, i32 2}
826 !1 = !{i32 0, i32 20, i32 171331627, !"simple_state_machine_no_openmp_attr", i32 66, i32 4}
827 !2 = !{i32 0, i32 20, i32 171331627, !"no_state_machine_needed", i32 14, i32 0}
828 !3 = !{i32 0, i32 20, i32 171331627, !"simple_state_machine_with_fallback", i32 55, i32 3}
829 !4 = !{i32 0, i32 20, i32 171331627, !"simple_state_machine_pure", i32 77, i32 5}
830 !5 = !{i32 0, i32 20, i32 171331627, !"simple_state_machine_interprocedural_nested_recursive", i32 92, i32 6}
831 !6 = !{i32 0, i32 20, i32 171331627, !"no_state_machine_weak_callee", i32 112, i32 7}
832 !7 = !{i32 0, i32 20, i32 171331627, !"simple_state_machine", i32 22, i32 1}
833 !8 = !{void ()* @__omp_offloading_14_a36502b_no_state_machine_needed_l14, !"kernel", i32 1}
834 !9 = !{void ()* @__omp_offloading_14_a36502b_simple_state_machine_l22, !"kernel", i32 1}
835 !10 = !{void ()* @__omp_offloading_14_a36502b_simple_state_machine_interprocedural_l39, !"kernel", i32 1}
836 !11 = !{void ()* @__omp_offloading_14_a36502b_simple_state_machine_with_fallback_l55, !"kernel", i32 1}
837 !12 = !{void ()* @__omp_offloading_14_a36502b_simple_state_machine_no_openmp_attr_l66, !"kernel", i32 1}
838 !13 = !{void ()* @__omp_offloading_14_a36502b_simple_state_machine_pure_l77, !"kernel", i32 1}
839 !14 = !{void ()* @__omp_offloading_14_a36502b_simple_state_machine_interprocedural_nested_recursive_l92, !"kernel", i32 1}
840 !15 = !{void ()* @__omp_offloading_14_a36502b_no_state_machine_weak_callee_l112, !"kernel", i32 1}
841 !16 = !{i32 1, !"wchar_size", i32 4}
842 !17 = !{i32 7, !"openmp", i32 50}
843 !18 = !{i32 7, !"openmp-device", i32 50}
844 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
845 ; CHECK-LABEL: define {{[^@]+}}@__omp_offloading_14_a36502b_no_state_machine_needed_l14
846 ; CHECK-SAME: () #[[ATTR0:[0-9]+]] {
848 ; CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
849 ; CHECK-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
850 ; CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i1 false, i1 false, i1 true)
851 ; CHECK-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
852 ; CHECK-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
853 ; CHECK: user_code.entry:
854 ; CHECK-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR3:[0-9]+]]
855 ; CHECK-NEXT: call void @__omp_outlined__(i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR3]]
856 ; CHECK-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true)
857 ; CHECK-NEXT: ret void
858 ; CHECK: worker.exit:
859 ; CHECK-NEXT: ret void
862 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
863 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__
864 ; CHECK-SAME: (i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
866 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
867 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
868 ; CHECK-NEXT: call void @no_parallel_region_in_here.internalized() #[[ATTR7:[0-9]+]]
869 ; CHECK-NEXT: call void @unknown_no_openmp() #[[ATTR8:[0-9]+]]
870 ; CHECK-NEXT: ret void
873 ; CHECK: Function Attrs: convergent noinline nounwind
874 ; CHECK-LABEL: define {{[^@]+}}@no_parallel_region_in_here.internalized
875 ; CHECK-SAME: () #[[ATTR1:[0-9]+]] {
877 ; CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* noundef @[[GLOB2:[0-9]+]]) #[[ATTR3]]
878 ; CHECK-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_single(%struct.ident_t* noundef @[[GLOB2]], i32 [[TMP0]]) #[[ATTR3]]
879 ; CHECK-NEXT: [[TMP2:%.*]] = icmp ne i32 [[TMP1]], 0
880 ; CHECK-NEXT: br i1 [[TMP2]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]]
881 ; CHECK: omp_if.then:
882 ; CHECK-NEXT: store i32 0, i32* @G, align 4
883 ; CHECK-NEXT: call void @__kmpc_end_single(%struct.ident_t* noundef @[[GLOB2]], i32 [[TMP0]]) #[[ATTR3]]
884 ; CHECK-NEXT: br label [[OMP_IF_END]]
886 ; CHECK-NEXT: call void @__kmpc_barrier(%struct.ident_t* noundef @[[GLOB3:[0-9]+]], i32 [[TMP0]]) #[[ATTR3]]
887 ; CHECK-NEXT: ret void
890 ; CHECK: Function Attrs: convergent noinline nounwind
891 ; CHECK-LABEL: define {{[^@]+}}@no_parallel_region_in_here
892 ; CHECK-SAME: () #[[ATTR1]] {
894 ; CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2]])
895 ; CHECK-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_single(%struct.ident_t* @[[GLOB2]], i32 [[TMP0]])
896 ; CHECK-NEXT: [[TMP2:%.*]] = icmp ne i32 [[TMP1]], 0
897 ; CHECK-NEXT: br i1 [[TMP2]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]]
898 ; CHECK: omp_if.then:
899 ; CHECK-NEXT: store i32 0, i32* @G, align 4
900 ; CHECK-NEXT: call void @__kmpc_end_single(%struct.ident_t* @[[GLOB2]], i32 [[TMP0]])
901 ; CHECK-NEXT: br label [[OMP_IF_END]]
903 ; CHECK-NEXT: call void @__kmpc_barrier(%struct.ident_t* @[[GLOB3]], i32 [[TMP0]])
904 ; CHECK-NEXT: ret void
907 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
908 ; CHECK-LABEL: define {{[^@]+}}@__omp_offloading_14_a36502b_simple_state_machine_l22
909 ; CHECK-SAME: () #[[ATTR0]] {
911 ; CHECK-NEXT: [[WORKER_WORK_FN_ADDR:%.*]] = alloca i8*, align 8
912 ; CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
913 ; CHECK-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
914 ; CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i1 false, i1 false, i1 true)
915 ; CHECK-NEXT: [[THREAD_IS_WORKER:%.*]] = icmp ne i32 [[TMP0]], -1
916 ; CHECK-NEXT: br i1 [[THREAD_IS_WORKER]], label [[WORKER_STATE_MACHINE_BEGIN:%.*]], label [[THREAD_USER_CODE_CHECK:%.*]]
917 ; CHECK: worker_state_machine.begin:
918 ; CHECK-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]])
919 ; CHECK-NEXT: [[WORKER_IS_ACTIVE:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORKER_WORK_FN_ADDR]])
920 ; CHECK-NEXT: [[WORKER_WORK_FN:%.*]] = load i8*, i8** [[WORKER_WORK_FN_ADDR]], align 8
921 ; CHECK-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)*
922 ; CHECK-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null
923 ; CHECK-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]]
924 ; CHECK: worker_state_machine.finished:
925 ; CHECK-NEXT: ret void
926 ; CHECK: worker_state_machine.is_active.check:
927 ; CHECK-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]]
928 ; CHECK: worker_state_machine.parallel_region.check:
929 ; CHECK-NEXT: [[WORKER_CHECK_PARALLEL_REGION:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], bitcast (i8* @__omp_outlined__2_wrapper.ID to void (i16, i32)*)
930 ; CHECK-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK1:%.*]]
931 ; CHECK: worker_state_machine.parallel_region.execute:
932 ; CHECK-NEXT: call void @__omp_outlined__2_wrapper(i16 0, i32 [[TMP0]])
933 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]]
934 ; CHECK: worker_state_machine.parallel_region.check1:
935 ; CHECK-NEXT: br i1 true, label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE2:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK3:%.*]]
936 ; CHECK: worker_state_machine.parallel_region.execute2:
937 ; CHECK-NEXT: call void @__omp_outlined__3_wrapper(i16 0, i32 [[TMP0]])
938 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]]
939 ; CHECK: worker_state_machine.parallel_region.check3:
940 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]]
941 ; CHECK: worker_state_machine.parallel_region.end:
942 ; CHECK-NEXT: call void @__kmpc_kernel_end_parallel()
943 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_DONE_BARRIER]]
944 ; CHECK: worker_state_machine.done.barrier:
945 ; CHECK-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]])
946 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_BEGIN]]
947 ; CHECK: thread.user_code.check:
948 ; CHECK-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
949 ; CHECK-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
950 ; CHECK: user_code.entry:
951 ; CHECK-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR3]]
952 ; CHECK-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4
953 ; CHECK-NEXT: call void @__omp_outlined__1(i32* noalias nocapture noundef nonnull readonly align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR3]]
954 ; CHECK-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true)
955 ; CHECK-NEXT: ret void
956 ; CHECK: worker.exit:
957 ; CHECK-NEXT: ret void
960 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
961 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__1
962 ; CHECK-SAME: (i32* noalias nocapture nofree noundef nonnull readonly align 4 dereferenceable(4) [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
964 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
965 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
966 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
967 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS1:%.*]] = alloca [0 x i8*], align 8
968 ; CHECK-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
969 ; CHECK-NEXT: call void @unknown_no_openmp() #[[ATTR8]]
970 ; CHECK-NEXT: [[TMP0:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4
971 ; CHECK-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
972 ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__2 to i8*), i8* noundef @__omp_outlined__2_wrapper.ID, i8** noundef [[TMP1]], i64 noundef 0)
973 ; CHECK-NEXT: call void @no_parallel_region_in_here.internalized() #[[ATTR7]]
974 ; CHECK-NEXT: [[TMP2:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS1]] to i8**
975 ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__3 to i8*), i8* noundef @__omp_outlined__3_wrapper.ID, i8** noundef [[TMP2]], i64 noundef 0)
976 ; CHECK-NEXT: ret void
979 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
980 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__2
981 ; CHECK-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
983 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
984 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
985 ; CHECK-NEXT: call void @p0() #[[ATTR9:[0-9]+]]
986 ; CHECK-NEXT: ret void
989 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
990 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__2_wrapper
991 ; CHECK-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
993 ; CHECK-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
994 ; CHECK-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
995 ; CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
996 ; CHECK-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
997 ; CHECK-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
998 ; CHECK-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
999 ; CHECK-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
1000 ; CHECK-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
1001 ; CHECK-NEXT: call void @__omp_outlined__2(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
1002 ; CHECK-NEXT: ret void
1005 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1006 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__3
1007 ; CHECK-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1008 ; CHECK-NEXT: entry:
1009 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1010 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1011 ; CHECK-NEXT: call void @p1() #[[ATTR9]]
1012 ; CHECK-NEXT: ret void
1015 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1016 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__3_wrapper
1017 ; CHECK-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
1018 ; CHECK-NEXT: entry:
1019 ; CHECK-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
1020 ; CHECK-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
1021 ; CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1022 ; CHECK-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
1023 ; CHECK-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
1024 ; CHECK-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
1025 ; CHECK-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
1026 ; CHECK-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
1027 ; CHECK-NEXT: call void @__omp_outlined__3(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
1028 ; CHECK-NEXT: ret void
1031 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1032 ; CHECK-LABEL: define {{[^@]+}}@__omp_offloading_14_a36502b_simple_state_machine_interprocedural_l39
1033 ; CHECK-SAME: () #[[ATTR0]] {
1034 ; CHECK-NEXT: entry:
1035 ; CHECK-NEXT: [[WORKER_WORK_FN_ADDR:%.*]] = alloca i8*, align 8
1036 ; CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1037 ; CHECK-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
1038 ; CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i1 false, i1 false, i1 true)
1039 ; CHECK-NEXT: [[THREAD_IS_WORKER:%.*]] = icmp ne i32 [[TMP0]], -1
1040 ; CHECK-NEXT: br i1 [[THREAD_IS_WORKER]], label [[WORKER_STATE_MACHINE_BEGIN:%.*]], label [[THREAD_USER_CODE_CHECK:%.*]]
1041 ; CHECK: worker_state_machine.begin:
1042 ; CHECK-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]])
1043 ; CHECK-NEXT: [[WORKER_IS_ACTIVE:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORKER_WORK_FN_ADDR]])
1044 ; CHECK-NEXT: [[WORKER_WORK_FN:%.*]] = load i8*, i8** [[WORKER_WORK_FN_ADDR]], align 8
1045 ; CHECK-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)*
1046 ; CHECK-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null
1047 ; CHECK-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]]
1048 ; CHECK: worker_state_machine.finished:
1049 ; CHECK-NEXT: ret void
1050 ; CHECK: worker_state_machine.is_active.check:
1051 ; CHECK-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]]
1052 ; CHECK: worker_state_machine.parallel_region.check:
1053 ; CHECK-NEXT: [[WORKER_CHECK_PARALLEL_REGION:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], @__omp_outlined__17_wrapper
1054 ; CHECK-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK1:%.*]]
1055 ; CHECK: worker_state_machine.parallel_region.execute:
1056 ; CHECK-NEXT: call void @__omp_outlined__17_wrapper(i16 0, i32 [[TMP0]])
1057 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]]
1058 ; CHECK: worker_state_machine.parallel_region.check1:
1059 ; CHECK-NEXT: [[WORKER_CHECK_PARALLEL_REGION4:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], bitcast (i8* @__omp_outlined__5_wrapper.ID to void (i16, i32)*)
1060 ; CHECK-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION4]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE2:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK3:%.*]]
1061 ; CHECK: worker_state_machine.parallel_region.execute2:
1062 ; CHECK-NEXT: call void @__omp_outlined__5_wrapper(i16 0, i32 [[TMP0]])
1063 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]]
1064 ; CHECK: worker_state_machine.parallel_region.check3:
1065 ; CHECK-NEXT: br i1 true, label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE5:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK6:%.*]]
1066 ; CHECK: worker_state_machine.parallel_region.execute5:
1067 ; CHECK-NEXT: call void @__omp_outlined__18_wrapper(i16 0, i32 [[TMP0]])
1068 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]]
1069 ; CHECK: worker_state_machine.parallel_region.check6:
1070 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]]
1071 ; CHECK: worker_state_machine.parallel_region.end:
1072 ; CHECK-NEXT: call void @__kmpc_kernel_end_parallel()
1073 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_DONE_BARRIER]]
1074 ; CHECK: worker_state_machine.done.barrier:
1075 ; CHECK-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]])
1076 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_BEGIN]]
1077 ; CHECK: thread.user_code.check:
1078 ; CHECK-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
1079 ; CHECK-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
1080 ; CHECK: user_code.entry:
1081 ; CHECK-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR3]]
1082 ; CHECK-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4
1083 ; CHECK-NEXT: call void @__omp_outlined__4(i32* noalias nocapture noundef nonnull readonly align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR3]]
1084 ; CHECK-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true)
1085 ; CHECK-NEXT: ret void
1086 ; CHECK: worker.exit:
1087 ; CHECK-NEXT: ret void
1090 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1091 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__4
1092 ; CHECK-SAME: (i32* noalias nocapture nofree noundef nonnull readonly align 4 dereferenceable(4) [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1093 ; CHECK-NEXT: entry:
1094 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1095 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1096 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
1097 ; CHECK-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
1098 ; CHECK-NEXT: call void @unknown_no_openmp() #[[ATTR8]]
1099 ; CHECK-NEXT: call void @simple_state_machine_interprocedural_before.internalized() #[[ATTR7]]
1100 ; CHECK-NEXT: call void @no_parallel_region_in_here.internalized() #[[ATTR7]]
1101 ; CHECK-NEXT: [[TMP0:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4
1102 ; CHECK-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
1103 ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__5 to i8*), i8* noundef @__omp_outlined__5_wrapper.ID, i8** noundef [[TMP1]], i64 noundef 0)
1104 ; CHECK-NEXT: call void @simple_state_machine_interprocedural_after.internalized() #[[ATTR7]]
1105 ; CHECK-NEXT: ret void
1108 ; CHECK: Function Attrs: convergent noinline nounwind
1109 ; CHECK-LABEL: define {{[^@]+}}@simple_state_machine_interprocedural_before.internalized
1110 ; CHECK-SAME: () #[[ATTR1]] {
1111 ; CHECK-NEXT: entry:
1112 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
1113 ; CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* noundef @[[GLOB2]]) #[[ATTR3]]
1114 ; CHECK-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
1115 ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB2]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__17 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__17_wrapper to i8*), i8** noundef [[TMP1]], i64 noundef 0)
1116 ; CHECK-NEXT: ret void
1119 ; CHECK: Function Attrs: convergent noinline nounwind
1120 ; CHECK-LABEL: define {{[^@]+}}@simple_state_machine_interprocedural_before
1121 ; CHECK-SAME: () #[[ATTR1]] {
1122 ; CHECK-NEXT: entry:
1123 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
1124 ; CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2]])
1125 ; CHECK-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
1126 ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP0]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__17 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__17_wrapper to i8*), i8** [[TMP1]], i64 0)
1127 ; CHECK-NEXT: ret void
1130 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1131 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__5
1132 ; CHECK-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1133 ; CHECK-NEXT: entry:
1134 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1135 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1136 ; CHECK-NEXT: call void @p1() #[[ATTR9]]
1137 ; CHECK-NEXT: ret void
1140 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1141 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__5_wrapper
1142 ; CHECK-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
1143 ; CHECK-NEXT: entry:
1144 ; CHECK-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
1145 ; CHECK-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
1146 ; CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1147 ; CHECK-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
1148 ; CHECK-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
1149 ; CHECK-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
1150 ; CHECK-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
1151 ; CHECK-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
1152 ; CHECK-NEXT: call void @__omp_outlined__5(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
1153 ; CHECK-NEXT: ret void
1156 ; CHECK: Function Attrs: convergent noinline nounwind
1157 ; CHECK-LABEL: define {{[^@]+}}@simple_state_machine_interprocedural_after.internalized
1158 ; CHECK-SAME: () #[[ATTR1]] {
1159 ; CHECK-NEXT: entry:
1160 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
1161 ; CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* noundef @[[GLOB2]]) #[[ATTR3]]
1162 ; CHECK-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
1163 ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB2]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__18 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__18_wrapper to i8*), i8** noundef [[TMP1]], i64 noundef 0)
1164 ; CHECK-NEXT: ret void
1167 ; CHECK: Function Attrs: convergent noinline nounwind
1168 ; CHECK-LABEL: define {{[^@]+}}@simple_state_machine_interprocedural_after
1169 ; CHECK-SAME: () #[[ATTR1]] {
1170 ; CHECK-NEXT: entry:
1171 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
1172 ; CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2]])
1173 ; CHECK-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
1174 ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP0]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__18 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__18_wrapper to i8*), i8** [[TMP1]], i64 0)
1175 ; CHECK-NEXT: ret void
1178 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1179 ; CHECK-LABEL: define {{[^@]+}}@__omp_offloading_14_a36502b_simple_state_machine_with_fallback_l55
1180 ; CHECK-SAME: () #[[ATTR0]] {
1181 ; CHECK-NEXT: entry:
1182 ; CHECK-NEXT: [[WORKER_WORK_FN_ADDR:%.*]] = alloca i8*, align 8
1183 ; CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1184 ; CHECK-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
1185 ; CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i1 false, i1 false, i1 true)
1186 ; CHECK-NEXT: [[THREAD_IS_WORKER:%.*]] = icmp ne i32 [[TMP0]], -1
1187 ; CHECK-NEXT: br i1 [[THREAD_IS_WORKER]], label [[WORKER_STATE_MACHINE_BEGIN:%.*]], label [[THREAD_USER_CODE_CHECK:%.*]]
1188 ; CHECK: worker_state_machine.begin:
1189 ; CHECK-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]])
1190 ; CHECK-NEXT: [[WORKER_IS_ACTIVE:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORKER_WORK_FN_ADDR]])
1191 ; CHECK-NEXT: [[WORKER_WORK_FN:%.*]] = load i8*, i8** [[WORKER_WORK_FN_ADDR]], align 8
1192 ; CHECK-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)*
1193 ; CHECK-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null
1194 ; CHECK-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]]
1195 ; CHECK: worker_state_machine.finished:
1196 ; CHECK-NEXT: ret void
1197 ; CHECK: worker_state_machine.is_active.check:
1198 ; CHECK-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]]
1199 ; CHECK: worker_state_machine.parallel_region.check:
1200 ; CHECK-NEXT: [[WORKER_CHECK_PARALLEL_REGION:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], bitcast (i8* @__omp_outlined__7_wrapper.ID to void (i16, i32)*)
1201 ; CHECK-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK1:%.*]]
1202 ; CHECK: worker_state_machine.parallel_region.execute:
1203 ; CHECK-NEXT: call void @__omp_outlined__7_wrapper(i16 0, i32 [[TMP0]])
1204 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]]
1205 ; CHECK: worker_state_machine.parallel_region.check1:
1206 ; CHECK-NEXT: [[WORKER_CHECK_PARALLEL_REGION4:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], bitcast (i8* @__omp_outlined__8_wrapper.ID to void (i16, i32)*)
1207 ; CHECK-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION4]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE2:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_FALLBACK_EXECUTE:%.*]]
1208 ; CHECK: worker_state_machine.parallel_region.execute2:
1209 ; CHECK-NEXT: call void @__omp_outlined__8_wrapper(i16 0, i32 [[TMP0]])
1210 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]]
1211 ; CHECK: worker_state_machine.parallel_region.fallback.execute:
1212 ; CHECK-NEXT: call void [[WORKER_WORK_FN_ADDR_CAST]](i16 0, i32 [[TMP0]])
1213 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]]
1214 ; CHECK: worker_state_machine.parallel_region.end:
1215 ; CHECK-NEXT: call void @__kmpc_kernel_end_parallel()
1216 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_DONE_BARRIER]]
1217 ; CHECK: worker_state_machine.done.barrier:
1218 ; CHECK-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]])
1219 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_BEGIN]]
1220 ; CHECK: thread.user_code.check:
1221 ; CHECK-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
1222 ; CHECK-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
1223 ; CHECK: user_code.entry:
1224 ; CHECK-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR3]]
1225 ; CHECK-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4
1226 ; CHECK-NEXT: call void @__omp_outlined__6(i32* noalias nocapture noundef nonnull readonly align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR3]]
1227 ; CHECK-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true)
1228 ; CHECK-NEXT: ret void
1229 ; CHECK: worker.exit:
1230 ; CHECK-NEXT: ret void
1233 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1234 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__6
1235 ; CHECK-SAME: (i32* noalias nocapture nofree noundef nonnull readonly align 4 dereferenceable(4) [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1236 ; CHECK-NEXT: entry:
1237 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1238 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1239 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
1240 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS1:%.*]] = alloca [0 x i8*], align 8
1241 ; CHECK-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
1242 ; CHECK-NEXT: [[TMP0:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4
1243 ; CHECK-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
1244 ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__7 to i8*), i8* noundef @__omp_outlined__7_wrapper.ID, i8** noundef [[TMP1]], i64 noundef 0)
1245 ; CHECK-NEXT: [[CALL:%.*]] = call i32 @unknown() #[[ATTR9]]
1246 ; CHECK-NEXT: [[TMP2:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS1]] to i8**
1247 ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__8 to i8*), i8* noundef @__omp_outlined__8_wrapper.ID, i8** noundef [[TMP2]], i64 noundef 0)
1248 ; CHECK-NEXT: ret void
1251 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1252 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__7
1253 ; CHECK-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1254 ; CHECK-NEXT: entry:
1255 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1256 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1257 ; CHECK-NEXT: call void @p0() #[[ATTR9]]
1258 ; CHECK-NEXT: ret void
1261 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1262 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__7_wrapper
1263 ; CHECK-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
1264 ; CHECK-NEXT: entry:
1265 ; CHECK-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
1266 ; CHECK-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
1267 ; CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1268 ; CHECK-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
1269 ; CHECK-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
1270 ; CHECK-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
1271 ; CHECK-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
1272 ; CHECK-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
1273 ; CHECK-NEXT: call void @__omp_outlined__7(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
1274 ; CHECK-NEXT: ret void
1277 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1278 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__8
1279 ; CHECK-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1280 ; CHECK-NEXT: entry:
1281 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1282 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1283 ; CHECK-NEXT: call void @p1() #[[ATTR9]]
1284 ; CHECK-NEXT: ret void
1287 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1288 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__8_wrapper
1289 ; CHECK-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
1290 ; CHECK-NEXT: entry:
1291 ; CHECK-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
1292 ; CHECK-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
1293 ; CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1294 ; CHECK-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
1295 ; CHECK-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
1296 ; CHECK-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
1297 ; CHECK-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
1298 ; CHECK-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
1299 ; CHECK-NEXT: call void @__omp_outlined__8(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
1300 ; CHECK-NEXT: ret void
1303 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1304 ; CHECK-LABEL: define {{[^@]+}}@__omp_offloading_14_a36502b_simple_state_machine_no_openmp_attr_l66
1305 ; CHECK-SAME: () #[[ATTR0]] {
1306 ; CHECK-NEXT: entry:
1307 ; CHECK-NEXT: [[WORKER_WORK_FN_ADDR:%.*]] = alloca i8*, align 8
1308 ; CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1309 ; CHECK-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
1310 ; CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i1 false, i1 false, i1 true)
1311 ; CHECK-NEXT: [[THREAD_IS_WORKER:%.*]] = icmp ne i32 [[TMP0]], -1
1312 ; CHECK-NEXT: br i1 [[THREAD_IS_WORKER]], label [[WORKER_STATE_MACHINE_BEGIN:%.*]], label [[THREAD_USER_CODE_CHECK:%.*]]
1313 ; CHECK: worker_state_machine.begin:
1314 ; CHECK-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]])
1315 ; CHECK-NEXT: [[WORKER_IS_ACTIVE:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORKER_WORK_FN_ADDR]])
1316 ; CHECK-NEXT: [[WORKER_WORK_FN:%.*]] = load i8*, i8** [[WORKER_WORK_FN_ADDR]], align 8
1317 ; CHECK-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)*
1318 ; CHECK-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null
1319 ; CHECK-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]]
1320 ; CHECK: worker_state_machine.finished:
1321 ; CHECK-NEXT: ret void
1322 ; CHECK: worker_state_machine.is_active.check:
1323 ; CHECK-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]]
1324 ; CHECK: worker_state_machine.parallel_region.check:
1325 ; CHECK-NEXT: [[WORKER_CHECK_PARALLEL_REGION:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], bitcast (i8* @__omp_outlined__10_wrapper.ID to void (i16, i32)*)
1326 ; CHECK-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK1:%.*]]
1327 ; CHECK: worker_state_machine.parallel_region.execute:
1328 ; CHECK-NEXT: call void @__omp_outlined__10_wrapper(i16 0, i32 [[TMP0]])
1329 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]]
1330 ; CHECK: worker_state_machine.parallel_region.check1:
1331 ; CHECK-NEXT: br i1 true, label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE2:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK3:%.*]]
1332 ; CHECK: worker_state_machine.parallel_region.execute2:
1333 ; CHECK-NEXT: call void @__omp_outlined__11_wrapper(i16 0, i32 [[TMP0]])
1334 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]]
1335 ; CHECK: worker_state_machine.parallel_region.check3:
1336 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]]
1337 ; CHECK: worker_state_machine.parallel_region.end:
1338 ; CHECK-NEXT: call void @__kmpc_kernel_end_parallel()
1339 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_DONE_BARRIER]]
1340 ; CHECK: worker_state_machine.done.barrier:
1341 ; CHECK-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]])
1342 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_BEGIN]]
1343 ; CHECK: thread.user_code.check:
1344 ; CHECK-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
1345 ; CHECK-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
1346 ; CHECK: user_code.entry:
1347 ; CHECK-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR3]]
1348 ; CHECK-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4
1349 ; CHECK-NEXT: call void @__omp_outlined__9(i32* noalias nocapture noundef nonnull readonly align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR3]]
1350 ; CHECK-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true)
1351 ; CHECK-NEXT: ret void
1352 ; CHECK: worker.exit:
1353 ; CHECK-NEXT: ret void
1356 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1357 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__9
1358 ; CHECK-SAME: (i32* noalias nocapture nofree noundef nonnull readonly align 4 dereferenceable(4) [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1359 ; CHECK-NEXT: entry:
1360 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1361 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1362 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
1363 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS1:%.*]] = alloca [0 x i8*], align 8
1364 ; CHECK-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
1365 ; CHECK-NEXT: [[TMP0:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4
1366 ; CHECK-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
1367 ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__10 to i8*), i8* noundef @__omp_outlined__10_wrapper.ID, i8** noundef [[TMP1]], i64 noundef 0)
1368 ; CHECK-NEXT: call void @unknown_no_openmp() #[[ATTR8]]
1369 ; CHECK-NEXT: [[TMP2:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS1]] to i8**
1370 ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__11 to i8*), i8* noundef @__omp_outlined__11_wrapper.ID, i8** noundef [[TMP2]], i64 noundef 0)
1371 ; CHECK-NEXT: ret void
1374 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1375 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__10
1376 ; CHECK-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1377 ; CHECK-NEXT: entry:
1378 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1379 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1380 ; CHECK-NEXT: call void @p0() #[[ATTR9]]
1381 ; CHECK-NEXT: ret void
1384 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1385 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__10_wrapper
1386 ; CHECK-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
1387 ; CHECK-NEXT: entry:
1388 ; CHECK-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
1389 ; CHECK-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
1390 ; CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1391 ; CHECK-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
1392 ; CHECK-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
1393 ; CHECK-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
1394 ; CHECK-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
1395 ; CHECK-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
1396 ; CHECK-NEXT: call void @__omp_outlined__10(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
1397 ; CHECK-NEXT: ret void
1400 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1401 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__11
1402 ; CHECK-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1403 ; CHECK-NEXT: entry:
1404 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1405 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1406 ; CHECK-NEXT: call void @p1() #[[ATTR9]]
1407 ; CHECK-NEXT: ret void
1410 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1411 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__11_wrapper
1412 ; CHECK-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
1413 ; CHECK-NEXT: entry:
1414 ; CHECK-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
1415 ; CHECK-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
1416 ; CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1417 ; CHECK-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
1418 ; CHECK-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
1419 ; CHECK-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
1420 ; CHECK-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
1421 ; CHECK-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
1422 ; CHECK-NEXT: call void @__omp_outlined__11(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
1423 ; CHECK-NEXT: ret void
1426 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1427 ; CHECK-LABEL: define {{[^@]+}}@__omp_offloading_14_a36502b_simple_state_machine_pure_l77
1428 ; CHECK-SAME: () #[[ATTR0]] {
1429 ; CHECK-NEXT: entry:
1430 ; CHECK-NEXT: [[WORKER_WORK_FN_ADDR:%.*]] = alloca i8*, align 8
1431 ; CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1432 ; CHECK-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
1433 ; CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i1 false, i1 false, i1 true)
1434 ; CHECK-NEXT: [[THREAD_IS_WORKER:%.*]] = icmp ne i32 [[TMP0]], -1
1435 ; CHECK-NEXT: br i1 [[THREAD_IS_WORKER]], label [[WORKER_STATE_MACHINE_BEGIN:%.*]], label [[THREAD_USER_CODE_CHECK:%.*]]
1436 ; CHECK: worker_state_machine.begin:
1437 ; CHECK-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]])
1438 ; CHECK-NEXT: [[WORKER_IS_ACTIVE:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORKER_WORK_FN_ADDR]])
1439 ; CHECK-NEXT: [[WORKER_WORK_FN:%.*]] = load i8*, i8** [[WORKER_WORK_FN_ADDR]], align 8
1440 ; CHECK-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)*
1441 ; CHECK-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null
1442 ; CHECK-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]]
1443 ; CHECK: worker_state_machine.finished:
1444 ; CHECK-NEXT: ret void
1445 ; CHECK: worker_state_machine.is_active.check:
1446 ; CHECK-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]]
1447 ; CHECK: worker_state_machine.parallel_region.check:
1448 ; CHECK-NEXT: [[WORKER_CHECK_PARALLEL_REGION:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], bitcast (i8* @__omp_outlined__13_wrapper.ID to void (i16, i32)*)
1449 ; CHECK-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK1:%.*]]
1450 ; CHECK: worker_state_machine.parallel_region.execute:
1451 ; CHECK-NEXT: call void @__omp_outlined__13_wrapper(i16 0, i32 [[TMP0]])
1452 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]]
1453 ; CHECK: worker_state_machine.parallel_region.check1:
1454 ; CHECK-NEXT: br i1 true, label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE2:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK3:%.*]]
1455 ; CHECK: worker_state_machine.parallel_region.execute2:
1456 ; CHECK-NEXT: call void @__omp_outlined__14_wrapper(i16 0, i32 [[TMP0]])
1457 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]]
1458 ; CHECK: worker_state_machine.parallel_region.check3:
1459 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]]
1460 ; CHECK: worker_state_machine.parallel_region.end:
1461 ; CHECK-NEXT: call void @__kmpc_kernel_end_parallel()
1462 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_DONE_BARRIER]]
1463 ; CHECK: worker_state_machine.done.barrier:
1464 ; CHECK-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]])
1465 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_BEGIN]]
1466 ; CHECK: thread.user_code.check:
1467 ; CHECK-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
1468 ; CHECK-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
1469 ; CHECK: user_code.entry:
1470 ; CHECK-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR3]]
1471 ; CHECK-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4
1472 ; CHECK-NEXT: call void @__omp_outlined__12(i32* noalias nocapture noundef nonnull readonly align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR3]]
1473 ; CHECK-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true)
1474 ; CHECK-NEXT: ret void
1475 ; CHECK: worker.exit:
1476 ; CHECK-NEXT: ret void
1479 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1480 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__12
1481 ; CHECK-SAME: (i32* noalias nocapture nofree noundef nonnull readonly align 4 dereferenceable(4) [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1482 ; CHECK-NEXT: entry:
1483 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1484 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1485 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
1486 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS1:%.*]] = alloca [0 x i8*], align 8
1487 ; CHECK-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
1488 ; CHECK-NEXT: call void @unknown_no_openmp() #[[ATTR8]]
1489 ; CHECK-NEXT: [[TMP0:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4
1490 ; CHECK-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
1491 ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__13 to i8*), i8* noundef @__omp_outlined__13_wrapper.ID, i8** noundef [[TMP1]], i64 noundef 0)
1492 ; CHECK-NEXT: [[TMP2:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS1]] to i8**
1493 ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__14 to i8*), i8* noundef @__omp_outlined__14_wrapper.ID, i8** noundef [[TMP2]], i64 noundef 0)
1494 ; CHECK-NEXT: ret void
1497 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1498 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__13
1499 ; CHECK-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1500 ; CHECK-NEXT: entry:
1501 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1502 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1503 ; CHECK-NEXT: call void @p0() #[[ATTR9]]
1504 ; CHECK-NEXT: ret void
1507 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1508 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__13_wrapper
1509 ; CHECK-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
1510 ; CHECK-NEXT: entry:
1511 ; CHECK-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
1512 ; CHECK-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
1513 ; CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1514 ; CHECK-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
1515 ; CHECK-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
1516 ; CHECK-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
1517 ; CHECK-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
1518 ; CHECK-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
1519 ; CHECK-NEXT: call void @__omp_outlined__13(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
1520 ; CHECK-NEXT: ret void
1523 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1524 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__14
1525 ; CHECK-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1526 ; CHECK-NEXT: entry:
1527 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1528 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1529 ; CHECK-NEXT: call void @p1() #[[ATTR9]]
1530 ; CHECK-NEXT: ret void
1533 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1534 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__14_wrapper
1535 ; CHECK-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
1536 ; CHECK-NEXT: entry:
1537 ; CHECK-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
1538 ; CHECK-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
1539 ; CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1540 ; CHECK-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
1541 ; CHECK-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
1542 ; CHECK-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
1543 ; CHECK-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
1544 ; CHECK-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
1545 ; CHECK-NEXT: call void @__omp_outlined__14(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
1546 ; CHECK-NEXT: ret void
1549 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1550 ; CHECK-LABEL: define {{[^@]+}}@__omp_offloading_14_a36502b_simple_state_machine_interprocedural_nested_recursive_l92
1551 ; CHECK-SAME: () #[[ATTR0]] {
1552 ; CHECK-NEXT: entry:
1553 ; CHECK-NEXT: [[WORKER_WORK_FN_ADDR:%.*]] = alloca i8*, align 8
1554 ; CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1555 ; CHECK-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
1556 ; CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i1 false, i1 false, i1 true)
1557 ; CHECK-NEXT: [[THREAD_IS_WORKER:%.*]] = icmp ne i32 [[TMP0]], -1
1558 ; CHECK-NEXT: br i1 [[THREAD_IS_WORKER]], label [[WORKER_STATE_MACHINE_BEGIN:%.*]], label [[THREAD_USER_CODE_CHECK:%.*]]
1559 ; CHECK: worker_state_machine.begin:
1560 ; CHECK-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]])
1561 ; CHECK-NEXT: [[WORKER_IS_ACTIVE:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORKER_WORK_FN_ADDR]])
1562 ; CHECK-NEXT: [[WORKER_WORK_FN:%.*]] = load i8*, i8** [[WORKER_WORK_FN_ADDR]], align 8
1563 ; CHECK-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)*
1564 ; CHECK-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null
1565 ; CHECK-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]]
1566 ; CHECK: worker_state_machine.finished:
1567 ; CHECK-NEXT: ret void
1568 ; CHECK: worker_state_machine.is_active.check:
1569 ; CHECK-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]]
1570 ; CHECK: worker_state_machine.parallel_region.check:
1571 ; CHECK-NEXT: [[WORKER_CHECK_PARALLEL_REGION:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], @__omp_outlined__19_wrapper
1572 ; CHECK-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_FALLBACK_EXECUTE:%.*]]
1573 ; CHECK: worker_state_machine.parallel_region.execute:
1574 ; CHECK-NEXT: call void @__omp_outlined__19_wrapper(i16 0, i32 [[TMP0]])
1575 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]]
1576 ; CHECK: worker_state_machine.parallel_region.fallback.execute:
1577 ; CHECK-NEXT: call void [[WORKER_WORK_FN_ADDR_CAST]](i16 0, i32 [[TMP0]])
1578 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]]
1579 ; CHECK: worker_state_machine.parallel_region.end:
1580 ; CHECK-NEXT: call void @__kmpc_kernel_end_parallel()
1581 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_DONE_BARRIER]]
1582 ; CHECK: worker_state_machine.done.barrier:
1583 ; CHECK-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]])
1584 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_BEGIN]]
1585 ; CHECK: thread.user_code.check:
1586 ; CHECK-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
1587 ; CHECK-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
1588 ; CHECK: user_code.entry:
1589 ; CHECK-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR3]]
1590 ; CHECK-NEXT: call void @__omp_outlined__15(i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR3]]
1591 ; CHECK-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true)
1592 ; CHECK-NEXT: ret void
1593 ; CHECK: worker.exit:
1594 ; CHECK-NEXT: ret void
1597 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1598 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__15
1599 ; CHECK-SAME: (i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1600 ; CHECK-NEXT: entry:
1601 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1602 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1603 ; CHECK-NEXT: [[CALL:%.*]] = call i32 bitcast (i32 (...)* @omp_get_thread_num to i32 ()*)() #[[ATTR9]]
1604 ; CHECK-NEXT: call void @simple_state_machine_interprocedural_nested_recursive_after.internalized(i32 [[CALL]]) #[[ATTR7]]
1605 ; CHECK-NEXT: ret void
1608 ; CHECK: Function Attrs: convergent noinline nounwind
1609 ; CHECK-LABEL: define {{[^@]+}}@simple_state_machine_interprocedural_nested_recursive_after.internalized
1610 ; CHECK-SAME: (i32 [[A:%.*]]) #[[ATTR1]] {
1611 ; CHECK-NEXT: entry:
1612 ; CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4
1613 ; CHECK-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4
1614 ; CHECK-NEXT: [[CMP:%.*]] = icmp eq i32 [[A]], 0
1615 ; CHECK-NEXT: br i1 [[CMP]], label [[IF_THEN:%.*]], label [[IF_END:%.*]]
1617 ; CHECK-NEXT: br label [[RETURN:%.*]]
1619 ; CHECK-NEXT: [[SUB:%.*]] = sub nsw i32 [[A]], 1
1620 ; CHECK-NEXT: call void @simple_state_machine_interprocedural_nested_recursive_after.internalized(i32 [[SUB]]) #[[ATTR7]]
1621 ; CHECK-NEXT: call void @simple_state_machine_interprocedural_nested_recursive_after_after.internalized() #[[ATTR7]]
1622 ; CHECK-NEXT: br label [[RETURN]]
1624 ; CHECK-NEXT: ret void
1627 ; CHECK: Function Attrs: convergent noinline nounwind
1628 ; CHECK-LABEL: define {{[^@]+}}@simple_state_machine_interprocedural_nested_recursive_after
1629 ; CHECK-SAME: (i32 [[A:%.*]]) #[[ATTR1]] {
1630 ; CHECK-NEXT: entry:
1631 ; CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4
1632 ; CHECK-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4
1633 ; CHECK-NEXT: [[TMP0:%.*]] = load i32, i32* [[A_ADDR]], align 4
1634 ; CHECK-NEXT: [[CMP:%.*]] = icmp eq i32 [[TMP0]], 0
1635 ; CHECK-NEXT: br i1 [[CMP]], label [[IF_THEN:%.*]], label [[IF_END:%.*]]
1637 ; CHECK-NEXT: br label [[RETURN:%.*]]
1639 ; CHECK-NEXT: [[TMP1:%.*]] = load i32, i32* [[A_ADDR]], align 4
1640 ; CHECK-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP1]], 1
1641 ; CHECK-NEXT: call void @simple_state_machine_interprocedural_nested_recursive_after(i32 [[SUB]]) #[[ATTR9]]
1642 ; CHECK-NEXT: call void @simple_state_machine_interprocedural_nested_recursive_after_after() #[[ATTR9]]
1643 ; CHECK-NEXT: br label [[RETURN]]
1645 ; CHECK-NEXT: ret void
1648 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1649 ; CHECK-LABEL: define {{[^@]+}}@__omp_offloading_14_a36502b_no_state_machine_weak_callee_l112
1650 ; CHECK-SAME: () #[[ATTR0]] {
1651 ; CHECK-NEXT: entry:
1652 ; CHECK-NEXT: [[WORKER_WORK_FN_ADDR:%.*]] = alloca i8*, align 8
1653 ; CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1654 ; CHECK-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
1655 ; CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i1 false, i1 false, i1 true)
1656 ; CHECK-NEXT: [[THREAD_IS_WORKER:%.*]] = icmp ne i32 [[TMP0]], -1
1657 ; CHECK-NEXT: br i1 [[THREAD_IS_WORKER]], label [[WORKER_STATE_MACHINE_BEGIN:%.*]], label [[THREAD_USER_CODE_CHECK:%.*]]
1658 ; CHECK: worker_state_machine.begin:
1659 ; CHECK-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]])
1660 ; CHECK-NEXT: [[WORKER_IS_ACTIVE:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORKER_WORK_FN_ADDR]])
1661 ; CHECK-NEXT: [[WORKER_WORK_FN:%.*]] = load i8*, i8** [[WORKER_WORK_FN_ADDR]], align 8
1662 ; CHECK-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)*
1663 ; CHECK-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null
1664 ; CHECK-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]]
1665 ; CHECK: worker_state_machine.finished:
1666 ; CHECK-NEXT: ret void
1667 ; CHECK: worker_state_machine.is_active.check:
1668 ; CHECK-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_FALLBACK_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]]
1669 ; CHECK: worker_state_machine.parallel_region.fallback.execute:
1670 ; CHECK-NEXT: call void [[WORKER_WORK_FN_ADDR_CAST]](i16 0, i32 [[TMP0]])
1671 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]]
1672 ; CHECK: worker_state_machine.parallel_region.end:
1673 ; CHECK-NEXT: call void @__kmpc_kernel_end_parallel()
1674 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_DONE_BARRIER]]
1675 ; CHECK: worker_state_machine.done.barrier:
1676 ; CHECK-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]])
1677 ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_BEGIN]]
1678 ; CHECK: thread.user_code.check:
1679 ; CHECK-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
1680 ; CHECK-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
1681 ; CHECK: user_code.entry:
1682 ; CHECK-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR3]]
1683 ; CHECK-NEXT: call void @__omp_outlined__16(i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR3]]
1684 ; CHECK-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true)
1685 ; CHECK-NEXT: ret void
1686 ; CHECK: worker.exit:
1687 ; CHECK-NEXT: ret void
1690 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1691 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__16
1692 ; CHECK-SAME: (i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1693 ; CHECK-NEXT: entry:
1694 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1695 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1696 ; CHECK-NEXT: call void @weak_callee_empty() #[[ATTR7]]
1697 ; CHECK-NEXT: ret void
1700 ; CHECK: Function Attrs: convergent noinline nounwind
1701 ; CHECK-LABEL: define {{[^@]+}}@weak_callee_empty
1702 ; CHECK-SAME: () #[[ATTR1]] {
1703 ; CHECK-NEXT: entry:
1704 ; CHECK-NEXT: ret void
1707 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1708 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__17
1709 ; CHECK-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1710 ; CHECK-NEXT: entry:
1711 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1712 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1713 ; CHECK-NEXT: call void @p0() #[[ATTR9]]
1714 ; CHECK-NEXT: ret void
1717 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1718 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__17_wrapper
1719 ; CHECK-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
1720 ; CHECK-NEXT: entry:
1721 ; CHECK-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
1722 ; CHECK-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
1723 ; CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1724 ; CHECK-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
1725 ; CHECK-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
1726 ; CHECK-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
1727 ; CHECK-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
1728 ; CHECK-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
1729 ; CHECK-NEXT: call void @__omp_outlined__17(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
1730 ; CHECK-NEXT: ret void
1733 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1734 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__18
1735 ; CHECK-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1736 ; CHECK-NEXT: entry:
1737 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1738 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1739 ; CHECK-NEXT: call void @p0() #[[ATTR9]]
1740 ; CHECK-NEXT: ret void
1743 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1744 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__18_wrapper
1745 ; CHECK-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
1746 ; CHECK-NEXT: entry:
1747 ; CHECK-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
1748 ; CHECK-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
1749 ; CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1750 ; CHECK-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
1751 ; CHECK-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
1752 ; CHECK-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
1753 ; CHECK-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
1754 ; CHECK-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
1755 ; CHECK-NEXT: call void @__omp_outlined__18(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
1756 ; CHECK-NEXT: ret void
1759 ; CHECK: Function Attrs: convergent noinline nounwind
1760 ; CHECK-LABEL: define {{[^@]+}}@simple_state_machine_interprocedural_nested_recursive_after_after.internalized
1761 ; CHECK-SAME: () #[[ATTR1]] {
1762 ; CHECK-NEXT: entry:
1763 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
1764 ; CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* noundef @[[GLOB2]]) #[[ATTR3]]
1765 ; CHECK-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
1766 ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB2]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__19 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__19_wrapper to i8*), i8** noundef [[TMP1]], i64 noundef 0)
1767 ; CHECK-NEXT: ret void
1770 ; CHECK: Function Attrs: convergent noinline nounwind
1771 ; CHECK-LABEL: define {{[^@]+}}@simple_state_machine_interprocedural_nested_recursive_after_after
1772 ; CHECK-SAME: () #[[ATTR1]] {
1773 ; CHECK-NEXT: entry:
1774 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
1775 ; CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2]])
1776 ; CHECK-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
1777 ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP0]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__19 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__19_wrapper to i8*), i8** [[TMP1]], i64 0)
1778 ; CHECK-NEXT: ret void
1781 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1782 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__19
1783 ; CHECK-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1784 ; CHECK-NEXT: entry:
1785 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1786 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1787 ; CHECK-NEXT: call void @p0() #[[ATTR9]]
1788 ; CHECK-NEXT: ret void
1791 ; CHECK: Function Attrs: convergent noinline norecurse nounwind
1792 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__19_wrapper
1793 ; CHECK-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
1794 ; CHECK-NEXT: entry:
1795 ; CHECK-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
1796 ; CHECK-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
1797 ; CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1798 ; CHECK-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
1799 ; CHECK-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
1800 ; CHECK-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
1801 ; CHECK-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
1802 ; CHECK-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
1803 ; CHECK-NEXT: call void @__omp_outlined__19(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
1804 ; CHECK-NEXT: ret void
1807 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
1808 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_offloading_14_a36502b_no_state_machine_needed_l14
1809 ; CHECK-DISABLED-SAME: () #[[ATTR0:[0-9]+]] {
1810 ; CHECK-DISABLED-NEXT: entry:
1811 ; CHECK-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1812 ; CHECK-DISABLED-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
1813 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i1 false, i1 true, i1 true)
1814 ; CHECK-DISABLED-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
1815 ; CHECK-DISABLED-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
1816 ; CHECK-DISABLED: user_code.entry:
1817 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR3:[0-9]+]]
1818 ; CHECK-DISABLED-NEXT: call void @__omp_outlined__(i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR3]]
1819 ; CHECK-DISABLED-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true)
1820 ; CHECK-DISABLED-NEXT: ret void
1821 ; CHECK-DISABLED: worker.exit:
1822 ; CHECK-DISABLED-NEXT: ret void
1825 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
1826 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__
1827 ; CHECK-DISABLED-SAME: (i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1828 ; CHECK-DISABLED-NEXT: entry:
1829 ; CHECK-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1830 ; CHECK-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1831 ; CHECK-DISABLED-NEXT: call void @no_parallel_region_in_here.internalized() #[[ATTR7:[0-9]+]]
1832 ; CHECK-DISABLED-NEXT: call void @unknown_no_openmp() #[[ATTR8:[0-9]+]]
1833 ; CHECK-DISABLED-NEXT: ret void
1836 ; CHECK-DISABLED: Function Attrs: convergent noinline nounwind
1837 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@no_parallel_region_in_here.internalized
1838 ; CHECK-DISABLED-SAME: () #[[ATTR1:[0-9]+]] {
1839 ; CHECK-DISABLED-NEXT: entry:
1840 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* noundef @[[GLOB2:[0-9]+]]) #[[ATTR3]]
1841 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_single(%struct.ident_t* noundef @[[GLOB2]], i32 [[TMP0]]) #[[ATTR3]]
1842 ; CHECK-DISABLED-NEXT: [[TMP2:%.*]] = icmp ne i32 [[TMP1]], 0
1843 ; CHECK-DISABLED-NEXT: br i1 [[TMP2]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]]
1844 ; CHECK-DISABLED: omp_if.then:
1845 ; CHECK-DISABLED-NEXT: store i32 0, i32* @G, align 4
1846 ; CHECK-DISABLED-NEXT: call void @__kmpc_end_single(%struct.ident_t* noundef @[[GLOB2]], i32 [[TMP0]]) #[[ATTR3]]
1847 ; CHECK-DISABLED-NEXT: br label [[OMP_IF_END]]
1848 ; CHECK-DISABLED: omp_if.end:
1849 ; CHECK-DISABLED-NEXT: call void @__kmpc_barrier(%struct.ident_t* noundef @[[GLOB3:[0-9]+]], i32 [[TMP0]]) #[[ATTR3]]
1850 ; CHECK-DISABLED-NEXT: ret void
1853 ; CHECK-DISABLED: Function Attrs: convergent noinline nounwind
1854 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@no_parallel_region_in_here
1855 ; CHECK-DISABLED-SAME: () #[[ATTR1]] {
1856 ; CHECK-DISABLED-NEXT: entry:
1857 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2]])
1858 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_single(%struct.ident_t* @[[GLOB2]], i32 [[TMP0]])
1859 ; CHECK-DISABLED-NEXT: [[TMP2:%.*]] = icmp ne i32 [[TMP1]], 0
1860 ; CHECK-DISABLED-NEXT: br i1 [[TMP2]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]]
1861 ; CHECK-DISABLED: omp_if.then:
1862 ; CHECK-DISABLED-NEXT: store i32 0, i32* @G, align 4
1863 ; CHECK-DISABLED-NEXT: call void @__kmpc_end_single(%struct.ident_t* @[[GLOB2]], i32 [[TMP0]])
1864 ; CHECK-DISABLED-NEXT: br label [[OMP_IF_END]]
1865 ; CHECK-DISABLED: omp_if.end:
1866 ; CHECK-DISABLED-NEXT: call void @__kmpc_barrier(%struct.ident_t* @[[GLOB3]], i32 [[TMP0]])
1867 ; CHECK-DISABLED-NEXT: ret void
1870 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
1871 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_offloading_14_a36502b_simple_state_machine_l22
1872 ; CHECK-DISABLED-SAME: () #[[ATTR0]] {
1873 ; CHECK-DISABLED-NEXT: entry:
1874 ; CHECK-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1875 ; CHECK-DISABLED-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
1876 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i1 false, i1 true, i1 true)
1877 ; CHECK-DISABLED-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
1878 ; CHECK-DISABLED-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
1879 ; CHECK-DISABLED: user_code.entry:
1880 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR3]]
1881 ; CHECK-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4
1882 ; CHECK-DISABLED-NEXT: call void @__omp_outlined__1(i32* noalias nocapture noundef nonnull readonly align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR3]]
1883 ; CHECK-DISABLED-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true)
1884 ; CHECK-DISABLED-NEXT: ret void
1885 ; CHECK-DISABLED: worker.exit:
1886 ; CHECK-DISABLED-NEXT: ret void
1889 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
1890 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__1
1891 ; CHECK-DISABLED-SAME: (i32* noalias nocapture nofree noundef nonnull readonly align 4 dereferenceable(4) [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1892 ; CHECK-DISABLED-NEXT: entry:
1893 ; CHECK-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1894 ; CHECK-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1895 ; CHECK-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
1896 ; CHECK-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS1:%.*]] = alloca [0 x i8*], align 8
1897 ; CHECK-DISABLED-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
1898 ; CHECK-DISABLED-NEXT: call void @unknown_no_openmp() #[[ATTR8]]
1899 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4
1900 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
1901 ; CHECK-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__2 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__2_wrapper to i8*), i8** noundef [[TMP1]], i64 noundef 0)
1902 ; CHECK-DISABLED-NEXT: call void @no_parallel_region_in_here.internalized() #[[ATTR7]]
1903 ; CHECK-DISABLED-NEXT: [[TMP2:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS1]] to i8**
1904 ; CHECK-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__3 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__3_wrapper to i8*), i8** noundef [[TMP2]], i64 noundef 0)
1905 ; CHECK-DISABLED-NEXT: ret void
1908 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
1909 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__2
1910 ; CHECK-DISABLED-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1911 ; CHECK-DISABLED-NEXT: entry:
1912 ; CHECK-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1913 ; CHECK-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1914 ; CHECK-DISABLED-NEXT: call void @p0() #[[ATTR9:[0-9]+]]
1915 ; CHECK-DISABLED-NEXT: ret void
1918 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
1919 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__2_wrapper
1920 ; CHECK-DISABLED-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
1921 ; CHECK-DISABLED-NEXT: entry:
1922 ; CHECK-DISABLED-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
1923 ; CHECK-DISABLED-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
1924 ; CHECK-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1925 ; CHECK-DISABLED-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
1926 ; CHECK-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
1927 ; CHECK-DISABLED-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
1928 ; CHECK-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
1929 ; CHECK-DISABLED-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
1930 ; CHECK-DISABLED-NEXT: call void @__omp_outlined__2(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
1931 ; CHECK-DISABLED-NEXT: ret void
1934 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
1935 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__3
1936 ; CHECK-DISABLED-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1937 ; CHECK-DISABLED-NEXT: entry:
1938 ; CHECK-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1939 ; CHECK-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1940 ; CHECK-DISABLED-NEXT: call void @p1() #[[ATTR9]]
1941 ; CHECK-DISABLED-NEXT: ret void
1944 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
1945 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__3_wrapper
1946 ; CHECK-DISABLED-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
1947 ; CHECK-DISABLED-NEXT: entry:
1948 ; CHECK-DISABLED-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
1949 ; CHECK-DISABLED-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
1950 ; CHECK-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1951 ; CHECK-DISABLED-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
1952 ; CHECK-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
1953 ; CHECK-DISABLED-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
1954 ; CHECK-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
1955 ; CHECK-DISABLED-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
1956 ; CHECK-DISABLED-NEXT: call void @__omp_outlined__3(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
1957 ; CHECK-DISABLED-NEXT: ret void
1960 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
1961 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_offloading_14_a36502b_simple_state_machine_interprocedural_l39
1962 ; CHECK-DISABLED-SAME: () #[[ATTR0]] {
1963 ; CHECK-DISABLED-NEXT: entry:
1964 ; CHECK-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
1965 ; CHECK-DISABLED-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
1966 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i1 false, i1 true, i1 true)
1967 ; CHECK-DISABLED-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
1968 ; CHECK-DISABLED-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
1969 ; CHECK-DISABLED: user_code.entry:
1970 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR3]]
1971 ; CHECK-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4
1972 ; CHECK-DISABLED-NEXT: call void @__omp_outlined__4(i32* noalias nocapture noundef nonnull readonly align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR3]]
1973 ; CHECK-DISABLED-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true)
1974 ; CHECK-DISABLED-NEXT: ret void
1975 ; CHECK-DISABLED: worker.exit:
1976 ; CHECK-DISABLED-NEXT: ret void
1979 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
1980 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__4
1981 ; CHECK-DISABLED-SAME: (i32* noalias nocapture nofree noundef nonnull readonly align 4 dereferenceable(4) [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
1982 ; CHECK-DISABLED-NEXT: entry:
1983 ; CHECK-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
1984 ; CHECK-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
1985 ; CHECK-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
1986 ; CHECK-DISABLED-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
1987 ; CHECK-DISABLED-NEXT: call void @unknown_no_openmp() #[[ATTR8]]
1988 ; CHECK-DISABLED-NEXT: call void @simple_state_machine_interprocedural_before.internalized() #[[ATTR7]]
1989 ; CHECK-DISABLED-NEXT: call void @no_parallel_region_in_here.internalized() #[[ATTR7]]
1990 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4
1991 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
1992 ; CHECK-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__5 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__5_wrapper to i8*), i8** noundef [[TMP1]], i64 noundef 0)
1993 ; CHECK-DISABLED-NEXT: call void @simple_state_machine_interprocedural_after.internalized() #[[ATTR7]]
1994 ; CHECK-DISABLED-NEXT: ret void
1997 ; CHECK-DISABLED: Function Attrs: convergent noinline nounwind
1998 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@simple_state_machine_interprocedural_before.internalized
1999 ; CHECK-DISABLED-SAME: () #[[ATTR1]] {
2000 ; CHECK-DISABLED-NEXT: entry:
2001 ; CHECK-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
2002 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* noundef @[[GLOB2]]) #[[ATTR3]]
2003 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
2004 ; CHECK-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB2]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__17 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__17_wrapper to i8*), i8** noundef [[TMP1]], i64 noundef 0)
2005 ; CHECK-DISABLED-NEXT: ret void
2008 ; CHECK-DISABLED: Function Attrs: convergent noinline nounwind
2009 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@simple_state_machine_interprocedural_before
2010 ; CHECK-DISABLED-SAME: () #[[ATTR1]] {
2011 ; CHECK-DISABLED-NEXT: entry:
2012 ; CHECK-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
2013 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2]])
2014 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
2015 ; CHECK-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP0]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__17 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__17_wrapper to i8*), i8** [[TMP1]], i64 0)
2016 ; CHECK-DISABLED-NEXT: ret void
2019 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2020 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__5
2021 ; CHECK-DISABLED-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
2022 ; CHECK-DISABLED-NEXT: entry:
2023 ; CHECK-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
2024 ; CHECK-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
2025 ; CHECK-DISABLED-NEXT: call void @p1() #[[ATTR9]]
2026 ; CHECK-DISABLED-NEXT: ret void
2029 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2030 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__5_wrapper
2031 ; CHECK-DISABLED-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
2032 ; CHECK-DISABLED-NEXT: entry:
2033 ; CHECK-DISABLED-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
2034 ; CHECK-DISABLED-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
2035 ; CHECK-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
2036 ; CHECK-DISABLED-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
2037 ; CHECK-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
2038 ; CHECK-DISABLED-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
2039 ; CHECK-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
2040 ; CHECK-DISABLED-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
2041 ; CHECK-DISABLED-NEXT: call void @__omp_outlined__5(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
2042 ; CHECK-DISABLED-NEXT: ret void
2045 ; CHECK-DISABLED: Function Attrs: convergent noinline nounwind
2046 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@simple_state_machine_interprocedural_after.internalized
2047 ; CHECK-DISABLED-SAME: () #[[ATTR1]] {
2048 ; CHECK-DISABLED-NEXT: entry:
2049 ; CHECK-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
2050 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* noundef @[[GLOB2]]) #[[ATTR3]]
2051 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
2052 ; CHECK-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB2]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__18 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__18_wrapper to i8*), i8** noundef [[TMP1]], i64 noundef 0)
2053 ; CHECK-DISABLED-NEXT: ret void
2056 ; CHECK-DISABLED: Function Attrs: convergent noinline nounwind
2057 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@simple_state_machine_interprocedural_after
2058 ; CHECK-DISABLED-SAME: () #[[ATTR1]] {
2059 ; CHECK-DISABLED-NEXT: entry:
2060 ; CHECK-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
2061 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2]])
2062 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
2063 ; CHECK-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP0]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__18 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__18_wrapper to i8*), i8** [[TMP1]], i64 0)
2064 ; CHECK-DISABLED-NEXT: ret void
2067 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2068 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_offloading_14_a36502b_simple_state_machine_with_fallback_l55
2069 ; CHECK-DISABLED-SAME: () #[[ATTR0]] {
2070 ; CHECK-DISABLED-NEXT: entry:
2071 ; CHECK-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
2072 ; CHECK-DISABLED-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
2073 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i1 false, i1 true, i1 true)
2074 ; CHECK-DISABLED-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
2075 ; CHECK-DISABLED-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
2076 ; CHECK-DISABLED: user_code.entry:
2077 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR3]]
2078 ; CHECK-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4
2079 ; CHECK-DISABLED-NEXT: call void @__omp_outlined__6(i32* noalias nocapture noundef nonnull readonly align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR3]]
2080 ; CHECK-DISABLED-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true)
2081 ; CHECK-DISABLED-NEXT: ret void
2082 ; CHECK-DISABLED: worker.exit:
2083 ; CHECK-DISABLED-NEXT: ret void
2086 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2087 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__6
2088 ; CHECK-DISABLED-SAME: (i32* noalias nocapture nofree noundef nonnull readonly align 4 dereferenceable(4) [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
2089 ; CHECK-DISABLED-NEXT: entry:
2090 ; CHECK-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
2091 ; CHECK-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
2092 ; CHECK-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
2093 ; CHECK-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS1:%.*]] = alloca [0 x i8*], align 8
2094 ; CHECK-DISABLED-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
2095 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4
2096 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
2097 ; CHECK-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__7 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__7_wrapper to i8*), i8** noundef [[TMP1]], i64 noundef 0)
2098 ; CHECK-DISABLED-NEXT: [[CALL:%.*]] = call i32 @unknown() #[[ATTR9]]
2099 ; CHECK-DISABLED-NEXT: [[TMP2:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS1]] to i8**
2100 ; CHECK-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__8 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__8_wrapper to i8*), i8** noundef [[TMP2]], i64 noundef 0)
2101 ; CHECK-DISABLED-NEXT: ret void
2104 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2105 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__7
2106 ; CHECK-DISABLED-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
2107 ; CHECK-DISABLED-NEXT: entry:
2108 ; CHECK-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
2109 ; CHECK-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
2110 ; CHECK-DISABLED-NEXT: call void @p0() #[[ATTR9]]
2111 ; CHECK-DISABLED-NEXT: ret void
2114 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2115 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__7_wrapper
2116 ; CHECK-DISABLED-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
2117 ; CHECK-DISABLED-NEXT: entry:
2118 ; CHECK-DISABLED-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
2119 ; CHECK-DISABLED-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
2120 ; CHECK-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
2121 ; CHECK-DISABLED-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
2122 ; CHECK-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
2123 ; CHECK-DISABLED-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
2124 ; CHECK-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
2125 ; CHECK-DISABLED-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
2126 ; CHECK-DISABLED-NEXT: call void @__omp_outlined__7(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
2127 ; CHECK-DISABLED-NEXT: ret void
2130 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2131 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__8
2132 ; CHECK-DISABLED-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
2133 ; CHECK-DISABLED-NEXT: entry:
2134 ; CHECK-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
2135 ; CHECK-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
2136 ; CHECK-DISABLED-NEXT: call void @p1() #[[ATTR9]]
2137 ; CHECK-DISABLED-NEXT: ret void
2140 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2141 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__8_wrapper
2142 ; CHECK-DISABLED-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
2143 ; CHECK-DISABLED-NEXT: entry:
2144 ; CHECK-DISABLED-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
2145 ; CHECK-DISABLED-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
2146 ; CHECK-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
2147 ; CHECK-DISABLED-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
2148 ; CHECK-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
2149 ; CHECK-DISABLED-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
2150 ; CHECK-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
2151 ; CHECK-DISABLED-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
2152 ; CHECK-DISABLED-NEXT: call void @__omp_outlined__8(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
2153 ; CHECK-DISABLED-NEXT: ret void
2156 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2157 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_offloading_14_a36502b_simple_state_machine_no_openmp_attr_l66
2158 ; CHECK-DISABLED-SAME: () #[[ATTR0]] {
2159 ; CHECK-DISABLED-NEXT: entry:
2160 ; CHECK-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
2161 ; CHECK-DISABLED-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
2162 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i1 false, i1 true, i1 true)
2163 ; CHECK-DISABLED-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
2164 ; CHECK-DISABLED-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
2165 ; CHECK-DISABLED: user_code.entry:
2166 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR3]]
2167 ; CHECK-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4
2168 ; CHECK-DISABLED-NEXT: call void @__omp_outlined__9(i32* noalias nocapture noundef nonnull readonly align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR3]]
2169 ; CHECK-DISABLED-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true)
2170 ; CHECK-DISABLED-NEXT: ret void
2171 ; CHECK-DISABLED: worker.exit:
2172 ; CHECK-DISABLED-NEXT: ret void
2175 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2176 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__9
2177 ; CHECK-DISABLED-SAME: (i32* noalias nocapture nofree noundef nonnull readonly align 4 dereferenceable(4) [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
2178 ; CHECK-DISABLED-NEXT: entry:
2179 ; CHECK-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
2180 ; CHECK-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
2181 ; CHECK-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
2182 ; CHECK-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS1:%.*]] = alloca [0 x i8*], align 8
2183 ; CHECK-DISABLED-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
2184 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4
2185 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
2186 ; CHECK-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__10 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__10_wrapper to i8*), i8** noundef [[TMP1]], i64 noundef 0)
2187 ; CHECK-DISABLED-NEXT: call void @unknown_no_openmp() #[[ATTR8]]
2188 ; CHECK-DISABLED-NEXT: [[TMP2:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS1]] to i8**
2189 ; CHECK-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__11 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__11_wrapper to i8*), i8** noundef [[TMP2]], i64 noundef 0)
2190 ; CHECK-DISABLED-NEXT: ret void
2193 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2194 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__10
2195 ; CHECK-DISABLED-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
2196 ; CHECK-DISABLED-NEXT: entry:
2197 ; CHECK-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
2198 ; CHECK-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
2199 ; CHECK-DISABLED-NEXT: call void @p0() #[[ATTR9]]
2200 ; CHECK-DISABLED-NEXT: ret void
2203 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2204 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__10_wrapper
2205 ; CHECK-DISABLED-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
2206 ; CHECK-DISABLED-NEXT: entry:
2207 ; CHECK-DISABLED-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
2208 ; CHECK-DISABLED-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
2209 ; CHECK-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
2210 ; CHECK-DISABLED-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
2211 ; CHECK-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
2212 ; CHECK-DISABLED-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
2213 ; CHECK-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
2214 ; CHECK-DISABLED-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
2215 ; CHECK-DISABLED-NEXT: call void @__omp_outlined__10(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
2216 ; CHECK-DISABLED-NEXT: ret void
2219 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2220 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__11
2221 ; CHECK-DISABLED-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
2222 ; CHECK-DISABLED-NEXT: entry:
2223 ; CHECK-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
2224 ; CHECK-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
2225 ; CHECK-DISABLED-NEXT: call void @p1() #[[ATTR9]]
2226 ; CHECK-DISABLED-NEXT: ret void
2229 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2230 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__11_wrapper
2231 ; CHECK-DISABLED-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
2232 ; CHECK-DISABLED-NEXT: entry:
2233 ; CHECK-DISABLED-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
2234 ; CHECK-DISABLED-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
2235 ; CHECK-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
2236 ; CHECK-DISABLED-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
2237 ; CHECK-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
2238 ; CHECK-DISABLED-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
2239 ; CHECK-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
2240 ; CHECK-DISABLED-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
2241 ; CHECK-DISABLED-NEXT: call void @__omp_outlined__11(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
2242 ; CHECK-DISABLED-NEXT: ret void
2245 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2246 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_offloading_14_a36502b_simple_state_machine_pure_l77
2247 ; CHECK-DISABLED-SAME: () #[[ATTR0]] {
2248 ; CHECK-DISABLED-NEXT: entry:
2249 ; CHECK-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
2250 ; CHECK-DISABLED-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
2251 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i1 false, i1 true, i1 true)
2252 ; CHECK-DISABLED-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
2253 ; CHECK-DISABLED-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
2254 ; CHECK-DISABLED: user_code.entry:
2255 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR3]]
2256 ; CHECK-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4
2257 ; CHECK-DISABLED-NEXT: call void @__omp_outlined__12(i32* noalias nocapture noundef nonnull readonly align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR3]]
2258 ; CHECK-DISABLED-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true)
2259 ; CHECK-DISABLED-NEXT: ret void
2260 ; CHECK-DISABLED: worker.exit:
2261 ; CHECK-DISABLED-NEXT: ret void
2264 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2265 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__12
2266 ; CHECK-DISABLED-SAME: (i32* noalias nocapture nofree noundef nonnull readonly align 4 dereferenceable(4) [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
2267 ; CHECK-DISABLED-NEXT: entry:
2268 ; CHECK-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
2269 ; CHECK-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
2270 ; CHECK-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
2271 ; CHECK-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS1:%.*]] = alloca [0 x i8*], align 8
2272 ; CHECK-DISABLED-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
2273 ; CHECK-DISABLED-NEXT: call void @unknown_no_openmp() #[[ATTR8]]
2274 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4
2275 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
2276 ; CHECK-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__13 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__13_wrapper to i8*), i8** noundef [[TMP1]], i64 noundef 0)
2277 ; CHECK-DISABLED-NEXT: [[TMP2:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS1]] to i8**
2278 ; CHECK-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__14 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__14_wrapper to i8*), i8** noundef [[TMP2]], i64 noundef 0)
2279 ; CHECK-DISABLED-NEXT: ret void
2282 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2283 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__13
2284 ; CHECK-DISABLED-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
2285 ; CHECK-DISABLED-NEXT: entry:
2286 ; CHECK-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
2287 ; CHECK-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
2288 ; CHECK-DISABLED-NEXT: call void @p0() #[[ATTR9]]
2289 ; CHECK-DISABLED-NEXT: ret void
2292 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2293 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__13_wrapper
2294 ; CHECK-DISABLED-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
2295 ; CHECK-DISABLED-NEXT: entry:
2296 ; CHECK-DISABLED-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
2297 ; CHECK-DISABLED-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
2298 ; CHECK-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
2299 ; CHECK-DISABLED-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
2300 ; CHECK-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
2301 ; CHECK-DISABLED-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
2302 ; CHECK-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
2303 ; CHECK-DISABLED-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
2304 ; CHECK-DISABLED-NEXT: call void @__omp_outlined__13(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
2305 ; CHECK-DISABLED-NEXT: ret void
2308 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2309 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__14
2310 ; CHECK-DISABLED-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
2311 ; CHECK-DISABLED-NEXT: entry:
2312 ; CHECK-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
2313 ; CHECK-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
2314 ; CHECK-DISABLED-NEXT: call void @p1() #[[ATTR9]]
2315 ; CHECK-DISABLED-NEXT: ret void
2318 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2319 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__14_wrapper
2320 ; CHECK-DISABLED-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
2321 ; CHECK-DISABLED-NEXT: entry:
2322 ; CHECK-DISABLED-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
2323 ; CHECK-DISABLED-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
2324 ; CHECK-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
2325 ; CHECK-DISABLED-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
2326 ; CHECK-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
2327 ; CHECK-DISABLED-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
2328 ; CHECK-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
2329 ; CHECK-DISABLED-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
2330 ; CHECK-DISABLED-NEXT: call void @__omp_outlined__14(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
2331 ; CHECK-DISABLED-NEXT: ret void
2334 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2335 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_offloading_14_a36502b_simple_state_machine_interprocedural_nested_recursive_l92
2336 ; CHECK-DISABLED-SAME: () #[[ATTR0]] {
2337 ; CHECK-DISABLED-NEXT: entry:
2338 ; CHECK-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
2339 ; CHECK-DISABLED-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
2340 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i1 false, i1 true, i1 true)
2341 ; CHECK-DISABLED-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
2342 ; CHECK-DISABLED-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
2343 ; CHECK-DISABLED: user_code.entry:
2344 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR3]]
2345 ; CHECK-DISABLED-NEXT: call void @__omp_outlined__15(i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR3]]
2346 ; CHECK-DISABLED-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true)
2347 ; CHECK-DISABLED-NEXT: ret void
2348 ; CHECK-DISABLED: worker.exit:
2349 ; CHECK-DISABLED-NEXT: ret void
2352 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2353 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__15
2354 ; CHECK-DISABLED-SAME: (i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
2355 ; CHECK-DISABLED-NEXT: entry:
2356 ; CHECK-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
2357 ; CHECK-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
2358 ; CHECK-DISABLED-NEXT: [[CALL:%.*]] = call i32 bitcast (i32 (...)* @omp_get_thread_num to i32 ()*)() #[[ATTR9]]
2359 ; CHECK-DISABLED-NEXT: call void @simple_state_machine_interprocedural_nested_recursive_after.internalized(i32 [[CALL]]) #[[ATTR7]]
2360 ; CHECK-DISABLED-NEXT: ret void
2363 ; CHECK-DISABLED: Function Attrs: convergent noinline nounwind
2364 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@simple_state_machine_interprocedural_nested_recursive_after.internalized
2365 ; CHECK-DISABLED-SAME: (i32 [[A:%.*]]) #[[ATTR1]] {
2366 ; CHECK-DISABLED-NEXT: entry:
2367 ; CHECK-DISABLED-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4
2368 ; CHECK-DISABLED-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4
2369 ; CHECK-DISABLED-NEXT: [[CMP:%.*]] = icmp eq i32 [[A]], 0
2370 ; CHECK-DISABLED-NEXT: br i1 [[CMP]], label [[IF_THEN:%.*]], label [[IF_END:%.*]]
2371 ; CHECK-DISABLED: if.then:
2372 ; CHECK-DISABLED-NEXT: br label [[RETURN:%.*]]
2373 ; CHECK-DISABLED: if.end:
2374 ; CHECK-DISABLED-NEXT: [[SUB:%.*]] = sub nsw i32 [[A]], 1
2375 ; CHECK-DISABLED-NEXT: call void @simple_state_machine_interprocedural_nested_recursive_after.internalized(i32 [[SUB]]) #[[ATTR7]]
2376 ; CHECK-DISABLED-NEXT: call void @simple_state_machine_interprocedural_nested_recursive_after_after.internalized() #[[ATTR7]]
2377 ; CHECK-DISABLED-NEXT: br label [[RETURN]]
2378 ; CHECK-DISABLED: return:
2379 ; CHECK-DISABLED-NEXT: ret void
2382 ; CHECK-DISABLED: Function Attrs: convergent noinline nounwind
2383 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@simple_state_machine_interprocedural_nested_recursive_after
2384 ; CHECK-DISABLED-SAME: (i32 [[A:%.*]]) #[[ATTR1]] {
2385 ; CHECK-DISABLED-NEXT: entry:
2386 ; CHECK-DISABLED-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4
2387 ; CHECK-DISABLED-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4
2388 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = load i32, i32* [[A_ADDR]], align 4
2389 ; CHECK-DISABLED-NEXT: [[CMP:%.*]] = icmp eq i32 [[TMP0]], 0
2390 ; CHECK-DISABLED-NEXT: br i1 [[CMP]], label [[IF_THEN:%.*]], label [[IF_END:%.*]]
2391 ; CHECK-DISABLED: if.then:
2392 ; CHECK-DISABLED-NEXT: br label [[RETURN:%.*]]
2393 ; CHECK-DISABLED: if.end:
2394 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = load i32, i32* [[A_ADDR]], align 4
2395 ; CHECK-DISABLED-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP1]], 1
2396 ; CHECK-DISABLED-NEXT: call void @simple_state_machine_interprocedural_nested_recursive_after(i32 [[SUB]]) #[[ATTR9]]
2397 ; CHECK-DISABLED-NEXT: call void @simple_state_machine_interprocedural_nested_recursive_after_after() #[[ATTR9]]
2398 ; CHECK-DISABLED-NEXT: br label [[RETURN]]
2399 ; CHECK-DISABLED: return:
2400 ; CHECK-DISABLED-NEXT: ret void
2403 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2404 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_offloading_14_a36502b_no_state_machine_weak_callee_l112
2405 ; CHECK-DISABLED-SAME: () #[[ATTR0]] {
2406 ; CHECK-DISABLED-NEXT: entry:
2407 ; CHECK-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
2408 ; CHECK-DISABLED-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
2409 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i1 false, i1 true, i1 true)
2410 ; CHECK-DISABLED-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
2411 ; CHECK-DISABLED-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
2412 ; CHECK-DISABLED: user_code.entry:
2413 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR3]]
2414 ; CHECK-DISABLED-NEXT: call void @__omp_outlined__16(i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR3]]
2415 ; CHECK-DISABLED-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true)
2416 ; CHECK-DISABLED-NEXT: ret void
2417 ; CHECK-DISABLED: worker.exit:
2418 ; CHECK-DISABLED-NEXT: ret void
2421 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2422 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__16
2423 ; CHECK-DISABLED-SAME: (i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree nonnull readnone align 4 dereferenceable(4) [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
2424 ; CHECK-DISABLED-NEXT: entry:
2425 ; CHECK-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
2426 ; CHECK-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
2427 ; CHECK-DISABLED-NEXT: call void @weak_callee_empty() #[[ATTR7]]
2428 ; CHECK-DISABLED-NEXT: ret void
2431 ; CHECK-DISABLED: Function Attrs: convergent noinline nounwind
2432 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@weak_callee_empty
2433 ; CHECK-DISABLED-SAME: () #[[ATTR1]] {
2434 ; CHECK-DISABLED-NEXT: entry:
2435 ; CHECK-DISABLED-NEXT: ret void
2438 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2439 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__17
2440 ; CHECK-DISABLED-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
2441 ; CHECK-DISABLED-NEXT: entry:
2442 ; CHECK-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
2443 ; CHECK-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
2444 ; CHECK-DISABLED-NEXT: call void @p0() #[[ATTR9]]
2445 ; CHECK-DISABLED-NEXT: ret void
2448 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2449 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__17_wrapper
2450 ; CHECK-DISABLED-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
2451 ; CHECK-DISABLED-NEXT: entry:
2452 ; CHECK-DISABLED-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
2453 ; CHECK-DISABLED-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
2454 ; CHECK-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
2455 ; CHECK-DISABLED-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
2456 ; CHECK-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
2457 ; CHECK-DISABLED-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
2458 ; CHECK-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
2459 ; CHECK-DISABLED-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
2460 ; CHECK-DISABLED-NEXT: call void @__omp_outlined__17(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
2461 ; CHECK-DISABLED-NEXT: ret void
2464 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2465 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__18
2466 ; CHECK-DISABLED-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
2467 ; CHECK-DISABLED-NEXT: entry:
2468 ; CHECK-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
2469 ; CHECK-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
2470 ; CHECK-DISABLED-NEXT: call void @p0() #[[ATTR9]]
2471 ; CHECK-DISABLED-NEXT: ret void
2474 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2475 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__18_wrapper
2476 ; CHECK-DISABLED-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
2477 ; CHECK-DISABLED-NEXT: entry:
2478 ; CHECK-DISABLED-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
2479 ; CHECK-DISABLED-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
2480 ; CHECK-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
2481 ; CHECK-DISABLED-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
2482 ; CHECK-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
2483 ; CHECK-DISABLED-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
2484 ; CHECK-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
2485 ; CHECK-DISABLED-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
2486 ; CHECK-DISABLED-NEXT: call void @__omp_outlined__18(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
2487 ; CHECK-DISABLED-NEXT: ret void
2490 ; CHECK-DISABLED: Function Attrs: convergent noinline nounwind
2491 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@simple_state_machine_interprocedural_nested_recursive_after_after.internalized
2492 ; CHECK-DISABLED-SAME: () #[[ATTR1]] {
2493 ; CHECK-DISABLED-NEXT: entry:
2494 ; CHECK-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
2495 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* noundef @[[GLOB2]]) #[[ATTR3]]
2496 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
2497 ; CHECK-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB2]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__19 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__19_wrapper to i8*), i8** noundef [[TMP1]], i64 noundef 0)
2498 ; CHECK-DISABLED-NEXT: ret void
2501 ; CHECK-DISABLED: Function Attrs: convergent noinline nounwind
2502 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@simple_state_machine_interprocedural_nested_recursive_after_after
2503 ; CHECK-DISABLED-SAME: () #[[ATTR1]] {
2504 ; CHECK-DISABLED-NEXT: entry:
2505 ; CHECK-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
2506 ; CHECK-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2]])
2507 ; CHECK-DISABLED-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
2508 ; CHECK-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP0]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__19 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__19_wrapper to i8*), i8** [[TMP1]], i64 0)
2509 ; CHECK-DISABLED-NEXT: ret void
2512 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2513 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__19
2514 ; CHECK-DISABLED-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
2515 ; CHECK-DISABLED-NEXT: entry:
2516 ; CHECK-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
2517 ; CHECK-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
2518 ; CHECK-DISABLED-NEXT: call void @p0() #[[ATTR9]]
2519 ; CHECK-DISABLED-NEXT: ret void
2522 ; CHECK-DISABLED: Function Attrs: convergent noinline norecurse nounwind
2523 ; CHECK-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__19_wrapper
2524 ; CHECK-DISABLED-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
2525 ; CHECK-DISABLED-NEXT: entry:
2526 ; CHECK-DISABLED-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
2527 ; CHECK-DISABLED-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
2528 ; CHECK-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
2529 ; CHECK-DISABLED-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
2530 ; CHECK-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
2531 ; CHECK-DISABLED-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
2532 ; CHECK-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
2533 ; CHECK-DISABLED-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
2534 ; CHECK-DISABLED-NEXT: call void @__omp_outlined__19(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
2535 ; CHECK-DISABLED-NEXT: ret void