1008 lines · cpp
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _2// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK13// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple x86_64-unknown-unknown -emit-pch -o %t %s4// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK15 6// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"7// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple x86_64-unknown-unknown -emit-pch -o %t %s8// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"9 10// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK111// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple x86_64-unknown-unknown -emit-pch -o %t %s12// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK113 14// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"15// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple x86_64-unknown-unknown -emit-pch -o %t %s16// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple x86_64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"17 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 target teams loop49 for(int i = 0 ; i < 100; i++) {}50 51#pragma omp target teams loop if (target: false)52 for(int i = 0 ; i < 100; i++) {53 gtid_test();54 }55}56 57 58template <typename T>59int tmain(T Arg) {60#pragma omp target teams loop if (true)61 for(int i = 0 ; i < 100; i++) {62 fn1();63 }64#pragma omp target teams loop if (false)65 for(int i = 0 ; i < 100; i++) {66 fn2();67 }68#pragma omp target teams loop if (target: Arg)69 for(int i = 0 ; i < 100; i++) {70 fn3();71 }72 return 0;73}74 75int main() {76#pragma omp target teams loop if (true)77 for(int i = 0 ; i < 100; i++) {78 79 80 fn4();81 }82 83#pragma omp target teams loop if (false)84 for(int i = 0 ; i < 100; i++) {85 86 87 fn5();88 }89 90#pragma omp target teams loop if (Arg)91 for(int i = 0 ; i < 100; i++) {92 93 94 fn6();95 }96 97 return tmain(Arg);98}99 100 101 102 103 104 105// call void [[T_OUTLINE_FUN_3:@.+]](106 107#endif108// CHECK1-LABEL: define {{[^@]+}}@_Z9gtid_testv109// CHECK1-SAME: () #[[ATTR0:[0-9]+]] {110// CHECK1-NEXT: entry:111// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4112// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8113// CHECK1-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0114// CHECK1-NEXT: store i32 3, ptr [[TMP0]], align 4115// CHECK1-NEXT: [[TMP1:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1116// CHECK1-NEXT: store i32 0, ptr [[TMP1]], align 4117// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2118// CHECK1-NEXT: store ptr null, ptr [[TMP2]], align 8119// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3120// CHECK1-NEXT: store ptr null, ptr [[TMP3]], align 8121// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4122// CHECK1-NEXT: store ptr null, ptr [[TMP4]], align 8123// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5124// CHECK1-NEXT: store ptr null, ptr [[TMP5]], align 8125// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6126// CHECK1-NEXT: store ptr null, ptr [[TMP6]], align 8127// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7128// CHECK1-NEXT: store ptr null, ptr [[TMP7]], align 8129// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8130// CHECK1-NEXT: store i64 100, ptr [[TMP8]], align 8131// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9132// CHECK1-NEXT: store i64 0, ptr [[TMP9]], align 8133// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10134// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP10]], align 4135// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11136// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP11]], align 4137// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12138// CHECK1-NEXT: store i32 0, ptr [[TMP12]], align 4139// 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]])140// CHECK1-NEXT: [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 0141// CHECK1-NEXT: br i1 [[TMP14]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]142// CHECK1: omp_offload.failed:143// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l48() #[[ATTR2:[0-9]+]]144// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]145// CHECK1: omp_offload.cont:146// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l51() #[[ATTR2]]147// CHECK1-NEXT: ret void148//149//150// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l48151// CHECK1-SAME: () #[[ATTR1:[0-9]+]] {152// CHECK1-NEXT: entry:153// 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)154// CHECK1-NEXT: ret void155//156//157// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l48.omp_outlined158// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {159// CHECK1-NEXT: entry:160// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8161// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8162// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4163// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4164// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4165// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4166// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4167// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4168// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4169// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8170// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8171// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4172// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_COMB_UB]], align 4173// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4174// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4175// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8176// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4177// 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)178// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4179// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 99180// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]181// CHECK1: cond.true:182// CHECK1-NEXT: br label [[COND_END:%.*]]183// CHECK1: cond.false:184// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4185// CHECK1-NEXT: br label [[COND_END]]186// CHECK1: cond.end:187// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]188// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4189// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4190// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4191// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]192// CHECK1: omp.inner.for.cond:193// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4194// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4195// CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]196// CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]197// CHECK1: omp.inner.for.body:198// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4199// CHECK1-NEXT: [[TMP8:%.*]] = zext i32 [[TMP7]] to i64200// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4201// CHECK1-NEXT: [[TMP10:%.*]] = zext i32 [[TMP9]] to i64202// 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]])203// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]204// CHECK1: omp.inner.for.inc:205// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4206// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4207// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP11]], [[TMP12]]208// CHECK1-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4209// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]210// CHECK1: omp.inner.for.end:211// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]212// CHECK1: omp.loop.exit:213// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])214// CHECK1-NEXT: ret void215//216//217// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l48.omp_outlined.omp_outlined218// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]]) #[[ATTR1]] {219// CHECK1-NEXT: entry:220// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8221// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8222// CHECK1-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8223// CHECK1-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8224// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4225// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4226// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4227// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4228// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4229// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4230// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4231// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8232// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8233// CHECK1-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8234// CHECK1-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8235// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4236// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_UB]], align 4237// CHECK1-NEXT: [[TMP0:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8238// CHECK1-NEXT: [[CONV:%.*]] = trunc i64 [[TMP0]] to i32239// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8240// CHECK1-NEXT: [[CONV1:%.*]] = trunc i64 [[TMP1]] to i32241// CHECK1-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 4242// CHECK1-NEXT: store i32 [[CONV1]], ptr [[DOTOMP_UB]], align 4243// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4244// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4245// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8246// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4247// 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)248// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4249// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 99250// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]251// CHECK1: cond.true:252// CHECK1-NEXT: br label [[COND_END:%.*]]253// CHECK1: cond.false:254// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4255// CHECK1-NEXT: br label [[COND_END]]256// CHECK1: cond.end:257// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]258// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4259// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4260// CHECK1-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 4261// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]262// CHECK1: omp.inner.for.cond:263// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4264// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4265// CHECK1-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]266// CHECK1-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]267// CHECK1: omp.inner.for.body:268// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4269// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP9]], 1270// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]271// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4272// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]273// CHECK1: omp.body.continue:274// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]275// CHECK1: omp.inner.for.inc:276// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4277// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP10]], 1278// CHECK1-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4279// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]280// CHECK1: omp.inner.for.end:281// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]282// CHECK1: omp.loop.exit:283// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP3]])284// CHECK1-NEXT: ret void285//286//287// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l51288// CHECK1-SAME: () #[[ATTR1]] {289// CHECK1-NEXT: entry:290// 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_l51.omp_outlined)291// CHECK1-NEXT: ret void292//293//294// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9gtid_testv_l51.omp_outlined295// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {296// CHECK1-NEXT: entry:297// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8298// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8299// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4300// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4301// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4302// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4303// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4304// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4305// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4306// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8307// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8308// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4309// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_COMB_UB]], align 4310// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4311// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4312// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8313// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4314// 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)315// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4316// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 99317// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]318// CHECK1: cond.true:319// CHECK1-NEXT: br label [[COND_END:%.*]]320// CHECK1: cond.false:321// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4322// CHECK1-NEXT: br label [[COND_END]]323// CHECK1: cond.end:324// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]325// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4326// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4327// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4328// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]329// CHECK1: omp.inner.for.cond:330// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4331// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4332// CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]333// CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]334// CHECK1: omp.inner.for.body:335// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4336// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1337// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]338// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4339// CHECK1-NEXT: call void @_Z9gtid_testv()340// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]341// CHECK1: omp.body.continue:342// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]343// CHECK1: omp.inner.for.inc:344// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4345// CHECK1-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP8]], 1346// CHECK1-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4347// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]348// CHECK1: omp.inner.for.end:349// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]350// CHECK1: omp.loop.exit:351// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])352// CHECK1-NEXT: ret void353//354//355// CHECK1-LABEL: define {{[^@]+}}@main356// CHECK1-SAME: () #[[ATTR3:[0-9]+]] {357// CHECK1-NEXT: entry:358// CHECK1-NEXT: [[RETVAL:%.*]] = alloca i32, align 4359// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4360// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8361// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i8, align 1362// CHECK1-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 8363// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8364// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8365// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8366// CHECK1-NEXT: [[_TMP4:%.*]] = alloca i32, align 4367// CHECK1-NEXT: [[KERNEL_ARGS5:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8368// CHECK1-NEXT: store i32 0, ptr [[RETVAL]], align 4369// CHECK1-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0370// CHECK1-NEXT: store i32 3, ptr [[TMP0]], align 4371// CHECK1-NEXT: [[TMP1:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1372// CHECK1-NEXT: store i32 0, ptr [[TMP1]], align 4373// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2374// CHECK1-NEXT: store ptr null, ptr [[TMP2]], align 8375// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3376// CHECK1-NEXT: store ptr null, ptr [[TMP3]], align 8377// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4378// CHECK1-NEXT: store ptr null, ptr [[TMP4]], align 8379// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5380// CHECK1-NEXT: store ptr null, ptr [[TMP5]], align 8381// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6382// CHECK1-NEXT: store ptr null, ptr [[TMP6]], align 8383// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7384// CHECK1-NEXT: store ptr null, ptr [[TMP7]], align 8385// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8386// CHECK1-NEXT: store i64 100, ptr [[TMP8]], align 8387// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9388// CHECK1-NEXT: store i64 0, ptr [[TMP9]], align 8389// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10390// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP10]], align 4391// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11392// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP11]], align 4393// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12394// CHECK1-NEXT: store i32 0, ptr [[TMP12]], align 4395// 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_l76.region_id, ptr [[KERNEL_ARGS]])396// CHECK1-NEXT: [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 0397// CHECK1-NEXT: br i1 [[TMP14]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]398// CHECK1: omp_offload.failed:399// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l76() #[[ATTR2]]400// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]401// CHECK1: omp_offload.cont:402// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l83() #[[ATTR2]]403// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr @Arg, align 4404// CHECK1-NEXT: [[TOBOOL:%.*]] = icmp ne i32 [[TMP15]], 0405// CHECK1-NEXT: [[STOREDV:%.*]] = zext i1 [[TOBOOL]] to i8406// CHECK1-NEXT: store i8 [[STOREDV]], ptr [[DOTCAPTURE_EXPR_]], align 1407// CHECK1-NEXT: [[TMP16:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 1408// CHECK1-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP16]] to i1409// CHECK1-NEXT: [[STOREDV1:%.*]] = zext i1 [[LOADEDV]] to i8410// CHECK1-NEXT: store i8 [[STOREDV1]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 1411// CHECK1-NEXT: [[TMP17:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 8412// CHECK1-NEXT: [[TMP18:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 1413// CHECK1-NEXT: [[LOADEDV2:%.*]] = trunc i8 [[TMP18]] to i1414// CHECK1-NEXT: br i1 [[LOADEDV2]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]415// CHECK1: omp_if.then:416// CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0417// CHECK1-NEXT: store i64 [[TMP17]], ptr [[TMP19]], align 8418// CHECK1-NEXT: [[TMP20:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0419// CHECK1-NEXT: store i64 [[TMP17]], ptr [[TMP20]], align 8420// CHECK1-NEXT: [[TMP21:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0421// CHECK1-NEXT: store ptr null, ptr [[TMP21]], align 8422// CHECK1-NEXT: [[TMP22:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0423// CHECK1-NEXT: [[TMP23:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0424// CHECK1-NEXT: [[TMP24:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 1425// CHECK1-NEXT: [[LOADEDV3:%.*]] = trunc i8 [[TMP24]] to i1426// CHECK1-NEXT: [[TMP25:%.*]] = select i1 [[LOADEDV3]], i32 0, i32 1427// CHECK1-NEXT: [[TMP26:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP25]], 0428// CHECK1-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 0429// CHECK1-NEXT: store i32 3, ptr [[TMP27]], align 4430// CHECK1-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 1431// CHECK1-NEXT: store i32 1, ptr [[TMP28]], align 4432// CHECK1-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 2433// CHECK1-NEXT: store ptr [[TMP22]], ptr [[TMP29]], align 8434// CHECK1-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 3435// CHECK1-NEXT: store ptr [[TMP23]], ptr [[TMP30]], align 8436// CHECK1-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 4437// CHECK1-NEXT: store ptr @.offload_sizes, ptr [[TMP31]], align 8438// CHECK1-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 5439// CHECK1-NEXT: store ptr @.offload_maptypes, ptr [[TMP32]], align 8440// CHECK1-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 6441// CHECK1-NEXT: store ptr null, ptr [[TMP33]], align 8442// CHECK1-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 7443// CHECK1-NEXT: store ptr null, ptr [[TMP34]], align 8444// CHECK1-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 8445// CHECK1-NEXT: store i64 100, ptr [[TMP35]], align 8446// CHECK1-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 9447// CHECK1-NEXT: store i64 0, ptr [[TMP36]], align 8448// CHECK1-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 10449// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP37]], align 4450// CHECK1-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 11451// CHECK1-NEXT: store [3 x i32] [[TMP26]], ptr [[TMP38]], align 4452// CHECK1-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 12453// CHECK1-NEXT: store i32 0, ptr [[TMP39]], align 4454// CHECK1-NEXT: [[TMP40:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 [[TMP25]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l90.region_id, ptr [[KERNEL_ARGS5]])455// CHECK1-NEXT: [[TMP41:%.*]] = icmp ne i32 [[TMP40]], 0456// CHECK1-NEXT: br i1 [[TMP41]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]]457// CHECK1: omp_offload.failed6:458// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l90(i64 [[TMP17]]) #[[ATTR2]]459// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT7]]460// CHECK1: omp_offload.cont7:461// CHECK1-NEXT: br label [[OMP_IF_END:%.*]]462// CHECK1: omp_if.else:463// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l90(i64 [[TMP17]]) #[[ATTR2]]464// CHECK1-NEXT: br label [[OMP_IF_END]]465// CHECK1: omp_if.end:466// CHECK1-NEXT: [[TMP42:%.*]] = load i32, ptr @Arg, align 4467// CHECK1-NEXT: [[CALL:%.*]] = call noundef i32 @_Z5tmainIiEiT_(i32 noundef [[TMP42]])468// CHECK1-NEXT: ret i32 [[CALL]]469//470//471// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l76472// CHECK1-SAME: () #[[ATTR1]] {473// CHECK1-NEXT: entry:474// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l76.omp_outlined)475// CHECK1-NEXT: ret void476//477//478// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l76.omp_outlined479// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {480// CHECK1-NEXT: entry:481// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8482// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8483// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4484// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4485// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4486// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4487// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4488// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4489// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4490// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8491// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8492// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4493// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_COMB_UB]], align 4494// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4495// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4496// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8497// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4498// 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)499// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4500// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 99501// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]502// CHECK1: cond.true:503// CHECK1-NEXT: br label [[COND_END:%.*]]504// CHECK1: cond.false:505// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4506// CHECK1-NEXT: br label [[COND_END]]507// CHECK1: cond.end:508// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]509// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4510// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4511// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4512// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]513// CHECK1: omp.inner.for.cond:514// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4515// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4516// CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]517// CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]518// CHECK1: omp.inner.for.body:519// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4520// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1521// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]522// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4523// CHECK1-NEXT: call void @_Z3fn4v()524// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]525// CHECK1: omp.body.continue:526// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]527// CHECK1: omp.inner.for.inc:528// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4529// CHECK1-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP8]], 1530// CHECK1-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4531// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]532// CHECK1: omp.inner.for.end:533// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]534// CHECK1: omp.loop.exit:535// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])536// CHECK1-NEXT: ret void537//538//539// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l83540// CHECK1-SAME: () #[[ATTR1]] {541// CHECK1-NEXT: entry:542// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l83.omp_outlined)543// CHECK1-NEXT: ret void544//545//546// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l83.omp_outlined547// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {548// CHECK1-NEXT: entry:549// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8550// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8551// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4552// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4553// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4554// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4555// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4556// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4557// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4558// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8559// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8560// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4561// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_COMB_UB]], align 4562// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4563// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4564// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8565// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4566// 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)567// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4568// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 99569// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]570// CHECK1: cond.true:571// CHECK1-NEXT: br label [[COND_END:%.*]]572// CHECK1: cond.false:573// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4574// CHECK1-NEXT: br label [[COND_END]]575// CHECK1: cond.end:576// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]577// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4578// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4579// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4580// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]581// CHECK1: omp.inner.for.cond:582// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4583// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4584// CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]585// CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]586// CHECK1: omp.inner.for.body:587// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4588// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1589// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]590// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4591// CHECK1-NEXT: call void @_Z3fn5v()592// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]593// CHECK1: omp.body.continue:594// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]595// CHECK1: omp.inner.for.inc:596// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4597// CHECK1-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP8]], 1598// CHECK1-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4599// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]600// CHECK1: omp.inner.for.end:601// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]602// CHECK1: omp.loop.exit:603// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])604// CHECK1-NEXT: ret void605//606//607// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l90608// CHECK1-SAME: (i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1]] {609// CHECK1-NEXT: entry:610// CHECK1-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 8611// CHECK1-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 8612// CHECK1-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 8613// CHECK1-NEXT: [[TMP0:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR__ADDR]], align 1614// CHECK1-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP0]] to i1615// CHECK1-NEXT: [[STOREDV:%.*]] = zext i1 [[LOADEDV]] to i8616// CHECK1-NEXT: store i8 [[STOREDV]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 1617// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 8618// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l90.omp_outlined, i64 [[TMP1]])619// CHECK1-NEXT: ret void620//621//622// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l90.omp_outlined623// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1]] {624// CHECK1-NEXT: entry:625// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8626// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8627// CHECK1-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 8628// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4629// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4630// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4631// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4632// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4633// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4634// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4635// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8636// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8637// CHECK1-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 8638// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4639// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_COMB_UB]], align 4640// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4641// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4642// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8643// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4644// 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)645// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4646// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 99647// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]648// CHECK1: cond.true:649// CHECK1-NEXT: br label [[COND_END:%.*]]650// CHECK1: cond.false:651// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4652// CHECK1-NEXT: br label [[COND_END]]653// CHECK1: cond.end:654// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]655// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4656// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4657// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4658// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]659// CHECK1: omp.inner.for.cond:660// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4661// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4662// CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]663// CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]664// CHECK1: omp.inner.for.body:665// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4666// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1667// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]668// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4669// CHECK1-NEXT: call void @_Z3fn6v()670// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]671// CHECK1: omp.body.continue:672// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]673// CHECK1: omp.inner.for.inc:674// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4675// CHECK1-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP8]], 1676// CHECK1-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4677// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]678// CHECK1: omp.inner.for.end:679// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]680// CHECK1: omp.loop.exit:681// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])682// CHECK1-NEXT: ret void683//684//685// CHECK1-LABEL: define {{[^@]+}}@_Z5tmainIiEiT_686// CHECK1-SAME: (i32 noundef [[ARG:%.*]]) #[[ATTR0]] comdat {687// CHECK1-NEXT: entry:688// CHECK1-NEXT: [[ARG_ADDR:%.*]] = alloca i32, align 4689// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4690// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8691// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i8, align 1692// CHECK1-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 8693// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8694// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8695// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8696// CHECK1-NEXT: [[_TMP3:%.*]] = alloca i32, align 4697// CHECK1-NEXT: [[KERNEL_ARGS4:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8698// CHECK1-NEXT: store i32 [[ARG]], ptr [[ARG_ADDR]], align 4699// CHECK1-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0700// CHECK1-NEXT: store i32 3, ptr [[TMP0]], align 4701// CHECK1-NEXT: [[TMP1:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1702// CHECK1-NEXT: store i32 0, ptr [[TMP1]], align 4703// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2704// CHECK1-NEXT: store ptr null, ptr [[TMP2]], align 8705// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3706// CHECK1-NEXT: store ptr null, ptr [[TMP3]], align 8707// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4708// CHECK1-NEXT: store ptr null, ptr [[TMP4]], align 8709// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5710// CHECK1-NEXT: store ptr null, ptr [[TMP5]], align 8711// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6712// CHECK1-NEXT: store ptr null, ptr [[TMP6]], align 8713// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7714// CHECK1-NEXT: store ptr null, ptr [[TMP7]], align 8715// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8716// CHECK1-NEXT: store i64 100, ptr [[TMP8]], align 8717// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9718// CHECK1-NEXT: store i64 0, ptr [[TMP9]], align 8719// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10720// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP10]], align 4721// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11722// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP11]], align 4723// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12724// CHECK1-NEXT: store i32 0, ptr [[TMP12]], align 4725// 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__l60.region_id, ptr [[KERNEL_ARGS]])726// CHECK1-NEXT: [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 0727// CHECK1-NEXT: br i1 [[TMP14]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]728// CHECK1: omp_offload.failed:729// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l60() #[[ATTR2]]730// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]731// CHECK1: omp_offload.cont:732// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l64() #[[ATTR2]]733// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[ARG_ADDR]], align 4734// CHECK1-NEXT: [[TOBOOL:%.*]] = icmp ne i32 [[TMP15]], 0735// CHECK1-NEXT: [[STOREDV:%.*]] = zext i1 [[TOBOOL]] to i8736// CHECK1-NEXT: store i8 [[STOREDV]], ptr [[DOTCAPTURE_EXPR_]], align 1737// CHECK1-NEXT: [[TMP16:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 1738// CHECK1-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP16]] to i1739// CHECK1-NEXT: [[STOREDV1:%.*]] = zext i1 [[LOADEDV]] to i8740// CHECK1-NEXT: store i8 [[STOREDV1]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 1741// CHECK1-NEXT: [[TMP17:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 8742// CHECK1-NEXT: [[TMP18:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 1743// CHECK1-NEXT: [[LOADEDV2:%.*]] = trunc i8 [[TMP18]] to i1744// CHECK1-NEXT: br i1 [[LOADEDV2]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]745// CHECK1: omp_if.then:746// CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0747// CHECK1-NEXT: store i64 [[TMP17]], ptr [[TMP19]], align 8748// CHECK1-NEXT: [[TMP20:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0749// CHECK1-NEXT: store i64 [[TMP17]], ptr [[TMP20]], align 8750// CHECK1-NEXT: [[TMP21:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0751// CHECK1-NEXT: store ptr null, ptr [[TMP21]], align 8752// CHECK1-NEXT: [[TMP22:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0753// CHECK1-NEXT: [[TMP23:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0754// CHECK1-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 0755// CHECK1-NEXT: store i32 3, ptr [[TMP24]], align 4756// CHECK1-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 1757// CHECK1-NEXT: store i32 1, ptr [[TMP25]], align 4758// CHECK1-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 2759// CHECK1-NEXT: store ptr [[TMP22]], ptr [[TMP26]], align 8760// CHECK1-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 3761// CHECK1-NEXT: store ptr [[TMP23]], ptr [[TMP27]], align 8762// CHECK1-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 4763// CHECK1-NEXT: store ptr @.offload_sizes.1, ptr [[TMP28]], align 8764// CHECK1-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 5765// CHECK1-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP29]], align 8766// CHECK1-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 6767// CHECK1-NEXT: store ptr null, ptr [[TMP30]], align 8768// CHECK1-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 7769// CHECK1-NEXT: store ptr null, ptr [[TMP31]], align 8770// CHECK1-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 8771// CHECK1-NEXT: store i64 100, ptr [[TMP32]], align 8772// CHECK1-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 9773// CHECK1-NEXT: store i64 0, ptr [[TMP33]], align 8774// CHECK1-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 10775// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP34]], align 4776// CHECK1-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 11777// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP35]], align 4778// CHECK1-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 12779// CHECK1-NEXT: store i32 0, ptr [[TMP36]], align 4780// CHECK1-NEXT: [[TMP37:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l68.region_id, ptr [[KERNEL_ARGS4]])781// CHECK1-NEXT: [[TMP38:%.*]] = icmp ne i32 [[TMP37]], 0782// CHECK1-NEXT: br i1 [[TMP38]], label [[OMP_OFFLOAD_FAILED5:%.*]], label [[OMP_OFFLOAD_CONT6:%.*]]783// CHECK1: omp_offload.failed5:784// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l68(i64 [[TMP17]]) #[[ATTR2]]785// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT6]]786// CHECK1: omp_offload.cont6:787// CHECK1-NEXT: br label [[OMP_IF_END:%.*]]788// CHECK1: omp_if.else:789// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l68(i64 [[TMP17]]) #[[ATTR2]]790// CHECK1-NEXT: br label [[OMP_IF_END]]791// CHECK1: omp_if.end:792// CHECK1-NEXT: ret i32 0793//794//795// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l60796// CHECK1-SAME: () #[[ATTR1]] {797// CHECK1-NEXT: entry:798// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l60.omp_outlined)799// CHECK1-NEXT: ret void800//801//802// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l60.omp_outlined803// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {804// CHECK1-NEXT: entry:805// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8806// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8807// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4808// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4809// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4810// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4811// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4812// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4813// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4814// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8815// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8816// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4817// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_COMB_UB]], align 4818// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4819// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4820// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8821// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4822// 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)823// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4824// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 99825// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]826// CHECK1: cond.true:827// CHECK1-NEXT: br label [[COND_END:%.*]]828// CHECK1: cond.false:829// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4830// CHECK1-NEXT: br label [[COND_END]]831// CHECK1: cond.end:832// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]833// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4834// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4835// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4836// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]837// CHECK1: omp.inner.for.cond:838// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4839// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4840// CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]841// CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]842// CHECK1: omp.inner.for.body:843// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4844// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1845// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]846// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4847// CHECK1-NEXT: call void @_Z3fn1v()848// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]849// CHECK1: omp.body.continue:850// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]851// CHECK1: omp.inner.for.inc:852// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4853// CHECK1-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP8]], 1854// CHECK1-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4855// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]856// CHECK1: omp.inner.for.end:857// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]858// CHECK1: omp.loop.exit:859// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])860// CHECK1-NEXT: ret void861//862//863// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l64864// CHECK1-SAME: () #[[ATTR1]] {865// CHECK1-NEXT: entry:866// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l64.omp_outlined)867// CHECK1-NEXT: ret void868//869//870// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l64.omp_outlined871// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {872// CHECK1-NEXT: entry:873// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8874// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8875// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4876// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4877// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4878// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4879// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4880// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4881// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4882// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8883// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8884// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4885// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_COMB_UB]], align 4886// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4887// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4888// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8889// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4890// 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)891// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4892// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 99893// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]894// CHECK1: cond.true:895// CHECK1-NEXT: br label [[COND_END:%.*]]896// CHECK1: cond.false:897// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4898// CHECK1-NEXT: br label [[COND_END]]899// CHECK1: cond.end:900// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]901// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4902// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4903// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4904// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]905// CHECK1: omp.inner.for.cond:906// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4907// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4908// CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]909// CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]910// CHECK1: omp.inner.for.body:911// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4912// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1913// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]914// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4915// CHECK1-NEXT: call void @_Z3fn2v()916// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]917// CHECK1: omp.body.continue:918// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]919// CHECK1: omp.inner.for.inc:920// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4921// CHECK1-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP8]], 1922// CHECK1-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4923// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]924// CHECK1: omp.inner.for.end:925// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]926// CHECK1: omp.loop.exit:927// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])928// CHECK1-NEXT: ret void929//930//931// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l68932// CHECK1-SAME: (i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1]] {933// CHECK1-NEXT: entry:934// CHECK1-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 8935// CHECK1-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 8936// CHECK1-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 8937// CHECK1-NEXT: [[TMP0:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR__ADDR]], align 1938// CHECK1-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP0]] to i1939// CHECK1-NEXT: [[STOREDV:%.*]] = zext i1 [[LOADEDV]] to i8940// CHECK1-NEXT: store i8 [[STOREDV]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 1941// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 8942// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l68.omp_outlined, i64 [[TMP1]])943// CHECK1-NEXT: ret void944//945//946// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiEiT__l68.omp_outlined947// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1]] {948// CHECK1-NEXT: entry:949// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8950// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8951// CHECK1-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 8952// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4953// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4954// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4955// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4956// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4957// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4958// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4959// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8960// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8961// CHECK1-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 8962// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4963// CHECK1-NEXT: store i32 99, ptr [[DOTOMP_COMB_UB]], align 4964// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4965// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4966// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8967// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4968// 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)969// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4970// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 99971// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]972// CHECK1: cond.true:973// CHECK1-NEXT: br label [[COND_END:%.*]]974// CHECK1: cond.false:975// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4976// CHECK1-NEXT: br label [[COND_END]]977// CHECK1: cond.end:978// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]979// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4980// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4981// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4982// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]983// CHECK1: omp.inner.for.cond:984// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4985// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4986// CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]987// CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]988// CHECK1: omp.inner.for.body:989// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4990// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1991// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]992// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4993// CHECK1-NEXT: call void @_Z3fn3v()994// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]995// CHECK1: omp.body.continue:996// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]997// CHECK1: omp.inner.for.inc:998// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4999// CHECK1-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP8]], 11000// CHECK1-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 41001// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]1002// CHECK1: omp.inner.for.end:1003// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1004// CHECK1: omp.loop.exit:1005// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])1006// CHECK1-NEXT: ret void1007//1008