1174 lines · cpp
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _2// 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 -DCHECK -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK510// 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 %s11// 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 --check-prefix=CHECK512// 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 --check-prefix=CHECK713// 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 %s14// 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 --check-prefix=CHECK715 16// 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=CHECK917// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s18// 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=CHECK919 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 --check-prefix=CHECK1121// 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 --check-prefix=CHECK1123 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 target33#pragma omp teams distribute simd reduction(+: t_var)34 for (int i = 0; i < 2; ++i) {35 t_var += (T) i;36 }37 return T();38}39 40int main() {41 static int sivar;42#ifdef LAMBDA43 44 [&]() {45#pragma omp target46#pragma omp teams distribute simd reduction(+: sivar)47 for (int i = 0; i < 2; ++i) {48 49 // Skip global and bound tid vars50 51 52 sivar += i;53 54 [&]() {55 56 sivar += 4;57 58 }();59 }60 }();61 return 0;62#else63#pragma omp target64#pragma omp teams distribute simd reduction(+: sivar)65 for (int i = 0; i < 2; ++i) {66 sivar += i;67 }68 return tmain<int>();69#endif70}71 72 73 74 75// Skip global and bound tid vars76 77 78 79 80 81// Skip global and bound tid vars82 83 84#endif85// CHECK1-LABEL: define {{[^@]+}}@main86// CHECK1-SAME: () #[[ATTR0:[0-9]+]] {87// CHECK1-NEXT: entry:88// CHECK1-NEXT: [[RETVAL:%.*]] = alloca i32, align 489// CHECK1-NEXT: [[SIVAR_CASTED:%.*]] = alloca i64, align 890// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 891// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 892// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 893// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 494// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 895// CHECK1-NEXT: store i32 0, ptr [[RETVAL]], align 496// CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr @_ZZ4mainE5sivar, align 497// CHECK1-NEXT: store i32 [[TMP0]], ptr [[SIVAR_CASTED]], align 498// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[SIVAR_CASTED]], align 899// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0100// CHECK1-NEXT: store i64 [[TMP1]], ptr [[TMP2]], align 8101// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0102// CHECK1-NEXT: store i64 [[TMP1]], ptr [[TMP3]], align 8103// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0104// CHECK1-NEXT: store ptr null, ptr [[TMP4]], align 8105// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0106// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0107// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0108// CHECK1-NEXT: store i32 3, ptr [[TMP7]], align 4109// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1110// CHECK1-NEXT: store i32 1, ptr [[TMP8]], align 4111// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2112// CHECK1-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 8113// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3114// CHECK1-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 8115// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4116// CHECK1-NEXT: store ptr @.offload_sizes, ptr [[TMP11]], align 8117// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5118// CHECK1-NEXT: store ptr @.offload_maptypes, ptr [[TMP12]], align 8119// CHECK1-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6120// CHECK1-NEXT: store ptr null, ptr [[TMP13]], align 8121// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7122// CHECK1-NEXT: store ptr null, ptr [[TMP14]], align 8123// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8124// CHECK1-NEXT: store i64 2, ptr [[TMP15]], align 8125// CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9126// CHECK1-NEXT: store i64 0, ptr [[TMP16]], align 8127// CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10128// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4129// CHECK1-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11130// CHECK1-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP18]], align 4131// CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12132// CHECK1-NEXT: store i32 0, ptr [[TMP19]], align 4133// CHECK1-NEXT: [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3:[0-9]+]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l63.region_id, ptr [[KERNEL_ARGS]])134// CHECK1-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0135// CHECK1-NEXT: br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]136// CHECK1: omp_offload.failed:137// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l63(i64 [[TMP1]]) #[[ATTR2:[0-9]+]]138// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]139// CHECK1: omp_offload.cont:140// CHECK1-NEXT: [[CALL:%.*]] = call noundef signext i32 @_Z5tmainIiET_v()141// CHECK1-NEXT: ret i32 [[CALL]]142//143//144// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l63145// CHECK1-SAME: (i64 noundef [[SIVAR:%.*]]) #[[ATTR1:[0-9]+]] {146// CHECK1-NEXT: entry:147// CHECK1-NEXT: [[SIVAR_ADDR:%.*]] = alloca i64, align 8148// CHECK1-NEXT: store i64 [[SIVAR]], ptr [[SIVAR_ADDR]], align 8149// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l63.omp_outlined, ptr [[SIVAR_ADDR]])150// CHECK1-NEXT: ret void151//152//153// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l63.omp_outlined154// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR1]] {155// CHECK1-NEXT: entry:156// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8157// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8158// CHECK1-NEXT: [[SIVAR_ADDR:%.*]] = alloca ptr, align 8159// CHECK1-NEXT: [[SIVAR1:%.*]] = alloca i32, align 4160// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4161// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4162// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4163// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4164// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4165// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4166// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4167// CHECK1-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 8168// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8169// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8170// CHECK1-NEXT: store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 8171// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 8172// CHECK1-NEXT: store i32 0, ptr [[SIVAR1]], align 4173// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4174// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 4175// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4176// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4177// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8178// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4179// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)180// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4181// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1182// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]183// CHECK1: cond.true:184// CHECK1-NEXT: br label [[COND_END:%.*]]185// CHECK1: cond.false:186// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4187// CHECK1-NEXT: br label [[COND_END]]188// CHECK1: cond.end:189// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]190// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4191// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4192// CHECK1-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4193// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]194// CHECK1: omp.inner.for.cond:195// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP5:![0-9]+]]196// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP5]]197// CHECK1-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]198// CHECK1-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]199// CHECK1: omp.inner.for.body:200// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP5]]201// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1202// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]203// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP5]]204// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP5]]205// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[SIVAR1]], align 4, !llvm.access.group [[ACC_GRP5]]206// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]207// CHECK1-NEXT: store i32 [[ADD3]], ptr [[SIVAR1]], align 4, !llvm.access.group [[ACC_GRP5]]208// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]209// CHECK1: omp.body.continue:210// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]211// CHECK1: omp.inner.for.inc:212// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP5]]213// CHECK1-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP11]], 1214// CHECK1-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP5]]215// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP6:![0-9]+]]216// CHECK1: omp.inner.for.end:217// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]218// CHECK1: omp.loop.exit:219// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])220// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 4221// CHECK1-NEXT: [[TMP13:%.*]] = icmp ne i32 [[TMP12]], 0222// CHECK1-NEXT: br i1 [[TMP13]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]223// CHECK1: .omp.final.then:224// CHECK1-NEXT: store i32 2, ptr [[I]], align 4225// CHECK1-NEXT: br label [[DOTOMP_FINAL_DONE]]226// CHECK1: .omp.final.done:227// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0228// CHECK1-NEXT: store ptr [[SIVAR1]], ptr [[TMP14]], align 8229// CHECK1-NEXT: [[TMP15:%.*]] = call i32 @__kmpc_reduce(ptr @[[GLOB2:[0-9]+]], i32 [[TMP2]], i32 1, i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l63.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)230// CHECK1-NEXT: switch i32 [[TMP15]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [231// CHECK1-NEXT: i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]232// CHECK1-NEXT: i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]233// CHECK1-NEXT: ]234// CHECK1: .omp.reduction.case1:235// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[TMP0]], align 4236// CHECK1-NEXT: [[TMP17:%.*]] = load i32, ptr [[SIVAR1]], align 4237// CHECK1-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP16]], [[TMP17]]238// CHECK1-NEXT: store i32 [[ADD5]], ptr [[TMP0]], align 4239// CHECK1-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)240// CHECK1-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]241// CHECK1: .omp.reduction.case2:242// CHECK1-NEXT: [[TMP18:%.*]] = load i32, ptr [[SIVAR1]], align 4243// CHECK1-NEXT: [[TMP19:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP18]] monotonic, align 4244// CHECK1-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)245// CHECK1-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]246// CHECK1: .omp.reduction.default:247// CHECK1-NEXT: ret void248//249//250// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l63.omp_outlined.omp.reduction.reduction_func251// CHECK1-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3:[0-9]+]] {252// CHECK1-NEXT: entry:253// CHECK1-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 8254// CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 8255// CHECK1-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 8256// CHECK1-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 8257// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 8258// CHECK1-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 8259// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i64 0, i64 0260// CHECK1-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 8261// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i64 0, i64 0262// CHECK1-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 8263// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4264// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4265// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]266// CHECK1-NEXT: store i32 [[ADD]], ptr [[TMP7]], align 4267// CHECK1-NEXT: ret void268//269//270// CHECK1-LABEL: define {{[^@]+}}@_Z5tmainIiET_v271// CHECK1-SAME: () #[[ATTR5:[0-9]+]] comdat {272// CHECK1-NEXT: entry:273// CHECK1-NEXT: [[T_VAR:%.*]] = alloca i32, align 4274// CHECK1-NEXT: [[VEC:%.*]] = alloca [2 x i32], align 4275// CHECK1-NEXT: [[T_VAR_CASTED:%.*]] = alloca i64, align 8276// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8277// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8278// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8279// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4280// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8281// CHECK1-NEXT: store i32 0, ptr [[T_VAR]], align 4282// CHECK1-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[VEC]], ptr align 4 @__const._Z5tmainIiET_v.vec, i64 8, i1 false)283// CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr [[T_VAR]], align 4284// CHECK1-NEXT: store i32 [[TMP0]], ptr [[T_VAR_CASTED]], align 4285// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[T_VAR_CASTED]], align 8286// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0287// CHECK1-NEXT: store i64 [[TMP1]], ptr [[TMP2]], align 8288// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0289// CHECK1-NEXT: store i64 [[TMP1]], ptr [[TMP3]], align 8290// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0291// CHECK1-NEXT: store ptr null, ptr [[TMP4]], align 8292// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0293// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0294// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0295// CHECK1-NEXT: store i32 3, ptr [[TMP7]], align 4296// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1297// CHECK1-NEXT: store i32 1, ptr [[TMP8]], align 4298// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2299// CHECK1-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 8300// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3301// CHECK1-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 8302// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4303// CHECK1-NEXT: store ptr @.offload_sizes.1, ptr [[TMP11]], align 8304// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5305// CHECK1-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP12]], align 8306// CHECK1-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6307// CHECK1-NEXT: store ptr null, ptr [[TMP13]], align 8308// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7309// CHECK1-NEXT: store ptr null, ptr [[TMP14]], align 8310// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8311// CHECK1-NEXT: store i64 2, ptr [[TMP15]], align 8312// CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9313// CHECK1-NEXT: store i64 0, ptr [[TMP16]], align 8314// CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10315// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4316// CHECK1-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11317// CHECK1-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP18]], align 4318// CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12319// CHECK1-NEXT: store i32 0, ptr [[TMP19]], align 4320// CHECK1-NEXT: [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.region_id, ptr [[KERNEL_ARGS]])321// CHECK1-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0322// CHECK1-NEXT: br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]323// CHECK1: omp_offload.failed:324// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32(i64 [[TMP1]]) #[[ATTR2]]325// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]326// CHECK1: omp_offload.cont:327// CHECK1-NEXT: ret i32 0328//329//330// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32331// CHECK1-SAME: (i64 noundef [[T_VAR:%.*]]) #[[ATTR1]] {332// CHECK1-NEXT: entry:333// CHECK1-NEXT: [[T_VAR_ADDR:%.*]] = alloca i64, align 8334// CHECK1-NEXT: store i64 [[T_VAR]], ptr [[T_VAR_ADDR]], align 8335// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined, ptr [[T_VAR_ADDR]])336// CHECK1-NEXT: ret void337//338//339// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined340// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[T_VAR:%.*]]) #[[ATTR1]] {341// CHECK1-NEXT: entry:342// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8343// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8344// CHECK1-NEXT: [[T_VAR_ADDR:%.*]] = alloca ptr, align 8345// CHECK1-NEXT: [[T_VAR1:%.*]] = alloca i32, align 4346// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4347// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4348// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4349// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4350// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4351// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4352// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4353// CHECK1-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 8354// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8355// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8356// CHECK1-NEXT: store ptr [[T_VAR]], ptr [[T_VAR_ADDR]], align 8357// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[T_VAR_ADDR]], align 8358// CHECK1-NEXT: store i32 0, ptr [[T_VAR1]], align 4359// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4360// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 4361// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4362// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4363// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8364// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4365// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)366// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4367// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1368// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]369// CHECK1: cond.true:370// CHECK1-NEXT: br label [[COND_END:%.*]]371// CHECK1: cond.false:372// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4373// CHECK1-NEXT: br label [[COND_END]]374// CHECK1: cond.end:375// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]376// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4377// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4378// CHECK1-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4379// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]380// CHECK1: omp.inner.for.cond:381// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11:![0-9]+]]382// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP11]]383// CHECK1-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]384// CHECK1-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]385// CHECK1: omp.inner.for.body:386// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]387// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1388// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]389// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP11]]390// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP11]]391// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP11]]392// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]393// CHECK1-NEXT: store i32 [[ADD3]], ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP11]]394// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]395// CHECK1: omp.body.continue:396// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]397// CHECK1: omp.inner.for.inc:398// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]399// CHECK1-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP11]], 1400// CHECK1-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]401// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP12:![0-9]+]]402// CHECK1: omp.inner.for.end:403// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]404// CHECK1: omp.loop.exit:405// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])406// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 4407// CHECK1-NEXT: [[TMP13:%.*]] = icmp ne i32 [[TMP12]], 0408// CHECK1-NEXT: br i1 [[TMP13]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]409// CHECK1: .omp.final.then:410// CHECK1-NEXT: store i32 2, ptr [[I]], align 4411// CHECK1-NEXT: br label [[DOTOMP_FINAL_DONE]]412// CHECK1: .omp.final.done:413// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0414// CHECK1-NEXT: store ptr [[T_VAR1]], ptr [[TMP14]], align 8415// CHECK1-NEXT: [[TMP15:%.*]] = call i32 @__kmpc_reduce(ptr @[[GLOB2]], 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)416// CHECK1-NEXT: switch i32 [[TMP15]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [417// CHECK1-NEXT: i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]418// CHECK1-NEXT: i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]419// CHECK1-NEXT: ]420// CHECK1: .omp.reduction.case1:421// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[TMP0]], align 4422// CHECK1-NEXT: [[TMP17:%.*]] = load i32, ptr [[T_VAR1]], align 4423// CHECK1-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP16]], [[TMP17]]424// CHECK1-NEXT: store i32 [[ADD5]], ptr [[TMP0]], align 4425// CHECK1-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)426// CHECK1-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]427// CHECK1: .omp.reduction.case2:428// CHECK1-NEXT: [[TMP18:%.*]] = load i32, ptr [[T_VAR1]], align 4429// CHECK1-NEXT: [[TMP19:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP18]] monotonic, align 4430// CHECK1-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)431// CHECK1-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]432// CHECK1: .omp.reduction.default:433// CHECK1-NEXT: ret void434//435//436// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp.reduction.reduction_func437// CHECK1-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3]] {438// CHECK1-NEXT: entry:439// CHECK1-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 8440// CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 8441// CHECK1-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 8442// CHECK1-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 8443// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 8444// CHECK1-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 8445// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i64 0, i64 0446// CHECK1-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 8447// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i64 0, i64 0448// CHECK1-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 8449// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4450// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4451// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]452// CHECK1-NEXT: store i32 [[ADD]], ptr [[TMP7]], align 4453// CHECK1-NEXT: ret void454//455//456// CHECK3-LABEL: define {{[^@]+}}@main457// CHECK3-SAME: () #[[ATTR0:[0-9]+]] {458// CHECK3-NEXT: entry:459// CHECK3-NEXT: [[RETVAL:%.*]] = alloca i32, align 4460// CHECK3-NEXT: [[SIVAR_CASTED:%.*]] = alloca i32, align 4461// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4462// CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4463// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4464// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4465// CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8466// CHECK3-NEXT: store i32 0, ptr [[RETVAL]], align 4467// CHECK3-NEXT: [[TMP0:%.*]] = load i32, ptr @_ZZ4mainE5sivar, align 4468// CHECK3-NEXT: store i32 [[TMP0]], ptr [[SIVAR_CASTED]], align 4469// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[SIVAR_CASTED]], align 4470// CHECK3-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0471// CHECK3-NEXT: store i32 [[TMP1]], ptr [[TMP2]], align 4472// CHECK3-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0473// CHECK3-NEXT: store i32 [[TMP1]], ptr [[TMP3]], align 4474// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0475// CHECK3-NEXT: store ptr null, ptr [[TMP4]], align 4476// CHECK3-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0477// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0478// CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0479// CHECK3-NEXT: store i32 3, ptr [[TMP7]], align 4480// CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1481// CHECK3-NEXT: store i32 1, ptr [[TMP8]], align 4482// CHECK3-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2483// CHECK3-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 4484// CHECK3-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3485// CHECK3-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 4486// CHECK3-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4487// CHECK3-NEXT: store ptr @.offload_sizes, ptr [[TMP11]], align 4488// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5489// CHECK3-NEXT: store ptr @.offload_maptypes, ptr [[TMP12]], align 4490// CHECK3-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6491// CHECK3-NEXT: store ptr null, ptr [[TMP13]], align 4492// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7493// CHECK3-NEXT: store ptr null, ptr [[TMP14]], align 4494// CHECK3-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8495// CHECK3-NEXT: store i64 2, ptr [[TMP15]], align 8496// CHECK3-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9497// CHECK3-NEXT: store i64 0, ptr [[TMP16]], align 8498// CHECK3-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10499// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4500// CHECK3-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11501// CHECK3-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP18]], align 4502// CHECK3-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12503// CHECK3-NEXT: store i32 0, ptr [[TMP19]], align 4504// CHECK3-NEXT: [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3:[0-9]+]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l63.region_id, ptr [[KERNEL_ARGS]])505// CHECK3-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0506// CHECK3-NEXT: br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]507// CHECK3: omp_offload.failed:508// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l63(i32 [[TMP1]]) #[[ATTR2:[0-9]+]]509// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]]510// CHECK3: omp_offload.cont:511// CHECK3-NEXT: [[CALL:%.*]] = call noundef i32 @_Z5tmainIiET_v()512// CHECK3-NEXT: ret i32 [[CALL]]513//514//515// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l63516// CHECK3-SAME: (i32 noundef [[SIVAR:%.*]]) #[[ATTR1:[0-9]+]] {517// CHECK3-NEXT: entry:518// CHECK3-NEXT: [[SIVAR_ADDR:%.*]] = alloca i32, align 4519// CHECK3-NEXT: store i32 [[SIVAR]], ptr [[SIVAR_ADDR]], align 4520// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l63.omp_outlined, ptr [[SIVAR_ADDR]])521// CHECK3-NEXT: ret void522//523//524// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l63.omp_outlined525// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR1]] {526// CHECK3-NEXT: entry:527// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4528// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4529// CHECK3-NEXT: [[SIVAR_ADDR:%.*]] = alloca ptr, align 4530// CHECK3-NEXT: [[SIVAR1:%.*]] = alloca i32, align 4531// CHECK3-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4532// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4533// CHECK3-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4534// CHECK3-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4535// CHECK3-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4536// CHECK3-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4537// CHECK3-NEXT: [[I:%.*]] = alloca i32, align 4538// CHECK3-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 4539// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4540// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4541// CHECK3-NEXT: store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 4542// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 4543// CHECK3-NEXT: store i32 0, ptr [[SIVAR1]], align 4544// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4545// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 4546// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4547// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4548// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4549// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4550// CHECK3-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)551// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4552// CHECK3-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1553// CHECK3-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]554// CHECK3: cond.true:555// CHECK3-NEXT: br label [[COND_END:%.*]]556// CHECK3: cond.false:557// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4558// CHECK3-NEXT: br label [[COND_END]]559// CHECK3: cond.end:560// CHECK3-NEXT: [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]561// CHECK3-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4562// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4563// CHECK3-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4564// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]565// CHECK3: omp.inner.for.cond:566// CHECK3-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6:![0-9]+]]567// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP6]]568// CHECK3-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]569// CHECK3-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]570// CHECK3: omp.inner.for.body:571// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6]]572// CHECK3-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1573// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]574// CHECK3-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP6]]575// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP6]]576// CHECK3-NEXT: [[TMP10:%.*]] = load i32, ptr [[SIVAR1]], align 4, !llvm.access.group [[ACC_GRP6]]577// CHECK3-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]578// CHECK3-NEXT: store i32 [[ADD3]], ptr [[SIVAR1]], align 4, !llvm.access.group [[ACC_GRP6]]579// CHECK3-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]580// CHECK3: omp.body.continue:581// CHECK3-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]582// CHECK3: omp.inner.for.inc:583// CHECK3-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6]]584// CHECK3-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP11]], 1585// CHECK3-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6]]586// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP7:![0-9]+]]587// CHECK3: omp.inner.for.end:588// CHECK3-NEXT: br label [[OMP_LOOP_EXIT:%.*]]589// CHECK3: omp.loop.exit:590// CHECK3-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])591// CHECK3-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 4592// CHECK3-NEXT: [[TMP13:%.*]] = icmp ne i32 [[TMP12]], 0593// CHECK3-NEXT: br i1 [[TMP13]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]594// CHECK3: .omp.final.then:595// CHECK3-NEXT: store i32 2, ptr [[I]], align 4596// CHECK3-NEXT: br label [[DOTOMP_FINAL_DONE]]597// CHECK3: .omp.final.done:598// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 0599// CHECK3-NEXT: store ptr [[SIVAR1]], ptr [[TMP14]], align 4600// CHECK3-NEXT: [[TMP15:%.*]] = call i32 @__kmpc_reduce(ptr @[[GLOB2:[0-9]+]], i32 [[TMP2]], i32 1, i32 4, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l63.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)601// CHECK3-NEXT: switch i32 [[TMP15]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [602// CHECK3-NEXT: i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]603// CHECK3-NEXT: i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]604// CHECK3-NEXT: ]605// CHECK3: .omp.reduction.case1:606// CHECK3-NEXT: [[TMP16:%.*]] = load i32, ptr [[TMP0]], align 4607// CHECK3-NEXT: [[TMP17:%.*]] = load i32, ptr [[SIVAR1]], align 4608// CHECK3-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP16]], [[TMP17]]609// CHECK3-NEXT: store i32 [[ADD5]], ptr [[TMP0]], align 4610// CHECK3-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)611// CHECK3-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]612// CHECK3: .omp.reduction.case2:613// CHECK3-NEXT: [[TMP18:%.*]] = load i32, ptr [[SIVAR1]], align 4614// CHECK3-NEXT: [[TMP19:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP18]] monotonic, align 4615// CHECK3-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)616// CHECK3-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]617// CHECK3: .omp.reduction.default:618// CHECK3-NEXT: ret void619//620//621// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l63.omp_outlined.omp.reduction.reduction_func622// CHECK3-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3:[0-9]+]] {623// CHECK3-NEXT: entry:624// CHECK3-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 4625// CHECK3-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 4626// CHECK3-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 4627// CHECK3-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 4628// CHECK3-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 4629// CHECK3-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 4630// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i32 0, i32 0631// CHECK3-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 4632// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i32 0, i32 0633// CHECK3-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 4634// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4635// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4636// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]637// CHECK3-NEXT: store i32 [[ADD]], ptr [[TMP7]], align 4638// CHECK3-NEXT: ret void639//640//641// CHECK3-LABEL: define {{[^@]+}}@_Z5tmainIiET_v642// CHECK3-SAME: () #[[ATTR5:[0-9]+]] comdat {643// CHECK3-NEXT: entry:644// CHECK3-NEXT: [[T_VAR:%.*]] = alloca i32, align 4645// CHECK3-NEXT: [[VEC:%.*]] = alloca [2 x i32], align 4646// CHECK3-NEXT: [[T_VAR_CASTED:%.*]] = alloca i32, align 4647// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4648// CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4649// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4650// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4651// CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8652// CHECK3-NEXT: store i32 0, ptr [[T_VAR]], align 4653// CHECK3-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[VEC]], ptr align 4 @__const._Z5tmainIiET_v.vec, i32 8, i1 false)654// CHECK3-NEXT: [[TMP0:%.*]] = load i32, ptr [[T_VAR]], align 4655// CHECK3-NEXT: store i32 [[TMP0]], ptr [[T_VAR_CASTED]], align 4656// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[T_VAR_CASTED]], align 4657// CHECK3-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0658// CHECK3-NEXT: store i32 [[TMP1]], ptr [[TMP2]], align 4659// CHECK3-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0660// CHECK3-NEXT: store i32 [[TMP1]], ptr [[TMP3]], align 4661// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0662// CHECK3-NEXT: store ptr null, ptr [[TMP4]], align 4663// CHECK3-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0664// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0665// CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0666// CHECK3-NEXT: store i32 3, ptr [[TMP7]], align 4667// CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1668// CHECK3-NEXT: store i32 1, ptr [[TMP8]], align 4669// CHECK3-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2670// CHECK3-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 4671// CHECK3-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3672// CHECK3-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 4673// CHECK3-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4674// CHECK3-NEXT: store ptr @.offload_sizes.1, ptr [[TMP11]], align 4675// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5676// CHECK3-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP12]], align 4677// CHECK3-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6678// CHECK3-NEXT: store ptr null, ptr [[TMP13]], align 4679// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7680// CHECK3-NEXT: store ptr null, ptr [[TMP14]], align 4681// CHECK3-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8682// CHECK3-NEXT: store i64 2, ptr [[TMP15]], align 8683// CHECK3-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9684// CHECK3-NEXT: store i64 0, ptr [[TMP16]], align 8685// CHECK3-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10686// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4687// CHECK3-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11688// CHECK3-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP18]], align 4689// CHECK3-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12690// CHECK3-NEXT: store i32 0, ptr [[TMP19]], align 4691// CHECK3-NEXT: [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.region_id, ptr [[KERNEL_ARGS]])692// CHECK3-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0693// CHECK3-NEXT: br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]694// CHECK3: omp_offload.failed:695// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32(i32 [[TMP1]]) #[[ATTR2]]696// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]]697// CHECK3: omp_offload.cont:698// CHECK3-NEXT: ret i32 0699//700//701// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32702// CHECK3-SAME: (i32 noundef [[T_VAR:%.*]]) #[[ATTR1]] {703// CHECK3-NEXT: entry:704// CHECK3-NEXT: [[T_VAR_ADDR:%.*]] = alloca i32, align 4705// CHECK3-NEXT: store i32 [[T_VAR]], ptr [[T_VAR_ADDR]], align 4706// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined, ptr [[T_VAR_ADDR]])707// CHECK3-NEXT: ret void708//709//710// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined711// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[T_VAR:%.*]]) #[[ATTR1]] {712// CHECK3-NEXT: entry:713// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4714// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4715// CHECK3-NEXT: [[T_VAR_ADDR:%.*]] = alloca ptr, align 4716// CHECK3-NEXT: [[T_VAR1:%.*]] = alloca i32, align 4717// CHECK3-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4718// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4719// CHECK3-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4720// CHECK3-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4721// CHECK3-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4722// CHECK3-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4723// CHECK3-NEXT: [[I:%.*]] = alloca i32, align 4724// CHECK3-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 4725// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4726// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4727// CHECK3-NEXT: store ptr [[T_VAR]], ptr [[T_VAR_ADDR]], align 4728// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[T_VAR_ADDR]], align 4729// CHECK3-NEXT: store i32 0, ptr [[T_VAR1]], align 4730// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4731// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 4732// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4733// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4734// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4735// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4736// CHECK3-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)737// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4738// CHECK3-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1739// CHECK3-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]740// CHECK3: cond.true:741// CHECK3-NEXT: br label [[COND_END:%.*]]742// CHECK3: cond.false:743// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4744// CHECK3-NEXT: br label [[COND_END]]745// CHECK3: cond.end:746// CHECK3-NEXT: [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]747// CHECK3-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4748// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4749// CHECK3-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4750// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]751// CHECK3: omp.inner.for.cond:752// CHECK3-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12:![0-9]+]]753// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP12]]754// CHECK3-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]755// CHECK3-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]756// CHECK3: omp.inner.for.body:757// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]758// CHECK3-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1759// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]760// CHECK3-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP12]]761// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP12]]762// CHECK3-NEXT: [[TMP10:%.*]] = load i32, ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP12]]763// CHECK3-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]764// CHECK3-NEXT: store i32 [[ADD3]], ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP12]]765// CHECK3-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]766// CHECK3: omp.body.continue:767// CHECK3-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]768// CHECK3: omp.inner.for.inc:769// CHECK3-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]770// CHECK3-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP11]], 1771// CHECK3-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]772// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP13:![0-9]+]]773// CHECK3: omp.inner.for.end:774// CHECK3-NEXT: br label [[OMP_LOOP_EXIT:%.*]]775// CHECK3: omp.loop.exit:776// CHECK3-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])777// CHECK3-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 4778// CHECK3-NEXT: [[TMP13:%.*]] = icmp ne i32 [[TMP12]], 0779// CHECK3-NEXT: br i1 [[TMP13]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]780// CHECK3: .omp.final.then:781// CHECK3-NEXT: store i32 2, ptr [[I]], align 4782// CHECK3-NEXT: br label [[DOTOMP_FINAL_DONE]]783// CHECK3: .omp.final.done:784// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 0785// CHECK3-NEXT: store ptr [[T_VAR1]], ptr [[TMP14]], align 4786// CHECK3-NEXT: [[TMP15:%.*]] = call i32 @__kmpc_reduce(ptr @[[GLOB2]], 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)787// CHECK3-NEXT: switch i32 [[TMP15]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [788// CHECK3-NEXT: i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]789// CHECK3-NEXT: i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]790// CHECK3-NEXT: ]791// CHECK3: .omp.reduction.case1:792// CHECK3-NEXT: [[TMP16:%.*]] = load i32, ptr [[TMP0]], align 4793// CHECK3-NEXT: [[TMP17:%.*]] = load i32, ptr [[T_VAR1]], align 4794// CHECK3-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP16]], [[TMP17]]795// CHECK3-NEXT: store i32 [[ADD5]], ptr [[TMP0]], align 4796// CHECK3-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)797// CHECK3-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]798// CHECK3: .omp.reduction.case2:799// CHECK3-NEXT: [[TMP18:%.*]] = load i32, ptr [[T_VAR1]], align 4800// CHECK3-NEXT: [[TMP19:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP18]] monotonic, align 4801// CHECK3-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)802// CHECK3-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]803// CHECK3: .omp.reduction.default:804// CHECK3-NEXT: ret void805//806//807// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp.reduction.reduction_func808// CHECK3-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3]] {809// CHECK3-NEXT: entry:810// CHECK3-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 4811// CHECK3-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 4812// CHECK3-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 4813// CHECK3-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 4814// CHECK3-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 4815// CHECK3-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 4816// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i32 0, i32 0817// CHECK3-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 4818// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i32 0, i32 0819// CHECK3-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 4820// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4821// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4822// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]823// CHECK3-NEXT: store i32 [[ADD]], ptr [[TMP7]], align 4824// CHECK3-NEXT: ret void825//826//827// CHECK5-LABEL: define {{[^@]+}}@main828// CHECK5-SAME: () #[[ATTR0:[0-9]+]] {829// CHECK5-NEXT: entry:830// CHECK5-NEXT: [[RETVAL:%.*]] = alloca i32, align 4831// CHECK5-NEXT: [[TMP:%.*]] = alloca i32, align 4832// CHECK5-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4833// CHECK5-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4834// CHECK5-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4835// CHECK5-NEXT: [[I:%.*]] = alloca i32, align 4836// CHECK5-NEXT: [[SIVAR:%.*]] = alloca i32, align 4837// CHECK5-NEXT: store i32 0, ptr [[RETVAL]], align 4838// CHECK5-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4839// CHECK5-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 4840// CHECK5-NEXT: [[TMP0:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4841// CHECK5-NEXT: store i32 [[TMP0]], ptr [[DOTOMP_IV]], align 4842// CHECK5-NEXT: store i32 0, ptr [[SIVAR]], align 4843// CHECK5-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]844// CHECK5: omp.inner.for.cond:845// CHECK5-NEXT: [[TMP1:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP2:![0-9]+]]846// CHECK5-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP2]]847// CHECK5-NEXT: [[CMP:%.*]] = icmp sle i32 [[TMP1]], [[TMP2]]848// CHECK5-NEXT: br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]849// CHECK5: omp.inner.for.body:850// CHECK5-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP2]]851// CHECK5-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1852// CHECK5-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]853// CHECK5-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP2]]854// CHECK5-NEXT: [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP2]]855// CHECK5-NEXT: [[TMP5:%.*]] = load i32, ptr [[SIVAR]], align 4, !llvm.access.group [[ACC_GRP2]]856// CHECK5-NEXT: [[ADD1:%.*]] = add nsw i32 [[TMP5]], [[TMP4]]857// CHECK5-NEXT: store i32 [[ADD1]], ptr [[SIVAR]], align 4, !llvm.access.group [[ACC_GRP2]]858// CHECK5-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]859// CHECK5: omp.body.continue:860// CHECK5-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]861// CHECK5: omp.inner.for.inc:862// CHECK5-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP2]]863// CHECK5-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1864// CHECK5-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP2]]865// CHECK5-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP3:![0-9]+]]866// CHECK5: omp.inner.for.end:867// CHECK5-NEXT: store i32 2, ptr [[I]], align 4868// CHECK5-NEXT: [[TMP7:%.*]] = load i32, ptr @_ZZ4mainE5sivar, align 4869// CHECK5-NEXT: [[TMP8:%.*]] = load i32, ptr [[SIVAR]], align 4870// CHECK5-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP7]], [[TMP8]]871// CHECK5-NEXT: store i32 [[ADD3]], ptr @_ZZ4mainE5sivar, align 4872// CHECK5-NEXT: [[CALL:%.*]] = call noundef signext i32 @_Z5tmainIiET_v()873// CHECK5-NEXT: ret i32 [[CALL]]874//875//876// CHECK5-LABEL: define {{[^@]+}}@_Z5tmainIiET_v877// CHECK5-SAME: () #[[ATTR1:[0-9]+]] comdat {878// CHECK5-NEXT: entry:879// CHECK5-NEXT: [[T_VAR:%.*]] = alloca i32, align 4880// CHECK5-NEXT: [[VEC:%.*]] = alloca [2 x i32], align 4881// CHECK5-NEXT: [[TMP:%.*]] = alloca i32, align 4882// CHECK5-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4883// CHECK5-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4884// CHECK5-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4885// CHECK5-NEXT: [[I:%.*]] = alloca i32, align 4886// CHECK5-NEXT: [[T_VAR1:%.*]] = alloca i32, align 4887// CHECK5-NEXT: store i32 0, ptr [[T_VAR]], align 4888// CHECK5-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[VEC]], ptr align 4 @__const._Z5tmainIiET_v.vec, i64 8, i1 false)889// CHECK5-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4890// CHECK5-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 4891// CHECK5-NEXT: [[TMP0:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4892// CHECK5-NEXT: store i32 [[TMP0]], ptr [[DOTOMP_IV]], align 4893// CHECK5-NEXT: store i32 0, ptr [[T_VAR1]], align 4894// CHECK5-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]895// CHECK5: omp.inner.for.cond:896// CHECK5-NEXT: [[TMP1:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6:![0-9]+]]897// CHECK5-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP6]]898// CHECK5-NEXT: [[CMP:%.*]] = icmp sle i32 [[TMP1]], [[TMP2]]899// CHECK5-NEXT: br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]900// CHECK5: omp.inner.for.body:901// CHECK5-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6]]902// CHECK5-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1903// CHECK5-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]904// CHECK5-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP6]]905// CHECK5-NEXT: [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP6]]906// CHECK5-NEXT: [[TMP5:%.*]] = load i32, ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP6]]907// CHECK5-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP5]], [[TMP4]]908// CHECK5-NEXT: store i32 [[ADD2]], ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP6]]909// CHECK5-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]910// CHECK5: omp.body.continue:911// CHECK5-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]912// CHECK5: omp.inner.for.inc:913// CHECK5-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6]]914// CHECK5-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP6]], 1915// CHECK5-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6]]916// CHECK5-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP7:![0-9]+]]917// CHECK5: omp.inner.for.end:918// CHECK5-NEXT: store i32 2, ptr [[I]], align 4919// CHECK5-NEXT: [[TMP7:%.*]] = load i32, ptr [[T_VAR]], align 4920// CHECK5-NEXT: [[TMP8:%.*]] = load i32, ptr [[T_VAR1]], align 4921// CHECK5-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP7]], [[TMP8]]922// CHECK5-NEXT: store i32 [[ADD4]], ptr [[T_VAR]], align 4923// CHECK5-NEXT: ret i32 0924//925//926// CHECK7-LABEL: define {{[^@]+}}@main927// CHECK7-SAME: () #[[ATTR0:[0-9]+]] {928// CHECK7-NEXT: entry:929// CHECK7-NEXT: [[RETVAL:%.*]] = alloca i32, align 4930// CHECK7-NEXT: [[TMP:%.*]] = alloca i32, align 4931// CHECK7-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4932// CHECK7-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4933// CHECK7-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4934// CHECK7-NEXT: [[I:%.*]] = alloca i32, align 4935// CHECK7-NEXT: [[SIVAR:%.*]] = alloca i32, align 4936// CHECK7-NEXT: store i32 0, ptr [[RETVAL]], align 4937// CHECK7-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4938// CHECK7-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 4939// CHECK7-NEXT: [[TMP0:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4940// CHECK7-NEXT: store i32 [[TMP0]], ptr [[DOTOMP_IV]], align 4941// CHECK7-NEXT: store i32 0, ptr [[SIVAR]], align 4942// CHECK7-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]943// CHECK7: omp.inner.for.cond:944// CHECK7-NEXT: [[TMP1:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP3:![0-9]+]]945// CHECK7-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP3]]946// CHECK7-NEXT: [[CMP:%.*]] = icmp sle i32 [[TMP1]], [[TMP2]]947// CHECK7-NEXT: br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]948// CHECK7: omp.inner.for.body:949// CHECK7-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP3]]950// CHECK7-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1951// CHECK7-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]952// CHECK7-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP3]]953// CHECK7-NEXT: [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP3]]954// CHECK7-NEXT: [[TMP5:%.*]] = load i32, ptr [[SIVAR]], align 4, !llvm.access.group [[ACC_GRP3]]955// CHECK7-NEXT: [[ADD1:%.*]] = add nsw i32 [[TMP5]], [[TMP4]]956// CHECK7-NEXT: store i32 [[ADD1]], ptr [[SIVAR]], align 4, !llvm.access.group [[ACC_GRP3]]957// CHECK7-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]958// CHECK7: omp.body.continue:959// CHECK7-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]960// CHECK7: omp.inner.for.inc:961// CHECK7-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP3]]962// CHECK7-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1963// CHECK7-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP3]]964// CHECK7-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP4:![0-9]+]]965// CHECK7: omp.inner.for.end:966// CHECK7-NEXT: store i32 2, ptr [[I]], align 4967// CHECK7-NEXT: [[TMP7:%.*]] = load i32, ptr @_ZZ4mainE5sivar, align 4968// CHECK7-NEXT: [[TMP8:%.*]] = load i32, ptr [[SIVAR]], align 4969// CHECK7-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP7]], [[TMP8]]970// CHECK7-NEXT: store i32 [[ADD3]], ptr @_ZZ4mainE5sivar, align 4971// CHECK7-NEXT: [[CALL:%.*]] = call noundef i32 @_Z5tmainIiET_v()972// CHECK7-NEXT: ret i32 [[CALL]]973//974//975// CHECK7-LABEL: define {{[^@]+}}@_Z5tmainIiET_v976// CHECK7-SAME: () #[[ATTR1:[0-9]+]] comdat {977// CHECK7-NEXT: entry:978// CHECK7-NEXT: [[T_VAR:%.*]] = alloca i32, align 4979// CHECK7-NEXT: [[VEC:%.*]] = alloca [2 x i32], align 4980// CHECK7-NEXT: [[TMP:%.*]] = alloca i32, align 4981// CHECK7-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4982// CHECK7-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4983// CHECK7-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4984// CHECK7-NEXT: [[I:%.*]] = alloca i32, align 4985// CHECK7-NEXT: [[T_VAR1:%.*]] = alloca i32, align 4986// CHECK7-NEXT: store i32 0, ptr [[T_VAR]], align 4987// CHECK7-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[VEC]], ptr align 4 @__const._Z5tmainIiET_v.vec, i32 8, i1 false)988// CHECK7-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4989// CHECK7-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 4990// CHECK7-NEXT: [[TMP0:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4991// CHECK7-NEXT: store i32 [[TMP0]], ptr [[DOTOMP_IV]], align 4992// CHECK7-NEXT: store i32 0, ptr [[T_VAR1]], align 4993// CHECK7-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]994// CHECK7: omp.inner.for.cond:995// CHECK7-NEXT: [[TMP1:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP7:![0-9]+]]996// CHECK7-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP7]]997// CHECK7-NEXT: [[CMP:%.*]] = icmp sle i32 [[TMP1]], [[TMP2]]998// CHECK7-NEXT: br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]999// CHECK7: omp.inner.for.body:1000// CHECK7-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP7]]1001// CHECK7-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP3]], 11002// CHECK7-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1003// CHECK7-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP7]]1004// CHECK7-NEXT: [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP7]]1005// CHECK7-NEXT: [[TMP5:%.*]] = load i32, ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP7]]1006// CHECK7-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP5]], [[TMP4]]1007// CHECK7-NEXT: store i32 [[ADD2]], ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP7]]1008// CHECK7-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1009// CHECK7: omp.body.continue:1010// CHECK7-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1011// CHECK7: omp.inner.for.inc:1012// CHECK7-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP7]]1013// CHECK7-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP6]], 11014// CHECK7-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP7]]1015// CHECK7-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP8:![0-9]+]]1016// CHECK7: omp.inner.for.end:1017// CHECK7-NEXT: store i32 2, ptr [[I]], align 41018// CHECK7-NEXT: [[TMP7:%.*]] = load i32, ptr [[T_VAR]], align 41019// CHECK7-NEXT: [[TMP8:%.*]] = load i32, ptr [[T_VAR1]], align 41020// CHECK7-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP7]], [[TMP8]]1021// CHECK7-NEXT: store i32 [[ADD4]], ptr [[T_VAR]], align 41022// CHECK7-NEXT: ret i32 01023//1024//1025// CHECK9-LABEL: define {{[^@]+}}@main1026// CHECK9-SAME: () #[[ATTR0:[0-9]+]] {1027// CHECK9-NEXT: entry:1028// CHECK9-NEXT: [[RETVAL:%.*]] = alloca i32, align 41029// CHECK9-NEXT: [[REF_TMP:%.*]] = alloca [[CLASS_ANON:%.*]], align 11030// CHECK9-NEXT: store i32 0, ptr [[RETVAL]], align 41031// CHECK9-NEXT: call void @"_ZZ4mainENK3$_0clEv"(ptr noundef nonnull align 1 dereferenceable(1) [[REF_TMP]])1032// CHECK9-NEXT: ret i32 01033//1034//1035// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l451036// CHECK9-SAME: (i64 noundef [[SIVAR:%.*]]) #[[ATTR2:[0-9]+]] {1037// CHECK9-NEXT: entry:1038// CHECK9-NEXT: [[SIVAR_ADDR:%.*]] = alloca i64, align 81039// CHECK9-NEXT: store i64 [[SIVAR]], ptr [[SIVAR_ADDR]], align 81040// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3:[0-9]+]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l45.omp_outlined, ptr [[SIVAR_ADDR]])1041// CHECK9-NEXT: ret void1042//1043//1044// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l45.omp_outlined1045// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR2]] {1046// CHECK9-NEXT: entry:1047// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81048// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81049// CHECK9-NEXT: [[SIVAR_ADDR:%.*]] = alloca ptr, align 81050// CHECK9-NEXT: [[SIVAR1:%.*]] = alloca i32, align 41051// CHECK9-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41052// CHECK9-NEXT: [[TMP:%.*]] = alloca i32, align 41053// CHECK9-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41054// CHECK9-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41055// CHECK9-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41056// CHECK9-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41057// CHECK9-NEXT: [[I:%.*]] = alloca i32, align 41058// CHECK9-NEXT: [[REF_TMP:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 81059// CHECK9-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 81060// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81061// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81062// CHECK9-NEXT: store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 81063// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 81064// CHECK9-NEXT: store i32 0, ptr [[SIVAR1]], align 41065// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41066// CHECK9-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 41067// CHECK9-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41068// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41069// CHECK9-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81070// CHECK9-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41071// CHECK9-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1072// CHECK9-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41073// CHECK9-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 11074// CHECK9-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1075// CHECK9: cond.true:1076// CHECK9-NEXT: br label [[COND_END:%.*]]1077// CHECK9: cond.false:1078// CHECK9-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41079// CHECK9-NEXT: br label [[COND_END]]1080// CHECK9: cond.end:1081// CHECK9-NEXT: [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]1082// CHECK9-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41083// CHECK9-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41084// CHECK9-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 41085// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1086// CHECK9: omp.inner.for.cond:1087// CHECK9-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP4:![0-9]+]]1088// CHECK9-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP4]]1089// CHECK9-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]1090// CHECK9-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1091// CHECK9: omp.inner.for.body:1092// CHECK9-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP4]]1093// CHECK9-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 11094// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1095// CHECK9-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP4]]1096// CHECK9-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP4]]1097// CHECK9-NEXT: [[TMP10:%.*]] = load i32, ptr [[SIVAR1]], align 4, !llvm.access.group [[ACC_GRP4]]1098// CHECK9-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]1099// CHECK9-NEXT: store i32 [[ADD3]], ptr [[SIVAR1]], align 4, !llvm.access.group [[ACC_GRP4]]1100// CHECK9-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 01101// CHECK9-NEXT: store ptr [[SIVAR1]], ptr [[TMP11]], align 8, !llvm.access.group [[ACC_GRP4]]1102// CHECK9-NEXT: call void @"_ZZZ4mainENK3$_0clEvENKUlvE_clEv"(ptr noundef nonnull align 8 dereferenceable(8) [[REF_TMP]]), !llvm.access.group [[ACC_GRP4]]1103// CHECK9-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1104// CHECK9: omp.body.continue:1105// CHECK9-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1106// CHECK9: omp.inner.for.inc:1107// CHECK9-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP4]]1108// CHECK9-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP12]], 11109// CHECK9-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP4]]1110// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP5:![0-9]+]]1111// CHECK9: omp.inner.for.end:1112// CHECK9-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1113// CHECK9: omp.loop.exit:1114// CHECK9-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])1115// CHECK9-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 41116// CHECK9-NEXT: [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 01117// CHECK9-NEXT: br i1 [[TMP14]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]1118// CHECK9: .omp.final.then:1119// CHECK9-NEXT: store i32 2, ptr [[I]], align 41120// CHECK9-NEXT: br label [[DOTOMP_FINAL_DONE]]1121// CHECK9: .omp.final.done:1122// CHECK9-NEXT: [[TMP15:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 01123// CHECK9-NEXT: store ptr [[SIVAR1]], ptr [[TMP15]], align 81124// CHECK9-NEXT: [[TMP16:%.*]] = call i32 @__kmpc_reduce(ptr @[[GLOB2:[0-9]+]], i32 [[TMP2]], i32 1, i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l45.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)1125// CHECK9-NEXT: switch i32 [[TMP16]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [1126// CHECK9-NEXT: i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]1127// CHECK9-NEXT: i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]1128// CHECK9-NEXT: ]1129// CHECK9: .omp.reduction.case1:1130// CHECK9-NEXT: [[TMP17:%.*]] = load i32, ptr [[TMP0]], align 41131// CHECK9-NEXT: [[TMP18:%.*]] = load i32, ptr [[SIVAR1]], align 41132// CHECK9-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP17]], [[TMP18]]1133// CHECK9-NEXT: store i32 [[ADD5]], ptr [[TMP0]], align 41134// CHECK9-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)1135// CHECK9-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]1136// CHECK9: .omp.reduction.case2:1137// CHECK9-NEXT: [[TMP19:%.*]] = load i32, ptr [[SIVAR1]], align 41138// CHECK9-NEXT: [[TMP20:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP19]] monotonic, align 41139// CHECK9-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)1140// CHECK9-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]1141// CHECK9: .omp.reduction.default:1142// CHECK9-NEXT: ret void1143//1144//1145// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l45.omp_outlined.omp.reduction.reduction_func1146// CHECK9-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR4:[0-9]+]] {1147// CHECK9-NEXT: entry:1148// CHECK9-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 81149// CHECK9-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 81150// CHECK9-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 81151// CHECK9-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 81152// CHECK9-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 81153// CHECK9-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 81154// CHECK9-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i64 0, i64 01155// CHECK9-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 81156// CHECK9-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i64 0, i64 01157// CHECK9-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 81158// CHECK9-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41159// CHECK9-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 41160// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]1161// CHECK9-NEXT: store i32 [[ADD]], ptr [[TMP7]], align 41162// CHECK9-NEXT: ret void1163//1164//1165// CHECK11-LABEL: define {{[^@]+}}@main1166// CHECK11-SAME: () #[[ATTR0:[0-9]+]] {1167// CHECK11-NEXT: entry:1168// CHECK11-NEXT: [[RETVAL:%.*]] = alloca i32, align 41169// CHECK11-NEXT: [[REF_TMP:%.*]] = alloca [[CLASS_ANON:%.*]], align 11170// CHECK11-NEXT: store i32 0, ptr [[RETVAL]], align 41171// CHECK11-NEXT: call void @"_ZZ4mainENK3$_0clEv"(ptr noundef nonnull align 1 dereferenceable(1) [[REF_TMP]])1172// CHECK11-NEXT: ret i32 01173//1174