brintos / llvm-project-archived public Read only

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// RUN: %clang_cc1 -DCHECK -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK13// RUN: %clang_cc1 -DCHECK -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s4// RUN: %clang_cc1 -DCHECK -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK15// RUN: %clang_cc1 -DCHECK -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK36// RUN: %clang_cc1 -DCHECK -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s7// RUN: %clang_cc1 -DCHECK -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK38 9// RUN: %clang_cc1 -DLAMBDA -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK510// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s11// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++  -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK512 13// RUN: %clang_cc1 -DCHECK -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"14// RUN: %clang_cc1 -DCHECK -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s15// RUN: %clang_cc1 -DCHECK -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"16// RUN: %clang_cc1 -DCHECK -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"17// RUN: %clang_cc1 -DCHECK -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s18// RUN: %clang_cc1 -DCHECK -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"19 20// RUN: %clang_cc1 -DLAMBDA -verify -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"21// RUN: %clang_cc1 -DLAMBDA -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s22// RUN: %clang_cc1 -DLAMBDA -fopenmp-simd -x c++  -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"23 24// expected-no-diagnostics25#ifndef HEADER26#define HEADER27 28template <typename T>29T tmain() {30  T t_var = T();31  T vec[] = {1, 2};32#pragma omp target teams distribute parallel for reduction(+: t_var)33  for (int i = 0; i < 2; ++i) {34    t_var += (T) i;35  }36  return T();37}38 39int main() {40  static int sivar;41#ifdef LAMBDA42 43  [&]() {44#pragma omp target teams distribute parallel for reduction(+: sivar)45  for (int i = 0; i < 2; ++i) {46 47    // Skip global and bound tid vars48 49 50 51    // Skip global and bound tid vars, and prev lb and ub vars52    // skip loop vars53 54 55    sivar += i;56 57    [&]() {58 59      sivar += 4;60 61    }();62  }63  }();64  return 0;65#else66#pragma omp target teams distribute parallel for reduction(+: sivar)67  for (int i = 0; i < 2; ++i) {68    sivar += i;69  }70  return tmain<int>();71#endif72}73 74 75 76 77// Skip global and bound tid vars78 79 80// Skip global and bound tid vars, and prev lb and ub81// skip loop vars82 83 84 85 86// Skip global and bound tid vars87 88 89// Skip global and bound tid vars, and prev lb and ub vars90// skip loop vars91 92#endif93// CHECK1-LABEL: define {{[^@]+}}@main94// CHECK1-SAME: () #[[ATTR0:[0-9]+]] {95// CHECK1-NEXT:  entry:96// CHECK1-NEXT:    [[RETVAL:%.*]] = alloca i32, align 497// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 898// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 899// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8100// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4101// CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8102// CHECK1-NEXT:    store i32 0, ptr [[RETVAL]], align 4103// CHECK1-NEXT:    [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0104// CHECK1-NEXT:    store ptr @_ZZ4mainE5sivar, ptr [[TMP0]], align 8105// CHECK1-NEXT:    [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0106// CHECK1-NEXT:    store ptr @_ZZ4mainE5sivar, ptr [[TMP1]], align 8107// CHECK1-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0108// CHECK1-NEXT:    store ptr null, ptr [[TMP2]], align 8109// CHECK1-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0110// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0111// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0112// CHECK1-NEXT:    store i32 3, ptr [[TMP5]], align 4113// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1114// CHECK1-NEXT:    store i32 1, ptr [[TMP6]], align 4115// CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2116// CHECK1-NEXT:    store ptr [[TMP3]], ptr [[TMP7]], align 8117// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3118// CHECK1-NEXT:    store ptr [[TMP4]], ptr [[TMP8]], align 8119// CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4120// CHECK1-NEXT:    store ptr @.offload_sizes, ptr [[TMP9]], align 8121// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5122// CHECK1-NEXT:    store ptr @.offload_maptypes, ptr [[TMP10]], align 8123// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6124// CHECK1-NEXT:    store ptr null, ptr [[TMP11]], align 8125// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7126// CHECK1-NEXT:    store ptr null, ptr [[TMP12]], align 8127// CHECK1-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8128// CHECK1-NEXT:    store i64 2, ptr [[TMP13]], align 8129// CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9130// CHECK1-NEXT:    store i64 0, ptr [[TMP14]], align 8131// CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10132// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4133// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11134// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4135// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12136// CHECK1-NEXT:    store i32 0, ptr [[TMP17]], align 4137// CHECK1-NEXT:    [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB4:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66.region_id, ptr [[KERNEL_ARGS]])138// CHECK1-NEXT:    [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0139// CHECK1-NEXT:    br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]140// CHECK1:       omp_offload.failed:141// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66(ptr @_ZZ4mainE5sivar) #[[ATTR2:[0-9]+]]142// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]143// CHECK1:       omp_offload.cont:144// CHECK1-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_Z5tmainIiET_v()145// CHECK1-NEXT:    ret i32 [[CALL]]146//147//148// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66149// CHECK1-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR1:[0-9]+]] {150// CHECK1-NEXT:  entry:151// CHECK1-NEXT:    [[SIVAR_ADDR:%.*]] = alloca ptr, align 8152// CHECK1-NEXT:    store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 8153// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 8154// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB4]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66.omp_outlined, ptr [[TMP0]])155// CHECK1-NEXT:    ret void156//157//158// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66.omp_outlined159// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR1]] {160// CHECK1-NEXT:  entry:161// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8162// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8163// CHECK1-NEXT:    [[SIVAR_ADDR:%.*]] = alloca ptr, align 8164// CHECK1-NEXT:    [[SIVAR1:%.*]] = alloca i32, align 4165// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4166// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4167// CHECK1-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4168// CHECK1-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4169// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4170// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4171// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4172// CHECK1-NEXT:    [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 8173// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8174// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8175// CHECK1-NEXT:    store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 8176// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 8177// CHECK1-NEXT:    store i32 0, ptr [[SIVAR1]], align 4178// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 4179// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_COMB_UB]], align 4180// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4181// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4182// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8183// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4184// CHECK1-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)185// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4186// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1187// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]188// CHECK1:       cond.true:189// CHECK1-NEXT:    br label [[COND_END:%.*]]190// CHECK1:       cond.false:191// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4192// CHECK1-NEXT:    br label [[COND_END]]193// CHECK1:       cond.end:194// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]195// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4196// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4197// CHECK1-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4198// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]199// CHECK1:       omp.inner.for.cond:200// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4201// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4202// CHECK1-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]203// CHECK1-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]204// CHECK1:       omp.inner.for.body:205// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4206// CHECK1-NEXT:    [[TMP9:%.*]] = zext i32 [[TMP8]] to i64207// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4208// CHECK1-NEXT:    [[TMP11:%.*]] = zext i32 [[TMP10]] to i64209// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB4]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66.omp_outlined.omp_outlined, i64 [[TMP9]], i64 [[TMP11]], ptr [[SIVAR1]])210// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]211// CHECK1:       omp.inner.for.inc:212// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4213// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4214// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP12]], [[TMP13]]215// CHECK1-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4216// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]217// CHECK1:       omp.inner.for.end:218// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]219// CHECK1:       omp.loop.exit:220// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])221// CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0222// CHECK1-NEXT:    store ptr [[SIVAR1]], ptr [[TMP14]], align 8223// CHECK1-NEXT:    [[TMP15:%.*]] = call i32 @__kmpc_reduce_nowait(ptr @[[GLOB3:[0-9]+]], i32 [[TMP2]], i32 1, i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)224// CHECK1-NEXT:    switch i32 [[TMP15]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [225// CHECK1-NEXT:      i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]226// CHECK1-NEXT:      i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]227// CHECK1-NEXT:    ]228// CHECK1:       .omp.reduction.case1:229// CHECK1-NEXT:    [[TMP16:%.*]] = load i32, ptr [[TMP0]], align 4230// CHECK1-NEXT:    [[TMP17:%.*]] = load i32, ptr [[SIVAR1]], align 4231// CHECK1-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP16]], [[TMP17]]232// CHECK1-NEXT:    store i32 [[ADD3]], ptr [[TMP0]], align 4233// CHECK1-NEXT:    call void @__kmpc_end_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)234// CHECK1-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]235// CHECK1:       .omp.reduction.case2:236// CHECK1-NEXT:    [[TMP18:%.*]] = load i32, ptr [[SIVAR1]], align 4237// CHECK1-NEXT:    [[TMP19:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP18]] monotonic, align 4238// CHECK1-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]239// CHECK1:       .omp.reduction.default:240// CHECK1-NEXT:    ret void241//242//243// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66.omp_outlined.omp_outlined244// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR1]] {245// CHECK1-NEXT:  entry:246// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8247// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8248// CHECK1-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8249// CHECK1-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8250// CHECK1-NEXT:    [[SIVAR_ADDR:%.*]] = alloca ptr, align 8251// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4252// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4253// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4254// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4255// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4256// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4257// CHECK1-NEXT:    [[SIVAR2:%.*]] = alloca i32, align 4258// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4259// CHECK1-NEXT:    [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 8260// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8261// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8262// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8263// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8264// CHECK1-NEXT:    store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 8265// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 8266// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4267// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 4268// CHECK1-NEXT:    [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8269// CHECK1-NEXT:    [[CONV:%.*]] = trunc i64 [[TMP1]] to i32270// CHECK1-NEXT:    [[TMP2:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8271// CHECK1-NEXT:    [[CONV1:%.*]] = trunc i64 [[TMP2]] to i32272// CHECK1-NEXT:    store i32 [[CONV]], ptr [[DOTOMP_LB]], align 4273// CHECK1-NEXT:    store i32 [[CONV1]], ptr [[DOTOMP_UB]], align 4274// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4275// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4276// CHECK1-NEXT:    store i32 0, ptr [[SIVAR2]], align 4277// CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8278// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[TMP3]], align 4279// CHECK1-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP4]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)280// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4281// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP5]], 1282// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]283// CHECK1:       cond.true:284// CHECK1-NEXT:    br label [[COND_END:%.*]]285// CHECK1:       cond.false:286// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4287// CHECK1-NEXT:    br label [[COND_END]]288// CHECK1:       cond.end:289// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP6]], [[COND_FALSE]] ]290// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4291// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4292// CHECK1-NEXT:    store i32 [[TMP7]], ptr [[DOTOMP_IV]], align 4293// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]294// CHECK1:       omp.inner.for.cond:295// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4296// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4297// CHECK1-NEXT:    [[CMP3:%.*]] = icmp sle i32 [[TMP8]], [[TMP9]]298// CHECK1-NEXT:    br i1 [[CMP3]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]299// CHECK1:       omp.inner.for.body:300// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4301// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP10]], 1302// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]303// CHECK1-NEXT:    store i32 [[ADD]], ptr [[I]], align 4304// CHECK1-NEXT:    [[TMP11:%.*]] = load i32, ptr [[I]], align 4305// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[SIVAR2]], align 4306// CHECK1-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP12]], [[TMP11]]307// CHECK1-NEXT:    store i32 [[ADD4]], ptr [[SIVAR2]], align 4308// CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]309// CHECK1:       omp.body.continue:310// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]311// CHECK1:       omp.inner.for.inc:312// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4313// CHECK1-NEXT:    [[ADD5:%.*]] = add nsw i32 [[TMP13]], 1314// CHECK1-NEXT:    store i32 [[ADD5]], ptr [[DOTOMP_IV]], align 4315// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]316// CHECK1:       omp.inner.for.end:317// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]318// CHECK1:       omp.loop.exit:319// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP4]])320// CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0321// CHECK1-NEXT:    store ptr [[SIVAR2]], ptr [[TMP14]], align 8322// CHECK1-NEXT:    [[TMP15:%.*]] = call i32 @__kmpc_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP4]], i32 1, i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66.omp_outlined.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)323// CHECK1-NEXT:    switch i32 [[TMP15]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [324// CHECK1-NEXT:      i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]325// CHECK1-NEXT:      i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]326// CHECK1-NEXT:    ]327// CHECK1:       .omp.reduction.case1:328// CHECK1-NEXT:    [[TMP16:%.*]] = load i32, ptr [[TMP0]], align 4329// CHECK1-NEXT:    [[TMP17:%.*]] = load i32, ptr [[SIVAR2]], align 4330// CHECK1-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP16]], [[TMP17]]331// CHECK1-NEXT:    store i32 [[ADD6]], ptr [[TMP0]], align 4332// CHECK1-NEXT:    call void @__kmpc_end_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP4]], ptr @.gomp_critical_user_.reduction.var)333// CHECK1-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]334// CHECK1:       .omp.reduction.case2:335// CHECK1-NEXT:    [[TMP18:%.*]] = load i32, ptr [[SIVAR2]], align 4336// CHECK1-NEXT:    [[TMP19:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP18]] monotonic, align 4337// CHECK1-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]338// CHECK1:       .omp.reduction.default:339// CHECK1-NEXT:    ret void340//341//342// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66.omp_outlined.omp_outlined.omp.reduction.reduction_func343// CHECK1-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3:[0-9]+]] {344// CHECK1-NEXT:  entry:345// CHECK1-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 8346// CHECK1-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 8347// CHECK1-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 8348// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 8349// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 8350// CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 8351// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i64 0, i64 0352// CHECK1-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 8353// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i64 0, i64 0354// CHECK1-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 8355// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4356// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4357// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]358// CHECK1-NEXT:    store i32 [[ADD]], ptr [[TMP7]], align 4359// CHECK1-NEXT:    ret void360//361//362// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66.omp_outlined.omp.reduction.reduction_func363// CHECK1-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3]] {364// CHECK1-NEXT:  entry:365// CHECK1-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 8366// CHECK1-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 8367// CHECK1-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 8368// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 8369// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 8370// CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 8371// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i64 0, i64 0372// CHECK1-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 8373// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i64 0, i64 0374// CHECK1-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 8375// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4376// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4377// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]378// CHECK1-NEXT:    store i32 [[ADD]], ptr [[TMP7]], align 4379// CHECK1-NEXT:    ret void380//381//382// CHECK1-LABEL: define {{[^@]+}}@_Z5tmainIiET_v383// CHECK1-SAME: () #[[ATTR5:[0-9]+]] comdat {384// CHECK1-NEXT:  entry:385// CHECK1-NEXT:    [[T_VAR:%.*]] = alloca i32, align 4386// CHECK1-NEXT:    [[VEC:%.*]] = alloca [2 x i32], align 4387// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8388// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8389// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8390// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4391// CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8392// CHECK1-NEXT:    store i32 0, ptr [[T_VAR]], align 4393// CHECK1-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[VEC]], ptr align 4 @__const._Z5tmainIiET_v.vec, i64 8, i1 false)394// CHECK1-NEXT:    [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0395// CHECK1-NEXT:    store ptr [[T_VAR]], ptr [[TMP0]], align 8396// CHECK1-NEXT:    [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0397// CHECK1-NEXT:    store ptr [[T_VAR]], ptr [[TMP1]], align 8398// CHECK1-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0399// CHECK1-NEXT:    store ptr null, ptr [[TMP2]], align 8400// CHECK1-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0401// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0402// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0403// CHECK1-NEXT:    store i32 3, ptr [[TMP5]], align 4404// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1405// CHECK1-NEXT:    store i32 1, ptr [[TMP6]], align 4406// CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2407// CHECK1-NEXT:    store ptr [[TMP3]], ptr [[TMP7]], align 8408// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3409// CHECK1-NEXT:    store ptr [[TMP4]], ptr [[TMP8]], align 8410// CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4411// CHECK1-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP9]], align 8412// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5413// CHECK1-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP10]], align 8414// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6415// CHECK1-NEXT:    store ptr null, ptr [[TMP11]], align 8416// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7417// CHECK1-NEXT:    store ptr null, ptr [[TMP12]], align 8418// CHECK1-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8419// CHECK1-NEXT:    store i64 2, ptr [[TMP13]], align 8420// CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9421// CHECK1-NEXT:    store i64 0, ptr [[TMP14]], align 8422// CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10423// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4424// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11425// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4426// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12427// CHECK1-NEXT:    store i32 0, ptr [[TMP17]], align 4428// CHECK1-NEXT:    [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB4]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.region_id, ptr [[KERNEL_ARGS]])429// CHECK1-NEXT:    [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0430// CHECK1-NEXT:    br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]431// CHECK1:       omp_offload.failed:432// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32(ptr [[T_VAR]]) #[[ATTR2]]433// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]434// CHECK1:       omp_offload.cont:435// CHECK1-NEXT:    ret i32 0436//437//438// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32439// CHECK1-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[T_VAR:%.*]]) #[[ATTR1]] {440// CHECK1-NEXT:  entry:441// CHECK1-NEXT:    [[T_VAR_ADDR:%.*]] = alloca ptr, align 8442// CHECK1-NEXT:    store ptr [[T_VAR]], ptr [[T_VAR_ADDR]], align 8443// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[T_VAR_ADDR]], align 8444// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB4]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined, ptr [[TMP0]])445// CHECK1-NEXT:    ret void446//447//448// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined449// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[T_VAR:%.*]]) #[[ATTR1]] {450// CHECK1-NEXT:  entry:451// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8452// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8453// CHECK1-NEXT:    [[T_VAR_ADDR:%.*]] = alloca ptr, align 8454// CHECK1-NEXT:    [[T_VAR1:%.*]] = alloca i32, align 4455// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4456// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4457// CHECK1-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4458// CHECK1-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4459// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4460// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4461// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4462// CHECK1-NEXT:    [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 8463// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8464// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8465// CHECK1-NEXT:    store ptr [[T_VAR]], ptr [[T_VAR_ADDR]], align 8466// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[T_VAR_ADDR]], align 8467// CHECK1-NEXT:    store i32 0, ptr [[T_VAR1]], align 4468// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 4469// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_COMB_UB]], align 4470// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4471// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4472// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8473// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4474// CHECK1-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)475// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4476// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1477// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]478// CHECK1:       cond.true:479// CHECK1-NEXT:    br label [[COND_END:%.*]]480// CHECK1:       cond.false:481// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4482// CHECK1-NEXT:    br label [[COND_END]]483// CHECK1:       cond.end:484// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]485// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4486// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4487// CHECK1-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4488// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]489// CHECK1:       omp.inner.for.cond:490// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4491// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4492// CHECK1-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]493// CHECK1-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]494// CHECK1:       omp.inner.for.body:495// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4496// CHECK1-NEXT:    [[TMP9:%.*]] = zext i32 [[TMP8]] to i64497// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4498// CHECK1-NEXT:    [[TMP11:%.*]] = zext i32 [[TMP10]] to i64499// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB4]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp_outlined, i64 [[TMP9]], i64 [[TMP11]], ptr [[T_VAR1]])500// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]501// CHECK1:       omp.inner.for.inc:502// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4503// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4504// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP12]], [[TMP13]]505// CHECK1-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4506// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]507// CHECK1:       omp.inner.for.end:508// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]509// CHECK1:       omp.loop.exit:510// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])511// CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0512// CHECK1-NEXT:    store ptr [[T_VAR1]], ptr [[TMP14]], align 8513// CHECK1-NEXT:    [[TMP15:%.*]] = call i32 @__kmpc_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP2]], i32 1, i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)514// CHECK1-NEXT:    switch i32 [[TMP15]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [515// CHECK1-NEXT:      i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]516// CHECK1-NEXT:      i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]517// CHECK1-NEXT:    ]518// CHECK1:       .omp.reduction.case1:519// CHECK1-NEXT:    [[TMP16:%.*]] = load i32, ptr [[TMP0]], align 4520// CHECK1-NEXT:    [[TMP17:%.*]] = load i32, ptr [[T_VAR1]], align 4521// CHECK1-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP16]], [[TMP17]]522// CHECK1-NEXT:    store i32 [[ADD3]], ptr [[TMP0]], align 4523// CHECK1-NEXT:    call void @__kmpc_end_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)524// CHECK1-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]525// CHECK1:       .omp.reduction.case2:526// CHECK1-NEXT:    [[TMP18:%.*]] = load i32, ptr [[T_VAR1]], align 4527// CHECK1-NEXT:    [[TMP19:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP18]] monotonic, align 4528// CHECK1-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]529// CHECK1:       .omp.reduction.default:530// CHECK1-NEXT:    ret void531//532//533// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp_outlined534// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[T_VAR:%.*]]) #[[ATTR1]] {535// CHECK1-NEXT:  entry:536// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8537// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8538// CHECK1-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8539// CHECK1-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8540// CHECK1-NEXT:    [[T_VAR_ADDR:%.*]] = alloca ptr, align 8541// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4542// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4543// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4544// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4545// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4546// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4547// CHECK1-NEXT:    [[T_VAR2:%.*]] = alloca i32, align 4548// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4549// CHECK1-NEXT:    [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 8550// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8551// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8552// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8553// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8554// CHECK1-NEXT:    store ptr [[T_VAR]], ptr [[T_VAR_ADDR]], align 8555// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[T_VAR_ADDR]], align 8556// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4557// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 4558// CHECK1-NEXT:    [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8559// CHECK1-NEXT:    [[CONV:%.*]] = trunc i64 [[TMP1]] to i32560// CHECK1-NEXT:    [[TMP2:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8561// CHECK1-NEXT:    [[CONV1:%.*]] = trunc i64 [[TMP2]] to i32562// CHECK1-NEXT:    store i32 [[CONV]], ptr [[DOTOMP_LB]], align 4563// CHECK1-NEXT:    store i32 [[CONV1]], ptr [[DOTOMP_UB]], align 4564// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4565// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4566// CHECK1-NEXT:    store i32 0, ptr [[T_VAR2]], align 4567// CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8568// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[TMP3]], align 4569// CHECK1-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB2]], i32 [[TMP4]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)570// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4571// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP5]], 1572// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]573// CHECK1:       cond.true:574// CHECK1-NEXT:    br label [[COND_END:%.*]]575// CHECK1:       cond.false:576// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4577// CHECK1-NEXT:    br label [[COND_END]]578// CHECK1:       cond.end:579// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP6]], [[COND_FALSE]] ]580// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4581// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4582// CHECK1-NEXT:    store i32 [[TMP7]], ptr [[DOTOMP_IV]], align 4583// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]584// CHECK1:       omp.inner.for.cond:585// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4586// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4587// CHECK1-NEXT:    [[CMP3:%.*]] = icmp sle i32 [[TMP8]], [[TMP9]]588// CHECK1-NEXT:    br i1 [[CMP3]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]589// CHECK1:       omp.inner.for.body:590// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4591// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP10]], 1592// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]593// CHECK1-NEXT:    store i32 [[ADD]], ptr [[I]], align 4594// CHECK1-NEXT:    [[TMP11:%.*]] = load i32, ptr [[I]], align 4595// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[T_VAR2]], align 4596// CHECK1-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP12]], [[TMP11]]597// CHECK1-NEXT:    store i32 [[ADD4]], ptr [[T_VAR2]], align 4598// CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]599// CHECK1:       omp.body.continue:600// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]601// CHECK1:       omp.inner.for.inc:602// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4603// CHECK1-NEXT:    [[ADD5:%.*]] = add nsw i32 [[TMP13]], 1604// CHECK1-NEXT:    store i32 [[ADD5]], ptr [[DOTOMP_IV]], align 4605// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]606// CHECK1:       omp.inner.for.end:607// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]608// CHECK1:       omp.loop.exit:609// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP4]])610// CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0611// CHECK1-NEXT:    store ptr [[T_VAR2]], ptr [[TMP14]], align 8612// CHECK1-NEXT:    [[TMP15:%.*]] = call i32 @__kmpc_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP4]], i32 1, i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)613// CHECK1-NEXT:    switch i32 [[TMP15]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [614// CHECK1-NEXT:      i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]615// CHECK1-NEXT:      i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]616// CHECK1-NEXT:    ]617// CHECK1:       .omp.reduction.case1:618// CHECK1-NEXT:    [[TMP16:%.*]] = load i32, ptr [[TMP0]], align 4619// CHECK1-NEXT:    [[TMP17:%.*]] = load i32, ptr [[T_VAR2]], align 4620// CHECK1-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP16]], [[TMP17]]621// CHECK1-NEXT:    store i32 [[ADD6]], ptr [[TMP0]], align 4622// CHECK1-NEXT:    call void @__kmpc_end_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP4]], ptr @.gomp_critical_user_.reduction.var)623// CHECK1-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]624// CHECK1:       .omp.reduction.case2:625// CHECK1-NEXT:    [[TMP18:%.*]] = load i32, ptr [[T_VAR2]], align 4626// CHECK1-NEXT:    [[TMP19:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP18]] monotonic, align 4627// CHECK1-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]628// CHECK1:       .omp.reduction.default:629// CHECK1-NEXT:    ret void630//631//632// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp_outlined.omp.reduction.reduction_func633// CHECK1-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3]] {634// CHECK1-NEXT:  entry:635// CHECK1-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 8636// CHECK1-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 8637// CHECK1-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 8638// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 8639// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 8640// CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 8641// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i64 0, i64 0642// CHECK1-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 8643// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i64 0, i64 0644// CHECK1-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 8645// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4646// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4647// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]648// CHECK1-NEXT:    store i32 [[ADD]], ptr [[TMP7]], align 4649// CHECK1-NEXT:    ret void650//651//652// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp.reduction.reduction_func653// CHECK1-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3]] {654// CHECK1-NEXT:  entry:655// CHECK1-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 8656// CHECK1-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 8657// CHECK1-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 8658// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 8659// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 8660// CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 8661// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i64 0, i64 0662// CHECK1-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 8663// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i64 0, i64 0664// CHECK1-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 8665// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4666// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4667// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]668// CHECK1-NEXT:    store i32 [[ADD]], ptr [[TMP7]], align 4669// CHECK1-NEXT:    ret void670//671//672// CHECK3-LABEL: define {{[^@]+}}@main673// CHECK3-SAME: () #[[ATTR0:[0-9]+]] {674// CHECK3-NEXT:  entry:675// CHECK3-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4676// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4677// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4678// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4679// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 4680// CHECK3-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8681// CHECK3-NEXT:    store i32 0, ptr [[RETVAL]], align 4682// CHECK3-NEXT:    [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0683// CHECK3-NEXT:    store ptr @_ZZ4mainE5sivar, ptr [[TMP0]], align 4684// CHECK3-NEXT:    [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0685// CHECK3-NEXT:    store ptr @_ZZ4mainE5sivar, ptr [[TMP1]], align 4686// CHECK3-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0687// CHECK3-NEXT:    store ptr null, ptr [[TMP2]], align 4688// CHECK3-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0689// CHECK3-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0690// CHECK3-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0691// CHECK3-NEXT:    store i32 3, ptr [[TMP5]], align 4692// CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1693// CHECK3-NEXT:    store i32 1, ptr [[TMP6]], align 4694// CHECK3-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2695// CHECK3-NEXT:    store ptr [[TMP3]], ptr [[TMP7]], align 4696// CHECK3-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3697// CHECK3-NEXT:    store ptr [[TMP4]], ptr [[TMP8]], align 4698// CHECK3-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4699// CHECK3-NEXT:    store ptr @.offload_sizes, ptr [[TMP9]], align 4700// CHECK3-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5701// CHECK3-NEXT:    store ptr @.offload_maptypes, ptr [[TMP10]], align 4702// CHECK3-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6703// CHECK3-NEXT:    store ptr null, ptr [[TMP11]], align 4704// CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7705// CHECK3-NEXT:    store ptr null, ptr [[TMP12]], align 4706// CHECK3-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8707// CHECK3-NEXT:    store i64 2, ptr [[TMP13]], align 8708// CHECK3-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9709// CHECK3-NEXT:    store i64 0, ptr [[TMP14]], align 8710// CHECK3-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10711// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4712// CHECK3-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11713// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4714// CHECK3-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12715// CHECK3-NEXT:    store i32 0, ptr [[TMP17]], align 4716// CHECK3-NEXT:    [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB4:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66.region_id, ptr [[KERNEL_ARGS]])717// CHECK3-NEXT:    [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0718// CHECK3-NEXT:    br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]719// CHECK3:       omp_offload.failed:720// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66(ptr @_ZZ4mainE5sivar) #[[ATTR2:[0-9]+]]721// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT]]722// CHECK3:       omp_offload.cont:723// CHECK3-NEXT:    [[CALL:%.*]] = call noundef i32 @_Z5tmainIiET_v()724// CHECK3-NEXT:    ret i32 [[CALL]]725//726//727// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66728// CHECK3-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR1:[0-9]+]] {729// CHECK3-NEXT:  entry:730// CHECK3-NEXT:    [[SIVAR_ADDR:%.*]] = alloca ptr, align 4731// CHECK3-NEXT:    store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 4732// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 4733// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB4]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66.omp_outlined, ptr [[TMP0]])734// CHECK3-NEXT:    ret void735//736//737// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66.omp_outlined738// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR1]] {739// CHECK3-NEXT:  entry:740// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4741// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4742// CHECK3-NEXT:    [[SIVAR_ADDR:%.*]] = alloca ptr, align 4743// CHECK3-NEXT:    [[SIVAR1:%.*]] = alloca i32, align 4744// CHECK3-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4745// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 4746// CHECK3-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4747// CHECK3-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4748// CHECK3-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4749// CHECK3-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4750// CHECK3-NEXT:    [[I:%.*]] = alloca i32, align 4751// CHECK3-NEXT:    [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 4752// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4753// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4754// CHECK3-NEXT:    store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 4755// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 4756// CHECK3-NEXT:    store i32 0, ptr [[SIVAR1]], align 4757// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 4758// CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_COMB_UB]], align 4759// CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4760// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4761// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4762// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4763// CHECK3-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)764// CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4765// CHECK3-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1766// CHECK3-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]767// CHECK3:       cond.true:768// CHECK3-NEXT:    br label [[COND_END:%.*]]769// CHECK3:       cond.false:770// CHECK3-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4771// CHECK3-NEXT:    br label [[COND_END]]772// CHECK3:       cond.end:773// CHECK3-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]774// CHECK3-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4775// CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4776// CHECK3-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4777// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]778// CHECK3:       omp.inner.for.cond:779// CHECK3-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4780// CHECK3-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4781// CHECK3-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]782// CHECK3-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]783// CHECK3:       omp.inner.for.body:784// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4785// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4786// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB4]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66.omp_outlined.omp_outlined, i32 [[TMP8]], i32 [[TMP9]], ptr [[SIVAR1]])787// CHECK3-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]788// CHECK3:       omp.inner.for.inc:789// CHECK3-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4790// CHECK3-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4791// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP10]], [[TMP11]]792// CHECK3-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4793// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND]]794// CHECK3:       omp.inner.for.end:795// CHECK3-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]796// CHECK3:       omp.loop.exit:797// CHECK3-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])798// CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 0799// CHECK3-NEXT:    store ptr [[SIVAR1]], ptr [[TMP12]], align 4800// CHECK3-NEXT:    [[TMP13:%.*]] = call i32 @__kmpc_reduce_nowait(ptr @[[GLOB3:[0-9]+]], i32 [[TMP2]], i32 1, i32 4, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)801// CHECK3-NEXT:    switch i32 [[TMP13]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [802// CHECK3-NEXT:      i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]803// CHECK3-NEXT:      i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]804// CHECK3-NEXT:    ]805// CHECK3:       .omp.reduction.case1:806// CHECK3-NEXT:    [[TMP14:%.*]] = load i32, ptr [[TMP0]], align 4807// CHECK3-NEXT:    [[TMP15:%.*]] = load i32, ptr [[SIVAR1]], align 4808// CHECK3-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP14]], [[TMP15]]809// CHECK3-NEXT:    store i32 [[ADD3]], ptr [[TMP0]], align 4810// CHECK3-NEXT:    call void @__kmpc_end_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)811// CHECK3-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]812// CHECK3:       .omp.reduction.case2:813// CHECK3-NEXT:    [[TMP16:%.*]] = load i32, ptr [[SIVAR1]], align 4814// CHECK3-NEXT:    [[TMP17:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP16]] monotonic, align 4815// CHECK3-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]816// CHECK3:       .omp.reduction.default:817// CHECK3-NEXT:    ret void818//819//820// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66.omp_outlined.omp_outlined821// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[DOTPREVIOUS_LB_:%.*]], i32 noundef [[DOTPREVIOUS_UB_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR1]] {822// CHECK3-NEXT:  entry:823// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4824// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4825// CHECK3-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i32, align 4826// CHECK3-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i32, align 4827// CHECK3-NEXT:    [[SIVAR_ADDR:%.*]] = alloca ptr, align 4828// CHECK3-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4829// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 4830// CHECK3-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4831// CHECK3-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4832// CHECK3-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4833// CHECK3-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4834// CHECK3-NEXT:    [[SIVAR1:%.*]] = alloca i32, align 4835// CHECK3-NEXT:    [[I:%.*]] = alloca i32, align 4836// CHECK3-NEXT:    [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 4837// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4838// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4839// CHECK3-NEXT:    store i32 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 4840// CHECK3-NEXT:    store i32 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 4841// CHECK3-NEXT:    store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 4842// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 4843// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4844// CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 4845// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTPREVIOUS_LB__ADDR]], align 4846// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTPREVIOUS_UB__ADDR]], align 4847// CHECK3-NEXT:    store i32 [[TMP1]], ptr [[DOTOMP_LB]], align 4848// CHECK3-NEXT:    store i32 [[TMP2]], ptr [[DOTOMP_UB]], align 4849// CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4850// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4851// CHECK3-NEXT:    store i32 0, ptr [[SIVAR1]], align 4852// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4853// CHECK3-NEXT:    [[TMP4:%.*]] = load i32, ptr [[TMP3]], align 4854// CHECK3-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP4]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)855// CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4856// CHECK3-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP5]], 1857// CHECK3-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]858// CHECK3:       cond.true:859// CHECK3-NEXT:    br label [[COND_END:%.*]]860// CHECK3:       cond.false:861// CHECK3-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4862// CHECK3-NEXT:    br label [[COND_END]]863// CHECK3:       cond.end:864// CHECK3-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP6]], [[COND_FALSE]] ]865// CHECK3-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4866// CHECK3-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4867// CHECK3-NEXT:    store i32 [[TMP7]], ptr [[DOTOMP_IV]], align 4868// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]869// CHECK3:       omp.inner.for.cond:870// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4871// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4872// CHECK3-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP8]], [[TMP9]]873// CHECK3-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]874// CHECK3:       omp.inner.for.body:875// CHECK3-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4876// CHECK3-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP10]], 1877// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]878// CHECK3-NEXT:    store i32 [[ADD]], ptr [[I]], align 4879// CHECK3-NEXT:    [[TMP11:%.*]] = load i32, ptr [[I]], align 4880// CHECK3-NEXT:    [[TMP12:%.*]] = load i32, ptr [[SIVAR1]], align 4881// CHECK3-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP12]], [[TMP11]]882// CHECK3-NEXT:    store i32 [[ADD3]], ptr [[SIVAR1]], align 4883// CHECK3-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]884// CHECK3:       omp.body.continue:885// CHECK3-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]886// CHECK3:       omp.inner.for.inc:887// CHECK3-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4888// CHECK3-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP13]], 1889// CHECK3-NEXT:    store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4890// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND]]891// CHECK3:       omp.inner.for.end:892// CHECK3-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]893// CHECK3:       omp.loop.exit:894// CHECK3-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP4]])895// CHECK3-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 0896// CHECK3-NEXT:    store ptr [[SIVAR1]], ptr [[TMP14]], align 4897// CHECK3-NEXT:    [[TMP15:%.*]] = call i32 @__kmpc_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP4]], i32 1, i32 4, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66.omp_outlined.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)898// CHECK3-NEXT:    switch i32 [[TMP15]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [899// CHECK3-NEXT:      i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]900// CHECK3-NEXT:      i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]901// CHECK3-NEXT:    ]902// CHECK3:       .omp.reduction.case1:903// CHECK3-NEXT:    [[TMP16:%.*]] = load i32, ptr [[TMP0]], align 4904// CHECK3-NEXT:    [[TMP17:%.*]] = load i32, ptr [[SIVAR1]], align 4905// CHECK3-NEXT:    [[ADD5:%.*]] = add nsw i32 [[TMP16]], [[TMP17]]906// CHECK3-NEXT:    store i32 [[ADD5]], ptr [[TMP0]], align 4907// CHECK3-NEXT:    call void @__kmpc_end_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP4]], ptr @.gomp_critical_user_.reduction.var)908// CHECK3-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]909// CHECK3:       .omp.reduction.case2:910// CHECK3-NEXT:    [[TMP18:%.*]] = load i32, ptr [[SIVAR1]], align 4911// CHECK3-NEXT:    [[TMP19:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP18]] monotonic, align 4912// CHECK3-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]913// CHECK3:       .omp.reduction.default:914// CHECK3-NEXT:    ret void915//916//917// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66.omp_outlined.omp_outlined.omp.reduction.reduction_func918// CHECK3-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3:[0-9]+]] {919// CHECK3-NEXT:  entry:920// CHECK3-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 4921// CHECK3-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 4922// CHECK3-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 4923// CHECK3-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 4924// CHECK3-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 4925// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 4926// CHECK3-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i32 0, i32 0927// CHECK3-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 4928// CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i32 0, i32 0929// CHECK3-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 4930// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4931// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4932// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]933// CHECK3-NEXT:    store i32 [[ADD]], ptr [[TMP7]], align 4934// CHECK3-NEXT:    ret void935//936//937// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l66.omp_outlined.omp.reduction.reduction_func938// CHECK3-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3]] {939// CHECK3-NEXT:  entry:940// CHECK3-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 4941// CHECK3-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 4942// CHECK3-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 4943// CHECK3-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 4944// CHECK3-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 4945// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 4946// CHECK3-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i32 0, i32 0947// CHECK3-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 4948// CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i32 0, i32 0949// CHECK3-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 4950// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4951// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4952// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]953// CHECK3-NEXT:    store i32 [[ADD]], ptr [[TMP7]], align 4954// CHECK3-NEXT:    ret void955//956//957// CHECK3-LABEL: define {{[^@]+}}@_Z5tmainIiET_v958// CHECK3-SAME: () #[[ATTR5:[0-9]+]] comdat {959// CHECK3-NEXT:  entry:960// CHECK3-NEXT:    [[T_VAR:%.*]] = alloca i32, align 4961// CHECK3-NEXT:    [[VEC:%.*]] = alloca [2 x i32], align 4962// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4963// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4964// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4965// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 4966// CHECK3-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8967// CHECK3-NEXT:    store i32 0, ptr [[T_VAR]], align 4968// CHECK3-NEXT:    call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[VEC]], ptr align 4 @__const._Z5tmainIiET_v.vec, i32 8, i1 false)969// CHECK3-NEXT:    [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0970// CHECK3-NEXT:    store ptr [[T_VAR]], ptr [[TMP0]], align 4971// CHECK3-NEXT:    [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0972// CHECK3-NEXT:    store ptr [[T_VAR]], ptr [[TMP1]], align 4973// CHECK3-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0974// CHECK3-NEXT:    store ptr null, ptr [[TMP2]], align 4975// CHECK3-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0976// CHECK3-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0977// CHECK3-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0978// CHECK3-NEXT:    store i32 3, ptr [[TMP5]], align 4979// CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1980// CHECK3-NEXT:    store i32 1, ptr [[TMP6]], align 4981// CHECK3-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2982// CHECK3-NEXT:    store ptr [[TMP3]], ptr [[TMP7]], align 4983// CHECK3-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3984// CHECK3-NEXT:    store ptr [[TMP4]], ptr [[TMP8]], align 4985// CHECK3-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4986// CHECK3-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP9]], align 4987// CHECK3-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5988// CHECK3-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP10]], align 4989// CHECK3-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6990// CHECK3-NEXT:    store ptr null, ptr [[TMP11]], align 4991// CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7992// CHECK3-NEXT:    store ptr null, ptr [[TMP12]], align 4993// CHECK3-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8994// CHECK3-NEXT:    store i64 2, ptr [[TMP13]], align 8995// CHECK3-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9996// CHECK3-NEXT:    store i64 0, ptr [[TMP14]], align 8997// CHECK3-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10998// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4999// CHECK3-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111000// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP16]], align 41001// CHECK3-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121002// CHECK3-NEXT:    store i32 0, ptr [[TMP17]], align 41003// CHECK3-NEXT:    [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB4]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.region_id, ptr [[KERNEL_ARGS]])1004// CHECK3-NEXT:    [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 01005// CHECK3-NEXT:    br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1006// CHECK3:       omp_offload.failed:1007// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32(ptr [[T_VAR]]) #[[ATTR2]]1008// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT]]1009// CHECK3:       omp_offload.cont:1010// CHECK3-NEXT:    ret i32 01011//1012//1013// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l321014// CHECK3-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[T_VAR:%.*]]) #[[ATTR1]] {1015// CHECK3-NEXT:  entry:1016// CHECK3-NEXT:    [[T_VAR_ADDR:%.*]] = alloca ptr, align 41017// CHECK3-NEXT:    store ptr [[T_VAR]], ptr [[T_VAR_ADDR]], align 41018// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[T_VAR_ADDR]], align 41019// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB4]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined, ptr [[TMP0]])1020// CHECK3-NEXT:    ret void1021//1022//1023// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined1024// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[T_VAR:%.*]]) #[[ATTR1]] {1025// CHECK3-NEXT:  entry:1026// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41027// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41028// CHECK3-NEXT:    [[T_VAR_ADDR:%.*]] = alloca ptr, align 41029// CHECK3-NEXT:    [[T_VAR1:%.*]] = alloca i32, align 41030// CHECK3-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41031// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 41032// CHECK3-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41033// CHECK3-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41034// CHECK3-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41035// CHECK3-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41036// CHECK3-NEXT:    [[I:%.*]] = alloca i32, align 41037// CHECK3-NEXT:    [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 41038// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41039// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41040// CHECK3-NEXT:    store ptr [[T_VAR]], ptr [[T_VAR_ADDR]], align 41041// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[T_VAR_ADDR]], align 41042// CHECK3-NEXT:    store i32 0, ptr [[T_VAR1]], align 41043// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 41044// CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_COMB_UB]], align 41045// CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41046// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41047// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41048// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41049// CHECK3-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1050// CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41051// CHECK3-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 11052// CHECK3-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1053// CHECK3:       cond.true:1054// CHECK3-NEXT:    br label [[COND_END:%.*]]1055// CHECK3:       cond.false:1056// CHECK3-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41057// CHECK3-NEXT:    br label [[COND_END]]1058// CHECK3:       cond.end:1059// CHECK3-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]1060// CHECK3-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41061// CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41062// CHECK3-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 41063// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1064// CHECK3:       omp.inner.for.cond:1065// CHECK3-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41066// CHECK3-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41067// CHECK3-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]1068// CHECK3-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1069// CHECK3:       omp.inner.for.body:1070// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41071// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41072// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB4]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp_outlined, i32 [[TMP8]], i32 [[TMP9]], ptr [[T_VAR1]])1073// CHECK3-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1074// CHECK3:       omp.inner.for.inc:1075// CHECK3-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41076// CHECK3-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41077// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP10]], [[TMP11]]1078// CHECK3-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 41079// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND]]1080// CHECK3:       omp.inner.for.end:1081// CHECK3-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1082// CHECK3:       omp.loop.exit:1083// CHECK3-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])1084// CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 01085// CHECK3-NEXT:    store ptr [[T_VAR1]], ptr [[TMP12]], align 41086// CHECK3-NEXT:    [[TMP13:%.*]] = call i32 @__kmpc_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP2]], i32 1, i32 4, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)1087// CHECK3-NEXT:    switch i32 [[TMP13]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [1088// CHECK3-NEXT:      i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]1089// CHECK3-NEXT:      i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]1090// CHECK3-NEXT:    ]1091// CHECK3:       .omp.reduction.case1:1092// CHECK3-NEXT:    [[TMP14:%.*]] = load i32, ptr [[TMP0]], align 41093// CHECK3-NEXT:    [[TMP15:%.*]] = load i32, ptr [[T_VAR1]], align 41094// CHECK3-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP14]], [[TMP15]]1095// CHECK3-NEXT:    store i32 [[ADD3]], ptr [[TMP0]], align 41096// CHECK3-NEXT:    call void @__kmpc_end_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)1097// CHECK3-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]1098// CHECK3:       .omp.reduction.case2:1099// CHECK3-NEXT:    [[TMP16:%.*]] = load i32, ptr [[T_VAR1]], align 41100// CHECK3-NEXT:    [[TMP17:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP16]] monotonic, align 41101// CHECK3-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]1102// CHECK3:       .omp.reduction.default:1103// CHECK3-NEXT:    ret void1104//1105//1106// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp_outlined1107// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[DOTPREVIOUS_LB_:%.*]], i32 noundef [[DOTPREVIOUS_UB_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[T_VAR:%.*]]) #[[ATTR1]] {1108// CHECK3-NEXT:  entry:1109// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41110// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41111// CHECK3-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i32, align 41112// CHECK3-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i32, align 41113// CHECK3-NEXT:    [[T_VAR_ADDR:%.*]] = alloca ptr, align 41114// CHECK3-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41115// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 41116// CHECK3-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 41117// CHECK3-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 41118// CHECK3-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41119// CHECK3-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41120// CHECK3-NEXT:    [[T_VAR1:%.*]] = alloca i32, align 41121// CHECK3-NEXT:    [[I:%.*]] = alloca i32, align 41122// CHECK3-NEXT:    [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 41123// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41124// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41125// CHECK3-NEXT:    store i32 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 41126// CHECK3-NEXT:    store i32 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 41127// CHECK3-NEXT:    store ptr [[T_VAR]], ptr [[T_VAR_ADDR]], align 41128// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[T_VAR_ADDR]], align 41129// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 41130// CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 41131// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTPREVIOUS_LB__ADDR]], align 41132// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTPREVIOUS_UB__ADDR]], align 41133// CHECK3-NEXT:    store i32 [[TMP1]], ptr [[DOTOMP_LB]], align 41134// CHECK3-NEXT:    store i32 [[TMP2]], ptr [[DOTOMP_UB]], align 41135// CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41136// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41137// CHECK3-NEXT:    store i32 0, ptr [[T_VAR1]], align 41138// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41139// CHECK3-NEXT:    [[TMP4:%.*]] = load i32, ptr [[TMP3]], align 41140// CHECK3-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB2]], i32 [[TMP4]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1141// CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41142// CHECK3-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP5]], 11143// CHECK3-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1144// CHECK3:       cond.true:1145// CHECK3-NEXT:    br label [[COND_END:%.*]]1146// CHECK3:       cond.false:1147// CHECK3-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41148// CHECK3-NEXT:    br label [[COND_END]]1149// CHECK3:       cond.end:1150// CHECK3-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP6]], [[COND_FALSE]] ]1151// CHECK3-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 41152// CHECK3-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41153// CHECK3-NEXT:    store i32 [[TMP7]], ptr [[DOTOMP_IV]], align 41154// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1155// CHECK3:       omp.inner.for.cond:1156// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41157// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41158// CHECK3-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP8]], [[TMP9]]1159// CHECK3-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1160// CHECK3:       omp.inner.for.body:1161// CHECK3-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41162// CHECK3-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP10]], 11163// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1164// CHECK3-NEXT:    store i32 [[ADD]], ptr [[I]], align 41165// CHECK3-NEXT:    [[TMP11:%.*]] = load i32, ptr [[I]], align 41166// CHECK3-NEXT:    [[TMP12:%.*]] = load i32, ptr [[T_VAR1]], align 41167// CHECK3-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP12]], [[TMP11]]1168// CHECK3-NEXT:    store i32 [[ADD3]], ptr [[T_VAR1]], align 41169// CHECK3-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1170// CHECK3:       omp.body.continue:1171// CHECK3-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1172// CHECK3:       omp.inner.for.inc:1173// CHECK3-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41174// CHECK3-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP13]], 11175// CHECK3-NEXT:    store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 41176// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND]]1177// CHECK3:       omp.inner.for.end:1178// CHECK3-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1179// CHECK3:       omp.loop.exit:1180// CHECK3-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP4]])1181// CHECK3-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 01182// CHECK3-NEXT:    store ptr [[T_VAR1]], ptr [[TMP14]], align 41183// CHECK3-NEXT:    [[TMP15:%.*]] = call i32 @__kmpc_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP4]], i32 1, i32 4, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)1184// CHECK3-NEXT:    switch i32 [[TMP15]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [1185// CHECK3-NEXT:      i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]1186// CHECK3-NEXT:      i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]1187// CHECK3-NEXT:    ]1188// CHECK3:       .omp.reduction.case1:1189// CHECK3-NEXT:    [[TMP16:%.*]] = load i32, ptr [[TMP0]], align 41190// CHECK3-NEXT:    [[TMP17:%.*]] = load i32, ptr [[T_VAR1]], align 41191// CHECK3-NEXT:    [[ADD5:%.*]] = add nsw i32 [[TMP16]], [[TMP17]]1192// CHECK3-NEXT:    store i32 [[ADD5]], ptr [[TMP0]], align 41193// CHECK3-NEXT:    call void @__kmpc_end_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP4]], ptr @.gomp_critical_user_.reduction.var)1194// CHECK3-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]1195// CHECK3:       .omp.reduction.case2:1196// CHECK3-NEXT:    [[TMP18:%.*]] = load i32, ptr [[T_VAR1]], align 41197// CHECK3-NEXT:    [[TMP19:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP18]] monotonic, align 41198// CHECK3-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]1199// CHECK3:       .omp.reduction.default:1200// CHECK3-NEXT:    ret void1201//1202//1203// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp_outlined.omp.reduction.reduction_func1204// CHECK3-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3]] {1205// CHECK3-NEXT:  entry:1206// CHECK3-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 41207// CHECK3-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 41208// CHECK3-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 41209// CHECK3-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 41210// CHECK3-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 41211// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 41212// CHECK3-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i32 0, i32 01213// CHECK3-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 41214// CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i32 0, i32 01215// CHECK3-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 41216// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41217// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 41218// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]1219// CHECK3-NEXT:    store i32 [[ADD]], ptr [[TMP7]], align 41220// CHECK3-NEXT:    ret void1221//1222//1223// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp.reduction.reduction_func1224// CHECK3-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3]] {1225// CHECK3-NEXT:  entry:1226// CHECK3-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 41227// CHECK3-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 41228// CHECK3-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 41229// CHECK3-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 41230// CHECK3-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 41231// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 41232// CHECK3-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i32 0, i32 01233// CHECK3-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 41234// CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i32 0, i32 01235// CHECK3-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 41236// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41237// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 41238// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]1239// CHECK3-NEXT:    store i32 [[ADD]], ptr [[TMP7]], align 41240// CHECK3-NEXT:    ret void1241//1242//1243// CHECK5-LABEL: define {{[^@]+}}@main1244// CHECK5-SAME: () #[[ATTR0:[0-9]+]] {1245// CHECK5-NEXT:  entry:1246// CHECK5-NEXT:    [[RETVAL:%.*]] = alloca i32, align 41247// CHECK5-NEXT:    [[REF_TMP:%.*]] = alloca [[CLASS_ANON:%.*]], align 11248// CHECK5-NEXT:    store i32 0, ptr [[RETVAL]], align 41249// CHECK5-NEXT:    call void @"_ZZ4mainENK3$_0clEv"(ptr noundef nonnull align 1 dereferenceable(1) [[REF_TMP]])1250// CHECK5-NEXT:    ret i32 01251//1252//1253// CHECK5-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l441254// CHECK5-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR2:[0-9]+]] {1255// CHECK5-NEXT:  entry:1256// CHECK5-NEXT:    [[SIVAR_ADDR:%.*]] = alloca ptr, align 81257// CHECK5-NEXT:    store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 81258// CHECK5-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 81259// CHECK5-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB4:[0-9]+]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l44.omp_outlined, ptr [[TMP0]])1260// CHECK5-NEXT:    ret void1261//1262//1263// CHECK5-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l44.omp_outlined1264// CHECK5-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR2]] {1265// CHECK5-NEXT:  entry:1266// CHECK5-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81267// CHECK5-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81268// CHECK5-NEXT:    [[SIVAR_ADDR:%.*]] = alloca ptr, align 81269// CHECK5-NEXT:    [[SIVAR1:%.*]] = alloca i32, align 41270// CHECK5-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41271// CHECK5-NEXT:    [[TMP:%.*]] = alloca i32, align 41272// CHECK5-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41273// CHECK5-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41274// CHECK5-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41275// CHECK5-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41276// CHECK5-NEXT:    [[I:%.*]] = alloca i32, align 41277// CHECK5-NEXT:    [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 81278// CHECK5-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81279// CHECK5-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81280// CHECK5-NEXT:    store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 81281// CHECK5-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 81282// CHECK5-NEXT:    store i32 0, ptr [[SIVAR1]], align 41283// CHECK5-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 41284// CHECK5-NEXT:    store i32 1, ptr [[DOTOMP_COMB_UB]], align 41285// CHECK5-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41286// CHECK5-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41287// CHECK5-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81288// CHECK5-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41289// CHECK5-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1290// CHECK5-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41291// CHECK5-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 11292// CHECK5-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1293// CHECK5:       cond.true:1294// CHECK5-NEXT:    br label [[COND_END:%.*]]1295// CHECK5:       cond.false:1296// CHECK5-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41297// CHECK5-NEXT:    br label [[COND_END]]1298// CHECK5:       cond.end:1299// CHECK5-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]1300// CHECK5-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41301// CHECK5-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41302// CHECK5-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 41303// CHECK5-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1304// CHECK5:       omp.inner.for.cond:1305// CHECK5-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41306// CHECK5-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41307// CHECK5-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]1308// CHECK5-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1309// CHECK5:       omp.inner.for.body:1310// CHECK5-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41311// CHECK5-NEXT:    [[TMP9:%.*]] = zext i32 [[TMP8]] to i641312// CHECK5-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41313// CHECK5-NEXT:    [[TMP11:%.*]] = zext i32 [[TMP10]] to i641314// CHECK5-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB4]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l44.omp_outlined.omp_outlined, i64 [[TMP9]], i64 [[TMP11]], ptr [[SIVAR1]])1315// CHECK5-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1316// CHECK5:       omp.inner.for.inc:1317// CHECK5-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41318// CHECK5-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41319// CHECK5-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP12]], [[TMP13]]1320// CHECK5-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 41321// CHECK5-NEXT:    br label [[OMP_INNER_FOR_COND]]1322// CHECK5:       omp.inner.for.end:1323// CHECK5-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1324// CHECK5:       omp.loop.exit:1325// CHECK5-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])1326// CHECK5-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 01327// CHECK5-NEXT:    store ptr [[SIVAR1]], ptr [[TMP14]], align 81328// CHECK5-NEXT:    [[TMP15:%.*]] = call i32 @__kmpc_reduce_nowait(ptr @[[GLOB3:[0-9]+]], i32 [[TMP2]], i32 1, i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l44.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)1329// CHECK5-NEXT:    switch i32 [[TMP15]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [1330// CHECK5-NEXT:      i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]1331// CHECK5-NEXT:      i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]1332// CHECK5-NEXT:    ]1333// CHECK5:       .omp.reduction.case1:1334// CHECK5-NEXT:    [[TMP16:%.*]] = load i32, ptr [[TMP0]], align 41335// CHECK5-NEXT:    [[TMP17:%.*]] = load i32, ptr [[SIVAR1]], align 41336// CHECK5-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP16]], [[TMP17]]1337// CHECK5-NEXT:    store i32 [[ADD3]], ptr [[TMP0]], align 41338// CHECK5-NEXT:    call void @__kmpc_end_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)1339// CHECK5-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]1340// CHECK5:       .omp.reduction.case2:1341// CHECK5-NEXT:    [[TMP18:%.*]] = load i32, ptr [[SIVAR1]], align 41342// CHECK5-NEXT:    [[TMP19:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP18]] monotonic, align 41343// CHECK5-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]1344// CHECK5:       .omp.reduction.default:1345// CHECK5-NEXT:    ret void1346//1347//1348// CHECK5-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l44.omp_outlined.omp_outlined1349// CHECK5-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR2]] {1350// CHECK5-NEXT:  entry:1351// CHECK5-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81352// CHECK5-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81353// CHECK5-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 81354// CHECK5-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 81355// CHECK5-NEXT:    [[SIVAR_ADDR:%.*]] = alloca ptr, align 81356// CHECK5-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41357// CHECK5-NEXT:    [[TMP:%.*]] = alloca i32, align 41358// CHECK5-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 41359// CHECK5-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 41360// CHECK5-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41361// CHECK5-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41362// CHECK5-NEXT:    [[SIVAR2:%.*]] = alloca i32, align 41363// CHECK5-NEXT:    [[I:%.*]] = alloca i32, align 41364// CHECK5-NEXT:    [[REF_TMP:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 81365// CHECK5-NEXT:    [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 81366// CHECK5-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81367// CHECK5-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81368// CHECK5-NEXT:    store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 81369// CHECK5-NEXT:    store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 81370// CHECK5-NEXT:    store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 81371// CHECK5-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 81372// CHECK5-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 41373// CHECK5-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 41374// CHECK5-NEXT:    [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 81375// CHECK5-NEXT:    [[CONV:%.*]] = trunc i64 [[TMP1]] to i321376// CHECK5-NEXT:    [[TMP2:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 81377// CHECK5-NEXT:    [[CONV1:%.*]] = trunc i64 [[TMP2]] to i321378// CHECK5-NEXT:    store i32 [[CONV]], ptr [[DOTOMP_LB]], align 41379// CHECK5-NEXT:    store i32 [[CONV1]], ptr [[DOTOMP_UB]], align 41380// CHECK5-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41381// CHECK5-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41382// CHECK5-NEXT:    store i32 0, ptr [[SIVAR2]], align 41383// CHECK5-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81384// CHECK5-NEXT:    [[TMP4:%.*]] = load i32, ptr [[TMP3]], align 41385// CHECK5-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP4]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1386// CHECK5-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41387// CHECK5-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP5]], 11388// CHECK5-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1389// CHECK5:       cond.true:1390// CHECK5-NEXT:    br label [[COND_END:%.*]]1391// CHECK5:       cond.false:1392// CHECK5-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41393// CHECK5-NEXT:    br label [[COND_END]]1394// CHECK5:       cond.end:1395// CHECK5-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP6]], [[COND_FALSE]] ]1396// CHECK5-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 41397// CHECK5-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41398// CHECK5-NEXT:    store i32 [[TMP7]], ptr [[DOTOMP_IV]], align 41399// CHECK5-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1400// CHECK5:       omp.inner.for.cond:1401// CHECK5-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41402// CHECK5-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41403// CHECK5-NEXT:    [[CMP3:%.*]] = icmp sle i32 [[TMP8]], [[TMP9]]1404// CHECK5-NEXT:    br i1 [[CMP3]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1405// CHECK5:       omp.inner.for.body:1406// CHECK5-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41407// CHECK5-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP10]], 11408// CHECK5-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1409// CHECK5-NEXT:    store i32 [[ADD]], ptr [[I]], align 41410// CHECK5-NEXT:    [[TMP11:%.*]] = load i32, ptr [[I]], align 41411// CHECK5-NEXT:    [[TMP12:%.*]] = load i32, ptr [[SIVAR2]], align 41412// CHECK5-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP12]], [[TMP11]]1413// CHECK5-NEXT:    store i32 [[ADD4]], ptr [[SIVAR2]], align 41414// CHECK5-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 01415// CHECK5-NEXT:    store ptr [[SIVAR2]], ptr [[TMP13]], align 81416// CHECK5-NEXT:    call void @"_ZZZ4mainENK3$_0clEvENKUlvE_clEv"(ptr noundef nonnull align 8 dereferenceable(8) [[REF_TMP]])1417// CHECK5-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1418// CHECK5:       omp.body.continue:1419// CHECK5-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1420// CHECK5:       omp.inner.for.inc:1421// CHECK5-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41422// CHECK5-NEXT:    [[ADD5:%.*]] = add nsw i32 [[TMP14]], 11423// CHECK5-NEXT:    store i32 [[ADD5]], ptr [[DOTOMP_IV]], align 41424// CHECK5-NEXT:    br label [[OMP_INNER_FOR_COND]]1425// CHECK5:       omp.inner.for.end:1426// CHECK5-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1427// CHECK5:       omp.loop.exit:1428// CHECK5-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP4]])1429// CHECK5-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 01430// CHECK5-NEXT:    store ptr [[SIVAR2]], ptr [[TMP15]], align 81431// CHECK5-NEXT:    [[TMP16:%.*]] = call i32 @__kmpc_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP4]], i32 1, i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l44.omp_outlined.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)1432// CHECK5-NEXT:    switch i32 [[TMP16]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [1433// CHECK5-NEXT:      i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]1434// CHECK5-NEXT:      i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]1435// CHECK5-NEXT:    ]1436// CHECK5:       .omp.reduction.case1:1437// CHECK5-NEXT:    [[TMP17:%.*]] = load i32, ptr [[TMP0]], align 41438// CHECK5-NEXT:    [[TMP18:%.*]] = load i32, ptr [[SIVAR2]], align 41439// CHECK5-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP17]], [[TMP18]]1440// CHECK5-NEXT:    store i32 [[ADD6]], ptr [[TMP0]], align 41441// CHECK5-NEXT:    call void @__kmpc_end_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP4]], ptr @.gomp_critical_user_.reduction.var)1442// CHECK5-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]1443// CHECK5:       .omp.reduction.case2:1444// CHECK5-NEXT:    [[TMP19:%.*]] = load i32, ptr [[SIVAR2]], align 41445// CHECK5-NEXT:    [[TMP20:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP19]] monotonic, align 41446// CHECK5-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]1447// CHECK5:       .omp.reduction.default:1448// CHECK5-NEXT:    ret void1449//1450//1451// CHECK5-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l44.omp_outlined.omp_outlined.omp.reduction.reduction_func1452// CHECK5-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR4:[0-9]+]] {1453// CHECK5-NEXT:  entry:1454// CHECK5-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 81455// CHECK5-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 81456// CHECK5-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 81457// CHECK5-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 81458// CHECK5-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 81459// CHECK5-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 81460// CHECK5-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i64 0, i64 01461// CHECK5-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 81462// CHECK5-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i64 0, i64 01463// CHECK5-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 81464// CHECK5-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41465// CHECK5-NEXT:    [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 41466// CHECK5-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]1467// CHECK5-NEXT:    store i32 [[ADD]], ptr [[TMP7]], align 41468// CHECK5-NEXT:    ret void1469//1470//1471// CHECK5-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l44.omp_outlined.omp.reduction.reduction_func1472// CHECK5-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR4]] {1473// CHECK5-NEXT:  entry:1474// CHECK5-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 81475// CHECK5-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 81476// CHECK5-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 81477// CHECK5-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 81478// CHECK5-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 81479// CHECK5-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 81480// CHECK5-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i64 0, i64 01481// CHECK5-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 81482// CHECK5-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i64 0, i64 01483// CHECK5-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 81484// CHECK5-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41485// CHECK5-NEXT:    [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 41486// CHECK5-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]1487// CHECK5-NEXT:    store i32 [[ADD]], ptr [[TMP7]], align 41488// CHECK5-NEXT:    ret void1489//1490