1; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --function-signature --check-globals 2; RUN: opt --mtriple=amdgcn-amd-amdhsa --data-layout=A5 -S -passes=openmp-opt < %s | FileCheck %s --check-prefixes=AMDGPU 3; RUN: opt --mtriple=nvptx64-- -S -passes=openmp-opt < %s | FileCheck %s --check-prefixes=NVPTX 4; RUN: opt --mtriple=amdgcn-amd-amdhsa --data-layout=A5 -S -passes=openmp-opt -openmp-opt-disable-spmdization < %s | FileCheck %s --check-prefix=AMDGPU-DISABLED 5; RUN: opt --mtriple=nvptx64-- -S -passes=openmp-opt -openmp-opt-disable-spmdization < %s | FileCheck %s --check-prefix=NVPTX-DISABLED 6 7;; void unknown(void); 8;; void spmd_amenable(void) __attribute__((assume("ompx_spmd_amenable"))); 9;; 10;; void sequential_loop() { 11;; #pragma omp target teams 12;; { 13;; for (int i = 0; i < 100; ++i) { 14;; #pragma omp parallel 15;; { 16;; unknown(); 17;; } 18;; } 19;; spmd_amenable(); 20;; } 21;; } 22;; 23;; void use(__attribute__((noescape)) int *) __attribute__((assume("ompx_spmd_amenable"))); 24;; 25;; void sequential_loop_to_stack_var() { 26;; #pragma omp target teams 27;; { 28;; int x; 29;; use(&x); 30;; for (int i = 0; i < 100; ++i) { 31;; #pragma omp parallel 32;; { 33;; unknown(); 34;; } 35;; } 36;; spmd_amenable(); 37;; } 38;; } 39;; 40;; void sequential_loop_to_shared_var() { 41;; #pragma omp target teams 42;; { 43;; int x; 44;; for (int i = 0; i < 100; ++i) { 45;; #pragma omp parallel 46;; { 47;; x++; 48;; unknown(); 49;; } 50;; } 51;; spmd_amenable(); 52;; } 53;; } 54;; 55;; void sequential_loop_to_shared_var_guarded() { 56;; #pragma omp target teams 57;; { 58;; int x = 42; 59;; for (int i = 0; i < 100; ++i) { 60;; #pragma omp parallel 61;; { 62;; x++; 63;; unknown(); 64;; } 65;; } 66;; spmd_amenable(); 67;; } 68;; } 69;; 70;; void do_not_spmdize_target() { 71;; #pragma omp target teams 72;; { 73;; // Incompatible parallel level, called both 74;; // from parallel and target regions 75;; unknown(); 76;; } 77;; } 78 79%struct.ident_t = type { i32, i32, i32, i32, i8* } 80 81@0 = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00", align 1 82@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 83@__omp_offloading_14_a34ca11_sequential_loop_l5_exec_mode = weak constant i8 1 84@__omp_offloading_14_a34ca11_sequential_loop_to_stack_var_l20_exec_mode = weak constant i8 1 85@__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_l35_exec_mode = weak constant i8 1 86@__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_guarded_l50_exec_mode = weak constant i8 1 87@__omp_offloading_14_a34ca11_do_not_spmdize_target_l65_exec_mode = weak constant i8 1 88@llvm.compiler.used = appending global [5 x i8*] [i8* @__omp_offloading_14_a34ca11_sequential_loop_l5_exec_mode, i8* @__omp_offloading_14_a34ca11_sequential_loop_to_stack_var_l20_exec_mode, i8* @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_l35_exec_mode, i8* @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_guarded_l50_exec_mode, i8* @__omp_offloading_14_a34ca11_do_not_spmdize_target_l65_exec_mode], section "llvm.metadata" 89 90 91;. 92; AMDGPU: @[[GLOB0:[0-9]+]] = private unnamed_addr constant [23 x i8] c" 93; AMDGPU: @[[GLOB1:[0-9]+]] = private unnamed_addr constant [[STRUCT_IDENT_T:%.*]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @[[GLOB0]], i32 0, i32 0) }, align 8 94; AMDGPU: @[[__OMP_OFFLOADING_14_A34CA11_SEQUENTIAL_LOOP_L5_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 3 95; AMDGPU: @[[__OMP_OFFLOADING_14_A34CA11_SEQUENTIAL_LOOP_TO_STACK_VAR_L20_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 3 96; AMDGPU: @[[__OMP_OFFLOADING_14_A34CA11_SEQUENTIAL_LOOP_TO_SHARED_VAR_L35_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 3 97; AMDGPU: @[[__OMP_OFFLOADING_14_A34CA11_SEQUENTIAL_LOOP_TO_SHARED_VAR_GUARDED_L50_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 3 98; AMDGPU: @[[__OMP_OFFLOADING_14_A34CA11_DO_NOT_SPMDIZE_TARGET_L65_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 1 99; AMDGPU: @[[LLVM_COMPILER_USED:[a-zA-Z0-9_$"\\.-]+]] = appending global [5 x i8*] [i8* @__omp_offloading_14_a34ca11_sequential_loop_l5_exec_mode, i8* @__omp_offloading_14_a34ca11_sequential_loop_to_stack_var_l20_exec_mode, i8* @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_l35_exec_mode, i8* @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_guarded_l50_exec_mode, i8* @__omp_offloading_14_a34ca11_do_not_spmdize_target_l65_exec_mode], section "llvm.metadata" 100; AMDGPU: @[[X:[a-zA-Z0-9_$"\\.-]+]] = internal addrspace(3) global [4 x i8] undef, align 32 101; AMDGPU: @[[X_1:[a-zA-Z0-9_$"\\.-]+]] = internal addrspace(3) global [4 x i8] undef, align 32 102;. 103; NVPTX: @[[GLOB0:[0-9]+]] = private unnamed_addr constant [23 x i8] c" 104; NVPTX: @[[GLOB1:[0-9]+]] = private unnamed_addr constant [[STRUCT_IDENT_T:%.*]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @[[GLOB0]], i32 0, i32 0) }, align 8 105; NVPTX: @[[__OMP_OFFLOADING_14_A34CA11_SEQUENTIAL_LOOP_L5_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 3 106; NVPTX: @[[__OMP_OFFLOADING_14_A34CA11_SEQUENTIAL_LOOP_TO_STACK_VAR_L20_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 3 107; NVPTX: @[[__OMP_OFFLOADING_14_A34CA11_SEQUENTIAL_LOOP_TO_SHARED_VAR_L35_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 3 108; NVPTX: @[[__OMP_OFFLOADING_14_A34CA11_SEQUENTIAL_LOOP_TO_SHARED_VAR_GUARDED_L50_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 3 109; NVPTX: @[[__OMP_OFFLOADING_14_A34CA11_DO_NOT_SPMDIZE_TARGET_L65_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 1 110; NVPTX: @[[LLVM_COMPILER_USED:[a-zA-Z0-9_$"\\.-]+]] = appending global [5 x i8*] [i8* @__omp_offloading_14_a34ca11_sequential_loop_l5_exec_mode, i8* @__omp_offloading_14_a34ca11_sequential_loop_to_stack_var_l20_exec_mode, i8* @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_l35_exec_mode, i8* @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_guarded_l50_exec_mode, i8* @__omp_offloading_14_a34ca11_do_not_spmdize_target_l65_exec_mode], section "llvm.metadata" 111; NVPTX: @[[X:[a-zA-Z0-9_$"\\.-]+]] = internal addrspace(3) global [4 x i8] undef, align 32 112; NVPTX: @[[X1:[a-zA-Z0-9_$"\\.-]+]] = internal addrspace(3) global [4 x i8] undef, align 32 113;. 114; AMDGPU-DISABLED: @[[GLOB0:[0-9]+]] = private unnamed_addr constant [23 x i8] c" 115; AMDGPU-DISABLED: @[[GLOB1:[0-9]+]] = private unnamed_addr constant [[STRUCT_IDENT_T:%.*]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @[[GLOB0]], i32 0, i32 0) }, align 8 116; AMDGPU-DISABLED: @[[__OMP_OFFLOADING_14_A34CA11_SEQUENTIAL_LOOP_L5_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 1 117; AMDGPU-DISABLED: @[[__OMP_OFFLOADING_14_A34CA11_SEQUENTIAL_LOOP_TO_STACK_VAR_L20_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 1 118; AMDGPU-DISABLED: @[[__OMP_OFFLOADING_14_A34CA11_SEQUENTIAL_LOOP_TO_SHARED_VAR_L35_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 1 119; AMDGPU-DISABLED: @[[__OMP_OFFLOADING_14_A34CA11_SEQUENTIAL_LOOP_TO_SHARED_VAR_GUARDED_L50_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 1 120; AMDGPU-DISABLED: @[[__OMP_OFFLOADING_14_A34CA11_DO_NOT_SPMDIZE_TARGET_L65_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 1 121; AMDGPU-DISABLED: @[[LLVM_COMPILER_USED:[a-zA-Z0-9_$"\\.-]+]] = appending global [5 x i8*] [i8* @__omp_offloading_14_a34ca11_sequential_loop_l5_exec_mode, i8* @__omp_offloading_14_a34ca11_sequential_loop_to_stack_var_l20_exec_mode, i8* @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_l35_exec_mode, i8* @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_guarded_l50_exec_mode, i8* @__omp_offloading_14_a34ca11_do_not_spmdize_target_l65_exec_mode], section "llvm.metadata" 122; AMDGPU-DISABLED: @[[X:[a-zA-Z0-9_$"\\.-]+]] = internal addrspace(3) global [4 x i8] undef, align 32 123; AMDGPU-DISABLED: @[[X_1:[a-zA-Z0-9_$"\\.-]+]] = internal addrspace(3) global [4 x i8] undef, align 32 124; AMDGPU-DISABLED: @[[__OMP_OUTLINED__1_WRAPPER_ID:[a-zA-Z0-9_$"\\.-]+]] = private constant i8 undef 125; AMDGPU-DISABLED: @[[__OMP_OUTLINED__3_WRAPPER_ID:[a-zA-Z0-9_$"\\.-]+]] = private constant i8 undef 126; AMDGPU-DISABLED: @[[__OMP_OUTLINED__5_WRAPPER_ID:[a-zA-Z0-9_$"\\.-]+]] = private constant i8 undef 127; AMDGPU-DISABLED: @[[__OMP_OUTLINED__7_WRAPPER_ID:[a-zA-Z0-9_$"\\.-]+]] = private constant i8 undef 128;. 129; NVPTX-DISABLED: @[[GLOB0:[0-9]+]] = private unnamed_addr constant [23 x i8] c" 130; NVPTX-DISABLED: @[[GLOB1:[0-9]+]] = private unnamed_addr constant [[STRUCT_IDENT_T:%.*]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @[[GLOB0]], i32 0, i32 0) }, align 8 131; NVPTX-DISABLED: @[[__OMP_OFFLOADING_14_A34CA11_SEQUENTIAL_LOOP_L5_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 1 132; NVPTX-DISABLED: @[[__OMP_OFFLOADING_14_A34CA11_SEQUENTIAL_LOOP_TO_STACK_VAR_L20_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 1 133; NVPTX-DISABLED: @[[__OMP_OFFLOADING_14_A34CA11_SEQUENTIAL_LOOP_TO_SHARED_VAR_L35_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 1 134; NVPTX-DISABLED: @[[__OMP_OFFLOADING_14_A34CA11_SEQUENTIAL_LOOP_TO_SHARED_VAR_GUARDED_L50_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 1 135; NVPTX-DISABLED: @[[__OMP_OFFLOADING_14_A34CA11_DO_NOT_SPMDIZE_TARGET_L65_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 1 136; NVPTX-DISABLED: @[[LLVM_COMPILER_USED:[a-zA-Z0-9_$"\\.-]+]] = appending global [5 x i8*] [i8* @__omp_offloading_14_a34ca11_sequential_loop_l5_exec_mode, i8* @__omp_offloading_14_a34ca11_sequential_loop_to_stack_var_l20_exec_mode, i8* @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_l35_exec_mode, i8* @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_guarded_l50_exec_mode, i8* @__omp_offloading_14_a34ca11_do_not_spmdize_target_l65_exec_mode], section "llvm.metadata" 137; NVPTX-DISABLED: @[[X:[a-zA-Z0-9_$"\\.-]+]] = internal addrspace(3) global [4 x i8] undef, align 32 138; NVPTX-DISABLED: @[[X1:[a-zA-Z0-9_$"\\.-]+]] = internal addrspace(3) global [4 x i8] undef, align 32 139; NVPTX-DISABLED: @[[__OMP_OUTLINED__1_WRAPPER_ID:[a-zA-Z0-9_$"\\.-]+]] = private constant i8 undef 140; NVPTX-DISABLED: @[[__OMP_OUTLINED__3_WRAPPER_ID:[a-zA-Z0-9_$"\\.-]+]] = private constant i8 undef 141; NVPTX-DISABLED: @[[__OMP_OUTLINED__5_WRAPPER_ID:[a-zA-Z0-9_$"\\.-]+]] = private constant i8 undef 142; NVPTX-DISABLED: @[[__OMP_OUTLINED__7_WRAPPER_ID:[a-zA-Z0-9_$"\\.-]+]] = private constant i8 undef 143;. 144define weak void @__omp_offloading_14_a34ca11_sequential_loop_l5() #0 { 145; AMDGPU-LABEL: define {{[^@]+}}@__omp_offloading_14_a34ca11_sequential_loop_l5 146; AMDGPU-SAME: () #[[ATTR0:[0-9]+]] { 147; AMDGPU-NEXT: entry: 148; AMDGPU-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 149; AMDGPU-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 150; AMDGPU-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) 151; AMDGPU-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 152; AMDGPU-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 153; AMDGPU: user_code.entry: 154; AMDGPU-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR4:[0-9]+]] 155; AMDGPU-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4 156; AMDGPU-NEXT: call void @__omp_outlined__(i32* noalias nocapture noundef nonnull readonly align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR4]] 157; AMDGPU-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) 158; AMDGPU-NEXT: ret void 159; AMDGPU: worker.exit: 160; AMDGPU-NEXT: ret void 161; 162; NVPTX-LABEL: define {{[^@]+}}@__omp_offloading_14_a34ca11_sequential_loop_l5 163; NVPTX-SAME: () #[[ATTR0:[0-9]+]] { 164; NVPTX-NEXT: entry: 165; NVPTX-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 166; NVPTX-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 167; NVPTX-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) 168; NVPTX-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 169; NVPTX-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 170; NVPTX: user_code.entry: 171; NVPTX-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR4:[0-9]+]] 172; NVPTX-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4 173; NVPTX-NEXT: call void @__omp_outlined__(i32* noalias nocapture noundef nonnull readonly align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR4]] 174; NVPTX-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) 175; NVPTX-NEXT: ret void 176; NVPTX: worker.exit: 177; NVPTX-NEXT: ret void 178; 179; AMDGPU-DISABLED-LABEL: define {{[^@]+}}@__omp_offloading_14_a34ca11_sequential_loop_l5 180; AMDGPU-DISABLED-SAME: () #[[ATTR0:[0-9]+]] { 181; AMDGPU-DISABLED-NEXT: entry: 182; AMDGPU-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR:%.*]] = alloca i8*, align 8, addrspace(5) 183; AMDGPU-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 184; AMDGPU-DISABLED-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 185; AMDGPU-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 186; AMDGPU-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 1, i1 false, i1 true) 187; AMDGPU-DISABLED-NEXT: [[BLOCK_HW_SIZE:%.*]] = call i32 @__kmpc_get_hardware_num_threads_in_block() 188; AMDGPU-DISABLED-NEXT: [[WARP_SIZE:%.*]] = call i32 @__kmpc_get_warp_size() 189; AMDGPU-DISABLED-NEXT: [[BLOCK_SIZE:%.*]] = sub i32 [[BLOCK_HW_SIZE]], [[WARP_SIZE]] 190; AMDGPU-DISABLED-NEXT: [[THREAD_IS_MAIN_OR_WORKER:%.*]] = icmp slt i32 [[TMP0]], [[BLOCK_SIZE]] 191; AMDGPU-DISABLED-NEXT: br i1 [[THREAD_IS_MAIN_OR_WORKER]], label [[IS_WORKER_CHECK:%.*]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]] 192; AMDGPU-DISABLED: is_worker_check: 193; AMDGPU-DISABLED-NEXT: [[THREAD_IS_WORKER:%.*]] = icmp ne i32 [[TMP0]], -1 194; AMDGPU-DISABLED-NEXT: br i1 [[THREAD_IS_WORKER]], label [[WORKER_STATE_MACHINE_BEGIN:%.*]], label [[THREAD_USER_CODE_CHECK:%.*]] 195; AMDGPU-DISABLED: worker_state_machine.begin: 196; AMDGPU-DISABLED-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 197; AMDGPU-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR_GENERIC:%.*]] = addrspacecast i8* addrspace(5)* [[WORKER_WORK_FN_ADDR]] to i8** 198; AMDGPU-DISABLED-NEXT: [[WORKER_IS_ACTIVE:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORKER_WORK_FN_ADDR_GENERIC]]) 199; AMDGPU-DISABLED-NEXT: [[WORKER_WORK_FN:%.*]] = load i8*, i8** [[WORKER_WORK_FN_ADDR_GENERIC]], align 8 200; AMDGPU-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)* 201; AMDGPU-DISABLED-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null 202; AMDGPU-DISABLED-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]] 203; AMDGPU-DISABLED: worker_state_machine.finished: 204; AMDGPU-DISABLED-NEXT: ret void 205; AMDGPU-DISABLED: worker_state_machine.is_active.check: 206; AMDGPU-DISABLED-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]] 207; AMDGPU-DISABLED: worker_state_machine.parallel_region.check: 208; AMDGPU-DISABLED-NEXT: br i1 true, label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK1:%.*]] 209; AMDGPU-DISABLED: worker_state_machine.parallel_region.execute: 210; AMDGPU-DISABLED-NEXT: call void @__omp_outlined__1_wrapper(i16 0, i32 [[TMP0]]) 211; AMDGPU-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]] 212; AMDGPU-DISABLED: worker_state_machine.parallel_region.check1: 213; AMDGPU-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]] 214; AMDGPU-DISABLED: worker_state_machine.parallel_region.end: 215; AMDGPU-DISABLED-NEXT: call void @__kmpc_kernel_end_parallel() 216; AMDGPU-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_DONE_BARRIER]] 217; AMDGPU-DISABLED: worker_state_machine.done.barrier: 218; AMDGPU-DISABLED-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 219; AMDGPU-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_BEGIN]] 220; AMDGPU-DISABLED: thread.user_code.check: 221; AMDGPU-DISABLED-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 222; AMDGPU-DISABLED-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 223; AMDGPU-DISABLED: user_code.entry: 224; AMDGPU-DISABLED-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR4:[0-9]+]] 225; AMDGPU-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4 226; AMDGPU-DISABLED-NEXT: call void @__omp_outlined__(i32* noalias nocapture noundef nonnull readonly align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR4]] 227; AMDGPU-DISABLED-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true) 228; AMDGPU-DISABLED-NEXT: ret void 229; AMDGPU-DISABLED: worker.exit: 230; AMDGPU-DISABLED-NEXT: ret void 231; 232; NVPTX-DISABLED-LABEL: define {{[^@]+}}@__omp_offloading_14_a34ca11_sequential_loop_l5 233; NVPTX-DISABLED-SAME: () #[[ATTR0:[0-9]+]] { 234; NVPTX-DISABLED-NEXT: entry: 235; NVPTX-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR:%.*]] = alloca i8*, align 8 236; NVPTX-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 237; NVPTX-DISABLED-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 238; NVPTX-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 239; NVPTX-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 1, i1 false, i1 true) 240; NVPTX-DISABLED-NEXT: [[BLOCK_HW_SIZE:%.*]] = call i32 @__kmpc_get_hardware_num_threads_in_block() 241; NVPTX-DISABLED-NEXT: [[WARP_SIZE:%.*]] = call i32 @__kmpc_get_warp_size() 242; NVPTX-DISABLED-NEXT: [[BLOCK_SIZE:%.*]] = sub i32 [[BLOCK_HW_SIZE]], [[WARP_SIZE]] 243; NVPTX-DISABLED-NEXT: [[THREAD_IS_MAIN_OR_WORKER:%.*]] = icmp slt i32 [[TMP0]], [[BLOCK_SIZE]] 244; NVPTX-DISABLED-NEXT: br i1 [[THREAD_IS_MAIN_OR_WORKER]], label [[IS_WORKER_CHECK:%.*]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]] 245; NVPTX-DISABLED: is_worker_check: 246; NVPTX-DISABLED-NEXT: [[THREAD_IS_WORKER:%.*]] = icmp ne i32 [[TMP0]], -1 247; NVPTX-DISABLED-NEXT: br i1 [[THREAD_IS_WORKER]], label [[WORKER_STATE_MACHINE_BEGIN:%.*]], label [[THREAD_USER_CODE_CHECK:%.*]] 248; NVPTX-DISABLED: worker_state_machine.begin: 249; NVPTX-DISABLED-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 250; NVPTX-DISABLED-NEXT: [[WORKER_IS_ACTIVE:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORKER_WORK_FN_ADDR]]) 251; NVPTX-DISABLED-NEXT: [[WORKER_WORK_FN:%.*]] = load i8*, i8** [[WORKER_WORK_FN_ADDR]], align 8 252; NVPTX-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)* 253; NVPTX-DISABLED-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null 254; NVPTX-DISABLED-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]] 255; NVPTX-DISABLED: worker_state_machine.finished: 256; NVPTX-DISABLED-NEXT: ret void 257; NVPTX-DISABLED: worker_state_machine.is_active.check: 258; NVPTX-DISABLED-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]] 259; NVPTX-DISABLED: worker_state_machine.parallel_region.check: 260; NVPTX-DISABLED-NEXT: br i1 true, label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK1:%.*]] 261; NVPTX-DISABLED: worker_state_machine.parallel_region.execute: 262; NVPTX-DISABLED-NEXT: call void @__omp_outlined__1_wrapper(i16 0, i32 [[TMP0]]) 263; NVPTX-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]] 264; NVPTX-DISABLED: worker_state_machine.parallel_region.check1: 265; NVPTX-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]] 266; NVPTX-DISABLED: worker_state_machine.parallel_region.end: 267; NVPTX-DISABLED-NEXT: call void @__kmpc_kernel_end_parallel() 268; NVPTX-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_DONE_BARRIER]] 269; NVPTX-DISABLED: worker_state_machine.done.barrier: 270; NVPTX-DISABLED-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 271; NVPTX-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_BEGIN]] 272; NVPTX-DISABLED: thread.user_code.check: 273; NVPTX-DISABLED-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 274; NVPTX-DISABLED-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 275; NVPTX-DISABLED: user_code.entry: 276; NVPTX-DISABLED-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR4:[0-9]+]] 277; NVPTX-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4 278; NVPTX-DISABLED-NEXT: call void @__omp_outlined__(i32* noalias nocapture noundef nonnull readonly align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR4]] 279; NVPTX-DISABLED-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true) 280; NVPTX-DISABLED-NEXT: ret void 281; NVPTX-DISABLED: worker.exit: 282; NVPTX-DISABLED-NEXT: ret void 283; 284entry: 285 %.zero.addr = alloca i32, align 4 286 %.threadid_temp. = alloca i32, align 4 287 store i32 0, i32* %.zero.addr, align 4 288 %0 = call i32 @__kmpc_target_init(%struct.ident_t* @1, i8 1, i1 true, i1 true) 289 %exec_user_code = icmp eq i32 %0, -1 290 br i1 %exec_user_code, label %user_code.entry, label %worker.exit 291 292user_code.entry: ; preds = %entry 293 %1 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @1) 294 store i32 %1, i32* %.threadid_temp., align 4 295 call void @__omp_outlined__(i32* %.threadid_temp., i32* %.zero.addr) #3 296 call void @__kmpc_target_deinit(%struct.ident_t* @1, i8 1, i1 true) 297 ret void 298 299worker.exit: ; preds = %entry 300 ret void 301} 302 303declare i32 @__kmpc_target_init(%struct.ident_t*, i8, i1, i1) 304 305define internal void @__omp_outlined__(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 { 306; 307; 308; AMDGPU-LABEL: define {{[^@]+}}@__omp_outlined__ 309; AMDGPU-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]] { 310; AMDGPU-NEXT: entry: 311; AMDGPU-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 312; AMDGPU-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 313; AMDGPU-NEXT: [[I:%.*]] = alloca i32, align 4 314; AMDGPU-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8 315; AMDGPU-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 316; AMDGPU-NEXT: store i32 0, i32* [[I]], align 4 317; AMDGPU-NEXT: br label [[FOR_COND:%.*]] 318; AMDGPU: for.cond: 319; AMDGPU-NEXT: [[TMP0:%.*]] = load i32, i32* [[I]], align 4 320; AMDGPU-NEXT: [[CMP:%.*]] = icmp slt i32 [[TMP0]], 100 321; AMDGPU-NEXT: br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]] 322; AMDGPU: for.body: 323; AMDGPU-NEXT: [[TMP1:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4 324; AMDGPU-NEXT: [[TMP2:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 325; AMDGPU-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP1]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__1 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*), i8** noundef [[TMP2]], i64 noundef 0) 326; AMDGPU-NEXT: br label [[FOR_INC:%.*]] 327; AMDGPU: for.inc: 328; AMDGPU-NEXT: [[TMP3:%.*]] = load i32, i32* [[I]], align 4 329; AMDGPU-NEXT: [[INC:%.*]] = add nsw i32 [[TMP3]], 1 330; AMDGPU-NEXT: store i32 [[INC]], i32* [[I]], align 4 331; AMDGPU-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP13:![0-9]+]] 332; AMDGPU: for.end: 333; AMDGPU-NEXT: call void @indirection() #[[ATTR7:[0-9]+]] 334; AMDGPU-NEXT: ret void 335; 336; NVPTX-LABEL: define {{[^@]+}}@__omp_outlined__ 337; NVPTX-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]] { 338; NVPTX-NEXT: entry: 339; NVPTX-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 340; NVPTX-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 341; NVPTX-NEXT: [[I:%.*]] = alloca i32, align 4 342; NVPTX-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8 343; NVPTX-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 344; NVPTX-NEXT: store i32 0, i32* [[I]], align 4 345; NVPTX-NEXT: br label [[FOR_COND:%.*]] 346; NVPTX: for.cond: 347; NVPTX-NEXT: [[TMP0:%.*]] = load i32, i32* [[I]], align 4 348; NVPTX-NEXT: [[CMP:%.*]] = icmp slt i32 [[TMP0]], 100 349; NVPTX-NEXT: br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]] 350; NVPTX: for.body: 351; NVPTX-NEXT: [[TMP1:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4 352; NVPTX-NEXT: [[TMP2:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 353; NVPTX-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP1]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__1 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*), i8** noundef [[TMP2]], i64 noundef 0) 354; NVPTX-NEXT: br label [[FOR_INC:%.*]] 355; NVPTX: for.inc: 356; NVPTX-NEXT: [[TMP3:%.*]] = load i32, i32* [[I]], align 4 357; NVPTX-NEXT: [[INC:%.*]] = add nsw i32 [[TMP3]], 1 358; NVPTX-NEXT: store i32 [[INC]], i32* [[I]], align 4 359; NVPTX-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP13:![0-9]+]] 360; NVPTX: for.end: 361; NVPTX-NEXT: call void @indirection() #[[ATTR7:[0-9]+]] 362; NVPTX-NEXT: ret void 363; 364; AMDGPU-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__ 365; AMDGPU-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]] { 366; AMDGPU-DISABLED-NEXT: entry: 367; AMDGPU-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 368; AMDGPU-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 369; AMDGPU-DISABLED-NEXT: [[I:%.*]] = alloca i32, align 4 370; AMDGPU-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8 371; AMDGPU-DISABLED-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 372; AMDGPU-DISABLED-NEXT: store i32 0, i32* [[I]], align 4 373; AMDGPU-DISABLED-NEXT: br label [[FOR_COND:%.*]] 374; AMDGPU-DISABLED: for.cond: 375; AMDGPU-DISABLED-NEXT: [[TMP0:%.*]] = load i32, i32* [[I]], align 4 376; AMDGPU-DISABLED-NEXT: [[CMP:%.*]] = icmp slt i32 [[TMP0]], 100 377; AMDGPU-DISABLED-NEXT: br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]] 378; AMDGPU-DISABLED: for.body: 379; AMDGPU-DISABLED-NEXT: [[TMP1:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4 380; AMDGPU-DISABLED-NEXT: [[TMP2:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 381; AMDGPU-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP1]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__1 to i8*), i8* noundef @__omp_outlined__1_wrapper.ID, i8** noundef [[TMP2]], i64 noundef 0) 382; AMDGPU-DISABLED-NEXT: br label [[FOR_INC:%.*]] 383; AMDGPU-DISABLED: for.inc: 384; AMDGPU-DISABLED-NEXT: [[TMP3:%.*]] = load i32, i32* [[I]], align 4 385; AMDGPU-DISABLED-NEXT: [[INC:%.*]] = add nsw i32 [[TMP3]], 1 386; AMDGPU-DISABLED-NEXT: store i32 [[INC]], i32* [[I]], align 4 387; AMDGPU-DISABLED-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP13:![0-9]+]] 388; AMDGPU-DISABLED: for.end: 389; AMDGPU-DISABLED-NEXT: call void @indirection() #[[ATTR7:[0-9]+]] 390; AMDGPU-DISABLED-NEXT: ret void 391; 392; NVPTX-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__ 393; NVPTX-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]] { 394; NVPTX-DISABLED-NEXT: entry: 395; NVPTX-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 396; NVPTX-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 397; NVPTX-DISABLED-NEXT: [[I:%.*]] = alloca i32, align 4 398; NVPTX-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8 399; NVPTX-DISABLED-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 400; NVPTX-DISABLED-NEXT: store i32 0, i32* [[I]], align 4 401; NVPTX-DISABLED-NEXT: br label [[FOR_COND:%.*]] 402; NVPTX-DISABLED: for.cond: 403; NVPTX-DISABLED-NEXT: [[TMP0:%.*]] = load i32, i32* [[I]], align 4 404; NVPTX-DISABLED-NEXT: [[CMP:%.*]] = icmp slt i32 [[TMP0]], 100 405; NVPTX-DISABLED-NEXT: br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]] 406; NVPTX-DISABLED: for.body: 407; NVPTX-DISABLED-NEXT: [[TMP1:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4 408; NVPTX-DISABLED-NEXT: [[TMP2:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 409; NVPTX-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP1]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__1 to i8*), i8* noundef @__omp_outlined__1_wrapper.ID, i8** noundef [[TMP2]], i64 noundef 0) 410; NVPTX-DISABLED-NEXT: br label [[FOR_INC:%.*]] 411; NVPTX-DISABLED: for.inc: 412; NVPTX-DISABLED-NEXT: [[TMP3:%.*]] = load i32, i32* [[I]], align 4 413; NVPTX-DISABLED-NEXT: [[INC:%.*]] = add nsw i32 [[TMP3]], 1 414; NVPTX-DISABLED-NEXT: store i32 [[INC]], i32* [[I]], align 4 415; NVPTX-DISABLED-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP13:![0-9]+]] 416; NVPTX-DISABLED: for.end: 417; NVPTX-DISABLED-NEXT: call void @indirection() #[[ATTR7:[0-9]+]] 418; NVPTX-DISABLED-NEXT: ret void 419; 420entry: 421 %.global_tid..addr = alloca i32*, align 8 422 %.bound_tid..addr = alloca i32*, align 8 423 %i = alloca i32, align 4 424 %captured_vars_addrs = alloca [0 x i8*], align 8 425 store i32* %.global_tid., i32** %.global_tid..addr, align 8 426 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8 427 store i32 0, i32* %i, align 4 428 br label %for.cond 429 430for.cond: ; preds = %for.inc, %entry 431 %0 = load i32, i32* %i, align 4 432 %cmp = icmp slt i32 %0, 100 433 br i1 %cmp, label %for.body, label %for.end 434 435for.body: ; preds = %for.cond 436 %1 = load i32*, i32** %.global_tid..addr, align 8 437 %2 = load i32, i32* %1, align 4 438 %3 = bitcast [0 x i8*]* %captured_vars_addrs to i8** 439 call void @__kmpc_parallel_51(%struct.ident_t* @1, i32 %2, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__1 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*), i8** %3, i64 0) 440 br label %for.inc 441 442for.inc: ; preds = %for.body 443 %4 = load i32, i32* %i, align 4 444 %inc = add nsw i32 %4, 1 445 store i32 %inc, i32* %i, align 4 446 br label %for.cond, !llvm.loop !13 447 448for.end: ; preds = %for.cond 449 call void @indirection() #4 450 ret void 451} 452 453define internal void @indirection() { 454; AMDGPU-LABEL: define {{[^@]+}}@indirection 455; AMDGPU-SAME: () #[[ATTR1:[0-9]+]] { 456; AMDGPU-NEXT: call void @spmd_amenable() #[[ATTR1]] 457; AMDGPU-NEXT: ret void 458; 459; NVPTX-LABEL: define {{[^@]+}}@indirection 460; NVPTX-SAME: () #[[ATTR1:[0-9]+]] { 461; NVPTX-NEXT: call void @spmd_amenable() #[[ATTR1]] 462; NVPTX-NEXT: ret void 463; 464; AMDGPU-DISABLED-LABEL: define {{[^@]+}}@indirection 465; AMDGPU-DISABLED-SAME: () #[[ATTR1:[0-9]+]] { 466; AMDGPU-DISABLED-NEXT: call void @spmd_amenable() #[[ATTR1]] 467; AMDGPU-DISABLED-NEXT: ret void 468; 469; NVPTX-DISABLED-LABEL: define {{[^@]+}}@indirection 470; NVPTX-DISABLED-SAME: () #[[ATTR1:[0-9]+]] { 471; NVPTX-DISABLED-NEXT: call void @spmd_amenable() #[[ATTR1]] 472; NVPTX-DISABLED-NEXT: ret void 473; 474 call void @spmd_amenable() 475 ret void 476} 477 478define internal void @__omp_outlined__1(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 { 479; 480; 481; AMDGPU-LABEL: define {{[^@]+}}@__omp_outlined__1 482; AMDGPU-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] { 483; AMDGPU-NEXT: entry: 484; AMDGPU-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 485; AMDGPU-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 486; AMDGPU-NEXT: call void @unknown() #[[ATTR8:[0-9]+]] 487; AMDGPU-NEXT: ret void 488; 489; NVPTX-LABEL: define {{[^@]+}}@__omp_outlined__1 490; NVPTX-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] { 491; NVPTX-NEXT: entry: 492; NVPTX-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 493; NVPTX-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 494; NVPTX-NEXT: call void @unknown() #[[ATTR8:[0-9]+]] 495; NVPTX-NEXT: ret void 496; 497; AMDGPU-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__1 498; AMDGPU-DISABLED-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] { 499; AMDGPU-DISABLED-NEXT: entry: 500; AMDGPU-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 501; AMDGPU-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 502; AMDGPU-DISABLED-NEXT: call void @unknown() #[[ATTR8:[0-9]+]] 503; AMDGPU-DISABLED-NEXT: ret void 504; 505; NVPTX-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__1 506; NVPTX-DISABLED-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] { 507; NVPTX-DISABLED-NEXT: entry: 508; NVPTX-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 509; NVPTX-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 510; NVPTX-DISABLED-NEXT: call void @unknown() #[[ATTR8:[0-9]+]] 511; NVPTX-DISABLED-NEXT: ret void 512; 513entry: 514 %.global_tid..addr = alloca i32*, align 8 515 %.bound_tid..addr = alloca i32*, align 8 516 store i32* %.global_tid., i32** %.global_tid..addr, align 8 517 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8 518 call void @unknown() #5 519 ret void 520} 521 522declare void @unknown() #1 523 524define internal void @__omp_outlined__1_wrapper(i16 zeroext %0, i32 %1) #0 { 525; 526; 527; AMDGPU-LABEL: define {{[^@]+}}@__omp_outlined__1_wrapper 528; AMDGPU-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 529; AMDGPU-NEXT: entry: 530; AMDGPU-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 531; AMDGPU-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 532; AMDGPU-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 533; AMDGPU-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8 534; AMDGPU-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 535; AMDGPU-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 536; AMDGPU-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 537; AMDGPU-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 538; AMDGPU-NEXT: call void @__omp_outlined__1(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR4]] 539; AMDGPU-NEXT: ret void 540; 541; NVPTX-LABEL: define {{[^@]+}}@__omp_outlined__1_wrapper 542; NVPTX-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 543; NVPTX-NEXT: entry: 544; NVPTX-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 545; NVPTX-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 546; NVPTX-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 547; NVPTX-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8 548; NVPTX-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 549; NVPTX-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 550; NVPTX-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 551; NVPTX-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 552; NVPTX-NEXT: call void @__omp_outlined__1(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR4]] 553; NVPTX-NEXT: ret void 554; 555; AMDGPU-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__1_wrapper 556; AMDGPU-DISABLED-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 557; AMDGPU-DISABLED-NEXT: entry: 558; AMDGPU-DISABLED-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 559; AMDGPU-DISABLED-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 560; AMDGPU-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 561; AMDGPU-DISABLED-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8 562; AMDGPU-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 563; AMDGPU-DISABLED-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 564; AMDGPU-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 565; AMDGPU-DISABLED-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 566; AMDGPU-DISABLED-NEXT: call void @__omp_outlined__1(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR4]] 567; AMDGPU-DISABLED-NEXT: ret void 568; 569; NVPTX-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__1_wrapper 570; NVPTX-DISABLED-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 571; NVPTX-DISABLED-NEXT: entry: 572; NVPTX-DISABLED-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 573; NVPTX-DISABLED-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 574; NVPTX-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 575; NVPTX-DISABLED-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8 576; NVPTX-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 577; NVPTX-DISABLED-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 578; NVPTX-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 579; NVPTX-DISABLED-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 580; NVPTX-DISABLED-NEXT: call void @__omp_outlined__1(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR4]] 581; NVPTX-DISABLED-NEXT: ret void 582; 583entry: 584 %.addr = alloca i16, align 2 585 %.addr1 = alloca i32, align 4 586 %.zero.addr = alloca i32, align 4 587 %global_args = alloca i8**, align 8 588 store i32 0, i32* %.zero.addr, align 4 589 store i16 %0, i16* %.addr, align 2 590 store i32 %1, i32* %.addr1, align 4 591 call void @__kmpc_get_shared_variables(i8*** %global_args) 592 call void @__omp_outlined__1(i32* %.addr1, i32* %.zero.addr) #3 593 ret void 594} 595 596declare void @__kmpc_get_shared_variables(i8***) 597 598declare void @__kmpc_parallel_51(%struct.ident_t*, i32, i32, i32, i32, i8*, i8*, i8**, i64) 599 600declare void @spmd_amenable() 601 602declare i32 @__kmpc_global_thread_num(%struct.ident_t*) #3 603 604declare void @__kmpc_target_deinit(%struct.ident_t*, i8, i1) 605 606define weak void @__omp_offloading_14_a34ca11_sequential_loop_to_stack_var_l20() #0 { 607; 608; 609; AMDGPU-LABEL: define {{[^@]+}}@__omp_offloading_14_a34ca11_sequential_loop_to_stack_var_l20 610; AMDGPU-SAME: () #[[ATTR0]] { 611; AMDGPU-NEXT: entry: 612; AMDGPU-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 613; AMDGPU-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 614; AMDGPU-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) 615; AMDGPU-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 616; AMDGPU-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 617; AMDGPU: user_code.entry: 618; AMDGPU-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR4]] 619; AMDGPU-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4 620; AMDGPU-NEXT: call void @__omp_outlined__2(i32* noalias nocapture noundef nonnull readonly align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR4]] 621; AMDGPU-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) 622; AMDGPU-NEXT: ret void 623; AMDGPU: worker.exit: 624; AMDGPU-NEXT: ret void 625; 626; NVPTX-LABEL: define {{[^@]+}}@__omp_offloading_14_a34ca11_sequential_loop_to_stack_var_l20 627; NVPTX-SAME: () #[[ATTR0]] { 628; NVPTX-NEXT: entry: 629; NVPTX-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 630; NVPTX-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 631; NVPTX-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) 632; NVPTX-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 633; NVPTX-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 634; NVPTX: user_code.entry: 635; NVPTX-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR4]] 636; NVPTX-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4 637; NVPTX-NEXT: call void @__omp_outlined__2(i32* noalias nocapture noundef nonnull readonly align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR4]] 638; NVPTX-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) 639; NVPTX-NEXT: ret void 640; NVPTX: worker.exit: 641; NVPTX-NEXT: ret void 642; 643; AMDGPU-DISABLED-LABEL: define {{[^@]+}}@__omp_offloading_14_a34ca11_sequential_loop_to_stack_var_l20 644; AMDGPU-DISABLED-SAME: () #[[ATTR0]] { 645; AMDGPU-DISABLED-NEXT: entry: 646; AMDGPU-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR:%.*]] = alloca i8*, align 8, addrspace(5) 647; AMDGPU-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 648; AMDGPU-DISABLED-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 649; AMDGPU-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 650; AMDGPU-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 1, i1 false, i1 true) 651; AMDGPU-DISABLED-NEXT: [[BLOCK_HW_SIZE:%.*]] = call i32 @__kmpc_get_hardware_num_threads_in_block() 652; AMDGPU-DISABLED-NEXT: [[WARP_SIZE:%.*]] = call i32 @__kmpc_get_warp_size() 653; AMDGPU-DISABLED-NEXT: [[BLOCK_SIZE:%.*]] = sub i32 [[BLOCK_HW_SIZE]], [[WARP_SIZE]] 654; AMDGPU-DISABLED-NEXT: [[THREAD_IS_MAIN_OR_WORKER:%.*]] = icmp slt i32 [[TMP0]], [[BLOCK_SIZE]] 655; AMDGPU-DISABLED-NEXT: br i1 [[THREAD_IS_MAIN_OR_WORKER]], label [[IS_WORKER_CHECK:%.*]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]] 656; AMDGPU-DISABLED: is_worker_check: 657; AMDGPU-DISABLED-NEXT: [[THREAD_IS_WORKER:%.*]] = icmp ne i32 [[TMP0]], -1 658; AMDGPU-DISABLED-NEXT: br i1 [[THREAD_IS_WORKER]], label [[WORKER_STATE_MACHINE_BEGIN:%.*]], label [[THREAD_USER_CODE_CHECK:%.*]] 659; AMDGPU-DISABLED: worker_state_machine.begin: 660; AMDGPU-DISABLED-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 661; AMDGPU-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR_GENERIC:%.*]] = addrspacecast i8* addrspace(5)* [[WORKER_WORK_FN_ADDR]] to i8** 662; AMDGPU-DISABLED-NEXT: [[WORKER_IS_ACTIVE:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORKER_WORK_FN_ADDR_GENERIC]]) 663; AMDGPU-DISABLED-NEXT: [[WORKER_WORK_FN:%.*]] = load i8*, i8** [[WORKER_WORK_FN_ADDR_GENERIC]], align 8 664; AMDGPU-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)* 665; AMDGPU-DISABLED-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null 666; AMDGPU-DISABLED-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]] 667; AMDGPU-DISABLED: worker_state_machine.finished: 668; AMDGPU-DISABLED-NEXT: ret void 669; AMDGPU-DISABLED: worker_state_machine.is_active.check: 670; AMDGPU-DISABLED-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]] 671; AMDGPU-DISABLED: worker_state_machine.parallel_region.check: 672; AMDGPU-DISABLED-NEXT: [[WORKER_CHECK_PARALLEL_REGION:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], bitcast (i8* @__omp_outlined__3_wrapper.ID to void (i16, i32)*) 673; AMDGPU-DISABLED-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_FALLBACK_EXECUTE:%.*]] 674; AMDGPU-DISABLED: worker_state_machine.parallel_region.execute: 675; AMDGPU-DISABLED-NEXT: call void @__omp_outlined__3_wrapper(i16 0, i32 [[TMP0]]) 676; AMDGPU-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]] 677; AMDGPU-DISABLED: worker_state_machine.parallel_region.fallback.execute: 678; AMDGPU-DISABLED-NEXT: call void [[WORKER_WORK_FN_ADDR_CAST]](i16 0, i32 [[TMP0]]) 679; AMDGPU-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]] 680; AMDGPU-DISABLED: worker_state_machine.parallel_region.end: 681; AMDGPU-DISABLED-NEXT: call void @__kmpc_kernel_end_parallel() 682; AMDGPU-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_DONE_BARRIER]] 683; AMDGPU-DISABLED: worker_state_machine.done.barrier: 684; AMDGPU-DISABLED-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 685; AMDGPU-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_BEGIN]] 686; AMDGPU-DISABLED: thread.user_code.check: 687; AMDGPU-DISABLED-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 688; AMDGPU-DISABLED-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 689; AMDGPU-DISABLED: user_code.entry: 690; AMDGPU-DISABLED-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR4]] 691; AMDGPU-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4 692; AMDGPU-DISABLED-NEXT: call void @__omp_outlined__2(i32* noalias nocapture noundef nonnull readonly align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR4]] 693; AMDGPU-DISABLED-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true) 694; AMDGPU-DISABLED-NEXT: ret void 695; AMDGPU-DISABLED: worker.exit: 696; AMDGPU-DISABLED-NEXT: ret void 697; 698; NVPTX-DISABLED-LABEL: define {{[^@]+}}@__omp_offloading_14_a34ca11_sequential_loop_to_stack_var_l20 699; NVPTX-DISABLED-SAME: () #[[ATTR0]] { 700; NVPTX-DISABLED-NEXT: entry: 701; NVPTX-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR:%.*]] = alloca i8*, align 8 702; NVPTX-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 703; NVPTX-DISABLED-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 704; NVPTX-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 705; NVPTX-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 1, i1 false, i1 true) 706; NVPTX-DISABLED-NEXT: [[BLOCK_HW_SIZE:%.*]] = call i32 @__kmpc_get_hardware_num_threads_in_block() 707; NVPTX-DISABLED-NEXT: [[WARP_SIZE:%.*]] = call i32 @__kmpc_get_warp_size() 708; NVPTX-DISABLED-NEXT: [[BLOCK_SIZE:%.*]] = sub i32 [[BLOCK_HW_SIZE]], [[WARP_SIZE]] 709; NVPTX-DISABLED-NEXT: [[THREAD_IS_MAIN_OR_WORKER:%.*]] = icmp slt i32 [[TMP0]], [[BLOCK_SIZE]] 710; NVPTX-DISABLED-NEXT: br i1 [[THREAD_IS_MAIN_OR_WORKER]], label [[IS_WORKER_CHECK:%.*]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]] 711; NVPTX-DISABLED: is_worker_check: 712; NVPTX-DISABLED-NEXT: [[THREAD_IS_WORKER:%.*]] = icmp ne i32 [[TMP0]], -1 713; NVPTX-DISABLED-NEXT: br i1 [[THREAD_IS_WORKER]], label [[WORKER_STATE_MACHINE_BEGIN:%.*]], label [[THREAD_USER_CODE_CHECK:%.*]] 714; NVPTX-DISABLED: worker_state_machine.begin: 715; NVPTX-DISABLED-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 716; NVPTX-DISABLED-NEXT: [[WORKER_IS_ACTIVE:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORKER_WORK_FN_ADDR]]) 717; NVPTX-DISABLED-NEXT: [[WORKER_WORK_FN:%.*]] = load i8*, i8** [[WORKER_WORK_FN_ADDR]], align 8 718; NVPTX-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)* 719; NVPTX-DISABLED-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null 720; NVPTX-DISABLED-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]] 721; NVPTX-DISABLED: worker_state_machine.finished: 722; NVPTX-DISABLED-NEXT: ret void 723; NVPTX-DISABLED: worker_state_machine.is_active.check: 724; NVPTX-DISABLED-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]] 725; NVPTX-DISABLED: worker_state_machine.parallel_region.check: 726; NVPTX-DISABLED-NEXT: [[WORKER_CHECK_PARALLEL_REGION:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], bitcast (i8* @__omp_outlined__3_wrapper.ID to void (i16, i32)*) 727; NVPTX-DISABLED-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_FALLBACK_EXECUTE:%.*]] 728; NVPTX-DISABLED: worker_state_machine.parallel_region.execute: 729; NVPTX-DISABLED-NEXT: call void @__omp_outlined__3_wrapper(i16 0, i32 [[TMP0]]) 730; NVPTX-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]] 731; NVPTX-DISABLED: worker_state_machine.parallel_region.fallback.execute: 732; NVPTX-DISABLED-NEXT: call void [[WORKER_WORK_FN_ADDR_CAST]](i16 0, i32 [[TMP0]]) 733; NVPTX-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]] 734; NVPTX-DISABLED: worker_state_machine.parallel_region.end: 735; NVPTX-DISABLED-NEXT: call void @__kmpc_kernel_end_parallel() 736; NVPTX-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_DONE_BARRIER]] 737; NVPTX-DISABLED: worker_state_machine.done.barrier: 738; NVPTX-DISABLED-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 739; NVPTX-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_BEGIN]] 740; NVPTX-DISABLED: thread.user_code.check: 741; NVPTX-DISABLED-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 742; NVPTX-DISABLED-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 743; NVPTX-DISABLED: user_code.entry: 744; NVPTX-DISABLED-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR4]] 745; NVPTX-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4 746; NVPTX-DISABLED-NEXT: call void @__omp_outlined__2(i32* noalias nocapture noundef nonnull readonly align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR4]] 747; NVPTX-DISABLED-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true) 748; NVPTX-DISABLED-NEXT: ret void 749; NVPTX-DISABLED: worker.exit: 750; NVPTX-DISABLED-NEXT: ret void 751; 752entry: 753 %.zero.addr = alloca i32, align 4 754 %.threadid_temp. = alloca i32, align 4 755 store i32 0, i32* %.zero.addr, align 4 756 %0 = call i32 @__kmpc_target_init(%struct.ident_t* @1, i8 1, i1 true, i1 true) 757 %exec_user_code = icmp eq i32 %0, -1 758 br i1 %exec_user_code, label %user_code.entry, label %worker.exit 759 760user_code.entry: ; preds = %entry 761 %1 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @1) 762 store i32 %1, i32* %.threadid_temp., align 4 763 call void @__omp_outlined__2(i32* %.threadid_temp., i32* %.zero.addr) #3 764 call void @__kmpc_target_deinit(%struct.ident_t* @1, i8 1, i1 true) 765 ret void 766 767worker.exit: ; preds = %entry 768 ret void 769} 770 771define internal void @__omp_outlined__2(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 { 772; AMDGPU-LABEL: define {{[^@]+}}@__omp_outlined__2 773; AMDGPU-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]] { 774; AMDGPU-NEXT: entry: 775; AMDGPU-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 776; AMDGPU-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 777; AMDGPU-NEXT: [[I:%.*]] = alloca i32, align 4 778; AMDGPU-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8 779; AMDGPU-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 780; AMDGPU-NEXT: [[TMP0:%.*]] = alloca i8, i64 4, align 1 781; AMDGPU-NEXT: [[X_ON_STACK:%.*]] = bitcast i8* [[TMP0]] to i32* 782; AMDGPU-NEXT: call void @use(i32* nocapture [[X_ON_STACK]]) #[[ATTR7]] 783; AMDGPU-NEXT: store i32 0, i32* [[I]], align 4 784; AMDGPU-NEXT: br label [[FOR_COND:%.*]] 785; AMDGPU: for.cond: 786; AMDGPU-NEXT: [[TMP1:%.*]] = load i32, i32* [[I]], align 4 787; AMDGPU-NEXT: [[CMP:%.*]] = icmp slt i32 [[TMP1]], 100 788; AMDGPU-NEXT: br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]] 789; AMDGPU: for.body: 790; AMDGPU-NEXT: [[TMP2:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4 791; AMDGPU-NEXT: [[TMP3:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 792; AMDGPU-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP2]], 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 [[TMP3]], i64 noundef 0) 793; AMDGPU-NEXT: br label [[FOR_INC:%.*]] 794; AMDGPU: for.inc: 795; AMDGPU-NEXT: [[TMP4:%.*]] = load i32, i32* [[I]], align 4 796; AMDGPU-NEXT: [[INC:%.*]] = add nsw i32 [[TMP4]], 1 797; AMDGPU-NEXT: store i32 [[INC]], i32* [[I]], align 4 798; AMDGPU-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP15:![0-9]+]] 799; AMDGPU: for.end: 800; AMDGPU-NEXT: call void @spmd_amenable() #[[ATTR7]] 801; AMDGPU-NEXT: ret void 802; 803; NVPTX-LABEL: define {{[^@]+}}@__omp_outlined__2 804; NVPTX-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]] { 805; NVPTX-NEXT: entry: 806; NVPTX-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 807; NVPTX-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 808; NVPTX-NEXT: [[I:%.*]] = alloca i32, align 4 809; NVPTX-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8 810; NVPTX-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 811; NVPTX-NEXT: [[TMP0:%.*]] = alloca i8, i64 4, align 1 812; NVPTX-NEXT: [[X_ON_STACK:%.*]] = bitcast i8* [[TMP0]] to i32* 813; NVPTX-NEXT: call void @use(i32* nocapture [[X_ON_STACK]]) #[[ATTR7]] 814; NVPTX-NEXT: store i32 0, i32* [[I]], align 4 815; NVPTX-NEXT: br label [[FOR_COND:%.*]] 816; NVPTX: for.cond: 817; NVPTX-NEXT: [[TMP1:%.*]] = load i32, i32* [[I]], align 4 818; NVPTX-NEXT: [[CMP:%.*]] = icmp slt i32 [[TMP1]], 100 819; NVPTX-NEXT: br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]] 820; NVPTX: for.body: 821; NVPTX-NEXT: [[TMP2:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4 822; NVPTX-NEXT: [[TMP3:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 823; NVPTX-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP2]], 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 [[TMP3]], i64 noundef 0) 824; NVPTX-NEXT: br label [[FOR_INC:%.*]] 825; NVPTX: for.inc: 826; NVPTX-NEXT: [[TMP4:%.*]] = load i32, i32* [[I]], align 4 827; NVPTX-NEXT: [[INC:%.*]] = add nsw i32 [[TMP4]], 1 828; NVPTX-NEXT: store i32 [[INC]], i32* [[I]], align 4 829; NVPTX-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP15:![0-9]+]] 830; NVPTX: for.end: 831; NVPTX-NEXT: call void @spmd_amenable() #[[ATTR7]] 832; NVPTX-NEXT: ret void 833; 834; AMDGPU-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__2 835; AMDGPU-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]] { 836; AMDGPU-DISABLED-NEXT: entry: 837; AMDGPU-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 838; AMDGPU-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 839; AMDGPU-DISABLED-NEXT: [[I:%.*]] = alloca i32, align 4 840; AMDGPU-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8 841; AMDGPU-DISABLED-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 842; AMDGPU-DISABLED-NEXT: [[TMP0:%.*]] = alloca i8, i64 4, align 1 843; AMDGPU-DISABLED-NEXT: [[X_ON_STACK:%.*]] = bitcast i8* [[TMP0]] to i32* 844; AMDGPU-DISABLED-NEXT: call void @use(i32* nocapture [[X_ON_STACK]]) #[[ATTR7]] 845; AMDGPU-DISABLED-NEXT: store i32 0, i32* [[I]], align 4 846; AMDGPU-DISABLED-NEXT: br label [[FOR_COND:%.*]] 847; AMDGPU-DISABLED: for.cond: 848; AMDGPU-DISABLED-NEXT: [[TMP1:%.*]] = load i32, i32* [[I]], align 4 849; AMDGPU-DISABLED-NEXT: [[CMP:%.*]] = icmp slt i32 [[TMP1]], 100 850; AMDGPU-DISABLED-NEXT: br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]] 851; AMDGPU-DISABLED: for.body: 852; AMDGPU-DISABLED-NEXT: [[TMP2:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4 853; AMDGPU-DISABLED-NEXT: [[TMP3:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 854; AMDGPU-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP2]], 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 [[TMP3]], i64 noundef 0) 855; AMDGPU-DISABLED-NEXT: br label [[FOR_INC:%.*]] 856; AMDGPU-DISABLED: for.inc: 857; AMDGPU-DISABLED-NEXT: [[TMP4:%.*]] = load i32, i32* [[I]], align 4 858; AMDGPU-DISABLED-NEXT: [[INC:%.*]] = add nsw i32 [[TMP4]], 1 859; AMDGPU-DISABLED-NEXT: store i32 [[INC]], i32* [[I]], align 4 860; AMDGPU-DISABLED-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP15:![0-9]+]] 861; AMDGPU-DISABLED: for.end: 862; AMDGPU-DISABLED-NEXT: call void @spmd_amenable() #[[ATTR7]] 863; AMDGPU-DISABLED-NEXT: ret void 864; 865; NVPTX-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__2 866; NVPTX-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]] { 867; NVPTX-DISABLED-NEXT: entry: 868; NVPTX-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 869; NVPTX-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 870; NVPTX-DISABLED-NEXT: [[I:%.*]] = alloca i32, align 4 871; NVPTX-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8 872; NVPTX-DISABLED-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 873; NVPTX-DISABLED-NEXT: [[TMP0:%.*]] = alloca i8, i64 4, align 1 874; NVPTX-DISABLED-NEXT: [[X_ON_STACK:%.*]] = bitcast i8* [[TMP0]] to i32* 875; NVPTX-DISABLED-NEXT: call void @use(i32* nocapture [[X_ON_STACK]]) #[[ATTR7]] 876; NVPTX-DISABLED-NEXT: store i32 0, i32* [[I]], align 4 877; NVPTX-DISABLED-NEXT: br label [[FOR_COND:%.*]] 878; NVPTX-DISABLED: for.cond: 879; NVPTX-DISABLED-NEXT: [[TMP1:%.*]] = load i32, i32* [[I]], align 4 880; NVPTX-DISABLED-NEXT: [[CMP:%.*]] = icmp slt i32 [[TMP1]], 100 881; NVPTX-DISABLED-NEXT: br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]] 882; NVPTX-DISABLED: for.body: 883; NVPTX-DISABLED-NEXT: [[TMP2:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4 884; NVPTX-DISABLED-NEXT: [[TMP3:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 885; NVPTX-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP2]], 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 [[TMP3]], i64 noundef 0) 886; NVPTX-DISABLED-NEXT: br label [[FOR_INC:%.*]] 887; NVPTX-DISABLED: for.inc: 888; NVPTX-DISABLED-NEXT: [[TMP4:%.*]] = load i32, i32* [[I]], align 4 889; NVPTX-DISABLED-NEXT: [[INC:%.*]] = add nsw i32 [[TMP4]], 1 890; NVPTX-DISABLED-NEXT: store i32 [[INC]], i32* [[I]], align 4 891; NVPTX-DISABLED-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP15:![0-9]+]] 892; NVPTX-DISABLED: for.end: 893; NVPTX-DISABLED-NEXT: call void @spmd_amenable() #[[ATTR7]] 894; NVPTX-DISABLED-NEXT: ret void 895; 896entry: 897 %.global_tid..addr = alloca i32*, align 8 898 %.bound_tid..addr = alloca i32*, align 8 899 %i = alloca i32, align 4 900 %captured_vars_addrs = alloca [0 x i8*], align 8 901 store i32* %.global_tid., i32** %.global_tid..addr, align 8 902 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8 903 %x = call i8* @__kmpc_alloc_shared(i64 4) 904 %x_on_stack = bitcast i8* %x to i32* 905 call void @use(i32* nocapture %x_on_stack) #4 906 store i32 0, i32* %i, align 4 907 br label %for.cond 908 909for.cond: ; preds = %for.inc, %entry 910 %0 = load i32, i32* %i, align 4 911 %cmp = icmp slt i32 %0, 100 912 br i1 %cmp, label %for.body, label %for.end 913 914for.body: ; preds = %for.cond 915 %1 = load i32*, i32** %.global_tid..addr, align 8 916 %2 = load i32, i32* %1, align 4 917 %3 = bitcast [0 x i8*]* %captured_vars_addrs to i8** 918 call void @__kmpc_parallel_51(%struct.ident_t* @1, i32 %2, 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) 919 br label %for.inc 920 921for.inc: ; preds = %for.body 922 %4 = load i32, i32* %i, align 4 923 %inc = add nsw i32 %4, 1 924 store i32 %inc, i32* %i, align 4 925 br label %for.cond, !llvm.loop !15 926 927for.end: ; preds = %for.cond 928 call void @spmd_amenable() #4 929 call void @__kmpc_free_shared(i8* %x, i64 4) 930 ret void 931} 932 933declare i8* @__kmpc_alloc_shared(i64) #3 934 935declare void @use(i32* nocapture) 936 937define internal void @__omp_outlined__3(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 { 938; 939; 940; AMDGPU-LABEL: define {{[^@]+}}@__omp_outlined__3 941; AMDGPU-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] { 942; AMDGPU-NEXT: entry: 943; AMDGPU-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 944; AMDGPU-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 945; AMDGPU-NEXT: call void @unknown() #[[ATTR8]] 946; AMDGPU-NEXT: ret void 947; 948; NVPTX-LABEL: define {{[^@]+}}@__omp_outlined__3 949; NVPTX-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] { 950; NVPTX-NEXT: entry: 951; NVPTX-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 952; NVPTX-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 953; NVPTX-NEXT: call void @unknown() #[[ATTR8]] 954; NVPTX-NEXT: ret void 955; 956; AMDGPU-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__3 957; AMDGPU-DISABLED-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] { 958; AMDGPU-DISABLED-NEXT: entry: 959; AMDGPU-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 960; AMDGPU-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 961; AMDGPU-DISABLED-NEXT: call void @unknown() #[[ATTR8]] 962; AMDGPU-DISABLED-NEXT: ret void 963; 964; NVPTX-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__3 965; NVPTX-DISABLED-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] { 966; NVPTX-DISABLED-NEXT: entry: 967; NVPTX-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 968; NVPTX-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 969; NVPTX-DISABLED-NEXT: call void @unknown() #[[ATTR8]] 970; NVPTX-DISABLED-NEXT: ret void 971; 972entry: 973 %.global_tid..addr = alloca i32*, align 8 974 %.bound_tid..addr = alloca i32*, align 8 975 store i32* %.global_tid., i32** %.global_tid..addr, align 8 976 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8 977 call void @unknown() #5 978 ret void 979} 980 981define internal void @__omp_outlined__3_wrapper(i16 zeroext %0, i32 %1) #0 { 982; 983; 984; AMDGPU-LABEL: define {{[^@]+}}@__omp_outlined__3_wrapper 985; AMDGPU-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 986; AMDGPU-NEXT: entry: 987; AMDGPU-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 988; AMDGPU-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 989; AMDGPU-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 990; AMDGPU-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8 991; AMDGPU-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 992; AMDGPU-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 993; AMDGPU-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 994; AMDGPU-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 995; AMDGPU-NEXT: call void @__omp_outlined__3(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR4]] 996; AMDGPU-NEXT: ret void 997; 998; NVPTX-LABEL: define {{[^@]+}}@__omp_outlined__3_wrapper 999; NVPTX-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 1000; NVPTX-NEXT: entry: 1001; NVPTX-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 1002; NVPTX-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 1003; NVPTX-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1004; NVPTX-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8 1005; NVPTX-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 1006; NVPTX-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 1007; NVPTX-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 1008; NVPTX-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 1009; NVPTX-NEXT: call void @__omp_outlined__3(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR4]] 1010; NVPTX-NEXT: ret void 1011; 1012; AMDGPU-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__3_wrapper 1013; AMDGPU-DISABLED-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 1014; AMDGPU-DISABLED-NEXT: entry: 1015; AMDGPU-DISABLED-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 1016; AMDGPU-DISABLED-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 1017; AMDGPU-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1018; AMDGPU-DISABLED-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8 1019; AMDGPU-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 1020; AMDGPU-DISABLED-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 1021; AMDGPU-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 1022; AMDGPU-DISABLED-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 1023; AMDGPU-DISABLED-NEXT: call void @__omp_outlined__3(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR4]] 1024; AMDGPU-DISABLED-NEXT: ret void 1025; 1026; NVPTX-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__3_wrapper 1027; NVPTX-DISABLED-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 1028; NVPTX-DISABLED-NEXT: entry: 1029; NVPTX-DISABLED-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 1030; NVPTX-DISABLED-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 1031; NVPTX-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1032; NVPTX-DISABLED-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8 1033; NVPTX-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 1034; NVPTX-DISABLED-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 1035; NVPTX-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 1036; NVPTX-DISABLED-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 1037; NVPTX-DISABLED-NEXT: call void @__omp_outlined__3(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR4]] 1038; NVPTX-DISABLED-NEXT: ret void 1039; 1040entry: 1041 %.addr = alloca i16, align 2 1042 %.addr1 = alloca i32, align 4 1043 %.zero.addr = alloca i32, align 4 1044 %global_args = alloca i8**, align 8 1045 store i32 0, i32* %.zero.addr, align 4 1046 store i16 %0, i16* %.addr, align 2 1047 store i32 %1, i32* %.addr1, align 4 1048 call void @__kmpc_get_shared_variables(i8*** %global_args) 1049 call void @__omp_outlined__3(i32* %.addr1, i32* %.zero.addr) #3 1050 ret void 1051} 1052 1053declare void @__kmpc_free_shared(i8* nocapture, i64) #3 1054 1055define weak void @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_l35() #0 { 1056; 1057; 1058; AMDGPU-LABEL: define {{[^@]+}}@__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_l35 1059; AMDGPU-SAME: () #[[ATTR0]] { 1060; AMDGPU-NEXT: entry: 1061; AMDGPU-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1062; AMDGPU-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 1063; AMDGPU-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) 1064; AMDGPU-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 1065; AMDGPU-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 1066; AMDGPU: user_code.entry: 1067; AMDGPU-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR4]] 1068; AMDGPU-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4 1069; AMDGPU-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]]) #[[ATTR4]] 1070; AMDGPU-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) 1071; AMDGPU-NEXT: ret void 1072; AMDGPU: worker.exit: 1073; AMDGPU-NEXT: ret void 1074; 1075; NVPTX-LABEL: define {{[^@]+}}@__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_l35 1076; NVPTX-SAME: () #[[ATTR0]] { 1077; NVPTX-NEXT: entry: 1078; NVPTX-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1079; NVPTX-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 1080; NVPTX-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) 1081; NVPTX-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 1082; NVPTX-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 1083; NVPTX: user_code.entry: 1084; NVPTX-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR4]] 1085; NVPTX-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4 1086; NVPTX-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]]) #[[ATTR4]] 1087; NVPTX-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) 1088; NVPTX-NEXT: ret void 1089; NVPTX: worker.exit: 1090; NVPTX-NEXT: ret void 1091; 1092; AMDGPU-DISABLED-LABEL: define {{[^@]+}}@__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_l35 1093; AMDGPU-DISABLED-SAME: () #[[ATTR0]] { 1094; AMDGPU-DISABLED-NEXT: entry: 1095; AMDGPU-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR:%.*]] = alloca i8*, align 8, addrspace(5) 1096; AMDGPU-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1097; AMDGPU-DISABLED-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 1098; AMDGPU-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 1099; AMDGPU-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 1, i1 false, i1 true) 1100; AMDGPU-DISABLED-NEXT: [[BLOCK_HW_SIZE:%.*]] = call i32 @__kmpc_get_hardware_num_threads_in_block() 1101; AMDGPU-DISABLED-NEXT: [[WARP_SIZE:%.*]] = call i32 @__kmpc_get_warp_size() 1102; AMDGPU-DISABLED-NEXT: [[BLOCK_SIZE:%.*]] = sub i32 [[BLOCK_HW_SIZE]], [[WARP_SIZE]] 1103; AMDGPU-DISABLED-NEXT: [[THREAD_IS_MAIN_OR_WORKER:%.*]] = icmp slt i32 [[TMP0]], [[BLOCK_SIZE]] 1104; AMDGPU-DISABLED-NEXT: br i1 [[THREAD_IS_MAIN_OR_WORKER]], label [[IS_WORKER_CHECK:%.*]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]] 1105; AMDGPU-DISABLED: is_worker_check: 1106; AMDGPU-DISABLED-NEXT: [[THREAD_IS_WORKER:%.*]] = icmp ne i32 [[TMP0]], -1 1107; AMDGPU-DISABLED-NEXT: br i1 [[THREAD_IS_WORKER]], label [[WORKER_STATE_MACHINE_BEGIN:%.*]], label [[THREAD_USER_CODE_CHECK:%.*]] 1108; AMDGPU-DISABLED: worker_state_machine.begin: 1109; AMDGPU-DISABLED-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 1110; AMDGPU-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR_GENERIC:%.*]] = addrspacecast i8* addrspace(5)* [[WORKER_WORK_FN_ADDR]] to i8** 1111; AMDGPU-DISABLED-NEXT: [[WORKER_IS_ACTIVE:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORKER_WORK_FN_ADDR_GENERIC]]) 1112; AMDGPU-DISABLED-NEXT: [[WORKER_WORK_FN:%.*]] = load i8*, i8** [[WORKER_WORK_FN_ADDR_GENERIC]], align 8 1113; AMDGPU-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)* 1114; AMDGPU-DISABLED-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null 1115; AMDGPU-DISABLED-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]] 1116; AMDGPU-DISABLED: worker_state_machine.finished: 1117; AMDGPU-DISABLED-NEXT: ret void 1118; AMDGPU-DISABLED: worker_state_machine.is_active.check: 1119; AMDGPU-DISABLED-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]] 1120; AMDGPU-DISABLED: worker_state_machine.parallel_region.check: 1121; AMDGPU-DISABLED-NEXT: [[WORKER_CHECK_PARALLEL_REGION:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], bitcast (i8* @__omp_outlined__5_wrapper.ID to void (i16, i32)*) 1122; AMDGPU-DISABLED-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_FALLBACK_EXECUTE:%.*]] 1123; AMDGPU-DISABLED: worker_state_machine.parallel_region.execute: 1124; AMDGPU-DISABLED-NEXT: call void @__omp_outlined__5_wrapper(i16 0, i32 [[TMP0]]) 1125; AMDGPU-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]] 1126; AMDGPU-DISABLED: worker_state_machine.parallel_region.fallback.execute: 1127; AMDGPU-DISABLED-NEXT: call void [[WORKER_WORK_FN_ADDR_CAST]](i16 0, i32 [[TMP0]]) 1128; AMDGPU-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]] 1129; AMDGPU-DISABLED: worker_state_machine.parallel_region.end: 1130; AMDGPU-DISABLED-NEXT: call void @__kmpc_kernel_end_parallel() 1131; AMDGPU-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_DONE_BARRIER]] 1132; AMDGPU-DISABLED: worker_state_machine.done.barrier: 1133; AMDGPU-DISABLED-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 1134; AMDGPU-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_BEGIN]] 1135; AMDGPU-DISABLED: thread.user_code.check: 1136; AMDGPU-DISABLED-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 1137; AMDGPU-DISABLED-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 1138; AMDGPU-DISABLED: user_code.entry: 1139; AMDGPU-DISABLED-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR4]] 1140; AMDGPU-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4 1141; AMDGPU-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]]) #[[ATTR4]] 1142; AMDGPU-DISABLED-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true) 1143; AMDGPU-DISABLED-NEXT: ret void 1144; AMDGPU-DISABLED: worker.exit: 1145; AMDGPU-DISABLED-NEXT: ret void 1146; 1147; NVPTX-DISABLED-LABEL: define {{[^@]+}}@__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_l35 1148; NVPTX-DISABLED-SAME: () #[[ATTR0]] { 1149; NVPTX-DISABLED-NEXT: entry: 1150; NVPTX-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR:%.*]] = alloca i8*, align 8 1151; NVPTX-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1152; NVPTX-DISABLED-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 1153; NVPTX-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 1154; NVPTX-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 1, i1 false, i1 true) 1155; NVPTX-DISABLED-NEXT: [[BLOCK_HW_SIZE:%.*]] = call i32 @__kmpc_get_hardware_num_threads_in_block() 1156; NVPTX-DISABLED-NEXT: [[WARP_SIZE:%.*]] = call i32 @__kmpc_get_warp_size() 1157; NVPTX-DISABLED-NEXT: [[BLOCK_SIZE:%.*]] = sub i32 [[BLOCK_HW_SIZE]], [[WARP_SIZE]] 1158; NVPTX-DISABLED-NEXT: [[THREAD_IS_MAIN_OR_WORKER:%.*]] = icmp slt i32 [[TMP0]], [[BLOCK_SIZE]] 1159; NVPTX-DISABLED-NEXT: br i1 [[THREAD_IS_MAIN_OR_WORKER]], label [[IS_WORKER_CHECK:%.*]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]] 1160; NVPTX-DISABLED: is_worker_check: 1161; NVPTX-DISABLED-NEXT: [[THREAD_IS_WORKER:%.*]] = icmp ne i32 [[TMP0]], -1 1162; NVPTX-DISABLED-NEXT: br i1 [[THREAD_IS_WORKER]], label [[WORKER_STATE_MACHINE_BEGIN:%.*]], label [[THREAD_USER_CODE_CHECK:%.*]] 1163; NVPTX-DISABLED: worker_state_machine.begin: 1164; NVPTX-DISABLED-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 1165; NVPTX-DISABLED-NEXT: [[WORKER_IS_ACTIVE:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORKER_WORK_FN_ADDR]]) 1166; NVPTX-DISABLED-NEXT: [[WORKER_WORK_FN:%.*]] = load i8*, i8** [[WORKER_WORK_FN_ADDR]], align 8 1167; NVPTX-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)* 1168; NVPTX-DISABLED-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null 1169; NVPTX-DISABLED-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]] 1170; NVPTX-DISABLED: worker_state_machine.finished: 1171; NVPTX-DISABLED-NEXT: ret void 1172; NVPTX-DISABLED: worker_state_machine.is_active.check: 1173; NVPTX-DISABLED-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]] 1174; NVPTX-DISABLED: worker_state_machine.parallel_region.check: 1175; NVPTX-DISABLED-NEXT: [[WORKER_CHECK_PARALLEL_REGION:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], bitcast (i8* @__omp_outlined__5_wrapper.ID to void (i16, i32)*) 1176; NVPTX-DISABLED-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_FALLBACK_EXECUTE:%.*]] 1177; NVPTX-DISABLED: worker_state_machine.parallel_region.execute: 1178; NVPTX-DISABLED-NEXT: call void @__omp_outlined__5_wrapper(i16 0, i32 [[TMP0]]) 1179; NVPTX-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]] 1180; NVPTX-DISABLED: worker_state_machine.parallel_region.fallback.execute: 1181; NVPTX-DISABLED-NEXT: call void [[WORKER_WORK_FN_ADDR_CAST]](i16 0, i32 [[TMP0]]) 1182; NVPTX-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]] 1183; NVPTX-DISABLED: worker_state_machine.parallel_region.end: 1184; NVPTX-DISABLED-NEXT: call void @__kmpc_kernel_end_parallel() 1185; NVPTX-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_DONE_BARRIER]] 1186; NVPTX-DISABLED: worker_state_machine.done.barrier: 1187; NVPTX-DISABLED-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 1188; NVPTX-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_BEGIN]] 1189; NVPTX-DISABLED: thread.user_code.check: 1190; NVPTX-DISABLED-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 1191; NVPTX-DISABLED-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 1192; NVPTX-DISABLED: user_code.entry: 1193; NVPTX-DISABLED-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR4]] 1194; NVPTX-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4 1195; NVPTX-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]]) #[[ATTR4]] 1196; NVPTX-DISABLED-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true) 1197; NVPTX-DISABLED-NEXT: ret void 1198; NVPTX-DISABLED: worker.exit: 1199; NVPTX-DISABLED-NEXT: ret void 1200; 1201entry: 1202 %.zero.addr = alloca i32, align 4 1203 %.threadid_temp. = alloca i32, align 4 1204 store i32 0, i32* %.zero.addr, align 4 1205 %0 = call i32 @__kmpc_target_init(%struct.ident_t* @1, i8 1, i1 true, i1 true) 1206 %exec_user_code = icmp eq i32 %0, -1 1207 br i1 %exec_user_code, label %user_code.entry, label %worker.exit 1208 1209user_code.entry: ; preds = %entry 1210 %1 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @1) 1211 store i32 %1, i32* %.threadid_temp., align 4 1212 call void @__omp_outlined__4(i32* %.threadid_temp., i32* %.zero.addr) #3 1213 call void @__kmpc_target_deinit(%struct.ident_t* @1, i8 1, i1 true) 1214 ret void 1215 1216worker.exit: ; preds = %entry 1217 ret void 1218} 1219 1220define internal void @__omp_outlined__4(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 { 1221; 1222; 1223; AMDGPU-LABEL: define {{[^@]+}}@__omp_outlined__4 1224; AMDGPU-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]] { 1225; AMDGPU-NEXT: entry: 1226; AMDGPU-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1227; AMDGPU-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1228; AMDGPU-NEXT: [[I:%.*]] = alloca i32, align 4 1229; AMDGPU-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 8 1230; AMDGPU-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 1231; AMDGPU-NEXT: store i32 0, i32* [[I]], align 4 1232; AMDGPU-NEXT: br label [[FOR_COND:%.*]] 1233; AMDGPU: for.cond: 1234; AMDGPU-NEXT: [[TMP0:%.*]] = load i32, i32* [[I]], align 4 1235; AMDGPU-NEXT: [[CMP:%.*]] = icmp slt i32 [[TMP0]], 100 1236; AMDGPU-NEXT: br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]] 1237; AMDGPU: for.body: 1238; AMDGPU-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0 1239; AMDGPU-NEXT: store i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([4 x i8], [4 x i8] addrspace(3)* @x, i32 0, i32 0) to i8*), i8** [[TMP1]], align 8 1240; AMDGPU-NEXT: [[TMP2:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4 1241; AMDGPU-NEXT: [[TMP3:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 1242; AMDGPU-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP2]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*, i32*)* @__omp_outlined__5 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__5_wrapper to i8*), i8** noundef [[TMP3]], i64 noundef 1) 1243; AMDGPU-NEXT: br label [[FOR_INC:%.*]] 1244; AMDGPU: for.inc: 1245; AMDGPU-NEXT: [[TMP4:%.*]] = load i32, i32* [[I]], align 4 1246; AMDGPU-NEXT: [[INC:%.*]] = add nsw i32 [[TMP4]], 1 1247; AMDGPU-NEXT: store i32 [[INC]], i32* [[I]], align 4 1248; AMDGPU-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP16:![0-9]+]] 1249; AMDGPU: for.end: 1250; AMDGPU-NEXT: call void @spmd_amenable() #[[ATTR7]] 1251; AMDGPU-NEXT: ret void 1252; 1253; NVPTX-LABEL: define {{[^@]+}}@__omp_outlined__4 1254; NVPTX-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]] { 1255; NVPTX-NEXT: entry: 1256; NVPTX-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1257; NVPTX-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1258; NVPTX-NEXT: [[I:%.*]] = alloca i32, align 4 1259; NVPTX-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 8 1260; NVPTX-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 1261; NVPTX-NEXT: store i32 0, i32* [[I]], align 4 1262; NVPTX-NEXT: br label [[FOR_COND:%.*]] 1263; NVPTX: for.cond: 1264; NVPTX-NEXT: [[TMP0:%.*]] = load i32, i32* [[I]], align 4 1265; NVPTX-NEXT: [[CMP:%.*]] = icmp slt i32 [[TMP0]], 100 1266; NVPTX-NEXT: br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]] 1267; NVPTX: for.body: 1268; NVPTX-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0 1269; NVPTX-NEXT: store i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([4 x i8], [4 x i8] addrspace(3)* @x, i32 0, i32 0) to i8*), i8** [[TMP1]], align 8 1270; NVPTX-NEXT: [[TMP2:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4 1271; NVPTX-NEXT: [[TMP3:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 1272; NVPTX-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP2]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*, i32*)* @__omp_outlined__5 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__5_wrapper to i8*), i8** noundef [[TMP3]], i64 noundef 1) 1273; NVPTX-NEXT: br label [[FOR_INC:%.*]] 1274; NVPTX: for.inc: 1275; NVPTX-NEXT: [[TMP4:%.*]] = load i32, i32* [[I]], align 4 1276; NVPTX-NEXT: [[INC:%.*]] = add nsw i32 [[TMP4]], 1 1277; NVPTX-NEXT: store i32 [[INC]], i32* [[I]], align 4 1278; NVPTX-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP16:![0-9]+]] 1279; NVPTX: for.end: 1280; NVPTX-NEXT: call void @spmd_amenable() #[[ATTR7]] 1281; NVPTX-NEXT: ret void 1282; 1283; AMDGPU-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__4 1284; AMDGPU-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]] { 1285; AMDGPU-DISABLED-NEXT: entry: 1286; AMDGPU-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1287; AMDGPU-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1288; AMDGPU-DISABLED-NEXT: [[I:%.*]] = alloca i32, align 4 1289; AMDGPU-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 8 1290; AMDGPU-DISABLED-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 1291; AMDGPU-DISABLED-NEXT: store i32 0, i32* [[I]], align 4 1292; AMDGPU-DISABLED-NEXT: br label [[FOR_COND:%.*]] 1293; AMDGPU-DISABLED: for.cond: 1294; AMDGPU-DISABLED-NEXT: [[TMP0:%.*]] = load i32, i32* [[I]], align 4 1295; AMDGPU-DISABLED-NEXT: [[CMP:%.*]] = icmp slt i32 [[TMP0]], 100 1296; AMDGPU-DISABLED-NEXT: br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]] 1297; AMDGPU-DISABLED: for.body: 1298; AMDGPU-DISABLED-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0 1299; AMDGPU-DISABLED-NEXT: store i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([4 x i8], [4 x i8] addrspace(3)* @x, i32 0, i32 0) to i8*), i8** [[TMP1]], align 8 1300; AMDGPU-DISABLED-NEXT: [[TMP2:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4 1301; AMDGPU-DISABLED-NEXT: [[TMP3:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 1302; AMDGPU-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP2]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*, i32*)* @__omp_outlined__5 to i8*), i8* noundef @__omp_outlined__5_wrapper.ID, i8** noundef [[TMP3]], i64 noundef 1) 1303; AMDGPU-DISABLED-NEXT: br label [[FOR_INC:%.*]] 1304; AMDGPU-DISABLED: for.inc: 1305; AMDGPU-DISABLED-NEXT: [[TMP4:%.*]] = load i32, i32* [[I]], align 4 1306; AMDGPU-DISABLED-NEXT: [[INC:%.*]] = add nsw i32 [[TMP4]], 1 1307; AMDGPU-DISABLED-NEXT: store i32 [[INC]], i32* [[I]], align 4 1308; AMDGPU-DISABLED-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP16:![0-9]+]] 1309; AMDGPU-DISABLED: for.end: 1310; AMDGPU-DISABLED-NEXT: call void @spmd_amenable() #[[ATTR7]] 1311; AMDGPU-DISABLED-NEXT: ret void 1312; 1313; NVPTX-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__4 1314; NVPTX-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]] { 1315; NVPTX-DISABLED-NEXT: entry: 1316; NVPTX-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1317; NVPTX-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1318; NVPTX-DISABLED-NEXT: [[I:%.*]] = alloca i32, align 4 1319; NVPTX-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 8 1320; NVPTX-DISABLED-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 1321; NVPTX-DISABLED-NEXT: store i32 0, i32* [[I]], align 4 1322; NVPTX-DISABLED-NEXT: br label [[FOR_COND:%.*]] 1323; NVPTX-DISABLED: for.cond: 1324; NVPTX-DISABLED-NEXT: [[TMP0:%.*]] = load i32, i32* [[I]], align 4 1325; NVPTX-DISABLED-NEXT: [[CMP:%.*]] = icmp slt i32 [[TMP0]], 100 1326; NVPTX-DISABLED-NEXT: br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]] 1327; NVPTX-DISABLED: for.body: 1328; NVPTX-DISABLED-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0 1329; NVPTX-DISABLED-NEXT: store i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([4 x i8], [4 x i8] addrspace(3)* @x, i32 0, i32 0) to i8*), i8** [[TMP1]], align 8 1330; NVPTX-DISABLED-NEXT: [[TMP2:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4 1331; NVPTX-DISABLED-NEXT: [[TMP3:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 1332; NVPTX-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP2]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*, i32*)* @__omp_outlined__5 to i8*), i8* noundef @__omp_outlined__5_wrapper.ID, i8** noundef [[TMP3]], i64 noundef 1) 1333; NVPTX-DISABLED-NEXT: br label [[FOR_INC:%.*]] 1334; NVPTX-DISABLED: for.inc: 1335; NVPTX-DISABLED-NEXT: [[TMP4:%.*]] = load i32, i32* [[I]], align 4 1336; NVPTX-DISABLED-NEXT: [[INC:%.*]] = add nsw i32 [[TMP4]], 1 1337; NVPTX-DISABLED-NEXT: store i32 [[INC]], i32* [[I]], align 4 1338; NVPTX-DISABLED-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP16:![0-9]+]] 1339; NVPTX-DISABLED: for.end: 1340; NVPTX-DISABLED-NEXT: call void @spmd_amenable() #[[ATTR7]] 1341; NVPTX-DISABLED-NEXT: ret void 1342; 1343entry: 1344 %.global_tid..addr = alloca i32*, align 8 1345 %.bound_tid..addr = alloca i32*, align 8 1346 %i = alloca i32, align 4 1347 %captured_vars_addrs = alloca [1 x i8*], align 8 1348 store i32* %.global_tid., i32** %.global_tid..addr, align 8 1349 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8 1350 %x = call i8* @__kmpc_alloc_shared(i64 4) 1351 %x_on_stack = bitcast i8* %x to i32* 1352 store i32 0, i32* %i, align 4 1353 br label %for.cond 1354 1355for.cond: ; preds = %for.inc, %entry 1356 %0 = load i32, i32* %i, align 4 1357 %cmp = icmp slt i32 %0, 100 1358 br i1 %cmp, label %for.body, label %for.end 1359 1360for.body: ; preds = %for.cond 1361 %1 = getelementptr inbounds [1 x i8*], [1 x i8*]* %captured_vars_addrs, i64 0, i64 0 1362 %2 = bitcast i32* %x_on_stack to i8* 1363 store i8* %2, i8** %1, align 8 1364 %3 = load i32*, i32** %.global_tid..addr, align 8 1365 %4 = load i32, i32* %3, align 4 1366 %5 = bitcast [1 x i8*]* %captured_vars_addrs to i8** 1367 call void @__kmpc_parallel_51(%struct.ident_t* @1, i32 %4, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i32*)* @__omp_outlined__5 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__5_wrapper to i8*), i8** %5, i64 1) 1368 br label %for.inc 1369 1370for.inc: ; preds = %for.body 1371 %6 = load i32, i32* %i, align 4 1372 %inc = add nsw i32 %6, 1 1373 store i32 %inc, i32* %i, align 4 1374 br label %for.cond, !llvm.loop !16 1375 1376for.end: ; preds = %for.cond 1377 call void @spmd_amenable() #4 1378 call void @__kmpc_free_shared(i8* %x, i64 4) 1379 ret void 1380} 1381 1382define internal void @__omp_outlined__5(i32* noalias %.global_tid., i32* noalias %.bound_tid., i32* nonnull align 4 dereferenceable(4) %x) #0 { 1383; 1384; 1385; AMDGPU-LABEL: define {{[^@]+}}@__omp_outlined__5 1386; AMDGPU-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]], i32* nocapture nofree nonnull align 4 dereferenceable(4) [[X:%.*]]) #[[ATTR0]] { 1387; AMDGPU-NEXT: entry: 1388; AMDGPU-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1389; AMDGPU-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1390; AMDGPU-NEXT: [[X_ADDR:%.*]] = alloca i32*, align 8 1391; AMDGPU-NEXT: store i32* [[X]], i32** [[X_ADDR]], align 8 1392; AMDGPU-NEXT: [[TMP0:%.*]] = load i32, i32* [[X]], align 4 1393; AMDGPU-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1 1394; AMDGPU-NEXT: store i32 [[INC]], i32* [[X]], align 4 1395; AMDGPU-NEXT: call void @unknown() #[[ATTR8]] 1396; AMDGPU-NEXT: ret void 1397; 1398; NVPTX-LABEL: define {{[^@]+}}@__omp_outlined__5 1399; NVPTX-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]], i32* nocapture nofree nonnull align 4 dereferenceable(4) [[X:%.*]]) #[[ATTR0]] { 1400; NVPTX-NEXT: entry: 1401; NVPTX-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1402; NVPTX-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1403; NVPTX-NEXT: [[X_ADDR:%.*]] = alloca i32*, align 8 1404; NVPTX-NEXT: store i32* [[X]], i32** [[X_ADDR]], align 8 1405; NVPTX-NEXT: [[TMP0:%.*]] = load i32, i32* [[X]], align 4 1406; NVPTX-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1 1407; NVPTX-NEXT: store i32 [[INC]], i32* [[X]], align 4 1408; NVPTX-NEXT: call void @unknown() #[[ATTR8]] 1409; NVPTX-NEXT: ret void 1410; 1411; AMDGPU-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__5 1412; AMDGPU-DISABLED-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]], i32* nocapture nofree nonnull align 4 dereferenceable(4) [[X:%.*]]) #[[ATTR0]] { 1413; AMDGPU-DISABLED-NEXT: entry: 1414; AMDGPU-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1415; AMDGPU-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1416; AMDGPU-DISABLED-NEXT: [[X_ADDR:%.*]] = alloca i32*, align 8 1417; AMDGPU-DISABLED-NEXT: store i32* [[X]], i32** [[X_ADDR]], align 8 1418; AMDGPU-DISABLED-NEXT: [[TMP0:%.*]] = load i32, i32* [[X]], align 4 1419; AMDGPU-DISABLED-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1 1420; AMDGPU-DISABLED-NEXT: store i32 [[INC]], i32* [[X]], align 4 1421; AMDGPU-DISABLED-NEXT: call void @unknown() #[[ATTR8]] 1422; AMDGPU-DISABLED-NEXT: ret void 1423; 1424; NVPTX-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__5 1425; NVPTX-DISABLED-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]], i32* nocapture nofree nonnull align 4 dereferenceable(4) [[X:%.*]]) #[[ATTR0]] { 1426; NVPTX-DISABLED-NEXT: entry: 1427; NVPTX-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1428; NVPTX-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1429; NVPTX-DISABLED-NEXT: [[X_ADDR:%.*]] = alloca i32*, align 8 1430; NVPTX-DISABLED-NEXT: store i32* [[X]], i32** [[X_ADDR]], align 8 1431; NVPTX-DISABLED-NEXT: [[TMP0:%.*]] = load i32, i32* [[X]], align 4 1432; NVPTX-DISABLED-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1 1433; NVPTX-DISABLED-NEXT: store i32 [[INC]], i32* [[X]], align 4 1434; NVPTX-DISABLED-NEXT: call void @unknown() #[[ATTR8]] 1435; NVPTX-DISABLED-NEXT: ret void 1436; 1437entry: 1438 %.global_tid..addr = alloca i32*, align 8 1439 %.bound_tid..addr = alloca i32*, align 8 1440 %x.addr = alloca i32*, align 8 1441 store i32* %.global_tid., i32** %.global_tid..addr, align 8 1442 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8 1443 store i32* %x, i32** %x.addr, align 8 1444 %0 = load i32*, i32** %x.addr, align 8 1445 %1 = load i32, i32* %0, align 4 1446 %inc = add nsw i32 %1, 1 1447 store i32 %inc, i32* %0, align 4 1448 call void @unknown() #5 1449 ret void 1450} 1451 1452define internal void @__omp_outlined__5_wrapper(i16 zeroext %0, i32 %1) #0 { 1453; 1454; 1455; AMDGPU-LABEL: define {{[^@]+}}@__omp_outlined__5_wrapper 1456; AMDGPU-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 1457; AMDGPU-NEXT: entry: 1458; AMDGPU-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 1459; AMDGPU-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 1460; AMDGPU-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1461; AMDGPU-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8 1462; AMDGPU-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 1463; AMDGPU-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 1464; AMDGPU-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 1465; AMDGPU-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 1466; AMDGPU-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 8 1467; AMDGPU-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i64 0 1468; AMDGPU-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32** 1469; AMDGPU-NEXT: [[TMP5:%.*]] = load i32*, i32** [[TMP4]], align 8 1470; AMDGPU-NEXT: call void @__omp_outlined__5(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32* [[TMP5]]) #[[ATTR4]] 1471; AMDGPU-NEXT: ret void 1472; 1473; NVPTX-LABEL: define {{[^@]+}}@__omp_outlined__5_wrapper 1474; NVPTX-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 1475; NVPTX-NEXT: entry: 1476; NVPTX-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 1477; NVPTX-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 1478; NVPTX-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1479; NVPTX-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8 1480; NVPTX-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 1481; NVPTX-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 1482; NVPTX-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 1483; NVPTX-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 1484; NVPTX-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 8 1485; NVPTX-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i64 0 1486; NVPTX-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32** 1487; NVPTX-NEXT: [[TMP5:%.*]] = load i32*, i32** [[TMP4]], align 8 1488; NVPTX-NEXT: call void @__omp_outlined__5(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32* [[TMP5]]) #[[ATTR4]] 1489; NVPTX-NEXT: ret void 1490; 1491; AMDGPU-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__5_wrapper 1492; AMDGPU-DISABLED-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 1493; AMDGPU-DISABLED-NEXT: entry: 1494; AMDGPU-DISABLED-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 1495; AMDGPU-DISABLED-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 1496; AMDGPU-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1497; AMDGPU-DISABLED-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8 1498; AMDGPU-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 1499; AMDGPU-DISABLED-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 1500; AMDGPU-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 1501; AMDGPU-DISABLED-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 1502; AMDGPU-DISABLED-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 8 1503; AMDGPU-DISABLED-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i64 0 1504; AMDGPU-DISABLED-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32** 1505; AMDGPU-DISABLED-NEXT: [[TMP5:%.*]] = load i32*, i32** [[TMP4]], align 8 1506; AMDGPU-DISABLED-NEXT: call void @__omp_outlined__5(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32* [[TMP5]]) #[[ATTR4]] 1507; AMDGPU-DISABLED-NEXT: ret void 1508; 1509; NVPTX-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__5_wrapper 1510; NVPTX-DISABLED-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 1511; NVPTX-DISABLED-NEXT: entry: 1512; NVPTX-DISABLED-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 1513; NVPTX-DISABLED-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 1514; NVPTX-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1515; NVPTX-DISABLED-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8 1516; NVPTX-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 1517; NVPTX-DISABLED-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 1518; NVPTX-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 1519; NVPTX-DISABLED-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 1520; NVPTX-DISABLED-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 8 1521; NVPTX-DISABLED-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i64 0 1522; NVPTX-DISABLED-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32** 1523; NVPTX-DISABLED-NEXT: [[TMP5:%.*]] = load i32*, i32** [[TMP4]], align 8 1524; NVPTX-DISABLED-NEXT: call void @__omp_outlined__5(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32* [[TMP5]]) #[[ATTR4]] 1525; NVPTX-DISABLED-NEXT: ret void 1526; 1527entry: 1528 %.addr = alloca i16, align 2 1529 %.addr1 = alloca i32, align 4 1530 %.zero.addr = alloca i32, align 4 1531 %global_args = alloca i8**, align 8 1532 store i32 0, i32* %.zero.addr, align 4 1533 store i16 %0, i16* %.addr, align 2 1534 store i32 %1, i32* %.addr1, align 4 1535 call void @__kmpc_get_shared_variables(i8*** %global_args) 1536 %2 = load i8**, i8*** %global_args, align 8 1537 %3 = getelementptr inbounds i8*, i8** %2, i64 0 1538 %4 = bitcast i8** %3 to i32** 1539 %5 = load i32*, i32** %4, align 8 1540 call void @__omp_outlined__5(i32* %.addr1, i32* %.zero.addr, i32* %5) #3 1541 ret void 1542} 1543 1544define weak void @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_guarded_l50() #0 { 1545; 1546; 1547; AMDGPU-LABEL: define {{[^@]+}}@__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_guarded_l50 1548; AMDGPU-SAME: () #[[ATTR0]] { 1549; AMDGPU-NEXT: entry: 1550; AMDGPU-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1551; AMDGPU-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 1552; AMDGPU-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) 1553; AMDGPU-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 1554; AMDGPU-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 1555; AMDGPU: user_code.entry: 1556; AMDGPU-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR4]] 1557; AMDGPU-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4 1558; AMDGPU-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]]) #[[ATTR4]] 1559; AMDGPU-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) 1560; AMDGPU-NEXT: ret void 1561; AMDGPU: worker.exit: 1562; AMDGPU-NEXT: ret void 1563; 1564; NVPTX-LABEL: define {{[^@]+}}@__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_guarded_l50 1565; NVPTX-SAME: () #[[ATTR0]] { 1566; NVPTX-NEXT: entry: 1567; NVPTX-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1568; NVPTX-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 1569; NVPTX-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) 1570; NVPTX-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 1571; NVPTX-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 1572; NVPTX: user_code.entry: 1573; NVPTX-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR4]] 1574; NVPTX-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4 1575; NVPTX-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]]) #[[ATTR4]] 1576; NVPTX-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) 1577; NVPTX-NEXT: ret void 1578; NVPTX: worker.exit: 1579; NVPTX-NEXT: ret void 1580; 1581; AMDGPU-DISABLED-LABEL: define {{[^@]+}}@__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_guarded_l50 1582; AMDGPU-DISABLED-SAME: () #[[ATTR0]] { 1583; AMDGPU-DISABLED-NEXT: entry: 1584; AMDGPU-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR:%.*]] = alloca i8*, align 8, addrspace(5) 1585; AMDGPU-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1586; AMDGPU-DISABLED-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 1587; AMDGPU-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 1588; AMDGPU-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 1, i1 false, i1 true) 1589; AMDGPU-DISABLED-NEXT: [[BLOCK_HW_SIZE:%.*]] = call i32 @__kmpc_get_hardware_num_threads_in_block() 1590; AMDGPU-DISABLED-NEXT: [[WARP_SIZE:%.*]] = call i32 @__kmpc_get_warp_size() 1591; AMDGPU-DISABLED-NEXT: [[BLOCK_SIZE:%.*]] = sub i32 [[BLOCK_HW_SIZE]], [[WARP_SIZE]] 1592; AMDGPU-DISABLED-NEXT: [[THREAD_IS_MAIN_OR_WORKER:%.*]] = icmp slt i32 [[TMP0]], [[BLOCK_SIZE]] 1593; AMDGPU-DISABLED-NEXT: br i1 [[THREAD_IS_MAIN_OR_WORKER]], label [[IS_WORKER_CHECK:%.*]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]] 1594; AMDGPU-DISABLED: is_worker_check: 1595; AMDGPU-DISABLED-NEXT: [[THREAD_IS_WORKER:%.*]] = icmp ne i32 [[TMP0]], -1 1596; AMDGPU-DISABLED-NEXT: br i1 [[THREAD_IS_WORKER]], label [[WORKER_STATE_MACHINE_BEGIN:%.*]], label [[THREAD_USER_CODE_CHECK:%.*]] 1597; AMDGPU-DISABLED: worker_state_machine.begin: 1598; AMDGPU-DISABLED-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 1599; AMDGPU-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR_GENERIC:%.*]] = addrspacecast i8* addrspace(5)* [[WORKER_WORK_FN_ADDR]] to i8** 1600; AMDGPU-DISABLED-NEXT: [[WORKER_IS_ACTIVE:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORKER_WORK_FN_ADDR_GENERIC]]) 1601; AMDGPU-DISABLED-NEXT: [[WORKER_WORK_FN:%.*]] = load i8*, i8** [[WORKER_WORK_FN_ADDR_GENERIC]], align 8 1602; AMDGPU-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)* 1603; AMDGPU-DISABLED-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null 1604; AMDGPU-DISABLED-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]] 1605; AMDGPU-DISABLED: worker_state_machine.finished: 1606; AMDGPU-DISABLED-NEXT: ret void 1607; AMDGPU-DISABLED: worker_state_machine.is_active.check: 1608; AMDGPU-DISABLED-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]] 1609; AMDGPU-DISABLED: worker_state_machine.parallel_region.check: 1610; AMDGPU-DISABLED-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)*) 1611; AMDGPU-DISABLED-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_FALLBACK_EXECUTE:%.*]] 1612; AMDGPU-DISABLED: worker_state_machine.parallel_region.execute: 1613; AMDGPU-DISABLED-NEXT: call void @__omp_outlined__7_wrapper(i16 0, i32 [[TMP0]]) 1614; AMDGPU-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]] 1615; AMDGPU-DISABLED: worker_state_machine.parallel_region.fallback.execute: 1616; AMDGPU-DISABLED-NEXT: call void [[WORKER_WORK_FN_ADDR_CAST]](i16 0, i32 [[TMP0]]) 1617; AMDGPU-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]] 1618; AMDGPU-DISABLED: worker_state_machine.parallel_region.end: 1619; AMDGPU-DISABLED-NEXT: call void @__kmpc_kernel_end_parallel() 1620; AMDGPU-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_DONE_BARRIER]] 1621; AMDGPU-DISABLED: worker_state_machine.done.barrier: 1622; AMDGPU-DISABLED-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 1623; AMDGPU-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_BEGIN]] 1624; AMDGPU-DISABLED: thread.user_code.check: 1625; AMDGPU-DISABLED-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 1626; AMDGPU-DISABLED-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 1627; AMDGPU-DISABLED: user_code.entry: 1628; AMDGPU-DISABLED-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR4]] 1629; AMDGPU-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4 1630; AMDGPU-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]]) #[[ATTR4]] 1631; AMDGPU-DISABLED-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true) 1632; AMDGPU-DISABLED-NEXT: ret void 1633; AMDGPU-DISABLED: worker.exit: 1634; AMDGPU-DISABLED-NEXT: ret void 1635; 1636; NVPTX-DISABLED-LABEL: define {{[^@]+}}@__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_guarded_l50 1637; NVPTX-DISABLED-SAME: () #[[ATTR0]] { 1638; NVPTX-DISABLED-NEXT: entry: 1639; NVPTX-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR:%.*]] = alloca i8*, align 8 1640; NVPTX-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1641; NVPTX-DISABLED-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 1642; NVPTX-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 1643; NVPTX-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 1, i1 false, i1 true) 1644; NVPTX-DISABLED-NEXT: [[BLOCK_HW_SIZE:%.*]] = call i32 @__kmpc_get_hardware_num_threads_in_block() 1645; NVPTX-DISABLED-NEXT: [[WARP_SIZE:%.*]] = call i32 @__kmpc_get_warp_size() 1646; NVPTX-DISABLED-NEXT: [[BLOCK_SIZE:%.*]] = sub i32 [[BLOCK_HW_SIZE]], [[WARP_SIZE]] 1647; NVPTX-DISABLED-NEXT: [[THREAD_IS_MAIN_OR_WORKER:%.*]] = icmp slt i32 [[TMP0]], [[BLOCK_SIZE]] 1648; NVPTX-DISABLED-NEXT: br i1 [[THREAD_IS_MAIN_OR_WORKER]], label [[IS_WORKER_CHECK:%.*]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]] 1649; NVPTX-DISABLED: is_worker_check: 1650; NVPTX-DISABLED-NEXT: [[THREAD_IS_WORKER:%.*]] = icmp ne i32 [[TMP0]], -1 1651; NVPTX-DISABLED-NEXT: br i1 [[THREAD_IS_WORKER]], label [[WORKER_STATE_MACHINE_BEGIN:%.*]], label [[THREAD_USER_CODE_CHECK:%.*]] 1652; NVPTX-DISABLED: worker_state_machine.begin: 1653; NVPTX-DISABLED-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 1654; NVPTX-DISABLED-NEXT: [[WORKER_IS_ACTIVE:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORKER_WORK_FN_ADDR]]) 1655; NVPTX-DISABLED-NEXT: [[WORKER_WORK_FN:%.*]] = load i8*, i8** [[WORKER_WORK_FN_ADDR]], align 8 1656; NVPTX-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)* 1657; NVPTX-DISABLED-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null 1658; NVPTX-DISABLED-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]] 1659; NVPTX-DISABLED: worker_state_machine.finished: 1660; NVPTX-DISABLED-NEXT: ret void 1661; NVPTX-DISABLED: worker_state_machine.is_active.check: 1662; NVPTX-DISABLED-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]] 1663; NVPTX-DISABLED: worker_state_machine.parallel_region.check: 1664; NVPTX-DISABLED-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)*) 1665; NVPTX-DISABLED-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_FALLBACK_EXECUTE:%.*]] 1666; NVPTX-DISABLED: worker_state_machine.parallel_region.execute: 1667; NVPTX-DISABLED-NEXT: call void @__omp_outlined__7_wrapper(i16 0, i32 [[TMP0]]) 1668; NVPTX-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]] 1669; NVPTX-DISABLED: worker_state_machine.parallel_region.fallback.execute: 1670; NVPTX-DISABLED-NEXT: call void [[WORKER_WORK_FN_ADDR_CAST]](i16 0, i32 [[TMP0]]) 1671; NVPTX-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]] 1672; NVPTX-DISABLED: worker_state_machine.parallel_region.end: 1673; NVPTX-DISABLED-NEXT: call void @__kmpc_kernel_end_parallel() 1674; NVPTX-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_DONE_BARRIER]] 1675; NVPTX-DISABLED: worker_state_machine.done.barrier: 1676; NVPTX-DISABLED-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 1677; NVPTX-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_BEGIN]] 1678; NVPTX-DISABLED: thread.user_code.check: 1679; NVPTX-DISABLED-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 1680; NVPTX-DISABLED-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 1681; NVPTX-DISABLED: user_code.entry: 1682; NVPTX-DISABLED-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR4]] 1683; NVPTX-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4 1684; NVPTX-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]]) #[[ATTR4]] 1685; NVPTX-DISABLED-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true) 1686; NVPTX-DISABLED-NEXT: ret void 1687; NVPTX-DISABLED: worker.exit: 1688; NVPTX-DISABLED-NEXT: ret void 1689; 1690entry: 1691 %.zero.addr = alloca i32, align 4 1692 %.threadid_temp. = alloca i32, align 4 1693 store i32 0, i32* %.zero.addr, align 4 1694 %0 = call i32 @__kmpc_target_init(%struct.ident_t* @1, i8 1, i1 true, i1 true) 1695 %exec_user_code = icmp eq i32 %0, -1 1696 br i1 %exec_user_code, label %user_code.entry, label %worker.exit 1697 1698user_code.entry: ; preds = %entry 1699 %1 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @1) 1700 store i32 %1, i32* %.threadid_temp., align 4 1701 call void @__omp_outlined__6(i32* %.threadid_temp., i32* %.zero.addr) #3 1702 call void @__kmpc_target_deinit(%struct.ident_t* @1, i8 1, i1 true) 1703 ret void 1704 1705worker.exit: ; preds = %entry 1706 ret void 1707} 1708 1709define internal void @__omp_outlined__6(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 { 1710; 1711; 1712; AMDGPU-LABEL: define {{[^@]+}}@__omp_outlined__6 1713; AMDGPU-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]] { 1714; AMDGPU-NEXT: entry: 1715; AMDGPU-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1716; AMDGPU-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1717; AMDGPU-NEXT: [[I:%.*]] = alloca i32, align 4 1718; AMDGPU-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 8 1719; AMDGPU-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 1720; AMDGPU-NEXT: [[X_ON_STACK:%.*]] = bitcast i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([4 x i8], [4 x i8] addrspace(3)* @x.1, i32 0, i32 0) to i8*) to i32* 1721; AMDGPU-NEXT: br label [[REGION_CHECK_TID:%.*]] 1722; AMDGPU: region.check.tid: 1723; AMDGPU-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block() 1724; AMDGPU-NEXT: [[TMP1:%.*]] = icmp eq i32 [[TMP0]], 0 1725; AMDGPU-NEXT: br i1 [[TMP1]], label [[REGION_GUARDED:%.*]], label [[REGION_BARRIER:%.*]] 1726; AMDGPU: region.guarded: 1727; AMDGPU-NEXT: store i32 42, i32* [[X_ON_STACK]], align 4 1728; AMDGPU-NEXT: br label [[REGION_GUARDED_END:%.*]] 1729; AMDGPU: region.guarded.end: 1730; AMDGPU-NEXT: br label [[REGION_BARRIER]] 1731; AMDGPU: region.barrier: 1732; AMDGPU-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 1733; AMDGPU-NEXT: br label [[REGION_EXIT:%.*]] 1734; AMDGPU: region.exit: 1735; AMDGPU-NEXT: store i32 0, i32* [[I]], align 4 1736; AMDGPU-NEXT: br label [[FOR_COND:%.*]] 1737; AMDGPU: for.cond: 1738; AMDGPU-NEXT: [[TMP2:%.*]] = load i32, i32* [[I]], align 4 1739; AMDGPU-NEXT: [[CMP:%.*]] = icmp slt i32 [[TMP2]], 100 1740; AMDGPU-NEXT: br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]] 1741; AMDGPU: for.body: 1742; AMDGPU-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0 1743; AMDGPU-NEXT: store i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([4 x i8], [4 x i8] addrspace(3)* @x.1, i32 0, i32 0) to i8*), i8** [[TMP3]], align 8 1744; AMDGPU-NEXT: [[TMP4:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4 1745; AMDGPU-NEXT: [[TMP5:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 1746; AMDGPU-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP4]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*, i32*)* @__omp_outlined__7 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__7_wrapper to i8*), i8** noundef [[TMP5]], i64 noundef 1) 1747; AMDGPU-NEXT: br label [[FOR_INC:%.*]] 1748; AMDGPU: for.inc: 1749; AMDGPU-NEXT: [[TMP6:%.*]] = load i32, i32* [[I]], align 4 1750; AMDGPU-NEXT: [[INC:%.*]] = add nsw i32 [[TMP6]], 1 1751; AMDGPU-NEXT: store i32 [[INC]], i32* [[I]], align 4 1752; AMDGPU-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP17:![0-9]+]] 1753; AMDGPU: for.end: 1754; AMDGPU-NEXT: call void @spmd_amenable() #[[ATTR7]] 1755; AMDGPU-NEXT: ret void 1756; 1757; NVPTX-LABEL: define {{[^@]+}}@__omp_outlined__6 1758; NVPTX-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]] { 1759; NVPTX-NEXT: entry: 1760; NVPTX-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1761; NVPTX-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1762; NVPTX-NEXT: [[I:%.*]] = alloca i32, align 4 1763; NVPTX-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 8 1764; NVPTX-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 1765; NVPTX-NEXT: [[X_ON_STACK:%.*]] = bitcast i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([4 x i8], [4 x i8] addrspace(3)* @x1, i32 0, i32 0) to i8*) to i32* 1766; NVPTX-NEXT: br label [[REGION_CHECK_TID:%.*]] 1767; NVPTX: region.check.tid: 1768; NVPTX-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block() 1769; NVPTX-NEXT: [[TMP1:%.*]] = icmp eq i32 [[TMP0]], 0 1770; NVPTX-NEXT: br i1 [[TMP1]], label [[REGION_GUARDED:%.*]], label [[REGION_BARRIER:%.*]] 1771; NVPTX: region.guarded: 1772; NVPTX-NEXT: store i32 42, i32* [[X_ON_STACK]], align 4 1773; NVPTX-NEXT: br label [[REGION_GUARDED_END:%.*]] 1774; NVPTX: region.guarded.end: 1775; NVPTX-NEXT: br label [[REGION_BARRIER]] 1776; NVPTX: region.barrier: 1777; NVPTX-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 1778; NVPTX-NEXT: br label [[REGION_EXIT:%.*]] 1779; NVPTX: region.exit: 1780; NVPTX-NEXT: store i32 0, i32* [[I]], align 4 1781; NVPTX-NEXT: br label [[FOR_COND:%.*]] 1782; NVPTX: for.cond: 1783; NVPTX-NEXT: [[TMP2:%.*]] = load i32, i32* [[I]], align 4 1784; NVPTX-NEXT: [[CMP:%.*]] = icmp slt i32 [[TMP2]], 100 1785; NVPTX-NEXT: br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]] 1786; NVPTX: for.body: 1787; NVPTX-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0 1788; NVPTX-NEXT: store i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([4 x i8], [4 x i8] addrspace(3)* @x1, i32 0, i32 0) to i8*), i8** [[TMP3]], align 8 1789; NVPTX-NEXT: [[TMP4:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4 1790; NVPTX-NEXT: [[TMP5:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 1791; NVPTX-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP4]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*, i32*)* @__omp_outlined__7 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__7_wrapper to i8*), i8** noundef [[TMP5]], i64 noundef 1) 1792; NVPTX-NEXT: br label [[FOR_INC:%.*]] 1793; NVPTX: for.inc: 1794; NVPTX-NEXT: [[TMP6:%.*]] = load i32, i32* [[I]], align 4 1795; NVPTX-NEXT: [[INC:%.*]] = add nsw i32 [[TMP6]], 1 1796; NVPTX-NEXT: store i32 [[INC]], i32* [[I]], align 4 1797; NVPTX-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP17:![0-9]+]] 1798; NVPTX: for.end: 1799; NVPTX-NEXT: call void @spmd_amenable() #[[ATTR7]] 1800; NVPTX-NEXT: ret void 1801; 1802; AMDGPU-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__6 1803; AMDGPU-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]] { 1804; AMDGPU-DISABLED-NEXT: entry: 1805; AMDGPU-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1806; AMDGPU-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1807; AMDGPU-DISABLED-NEXT: [[I:%.*]] = alloca i32, align 4 1808; AMDGPU-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 8 1809; AMDGPU-DISABLED-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 1810; AMDGPU-DISABLED-NEXT: [[X_ON_STACK:%.*]] = bitcast i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([4 x i8], [4 x i8] addrspace(3)* @x.1, i32 0, i32 0) to i8*) to i32* 1811; AMDGPU-DISABLED-NEXT: store i32 42, i32* [[X_ON_STACK]], align 4 1812; AMDGPU-DISABLED-NEXT: store i32 0, i32* [[I]], align 4 1813; AMDGPU-DISABLED-NEXT: br label [[FOR_COND:%.*]] 1814; AMDGPU-DISABLED: for.cond: 1815; AMDGPU-DISABLED-NEXT: [[TMP0:%.*]] = load i32, i32* [[I]], align 4 1816; AMDGPU-DISABLED-NEXT: [[CMP:%.*]] = icmp slt i32 [[TMP0]], 100 1817; AMDGPU-DISABLED-NEXT: br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]] 1818; AMDGPU-DISABLED: for.body: 1819; AMDGPU-DISABLED-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0 1820; AMDGPU-DISABLED-NEXT: store i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([4 x i8], [4 x i8] addrspace(3)* @x.1, i32 0, i32 0) to i8*), i8** [[TMP1]], align 8 1821; AMDGPU-DISABLED-NEXT: [[TMP2:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4 1822; AMDGPU-DISABLED-NEXT: [[TMP3:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 1823; AMDGPU-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP2]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*, i32*)* @__omp_outlined__7 to i8*), i8* noundef @__omp_outlined__7_wrapper.ID, i8** noundef [[TMP3]], i64 noundef 1) 1824; AMDGPU-DISABLED-NEXT: br label [[FOR_INC:%.*]] 1825; AMDGPU-DISABLED: for.inc: 1826; AMDGPU-DISABLED-NEXT: [[TMP4:%.*]] = load i32, i32* [[I]], align 4 1827; AMDGPU-DISABLED-NEXT: [[INC:%.*]] = add nsw i32 [[TMP4]], 1 1828; AMDGPU-DISABLED-NEXT: store i32 [[INC]], i32* [[I]], align 4 1829; AMDGPU-DISABLED-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP17:![0-9]+]] 1830; AMDGPU-DISABLED: for.end: 1831; AMDGPU-DISABLED-NEXT: call void @spmd_amenable() #[[ATTR7]] 1832; AMDGPU-DISABLED-NEXT: ret void 1833; 1834; NVPTX-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__6 1835; NVPTX-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]] { 1836; NVPTX-DISABLED-NEXT: entry: 1837; NVPTX-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1838; NVPTX-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1839; NVPTX-DISABLED-NEXT: [[I:%.*]] = alloca i32, align 4 1840; NVPTX-DISABLED-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 8 1841; NVPTX-DISABLED-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 1842; NVPTX-DISABLED-NEXT: [[X_ON_STACK:%.*]] = bitcast i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([4 x i8], [4 x i8] addrspace(3)* @x1, i32 0, i32 0) to i8*) to i32* 1843; NVPTX-DISABLED-NEXT: store i32 42, i32* [[X_ON_STACK]], align 4 1844; NVPTX-DISABLED-NEXT: store i32 0, i32* [[I]], align 4 1845; NVPTX-DISABLED-NEXT: br label [[FOR_COND:%.*]] 1846; NVPTX-DISABLED: for.cond: 1847; NVPTX-DISABLED-NEXT: [[TMP0:%.*]] = load i32, i32* [[I]], align 4 1848; NVPTX-DISABLED-NEXT: [[CMP:%.*]] = icmp slt i32 [[TMP0]], 100 1849; NVPTX-DISABLED-NEXT: br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]] 1850; NVPTX-DISABLED: for.body: 1851; NVPTX-DISABLED-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0 1852; NVPTX-DISABLED-NEXT: store i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([4 x i8], [4 x i8] addrspace(3)* @x1, i32 0, i32 0) to i8*), i8** [[TMP1]], align 8 1853; NVPTX-DISABLED-NEXT: [[TMP2:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4 1854; NVPTX-DISABLED-NEXT: [[TMP3:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 1855; NVPTX-DISABLED-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP2]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*, i32*)* @__omp_outlined__7 to i8*), i8* noundef @__omp_outlined__7_wrapper.ID, i8** noundef [[TMP3]], i64 noundef 1) 1856; NVPTX-DISABLED-NEXT: br label [[FOR_INC:%.*]] 1857; NVPTX-DISABLED: for.inc: 1858; NVPTX-DISABLED-NEXT: [[TMP4:%.*]] = load i32, i32* [[I]], align 4 1859; NVPTX-DISABLED-NEXT: [[INC:%.*]] = add nsw i32 [[TMP4]], 1 1860; NVPTX-DISABLED-NEXT: store i32 [[INC]], i32* [[I]], align 4 1861; NVPTX-DISABLED-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP17:![0-9]+]] 1862; NVPTX-DISABLED: for.end: 1863; NVPTX-DISABLED-NEXT: call void @spmd_amenable() #[[ATTR7]] 1864; NVPTX-DISABLED-NEXT: ret void 1865; 1866entry: 1867 %.global_tid..addr = alloca i32*, align 8 1868 %.bound_tid..addr = alloca i32*, align 8 1869 %i = alloca i32, align 4 1870 %captured_vars_addrs = alloca [1 x i8*], align 8 1871 store i32* %.global_tid., i32** %.global_tid..addr, align 8 1872 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8 1873 %x = call i8* @__kmpc_alloc_shared(i64 4) 1874 %x_on_stack = bitcast i8* %x to i32* 1875 store i32 42, i32* %x_on_stack, align 4 1876 store i32 0, i32* %i, align 4 1877 br label %for.cond 1878 1879for.cond: ; preds = %for.inc, %entry 1880 %0 = load i32, i32* %i, align 4 1881 %cmp = icmp slt i32 %0, 100 1882 br i1 %cmp, label %for.body, label %for.end 1883 1884for.body: ; preds = %for.cond 1885 %1 = getelementptr inbounds [1 x i8*], [1 x i8*]* %captured_vars_addrs, i64 0, i64 0 1886 %2 = bitcast i32* %x_on_stack to i8* 1887 store i8* %2, i8** %1, align 8 1888 %3 = load i32*, i32** %.global_tid..addr, align 8 1889 %4 = load i32, i32* %3, align 4 1890 %5 = bitcast [1 x i8*]* %captured_vars_addrs to i8** 1891 call void @__kmpc_parallel_51(%struct.ident_t* @1, i32 %4, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i32*)* @__omp_outlined__7 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__7_wrapper to i8*), i8** %5, i64 1) 1892 br label %for.inc 1893 1894for.inc: ; preds = %for.body 1895 %6 = load i32, i32* %i, align 4 1896 %inc = add nsw i32 %6, 1 1897 store i32 %inc, i32* %i, align 4 1898 br label %for.cond, !llvm.loop !17 1899 1900for.end: ; preds = %for.cond 1901 call void @spmd_amenable() #4 1902 call void @__kmpc_free_shared(i8* %x, i64 4) 1903 ret void 1904} 1905 1906define internal void @__omp_outlined__7(i32* noalias %.global_tid., i32* noalias %.bound_tid., i32* nonnull align 4 dereferenceable(4) %x) #0 { 1907; 1908; 1909; AMDGPU-LABEL: define {{[^@]+}}@__omp_outlined__7 1910; AMDGPU-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]], i32* nocapture nofree nonnull align 4 dereferenceable(4) [[X:%.*]]) #[[ATTR0]] { 1911; AMDGPU-NEXT: entry: 1912; AMDGPU-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1913; AMDGPU-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1914; AMDGPU-NEXT: [[X_ADDR:%.*]] = alloca i32*, align 8 1915; AMDGPU-NEXT: store i32* [[X]], i32** [[X_ADDR]], align 8 1916; AMDGPU-NEXT: [[TMP0:%.*]] = load i32, i32* [[X]], align 4 1917; AMDGPU-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1 1918; AMDGPU-NEXT: store i32 [[INC]], i32* [[X]], align 4 1919; AMDGPU-NEXT: call void @unknown() #[[ATTR8]] 1920; AMDGPU-NEXT: ret void 1921; 1922; NVPTX-LABEL: define {{[^@]+}}@__omp_outlined__7 1923; NVPTX-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]], i32* nocapture nofree nonnull align 4 dereferenceable(4) [[X:%.*]]) #[[ATTR0]] { 1924; NVPTX-NEXT: entry: 1925; NVPTX-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1926; NVPTX-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1927; NVPTX-NEXT: [[X_ADDR:%.*]] = alloca i32*, align 8 1928; NVPTX-NEXT: store i32* [[X]], i32** [[X_ADDR]], align 8 1929; NVPTX-NEXT: [[TMP0:%.*]] = load i32, i32* [[X]], align 4 1930; NVPTX-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1 1931; NVPTX-NEXT: store i32 [[INC]], i32* [[X]], align 4 1932; NVPTX-NEXT: call void @unknown() #[[ATTR8]] 1933; NVPTX-NEXT: ret void 1934; 1935; AMDGPU-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__7 1936; AMDGPU-DISABLED-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]], i32* nocapture nofree nonnull align 4 dereferenceable(4) [[X:%.*]]) #[[ATTR0]] { 1937; AMDGPU-DISABLED-NEXT: entry: 1938; AMDGPU-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1939; AMDGPU-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1940; AMDGPU-DISABLED-NEXT: [[X_ADDR:%.*]] = alloca i32*, align 8 1941; AMDGPU-DISABLED-NEXT: store i32* [[X]], i32** [[X_ADDR]], align 8 1942; AMDGPU-DISABLED-NEXT: [[TMP0:%.*]] = load i32, i32* [[X]], align 4 1943; AMDGPU-DISABLED-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1 1944; AMDGPU-DISABLED-NEXT: store i32 [[INC]], i32* [[X]], align 4 1945; AMDGPU-DISABLED-NEXT: call void @unknown() #[[ATTR8]] 1946; AMDGPU-DISABLED-NEXT: ret void 1947; 1948; NVPTX-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__7 1949; NVPTX-DISABLED-SAME: (i32* noalias nocapture nofree readnone [[DOTGLOBAL_TID_:%.*]], i32* noalias nocapture nofree readnone [[DOTBOUND_TID_:%.*]], i32* nocapture nofree nonnull align 4 dereferenceable(4) [[X:%.*]]) #[[ATTR0]] { 1950; NVPTX-DISABLED-NEXT: entry: 1951; NVPTX-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1952; NVPTX-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1953; NVPTX-DISABLED-NEXT: [[X_ADDR:%.*]] = alloca i32*, align 8 1954; NVPTX-DISABLED-NEXT: store i32* [[X]], i32** [[X_ADDR]], align 8 1955; NVPTX-DISABLED-NEXT: [[TMP0:%.*]] = load i32, i32* [[X]], align 4 1956; NVPTX-DISABLED-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1 1957; NVPTX-DISABLED-NEXT: store i32 [[INC]], i32* [[X]], align 4 1958; NVPTX-DISABLED-NEXT: call void @unknown() #[[ATTR8]] 1959; NVPTX-DISABLED-NEXT: ret void 1960; 1961entry: 1962 %.global_tid..addr = alloca i32*, align 8 1963 %.bound_tid..addr = alloca i32*, align 8 1964 %x.addr = alloca i32*, align 8 1965 store i32* %.global_tid., i32** %.global_tid..addr, align 8 1966 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8 1967 store i32* %x, i32** %x.addr, align 8 1968 %0 = load i32*, i32** %x.addr, align 8 1969 %1 = load i32, i32* %0, align 4 1970 %inc = add nsw i32 %1, 1 1971 store i32 %inc, i32* %0, align 4 1972 call void @unknown() #5 1973 ret void 1974} 1975 1976define internal void @__omp_outlined__7_wrapper(i16 zeroext %0, i32 %1) #0 { 1977; 1978; 1979; AMDGPU-LABEL: define {{[^@]+}}@__omp_outlined__7_wrapper 1980; AMDGPU-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 1981; AMDGPU-NEXT: entry: 1982; AMDGPU-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 1983; AMDGPU-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 1984; AMDGPU-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1985; AMDGPU-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8 1986; AMDGPU-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 1987; AMDGPU-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 1988; AMDGPU-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 1989; AMDGPU-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 1990; AMDGPU-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 8 1991; AMDGPU-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i64 0 1992; AMDGPU-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32** 1993; AMDGPU-NEXT: [[TMP5:%.*]] = load i32*, i32** [[TMP4]], align 8 1994; AMDGPU-NEXT: call void @__omp_outlined__7(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32* [[TMP5]]) #[[ATTR4]] 1995; AMDGPU-NEXT: ret void 1996; 1997; NVPTX-LABEL: define {{[^@]+}}@__omp_outlined__7_wrapper 1998; NVPTX-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 1999; NVPTX-NEXT: entry: 2000; NVPTX-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 2001; NVPTX-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 2002; NVPTX-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 2003; NVPTX-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8 2004; NVPTX-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 2005; NVPTX-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 2006; NVPTX-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 2007; NVPTX-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 2008; NVPTX-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 8 2009; NVPTX-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i64 0 2010; NVPTX-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32** 2011; NVPTX-NEXT: [[TMP5:%.*]] = load i32*, i32** [[TMP4]], align 8 2012; NVPTX-NEXT: call void @__omp_outlined__7(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32* [[TMP5]]) #[[ATTR4]] 2013; NVPTX-NEXT: ret void 2014; 2015; AMDGPU-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__7_wrapper 2016; AMDGPU-DISABLED-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 2017; AMDGPU-DISABLED-NEXT: entry: 2018; AMDGPU-DISABLED-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 2019; AMDGPU-DISABLED-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 2020; AMDGPU-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 2021; AMDGPU-DISABLED-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8 2022; AMDGPU-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 2023; AMDGPU-DISABLED-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 2024; AMDGPU-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 2025; AMDGPU-DISABLED-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 2026; AMDGPU-DISABLED-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 8 2027; AMDGPU-DISABLED-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i64 0 2028; AMDGPU-DISABLED-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32** 2029; AMDGPU-DISABLED-NEXT: [[TMP5:%.*]] = load i32*, i32** [[TMP4]], align 8 2030; AMDGPU-DISABLED-NEXT: call void @__omp_outlined__7(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32* [[TMP5]]) #[[ATTR4]] 2031; AMDGPU-DISABLED-NEXT: ret void 2032; 2033; NVPTX-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__7_wrapper 2034; NVPTX-DISABLED-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 2035; NVPTX-DISABLED-NEXT: entry: 2036; NVPTX-DISABLED-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 2037; NVPTX-DISABLED-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 2038; NVPTX-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 2039; NVPTX-DISABLED-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8 2040; NVPTX-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 2041; NVPTX-DISABLED-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 2042; NVPTX-DISABLED-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 2043; NVPTX-DISABLED-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 2044; NVPTX-DISABLED-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 8 2045; NVPTX-DISABLED-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i64 0 2046; NVPTX-DISABLED-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32** 2047; NVPTX-DISABLED-NEXT: [[TMP5:%.*]] = load i32*, i32** [[TMP4]], align 8 2048; NVPTX-DISABLED-NEXT: call void @__omp_outlined__7(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32* [[TMP5]]) #[[ATTR4]] 2049; NVPTX-DISABLED-NEXT: ret void 2050; 2051entry: 2052 %.addr = alloca i16, align 2 2053 %.addr1 = alloca i32, align 4 2054 %.zero.addr = alloca i32, align 4 2055 %global_args = alloca i8**, align 8 2056 store i32 0, i32* %.zero.addr, align 4 2057 store i16 %0, i16* %.addr, align 2 2058 store i32 %1, i32* %.addr1, align 4 2059 call void @__kmpc_get_shared_variables(i8*** %global_args) 2060 %2 = load i8**, i8*** %global_args, align 8 2061 %3 = getelementptr inbounds i8*, i8** %2, i64 0 2062 %4 = bitcast i8** %3 to i32** 2063 %5 = load i32*, i32** %4, align 8 2064 call void @__omp_outlined__7(i32* %.addr1, i32* %.zero.addr, i32* %5) #3 2065 ret void 2066} 2067 2068define weak void @__omp_offloading_14_a34ca11_do_not_spmdize_target_l65() #0 { 2069; 2070; 2071; AMDGPU-LABEL: define {{[^@]+}}@__omp_offloading_14_a34ca11_do_not_spmdize_target_l65 2072; AMDGPU-SAME: () #[[ATTR0]] { 2073; AMDGPU-NEXT: entry: 2074; AMDGPU-NEXT: [[WORKER_WORK_FN_ADDR:%.*]] = alloca i8*, align 8, addrspace(5) 2075; AMDGPU-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 2076; AMDGPU-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 2077; AMDGPU-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 1, i1 false, i1 true) 2078; AMDGPU-NEXT: [[BLOCK_HW_SIZE:%.*]] = call i32 @__kmpc_get_hardware_num_threads_in_block() 2079; AMDGPU-NEXT: [[WARP_SIZE:%.*]] = call i32 @__kmpc_get_warp_size() 2080; AMDGPU-NEXT: [[BLOCK_SIZE:%.*]] = sub i32 [[BLOCK_HW_SIZE]], [[WARP_SIZE]] 2081; AMDGPU-NEXT: [[THREAD_IS_MAIN_OR_WORKER:%.*]] = icmp slt i32 [[TMP0]], [[BLOCK_SIZE]] 2082; AMDGPU-NEXT: br i1 [[THREAD_IS_MAIN_OR_WORKER]], label [[IS_WORKER_CHECK:%.*]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]] 2083; AMDGPU: is_worker_check: 2084; AMDGPU-NEXT: [[THREAD_IS_WORKER:%.*]] = icmp ne i32 [[TMP0]], -1 2085; AMDGPU-NEXT: br i1 [[THREAD_IS_WORKER]], label [[WORKER_STATE_MACHINE_BEGIN:%.*]], label [[THREAD_USER_CODE_CHECK:%.*]] 2086; AMDGPU: worker_state_machine.begin: 2087; AMDGPU-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 2088; AMDGPU-NEXT: [[WORKER_WORK_FN_ADDR_GENERIC:%.*]] = addrspacecast i8* addrspace(5)* [[WORKER_WORK_FN_ADDR]] to i8** 2089; AMDGPU-NEXT: [[WORKER_IS_ACTIVE:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORKER_WORK_FN_ADDR_GENERIC]]) 2090; AMDGPU-NEXT: [[WORKER_WORK_FN:%.*]] = load i8*, i8** [[WORKER_WORK_FN_ADDR_GENERIC]], align 8 2091; AMDGPU-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)* 2092; AMDGPU-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null 2093; AMDGPU-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]] 2094; AMDGPU: worker_state_machine.finished: 2095; AMDGPU-NEXT: ret void 2096; AMDGPU: worker_state_machine.is_active.check: 2097; AMDGPU-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_FALLBACK_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]] 2098; AMDGPU: worker_state_machine.parallel_region.fallback.execute: 2099; AMDGPU-NEXT: call void [[WORKER_WORK_FN_ADDR_CAST]](i16 0, i32 [[TMP0]]) 2100; AMDGPU-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]] 2101; AMDGPU: worker_state_machine.parallel_region.end: 2102; AMDGPU-NEXT: call void @__kmpc_kernel_end_parallel() 2103; AMDGPU-NEXT: br label [[WORKER_STATE_MACHINE_DONE_BARRIER]] 2104; AMDGPU: worker_state_machine.done.barrier: 2105; AMDGPU-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 2106; AMDGPU-NEXT: br label [[WORKER_STATE_MACHINE_BEGIN]] 2107; AMDGPU: thread.user_code.check: 2108; AMDGPU-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 2109; AMDGPU-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 2110; AMDGPU: user_code.entry: 2111; AMDGPU-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR4]] 2112; AMDGPU-NEXT: call void @__omp_outlined__8(i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR4]] 2113; AMDGPU-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true) 2114; AMDGPU-NEXT: ret void 2115; AMDGPU: worker.exit: 2116; AMDGPU-NEXT: ret void 2117; 2118; NVPTX-LABEL: define {{[^@]+}}@__omp_offloading_14_a34ca11_do_not_spmdize_target_l65 2119; NVPTX-SAME: () #[[ATTR0]] { 2120; NVPTX-NEXT: entry: 2121; NVPTX-NEXT: [[WORKER_WORK_FN_ADDR:%.*]] = alloca i8*, align 8 2122; NVPTX-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 2123; NVPTX-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 2124; NVPTX-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 1, i1 false, i1 true) 2125; NVPTX-NEXT: [[BLOCK_HW_SIZE:%.*]] = call i32 @__kmpc_get_hardware_num_threads_in_block() 2126; NVPTX-NEXT: [[WARP_SIZE:%.*]] = call i32 @__kmpc_get_warp_size() 2127; NVPTX-NEXT: [[BLOCK_SIZE:%.*]] = sub i32 [[BLOCK_HW_SIZE]], [[WARP_SIZE]] 2128; NVPTX-NEXT: [[THREAD_IS_MAIN_OR_WORKER:%.*]] = icmp slt i32 [[TMP0]], [[BLOCK_SIZE]] 2129; NVPTX-NEXT: br i1 [[THREAD_IS_MAIN_OR_WORKER]], label [[IS_WORKER_CHECK:%.*]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]] 2130; NVPTX: is_worker_check: 2131; NVPTX-NEXT: [[THREAD_IS_WORKER:%.*]] = icmp ne i32 [[TMP0]], -1 2132; NVPTX-NEXT: br i1 [[THREAD_IS_WORKER]], label [[WORKER_STATE_MACHINE_BEGIN:%.*]], label [[THREAD_USER_CODE_CHECK:%.*]] 2133; NVPTX: worker_state_machine.begin: 2134; NVPTX-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 2135; NVPTX-NEXT: [[WORKER_IS_ACTIVE:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORKER_WORK_FN_ADDR]]) 2136; NVPTX-NEXT: [[WORKER_WORK_FN:%.*]] = load i8*, i8** [[WORKER_WORK_FN_ADDR]], align 8 2137; NVPTX-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)* 2138; NVPTX-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null 2139; NVPTX-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]] 2140; NVPTX: worker_state_machine.finished: 2141; NVPTX-NEXT: ret void 2142; NVPTX: worker_state_machine.is_active.check: 2143; NVPTX-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_FALLBACK_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]] 2144; NVPTX: worker_state_machine.parallel_region.fallback.execute: 2145; NVPTX-NEXT: call void [[WORKER_WORK_FN_ADDR_CAST]](i16 0, i32 [[TMP0]]) 2146; NVPTX-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]] 2147; NVPTX: worker_state_machine.parallel_region.end: 2148; NVPTX-NEXT: call void @__kmpc_kernel_end_parallel() 2149; NVPTX-NEXT: br label [[WORKER_STATE_MACHINE_DONE_BARRIER]] 2150; NVPTX: worker_state_machine.done.barrier: 2151; NVPTX-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 2152; NVPTX-NEXT: br label [[WORKER_STATE_MACHINE_BEGIN]] 2153; NVPTX: thread.user_code.check: 2154; NVPTX-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 2155; NVPTX-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 2156; NVPTX: user_code.entry: 2157; NVPTX-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR4]] 2158; NVPTX-NEXT: call void @__omp_outlined__8(i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR4]] 2159; NVPTX-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true) 2160; NVPTX-NEXT: ret void 2161; NVPTX: worker.exit: 2162; NVPTX-NEXT: ret void 2163; 2164; AMDGPU-DISABLED-LABEL: define {{[^@]+}}@__omp_offloading_14_a34ca11_do_not_spmdize_target_l65 2165; AMDGPU-DISABLED-SAME: () #[[ATTR0]] { 2166; AMDGPU-DISABLED-NEXT: entry: 2167; AMDGPU-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR:%.*]] = alloca i8*, align 8, addrspace(5) 2168; AMDGPU-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 2169; AMDGPU-DISABLED-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 2170; AMDGPU-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 2171; AMDGPU-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 1, i1 false, i1 true) 2172; AMDGPU-DISABLED-NEXT: [[BLOCK_HW_SIZE:%.*]] = call i32 @__kmpc_get_hardware_num_threads_in_block() 2173; AMDGPU-DISABLED-NEXT: [[WARP_SIZE:%.*]] = call i32 @__kmpc_get_warp_size() 2174; AMDGPU-DISABLED-NEXT: [[BLOCK_SIZE:%.*]] = sub i32 [[BLOCK_HW_SIZE]], [[WARP_SIZE]] 2175; AMDGPU-DISABLED-NEXT: [[THREAD_IS_MAIN_OR_WORKER:%.*]] = icmp slt i32 [[TMP0]], [[BLOCK_SIZE]] 2176; AMDGPU-DISABLED-NEXT: br i1 [[THREAD_IS_MAIN_OR_WORKER]], label [[IS_WORKER_CHECK:%.*]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]] 2177; AMDGPU-DISABLED: is_worker_check: 2178; AMDGPU-DISABLED-NEXT: [[THREAD_IS_WORKER:%.*]] = icmp ne i32 [[TMP0]], -1 2179; AMDGPU-DISABLED-NEXT: br i1 [[THREAD_IS_WORKER]], label [[WORKER_STATE_MACHINE_BEGIN:%.*]], label [[THREAD_USER_CODE_CHECK:%.*]] 2180; AMDGPU-DISABLED: worker_state_machine.begin: 2181; AMDGPU-DISABLED-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 2182; AMDGPU-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR_GENERIC:%.*]] = addrspacecast i8* addrspace(5)* [[WORKER_WORK_FN_ADDR]] to i8** 2183; AMDGPU-DISABLED-NEXT: [[WORKER_IS_ACTIVE:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORKER_WORK_FN_ADDR_GENERIC]]) 2184; AMDGPU-DISABLED-NEXT: [[WORKER_WORK_FN:%.*]] = load i8*, i8** [[WORKER_WORK_FN_ADDR_GENERIC]], align 8 2185; AMDGPU-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)* 2186; AMDGPU-DISABLED-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null 2187; AMDGPU-DISABLED-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]] 2188; AMDGPU-DISABLED: worker_state_machine.finished: 2189; AMDGPU-DISABLED-NEXT: ret void 2190; AMDGPU-DISABLED: worker_state_machine.is_active.check: 2191; AMDGPU-DISABLED-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_FALLBACK_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]] 2192; AMDGPU-DISABLED: worker_state_machine.parallel_region.fallback.execute: 2193; AMDGPU-DISABLED-NEXT: call void [[WORKER_WORK_FN_ADDR_CAST]](i16 0, i32 [[TMP0]]) 2194; AMDGPU-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]] 2195; AMDGPU-DISABLED: worker_state_machine.parallel_region.end: 2196; AMDGPU-DISABLED-NEXT: call void @__kmpc_kernel_end_parallel() 2197; AMDGPU-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_DONE_BARRIER]] 2198; AMDGPU-DISABLED: worker_state_machine.done.barrier: 2199; AMDGPU-DISABLED-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 2200; AMDGPU-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_BEGIN]] 2201; AMDGPU-DISABLED: thread.user_code.check: 2202; AMDGPU-DISABLED-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 2203; AMDGPU-DISABLED-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 2204; AMDGPU-DISABLED: user_code.entry: 2205; AMDGPU-DISABLED-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR4]] 2206; AMDGPU-DISABLED-NEXT: call void @__omp_outlined__8(i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR4]] 2207; AMDGPU-DISABLED-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true) 2208; AMDGPU-DISABLED-NEXT: ret void 2209; AMDGPU-DISABLED: worker.exit: 2210; AMDGPU-DISABLED-NEXT: ret void 2211; 2212; NVPTX-DISABLED-LABEL: define {{[^@]+}}@__omp_offloading_14_a34ca11_do_not_spmdize_target_l65 2213; NVPTX-DISABLED-SAME: () #[[ATTR0]] { 2214; NVPTX-DISABLED-NEXT: entry: 2215; NVPTX-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR:%.*]] = alloca i8*, align 8 2216; NVPTX-DISABLED-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 2217; NVPTX-DISABLED-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 2218; NVPTX-DISABLED-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 2219; NVPTX-DISABLED-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 1, i1 false, i1 true) 2220; NVPTX-DISABLED-NEXT: [[BLOCK_HW_SIZE:%.*]] = call i32 @__kmpc_get_hardware_num_threads_in_block() 2221; NVPTX-DISABLED-NEXT: [[WARP_SIZE:%.*]] = call i32 @__kmpc_get_warp_size() 2222; NVPTX-DISABLED-NEXT: [[BLOCK_SIZE:%.*]] = sub i32 [[BLOCK_HW_SIZE]], [[WARP_SIZE]] 2223; NVPTX-DISABLED-NEXT: [[THREAD_IS_MAIN_OR_WORKER:%.*]] = icmp slt i32 [[TMP0]], [[BLOCK_SIZE]] 2224; NVPTX-DISABLED-NEXT: br i1 [[THREAD_IS_MAIN_OR_WORKER]], label [[IS_WORKER_CHECK:%.*]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]] 2225; NVPTX-DISABLED: is_worker_check: 2226; NVPTX-DISABLED-NEXT: [[THREAD_IS_WORKER:%.*]] = icmp ne i32 [[TMP0]], -1 2227; NVPTX-DISABLED-NEXT: br i1 [[THREAD_IS_WORKER]], label [[WORKER_STATE_MACHINE_BEGIN:%.*]], label [[THREAD_USER_CODE_CHECK:%.*]] 2228; NVPTX-DISABLED: worker_state_machine.begin: 2229; NVPTX-DISABLED-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 2230; NVPTX-DISABLED-NEXT: [[WORKER_IS_ACTIVE:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORKER_WORK_FN_ADDR]]) 2231; NVPTX-DISABLED-NEXT: [[WORKER_WORK_FN:%.*]] = load i8*, i8** [[WORKER_WORK_FN_ADDR]], align 8 2232; NVPTX-DISABLED-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)* 2233; NVPTX-DISABLED-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null 2234; NVPTX-DISABLED-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]] 2235; NVPTX-DISABLED: worker_state_machine.finished: 2236; NVPTX-DISABLED-NEXT: ret void 2237; NVPTX-DISABLED: worker_state_machine.is_active.check: 2238; NVPTX-DISABLED-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_FALLBACK_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]] 2239; NVPTX-DISABLED: worker_state_machine.parallel_region.fallback.execute: 2240; NVPTX-DISABLED-NEXT: call void [[WORKER_WORK_FN_ADDR_CAST]](i16 0, i32 [[TMP0]]) 2241; NVPTX-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]] 2242; NVPTX-DISABLED: worker_state_machine.parallel_region.end: 2243; NVPTX-DISABLED-NEXT: call void @__kmpc_kernel_end_parallel() 2244; NVPTX-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_DONE_BARRIER]] 2245; NVPTX-DISABLED: worker_state_machine.done.barrier: 2246; NVPTX-DISABLED-NEXT: call void @__kmpc_barrier_simple_generic(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]]) 2247; NVPTX-DISABLED-NEXT: br label [[WORKER_STATE_MACHINE_BEGIN]] 2248; NVPTX-DISABLED: thread.user_code.check: 2249; NVPTX-DISABLED-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 2250; NVPTX-DISABLED-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 2251; NVPTX-DISABLED: user_code.entry: 2252; NVPTX-DISABLED-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) #[[ATTR4]] 2253; NVPTX-DISABLED-NEXT: call void @__omp_outlined__8(i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTTHREADID_TEMP_]], i32* noalias nocapture noundef nonnull readnone align 4 dereferenceable(4) [[DOTZERO_ADDR]]) #[[ATTR4]] 2254; NVPTX-DISABLED-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true) 2255; NVPTX-DISABLED-NEXT: ret void 2256; NVPTX-DISABLED: worker.exit: 2257; NVPTX-DISABLED-NEXT: ret void 2258; 2259entry: 2260 %.zero.addr = alloca i32, align 4 2261 %.threadid_temp. = alloca i32, align 4 2262 store i32 0, i32* %.zero.addr, align 4 2263 %0 = call i32 @__kmpc_target_init(%struct.ident_t* @1, i8 1, i1 true, i1 true) 2264 %exec_user_code = icmp eq i32 %0, -1 2265 br i1 %exec_user_code, label %user_code.entry, label %worker.exit 2266 2267user_code.entry: ; preds = %entry 2268 %1 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @1) 2269 store i32 %1, i32* %.threadid_temp., align 4 2270 call void @__omp_outlined__8(i32* %.threadid_temp., i32* %.zero.addr) #3 2271 call void @__kmpc_target_deinit(%struct.ident_t* @1, i8 1, i1 true) 2272 ret void 2273 2274worker.exit: ; preds = %entry 2275 ret void 2276} 2277 2278define internal void @__omp_outlined__8(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #0 { 2279; 2280; 2281; AMDGPU-LABEL: define {{[^@]+}}@__omp_outlined__8 2282; AMDGPU-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]] { 2283; AMDGPU-NEXT: entry: 2284; AMDGPU-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 2285; AMDGPU-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 2286; AMDGPU-NEXT: call void @unknown() #[[ATTR8]] 2287; AMDGPU-NEXT: ret void 2288; 2289; NVPTX-LABEL: define {{[^@]+}}@__omp_outlined__8 2290; NVPTX-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]] { 2291; NVPTX-NEXT: entry: 2292; NVPTX-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 2293; NVPTX-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 2294; NVPTX-NEXT: call void @unknown() #[[ATTR8]] 2295; NVPTX-NEXT: ret void 2296; 2297; AMDGPU-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__8 2298; AMDGPU-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]] { 2299; AMDGPU-DISABLED-NEXT: entry: 2300; AMDGPU-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 2301; AMDGPU-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 2302; AMDGPU-DISABLED-NEXT: call void @unknown() #[[ATTR8]] 2303; AMDGPU-DISABLED-NEXT: ret void 2304; 2305; NVPTX-DISABLED-LABEL: define {{[^@]+}}@__omp_outlined__8 2306; NVPTX-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]] { 2307; NVPTX-DISABLED-NEXT: entry: 2308; NVPTX-DISABLED-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 2309; NVPTX-DISABLED-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 2310; NVPTX-DISABLED-NEXT: call void @unknown() #[[ATTR8]] 2311; NVPTX-DISABLED-NEXT: ret void 2312; 2313entry: 2314 %.global_tid..addr = alloca i32*, align 8 2315 %.bound_tid..addr = alloca i32*, align 8 2316 store i32* %.global_tid., i32** %.global_tid..addr, align 8 2317 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8 2318 call void @unknown() #5 2319 ret void 2320} 2321 2322attributes #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" } 2323attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32,+sm_20" } 2324attributes #2 = { convergent "frame-pointer"="none" "llvm.assume"="ompx_spmd_amenable" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32,+sm_20" } 2325attributes #3 = { nounwind } 2326attributes #4 = { convergent "llvm.assume"="ompx_spmd_amenable" } 2327attributes #5 = { convergent } 2328 2329!omp_offload.info = !{!0, !1, !2, !3, !4} 2330!nvvm.annotations = !{!5, !6, !7, !8, !9} 2331!llvm.module.flags = !{!10, !11, !12} 2332 2333!0 = !{i32 0, i32 20, i32 171231761, !"sequential_loop_to_stack_var", i32 20, i32 1} 2334!1 = !{i32 0, i32 20, i32 171231761, !"sequential_loop", i32 5, i32 0} 2335!2 = !{i32 0, i32 20, i32 171231761, !"sequential_loop_to_shared_var", i32 35, i32 2} 2336!3 = !{i32 0, i32 20, i32 171231761, !"do_not_spmdize_target", i32 65, i32 4} 2337!4 = !{i32 0, i32 20, i32 171231761, !"sequential_loop_to_shared_var_guarded", i32 50, i32 3} 2338!5 = !{void ()* @__omp_offloading_14_a34ca11_sequential_loop_l5, !"kernel", i32 1} 2339!6 = !{void ()* @__omp_offloading_14_a34ca11_sequential_loop_to_stack_var_l20, !"kernel", i32 1} 2340!7 = !{void ()* @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_l35, !"kernel", i32 1} 2341!8 = !{void ()* @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_guarded_l50, !"kernel", i32 1} 2342!9 = !{void ()* @__omp_offloading_14_a34ca11_do_not_spmdize_target_l65, !"kernel", i32 1} 2343!10 = !{i32 1, !"wchar_size", i32 4} 2344!11 = !{i32 7, !"openmp", i32 50} 2345!12 = !{i32 7, !"openmp-device", i32 50} 2346!13 = distinct !{!13, !14} 2347!14 = !{!"llvm.loop.mustprogress"} 2348!15 = distinct !{!15, !14} 2349!16 = distinct !{!16, !14} 2350!17 = distinct !{!17, !14} 2351;. 2352; AMDGPU: attributes #[[ATTR0]] = { 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" } 2353; AMDGPU: attributes #[[ATTR1]] = { "llvm.assume"="ompx_spmd_amenable" } 2354; AMDGPU: attributes #[[ATTR2:[0-9]+]] = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32,+sm_20" } 2355; AMDGPU: attributes #[[ATTR3:[0-9]+]] = { alwaysinline } 2356; AMDGPU: attributes #[[ATTR4]] = { nounwind } 2357; AMDGPU: attributes #[[ATTR5:[0-9]+]] = { nosync nounwind } 2358; AMDGPU: attributes #[[ATTR6:[0-9]+]] = { convergent nounwind } 2359; AMDGPU: attributes #[[ATTR7]] = { convergent "llvm.assume"="ompx_spmd_amenable" } 2360; AMDGPU: attributes #[[ATTR8]] = { convergent } 2361;. 2362; NVPTX: attributes #[[ATTR0]] = { 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" } 2363; NVPTX: attributes #[[ATTR1]] = { "llvm.assume"="ompx_spmd_amenable" } 2364; NVPTX: attributes #[[ATTR2:[0-9]+]] = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32,+sm_20" } 2365; NVPTX: attributes #[[ATTR3:[0-9]+]] = { alwaysinline } 2366; NVPTX: attributes #[[ATTR4]] = { nounwind } 2367; NVPTX: attributes #[[ATTR5:[0-9]+]] = { nosync nounwind } 2368; NVPTX: attributes #[[ATTR6:[0-9]+]] = { convergent nounwind } 2369; NVPTX: attributes #[[ATTR7]] = { convergent "llvm.assume"="ompx_spmd_amenable" } 2370; NVPTX: attributes #[[ATTR8]] = { convergent } 2371;. 2372; AMDGPU-DISABLED: attributes #[[ATTR0]] = { 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" } 2373; AMDGPU-DISABLED: attributes #[[ATTR1]] = { "llvm.assume"="ompx_spmd_amenable" } 2374; AMDGPU-DISABLED: attributes #[[ATTR2:[0-9]+]] = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32,+sm_20" } 2375; AMDGPU-DISABLED: attributes #[[ATTR3:[0-9]+]] = { alwaysinline } 2376; AMDGPU-DISABLED: attributes #[[ATTR4]] = { nounwind } 2377; AMDGPU-DISABLED: attributes #[[ATTR5:[0-9]+]] = { nosync nounwind } 2378; AMDGPU-DISABLED: attributes #[[ATTR6:[0-9]+]] = { convergent nounwind } 2379; AMDGPU-DISABLED: attributes #[[ATTR7]] = { convergent "llvm.assume"="ompx_spmd_amenable" } 2380; AMDGPU-DISABLED: attributes #[[ATTR8]] = { convergent } 2381;. 2382; NVPTX-DISABLED: attributes #[[ATTR0]] = { 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" } 2383; NVPTX-DISABLED: attributes #[[ATTR1]] = { "llvm.assume"="ompx_spmd_amenable" } 2384; NVPTX-DISABLED: attributes #[[ATTR2:[0-9]+]] = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32,+sm_20" } 2385; NVPTX-DISABLED: attributes #[[ATTR3:[0-9]+]] = { alwaysinline } 2386; NVPTX-DISABLED: attributes #[[ATTR4]] = { nounwind } 2387; NVPTX-DISABLED: attributes #[[ATTR5:[0-9]+]] = { nosync nounwind } 2388; NVPTX-DISABLED: attributes #[[ATTR6:[0-9]+]] = { convergent nounwind } 2389; NVPTX-DISABLED: attributes #[[ATTR7]] = { convergent "llvm.assume"="ompx_spmd_amenable" } 2390; NVPTX-DISABLED: attributes #[[ATTR8]] = { convergent } 2391;. 2392; AMDGPU: [[META0:![0-9]+]] = !{i32 0, i32 20, i32 171231761, !"sequential_loop_to_stack_var", i32 20, i32 1} 2393; AMDGPU: [[META1:![0-9]+]] = !{i32 0, i32 20, i32 171231761, !"sequential_loop", i32 5, i32 0} 2394; AMDGPU: [[META2:![0-9]+]] = !{i32 0, i32 20, i32 171231761, !"sequential_loop_to_shared_var", i32 35, i32 2} 2395; AMDGPU: [[META3:![0-9]+]] = !{i32 0, i32 20, i32 171231761, !"do_not_spmdize_target", i32 65, i32 4} 2396; AMDGPU: [[META4:![0-9]+]] = !{i32 0, i32 20, i32 171231761, !"sequential_loop_to_shared_var_guarded", i32 50, i32 3} 2397; AMDGPU: [[META5:![0-9]+]] = !{void ()* @__omp_offloading_14_a34ca11_sequential_loop_l5, !"kernel", i32 1} 2398; AMDGPU: [[META6:![0-9]+]] = !{void ()* @__omp_offloading_14_a34ca11_sequential_loop_to_stack_var_l20, !"kernel", i32 1} 2399; AMDGPU: [[META7:![0-9]+]] = !{void ()* @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_l35, !"kernel", i32 1} 2400; AMDGPU: [[META8:![0-9]+]] = !{void ()* @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_guarded_l50, !"kernel", i32 1} 2401; AMDGPU: [[META9:![0-9]+]] = !{void ()* @__omp_offloading_14_a34ca11_do_not_spmdize_target_l65, !"kernel", i32 1} 2402; AMDGPU: [[META10:![0-9]+]] = !{i32 1, !"wchar_size", i32 4} 2403; AMDGPU: [[META11:![0-9]+]] = !{i32 7, !"openmp", i32 50} 2404; AMDGPU: [[META12:![0-9]+]] = !{i32 7, !"openmp-device", i32 50} 2405; AMDGPU: [[LOOP13]] = distinct !{!13, !14} 2406; AMDGPU: [[META14:![0-9]+]] = !{!"llvm.loop.mustprogress"} 2407; AMDGPU: [[LOOP15]] = distinct !{!15, !14} 2408; AMDGPU: [[LOOP16]] = distinct !{!16, !14} 2409; AMDGPU: [[LOOP17]] = distinct !{!17, !14} 2410;. 2411; NVPTX: [[META0:![0-9]+]] = !{i32 0, i32 20, i32 171231761, !"sequential_loop_to_stack_var", i32 20, i32 1} 2412; NVPTX: [[META1:![0-9]+]] = !{i32 0, i32 20, i32 171231761, !"sequential_loop", i32 5, i32 0} 2413; NVPTX: [[META2:![0-9]+]] = !{i32 0, i32 20, i32 171231761, !"sequential_loop_to_shared_var", i32 35, i32 2} 2414; NVPTX: [[META3:![0-9]+]] = !{i32 0, i32 20, i32 171231761, !"do_not_spmdize_target", i32 65, i32 4} 2415; NVPTX: [[META4:![0-9]+]] = !{i32 0, i32 20, i32 171231761, !"sequential_loop_to_shared_var_guarded", i32 50, i32 3} 2416; NVPTX: [[META5:![0-9]+]] = !{void ()* @__omp_offloading_14_a34ca11_sequential_loop_l5, !"kernel", i32 1} 2417; NVPTX: [[META6:![0-9]+]] = !{void ()* @__omp_offloading_14_a34ca11_sequential_loop_to_stack_var_l20, !"kernel", i32 1} 2418; NVPTX: [[META7:![0-9]+]] = !{void ()* @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_l35, !"kernel", i32 1} 2419; NVPTX: [[META8:![0-9]+]] = !{void ()* @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_guarded_l50, !"kernel", i32 1} 2420; NVPTX: [[META9:![0-9]+]] = !{void ()* @__omp_offloading_14_a34ca11_do_not_spmdize_target_l65, !"kernel", i32 1} 2421; NVPTX: [[META10:![0-9]+]] = !{i32 1, !"wchar_size", i32 4} 2422; NVPTX: [[META11:![0-9]+]] = !{i32 7, !"openmp", i32 50} 2423; NVPTX: [[META12:![0-9]+]] = !{i32 7, !"openmp-device", i32 50} 2424; NVPTX: [[LOOP13]] = distinct !{!13, !14} 2425; NVPTX: [[META14:![0-9]+]] = !{!"llvm.loop.mustprogress"} 2426; NVPTX: [[LOOP15]] = distinct !{!15, !14} 2427; NVPTX: [[LOOP16]] = distinct !{!16, !14} 2428; NVPTX: [[LOOP17]] = distinct !{!17, !14} 2429;. 2430; AMDGPU-DISABLED: [[META0:![0-9]+]] = !{i32 0, i32 20, i32 171231761, !"sequential_loop_to_stack_var", i32 20, i32 1} 2431; AMDGPU-DISABLED: [[META1:![0-9]+]] = !{i32 0, i32 20, i32 171231761, !"sequential_loop", i32 5, i32 0} 2432; AMDGPU-DISABLED: [[META2:![0-9]+]] = !{i32 0, i32 20, i32 171231761, !"sequential_loop_to_shared_var", i32 35, i32 2} 2433; AMDGPU-DISABLED: [[META3:![0-9]+]] = !{i32 0, i32 20, i32 171231761, !"do_not_spmdize_target", i32 65, i32 4} 2434; AMDGPU-DISABLED: [[META4:![0-9]+]] = !{i32 0, i32 20, i32 171231761, !"sequential_loop_to_shared_var_guarded", i32 50, i32 3} 2435; AMDGPU-DISABLED: [[META5:![0-9]+]] = !{void ()* @__omp_offloading_14_a34ca11_sequential_loop_l5, !"kernel", i32 1} 2436; AMDGPU-DISABLED: [[META6:![0-9]+]] = !{void ()* @__omp_offloading_14_a34ca11_sequential_loop_to_stack_var_l20, !"kernel", i32 1} 2437; AMDGPU-DISABLED: [[META7:![0-9]+]] = !{void ()* @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_l35, !"kernel", i32 1} 2438; AMDGPU-DISABLED: [[META8:![0-9]+]] = !{void ()* @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_guarded_l50, !"kernel", i32 1} 2439; AMDGPU-DISABLED: [[META9:![0-9]+]] = !{void ()* @__omp_offloading_14_a34ca11_do_not_spmdize_target_l65, !"kernel", i32 1} 2440; AMDGPU-DISABLED: [[META10:![0-9]+]] = !{i32 1, !"wchar_size", i32 4} 2441; AMDGPU-DISABLED: [[META11:![0-9]+]] = !{i32 7, !"openmp", i32 50} 2442; AMDGPU-DISABLED: [[META12:![0-9]+]] = !{i32 7, !"openmp-device", i32 50} 2443; AMDGPU-DISABLED: [[LOOP13]] = distinct !{!13, !14} 2444; AMDGPU-DISABLED: [[META14:![0-9]+]] = !{!"llvm.loop.mustprogress"} 2445; AMDGPU-DISABLED: [[LOOP15]] = distinct !{!15, !14} 2446; AMDGPU-DISABLED: [[LOOP16]] = distinct !{!16, !14} 2447; AMDGPU-DISABLED: [[LOOP17]] = distinct !{!17, !14} 2448;. 2449; NVPTX-DISABLED: [[META0:![0-9]+]] = !{i32 0, i32 20, i32 171231761, !"sequential_loop_to_stack_var", i32 20, i32 1} 2450; NVPTX-DISABLED: [[META1:![0-9]+]] = !{i32 0, i32 20, i32 171231761, !"sequential_loop", i32 5, i32 0} 2451; NVPTX-DISABLED: [[META2:![0-9]+]] = !{i32 0, i32 20, i32 171231761, !"sequential_loop_to_shared_var", i32 35, i32 2} 2452; NVPTX-DISABLED: [[META3:![0-9]+]] = !{i32 0, i32 20, i32 171231761, !"do_not_spmdize_target", i32 65, i32 4} 2453; NVPTX-DISABLED: [[META4:![0-9]+]] = !{i32 0, i32 20, i32 171231761, !"sequential_loop_to_shared_var_guarded", i32 50, i32 3} 2454; NVPTX-DISABLED: [[META5:![0-9]+]] = !{void ()* @__omp_offloading_14_a34ca11_sequential_loop_l5, !"kernel", i32 1} 2455; NVPTX-DISABLED: [[META6:![0-9]+]] = !{void ()* @__omp_offloading_14_a34ca11_sequential_loop_to_stack_var_l20, !"kernel", i32 1} 2456; NVPTX-DISABLED: [[META7:![0-9]+]] = !{void ()* @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_l35, !"kernel", i32 1} 2457; NVPTX-DISABLED: [[META8:![0-9]+]] = !{void ()* @__omp_offloading_14_a34ca11_sequential_loop_to_shared_var_guarded_l50, !"kernel", i32 1} 2458; NVPTX-DISABLED: [[META9:![0-9]+]] = !{void ()* @__omp_offloading_14_a34ca11_do_not_spmdize_target_l65, !"kernel", i32 1} 2459; NVPTX-DISABLED: [[META10:![0-9]+]] = !{i32 1, !"wchar_size", i32 4} 2460; NVPTX-DISABLED: [[META11:![0-9]+]] = !{i32 7, !"openmp", i32 50} 2461; NVPTX-DISABLED: [[META12:![0-9]+]] = !{i32 7, !"openmp-device", i32 50} 2462; NVPTX-DISABLED: [[LOOP13]] = distinct !{!13, !14} 2463; NVPTX-DISABLED: [[META14:![0-9]+]] = !{!"llvm.loop.mustprogress"} 2464; NVPTX-DISABLED: [[LOOP15]] = distinct !{!15, !14} 2465; NVPTX-DISABLED: [[LOOP16]] = distinct !{!16, !14} 2466; NVPTX-DISABLED: [[LOOP17]] = distinct !{!17, !14} 2467;. 2468