1273 lines · cpp
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _2// REQUIRES: amdgpu-registered-target3 4// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc5// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s6// expected-no-diagnostics7#ifndef HEADER8#define HEADER9 10int foo1() {11 int sum = 0.0;12 #pragma omp target map(tofrom: sum)13 {14 int N = 10;15 int A[N];16 17 for (int i = 0; i < N; i++)18 A[i] = i;19 20 for (int i = 0; i < N; i++)21 sum += A[i];22 }23 return sum;24}25 26int foo2() {27 int sum = 0.0;28 int M = 12;29 int result[M];30 #pragma omp target teams distribute parallel for map(from: result[:M])31 for (int i = 0; i < M; i++) {32 int N = 10;33 int A[N];34 result[i] = i;35 36 for (int j = 0; j < N; j++)37 A[j] = j;38 39 for (int j = 0; j < N; j++)40 result[i] += A[j];41 }42 43 for (int i = 0; i < M; i++)44 sum += result[i];45 return sum;46}47 48int foo3() {49 int sum = 0.0;50 int M = 12;51 int result[M];52 #pragma omp target teams distribute map(from: result[:M])53 for (int i = 0; i < M; i++) {54 int N = 10;55 int A[N];56 result[i] = i;57 58 #pragma omp parallel for59 for (int j = 0; j < N; j++)60 A[j] = j;61 62 for (int j = 0; j < N; j++)63 result[i] += A[j];64 }65 66 for (int i = 0; i < M; i++)67 sum += result[i];68 return sum;69}70 71int foo4() {72 int sum = 0.0;73 int M = 12;74 int result[M];75 int N = 10;76 #pragma omp target teams distribute map(from: result[:M])77 for (int i = 0; i < M; i++) {78 int A[N];79 result[i] = i;80 81 #pragma omp parallel for82 for (int j = 0; j < N; j++)83 A[j] = j;84 85 for (int j = 0; j < N; j++)86 result[i] += A[j];87 }88 89 for (int i = 0; i < M; i++)90 sum += result[i];91 return sum;92}93 94int main() {95 return foo1() + foo2() + foo3() + foo4();96}97 98#endif99// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo1v_l12100// CHECK-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SUM:%.*]]) #[[ATTR0:[0-9]+]] {101// CHECK-NEXT: entry:102// CHECK-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)103// CHECK-NEXT: [[SUM_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)104// CHECK-NEXT: [[N:%.*]] = alloca i32, align 4, addrspace(5)105// CHECK-NEXT: [[__VLA_EXPR0:%.*]] = alloca i64, align 8, addrspace(5)106// CHECK-NEXT: [[I:%.*]] = alloca i32, align 4, addrspace(5)107// CHECK-NEXT: [[I1:%.*]] = alloca i32, align 4, addrspace(5)108// CHECK-NEXT: [[DYN_PTR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DYN_PTR_ADDR]] to ptr109// CHECK-NEXT: [[SUM_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SUM_ADDR]] to ptr110// CHECK-NEXT: [[N_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[N]] to ptr111// CHECK-NEXT: [[__VLA_EXPR0_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__VLA_EXPR0]] to ptr112// CHECK-NEXT: [[I_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[I]] to ptr113// CHECK-NEXT: [[I1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[I1]] to ptr114// CHECK-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR_ASCAST]], align 8115// CHECK-NEXT: store ptr [[SUM]], ptr [[SUM_ADDR_ASCAST]], align 8116// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SUM_ADDR_ASCAST]], align 8117// CHECK-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo1v_l12_kernel_environment to ptr), ptr [[DYN_PTR]])118// CHECK-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1119// CHECK-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]120// CHECK: user_code.entry:121// CHECK-NEXT: store i32 10, ptr [[N_ASCAST]], align 4122// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ASCAST]], align 4123// CHECK-NEXT: [[TMP3:%.*]] = zext i32 [[TMP2]] to i64124// CHECK-NEXT: [[TMP4:%.*]] = mul nuw i64 [[TMP3]], 4125// CHECK-NEXT: [[TMP5:%.*]] = add nuw i64 [[TMP4]], 3126// CHECK-NEXT: [[TMP6:%.*]] = udiv i64 [[TMP5]], 4127// CHECK-NEXT: [[TMP7:%.*]] = mul nuw i64 [[TMP6]], 4128// CHECK-NEXT: [[A:%.*]] = call align 4 ptr @__kmpc_alloc_shared(i64 [[TMP7]])129// CHECK-NEXT: store i64 [[TMP3]], ptr [[__VLA_EXPR0_ASCAST]], align 8130// CHECK-NEXT: store i32 0, ptr [[I_ASCAST]], align 4131// CHECK-NEXT: br label [[FOR_COND:%.*]]132// CHECK: for.cond:133// CHECK-NEXT: [[TMP8:%.*]] = load i32, ptr [[I_ASCAST]], align 4134// CHECK-NEXT: [[TMP9:%.*]] = load i32, ptr [[N_ASCAST]], align 4135// CHECK-NEXT: [[CMP:%.*]] = icmp slt i32 [[TMP8]], [[TMP9]]136// CHECK-NEXT: br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]]137// CHECK: for.body:138// CHECK-NEXT: [[TMP10:%.*]] = load i32, ptr [[I_ASCAST]], align 4139// CHECK-NEXT: [[TMP11:%.*]] = load i32, ptr [[I_ASCAST]], align 4140// CHECK-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP11]] to i64141// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[A]], i64 [[IDXPROM]]142// CHECK-NEXT: store i32 [[TMP10]], ptr [[ARRAYIDX]], align 4143// CHECK-NEXT: br label [[FOR_INC:%.*]]144// CHECK: for.inc:145// CHECK-NEXT: [[TMP12:%.*]] = load i32, ptr [[I_ASCAST]], align 4146// CHECK-NEXT: [[INC:%.*]] = add nsw i32 [[TMP12]], 1147// CHECK-NEXT: store i32 [[INC]], ptr [[I_ASCAST]], align 4148// CHECK-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP13:![0-9]+]]149// CHECK: worker.exit:150// CHECK-NEXT: ret void151// CHECK: for.end:152// CHECK-NEXT: store i32 0, ptr [[I1_ASCAST]], align 4153// CHECK-NEXT: br label [[FOR_COND2:%.*]]154// CHECK: for.cond2:155// CHECK-NEXT: [[TMP13:%.*]] = load i32, ptr [[I1_ASCAST]], align 4156// CHECK-NEXT: [[TMP14:%.*]] = load i32, ptr [[N_ASCAST]], align 4157// CHECK-NEXT: [[CMP3:%.*]] = icmp slt i32 [[TMP13]], [[TMP14]]158// CHECK-NEXT: br i1 [[CMP3]], label [[FOR_BODY4:%.*]], label [[FOR_END9:%.*]]159// CHECK: for.body4:160// CHECK-NEXT: [[TMP15:%.*]] = load i32, ptr [[I1_ASCAST]], align 4161// CHECK-NEXT: [[IDXPROM5:%.*]] = sext i32 [[TMP15]] to i64162// CHECK-NEXT: [[ARRAYIDX6:%.*]] = getelementptr inbounds i32, ptr [[A]], i64 [[IDXPROM5]]163// CHECK-NEXT: [[TMP16:%.*]] = load i32, ptr [[ARRAYIDX6]], align 4164// CHECK-NEXT: [[TMP17:%.*]] = load i32, ptr [[TMP0]], align 4165// CHECK-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP17]], [[TMP16]]166// CHECK-NEXT: store i32 [[ADD]], ptr [[TMP0]], align 4167// CHECK-NEXT: br label [[FOR_INC7:%.*]]168// CHECK: for.inc7:169// CHECK-NEXT: [[TMP18:%.*]] = load i32, ptr [[I1_ASCAST]], align 4170// CHECK-NEXT: [[INC8:%.*]] = add nsw i32 [[TMP18]], 1171// CHECK-NEXT: store i32 [[INC8]], ptr [[I1_ASCAST]], align 4172// CHECK-NEXT: br label [[FOR_COND2]], !llvm.loop [[LOOP15:![0-9]+]]173// CHECK: for.end9:174// CHECK-NEXT: call void @__kmpc_free_shared(ptr [[A]], i64 [[TMP7]])175// CHECK-NEXT: call void @__kmpc_target_deinit()176// CHECK-NEXT: ret void177//178//179// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo2v_l30180// CHECK-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[M:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[RESULT:%.*]]) #[[ATTR0]] {181// CHECK-NEXT: entry:182// CHECK-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)183// CHECK-NEXT: [[M_ADDR:%.*]] = alloca i64, align 8, addrspace(5)184// CHECK-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 8, addrspace(5)185// CHECK-NEXT: [[RESULT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)186// CHECK-NEXT: [[M_CASTED:%.*]] = alloca i64, align 8, addrspace(5)187// CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4, addrspace(5)188// CHECK-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4, addrspace(5)189// CHECK-NEXT: [[DYN_PTR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DYN_PTR_ADDR]] to ptr190// CHECK-NEXT: [[M_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[M_ADDR]] to ptr191// CHECK-NEXT: [[VLA_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[VLA_ADDR]] to ptr192// CHECK-NEXT: [[RESULT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RESULT_ADDR]] to ptr193// CHECK-NEXT: [[M_CASTED_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[M_CASTED]] to ptr194// CHECK-NEXT: [[DOTZERO_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTZERO_ADDR]] to ptr195// CHECK-NEXT: [[DOTTHREADID_TEMP__ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTTHREADID_TEMP_]] to ptr196// CHECK-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR_ASCAST]], align 8197// CHECK-NEXT: store i64 [[M]], ptr [[M_ADDR_ASCAST]], align 8198// CHECK-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR_ASCAST]], align 8199// CHECK-NEXT: store ptr [[RESULT]], ptr [[RESULT_ADDR_ASCAST]], align 8200// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR_ASCAST]], align 8201// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[RESULT_ADDR_ASCAST]], align 8202// CHECK-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo2v_l30_kernel_environment to ptr), ptr [[DYN_PTR]])203// CHECK-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -1204// CHECK-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]205// CHECK: user_code.entry:206// CHECK-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_global_thread_num(ptr addrspacecast (ptr addrspace(1) @[[GLOB1:[0-9]+]] to ptr))207// CHECK-NEXT: [[TMP4:%.*]] = load i32, ptr [[M_ADDR_ASCAST]], align 4208// CHECK-NEXT: store i32 [[TMP4]], ptr [[M_CASTED_ASCAST]], align 4209// CHECK-NEXT: [[TMP5:%.*]] = load i64, ptr [[M_CASTED_ASCAST]], align 8210// CHECK-NEXT: store i32 0, ptr [[DOTZERO_ADDR_ASCAST]], align 4211// CHECK-NEXT: store i32 [[TMP3]], ptr [[DOTTHREADID_TEMP__ASCAST]], align 4212// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo2v_l30_omp_outlined(ptr [[DOTTHREADID_TEMP__ASCAST]], ptr [[DOTZERO_ADDR_ASCAST]], i64 [[TMP5]], i64 [[TMP0]], ptr [[TMP1]]) #[[ATTR4:[0-9]+]]213// CHECK-NEXT: call void @__kmpc_target_deinit()214// CHECK-NEXT: ret void215// CHECK: worker.exit:216// CHECK-NEXT: ret void217//218//219// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo2v_l30_omp_outlined220// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[M:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[RESULT:%.*]]) #[[ATTR3:[0-9]+]] {221// CHECK-NEXT: entry:222// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)223// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)224// CHECK-NEXT: [[M_ADDR:%.*]] = alloca i64, align 8, addrspace(5)225// CHECK-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 8, addrspace(5)226// CHECK-NEXT: [[RESULT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)227// CHECK-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4, addrspace(5)228// CHECK-NEXT: [[TMP:%.*]] = alloca i32, align 4, addrspace(5)229// CHECK-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4, addrspace(5)230// CHECK-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4, addrspace(5)231// CHECK-NEXT: [[I:%.*]] = alloca i32, align 4, addrspace(5)232// CHECK-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4, addrspace(5)233// CHECK-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4, addrspace(5)234// CHECK-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4, addrspace(5)235// CHECK-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4, addrspace(5)236// CHECK-NEXT: [[I3:%.*]] = alloca i32, align 4, addrspace(5)237// CHECK-NEXT: [[M_CASTED:%.*]] = alloca i64, align 8, addrspace(5)238// CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [5 x ptr], align 8, addrspace(5)239// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTGLOBAL_TID__ADDR]] to ptr240// CHECK-NEXT: [[DOTBOUND_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTBOUND_TID__ADDR]] to ptr241// CHECK-NEXT: [[M_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[M_ADDR]] to ptr242// CHECK-NEXT: [[VLA_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[VLA_ADDR]] to ptr243// CHECK-NEXT: [[RESULT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RESULT_ADDR]] to ptr244// CHECK-NEXT: [[DOTOMP_IV_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IV]] to ptr245// CHECK-NEXT: [[TMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[TMP]] to ptr246// CHECK-NEXT: [[DOTCAPTURE_EXPR__ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTCAPTURE_EXPR_]] to ptr247// CHECK-NEXT: [[DOTCAPTURE_EXPR_1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTCAPTURE_EXPR_1]] to ptr248// CHECK-NEXT: [[I_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[I]] to ptr249// CHECK-NEXT: [[DOTOMP_COMB_LB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_COMB_LB]] to ptr250// CHECK-NEXT: [[DOTOMP_COMB_UB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_COMB_UB]] to ptr251// CHECK-NEXT: [[DOTOMP_STRIDE_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_STRIDE]] to ptr252// CHECK-NEXT: [[DOTOMP_IS_LAST_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IS_LAST]] to ptr253// CHECK-NEXT: [[I3_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[I3]] to ptr254// CHECK-NEXT: [[M_CASTED_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[M_CASTED]] to ptr255// CHECK-NEXT: [[CAPTURED_VARS_ADDRS_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[CAPTURED_VARS_ADDRS]] to ptr256// CHECK-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8257// CHECK-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR_ASCAST]], align 8258// CHECK-NEXT: store i64 [[M]], ptr [[M_ADDR_ASCAST]], align 8259// CHECK-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR_ASCAST]], align 8260// CHECK-NEXT: store ptr [[RESULT]], ptr [[RESULT_ADDR_ASCAST]], align 8261// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR_ASCAST]], align 8262// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[RESULT_ADDR_ASCAST]], align 8263// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[M_ADDR_ASCAST]], align 4264// CHECK-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR__ASCAST]], align 4265// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ASCAST]], align 4266// CHECK-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0267// CHECK-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1268// CHECK-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1269// CHECK-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1_ASCAST]], align 4270// CHECK-NEXT: store i32 0, ptr [[I_ASCAST]], align 4271// CHECK-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ASCAST]], align 4272// CHECK-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]273// CHECK-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]274// CHECK: omp.precond.then:275// CHECK-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB_ASCAST]], align 4276// CHECK-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1_ASCAST]], align 4277// CHECK-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_COMB_UB_ASCAST]], align 4278// CHECK-NEXT: store i32 1, ptr [[DOTOMP_STRIDE_ASCAST]], align 4279// CHECK-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST_ASCAST]], align 4280// CHECK-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @__kmpc_get_hardware_num_threads_in_block()281// CHECK-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8282// CHECK-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 4283// CHECK-NEXT: call void @__kmpc_distribute_static_init_4(ptr addrspacecast (ptr addrspace(1) @[[GLOB2:[0-9]+]] to ptr), i32 [[TMP7]], i32 91, ptr [[DOTOMP_IS_LAST_ASCAST]], ptr [[DOTOMP_COMB_LB_ASCAST]], ptr [[DOTOMP_COMB_UB_ASCAST]], ptr [[DOTOMP_STRIDE_ASCAST]], i32 1, i32 [[NVPTX_NUM_THREADS]])284// CHECK-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_COMB_UB_ASCAST]], align 4285// CHECK-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1_ASCAST]], align 4286// CHECK-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]287// CHECK-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]288// CHECK: cond.true:289// CHECK-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1_ASCAST]], align 4290// CHECK-NEXT: br label [[COND_END:%.*]]291// CHECK: cond.false:292// CHECK-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_UB_ASCAST]], align 4293// CHECK-NEXT: br label [[COND_END]]294// CHECK: cond.end:295// CHECK-NEXT: [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]296// CHECK-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB_ASCAST]], align 4297// CHECK-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_COMB_LB_ASCAST]], align 4298// CHECK-NEXT: store i32 [[TMP12]], ptr [[DOTOMP_IV_ASCAST]], align 4299// CHECK-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]300// CHECK: omp.inner.for.cond:301// CHECK-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4302// CHECK-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1_ASCAST]], align 4303// CHECK-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP14]], 1304// CHECK-NEXT: [[CMP5:%.*]] = icmp slt i32 [[TMP13]], [[ADD]]305// CHECK-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]306// CHECK: omp.inner.for.body:307// CHECK-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_COMB_LB_ASCAST]], align 4308// CHECK-NEXT: [[TMP16:%.*]] = zext i32 [[TMP15]] to i64309// CHECK-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_COMB_UB_ASCAST]], align 4310// CHECK-NEXT: [[TMP18:%.*]] = zext i32 [[TMP17]] to i64311// CHECK-NEXT: [[TMP19:%.*]] = load i32, ptr [[M_ADDR_ASCAST]], align 4312// CHECK-NEXT: store i32 [[TMP19]], ptr [[M_CASTED_ASCAST]], align 4313// CHECK-NEXT: [[TMP20:%.*]] = load i64, ptr [[M_CASTED_ASCAST]], align 8314// CHECK-NEXT: [[TMP21:%.*]] = getelementptr inbounds [5 x ptr], ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 0, i64 0315// CHECK-NEXT: [[TMP22:%.*]] = inttoptr i64 [[TMP16]] to ptr316// CHECK-NEXT: store ptr [[TMP22]], ptr [[TMP21]], align 8317// CHECK-NEXT: [[TMP23:%.*]] = getelementptr inbounds [5 x ptr], ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 0, i64 1318// CHECK-NEXT: [[TMP24:%.*]] = inttoptr i64 [[TMP18]] to ptr319// CHECK-NEXT: store ptr [[TMP24]], ptr [[TMP23]], align 8320// CHECK-NEXT: [[TMP25:%.*]] = getelementptr inbounds [5 x ptr], ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 0, i64 2321// CHECK-NEXT: [[TMP26:%.*]] = inttoptr i64 [[TMP20]] to ptr322// CHECK-NEXT: store ptr [[TMP26]], ptr [[TMP25]], align 8323// CHECK-NEXT: [[TMP27:%.*]] = getelementptr inbounds [5 x ptr], ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 0, i64 3324// CHECK-NEXT: [[TMP28:%.*]] = inttoptr i64 [[TMP0]] to ptr325// CHECK-NEXT: store ptr [[TMP28]], ptr [[TMP27]], align 8326// CHECK-NEXT: [[TMP29:%.*]] = getelementptr inbounds [5 x ptr], ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 0, i64 4327// CHECK-NEXT: store ptr [[TMP1]], ptr [[TMP29]], align 8328// CHECK-NEXT: [[TMP30:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8329// CHECK-NEXT: [[TMP31:%.*]] = load i32, ptr [[TMP30]], align 4330// CHECK-NEXT: call void @__kmpc_parallel_51(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i32 [[TMP31]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo2v_l30_omp_outlined_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 5)331// CHECK-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]332// CHECK: omp.inner.for.inc:333// CHECK-NEXT: [[TMP32:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4334// CHECK-NEXT: [[TMP33:%.*]] = load i32, ptr [[DOTOMP_STRIDE_ASCAST]], align 4335// CHECK-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP32]], [[TMP33]]336// CHECK-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV_ASCAST]], align 4337// CHECK-NEXT: [[TMP34:%.*]] = load i32, ptr [[DOTOMP_COMB_LB_ASCAST]], align 4338// CHECK-NEXT: [[TMP35:%.*]] = load i32, ptr [[DOTOMP_STRIDE_ASCAST]], align 4339// CHECK-NEXT: [[ADD7:%.*]] = add nsw i32 [[TMP34]], [[TMP35]]340// CHECK-NEXT: store i32 [[ADD7]], ptr [[DOTOMP_COMB_LB_ASCAST]], align 4341// CHECK-NEXT: [[TMP36:%.*]] = load i32, ptr [[DOTOMP_COMB_UB_ASCAST]], align 4342// CHECK-NEXT: [[TMP37:%.*]] = load i32, ptr [[DOTOMP_STRIDE_ASCAST]], align 4343// CHECK-NEXT: [[ADD8:%.*]] = add nsw i32 [[TMP36]], [[TMP37]]344// CHECK-NEXT: store i32 [[ADD8]], ptr [[DOTOMP_COMB_UB_ASCAST]], align 4345// CHECK-NEXT: [[TMP38:%.*]] = load i32, ptr [[DOTOMP_COMB_UB_ASCAST]], align 4346// CHECK-NEXT: [[TMP39:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1_ASCAST]], align 4347// CHECK-NEXT: [[CMP9:%.*]] = icmp sgt i32 [[TMP38]], [[TMP39]]348// CHECK-NEXT: br i1 [[CMP9]], label [[COND_TRUE10:%.*]], label [[COND_FALSE11:%.*]]349// CHECK: cond.true10:350// CHECK-NEXT: [[TMP40:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1_ASCAST]], align 4351// CHECK-NEXT: br label [[COND_END12:%.*]]352// CHECK: cond.false11:353// CHECK-NEXT: [[TMP41:%.*]] = load i32, ptr [[DOTOMP_COMB_UB_ASCAST]], align 4354// CHECK-NEXT: br label [[COND_END12]]355// CHECK: cond.end12:356// CHECK-NEXT: [[COND13:%.*]] = phi i32 [ [[TMP40]], [[COND_TRUE10]] ], [ [[TMP41]], [[COND_FALSE11]] ]357// CHECK-NEXT: store i32 [[COND13]], ptr [[DOTOMP_COMB_UB_ASCAST]], align 4358// CHECK-NEXT: [[TMP42:%.*]] = load i32, ptr [[DOTOMP_COMB_LB_ASCAST]], align 4359// CHECK-NEXT: store i32 [[TMP42]], ptr [[DOTOMP_IV_ASCAST]], align 4360// CHECK-NEXT: br label [[OMP_INNER_FOR_COND]]361// CHECK: omp.inner.for.end:362// CHECK-NEXT: br label [[OMP_LOOP_EXIT:%.*]]363// CHECK: omp.loop.exit:364// CHECK-NEXT: [[TMP43:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8365// CHECK-NEXT: [[TMP44:%.*]] = load i32, ptr [[TMP43]], align 4366// CHECK-NEXT: call void @__kmpc_distribute_static_fini(ptr addrspacecast (ptr addrspace(1) @[[GLOB2]] to ptr), i32 [[TMP44]])367// CHECK-NEXT: br label [[OMP_PRECOND_END]]368// CHECK: omp.precond.end:369// CHECK-NEXT: ret void370//371//372// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo2v_l30_omp_outlined_omp_outlined373// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]], i64 noundef [[M:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[RESULT:%.*]]) #[[ATTR3]] {374// CHECK-NEXT: entry:375// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)376// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)377// CHECK-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8, addrspace(5)378// CHECK-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8, addrspace(5)379// CHECK-NEXT: [[M_ADDR:%.*]] = alloca i64, align 8, addrspace(5)380// CHECK-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 8, addrspace(5)381// CHECK-NEXT: [[RESULT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)382// CHECK-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4, addrspace(5)383// CHECK-NEXT: [[TMP:%.*]] = alloca i32, align 4, addrspace(5)384// CHECK-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4, addrspace(5)385// CHECK-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4, addrspace(5)386// CHECK-NEXT: [[I:%.*]] = alloca i32, align 4, addrspace(5)387// CHECK-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4, addrspace(5)388// CHECK-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4, addrspace(5)389// CHECK-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4, addrspace(5)390// CHECK-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4, addrspace(5)391// CHECK-NEXT: [[I4:%.*]] = alloca i32, align 4, addrspace(5)392// CHECK-NEXT: [[N:%.*]] = alloca i32, align 4, addrspace(5)393// CHECK-NEXT: [[SAVED_STACK:%.*]] = alloca ptr addrspace(5), align 4, addrspace(5)394// CHECK-NEXT: [[__VLA_EXPR0:%.*]] = alloca i64, align 8, addrspace(5)395// CHECK-NEXT: [[J:%.*]] = alloca i32, align 4, addrspace(5)396// CHECK-NEXT: [[J11:%.*]] = alloca i32, align 4, addrspace(5)397// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTGLOBAL_TID__ADDR]] to ptr398// CHECK-NEXT: [[DOTBOUND_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTBOUND_TID__ADDR]] to ptr399// CHECK-NEXT: [[DOTPREVIOUS_LB__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTPREVIOUS_LB__ADDR]] to ptr400// CHECK-NEXT: [[DOTPREVIOUS_UB__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTPREVIOUS_UB__ADDR]] to ptr401// CHECK-NEXT: [[M_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[M_ADDR]] to ptr402// CHECK-NEXT: [[VLA_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[VLA_ADDR]] to ptr403// CHECK-NEXT: [[RESULT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RESULT_ADDR]] to ptr404// CHECK-NEXT: [[DOTOMP_IV_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IV]] to ptr405// CHECK-NEXT: [[TMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[TMP]] to ptr406// CHECK-NEXT: [[DOTCAPTURE_EXPR__ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTCAPTURE_EXPR_]] to ptr407// CHECK-NEXT: [[DOTCAPTURE_EXPR_1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTCAPTURE_EXPR_1]] to ptr408// CHECK-NEXT: [[I_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[I]] to ptr409// CHECK-NEXT: [[DOTOMP_LB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_LB]] to ptr410// CHECK-NEXT: [[DOTOMP_UB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_UB]] to ptr411// CHECK-NEXT: [[DOTOMP_STRIDE_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_STRIDE]] to ptr412// CHECK-NEXT: [[DOTOMP_IS_LAST_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IS_LAST]] to ptr413// CHECK-NEXT: [[I4_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[I4]] to ptr414// CHECK-NEXT: [[N_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[N]] to ptr415// CHECK-NEXT: [[SAVED_STACK_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SAVED_STACK]] to ptr416// CHECK-NEXT: [[__VLA_EXPR0_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__VLA_EXPR0]] to ptr417// CHECK-NEXT: [[J_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[J]] to ptr418// CHECK-NEXT: [[J11_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[J11]] to ptr419// CHECK-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8420// CHECK-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR_ASCAST]], align 8421// CHECK-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR_ASCAST]], align 8422// CHECK-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR_ASCAST]], align 8423// CHECK-NEXT: store i64 [[M]], ptr [[M_ADDR_ASCAST]], align 8424// CHECK-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR_ASCAST]], align 8425// CHECK-NEXT: store ptr [[RESULT]], ptr [[RESULT_ADDR_ASCAST]], align 8426// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR_ASCAST]], align 8427// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[RESULT_ADDR_ASCAST]], align 8428// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[M_ADDR_ASCAST]], align 4429// CHECK-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR__ASCAST]], align 4430// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ASCAST]], align 4431// CHECK-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0432// CHECK-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1433// CHECK-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1434// CHECK-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1_ASCAST]], align 4435// CHECK-NEXT: store i32 0, ptr [[I_ASCAST]], align 4436// CHECK-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ASCAST]], align 4437// CHECK-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]438// CHECK-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]439// CHECK: omp.precond.then:440// CHECK-NEXT: store i32 0, ptr [[DOTOMP_LB_ASCAST]], align 4441// CHECK-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1_ASCAST]], align 4442// CHECK-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_UB_ASCAST]], align 4443// CHECK-NEXT: [[TMP6:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR_ASCAST]], align 8444// CHECK-NEXT: [[CONV:%.*]] = trunc i64 [[TMP6]] to i32445// CHECK-NEXT: [[TMP7:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR_ASCAST]], align 8446// CHECK-NEXT: [[CONV3:%.*]] = trunc i64 [[TMP7]] to i32447// CHECK-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB_ASCAST]], align 4448// CHECK-NEXT: store i32 [[CONV3]], ptr [[DOTOMP_UB_ASCAST]], align 4449// CHECK-NEXT: store i32 1, ptr [[DOTOMP_STRIDE_ASCAST]], align 4450// CHECK-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST_ASCAST]], align 4451// CHECK-NEXT: [[TMP8:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8452// CHECK-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP8]], align 4453// CHECK-NEXT: call void @__kmpc_for_static_init_4(ptr addrspacecast (ptr addrspace(1) @[[GLOB3:[0-9]+]] to ptr), i32 [[TMP9]], i32 33, ptr [[DOTOMP_IS_LAST_ASCAST]], ptr [[DOTOMP_LB_ASCAST]], ptr [[DOTOMP_UB_ASCAST]], ptr [[DOTOMP_STRIDE_ASCAST]], i32 1, i32 1)454// CHECK-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_LB_ASCAST]], align 4455// CHECK-NEXT: store i32 [[TMP10]], ptr [[DOTOMP_IV_ASCAST]], align 4456// CHECK-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]457// CHECK: omp.inner.for.cond:458// CHECK-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4459// CHECK-NEXT: [[CONV5:%.*]] = sext i32 [[TMP11]] to i64460// CHECK-NEXT: [[TMP12:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR_ASCAST]], align 8461// CHECK-NEXT: [[CMP6:%.*]] = icmp ule i64 [[CONV5]], [[TMP12]]462// CHECK-NEXT: br i1 [[CMP6]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]463// CHECK: omp.inner.for.body:464// CHECK-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4465// CHECK-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP13]], 1466// CHECK-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]467// CHECK-NEXT: store i32 [[ADD]], ptr [[I4_ASCAST]], align 4468// CHECK-NEXT: store i32 10, ptr [[N_ASCAST]], align 4469// CHECK-NEXT: [[TMP14:%.*]] = load i32, ptr [[N_ASCAST]], align 4470// CHECK-NEXT: [[TMP15:%.*]] = zext i32 [[TMP14]] to i64471// CHECK-NEXT: [[TMP16:%.*]] = call ptr addrspace(5) @llvm.stacksave.p5()472// CHECK-NEXT: store ptr addrspace(5) [[TMP16]], ptr [[SAVED_STACK_ASCAST]], align 4473// CHECK-NEXT: [[VLA7:%.*]] = alloca i32, i64 [[TMP15]], align 4, addrspace(5)474// CHECK-NEXT: [[VLA7_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[VLA7]] to ptr475// CHECK-NEXT: store i64 [[TMP15]], ptr [[__VLA_EXPR0_ASCAST]], align 8476// CHECK-NEXT: [[TMP17:%.*]] = load i32, ptr [[I4_ASCAST]], align 4477// CHECK-NEXT: [[TMP18:%.*]] = load i32, ptr [[I4_ASCAST]], align 4478// CHECK-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP18]] to i64479// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 [[IDXPROM]]480// CHECK-NEXT: store i32 [[TMP17]], ptr [[ARRAYIDX]], align 4481// CHECK-NEXT: store i32 0, ptr [[J_ASCAST]], align 4482// CHECK-NEXT: br label [[FOR_COND:%.*]]483// CHECK: for.cond:484// CHECK-NEXT: [[TMP19:%.*]] = load i32, ptr [[J_ASCAST]], align 4485// CHECK-NEXT: [[TMP20:%.*]] = load i32, ptr [[N_ASCAST]], align 4486// CHECK-NEXT: [[CMP8:%.*]] = icmp slt i32 [[TMP19]], [[TMP20]]487// CHECK-NEXT: br i1 [[CMP8]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]]488// CHECK: for.body:489// CHECK-NEXT: [[TMP21:%.*]] = load i32, ptr [[J_ASCAST]], align 4490// CHECK-NEXT: [[TMP22:%.*]] = load i32, ptr [[J_ASCAST]], align 4491// CHECK-NEXT: [[IDXPROM9:%.*]] = sext i32 [[TMP22]] to i64492// CHECK-NEXT: [[ARRAYIDX10:%.*]] = getelementptr inbounds i32, ptr [[VLA7_ASCAST]], i64 [[IDXPROM9]]493// CHECK-NEXT: store i32 [[TMP21]], ptr [[ARRAYIDX10]], align 4494// CHECK-NEXT: br label [[FOR_INC:%.*]]495// CHECK: for.inc:496// CHECK-NEXT: [[TMP23:%.*]] = load i32, ptr [[J_ASCAST]], align 4497// CHECK-NEXT: [[INC:%.*]] = add nsw i32 [[TMP23]], 1498// CHECK-NEXT: store i32 [[INC]], ptr [[J_ASCAST]], align 4499// CHECK-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP16:![0-9]+]]500// CHECK: for.end:501// CHECK-NEXT: store i32 0, ptr [[J11_ASCAST]], align 4502// CHECK-NEXT: br label [[FOR_COND12:%.*]]503// CHECK: for.cond12:504// CHECK-NEXT: [[TMP24:%.*]] = load i32, ptr [[J11_ASCAST]], align 4505// CHECK-NEXT: [[TMP25:%.*]] = load i32, ptr [[N_ASCAST]], align 4506// CHECK-NEXT: [[CMP13:%.*]] = icmp slt i32 [[TMP24]], [[TMP25]]507// CHECK-NEXT: br i1 [[CMP13]], label [[FOR_BODY14:%.*]], label [[FOR_END22:%.*]]508// CHECK: for.body14:509// CHECK-NEXT: [[TMP26:%.*]] = load i32, ptr [[J11_ASCAST]], align 4510// CHECK-NEXT: [[IDXPROM15:%.*]] = sext i32 [[TMP26]] to i64511// CHECK-NEXT: [[ARRAYIDX16:%.*]] = getelementptr inbounds i32, ptr [[VLA7_ASCAST]], i64 [[IDXPROM15]]512// CHECK-NEXT: [[TMP27:%.*]] = load i32, ptr [[ARRAYIDX16]], align 4513// CHECK-NEXT: [[TMP28:%.*]] = load i32, ptr [[I4_ASCAST]], align 4514// CHECK-NEXT: [[IDXPROM17:%.*]] = sext i32 [[TMP28]] to i64515// CHECK-NEXT: [[ARRAYIDX18:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 [[IDXPROM17]]516// CHECK-NEXT: [[TMP29:%.*]] = load i32, ptr [[ARRAYIDX18]], align 4517// CHECK-NEXT: [[ADD19:%.*]] = add nsw i32 [[TMP29]], [[TMP27]]518// CHECK-NEXT: store i32 [[ADD19]], ptr [[ARRAYIDX18]], align 4519// CHECK-NEXT: br label [[FOR_INC20:%.*]]520// CHECK: for.inc20:521// CHECK-NEXT: [[TMP30:%.*]] = load i32, ptr [[J11_ASCAST]], align 4522// CHECK-NEXT: [[INC21:%.*]] = add nsw i32 [[TMP30]], 1523// CHECK-NEXT: store i32 [[INC21]], ptr [[J11_ASCAST]], align 4524// CHECK-NEXT: br label [[FOR_COND12]], !llvm.loop [[LOOP17:![0-9]+]]525// CHECK: for.end22:526// CHECK-NEXT: [[TMP31:%.*]] = load ptr addrspace(5), ptr [[SAVED_STACK_ASCAST]], align 4527// CHECK-NEXT: call void @llvm.stackrestore.p5(ptr addrspace(5) [[TMP31]])528// CHECK-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]529// CHECK: omp.body.continue:530// CHECK-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]531// CHECK: omp.inner.for.inc:532// CHECK-NEXT: [[TMP32:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4533// CHECK-NEXT: [[TMP33:%.*]] = load i32, ptr [[DOTOMP_STRIDE_ASCAST]], align 4534// CHECK-NEXT: [[ADD23:%.*]] = add nsw i32 [[TMP32]], [[TMP33]]535// CHECK-NEXT: store i32 [[ADD23]], ptr [[DOTOMP_IV_ASCAST]], align 4536// CHECK-NEXT: br label [[OMP_INNER_FOR_COND]]537// CHECK: omp.inner.for.end:538// CHECK-NEXT: br label [[OMP_LOOP_EXIT:%.*]]539// CHECK: omp.loop.exit:540// CHECK-NEXT: [[TMP34:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8541// CHECK-NEXT: [[TMP35:%.*]] = load i32, ptr [[TMP34]], align 4542// CHECK-NEXT: call void @__kmpc_for_static_fini(ptr addrspacecast (ptr addrspace(1) @[[GLOB3]] to ptr), i32 [[TMP35]])543// CHECK-NEXT: br label [[OMP_PRECOND_END]]544// CHECK: omp.precond.end:545// CHECK-NEXT: ret void546//547//548// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo3v_l52549// CHECK-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[M:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[RESULT:%.*]]) #[[ATTR0]] {550// CHECK-NEXT: entry:551// CHECK-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)552// CHECK-NEXT: [[M_ADDR:%.*]] = alloca i64, align 8, addrspace(5)553// CHECK-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 8, addrspace(5)554// CHECK-NEXT: [[RESULT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)555// CHECK-NEXT: [[M_CASTED:%.*]] = alloca i64, align 8, addrspace(5)556// CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4, addrspace(5)557// CHECK-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4, addrspace(5)558// CHECK-NEXT: [[DYN_PTR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DYN_PTR_ADDR]] to ptr559// CHECK-NEXT: [[M_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[M_ADDR]] to ptr560// CHECK-NEXT: [[VLA_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[VLA_ADDR]] to ptr561// CHECK-NEXT: [[RESULT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RESULT_ADDR]] to ptr562// CHECK-NEXT: [[M_CASTED_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[M_CASTED]] to ptr563// CHECK-NEXT: [[DOTZERO_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTZERO_ADDR]] to ptr564// CHECK-NEXT: [[DOTTHREADID_TEMP__ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTTHREADID_TEMP_]] to ptr565// CHECK-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR_ASCAST]], align 8566// CHECK-NEXT: store i64 [[M]], ptr [[M_ADDR_ASCAST]], align 8567// CHECK-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR_ASCAST]], align 8568// CHECK-NEXT: store ptr [[RESULT]], ptr [[RESULT_ADDR_ASCAST]], align 8569// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR_ASCAST]], align 8570// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[RESULT_ADDR_ASCAST]], align 8571// CHECK-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo3v_l52_kernel_environment to ptr), ptr [[DYN_PTR]])572// CHECK-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -1573// CHECK-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]574// CHECK: user_code.entry:575// CHECK-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_global_thread_num(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr))576// CHECK-NEXT: [[TMP4:%.*]] = load i32, ptr [[M_ADDR_ASCAST]], align 4577// CHECK-NEXT: store i32 [[TMP4]], ptr [[M_CASTED_ASCAST]], align 4578// CHECK-NEXT: [[TMP5:%.*]] = load i64, ptr [[M_CASTED_ASCAST]], align 8579// CHECK-NEXT: store i32 0, ptr [[DOTZERO_ADDR_ASCAST]], align 4580// CHECK-NEXT: store i32 [[TMP3]], ptr [[DOTTHREADID_TEMP__ASCAST]], align 4581// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo3v_l52_omp_outlined(ptr [[DOTTHREADID_TEMP__ASCAST]], ptr [[DOTZERO_ADDR_ASCAST]], i64 [[TMP5]], i64 [[TMP0]], ptr [[TMP1]]) #[[ATTR4]]582// CHECK-NEXT: call void @__kmpc_target_deinit()583// CHECK-NEXT: ret void584// CHECK: worker.exit:585// CHECK-NEXT: ret void586//587//588// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo3v_l52_omp_outlined589// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[M:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[RESULT:%.*]]) #[[ATTR3]] {590// CHECK-NEXT: entry:591// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)592// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)593// CHECK-NEXT: [[M_ADDR:%.*]] = alloca i64, align 8, addrspace(5)594// CHECK-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 8, addrspace(5)595// CHECK-NEXT: [[RESULT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)596// CHECK-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4, addrspace(5)597// CHECK-NEXT: [[TMP:%.*]] = alloca i32, align 4, addrspace(5)598// CHECK-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4, addrspace(5)599// CHECK-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4, addrspace(5)600// CHECK-NEXT: [[I:%.*]] = alloca i32, align 4, addrspace(5)601// CHECK-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4, addrspace(5)602// CHECK-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4, addrspace(5)603// CHECK-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4, addrspace(5)604// CHECK-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4, addrspace(5)605// CHECK-NEXT: [[I3:%.*]] = alloca i32, align 4, addrspace(5)606// CHECK-NEXT: [[__VLA_EXPR0:%.*]] = alloca i64, align 8, addrspace(5)607// CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [3 x ptr], align 8, addrspace(5)608// CHECK-NEXT: [[J:%.*]] = alloca i32, align 4, addrspace(5)609// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTGLOBAL_TID__ADDR]] to ptr610// CHECK-NEXT: [[DOTBOUND_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTBOUND_TID__ADDR]] to ptr611// CHECK-NEXT: [[M_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[M_ADDR]] to ptr612// CHECK-NEXT: [[VLA_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[VLA_ADDR]] to ptr613// CHECK-NEXT: [[RESULT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RESULT_ADDR]] to ptr614// CHECK-NEXT: [[DOTOMP_IV_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IV]] to ptr615// CHECK-NEXT: [[TMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[TMP]] to ptr616// CHECK-NEXT: [[DOTCAPTURE_EXPR__ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTCAPTURE_EXPR_]] to ptr617// CHECK-NEXT: [[DOTCAPTURE_EXPR_1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTCAPTURE_EXPR_1]] to ptr618// CHECK-NEXT: [[I_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[I]] to ptr619// CHECK-NEXT: [[DOTOMP_LB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_LB]] to ptr620// CHECK-NEXT: [[DOTOMP_UB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_UB]] to ptr621// CHECK-NEXT: [[DOTOMP_STRIDE_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_STRIDE]] to ptr622// CHECK-NEXT: [[DOTOMP_IS_LAST_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IS_LAST]] to ptr623// CHECK-NEXT: [[I3_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[I3]] to ptr624// CHECK-NEXT: [[__VLA_EXPR0_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__VLA_EXPR0]] to ptr625// CHECK-NEXT: [[CAPTURED_VARS_ADDRS_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[CAPTURED_VARS_ADDRS]] to ptr626// CHECK-NEXT: [[J_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[J]] to ptr627// CHECK-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8628// CHECK-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR_ASCAST]], align 8629// CHECK-NEXT: store i64 [[M]], ptr [[M_ADDR_ASCAST]], align 8630// CHECK-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR_ASCAST]], align 8631// CHECK-NEXT: store ptr [[RESULT]], ptr [[RESULT_ADDR_ASCAST]], align 8632// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR_ASCAST]], align 8633// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[RESULT_ADDR_ASCAST]], align 8634// CHECK-NEXT: [[N:%.*]] = call align 8 ptr @__kmpc_alloc_shared(i64 4)635// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[M_ADDR_ASCAST]], align 4636// CHECK-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR__ASCAST]], align 4637// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ASCAST]], align 4638// CHECK-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0639// CHECK-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1640// CHECK-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1641// CHECK-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1_ASCAST]], align 4642// CHECK-NEXT: store i32 0, ptr [[I_ASCAST]], align 4643// CHECK-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ASCAST]], align 4644// CHECK-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]645// CHECK-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]646// CHECK: omp.precond.then:647// CHECK-NEXT: store i32 0, ptr [[DOTOMP_LB_ASCAST]], align 4648// CHECK-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1_ASCAST]], align 4649// CHECK-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_UB_ASCAST]], align 4650// CHECK-NEXT: store i32 1, ptr [[DOTOMP_STRIDE_ASCAST]], align 4651// CHECK-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST_ASCAST]], align 4652// CHECK-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8653// CHECK-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 4654// CHECK-NEXT: call void @__kmpc_distribute_static_init_4(ptr addrspacecast (ptr addrspace(1) @[[GLOB2]] to ptr), i32 [[TMP7]], i32 92, ptr [[DOTOMP_IS_LAST_ASCAST]], ptr [[DOTOMP_LB_ASCAST]], ptr [[DOTOMP_UB_ASCAST]], ptr [[DOTOMP_STRIDE_ASCAST]], i32 1, i32 1)655// CHECK-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 4656// CHECK-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1_ASCAST]], align 4657// CHECK-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]658// CHECK-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]659// CHECK: cond.true:660// CHECK-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1_ASCAST]], align 4661// CHECK-NEXT: br label [[COND_END:%.*]]662// CHECK: cond.false:663// CHECK-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 4664// CHECK-NEXT: br label [[COND_END]]665// CHECK: cond.end:666// CHECK-NEXT: [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]667// CHECK-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB_ASCAST]], align 4668// CHECK-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_LB_ASCAST]], align 4669// CHECK-NEXT: store i32 [[TMP12]], ptr [[DOTOMP_IV_ASCAST]], align 4670// CHECK-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]671// CHECK: omp.inner.for.cond:672// CHECK-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4673// CHECK-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 4674// CHECK-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]675// CHECK-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]676// CHECK: omp.inner.for.body:677// CHECK-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4678// CHECK-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP15]], 1679// CHECK-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]680// CHECK-NEXT: store i32 [[ADD]], ptr [[I3_ASCAST]], align 4681// CHECK-NEXT: store i32 10, ptr [[N]], align 4682// CHECK-NEXT: [[TMP16:%.*]] = load i32, ptr [[N]], align 4683// CHECK-NEXT: [[TMP17:%.*]] = zext i32 [[TMP16]] to i64684// CHECK-NEXT: [[TMP18:%.*]] = mul nuw i64 [[TMP17]], 4685// CHECK-NEXT: [[TMP19:%.*]] = add nuw i64 [[TMP18]], 3686// CHECK-NEXT: [[TMP20:%.*]] = udiv i64 [[TMP19]], 4687// CHECK-NEXT: [[TMP21:%.*]] = mul nuw i64 [[TMP20]], 4688// CHECK-NEXT: [[A:%.*]] = call align 4 ptr @__kmpc_alloc_shared(i64 [[TMP21]])689// CHECK-NEXT: store i64 [[TMP17]], ptr [[__VLA_EXPR0_ASCAST]], align 8690// CHECK-NEXT: [[TMP22:%.*]] = load i32, ptr [[I3_ASCAST]], align 4691// CHECK-NEXT: [[TMP23:%.*]] = load i32, ptr [[I3_ASCAST]], align 4692// CHECK-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP23]] to i64693// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 [[IDXPROM]]694// CHECK-NEXT: store i32 [[TMP22]], ptr [[ARRAYIDX]], align 4695// CHECK-NEXT: [[TMP24:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 0, i64 0696// CHECK-NEXT: store ptr [[N]], ptr [[TMP24]], align 8697// CHECK-NEXT: [[TMP25:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 0, i64 1698// CHECK-NEXT: [[TMP26:%.*]] = inttoptr i64 [[TMP17]] to ptr699// CHECK-NEXT: store ptr [[TMP26]], ptr [[TMP25]], align 8700// CHECK-NEXT: [[TMP27:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 0, i64 2701// CHECK-NEXT: store ptr [[A]], ptr [[TMP27]], align 8702// CHECK-NEXT: [[TMP28:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8703// CHECK-NEXT: [[TMP29:%.*]] = load i32, ptr [[TMP28]], align 4704// CHECK-NEXT: call void @__kmpc_parallel_51(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i32 [[TMP29]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo3v_l52_omp_outlined_omp_outlined, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo3v_l52_omp_outlined_omp_outlined_wrapper, ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 3)705// CHECK-NEXT: store i32 0, ptr [[J_ASCAST]], align 4706// CHECK-NEXT: br label [[FOR_COND:%.*]]707// CHECK: for.cond:708// CHECK-NEXT: [[TMP30:%.*]] = load i32, ptr [[J_ASCAST]], align 4709// CHECK-NEXT: [[TMP31:%.*]] = load i32, ptr [[N]], align 4710// CHECK-NEXT: [[CMP6:%.*]] = icmp slt i32 [[TMP30]], [[TMP31]]711// CHECK-NEXT: br i1 [[CMP6]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]]712// CHECK: for.body:713// CHECK-NEXT: [[TMP32:%.*]] = load i32, ptr [[J_ASCAST]], align 4714// CHECK-NEXT: [[IDXPROM7:%.*]] = sext i32 [[TMP32]] to i64715// CHECK-NEXT: [[ARRAYIDX8:%.*]] = getelementptr inbounds i32, ptr [[A]], i64 [[IDXPROM7]]716// CHECK-NEXT: [[TMP33:%.*]] = load i32, ptr [[ARRAYIDX8]], align 4717// CHECK-NEXT: [[TMP34:%.*]] = load i32, ptr [[I3_ASCAST]], align 4718// CHECK-NEXT: [[IDXPROM9:%.*]] = sext i32 [[TMP34]] to i64719// CHECK-NEXT: [[ARRAYIDX10:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 [[IDXPROM9]]720// CHECK-NEXT: [[TMP35:%.*]] = load i32, ptr [[ARRAYIDX10]], align 4721// CHECK-NEXT: [[ADD11:%.*]] = add nsw i32 [[TMP35]], [[TMP33]]722// CHECK-NEXT: store i32 [[ADD11]], ptr [[ARRAYIDX10]], align 4723// CHECK-NEXT: br label [[FOR_INC:%.*]]724// CHECK: for.inc:725// CHECK-NEXT: [[TMP36:%.*]] = load i32, ptr [[J_ASCAST]], align 4726// CHECK-NEXT: [[INC:%.*]] = add nsw i32 [[TMP36]], 1727// CHECK-NEXT: store i32 [[INC]], ptr [[J_ASCAST]], align 4728// CHECK-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP18:![0-9]+]]729// CHECK: for.end:730// CHECK-NEXT: call void @__kmpc_free_shared(ptr [[A]], i64 [[TMP21]])731// CHECK-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]732// CHECK: omp.body.continue:733// CHECK-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]734// CHECK: omp.inner.for.inc:735// CHECK-NEXT: [[TMP37:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4736// CHECK-NEXT: [[ADD12:%.*]] = add nsw i32 [[TMP37]], 1737// CHECK-NEXT: store i32 [[ADD12]], ptr [[DOTOMP_IV_ASCAST]], align 4738// CHECK-NEXT: br label [[OMP_INNER_FOR_COND]]739// CHECK: omp.inner.for.end:740// CHECK-NEXT: br label [[OMP_LOOP_EXIT:%.*]]741// CHECK: omp.loop.exit:742// CHECK-NEXT: [[TMP38:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8743// CHECK-NEXT: [[TMP39:%.*]] = load i32, ptr [[TMP38]], align 4744// CHECK-NEXT: call void @__kmpc_distribute_static_fini(ptr addrspacecast (ptr addrspace(1) @[[GLOB2]] to ptr), i32 [[TMP39]])745// CHECK-NEXT: br label [[OMP_PRECOND_END]]746// CHECK: omp.precond.end:747// CHECK-NEXT: call void @__kmpc_free_shared(ptr [[N]], i64 4)748// CHECK-NEXT: ret void749//750//751// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo3v_l52_omp_outlined_omp_outlined752// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR3]] {753// CHECK-NEXT: entry:754// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)755// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)756// CHECK-NEXT: [[N_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)757// CHECK-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 8, addrspace(5)758// CHECK-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)759// CHECK-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4, addrspace(5)760// CHECK-NEXT: [[TMP:%.*]] = alloca i32, align 4, addrspace(5)761// CHECK-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4, addrspace(5)762// CHECK-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4, addrspace(5)763// CHECK-NEXT: [[J:%.*]] = alloca i32, align 4, addrspace(5)764// CHECK-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4, addrspace(5)765// CHECK-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4, addrspace(5)766// CHECK-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4, addrspace(5)767// CHECK-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4, addrspace(5)768// CHECK-NEXT: [[J3:%.*]] = alloca i32, align 4, addrspace(5)769// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTGLOBAL_TID__ADDR]] to ptr770// CHECK-NEXT: [[DOTBOUND_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTBOUND_TID__ADDR]] to ptr771// CHECK-NEXT: [[N_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[N_ADDR]] to ptr772// CHECK-NEXT: [[VLA_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[VLA_ADDR]] to ptr773// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr774// CHECK-NEXT: [[DOTOMP_IV_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IV]] to ptr775// CHECK-NEXT: [[TMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[TMP]] to ptr776// CHECK-NEXT: [[DOTCAPTURE_EXPR__ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTCAPTURE_EXPR_]] to ptr777// CHECK-NEXT: [[DOTCAPTURE_EXPR_1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTCAPTURE_EXPR_1]] to ptr778// CHECK-NEXT: [[J_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[J]] to ptr779// CHECK-NEXT: [[DOTOMP_LB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_LB]] to ptr780// CHECK-NEXT: [[DOTOMP_UB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_UB]] to ptr781// CHECK-NEXT: [[DOTOMP_STRIDE_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_STRIDE]] to ptr782// CHECK-NEXT: [[DOTOMP_IS_LAST_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IS_LAST]] to ptr783// CHECK-NEXT: [[J3_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[J3]] to ptr784// CHECK-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8785// CHECK-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR_ASCAST]], align 8786// CHECK-NEXT: store ptr [[N]], ptr [[N_ADDR_ASCAST]], align 8787// CHECK-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR_ASCAST]], align 8788// CHECK-NEXT: store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8789// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[N_ADDR_ASCAST]], align 8790// CHECK-NEXT: [[TMP1:%.*]] = load i64, ptr [[VLA_ADDR_ASCAST]], align 8791// CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8792// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 4793// CHECK-NEXT: store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR__ASCAST]], align 4794// CHECK-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ASCAST]], align 4795// CHECK-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP4]], 0796// CHECK-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1797// CHECK-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1798// CHECK-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1_ASCAST]], align 4799// CHECK-NEXT: store i32 0, ptr [[J_ASCAST]], align 4800// CHECK-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ASCAST]], align 4801// CHECK-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP5]]802// CHECK-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]803// CHECK: omp.precond.then:804// CHECK-NEXT: store i32 0, ptr [[DOTOMP_LB_ASCAST]], align 4805// CHECK-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1_ASCAST]], align 4806// CHECK-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_UB_ASCAST]], align 4807// CHECK-NEXT: store i32 1, ptr [[DOTOMP_STRIDE_ASCAST]], align 4808// CHECK-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST_ASCAST]], align 4809// CHECK-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8810// CHECK-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4811// CHECK-NEXT: call void @__kmpc_for_static_init_4(ptr addrspacecast (ptr addrspace(1) @[[GLOB3]] to ptr), i32 [[TMP8]], i32 33, ptr [[DOTOMP_IS_LAST_ASCAST]], ptr [[DOTOMP_LB_ASCAST]], ptr [[DOTOMP_UB_ASCAST]], ptr [[DOTOMP_STRIDE_ASCAST]], i32 1, i32 1)812// CHECK-NEXT: br label [[OMP_DISPATCH_COND:%.*]]813// CHECK: omp.dispatch.cond:814// CHECK-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 4815// CHECK-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1_ASCAST]], align 4816// CHECK-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]817// CHECK-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]818// CHECK: cond.true:819// CHECK-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1_ASCAST]], align 4820// CHECK-NEXT: br label [[COND_END:%.*]]821// CHECK: cond.false:822// CHECK-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 4823// CHECK-NEXT: br label [[COND_END]]824// CHECK: cond.end:825// CHECK-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]826// CHECK-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB_ASCAST]], align 4827// CHECK-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB_ASCAST]], align 4828// CHECK-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV_ASCAST]], align 4829// CHECK-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4830// CHECK-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 4831// CHECK-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]832// CHECK-NEXT: br i1 [[CMP5]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]833// CHECK: omp.dispatch.body:834// CHECK-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]835// CHECK: omp.inner.for.cond:836// CHECK-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4837// CHECK-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 4838// CHECK-NEXT: [[CMP6:%.*]] = icmp sle i32 [[TMP16]], [[TMP17]]839// CHECK-NEXT: br i1 [[CMP6]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]840// CHECK: omp.inner.for.body:841// CHECK-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4842// CHECK-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP18]], 1843// CHECK-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]844// CHECK-NEXT: store i32 [[ADD]], ptr [[J3_ASCAST]], align 4845// CHECK-NEXT: [[TMP19:%.*]] = load i32, ptr [[J3_ASCAST]], align 4846// CHECK-NEXT: [[TMP20:%.*]] = load i32, ptr [[J3_ASCAST]], align 4847// CHECK-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP20]] to i64848// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i64 [[IDXPROM]]849// CHECK-NEXT: store i32 [[TMP19]], ptr [[ARRAYIDX]], align 4850// CHECK-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]851// CHECK: omp.body.continue:852// CHECK-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]853// CHECK: omp.inner.for.inc:854// CHECK-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4855// CHECK-NEXT: [[ADD7:%.*]] = add nsw i32 [[TMP21]], 1856// CHECK-NEXT: store i32 [[ADD7]], ptr [[DOTOMP_IV_ASCAST]], align 4857// CHECK-NEXT: br label [[OMP_INNER_FOR_COND]]858// CHECK: omp.inner.for.end:859// CHECK-NEXT: br label [[OMP_DISPATCH_INC:%.*]]860// CHECK: omp.dispatch.inc:861// CHECK-NEXT: [[TMP22:%.*]] = load i32, ptr [[DOTOMP_LB_ASCAST]], align 4862// CHECK-NEXT: [[TMP23:%.*]] = load i32, ptr [[DOTOMP_STRIDE_ASCAST]], align 4863// CHECK-NEXT: [[ADD8:%.*]] = add nsw i32 [[TMP22]], [[TMP23]]864// CHECK-NEXT: store i32 [[ADD8]], ptr [[DOTOMP_LB_ASCAST]], align 4865// CHECK-NEXT: [[TMP24:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 4866// CHECK-NEXT: [[TMP25:%.*]] = load i32, ptr [[DOTOMP_STRIDE_ASCAST]], align 4867// CHECK-NEXT: [[ADD9:%.*]] = add nsw i32 [[TMP24]], [[TMP25]]868// CHECK-NEXT: store i32 [[ADD9]], ptr [[DOTOMP_UB_ASCAST]], align 4869// CHECK-NEXT: br label [[OMP_DISPATCH_COND]]870// CHECK: omp.dispatch.end:871// CHECK-NEXT: [[TMP26:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8872// CHECK-NEXT: [[TMP27:%.*]] = load i32, ptr [[TMP26]], align 4873// CHECK-NEXT: call void @__kmpc_for_static_fini(ptr addrspacecast (ptr addrspace(1) @[[GLOB3]] to ptr), i32 [[TMP27]])874// CHECK-NEXT: br label [[OMP_PRECOND_END]]875// CHECK: omp.precond.end:876// CHECK-NEXT: ret void877//878//879// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo3v_l52_omp_outlined_omp_outlined_wrapper880// CHECK-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR7:[0-9]+]] {881// CHECK-NEXT: entry:882// CHECK-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2, addrspace(5)883// CHECK-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4, addrspace(5)884// CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4, addrspace(5)885// CHECK-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 8, addrspace(5)886// CHECK-NEXT: [[DOTADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTADDR]] to ptr887// CHECK-NEXT: [[DOTADDR1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTADDR1]] to ptr888// CHECK-NEXT: [[DOTZERO_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTZERO_ADDR]] to ptr889// CHECK-NEXT: [[GLOBAL_ARGS_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[GLOBAL_ARGS]] to ptr890// CHECK-NEXT: store i16 [[TMP0]], ptr [[DOTADDR_ASCAST]], align 2891// CHECK-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1_ASCAST]], align 4892// CHECK-NEXT: store i32 0, ptr [[DOTZERO_ADDR_ASCAST]], align 4893// CHECK-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS_ASCAST]])894// CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[GLOBAL_ARGS_ASCAST]], align 8895// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 0896// CHECK-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP3]], align 8897// CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 1898// CHECK-NEXT: [[TMP6:%.*]] = load i64, ptr [[TMP5]], align 8899// CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 2900// CHECK-NEXT: [[TMP8:%.*]] = load ptr, ptr [[TMP7]], align 8901// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo3v_l52_omp_outlined_omp_outlined(ptr [[DOTADDR1_ASCAST]], ptr [[DOTZERO_ADDR_ASCAST]], ptr [[TMP4]], i64 [[TMP6]], ptr [[TMP8]]) #[[ATTR4]]902// CHECK-NEXT: ret void903//904//905// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l76906// CHECK-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[M:%.*]], i64 noundef [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[RESULT:%.*]]) #[[ATTR0]] {907// CHECK-NEXT: entry:908// CHECK-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)909// CHECK-NEXT: [[M_ADDR:%.*]] = alloca i64, align 8, addrspace(5)910// CHECK-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8, addrspace(5)911// CHECK-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 8, addrspace(5)912// CHECK-NEXT: [[RESULT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)913// CHECK-NEXT: [[M_CASTED:%.*]] = alloca i64, align 8, addrspace(5)914// CHECK-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8, addrspace(5)915// CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4, addrspace(5)916// CHECK-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4, addrspace(5)917// CHECK-NEXT: [[DYN_PTR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DYN_PTR_ADDR]] to ptr918// CHECK-NEXT: [[M_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[M_ADDR]] to ptr919// CHECK-NEXT: [[N_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[N_ADDR]] to ptr920// CHECK-NEXT: [[VLA_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[VLA_ADDR]] to ptr921// CHECK-NEXT: [[RESULT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RESULT_ADDR]] to ptr922// CHECK-NEXT: [[M_CASTED_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[M_CASTED]] to ptr923// CHECK-NEXT: [[N_CASTED_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[N_CASTED]] to ptr924// CHECK-NEXT: [[DOTZERO_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTZERO_ADDR]] to ptr925// CHECK-NEXT: [[DOTTHREADID_TEMP__ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTTHREADID_TEMP_]] to ptr926// CHECK-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR_ASCAST]], align 8927// CHECK-NEXT: store i64 [[M]], ptr [[M_ADDR_ASCAST]], align 8928// CHECK-NEXT: store i64 [[N]], ptr [[N_ADDR_ASCAST]], align 8929// CHECK-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR_ASCAST]], align 8930// CHECK-NEXT: store ptr [[RESULT]], ptr [[RESULT_ADDR_ASCAST]], align 8931// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR_ASCAST]], align 8932// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[RESULT_ADDR_ASCAST]], align 8933// CHECK-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l76_kernel_environment to ptr), ptr [[DYN_PTR]])934// CHECK-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -1935// CHECK-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]936// CHECK: user_code.entry:937// CHECK-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_global_thread_num(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr))938// CHECK-NEXT: [[TMP4:%.*]] = load i32, ptr [[M_ADDR_ASCAST]], align 4939// CHECK-NEXT: store i32 [[TMP4]], ptr [[M_CASTED_ASCAST]], align 4940// CHECK-NEXT: [[TMP5:%.*]] = load i64, ptr [[M_CASTED_ASCAST]], align 8941// CHECK-NEXT: [[TMP6:%.*]] = load i32, ptr [[N_ADDR_ASCAST]], align 4942// CHECK-NEXT: store i32 [[TMP6]], ptr [[N_CASTED_ASCAST]], align 4943// CHECK-NEXT: [[TMP7:%.*]] = load i64, ptr [[N_CASTED_ASCAST]], align 8944// CHECK-NEXT: store i32 0, ptr [[DOTZERO_ADDR_ASCAST]], align 4945// CHECK-NEXT: store i32 [[TMP3]], ptr [[DOTTHREADID_TEMP__ASCAST]], align 4946// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l76_omp_outlined(ptr [[DOTTHREADID_TEMP__ASCAST]], ptr [[DOTZERO_ADDR_ASCAST]], i64 [[TMP5]], i64 [[TMP7]], i64 [[TMP0]], ptr [[TMP1]]) #[[ATTR4]]947// CHECK-NEXT: call void @__kmpc_target_deinit()948// CHECK-NEXT: ret void949// CHECK: worker.exit:950// CHECK-NEXT: ret void951//952//953// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l76_omp_outlined954// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[M:%.*]], i64 noundef [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[RESULT:%.*]]) #[[ATTR3]] {955// CHECK-NEXT: entry:956// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)957// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)958// CHECK-NEXT: [[M_ADDR:%.*]] = alloca i64, align 8, addrspace(5)959// CHECK-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8, addrspace(5)960// CHECK-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 8, addrspace(5)961// CHECK-NEXT: [[RESULT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)962// CHECK-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4, addrspace(5)963// CHECK-NEXT: [[TMP:%.*]] = alloca i32, align 4, addrspace(5)964// CHECK-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4, addrspace(5)965// CHECK-NEXT: [[DOTCAPTURE_EXPR_2:%.*]] = alloca i32, align 4, addrspace(5)966// CHECK-NEXT: [[I:%.*]] = alloca i32, align 4, addrspace(5)967// CHECK-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4, addrspace(5)968// CHECK-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4, addrspace(5)969// CHECK-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4, addrspace(5)970// CHECK-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4, addrspace(5)971// CHECK-NEXT: [[I4:%.*]] = alloca i32, align 4, addrspace(5)972// CHECK-NEXT: [[__VLA_EXPR0:%.*]] = alloca i64, align 8, addrspace(5)973// CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [3 x ptr], align 8, addrspace(5)974// CHECK-NEXT: [[J:%.*]] = alloca i32, align 4, addrspace(5)975// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTGLOBAL_TID__ADDR]] to ptr976// CHECK-NEXT: [[DOTBOUND_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTBOUND_TID__ADDR]] to ptr977// CHECK-NEXT: [[M_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[M_ADDR]] to ptr978// CHECK-NEXT: [[N_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[N_ADDR]] to ptr979// CHECK-NEXT: [[VLA_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[VLA_ADDR]] to ptr980// CHECK-NEXT: [[RESULT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RESULT_ADDR]] to ptr981// CHECK-NEXT: [[DOTOMP_IV_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IV]] to ptr982// CHECK-NEXT: [[TMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[TMP]] to ptr983// CHECK-NEXT: [[DOTCAPTURE_EXPR__ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTCAPTURE_EXPR_]] to ptr984// CHECK-NEXT: [[DOTCAPTURE_EXPR_2_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTCAPTURE_EXPR_2]] to ptr985// CHECK-NEXT: [[I_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[I]] to ptr986// CHECK-NEXT: [[DOTOMP_LB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_LB]] to ptr987// CHECK-NEXT: [[DOTOMP_UB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_UB]] to ptr988// CHECK-NEXT: [[DOTOMP_STRIDE_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_STRIDE]] to ptr989// CHECK-NEXT: [[DOTOMP_IS_LAST_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IS_LAST]] to ptr990// CHECK-NEXT: [[I4_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[I4]] to ptr991// CHECK-NEXT: [[__VLA_EXPR0_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__VLA_EXPR0]] to ptr992// CHECK-NEXT: [[CAPTURED_VARS_ADDRS_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[CAPTURED_VARS_ADDRS]] to ptr993// CHECK-NEXT: [[J_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[J]] to ptr994// CHECK-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8995// CHECK-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR_ASCAST]], align 8996// CHECK-NEXT: store i64 [[M]], ptr [[M_ADDR_ASCAST]], align 8997// CHECK-NEXT: store i64 [[N]], ptr [[N_ADDR_ASCAST]], align 8998// CHECK-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR_ASCAST]], align 8999// CHECK-NEXT: store ptr [[RESULT]], ptr [[RESULT_ADDR_ASCAST]], align 81000// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR_ASCAST]], align 81001// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[RESULT_ADDR_ASCAST]], align 81002// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR_ASCAST]], align 41003// CHECK-NEXT: [[N1:%.*]] = call align 8 ptr @__kmpc_alloc_shared(i64 4)1004// CHECK-NEXT: store i32 [[TMP2]], ptr [[N1]], align 41005// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[M_ADDR_ASCAST]], align 41006// CHECK-NEXT: store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR__ASCAST]], align 41007// CHECK-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ASCAST]], align 41008// CHECK-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP4]], 01009// CHECK-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11010// CHECK-NEXT: [[SUB3:%.*]] = sub nsw i32 [[DIV]], 11011// CHECK-NEXT: store i32 [[SUB3]], ptr [[DOTCAPTURE_EXPR_2_ASCAST]], align 41012// CHECK-NEXT: store i32 0, ptr [[I_ASCAST]], align 41013// CHECK-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ASCAST]], align 41014// CHECK-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP5]]1015// CHECK-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1016// CHECK: omp.precond.then:1017// CHECK-NEXT: store i32 0, ptr [[DOTOMP_LB_ASCAST]], align 41018// CHECK-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2_ASCAST]], align 41019// CHECK-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_UB_ASCAST]], align 41020// CHECK-NEXT: store i32 1, ptr [[DOTOMP_STRIDE_ASCAST]], align 41021// CHECK-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST_ASCAST]], align 41022// CHECK-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 81023// CHECK-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41024// CHECK-NEXT: call void @__kmpc_distribute_static_init_4(ptr addrspacecast (ptr addrspace(1) @[[GLOB2]] to ptr), i32 [[TMP8]], i32 92, ptr [[DOTOMP_IS_LAST_ASCAST]], ptr [[DOTOMP_LB_ASCAST]], ptr [[DOTOMP_UB_ASCAST]], ptr [[DOTOMP_STRIDE_ASCAST]], i32 1, i32 1)1025// CHECK-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 41026// CHECK-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2_ASCAST]], align 41027// CHECK-NEXT: [[CMP5:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]1028// CHECK-NEXT: br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1029// CHECK: cond.true:1030// CHECK-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2_ASCAST]], align 41031// CHECK-NEXT: br label [[COND_END:%.*]]1032// CHECK: cond.false:1033// CHECK-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 41034// CHECK-NEXT: br label [[COND_END]]1035// CHECK: cond.end:1036// CHECK-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]1037// CHECK-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB_ASCAST]], align 41038// CHECK-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB_ASCAST]], align 41039// CHECK-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV_ASCAST]], align 41040// CHECK-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1041// CHECK: omp.inner.for.cond:1042// CHECK-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 41043// CHECK-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 41044// CHECK-NEXT: [[CMP6:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]1045// CHECK-NEXT: br i1 [[CMP6]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1046// CHECK: omp.inner.for.body:1047// CHECK-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 41048// CHECK-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 11049// CHECK-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1050// CHECK-NEXT: store i32 [[ADD]], ptr [[I4_ASCAST]], align 41051// CHECK-NEXT: [[TMP17:%.*]] = load i32, ptr [[N1]], align 41052// CHECK-NEXT: [[TMP18:%.*]] = zext i32 [[TMP17]] to i641053// CHECK-NEXT: [[TMP19:%.*]] = mul nuw i64 [[TMP18]], 41054// CHECK-NEXT: [[TMP20:%.*]] = add nuw i64 [[TMP19]], 31055// CHECK-NEXT: [[TMP21:%.*]] = udiv i64 [[TMP20]], 41056// CHECK-NEXT: [[TMP22:%.*]] = mul nuw i64 [[TMP21]], 41057// CHECK-NEXT: [[A:%.*]] = call align 4 ptr @__kmpc_alloc_shared(i64 [[TMP22]])1058// CHECK-NEXT: store i64 [[TMP18]], ptr [[__VLA_EXPR0_ASCAST]], align 81059// CHECK-NEXT: [[TMP23:%.*]] = load i32, ptr [[I4_ASCAST]], align 41060// CHECK-NEXT: [[TMP24:%.*]] = load i32, ptr [[I4_ASCAST]], align 41061// CHECK-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP24]] to i641062// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 [[IDXPROM]]1063// CHECK-NEXT: store i32 [[TMP23]], ptr [[ARRAYIDX]], align 41064// CHECK-NEXT: [[TMP25:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 0, i64 01065// CHECK-NEXT: store ptr [[N1]], ptr [[TMP25]], align 81066// CHECK-NEXT: [[TMP26:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 0, i64 11067// CHECK-NEXT: [[TMP27:%.*]] = inttoptr i64 [[TMP18]] to ptr1068// CHECK-NEXT: store ptr [[TMP27]], ptr [[TMP26]], align 81069// CHECK-NEXT: [[TMP28:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 0, i64 21070// CHECK-NEXT: store ptr [[A]], ptr [[TMP28]], align 81071// CHECK-NEXT: [[TMP29:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 81072// CHECK-NEXT: [[TMP30:%.*]] = load i32, ptr [[TMP29]], align 41073// CHECK-NEXT: call void @__kmpc_parallel_51(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i32 [[TMP30]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l76_omp_outlined_omp_outlined, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l76_omp_outlined_omp_outlined_wrapper, ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 3)1074// CHECK-NEXT: store i32 0, ptr [[J_ASCAST]], align 41075// CHECK-NEXT: br label [[FOR_COND:%.*]]1076// CHECK: for.cond:1077// CHECK-NEXT: [[TMP31:%.*]] = load i32, ptr [[J_ASCAST]], align 41078// CHECK-NEXT: [[TMP32:%.*]] = load i32, ptr [[N1]], align 41079// CHECK-NEXT: [[CMP7:%.*]] = icmp slt i32 [[TMP31]], [[TMP32]]1080// CHECK-NEXT: br i1 [[CMP7]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]]1081// CHECK: for.body:1082// CHECK-NEXT: [[TMP33:%.*]] = load i32, ptr [[J_ASCAST]], align 41083// CHECK-NEXT: [[IDXPROM8:%.*]] = sext i32 [[TMP33]] to i641084// CHECK-NEXT: [[ARRAYIDX9:%.*]] = getelementptr inbounds i32, ptr [[A]], i64 [[IDXPROM8]]1085// CHECK-NEXT: [[TMP34:%.*]] = load i32, ptr [[ARRAYIDX9]], align 41086// CHECK-NEXT: [[TMP35:%.*]] = load i32, ptr [[I4_ASCAST]], align 41087// CHECK-NEXT: [[IDXPROM10:%.*]] = sext i32 [[TMP35]] to i641088// CHECK-NEXT: [[ARRAYIDX11:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 [[IDXPROM10]]1089// CHECK-NEXT: [[TMP36:%.*]] = load i32, ptr [[ARRAYIDX11]], align 41090// CHECK-NEXT: [[ADD12:%.*]] = add nsw i32 [[TMP36]], [[TMP34]]1091// CHECK-NEXT: store i32 [[ADD12]], ptr [[ARRAYIDX11]], align 41092// CHECK-NEXT: br label [[FOR_INC:%.*]]1093// CHECK: for.inc:1094// CHECK-NEXT: [[TMP37:%.*]] = load i32, ptr [[J_ASCAST]], align 41095// CHECK-NEXT: [[INC:%.*]] = add nsw i32 [[TMP37]], 11096// CHECK-NEXT: store i32 [[INC]], ptr [[J_ASCAST]], align 41097// CHECK-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP19:![0-9]+]]1098// CHECK: for.end:1099// CHECK-NEXT: call void @__kmpc_free_shared(ptr [[A]], i64 [[TMP22]])1100// CHECK-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1101// CHECK: omp.body.continue:1102// CHECK-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1103// CHECK: omp.inner.for.inc:1104// CHECK-NEXT: [[TMP38:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 41105// CHECK-NEXT: [[ADD13:%.*]] = add nsw i32 [[TMP38]], 11106// CHECK-NEXT: store i32 [[ADD13]], ptr [[DOTOMP_IV_ASCAST]], align 41107// CHECK-NEXT: br label [[OMP_INNER_FOR_COND]]1108// CHECK: omp.inner.for.end:1109// CHECK-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1110// CHECK: omp.loop.exit:1111// CHECK-NEXT: [[TMP39:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 81112// CHECK-NEXT: [[TMP40:%.*]] = load i32, ptr [[TMP39]], align 41113// CHECK-NEXT: call void @__kmpc_distribute_static_fini(ptr addrspacecast (ptr addrspace(1) @[[GLOB2]] to ptr), i32 [[TMP40]])1114// CHECK-NEXT: br label [[OMP_PRECOND_END]]1115// CHECK: omp.precond.end:1116// CHECK-NEXT: call void @__kmpc_free_shared(ptr [[N1]], i64 4)1117// CHECK-NEXT: ret void1118//1119//1120// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l76_omp_outlined_omp_outlined1121// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR3]] {1122// CHECK-NEXT: entry:1123// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)1124// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)1125// CHECK-NEXT: [[N_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)1126// CHECK-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 8, addrspace(5)1127// CHECK-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)1128// CHECK-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4, addrspace(5)1129// CHECK-NEXT: [[TMP:%.*]] = alloca i32, align 4, addrspace(5)1130// CHECK-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4, addrspace(5)1131// CHECK-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4, addrspace(5)1132// CHECK-NEXT: [[J:%.*]] = alloca i32, align 4, addrspace(5)1133// CHECK-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4, addrspace(5)1134// CHECK-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4, addrspace(5)1135// CHECK-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4, addrspace(5)1136// CHECK-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4, addrspace(5)1137// CHECK-NEXT: [[J3:%.*]] = alloca i32, align 4, addrspace(5)1138// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTGLOBAL_TID__ADDR]] to ptr1139// CHECK-NEXT: [[DOTBOUND_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTBOUND_TID__ADDR]] to ptr1140// CHECK-NEXT: [[N_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[N_ADDR]] to ptr1141// CHECK-NEXT: [[VLA_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[VLA_ADDR]] to ptr1142// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr1143// CHECK-NEXT: [[DOTOMP_IV_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IV]] to ptr1144// CHECK-NEXT: [[TMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[TMP]] to ptr1145// CHECK-NEXT: [[DOTCAPTURE_EXPR__ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTCAPTURE_EXPR_]] to ptr1146// CHECK-NEXT: [[DOTCAPTURE_EXPR_1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTCAPTURE_EXPR_1]] to ptr1147// CHECK-NEXT: [[J_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[J]] to ptr1148// CHECK-NEXT: [[DOTOMP_LB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_LB]] to ptr1149// CHECK-NEXT: [[DOTOMP_UB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_UB]] to ptr1150// CHECK-NEXT: [[DOTOMP_STRIDE_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_STRIDE]] to ptr1151// CHECK-NEXT: [[DOTOMP_IS_LAST_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IS_LAST]] to ptr1152// CHECK-NEXT: [[J3_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[J3]] to ptr1153// CHECK-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 81154// CHECK-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR_ASCAST]], align 81155// CHECK-NEXT: store ptr [[N]], ptr [[N_ADDR_ASCAST]], align 81156// CHECK-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR_ASCAST]], align 81157// CHECK-NEXT: store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 81158// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[N_ADDR_ASCAST]], align 81159// CHECK-NEXT: [[TMP1:%.*]] = load i64, ptr [[VLA_ADDR_ASCAST]], align 81160// CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 81161// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 41162// CHECK-NEXT: store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR__ASCAST]], align 41163// CHECK-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ASCAST]], align 41164// CHECK-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP4]], 01165// CHECK-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11166// CHECK-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11167// CHECK-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1_ASCAST]], align 41168// CHECK-NEXT: store i32 0, ptr [[J_ASCAST]], align 41169// CHECK-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ASCAST]], align 41170// CHECK-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP5]]1171// CHECK-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1172// CHECK: omp.precond.then:1173// CHECK-NEXT: store i32 0, ptr [[DOTOMP_LB_ASCAST]], align 41174// CHECK-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1_ASCAST]], align 41175// CHECK-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_UB_ASCAST]], align 41176// CHECK-NEXT: store i32 1, ptr [[DOTOMP_STRIDE_ASCAST]], align 41177// CHECK-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST_ASCAST]], align 41178// CHECK-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 81179// CHECK-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41180// CHECK-NEXT: call void @__kmpc_for_static_init_4(ptr addrspacecast (ptr addrspace(1) @[[GLOB3]] to ptr), i32 [[TMP8]], i32 33, ptr [[DOTOMP_IS_LAST_ASCAST]], ptr [[DOTOMP_LB_ASCAST]], ptr [[DOTOMP_UB_ASCAST]], ptr [[DOTOMP_STRIDE_ASCAST]], i32 1, i32 1)1181// CHECK-NEXT: br label [[OMP_DISPATCH_COND:%.*]]1182// CHECK: omp.dispatch.cond:1183// CHECK-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 41184// CHECK-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1_ASCAST]], align 41185// CHECK-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]1186// CHECK-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1187// CHECK: cond.true:1188// CHECK-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1_ASCAST]], align 41189// CHECK-NEXT: br label [[COND_END:%.*]]1190// CHECK: cond.false:1191// CHECK-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 41192// CHECK-NEXT: br label [[COND_END]]1193// CHECK: cond.end:1194// CHECK-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]1195// CHECK-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB_ASCAST]], align 41196// CHECK-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB_ASCAST]], align 41197// CHECK-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV_ASCAST]], align 41198// CHECK-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 41199// CHECK-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 41200// CHECK-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]1201// CHECK-NEXT: br i1 [[CMP5]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]1202// CHECK: omp.dispatch.body:1203// CHECK-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1204// CHECK: omp.inner.for.cond:1205// CHECK-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 41206// CHECK-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 41207// CHECK-NEXT: [[CMP6:%.*]] = icmp sle i32 [[TMP16]], [[TMP17]]1208// CHECK-NEXT: br i1 [[CMP6]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1209// CHECK: omp.inner.for.body:1210// CHECK-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 41211// CHECK-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP18]], 11212// CHECK-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1213// CHECK-NEXT: store i32 [[ADD]], ptr [[J3_ASCAST]], align 41214// CHECK-NEXT: [[TMP19:%.*]] = load i32, ptr [[J3_ASCAST]], align 41215// CHECK-NEXT: [[TMP20:%.*]] = load i32, ptr [[J3_ASCAST]], align 41216// CHECK-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP20]] to i641217// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i64 [[IDXPROM]]1218// CHECK-NEXT: store i32 [[TMP19]], ptr [[ARRAYIDX]], align 41219// CHECK-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1220// CHECK: omp.body.continue:1221// CHECK-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1222// CHECK: omp.inner.for.inc:1223// CHECK-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 41224// CHECK-NEXT: [[ADD7:%.*]] = add nsw i32 [[TMP21]], 11225// CHECK-NEXT: store i32 [[ADD7]], ptr [[DOTOMP_IV_ASCAST]], align 41226// CHECK-NEXT: br label [[OMP_INNER_FOR_COND]]1227// CHECK: omp.inner.for.end:1228// CHECK-NEXT: br label [[OMP_DISPATCH_INC:%.*]]1229// CHECK: omp.dispatch.inc:1230// CHECK-NEXT: [[TMP22:%.*]] = load i32, ptr [[DOTOMP_LB_ASCAST]], align 41231// CHECK-NEXT: [[TMP23:%.*]] = load i32, ptr [[DOTOMP_STRIDE_ASCAST]], align 41232// CHECK-NEXT: [[ADD8:%.*]] = add nsw i32 [[TMP22]], [[TMP23]]1233// CHECK-NEXT: store i32 [[ADD8]], ptr [[DOTOMP_LB_ASCAST]], align 41234// CHECK-NEXT: [[TMP24:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 41235// CHECK-NEXT: [[TMP25:%.*]] = load i32, ptr [[DOTOMP_STRIDE_ASCAST]], align 41236// CHECK-NEXT: [[ADD9:%.*]] = add nsw i32 [[TMP24]], [[TMP25]]1237// CHECK-NEXT: store i32 [[ADD9]], ptr [[DOTOMP_UB_ASCAST]], align 41238// CHECK-NEXT: br label [[OMP_DISPATCH_COND]]1239// CHECK: omp.dispatch.end:1240// CHECK-NEXT: [[TMP26:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 81241// CHECK-NEXT: [[TMP27:%.*]] = load i32, ptr [[TMP26]], align 41242// CHECK-NEXT: call void @__kmpc_for_static_fini(ptr addrspacecast (ptr addrspace(1) @[[GLOB3]] to ptr), i32 [[TMP27]])1243// CHECK-NEXT: br label [[OMP_PRECOND_END]]1244// CHECK: omp.precond.end:1245// CHECK-NEXT: ret void1246//1247//1248// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l76_omp_outlined_omp_outlined_wrapper1249// CHECK-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR7]] {1250// CHECK-NEXT: entry:1251// CHECK-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2, addrspace(5)1252// CHECK-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4, addrspace(5)1253// CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4, addrspace(5)1254// CHECK-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 8, addrspace(5)1255// CHECK-NEXT: [[DOTADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTADDR]] to ptr1256// CHECK-NEXT: [[DOTADDR1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTADDR1]] to ptr1257// CHECK-NEXT: [[DOTZERO_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTZERO_ADDR]] to ptr1258// CHECK-NEXT: [[GLOBAL_ARGS_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[GLOBAL_ARGS]] to ptr1259// CHECK-NEXT: store i16 [[TMP0]], ptr [[DOTADDR_ASCAST]], align 21260// CHECK-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1_ASCAST]], align 41261// CHECK-NEXT: store i32 0, ptr [[DOTZERO_ADDR_ASCAST]], align 41262// CHECK-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS_ASCAST]])1263// CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[GLOBAL_ARGS_ASCAST]], align 81264// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 01265// CHECK-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP3]], align 81266// CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 11267// CHECK-NEXT: [[TMP6:%.*]] = load i64, ptr [[TMP5]], align 81268// CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 21269// CHECK-NEXT: [[TMP8:%.*]] = load ptr, ptr [[TMP7]], align 81270// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l76_omp_outlined_omp_outlined(ptr [[DOTADDR1_ASCAST]], ptr [[DOTZERO_ADDR_ASCAST]], ptr [[TMP4]], i64 [[TMP6]], ptr [[TMP8]]) #[[ATTR4]]1271// CHECK-NEXT: ret void1272//1273