937 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 loop 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 loop reduction(+: sivar)47 for (int i = 0; i < 2; ++i) {48 49 // Skip global and bound tid vars50 51 52 53 // Skip global and bound tid vars, and prev lb and ub vars54 // skip loop vars55 56 57 sivar += i;58 59 [&]() {60 61 sivar += 4;62 63 }();64 }65 }();66 return 0;67#else68#pragma omp target69#pragma omp teams loop reduction(+: sivar)70 for (int i = 0; i < 2; ++i) {71 sivar += i;72 }73 return tmain<int>();74#endif75}76 77 78 79 80// Skip global and bound tid vars81 82 83// Skip global and bound tid vars, and prev lb and ub84// skip loop vars85 86 87 88 89// Skip global and bound tid vars90 91 92// Skip global and bound tid vars, and prev lb and ub vars93// skip loop vars94 95#endif96// CHECK1-LABEL: define {{[^@]+}}@main97// CHECK1-SAME: () #[[ATTR0:[0-9]+]] {98// CHECK1-NEXT: entry:99// CHECK1-NEXT: [[RETVAL:%.*]] = alloca i32, align 4100// CHECK1-NEXT: [[SIVAR_CASTED:%.*]] = alloca i64, align 8101// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8102// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8103// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8104// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4105// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8106// CHECK1-NEXT: store i32 0, ptr [[RETVAL]], align 4107// CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr @_ZZ4mainE5sivar, align 4108// CHECK1-NEXT: store i32 [[TMP0]], ptr [[SIVAR_CASTED]], align 4109// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[SIVAR_CASTED]], align 8110// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0111// CHECK1-NEXT: store i64 [[TMP1]], ptr [[TMP2]], align 8112// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0113// CHECK1-NEXT: store i64 [[TMP1]], ptr [[TMP3]], align 8114// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0115// CHECK1-NEXT: store ptr null, ptr [[TMP4]], align 8116// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0117// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0118// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0119// CHECK1-NEXT: store i32 3, ptr [[TMP7]], align 4120// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1121// CHECK1-NEXT: store i32 1, ptr [[TMP8]], align 4122// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2123// CHECK1-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 8124// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3125// CHECK1-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 8126// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4127// CHECK1-NEXT: store ptr @.offload_sizes, ptr [[TMP11]], align 8128// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5129// CHECK1-NEXT: store ptr @.offload_maptypes, ptr [[TMP12]], align 8130// CHECK1-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6131// CHECK1-NEXT: store ptr null, ptr [[TMP13]], align 8132// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7133// CHECK1-NEXT: store ptr null, ptr [[TMP14]], align 8134// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8135// CHECK1-NEXT: store i64 2, ptr [[TMP15]], align 8136// CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9137// CHECK1-NEXT: store i64 0, ptr [[TMP16]], align 8138// CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10139// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4140// CHECK1-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11141// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP18]], align 4142// CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12143// CHECK1-NEXT: store i32 0, ptr [[TMP19]], align 4144// 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_l68.region_id, ptr [[KERNEL_ARGS]])145// CHECK1-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0146// CHECK1-NEXT: br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]147// CHECK1: omp_offload.failed:148// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68(i64 [[TMP1]]) #[[ATTR2:[0-9]+]]149// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]150// CHECK1: omp_offload.cont:151// CHECK1-NEXT: [[CALL:%.*]] = call noundef signext i32 @_Z5tmainIiET_v()152// CHECK1-NEXT: ret i32 [[CALL]]153//154//155// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68156// CHECK1-SAME: (i64 noundef [[SIVAR:%.*]]) #[[ATTR1:[0-9]+]] {157// CHECK1-NEXT: entry:158// CHECK1-NEXT: [[SIVAR_ADDR:%.*]] = alloca i64, align 8159// CHECK1-NEXT: store i64 [[SIVAR]], ptr [[SIVAR_ADDR]], align 8160// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68.omp_outlined, ptr [[SIVAR_ADDR]])161// CHECK1-NEXT: ret void162//163//164// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68.omp_outlined165// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR1]] {166// CHECK1-NEXT: entry:167// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8168// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8169// CHECK1-NEXT: [[SIVAR_ADDR:%.*]] = alloca ptr, align 8170// CHECK1-NEXT: [[SIVAR1:%.*]] = alloca i32, align 4171// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4172// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4173// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4174// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4175// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4176// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4177// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4178// CHECK1-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 8179// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8180// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8181// CHECK1-NEXT: store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 8182// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 8183// CHECK1-NEXT: store i32 0, ptr [[SIVAR1]], align 4184// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4185// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_COMB_UB]], align 4186// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4187// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4188// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8189// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4190// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)191// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4192// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1193// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]194// CHECK1: cond.true:195// CHECK1-NEXT: br label [[COND_END:%.*]]196// CHECK1: cond.false:197// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4198// CHECK1-NEXT: br label [[COND_END]]199// CHECK1: cond.end:200// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]201// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4202// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4203// CHECK1-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4204// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]205// CHECK1: omp.inner.for.cond:206// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4207// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4208// CHECK1-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]209// CHECK1-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]210// CHECK1: omp.inner.for.body:211// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4212// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1213// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]214// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4215// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4216// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[SIVAR1]], align 4217// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]218// CHECK1-NEXT: store i32 [[ADD3]], ptr [[SIVAR1]], align 4219// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]220// CHECK1: omp.body.continue:221// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]222// CHECK1: omp.inner.for.inc:223// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4224// CHECK1-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP11]], 1225// CHECK1-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4226// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]227// CHECK1: omp.inner.for.end:228// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]229// CHECK1: omp.loop.exit:230// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])231// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0232// CHECK1-NEXT: store ptr [[SIVAR1]], ptr [[TMP12]], align 8233// CHECK1-NEXT: [[TMP13:%.*]] = call i32 @__kmpc_reduce_nowait(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_l68.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)234// CHECK1-NEXT: switch i32 [[TMP13]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [235// CHECK1-NEXT: i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]236// CHECK1-NEXT: i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]237// CHECK1-NEXT: ]238// CHECK1: .omp.reduction.case1:239// CHECK1-NEXT: [[TMP14:%.*]] = load i32, ptr [[TMP0]], align 4240// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[SIVAR1]], align 4241// CHECK1-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP14]], [[TMP15]]242// CHECK1-NEXT: store i32 [[ADD5]], ptr [[TMP0]], align 4243// CHECK1-NEXT: call void @__kmpc_end_reduce_nowait(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)244// CHECK1-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]245// CHECK1: .omp.reduction.case2:246// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[SIVAR1]], align 4247// CHECK1-NEXT: [[TMP17:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP16]] monotonic, align 4248// CHECK1-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]249// CHECK1: .omp.reduction.default:250// CHECK1-NEXT: ret void251//252//253// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68.omp_outlined.omp.reduction.reduction_func254// CHECK1-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3:[0-9]+]] {255// CHECK1-NEXT: entry:256// CHECK1-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 8257// CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 8258// CHECK1-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 8259// CHECK1-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 8260// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 8261// CHECK1-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 8262// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i64 0, i64 0263// CHECK1-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 8264// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i64 0, i64 0265// CHECK1-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 8266// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4267// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4268// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]269// CHECK1-NEXT: store i32 [[ADD]], ptr [[TMP7]], align 4270// CHECK1-NEXT: ret void271//272//273// CHECK1-LABEL: define {{[^@]+}}@_Z5tmainIiET_v274// CHECK1-SAME: () #[[ATTR5:[0-9]+]] comdat {275// CHECK1-NEXT: entry:276// CHECK1-NEXT: [[T_VAR:%.*]] = alloca i32, align 4277// CHECK1-NEXT: [[VEC:%.*]] = alloca [2 x i32], align 4278// CHECK1-NEXT: [[T_VAR_CASTED:%.*]] = alloca i64, align 8279// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8280// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8281// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8282// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4283// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8284// CHECK1-NEXT: store i32 0, ptr [[T_VAR]], align 4285// CHECK1-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[VEC]], ptr align 4 @__const._Z5tmainIiET_v.vec, i64 8, i1 false)286// CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr [[T_VAR]], align 4287// CHECK1-NEXT: store i32 [[TMP0]], ptr [[T_VAR_CASTED]], align 4288// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[T_VAR_CASTED]], align 8289// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0290// CHECK1-NEXT: store i64 [[TMP1]], ptr [[TMP2]], align 8291// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0292// CHECK1-NEXT: store i64 [[TMP1]], ptr [[TMP3]], align 8293// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0294// CHECK1-NEXT: store ptr null, ptr [[TMP4]], align 8295// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0296// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0297// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0298// CHECK1-NEXT: store i32 3, ptr [[TMP7]], align 4299// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1300// CHECK1-NEXT: store i32 1, ptr [[TMP8]], align 4301// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2302// CHECK1-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 8303// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3304// CHECK1-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 8305// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4306// CHECK1-NEXT: store ptr @.offload_sizes.1, ptr [[TMP11]], align 8307// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5308// CHECK1-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP12]], align 8309// CHECK1-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6310// CHECK1-NEXT: store ptr null, ptr [[TMP13]], align 8311// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7312// CHECK1-NEXT: store ptr null, ptr [[TMP14]], align 8313// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8314// CHECK1-NEXT: store i64 2, ptr [[TMP15]], align 8315// CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9316// CHECK1-NEXT: store i64 0, ptr [[TMP16]], align 8317// CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10318// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4319// CHECK1-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11320// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP18]], align 4321// CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12322// CHECK1-NEXT: store i32 0, ptr [[TMP19]], align 4323// 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]])324// CHECK1-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0325// CHECK1-NEXT: br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]326// CHECK1: omp_offload.failed:327// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32(i64 [[TMP1]]) #[[ATTR2]]328// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]329// CHECK1: omp_offload.cont:330// CHECK1-NEXT: ret i32 0331//332//333// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32334// CHECK1-SAME: (i64 noundef [[T_VAR:%.*]]) #[[ATTR1]] {335// CHECK1-NEXT: entry:336// CHECK1-NEXT: [[T_VAR_ADDR:%.*]] = alloca i64, align 8337// CHECK1-NEXT: store i64 [[T_VAR]], ptr [[T_VAR_ADDR]], align 8338// 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]])339// CHECK1-NEXT: ret void340//341//342// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined343// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[T_VAR:%.*]]) #[[ATTR1]] {344// CHECK1-NEXT: entry:345// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8346// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8347// CHECK1-NEXT: [[T_VAR_ADDR:%.*]] = alloca ptr, align 8348// CHECK1-NEXT: [[T_VAR1:%.*]] = alloca i32, align 4349// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4350// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4351// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4352// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4353// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4354// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4355// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4356// CHECK1-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 8357// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8358// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8359// CHECK1-NEXT: store ptr [[T_VAR]], ptr [[T_VAR_ADDR]], align 8360// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[T_VAR_ADDR]], align 8361// CHECK1-NEXT: store i32 0, ptr [[T_VAR1]], align 4362// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4363// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_COMB_UB]], align 4364// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4365// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4366// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8367// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4368// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)369// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4370// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1371// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]372// CHECK1: cond.true:373// CHECK1-NEXT: br label [[COND_END:%.*]]374// CHECK1: cond.false:375// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4376// CHECK1-NEXT: br label [[COND_END]]377// CHECK1: cond.end:378// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]379// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4380// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4381// CHECK1-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4382// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]383// CHECK1: omp.inner.for.cond:384// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4385// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4386// CHECK1-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]387// CHECK1-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]388// CHECK1: omp.inner.for.body:389// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4390// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1391// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]392// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4393// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4394// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[T_VAR1]], align 4395// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]396// CHECK1-NEXT: store i32 [[ADD3]], ptr [[T_VAR1]], align 4397// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]398// CHECK1: omp.body.continue:399// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]400// CHECK1: omp.inner.for.inc:401// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4402// CHECK1-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP11]], 1403// CHECK1-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4404// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]405// CHECK1: omp.inner.for.end:406// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]407// CHECK1: omp.loop.exit:408// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])409// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0410// CHECK1-NEXT: store ptr [[T_VAR1]], ptr [[TMP12]], align 8411// CHECK1-NEXT: [[TMP13:%.*]] = call i32 @__kmpc_reduce_nowait(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)412// CHECK1-NEXT: switch i32 [[TMP13]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [413// CHECK1-NEXT: i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]414// CHECK1-NEXT: i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]415// CHECK1-NEXT: ]416// CHECK1: .omp.reduction.case1:417// CHECK1-NEXT: [[TMP14:%.*]] = load i32, ptr [[TMP0]], align 4418// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[T_VAR1]], align 4419// CHECK1-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP14]], [[TMP15]]420// CHECK1-NEXT: store i32 [[ADD5]], ptr [[TMP0]], align 4421// CHECK1-NEXT: call void @__kmpc_end_reduce_nowait(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)422// CHECK1-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]423// CHECK1: .omp.reduction.case2:424// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[T_VAR1]], align 4425// CHECK1-NEXT: [[TMP17:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP16]] monotonic, align 4426// CHECK1-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]427// CHECK1: .omp.reduction.default:428// CHECK1-NEXT: ret void429//430//431// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp.reduction.reduction_func432// CHECK1-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3]] {433// CHECK1-NEXT: entry:434// CHECK1-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 8435// CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 8436// CHECK1-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 8437// CHECK1-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 8438// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 8439// CHECK1-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 8440// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i64 0, i64 0441// CHECK1-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 8442// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i64 0, i64 0443// CHECK1-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 8444// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4445// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4446// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]447// CHECK1-NEXT: store i32 [[ADD]], ptr [[TMP7]], align 4448// CHECK1-NEXT: ret void449//450//451// CHECK3-LABEL: define {{[^@]+}}@main452// CHECK3-SAME: () #[[ATTR0:[0-9]+]] {453// CHECK3-NEXT: entry:454// CHECK3-NEXT: [[RETVAL:%.*]] = alloca i32, align 4455// CHECK3-NEXT: [[SIVAR_CASTED:%.*]] = alloca i32, align 4456// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4457// CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4458// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4459// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4460// CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8461// CHECK3-NEXT: store i32 0, ptr [[RETVAL]], align 4462// CHECK3-NEXT: [[TMP0:%.*]] = load i32, ptr @_ZZ4mainE5sivar, align 4463// CHECK3-NEXT: store i32 [[TMP0]], ptr [[SIVAR_CASTED]], align 4464// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[SIVAR_CASTED]], align 4465// CHECK3-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0466// CHECK3-NEXT: store i32 [[TMP1]], ptr [[TMP2]], align 4467// CHECK3-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0468// CHECK3-NEXT: store i32 [[TMP1]], ptr [[TMP3]], align 4469// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0470// CHECK3-NEXT: store ptr null, ptr [[TMP4]], align 4471// CHECK3-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0472// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0473// CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0474// CHECK3-NEXT: store i32 3, ptr [[TMP7]], align 4475// CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1476// CHECK3-NEXT: store i32 1, ptr [[TMP8]], align 4477// CHECK3-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2478// CHECK3-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 4479// CHECK3-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3480// CHECK3-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 4481// CHECK3-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4482// CHECK3-NEXT: store ptr @.offload_sizes, ptr [[TMP11]], align 4483// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5484// CHECK3-NEXT: store ptr @.offload_maptypes, ptr [[TMP12]], align 4485// CHECK3-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6486// CHECK3-NEXT: store ptr null, ptr [[TMP13]], align 4487// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7488// CHECK3-NEXT: store ptr null, ptr [[TMP14]], align 4489// CHECK3-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8490// CHECK3-NEXT: store i64 2, ptr [[TMP15]], align 8491// CHECK3-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9492// CHECK3-NEXT: store i64 0, ptr [[TMP16]], align 8493// CHECK3-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10494// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4495// CHECK3-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11496// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP18]], align 4497// CHECK3-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12498// CHECK3-NEXT: store i32 0, ptr [[TMP19]], align 4499// 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_l68.region_id, ptr [[KERNEL_ARGS]])500// CHECK3-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0501// CHECK3-NEXT: br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]502// CHECK3: omp_offload.failed:503// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68(i32 [[TMP1]]) #[[ATTR2:[0-9]+]]504// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]]505// CHECK3: omp_offload.cont:506// CHECK3-NEXT: [[CALL:%.*]] = call noundef i32 @_Z5tmainIiET_v()507// CHECK3-NEXT: ret i32 [[CALL]]508//509//510// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68511// CHECK3-SAME: (i32 noundef [[SIVAR:%.*]]) #[[ATTR1:[0-9]+]] {512// CHECK3-NEXT: entry:513// CHECK3-NEXT: [[SIVAR_ADDR:%.*]] = alloca i32, align 4514// CHECK3-NEXT: store i32 [[SIVAR]], ptr [[SIVAR_ADDR]], align 4515// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68.omp_outlined, ptr [[SIVAR_ADDR]])516// CHECK3-NEXT: ret void517//518//519// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68.omp_outlined520// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR1]] {521// CHECK3-NEXT: entry:522// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4523// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4524// CHECK3-NEXT: [[SIVAR_ADDR:%.*]] = alloca ptr, align 4525// CHECK3-NEXT: [[SIVAR1:%.*]] = alloca i32, align 4526// CHECK3-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4527// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4528// CHECK3-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4529// CHECK3-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4530// CHECK3-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4531// CHECK3-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4532// CHECK3-NEXT: [[I:%.*]] = alloca i32, align 4533// CHECK3-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 4534// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4535// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4536// CHECK3-NEXT: store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 4537// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 4538// CHECK3-NEXT: store i32 0, ptr [[SIVAR1]], align 4539// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4540// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_COMB_UB]], align 4541// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4542// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4543// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4544// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4545// CHECK3-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)546// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4547// CHECK3-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1548// CHECK3-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]549// CHECK3: cond.true:550// CHECK3-NEXT: br label [[COND_END:%.*]]551// CHECK3: cond.false:552// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4553// CHECK3-NEXT: br label [[COND_END]]554// CHECK3: cond.end:555// CHECK3-NEXT: [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]556// CHECK3-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4557// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4558// CHECK3-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4559// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]560// CHECK3: omp.inner.for.cond:561// CHECK3-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4562// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4563// CHECK3-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]564// CHECK3-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]565// CHECK3: omp.inner.for.body:566// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4567// CHECK3-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1568// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]569// CHECK3-NEXT: store i32 [[ADD]], ptr [[I]], align 4570// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4571// CHECK3-NEXT: [[TMP10:%.*]] = load i32, ptr [[SIVAR1]], align 4572// CHECK3-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]573// CHECK3-NEXT: store i32 [[ADD3]], ptr [[SIVAR1]], align 4574// CHECK3-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]575// CHECK3: omp.body.continue:576// CHECK3-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]577// CHECK3: omp.inner.for.inc:578// CHECK3-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4579// CHECK3-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP11]], 1580// CHECK3-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4581// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND]]582// CHECK3: omp.inner.for.end:583// CHECK3-NEXT: br label [[OMP_LOOP_EXIT:%.*]]584// CHECK3: omp.loop.exit:585// CHECK3-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])586// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 0587// CHECK3-NEXT: store ptr [[SIVAR1]], ptr [[TMP12]], align 4588// CHECK3-NEXT: [[TMP13:%.*]] = call i32 @__kmpc_reduce_nowait(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_l68.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)589// CHECK3-NEXT: switch i32 [[TMP13]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [590// CHECK3-NEXT: i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]591// CHECK3-NEXT: i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]592// CHECK3-NEXT: ]593// CHECK3: .omp.reduction.case1:594// CHECK3-NEXT: [[TMP14:%.*]] = load i32, ptr [[TMP0]], align 4595// CHECK3-NEXT: [[TMP15:%.*]] = load i32, ptr [[SIVAR1]], align 4596// CHECK3-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP14]], [[TMP15]]597// CHECK3-NEXT: store i32 [[ADD5]], ptr [[TMP0]], align 4598// CHECK3-NEXT: call void @__kmpc_end_reduce_nowait(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)599// CHECK3-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]600// CHECK3: .omp.reduction.case2:601// CHECK3-NEXT: [[TMP16:%.*]] = load i32, ptr [[SIVAR1]], align 4602// CHECK3-NEXT: [[TMP17:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP16]] monotonic, align 4603// CHECK3-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]604// CHECK3: .omp.reduction.default:605// CHECK3-NEXT: ret void606//607//608// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68.omp_outlined.omp.reduction.reduction_func609// CHECK3-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3:[0-9]+]] {610// CHECK3-NEXT: entry:611// CHECK3-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 4612// CHECK3-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 4613// CHECK3-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 4614// CHECK3-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 4615// CHECK3-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 4616// CHECK3-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 4617// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i32 0, i32 0618// CHECK3-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 4619// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i32 0, i32 0620// CHECK3-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 4621// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4622// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4623// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]624// CHECK3-NEXT: store i32 [[ADD]], ptr [[TMP7]], align 4625// CHECK3-NEXT: ret void626//627//628// CHECK3-LABEL: define {{[^@]+}}@_Z5tmainIiET_v629// CHECK3-SAME: () #[[ATTR5:[0-9]+]] comdat {630// CHECK3-NEXT: entry:631// CHECK3-NEXT: [[T_VAR:%.*]] = alloca i32, align 4632// CHECK3-NEXT: [[VEC:%.*]] = alloca [2 x i32], align 4633// CHECK3-NEXT: [[T_VAR_CASTED:%.*]] = alloca i32, align 4634// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4635// CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4636// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4637// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4638// CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8639// CHECK3-NEXT: store i32 0, ptr [[T_VAR]], align 4640// CHECK3-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[VEC]], ptr align 4 @__const._Z5tmainIiET_v.vec, i32 8, i1 false)641// CHECK3-NEXT: [[TMP0:%.*]] = load i32, ptr [[T_VAR]], align 4642// CHECK3-NEXT: store i32 [[TMP0]], ptr [[T_VAR_CASTED]], align 4643// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[T_VAR_CASTED]], align 4644// CHECK3-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0645// CHECK3-NEXT: store i32 [[TMP1]], ptr [[TMP2]], align 4646// CHECK3-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0647// CHECK3-NEXT: store i32 [[TMP1]], ptr [[TMP3]], align 4648// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0649// CHECK3-NEXT: store ptr null, ptr [[TMP4]], align 4650// CHECK3-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0651// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0652// CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0653// CHECK3-NEXT: store i32 3, ptr [[TMP7]], align 4654// CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1655// CHECK3-NEXT: store i32 1, ptr [[TMP8]], align 4656// CHECK3-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2657// CHECK3-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 4658// CHECK3-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3659// CHECK3-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 4660// CHECK3-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4661// CHECK3-NEXT: store ptr @.offload_sizes.1, ptr [[TMP11]], align 4662// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5663// CHECK3-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP12]], align 4664// CHECK3-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6665// CHECK3-NEXT: store ptr null, ptr [[TMP13]], align 4666// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7667// CHECK3-NEXT: store ptr null, ptr [[TMP14]], align 4668// CHECK3-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8669// CHECK3-NEXT: store i64 2, ptr [[TMP15]], align 8670// CHECK3-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9671// CHECK3-NEXT: store i64 0, ptr [[TMP16]], align 8672// CHECK3-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10673// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4674// CHECK3-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11675// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP18]], align 4676// CHECK3-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12677// CHECK3-NEXT: store i32 0, ptr [[TMP19]], align 4678// 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]])679// CHECK3-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0680// CHECK3-NEXT: br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]681// CHECK3: omp_offload.failed:682// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32(i32 [[TMP1]]) #[[ATTR2]]683// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]]684// CHECK3: omp_offload.cont:685// CHECK3-NEXT: ret i32 0686//687//688// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32689// CHECK3-SAME: (i32 noundef [[T_VAR:%.*]]) #[[ATTR1]] {690// CHECK3-NEXT: entry:691// CHECK3-NEXT: [[T_VAR_ADDR:%.*]] = alloca i32, align 4692// CHECK3-NEXT: store i32 [[T_VAR]], ptr [[T_VAR_ADDR]], align 4693// 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]])694// CHECK3-NEXT: ret void695//696//697// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined698// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[T_VAR:%.*]]) #[[ATTR1]] {699// CHECK3-NEXT: entry:700// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4701// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4702// CHECK3-NEXT: [[T_VAR_ADDR:%.*]] = alloca ptr, align 4703// CHECK3-NEXT: [[T_VAR1:%.*]] = alloca i32, align 4704// CHECK3-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4705// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4706// CHECK3-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4707// CHECK3-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4708// CHECK3-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4709// CHECK3-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4710// CHECK3-NEXT: [[I:%.*]] = alloca i32, align 4711// CHECK3-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 4712// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4713// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4714// CHECK3-NEXT: store ptr [[T_VAR]], ptr [[T_VAR_ADDR]], align 4715// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[T_VAR_ADDR]], align 4716// CHECK3-NEXT: store i32 0, ptr [[T_VAR1]], align 4717// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4718// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_COMB_UB]], align 4719// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4720// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4721// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4722// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4723// CHECK3-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)724// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4725// CHECK3-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1726// CHECK3-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]727// CHECK3: cond.true:728// CHECK3-NEXT: br label [[COND_END:%.*]]729// CHECK3: cond.false:730// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4731// CHECK3-NEXT: br label [[COND_END]]732// CHECK3: cond.end:733// CHECK3-NEXT: [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]734// CHECK3-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4735// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4736// CHECK3-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4737// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]738// CHECK3: omp.inner.for.cond:739// CHECK3-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4740// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4741// CHECK3-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]742// CHECK3-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]743// CHECK3: omp.inner.for.body:744// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4745// CHECK3-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1746// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]747// CHECK3-NEXT: store i32 [[ADD]], ptr [[I]], align 4748// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4749// CHECK3-NEXT: [[TMP10:%.*]] = load i32, ptr [[T_VAR1]], align 4750// CHECK3-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]751// CHECK3-NEXT: store i32 [[ADD3]], ptr [[T_VAR1]], align 4752// CHECK3-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]753// CHECK3: omp.body.continue:754// CHECK3-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]755// CHECK3: omp.inner.for.inc:756// CHECK3-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4757// CHECK3-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP11]], 1758// CHECK3-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4759// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND]]760// CHECK3: omp.inner.for.end:761// CHECK3-NEXT: br label [[OMP_LOOP_EXIT:%.*]]762// CHECK3: omp.loop.exit:763// CHECK3-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])764// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 0765// CHECK3-NEXT: store ptr [[T_VAR1]], ptr [[TMP12]], align 4766// CHECK3-NEXT: [[TMP13:%.*]] = call i32 @__kmpc_reduce_nowait(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)767// CHECK3-NEXT: switch i32 [[TMP13]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [768// CHECK3-NEXT: i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]769// CHECK3-NEXT: i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]770// CHECK3-NEXT: ]771// CHECK3: .omp.reduction.case1:772// CHECK3-NEXT: [[TMP14:%.*]] = load i32, ptr [[TMP0]], align 4773// CHECK3-NEXT: [[TMP15:%.*]] = load i32, ptr [[T_VAR1]], align 4774// CHECK3-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP14]], [[TMP15]]775// CHECK3-NEXT: store i32 [[ADD5]], ptr [[TMP0]], align 4776// CHECK3-NEXT: call void @__kmpc_end_reduce_nowait(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)777// CHECK3-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]778// CHECK3: .omp.reduction.case2:779// CHECK3-NEXT: [[TMP16:%.*]] = load i32, ptr [[T_VAR1]], align 4780// CHECK3-NEXT: [[TMP17:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP16]] monotonic, align 4781// CHECK3-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]782// CHECK3: .omp.reduction.default:783// CHECK3-NEXT: ret void784//785//786// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp.reduction.reduction_func787// CHECK3-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3]] {788// CHECK3-NEXT: entry:789// CHECK3-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 4790// CHECK3-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 4791// CHECK3-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 4792// CHECK3-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 4793// CHECK3-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 4794// CHECK3-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 4795// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i32 0, i32 0796// CHECK3-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 4797// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i32 0, i32 0798// CHECK3-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 4799// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4800// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4801// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]802// CHECK3-NEXT: store i32 [[ADD]], ptr [[TMP7]], align 4803// CHECK3-NEXT: ret void804//805//806// CHECK9-LABEL: define {{[^@]+}}@main807// CHECK9-SAME: () #[[ATTR0:[0-9]+]] {808// CHECK9-NEXT: entry:809// CHECK9-NEXT: [[RETVAL:%.*]] = alloca i32, align 4810// CHECK9-NEXT: [[REF_TMP:%.*]] = alloca [[CLASS_ANON:%.*]], align 1811// CHECK9-NEXT: store i32 0, ptr [[RETVAL]], align 4812// CHECK9-NEXT: call void @"_ZZ4mainENK3$_0clEv"(ptr noundef nonnull align 1 dereferenceable(1) [[REF_TMP]])813// CHECK9-NEXT: ret i32 0814//815//816// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l45817// CHECK9-SAME: (i64 noundef [[SIVAR:%.*]]) #[[ATTR2:[0-9]+]] {818// CHECK9-NEXT: entry:819// CHECK9-NEXT: [[SIVAR_ADDR:%.*]] = alloca i64, align 8820// CHECK9-NEXT: store i64 [[SIVAR]], ptr [[SIVAR_ADDR]], align 8821// 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]])822// CHECK9-NEXT: ret void823//824//825// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l45.omp_outlined826// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR2]] {827// CHECK9-NEXT: entry:828// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8829// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8830// CHECK9-NEXT: [[SIVAR_ADDR:%.*]] = alloca ptr, align 8831// CHECK9-NEXT: [[SIVAR1:%.*]] = alloca i32, align 4832// CHECK9-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4833// CHECK9-NEXT: [[TMP:%.*]] = alloca i32, align 4834// CHECK9-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4835// CHECK9-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4836// CHECK9-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4837// CHECK9-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4838// CHECK9-NEXT: [[I:%.*]] = alloca i32, align 4839// CHECK9-NEXT: [[REF_TMP:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 8840// CHECK9-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 8841// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8842// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8843// CHECK9-NEXT: store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 8844// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 8845// CHECK9-NEXT: store i32 0, ptr [[SIVAR1]], align 4846// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4847// CHECK9-NEXT: store i32 1, ptr [[DOTOMP_COMB_UB]], align 4848// CHECK9-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4849// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4850// CHECK9-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8851// CHECK9-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4852// CHECK9-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)853// CHECK9-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4854// CHECK9-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1855// CHECK9-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]856// CHECK9: cond.true:857// CHECK9-NEXT: br label [[COND_END:%.*]]858// CHECK9: cond.false:859// CHECK9-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4860// CHECK9-NEXT: br label [[COND_END]]861// CHECK9: cond.end:862// CHECK9-NEXT: [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]863// CHECK9-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4864// CHECK9-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4865// CHECK9-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4866// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]867// CHECK9: omp.inner.for.cond:868// CHECK9-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4869// CHECK9-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4870// CHECK9-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]871// CHECK9-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]872// CHECK9: omp.inner.for.body:873// CHECK9-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4874// CHECK9-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1875// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]876// CHECK9-NEXT: store i32 [[ADD]], ptr [[I]], align 4877// CHECK9-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4878// CHECK9-NEXT: [[TMP10:%.*]] = load i32, ptr [[SIVAR1]], align 4879// CHECK9-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]880// CHECK9-NEXT: store i32 [[ADD3]], ptr [[SIVAR1]], align 4881// CHECK9-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 0882// CHECK9-NEXT: store ptr [[SIVAR1]], ptr [[TMP11]], align 8883// CHECK9-NEXT: call void @"_ZZZ4mainENK3$_0clEvENKUlvE_clEv"(ptr noundef nonnull align 8 dereferenceable(8) [[REF_TMP]])884// CHECK9-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]885// CHECK9: omp.body.continue:886// CHECK9-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]887// CHECK9: omp.inner.for.inc:888// CHECK9-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4889// CHECK9-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP12]], 1890// CHECK9-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4891// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND]]892// CHECK9: omp.inner.for.end:893// CHECK9-NEXT: br label [[OMP_LOOP_EXIT:%.*]]894// CHECK9: omp.loop.exit:895// CHECK9-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])896// CHECK9-NEXT: [[TMP13:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0897// CHECK9-NEXT: store ptr [[SIVAR1]], ptr [[TMP13]], align 8898// CHECK9-NEXT: [[TMP14:%.*]] = call i32 @__kmpc_reduce_nowait(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)899// CHECK9-NEXT: switch i32 [[TMP14]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [900// CHECK9-NEXT: i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]901// CHECK9-NEXT: i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]902// CHECK9-NEXT: ]903// CHECK9: .omp.reduction.case1:904// CHECK9-NEXT: [[TMP15:%.*]] = load i32, ptr [[TMP0]], align 4905// CHECK9-NEXT: [[TMP16:%.*]] = load i32, ptr [[SIVAR1]], align 4906// CHECK9-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP15]], [[TMP16]]907// CHECK9-NEXT: store i32 [[ADD5]], ptr [[TMP0]], align 4908// CHECK9-NEXT: call void @__kmpc_end_reduce_nowait(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)909// CHECK9-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]910// CHECK9: .omp.reduction.case2:911// CHECK9-NEXT: [[TMP17:%.*]] = load i32, ptr [[SIVAR1]], align 4912// CHECK9-NEXT: [[TMP18:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP17]] monotonic, align 4913// CHECK9-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]914// CHECK9: .omp.reduction.default:915// CHECK9-NEXT: ret void916//917//918// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l45.omp_outlined.omp.reduction.reduction_func919// CHECK9-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR4:[0-9]+]] {920// CHECK9-NEXT: entry:921// CHECK9-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 8922// CHECK9-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 8923// CHECK9-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 8924// CHECK9-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 8925// CHECK9-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 8926// CHECK9-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 8927// CHECK9-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i64 0, i64 0928// CHECK9-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 8929// CHECK9-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i64 0, i64 0930// CHECK9-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 8931// CHECK9-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4932// CHECK9-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4933// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]934// CHECK9-NEXT: store i32 [[ADD]], ptr [[TMP7]], align 4935// CHECK9-NEXT: ret void936//937