1159 lines · cpp
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _2// RUN: %clang_cc1 -DCHECK -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK13// RUN: %clang_cc1 -DCHECK -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s4// RUN: %clang_cc1 -DCHECK -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK15// RUN: %clang_cc1 -DCHECK -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK36// RUN: %clang_cc1 -DCHECK -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s7// RUN: %clang_cc1 -DCHECK -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK38 9// RUN: %clang_cc1 -DCHECK -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK510// RUN: %clang_cc1 -DCHECK -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s11// RUN: %clang_cc1 -DCHECK -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK512// RUN: %clang_cc1 -DCHECK -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK713// RUN: %clang_cc1 -DCHECK -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s14// RUN: %clang_cc1 -DCHECK -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK715 16// RUN: %clang_cc1 -DLAMBDA -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK917// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s18// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK919 20// RUN: %clang_cc1 -DLAMBDA -verify -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK1121// RUN: %clang_cc1 -DLAMBDA -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s22// RUN: %clang_cc1 -DLAMBDA -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK1123 24// expected-no-diagnostics25#ifndef HEADER26#define HEADER27 28template <typename T>29T tmain() {30 T t_var = T();31 T vec[] = {1, 2};32#pragma omp target teams distribute simd reduction(+: t_var)33 for (int i = 0; i < 2; ++i) {34 t_var += (T) i;35 }36 return T();37}38 39int main() {40 static int sivar;41#ifdef LAMBDA42 43 [&]() {44#pragma omp target teams distribute simd reduction(+: sivar)45 for (int i = 0; i < 2; ++i) {46 47 // Skip global and bound tid vars48 49 sivar += i;50 51 [&]() {52 53 sivar += 4;54 55 }();56 }57 }();58 return 0;59#else60#pragma omp target teams distribute simd reduction(+: sivar)61 for (int i = 0; i < 2; ++i) {62 sivar += i;63 }64 return tmain<int>();65#endif66}67 68 69 70 71// Skip global and bound tid vars72 73 74 75 76 77// Skip global and bound tid vars78 79 80#endif81// CHECK1-LABEL: define {{[^@]+}}@main82// CHECK1-SAME: () #[[ATTR0:[0-9]+]] {83// CHECK1-NEXT: entry:84// CHECK1-NEXT: [[RETVAL:%.*]] = alloca i32, align 485// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 886// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 887// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 888// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 489// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 890// CHECK1-NEXT: store i32 0, ptr [[RETVAL]], align 491// CHECK1-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 092// CHECK1-NEXT: store ptr @_ZZ4mainE5sivar, ptr [[TMP0]], align 893// CHECK1-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 094// CHECK1-NEXT: store ptr @_ZZ4mainE5sivar, ptr [[TMP1]], align 895// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 096// CHECK1-NEXT: store ptr null, ptr [[TMP2]], align 897// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 098// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 099// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0100// CHECK1-NEXT: store i32 3, ptr [[TMP5]], align 4101// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1102// CHECK1-NEXT: store i32 1, ptr [[TMP6]], align 4103// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2104// CHECK1-NEXT: store ptr [[TMP3]], ptr [[TMP7]], align 8105// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3106// CHECK1-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 8107// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4108// CHECK1-NEXT: store ptr @.offload_sizes, ptr [[TMP9]], align 8109// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5110// CHECK1-NEXT: store ptr @.offload_maptypes, ptr [[TMP10]], align 8111// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6112// CHECK1-NEXT: store ptr null, ptr [[TMP11]], align 8113// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7114// CHECK1-NEXT: store ptr null, ptr [[TMP12]], align 8115// CHECK1-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8116// CHECK1-NEXT: store i64 2, ptr [[TMP13]], align 8117// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9118// CHECK1-NEXT: store i64 0, ptr [[TMP14]], align 8119// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10120// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4121// CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11122// CHECK1-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP16]], align 4123// CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12124// CHECK1-NEXT: store i32 0, ptr [[TMP17]], align 4125// CHECK1-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3:[0-9]+]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l60.region_id, ptr [[KERNEL_ARGS]])126// CHECK1-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0127// CHECK1-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]128// CHECK1: omp_offload.failed:129// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l60(ptr @_ZZ4mainE5sivar) #[[ATTR2:[0-9]+]]130// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]131// CHECK1: omp_offload.cont:132// CHECK1-NEXT: [[CALL:%.*]] = call noundef signext i32 @_Z5tmainIiET_v()133// CHECK1-NEXT: ret i32 [[CALL]]134//135//136// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l60137// CHECK1-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR1:[0-9]+]] {138// CHECK1-NEXT: entry:139// CHECK1-NEXT: [[SIVAR_ADDR:%.*]] = alloca ptr, align 8140// CHECK1-NEXT: store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 8141// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 8142// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l60.omp_outlined, ptr [[TMP0]])143// CHECK1-NEXT: ret void144//145//146// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l60.omp_outlined147// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR1]] {148// CHECK1-NEXT: entry:149// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8150// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8151// CHECK1-NEXT: [[SIVAR_ADDR:%.*]] = alloca ptr, align 8152// CHECK1-NEXT: [[SIVAR1:%.*]] = alloca i32, align 4153// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4154// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4155// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4156// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4157// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4158// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4159// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4160// CHECK1-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 8161// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8162// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8163// CHECK1-NEXT: store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 8164// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 8165// CHECK1-NEXT: store i32 0, ptr [[SIVAR1]], align 4166// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4167// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 4168// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4169// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4170// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8171// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4172// 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)173// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4174// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1175// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]176// CHECK1: cond.true:177// CHECK1-NEXT: br label [[COND_END:%.*]]178// CHECK1: cond.false:179// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4180// CHECK1-NEXT: br label [[COND_END]]181// CHECK1: cond.end:182// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]183// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4184// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4185// CHECK1-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4186// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]187// CHECK1: omp.inner.for.cond:188// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP5:![0-9]+]]189// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP5]]190// CHECK1-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]191// CHECK1-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]192// CHECK1: omp.inner.for.body:193// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP5]]194// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1195// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]196// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP5]]197// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP5]]198// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[SIVAR1]], align 4, !llvm.access.group [[ACC_GRP5]]199// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]200// CHECK1-NEXT: store i32 [[ADD3]], ptr [[SIVAR1]], align 4, !llvm.access.group [[ACC_GRP5]]201// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]202// CHECK1: omp.body.continue:203// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]204// CHECK1: omp.inner.for.inc:205// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP5]]206// CHECK1-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP11]], 1207// CHECK1-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP5]]208// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP6:![0-9]+]]209// CHECK1: omp.inner.for.end:210// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]211// CHECK1: omp.loop.exit:212// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])213// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 4214// CHECK1-NEXT: [[TMP13:%.*]] = icmp ne i32 [[TMP12]], 0215// CHECK1-NEXT: br i1 [[TMP13]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]216// CHECK1: .omp.final.then:217// CHECK1-NEXT: store i32 2, ptr [[I]], align 4218// CHECK1-NEXT: br label [[DOTOMP_FINAL_DONE]]219// CHECK1: .omp.final.done:220// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0221// CHECK1-NEXT: store ptr [[SIVAR1]], ptr [[TMP14]], align 8222// CHECK1-NEXT: [[TMP15:%.*]] = call i32 @__kmpc_reduce(ptr @[[GLOB2:[0-9]+]], i32 [[TMP2]], i32 1, i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l60.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)223// CHECK1-NEXT: switch i32 [[TMP15]], 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: [[TMP16:%.*]] = load i32, ptr [[TMP0]], align 4229// CHECK1-NEXT: [[TMP17:%.*]] = load i32, ptr [[SIVAR1]], align 4230// CHECK1-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP16]], [[TMP17]]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: [[TMP18:%.*]] = load i32, ptr [[SIVAR1]], align 4236// CHECK1-NEXT: [[TMP19:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP18]] 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_l60.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: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8269// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8270// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8271// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4272// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8273// CHECK1-NEXT: store i32 0, ptr [[T_VAR]], align 4274// CHECK1-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[VEC]], ptr align 4 @__const._Z5tmainIiET_v.vec, i64 8, i1 false)275// CHECK1-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0276// CHECK1-NEXT: store ptr [[T_VAR]], ptr [[TMP0]], align 8277// CHECK1-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0278// CHECK1-NEXT: store ptr [[T_VAR]], ptr [[TMP1]], align 8279// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0280// CHECK1-NEXT: store ptr null, ptr [[TMP2]], align 8281// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0282// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0283// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0284// CHECK1-NEXT: store i32 3, ptr [[TMP5]], align 4285// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1286// CHECK1-NEXT: store i32 1, ptr [[TMP6]], align 4287// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2288// CHECK1-NEXT: store ptr [[TMP3]], ptr [[TMP7]], align 8289// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3290// CHECK1-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 8291// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4292// CHECK1-NEXT: store ptr @.offload_sizes.1, ptr [[TMP9]], align 8293// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5294// CHECK1-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP10]], align 8295// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6296// CHECK1-NEXT: store ptr null, ptr [[TMP11]], align 8297// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7298// CHECK1-NEXT: store ptr null, ptr [[TMP12]], align 8299// CHECK1-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8300// CHECK1-NEXT: store i64 2, ptr [[TMP13]], align 8301// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9302// CHECK1-NEXT: store i64 0, ptr [[TMP14]], align 8303// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10304// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4305// CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11306// CHECK1-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP16]], align 4307// CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12308// CHECK1-NEXT: store i32 0, ptr [[TMP17]], align 4309// CHECK1-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.region_id, ptr [[KERNEL_ARGS]])310// CHECK1-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0311// CHECK1-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]312// CHECK1: omp_offload.failed:313// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32(ptr [[T_VAR]]) #[[ATTR2]]314// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]315// CHECK1: omp_offload.cont:316// CHECK1-NEXT: ret i32 0317//318//319// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32320// CHECK1-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[T_VAR:%.*]]) #[[ATTR1]] {321// CHECK1-NEXT: entry:322// CHECK1-NEXT: [[T_VAR_ADDR:%.*]] = alloca ptr, align 8323// CHECK1-NEXT: store ptr [[T_VAR]], ptr [[T_VAR_ADDR]], align 8324// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[T_VAR_ADDR]], align 8325// 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 [[TMP0]])326// CHECK1-NEXT: ret void327//328//329// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined330// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[T_VAR:%.*]]) #[[ATTR1]] {331// CHECK1-NEXT: entry:332// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8333// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8334// CHECK1-NEXT: [[T_VAR_ADDR:%.*]] = alloca ptr, align 8335// CHECK1-NEXT: [[T_VAR1:%.*]] = alloca i32, align 4336// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4337// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4338// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4339// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4340// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4341// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4342// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4343// CHECK1-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 8344// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8345// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8346// CHECK1-NEXT: store ptr [[T_VAR]], ptr [[T_VAR_ADDR]], align 8347// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[T_VAR_ADDR]], align 8348// CHECK1-NEXT: store i32 0, ptr [[T_VAR1]], align 4349// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4350// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 4351// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4352// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4353// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8354// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4355// 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)356// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4357// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1358// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]359// CHECK1: cond.true:360// CHECK1-NEXT: br label [[COND_END:%.*]]361// CHECK1: cond.false:362// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4363// CHECK1-NEXT: br label [[COND_END]]364// CHECK1: cond.end:365// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]366// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4367// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4368// CHECK1-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4369// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]370// CHECK1: omp.inner.for.cond:371// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11:![0-9]+]]372// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP11]]373// CHECK1-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]374// CHECK1-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]375// CHECK1: omp.inner.for.body:376// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]377// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1378// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]379// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP11]]380// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP11]]381// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP11]]382// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]383// CHECK1-NEXT: store i32 [[ADD3]], ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP11]]384// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]385// CHECK1: omp.body.continue:386// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]387// CHECK1: omp.inner.for.inc:388// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]389// CHECK1-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP11]], 1390// CHECK1-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]391// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP12:![0-9]+]]392// CHECK1: omp.inner.for.end:393// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]394// CHECK1: omp.loop.exit:395// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])396// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 4397// CHECK1-NEXT: [[TMP13:%.*]] = icmp ne i32 [[TMP12]], 0398// CHECK1-NEXT: br i1 [[TMP13]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]399// CHECK1: .omp.final.then:400// CHECK1-NEXT: store i32 2, ptr [[I]], align 4401// CHECK1-NEXT: br label [[DOTOMP_FINAL_DONE]]402// CHECK1: .omp.final.done:403// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0404// CHECK1-NEXT: store ptr [[T_VAR1]], ptr [[TMP14]], align 8405// CHECK1-NEXT: [[TMP15:%.*]] = call i32 @__kmpc_reduce(ptr @[[GLOB2]], i32 [[TMP2]], i32 1, i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)406// CHECK1-NEXT: switch i32 [[TMP15]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [407// CHECK1-NEXT: i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]408// CHECK1-NEXT: i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]409// CHECK1-NEXT: ]410// CHECK1: .omp.reduction.case1:411// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[TMP0]], align 4412// CHECK1-NEXT: [[TMP17:%.*]] = load i32, ptr [[T_VAR1]], align 4413// CHECK1-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP16]], [[TMP17]]414// CHECK1-NEXT: store i32 [[ADD5]], ptr [[TMP0]], align 4415// CHECK1-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)416// CHECK1-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]417// CHECK1: .omp.reduction.case2:418// CHECK1-NEXT: [[TMP18:%.*]] = load i32, ptr [[T_VAR1]], align 4419// CHECK1-NEXT: [[TMP19:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP18]] monotonic, align 4420// CHECK1-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)421// CHECK1-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]422// CHECK1: .omp.reduction.default:423// CHECK1-NEXT: ret void424//425//426// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp.reduction.reduction_func427// CHECK1-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3]] {428// CHECK1-NEXT: entry:429// CHECK1-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 8430// CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 8431// CHECK1-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 8432// CHECK1-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 8433// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 8434// CHECK1-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 8435// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i64 0, i64 0436// CHECK1-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 8437// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i64 0, i64 0438// CHECK1-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 8439// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4440// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4441// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]442// CHECK1-NEXT: store i32 [[ADD]], ptr [[TMP7]], align 4443// CHECK1-NEXT: ret void444//445//446// CHECK3-LABEL: define {{[^@]+}}@main447// CHECK3-SAME: () #[[ATTR0:[0-9]+]] {448// CHECK3-NEXT: entry:449// CHECK3-NEXT: [[RETVAL:%.*]] = alloca i32, align 4450// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4451// CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4452// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4453// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4454// CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8455// CHECK3-NEXT: store i32 0, ptr [[RETVAL]], align 4456// CHECK3-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0457// CHECK3-NEXT: store ptr @_ZZ4mainE5sivar, ptr [[TMP0]], align 4458// CHECK3-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0459// CHECK3-NEXT: store ptr @_ZZ4mainE5sivar, ptr [[TMP1]], align 4460// CHECK3-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0461// CHECK3-NEXT: store ptr null, ptr [[TMP2]], align 4462// CHECK3-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0463// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0464// CHECK3-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0465// CHECK3-NEXT: store i32 3, ptr [[TMP5]], align 4466// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1467// CHECK3-NEXT: store i32 1, ptr [[TMP6]], align 4468// CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2469// CHECK3-NEXT: store ptr [[TMP3]], ptr [[TMP7]], align 4470// CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3471// CHECK3-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 4472// CHECK3-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4473// CHECK3-NEXT: store ptr @.offload_sizes, ptr [[TMP9]], align 4474// CHECK3-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5475// CHECK3-NEXT: store ptr @.offload_maptypes, ptr [[TMP10]], align 4476// CHECK3-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6477// CHECK3-NEXT: store ptr null, ptr [[TMP11]], align 4478// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7479// CHECK3-NEXT: store ptr null, ptr [[TMP12]], align 4480// CHECK3-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8481// CHECK3-NEXT: store i64 2, ptr [[TMP13]], align 8482// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9483// CHECK3-NEXT: store i64 0, ptr [[TMP14]], align 8484// CHECK3-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10485// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4486// CHECK3-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11487// CHECK3-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP16]], align 4488// CHECK3-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12489// CHECK3-NEXT: store i32 0, ptr [[TMP17]], align 4490// CHECK3-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3:[0-9]+]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l60.region_id, ptr [[KERNEL_ARGS]])491// CHECK3-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0492// CHECK3-NEXT: br i1 [[TMP19]], 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_l60(ptr @_ZZ4mainE5sivar) #[[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_l60502// CHECK3-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR1:[0-9]+]] {503// CHECK3-NEXT: entry:504// CHECK3-NEXT: [[SIVAR_ADDR:%.*]] = alloca ptr, align 4505// CHECK3-NEXT: store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 4506// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 4507// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l60.omp_outlined, ptr [[TMP0]])508// CHECK3-NEXT: ret void509//510//511// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l60.omp_outlined512// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR1]] {513// CHECK3-NEXT: entry:514// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4515// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4516// CHECK3-NEXT: [[SIVAR_ADDR:%.*]] = alloca ptr, align 4517// CHECK3-NEXT: [[SIVAR1:%.*]] = alloca i32, align 4518// CHECK3-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4519// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4520// CHECK3-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4521// CHECK3-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4522// CHECK3-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4523// CHECK3-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4524// CHECK3-NEXT: [[I:%.*]] = alloca i32, align 4525// CHECK3-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 4526// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4527// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4528// CHECK3-NEXT: store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 4529// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 4530// CHECK3-NEXT: store i32 0, ptr [[SIVAR1]], align 4531// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4532// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 4533// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4534// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4535// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4536// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4537// 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)538// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4539// CHECK3-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1540// CHECK3-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]541// CHECK3: cond.true:542// CHECK3-NEXT: br label [[COND_END:%.*]]543// CHECK3: cond.false:544// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4545// CHECK3-NEXT: br label [[COND_END]]546// CHECK3: cond.end:547// CHECK3-NEXT: [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]548// CHECK3-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4549// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4550// CHECK3-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4551// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]552// CHECK3: omp.inner.for.cond:553// CHECK3-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6:![0-9]+]]554// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP6]]555// CHECK3-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]556// CHECK3-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]557// CHECK3: omp.inner.for.body:558// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6]]559// CHECK3-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1560// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]561// CHECK3-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP6]]562// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP6]]563// CHECK3-NEXT: [[TMP10:%.*]] = load i32, ptr [[SIVAR1]], align 4, !llvm.access.group [[ACC_GRP6]]564// CHECK3-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]565// CHECK3-NEXT: store i32 [[ADD3]], ptr [[SIVAR1]], align 4, !llvm.access.group [[ACC_GRP6]]566// CHECK3-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]567// CHECK3: omp.body.continue:568// CHECK3-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]569// CHECK3: omp.inner.for.inc:570// CHECK3-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6]]571// CHECK3-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP11]], 1572// CHECK3-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6]]573// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP7:![0-9]+]]574// CHECK3: omp.inner.for.end:575// CHECK3-NEXT: br label [[OMP_LOOP_EXIT:%.*]]576// CHECK3: omp.loop.exit:577// CHECK3-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])578// CHECK3-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 4579// CHECK3-NEXT: [[TMP13:%.*]] = icmp ne i32 [[TMP12]], 0580// CHECK3-NEXT: br i1 [[TMP13]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]581// CHECK3: .omp.final.then:582// CHECK3-NEXT: store i32 2, ptr [[I]], align 4583// CHECK3-NEXT: br label [[DOTOMP_FINAL_DONE]]584// CHECK3: .omp.final.done:585// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 0586// CHECK3-NEXT: store ptr [[SIVAR1]], ptr [[TMP14]], align 4587// CHECK3-NEXT: [[TMP15:%.*]] = call i32 @__kmpc_reduce(ptr @[[GLOB2:[0-9]+]], i32 [[TMP2]], i32 1, i32 4, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l60.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)588// CHECK3-NEXT: switch i32 [[TMP15]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [589// CHECK3-NEXT: i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]590// CHECK3-NEXT: i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]591// CHECK3-NEXT: ]592// CHECK3: .omp.reduction.case1:593// CHECK3-NEXT: [[TMP16:%.*]] = load i32, ptr [[TMP0]], align 4594// CHECK3-NEXT: [[TMP17:%.*]] = load i32, ptr [[SIVAR1]], align 4595// CHECK3-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP16]], [[TMP17]]596// CHECK3-NEXT: store i32 [[ADD5]], ptr [[TMP0]], align 4597// CHECK3-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)598// CHECK3-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]599// CHECK3: .omp.reduction.case2:600// CHECK3-NEXT: [[TMP18:%.*]] = load i32, ptr [[SIVAR1]], align 4601// CHECK3-NEXT: [[TMP19:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP18]] monotonic, align 4602// CHECK3-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)603// 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_l60.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: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4634// CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4635// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4636// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4637// CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8638// CHECK3-NEXT: store i32 0, ptr [[T_VAR]], align 4639// CHECK3-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[VEC]], ptr align 4 @__const._Z5tmainIiET_v.vec, i32 8, i1 false)640// CHECK3-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0641// CHECK3-NEXT: store ptr [[T_VAR]], ptr [[TMP0]], align 4642// CHECK3-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0643// CHECK3-NEXT: store ptr [[T_VAR]], ptr [[TMP1]], align 4644// CHECK3-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0645// CHECK3-NEXT: store ptr null, ptr [[TMP2]], align 4646// CHECK3-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0647// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0648// CHECK3-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0649// CHECK3-NEXT: store i32 3, ptr [[TMP5]], align 4650// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1651// CHECK3-NEXT: store i32 1, ptr [[TMP6]], align 4652// CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2653// CHECK3-NEXT: store ptr [[TMP3]], ptr [[TMP7]], align 4654// CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3655// CHECK3-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 4656// CHECK3-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4657// CHECK3-NEXT: store ptr @.offload_sizes.1, ptr [[TMP9]], align 4658// CHECK3-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5659// CHECK3-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP10]], align 4660// CHECK3-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6661// CHECK3-NEXT: store ptr null, ptr [[TMP11]], align 4662// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7663// CHECK3-NEXT: store ptr null, ptr [[TMP12]], align 4664// CHECK3-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8665// CHECK3-NEXT: store i64 2, ptr [[TMP13]], align 8666// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9667// CHECK3-NEXT: store i64 0, ptr [[TMP14]], align 8668// CHECK3-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10669// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4670// CHECK3-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11671// CHECK3-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP16]], align 4672// CHECK3-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12673// CHECK3-NEXT: store i32 0, ptr [[TMP17]], align 4674// CHECK3-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.region_id, ptr [[KERNEL_ARGS]])675// CHECK3-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0676// CHECK3-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]677// CHECK3: omp_offload.failed:678// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32(ptr [[T_VAR]]) #[[ATTR2]]679// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]]680// CHECK3: omp_offload.cont:681// CHECK3-NEXT: ret i32 0682//683//684// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32685// CHECK3-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[T_VAR:%.*]]) #[[ATTR1]] {686// CHECK3-NEXT: entry:687// CHECK3-NEXT: [[T_VAR_ADDR:%.*]] = alloca ptr, align 4688// CHECK3-NEXT: store ptr [[T_VAR]], ptr [[T_VAR_ADDR]], align 4689// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[T_VAR_ADDR]], align 4690// 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 [[TMP0]])691// CHECK3-NEXT: ret void692//693//694// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined695// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[T_VAR:%.*]]) #[[ATTR1]] {696// CHECK3-NEXT: entry:697// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4698// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4699// CHECK3-NEXT: [[T_VAR_ADDR:%.*]] = alloca ptr, align 4700// CHECK3-NEXT: [[T_VAR1:%.*]] = alloca i32, align 4701// CHECK3-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4702// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4703// CHECK3-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4704// CHECK3-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4705// CHECK3-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4706// CHECK3-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4707// CHECK3-NEXT: [[I:%.*]] = alloca i32, align 4708// CHECK3-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 4709// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4710// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4711// CHECK3-NEXT: store ptr [[T_VAR]], ptr [[T_VAR_ADDR]], align 4712// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[T_VAR_ADDR]], align 4713// CHECK3-NEXT: store i32 0, ptr [[T_VAR1]], align 4714// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4715// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 4716// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4717// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4718// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4719// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4720// 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)721// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4722// CHECK3-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1723// CHECK3-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]724// CHECK3: cond.true:725// CHECK3-NEXT: br label [[COND_END:%.*]]726// CHECK3: cond.false:727// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4728// CHECK3-NEXT: br label [[COND_END]]729// CHECK3: cond.end:730// CHECK3-NEXT: [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]731// CHECK3-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4732// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4733// CHECK3-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4734// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]735// CHECK3: omp.inner.for.cond:736// CHECK3-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12:![0-9]+]]737// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP12]]738// CHECK3-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]739// CHECK3-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]740// CHECK3: omp.inner.for.body:741// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]742// CHECK3-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1743// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]744// CHECK3-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP12]]745// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP12]]746// CHECK3-NEXT: [[TMP10:%.*]] = load i32, ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP12]]747// CHECK3-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]748// CHECK3-NEXT: store i32 [[ADD3]], ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP12]]749// CHECK3-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]750// CHECK3: omp.body.continue:751// CHECK3-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]752// CHECK3: omp.inner.for.inc:753// CHECK3-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]754// CHECK3-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP11]], 1755// CHECK3-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]756// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP13:![0-9]+]]757// CHECK3: omp.inner.for.end:758// CHECK3-NEXT: br label [[OMP_LOOP_EXIT:%.*]]759// CHECK3: omp.loop.exit:760// CHECK3-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])761// CHECK3-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 4762// CHECK3-NEXT: [[TMP13:%.*]] = icmp ne i32 [[TMP12]], 0763// CHECK3-NEXT: br i1 [[TMP13]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]764// CHECK3: .omp.final.then:765// CHECK3-NEXT: store i32 2, ptr [[I]], align 4766// CHECK3-NEXT: br label [[DOTOMP_FINAL_DONE]]767// CHECK3: .omp.final.done:768// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 0769// CHECK3-NEXT: store ptr [[T_VAR1]], ptr [[TMP14]], align 4770// CHECK3-NEXT: [[TMP15:%.*]] = call i32 @__kmpc_reduce(ptr @[[GLOB2]], i32 [[TMP2]], i32 1, i32 4, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)771// CHECK3-NEXT: switch i32 [[TMP15]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [772// CHECK3-NEXT: i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]773// CHECK3-NEXT: i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]774// CHECK3-NEXT: ]775// CHECK3: .omp.reduction.case1:776// CHECK3-NEXT: [[TMP16:%.*]] = load i32, ptr [[TMP0]], align 4777// CHECK3-NEXT: [[TMP17:%.*]] = load i32, ptr [[T_VAR1]], align 4778// CHECK3-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP16]], [[TMP17]]779// CHECK3-NEXT: store i32 [[ADD5]], ptr [[TMP0]], align 4780// CHECK3-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)781// CHECK3-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]782// CHECK3: .omp.reduction.case2:783// CHECK3-NEXT: [[TMP18:%.*]] = load i32, ptr [[T_VAR1]], align 4784// CHECK3-NEXT: [[TMP19:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP18]] monotonic, align 4785// CHECK3-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)786// CHECK3-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]787// CHECK3: .omp.reduction.default:788// CHECK3-NEXT: ret void789//790//791// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined.omp.reduction.reduction_func792// CHECK3-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3]] {793// CHECK3-NEXT: entry:794// CHECK3-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 4795// CHECK3-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 4796// CHECK3-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 4797// CHECK3-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 4798// CHECK3-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 4799// CHECK3-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 4800// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i32 0, i32 0801// CHECK3-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 4802// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i32 0, i32 0803// CHECK3-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 4804// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4805// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4806// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]807// CHECK3-NEXT: store i32 [[ADD]], ptr [[TMP7]], align 4808// CHECK3-NEXT: ret void809//810//811// CHECK5-LABEL: define {{[^@]+}}@main812// CHECK5-SAME: () #[[ATTR0:[0-9]+]] {813// CHECK5-NEXT: entry:814// CHECK5-NEXT: [[RETVAL:%.*]] = alloca i32, align 4815// CHECK5-NEXT: [[TMP:%.*]] = alloca i32, align 4816// CHECK5-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4817// CHECK5-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4818// CHECK5-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4819// CHECK5-NEXT: [[I:%.*]] = alloca i32, align 4820// CHECK5-NEXT: [[SIVAR:%.*]] = alloca i32, align 4821// CHECK5-NEXT: store i32 0, ptr [[RETVAL]], align 4822// CHECK5-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4823// CHECK5-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 4824// CHECK5-NEXT: [[TMP0:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4825// CHECK5-NEXT: store i32 [[TMP0]], ptr [[DOTOMP_IV]], align 4826// CHECK5-NEXT: store i32 0, ptr [[SIVAR]], align 4827// CHECK5-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]828// CHECK5: omp.inner.for.cond:829// CHECK5-NEXT: [[TMP1:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP2:![0-9]+]]830// CHECK5-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP2]]831// CHECK5-NEXT: [[CMP:%.*]] = icmp sle i32 [[TMP1]], [[TMP2]]832// CHECK5-NEXT: br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]833// CHECK5: omp.inner.for.body:834// CHECK5-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP2]]835// CHECK5-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1836// CHECK5-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]837// CHECK5-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP2]]838// CHECK5-NEXT: [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP2]]839// CHECK5-NEXT: [[TMP5:%.*]] = load i32, ptr [[SIVAR]], align 4, !llvm.access.group [[ACC_GRP2]]840// CHECK5-NEXT: [[ADD1:%.*]] = add nsw i32 [[TMP5]], [[TMP4]]841// CHECK5-NEXT: store i32 [[ADD1]], ptr [[SIVAR]], align 4, !llvm.access.group [[ACC_GRP2]]842// CHECK5-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]843// CHECK5: omp.body.continue:844// CHECK5-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]845// CHECK5: omp.inner.for.inc:846// CHECK5-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP2]]847// CHECK5-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1848// CHECK5-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP2]]849// CHECK5-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP3:![0-9]+]]850// CHECK5: omp.inner.for.end:851// CHECK5-NEXT: store i32 2, ptr [[I]], align 4852// CHECK5-NEXT: [[TMP7:%.*]] = load i32, ptr @_ZZ4mainE5sivar, align 4853// CHECK5-NEXT: [[TMP8:%.*]] = load i32, ptr [[SIVAR]], align 4854// CHECK5-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP7]], [[TMP8]]855// CHECK5-NEXT: store i32 [[ADD3]], ptr @_ZZ4mainE5sivar, align 4856// CHECK5-NEXT: [[CALL:%.*]] = call noundef signext i32 @_Z5tmainIiET_v()857// CHECK5-NEXT: ret i32 [[CALL]]858//859//860// CHECK5-LABEL: define {{[^@]+}}@_Z5tmainIiET_v861// CHECK5-SAME: () #[[ATTR1:[0-9]+]] comdat {862// CHECK5-NEXT: entry:863// CHECK5-NEXT: [[T_VAR:%.*]] = alloca i32, align 4864// CHECK5-NEXT: [[VEC:%.*]] = alloca [2 x i32], align 4865// CHECK5-NEXT: [[TMP:%.*]] = alloca i32, align 4866// CHECK5-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4867// CHECK5-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4868// CHECK5-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4869// CHECK5-NEXT: [[I:%.*]] = alloca i32, align 4870// CHECK5-NEXT: [[T_VAR1:%.*]] = alloca i32, align 4871// CHECK5-NEXT: store i32 0, ptr [[T_VAR]], align 4872// CHECK5-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[VEC]], ptr align 4 @__const._Z5tmainIiET_v.vec, i64 8, i1 false)873// CHECK5-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4874// CHECK5-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 4875// CHECK5-NEXT: [[TMP0:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4876// CHECK5-NEXT: store i32 [[TMP0]], ptr [[DOTOMP_IV]], align 4877// CHECK5-NEXT: store i32 0, ptr [[T_VAR1]], align 4878// CHECK5-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]879// CHECK5: omp.inner.for.cond:880// CHECK5-NEXT: [[TMP1:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6:![0-9]+]]881// CHECK5-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP6]]882// CHECK5-NEXT: [[CMP:%.*]] = icmp sle i32 [[TMP1]], [[TMP2]]883// CHECK5-NEXT: br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]884// CHECK5: omp.inner.for.body:885// CHECK5-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6]]886// CHECK5-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1887// CHECK5-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]888// CHECK5-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP6]]889// CHECK5-NEXT: [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP6]]890// CHECK5-NEXT: [[TMP5:%.*]] = load i32, ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP6]]891// CHECK5-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP5]], [[TMP4]]892// CHECK5-NEXT: store i32 [[ADD2]], ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP6]]893// CHECK5-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]894// CHECK5: omp.body.continue:895// CHECK5-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]896// CHECK5: omp.inner.for.inc:897// CHECK5-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6]]898// CHECK5-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP6]], 1899// CHECK5-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6]]900// CHECK5-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP7:![0-9]+]]901// CHECK5: omp.inner.for.end:902// CHECK5-NEXT: store i32 2, ptr [[I]], align 4903// CHECK5-NEXT: [[TMP7:%.*]] = load i32, ptr [[T_VAR]], align 4904// CHECK5-NEXT: [[TMP8:%.*]] = load i32, ptr [[T_VAR1]], align 4905// CHECK5-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP7]], [[TMP8]]906// CHECK5-NEXT: store i32 [[ADD4]], ptr [[T_VAR]], align 4907// CHECK5-NEXT: ret i32 0908//909//910// CHECK7-LABEL: define {{[^@]+}}@main911// CHECK7-SAME: () #[[ATTR0:[0-9]+]] {912// CHECK7-NEXT: entry:913// CHECK7-NEXT: [[RETVAL:%.*]] = alloca i32, align 4914// CHECK7-NEXT: [[TMP:%.*]] = alloca i32, align 4915// CHECK7-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4916// CHECK7-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4917// CHECK7-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4918// CHECK7-NEXT: [[I:%.*]] = alloca i32, align 4919// CHECK7-NEXT: [[SIVAR:%.*]] = alloca i32, align 4920// CHECK7-NEXT: store i32 0, ptr [[RETVAL]], align 4921// CHECK7-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4922// CHECK7-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 4923// CHECK7-NEXT: [[TMP0:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4924// CHECK7-NEXT: store i32 [[TMP0]], ptr [[DOTOMP_IV]], align 4925// CHECK7-NEXT: store i32 0, ptr [[SIVAR]], align 4926// CHECK7-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]927// CHECK7: omp.inner.for.cond:928// CHECK7-NEXT: [[TMP1:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP3:![0-9]+]]929// CHECK7-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP3]]930// CHECK7-NEXT: [[CMP:%.*]] = icmp sle i32 [[TMP1]], [[TMP2]]931// CHECK7-NEXT: br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]932// CHECK7: omp.inner.for.body:933// CHECK7-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP3]]934// CHECK7-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1935// CHECK7-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]936// CHECK7-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP3]]937// CHECK7-NEXT: [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP3]]938// CHECK7-NEXT: [[TMP5:%.*]] = load i32, ptr [[SIVAR]], align 4, !llvm.access.group [[ACC_GRP3]]939// CHECK7-NEXT: [[ADD1:%.*]] = add nsw i32 [[TMP5]], [[TMP4]]940// CHECK7-NEXT: store i32 [[ADD1]], ptr [[SIVAR]], align 4, !llvm.access.group [[ACC_GRP3]]941// CHECK7-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]942// CHECK7: omp.body.continue:943// CHECK7-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]944// CHECK7: omp.inner.for.inc:945// CHECK7-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP3]]946// CHECK7-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1947// CHECK7-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP3]]948// CHECK7-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP4:![0-9]+]]949// CHECK7: omp.inner.for.end:950// CHECK7-NEXT: store i32 2, ptr [[I]], align 4951// CHECK7-NEXT: [[TMP7:%.*]] = load i32, ptr @_ZZ4mainE5sivar, align 4952// CHECK7-NEXT: [[TMP8:%.*]] = load i32, ptr [[SIVAR]], align 4953// CHECK7-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP7]], [[TMP8]]954// CHECK7-NEXT: store i32 [[ADD3]], ptr @_ZZ4mainE5sivar, align 4955// CHECK7-NEXT: [[CALL:%.*]] = call noundef i32 @_Z5tmainIiET_v()956// CHECK7-NEXT: ret i32 [[CALL]]957//958//959// CHECK7-LABEL: define {{[^@]+}}@_Z5tmainIiET_v960// CHECK7-SAME: () #[[ATTR1:[0-9]+]] comdat {961// CHECK7-NEXT: entry:962// CHECK7-NEXT: [[T_VAR:%.*]] = alloca i32, align 4963// CHECK7-NEXT: [[VEC:%.*]] = alloca [2 x i32], align 4964// CHECK7-NEXT: [[TMP:%.*]] = alloca i32, align 4965// CHECK7-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4966// CHECK7-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4967// CHECK7-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4968// CHECK7-NEXT: [[I:%.*]] = alloca i32, align 4969// CHECK7-NEXT: [[T_VAR1:%.*]] = alloca i32, align 4970// CHECK7-NEXT: store i32 0, ptr [[T_VAR]], align 4971// CHECK7-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[VEC]], ptr align 4 @__const._Z5tmainIiET_v.vec, i32 8, i1 false)972// CHECK7-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4973// CHECK7-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 4974// CHECK7-NEXT: [[TMP0:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4975// CHECK7-NEXT: store i32 [[TMP0]], ptr [[DOTOMP_IV]], align 4976// CHECK7-NEXT: store i32 0, ptr [[T_VAR1]], align 4977// CHECK7-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]978// CHECK7: omp.inner.for.cond:979// CHECK7-NEXT: [[TMP1:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP7:![0-9]+]]980// CHECK7-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP7]]981// CHECK7-NEXT: [[CMP:%.*]] = icmp sle i32 [[TMP1]], [[TMP2]]982// CHECK7-NEXT: br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]983// CHECK7: omp.inner.for.body:984// CHECK7-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP7]]985// CHECK7-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1986// CHECK7-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]987// CHECK7-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP7]]988// CHECK7-NEXT: [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP7]]989// CHECK7-NEXT: [[TMP5:%.*]] = load i32, ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP7]]990// CHECK7-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP5]], [[TMP4]]991// CHECK7-NEXT: store i32 [[ADD2]], ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP7]]992// CHECK7-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]993// CHECK7: omp.body.continue:994// CHECK7-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]995// CHECK7: omp.inner.for.inc:996// CHECK7-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP7]]997// CHECK7-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP6]], 1998// CHECK7-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP7]]999// CHECK7-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP8:![0-9]+]]1000// CHECK7: omp.inner.for.end:1001// CHECK7-NEXT: store i32 2, ptr [[I]], align 41002// CHECK7-NEXT: [[TMP7:%.*]] = load i32, ptr [[T_VAR]], align 41003// CHECK7-NEXT: [[TMP8:%.*]] = load i32, ptr [[T_VAR1]], align 41004// CHECK7-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP7]], [[TMP8]]1005// CHECK7-NEXT: store i32 [[ADD4]], ptr [[T_VAR]], align 41006// CHECK7-NEXT: ret i32 01007//1008//1009// CHECK9-LABEL: define {{[^@]+}}@main1010// CHECK9-SAME: () #[[ATTR0:[0-9]+]] {1011// CHECK9-NEXT: entry:1012// CHECK9-NEXT: [[RETVAL:%.*]] = alloca i32, align 41013// CHECK9-NEXT: [[REF_TMP:%.*]] = alloca [[CLASS_ANON:%.*]], align 11014// CHECK9-NEXT: store i32 0, ptr [[RETVAL]], align 41015// CHECK9-NEXT: call void @"_ZZ4mainENK3$_0clEv"(ptr noundef nonnull align 1 dereferenceable(1) [[REF_TMP]])1016// CHECK9-NEXT: ret i32 01017//1018//1019// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l441020// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR2:[0-9]+]] {1021// CHECK9-NEXT: entry:1022// CHECK9-NEXT: [[SIVAR_ADDR:%.*]] = alloca ptr, align 81023// CHECK9-NEXT: store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 81024// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 81025// 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_l44.omp_outlined, ptr [[TMP0]])1026// CHECK9-NEXT: ret void1027//1028//1029// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l44.omp_outlined1030// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR2]] {1031// CHECK9-NEXT: entry:1032// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81033// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81034// CHECK9-NEXT: [[SIVAR_ADDR:%.*]] = alloca ptr, align 81035// CHECK9-NEXT: [[SIVAR1:%.*]] = alloca i32, align 41036// CHECK9-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41037// CHECK9-NEXT: [[TMP:%.*]] = alloca i32, align 41038// CHECK9-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41039// CHECK9-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41040// CHECK9-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41041// CHECK9-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41042// CHECK9-NEXT: [[I:%.*]] = alloca i32, align 41043// CHECK9-NEXT: [[REF_TMP:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 81044// CHECK9-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 81045// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81046// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81047// CHECK9-NEXT: store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 81048// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 81049// CHECK9-NEXT: store i32 0, ptr [[SIVAR1]], align 41050// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41051// CHECK9-NEXT: store i32 1, ptr [[DOTOMP_UB]], align 41052// CHECK9-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41053// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41054// CHECK9-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81055// CHECK9-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41056// 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)1057// CHECK9-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41058// CHECK9-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 11059// CHECK9-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1060// CHECK9: cond.true:1061// CHECK9-NEXT: br label [[COND_END:%.*]]1062// CHECK9: cond.false:1063// CHECK9-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41064// CHECK9-NEXT: br label [[COND_END]]1065// CHECK9: cond.end:1066// CHECK9-NEXT: [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]1067// CHECK9-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41068// CHECK9-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41069// CHECK9-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 41070// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1071// CHECK9: omp.inner.for.cond:1072// CHECK9-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP4:![0-9]+]]1073// CHECK9-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP4]]1074// CHECK9-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]1075// CHECK9-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1076// CHECK9: omp.inner.for.body:1077// CHECK9-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP4]]1078// CHECK9-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 11079// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1080// CHECK9-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP4]]1081// CHECK9-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP4]]1082// CHECK9-NEXT: [[TMP10:%.*]] = load i32, ptr [[SIVAR1]], align 4, !llvm.access.group [[ACC_GRP4]]1083// CHECK9-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]1084// CHECK9-NEXT: store i32 [[ADD3]], ptr [[SIVAR1]], align 4, !llvm.access.group [[ACC_GRP4]]1085// CHECK9-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 01086// CHECK9-NEXT: store ptr [[SIVAR1]], ptr [[TMP11]], align 8, !llvm.access.group [[ACC_GRP4]]1087// CHECK9-NEXT: call void @"_ZZZ4mainENK3$_0clEvENKUlvE_clEv"(ptr noundef nonnull align 8 dereferenceable(8) [[REF_TMP]]), !llvm.access.group [[ACC_GRP4]]1088// CHECK9-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1089// CHECK9: omp.body.continue:1090// CHECK9-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1091// CHECK9: omp.inner.for.inc:1092// CHECK9-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP4]]1093// CHECK9-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP12]], 11094// CHECK9-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP4]]1095// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP5:![0-9]+]]1096// CHECK9: omp.inner.for.end:1097// CHECK9-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1098// CHECK9: omp.loop.exit:1099// CHECK9-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])1100// CHECK9-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 41101// CHECK9-NEXT: [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 01102// CHECK9-NEXT: br i1 [[TMP14]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]1103// CHECK9: .omp.final.then:1104// CHECK9-NEXT: store i32 2, ptr [[I]], align 41105// CHECK9-NEXT: br label [[DOTOMP_FINAL_DONE]]1106// CHECK9: .omp.final.done:1107// CHECK9-NEXT: [[TMP15:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 01108// CHECK9-NEXT: store ptr [[SIVAR1]], ptr [[TMP15]], align 81109// CHECK9-NEXT: [[TMP16:%.*]] = call i32 @__kmpc_reduce(ptr @[[GLOB2:[0-9]+]], i32 [[TMP2]], i32 1, i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l44.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)1110// CHECK9-NEXT: switch i32 [[TMP16]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [1111// CHECK9-NEXT: i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]1112// CHECK9-NEXT: i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]1113// CHECK9-NEXT: ]1114// CHECK9: .omp.reduction.case1:1115// CHECK9-NEXT: [[TMP17:%.*]] = load i32, ptr [[TMP0]], align 41116// CHECK9-NEXT: [[TMP18:%.*]] = load i32, ptr [[SIVAR1]], align 41117// CHECK9-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP17]], [[TMP18]]1118// CHECK9-NEXT: store i32 [[ADD5]], ptr [[TMP0]], align 41119// CHECK9-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)1120// CHECK9-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]1121// CHECK9: .omp.reduction.case2:1122// CHECK9-NEXT: [[TMP19:%.*]] = load i32, ptr [[SIVAR1]], align 41123// CHECK9-NEXT: [[TMP20:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP19]] monotonic, align 41124// CHECK9-NEXT: call void @__kmpc_end_reduce(ptr @[[GLOB2]], i32 [[TMP2]], ptr @.gomp_critical_user_.reduction.var)1125// CHECK9-NEXT: br label [[DOTOMP_REDUCTION_DEFAULT]]1126// CHECK9: .omp.reduction.default:1127// CHECK9-NEXT: ret void1128//1129//1130// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l44.omp_outlined.omp.reduction.reduction_func1131// CHECK9-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR4:[0-9]+]] {1132// CHECK9-NEXT: entry:1133// CHECK9-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 81134// CHECK9-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 81135// CHECK9-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 81136// CHECK9-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 81137// CHECK9-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 81138// CHECK9-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 81139// CHECK9-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i64 0, i64 01140// CHECK9-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 81141// CHECK9-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i64 0, i64 01142// CHECK9-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 81143// CHECK9-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41144// CHECK9-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 41145// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]1146// CHECK9-NEXT: store i32 [[ADD]], ptr [[TMP7]], align 41147// CHECK9-NEXT: ret void1148//1149//1150// CHECK11-LABEL: define {{[^@]+}}@main1151// CHECK11-SAME: () #[[ATTR0:[0-9]+]] {1152// CHECK11-NEXT: entry:1153// CHECK11-NEXT: [[RETVAL:%.*]] = alloca i32, align 41154// CHECK11-NEXT: [[REF_TMP:%.*]] = alloca [[CLASS_ANON:%.*]], align 11155// CHECK11-NEXT: store i32 0, ptr [[RETVAL]], align 41156// CHECK11-NEXT: call void @"_ZZ4mainENK3$_0clEv"(ptr noundef nonnull align 1 dereferenceable(1) [[REF_TMP]])1157// CHECK11-NEXT: ret i32 01158//1159