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