brintos / llvm-project-archived public Read only

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 18// 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=CHECK119// 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 %s20// 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=CHECK121 22// 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}}"23// 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 %s24// 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}}"25 26// 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=CHECK127// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple aarch64-unknown-unknown -emit-pch -o %t %s28// 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=CHECK129 30// 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}}"31// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple aarch64-unknown-unknown -emit-pch -o %t %s32// 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}}"33 34// expected-no-diagnostics35#ifndef HEADER36#define HEADER37 38void fn1();39void fn2();40void fn3();41void fn4();42void fn5();43void fn6();44 45int Arg;46 47void gtid_test() {48#pragma omp target49#pragma omp teams distribute parallel for50  for(int i = 0 ; i < 100; i++) {}51 52#pragma omp target53#pragma omp teams distribute parallel for if (parallel: false)54  for(int i = 0 ; i < 100; i++) {55    gtid_test();56  }57}58 59 60template <typename T>61int tmain(T Arg) {62#pragma omp target63#pragma omp teams distribute parallel for if (true)64  for(int i = 0 ; i < 100; i++) {65    fn1();66  }67#pragma omp target68#pragma omp teams distribute parallel for if (false)69  for(int i = 0 ; i < 100; i++) {70    fn2();71  }72#pragma omp target73#pragma omp teams distribute parallel for if (parallel: Arg)74  for(int i = 0 ; i < 100; i++) {75    fn3();76  }77  return 0;78}79 80int main() {81#pragma omp target82#pragma omp teams distribute parallel for if (true)83  for(int i = 0 ; i < 100; i++) {84 85 86    fn4();87  }88 89#pragma omp target90#pragma omp teams distribute parallel for if (false)91  for(int i = 0 ; i < 100; i++) {92 93 94    fn5();95  }96 97#pragma omp target98#pragma omp teams distribute parallel for if (Arg)99  for(int i = 0 ; i < 100; i++) {100 101 102    fn6();103  }104 105  return tmain(Arg);106}107 108 109 110 111 112 113// call void [[T_OUTLINE_FUN_3:@.+]](114 115#endif116// CHECK1-LABEL: define {{[^@]+}}@_Z9gtid_testv117// CHECK1-SAME: () #[[ATTR0:[0-9]+]] {118// CHECK1-NEXT:  entry:119// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4120// CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8121// CHECK1-NEXT:    [[_TMP1:%.*]] = alloca i32, align 4122// CHECK1-NEXT:    [[KERNEL_ARGS2:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8123// CHECK1-NEXT:    [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0124// CHECK1-NEXT:    store i32 3, ptr [[TMP0]], align 4125// CHECK1-NEXT:    [[TMP1:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1126// CHECK1-NEXT:    store i32 0, ptr [[TMP1]], align 4127// CHECK1-NEXT:    [[TMP2:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2128// CHECK1-NEXT:    store ptr null, ptr [[TMP2]], align 8129// CHECK1-NEXT:    [[TMP3:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3130// CHECK1-NEXT:    store ptr null, ptr [[TMP3]], align 8131// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4132// CHECK1-NEXT:    store ptr null, ptr [[TMP4]], align 8133// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5134// CHECK1-NEXT:    store ptr null, ptr [[TMP5]], align 8135// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6136// CHECK1-NEXT:    store ptr null, ptr [[TMP6]], align 8137// CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7138// CHECK1-NEXT:    store ptr null, ptr [[TMP7]], align 8139// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8140// CHECK1-NEXT:    store i64 100, ptr [[TMP8]], align 8141// CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9142// CHECK1-NEXT:    store i64 0, ptr [[TMP9]], align 8143// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10144// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP10]], align 4145// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11146// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP11]], align 4147// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12148// CHECK1-NEXT:    store i32 0, ptr [[TMP12]], align 4149// 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_l48.region_id, ptr [[KERNEL_ARGS]])150// CHECK1-NEXT:    [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 0151// CHECK1-NEXT:    br i1 [[TMP14]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]152// CHECK1:       omp_offload.failed:153// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l48() #[[ATTR2:[0-9]+]]154// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]155// CHECK1:       omp_offload.cont:156// CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 0157// CHECK1-NEXT:    store i32 3, ptr [[TMP15]], align 4158// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 1159// CHECK1-NEXT:    store i32 0, ptr [[TMP16]], align 4160// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 2161// CHECK1-NEXT:    store ptr null, ptr [[TMP17]], align 8162// CHECK1-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 3163// CHECK1-NEXT:    store ptr null, ptr [[TMP18]], align 8164// CHECK1-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 4165// CHECK1-NEXT:    store ptr null, ptr [[TMP19]], align 8166// CHECK1-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 5167// CHECK1-NEXT:    store ptr null, ptr [[TMP20]], align 8168// CHECK1-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 6169// CHECK1-NEXT:    store ptr null, ptr [[TMP21]], align 8170// CHECK1-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 7171// CHECK1-NEXT:    store ptr null, ptr [[TMP22]], align 8172// CHECK1-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 8173// CHECK1-NEXT:    store i64 100, ptr [[TMP23]], align 8174// CHECK1-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 9175// CHECK1-NEXT:    store i64 0, ptr [[TMP24]], align 8176// CHECK1-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 10177// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP25]], align 4178// CHECK1-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 11179// CHECK1-NEXT:    store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP26]], align 4180// CHECK1-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 12181// CHECK1-NEXT:    store i32 0, ptr [[TMP27]], align 4182// CHECK1-NEXT:    [[TMP28:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l52.region_id, ptr [[KERNEL_ARGS2]])183// CHECK1-NEXT:    [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 0184// CHECK1-NEXT:    br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED3:%.*]], label [[OMP_OFFLOAD_CONT4:%.*]]185// CHECK1:       omp_offload.failed3:186// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l52() #[[ATTR2]]187// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT4]]188// CHECK1:       omp_offload.cont4:189// CHECK1-NEXT:    ret void190//191//192// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l48193// CHECK1-SAME: () #[[ATTR1:[0-9]+]] {194// CHECK1-NEXT:  entry:195// 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_l48.omp_outlined)196// CHECK1-NEXT:    ret void197//198//199// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l48.omp_outlined200// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {201// CHECK1-NEXT:  entry:202// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8203// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8204// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4205// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4206// CHECK1-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4207// CHECK1-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4208// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4209// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4210// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4211// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8212// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8213// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 4214// CHECK1-NEXT:    store i32 99, ptr [[DOTOMP_COMB_UB]], align 4215// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4216// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4217// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8218// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4219// 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)220// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4221// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 99222// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]223// CHECK1:       cond.true:224// CHECK1-NEXT:    br label [[COND_END:%.*]]225// CHECK1:       cond.false:226// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4227// CHECK1-NEXT:    br label [[COND_END]]228// CHECK1:       cond.end:229// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]230// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4231// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4232// CHECK1-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4233// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]234// CHECK1:       omp.inner.for.cond:235// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4236// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4237// CHECK1-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]238// CHECK1-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]239// CHECK1:       omp.inner.for.body:240// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4241// CHECK1-NEXT:    [[TMP8:%.*]] = zext i32 [[TMP7]] to i64242// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4243// CHECK1-NEXT:    [[TMP10:%.*]] = zext i32 [[TMP9]] to i64244// 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_l48.omp_outlined.omp_outlined, i64 [[TMP8]], i64 [[TMP10]])245// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]246// CHECK1:       omp.inner.for.inc:247// CHECK1-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4248// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4249// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP11]], [[TMP12]]250// CHECK1-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4251// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]252// CHECK1:       omp.inner.for.end:253// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]254// CHECK1:       omp.loop.exit:255// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])256// CHECK1-NEXT:    ret void257//258//259// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l48.omp_outlined.omp_outlined260// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]]) #[[ATTR1]] {261// CHECK1-NEXT:  entry:262// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8263// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8264// CHECK1-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8265// CHECK1-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8266// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4267// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4268// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4269// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4270// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4271// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4272// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4273// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8274// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8275// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8276// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8277// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4278// CHECK1-NEXT:    store i32 99, ptr [[DOTOMP_UB]], align 4279// CHECK1-NEXT:    [[TMP0:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8280// CHECK1-NEXT:    [[CONV:%.*]] = trunc i64 [[TMP0]] to i32281// CHECK1-NEXT:    [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8282// CHECK1-NEXT:    [[CONV1:%.*]] = trunc i64 [[TMP1]] to i32283// CHECK1-NEXT:    store i32 [[CONV]], ptr [[DOTOMP_LB]], align 4284// CHECK1-NEXT:    store i32 [[CONV1]], ptr [[DOTOMP_UB]], align 4285// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4286// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4287// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8288// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4289// 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)290// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4291// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 99292// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]293// CHECK1:       cond.true:294// CHECK1-NEXT:    br label [[COND_END:%.*]]295// CHECK1:       cond.false:296// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4297// CHECK1-NEXT:    br label [[COND_END]]298// CHECK1:       cond.end:299// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]300// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4301// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4302// CHECK1-NEXT:    store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 4303// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]304// CHECK1:       omp.inner.for.cond:305// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4306// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4307// CHECK1-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]308// CHECK1-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]309// CHECK1:       omp.inner.for.body:310// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4311// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP9]], 1312// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]313// CHECK1-NEXT:    store i32 [[ADD]], ptr [[I]], align 4314// CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]315// CHECK1:       omp.body.continue:316// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]317// CHECK1:       omp.inner.for.inc:318// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4319// CHECK1-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP10]], 1320// CHECK1-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4321// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]322// CHECK1:       omp.inner.for.end:323// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]324// CHECK1:       omp.loop.exit:325// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP3]])326// CHECK1-NEXT:    ret void327//328//329// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l52330// CHECK1-SAME: () #[[ATTR1]] {331// CHECK1-NEXT:  entry:332// 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)333// CHECK1-NEXT:    ret void334//335//336// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l52.omp_outlined337// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {338// CHECK1-NEXT:  entry:339// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8340// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8341// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4342// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4343// CHECK1-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4344// CHECK1-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4345// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4346// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4347// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4348// CHECK1-NEXT:    [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 4349// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8350// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8351// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 4352// CHECK1-NEXT:    store i32 99, ptr [[DOTOMP_COMB_UB]], align 4353// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4354// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4355// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8356// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4357// 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)358// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4359// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 99360// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]361// CHECK1:       cond.true:362// CHECK1-NEXT:    br label [[COND_END:%.*]]363// CHECK1:       cond.false:364// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4365// CHECK1-NEXT:    br label [[COND_END]]366// CHECK1:       cond.end:367// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]368// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4369// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4370// CHECK1-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4371// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]372// CHECK1:       omp.inner.for.cond:373// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4374// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4375// CHECK1-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]376// CHECK1-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]377// CHECK1:       omp.inner.for.body:378// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4379// CHECK1-NEXT:    [[TMP8:%.*]] = zext i32 [[TMP7]] to i64380// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4381// CHECK1-NEXT:    [[TMP10:%.*]] = zext i32 [[TMP9]] to i64382// CHECK1-NEXT:    call void @__kmpc_serialized_parallel(ptr @[[GLOB3]], i32 [[TMP1]])383// CHECK1-NEXT:    [[TMP11:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8384// CHECK1-NEXT:    store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 4385// 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]]386// CHECK1-NEXT:    call void @__kmpc_end_serialized_parallel(ptr @[[GLOB3]], i32 [[TMP1]])387// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]388// CHECK1:       omp.inner.for.inc:389// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4390// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4391// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP12]], [[TMP13]]392// CHECK1-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4393// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]394// CHECK1:       omp.inner.for.end:395// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]396// CHECK1:       omp.loop.exit:397// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])398// CHECK1-NEXT:    ret void399//400//401// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l52.omp_outlined.omp_outlined402// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]]) #[[ATTR1]] {403// CHECK1-NEXT:  entry:404// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8405// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8406// CHECK1-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8407// CHECK1-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8408// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4409// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4410// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4411// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4412// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4413// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4414// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4415// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8416// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8417// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8418// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8419// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4420// CHECK1-NEXT:    store i32 99, ptr [[DOTOMP_UB]], align 4421// CHECK1-NEXT:    [[TMP0:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8422// CHECK1-NEXT:    [[CONV:%.*]] = trunc i64 [[TMP0]] to i32423// CHECK1-NEXT:    [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8424// CHECK1-NEXT:    [[CONV1:%.*]] = trunc i64 [[TMP1]] to i32425// CHECK1-NEXT:    store i32 [[CONV]], ptr [[DOTOMP_LB]], align 4426// CHECK1-NEXT:    store i32 [[CONV1]], ptr [[DOTOMP_UB]], align 4427// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4428// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4429// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8430// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4431// 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)432// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4433// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 99434// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]435// CHECK1:       cond.true:436// CHECK1-NEXT:    br label [[COND_END:%.*]]437// CHECK1:       cond.false:438// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4439// CHECK1-NEXT:    br label [[COND_END]]440// CHECK1:       cond.end:441// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]442// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4443// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4444// CHECK1-NEXT:    store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 4445// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]446// CHECK1:       omp.inner.for.cond:447// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4448// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4449// CHECK1-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]450// CHECK1-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]451// CHECK1:       omp.inner.for.body:452// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4453// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP9]], 1454// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]455// CHECK1-NEXT:    store i32 [[ADD]], ptr [[I]], align 4456// CHECK1-NEXT:    call void @_Z9gtid_testv()457// CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]458// CHECK1:       omp.body.continue:459// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]460// CHECK1:       omp.inner.for.inc:461// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4462// CHECK1-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP10]], 1463// CHECK1-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4464// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]465// CHECK1:       omp.inner.for.end:466// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]467// CHECK1:       omp.loop.exit:468// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP3]])469// CHECK1-NEXT:    ret void470//471//472// CHECK1-LABEL: define {{[^@]+}}@main473// CHECK1-SAME: () #[[ATTR3:[0-9]+]] {474// CHECK1-NEXT:  entry:475// CHECK1-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4476// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4477// CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8478// CHECK1-NEXT:    [[_TMP1:%.*]] = alloca i32, align 4479// CHECK1-NEXT:    [[KERNEL_ARGS2:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8480// CHECK1-NEXT:    [[ARG_CASTED:%.*]] = alloca i64, align 8481// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8482// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8483// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8484// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i8, align 1485// CHECK1-NEXT:    [[_TMP5:%.*]] = alloca i32, align 4486// CHECK1-NEXT:    [[KERNEL_ARGS6:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8487// CHECK1-NEXT:    store i32 0, ptr [[RETVAL]], align 4488// CHECK1-NEXT:    [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0489// CHECK1-NEXT:    store i32 3, ptr [[TMP0]], align 4490// CHECK1-NEXT:    [[TMP1:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1491// CHECK1-NEXT:    store i32 0, ptr [[TMP1]], align 4492// CHECK1-NEXT:    [[TMP2:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2493// CHECK1-NEXT:    store ptr null, ptr [[TMP2]], align 8494// CHECK1-NEXT:    [[TMP3:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3495// CHECK1-NEXT:    store ptr null, ptr [[TMP3]], align 8496// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4497// CHECK1-NEXT:    store ptr null, ptr [[TMP4]], align 8498// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5499// CHECK1-NEXT:    store ptr null, ptr [[TMP5]], align 8500// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6501// CHECK1-NEXT:    store ptr null, ptr [[TMP6]], align 8502// CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7503// CHECK1-NEXT:    store ptr null, ptr [[TMP7]], align 8504// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8505// CHECK1-NEXT:    store i64 100, ptr [[TMP8]], align 8506// CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9507// CHECK1-NEXT:    store i64 0, ptr [[TMP9]], align 8508// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10509// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP10]], align 4510// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11511// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP11]], align 4512// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12513// CHECK1-NEXT:    store i32 0, ptr [[TMP12]], align 4514// 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_l81.region_id, ptr [[KERNEL_ARGS]])515// CHECK1-NEXT:    [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 0516// CHECK1-NEXT:    br i1 [[TMP14]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]517// CHECK1:       omp_offload.failed:518// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l81() #[[ATTR2]]519// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]520// CHECK1:       omp_offload.cont:521// CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 0522// CHECK1-NEXT:    store i32 3, ptr [[TMP15]], align 4523// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 1524// CHECK1-NEXT:    store i32 0, ptr [[TMP16]], align 4525// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 2526// CHECK1-NEXT:    store ptr null, ptr [[TMP17]], align 8527// CHECK1-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 3528// CHECK1-NEXT:    store ptr null, ptr [[TMP18]], align 8529// CHECK1-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 4530// CHECK1-NEXT:    store ptr null, ptr [[TMP19]], align 8531// CHECK1-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 5532// CHECK1-NEXT:    store ptr null, ptr [[TMP20]], align 8533// CHECK1-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 6534// CHECK1-NEXT:    store ptr null, ptr [[TMP21]], align 8535// CHECK1-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 7536// CHECK1-NEXT:    store ptr null, ptr [[TMP22]], align 8537// CHECK1-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 8538// CHECK1-NEXT:    store i64 100, ptr [[TMP23]], align 8539// CHECK1-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 9540// CHECK1-NEXT:    store i64 0, ptr [[TMP24]], align 8541// CHECK1-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 10542// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP25]], align 4543// CHECK1-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 11544// CHECK1-NEXT:    store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP26]], align 4545// CHECK1-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 12546// CHECK1-NEXT:    store i32 0, ptr [[TMP27]], align 4547// CHECK1-NEXT:    [[TMP28:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l89.region_id, ptr [[KERNEL_ARGS2]])548// CHECK1-NEXT:    [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 0549// CHECK1-NEXT:    br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED3:%.*]], label [[OMP_OFFLOAD_CONT4:%.*]]550// CHECK1:       omp_offload.failed3:551// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l89() #[[ATTR2]]552// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT4]]553// CHECK1:       omp_offload.cont4:554// CHECK1-NEXT:    [[TMP30:%.*]] = load i32, ptr @Arg, align 4555// CHECK1-NEXT:    store i32 [[TMP30]], ptr [[ARG_CASTED]], align 4556// CHECK1-NEXT:    [[TMP31:%.*]] = load i64, ptr [[ARG_CASTED]], align 8557// CHECK1-NEXT:    [[TMP32:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0558// CHECK1-NEXT:    store i64 [[TMP31]], ptr [[TMP32]], align 8559// CHECK1-NEXT:    [[TMP33:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0560// CHECK1-NEXT:    store i64 [[TMP31]], ptr [[TMP33]], align 8561// CHECK1-NEXT:    [[TMP34:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0562// CHECK1-NEXT:    store ptr null, ptr [[TMP34]], align 8563// CHECK1-NEXT:    [[TMP35:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0564// CHECK1-NEXT:    [[TMP36:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0565// CHECK1-NEXT:    [[TMP37:%.*]] = load i32, ptr @Arg, align 4566// CHECK1-NEXT:    [[TOBOOL:%.*]] = icmp ne i32 [[TMP37]], 0567// CHECK1-NEXT:    [[STOREDV:%.*]] = zext i1 [[TOBOOL]] to i8568// CHECK1-NEXT:    store i8 [[STOREDV]], ptr [[DOTCAPTURE_EXPR_]], align 1569// CHECK1-NEXT:    [[TMP38:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 1570// CHECK1-NEXT:    [[LOADEDV:%.*]] = trunc i8 [[TMP38]] to i1571// CHECK1-NEXT:    [[TMP39:%.*]] = select i1 [[LOADEDV]], i32 0, i32 1572// CHECK1-NEXT:    [[TMP40:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP39]], 0573// CHECK1-NEXT:    [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 0574// CHECK1-NEXT:    store i32 3, ptr [[TMP41]], align 4575// CHECK1-NEXT:    [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 1576// CHECK1-NEXT:    store i32 1, ptr [[TMP42]], align 4577// CHECK1-NEXT:    [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 2578// CHECK1-NEXT:    store ptr [[TMP35]], ptr [[TMP43]], align 8579// CHECK1-NEXT:    [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 3580// CHECK1-NEXT:    store ptr [[TMP36]], ptr [[TMP44]], align 8581// CHECK1-NEXT:    [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 4582// CHECK1-NEXT:    store ptr @.offload_sizes, ptr [[TMP45]], align 8583// CHECK1-NEXT:    [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 5584// CHECK1-NEXT:    store ptr @.offload_maptypes, ptr [[TMP46]], align 8585// CHECK1-NEXT:    [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 6586// CHECK1-NEXT:    store ptr null, ptr [[TMP47]], align 8587// CHECK1-NEXT:    [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 7588// CHECK1-NEXT:    store ptr null, ptr [[TMP48]], align 8589// CHECK1-NEXT:    [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 8590// CHECK1-NEXT:    store i64 100, ptr [[TMP49]], align 8591// CHECK1-NEXT:    [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 9592// CHECK1-NEXT:    store i64 0, ptr [[TMP50]], align 8593// CHECK1-NEXT:    [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 10594// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP51]], align 4595// CHECK1-NEXT:    [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 11596// CHECK1-NEXT:    store [3 x i32] [[TMP40]], ptr [[TMP52]], align 4597// CHECK1-NEXT:    [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 12598// CHECK1-NEXT:    store i32 0, ptr [[TMP53]], align 4599// CHECK1-NEXT:    [[TMP54:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 [[TMP39]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l97.region_id, ptr [[KERNEL_ARGS6]])600// CHECK1-NEXT:    [[TMP55:%.*]] = icmp ne i32 [[TMP54]], 0601// CHECK1-NEXT:    br i1 [[TMP55]], label [[OMP_OFFLOAD_FAILED7:%.*]], label [[OMP_OFFLOAD_CONT8:%.*]]602// CHECK1:       omp_offload.failed7:603// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l97(i64 [[TMP31]]) #[[ATTR2]]604// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT8]]605// CHECK1:       omp_offload.cont8:606// CHECK1-NEXT:    [[TMP56:%.*]] = load i32, ptr @Arg, align 4607// CHECK1-NEXT:    [[CALL:%.*]] = call noundef i32 @_Z5tmainIiEiT_(i32 noundef [[TMP56]])608// CHECK1-NEXT:    ret i32 [[CALL]]609//610//611// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l81612// CHECK1-SAME: () #[[ATTR1]] {613// CHECK1-NEXT:  entry:614// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l81.omp_outlined)615// CHECK1-NEXT:    ret void616//617//618// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l81.omp_outlined619// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {620// CHECK1-NEXT:  entry:621// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8622// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8623// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4624// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4625// CHECK1-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4626// CHECK1-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4627// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4628// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4629// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4630// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8631// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8632// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 4633// CHECK1-NEXT:    store i32 99, ptr [[DOTOMP_COMB_UB]], align 4634// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4635// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4636// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8637// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4638// 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)639// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4640// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 99641// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]642// CHECK1:       cond.true:643// CHECK1-NEXT:    br label [[COND_END:%.*]]644// CHECK1:       cond.false:645// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4646// CHECK1-NEXT:    br label [[COND_END]]647// CHECK1:       cond.end:648// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]649// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4650// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4651// CHECK1-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4652// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]653// CHECK1:       omp.inner.for.cond:654// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4655// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4656// CHECK1-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]657// CHECK1-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]658// CHECK1:       omp.inner.for.body:659// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4660// CHECK1-NEXT:    [[TMP8:%.*]] = zext i32 [[TMP7]] to i64661// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4662// CHECK1-NEXT:    [[TMP10:%.*]] = zext i32 [[TMP9]] to i64663// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l81.omp_outlined.omp_outlined, i64 [[TMP8]], i64 [[TMP10]])664// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]665// CHECK1:       omp.inner.for.inc:666// CHECK1-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4667// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4668// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP11]], [[TMP12]]669// CHECK1-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4670// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]671// CHECK1:       omp.inner.for.end:672// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]673// CHECK1:       omp.loop.exit:674// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])675// CHECK1-NEXT:    ret void676//677//678// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l81.omp_outlined.omp_outlined679// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]]) #[[ATTR1]] {680// CHECK1-NEXT:  entry:681// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8682// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8683// CHECK1-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8684// CHECK1-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8685// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4686// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4687// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4688// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4689// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4690// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4691// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4692// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8693// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8694// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8695// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8696// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4697// CHECK1-NEXT:    store i32 99, ptr [[DOTOMP_UB]], align 4698// CHECK1-NEXT:    [[TMP0:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8699// CHECK1-NEXT:    [[CONV:%.*]] = trunc i64 [[TMP0]] to i32700// CHECK1-NEXT:    [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8701// CHECK1-NEXT:    [[CONV1:%.*]] = trunc i64 [[TMP1]] to i32702// CHECK1-NEXT:    store i32 [[CONV]], ptr [[DOTOMP_LB]], align 4703// CHECK1-NEXT:    store i32 [[CONV1]], ptr [[DOTOMP_UB]], align 4704// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4705// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4706// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8707// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4708// 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)709// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4710// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 99711// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]712// CHECK1:       cond.true:713// CHECK1-NEXT:    br label [[COND_END:%.*]]714// CHECK1:       cond.false:715// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4716// CHECK1-NEXT:    br label [[COND_END]]717// CHECK1:       cond.end:718// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]719// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4720// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4721// CHECK1-NEXT:    store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 4722// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]723// CHECK1:       omp.inner.for.cond:724// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4725// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4726// CHECK1-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]727// CHECK1-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]728// CHECK1:       omp.inner.for.body:729// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4730// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP9]], 1731// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]732// CHECK1-NEXT:    store i32 [[ADD]], ptr [[I]], align 4733// CHECK1-NEXT:    call void @_Z3fn4v()734// CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]735// CHECK1:       omp.body.continue:736// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]737// CHECK1:       omp.inner.for.inc:738// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4739// CHECK1-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP10]], 1740// CHECK1-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4741// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]742// CHECK1:       omp.inner.for.end:743// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]744// CHECK1:       omp.loop.exit:745// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP3]])746// CHECK1-NEXT:    ret void747//748//749// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l89750// CHECK1-SAME: () #[[ATTR1]] {751// CHECK1-NEXT:  entry:752// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l89.omp_outlined)753// CHECK1-NEXT:    ret void754//755//756// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l89.omp_outlined757// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {758// CHECK1-NEXT:  entry:759// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8760// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8761// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4762// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4763// CHECK1-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4764// CHECK1-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4765// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4766// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4767// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4768// CHECK1-NEXT:    [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 4769// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8770// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8771// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 4772// CHECK1-NEXT:    store i32 99, ptr [[DOTOMP_COMB_UB]], align 4773// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4774// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4775// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8776// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4777// 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)778// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4779// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 99780// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]781// CHECK1:       cond.true:782// CHECK1-NEXT:    br label [[COND_END:%.*]]783// CHECK1:       cond.false:784// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4785// CHECK1-NEXT:    br label [[COND_END]]786// CHECK1:       cond.end:787// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]788// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4789// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4790// CHECK1-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4791// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]792// CHECK1:       omp.inner.for.cond:793// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4794// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4795// CHECK1-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]796// CHECK1-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]797// CHECK1:       omp.inner.for.body:798// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4799// CHECK1-NEXT:    [[TMP8:%.*]] = zext i32 [[TMP7]] to i64800// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4801// CHECK1-NEXT:    [[TMP10:%.*]] = zext i32 [[TMP9]] to i64802// CHECK1-NEXT:    call void @__kmpc_serialized_parallel(ptr @[[GLOB3]], i32 [[TMP1]])803// CHECK1-NEXT:    [[TMP11:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8804// CHECK1-NEXT:    store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 4805// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l89.omp_outlined.omp_outlined(ptr [[TMP11]], ptr [[DOTBOUND_ZERO_ADDR]], i64 [[TMP8]], i64 [[TMP10]]) #[[ATTR2]]806// CHECK1-NEXT:    call void @__kmpc_end_serialized_parallel(ptr @[[GLOB3]], i32 [[TMP1]])807// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]808// CHECK1:       omp.inner.for.inc:809// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4810// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4811// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP12]], [[TMP13]]812// CHECK1-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4813// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]814// CHECK1:       omp.inner.for.end:815// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]816// CHECK1:       omp.loop.exit:817// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])818// CHECK1-NEXT:    ret void819//820//821// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l89.omp_outlined.omp_outlined822// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]]) #[[ATTR1]] {823// CHECK1-NEXT:  entry:824// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8825// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8826// CHECK1-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8827// CHECK1-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8828// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4829// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4830// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4831// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4832// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4833// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4834// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4835// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8836// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8837// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8838// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8839// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4840// CHECK1-NEXT:    store i32 99, ptr [[DOTOMP_UB]], align 4841// CHECK1-NEXT:    [[TMP0:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8842// CHECK1-NEXT:    [[CONV:%.*]] = trunc i64 [[TMP0]] to i32843// CHECK1-NEXT:    [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8844// CHECK1-NEXT:    [[CONV1:%.*]] = trunc i64 [[TMP1]] to i32845// CHECK1-NEXT:    store i32 [[CONV]], ptr [[DOTOMP_LB]], align 4846// CHECK1-NEXT:    store i32 [[CONV1]], ptr [[DOTOMP_UB]], align 4847// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4848// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4849// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8850// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4851// 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)852// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4853// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 99854// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]855// CHECK1:       cond.true:856// CHECK1-NEXT:    br label [[COND_END:%.*]]857// CHECK1:       cond.false:858// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4859// CHECK1-NEXT:    br label [[COND_END]]860// CHECK1:       cond.end:861// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]862// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4863// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4864// CHECK1-NEXT:    store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 4865// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]866// CHECK1:       omp.inner.for.cond:867// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4868// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4869// CHECK1-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]870// CHECK1-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]871// CHECK1:       omp.inner.for.body:872// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4873// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP9]], 1874// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]875// CHECK1-NEXT:    store i32 [[ADD]], ptr [[I]], align 4876// CHECK1-NEXT:    call void @_Z3fn5v()877// CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]878// CHECK1:       omp.body.continue:879// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]880// CHECK1:       omp.inner.for.inc:881// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4882// CHECK1-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP10]], 1883// CHECK1-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4884// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]885// CHECK1:       omp.inner.for.end:886// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]887// CHECK1:       omp.loop.exit:888// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP3]])889// CHECK1-NEXT:    ret void890//891//892// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l97893// CHECK1-SAME: (i64 noundef [[ARG:%.*]]) #[[ATTR1]] {894// CHECK1-NEXT:  entry:895// CHECK1-NEXT:    [[ARG_ADDR:%.*]] = alloca i64, align 8896// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i8, align 1897// CHECK1-NEXT:    [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 8898// CHECK1-NEXT:    store i64 [[ARG]], ptr [[ARG_ADDR]], align 8899// CHECK1-NEXT:    [[TMP0:%.*]] = load i32, ptr [[ARG_ADDR]], align 4900// CHECK1-NEXT:    [[TOBOOL:%.*]] = icmp ne i32 [[TMP0]], 0901// CHECK1-NEXT:    [[STOREDV:%.*]] = zext i1 [[TOBOOL]] to i8902// CHECK1-NEXT:    store i8 [[STOREDV]], ptr [[DOTCAPTURE_EXPR_]], align 1903// CHECK1-NEXT:    [[TMP1:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 1904// CHECK1-NEXT:    [[LOADEDV:%.*]] = trunc i8 [[TMP1]] to i1905// CHECK1-NEXT:    [[STOREDV1:%.*]] = zext i1 [[LOADEDV]] to i8906// CHECK1-NEXT:    store i8 [[STOREDV1]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 1907// CHECK1-NEXT:    [[TMP2:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 8908// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l97.omp_outlined, i64 [[TMP2]])909// CHECK1-NEXT:    ret void910//911//912// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l97.omp_outlined913// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1]] {914// CHECK1-NEXT:  entry:915// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8916// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8917// CHECK1-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 8918// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4919// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4920// CHECK1-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4921// CHECK1-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4922// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4923// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4924// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4925// CHECK1-NEXT:    [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 4926// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8927// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8928// CHECK1-NEXT:    store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 8929// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 4930// CHECK1-NEXT:    store i32 99, ptr [[DOTOMP_COMB_UB]], align 4931// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4932// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4933// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8934// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4935// 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)936// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4937// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 99938// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]939// CHECK1:       cond.true:940// CHECK1-NEXT:    br label [[COND_END:%.*]]941// CHECK1:       cond.false:942// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4943// CHECK1-NEXT:    br label [[COND_END]]944// CHECK1:       cond.end:945// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]946// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4947// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4948// CHECK1-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4949// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]950// CHECK1:       omp.inner.for.cond:951// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4952// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4953// CHECK1-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]954// CHECK1-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]955// CHECK1:       omp.inner.for.body:956// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4957// CHECK1-NEXT:    [[TMP8:%.*]] = zext i32 [[TMP7]] to i64958// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4959// CHECK1-NEXT:    [[TMP10:%.*]] = zext i32 [[TMP9]] to i64960// CHECK1-NEXT:    [[TMP11:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR__ADDR]], align 1961// CHECK1-NEXT:    [[LOADEDV:%.*]] = trunc i8 [[TMP11]] to i1962// CHECK1-NEXT:    br i1 [[LOADEDV]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]963// CHECK1:       omp_if.then:964// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l97.omp_outlined.omp_outlined, i64 [[TMP8]], i64 [[TMP10]])965// CHECK1-NEXT:    br label [[OMP_IF_END:%.*]]966// CHECK1:       omp_if.else:967// CHECK1-NEXT:    call void @__kmpc_serialized_parallel(ptr @[[GLOB3]], i32 [[TMP1]])968// CHECK1-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8969// CHECK1-NEXT:    store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 4970// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l97.omp_outlined.omp_outlined(ptr [[TMP12]], ptr [[DOTBOUND_ZERO_ADDR]], i64 [[TMP8]], i64 [[TMP10]]) #[[ATTR2]]971// CHECK1-NEXT:    call void @__kmpc_end_serialized_parallel(ptr @[[GLOB3]], i32 [[TMP1]])972// CHECK1-NEXT:    br label [[OMP_IF_END]]973// CHECK1:       omp_if.end:974// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]975// CHECK1:       omp.inner.for.inc:976// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4977// CHECK1-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4978// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP13]], [[TMP14]]979// CHECK1-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4980// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]981// CHECK1:       omp.inner.for.end:982// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]983// CHECK1:       omp.loop.exit:984// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])985// CHECK1-NEXT:    ret void986//987//988// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l97.omp_outlined.omp_outlined989// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]]) #[[ATTR1]] {990// CHECK1-NEXT:  entry:991// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8992// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8993// CHECK1-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8994// CHECK1-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8995// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4996// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4997// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4998// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4999// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41000// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41001// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 41002// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81003// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81004// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 81005// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 81006// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 41007// CHECK1-NEXT:    store i32 99, ptr [[DOTOMP_UB]], align 41008// CHECK1-NEXT:    [[TMP0:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 81009// CHECK1-NEXT:    [[CONV:%.*]] = trunc i64 [[TMP0]] to i321010// CHECK1-NEXT:    [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 81011// CHECK1-NEXT:    [[CONV1:%.*]] = trunc i64 [[TMP1]] to i321012// CHECK1-NEXT:    store i32 [[CONV]], ptr [[DOTOMP_LB]], align 41013// CHECK1-NEXT:    store i32 [[CONV1]], ptr [[DOTOMP_UB]], align 41014// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41015// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41016// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81017// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 41018// 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)1019// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41020// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 991021// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1022// CHECK1:       cond.true:1023// CHECK1-NEXT:    br label [[COND_END:%.*]]1024// CHECK1:       cond.false:1025// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41026// CHECK1-NEXT:    br label [[COND_END]]1027// CHECK1:       cond.end:1028// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]1029// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 41030// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41031// CHECK1-NEXT:    store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 41032// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1033// CHECK1:       omp.inner.for.cond:1034// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41035// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41036// CHECK1-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]1037// CHECK1-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1038// CHECK1:       omp.inner.for.body:1039// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41040// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP9]], 11041// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1042// CHECK1-NEXT:    store i32 [[ADD]], ptr [[I]], align 41043// CHECK1-NEXT:    call void @_Z3fn6v()1044// CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1045// CHECK1:       omp.body.continue:1046// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1047// CHECK1:       omp.inner.for.inc:1048// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41049// CHECK1-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP10]], 11050// CHECK1-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 41051// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]1052// CHECK1:       omp.inner.for.end:1053// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1054// CHECK1:       omp.loop.exit:1055// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP3]])1056// CHECK1-NEXT:    ret void1057//1058//1059// CHECK1-LABEL: define {{[^@]+}}@_Z5tmainIiEiT_1060// CHECK1-SAME: (i32 noundef [[ARG:%.*]]) #[[ATTR0]] comdat {1061// CHECK1-NEXT:  entry:1062// CHECK1-NEXT:    [[ARG_ADDR:%.*]] = alloca i32, align 41063// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 41064// CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81065// CHECK1-NEXT:    [[_TMP1:%.*]] = alloca i32, align 41066// CHECK1-NEXT:    [[KERNEL_ARGS2:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 81067// CHECK1-NEXT:    [[ARG_CASTED:%.*]] = alloca i64, align 81068// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 81069// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 81070// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 81071// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i8, align 11072// CHECK1-NEXT:    [[_TMP5:%.*]] = alloca i32, align 41073// CHECK1-NEXT:    [[KERNEL_ARGS6:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 81074// CHECK1-NEXT:    store i32 [[ARG]], ptr [[ARG_ADDR]], align 41075// CHECK1-NEXT:    [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01076// CHECK1-NEXT:    store i32 3, ptr [[TMP0]], align 41077// CHECK1-NEXT:    [[TMP1:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11078// CHECK1-NEXT:    store i32 0, ptr [[TMP1]], align 41079// CHECK1-NEXT:    [[TMP2:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21080// CHECK1-NEXT:    store ptr null, ptr [[TMP2]], align 81081// CHECK1-NEXT:    [[TMP3:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31082// CHECK1-NEXT:    store ptr null, ptr [[TMP3]], align 81083// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41084// CHECK1-NEXT:    store ptr null, ptr [[TMP4]], align 81085// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51086// CHECK1-NEXT:    store ptr null, ptr [[TMP5]], align 81087// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61088// CHECK1-NEXT:    store ptr null, ptr [[TMP6]], align 81089// CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71090// CHECK1-NEXT:    store ptr null, ptr [[TMP7]], align 81091// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81092// CHECK1-NEXT:    store i64 100, ptr [[TMP8]], align 81093// CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91094// CHECK1-NEXT:    store i64 0, ptr [[TMP9]], align 81095// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101096// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP10]], align 41097// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111098// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP11]], align 41099// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121100// CHECK1-NEXT:    store i32 0, ptr [[TMP12]], align 41101// 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__l62.region_id, ptr [[KERNEL_ARGS]])1102// CHECK1-NEXT:    [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 01103// CHECK1-NEXT:    br i1 [[TMP14]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1104// CHECK1:       omp_offload.failed:1105// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l62() #[[ATTR2]]1106// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]1107// CHECK1:       omp_offload.cont:1108// CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 01109// CHECK1-NEXT:    store i32 3, ptr [[TMP15]], align 41110// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 11111// CHECK1-NEXT:    store i32 0, ptr [[TMP16]], align 41112// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 21113// CHECK1-NEXT:    store ptr null, ptr [[TMP17]], align 81114// CHECK1-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 31115// CHECK1-NEXT:    store ptr null, ptr [[TMP18]], align 81116// CHECK1-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 41117// CHECK1-NEXT:    store ptr null, ptr [[TMP19]], align 81118// CHECK1-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 51119// CHECK1-NEXT:    store ptr null, ptr [[TMP20]], align 81120// CHECK1-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 61121// CHECK1-NEXT:    store ptr null, ptr [[TMP21]], align 81122// CHECK1-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 71123// CHECK1-NEXT:    store ptr null, ptr [[TMP22]], align 81124// CHECK1-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 81125// CHECK1-NEXT:    store i64 100, ptr [[TMP23]], align 81126// CHECK1-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 91127// CHECK1-NEXT:    store i64 0, ptr [[TMP24]], align 81128// CHECK1-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 101129// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP25]], align 41130// CHECK1-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 111131// CHECK1-NEXT:    store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP26]], align 41132// CHECK1-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS2]], i32 0, i32 121133// CHECK1-NEXT:    store i32 0, ptr [[TMP27]], align 41134// CHECK1-NEXT:    [[TMP28:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l67.region_id, ptr [[KERNEL_ARGS2]])1135// CHECK1-NEXT:    [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 01136// CHECK1-NEXT:    br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED3:%.*]], label [[OMP_OFFLOAD_CONT4:%.*]]1137// CHECK1:       omp_offload.failed3:1138// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l67() #[[ATTR2]]1139// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT4]]1140// CHECK1:       omp_offload.cont4:1141// CHECK1-NEXT:    [[TMP30:%.*]] = load i32, ptr [[ARG_ADDR]], align 41142// CHECK1-NEXT:    store i32 [[TMP30]], ptr [[ARG_CASTED]], align 41143// CHECK1-NEXT:    [[TMP31:%.*]] = load i64, ptr [[ARG_CASTED]], align 81144// CHECK1-NEXT:    [[TMP32:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01145// CHECK1-NEXT:    store i64 [[TMP31]], ptr [[TMP32]], align 81146// CHECK1-NEXT:    [[TMP33:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01147// CHECK1-NEXT:    store i64 [[TMP31]], ptr [[TMP33]], align 81148// CHECK1-NEXT:    [[TMP34:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 01149// CHECK1-NEXT:    store ptr null, ptr [[TMP34]], align 81150// CHECK1-NEXT:    [[TMP35:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01151// CHECK1-NEXT:    [[TMP36:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01152// CHECK1-NEXT:    [[TMP37:%.*]] = load i32, ptr [[ARG_ADDR]], align 41153// CHECK1-NEXT:    [[TOBOOL:%.*]] = icmp ne i32 [[TMP37]], 01154// CHECK1-NEXT:    [[STOREDV:%.*]] = zext i1 [[TOBOOL]] to i81155// CHECK1-NEXT:    store i8 [[STOREDV]], ptr [[DOTCAPTURE_EXPR_]], align 11156// CHECK1-NEXT:    [[TMP38:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 11157// CHECK1-NEXT:    [[LOADEDV:%.*]] = trunc i8 [[TMP38]] to i11158// CHECK1-NEXT:    [[TMP39:%.*]] = select i1 [[LOADEDV]], i32 0, i32 11159// CHECK1-NEXT:    [[TMP40:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP39]], 01160// CHECK1-NEXT:    [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 01161// CHECK1-NEXT:    store i32 3, ptr [[TMP41]], align 41162// CHECK1-NEXT:    [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 11163// CHECK1-NEXT:    store i32 1, ptr [[TMP42]], align 41164// CHECK1-NEXT:    [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 21165// CHECK1-NEXT:    store ptr [[TMP35]], ptr [[TMP43]], align 81166// CHECK1-NEXT:    [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 31167// CHECK1-NEXT:    store ptr [[TMP36]], ptr [[TMP44]], align 81168// CHECK1-NEXT:    [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 41169// CHECK1-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP45]], align 81170// CHECK1-NEXT:    [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 51171// CHECK1-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP46]], align 81172// CHECK1-NEXT:    [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 61173// CHECK1-NEXT:    store ptr null, ptr [[TMP47]], align 81174// CHECK1-NEXT:    [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 71175// CHECK1-NEXT:    store ptr null, ptr [[TMP48]], align 81176// CHECK1-NEXT:    [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 81177// CHECK1-NEXT:    store i64 100, ptr [[TMP49]], align 81178// CHECK1-NEXT:    [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 91179// CHECK1-NEXT:    store i64 0, ptr [[TMP50]], align 81180// CHECK1-NEXT:    [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 101181// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP51]], align 41182// CHECK1-NEXT:    [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 111183// CHECK1-NEXT:    store [3 x i32] [[TMP40]], ptr [[TMP52]], align 41184// CHECK1-NEXT:    [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 121185// CHECK1-NEXT:    store i32 0, ptr [[TMP53]], align 41186// CHECK1-NEXT:    [[TMP54:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 [[TMP39]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l72.region_id, ptr [[KERNEL_ARGS6]])1187// CHECK1-NEXT:    [[TMP55:%.*]] = icmp ne i32 [[TMP54]], 01188// CHECK1-NEXT:    br i1 [[TMP55]], label [[OMP_OFFLOAD_FAILED7:%.*]], label [[OMP_OFFLOAD_CONT8:%.*]]1189// CHECK1:       omp_offload.failed7:1190// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l72(i64 [[TMP31]]) #[[ATTR2]]1191// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT8]]1192// CHECK1:       omp_offload.cont8:1193// CHECK1-NEXT:    ret i32 01194//1195//1196// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l621197// CHECK1-SAME: () #[[ATTR1]] {1198// CHECK1-NEXT:  entry:1199// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l62.omp_outlined)1200// CHECK1-NEXT:    ret void1201//1202//1203// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l62.omp_outlined1204// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {1205// CHECK1-NEXT:  entry:1206// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81207// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81208// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41209// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 41210// CHECK1-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41211// CHECK1-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41212// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41213// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41214// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 41215// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81216// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81217// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 41218// CHECK1-NEXT:    store i32 99, ptr [[DOTOMP_COMB_UB]], align 41219// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41220// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41221// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81222// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 41223// 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)1224// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41225// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 991226// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1227// CHECK1:       cond.true:1228// CHECK1-NEXT:    br label [[COND_END:%.*]]1229// CHECK1:       cond.false:1230// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41231// CHECK1-NEXT:    br label [[COND_END]]1232// CHECK1:       cond.end:1233// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]1234// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41235// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41236// CHECK1-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 41237// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1238// CHECK1:       omp.inner.for.cond:1239// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41240// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41241// CHECK1-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]1242// CHECK1-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1243// CHECK1:       omp.inner.for.body:1244// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41245// CHECK1-NEXT:    [[TMP8:%.*]] = zext i32 [[TMP7]] to i641246// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41247// CHECK1-NEXT:    [[TMP10:%.*]] = zext i32 [[TMP9]] to i641248// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l62.omp_outlined.omp_outlined, i64 [[TMP8]], i64 [[TMP10]])1249// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1250// CHECK1:       omp.inner.for.inc:1251// CHECK1-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41252// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41253// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP11]], [[TMP12]]1254// CHECK1-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 41255// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]1256// CHECK1:       omp.inner.for.end:1257// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1258// CHECK1:       omp.loop.exit:1259// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])1260// CHECK1-NEXT:    ret void1261//1262//1263// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l62.omp_outlined.omp_outlined1264// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]]) #[[ATTR1]] {1265// CHECK1-NEXT:  entry:1266// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81267// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81268// CHECK1-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 81269// CHECK1-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 81270// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41271// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 41272// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 41273// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 41274// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41275// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41276// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 41277// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81278// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81279// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 81280// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 81281// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 41282// CHECK1-NEXT:    store i32 99, ptr [[DOTOMP_UB]], align 41283// CHECK1-NEXT:    [[TMP0:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 81284// CHECK1-NEXT:    [[CONV:%.*]] = trunc i64 [[TMP0]] to i321285// CHECK1-NEXT:    [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 81286// CHECK1-NEXT:    [[CONV1:%.*]] = trunc i64 [[TMP1]] to i321287// CHECK1-NEXT:    store i32 [[CONV]], ptr [[DOTOMP_LB]], align 41288// CHECK1-NEXT:    store i32 [[CONV1]], ptr [[DOTOMP_UB]], align 41289// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41290// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41291// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81292// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 41293// 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)1294// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41295// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 991296// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1297// CHECK1:       cond.true:1298// CHECK1-NEXT:    br label [[COND_END:%.*]]1299// CHECK1:       cond.false:1300// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41301// CHECK1-NEXT:    br label [[COND_END]]1302// CHECK1:       cond.end:1303// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]1304// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 41305// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41306// CHECK1-NEXT:    store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 41307// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1308// CHECK1:       omp.inner.for.cond:1309// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41310// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41311// CHECK1-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]1312// CHECK1-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1313// CHECK1:       omp.inner.for.body:1314// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41315// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP9]], 11316// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1317// CHECK1-NEXT:    store i32 [[ADD]], ptr [[I]], align 41318// CHECK1-NEXT:    call void @_Z3fn1v()1319// CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1320// CHECK1:       omp.body.continue:1321// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1322// CHECK1:       omp.inner.for.inc:1323// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41324// CHECK1-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP10]], 11325// CHECK1-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 41326// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]1327// CHECK1:       omp.inner.for.end:1328// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1329// CHECK1:       omp.loop.exit:1330// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP3]])1331// CHECK1-NEXT:    ret void1332//1333//1334// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l671335// CHECK1-SAME: () #[[ATTR1]] {1336// CHECK1-NEXT:  entry:1337// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l67.omp_outlined)1338// CHECK1-NEXT:    ret void1339//1340//1341// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l67.omp_outlined1342// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {1343// CHECK1-NEXT:  entry:1344// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81345// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81346// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41347// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 41348// CHECK1-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41349// CHECK1-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41350// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41351// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41352// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 41353// CHECK1-NEXT:    [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 41354// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81355// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81356// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 41357// CHECK1-NEXT:    store i32 99, ptr [[DOTOMP_COMB_UB]], align 41358// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41359// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41360// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81361// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 41362// 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)1363// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41364// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 991365// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1366// CHECK1:       cond.true:1367// CHECK1-NEXT:    br label [[COND_END:%.*]]1368// CHECK1:       cond.false:1369// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41370// CHECK1-NEXT:    br label [[COND_END]]1371// CHECK1:       cond.end:1372// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]1373// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41374// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41375// CHECK1-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 41376// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1377// CHECK1:       omp.inner.for.cond:1378// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41379// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41380// CHECK1-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]1381// CHECK1-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1382// CHECK1:       omp.inner.for.body:1383// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41384// CHECK1-NEXT:    [[TMP8:%.*]] = zext i32 [[TMP7]] to i641385// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41386// CHECK1-NEXT:    [[TMP10:%.*]] = zext i32 [[TMP9]] to i641387// CHECK1-NEXT:    call void @__kmpc_serialized_parallel(ptr @[[GLOB3]], i32 [[TMP1]])1388// CHECK1-NEXT:    [[TMP11:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81389// CHECK1-NEXT:    store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 41390// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l67.omp_outlined.omp_outlined(ptr [[TMP11]], ptr [[DOTBOUND_ZERO_ADDR]], i64 [[TMP8]], i64 [[TMP10]]) #[[ATTR2]]1391// CHECK1-NEXT:    call void @__kmpc_end_serialized_parallel(ptr @[[GLOB3]], i32 [[TMP1]])1392// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1393// CHECK1:       omp.inner.for.inc:1394// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41395// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41396// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP12]], [[TMP13]]1397// CHECK1-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 41398// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]1399// CHECK1:       omp.inner.for.end:1400// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1401// CHECK1:       omp.loop.exit:1402// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])1403// CHECK1-NEXT:    ret void1404//1405//1406// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l67.omp_outlined.omp_outlined1407// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]]) #[[ATTR1]] {1408// CHECK1-NEXT:  entry:1409// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81410// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81411// CHECK1-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 81412// CHECK1-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 81413// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41414// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 41415// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 41416// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 41417// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41418// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41419// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 41420// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81421// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81422// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 81423// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 81424// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 41425// CHECK1-NEXT:    store i32 99, ptr [[DOTOMP_UB]], align 41426// CHECK1-NEXT:    [[TMP0:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 81427// CHECK1-NEXT:    [[CONV:%.*]] = trunc i64 [[TMP0]] to i321428// CHECK1-NEXT:    [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 81429// CHECK1-NEXT:    [[CONV1:%.*]] = trunc i64 [[TMP1]] to i321430// CHECK1-NEXT:    store i32 [[CONV]], ptr [[DOTOMP_LB]], align 41431// CHECK1-NEXT:    store i32 [[CONV1]], ptr [[DOTOMP_UB]], align 41432// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41433// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41434// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81435// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 41436// 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)1437// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41438// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 991439// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1440// CHECK1:       cond.true:1441// CHECK1-NEXT:    br label [[COND_END:%.*]]1442// CHECK1:       cond.false:1443// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41444// CHECK1-NEXT:    br label [[COND_END]]1445// CHECK1:       cond.end:1446// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]1447// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 41448// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41449// CHECK1-NEXT:    store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 41450// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1451// CHECK1:       omp.inner.for.cond:1452// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41453// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41454// CHECK1-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]1455// CHECK1-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1456// CHECK1:       omp.inner.for.body:1457// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41458// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP9]], 11459// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1460// CHECK1-NEXT:    store i32 [[ADD]], ptr [[I]], align 41461// CHECK1-NEXT:    call void @_Z3fn2v()1462// CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1463// CHECK1:       omp.body.continue:1464// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1465// CHECK1:       omp.inner.for.inc:1466// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41467// CHECK1-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP10]], 11468// CHECK1-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 41469// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]1470// CHECK1:       omp.inner.for.end:1471// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1472// CHECK1:       omp.loop.exit:1473// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP3]])1474// CHECK1-NEXT:    ret void1475//1476//1477// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l721478// CHECK1-SAME: (i64 noundef [[ARG:%.*]]) #[[ATTR1]] {1479// CHECK1-NEXT:  entry:1480// CHECK1-NEXT:    [[ARG_ADDR:%.*]] = alloca i64, align 81481// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i8, align 11482// CHECK1-NEXT:    [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 81483// CHECK1-NEXT:    store i64 [[ARG]], ptr [[ARG_ADDR]], align 81484// CHECK1-NEXT:    [[TMP0:%.*]] = load i32, ptr [[ARG_ADDR]], align 41485// CHECK1-NEXT:    [[TOBOOL:%.*]] = icmp ne i32 [[TMP0]], 01486// CHECK1-NEXT:    [[STOREDV:%.*]] = zext i1 [[TOBOOL]] to i81487// CHECK1-NEXT:    store i8 [[STOREDV]], ptr [[DOTCAPTURE_EXPR_]], align 11488// CHECK1-NEXT:    [[TMP1:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 11489// CHECK1-NEXT:    [[LOADEDV:%.*]] = trunc i8 [[TMP1]] to i11490// CHECK1-NEXT:    [[STOREDV1:%.*]] = zext i1 [[LOADEDV]] to i81491// CHECK1-NEXT:    store i8 [[STOREDV1]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 11492// CHECK1-NEXT:    [[TMP2:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 81493// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l72.omp_outlined, i64 [[TMP2]])1494// CHECK1-NEXT:    ret void1495//1496//1497// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l72.omp_outlined1498// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1]] {1499// CHECK1-NEXT:  entry:1500// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81501// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81502// CHECK1-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 81503// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41504// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 41505// CHECK1-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41506// CHECK1-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41507// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41508// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41509// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 41510// CHECK1-NEXT:    [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 41511// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81512// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81513// CHECK1-NEXT:    store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 81514// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 41515// CHECK1-NEXT:    store i32 99, ptr [[DOTOMP_COMB_UB]], align 41516// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41517// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41518// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81519// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 41520// 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)1521// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41522// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 991523// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1524// CHECK1:       cond.true:1525// CHECK1-NEXT:    br label [[COND_END:%.*]]1526// CHECK1:       cond.false:1527// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41528// CHECK1-NEXT:    br label [[COND_END]]1529// CHECK1:       cond.end:1530// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]1531// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41532// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41533// CHECK1-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 41534// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1535// CHECK1:       omp.inner.for.cond:1536// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41537// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41538// CHECK1-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]1539// CHECK1-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1540// CHECK1:       omp.inner.for.body:1541// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41542// CHECK1-NEXT:    [[TMP8:%.*]] = zext i32 [[TMP7]] to i641543// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41544// CHECK1-NEXT:    [[TMP10:%.*]] = zext i32 [[TMP9]] to i641545// CHECK1-NEXT:    [[TMP11:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR__ADDR]], align 11546// CHECK1-NEXT:    [[LOADEDV:%.*]] = trunc i8 [[TMP11]] to i11547// CHECK1-NEXT:    br i1 [[LOADEDV]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]1548// CHECK1:       omp_if.then:1549// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l72.omp_outlined.omp_outlined, i64 [[TMP8]], i64 [[TMP10]])1550// CHECK1-NEXT:    br label [[OMP_IF_END:%.*]]1551// CHECK1:       omp_if.else:1552// CHECK1-NEXT:    call void @__kmpc_serialized_parallel(ptr @[[GLOB3]], i32 [[TMP1]])1553// CHECK1-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81554// CHECK1-NEXT:    store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 41555// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l72.omp_outlined.omp_outlined(ptr [[TMP12]], ptr [[DOTBOUND_ZERO_ADDR]], i64 [[TMP8]], i64 [[TMP10]]) #[[ATTR2]]1556// CHECK1-NEXT:    call void @__kmpc_end_serialized_parallel(ptr @[[GLOB3]], i32 [[TMP1]])1557// CHECK1-NEXT:    br label [[OMP_IF_END]]1558// CHECK1:       omp_if.end:1559// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1560// CHECK1:       omp.inner.for.inc:1561// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41562// CHECK1-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41563// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP13]], [[TMP14]]1564// CHECK1-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 41565// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]1566// CHECK1:       omp.inner.for.end:1567// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1568// CHECK1:       omp.loop.exit:1569// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])1570// CHECK1-NEXT:    ret void1571//1572//1573// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l72.omp_outlined.omp_outlined1574// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]]) #[[ATTR1]] {1575// CHECK1-NEXT:  entry:1576// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81577// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81578// CHECK1-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 81579// CHECK1-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 81580// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41581// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 41582// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 41583// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 41584// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41585// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41586// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 41587// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81588// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81589// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 81590// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 81591// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 41592// CHECK1-NEXT:    store i32 99, ptr [[DOTOMP_UB]], align 41593// CHECK1-NEXT:    [[TMP0:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 81594// CHECK1-NEXT:    [[CONV:%.*]] = trunc i64 [[TMP0]] to i321595// CHECK1-NEXT:    [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 81596// CHECK1-NEXT:    [[CONV1:%.*]] = trunc i64 [[TMP1]] to i321597// CHECK1-NEXT:    store i32 [[CONV]], ptr [[DOTOMP_LB]], align 41598// CHECK1-NEXT:    store i32 [[CONV1]], ptr [[DOTOMP_UB]], align 41599// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41600// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41601// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81602// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 41603// 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)1604// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41605// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 991606// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1607// CHECK1:       cond.true:1608// CHECK1-NEXT:    br label [[COND_END:%.*]]1609// CHECK1:       cond.false:1610// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41611// CHECK1-NEXT:    br label [[COND_END]]1612// CHECK1:       cond.end:1613// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]1614// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 41615// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41616// CHECK1-NEXT:    store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 41617// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1618// CHECK1:       omp.inner.for.cond:1619// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41620// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41621// CHECK1-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]1622// CHECK1-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1623// CHECK1:       omp.inner.for.body:1624// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41625// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP9]], 11626// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1627// CHECK1-NEXT:    store i32 [[ADD]], ptr [[I]], align 41628// CHECK1-NEXT:    call void @_Z3fn3v()1629// CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1630// CHECK1:       omp.body.continue:1631// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1632// CHECK1:       omp.inner.for.inc:1633// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41634// CHECK1-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP10]], 11635// CHECK1-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 41636// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]1637// CHECK1:       omp.inner.for.end:1638// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1639// CHECK1:       omp.loop.exit:1640// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP3]])1641// CHECK1-NEXT:    ret void1642//1643