1612 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 -verify -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK13// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple x86_64-unknown-unknown -emit-pch -o %t %s4// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK15 6// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"7// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple x86_64-unknown-unknown -emit-pch -o %t %s8// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"9 10// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK111// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple x86_64-unknown-unknown -emit-pch -o %t %s12// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK113 14// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"15// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple x86_64-unknown-unknown -emit-pch -o %t %s16// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"17// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple aarch64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK118// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple aarch64-unknown-unknown -emit-pch -o %t %s19// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple aarch64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK120 21// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple aarch64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"22// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple aarch64-unknown-unknown -emit-pch -o %t %s23// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple aarch64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"24 25// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple aarch64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK126// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple aarch64-unknown-unknown -emit-pch -o %t %s27// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple aarch64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK128 29// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple aarch64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"30// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple aarch64-unknown-unknown -emit-pch -o %t %s31// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple aarch64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"32 33// expected-no-diagnostics34#ifndef HEADER35#define HEADER36 37void fn1();38void fn2();39void fn3();40void fn4();41void fn5();42void fn6();43 44int Arg;45 46void gtid_test() {47#pragma omp target48#pragma omp teams49#pragma omp distribute parallel for50 for(int i = 0 ; i < 100; i++) {}51 52#pragma omp target53#pragma omp teams54#pragma omp distribute parallel for if (parallel: false)55 for(int i = 0 ; i < 100; i++) {56 gtid_test();57 }58}59 60 61template <typename T>62int tmain(T Arg) {63#pragma omp target64#pragma omp teams65#pragma omp distribute parallel for if (true)66 for(int i = 0 ; i < 100; i++) {67 fn1();68 }69#pragma omp target70#pragma omp teams71#pragma omp distribute parallel for if (false)72 for(int i = 0 ; i < 100; i++) {73 fn2();74 }75#pragma omp target76#pragma omp teams77#pragma omp distribute parallel for if (parallel: Arg)78 for(int i = 0 ; i < 100; i++) {79 fn3();80 }81 return 0;82}83 84int main() {85#pragma omp target86#pragma omp teams87#pragma omp distribute parallel for if (true)88 for(int i = 0 ; i < 100; i++) {89 90 91 fn4();92 }93 94#pragma omp target95#pragma omp teams96#pragma omp distribute parallel for if (false)97 for(int i = 0 ; i < 100; i++) {98 99 100 fn5();101 }102 103#pragma omp target104#pragma omp teams105#pragma omp distribute parallel for if (Arg)106 for(int i = 0 ; i < 100; i++) {107 108 109 fn6();110 }111 112 return tmain(Arg);113}114 115 116 117 118 119 120// call void [[T_OUTLINE_FUN_3:@.+]](121 122#endif123// CHECK1-LABEL: define {{[^@]+}}@_Z9gtid_testv124// CHECK1-SAME: () #[[ATTR0:[0-9]+]] {125// CHECK1-NEXT: entry:126// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4127// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8128// CHECK1-NEXT: [[_TMP1:%.*]] = alloca i32, align 4129// CHECK1-NEXT: [[KERNEL_ARGS2:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8130// CHECK1-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0131// CHECK1-NEXT: store i32 3, ptr [[TMP0]], align 4132// CHECK1-NEXT: [[TMP1:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1133// CHECK1-NEXT: store i32 0, ptr [[TMP1]], align 4134// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2135// CHECK1-NEXT: store ptr null, ptr [[TMP2]], align 8136// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3137// CHECK1-NEXT: store ptr null, ptr [[TMP3]], align 8138// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4139// CHECK1-NEXT: store ptr null, ptr [[TMP4]], align 8140// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5141// CHECK1-NEXT: store ptr null, ptr [[TMP5]], align 8142// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6143// CHECK1-NEXT: store ptr null, ptr [[TMP6]], align 8144// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7145// CHECK1-NEXT: store ptr null, ptr [[TMP7]], align 8146// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8147// CHECK1-NEXT: store i64 100, ptr [[TMP8]], align 8148// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9149// CHECK1-NEXT: store i64 0, ptr [[TMP9]], align 8150// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10151// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP10]], align 4152// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11153// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP11]], align 4154// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12155// CHECK1-NEXT: store i32 0, ptr [[TMP12]], align 4156// CHECK1-NEXT: [[TMP13:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l47.region_id, ptr [[KERNEL_ARGS]])157// CHECK1-NEXT: [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 0158// CHECK1-NEXT: br i1 [[TMP14]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]159// CHECK1: omp_offload.failed:160// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l47() #[[ATTR2:[0-9]+]]161// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]162// CHECK1: omp_offload.cont:163// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 0164// CHECK1-NEXT: store i32 3, ptr [[TMP15]], align 4165// CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 1166// CHECK1-NEXT: store i32 0, ptr [[TMP16]], align 4167// CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 2168// CHECK1-NEXT: store ptr null, ptr [[TMP17]], align 8169// CHECK1-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 3170// CHECK1-NEXT: store ptr null, ptr [[TMP18]], align 8171// CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 4172// CHECK1-NEXT: store ptr null, ptr [[TMP19]], align 8173// CHECK1-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 5174// CHECK1-NEXT: store ptr null, ptr [[TMP20]], align 8175// CHECK1-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 6176// CHECK1-NEXT: store ptr null, ptr [[TMP21]], align 8177// CHECK1-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 7178// CHECK1-NEXT: store ptr null, ptr [[TMP22]], align 8179// CHECK1-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 8180// CHECK1-NEXT: store i64 100, ptr [[TMP23]], align 8181// CHECK1-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 9182// CHECK1-NEXT: store i64 0, ptr [[TMP24]], align 8183// CHECK1-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 10184// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP25]], align 4185// CHECK1-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 11186// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP26]], align 4187// CHECK1-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 12188// CHECK1-NEXT: store i32 0, ptr [[TMP27]], align 4189// CHECK1-NEXT: [[TMP28:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l52.region_id, ptr [[KERNEL_ARGS2]])190// CHECK1-NEXT: [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 0191// CHECK1-NEXT: br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED3:%.*]], label [[OMP_OFFLOAD_CONT4:%.*]]192// CHECK1: omp_offload.failed3:193// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l52() #[[ATTR2]]194// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT4]]195// CHECK1: omp_offload.cont4:196// CHECK1-NEXT: ret void197//198//199// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l47200// CHECK1-SAME: () #[[ATTR1:[0-9]+]] {201// CHECK1-NEXT: entry:202// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l47.omp_outlined)203// CHECK1-NEXT: ret void204//205//206// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l47.omp_outlined207// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {208// CHECK1-NEXT: entry:209// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8210// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8211// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4212// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4213// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4214// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4215// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4216// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4217// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4218// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8219// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8220// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4221// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_COMB_UB]], align 4222// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4223// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4224// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8225// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4226// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP1]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)227// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4228// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 99229// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]230// CHECK1: cond.true:231// CHECK1-NEXT: br label [[COND_END:%.*]]232// CHECK1: cond.false:233// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4234// CHECK1-NEXT: br label [[COND_END]]235// CHECK1: cond.end:236// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]237// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4238// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4239// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4240// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]241// CHECK1: omp.inner.for.cond:242// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4243// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4244// CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]245// CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]246// CHECK1: omp.inner.for.body:247// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4248// CHECK1-NEXT: [[TMP8:%.*]] = zext i32 [[TMP7]] to i64249// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4250// CHECK1-NEXT: [[TMP10:%.*]] = zext i32 [[TMP9]] to i64251// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l47.omp_outlined.omp_outlined, i64 [[TMP8]], i64 [[TMP10]])252// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]253// CHECK1: omp.inner.for.inc:254// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4255// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4256// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP11]], [[TMP12]]257// CHECK1-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4258// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]259// CHECK1: omp.inner.for.end:260// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]261// CHECK1: omp.loop.exit:262// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])263// CHECK1-NEXT: ret void264//265//266// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l47.omp_outlined.omp_outlined267// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]]) #[[ATTR1]] {268// CHECK1-NEXT: entry:269// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8270// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8271// CHECK1-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8272// CHECK1-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8273// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4274// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4275// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4276// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4277// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4278// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4279// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4280// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8281// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8282// CHECK1-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8283// CHECK1-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8284// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4285// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_UB]], align 4286// CHECK1-NEXT: [[TMP0:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8287// CHECK1-NEXT: [[CONV:%.*]] = trunc i64 [[TMP0]] to i32288// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8289// CHECK1-NEXT: [[CONV1:%.*]] = trunc i64 [[TMP1]] to i32290// CHECK1-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 4291// CHECK1-NEXT: store i32 [[CONV1]], ptr [[DOTOMP_UB]], align 4292// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4293// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4294// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8295// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4296// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP3]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)297// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4298// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 99299// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]300// CHECK1: cond.true:301// CHECK1-NEXT: br label [[COND_END:%.*]]302// CHECK1: cond.false:303// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4304// CHECK1-NEXT: br label [[COND_END]]305// CHECK1: cond.end:306// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]307// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4308// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4309// CHECK1-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 4310// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]311// CHECK1: omp.inner.for.cond:312// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4313// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4314// CHECK1-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]315// CHECK1-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]316// CHECK1: omp.inner.for.body:317// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4318// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP9]], 1319// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]320// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4321// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]322// CHECK1: omp.body.continue:323// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]324// CHECK1: omp.inner.for.inc:325// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4326// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], 1327// CHECK1-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4328// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]329// CHECK1: omp.inner.for.end:330// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]331// CHECK1: omp.loop.exit:332// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP3]])333// CHECK1-NEXT: ret void334//335//336// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l52337// CHECK1-SAME: () #[[ATTR1]] {338// CHECK1-NEXT: entry:339// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l52.omp_outlined)340// CHECK1-NEXT: ret void341//342//343// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l52.omp_outlined344// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {345// CHECK1-NEXT: entry:346// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8347// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8348// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4349// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4350// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4351// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4352// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4353// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4354// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4355// CHECK1-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 4356// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8357// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8358// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4359// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_COMB_UB]], align 4360// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4361// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4362// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8363// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4364// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP1]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)365// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4366// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 99367// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]368// CHECK1: cond.true:369// CHECK1-NEXT: br label [[COND_END:%.*]]370// CHECK1: cond.false:371// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4372// CHECK1-NEXT: br label [[COND_END]]373// CHECK1: cond.end:374// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]375// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4376// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4377// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4378// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]379// CHECK1: omp.inner.for.cond:380// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4381// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4382// CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]383// CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]384// CHECK1: omp.inner.for.body:385// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4386// CHECK1-NEXT: [[TMP8:%.*]] = zext i32 [[TMP7]] to i64387// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4388// CHECK1-NEXT: [[TMP10:%.*]] = zext i32 [[TMP9]] to i64389// CHECK1-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB3]], i32 [[TMP1]])390// CHECK1-NEXT: [[TMP11:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8391// CHECK1-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 4392// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l52.omp_outlined.omp_outlined(ptr [[TMP11]], ptr [[DOTBOUND_ZERO_ADDR]], i64 [[TMP8]], i64 [[TMP10]]) #[[ATTR2]]393// CHECK1-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB3]], i32 [[TMP1]])394// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]395// CHECK1: omp.inner.for.inc:396// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4397// CHECK1-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4398// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP12]], [[TMP13]]399// CHECK1-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4400// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]401// CHECK1: omp.inner.for.end:402// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]403// CHECK1: omp.loop.exit:404// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])405// CHECK1-NEXT: ret void406//407//408// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l52.omp_outlined.omp_outlined409// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]]) #[[ATTR1]] {410// CHECK1-NEXT: entry:411// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8412// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8413// CHECK1-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8414// CHECK1-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8415// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4416// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4417// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4418// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4419// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4420// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4421// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4422// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8423// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8424// CHECK1-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8425// CHECK1-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8426// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4427// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_UB]], align 4428// CHECK1-NEXT: [[TMP0:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8429// CHECK1-NEXT: [[CONV:%.*]] = trunc i64 [[TMP0]] to i32430// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8431// CHECK1-NEXT: [[CONV1:%.*]] = trunc i64 [[TMP1]] to i32432// CHECK1-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 4433// CHECK1-NEXT: store i32 [[CONV1]], ptr [[DOTOMP_UB]], align 4434// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4435// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4436// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8437// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4438// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2]], i32 [[TMP3]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)439// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4440// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 99441// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]442// CHECK1: cond.true:443// CHECK1-NEXT: br label [[COND_END:%.*]]444// CHECK1: cond.false:445// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4446// CHECK1-NEXT: br label [[COND_END]]447// CHECK1: cond.end:448// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]449// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4450// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4451// CHECK1-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 4452// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]453// CHECK1: omp.inner.for.cond:454// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4455// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4456// CHECK1-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]457// CHECK1-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]458// CHECK1: omp.inner.for.body:459// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4460// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP9]], 1461// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]462// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4463// CHECK1-NEXT: call void @_Z9gtid_testv()464// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]465// CHECK1: omp.body.continue:466// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]467// CHECK1: omp.inner.for.inc:468// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4469// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], 1470// CHECK1-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4471// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]472// CHECK1: omp.inner.for.end:473// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]474// CHECK1: omp.loop.exit:475// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP3]])476// CHECK1-NEXT: ret void477//478//479// CHECK1-LABEL: define {{[^@]+}}@main480// CHECK1-SAME: () #[[ATTR3:[0-9]+]] {481// CHECK1-NEXT: entry:482// CHECK1-NEXT: [[RETVAL:%.*]] = alloca i32, align 4483// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4484// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8485// CHECK1-NEXT: [[_TMP1:%.*]] = alloca i32, align 4486// CHECK1-NEXT: [[KERNEL_ARGS2:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8487// CHECK1-NEXT: [[ARG_CASTED:%.*]] = alloca i64, align 8488// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8489// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8490// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8491// CHECK1-NEXT: [[_TMP5:%.*]] = alloca i32, align 4492// CHECK1-NEXT: [[KERNEL_ARGS6:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8493// CHECK1-NEXT: store i32 0, ptr [[RETVAL]], align 4494// CHECK1-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0495// CHECK1-NEXT: store i32 3, ptr [[TMP0]], align 4496// CHECK1-NEXT: [[TMP1:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1497// CHECK1-NEXT: store i32 0, ptr [[TMP1]], align 4498// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2499// CHECK1-NEXT: store ptr null, ptr [[TMP2]], align 8500// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3501// CHECK1-NEXT: store ptr null, ptr [[TMP3]], align 8502// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4503// CHECK1-NEXT: store ptr null, ptr [[TMP4]], align 8504// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5505// CHECK1-NEXT: store ptr null, ptr [[TMP5]], align 8506// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6507// CHECK1-NEXT: store ptr null, ptr [[TMP6]], align 8508// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7509// CHECK1-NEXT: store ptr null, ptr [[TMP7]], align 8510// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8511// CHECK1-NEXT: store i64 100, ptr [[TMP8]], align 8512// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9513// CHECK1-NEXT: store i64 0, ptr [[TMP9]], align 8514// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10515// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP10]], align 4516// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11517// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP11]], align 4518// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12519// CHECK1-NEXT: store i32 0, ptr [[TMP12]], align 4520// CHECK1-NEXT: [[TMP13:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l85.region_id, ptr [[KERNEL_ARGS]])521// CHECK1-NEXT: [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 0522// CHECK1-NEXT: br i1 [[TMP14]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]523// CHECK1: omp_offload.failed:524// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l85() #[[ATTR2]]525// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]526// CHECK1: omp_offload.cont:527// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 0528// CHECK1-NEXT: store i32 3, ptr [[TMP15]], align 4529// CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 1530// CHECK1-NEXT: store i32 0, ptr [[TMP16]], align 4531// CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 2532// CHECK1-NEXT: store ptr null, ptr [[TMP17]], align 8533// CHECK1-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 3534// CHECK1-NEXT: store ptr null, ptr [[TMP18]], align 8535// CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 4536// CHECK1-NEXT: store ptr null, ptr [[TMP19]], align 8537// CHECK1-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 5538// CHECK1-NEXT: store ptr null, ptr [[TMP20]], align 8539// CHECK1-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 6540// CHECK1-NEXT: store ptr null, ptr [[TMP21]], align 8541// CHECK1-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 7542// CHECK1-NEXT: store ptr null, ptr [[TMP22]], align 8543// CHECK1-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 8544// CHECK1-NEXT: store i64 100, ptr [[TMP23]], align 8545// CHECK1-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 9546// CHECK1-NEXT: store i64 0, ptr [[TMP24]], align 8547// CHECK1-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 10548// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP25]], align 4549// CHECK1-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 11550// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP26]], align 4551// CHECK1-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 12552// CHECK1-NEXT: store i32 0, ptr [[TMP27]], align 4553// CHECK1-NEXT: [[TMP28:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l94.region_id, ptr [[KERNEL_ARGS2]])554// CHECK1-NEXT: [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 0555// CHECK1-NEXT: br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED3:%.*]], label [[OMP_OFFLOAD_CONT4:%.*]]556// CHECK1: omp_offload.failed3:557// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l94() #[[ATTR2]]558// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT4]]559// CHECK1: omp_offload.cont4:560// CHECK1-NEXT: [[TMP30:%.*]] = load i32, ptr @Arg, align 4561// CHECK1-NEXT: store i32 [[TMP30]], ptr [[ARG_CASTED]], align 4562// CHECK1-NEXT: [[TMP31:%.*]] = load i64, ptr [[ARG_CASTED]], align 8563// CHECK1-NEXT: [[TMP32:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0564// CHECK1-NEXT: store i64 [[TMP31]], ptr [[TMP32]], align 8565// CHECK1-NEXT: [[TMP33:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0566// CHECK1-NEXT: store i64 [[TMP31]], ptr [[TMP33]], align 8567// CHECK1-NEXT: [[TMP34:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0568// CHECK1-NEXT: store ptr null, ptr [[TMP34]], align 8569// CHECK1-NEXT: [[TMP35:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0570// CHECK1-NEXT: [[TMP36:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0571// CHECK1-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 0572// CHECK1-NEXT: store i32 3, ptr [[TMP37]], align 4573// CHECK1-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 1574// CHECK1-NEXT: store i32 1, ptr [[TMP38]], align 4575// CHECK1-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 2576// CHECK1-NEXT: store ptr [[TMP35]], ptr [[TMP39]], align 8577// CHECK1-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 3578// CHECK1-NEXT: store ptr [[TMP36]], ptr [[TMP40]], align 8579// CHECK1-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 4580// CHECK1-NEXT: store ptr @.offload_sizes, ptr [[TMP41]], align 8581// CHECK1-NEXT: [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 5582// CHECK1-NEXT: store ptr @.offload_maptypes, ptr [[TMP42]], align 8583// CHECK1-NEXT: [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 6584// CHECK1-NEXT: store ptr null, ptr [[TMP43]], align 8585// CHECK1-NEXT: [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 7586// CHECK1-NEXT: store ptr null, ptr [[TMP44]], align 8587// CHECK1-NEXT: [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 8588// CHECK1-NEXT: store i64 100, ptr [[TMP45]], align 8589// CHECK1-NEXT: [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 9590// CHECK1-NEXT: store i64 0, ptr [[TMP46]], align 8591// CHECK1-NEXT: [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 10592// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP47]], align 4593// CHECK1-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 11594// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP48]], align 4595// CHECK1-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 12596// CHECK1-NEXT: store i32 0, ptr [[TMP49]], align 4597// CHECK1-NEXT: [[TMP50:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l103.region_id, ptr [[KERNEL_ARGS6]])598// CHECK1-NEXT: [[TMP51:%.*]] = icmp ne i32 [[TMP50]], 0599// CHECK1-NEXT: br i1 [[TMP51]], label [[OMP_OFFLOAD_FAILED7:%.*]], label [[OMP_OFFLOAD_CONT8:%.*]]600// CHECK1: omp_offload.failed7:601// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l103(i64 [[TMP31]]) #[[ATTR2]]602// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT8]]603// CHECK1: omp_offload.cont8:604// CHECK1-NEXT: [[TMP52:%.*]] = load i32, ptr @Arg, align 4605// CHECK1-NEXT: [[CALL:%.*]] = call noundef i32 @_Z5tmainIiEiT_(i32 noundef [[TMP52]])606// CHECK1-NEXT: ret i32 [[CALL]]607//608//609// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l85610// CHECK1-SAME: () #[[ATTR1]] {611// CHECK1-NEXT: entry:612// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l85.omp_outlined)613// CHECK1-NEXT: ret void614//615//616// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l85.omp_outlined617// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {618// CHECK1-NEXT: entry:619// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8620// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8621// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4622// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4623// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4624// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4625// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4626// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4627// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4628// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8629// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8630// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4631// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_COMB_UB]], align 4632// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4633// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4634// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8635// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4636// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP1]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)637// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4638// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 99639// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]640// CHECK1: cond.true:641// CHECK1-NEXT: br label [[COND_END:%.*]]642// CHECK1: cond.false:643// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4644// CHECK1-NEXT: br label [[COND_END]]645// CHECK1: cond.end:646// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]647// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4648// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4649// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4650// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]651// CHECK1: omp.inner.for.cond:652// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4653// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4654// CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]655// CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]656// CHECK1: omp.inner.for.body:657// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4658// CHECK1-NEXT: [[TMP8:%.*]] = zext i32 [[TMP7]] to i64659// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4660// CHECK1-NEXT: [[TMP10:%.*]] = zext i32 [[TMP9]] to i64661// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l85.omp_outlined.omp_outlined, i64 [[TMP8]], i64 [[TMP10]])662// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]663// CHECK1: omp.inner.for.inc:664// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4665// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4666// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP11]], [[TMP12]]667// CHECK1-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4668// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]669// CHECK1: omp.inner.for.end:670// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]671// CHECK1: omp.loop.exit:672// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])673// CHECK1-NEXT: ret void674//675//676// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l85.omp_outlined.omp_outlined677// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]]) #[[ATTR1]] {678// CHECK1-NEXT: entry:679// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8680// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8681// CHECK1-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8682// CHECK1-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8683// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4684// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4685// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4686// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4687// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4688// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4689// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4690// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8691// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8692// CHECK1-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8693// CHECK1-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8694// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4695// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_UB]], align 4696// CHECK1-NEXT: [[TMP0:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8697// CHECK1-NEXT: [[CONV:%.*]] = trunc i64 [[TMP0]] to i32698// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8699// CHECK1-NEXT: [[CONV1:%.*]] = trunc i64 [[TMP1]] to i32700// CHECK1-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 4701// CHECK1-NEXT: store i32 [[CONV1]], ptr [[DOTOMP_UB]], align 4702// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4703// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4704// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8705// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4706// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2]], i32 [[TMP3]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)707// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4708// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 99709// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]710// CHECK1: cond.true:711// CHECK1-NEXT: br label [[COND_END:%.*]]712// CHECK1: cond.false:713// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4714// CHECK1-NEXT: br label [[COND_END]]715// CHECK1: cond.end:716// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]717// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4718// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4719// CHECK1-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 4720// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]721// CHECK1: omp.inner.for.cond:722// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4723// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4724// CHECK1-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]725// CHECK1-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]726// CHECK1: omp.inner.for.body:727// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4728// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP9]], 1729// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]730// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4731// CHECK1-NEXT: call void @_Z3fn4v()732// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]733// CHECK1: omp.body.continue:734// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]735// CHECK1: omp.inner.for.inc:736// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4737// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], 1738// CHECK1-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4739// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]740// CHECK1: omp.inner.for.end:741// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]742// CHECK1: omp.loop.exit:743// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP3]])744// CHECK1-NEXT: ret void745//746//747// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l94748// CHECK1-SAME: () #[[ATTR1]] {749// CHECK1-NEXT: entry:750// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l94.omp_outlined)751// CHECK1-NEXT: ret void752//753//754// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l94.omp_outlined755// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {756// CHECK1-NEXT: entry:757// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8758// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8759// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4760// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4761// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4762// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4763// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4764// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4765// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4766// CHECK1-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 4767// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8768// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8769// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4770// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_COMB_UB]], align 4771// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4772// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4773// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8774// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4775// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP1]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)776// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4777// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 99778// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]779// CHECK1: cond.true:780// CHECK1-NEXT: br label [[COND_END:%.*]]781// CHECK1: cond.false:782// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4783// CHECK1-NEXT: br label [[COND_END]]784// CHECK1: cond.end:785// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]786// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4787// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4788// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4789// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]790// CHECK1: omp.inner.for.cond:791// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4792// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4793// CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]794// CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]795// CHECK1: omp.inner.for.body:796// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4797// CHECK1-NEXT: [[TMP8:%.*]] = zext i32 [[TMP7]] to i64798// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4799// CHECK1-NEXT: [[TMP10:%.*]] = zext i32 [[TMP9]] to i64800// CHECK1-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB3]], i32 [[TMP1]])801// CHECK1-NEXT: [[TMP11:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8802// CHECK1-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 4803// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l94.omp_outlined.omp_outlined(ptr [[TMP11]], ptr [[DOTBOUND_ZERO_ADDR]], i64 [[TMP8]], i64 [[TMP10]]) #[[ATTR2]]804// CHECK1-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB3]], i32 [[TMP1]])805// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]806// CHECK1: omp.inner.for.inc:807// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4808// CHECK1-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4809// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP12]], [[TMP13]]810// CHECK1-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4811// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]812// CHECK1: omp.inner.for.end:813// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]814// CHECK1: omp.loop.exit:815// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])816// CHECK1-NEXT: ret void817//818//819// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l94.omp_outlined.omp_outlined820// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]]) #[[ATTR1]] {821// CHECK1-NEXT: entry:822// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8823// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8824// CHECK1-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8825// CHECK1-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8826// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4827// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4828// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4829// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4830// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4831// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4832// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4833// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8834// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8835// CHECK1-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8836// CHECK1-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8837// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4838// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_UB]], align 4839// CHECK1-NEXT: [[TMP0:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8840// CHECK1-NEXT: [[CONV:%.*]] = trunc i64 [[TMP0]] to i32841// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8842// CHECK1-NEXT: [[CONV1:%.*]] = trunc i64 [[TMP1]] to i32843// CHECK1-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 4844// CHECK1-NEXT: store i32 [[CONV1]], ptr [[DOTOMP_UB]], align 4845// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4846// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4847// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8848// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4849// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2]], i32 [[TMP3]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)850// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4851// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 99852// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]853// CHECK1: cond.true:854// CHECK1-NEXT: br label [[COND_END:%.*]]855// CHECK1: cond.false:856// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4857// CHECK1-NEXT: br label [[COND_END]]858// CHECK1: cond.end:859// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]860// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4861// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4862// CHECK1-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 4863// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]864// CHECK1: omp.inner.for.cond:865// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4866// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4867// CHECK1-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]868// CHECK1-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]869// CHECK1: omp.inner.for.body:870// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4871// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP9]], 1872// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]873// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4874// CHECK1-NEXT: call void @_Z3fn5v()875// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]876// CHECK1: omp.body.continue:877// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]878// CHECK1: omp.inner.for.inc:879// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4880// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], 1881// CHECK1-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4882// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]883// CHECK1: omp.inner.for.end:884// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]885// CHECK1: omp.loop.exit:886// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP3]])887// CHECK1-NEXT: ret void888//889//890// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l103891// CHECK1-SAME: (i64 noundef [[ARG:%.*]]) #[[ATTR1]] {892// CHECK1-NEXT: entry:893// CHECK1-NEXT: [[ARG_ADDR:%.*]] = alloca i64, align 8894// CHECK1-NEXT: store i64 [[ARG]], ptr [[ARG_ADDR]], align 8895// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l103.omp_outlined, ptr [[ARG_ADDR]])896// CHECK1-NEXT: ret void897//898//899// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l103.omp_outlined900// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARG:%.*]]) #[[ATTR1]] {901// CHECK1-NEXT: entry:902// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8903// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8904// CHECK1-NEXT: [[ARG_ADDR:%.*]] = alloca ptr, align 8905// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4906// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4907// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4908// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4909// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4910// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4911// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4912// CHECK1-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 4913// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8914// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8915// CHECK1-NEXT: store ptr [[ARG]], ptr [[ARG_ADDR]], align 8916// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[ARG_ADDR]], align 8917// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4918// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_COMB_UB]], align 4919// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4920// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4921// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8922// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4923// 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)924// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4925// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 99926// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]927// CHECK1: cond.true:928// CHECK1-NEXT: br label [[COND_END:%.*]]929// CHECK1: cond.false:930// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4931// CHECK1-NEXT: br label [[COND_END]]932// CHECK1: cond.end:933// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]934// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4935// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4936// CHECK1-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4937// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]938// CHECK1: omp.inner.for.cond:939// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4940// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4941// CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]942// CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]943// CHECK1: omp.inner.for.body:944// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4945// CHECK1-NEXT: [[TMP9:%.*]] = zext i32 [[TMP8]] to i64946// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4947// CHECK1-NEXT: [[TMP11:%.*]] = zext i32 [[TMP10]] to i64948// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[TMP0]], align 4949// CHECK1-NEXT: [[TOBOOL:%.*]] = icmp ne i32 [[TMP12]], 0950// CHECK1-NEXT: br i1 [[TOBOOL]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]951// CHECK1: omp_if.then:952// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l103.omp_outlined.omp_outlined, i64 [[TMP9]], i64 [[TMP11]])953// CHECK1-NEXT: br label [[OMP_IF_END:%.*]]954// CHECK1: omp_if.else:955// CHECK1-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB3]], i32 [[TMP2]])956// CHECK1-NEXT: [[TMP13:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8957// CHECK1-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 4958// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l103.omp_outlined.omp_outlined(ptr [[TMP13]], ptr [[DOTBOUND_ZERO_ADDR]], i64 [[TMP9]], i64 [[TMP11]]) #[[ATTR2]]959// CHECK1-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB3]], i32 [[TMP2]])960// CHECK1-NEXT: br label [[OMP_IF_END]]961// CHECK1: omp_if.end:962// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]963// CHECK1: omp.inner.for.inc:964// CHECK1-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4965// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4966// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP14]], [[TMP15]]967// CHECK1-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4968// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]969// CHECK1: omp.inner.for.end:970// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]971// CHECK1: omp.loop.exit:972// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])973// CHECK1-NEXT: ret void974//975//976// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l103.omp_outlined.omp_outlined977// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]]) #[[ATTR1]] {978// CHECK1-NEXT: entry:979// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8980// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8981// CHECK1-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8982// CHECK1-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8983// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4984// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4985// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4986// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4987// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4988// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4989// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4990// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8991// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8992// CHECK1-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8993// CHECK1-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8994// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4995// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_UB]], align 4996// CHECK1-NEXT: [[TMP0:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8997// CHECK1-NEXT: [[CONV:%.*]] = trunc i64 [[TMP0]] to i32998// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8999// CHECK1-NEXT: [[CONV1:%.*]] = trunc i64 [[TMP1]] to i321000// CHECK1-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 41001// CHECK1-NEXT: store i32 [[CONV1]], ptr [[DOTOMP_UB]], align 41002// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41003// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41004// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81005// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 41006// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2]], i32 [[TMP3]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1007// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41008// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 991009// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1010// CHECK1: cond.true:1011// CHECK1-NEXT: br label [[COND_END:%.*]]1012// CHECK1: cond.false:1013// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41014// CHECK1-NEXT: br label [[COND_END]]1015// CHECK1: cond.end:1016// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]1017// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41018// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41019// CHECK1-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 41020// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1021// CHECK1: omp.inner.for.cond:1022// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41023// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41024// CHECK1-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]1025// CHECK1-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1026// CHECK1: omp.inner.for.body:1027// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41028// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP9]], 11029// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1030// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 41031// CHECK1-NEXT: call void @_Z3fn6v()1032// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1033// CHECK1: omp.body.continue:1034// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1035// CHECK1: omp.inner.for.inc:1036// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41037// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], 11038// CHECK1-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 41039// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]1040// CHECK1: omp.inner.for.end:1041// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1042// CHECK1: omp.loop.exit:1043// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP3]])1044// CHECK1-NEXT: ret void1045//1046//1047// CHECK1-LABEL: define {{[^@]+}}@_Z5tmainIiEiT_1048// CHECK1-SAME: (i32 noundef [[ARG:%.*]]) #[[ATTR0]] comdat {1049// CHECK1-NEXT: entry:1050// CHECK1-NEXT: [[ARG_ADDR:%.*]] = alloca i32, align 41051// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 41052// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81053// CHECK1-NEXT: [[_TMP1:%.*]] = alloca i32, align 41054// CHECK1-NEXT: [[KERNEL_ARGS2:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 81055// CHECK1-NEXT: [[ARG_CASTED:%.*]] = alloca i64, align 81056// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 81057// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 81058// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 81059// CHECK1-NEXT: [[_TMP5:%.*]] = alloca i32, align 41060// CHECK1-NEXT: [[KERNEL_ARGS6:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 81061// CHECK1-NEXT: store i32 [[ARG]], ptr [[ARG_ADDR]], align 41062// CHECK1-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01063// CHECK1-NEXT: store i32 3, ptr [[TMP0]], align 41064// CHECK1-NEXT: [[TMP1:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11065// CHECK1-NEXT: store i32 0, ptr [[TMP1]], align 41066// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21067// CHECK1-NEXT: store ptr null, ptr [[TMP2]], align 81068// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31069// CHECK1-NEXT: store ptr null, ptr [[TMP3]], align 81070// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41071// CHECK1-NEXT: store ptr null, ptr [[TMP4]], align 81072// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51073// CHECK1-NEXT: store ptr null, ptr [[TMP5]], align 81074// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61075// CHECK1-NEXT: store ptr null, ptr [[TMP6]], align 81076// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71077// CHECK1-NEXT: store ptr null, ptr [[TMP7]], align 81078// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81079// CHECK1-NEXT: store i64 100, ptr [[TMP8]], align 81080// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91081// CHECK1-NEXT: store i64 0, ptr [[TMP9]], align 81082// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101083// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP10]], align 41084// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111085// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP11]], align 41086// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121087// CHECK1-NEXT: store i32 0, ptr [[TMP12]], align 41088// CHECK1-NEXT: [[TMP13:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l63.region_id, ptr [[KERNEL_ARGS]])1089// CHECK1-NEXT: [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 01090// CHECK1-NEXT: br i1 [[TMP14]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1091// CHECK1: omp_offload.failed:1092// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l63() #[[ATTR2]]1093// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]1094// CHECK1: omp_offload.cont:1095// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 01096// CHECK1-NEXT: store i32 3, ptr [[TMP15]], align 41097// CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 11098// CHECK1-NEXT: store i32 0, ptr [[TMP16]], align 41099// CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 21100// CHECK1-NEXT: store ptr null, ptr [[TMP17]], align 81101// CHECK1-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 31102// CHECK1-NEXT: store ptr null, ptr [[TMP18]], align 81103// CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 41104// CHECK1-NEXT: store ptr null, ptr [[TMP19]], align 81105// CHECK1-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 51106// CHECK1-NEXT: store ptr null, ptr [[TMP20]], align 81107// CHECK1-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 61108// CHECK1-NEXT: store ptr null, ptr [[TMP21]], align 81109// CHECK1-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 71110// CHECK1-NEXT: store ptr null, ptr [[TMP22]], align 81111// CHECK1-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 81112// CHECK1-NEXT: store i64 100, ptr [[TMP23]], align 81113// CHECK1-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 91114// CHECK1-NEXT: store i64 0, ptr [[TMP24]], align 81115// CHECK1-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 101116// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP25]], align 41117// CHECK1-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 111118// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP26]], align 41119// CHECK1-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 121120// CHECK1-NEXT: store i32 0, ptr [[TMP27]], align 41121// CHECK1-NEXT: [[TMP28:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l69.region_id, ptr [[KERNEL_ARGS2]])1122// CHECK1-NEXT: [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 01123// CHECK1-NEXT: br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED3:%.*]], label [[OMP_OFFLOAD_CONT4:%.*]]1124// CHECK1: omp_offload.failed3:1125// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l69() #[[ATTR2]]1126// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT4]]1127// CHECK1: omp_offload.cont4:1128// CHECK1-NEXT: [[TMP30:%.*]] = load i32, ptr [[ARG_ADDR]], align 41129// CHECK1-NEXT: store i32 [[TMP30]], ptr [[ARG_CASTED]], align 41130// CHECK1-NEXT: [[TMP31:%.*]] = load i64, ptr [[ARG_CASTED]], align 81131// CHECK1-NEXT: [[TMP32:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01132// CHECK1-NEXT: store i64 [[TMP31]], ptr [[TMP32]], align 81133// CHECK1-NEXT: [[TMP33:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01134// CHECK1-NEXT: store i64 [[TMP31]], ptr [[TMP33]], align 81135// CHECK1-NEXT: [[TMP34:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 01136// CHECK1-NEXT: store ptr null, ptr [[TMP34]], align 81137// CHECK1-NEXT: [[TMP35:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01138// CHECK1-NEXT: [[TMP36:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01139// CHECK1-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 01140// CHECK1-NEXT: store i32 3, ptr [[TMP37]], align 41141// CHECK1-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 11142// CHECK1-NEXT: store i32 1, ptr [[TMP38]], align 41143// CHECK1-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 21144// CHECK1-NEXT: store ptr [[TMP35]], ptr [[TMP39]], align 81145// CHECK1-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 31146// CHECK1-NEXT: store ptr [[TMP36]], ptr [[TMP40]], align 81147// CHECK1-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 41148// CHECK1-NEXT: store ptr @.offload_sizes.1, ptr [[TMP41]], align 81149// CHECK1-NEXT: [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 51150// CHECK1-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP42]], align 81151// CHECK1-NEXT: [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 61152// CHECK1-NEXT: store ptr null, ptr [[TMP43]], align 81153// CHECK1-NEXT: [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 71154// CHECK1-NEXT: store ptr null, ptr [[TMP44]], align 81155// CHECK1-NEXT: [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 81156// CHECK1-NEXT: store i64 100, ptr [[TMP45]], align 81157// CHECK1-NEXT: [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 91158// CHECK1-NEXT: store i64 0, ptr [[TMP46]], align 81159// CHECK1-NEXT: [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 101160// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP47]], align 41161// CHECK1-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 111162// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP48]], align 41163// CHECK1-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 121164// CHECK1-NEXT: store i32 0, ptr [[TMP49]], align 41165// CHECK1-NEXT: [[TMP50:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l75.region_id, ptr [[KERNEL_ARGS6]])1166// CHECK1-NEXT: [[TMP51:%.*]] = icmp ne i32 [[TMP50]], 01167// CHECK1-NEXT: br i1 [[TMP51]], label [[OMP_OFFLOAD_FAILED7:%.*]], label [[OMP_OFFLOAD_CONT8:%.*]]1168// CHECK1: omp_offload.failed7:1169// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l75(i64 [[TMP31]]) #[[ATTR2]]1170// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT8]]1171// CHECK1: omp_offload.cont8:1172// CHECK1-NEXT: ret i32 01173//1174//1175// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l631176// CHECK1-SAME: () #[[ATTR1]] {1177// CHECK1-NEXT: entry:1178// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l63.omp_outlined)1179// CHECK1-NEXT: ret void1180//1181//1182// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l63.omp_outlined1183// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {1184// CHECK1-NEXT: entry:1185// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81186// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81187// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41188// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 41189// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41190// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41191// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41192// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41193// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 41194// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81195// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81196// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 41197// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_COMB_UB]], align 41198// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41199// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41200// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81201// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 41202// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP1]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1203// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41204// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 991205// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1206// CHECK1: cond.true:1207// CHECK1-NEXT: br label [[COND_END:%.*]]1208// CHECK1: cond.false:1209// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41210// CHECK1-NEXT: br label [[COND_END]]1211// CHECK1: cond.end:1212// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]1213// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41214// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41215// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 41216// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1217// CHECK1: omp.inner.for.cond:1218// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41219// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41220// CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]1221// CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1222// CHECK1: omp.inner.for.body:1223// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41224// CHECK1-NEXT: [[TMP8:%.*]] = zext i32 [[TMP7]] to i641225// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41226// CHECK1-NEXT: [[TMP10:%.*]] = zext i32 [[TMP9]] to i641227// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l63.omp_outlined.omp_outlined, i64 [[TMP8]], i64 [[TMP10]])1228// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1229// CHECK1: omp.inner.for.inc:1230// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41231// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41232// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP11]], [[TMP12]]1233// CHECK1-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 41234// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]1235// CHECK1: omp.inner.for.end:1236// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1237// CHECK1: omp.loop.exit:1238// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])1239// CHECK1-NEXT: ret void1240//1241//1242// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l63.omp_outlined.omp_outlined1243// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]]) #[[ATTR1]] {1244// CHECK1-NEXT: entry:1245// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81246// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81247// CHECK1-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 81248// CHECK1-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 81249// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41250// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 41251// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41252// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41253// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41254// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41255// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 41256// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81257// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81258// CHECK1-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 81259// CHECK1-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 81260// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41261// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_UB]], align 41262// CHECK1-NEXT: [[TMP0:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 81263// CHECK1-NEXT: [[CONV:%.*]] = trunc i64 [[TMP0]] to i321264// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 81265// CHECK1-NEXT: [[CONV1:%.*]] = trunc i64 [[TMP1]] to i321266// CHECK1-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 41267// CHECK1-NEXT: store i32 [[CONV1]], ptr [[DOTOMP_UB]], align 41268// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41269// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41270// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81271// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 41272// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2]], i32 [[TMP3]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1273// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41274// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 991275// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1276// CHECK1: cond.true:1277// CHECK1-NEXT: br label [[COND_END:%.*]]1278// CHECK1: cond.false:1279// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41280// CHECK1-NEXT: br label [[COND_END]]1281// CHECK1: cond.end:1282// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]1283// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41284// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41285// CHECK1-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 41286// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1287// CHECK1: omp.inner.for.cond:1288// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41289// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41290// CHECK1-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]1291// CHECK1-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1292// CHECK1: omp.inner.for.body:1293// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41294// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP9]], 11295// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1296// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 41297// CHECK1-NEXT: call void @_Z3fn1v()1298// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1299// CHECK1: omp.body.continue:1300// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1301// CHECK1: omp.inner.for.inc:1302// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41303// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], 11304// CHECK1-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 41305// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]1306// CHECK1: omp.inner.for.end:1307// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1308// CHECK1: omp.loop.exit:1309// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP3]])1310// CHECK1-NEXT: ret void1311//1312//1313// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l691314// CHECK1-SAME: () #[[ATTR1]] {1315// CHECK1-NEXT: entry:1316// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l69.omp_outlined)1317// CHECK1-NEXT: ret void1318//1319//1320// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l69.omp_outlined1321// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {1322// CHECK1-NEXT: entry:1323// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81324// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81325// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41326// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 41327// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41328// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41329// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41330// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41331// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 41332// CHECK1-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 41333// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81334// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81335// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 41336// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_COMB_UB]], align 41337// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41338// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41339// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81340// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 41341// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP1]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1342// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41343// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 991344// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1345// CHECK1: cond.true:1346// CHECK1-NEXT: br label [[COND_END:%.*]]1347// CHECK1: cond.false:1348// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41349// CHECK1-NEXT: br label [[COND_END]]1350// CHECK1: cond.end:1351// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]1352// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41353// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41354// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 41355// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1356// CHECK1: omp.inner.for.cond:1357// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41358// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41359// CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]1360// CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1361// CHECK1: omp.inner.for.body:1362// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41363// CHECK1-NEXT: [[TMP8:%.*]] = zext i32 [[TMP7]] to i641364// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41365// CHECK1-NEXT: [[TMP10:%.*]] = zext i32 [[TMP9]] to i641366// CHECK1-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB3]], i32 [[TMP1]])1367// CHECK1-NEXT: [[TMP11:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81368// CHECK1-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 41369// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l69.omp_outlined.omp_outlined(ptr [[TMP11]], ptr [[DOTBOUND_ZERO_ADDR]], i64 [[TMP8]], i64 [[TMP10]]) #[[ATTR2]]1370// CHECK1-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB3]], i32 [[TMP1]])1371// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1372// CHECK1: omp.inner.for.inc:1373// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41374// CHECK1-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41375// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP12]], [[TMP13]]1376// CHECK1-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 41377// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]1378// CHECK1: omp.inner.for.end:1379// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1380// CHECK1: omp.loop.exit:1381// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])1382// CHECK1-NEXT: ret void1383//1384//1385// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l69.omp_outlined.omp_outlined1386// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]]) #[[ATTR1]] {1387// CHECK1-NEXT: entry:1388// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81389// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81390// CHECK1-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 81391// CHECK1-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 81392// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41393// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 41394// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41395// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41396// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41397// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41398// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 41399// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81400// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81401// CHECK1-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 81402// CHECK1-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 81403// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41404// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_UB]], align 41405// CHECK1-NEXT: [[TMP0:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 81406// CHECK1-NEXT: [[CONV:%.*]] = trunc i64 [[TMP0]] to i321407// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 81408// CHECK1-NEXT: [[CONV1:%.*]] = trunc i64 [[TMP1]] to i321409// CHECK1-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 41410// CHECK1-NEXT: store i32 [[CONV1]], ptr [[DOTOMP_UB]], align 41411// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41412// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41413// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81414// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 41415// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2]], i32 [[TMP3]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1416// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41417// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 991418// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1419// CHECK1: cond.true:1420// CHECK1-NEXT: br label [[COND_END:%.*]]1421// CHECK1: cond.false:1422// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41423// CHECK1-NEXT: br label [[COND_END]]1424// CHECK1: cond.end:1425// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]1426// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41427// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41428// CHECK1-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 41429// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1430// CHECK1: omp.inner.for.cond:1431// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41432// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41433// CHECK1-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]1434// CHECK1-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1435// CHECK1: omp.inner.for.body:1436// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41437// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP9]], 11438// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1439// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 41440// CHECK1-NEXT: call void @_Z3fn2v()1441// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1442// CHECK1: omp.body.continue:1443// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1444// CHECK1: omp.inner.for.inc:1445// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41446// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], 11447// CHECK1-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 41448// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]1449// CHECK1: omp.inner.for.end:1450// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1451// CHECK1: omp.loop.exit:1452// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP3]])1453// CHECK1-NEXT: ret void1454//1455//1456// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l751457// CHECK1-SAME: (i64 noundef [[ARG:%.*]]) #[[ATTR1]] {1458// CHECK1-NEXT: entry:1459// CHECK1-NEXT: [[ARG_ADDR:%.*]] = alloca i64, align 81460// CHECK1-NEXT: store i64 [[ARG]], ptr [[ARG_ADDR]], align 81461// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l75.omp_outlined, ptr [[ARG_ADDR]])1462// CHECK1-NEXT: ret void1463//1464//1465// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l75.omp_outlined1466// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARG:%.*]]) #[[ATTR1]] {1467// CHECK1-NEXT: entry:1468// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81469// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81470// CHECK1-NEXT: [[ARG_ADDR:%.*]] = alloca ptr, align 81471// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41472// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 41473// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41474// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41475// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41476// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41477// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 41478// CHECK1-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 41479// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81480// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81481// CHECK1-NEXT: store ptr [[ARG]], ptr [[ARG_ADDR]], align 81482// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[ARG_ADDR]], align 81483// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 41484// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_COMB_UB]], align 41485// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41486// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41487// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81488// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41489// 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)1490// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41491// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 991492// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1493// CHECK1: cond.true:1494// CHECK1-NEXT: br label [[COND_END:%.*]]1495// CHECK1: cond.false:1496// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41497// CHECK1-NEXT: br label [[COND_END]]1498// CHECK1: cond.end:1499// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]1500// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41501// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41502// CHECK1-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 41503// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1504// CHECK1: omp.inner.for.cond:1505// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41506// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41507// CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]1508// CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1509// CHECK1: omp.inner.for.body:1510// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41511// CHECK1-NEXT: [[TMP9:%.*]] = zext i32 [[TMP8]] to i641512// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41513// CHECK1-NEXT: [[TMP11:%.*]] = zext i32 [[TMP10]] to i641514// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[TMP0]], align 41515// CHECK1-NEXT: [[TOBOOL:%.*]] = icmp ne i32 [[TMP12]], 01516// CHECK1-NEXT: br i1 [[TOBOOL]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]1517// CHECK1: omp_if.then:1518// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l75.omp_outlined.omp_outlined, i64 [[TMP9]], i64 [[TMP11]])1519// CHECK1-NEXT: br label [[OMP_IF_END:%.*]]1520// CHECK1: omp_if.else:1521// CHECK1-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB3]], i32 [[TMP2]])1522// CHECK1-NEXT: [[TMP13:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81523// CHECK1-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 41524// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l75.omp_outlined.omp_outlined(ptr [[TMP13]], ptr [[DOTBOUND_ZERO_ADDR]], i64 [[TMP9]], i64 [[TMP11]]) #[[ATTR2]]1525// CHECK1-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB3]], i32 [[TMP2]])1526// CHECK1-NEXT: br label [[OMP_IF_END]]1527// CHECK1: omp_if.end:1528// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1529// CHECK1: omp.inner.for.inc:1530// CHECK1-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41531// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41532// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP14]], [[TMP15]]1533// CHECK1-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 41534// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]1535// CHECK1: omp.inner.for.end:1536// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1537// CHECK1: omp.loop.exit:1538// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])1539// CHECK1-NEXT: ret void1540//1541//1542// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l75.omp_outlined.omp_outlined1543// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]]) #[[ATTR1]] {1544// CHECK1-NEXT: entry:1545// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81546// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81547// CHECK1-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 81548// CHECK1-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 81549// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41550// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 41551// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41552// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41553// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41554// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41555// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 41556// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81557// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81558// CHECK1-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 81559// CHECK1-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 81560// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41561// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_UB]], align 41562// CHECK1-NEXT: [[TMP0:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 81563// CHECK1-NEXT: [[CONV:%.*]] = trunc i64 [[TMP0]] to i321564// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 81565// CHECK1-NEXT: [[CONV1:%.*]] = trunc i64 [[TMP1]] to i321566// CHECK1-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 41567// CHECK1-NEXT: store i32 [[CONV1]], ptr [[DOTOMP_UB]], align 41568// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41569// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41570// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81571// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 41572// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2]], i32 [[TMP3]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1573// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41574// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 991575// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1576// CHECK1: cond.true:1577// CHECK1-NEXT: br label [[COND_END:%.*]]1578// CHECK1: cond.false:1579// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41580// CHECK1-NEXT: br label [[COND_END]]1581// CHECK1: cond.end:1582// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]1583// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41584// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41585// CHECK1-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 41586// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1587// CHECK1: omp.inner.for.cond:1588// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41589// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41590// CHECK1-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]1591// CHECK1-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1592// CHECK1: omp.inner.for.body:1593// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41594// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP9]], 11595// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1596// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 41597// CHECK1-NEXT: call void @_Z3fn3v()1598// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1599// CHECK1: omp.body.continue:1600// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1601// CHECK1: omp.inner.for.inc:1602// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41603// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], 11604// CHECK1-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 41605// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]1606// CHECK1: omp.inner.for.end:1607// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1608// CHECK1: omp.loop.exit:1609// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP3]])1610// CHECK1-NEXT: ret void1611//1612