2659 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// expected-no-diagnostics3#ifndef HEADER4#define HEADER5// Test host codegen.6// RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK17// RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK28// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s9// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK210// RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK411// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s12// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK413 14// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"15// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s16// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"17// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"18// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s19// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"20 21// Test target codegen - host bc file has to be created first. (no significant differences with host version of target region)22// RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc23// RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK1024// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s25// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK1026// RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc27// RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix=CHECK1228// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s29// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK1230 31// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc32// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"33// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s34// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"35// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc36// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"37// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s38// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"39 40#ifdef CK141 42 43int target_teams_fun(int *g){44 int n = 1000;45 int a[1000];46 int te = n / 128;47 int th = 128;48 // discard n_addr49 // discard capture expressions for te and th50 51 #pragma omp target teams loop num_teams(te), thread_limit(th)52 for(int i = 0; i < n; i++) {53 a[i] = 0;54 }55 56 {{{57 #pragma omp target teams loop is_device_ptr(g)58 for(int i = 0; i < n; i++) {59 a[i] = g[0];60 }61 }}}62 63 // outlined target regions64 65 66 67 68 return a[0];69}70 71#endif // CK172#endif // HEADER73// CHECK1-LABEL: define {{[^@]+}}@_Z16target_teams_funPi74// CHECK1-SAME: (ptr noundef [[G:%.*]]) #[[ATTR0:[0-9]+]] {75// CHECK1-NEXT: entry:76// CHECK1-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 877// CHECK1-NEXT: [[N:%.*]] = alloca i32, align 478// CHECK1-NEXT: [[A:%.*]] = alloca [1000 x i32], align 479// CHECK1-NEXT: [[TE:%.*]] = alloca i32, align 480// CHECK1-NEXT: [[TH:%.*]] = alloca i32, align 481// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 482// CHECK1-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 483// CHECK1-NEXT: [[N_CASTED:%.*]] = alloca i64, align 884// CHECK1-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 885// CHECK1-NEXT: [[DOTCAPTURE_EXPR__CASTED2:%.*]] = alloca i64, align 886// CHECK1-NEXT: [[N_CASTED3:%.*]] = alloca i64, align 887// CHECK1-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 888// CHECK1-NEXT: store i32 1000, ptr [[N]], align 489// CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr [[N]], align 490// CHECK1-NEXT: [[DIV:%.*]] = sdiv i32 [[TMP0]], 12891// CHECK1-NEXT: store i32 [[DIV]], ptr [[TE]], align 492// CHECK1-NEXT: store i32 128, ptr [[TH]], align 493// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TE]], align 494// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 495// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[TH]], align 496// CHECK1-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_1]], align 497// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[N]], align 498// CHECK1-NEXT: store i32 [[TMP3]], ptr [[N_CASTED]], align 499// CHECK1-NEXT: [[TMP4:%.*]] = load i64, ptr [[N_CASTED]], align 8100// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4101// CHECK1-NEXT: store i32 [[TMP5]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 4102// CHECK1-NEXT: [[TMP6:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 8103// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4104// CHECK1-NEXT: store i32 [[TMP7]], ptr [[DOTCAPTURE_EXPR__CASTED2]], align 4105// CHECK1-NEXT: [[TMP8:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED2]], align 8106// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51(i64 [[TMP4]], ptr [[A]], i64 [[TMP6]], i64 [[TMP8]]) #[[ATTR2:[0-9]+]]107// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[N]], align 4108// CHECK1-NEXT: store i32 [[TMP9]], ptr [[N_CASTED3]], align 4109// CHECK1-NEXT: [[TMP10:%.*]] = load i64, ptr [[N_CASTED3]], align 8110// CHECK1-NEXT: [[TMP11:%.*]] = load ptr, ptr [[G_ADDR]], align 8111// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57(i64 [[TMP10]], ptr [[A]], ptr [[TMP11]]) #[[ATTR2]]112// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[A]], i64 0, i64 0113// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[ARRAYIDX]], align 4114// CHECK1-NEXT: ret i32 [[TMP12]]115//116//117// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51118// CHECK1-SAME: (i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]], i64 noundef [[DOTCAPTURE_EXPR_1:%.*]]) #[[ATTR1:[0-9]+]] {119// CHECK1-NEXT: entry:120// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8121// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8122// CHECK1-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 8123// CHECK1-NEXT: [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca i64, align 8124// CHECK1-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8125// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB3:[0-9]+]])126// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8127// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8128// CHECK1-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 8129// CHECK1-NEXT: store i64 [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 8130// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 8131// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 4132// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR2]], align 4133// CHECK1-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB3]], i32 [[TMP0]], i32 [[TMP2]], i32 [[TMP3]])134// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[N_ADDR]], align 4135// CHECK1-NEXT: store i32 [[TMP4]], ptr [[N_CASTED]], align 4136// CHECK1-NEXT: [[TMP5:%.*]] = load i64, ptr [[N_CASTED]], align 8137// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined, i64 [[TMP5]], ptr [[TMP1]])138// CHECK1-NEXT: ret void139//140//141// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined142// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR1]] {143// CHECK1-NEXT: entry:144// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8145// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8146// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8147// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8148// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4149// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4150// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4151// CHECK1-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4152// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4153// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4154// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4155// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4156// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4157// CHECK1-NEXT: [[I3:%.*]] = alloca i32, align 4158// CHECK1-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8159// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8160// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8161// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8162// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8163// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8164// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 4165// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 4166// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4167// CHECK1-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 0168// CHECK1-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1169// CHECK1-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1170// CHECK1-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4171// CHECK1-NEXT: store i32 0, ptr [[I]], align 4172// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4173// CHECK1-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]174// CHECK1-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]175// CHECK1: omp.precond.then:176// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4177// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4178// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_COMB_UB]], align 4179// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4180// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4181// CHECK1-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8182// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 4183// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP6]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)184// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4185// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4186// CHECK1-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP7]], [[TMP8]]187// CHECK1-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]188// CHECK1: cond.true:189// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4190// CHECK1-NEXT: br label [[COND_END:%.*]]191// CHECK1: cond.false:192// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4193// CHECK1-NEXT: br label [[COND_END]]194// CHECK1: cond.end:195// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ [[TMP9]], [[COND_TRUE]] ], [ [[TMP10]], [[COND_FALSE]] ]196// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4197// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4198// CHECK1-NEXT: store i32 [[TMP11]], ptr [[DOTOMP_IV]], align 4199// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]200// CHECK1: omp.inner.for.cond:201// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4202// CHECK1-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4203// CHECK1-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP12]], [[TMP13]]204// CHECK1-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]205// CHECK1: omp.inner.for.body:206// CHECK1-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4207// CHECK1-NEXT: [[TMP15:%.*]] = zext i32 [[TMP14]] to i64208// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4209// CHECK1-NEXT: [[TMP17:%.*]] = zext i32 [[TMP16]] to i64210// CHECK1-NEXT: [[TMP18:%.*]] = load i32, ptr [[N_ADDR]], align 4211// CHECK1-NEXT: store i32 [[TMP18]], ptr [[N_CASTED]], align 4212// CHECK1-NEXT: [[TMP19:%.*]] = load i64, ptr [[N_CASTED]], align 8213// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 4, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined.omp_outlined, i64 [[TMP15]], i64 [[TMP17]], i64 [[TMP19]], ptr [[TMP0]])214// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]215// CHECK1: omp.inner.for.inc:216// CHECK1-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4217// CHECK1-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4218// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP20]], [[TMP21]]219// CHECK1-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4220// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]221// CHECK1: omp.inner.for.end:222// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]223// CHECK1: omp.loop.exit:224// CHECK1-NEXT: [[TMP22:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8225// CHECK1-NEXT: [[TMP23:%.*]] = load i32, ptr [[TMP22]], align 4226// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP23]])227// CHECK1-NEXT: br label [[OMP_PRECOND_END]]228// CHECK1: omp.precond.end:229// CHECK1-NEXT: ret void230//231//232// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined.omp_outlined233// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR1]] {234// CHECK1-NEXT: entry:235// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8236// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8237// CHECK1-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8238// CHECK1-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8239// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8240// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8241// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4242// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4243// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4244// CHECK1-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4245// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4246// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4247// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4248// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4249// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4250// CHECK1-NEXT: [[I4:%.*]] = alloca i32, align 4251// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8252// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8253// CHECK1-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8254// CHECK1-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8255// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8256// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8257// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8258// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 4259// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 4260// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4261// CHECK1-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 0262// CHECK1-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1263// CHECK1-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1264// CHECK1-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4265// CHECK1-NEXT: store i32 0, ptr [[I]], align 4266// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4267// CHECK1-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]268// CHECK1-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]269// CHECK1: omp.precond.then:270// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4271// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4272// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_UB]], align 4273// CHECK1-NEXT: [[TMP5:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8274// CHECK1-NEXT: [[CONV:%.*]] = trunc i64 [[TMP5]] to i32275// CHECK1-NEXT: [[TMP6:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8276// CHECK1-NEXT: [[CONV3:%.*]] = trunc i64 [[TMP6]] to i32277// CHECK1-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 4278// CHECK1-NEXT: store i32 [[CONV3]], ptr [[DOTOMP_UB]], align 4279// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4280// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4281// CHECK1-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8282// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4283// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP8]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)284// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4285// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4286// CHECK1-NEXT: [[CMP5:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]287// CHECK1-NEXT: br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]288// CHECK1: cond.true:289// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4290// CHECK1-NEXT: br label [[COND_END:%.*]]291// CHECK1: cond.false:292// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4293// CHECK1-NEXT: br label [[COND_END]]294// CHECK1: cond.end:295// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]296// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4297// CHECK1-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4298// CHECK1-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 4299// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]300// CHECK1: omp.inner.for.cond:301// CHECK1-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4302// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4303// CHECK1-NEXT: [[CMP6:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]304// CHECK1-NEXT: br i1 [[CMP6]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]305// CHECK1: omp.inner.for.body:306// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4307// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 1308// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]309// CHECK1-NEXT: store i32 [[ADD]], ptr [[I4]], align 4310// CHECK1-NEXT: [[TMP17:%.*]] = load i32, ptr [[I4]], align 4311// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP17]] to i64312// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]313// CHECK1-NEXT: store i32 0, ptr [[ARRAYIDX]], 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: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4319// CHECK1-NEXT: [[ADD7:%.*]] = add nsw i32 [[TMP18]], 1320// CHECK1-NEXT: store i32 [[ADD7]], 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: [[TMP19:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8326// CHECK1-NEXT: [[TMP20:%.*]] = load i32, ptr [[TMP19]], align 4327// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP20]])328// CHECK1-NEXT: br label [[OMP_PRECOND_END]]329// CHECK1: omp.precond.end:330// CHECK1-NEXT: ret void331//332//333// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57334// CHECK1-SAME: (i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], ptr noundef [[G:%.*]]) #[[ATTR1]] {335// CHECK1-NEXT: entry:336// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8337// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8338// CHECK1-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 8339// CHECK1-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8340// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8341// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8342// CHECK1-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 8343// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8344// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 4345// CHECK1-NEXT: store i32 [[TMP1]], ptr [[N_CASTED]], align 4346// CHECK1-NEXT: [[TMP2:%.*]] = load i64, ptr [[N_CASTED]], align 8347// CHECK1-NEXT: [[TMP3:%.*]] = load ptr, ptr [[G_ADDR]], align 8348// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.omp_outlined, i64 [[TMP2]], ptr [[TMP0]], ptr [[TMP3]])349// CHECK1-NEXT: ret void350//351//352// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.omp_outlined353// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], ptr noundef [[G:%.*]]) #[[ATTR1]] {354// CHECK1-NEXT: entry:355// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8356// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8357// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8358// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8359// CHECK1-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 8360// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4361// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4362// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4363// CHECK1-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4364// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4365// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4366// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4367// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4368// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4369// CHECK1-NEXT: [[I3:%.*]] = alloca i32, align 4370// CHECK1-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8371// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8372// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8373// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8374// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8375// CHECK1-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 8376// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8377// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 4378// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 4379// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4380// CHECK1-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 0381// CHECK1-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1382// CHECK1-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1383// CHECK1-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4384// CHECK1-NEXT: store i32 0, ptr [[I]], align 4385// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4386// CHECK1-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]387// CHECK1-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]388// CHECK1: omp.precond.then:389// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4390// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4391// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_COMB_UB]], align 4392// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4393// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4394// CHECK1-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8395// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 4396// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP6]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)397// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4398// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4399// CHECK1-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP7]], [[TMP8]]400// CHECK1-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]401// CHECK1: cond.true:402// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4403// CHECK1-NEXT: br label [[COND_END:%.*]]404// CHECK1: cond.false:405// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4406// CHECK1-NEXT: br label [[COND_END]]407// CHECK1: cond.end:408// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ [[TMP9]], [[COND_TRUE]] ], [ [[TMP10]], [[COND_FALSE]] ]409// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4410// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4411// CHECK1-NEXT: store i32 [[TMP11]], ptr [[DOTOMP_IV]], align 4412// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]413// CHECK1: omp.inner.for.cond:414// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4415// CHECK1-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4416// CHECK1-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP12]], [[TMP13]]417// CHECK1-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]418// CHECK1: omp.inner.for.body:419// CHECK1-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4420// CHECK1-NEXT: [[TMP15:%.*]] = zext i32 [[TMP14]] to i64421// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4422// CHECK1-NEXT: [[TMP17:%.*]] = zext i32 [[TMP16]] to i64423// CHECK1-NEXT: [[TMP18:%.*]] = load i32, ptr [[N_ADDR]], align 4424// CHECK1-NEXT: store i32 [[TMP18]], ptr [[N_CASTED]], align 4425// CHECK1-NEXT: [[TMP19:%.*]] = load i64, ptr [[N_CASTED]], align 8426// CHECK1-NEXT: [[TMP20:%.*]] = load ptr, ptr [[G_ADDR]], align 8427// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 5, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.omp_outlined.omp_outlined, i64 [[TMP15]], i64 [[TMP17]], i64 [[TMP19]], ptr [[TMP0]], ptr [[TMP20]])428// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]429// CHECK1: omp.inner.for.inc:430// CHECK1-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4431// CHECK1-NEXT: [[TMP22:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4432// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP21]], [[TMP22]]433// CHECK1-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4434// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]435// CHECK1: omp.inner.for.end:436// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]437// CHECK1: omp.loop.exit:438// CHECK1-NEXT: [[TMP23:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8439// CHECK1-NEXT: [[TMP24:%.*]] = load i32, ptr [[TMP23]], align 4440// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP24]])441// CHECK1-NEXT: br label [[OMP_PRECOND_END]]442// CHECK1: omp.precond.end:443// CHECK1-NEXT: ret void444//445//446// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.omp_outlined.omp_outlined447// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], ptr noundef [[G:%.*]]) #[[ATTR1]] {448// CHECK1-NEXT: entry:449// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8450// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8451// CHECK1-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8452// CHECK1-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8453// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8454// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8455// CHECK1-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 8456// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4457// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4458// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4459// CHECK1-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4460// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4461// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4462// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4463// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4464// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4465// CHECK1-NEXT: [[I4:%.*]] = alloca i32, align 4466// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8467// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8468// CHECK1-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8469// CHECK1-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8470// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8471// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8472// CHECK1-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 8473// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8474// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 4475// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 4476// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4477// CHECK1-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 0478// CHECK1-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1479// CHECK1-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1480// CHECK1-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4481// CHECK1-NEXT: store i32 0, ptr [[I]], align 4482// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4483// CHECK1-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]484// CHECK1-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]485// CHECK1: omp.precond.then:486// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4487// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4488// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_UB]], align 4489// CHECK1-NEXT: [[TMP5:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8490// CHECK1-NEXT: [[CONV:%.*]] = trunc i64 [[TMP5]] to i32491// CHECK1-NEXT: [[TMP6:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8492// CHECK1-NEXT: [[CONV3:%.*]] = trunc i64 [[TMP6]] to i32493// CHECK1-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 4494// CHECK1-NEXT: store i32 [[CONV3]], ptr [[DOTOMP_UB]], align 4495// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4496// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4497// CHECK1-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8498// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4499// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2]], i32 [[TMP8]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)500// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4501// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4502// CHECK1-NEXT: [[CMP5:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]503// CHECK1-NEXT: br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]504// CHECK1: cond.true:505// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4506// CHECK1-NEXT: br label [[COND_END:%.*]]507// CHECK1: cond.false:508// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4509// CHECK1-NEXT: br label [[COND_END]]510// CHECK1: cond.end:511// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]512// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4513// CHECK1-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4514// CHECK1-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 4515// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]516// CHECK1: omp.inner.for.cond:517// CHECK1-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4518// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4519// CHECK1-NEXT: [[CMP6:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]520// CHECK1-NEXT: br i1 [[CMP6]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]521// CHECK1: omp.inner.for.body:522// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4523// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 1524// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]525// CHECK1-NEXT: store i32 [[ADD]], ptr [[I4]], align 4526// CHECK1-NEXT: [[TMP17:%.*]] = load ptr, ptr [[G_ADDR]], align 8527// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP17]], i64 0528// CHECK1-NEXT: [[TMP18:%.*]] = load i32, ptr [[ARRAYIDX]], align 4529// CHECK1-NEXT: [[TMP19:%.*]] = load i32, ptr [[I4]], align 4530// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP19]] to i64531// CHECK1-NEXT: [[ARRAYIDX7:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]532// CHECK1-NEXT: store i32 [[TMP18]], ptr [[ARRAYIDX7]], align 4533// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]534// CHECK1: omp.body.continue:535// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]536// CHECK1: omp.inner.for.inc:537// CHECK1-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4538// CHECK1-NEXT: [[ADD8:%.*]] = add nsw i32 [[TMP20]], 1539// CHECK1-NEXT: store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 4540// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]541// CHECK1: omp.inner.for.end:542// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]543// CHECK1: omp.loop.exit:544// CHECK1-NEXT: [[TMP21:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8545// CHECK1-NEXT: [[TMP22:%.*]] = load i32, ptr [[TMP21]], align 4546// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP22]])547// CHECK1-NEXT: br label [[OMP_PRECOND_END]]548// CHECK1: omp.precond.end:549// CHECK1-NEXT: ret void550//551//552// CHECK2-LABEL: define {{[^@]+}}@_Z16target_teams_funPi553// CHECK2-SAME: (ptr noundef [[G:%.*]]) #[[ATTR0:[0-9]+]] {554// CHECK2-NEXT: entry:555// CHECK2-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 8556// CHECK2-NEXT: [[N:%.*]] = alloca i32, align 4557// CHECK2-NEXT: [[A:%.*]] = alloca [1000 x i32], align 4558// CHECK2-NEXT: [[TE:%.*]] = alloca i32, align 4559// CHECK2-NEXT: [[TH:%.*]] = alloca i32, align 4560// CHECK2-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4561// CHECK2-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4562// CHECK2-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8563// CHECK2-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 8564// CHECK2-NEXT: [[DOTCAPTURE_EXPR__CASTED2:%.*]] = alloca i64, align 8565// CHECK2-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [4 x ptr], align 8566// CHECK2-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [4 x ptr], align 8567// CHECK2-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [4 x ptr], align 8568// CHECK2-NEXT: [[TMP:%.*]] = alloca i32, align 4569// CHECK2-NEXT: [[DOTCAPTURE_EXPR_3:%.*]] = alloca i32, align 4570// CHECK2-NEXT: [[DOTCAPTURE_EXPR_4:%.*]] = alloca i32, align 4571// CHECK2-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8572// CHECK2-NEXT: [[N_CASTED7:%.*]] = alloca i64, align 8573// CHECK2-NEXT: [[DOTOFFLOAD_BASEPTRS8:%.*]] = alloca [3 x ptr], align 8574// CHECK2-NEXT: [[DOTOFFLOAD_PTRS9:%.*]] = alloca [3 x ptr], align 8575// CHECK2-NEXT: [[DOTOFFLOAD_MAPPERS10:%.*]] = alloca [3 x ptr], align 8576// CHECK2-NEXT: [[_TMP11:%.*]] = alloca i32, align 4577// CHECK2-NEXT: [[DOTCAPTURE_EXPR_12:%.*]] = alloca i32, align 4578// CHECK2-NEXT: [[DOTCAPTURE_EXPR_13:%.*]] = alloca i32, align 4579// CHECK2-NEXT: [[KERNEL_ARGS18:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8580// CHECK2-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 8581// CHECK2-NEXT: store i32 1000, ptr [[N]], align 4582// CHECK2-NEXT: [[TMP0:%.*]] = load i32, ptr [[N]], align 4583// CHECK2-NEXT: [[DIV:%.*]] = sdiv i32 [[TMP0]], 128584// CHECK2-NEXT: store i32 [[DIV]], ptr [[TE]], align 4585// CHECK2-NEXT: store i32 128, ptr [[TH]], align 4586// CHECK2-NEXT: [[TMP1:%.*]] = load i32, ptr [[TE]], align 4587// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 4588// CHECK2-NEXT: [[TMP2:%.*]] = load i32, ptr [[TH]], align 4589// CHECK2-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_1]], align 4590// CHECK2-NEXT: [[TMP3:%.*]] = load i32, ptr [[N]], align 4591// CHECK2-NEXT: store i32 [[TMP3]], ptr [[N_CASTED]], align 4592// CHECK2-NEXT: [[TMP4:%.*]] = load i64, ptr [[N_CASTED]], align 8593// CHECK2-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4594// CHECK2-NEXT: store i32 [[TMP5]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 4595// CHECK2-NEXT: [[TMP6:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 8596// CHECK2-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4597// CHECK2-NEXT: store i32 [[TMP7]], ptr [[DOTCAPTURE_EXPR__CASTED2]], align 4598// CHECK2-NEXT: [[TMP8:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED2]], align 8599// CHECK2-NEXT: [[TMP9:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0600// CHECK2-NEXT: store i64 [[TMP4]], ptr [[TMP9]], align 8601// CHECK2-NEXT: [[TMP10:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0602// CHECK2-NEXT: store i64 [[TMP4]], ptr [[TMP10]], align 8603// CHECK2-NEXT: [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0604// CHECK2-NEXT: store ptr null, ptr [[TMP11]], align 8605// CHECK2-NEXT: [[TMP12:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1606// CHECK2-NEXT: store ptr [[A]], ptr [[TMP12]], align 8607// CHECK2-NEXT: [[TMP13:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1608// CHECK2-NEXT: store ptr [[A]], ptr [[TMP13]], align 8609// CHECK2-NEXT: [[TMP14:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1610// CHECK2-NEXT: store ptr null, ptr [[TMP14]], align 8611// CHECK2-NEXT: [[TMP15:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2612// CHECK2-NEXT: store i64 [[TMP6]], ptr [[TMP15]], align 8613// CHECK2-NEXT: [[TMP16:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2614// CHECK2-NEXT: store i64 [[TMP6]], ptr [[TMP16]], align 8615// CHECK2-NEXT: [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2616// CHECK2-NEXT: store ptr null, ptr [[TMP17]], align 8617// CHECK2-NEXT: [[TMP18:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3618// CHECK2-NEXT: store i64 [[TMP8]], ptr [[TMP18]], align 8619// CHECK2-NEXT: [[TMP19:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3620// CHECK2-NEXT: store i64 [[TMP8]], ptr [[TMP19]], align 8621// CHECK2-NEXT: [[TMP20:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3622// CHECK2-NEXT: store ptr null, ptr [[TMP20]], align 8623// CHECK2-NEXT: [[TMP21:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0624// CHECK2-NEXT: [[TMP22:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0625// CHECK2-NEXT: [[TMP23:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4626// CHECK2-NEXT: [[TMP24:%.*]] = load i32, ptr [[N]], align 4627// CHECK2-NEXT: store i32 [[TMP24]], ptr [[DOTCAPTURE_EXPR_3]], align 4628// CHECK2-NEXT: [[TMP25:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_3]], align 4629// CHECK2-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP25]], 0630// CHECK2-NEXT: [[DIV5:%.*]] = sdiv i32 [[SUB]], 1631// CHECK2-NEXT: [[SUB6:%.*]] = sub nsw i32 [[DIV5]], 1632// CHECK2-NEXT: store i32 [[SUB6]], ptr [[DOTCAPTURE_EXPR_4]], align 4633// CHECK2-NEXT: [[TMP26:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 4634// CHECK2-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP26]], 1635// CHECK2-NEXT: [[TMP27:%.*]] = zext i32 [[ADD]] to i64636// CHECK2-NEXT: [[TMP28:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP23]], 0637// CHECK2-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0638// CHECK2-NEXT: store i32 3, ptr [[TMP29]], align 4639// CHECK2-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1640// CHECK2-NEXT: store i32 4, ptr [[TMP30]], align 4641// CHECK2-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2642// CHECK2-NEXT: store ptr [[TMP21]], ptr [[TMP31]], align 8643// CHECK2-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3644// CHECK2-NEXT: store ptr [[TMP22]], ptr [[TMP32]], align 8645// CHECK2-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4646// CHECK2-NEXT: store ptr @.offload_sizes, ptr [[TMP33]], align 8647// CHECK2-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5648// CHECK2-NEXT: store ptr @.offload_maptypes, ptr [[TMP34]], align 8649// CHECK2-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6650// CHECK2-NEXT: store ptr null, ptr [[TMP35]], align 8651// CHECK2-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7652// CHECK2-NEXT: store ptr null, ptr [[TMP36]], align 8653// CHECK2-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8654// CHECK2-NEXT: store i64 [[TMP27]], ptr [[TMP37]], align 8655// CHECK2-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9656// CHECK2-NEXT: store i64 0, ptr [[TMP38]], align 8657// CHECK2-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10658// CHECK2-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP39]], align 4659// CHECK2-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11660// CHECK2-NEXT: store [3 x i32] [[TMP28]], ptr [[TMP40]], align 4661// CHECK2-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12662// CHECK2-NEXT: store i32 0, ptr [[TMP41]], align 4663// CHECK2-NEXT: [[TMP42:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3:[0-9]+]], i64 -1, i32 -1, i32 [[TMP23]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.region_id, ptr [[KERNEL_ARGS]])664// CHECK2-NEXT: [[TMP43:%.*]] = icmp ne i32 [[TMP42]], 0665// CHECK2-NEXT: br i1 [[TMP43]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]666// CHECK2: omp_offload.failed:667// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51(i64 [[TMP4]], ptr [[A]], i64 [[TMP6]], i64 [[TMP8]]) #[[ATTR2:[0-9]+]]668// CHECK2-NEXT: br label [[OMP_OFFLOAD_CONT]]669// CHECK2: omp_offload.cont:670// CHECK2-NEXT: [[TMP44:%.*]] = load i32, ptr [[N]], align 4671// CHECK2-NEXT: store i32 [[TMP44]], ptr [[N_CASTED7]], align 4672// CHECK2-NEXT: [[TMP45:%.*]] = load i64, ptr [[N_CASTED7]], align 8673// CHECK2-NEXT: [[TMP46:%.*]] = load ptr, ptr [[G_ADDR]], align 8674// CHECK2-NEXT: [[TMP47:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 0675// CHECK2-NEXT: store i64 [[TMP45]], ptr [[TMP47]], align 8676// CHECK2-NEXT: [[TMP48:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 0677// CHECK2-NEXT: store i64 [[TMP45]], ptr [[TMP48]], align 8678// CHECK2-NEXT: [[TMP49:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS10]], i64 0, i64 0679// CHECK2-NEXT: store ptr null, ptr [[TMP49]], align 8680// CHECK2-NEXT: [[TMP50:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 1681// CHECK2-NEXT: store ptr [[A]], ptr [[TMP50]], align 8682// CHECK2-NEXT: [[TMP51:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 1683// CHECK2-NEXT: store ptr [[A]], ptr [[TMP51]], align 8684// CHECK2-NEXT: [[TMP52:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS10]], i64 0, i64 1685// CHECK2-NEXT: store ptr null, ptr [[TMP52]], align 8686// CHECK2-NEXT: [[TMP53:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 2687// CHECK2-NEXT: store ptr [[TMP46]], ptr [[TMP53]], align 8688// CHECK2-NEXT: [[TMP54:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 2689// CHECK2-NEXT: store ptr [[TMP46]], ptr [[TMP54]], align 8690// CHECK2-NEXT: [[TMP55:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS10]], i64 0, i64 2691// CHECK2-NEXT: store ptr null, ptr [[TMP55]], align 8692// CHECK2-NEXT: [[TMP56:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 0693// CHECK2-NEXT: [[TMP57:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 0694// CHECK2-NEXT: [[TMP58:%.*]] = load i32, ptr [[N]], align 4695// CHECK2-NEXT: store i32 [[TMP58]], ptr [[DOTCAPTURE_EXPR_12]], align 4696// CHECK2-NEXT: [[TMP59:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_12]], align 4697// CHECK2-NEXT: [[SUB14:%.*]] = sub nsw i32 [[TMP59]], 0698// CHECK2-NEXT: [[DIV15:%.*]] = sdiv i32 [[SUB14]], 1699// CHECK2-NEXT: [[SUB16:%.*]] = sub nsw i32 [[DIV15]], 1700// CHECK2-NEXT: store i32 [[SUB16]], ptr [[DOTCAPTURE_EXPR_13]], align 4701// CHECK2-NEXT: [[TMP60:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_13]], align 4702// CHECK2-NEXT: [[ADD17:%.*]] = add nsw i32 [[TMP60]], 1703// CHECK2-NEXT: [[TMP61:%.*]] = zext i32 [[ADD17]] to i64704// CHECK2-NEXT: [[TMP62:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 0705// CHECK2-NEXT: store i32 3, ptr [[TMP62]], align 4706// CHECK2-NEXT: [[TMP63:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 1707// CHECK2-NEXT: store i32 3, ptr [[TMP63]], align 4708// CHECK2-NEXT: [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 2709// CHECK2-NEXT: store ptr [[TMP56]], ptr [[TMP64]], align 8710// CHECK2-NEXT: [[TMP65:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 3711// CHECK2-NEXT: store ptr [[TMP57]], ptr [[TMP65]], align 8712// CHECK2-NEXT: [[TMP66:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 4713// CHECK2-NEXT: store ptr @.offload_sizes.1, ptr [[TMP66]], align 8714// CHECK2-NEXT: [[TMP67:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 5715// CHECK2-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP67]], align 8716// CHECK2-NEXT: [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 6717// CHECK2-NEXT: store ptr null, ptr [[TMP68]], align 8718// CHECK2-NEXT: [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 7719// CHECK2-NEXT: store ptr null, ptr [[TMP69]], align 8720// CHECK2-NEXT: [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 8721// CHECK2-NEXT: store i64 [[TMP61]], ptr [[TMP70]], align 8722// CHECK2-NEXT: [[TMP71:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 9723// CHECK2-NEXT: store i64 0, ptr [[TMP71]], align 8724// CHECK2-NEXT: [[TMP72:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 10725// CHECK2-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP72]], align 4726// CHECK2-NEXT: [[TMP73:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 11727// CHECK2-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP73]], align 4728// CHECK2-NEXT: [[TMP74:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 12729// CHECK2-NEXT: store i32 0, ptr [[TMP74]], align 4730// CHECK2-NEXT: [[TMP75:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.region_id, ptr [[KERNEL_ARGS18]])731// CHECK2-NEXT: [[TMP76:%.*]] = icmp ne i32 [[TMP75]], 0732// CHECK2-NEXT: br i1 [[TMP76]], label [[OMP_OFFLOAD_FAILED19:%.*]], label [[OMP_OFFLOAD_CONT20:%.*]]733// CHECK2: omp_offload.failed19:734// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57(i64 [[TMP45]], ptr [[A]], ptr [[TMP46]]) #[[ATTR2]]735// CHECK2-NEXT: br label [[OMP_OFFLOAD_CONT20]]736// CHECK2: omp_offload.cont20:737// CHECK2-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[A]], i64 0, i64 0738// CHECK2-NEXT: [[TMP77:%.*]] = load i32, ptr [[ARRAYIDX]], align 4739// CHECK2-NEXT: ret i32 [[TMP77]]740//741//742// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51743// CHECK2-SAME: (i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]], i64 noundef [[DOTCAPTURE_EXPR_1:%.*]]) #[[ATTR1:[0-9]+]] {744// CHECK2-NEXT: entry:745// CHECK2-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8746// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8747// CHECK2-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 8748// CHECK2-NEXT: [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca i64, align 8749// CHECK2-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8750// CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB3]])751// CHECK2-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8752// CHECK2-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8753// CHECK2-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 8754// CHECK2-NEXT: store i64 [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 8755// CHECK2-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 8756// CHECK2-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 4757// CHECK2-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR2]], align 4758// CHECK2-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB3]], i32 [[TMP0]], i32 [[TMP2]], i32 [[TMP3]])759// CHECK2-NEXT: [[TMP4:%.*]] = load i32, ptr [[N_ADDR]], align 4760// CHECK2-NEXT: store i32 [[TMP4]], ptr [[N_CASTED]], align 4761// CHECK2-NEXT: [[TMP5:%.*]] = load i64, ptr [[N_CASTED]], align 8762// CHECK2-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined, i64 [[TMP5]], ptr [[TMP1]])763// CHECK2-NEXT: ret void764//765//766// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined767// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR1]] {768// CHECK2-NEXT: entry:769// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8770// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8771// CHECK2-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8772// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8773// CHECK2-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4774// CHECK2-NEXT: [[TMP:%.*]] = alloca i32, align 4775// CHECK2-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4776// CHECK2-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4777// CHECK2-NEXT: [[I:%.*]] = alloca i32, align 4778// CHECK2-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4779// CHECK2-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4780// CHECK2-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4781// CHECK2-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4782// CHECK2-NEXT: [[I3:%.*]] = alloca i32, align 4783// CHECK2-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8784// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8785// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8786// CHECK2-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8787// CHECK2-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8788// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8789// CHECK2-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 4790// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 4791// CHECK2-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4792// CHECK2-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 0793// CHECK2-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1794// CHECK2-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1795// CHECK2-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4796// CHECK2-NEXT: store i32 0, ptr [[I]], align 4797// CHECK2-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4798// CHECK2-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]799// CHECK2-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]800// CHECK2: omp.precond.then:801// CHECK2-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4802// CHECK2-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4803// CHECK2-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_COMB_UB]], align 4804// CHECK2-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4805// CHECK2-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4806// CHECK2-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8807// CHECK2-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 4808// CHECK2-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP6]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)809// CHECK2-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4810// CHECK2-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4811// CHECK2-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP7]], [[TMP8]]812// CHECK2-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]813// CHECK2: cond.true:814// CHECK2-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4815// CHECK2-NEXT: br label [[COND_END:%.*]]816// CHECK2: cond.false:817// CHECK2-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4818// CHECK2-NEXT: br label [[COND_END]]819// CHECK2: cond.end:820// CHECK2-NEXT: [[COND:%.*]] = phi i32 [ [[TMP9]], [[COND_TRUE]] ], [ [[TMP10]], [[COND_FALSE]] ]821// CHECK2-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4822// CHECK2-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4823// CHECK2-NEXT: store i32 [[TMP11]], ptr [[DOTOMP_IV]], align 4824// CHECK2-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]825// CHECK2: omp.inner.for.cond:826// CHECK2-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4827// CHECK2-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4828// CHECK2-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP12]], [[TMP13]]829// CHECK2-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]830// CHECK2: omp.inner.for.body:831// CHECK2-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4832// CHECK2-NEXT: [[TMP15:%.*]] = zext i32 [[TMP14]] to i64833// CHECK2-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4834// CHECK2-NEXT: [[TMP17:%.*]] = zext i32 [[TMP16]] to i64835// CHECK2-NEXT: [[TMP18:%.*]] = load i32, ptr [[N_ADDR]], align 4836// CHECK2-NEXT: store i32 [[TMP18]], ptr [[N_CASTED]], align 4837// CHECK2-NEXT: [[TMP19:%.*]] = load i64, ptr [[N_CASTED]], align 8838// CHECK2-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 4, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined.omp_outlined, i64 [[TMP15]], i64 [[TMP17]], i64 [[TMP19]], ptr [[TMP0]])839// CHECK2-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]840// CHECK2: omp.inner.for.inc:841// CHECK2-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4842// CHECK2-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4843// CHECK2-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP20]], [[TMP21]]844// CHECK2-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4845// CHECK2-NEXT: br label [[OMP_INNER_FOR_COND]]846// CHECK2: omp.inner.for.end:847// CHECK2-NEXT: br label [[OMP_LOOP_EXIT:%.*]]848// CHECK2: omp.loop.exit:849// CHECK2-NEXT: [[TMP22:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8850// CHECK2-NEXT: [[TMP23:%.*]] = load i32, ptr [[TMP22]], align 4851// CHECK2-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP23]])852// CHECK2-NEXT: br label [[OMP_PRECOND_END]]853// CHECK2: omp.precond.end:854// CHECK2-NEXT: ret void855//856//857// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined.omp_outlined858// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR1]] {859// CHECK2-NEXT: entry:860// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8861// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8862// CHECK2-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8863// CHECK2-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8864// CHECK2-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8865// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8866// CHECK2-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4867// CHECK2-NEXT: [[TMP:%.*]] = alloca i32, align 4868// CHECK2-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4869// CHECK2-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4870// CHECK2-NEXT: [[I:%.*]] = alloca i32, align 4871// CHECK2-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4872// CHECK2-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4873// CHECK2-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4874// CHECK2-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4875// CHECK2-NEXT: [[I4:%.*]] = alloca i32, align 4876// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8877// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8878// CHECK2-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8879// CHECK2-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8880// CHECK2-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8881// CHECK2-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8882// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8883// CHECK2-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 4884// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 4885// CHECK2-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4886// CHECK2-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 0887// CHECK2-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1888// CHECK2-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1889// CHECK2-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4890// CHECK2-NEXT: store i32 0, ptr [[I]], align 4891// CHECK2-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4892// CHECK2-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]893// CHECK2-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]894// CHECK2: omp.precond.then:895// CHECK2-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4896// CHECK2-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4897// CHECK2-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_UB]], align 4898// CHECK2-NEXT: [[TMP5:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8899// CHECK2-NEXT: [[CONV:%.*]] = trunc i64 [[TMP5]] to i32900// CHECK2-NEXT: [[TMP6:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8901// CHECK2-NEXT: [[CONV3:%.*]] = trunc i64 [[TMP6]] to i32902// CHECK2-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 4903// CHECK2-NEXT: store i32 [[CONV3]], ptr [[DOTOMP_UB]], align 4904// CHECK2-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4905// CHECK2-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4906// CHECK2-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8907// CHECK2-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4908// CHECK2-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP8]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)909// CHECK2-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4910// CHECK2-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4911// CHECK2-NEXT: [[CMP5:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]912// CHECK2-NEXT: br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]913// CHECK2: cond.true:914// CHECK2-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4915// CHECK2-NEXT: br label [[COND_END:%.*]]916// CHECK2: cond.false:917// CHECK2-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4918// CHECK2-NEXT: br label [[COND_END]]919// CHECK2: cond.end:920// CHECK2-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]921// CHECK2-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4922// CHECK2-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4923// CHECK2-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 4924// CHECK2-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]925// CHECK2: omp.inner.for.cond:926// CHECK2-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4927// CHECK2-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4928// CHECK2-NEXT: [[CMP6:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]929// CHECK2-NEXT: br i1 [[CMP6]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]930// CHECK2: omp.inner.for.body:931// CHECK2-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4932// CHECK2-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 1933// CHECK2-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]934// CHECK2-NEXT: store i32 [[ADD]], ptr [[I4]], align 4935// CHECK2-NEXT: [[TMP17:%.*]] = load i32, ptr [[I4]], align 4936// CHECK2-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP17]] to i64937// CHECK2-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]938// CHECK2-NEXT: store i32 0, ptr [[ARRAYIDX]], align 4939// CHECK2-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]940// CHECK2: omp.body.continue:941// CHECK2-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]942// CHECK2: omp.inner.for.inc:943// CHECK2-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4944// CHECK2-NEXT: [[ADD7:%.*]] = add nsw i32 [[TMP18]], 1945// CHECK2-NEXT: store i32 [[ADD7]], ptr [[DOTOMP_IV]], align 4946// CHECK2-NEXT: br label [[OMP_INNER_FOR_COND]]947// CHECK2: omp.inner.for.end:948// CHECK2-NEXT: br label [[OMP_LOOP_EXIT:%.*]]949// CHECK2: omp.loop.exit:950// CHECK2-NEXT: [[TMP19:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8951// CHECK2-NEXT: [[TMP20:%.*]] = load i32, ptr [[TMP19]], align 4952// CHECK2-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP20]])953// CHECK2-NEXT: br label [[OMP_PRECOND_END]]954// CHECK2: omp.precond.end:955// CHECK2-NEXT: ret void956//957//958// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57959// CHECK2-SAME: (i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], ptr noundef [[G:%.*]]) #[[ATTR1]] {960// CHECK2-NEXT: entry:961// CHECK2-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8962// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8963// CHECK2-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 8964// CHECK2-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8965// CHECK2-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8966// CHECK2-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8967// CHECK2-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 8968// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8969// CHECK2-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 4970// CHECK2-NEXT: store i32 [[TMP1]], ptr [[N_CASTED]], align 4971// CHECK2-NEXT: [[TMP2:%.*]] = load i64, ptr [[N_CASTED]], align 8972// CHECK2-NEXT: [[TMP3:%.*]] = load ptr, ptr [[G_ADDR]], align 8973// CHECK2-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.omp_outlined, i64 [[TMP2]], ptr [[TMP0]], ptr [[TMP3]])974// CHECK2-NEXT: ret void975//976//977// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.omp_outlined978// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], ptr noundef [[G:%.*]]) #[[ATTR1]] {979// CHECK2-NEXT: entry:980// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8981// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8982// CHECK2-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8983// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8984// CHECK2-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 8985// CHECK2-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4986// CHECK2-NEXT: [[TMP:%.*]] = alloca i32, align 4987// CHECK2-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4988// CHECK2-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4989// CHECK2-NEXT: [[I:%.*]] = alloca i32, align 4990// CHECK2-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4991// CHECK2-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4992// CHECK2-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4993// CHECK2-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4994// CHECK2-NEXT: [[I3:%.*]] = alloca i32, align 4995// CHECK2-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8996// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8997// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8998// CHECK2-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8999// CHECK2-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81000// CHECK2-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 81001// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81002// CHECK2-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41003// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 41004// CHECK2-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41005// CHECK2-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 01006// CHECK2-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11007// CHECK2-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11008// CHECK2-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41009// CHECK2-NEXT: store i32 0, ptr [[I]], align 41010// CHECK2-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41011// CHECK2-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]1012// CHECK2-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1013// CHECK2: omp.precond.then:1014// CHECK2-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 41015// CHECK2-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41016// CHECK2-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_COMB_UB]], align 41017// CHECK2-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41018// CHECK2-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41019// CHECK2-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81020// CHECK2-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 41021// CHECK2-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP6]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1022// CHECK2-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41023// CHECK2-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41024// CHECK2-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP7]], [[TMP8]]1025// CHECK2-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1026// CHECK2: cond.true:1027// CHECK2-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41028// CHECK2-NEXT: br label [[COND_END:%.*]]1029// CHECK2: cond.false:1030// CHECK2-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41031// CHECK2-NEXT: br label [[COND_END]]1032// CHECK2: cond.end:1033// CHECK2-NEXT: [[COND:%.*]] = phi i32 [ [[TMP9]], [[COND_TRUE]] ], [ [[TMP10]], [[COND_FALSE]] ]1034// CHECK2-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41035// CHECK2-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41036// CHECK2-NEXT: store i32 [[TMP11]], ptr [[DOTOMP_IV]], align 41037// CHECK2-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1038// CHECK2: omp.inner.for.cond:1039// CHECK2-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41040// CHECK2-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41041// CHECK2-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP12]], [[TMP13]]1042// CHECK2-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1043// CHECK2: omp.inner.for.body:1044// CHECK2-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41045// CHECK2-NEXT: [[TMP15:%.*]] = zext i32 [[TMP14]] to i641046// CHECK2-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41047// CHECK2-NEXT: [[TMP17:%.*]] = zext i32 [[TMP16]] to i641048// CHECK2-NEXT: [[TMP18:%.*]] = load i32, ptr [[N_ADDR]], align 41049// CHECK2-NEXT: store i32 [[TMP18]], ptr [[N_CASTED]], align 41050// CHECK2-NEXT: [[TMP19:%.*]] = load i64, ptr [[N_CASTED]], align 81051// CHECK2-NEXT: [[TMP20:%.*]] = load ptr, ptr [[G_ADDR]], align 81052// CHECK2-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 5, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.omp_outlined.omp_outlined, i64 [[TMP15]], i64 [[TMP17]], i64 [[TMP19]], ptr [[TMP0]], ptr [[TMP20]])1053// CHECK2-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1054// CHECK2: omp.inner.for.inc:1055// CHECK2-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41056// CHECK2-NEXT: [[TMP22:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41057// CHECK2-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP21]], [[TMP22]]1058// CHECK2-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 41059// CHECK2-NEXT: br label [[OMP_INNER_FOR_COND]]1060// CHECK2: omp.inner.for.end:1061// CHECK2-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1062// CHECK2: omp.loop.exit:1063// CHECK2-NEXT: [[TMP23:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81064// CHECK2-NEXT: [[TMP24:%.*]] = load i32, ptr [[TMP23]], align 41065// CHECK2-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP24]])1066// CHECK2-NEXT: br label [[OMP_PRECOND_END]]1067// CHECK2: omp.precond.end:1068// CHECK2-NEXT: ret void1069//1070//1071// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.omp_outlined.omp_outlined1072// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], ptr noundef [[G:%.*]]) #[[ATTR1]] {1073// CHECK2-NEXT: entry:1074// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81075// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81076// CHECK2-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 81077// CHECK2-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 81078// CHECK2-NEXT: [[N_ADDR:%.*]] = alloca i64, align 81079// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81080// CHECK2-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 81081// CHECK2-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41082// CHECK2-NEXT: [[TMP:%.*]] = alloca i32, align 41083// CHECK2-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41084// CHECK2-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41085// CHECK2-NEXT: [[I:%.*]] = alloca i32, align 41086// CHECK2-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41087// CHECK2-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41088// CHECK2-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41089// CHECK2-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41090// CHECK2-NEXT: [[I4:%.*]] = alloca i32, align 41091// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81092// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81093// CHECK2-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 81094// CHECK2-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 81095// CHECK2-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 81096// CHECK2-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81097// CHECK2-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 81098// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81099// CHECK2-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41100// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 41101// CHECK2-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41102// CHECK2-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 01103// CHECK2-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11104// CHECK2-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11105// CHECK2-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41106// CHECK2-NEXT: store i32 0, ptr [[I]], align 41107// CHECK2-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41108// CHECK2-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]1109// CHECK2-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1110// CHECK2: omp.precond.then:1111// CHECK2-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41112// CHECK2-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41113// CHECK2-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_UB]], align 41114// CHECK2-NEXT: [[TMP5:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 81115// CHECK2-NEXT: [[CONV:%.*]] = trunc i64 [[TMP5]] to i321116// CHECK2-NEXT: [[TMP6:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 81117// CHECK2-NEXT: [[CONV3:%.*]] = trunc i64 [[TMP6]] to i321118// CHECK2-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 41119// CHECK2-NEXT: store i32 [[CONV3]], ptr [[DOTOMP_UB]], align 41120// CHECK2-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41121// CHECK2-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41122// CHECK2-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81123// CHECK2-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41124// CHECK2-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2]], i32 [[TMP8]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1125// CHECK2-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41126// CHECK2-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41127// CHECK2-NEXT: [[CMP5:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]1128// CHECK2-NEXT: br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1129// CHECK2: cond.true:1130// CHECK2-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41131// CHECK2-NEXT: br label [[COND_END:%.*]]1132// CHECK2: cond.false:1133// CHECK2-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41134// CHECK2-NEXT: br label [[COND_END]]1135// CHECK2: cond.end:1136// CHECK2-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]1137// CHECK2-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41138// CHECK2-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41139// CHECK2-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 41140// CHECK2-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1141// CHECK2: omp.inner.for.cond:1142// CHECK2-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41143// CHECK2-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41144// CHECK2-NEXT: [[CMP6:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]1145// CHECK2-NEXT: br i1 [[CMP6]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1146// CHECK2: omp.inner.for.body:1147// CHECK2-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41148// CHECK2-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 11149// CHECK2-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1150// CHECK2-NEXT: store i32 [[ADD]], ptr [[I4]], align 41151// CHECK2-NEXT: [[TMP17:%.*]] = load ptr, ptr [[G_ADDR]], align 81152// CHECK2-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP17]], i64 01153// CHECK2-NEXT: [[TMP18:%.*]] = load i32, ptr [[ARRAYIDX]], align 41154// CHECK2-NEXT: [[TMP19:%.*]] = load i32, ptr [[I4]], align 41155// CHECK2-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP19]] to i641156// CHECK2-NEXT: [[ARRAYIDX7:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]1157// CHECK2-NEXT: store i32 [[TMP18]], ptr [[ARRAYIDX7]], align 41158// CHECK2-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1159// CHECK2: omp.body.continue:1160// CHECK2-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1161// CHECK2: omp.inner.for.inc:1162// CHECK2-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41163// CHECK2-NEXT: [[ADD8:%.*]] = add nsw i32 [[TMP20]], 11164// CHECK2-NEXT: store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 41165// CHECK2-NEXT: br label [[OMP_INNER_FOR_COND]]1166// CHECK2: omp.inner.for.end:1167// CHECK2-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1168// CHECK2: omp.loop.exit:1169// CHECK2-NEXT: [[TMP21:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81170// CHECK2-NEXT: [[TMP22:%.*]] = load i32, ptr [[TMP21]], align 41171// CHECK2-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP22]])1172// CHECK2-NEXT: br label [[OMP_PRECOND_END]]1173// CHECK2: omp.precond.end:1174// CHECK2-NEXT: ret void1175//1176//1177// CHECK4-LABEL: define {{[^@]+}}@_Z16target_teams_funPi1178// CHECK4-SAME: (ptr noundef [[G:%.*]]) #[[ATTR0:[0-9]+]] {1179// CHECK4-NEXT: entry:1180// CHECK4-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 41181// CHECK4-NEXT: [[N:%.*]] = alloca i32, align 41182// CHECK4-NEXT: [[A:%.*]] = alloca [1000 x i32], align 41183// CHECK4-NEXT: [[TE:%.*]] = alloca i32, align 41184// CHECK4-NEXT: [[TH:%.*]] = alloca i32, align 41185// CHECK4-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41186// CHECK4-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41187// CHECK4-NEXT: [[N_CASTED:%.*]] = alloca i32, align 41188// CHECK4-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i32, align 41189// CHECK4-NEXT: [[DOTCAPTURE_EXPR__CASTED2:%.*]] = alloca i32, align 41190// CHECK4-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [4 x ptr], align 41191// CHECK4-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [4 x ptr], align 41192// CHECK4-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [4 x ptr], align 41193// CHECK4-NEXT: [[TMP:%.*]] = alloca i32, align 41194// CHECK4-NEXT: [[DOTCAPTURE_EXPR_3:%.*]] = alloca i32, align 41195// CHECK4-NEXT: [[DOTCAPTURE_EXPR_4:%.*]] = alloca i32, align 41196// CHECK4-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81197// CHECK4-NEXT: [[N_CASTED7:%.*]] = alloca i32, align 41198// CHECK4-NEXT: [[DOTOFFLOAD_BASEPTRS8:%.*]] = alloca [3 x ptr], align 41199// CHECK4-NEXT: [[DOTOFFLOAD_PTRS9:%.*]] = alloca [3 x ptr], align 41200// CHECK4-NEXT: [[DOTOFFLOAD_MAPPERS10:%.*]] = alloca [3 x ptr], align 41201// CHECK4-NEXT: [[_TMP11:%.*]] = alloca i32, align 41202// CHECK4-NEXT: [[DOTCAPTURE_EXPR_12:%.*]] = alloca i32, align 41203// CHECK4-NEXT: [[DOTCAPTURE_EXPR_13:%.*]] = alloca i32, align 41204// CHECK4-NEXT: [[KERNEL_ARGS18:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 81205// CHECK4-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 41206// CHECK4-NEXT: store i32 1000, ptr [[N]], align 41207// CHECK4-NEXT: [[TMP0:%.*]] = load i32, ptr [[N]], align 41208// CHECK4-NEXT: [[DIV:%.*]] = sdiv i32 [[TMP0]], 1281209// CHECK4-NEXT: store i32 [[DIV]], ptr [[TE]], align 41210// CHECK4-NEXT: store i32 128, ptr [[TH]], align 41211// CHECK4-NEXT: [[TMP1:%.*]] = load i32, ptr [[TE]], align 41212// CHECK4-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 41213// CHECK4-NEXT: [[TMP2:%.*]] = load i32, ptr [[TH]], align 41214// CHECK4-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_1]], align 41215// CHECK4-NEXT: [[TMP3:%.*]] = load i32, ptr [[N]], align 41216// CHECK4-NEXT: store i32 [[TMP3]], ptr [[N_CASTED]], align 41217// CHECK4-NEXT: [[TMP4:%.*]] = load i32, ptr [[N_CASTED]], align 41218// CHECK4-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41219// CHECK4-NEXT: store i32 [[TMP5]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 41220// CHECK4-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED]], align 41221// CHECK4-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41222// CHECK4-NEXT: store i32 [[TMP7]], ptr [[DOTCAPTURE_EXPR__CASTED2]], align 41223// CHECK4-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED2]], align 41224// CHECK4-NEXT: [[TMP9:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01225// CHECK4-NEXT: store i32 [[TMP4]], ptr [[TMP9]], align 41226// CHECK4-NEXT: [[TMP10:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01227// CHECK4-NEXT: store i32 [[TMP4]], ptr [[TMP10]], align 41228// CHECK4-NEXT: [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01229// CHECK4-NEXT: store ptr null, ptr [[TMP11]], align 41230// CHECK4-NEXT: [[TMP12:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11231// CHECK4-NEXT: store ptr [[A]], ptr [[TMP12]], align 41232// CHECK4-NEXT: [[TMP13:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11233// CHECK4-NEXT: store ptr [[A]], ptr [[TMP13]], align 41234// CHECK4-NEXT: [[TMP14:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 11235// CHECK4-NEXT: store ptr null, ptr [[TMP14]], align 41236// CHECK4-NEXT: [[TMP15:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21237// CHECK4-NEXT: store i32 [[TMP6]], ptr [[TMP15]], align 41238// CHECK4-NEXT: [[TMP16:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21239// CHECK4-NEXT: store i32 [[TMP6]], ptr [[TMP16]], align 41240// CHECK4-NEXT: [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 21241// CHECK4-NEXT: store ptr null, ptr [[TMP17]], align 41242// CHECK4-NEXT: [[TMP18:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 31243// CHECK4-NEXT: store i32 [[TMP8]], ptr [[TMP18]], align 41244// CHECK4-NEXT: [[TMP19:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 31245// CHECK4-NEXT: store i32 [[TMP8]], ptr [[TMP19]], align 41246// CHECK4-NEXT: [[TMP20:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 31247// CHECK4-NEXT: store ptr null, ptr [[TMP20]], align 41248// CHECK4-NEXT: [[TMP21:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01249// CHECK4-NEXT: [[TMP22:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01250// CHECK4-NEXT: [[TMP23:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41251// CHECK4-NEXT: [[TMP24:%.*]] = load i32, ptr [[N]], align 41252// CHECK4-NEXT: store i32 [[TMP24]], ptr [[DOTCAPTURE_EXPR_3]], align 41253// CHECK4-NEXT: [[TMP25:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_3]], align 41254// CHECK4-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP25]], 01255// CHECK4-NEXT: [[DIV5:%.*]] = sdiv i32 [[SUB]], 11256// CHECK4-NEXT: [[SUB6:%.*]] = sub nsw i32 [[DIV5]], 11257// CHECK4-NEXT: store i32 [[SUB6]], ptr [[DOTCAPTURE_EXPR_4]], align 41258// CHECK4-NEXT: [[TMP26:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 41259// CHECK4-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP26]], 11260// CHECK4-NEXT: [[TMP27:%.*]] = zext i32 [[ADD]] to i641261// CHECK4-NEXT: [[TMP28:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP23]], 01262// CHECK4-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01263// CHECK4-NEXT: store i32 3, ptr [[TMP29]], align 41264// CHECK4-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11265// CHECK4-NEXT: store i32 4, ptr [[TMP30]], align 41266// CHECK4-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21267// CHECK4-NEXT: store ptr [[TMP21]], ptr [[TMP31]], align 41268// CHECK4-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31269// CHECK4-NEXT: store ptr [[TMP22]], ptr [[TMP32]], align 41270// CHECK4-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41271// CHECK4-NEXT: store ptr @.offload_sizes, ptr [[TMP33]], align 41272// CHECK4-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51273// CHECK4-NEXT: store ptr @.offload_maptypes, ptr [[TMP34]], align 41274// CHECK4-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61275// CHECK4-NEXT: store ptr null, ptr [[TMP35]], align 41276// CHECK4-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71277// CHECK4-NEXT: store ptr null, ptr [[TMP36]], align 41278// CHECK4-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81279// CHECK4-NEXT: store i64 [[TMP27]], ptr [[TMP37]], align 81280// CHECK4-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91281// CHECK4-NEXT: store i64 0, ptr [[TMP38]], align 81282// CHECK4-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101283// CHECK4-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP39]], align 41284// CHECK4-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111285// CHECK4-NEXT: store [3 x i32] [[TMP28]], ptr [[TMP40]], align 41286// CHECK4-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121287// CHECK4-NEXT: store i32 0, ptr [[TMP41]], align 41288// CHECK4-NEXT: [[TMP42:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3:[0-9]+]], i64 -1, i32 -1, i32 [[TMP23]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.region_id, ptr [[KERNEL_ARGS]])1289// CHECK4-NEXT: [[TMP43:%.*]] = icmp ne i32 [[TMP42]], 01290// CHECK4-NEXT: br i1 [[TMP43]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1291// CHECK4: omp_offload.failed:1292// CHECK4-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51(i32 [[TMP4]], ptr [[A]], i32 [[TMP6]], i32 [[TMP8]]) #[[ATTR2:[0-9]+]]1293// CHECK4-NEXT: br label [[OMP_OFFLOAD_CONT]]1294// CHECK4: omp_offload.cont:1295// CHECK4-NEXT: [[TMP44:%.*]] = load i32, ptr [[N]], align 41296// CHECK4-NEXT: store i32 [[TMP44]], ptr [[N_CASTED7]], align 41297// CHECK4-NEXT: [[TMP45:%.*]] = load i32, ptr [[N_CASTED7]], align 41298// CHECK4-NEXT: [[TMP46:%.*]] = load ptr, ptr [[G_ADDR]], align 41299// CHECK4-NEXT: [[TMP47:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 01300// CHECK4-NEXT: store i32 [[TMP45]], ptr [[TMP47]], align 41301// CHECK4-NEXT: [[TMP48:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 01302// CHECK4-NEXT: store i32 [[TMP45]], ptr [[TMP48]], align 41303// CHECK4-NEXT: [[TMP49:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS10]], i32 0, i32 01304// CHECK4-NEXT: store ptr null, ptr [[TMP49]], align 41305// CHECK4-NEXT: [[TMP50:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 11306// CHECK4-NEXT: store ptr [[A]], ptr [[TMP50]], align 41307// CHECK4-NEXT: [[TMP51:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 11308// CHECK4-NEXT: store ptr [[A]], ptr [[TMP51]], align 41309// CHECK4-NEXT: [[TMP52:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS10]], i32 0, i32 11310// CHECK4-NEXT: store ptr null, ptr [[TMP52]], align 41311// CHECK4-NEXT: [[TMP53:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 21312// CHECK4-NEXT: store ptr [[TMP46]], ptr [[TMP53]], align 41313// CHECK4-NEXT: [[TMP54:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 21314// CHECK4-NEXT: store ptr [[TMP46]], ptr [[TMP54]], align 41315// CHECK4-NEXT: [[TMP55:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS10]], i32 0, i32 21316// CHECK4-NEXT: store ptr null, ptr [[TMP55]], align 41317// CHECK4-NEXT: [[TMP56:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 01318// CHECK4-NEXT: [[TMP57:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 01319// CHECK4-NEXT: [[TMP58:%.*]] = load i32, ptr [[N]], align 41320// CHECK4-NEXT: store i32 [[TMP58]], ptr [[DOTCAPTURE_EXPR_12]], align 41321// CHECK4-NEXT: [[TMP59:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_12]], align 41322// CHECK4-NEXT: [[SUB14:%.*]] = sub nsw i32 [[TMP59]], 01323// CHECK4-NEXT: [[DIV15:%.*]] = sdiv i32 [[SUB14]], 11324// CHECK4-NEXT: [[SUB16:%.*]] = sub nsw i32 [[DIV15]], 11325// CHECK4-NEXT: store i32 [[SUB16]], ptr [[DOTCAPTURE_EXPR_13]], align 41326// CHECK4-NEXT: [[TMP60:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_13]], align 41327// CHECK4-NEXT: [[ADD17:%.*]] = add nsw i32 [[TMP60]], 11328// CHECK4-NEXT: [[TMP61:%.*]] = zext i32 [[ADD17]] to i641329// CHECK4-NEXT: [[TMP62:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 01330// CHECK4-NEXT: store i32 3, ptr [[TMP62]], align 41331// CHECK4-NEXT: [[TMP63:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 11332// CHECK4-NEXT: store i32 3, ptr [[TMP63]], align 41333// CHECK4-NEXT: [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 21334// CHECK4-NEXT: store ptr [[TMP56]], ptr [[TMP64]], align 41335// CHECK4-NEXT: [[TMP65:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 31336// CHECK4-NEXT: store ptr [[TMP57]], ptr [[TMP65]], align 41337// CHECK4-NEXT: [[TMP66:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 41338// CHECK4-NEXT: store ptr @.offload_sizes.1, ptr [[TMP66]], align 41339// CHECK4-NEXT: [[TMP67:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 51340// CHECK4-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP67]], align 41341// CHECK4-NEXT: [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 61342// CHECK4-NEXT: store ptr null, ptr [[TMP68]], align 41343// CHECK4-NEXT: [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 71344// CHECK4-NEXT: store ptr null, ptr [[TMP69]], align 41345// CHECK4-NEXT: [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 81346// CHECK4-NEXT: store i64 [[TMP61]], ptr [[TMP70]], align 81347// CHECK4-NEXT: [[TMP71:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 91348// CHECK4-NEXT: store i64 0, ptr [[TMP71]], align 81349// CHECK4-NEXT: [[TMP72:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 101350// CHECK4-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP72]], align 41351// CHECK4-NEXT: [[TMP73:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 111352// CHECK4-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP73]], align 41353// CHECK4-NEXT: [[TMP74:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 121354// CHECK4-NEXT: store i32 0, ptr [[TMP74]], align 41355// CHECK4-NEXT: [[TMP75:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.region_id, ptr [[KERNEL_ARGS18]])1356// CHECK4-NEXT: [[TMP76:%.*]] = icmp ne i32 [[TMP75]], 01357// CHECK4-NEXT: br i1 [[TMP76]], label [[OMP_OFFLOAD_FAILED19:%.*]], label [[OMP_OFFLOAD_CONT20:%.*]]1358// CHECK4: omp_offload.failed19:1359// CHECK4-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57(i32 [[TMP45]], ptr [[A]], ptr [[TMP46]]) #[[ATTR2]]1360// CHECK4-NEXT: br label [[OMP_OFFLOAD_CONT20]]1361// CHECK4: omp_offload.cont20:1362// CHECK4-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[A]], i32 0, i32 01363// CHECK4-NEXT: [[TMP77:%.*]] = load i32, ptr [[ARRAYIDX]], align 41364// CHECK4-NEXT: ret i32 [[TMP77]]1365//1366//1367// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l511368// CHECK4-SAME: (i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]], i32 noundef [[DOTCAPTURE_EXPR_1:%.*]]) #[[ATTR1:[0-9]+]] {1369// CHECK4-NEXT: entry:1370// CHECK4-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41371// CHECK4-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41372// CHECK4-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 41373// CHECK4-NEXT: [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca i32, align 41374// CHECK4-NEXT: [[N_CASTED:%.*]] = alloca i32, align 41375// CHECK4-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB3]])1376// CHECK4-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41377// CHECK4-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41378// CHECK4-NEXT: store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 41379// CHECK4-NEXT: store i32 [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 41380// CHECK4-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 41381// CHECK4-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 41382// CHECK4-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR2]], align 41383// CHECK4-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB3]], i32 [[TMP0]], i32 [[TMP2]], i32 [[TMP3]])1384// CHECK4-NEXT: [[TMP4:%.*]] = load i32, ptr [[N_ADDR]], align 41385// CHECK4-NEXT: store i32 [[TMP4]], ptr [[N_CASTED]], align 41386// CHECK4-NEXT: [[TMP5:%.*]] = load i32, ptr [[N_CASTED]], align 41387// CHECK4-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined, i32 [[TMP5]], ptr [[TMP1]])1388// CHECK4-NEXT: ret void1389//1390//1391// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined1392// CHECK4-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR1]] {1393// CHECK4-NEXT: entry:1394// CHECK4-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41395// CHECK4-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41396// CHECK4-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41397// CHECK4-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41398// CHECK4-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41399// CHECK4-NEXT: [[TMP:%.*]] = alloca i32, align 41400// CHECK4-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41401// CHECK4-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41402// CHECK4-NEXT: [[I:%.*]] = alloca i32, align 41403// CHECK4-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41404// CHECK4-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41405// CHECK4-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41406// CHECK4-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41407// CHECK4-NEXT: [[I3:%.*]] = alloca i32, align 41408// CHECK4-NEXT: [[N_CASTED:%.*]] = alloca i32, align 41409// CHECK4-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41410// CHECK4-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41411// CHECK4-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41412// CHECK4-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41413// CHECK4-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41414// CHECK4-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41415// CHECK4-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 41416// CHECK4-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41417// CHECK4-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 01418// CHECK4-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11419// CHECK4-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11420// CHECK4-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41421// CHECK4-NEXT: store i32 0, ptr [[I]], align 41422// CHECK4-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41423// CHECK4-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]1424// CHECK4-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1425// CHECK4: omp.precond.then:1426// CHECK4-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 41427// CHECK4-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41428// CHECK4-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_COMB_UB]], align 41429// CHECK4-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41430// CHECK4-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41431// CHECK4-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41432// CHECK4-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 41433// CHECK4-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP6]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1434// CHECK4-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41435// CHECK4-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41436// CHECK4-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP7]], [[TMP8]]1437// CHECK4-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1438// CHECK4: cond.true:1439// CHECK4-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41440// CHECK4-NEXT: br label [[COND_END:%.*]]1441// CHECK4: cond.false:1442// CHECK4-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41443// CHECK4-NEXT: br label [[COND_END]]1444// CHECK4: cond.end:1445// CHECK4-NEXT: [[COND:%.*]] = phi i32 [ [[TMP9]], [[COND_TRUE]] ], [ [[TMP10]], [[COND_FALSE]] ]1446// CHECK4-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41447// CHECK4-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41448// CHECK4-NEXT: store i32 [[TMP11]], ptr [[DOTOMP_IV]], align 41449// CHECK4-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1450// CHECK4: omp.inner.for.cond:1451// CHECK4-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41452// CHECK4-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41453// CHECK4-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP12]], [[TMP13]]1454// CHECK4-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1455// CHECK4: omp.inner.for.body:1456// CHECK4-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41457// CHECK4-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41458// CHECK4-NEXT: [[TMP16:%.*]] = load i32, ptr [[N_ADDR]], align 41459// CHECK4-NEXT: store i32 [[TMP16]], ptr [[N_CASTED]], align 41460// CHECK4-NEXT: [[TMP17:%.*]] = load i32, ptr [[N_CASTED]], align 41461// CHECK4-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 4, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined.omp_outlined, i32 [[TMP14]], i32 [[TMP15]], i32 [[TMP17]], ptr [[TMP0]])1462// CHECK4-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1463// CHECK4: omp.inner.for.inc:1464// CHECK4-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41465// CHECK4-NEXT: [[TMP19:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41466// CHECK4-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP18]], [[TMP19]]1467// CHECK4-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 41468// CHECK4-NEXT: br label [[OMP_INNER_FOR_COND]]1469// CHECK4: omp.inner.for.end:1470// CHECK4-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1471// CHECK4: omp.loop.exit:1472// CHECK4-NEXT: [[TMP20:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41473// CHECK4-NEXT: [[TMP21:%.*]] = load i32, ptr [[TMP20]], align 41474// CHECK4-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP21]])1475// CHECK4-NEXT: br label [[OMP_PRECOND_END]]1476// CHECK4: omp.precond.end:1477// CHECK4-NEXT: ret void1478//1479//1480// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined.omp_outlined1481// CHECK4-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[DOTPREVIOUS_LB_:%.*]], i32 noundef [[DOTPREVIOUS_UB_:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR1]] {1482// CHECK4-NEXT: entry:1483// CHECK4-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41484// CHECK4-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41485// CHECK4-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i32, align 41486// CHECK4-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i32, align 41487// CHECK4-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41488// CHECK4-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41489// CHECK4-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41490// CHECK4-NEXT: [[TMP:%.*]] = alloca i32, align 41491// CHECK4-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41492// CHECK4-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41493// CHECK4-NEXT: [[I:%.*]] = alloca i32, align 41494// CHECK4-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41495// CHECK4-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41496// CHECK4-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41497// CHECK4-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41498// CHECK4-NEXT: [[I3:%.*]] = alloca i32, align 41499// CHECK4-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41500// CHECK4-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41501// CHECK4-NEXT: store i32 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 41502// CHECK4-NEXT: store i32 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 41503// CHECK4-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41504// CHECK4-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41505// CHECK4-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41506// CHECK4-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41507// CHECK4-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 41508// CHECK4-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41509// CHECK4-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 01510// CHECK4-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11511// CHECK4-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11512// CHECK4-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41513// CHECK4-NEXT: store i32 0, ptr [[I]], align 41514// CHECK4-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41515// CHECK4-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]1516// CHECK4-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1517// CHECK4: omp.precond.then:1518// CHECK4-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41519// CHECK4-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41520// CHECK4-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_UB]], align 41521// CHECK4-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTPREVIOUS_LB__ADDR]], align 41522// CHECK4-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTPREVIOUS_UB__ADDR]], align 41523// CHECK4-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_LB]], align 41524// CHECK4-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_UB]], align 41525// CHECK4-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41526// CHECK4-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41527// CHECK4-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41528// CHECK4-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41529// CHECK4-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP8]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1530// CHECK4-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41531// CHECK4-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41532// CHECK4-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]1533// CHECK4-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1534// CHECK4: cond.true:1535// CHECK4-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41536// CHECK4-NEXT: br label [[COND_END:%.*]]1537// CHECK4: cond.false:1538// CHECK4-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41539// CHECK4-NEXT: br label [[COND_END]]1540// CHECK4: cond.end:1541// CHECK4-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]1542// CHECK4-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41543// CHECK4-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41544// CHECK4-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 41545// CHECK4-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1546// CHECK4: omp.inner.for.cond:1547// CHECK4-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41548// CHECK4-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41549// CHECK4-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]1550// CHECK4-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1551// CHECK4: omp.inner.for.body:1552// CHECK4-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41553// CHECK4-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 11554// CHECK4-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1555// CHECK4-NEXT: store i32 [[ADD]], ptr [[I3]], align 41556// CHECK4-NEXT: [[TMP17:%.*]] = load i32, ptr [[I3]], align 41557// CHECK4-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i32 0, i32 [[TMP17]]1558// CHECK4-NEXT: store i32 0, ptr [[ARRAYIDX]], align 41559// CHECK4-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1560// CHECK4: omp.body.continue:1561// CHECK4-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1562// CHECK4: omp.inner.for.inc:1563// CHECK4-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41564// CHECK4-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP18]], 11565// CHECK4-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 41566// CHECK4-NEXT: br label [[OMP_INNER_FOR_COND]]1567// CHECK4: omp.inner.for.end:1568// CHECK4-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1569// CHECK4: omp.loop.exit:1570// CHECK4-NEXT: [[TMP19:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41571// CHECK4-NEXT: [[TMP20:%.*]] = load i32, ptr [[TMP19]], align 41572// CHECK4-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP20]])1573// CHECK4-NEXT: br label [[OMP_PRECOND_END]]1574// CHECK4: omp.precond.end:1575// CHECK4-NEXT: ret void1576//1577//1578// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l571579// CHECK4-SAME: (i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], ptr noundef [[G:%.*]]) #[[ATTR1]] {1580// CHECK4-NEXT: entry:1581// CHECK4-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41582// CHECK4-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41583// CHECK4-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 41584// CHECK4-NEXT: [[N_CASTED:%.*]] = alloca i32, align 41585// CHECK4-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41586// CHECK4-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41587// CHECK4-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 41588// CHECK4-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41589// CHECK4-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41590// CHECK4-NEXT: store i32 [[TMP1]], ptr [[N_CASTED]], align 41591// CHECK4-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_CASTED]], align 41592// CHECK4-NEXT: [[TMP3:%.*]] = load ptr, ptr [[G_ADDR]], align 41593// CHECK4-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.omp_outlined, i32 [[TMP2]], ptr [[TMP0]], ptr [[TMP3]])1594// CHECK4-NEXT: ret void1595//1596//1597// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.omp_outlined1598// CHECK4-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], ptr noundef [[G:%.*]]) #[[ATTR1]] {1599// CHECK4-NEXT: entry:1600// CHECK4-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41601// CHECK4-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41602// CHECK4-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41603// CHECK4-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41604// CHECK4-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 41605// CHECK4-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41606// CHECK4-NEXT: [[TMP:%.*]] = alloca i32, align 41607// CHECK4-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41608// CHECK4-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41609// CHECK4-NEXT: [[I:%.*]] = alloca i32, align 41610// CHECK4-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41611// CHECK4-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41612// CHECK4-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41613// CHECK4-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41614// CHECK4-NEXT: [[I3:%.*]] = alloca i32, align 41615// CHECK4-NEXT: [[N_CASTED:%.*]] = alloca i32, align 41616// CHECK4-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41617// CHECK4-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41618// CHECK4-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41619// CHECK4-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41620// CHECK4-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 41621// CHECK4-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41622// CHECK4-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41623// CHECK4-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 41624// CHECK4-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41625// CHECK4-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 01626// CHECK4-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11627// CHECK4-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11628// CHECK4-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41629// CHECK4-NEXT: store i32 0, ptr [[I]], align 41630// CHECK4-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41631// CHECK4-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]1632// CHECK4-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1633// CHECK4: omp.precond.then:1634// CHECK4-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 41635// CHECK4-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41636// CHECK4-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_COMB_UB]], align 41637// CHECK4-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41638// CHECK4-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41639// CHECK4-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41640// CHECK4-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 41641// CHECK4-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP6]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1642// CHECK4-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41643// CHECK4-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41644// CHECK4-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP7]], [[TMP8]]1645// CHECK4-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1646// CHECK4: cond.true:1647// CHECK4-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41648// CHECK4-NEXT: br label [[COND_END:%.*]]1649// CHECK4: cond.false:1650// CHECK4-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41651// CHECK4-NEXT: br label [[COND_END]]1652// CHECK4: cond.end:1653// CHECK4-NEXT: [[COND:%.*]] = phi i32 [ [[TMP9]], [[COND_TRUE]] ], [ [[TMP10]], [[COND_FALSE]] ]1654// CHECK4-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41655// CHECK4-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41656// CHECK4-NEXT: store i32 [[TMP11]], ptr [[DOTOMP_IV]], align 41657// CHECK4-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1658// CHECK4: omp.inner.for.cond:1659// CHECK4-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41660// CHECK4-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41661// CHECK4-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP12]], [[TMP13]]1662// CHECK4-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1663// CHECK4: omp.inner.for.body:1664// CHECK4-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41665// CHECK4-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41666// CHECK4-NEXT: [[TMP16:%.*]] = load i32, ptr [[N_ADDR]], align 41667// CHECK4-NEXT: store i32 [[TMP16]], ptr [[N_CASTED]], align 41668// CHECK4-NEXT: [[TMP17:%.*]] = load i32, ptr [[N_CASTED]], align 41669// CHECK4-NEXT: [[TMP18:%.*]] = load ptr, ptr [[G_ADDR]], align 41670// CHECK4-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 5, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.omp_outlined.omp_outlined, i32 [[TMP14]], i32 [[TMP15]], i32 [[TMP17]], ptr [[TMP0]], ptr [[TMP18]])1671// CHECK4-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1672// CHECK4: omp.inner.for.inc:1673// CHECK4-NEXT: [[TMP19:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41674// CHECK4-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41675// CHECK4-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP19]], [[TMP20]]1676// CHECK4-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 41677// CHECK4-NEXT: br label [[OMP_INNER_FOR_COND]]1678// CHECK4: omp.inner.for.end:1679// CHECK4-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1680// CHECK4: omp.loop.exit:1681// CHECK4-NEXT: [[TMP21:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41682// CHECK4-NEXT: [[TMP22:%.*]] = load i32, ptr [[TMP21]], align 41683// CHECK4-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP22]])1684// CHECK4-NEXT: br label [[OMP_PRECOND_END]]1685// CHECK4: omp.precond.end:1686// CHECK4-NEXT: ret void1687//1688//1689// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.omp_outlined.omp_outlined1690// CHECK4-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[DOTPREVIOUS_LB_:%.*]], i32 noundef [[DOTPREVIOUS_UB_:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], ptr noundef [[G:%.*]]) #[[ATTR1]] {1691// CHECK4-NEXT: entry:1692// CHECK4-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41693// CHECK4-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41694// CHECK4-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i32, align 41695// CHECK4-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i32, align 41696// CHECK4-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41697// CHECK4-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41698// CHECK4-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 41699// CHECK4-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41700// CHECK4-NEXT: [[TMP:%.*]] = alloca i32, align 41701// CHECK4-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41702// CHECK4-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41703// CHECK4-NEXT: [[I:%.*]] = alloca i32, align 41704// CHECK4-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41705// CHECK4-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41706// CHECK4-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41707// CHECK4-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41708// CHECK4-NEXT: [[I3:%.*]] = alloca i32, align 41709// CHECK4-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41710// CHECK4-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41711// CHECK4-NEXT: store i32 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 41712// CHECK4-NEXT: store i32 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 41713// CHECK4-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41714// CHECK4-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41715// CHECK4-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 41716// CHECK4-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41717// CHECK4-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41718// CHECK4-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 41719// CHECK4-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41720// CHECK4-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 01721// CHECK4-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11722// CHECK4-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11723// CHECK4-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41724// CHECK4-NEXT: store i32 0, ptr [[I]], align 41725// CHECK4-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41726// CHECK4-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]1727// CHECK4-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1728// CHECK4: omp.precond.then:1729// CHECK4-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41730// CHECK4-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41731// CHECK4-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_UB]], align 41732// CHECK4-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTPREVIOUS_LB__ADDR]], align 41733// CHECK4-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTPREVIOUS_UB__ADDR]], align 41734// CHECK4-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_LB]], align 41735// CHECK4-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_UB]], align 41736// CHECK4-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41737// CHECK4-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41738// CHECK4-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41739// CHECK4-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41740// CHECK4-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2]], i32 [[TMP8]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1741// CHECK4-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41742// CHECK4-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41743// CHECK4-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]1744// CHECK4-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1745// CHECK4: cond.true:1746// CHECK4-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41747// CHECK4-NEXT: br label [[COND_END:%.*]]1748// CHECK4: cond.false:1749// CHECK4-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41750// CHECK4-NEXT: br label [[COND_END]]1751// CHECK4: cond.end:1752// CHECK4-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]1753// CHECK4-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41754// CHECK4-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41755// CHECK4-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 41756// CHECK4-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1757// CHECK4: omp.inner.for.cond:1758// CHECK4-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41759// CHECK4-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41760// CHECK4-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]1761// CHECK4-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1762// CHECK4: omp.inner.for.body:1763// CHECK4-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41764// CHECK4-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 11765// CHECK4-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1766// CHECK4-NEXT: store i32 [[ADD]], ptr [[I3]], align 41767// CHECK4-NEXT: [[TMP17:%.*]] = load ptr, ptr [[G_ADDR]], align 41768// CHECK4-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP17]], i32 01769// CHECK4-NEXT: [[TMP18:%.*]] = load i32, ptr [[ARRAYIDX]], align 41770// CHECK4-NEXT: [[TMP19:%.*]] = load i32, ptr [[I3]], align 41771// CHECK4-NEXT: [[ARRAYIDX6:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i32 0, i32 [[TMP19]]1772// CHECK4-NEXT: store i32 [[TMP18]], ptr [[ARRAYIDX6]], align 41773// CHECK4-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1774// CHECK4: omp.body.continue:1775// CHECK4-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1776// CHECK4: omp.inner.for.inc:1777// CHECK4-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41778// CHECK4-NEXT: [[ADD7:%.*]] = add nsw i32 [[TMP20]], 11779// CHECK4-NEXT: store i32 [[ADD7]], ptr [[DOTOMP_IV]], align 41780// CHECK4-NEXT: br label [[OMP_INNER_FOR_COND]]1781// CHECK4: omp.inner.for.end:1782// CHECK4-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1783// CHECK4: omp.loop.exit:1784// CHECK4-NEXT: [[TMP21:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41785// CHECK4-NEXT: [[TMP22:%.*]] = load i32, ptr [[TMP21]], align 41786// CHECK4-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP22]])1787// CHECK4-NEXT: br label [[OMP_PRECOND_END]]1788// CHECK4: omp.precond.end:1789// CHECK4-NEXT: ret void1790//1791//1792// CHECK10-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l511793// CHECK10-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]], i64 noundef [[DOTCAPTURE_EXPR_1:%.*]]) #[[ATTR0:[0-9]+]] {1794// CHECK10-NEXT: entry:1795// CHECK10-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81796// CHECK10-NEXT: [[N_ADDR:%.*]] = alloca i64, align 81797// CHECK10-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81798// CHECK10-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 81799// CHECK10-NEXT: [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca i64, align 81800// CHECK10-NEXT: [[N_CASTED:%.*]] = alloca i64, align 81801// CHECK10-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB3:[0-9]+]])1802// CHECK10-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81803// CHECK10-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 81804// CHECK10-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81805// CHECK10-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 81806// CHECK10-NEXT: store i64 [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 81807// CHECK10-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 81808// CHECK10-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 41809// CHECK10-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR2]], align 41810// CHECK10-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB3]], i32 [[TMP0]], i32 [[TMP2]], i32 [[TMP3]])1811// CHECK10-NEXT: [[TMP4:%.*]] = load i32, ptr [[N_ADDR]], align 41812// CHECK10-NEXT: store i32 [[TMP4]], ptr [[N_CASTED]], align 41813// CHECK10-NEXT: [[TMP5:%.*]] = load i64, ptr [[N_CASTED]], align 81814// CHECK10-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined, i64 [[TMP5]], ptr [[TMP1]])1815// CHECK10-NEXT: ret void1816//1817//1818// CHECK10-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined1819// CHECK10-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR0]] {1820// CHECK10-NEXT: entry:1821// CHECK10-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81822// CHECK10-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81823// CHECK10-NEXT: [[N_ADDR:%.*]] = alloca i64, align 81824// CHECK10-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81825// CHECK10-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41826// CHECK10-NEXT: [[TMP:%.*]] = alloca i32, align 41827// CHECK10-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41828// CHECK10-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41829// CHECK10-NEXT: [[I:%.*]] = alloca i32, align 41830// CHECK10-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41831// CHECK10-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41832// CHECK10-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41833// CHECK10-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41834// CHECK10-NEXT: [[I3:%.*]] = alloca i32, align 41835// CHECK10-NEXT: [[N_CASTED:%.*]] = alloca i64, align 81836// CHECK10-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81837// CHECK10-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81838// CHECK10-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 81839// CHECK10-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81840// CHECK10-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81841// CHECK10-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41842// CHECK10-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 41843// CHECK10-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41844// CHECK10-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 01845// CHECK10-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11846// CHECK10-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11847// CHECK10-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41848// CHECK10-NEXT: store i32 0, ptr [[I]], align 41849// CHECK10-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41850// CHECK10-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]1851// CHECK10-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1852// CHECK10: omp.precond.then:1853// CHECK10-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 41854// CHECK10-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41855// CHECK10-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_COMB_UB]], align 41856// CHECK10-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41857// CHECK10-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41858// CHECK10-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81859// CHECK10-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 41860// CHECK10-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP6]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1861// CHECK10-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41862// CHECK10-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41863// CHECK10-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP7]], [[TMP8]]1864// CHECK10-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1865// CHECK10: cond.true:1866// CHECK10-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41867// CHECK10-NEXT: br label [[COND_END:%.*]]1868// CHECK10: cond.false:1869// CHECK10-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41870// CHECK10-NEXT: br label [[COND_END]]1871// CHECK10: cond.end:1872// CHECK10-NEXT: [[COND:%.*]] = phi i32 [ [[TMP9]], [[COND_TRUE]] ], [ [[TMP10]], [[COND_FALSE]] ]1873// CHECK10-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41874// CHECK10-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41875// CHECK10-NEXT: store i32 [[TMP11]], ptr [[DOTOMP_IV]], align 41876// CHECK10-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1877// CHECK10: omp.inner.for.cond:1878// CHECK10-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41879// CHECK10-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41880// CHECK10-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP12]], [[TMP13]]1881// CHECK10-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1882// CHECK10: omp.inner.for.body:1883// CHECK10-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41884// CHECK10-NEXT: [[TMP15:%.*]] = zext i32 [[TMP14]] to i641885// CHECK10-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41886// CHECK10-NEXT: [[TMP17:%.*]] = zext i32 [[TMP16]] to i641887// CHECK10-NEXT: [[TMP18:%.*]] = load i32, ptr [[N_ADDR]], align 41888// CHECK10-NEXT: store i32 [[TMP18]], ptr [[N_CASTED]], align 41889// CHECK10-NEXT: [[TMP19:%.*]] = load i64, ptr [[N_CASTED]], align 81890// CHECK10-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 4, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined.omp_outlined, i64 [[TMP15]], i64 [[TMP17]], i64 [[TMP19]], ptr [[TMP0]])1891// CHECK10-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1892// CHECK10: omp.inner.for.inc:1893// CHECK10-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41894// CHECK10-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41895// CHECK10-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP20]], [[TMP21]]1896// CHECK10-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 41897// CHECK10-NEXT: br label [[OMP_INNER_FOR_COND]]1898// CHECK10: omp.inner.for.end:1899// CHECK10-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1900// CHECK10: omp.loop.exit:1901// CHECK10-NEXT: [[TMP22:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81902// CHECK10-NEXT: [[TMP23:%.*]] = load i32, ptr [[TMP22]], align 41903// CHECK10-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP23]])1904// CHECK10-NEXT: br label [[OMP_PRECOND_END]]1905// CHECK10: omp.precond.end:1906// CHECK10-NEXT: ret void1907//1908//1909// CHECK10-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined.omp_outlined1910// CHECK10-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR0]] {1911// CHECK10-NEXT: entry:1912// CHECK10-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81913// CHECK10-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81914// CHECK10-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 81915// CHECK10-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 81916// CHECK10-NEXT: [[N_ADDR:%.*]] = alloca i64, align 81917// CHECK10-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81918// CHECK10-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41919// CHECK10-NEXT: [[TMP:%.*]] = alloca i32, align 41920// CHECK10-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41921// CHECK10-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41922// CHECK10-NEXT: [[I:%.*]] = alloca i32, align 41923// CHECK10-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41924// CHECK10-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41925// CHECK10-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41926// CHECK10-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41927// CHECK10-NEXT: [[I4:%.*]] = alloca i32, align 41928// CHECK10-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81929// CHECK10-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81930// CHECK10-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 81931// CHECK10-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 81932// CHECK10-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 81933// CHECK10-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81934// CHECK10-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81935// CHECK10-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41936// CHECK10-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 41937// CHECK10-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41938// CHECK10-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 01939// CHECK10-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11940// CHECK10-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11941// CHECK10-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41942// CHECK10-NEXT: store i32 0, ptr [[I]], align 41943// CHECK10-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41944// CHECK10-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]1945// CHECK10-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1946// CHECK10: omp.precond.then:1947// CHECK10-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41948// CHECK10-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41949// CHECK10-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_UB]], align 41950// CHECK10-NEXT: [[TMP5:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 81951// CHECK10-NEXT: [[CONV:%.*]] = trunc i64 [[TMP5]] to i321952// CHECK10-NEXT: [[TMP6:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 81953// CHECK10-NEXT: [[CONV3:%.*]] = trunc i64 [[TMP6]] to i321954// CHECK10-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 41955// CHECK10-NEXT: store i32 [[CONV3]], ptr [[DOTOMP_UB]], align 41956// CHECK10-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41957// CHECK10-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41958// CHECK10-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81959// CHECK10-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41960// CHECK10-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP8]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1961// CHECK10-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41962// CHECK10-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41963// CHECK10-NEXT: [[CMP5:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]1964// CHECK10-NEXT: br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1965// CHECK10: cond.true:1966// CHECK10-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41967// CHECK10-NEXT: br label [[COND_END:%.*]]1968// CHECK10: cond.false:1969// CHECK10-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41970// CHECK10-NEXT: br label [[COND_END]]1971// CHECK10: cond.end:1972// CHECK10-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]1973// CHECK10-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41974// CHECK10-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41975// CHECK10-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 41976// CHECK10-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1977// CHECK10: omp.inner.for.cond:1978// CHECK10-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41979// CHECK10-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41980// CHECK10-NEXT: [[CMP6:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]1981// CHECK10-NEXT: br i1 [[CMP6]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1982// CHECK10: omp.inner.for.body:1983// CHECK10-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41984// CHECK10-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 11985// CHECK10-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1986// CHECK10-NEXT: store i32 [[ADD]], ptr [[I4]], align 41987// CHECK10-NEXT: [[TMP17:%.*]] = load i32, ptr [[I4]], align 41988// CHECK10-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP17]] to i641989// CHECK10-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]1990// CHECK10-NEXT: store i32 0, ptr [[ARRAYIDX]], align 41991// CHECK10-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1992// CHECK10: omp.body.continue:1993// CHECK10-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1994// CHECK10: omp.inner.for.inc:1995// CHECK10-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41996// CHECK10-NEXT: [[ADD7:%.*]] = add nsw i32 [[TMP18]], 11997// CHECK10-NEXT: store i32 [[ADD7]], ptr [[DOTOMP_IV]], align 41998// CHECK10-NEXT: br label [[OMP_INNER_FOR_COND]]1999// CHECK10: omp.inner.for.end:2000// CHECK10-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2001// CHECK10: omp.loop.exit:2002// CHECK10-NEXT: [[TMP19:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82003// CHECK10-NEXT: [[TMP20:%.*]] = load i32, ptr [[TMP19]], align 42004// CHECK10-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP20]])2005// CHECK10-NEXT: br label [[OMP_PRECOND_END]]2006// CHECK10: omp.precond.end:2007// CHECK10-NEXT: ret void2008//2009//2010// CHECK10-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l572011// CHECK10-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], ptr noundef [[G:%.*]]) #[[ATTR0]] {2012// CHECK10-NEXT: entry:2013// CHECK10-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 82014// CHECK10-NEXT: [[N_ADDR:%.*]] = alloca i64, align 82015// CHECK10-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 82016// CHECK10-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 82017// CHECK10-NEXT: [[N_CASTED:%.*]] = alloca i64, align 82018// CHECK10-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 82019// CHECK10-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 82020// CHECK10-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 82021// CHECK10-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 82022// CHECK10-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 82023// CHECK10-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 42024// CHECK10-NEXT: store i32 [[TMP1]], ptr [[N_CASTED]], align 42025// CHECK10-NEXT: [[TMP2:%.*]] = load i64, ptr [[N_CASTED]], align 82026// CHECK10-NEXT: [[TMP3:%.*]] = load ptr, ptr [[G_ADDR]], align 82027// CHECK10-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.omp_outlined, i64 [[TMP2]], ptr [[TMP0]], ptr [[TMP3]])2028// CHECK10-NEXT: ret void2029//2030//2031// CHECK10-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.omp_outlined2032// CHECK10-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], ptr noundef [[G:%.*]]) #[[ATTR0]] {2033// CHECK10-NEXT: entry:2034// CHECK10-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82035// CHECK10-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82036// CHECK10-NEXT: [[N_ADDR:%.*]] = alloca i64, align 82037// CHECK10-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 82038// CHECK10-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 82039// CHECK10-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42040// CHECK10-NEXT: [[TMP:%.*]] = alloca i32, align 42041// CHECK10-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42042// CHECK10-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42043// CHECK10-NEXT: [[I:%.*]] = alloca i32, align 42044// CHECK10-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 42045// CHECK10-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 42046// CHECK10-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42047// CHECK10-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42048// CHECK10-NEXT: [[I3:%.*]] = alloca i32, align 42049// CHECK10-NEXT: [[N_CASTED:%.*]] = alloca i64, align 82050// CHECK10-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82051// CHECK10-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82052// CHECK10-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 82053// CHECK10-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 82054// CHECK10-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 82055// CHECK10-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 82056// CHECK10-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 42057// CHECK10-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 42058// CHECK10-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42059// CHECK10-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 02060// CHECK10-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12061// CHECK10-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12062// CHECK10-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42063// CHECK10-NEXT: store i32 0, ptr [[I]], align 42064// CHECK10-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42065// CHECK10-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]2066// CHECK10-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2067// CHECK10: omp.precond.then:2068// CHECK10-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 42069// CHECK10-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42070// CHECK10-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_COMB_UB]], align 42071// CHECK10-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42072// CHECK10-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42073// CHECK10-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82074// CHECK10-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 42075// CHECK10-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP6]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2076// CHECK10-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42077// CHECK10-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42078// CHECK10-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP7]], [[TMP8]]2079// CHECK10-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2080// CHECK10: cond.true:2081// CHECK10-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42082// CHECK10-NEXT: br label [[COND_END:%.*]]2083// CHECK10: cond.false:2084// CHECK10-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42085// CHECK10-NEXT: br label [[COND_END]]2086// CHECK10: cond.end:2087// CHECK10-NEXT: [[COND:%.*]] = phi i32 [ [[TMP9]], [[COND_TRUE]] ], [ [[TMP10]], [[COND_FALSE]] ]2088// CHECK10-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 42089// CHECK10-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 42090// CHECK10-NEXT: store i32 [[TMP11]], ptr [[DOTOMP_IV]], align 42091// CHECK10-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2092// CHECK10: omp.inner.for.cond:2093// CHECK10-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42094// CHECK10-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42095// CHECK10-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP12]], [[TMP13]]2096// CHECK10-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2097// CHECK10: omp.inner.for.body:2098// CHECK10-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 42099// CHECK10-NEXT: [[TMP15:%.*]] = zext i32 [[TMP14]] to i642100// CHECK10-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42101// CHECK10-NEXT: [[TMP17:%.*]] = zext i32 [[TMP16]] to i642102// CHECK10-NEXT: [[TMP18:%.*]] = load i32, ptr [[N_ADDR]], align 42103// CHECK10-NEXT: store i32 [[TMP18]], ptr [[N_CASTED]], align 42104// CHECK10-NEXT: [[TMP19:%.*]] = load i64, ptr [[N_CASTED]], align 82105// CHECK10-NEXT: [[TMP20:%.*]] = load ptr, ptr [[G_ADDR]], align 82106// CHECK10-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 5, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.omp_outlined.omp_outlined, i64 [[TMP15]], i64 [[TMP17]], i64 [[TMP19]], ptr [[TMP0]], ptr [[TMP20]])2107// CHECK10-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2108// CHECK10: omp.inner.for.inc:2109// CHECK10-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42110// CHECK10-NEXT: [[TMP22:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 42111// CHECK10-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP21]], [[TMP22]]2112// CHECK10-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 42113// CHECK10-NEXT: br label [[OMP_INNER_FOR_COND]]2114// CHECK10: omp.inner.for.end:2115// CHECK10-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2116// CHECK10: omp.loop.exit:2117// CHECK10-NEXT: [[TMP23:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82118// CHECK10-NEXT: [[TMP24:%.*]] = load i32, ptr [[TMP23]], align 42119// CHECK10-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP24]])2120// CHECK10-NEXT: br label [[OMP_PRECOND_END]]2121// CHECK10: omp.precond.end:2122// CHECK10-NEXT: ret void2123//2124//2125// CHECK10-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.omp_outlined.omp_outlined2126// CHECK10-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], ptr noundef [[G:%.*]]) #[[ATTR0]] {2127// CHECK10-NEXT: entry:2128// CHECK10-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82129// CHECK10-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82130// CHECK10-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 82131// CHECK10-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 82132// CHECK10-NEXT: [[N_ADDR:%.*]] = alloca i64, align 82133// CHECK10-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 82134// CHECK10-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 82135// CHECK10-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42136// CHECK10-NEXT: [[TMP:%.*]] = alloca i32, align 42137// CHECK10-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42138// CHECK10-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42139// CHECK10-NEXT: [[I:%.*]] = alloca i32, align 42140// CHECK10-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 42141// CHECK10-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 42142// CHECK10-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42143// CHECK10-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42144// CHECK10-NEXT: [[I4:%.*]] = alloca i32, align 42145// CHECK10-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82146// CHECK10-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82147// CHECK10-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 82148// CHECK10-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 82149// CHECK10-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 82150// CHECK10-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 82151// CHECK10-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 82152// CHECK10-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 82153// CHECK10-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 42154// CHECK10-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 42155// CHECK10-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42156// CHECK10-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 02157// CHECK10-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12158// CHECK10-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12159// CHECK10-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42160// CHECK10-NEXT: store i32 0, ptr [[I]], align 42161// CHECK10-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42162// CHECK10-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]2163// CHECK10-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2164// CHECK10: omp.precond.then:2165// CHECK10-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 42166// CHECK10-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42167// CHECK10-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_UB]], align 42168// CHECK10-NEXT: [[TMP5:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 82169// CHECK10-NEXT: [[CONV:%.*]] = trunc i64 [[TMP5]] to i322170// CHECK10-NEXT: [[TMP6:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 82171// CHECK10-NEXT: [[CONV3:%.*]] = trunc i64 [[TMP6]] to i322172// CHECK10-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 42173// CHECK10-NEXT: store i32 [[CONV3]], ptr [[DOTOMP_UB]], align 42174// CHECK10-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42175// CHECK10-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42176// CHECK10-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82177// CHECK10-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 42178// CHECK10-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2]], i32 [[TMP8]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2179// CHECK10-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42180// CHECK10-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42181// CHECK10-NEXT: [[CMP5:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]2182// CHECK10-NEXT: br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2183// CHECK10: cond.true:2184// CHECK10-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42185// CHECK10-NEXT: br label [[COND_END:%.*]]2186// CHECK10: cond.false:2187// CHECK10-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42188// CHECK10-NEXT: br label [[COND_END]]2189// CHECK10: cond.end:2190// CHECK10-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]2191// CHECK10-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 42192// CHECK10-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42193// CHECK10-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 42194// CHECK10-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2195// CHECK10: omp.inner.for.cond:2196// CHECK10-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42197// CHECK10-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42198// CHECK10-NEXT: [[CMP6:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]2199// CHECK10-NEXT: br i1 [[CMP6]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2200// CHECK10: omp.inner.for.body:2201// CHECK10-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42202// CHECK10-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 12203// CHECK10-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]2204// CHECK10-NEXT: store i32 [[ADD]], ptr [[I4]], align 42205// CHECK10-NEXT: [[TMP17:%.*]] = load ptr, ptr [[G_ADDR]], align 82206// CHECK10-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP17]], i64 02207// CHECK10-NEXT: [[TMP18:%.*]] = load i32, ptr [[ARRAYIDX]], align 42208// CHECK10-NEXT: [[TMP19:%.*]] = load i32, ptr [[I4]], align 42209// CHECK10-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP19]] to i642210// CHECK10-NEXT: [[ARRAYIDX7:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]2211// CHECK10-NEXT: store i32 [[TMP18]], ptr [[ARRAYIDX7]], align 42212// CHECK10-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]2213// CHECK10: omp.body.continue:2214// CHECK10-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2215// CHECK10: omp.inner.for.inc:2216// CHECK10-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42217// CHECK10-NEXT: [[ADD8:%.*]] = add nsw i32 [[TMP20]], 12218// CHECK10-NEXT: store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 42219// CHECK10-NEXT: br label [[OMP_INNER_FOR_COND]]2220// CHECK10: omp.inner.for.end:2221// CHECK10-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2222// CHECK10: omp.loop.exit:2223// CHECK10-NEXT: [[TMP21:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82224// CHECK10-NEXT: [[TMP22:%.*]] = load i32, ptr [[TMP21]], align 42225// CHECK10-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP22]])2226// CHECK10-NEXT: br label [[OMP_PRECOND_END]]2227// CHECK10: omp.precond.end:2228// CHECK10-NEXT: ret void2229//2230//2231// CHECK12-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l512232// CHECK12-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]], i32 noundef [[DOTCAPTURE_EXPR_1:%.*]]) #[[ATTR0:[0-9]+]] {2233// CHECK12-NEXT: entry:2234// CHECK12-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42235// CHECK12-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42236// CHECK12-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42237// CHECK12-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 42238// CHECK12-NEXT: [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca i32, align 42239// CHECK12-NEXT: [[N_CASTED:%.*]] = alloca i32, align 42240// CHECK12-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB3:[0-9]+]])2241// CHECK12-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42242// CHECK12-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42243// CHECK12-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42244// CHECK12-NEXT: store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 42245// CHECK12-NEXT: store i32 [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 42246// CHECK12-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 42247// CHECK12-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 42248// CHECK12-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR2]], align 42249// CHECK12-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB3]], i32 [[TMP0]], i32 [[TMP2]], i32 [[TMP3]])2250// CHECK12-NEXT: [[TMP4:%.*]] = load i32, ptr [[N_ADDR]], align 42251// CHECK12-NEXT: store i32 [[TMP4]], ptr [[N_CASTED]], align 42252// CHECK12-NEXT: [[TMP5:%.*]] = load i32, ptr [[N_CASTED]], align 42253// CHECK12-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined, i32 [[TMP5]], ptr [[TMP1]])2254// CHECK12-NEXT: ret void2255//2256//2257// CHECK12-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined2258// CHECK12-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR0]] {2259// CHECK12-NEXT: entry:2260// CHECK12-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42261// CHECK12-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42262// CHECK12-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42263// CHECK12-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42264// CHECK12-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42265// CHECK12-NEXT: [[TMP:%.*]] = alloca i32, align 42266// CHECK12-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42267// CHECK12-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42268// CHECK12-NEXT: [[I:%.*]] = alloca i32, align 42269// CHECK12-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 42270// CHECK12-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 42271// CHECK12-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42272// CHECK12-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42273// CHECK12-NEXT: [[I3:%.*]] = alloca i32, align 42274// CHECK12-NEXT: [[N_CASTED:%.*]] = alloca i32, align 42275// CHECK12-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42276// CHECK12-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42277// CHECK12-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42278// CHECK12-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42279// CHECK12-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42280// CHECK12-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 42281// CHECK12-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 42282// CHECK12-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42283// CHECK12-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 02284// CHECK12-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12285// CHECK12-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12286// CHECK12-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42287// CHECK12-NEXT: store i32 0, ptr [[I]], align 42288// CHECK12-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42289// CHECK12-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]2290// CHECK12-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2291// CHECK12: omp.precond.then:2292// CHECK12-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 42293// CHECK12-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42294// CHECK12-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_COMB_UB]], align 42295// CHECK12-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42296// CHECK12-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42297// CHECK12-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42298// CHECK12-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 42299// CHECK12-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP6]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2300// CHECK12-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42301// CHECK12-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42302// CHECK12-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP7]], [[TMP8]]2303// CHECK12-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2304// CHECK12: cond.true:2305// CHECK12-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42306// CHECK12-NEXT: br label [[COND_END:%.*]]2307// CHECK12: cond.false:2308// CHECK12-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42309// CHECK12-NEXT: br label [[COND_END]]2310// CHECK12: cond.end:2311// CHECK12-NEXT: [[COND:%.*]] = phi i32 [ [[TMP9]], [[COND_TRUE]] ], [ [[TMP10]], [[COND_FALSE]] ]2312// CHECK12-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 42313// CHECK12-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 42314// CHECK12-NEXT: store i32 [[TMP11]], ptr [[DOTOMP_IV]], align 42315// CHECK12-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2316// CHECK12: omp.inner.for.cond:2317// CHECK12-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42318// CHECK12-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42319// CHECK12-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP12]], [[TMP13]]2320// CHECK12-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2321// CHECK12: omp.inner.for.body:2322// CHECK12-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 42323// CHECK12-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42324// CHECK12-NEXT: [[TMP16:%.*]] = load i32, ptr [[N_ADDR]], align 42325// CHECK12-NEXT: store i32 [[TMP16]], ptr [[N_CASTED]], align 42326// CHECK12-NEXT: [[TMP17:%.*]] = load i32, ptr [[N_CASTED]], align 42327// CHECK12-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 4, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined.omp_outlined, i32 [[TMP14]], i32 [[TMP15]], i32 [[TMP17]], ptr [[TMP0]])2328// CHECK12-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2329// CHECK12: omp.inner.for.inc:2330// CHECK12-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42331// CHECK12-NEXT: [[TMP19:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 42332// CHECK12-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP18]], [[TMP19]]2333// CHECK12-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 42334// CHECK12-NEXT: br label [[OMP_INNER_FOR_COND]]2335// CHECK12: omp.inner.for.end:2336// CHECK12-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2337// CHECK12: omp.loop.exit:2338// CHECK12-NEXT: [[TMP20:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42339// CHECK12-NEXT: [[TMP21:%.*]] = load i32, ptr [[TMP20]], align 42340// CHECK12-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP21]])2341// CHECK12-NEXT: br label [[OMP_PRECOND_END]]2342// CHECK12: omp.precond.end:2343// CHECK12-NEXT: ret void2344//2345//2346// CHECK12-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined.omp_outlined2347// CHECK12-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[DOTPREVIOUS_LB_:%.*]], i32 noundef [[DOTPREVIOUS_UB_:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR0]] {2348// CHECK12-NEXT: entry:2349// CHECK12-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42350// CHECK12-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42351// CHECK12-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i32, align 42352// CHECK12-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i32, align 42353// CHECK12-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42354// CHECK12-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42355// CHECK12-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42356// CHECK12-NEXT: [[TMP:%.*]] = alloca i32, align 42357// CHECK12-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42358// CHECK12-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42359// CHECK12-NEXT: [[I:%.*]] = alloca i32, align 42360// CHECK12-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 42361// CHECK12-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 42362// CHECK12-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42363// CHECK12-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42364// CHECK12-NEXT: [[I3:%.*]] = alloca i32, align 42365// CHECK12-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42366// CHECK12-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42367// CHECK12-NEXT: store i32 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 42368// CHECK12-NEXT: store i32 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 42369// CHECK12-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42370// CHECK12-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42371// CHECK12-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42372// CHECK12-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 42373// CHECK12-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 42374// CHECK12-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42375// CHECK12-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 02376// CHECK12-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12377// CHECK12-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12378// CHECK12-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42379// CHECK12-NEXT: store i32 0, ptr [[I]], align 42380// CHECK12-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42381// CHECK12-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]2382// CHECK12-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2383// CHECK12: omp.precond.then:2384// CHECK12-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 42385// CHECK12-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42386// CHECK12-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_UB]], align 42387// CHECK12-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTPREVIOUS_LB__ADDR]], align 42388// CHECK12-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTPREVIOUS_UB__ADDR]], align 42389// CHECK12-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_LB]], align 42390// CHECK12-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_UB]], align 42391// CHECK12-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42392// CHECK12-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42393// CHECK12-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42394// CHECK12-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 42395// CHECK12-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP8]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2396// CHECK12-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42397// CHECK12-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42398// CHECK12-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]2399// CHECK12-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2400// CHECK12: cond.true:2401// CHECK12-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42402// CHECK12-NEXT: br label [[COND_END:%.*]]2403// CHECK12: cond.false:2404// CHECK12-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42405// CHECK12-NEXT: br label [[COND_END]]2406// CHECK12: cond.end:2407// CHECK12-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]2408// CHECK12-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 42409// CHECK12-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42410// CHECK12-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 42411// CHECK12-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2412// CHECK12: omp.inner.for.cond:2413// CHECK12-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42414// CHECK12-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42415// CHECK12-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]2416// CHECK12-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2417// CHECK12: omp.inner.for.body:2418// CHECK12-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42419// CHECK12-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 12420// CHECK12-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]2421// CHECK12-NEXT: store i32 [[ADD]], ptr [[I3]], align 42422// CHECK12-NEXT: [[TMP17:%.*]] = load i32, ptr [[I3]], align 42423// CHECK12-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i32 0, i32 [[TMP17]]2424// CHECK12-NEXT: store i32 0, ptr [[ARRAYIDX]], align 42425// CHECK12-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]2426// CHECK12: omp.body.continue:2427// CHECK12-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2428// CHECK12: omp.inner.for.inc:2429// CHECK12-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42430// CHECK12-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP18]], 12431// CHECK12-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 42432// CHECK12-NEXT: br label [[OMP_INNER_FOR_COND]]2433// CHECK12: omp.inner.for.end:2434// CHECK12-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2435// CHECK12: omp.loop.exit:2436// CHECK12-NEXT: [[TMP19:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42437// CHECK12-NEXT: [[TMP20:%.*]] = load i32, ptr [[TMP19]], align 42438// CHECK12-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP20]])2439// CHECK12-NEXT: br label [[OMP_PRECOND_END]]2440// CHECK12: omp.precond.end:2441// CHECK12-NEXT: ret void2442//2443//2444// CHECK12-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l572445// CHECK12-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], ptr noundef [[G:%.*]]) #[[ATTR0]] {2446// CHECK12-NEXT: entry:2447// CHECK12-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42448// CHECK12-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42449// CHECK12-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42450// CHECK12-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 42451// CHECK12-NEXT: [[N_CASTED:%.*]] = alloca i32, align 42452// CHECK12-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42453// CHECK12-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42454// CHECK12-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42455// CHECK12-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 42456// CHECK12-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42457// CHECK12-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 42458// CHECK12-NEXT: store i32 [[TMP1]], ptr [[N_CASTED]], align 42459// CHECK12-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_CASTED]], align 42460// CHECK12-NEXT: [[TMP3:%.*]] = load ptr, ptr [[G_ADDR]], align 42461// CHECK12-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.omp_outlined, i32 [[TMP2]], ptr [[TMP0]], ptr [[TMP3]])2462// CHECK12-NEXT: ret void2463//2464//2465// CHECK12-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.omp_outlined2466// CHECK12-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], ptr noundef [[G:%.*]]) #[[ATTR0]] {2467// CHECK12-NEXT: entry:2468// CHECK12-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42469// CHECK12-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42470// CHECK12-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42471// CHECK12-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42472// CHECK12-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 42473// CHECK12-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42474// CHECK12-NEXT: [[TMP:%.*]] = alloca i32, align 42475// CHECK12-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42476// CHECK12-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42477// CHECK12-NEXT: [[I:%.*]] = alloca i32, align 42478// CHECK12-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 42479// CHECK12-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 42480// CHECK12-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42481// CHECK12-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42482// CHECK12-NEXT: [[I3:%.*]] = alloca i32, align 42483// CHECK12-NEXT: [[N_CASTED:%.*]] = alloca i32, align 42484// CHECK12-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42485// CHECK12-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42486// CHECK12-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42487// CHECK12-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42488// CHECK12-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 42489// CHECK12-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42490// CHECK12-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 42491// CHECK12-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 42492// CHECK12-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42493// CHECK12-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 02494// CHECK12-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12495// CHECK12-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12496// CHECK12-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42497// CHECK12-NEXT: store i32 0, ptr [[I]], align 42498// CHECK12-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42499// CHECK12-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]2500// CHECK12-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2501// CHECK12: omp.precond.then:2502// CHECK12-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 42503// CHECK12-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42504// CHECK12-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_COMB_UB]], align 42505// CHECK12-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42506// CHECK12-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42507// CHECK12-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42508// CHECK12-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 42509// CHECK12-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP6]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2510// CHECK12-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42511// CHECK12-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42512// CHECK12-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP7]], [[TMP8]]2513// CHECK12-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2514// CHECK12: cond.true:2515// CHECK12-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42516// CHECK12-NEXT: br label [[COND_END:%.*]]2517// CHECK12: cond.false:2518// CHECK12-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42519// CHECK12-NEXT: br label [[COND_END]]2520// CHECK12: cond.end:2521// CHECK12-NEXT: [[COND:%.*]] = phi i32 [ [[TMP9]], [[COND_TRUE]] ], [ [[TMP10]], [[COND_FALSE]] ]2522// CHECK12-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 42523// CHECK12-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 42524// CHECK12-NEXT: store i32 [[TMP11]], ptr [[DOTOMP_IV]], align 42525// CHECK12-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2526// CHECK12: omp.inner.for.cond:2527// CHECK12-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42528// CHECK12-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42529// CHECK12-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP12]], [[TMP13]]2530// CHECK12-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2531// CHECK12: omp.inner.for.body:2532// CHECK12-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 42533// CHECK12-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42534// CHECK12-NEXT: [[TMP16:%.*]] = load i32, ptr [[N_ADDR]], align 42535// CHECK12-NEXT: store i32 [[TMP16]], ptr [[N_CASTED]], align 42536// CHECK12-NEXT: [[TMP17:%.*]] = load i32, ptr [[N_CASTED]], align 42537// CHECK12-NEXT: [[TMP18:%.*]] = load ptr, ptr [[G_ADDR]], align 42538// CHECK12-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 5, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.omp_outlined.omp_outlined, i32 [[TMP14]], i32 [[TMP15]], i32 [[TMP17]], ptr [[TMP0]], ptr [[TMP18]])2539// CHECK12-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2540// CHECK12: omp.inner.for.inc:2541// CHECK12-NEXT: [[TMP19:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42542// CHECK12-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 42543// CHECK12-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP19]], [[TMP20]]2544// CHECK12-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 42545// CHECK12-NEXT: br label [[OMP_INNER_FOR_COND]]2546// CHECK12: omp.inner.for.end:2547// CHECK12-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2548// CHECK12: omp.loop.exit:2549// CHECK12-NEXT: [[TMP21:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42550// CHECK12-NEXT: [[TMP22:%.*]] = load i32, ptr [[TMP21]], align 42551// CHECK12-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP22]])2552// CHECK12-NEXT: br label [[OMP_PRECOND_END]]2553// CHECK12: omp.precond.end:2554// CHECK12-NEXT: ret void2555//2556//2557// CHECK12-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l57.omp_outlined.omp_outlined2558// CHECK12-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[DOTPREVIOUS_LB_:%.*]], i32 noundef [[DOTPREVIOUS_UB_:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], ptr noundef [[G:%.*]]) #[[ATTR0]] {2559// CHECK12-NEXT: entry:2560// CHECK12-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42561// CHECK12-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42562// CHECK12-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i32, align 42563// CHECK12-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i32, align 42564// CHECK12-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42565// CHECK12-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42566// CHECK12-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 42567// CHECK12-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42568// CHECK12-NEXT: [[TMP:%.*]] = alloca i32, align 42569// CHECK12-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42570// CHECK12-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42571// CHECK12-NEXT: [[I:%.*]] = alloca i32, align 42572// CHECK12-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 42573// CHECK12-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 42574// CHECK12-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42575// CHECK12-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42576// CHECK12-NEXT: [[I3:%.*]] = alloca i32, align 42577// CHECK12-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42578// CHECK12-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42579// CHECK12-NEXT: store i32 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 42580// CHECK12-NEXT: store i32 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 42581// CHECK12-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42582// CHECK12-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42583// CHECK12-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 42584// CHECK12-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42585// CHECK12-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 42586// CHECK12-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 42587// CHECK12-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42588// CHECK12-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 02589// CHECK12-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12590// CHECK12-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12591// CHECK12-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42592// CHECK12-NEXT: store i32 0, ptr [[I]], align 42593// CHECK12-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42594// CHECK12-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]2595// CHECK12-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2596// CHECK12: omp.precond.then:2597// CHECK12-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 42598// CHECK12-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42599// CHECK12-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_UB]], align 42600// CHECK12-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTPREVIOUS_LB__ADDR]], align 42601// CHECK12-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTPREVIOUS_UB__ADDR]], align 42602// CHECK12-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_LB]], align 42603// CHECK12-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_UB]], align 42604// CHECK12-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42605// CHECK12-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42606// CHECK12-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42607// CHECK12-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 42608// CHECK12-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2]], i32 [[TMP8]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2609// CHECK12-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42610// CHECK12-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42611// CHECK12-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]2612// CHECK12-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2613// CHECK12: cond.true:2614// CHECK12-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42615// CHECK12-NEXT: br label [[COND_END:%.*]]2616// CHECK12: cond.false:2617// CHECK12-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42618// CHECK12-NEXT: br label [[COND_END]]2619// CHECK12: cond.end:2620// CHECK12-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]2621// CHECK12-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 42622// CHECK12-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42623// CHECK12-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 42624// CHECK12-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2625// CHECK12: omp.inner.for.cond:2626// CHECK12-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42627// CHECK12-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42628// CHECK12-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]2629// CHECK12-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2630// CHECK12: omp.inner.for.body:2631// CHECK12-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42632// CHECK12-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 12633// CHECK12-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]2634// CHECK12-NEXT: store i32 [[ADD]], ptr [[I3]], align 42635// CHECK12-NEXT: [[TMP17:%.*]] = load ptr, ptr [[G_ADDR]], align 42636// CHECK12-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP17]], i32 02637// CHECK12-NEXT: [[TMP18:%.*]] = load i32, ptr [[ARRAYIDX]], align 42638// CHECK12-NEXT: [[TMP19:%.*]] = load i32, ptr [[I3]], align 42639// CHECK12-NEXT: [[ARRAYIDX6:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i32 0, i32 [[TMP19]]2640// CHECK12-NEXT: store i32 [[TMP18]], ptr [[ARRAYIDX6]], align 42641// CHECK12-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]2642// CHECK12: omp.body.continue:2643// CHECK12-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2644// CHECK12: omp.inner.for.inc:2645// CHECK12-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42646// CHECK12-NEXT: [[ADD7:%.*]] = add nsw i32 [[TMP20]], 12647// CHECK12-NEXT: store i32 [[ADD7]], ptr [[DOTOMP_IV]], align 42648// CHECK12-NEXT: br label [[OMP_INNER_FOR_COND]]2649// CHECK12: omp.inner.for.end:2650// CHECK12-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2651// CHECK12: omp.loop.exit:2652// CHECK12-NEXT: [[TMP21:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42653// CHECK12-NEXT: [[TMP22:%.*]] = load i32, ptr [[TMP21]], align 42654// CHECK12-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP22]])2655// CHECK12-NEXT: br label [[OMP_PRECOND_END]]2656// CHECK12: omp.precond.end:2657// CHECK12-NEXT: ret void2658//2659