931 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 --implicit-check-not="{{__kmpc|__tgt}}"10// 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 --implicit-check-not="{{__kmpc|__tgt}}"12// 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}}"13// 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 --implicit-check-not="{{__kmpc|__tgt}}"15 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 --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 target33#pragma omp teams distribute 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 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 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] zeroinitializer, 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 0, 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 4196// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4197// 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 4201// 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 4204// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4205// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[SIVAR1]], align 4206// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]207// CHECK1-NEXT: store i32 [[ADD3]], ptr [[SIVAR1]], align 4208// 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 4213// CHECK1-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP11]], 1214// CHECK1-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4215// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]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:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0221// CHECK1-NEXT: store ptr [[SIVAR1]], ptr [[TMP12]], align 8222// CHECK1-NEXT: [[TMP13:%.*]] = 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)223// CHECK1-NEXT: switch i32 [[TMP13]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [224// CHECK1-NEXT: i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]225// CHECK1-NEXT: i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]226// CHECK1-NEXT: ]227// CHECK1: .omp.reduction.case1:228// CHECK1-NEXT: [[TMP14:%.*]] = load i32, ptr [[TMP0]], align 4229// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[SIVAR1]], align 4230// CHECK1-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP14]], [[TMP15]]231// CHECK1-NEXT: store i32 [[ADD5]], ptr [[TMP0]], align 4232// CHECK1-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)233// CHECK1-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]234// CHECK1: .omp.reduction.case2:235// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[SIVAR1]], align 4236// CHECK1-NEXT: [[TMP17:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP16]] monotonic, align 4237// CHECK1-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)238// 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_l63.omp_outlined.omp.reduction.reduction_func244// CHECK1-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3:[0-9]+]] {245// CHECK1-NEXT: entry:246// CHECK1-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 8247// CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 8248// CHECK1-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 8249// CHECK1-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 8250// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 8251// CHECK1-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 8252// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i64 0, i64 0253// CHECK1-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 8254// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i64 0, i64 0255// CHECK1-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 8256// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4257// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4258// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]259// CHECK1-NEXT: store i32 [[ADD]], ptr [[TMP7]], align 4260// CHECK1-NEXT: ret void261//262//263// CHECK1-LABEL: define {{[^@]+}}@_Z5tmainIiET_v264// CHECK1-SAME: () #[[ATTR5:[0-9]+]] comdat {265// CHECK1-NEXT: entry:266// CHECK1-NEXT: [[T_VAR:%.*]] = alloca i32, align 4267// CHECK1-NEXT: [[VEC:%.*]] = alloca [2 x i32], align 4268// CHECK1-NEXT: [[T_VAR_CASTED:%.*]] = alloca i64, align 8269// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8270// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8271// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8272// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4273// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8274// CHECK1-NEXT: store i32 0, ptr [[T_VAR]], align 4275// CHECK1-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[VEC]], ptr align 4 @__const._Z5tmainIiET_v.vec, i64 8, i1 false)276// CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr [[T_VAR]], align 4277// CHECK1-NEXT: store i32 [[TMP0]], ptr [[T_VAR_CASTED]], align 4278// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[T_VAR_CASTED]], align 8279// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0280// CHECK1-NEXT: store i64 [[TMP1]], ptr [[TMP2]], align 8281// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0282// CHECK1-NEXT: store i64 [[TMP1]], ptr [[TMP3]], align 8283// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0284// CHECK1-NEXT: store ptr null, ptr [[TMP4]], align 8285// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0286// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0287// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0288// CHECK1-NEXT: store i32 3, ptr [[TMP7]], align 4289// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1290// CHECK1-NEXT: store i32 1, ptr [[TMP8]], align 4291// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2292// CHECK1-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 8293// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3294// CHECK1-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 8295// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4296// CHECK1-NEXT: store ptr @.offload_sizes.1, ptr [[TMP11]], align 8297// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5298// CHECK1-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP12]], align 8299// CHECK1-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6300// CHECK1-NEXT: store ptr null, ptr [[TMP13]], align 8301// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7302// CHECK1-NEXT: store ptr null, ptr [[TMP14]], align 8303// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8304// CHECK1-NEXT: store i64 2, ptr [[TMP15]], align 8305// CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9306// CHECK1-NEXT: store i64 0, ptr [[TMP16]], align 8307// CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10308// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4309// CHECK1-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11310// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP18]], align 4311// CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12312// CHECK1-NEXT: store i32 0, ptr [[TMP19]], align 4313// CHECK1-NEXT: [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.region_id, ptr [[KERNEL_ARGS]])314// CHECK1-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0315// CHECK1-NEXT: br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]316// CHECK1: omp_offload.failed:317// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32(i64 [[TMP1]]) #[[ATTR2]]318// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]319// CHECK1: omp_offload.cont:320// CHECK1-NEXT: ret i32 0321//322//323// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32324// CHECK1-SAME: (i64 noundef [[T_VAR:%.*]]) #[[ATTR1]] {325// CHECK1-NEXT: entry:326// CHECK1-NEXT: [[T_VAR_ADDR:%.*]] = alloca i64, align 8327// CHECK1-NEXT: store i64 [[T_VAR]], ptr [[T_VAR_ADDR]], align 8328// 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]])329// CHECK1-NEXT: ret void330//331//332// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined333// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[T_VAR:%.*]]) #[[ATTR1]] {334// CHECK1-NEXT: entry:335// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8336// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8337// CHECK1-NEXT: [[T_VAR_ADDR:%.*]] = alloca ptr, align 8338// CHECK1-NEXT: [[T_VAR1:%.*]] = alloca i32, align 4339// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4340// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4341// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4342// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4343// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4344// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4345// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4346// CHECK1-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 8347// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8348// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8349// CHECK1-NEXT: store ptr [[T_VAR]], ptr [[T_VAR_ADDR]], align 8350// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[T_VAR_ADDR]], align 8351// CHECK1-NEXT: store i32 0, ptr [[T_VAR1]], align 4352// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4353// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 4354// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4355// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4356// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8357// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4358// 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)359// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4360// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1361// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]362// CHECK1: cond.true:363// CHECK1-NEXT: br label [[COND_END:%.*]]364// CHECK1: cond.false:365// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4366// CHECK1-NEXT: br label [[COND_END]]367// CHECK1: cond.end:368// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]369// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4370// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4371// CHECK1-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4372// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]373// CHECK1: omp.inner.for.cond:374// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4375// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4376// CHECK1-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]377// CHECK1-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]378// CHECK1: omp.inner.for.body:379// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4380// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1381// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]382// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4383// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4384// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[T_VAR1]], align 4385// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]386// CHECK1-NEXT: store i32 [[ADD3]], ptr [[T_VAR1]], align 4387// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]388// CHECK1: omp.body.continue:389// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]390// CHECK1: omp.inner.for.inc:391// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4392// CHECK1-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP11]], 1393// CHECK1-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4394// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]395// CHECK1: omp.inner.for.end:396// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]397// CHECK1: omp.loop.exit:398// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])399// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0400// CHECK1-NEXT: store ptr [[T_VAR1]], ptr [[TMP12]], align 8401// CHECK1-NEXT: [[TMP13:%.*]] = 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)402// CHECK1-NEXT: switch i32 [[TMP13]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [403// CHECK1-NEXT: i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]404// CHECK1-NEXT: i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]405// CHECK1-NEXT: ]406// CHECK1: .omp.reduction.case1:407// CHECK1-NEXT: [[TMP14:%.*]] = load i32, ptr [[TMP0]], align 4408// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[T_VAR1]], align 4409// CHECK1-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP14]], [[TMP15]]410// CHECK1-NEXT: store i32 [[ADD5]], ptr [[TMP0]], align 4411// CHECK1-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)412// CHECK1-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]413// CHECK1: .omp.reduction.case2:414// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[T_VAR1]], align 4415// CHECK1-NEXT: [[TMP17:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP16]] monotonic, align 4416// CHECK1-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)417// CHECK1-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]418// CHECK1: .omp.reduction.default:419// CHECK1-NEXT: ret void420//421//422// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp.reduction.reduction_func423// CHECK1-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3]] {424// CHECK1-NEXT: entry:425// CHECK1-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 8426// CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 8427// CHECK1-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 8428// CHECK1-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 8429// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 8430// CHECK1-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 8431// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i64 0, i64 0432// CHECK1-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 8433// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i64 0, i64 0434// CHECK1-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 8435// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4436// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4437// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]438// CHECK1-NEXT: store i32 [[ADD]], ptr [[TMP7]], align 4439// CHECK1-NEXT: ret void440//441//442// CHECK3-LABEL: define {{[^@]+}}@main443// CHECK3-SAME: () #[[ATTR0:[0-9]+]] {444// CHECK3-NEXT: entry:445// CHECK3-NEXT: [[RETVAL:%.*]] = alloca i32, align 4446// CHECK3-NEXT: [[SIVAR_CASTED:%.*]] = alloca i32, align 4447// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4448// CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4449// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4450// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4451// CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8452// CHECK3-NEXT: store i32 0, ptr [[RETVAL]], align 4453// CHECK3-NEXT: [[TMP0:%.*]] = load i32, ptr @_ZZ4mainE5sivar, align 4454// CHECK3-NEXT: store i32 [[TMP0]], ptr [[SIVAR_CASTED]], align 4455// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[SIVAR_CASTED]], align 4456// CHECK3-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0457// CHECK3-NEXT: store i32 [[TMP1]], ptr [[TMP2]], align 4458// CHECK3-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0459// CHECK3-NEXT: store i32 [[TMP1]], ptr [[TMP3]], align 4460// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0461// CHECK3-NEXT: store ptr null, ptr [[TMP4]], align 4462// CHECK3-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0463// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0464// CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0465// CHECK3-NEXT: store i32 3, ptr [[TMP7]], align 4466// CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1467// CHECK3-NEXT: store i32 1, ptr [[TMP8]], align 4468// CHECK3-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2469// CHECK3-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 4470// CHECK3-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3471// CHECK3-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 4472// CHECK3-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4473// CHECK3-NEXT: store ptr @.offload_sizes, ptr [[TMP11]], align 4474// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5475// CHECK3-NEXT: store ptr @.offload_maptypes, ptr [[TMP12]], align 4476// CHECK3-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6477// CHECK3-NEXT: store ptr null, ptr [[TMP13]], align 4478// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7479// CHECK3-NEXT: store ptr null, ptr [[TMP14]], align 4480// CHECK3-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8481// CHECK3-NEXT: store i64 2, ptr [[TMP15]], align 8482// CHECK3-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9483// CHECK3-NEXT: store i64 0, ptr [[TMP16]], align 8484// CHECK3-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10485// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4486// CHECK3-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11487// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP18]], align 4488// CHECK3-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12489// CHECK3-NEXT: store i32 0, ptr [[TMP19]], align 4490// CHECK3-NEXT: [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l63.region_id, ptr [[KERNEL_ARGS]])491// CHECK3-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0492// CHECK3-NEXT: br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]493// CHECK3: omp_offload.failed:494// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l63(i32 [[TMP1]]) #[[ATTR2:[0-9]+]]495// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]]496// CHECK3: omp_offload.cont:497// CHECK3-NEXT: [[CALL:%.*]] = call noundef i32 @_Z5tmainIiET_v()498// CHECK3-NEXT: ret i32 [[CALL]]499//500//501// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l63502// CHECK3-SAME: (i32 noundef [[SIVAR:%.*]]) #[[ATTR1:[0-9]+]] {503// CHECK3-NEXT: entry:504// CHECK3-NEXT: [[SIVAR_ADDR:%.*]] = alloca i32, align 4505// CHECK3-NEXT: store i32 [[SIVAR]], ptr [[SIVAR_ADDR]], align 4506// 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]])507// CHECK3-NEXT: ret void508//509//510// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l63.omp_outlined511// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR1]] {512// CHECK3-NEXT: entry:513// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4514// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4515// CHECK3-NEXT: [[SIVAR_ADDR:%.*]] = alloca ptr, align 4516// CHECK3-NEXT: [[SIVAR1:%.*]] = alloca i32, align 4517// CHECK3-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4518// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4519// CHECK3-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4520// CHECK3-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4521// CHECK3-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4522// CHECK3-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4523// CHECK3-NEXT: [[I:%.*]] = alloca i32, align 4524// CHECK3-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 4525// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4526// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4527// CHECK3-NEXT: store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 4528// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 4529// CHECK3-NEXT: store i32 0, ptr [[SIVAR1]], align 4530// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4531// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 4532// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4533// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4534// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4535// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4536// 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)537// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4538// CHECK3-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1539// CHECK3-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]540// CHECK3: cond.true:541// CHECK3-NEXT: br label [[COND_END:%.*]]542// CHECK3: cond.false:543// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4544// CHECK3-NEXT: br label [[COND_END]]545// CHECK3: cond.end:546// CHECK3-NEXT: [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]547// CHECK3-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4548// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4549// CHECK3-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4550// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]551// CHECK3: omp.inner.for.cond:552// CHECK3-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4553// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4554// CHECK3-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]555// CHECK3-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]556// CHECK3: omp.inner.for.body:557// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4558// CHECK3-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1559// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]560// CHECK3-NEXT: store i32 [[ADD]], ptr [[I]], align 4561// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4562// CHECK3-NEXT: [[TMP10:%.*]] = load i32, ptr [[SIVAR1]], align 4563// CHECK3-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]564// CHECK3-NEXT: store i32 [[ADD3]], ptr [[SIVAR1]], align 4565// CHECK3-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]566// CHECK3: omp.body.continue:567// CHECK3-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]568// CHECK3: omp.inner.for.inc:569// CHECK3-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4570// CHECK3-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP11]], 1571// CHECK3-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4572// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND]]573// CHECK3: omp.inner.for.end:574// CHECK3-NEXT: br label [[OMP_LOOP_EXIT:%.*]]575// CHECK3: omp.loop.exit:576// CHECK3-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])577// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 0578// CHECK3-NEXT: store ptr [[SIVAR1]], ptr [[TMP12]], align 4579// CHECK3-NEXT: [[TMP13:%.*]] = 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)580// CHECK3-NEXT: switch i32 [[TMP13]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [581// CHECK3-NEXT: i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]582// CHECK3-NEXT: i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]583// CHECK3-NEXT: ]584// CHECK3: .omp.reduction.case1:585// CHECK3-NEXT: [[TMP14:%.*]] = load i32, ptr [[TMP0]], align 4586// CHECK3-NEXT: [[TMP15:%.*]] = load i32, ptr [[SIVAR1]], align 4587// CHECK3-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP14]], [[TMP15]]588// CHECK3-NEXT: store i32 [[ADD5]], ptr [[TMP0]], align 4589// CHECK3-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)590// CHECK3-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]591// CHECK3: .omp.reduction.case2:592// CHECK3-NEXT: [[TMP16:%.*]] = load i32, ptr [[SIVAR1]], align 4593// CHECK3-NEXT: [[TMP17:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP16]] monotonic, align 4594// CHECK3-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)595// CHECK3-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]596// CHECK3: .omp.reduction.default:597// CHECK3-NEXT: ret void598//599//600// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l63.omp_outlined.omp.reduction.reduction_func601// CHECK3-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3:[0-9]+]] {602// CHECK3-NEXT: entry:603// CHECK3-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 4604// CHECK3-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 4605// CHECK3-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 4606// CHECK3-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 4607// CHECK3-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 4608// CHECK3-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 4609// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i32 0, i32 0610// CHECK3-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 4611// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i32 0, i32 0612// CHECK3-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 4613// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4614// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4615// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]616// CHECK3-NEXT: store i32 [[ADD]], ptr [[TMP7]], align 4617// CHECK3-NEXT: ret void618//619//620// CHECK3-LABEL: define {{[^@]+}}@_Z5tmainIiET_v621// CHECK3-SAME: () #[[ATTR5:[0-9]+]] comdat {622// CHECK3-NEXT: entry:623// CHECK3-NEXT: [[T_VAR:%.*]] = alloca i32, align 4624// CHECK3-NEXT: [[VEC:%.*]] = alloca [2 x i32], align 4625// CHECK3-NEXT: [[T_VAR_CASTED:%.*]] = alloca i32, align 4626// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4627// CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4628// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4629// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4630// CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8631// CHECK3-NEXT: store i32 0, ptr [[T_VAR]], align 4632// CHECK3-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[VEC]], ptr align 4 @__const._Z5tmainIiET_v.vec, i32 8, i1 false)633// CHECK3-NEXT: [[TMP0:%.*]] = load i32, ptr [[T_VAR]], align 4634// CHECK3-NEXT: store i32 [[TMP0]], ptr [[T_VAR_CASTED]], align 4635// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[T_VAR_CASTED]], align 4636// CHECK3-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0637// CHECK3-NEXT: store i32 [[TMP1]], ptr [[TMP2]], align 4638// CHECK3-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0639// CHECK3-NEXT: store i32 [[TMP1]], ptr [[TMP3]], align 4640// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0641// CHECK3-NEXT: store ptr null, ptr [[TMP4]], align 4642// CHECK3-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0643// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0644// CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0645// CHECK3-NEXT: store i32 3, ptr [[TMP7]], align 4646// CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1647// CHECK3-NEXT: store i32 1, ptr [[TMP8]], align 4648// CHECK3-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2649// CHECK3-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 4650// CHECK3-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3651// CHECK3-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 4652// CHECK3-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4653// CHECK3-NEXT: store ptr @.offload_sizes.1, ptr [[TMP11]], align 4654// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5655// CHECK3-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP12]], align 4656// CHECK3-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6657// CHECK3-NEXT: store ptr null, ptr [[TMP13]], align 4658// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7659// CHECK3-NEXT: store ptr null, ptr [[TMP14]], align 4660// CHECK3-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8661// CHECK3-NEXT: store i64 2, ptr [[TMP15]], align 8662// CHECK3-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9663// CHECK3-NEXT: store i64 0, ptr [[TMP16]], align 8664// CHECK3-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10665// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4666// CHECK3-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11667// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP18]], align 4668// CHECK3-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12669// CHECK3-NEXT: store i32 0, ptr [[TMP19]], align 4670// CHECK3-NEXT: [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.region_id, ptr [[KERNEL_ARGS]])671// CHECK3-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0672// CHECK3-NEXT: br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]673// CHECK3: omp_offload.failed:674// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32(i32 [[TMP1]]) #[[ATTR2]]675// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]]676// CHECK3: omp_offload.cont:677// CHECK3-NEXT: ret i32 0678//679//680// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32681// CHECK3-SAME: (i32 noundef [[T_VAR:%.*]]) #[[ATTR1]] {682// CHECK3-NEXT: entry:683// CHECK3-NEXT: [[T_VAR_ADDR:%.*]] = alloca i32, align 4684// CHECK3-NEXT: store i32 [[T_VAR]], ptr [[T_VAR_ADDR]], align 4685// 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]])686// CHECK3-NEXT: ret void687//688//689// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined690// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[T_VAR:%.*]]) #[[ATTR1]] {691// CHECK3-NEXT: entry:692// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4693// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4694// CHECK3-NEXT: [[T_VAR_ADDR:%.*]] = alloca ptr, align 4695// CHECK3-NEXT: [[T_VAR1:%.*]] = alloca i32, align 4696// CHECK3-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4697// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4698// CHECK3-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4699// CHECK3-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4700// CHECK3-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4701// CHECK3-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4702// CHECK3-NEXT: [[I:%.*]] = alloca i32, align 4703// CHECK3-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 4704// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4705// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4706// CHECK3-NEXT: store ptr [[T_VAR]], ptr [[T_VAR_ADDR]], align 4707// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[T_VAR_ADDR]], align 4708// CHECK3-NEXT: store i32 0, ptr [[T_VAR1]], align 4709// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4710// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 4711// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4712// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4713// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4714// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4715// 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)716// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4717// CHECK3-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1718// CHECK3-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]719// CHECK3: cond.true:720// CHECK3-NEXT: br label [[COND_END:%.*]]721// CHECK3: cond.false:722// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4723// CHECK3-NEXT: br label [[COND_END]]724// CHECK3: cond.end:725// CHECK3-NEXT: [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]726// CHECK3-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4727// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4728// CHECK3-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4729// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]730// CHECK3: omp.inner.for.cond:731// CHECK3-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4732// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4733// CHECK3-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]734// CHECK3-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]735// CHECK3: omp.inner.for.body:736// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4737// CHECK3-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1738// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]739// CHECK3-NEXT: store i32 [[ADD]], ptr [[I]], align 4740// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4741// CHECK3-NEXT: [[TMP10:%.*]] = load i32, ptr [[T_VAR1]], align 4742// CHECK3-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]743// CHECK3-NEXT: store i32 [[ADD3]], ptr [[T_VAR1]], align 4744// CHECK3-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]745// CHECK3: omp.body.continue:746// CHECK3-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]747// CHECK3: omp.inner.for.inc:748// CHECK3-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4749// CHECK3-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP11]], 1750// CHECK3-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4751// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND]]752// CHECK3: omp.inner.for.end:753// CHECK3-NEXT: br label [[OMP_LOOP_EXIT:%.*]]754// CHECK3: omp.loop.exit:755// CHECK3-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])756// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 0757// CHECK3-NEXT: store ptr [[T_VAR1]], ptr [[TMP12]], align 4758// CHECK3-NEXT: [[TMP13:%.*]] = 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)759// CHECK3-NEXT: switch i32 [[TMP13]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [760// CHECK3-NEXT: i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]761// CHECK3-NEXT: i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]762// CHECK3-NEXT: ]763// CHECK3: .omp.reduction.case1:764// CHECK3-NEXT: [[TMP14:%.*]] = load i32, ptr [[TMP0]], align 4765// CHECK3-NEXT: [[TMP15:%.*]] = load i32, ptr [[T_VAR1]], align 4766// CHECK3-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP14]], [[TMP15]]767// CHECK3-NEXT: store i32 [[ADD5]], ptr [[TMP0]], align 4768// CHECK3-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)769// CHECK3-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]770// CHECK3: .omp.reduction.case2:771// CHECK3-NEXT: [[TMP16:%.*]] = load i32, ptr [[T_VAR1]], align 4772// CHECK3-NEXT: [[TMP17:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP16]] monotonic, align 4773// CHECK3-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)774// CHECK3-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]775// CHECK3: .omp.reduction.default:776// CHECK3-NEXT: ret void777//778//779// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp.reduction.reduction_func780// CHECK3-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3]] {781// CHECK3-NEXT: entry:782// CHECK3-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 4783// CHECK3-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 4784// CHECK3-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 4785// CHECK3-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 4786// CHECK3-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 4787// CHECK3-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 4788// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i32 0, i32 0789// CHECK3-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 4790// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i32 0, i32 0791// CHECK3-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 4792// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4793// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4794// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]795// CHECK3-NEXT: store i32 [[ADD]], ptr [[TMP7]], align 4796// CHECK3-NEXT: ret void797//798//799// CHECK9-LABEL: define {{[^@]+}}@main800// CHECK9-SAME: () #[[ATTR0:[0-9]+]] {801// CHECK9-NEXT: entry:802// CHECK9-NEXT: [[RETVAL:%.*]] = alloca i32, align 4803// CHECK9-NEXT: [[REF_TMP:%.*]] = alloca [[CLASS_ANON:%.*]], align 1804// CHECK9-NEXT: store i32 0, ptr [[RETVAL]], align 4805// CHECK9-NEXT: call void @"_ZZ4mainENK3$_0clEv"(ptr noundef nonnull align 1 dereferenceable(1) [[REF_TMP]])806// CHECK9-NEXT: ret i32 0807//808//809// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l45810// CHECK9-SAME: (i64 noundef [[SIVAR:%.*]]) #[[ATTR2:[0-9]+]] {811// CHECK9-NEXT: entry:812// CHECK9-NEXT: [[SIVAR_ADDR:%.*]] = alloca i64, align 8813// CHECK9-NEXT: store i64 [[SIVAR]], ptr [[SIVAR_ADDR]], align 8814// 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]])815// CHECK9-NEXT: ret void816//817//818// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l45.omp_outlined819// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR2]] {820// CHECK9-NEXT: entry:821// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8822// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8823// CHECK9-NEXT: [[SIVAR_ADDR:%.*]] = alloca ptr, align 8824// CHECK9-NEXT: [[SIVAR1:%.*]] = alloca i32, align 4825// CHECK9-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4826// CHECK9-NEXT: [[TMP:%.*]] = alloca i32, align 4827// CHECK9-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4828// CHECK9-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4829// CHECK9-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4830// CHECK9-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4831// CHECK9-NEXT: [[I:%.*]] = alloca i32, align 4832// CHECK9-NEXT: [[REF_TMP:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 8833// CHECK9-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 8834// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8835// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8836// CHECK9-NEXT: store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 8837// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 8838// CHECK9-NEXT: store i32 0, ptr [[SIVAR1]], align 4839// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4840// CHECK9-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 4841// CHECK9-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4842// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4843// CHECK9-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8844// CHECK9-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4845// 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)846// CHECK9-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4847// CHECK9-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1848// CHECK9-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]849// CHECK9: cond.true:850// CHECK9-NEXT: br label [[COND_END:%.*]]851// CHECK9: cond.false:852// CHECK9-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4853// CHECK9-NEXT: br label [[COND_END]]854// CHECK9: cond.end:855// CHECK9-NEXT: [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]856// CHECK9-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4857// CHECK9-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4858// CHECK9-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4859// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]860// CHECK9: omp.inner.for.cond:861// CHECK9-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4862// CHECK9-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4863// CHECK9-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]864// CHECK9-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]865// CHECK9: omp.inner.for.body:866// CHECK9-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4867// CHECK9-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1868// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]869// CHECK9-NEXT: store i32 [[ADD]], ptr [[I]], align 4870// CHECK9-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4871// CHECK9-NEXT: [[TMP10:%.*]] = load i32, ptr [[SIVAR1]], align 4872// CHECK9-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]873// CHECK9-NEXT: store i32 [[ADD3]], ptr [[SIVAR1]], align 4874// CHECK9-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 0875// CHECK9-NEXT: store ptr [[SIVAR1]], ptr [[TMP11]], align 8876// CHECK9-NEXT: call void @"_ZZZ4mainENK3$_0clEvENKUlvE_clEv"(ptr noundef nonnull align 8 dereferenceable(8) [[REF_TMP]])877// CHECK9-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]878// CHECK9: omp.body.continue:879// CHECK9-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]880// CHECK9: omp.inner.for.inc:881// CHECK9-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4882// CHECK9-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP12]], 1883// CHECK9-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4884// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND]]885// CHECK9: omp.inner.for.end:886// CHECK9-NEXT: br label [[OMP_LOOP_EXIT:%.*]]887// CHECK9: omp.loop.exit:888// CHECK9-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])889// CHECK9-NEXT: [[TMP13:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0890// CHECK9-NEXT: store ptr [[SIVAR1]], ptr [[TMP13]], align 8891// CHECK9-NEXT: [[TMP14:%.*]] = 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)892// CHECK9-NEXT: switch i32 [[TMP14]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [893// CHECK9-NEXT: i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]894// CHECK9-NEXT: i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]895// CHECK9-NEXT: ]896// CHECK9: .omp.reduction.case1:897// CHECK9-NEXT: [[TMP15:%.*]] = load i32, ptr [[TMP0]], align 4898// CHECK9-NEXT: [[TMP16:%.*]] = load i32, ptr [[SIVAR1]], align 4899// CHECK9-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP15]], [[TMP16]]900// CHECK9-NEXT: store i32 [[ADD5]], ptr [[TMP0]], align 4901// CHECK9-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)902// CHECK9-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]903// CHECK9: .omp.reduction.case2:904// CHECK9-NEXT: [[TMP17:%.*]] = load i32, ptr [[SIVAR1]], align 4905// CHECK9-NEXT: [[TMP18:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP17]] monotonic, align 4906// CHECK9-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)907// CHECK9-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]908// CHECK9: .omp.reduction.default:909// CHECK9-NEXT: ret void910//911//912// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l45.omp_outlined.omp.reduction.reduction_func913// CHECK9-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR4:[0-9]+]] {914// CHECK9-NEXT: entry:915// CHECK9-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 8916// CHECK9-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 8917// CHECK9-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 8918// CHECK9-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 8919// CHECK9-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 8920// CHECK9-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 8921// CHECK9-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i64 0, i64 0922// CHECK9-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 8923// CHECK9-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i64 0, i64 0924// CHECK9-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 8925// CHECK9-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4926// CHECK9-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4927// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]928// CHECK9-NEXT: store i32 [[ADD]], ptr [[TMP7]], align 4929// CHECK9-NEXT: ret void930//931