2739 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 distribute parallel for num_teams(te), thread_limit(th)52 for(int i = 0; i < n; i++) {53 a[i] = 0;54 #pragma omp cancel for55 }56 57 {{{58 #pragma omp target teams distribute parallel for is_device_ptr(g)59 for(int i = 0; i < n; i++) {60 a[i] = g[0];61 }62 }}}63 64 // outlined target regions65 66 67 68 69 return a[0];70}71 72#endif // CK173#endif // HEADER74// CHECK1-LABEL: define {{[^@]+}}@_Z16target_teams_funPi75// CHECK1-SAME: (ptr noundef [[G:%.*]]) #[[ATTR0:[0-9]+]] {76// CHECK1-NEXT: entry:77// CHECK1-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 878// CHECK1-NEXT: [[N:%.*]] = alloca i32, align 479// CHECK1-NEXT: [[A:%.*]] = alloca [1000 x i32], align 480// CHECK1-NEXT: [[TE:%.*]] = alloca i32, align 481// CHECK1-NEXT: [[TH:%.*]] = alloca i32, align 482// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 483// CHECK1-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 484// CHECK1-NEXT: [[N_CASTED:%.*]] = alloca i64, align 885// CHECK1-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 886// CHECK1-NEXT: [[DOTCAPTURE_EXPR__CASTED2:%.*]] = alloca i64, align 887// CHECK1-NEXT: [[N_CASTED3:%.*]] = alloca i64, align 888// CHECK1-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 889// CHECK1-NEXT: store i32 1000, ptr [[N]], align 490// CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr [[N]], align 491// CHECK1-NEXT: [[DIV:%.*]] = sdiv i32 [[TMP0]], 12892// CHECK1-NEXT: store i32 [[DIV]], ptr [[TE]], align 493// CHECK1-NEXT: store i32 128, ptr [[TH]], align 494// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TE]], align 495// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 496// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[TH]], align 497// CHECK1-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_1]], align 498// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[N]], align 499// CHECK1-NEXT: store i32 [[TMP3]], ptr [[N_CASTED]], align 4100// CHECK1-NEXT: [[TMP4:%.*]] = load i64, ptr [[N_CASTED]], align 8101// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4102// CHECK1-NEXT: store i32 [[TMP5]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 4103// CHECK1-NEXT: [[TMP6:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 8104// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4105// CHECK1-NEXT: store i32 [[TMP7]], ptr [[DOTCAPTURE_EXPR__CASTED2]], align 4106// CHECK1-NEXT: [[TMP8:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED2]], align 8107// 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]+]]108// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[N]], align 4109// CHECK1-NEXT: store i32 [[TMP9]], ptr [[N_CASTED3]], align 4110// CHECK1-NEXT: [[TMP10:%.*]] = load i64, ptr [[N_CASTED3]], align 8111// CHECK1-NEXT: [[TMP11:%.*]] = load ptr, ptr [[G_ADDR]], align 8112// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l58(i64 [[TMP10]], ptr [[A]], ptr [[TMP11]]) #[[ATTR2]]113// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[A]], i64 0, i64 0114// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[ARRAYIDX]], align 4115// CHECK1-NEXT: ret i32 [[TMP12]]116//117//118// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51119// 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]+]] {120// CHECK1-NEXT: entry:121// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8122// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8123// CHECK1-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 8124// CHECK1-NEXT: [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca i64, align 8125// CHECK1-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8126// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB3:[0-9]+]])127// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8128// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8129// CHECK1-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 8130// CHECK1-NEXT: store i64 [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 8131// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 8132// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 4133// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR2]], align 4134// CHECK1-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB3]], i32 [[TMP0]], i32 [[TMP2]], i32 [[TMP3]])135// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[N_ADDR]], align 4136// CHECK1-NEXT: store i32 [[TMP4]], ptr [[N_CASTED]], align 4137// CHECK1-NEXT: [[TMP5:%.*]] = load i64, ptr [[N_CASTED]], align 8138// 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]])139// CHECK1-NEXT: ret void140//141//142// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined143// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR1]] {144// CHECK1-NEXT: entry:145// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8146// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8147// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8148// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8149// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4150// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4151// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4152// CHECK1-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4153// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4154// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4155// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4156// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4157// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4158// CHECK1-NEXT: [[I3:%.*]] = alloca i32, align 4159// CHECK1-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8160// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8161// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8162// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8163// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8164// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8165// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 4166// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 4167// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4168// CHECK1-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 0169// CHECK1-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1170// CHECK1-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1171// CHECK1-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4172// CHECK1-NEXT: store i32 0, ptr [[I]], align 4173// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4174// CHECK1-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]175// CHECK1-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]176// CHECK1: omp.precond.then:177// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4178// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4179// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_COMB_UB]], align 4180// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4181// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4182// CHECK1-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8183// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 4184// 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)185// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4186// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4187// CHECK1-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP7]], [[TMP8]]188// CHECK1-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]189// CHECK1: cond.true:190// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4191// CHECK1-NEXT: br label [[COND_END:%.*]]192// CHECK1: cond.false:193// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4194// CHECK1-NEXT: br label [[COND_END]]195// CHECK1: cond.end:196// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ [[TMP9]], [[COND_TRUE]] ], [ [[TMP10]], [[COND_FALSE]] ]197// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4198// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4199// CHECK1-NEXT: store i32 [[TMP11]], ptr [[DOTOMP_IV]], align 4200// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]201// CHECK1: omp.inner.for.cond:202// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4203// CHECK1-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4204// CHECK1-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP12]], [[TMP13]]205// CHECK1-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]206// CHECK1: omp.inner.for.body:207// CHECK1-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4208// CHECK1-NEXT: [[TMP15:%.*]] = zext i32 [[TMP14]] to i64209// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4210// CHECK1-NEXT: [[TMP17:%.*]] = zext i32 [[TMP16]] to i64211// CHECK1-NEXT: [[TMP18:%.*]] = load i32, ptr [[N_ADDR]], align 4212// CHECK1-NEXT: store i32 [[TMP18]], ptr [[N_CASTED]], align 4213// CHECK1-NEXT: [[TMP19:%.*]] = load i64, ptr [[N_CASTED]], align 8214// 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]])215// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]216// CHECK1: omp.inner.for.inc:217// CHECK1-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4218// CHECK1-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4219// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP20]], [[TMP21]]220// CHECK1-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4221// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]222// CHECK1: omp.inner.for.end:223// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]224// CHECK1: omp.loop.exit:225// CHECK1-NEXT: [[TMP22:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8226// CHECK1-NEXT: [[TMP23:%.*]] = load i32, ptr [[TMP22]], align 4227// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP23]])228// CHECK1-NEXT: br label [[OMP_PRECOND_END]]229// CHECK1: omp.precond.end:230// CHECK1-NEXT: ret void231//232//233// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined.omp_outlined234// 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]] {235// CHECK1-NEXT: entry:236// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8237// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8238// CHECK1-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8239// CHECK1-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8240// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8241// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8242// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4243// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4244// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4245// CHECK1-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4246// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4247// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4248// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4249// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4250// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4251// CHECK1-NEXT: [[I4:%.*]] = alloca i32, align 4252// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8253// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8254// CHECK1-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8255// CHECK1-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8256// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8257// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8258// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8259// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 4260// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 4261// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4262// CHECK1-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 0263// CHECK1-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1264// CHECK1-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1265// CHECK1-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4266// CHECK1-NEXT: store i32 0, ptr [[I]], align 4267// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4268// CHECK1-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]269// CHECK1-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]270// CHECK1: omp.precond.then:271// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4272// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4273// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_UB]], align 4274// CHECK1-NEXT: [[TMP5:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8275// CHECK1-NEXT: [[CONV:%.*]] = trunc i64 [[TMP5]] to i32276// CHECK1-NEXT: [[TMP6:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8277// CHECK1-NEXT: [[CONV3:%.*]] = trunc i64 [[TMP6]] to i32278// CHECK1-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 4279// CHECK1-NEXT: store i32 [[CONV3]], ptr [[DOTOMP_UB]], align 4280// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4281// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4282// CHECK1-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8283// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4284// 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)285// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4286// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4287// CHECK1-NEXT: [[CMP5:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]288// CHECK1-NEXT: br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]289// CHECK1: cond.true:290// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4291// CHECK1-NEXT: br label [[COND_END:%.*]]292// CHECK1: cond.false:293// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4294// CHECK1-NEXT: br label [[COND_END]]295// CHECK1: cond.end:296// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]297// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4298// CHECK1-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4299// CHECK1-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 4300// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]301// CHECK1: omp.inner.for.cond:302// CHECK1-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4303// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4304// CHECK1-NEXT: [[CMP6:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]305// CHECK1-NEXT: br i1 [[CMP6]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]306// CHECK1: omp.inner.for.body:307// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4308// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 1309// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]310// CHECK1-NEXT: store i32 [[ADD]], ptr [[I4]], align 4311// CHECK1-NEXT: [[TMP17:%.*]] = load i32, ptr [[I4]], align 4312// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP17]] to i64313// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]314// CHECK1-NEXT: store i32 0, ptr [[ARRAYIDX]], align 4315// CHECK1-NEXT: [[TMP18:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8316// CHECK1-NEXT: [[TMP19:%.*]] = load i32, ptr [[TMP18]], align 4317// CHECK1-NEXT: [[TMP20:%.*]] = call i32 @__kmpc_cancel(ptr @[[GLOB3]], i32 [[TMP19]], i32 2)318// CHECK1-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0319// CHECK1-NEXT: br i1 [[TMP21]], label [[DOTCANCEL_EXIT:%.*]], label [[DOTCANCEL_CONTINUE:%.*]]320// CHECK1: .cancel.exit:321// CHECK1-NEXT: br label [[CANCEL_EXIT:%.*]]322// CHECK1: .cancel.continue:323// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]324// CHECK1: omp.body.continue:325// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]326// CHECK1: omp.inner.for.inc:327// CHECK1-NEXT: [[TMP22:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4328// CHECK1-NEXT: [[ADD7:%.*]] = add nsw i32 [[TMP22]], 1329// CHECK1-NEXT: store i32 [[ADD7]], ptr [[DOTOMP_IV]], align 4330// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]331// CHECK1: omp.inner.for.end:332// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]333// CHECK1: omp.loop.exit:334// CHECK1-NEXT: [[TMP23:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8335// CHECK1-NEXT: [[TMP24:%.*]] = load i32, ptr [[TMP23]], align 4336// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP24]])337// CHECK1-NEXT: br label [[OMP_PRECOND_END]]338// CHECK1: cancel.exit:339// CHECK1-NEXT: [[TMP25:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8340// CHECK1-NEXT: [[TMP26:%.*]] = load i32, ptr [[TMP25]], align 4341// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP26]])342// CHECK1-NEXT: br label [[CANCEL_CONT:%.*]]343// CHECK1: omp.precond.end:344// CHECK1-NEXT: br label [[CANCEL_CONT]]345// CHECK1: cancel.cont:346// CHECK1-NEXT: ret void347//348//349// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l58350// CHECK1-SAME: (i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], ptr noundef [[G:%.*]]) #[[ATTR1]] {351// CHECK1-NEXT: entry:352// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8353// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8354// CHECK1-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 8355// CHECK1-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8356// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8357// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8358// CHECK1-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 8359// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8360// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 4361// CHECK1-NEXT: store i32 [[TMP1]], ptr [[N_CASTED]], align 4362// CHECK1-NEXT: [[TMP2:%.*]] = load i64, ptr [[N_CASTED]], align 8363// CHECK1-NEXT: [[TMP3:%.*]] = load ptr, ptr [[G_ADDR]], align 8364// 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_l58.omp_outlined, i64 [[TMP2]], ptr [[TMP0]], ptr [[TMP3]])365// CHECK1-NEXT: ret void366//367//368// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l58.omp_outlined369// 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]] {370// CHECK1-NEXT: entry:371// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8372// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8373// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8374// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8375// CHECK1-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 8376// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4377// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4378// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4379// CHECK1-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4380// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4381// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4382// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4383// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4384// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4385// CHECK1-NEXT: [[I3:%.*]] = alloca i32, align 4386// CHECK1-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8387// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8388// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8389// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8390// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8391// CHECK1-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 8392// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8393// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 4394// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 4395// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4396// CHECK1-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 0397// CHECK1-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1398// CHECK1-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1399// CHECK1-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4400// CHECK1-NEXT: store i32 0, ptr [[I]], align 4401// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4402// CHECK1-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]403// CHECK1-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]404// CHECK1: omp.precond.then:405// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4406// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4407// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_COMB_UB]], align 4408// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4409// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4410// CHECK1-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8411// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 4412// 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)413// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4414// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4415// CHECK1-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP7]], [[TMP8]]416// CHECK1-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]417// CHECK1: cond.true:418// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4419// CHECK1-NEXT: br label [[COND_END:%.*]]420// CHECK1: cond.false:421// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4422// CHECK1-NEXT: br label [[COND_END]]423// CHECK1: cond.end:424// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ [[TMP9]], [[COND_TRUE]] ], [ [[TMP10]], [[COND_FALSE]] ]425// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4426// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4427// CHECK1-NEXT: store i32 [[TMP11]], ptr [[DOTOMP_IV]], align 4428// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]429// CHECK1: omp.inner.for.cond:430// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4431// CHECK1-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4432// CHECK1-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP12]], [[TMP13]]433// CHECK1-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]434// CHECK1: omp.inner.for.body:435// CHECK1-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4436// CHECK1-NEXT: [[TMP15:%.*]] = zext i32 [[TMP14]] to i64437// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4438// CHECK1-NEXT: [[TMP17:%.*]] = zext i32 [[TMP16]] to i64439// CHECK1-NEXT: [[TMP18:%.*]] = load i32, ptr [[N_ADDR]], align 4440// CHECK1-NEXT: store i32 [[TMP18]], ptr [[N_CASTED]], align 4441// CHECK1-NEXT: [[TMP19:%.*]] = load i64, ptr [[N_CASTED]], align 8442// CHECK1-NEXT: [[TMP20:%.*]] = load ptr, ptr [[G_ADDR]], align 8443// 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_l58.omp_outlined.omp_outlined, i64 [[TMP15]], i64 [[TMP17]], i64 [[TMP19]], ptr [[TMP0]], ptr [[TMP20]])444// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]445// CHECK1: omp.inner.for.inc:446// CHECK1-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4447// CHECK1-NEXT: [[TMP22:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4448// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP21]], [[TMP22]]449// CHECK1-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4450// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]451// CHECK1: omp.inner.for.end:452// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]453// CHECK1: omp.loop.exit:454// CHECK1-NEXT: [[TMP23:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8455// CHECK1-NEXT: [[TMP24:%.*]] = load i32, ptr [[TMP23]], align 4456// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP24]])457// CHECK1-NEXT: br label [[OMP_PRECOND_END]]458// CHECK1: omp.precond.end:459// CHECK1-NEXT: ret void460//461//462// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l58.omp_outlined.omp_outlined463// 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]] {464// CHECK1-NEXT: entry:465// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8466// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8467// CHECK1-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8468// CHECK1-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8469// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8470// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8471// CHECK1-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 8472// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4473// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4474// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4475// CHECK1-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4476// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4477// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4478// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4479// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4480// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4481// CHECK1-NEXT: [[I4:%.*]] = alloca i32, align 4482// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8483// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8484// CHECK1-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8485// CHECK1-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8486// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8487// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8488// CHECK1-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 8489// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8490// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 4491// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 4492// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4493// CHECK1-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 0494// CHECK1-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1495// CHECK1-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1496// CHECK1-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4497// CHECK1-NEXT: store i32 0, ptr [[I]], align 4498// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4499// CHECK1-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]500// CHECK1-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]501// CHECK1: omp.precond.then:502// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4503// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4504// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_UB]], align 4505// CHECK1-NEXT: [[TMP5:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8506// CHECK1-NEXT: [[CONV:%.*]] = trunc i64 [[TMP5]] to i32507// CHECK1-NEXT: [[TMP6:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8508// CHECK1-NEXT: [[CONV3:%.*]] = trunc i64 [[TMP6]] to i32509// CHECK1-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 4510// CHECK1-NEXT: store i32 [[CONV3]], ptr [[DOTOMP_UB]], align 4511// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4512// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4513// CHECK1-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8514// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4515// 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)516// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4517// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4518// CHECK1-NEXT: [[CMP5:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]519// CHECK1-NEXT: br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]520// CHECK1: cond.true:521// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4522// CHECK1-NEXT: br label [[COND_END:%.*]]523// CHECK1: cond.false:524// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4525// CHECK1-NEXT: br label [[COND_END]]526// CHECK1: cond.end:527// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]528// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4529// CHECK1-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4530// CHECK1-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 4531// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]532// CHECK1: omp.inner.for.cond:533// CHECK1-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4534// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4535// CHECK1-NEXT: [[CMP6:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]536// CHECK1-NEXT: br i1 [[CMP6]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]537// CHECK1: omp.inner.for.body:538// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4539// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 1540// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]541// CHECK1-NEXT: store i32 [[ADD]], ptr [[I4]], align 4542// CHECK1-NEXT: [[TMP17:%.*]] = load ptr, ptr [[G_ADDR]], align 8543// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP17]], i64 0544// CHECK1-NEXT: [[TMP18:%.*]] = load i32, ptr [[ARRAYIDX]], align 4545// CHECK1-NEXT: [[TMP19:%.*]] = load i32, ptr [[I4]], align 4546// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP19]] to i64547// CHECK1-NEXT: [[ARRAYIDX7:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]548// CHECK1-NEXT: store i32 [[TMP18]], ptr [[ARRAYIDX7]], align 4549// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]550// CHECK1: omp.body.continue:551// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]552// CHECK1: omp.inner.for.inc:553// CHECK1-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4554// CHECK1-NEXT: [[ADD8:%.*]] = add nsw i32 [[TMP20]], 1555// CHECK1-NEXT: store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 4556// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]557// CHECK1: omp.inner.for.end:558// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]559// CHECK1: omp.loop.exit:560// CHECK1-NEXT: [[TMP21:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8561// CHECK1-NEXT: [[TMP22:%.*]] = load i32, ptr [[TMP21]], align 4562// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP22]])563// CHECK1-NEXT: br label [[OMP_PRECOND_END]]564// CHECK1: omp.precond.end:565// CHECK1-NEXT: ret void566//567//568// CHECK2-LABEL: define {{[^@]+}}@_Z16target_teams_funPi569// CHECK2-SAME: (ptr noundef [[G:%.*]]) #[[ATTR0:[0-9]+]] {570// CHECK2-NEXT: entry:571// CHECK2-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 8572// CHECK2-NEXT: [[N:%.*]] = alloca i32, align 4573// CHECK2-NEXT: [[A:%.*]] = alloca [1000 x i32], align 4574// CHECK2-NEXT: [[TE:%.*]] = alloca i32, align 4575// CHECK2-NEXT: [[TH:%.*]] = alloca i32, align 4576// CHECK2-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4577// CHECK2-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4578// CHECK2-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8579// CHECK2-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 8580// CHECK2-NEXT: [[DOTCAPTURE_EXPR__CASTED2:%.*]] = alloca i64, align 8581// CHECK2-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [4 x ptr], align 8582// CHECK2-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [4 x ptr], align 8583// CHECK2-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [4 x ptr], align 8584// CHECK2-NEXT: [[TMP:%.*]] = alloca i32, align 4585// CHECK2-NEXT: [[DOTCAPTURE_EXPR_3:%.*]] = alloca i32, align 4586// CHECK2-NEXT: [[DOTCAPTURE_EXPR_4:%.*]] = alloca i32, align 4587// CHECK2-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8588// CHECK2-NEXT: [[N_CASTED7:%.*]] = alloca i64, align 8589// CHECK2-NEXT: [[DOTOFFLOAD_BASEPTRS8:%.*]] = alloca [3 x ptr], align 8590// CHECK2-NEXT: [[DOTOFFLOAD_PTRS9:%.*]] = alloca [3 x ptr], align 8591// CHECK2-NEXT: [[DOTOFFLOAD_MAPPERS10:%.*]] = alloca [3 x ptr], align 8592// CHECK2-NEXT: [[_TMP11:%.*]] = alloca i32, align 4593// CHECK2-NEXT: [[DOTCAPTURE_EXPR_12:%.*]] = alloca i32, align 4594// CHECK2-NEXT: [[DOTCAPTURE_EXPR_13:%.*]] = alloca i32, align 4595// CHECK2-NEXT: [[KERNEL_ARGS18:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8596// CHECK2-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 8597// CHECK2-NEXT: store i32 1000, ptr [[N]], align 4598// CHECK2-NEXT: [[TMP0:%.*]] = load i32, ptr [[N]], align 4599// CHECK2-NEXT: [[DIV:%.*]] = sdiv i32 [[TMP0]], 128600// CHECK2-NEXT: store i32 [[DIV]], ptr [[TE]], align 4601// CHECK2-NEXT: store i32 128, ptr [[TH]], align 4602// CHECK2-NEXT: [[TMP1:%.*]] = load i32, ptr [[TE]], align 4603// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 4604// CHECK2-NEXT: [[TMP2:%.*]] = load i32, ptr [[TH]], align 4605// CHECK2-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_1]], align 4606// CHECK2-NEXT: [[TMP3:%.*]] = load i32, ptr [[N]], align 4607// CHECK2-NEXT: store i32 [[TMP3]], ptr [[N_CASTED]], align 4608// CHECK2-NEXT: [[TMP4:%.*]] = load i64, ptr [[N_CASTED]], align 8609// CHECK2-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4610// CHECK2-NEXT: store i32 [[TMP5]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 4611// CHECK2-NEXT: [[TMP6:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 8612// CHECK2-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4613// CHECK2-NEXT: store i32 [[TMP7]], ptr [[DOTCAPTURE_EXPR__CASTED2]], align 4614// CHECK2-NEXT: [[TMP8:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED2]], align 8615// CHECK2-NEXT: [[TMP9:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0616// CHECK2-NEXT: store i64 [[TMP4]], ptr [[TMP9]], align 8617// CHECK2-NEXT: [[TMP10:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0618// CHECK2-NEXT: store i64 [[TMP4]], ptr [[TMP10]], align 8619// CHECK2-NEXT: [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0620// CHECK2-NEXT: store ptr null, ptr [[TMP11]], align 8621// CHECK2-NEXT: [[TMP12:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1622// CHECK2-NEXT: store ptr [[A]], ptr [[TMP12]], align 8623// CHECK2-NEXT: [[TMP13:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1624// CHECK2-NEXT: store ptr [[A]], ptr [[TMP13]], align 8625// CHECK2-NEXT: [[TMP14:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1626// CHECK2-NEXT: store ptr null, ptr [[TMP14]], align 8627// CHECK2-NEXT: [[TMP15:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2628// CHECK2-NEXT: store i64 [[TMP6]], ptr [[TMP15]], align 8629// CHECK2-NEXT: [[TMP16:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2630// CHECK2-NEXT: store i64 [[TMP6]], ptr [[TMP16]], align 8631// CHECK2-NEXT: [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2632// CHECK2-NEXT: store ptr null, ptr [[TMP17]], align 8633// CHECK2-NEXT: [[TMP18:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3634// CHECK2-NEXT: store i64 [[TMP8]], ptr [[TMP18]], align 8635// CHECK2-NEXT: [[TMP19:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3636// CHECK2-NEXT: store i64 [[TMP8]], ptr [[TMP19]], align 8637// CHECK2-NEXT: [[TMP20:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3638// CHECK2-NEXT: store ptr null, ptr [[TMP20]], align 8639// CHECK2-NEXT: [[TMP21:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0640// CHECK2-NEXT: [[TMP22:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0641// CHECK2-NEXT: [[TMP23:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4642// CHECK2-NEXT: [[TMP24:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4643// CHECK2-NEXT: [[TMP25:%.*]] = load i32, ptr [[N]], align 4644// CHECK2-NEXT: store i32 [[TMP25]], ptr [[DOTCAPTURE_EXPR_3]], align 4645// CHECK2-NEXT: [[TMP26:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_3]], align 4646// CHECK2-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP26]], 0647// CHECK2-NEXT: [[DIV5:%.*]] = sdiv i32 [[SUB]], 1648// CHECK2-NEXT: [[SUB6:%.*]] = sub nsw i32 [[DIV5]], 1649// CHECK2-NEXT: store i32 [[SUB6]], ptr [[DOTCAPTURE_EXPR_4]], align 4650// CHECK2-NEXT: [[TMP27:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 4651// CHECK2-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP27]], 1652// CHECK2-NEXT: [[TMP28:%.*]] = zext i32 [[ADD]] to i64653// CHECK2-NEXT: [[TMP29:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP23]], 0654// CHECK2-NEXT: [[TMP30:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP24]], 0655// CHECK2-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0656// CHECK2-NEXT: store i32 3, ptr [[TMP31]], align 4657// CHECK2-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1658// CHECK2-NEXT: store i32 4, ptr [[TMP32]], align 4659// CHECK2-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2660// CHECK2-NEXT: store ptr [[TMP21]], ptr [[TMP33]], align 8661// CHECK2-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3662// CHECK2-NEXT: store ptr [[TMP22]], ptr [[TMP34]], align 8663// CHECK2-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4664// CHECK2-NEXT: store ptr @.offload_sizes, ptr [[TMP35]], align 8665// CHECK2-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5666// CHECK2-NEXT: store ptr @.offload_maptypes, ptr [[TMP36]], align 8667// CHECK2-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6668// CHECK2-NEXT: store ptr null, ptr [[TMP37]], align 8669// CHECK2-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7670// CHECK2-NEXT: store ptr null, ptr [[TMP38]], align 8671// CHECK2-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8672// CHECK2-NEXT: store i64 [[TMP28]], ptr [[TMP39]], align 8673// CHECK2-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9674// CHECK2-NEXT: store i64 0, ptr [[TMP40]], align 8675// CHECK2-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10676// CHECK2-NEXT: store [3 x i32] [[TMP29]], ptr [[TMP41]], align 4677// CHECK2-NEXT: [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11678// CHECK2-NEXT: store [3 x i32] [[TMP30]], ptr [[TMP42]], align 4679// CHECK2-NEXT: [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12680// CHECK2-NEXT: store i32 0, ptr [[TMP43]], align 4681// CHECK2-NEXT: [[TMP44:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3:[0-9]+]], i64 -1, i32 [[TMP23]], i32 [[TMP24]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.region_id, ptr [[KERNEL_ARGS]])682// CHECK2-NEXT: [[TMP45:%.*]] = icmp ne i32 [[TMP44]], 0683// CHECK2-NEXT: br i1 [[TMP45]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]684// CHECK2: omp_offload.failed:685// 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]+]]686// CHECK2-NEXT: br label [[OMP_OFFLOAD_CONT]]687// CHECK2: omp_offload.cont:688// CHECK2-NEXT: [[TMP46:%.*]] = load i32, ptr [[N]], align 4689// CHECK2-NEXT: store i32 [[TMP46]], ptr [[N_CASTED7]], align 4690// CHECK2-NEXT: [[TMP47:%.*]] = load i64, ptr [[N_CASTED7]], align 8691// CHECK2-NEXT: [[TMP48:%.*]] = load ptr, ptr [[G_ADDR]], align 8692// CHECK2-NEXT: [[TMP49:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 0693// CHECK2-NEXT: store i64 [[TMP47]], ptr [[TMP49]], align 8694// CHECK2-NEXT: [[TMP50:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 0695// CHECK2-NEXT: store i64 [[TMP47]], ptr [[TMP50]], align 8696// CHECK2-NEXT: [[TMP51:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS10]], i64 0, i64 0697// CHECK2-NEXT: store ptr null, ptr [[TMP51]], align 8698// CHECK2-NEXT: [[TMP52:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 1699// CHECK2-NEXT: store ptr [[A]], ptr [[TMP52]], align 8700// CHECK2-NEXT: [[TMP53:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 1701// CHECK2-NEXT: store ptr [[A]], ptr [[TMP53]], align 8702// CHECK2-NEXT: [[TMP54:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS10]], i64 0, i64 1703// CHECK2-NEXT: store ptr null, ptr [[TMP54]], align 8704// CHECK2-NEXT: [[TMP55:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 2705// CHECK2-NEXT: store ptr [[TMP48]], ptr [[TMP55]], align 8706// CHECK2-NEXT: [[TMP56:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 2707// CHECK2-NEXT: store ptr [[TMP48]], ptr [[TMP56]], align 8708// CHECK2-NEXT: [[TMP57:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS10]], i64 0, i64 2709// CHECK2-NEXT: store ptr null, ptr [[TMP57]], align 8710// CHECK2-NEXT: [[TMP58:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 0711// CHECK2-NEXT: [[TMP59:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 0712// CHECK2-NEXT: [[TMP60:%.*]] = load i32, ptr [[N]], align 4713// CHECK2-NEXT: store i32 [[TMP60]], ptr [[DOTCAPTURE_EXPR_12]], align 4714// CHECK2-NEXT: [[TMP61:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_12]], align 4715// CHECK2-NEXT: [[SUB14:%.*]] = sub nsw i32 [[TMP61]], 0716// CHECK2-NEXT: [[DIV15:%.*]] = sdiv i32 [[SUB14]], 1717// CHECK2-NEXT: [[SUB16:%.*]] = sub nsw i32 [[DIV15]], 1718// CHECK2-NEXT: store i32 [[SUB16]], ptr [[DOTCAPTURE_EXPR_13]], align 4719// CHECK2-NEXT: [[TMP62:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_13]], align 4720// CHECK2-NEXT: [[ADD17:%.*]] = add nsw i32 [[TMP62]], 1721// CHECK2-NEXT: [[TMP63:%.*]] = zext i32 [[ADD17]] to i64722// CHECK2-NEXT: [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 0723// CHECK2-NEXT: store i32 3, ptr [[TMP64]], align 4724// CHECK2-NEXT: [[TMP65:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 1725// CHECK2-NEXT: store i32 3, ptr [[TMP65]], align 4726// CHECK2-NEXT: [[TMP66:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 2727// CHECK2-NEXT: store ptr [[TMP58]], ptr [[TMP66]], align 8728// CHECK2-NEXT: [[TMP67:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 3729// CHECK2-NEXT: store ptr [[TMP59]], ptr [[TMP67]], align 8730// CHECK2-NEXT: [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 4731// CHECK2-NEXT: store ptr @.offload_sizes.1, ptr [[TMP68]], align 8732// CHECK2-NEXT: [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 5733// CHECK2-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP69]], align 8734// CHECK2-NEXT: [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 6735// CHECK2-NEXT: store ptr null, ptr [[TMP70]], align 8736// CHECK2-NEXT: [[TMP71:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 7737// CHECK2-NEXT: store ptr null, ptr [[TMP71]], align 8738// CHECK2-NEXT: [[TMP72:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 8739// CHECK2-NEXT: store i64 [[TMP63]], ptr [[TMP72]], align 8740// CHECK2-NEXT: [[TMP73:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 9741// CHECK2-NEXT: store i64 0, ptr [[TMP73]], align 8742// CHECK2-NEXT: [[TMP74:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 10743// CHECK2-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP74]], align 4744// CHECK2-NEXT: [[TMP75:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 11745// CHECK2-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP75]], align 4746// CHECK2-NEXT: [[TMP76:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 12747// CHECK2-NEXT: store i32 0, ptr [[TMP76]], align 4748// CHECK2-NEXT: [[TMP77:%.*]] = 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_l58.region_id, ptr [[KERNEL_ARGS18]])749// CHECK2-NEXT: [[TMP78:%.*]] = icmp ne i32 [[TMP77]], 0750// CHECK2-NEXT: br i1 [[TMP78]], label [[OMP_OFFLOAD_FAILED19:%.*]], label [[OMP_OFFLOAD_CONT20:%.*]]751// CHECK2: omp_offload.failed19:752// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l58(i64 [[TMP47]], ptr [[A]], ptr [[TMP48]]) #[[ATTR2]]753// CHECK2-NEXT: br label [[OMP_OFFLOAD_CONT20]]754// CHECK2: omp_offload.cont20:755// CHECK2-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[A]], i64 0, i64 0756// CHECK2-NEXT: [[TMP79:%.*]] = load i32, ptr [[ARRAYIDX]], align 4757// CHECK2-NEXT: ret i32 [[TMP79]]758//759//760// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51761// 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]+]] {762// CHECK2-NEXT: entry:763// CHECK2-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8764// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8765// CHECK2-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 8766// CHECK2-NEXT: [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca i64, align 8767// CHECK2-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8768// CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB3]])769// CHECK2-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8770// CHECK2-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8771// CHECK2-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 8772// CHECK2-NEXT: store i64 [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 8773// CHECK2-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 8774// CHECK2-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 4775// CHECK2-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR2]], align 4776// CHECK2-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB3]], i32 [[TMP0]], i32 [[TMP2]], i32 [[TMP3]])777// CHECK2-NEXT: [[TMP4:%.*]] = load i32, ptr [[N_ADDR]], align 4778// CHECK2-NEXT: store i32 [[TMP4]], ptr [[N_CASTED]], align 4779// CHECK2-NEXT: [[TMP5:%.*]] = load i64, ptr [[N_CASTED]], align 8780// 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]])781// CHECK2-NEXT: ret void782//783//784// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined785// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR1]] {786// CHECK2-NEXT: entry:787// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8788// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8789// CHECK2-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8790// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8791// CHECK2-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4792// CHECK2-NEXT: [[TMP:%.*]] = alloca i32, align 4793// CHECK2-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4794// CHECK2-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4795// CHECK2-NEXT: [[I:%.*]] = alloca i32, align 4796// CHECK2-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4797// CHECK2-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4798// CHECK2-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4799// CHECK2-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4800// CHECK2-NEXT: [[I3:%.*]] = alloca i32, align 4801// CHECK2-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8802// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8803// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8804// CHECK2-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8805// CHECK2-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8806// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8807// CHECK2-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 4808// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 4809// CHECK2-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4810// CHECK2-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 0811// CHECK2-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1812// CHECK2-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1813// CHECK2-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4814// CHECK2-NEXT: store i32 0, ptr [[I]], align 4815// CHECK2-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4816// CHECK2-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]817// CHECK2-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]818// CHECK2: omp.precond.then:819// CHECK2-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4820// CHECK2-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4821// CHECK2-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_COMB_UB]], align 4822// CHECK2-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4823// CHECK2-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4824// CHECK2-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8825// CHECK2-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 4826// 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)827// CHECK2-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4828// CHECK2-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4829// CHECK2-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP7]], [[TMP8]]830// CHECK2-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]831// CHECK2: cond.true:832// CHECK2-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4833// CHECK2-NEXT: br label [[COND_END:%.*]]834// CHECK2: cond.false:835// CHECK2-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4836// CHECK2-NEXT: br label [[COND_END]]837// CHECK2: cond.end:838// CHECK2-NEXT: [[COND:%.*]] = phi i32 [ [[TMP9]], [[COND_TRUE]] ], [ [[TMP10]], [[COND_FALSE]] ]839// CHECK2-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4840// CHECK2-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4841// CHECK2-NEXT: store i32 [[TMP11]], ptr [[DOTOMP_IV]], align 4842// CHECK2-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]843// CHECK2: omp.inner.for.cond:844// CHECK2-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4845// CHECK2-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4846// CHECK2-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP12]], [[TMP13]]847// CHECK2-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]848// CHECK2: omp.inner.for.body:849// CHECK2-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4850// CHECK2-NEXT: [[TMP15:%.*]] = zext i32 [[TMP14]] to i64851// CHECK2-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4852// CHECK2-NEXT: [[TMP17:%.*]] = zext i32 [[TMP16]] to i64853// CHECK2-NEXT: [[TMP18:%.*]] = load i32, ptr [[N_ADDR]], align 4854// CHECK2-NEXT: store i32 [[TMP18]], ptr [[N_CASTED]], align 4855// CHECK2-NEXT: [[TMP19:%.*]] = load i64, ptr [[N_CASTED]], align 8856// 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]])857// CHECK2-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]858// CHECK2: omp.inner.for.inc:859// CHECK2-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4860// CHECK2-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4861// CHECK2-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP20]], [[TMP21]]862// CHECK2-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4863// CHECK2-NEXT: br label [[OMP_INNER_FOR_COND]]864// CHECK2: omp.inner.for.end:865// CHECK2-NEXT: br label [[OMP_LOOP_EXIT:%.*]]866// CHECK2: omp.loop.exit:867// CHECK2-NEXT: [[TMP22:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8868// CHECK2-NEXT: [[TMP23:%.*]] = load i32, ptr [[TMP22]], align 4869// CHECK2-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP23]])870// CHECK2-NEXT: br label [[OMP_PRECOND_END]]871// CHECK2: omp.precond.end:872// CHECK2-NEXT: ret void873//874//875// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined.omp_outlined876// 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]] {877// CHECK2-NEXT: entry:878// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8879// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8880// CHECK2-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8881// CHECK2-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8882// CHECK2-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8883// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8884// CHECK2-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4885// CHECK2-NEXT: [[TMP:%.*]] = alloca i32, align 4886// CHECK2-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4887// CHECK2-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4888// CHECK2-NEXT: [[I:%.*]] = alloca i32, align 4889// CHECK2-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4890// CHECK2-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4891// CHECK2-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4892// CHECK2-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4893// CHECK2-NEXT: [[I4:%.*]] = alloca i32, align 4894// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8895// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8896// CHECK2-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8897// CHECK2-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8898// CHECK2-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8899// CHECK2-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8900// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8901// CHECK2-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 4902// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 4903// CHECK2-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4904// CHECK2-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 0905// CHECK2-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1906// CHECK2-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1907// CHECK2-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4908// CHECK2-NEXT: store i32 0, ptr [[I]], align 4909// CHECK2-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4910// CHECK2-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]911// CHECK2-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]912// CHECK2: omp.precond.then:913// CHECK2-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4914// CHECK2-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4915// CHECK2-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_UB]], align 4916// CHECK2-NEXT: [[TMP5:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8917// CHECK2-NEXT: [[CONV:%.*]] = trunc i64 [[TMP5]] to i32918// CHECK2-NEXT: [[TMP6:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8919// CHECK2-NEXT: [[CONV3:%.*]] = trunc i64 [[TMP6]] to i32920// CHECK2-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 4921// CHECK2-NEXT: store i32 [[CONV3]], ptr [[DOTOMP_UB]], align 4922// CHECK2-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4923// CHECK2-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4924// CHECK2-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8925// CHECK2-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4926// 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)927// CHECK2-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4928// CHECK2-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4929// CHECK2-NEXT: [[CMP5:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]930// CHECK2-NEXT: br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]931// CHECK2: cond.true:932// CHECK2-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4933// CHECK2-NEXT: br label [[COND_END:%.*]]934// CHECK2: cond.false:935// CHECK2-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4936// CHECK2-NEXT: br label [[COND_END]]937// CHECK2: cond.end:938// CHECK2-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]939// CHECK2-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4940// CHECK2-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4941// CHECK2-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 4942// CHECK2-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]943// CHECK2: omp.inner.for.cond:944// CHECK2-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4945// CHECK2-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4946// CHECK2-NEXT: [[CMP6:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]947// CHECK2-NEXT: br i1 [[CMP6]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]948// CHECK2: omp.inner.for.body:949// CHECK2-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4950// CHECK2-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 1951// CHECK2-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]952// CHECK2-NEXT: store i32 [[ADD]], ptr [[I4]], align 4953// CHECK2-NEXT: [[TMP17:%.*]] = load i32, ptr [[I4]], align 4954// CHECK2-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP17]] to i64955// CHECK2-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]956// CHECK2-NEXT: store i32 0, ptr [[ARRAYIDX]], align 4957// CHECK2-NEXT: [[TMP18:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8958// CHECK2-NEXT: [[TMP19:%.*]] = load i32, ptr [[TMP18]], align 4959// CHECK2-NEXT: [[TMP20:%.*]] = call i32 @__kmpc_cancel(ptr @[[GLOB3]], i32 [[TMP19]], i32 2)960// CHECK2-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0961// CHECK2-NEXT: br i1 [[TMP21]], label [[DOTCANCEL_EXIT:%.*]], label [[DOTCANCEL_CONTINUE:%.*]]962// CHECK2: .cancel.exit:963// CHECK2-NEXT: br label [[CANCEL_EXIT:%.*]]964// CHECK2: .cancel.continue:965// CHECK2-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]966// CHECK2: omp.body.continue:967// CHECK2-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]968// CHECK2: omp.inner.for.inc:969// CHECK2-NEXT: [[TMP22:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4970// CHECK2-NEXT: [[ADD7:%.*]] = add nsw i32 [[TMP22]], 1971// CHECK2-NEXT: store i32 [[ADD7]], ptr [[DOTOMP_IV]], align 4972// CHECK2-NEXT: br label [[OMP_INNER_FOR_COND]]973// CHECK2: omp.inner.for.end:974// CHECK2-NEXT: br label [[OMP_LOOP_EXIT:%.*]]975// CHECK2: omp.loop.exit:976// CHECK2-NEXT: [[TMP23:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8977// CHECK2-NEXT: [[TMP24:%.*]] = load i32, ptr [[TMP23]], align 4978// CHECK2-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP24]])979// CHECK2-NEXT: br label [[OMP_PRECOND_END]]980// CHECK2: cancel.exit:981// CHECK2-NEXT: [[TMP25:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8982// CHECK2-NEXT: [[TMP26:%.*]] = load i32, ptr [[TMP25]], align 4983// CHECK2-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP26]])984// CHECK2-NEXT: br label [[CANCEL_CONT:%.*]]985// CHECK2: omp.precond.end:986// CHECK2-NEXT: br label [[CANCEL_CONT]]987// CHECK2: cancel.cont:988// CHECK2-NEXT: ret void989//990//991// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l58992// CHECK2-SAME: (i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], ptr noundef [[G:%.*]]) #[[ATTR1]] {993// CHECK2-NEXT: entry:994// CHECK2-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8995// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8996// CHECK2-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 8997// CHECK2-NEXT: [[N_CASTED:%.*]] = alloca i64, 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 [[N_CASTED]], align 41004// CHECK2-NEXT: [[TMP2:%.*]] = load i64, ptr [[N_CASTED]], align 81005// CHECK2-NEXT: [[TMP3:%.*]] = load ptr, ptr [[G_ADDR]], align 81006// 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_l58.omp_outlined, i64 [[TMP2]], ptr [[TMP0]], ptr [[TMP3]])1007// CHECK2-NEXT: ret void1008//1009//1010// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l58.omp_outlined1011// 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]] {1012// CHECK2-NEXT: entry:1013// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81014// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81015// CHECK2-NEXT: [[N_ADDR:%.*]] = alloca i64, align 81016// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81017// CHECK2-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 81018// CHECK2-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41019// CHECK2-NEXT: [[TMP:%.*]] = alloca i32, align 41020// CHECK2-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41021// CHECK2-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41022// CHECK2-NEXT: [[I:%.*]] = alloca i32, align 41023// CHECK2-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41024// CHECK2-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41025// CHECK2-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41026// CHECK2-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41027// CHECK2-NEXT: [[I3:%.*]] = alloca i32, align 41028// CHECK2-NEXT: [[N_CASTED:%.*]] = alloca i64, align 81029// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81030// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81031// CHECK2-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 81032// CHECK2-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81033// CHECK2-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 81034// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81035// CHECK2-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41036// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 41037// CHECK2-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41038// CHECK2-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 01039// CHECK2-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11040// CHECK2-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11041// CHECK2-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41042// CHECK2-NEXT: store i32 0, ptr [[I]], align 41043// CHECK2-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41044// CHECK2-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]1045// CHECK2-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1046// CHECK2: omp.precond.then:1047// CHECK2-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 41048// CHECK2-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41049// CHECK2-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_COMB_UB]], align 41050// CHECK2-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41051// CHECK2-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41052// CHECK2-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81053// CHECK2-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 41054// 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)1055// CHECK2-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41056// CHECK2-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41057// CHECK2-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP7]], [[TMP8]]1058// CHECK2-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1059// CHECK2: cond.true:1060// CHECK2-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41061// CHECK2-NEXT: br label [[COND_END:%.*]]1062// CHECK2: cond.false:1063// CHECK2-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41064// CHECK2-NEXT: br label [[COND_END]]1065// CHECK2: cond.end:1066// CHECK2-NEXT: [[COND:%.*]] = phi i32 [ [[TMP9]], [[COND_TRUE]] ], [ [[TMP10]], [[COND_FALSE]] ]1067// CHECK2-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41068// CHECK2-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41069// CHECK2-NEXT: store i32 [[TMP11]], ptr [[DOTOMP_IV]], align 41070// CHECK2-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1071// CHECK2: omp.inner.for.cond:1072// CHECK2-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41073// CHECK2-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41074// CHECK2-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP12]], [[TMP13]]1075// CHECK2-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1076// CHECK2: omp.inner.for.body:1077// CHECK2-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41078// CHECK2-NEXT: [[TMP15:%.*]] = zext i32 [[TMP14]] to i641079// CHECK2-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41080// CHECK2-NEXT: [[TMP17:%.*]] = zext i32 [[TMP16]] to i641081// CHECK2-NEXT: [[TMP18:%.*]] = load i32, ptr [[N_ADDR]], align 41082// CHECK2-NEXT: store i32 [[TMP18]], ptr [[N_CASTED]], align 41083// CHECK2-NEXT: [[TMP19:%.*]] = load i64, ptr [[N_CASTED]], align 81084// CHECK2-NEXT: [[TMP20:%.*]] = load ptr, ptr [[G_ADDR]], align 81085// 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_l58.omp_outlined.omp_outlined, i64 [[TMP15]], i64 [[TMP17]], i64 [[TMP19]], ptr [[TMP0]], ptr [[TMP20]])1086// CHECK2-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1087// CHECK2: omp.inner.for.inc:1088// CHECK2-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41089// CHECK2-NEXT: [[TMP22:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41090// CHECK2-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP21]], [[TMP22]]1091// CHECK2-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 41092// CHECK2-NEXT: br label [[OMP_INNER_FOR_COND]]1093// CHECK2: omp.inner.for.end:1094// CHECK2-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1095// CHECK2: omp.loop.exit:1096// CHECK2-NEXT: [[TMP23:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81097// CHECK2-NEXT: [[TMP24:%.*]] = load i32, ptr [[TMP23]], align 41098// CHECK2-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP24]])1099// CHECK2-NEXT: br label [[OMP_PRECOND_END]]1100// CHECK2: omp.precond.end:1101// CHECK2-NEXT: ret void1102//1103//1104// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l58.omp_outlined.omp_outlined1105// 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]] {1106// CHECK2-NEXT: entry:1107// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81108// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81109// CHECK2-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 81110// CHECK2-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 81111// CHECK2-NEXT: [[N_ADDR:%.*]] = alloca i64, align 81112// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81113// CHECK2-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 81114// CHECK2-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41115// CHECK2-NEXT: [[TMP:%.*]] = alloca i32, align 41116// CHECK2-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41117// CHECK2-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41118// CHECK2-NEXT: [[I:%.*]] = alloca i32, align 41119// CHECK2-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41120// CHECK2-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41121// CHECK2-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41122// CHECK2-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41123// CHECK2-NEXT: [[I4:%.*]] = alloca i32, align 41124// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81125// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81126// CHECK2-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 81127// CHECK2-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 81128// CHECK2-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 81129// CHECK2-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81130// CHECK2-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 81131// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81132// CHECK2-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41133// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 41134// CHECK2-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41135// CHECK2-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 01136// CHECK2-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11137// CHECK2-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11138// CHECK2-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41139// CHECK2-NEXT: store i32 0, ptr [[I]], align 41140// CHECK2-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41141// CHECK2-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]1142// CHECK2-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1143// CHECK2: omp.precond.then:1144// CHECK2-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41145// CHECK2-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41146// CHECK2-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_UB]], align 41147// CHECK2-NEXT: [[TMP5:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 81148// CHECK2-NEXT: [[CONV:%.*]] = trunc i64 [[TMP5]] to i321149// CHECK2-NEXT: [[TMP6:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 81150// CHECK2-NEXT: [[CONV3:%.*]] = trunc i64 [[TMP6]] to i321151// CHECK2-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 41152// CHECK2-NEXT: store i32 [[CONV3]], ptr [[DOTOMP_UB]], align 41153// CHECK2-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41154// CHECK2-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41155// CHECK2-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81156// CHECK2-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41157// 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)1158// CHECK2-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41159// CHECK2-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41160// CHECK2-NEXT: [[CMP5:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]1161// CHECK2-NEXT: br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1162// CHECK2: cond.true:1163// CHECK2-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41164// CHECK2-NEXT: br label [[COND_END:%.*]]1165// CHECK2: cond.false:1166// CHECK2-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41167// CHECK2-NEXT: br label [[COND_END]]1168// CHECK2: cond.end:1169// CHECK2-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]1170// CHECK2-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41171// CHECK2-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41172// CHECK2-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 41173// CHECK2-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1174// CHECK2: omp.inner.for.cond:1175// CHECK2-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41176// CHECK2-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41177// CHECK2-NEXT: [[CMP6:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]1178// CHECK2-NEXT: br i1 [[CMP6]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1179// CHECK2: omp.inner.for.body:1180// CHECK2-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41181// CHECK2-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 11182// CHECK2-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1183// CHECK2-NEXT: store i32 [[ADD]], ptr [[I4]], align 41184// CHECK2-NEXT: [[TMP17:%.*]] = load ptr, ptr [[G_ADDR]], align 81185// CHECK2-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP17]], i64 01186// CHECK2-NEXT: [[TMP18:%.*]] = load i32, ptr [[ARRAYIDX]], align 41187// CHECK2-NEXT: [[TMP19:%.*]] = load i32, ptr [[I4]], align 41188// CHECK2-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP19]] to i641189// CHECK2-NEXT: [[ARRAYIDX7:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]1190// CHECK2-NEXT: store i32 [[TMP18]], ptr [[ARRAYIDX7]], align 41191// CHECK2-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1192// CHECK2: omp.body.continue:1193// CHECK2-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1194// CHECK2: omp.inner.for.inc:1195// CHECK2-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41196// CHECK2-NEXT: [[ADD8:%.*]] = add nsw i32 [[TMP20]], 11197// CHECK2-NEXT: store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 41198// CHECK2-NEXT: br label [[OMP_INNER_FOR_COND]]1199// CHECK2: omp.inner.for.end:1200// CHECK2-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1201// CHECK2: omp.loop.exit:1202// CHECK2-NEXT: [[TMP21:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81203// CHECK2-NEXT: [[TMP22:%.*]] = load i32, ptr [[TMP21]], align 41204// CHECK2-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP22]])1205// CHECK2-NEXT: br label [[OMP_PRECOND_END]]1206// CHECK2: omp.precond.end:1207// CHECK2-NEXT: ret void1208//1209//1210// CHECK4-LABEL: define {{[^@]+}}@_Z16target_teams_funPi1211// CHECK4-SAME: (ptr noundef [[G:%.*]]) #[[ATTR0:[0-9]+]] {1212// CHECK4-NEXT: entry:1213// CHECK4-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 41214// CHECK4-NEXT: [[N:%.*]] = alloca i32, align 41215// CHECK4-NEXT: [[A:%.*]] = alloca [1000 x i32], align 41216// CHECK4-NEXT: [[TE:%.*]] = alloca i32, align 41217// CHECK4-NEXT: [[TH:%.*]] = alloca i32, align 41218// CHECK4-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41219// CHECK4-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41220// CHECK4-NEXT: [[N_CASTED:%.*]] = alloca i32, align 41221// CHECK4-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i32, align 41222// CHECK4-NEXT: [[DOTCAPTURE_EXPR__CASTED2:%.*]] = alloca i32, align 41223// CHECK4-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [4 x ptr], align 41224// CHECK4-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [4 x ptr], align 41225// CHECK4-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [4 x ptr], align 41226// CHECK4-NEXT: [[TMP:%.*]] = alloca i32, align 41227// CHECK4-NEXT: [[DOTCAPTURE_EXPR_3:%.*]] = alloca i32, align 41228// CHECK4-NEXT: [[DOTCAPTURE_EXPR_4:%.*]] = alloca i32, align 41229// CHECK4-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81230// CHECK4-NEXT: [[N_CASTED7:%.*]] = alloca i32, align 41231// CHECK4-NEXT: [[DOTOFFLOAD_BASEPTRS8:%.*]] = alloca [3 x ptr], align 41232// CHECK4-NEXT: [[DOTOFFLOAD_PTRS9:%.*]] = alloca [3 x ptr], align 41233// CHECK4-NEXT: [[DOTOFFLOAD_MAPPERS10:%.*]] = alloca [3 x ptr], align 41234// CHECK4-NEXT: [[_TMP11:%.*]] = alloca i32, align 41235// CHECK4-NEXT: [[DOTCAPTURE_EXPR_12:%.*]] = alloca i32, align 41236// CHECK4-NEXT: [[DOTCAPTURE_EXPR_13:%.*]] = alloca i32, align 41237// CHECK4-NEXT: [[KERNEL_ARGS18:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 81238// CHECK4-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 41239// CHECK4-NEXT: store i32 1000, ptr [[N]], align 41240// CHECK4-NEXT: [[TMP0:%.*]] = load i32, ptr [[N]], align 41241// CHECK4-NEXT: [[DIV:%.*]] = sdiv i32 [[TMP0]], 1281242// CHECK4-NEXT: store i32 [[DIV]], ptr [[TE]], align 41243// CHECK4-NEXT: store i32 128, ptr [[TH]], align 41244// CHECK4-NEXT: [[TMP1:%.*]] = load i32, ptr [[TE]], align 41245// CHECK4-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 41246// CHECK4-NEXT: [[TMP2:%.*]] = load i32, ptr [[TH]], align 41247// CHECK4-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_1]], align 41248// CHECK4-NEXT: [[TMP3:%.*]] = load i32, ptr [[N]], align 41249// CHECK4-NEXT: store i32 [[TMP3]], ptr [[N_CASTED]], align 41250// CHECK4-NEXT: [[TMP4:%.*]] = load i32, ptr [[N_CASTED]], align 41251// CHECK4-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41252// CHECK4-NEXT: store i32 [[TMP5]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 41253// CHECK4-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED]], align 41254// CHECK4-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41255// CHECK4-NEXT: store i32 [[TMP7]], ptr [[DOTCAPTURE_EXPR__CASTED2]], align 41256// CHECK4-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED2]], align 41257// CHECK4-NEXT: [[TMP9:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01258// CHECK4-NEXT: store i32 [[TMP4]], ptr [[TMP9]], align 41259// CHECK4-NEXT: [[TMP10:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01260// CHECK4-NEXT: store i32 [[TMP4]], ptr [[TMP10]], align 41261// CHECK4-NEXT: [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01262// CHECK4-NEXT: store ptr null, ptr [[TMP11]], align 41263// CHECK4-NEXT: [[TMP12:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11264// CHECK4-NEXT: store ptr [[A]], ptr [[TMP12]], align 41265// CHECK4-NEXT: [[TMP13:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11266// CHECK4-NEXT: store ptr [[A]], ptr [[TMP13]], align 41267// CHECK4-NEXT: [[TMP14:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 11268// CHECK4-NEXT: store ptr null, ptr [[TMP14]], align 41269// CHECK4-NEXT: [[TMP15:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21270// CHECK4-NEXT: store i32 [[TMP6]], ptr [[TMP15]], align 41271// CHECK4-NEXT: [[TMP16:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21272// CHECK4-NEXT: store i32 [[TMP6]], ptr [[TMP16]], align 41273// CHECK4-NEXT: [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 21274// CHECK4-NEXT: store ptr null, ptr [[TMP17]], align 41275// CHECK4-NEXT: [[TMP18:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 31276// CHECK4-NEXT: store i32 [[TMP8]], ptr [[TMP18]], align 41277// CHECK4-NEXT: [[TMP19:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 31278// CHECK4-NEXT: store i32 [[TMP8]], ptr [[TMP19]], align 41279// CHECK4-NEXT: [[TMP20:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 31280// CHECK4-NEXT: store ptr null, ptr [[TMP20]], align 41281// CHECK4-NEXT: [[TMP21:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01282// CHECK4-NEXT: [[TMP22:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01283// CHECK4-NEXT: [[TMP23:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41284// CHECK4-NEXT: [[TMP24:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41285// CHECK4-NEXT: [[TMP25:%.*]] = load i32, ptr [[N]], align 41286// CHECK4-NEXT: store i32 [[TMP25]], ptr [[DOTCAPTURE_EXPR_3]], align 41287// CHECK4-NEXT: [[TMP26:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_3]], align 41288// CHECK4-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP26]], 01289// CHECK4-NEXT: [[DIV5:%.*]] = sdiv i32 [[SUB]], 11290// CHECK4-NEXT: [[SUB6:%.*]] = sub nsw i32 [[DIV5]], 11291// CHECK4-NEXT: store i32 [[SUB6]], ptr [[DOTCAPTURE_EXPR_4]], align 41292// CHECK4-NEXT: [[TMP27:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 41293// CHECK4-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP27]], 11294// CHECK4-NEXT: [[TMP28:%.*]] = zext i32 [[ADD]] to i641295// CHECK4-NEXT: [[TMP29:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP23]], 01296// CHECK4-NEXT: [[TMP30:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP24]], 01297// CHECK4-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01298// CHECK4-NEXT: store i32 3, ptr [[TMP31]], align 41299// CHECK4-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11300// CHECK4-NEXT: store i32 4, ptr [[TMP32]], align 41301// CHECK4-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21302// CHECK4-NEXT: store ptr [[TMP21]], ptr [[TMP33]], align 41303// CHECK4-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31304// CHECK4-NEXT: store ptr [[TMP22]], ptr [[TMP34]], align 41305// CHECK4-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41306// CHECK4-NEXT: store ptr @.offload_sizes, ptr [[TMP35]], align 41307// CHECK4-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51308// CHECK4-NEXT: store ptr @.offload_maptypes, ptr [[TMP36]], align 41309// CHECK4-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61310// CHECK4-NEXT: store ptr null, ptr [[TMP37]], align 41311// CHECK4-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71312// CHECK4-NEXT: store ptr null, ptr [[TMP38]], align 41313// CHECK4-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81314// CHECK4-NEXT: store i64 [[TMP28]], ptr [[TMP39]], align 81315// CHECK4-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91316// CHECK4-NEXT: store i64 0, ptr [[TMP40]], align 81317// CHECK4-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101318// CHECK4-NEXT: store [3 x i32] [[TMP29]], ptr [[TMP41]], align 41319// CHECK4-NEXT: [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111320// CHECK4-NEXT: store [3 x i32] [[TMP30]], ptr [[TMP42]], align 41321// CHECK4-NEXT: [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121322// CHECK4-NEXT: store i32 0, ptr [[TMP43]], align 41323// CHECK4-NEXT: [[TMP44:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3:[0-9]+]], i64 -1, i32 [[TMP23]], i32 [[TMP24]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.region_id, ptr [[KERNEL_ARGS]])1324// CHECK4-NEXT: [[TMP45:%.*]] = icmp ne i32 [[TMP44]], 01325// CHECK4-NEXT: br i1 [[TMP45]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1326// CHECK4: omp_offload.failed:1327// 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]+]]1328// CHECK4-NEXT: br label [[OMP_OFFLOAD_CONT]]1329// CHECK4: omp_offload.cont:1330// CHECK4-NEXT: [[TMP46:%.*]] = load i32, ptr [[N]], align 41331// CHECK4-NEXT: store i32 [[TMP46]], ptr [[N_CASTED7]], align 41332// CHECK4-NEXT: [[TMP47:%.*]] = load i32, ptr [[N_CASTED7]], align 41333// CHECK4-NEXT: [[TMP48:%.*]] = load ptr, ptr [[G_ADDR]], align 41334// CHECK4-NEXT: [[TMP49:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 01335// CHECK4-NEXT: store i32 [[TMP47]], ptr [[TMP49]], align 41336// CHECK4-NEXT: [[TMP50:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 01337// CHECK4-NEXT: store i32 [[TMP47]], ptr [[TMP50]], align 41338// CHECK4-NEXT: [[TMP51:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS10]], i32 0, i32 01339// CHECK4-NEXT: store ptr null, ptr [[TMP51]], align 41340// CHECK4-NEXT: [[TMP52:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 11341// CHECK4-NEXT: store ptr [[A]], ptr [[TMP52]], align 41342// CHECK4-NEXT: [[TMP53:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 11343// CHECK4-NEXT: store ptr [[A]], ptr [[TMP53]], align 41344// CHECK4-NEXT: [[TMP54:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS10]], i32 0, i32 11345// CHECK4-NEXT: store ptr null, ptr [[TMP54]], align 41346// CHECK4-NEXT: [[TMP55:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 21347// CHECK4-NEXT: store ptr [[TMP48]], ptr [[TMP55]], align 41348// CHECK4-NEXT: [[TMP56:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 21349// CHECK4-NEXT: store ptr [[TMP48]], ptr [[TMP56]], align 41350// CHECK4-NEXT: [[TMP57:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS10]], i32 0, i32 21351// CHECK4-NEXT: store ptr null, ptr [[TMP57]], align 41352// CHECK4-NEXT: [[TMP58:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 01353// CHECK4-NEXT: [[TMP59:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 01354// CHECK4-NEXT: [[TMP60:%.*]] = load i32, ptr [[N]], align 41355// CHECK4-NEXT: store i32 [[TMP60]], ptr [[DOTCAPTURE_EXPR_12]], align 41356// CHECK4-NEXT: [[TMP61:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_12]], align 41357// CHECK4-NEXT: [[SUB14:%.*]] = sub nsw i32 [[TMP61]], 01358// CHECK4-NEXT: [[DIV15:%.*]] = sdiv i32 [[SUB14]], 11359// CHECK4-NEXT: [[SUB16:%.*]] = sub nsw i32 [[DIV15]], 11360// CHECK4-NEXT: store i32 [[SUB16]], ptr [[DOTCAPTURE_EXPR_13]], align 41361// CHECK4-NEXT: [[TMP62:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_13]], align 41362// CHECK4-NEXT: [[ADD17:%.*]] = add nsw i32 [[TMP62]], 11363// CHECK4-NEXT: [[TMP63:%.*]] = zext i32 [[ADD17]] to i641364// CHECK4-NEXT: [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 01365// CHECK4-NEXT: store i32 3, ptr [[TMP64]], align 41366// CHECK4-NEXT: [[TMP65:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 11367// CHECK4-NEXT: store i32 3, ptr [[TMP65]], align 41368// CHECK4-NEXT: [[TMP66:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 21369// CHECK4-NEXT: store ptr [[TMP58]], ptr [[TMP66]], align 41370// CHECK4-NEXT: [[TMP67:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 31371// CHECK4-NEXT: store ptr [[TMP59]], ptr [[TMP67]], align 41372// CHECK4-NEXT: [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 41373// CHECK4-NEXT: store ptr @.offload_sizes.1, ptr [[TMP68]], align 41374// CHECK4-NEXT: [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 51375// CHECK4-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP69]], align 41376// CHECK4-NEXT: [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 61377// CHECK4-NEXT: store ptr null, ptr [[TMP70]], align 41378// CHECK4-NEXT: [[TMP71:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 71379// CHECK4-NEXT: store ptr null, ptr [[TMP71]], align 41380// CHECK4-NEXT: [[TMP72:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 81381// CHECK4-NEXT: store i64 [[TMP63]], ptr [[TMP72]], align 81382// CHECK4-NEXT: [[TMP73:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 91383// CHECK4-NEXT: store i64 0, ptr [[TMP73]], align 81384// CHECK4-NEXT: [[TMP74:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 101385// CHECK4-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP74]], align 41386// CHECK4-NEXT: [[TMP75:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 111387// CHECK4-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP75]], align 41388// CHECK4-NEXT: [[TMP76:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 121389// CHECK4-NEXT: store i32 0, ptr [[TMP76]], align 41390// CHECK4-NEXT: [[TMP77:%.*]] = 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_l58.region_id, ptr [[KERNEL_ARGS18]])1391// CHECK4-NEXT: [[TMP78:%.*]] = icmp ne i32 [[TMP77]], 01392// CHECK4-NEXT: br i1 [[TMP78]], label [[OMP_OFFLOAD_FAILED19:%.*]], label [[OMP_OFFLOAD_CONT20:%.*]]1393// CHECK4: omp_offload.failed19:1394// CHECK4-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l58(i32 [[TMP47]], ptr [[A]], ptr [[TMP48]]) #[[ATTR2]]1395// CHECK4-NEXT: br label [[OMP_OFFLOAD_CONT20]]1396// CHECK4: omp_offload.cont20:1397// CHECK4-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[A]], i32 0, i32 01398// CHECK4-NEXT: [[TMP79:%.*]] = load i32, ptr [[ARRAYIDX]], align 41399// CHECK4-NEXT: ret i32 [[TMP79]]1400//1401//1402// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l511403// 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]+]] {1404// CHECK4-NEXT: entry:1405// CHECK4-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41406// CHECK4-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41407// CHECK4-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 41408// CHECK4-NEXT: [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca i32, align 41409// CHECK4-NEXT: [[N_CASTED:%.*]] = alloca i32, align 41410// CHECK4-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB3]])1411// CHECK4-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41412// CHECK4-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41413// CHECK4-NEXT: store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 41414// CHECK4-NEXT: store i32 [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 41415// CHECK4-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 41416// CHECK4-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 41417// CHECK4-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR2]], align 41418// CHECK4-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB3]], i32 [[TMP0]], i32 [[TMP2]], i32 [[TMP3]])1419// CHECK4-NEXT: [[TMP4:%.*]] = load i32, ptr [[N_ADDR]], align 41420// CHECK4-NEXT: store i32 [[TMP4]], ptr [[N_CASTED]], align 41421// CHECK4-NEXT: [[TMP5:%.*]] = load i32, ptr [[N_CASTED]], align 41422// 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]])1423// CHECK4-NEXT: ret void1424//1425//1426// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined1427// CHECK4-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR1]] {1428// CHECK4-NEXT: entry:1429// CHECK4-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41430// CHECK4-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41431// CHECK4-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41432// CHECK4-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41433// CHECK4-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41434// CHECK4-NEXT: [[TMP:%.*]] = alloca i32, align 41435// CHECK4-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41436// CHECK4-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41437// CHECK4-NEXT: [[I:%.*]] = alloca i32, align 41438// CHECK4-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41439// CHECK4-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41440// CHECK4-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41441// CHECK4-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41442// CHECK4-NEXT: [[I3:%.*]] = alloca i32, align 41443// CHECK4-NEXT: [[N_CASTED:%.*]] = alloca i32, align 41444// CHECK4-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41445// CHECK4-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41446// CHECK4-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41447// CHECK4-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41448// CHECK4-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41449// CHECK4-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41450// CHECK4-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 41451// CHECK4-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41452// CHECK4-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 01453// CHECK4-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11454// CHECK4-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11455// CHECK4-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41456// CHECK4-NEXT: store i32 0, ptr [[I]], align 41457// CHECK4-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41458// CHECK4-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]1459// CHECK4-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1460// CHECK4: omp.precond.then:1461// CHECK4-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 41462// CHECK4-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41463// CHECK4-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_COMB_UB]], align 41464// CHECK4-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41465// CHECK4-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41466// CHECK4-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41467// CHECK4-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 41468// 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)1469// CHECK4-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41470// CHECK4-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41471// CHECK4-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP7]], [[TMP8]]1472// CHECK4-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1473// CHECK4: cond.true:1474// CHECK4-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41475// CHECK4-NEXT: br label [[COND_END:%.*]]1476// CHECK4: cond.false:1477// CHECK4-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41478// CHECK4-NEXT: br label [[COND_END]]1479// CHECK4: cond.end:1480// CHECK4-NEXT: [[COND:%.*]] = phi i32 [ [[TMP9]], [[COND_TRUE]] ], [ [[TMP10]], [[COND_FALSE]] ]1481// CHECK4-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41482// CHECK4-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41483// CHECK4-NEXT: store i32 [[TMP11]], ptr [[DOTOMP_IV]], align 41484// CHECK4-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1485// CHECK4: omp.inner.for.cond:1486// CHECK4-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41487// CHECK4-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41488// CHECK4-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP12]], [[TMP13]]1489// CHECK4-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1490// CHECK4: omp.inner.for.body:1491// CHECK4-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41492// CHECK4-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41493// CHECK4-NEXT: [[TMP16:%.*]] = load i32, ptr [[N_ADDR]], align 41494// CHECK4-NEXT: store i32 [[TMP16]], ptr [[N_CASTED]], align 41495// CHECK4-NEXT: [[TMP17:%.*]] = load i32, ptr [[N_CASTED]], align 41496// 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]])1497// CHECK4-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1498// CHECK4: omp.inner.for.inc:1499// CHECK4-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41500// CHECK4-NEXT: [[TMP19:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41501// CHECK4-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP18]], [[TMP19]]1502// CHECK4-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 41503// CHECK4-NEXT: br label [[OMP_INNER_FOR_COND]]1504// CHECK4: omp.inner.for.end:1505// CHECK4-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1506// CHECK4: omp.loop.exit:1507// CHECK4-NEXT: [[TMP20:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41508// CHECK4-NEXT: [[TMP21:%.*]] = load i32, ptr [[TMP20]], align 41509// CHECK4-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP21]])1510// CHECK4-NEXT: br label [[OMP_PRECOND_END]]1511// CHECK4: omp.precond.end:1512// CHECK4-NEXT: ret void1513//1514//1515// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined.omp_outlined1516// 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]] {1517// CHECK4-NEXT: entry:1518// CHECK4-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41519// CHECK4-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41520// CHECK4-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i32, align 41521// CHECK4-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i32, align 41522// CHECK4-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41523// CHECK4-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41524// CHECK4-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41525// CHECK4-NEXT: [[TMP:%.*]] = alloca i32, align 41526// CHECK4-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41527// CHECK4-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41528// CHECK4-NEXT: [[I:%.*]] = alloca i32, align 41529// CHECK4-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41530// CHECK4-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41531// CHECK4-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41532// CHECK4-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41533// CHECK4-NEXT: [[I3:%.*]] = alloca i32, align 41534// CHECK4-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41535// CHECK4-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41536// CHECK4-NEXT: store i32 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 41537// CHECK4-NEXT: store i32 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 41538// CHECK4-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41539// CHECK4-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41540// CHECK4-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41541// CHECK4-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41542// CHECK4-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 41543// CHECK4-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41544// CHECK4-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 01545// CHECK4-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11546// CHECK4-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11547// CHECK4-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41548// CHECK4-NEXT: store i32 0, ptr [[I]], align 41549// CHECK4-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41550// CHECK4-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]1551// CHECK4-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1552// CHECK4: omp.precond.then:1553// CHECK4-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41554// CHECK4-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41555// CHECK4-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_UB]], align 41556// CHECK4-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTPREVIOUS_LB__ADDR]], align 41557// CHECK4-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTPREVIOUS_UB__ADDR]], align 41558// CHECK4-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_LB]], align 41559// CHECK4-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_UB]], align 41560// CHECK4-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41561// CHECK4-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41562// CHECK4-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41563// CHECK4-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41564// 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)1565// CHECK4-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41566// CHECK4-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41567// CHECK4-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]1568// CHECK4-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1569// CHECK4: cond.true:1570// CHECK4-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41571// CHECK4-NEXT: br label [[COND_END:%.*]]1572// CHECK4: cond.false:1573// CHECK4-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41574// CHECK4-NEXT: br label [[COND_END]]1575// CHECK4: cond.end:1576// CHECK4-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]1577// CHECK4-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41578// CHECK4-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41579// CHECK4-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 41580// CHECK4-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1581// CHECK4: omp.inner.for.cond:1582// CHECK4-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41583// CHECK4-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41584// CHECK4-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]1585// CHECK4-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1586// CHECK4: omp.inner.for.body:1587// CHECK4-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41588// CHECK4-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 11589// CHECK4-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1590// CHECK4-NEXT: store i32 [[ADD]], ptr [[I3]], align 41591// CHECK4-NEXT: [[TMP17:%.*]] = load i32, ptr [[I3]], align 41592// CHECK4-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i32 0, i32 [[TMP17]]1593// CHECK4-NEXT: store i32 0, ptr [[ARRAYIDX]], align 41594// CHECK4-NEXT: [[TMP18:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41595// CHECK4-NEXT: [[TMP19:%.*]] = load i32, ptr [[TMP18]], align 41596// CHECK4-NEXT: [[TMP20:%.*]] = call i32 @__kmpc_cancel(ptr @[[GLOB3]], i32 [[TMP19]], i32 2)1597// CHECK4-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 01598// CHECK4-NEXT: br i1 [[TMP21]], label [[DOTCANCEL_EXIT:%.*]], label [[DOTCANCEL_CONTINUE:%.*]]1599// CHECK4: .cancel.exit:1600// CHECK4-NEXT: br label [[CANCEL_EXIT:%.*]]1601// CHECK4: .cancel.continue:1602// CHECK4-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1603// CHECK4: omp.body.continue:1604// CHECK4-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1605// CHECK4: omp.inner.for.inc:1606// CHECK4-NEXT: [[TMP22:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41607// CHECK4-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP22]], 11608// CHECK4-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 41609// CHECK4-NEXT: br label [[OMP_INNER_FOR_COND]]1610// CHECK4: omp.inner.for.end:1611// CHECK4-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1612// CHECK4: omp.loop.exit:1613// CHECK4-NEXT: [[TMP23:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41614// CHECK4-NEXT: [[TMP24:%.*]] = load i32, ptr [[TMP23]], align 41615// CHECK4-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP24]])1616// CHECK4-NEXT: br label [[OMP_PRECOND_END]]1617// CHECK4: cancel.exit:1618// CHECK4-NEXT: [[TMP25:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41619// CHECK4-NEXT: [[TMP26:%.*]] = load i32, ptr [[TMP25]], align 41620// CHECK4-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP26]])1621// CHECK4-NEXT: br label [[CANCEL_CONT:%.*]]1622// CHECK4: omp.precond.end:1623// CHECK4-NEXT: br label [[CANCEL_CONT]]1624// CHECK4: cancel.cont:1625// CHECK4-NEXT: ret void1626//1627//1628// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l581629// CHECK4-SAME: (i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], ptr noundef [[G:%.*]]) #[[ATTR1]] {1630// CHECK4-NEXT: entry:1631// CHECK4-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41632// CHECK4-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41633// CHECK4-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 41634// CHECK4-NEXT: [[N_CASTED:%.*]] = alloca i32, align 41635// CHECK4-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41636// CHECK4-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41637// CHECK4-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 41638// CHECK4-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41639// CHECK4-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41640// CHECK4-NEXT: store i32 [[TMP1]], ptr [[N_CASTED]], align 41641// CHECK4-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_CASTED]], align 41642// CHECK4-NEXT: [[TMP3:%.*]] = load ptr, ptr [[G_ADDR]], align 41643// 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_l58.omp_outlined, i32 [[TMP2]], ptr [[TMP0]], ptr [[TMP3]])1644// CHECK4-NEXT: ret void1645//1646//1647// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l58.omp_outlined1648// 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]] {1649// CHECK4-NEXT: entry:1650// CHECK4-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41651// CHECK4-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41652// CHECK4-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41653// CHECK4-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41654// CHECK4-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 41655// CHECK4-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41656// CHECK4-NEXT: [[TMP:%.*]] = alloca i32, align 41657// CHECK4-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41658// CHECK4-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41659// CHECK4-NEXT: [[I:%.*]] = alloca i32, align 41660// CHECK4-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41661// CHECK4-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41662// CHECK4-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41663// CHECK4-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41664// CHECK4-NEXT: [[I3:%.*]] = alloca i32, align 41665// CHECK4-NEXT: [[N_CASTED:%.*]] = alloca i32, align 41666// CHECK4-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41667// CHECK4-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41668// CHECK4-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41669// CHECK4-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41670// CHECK4-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 41671// CHECK4-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41672// CHECK4-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41673// CHECK4-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 41674// CHECK4-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41675// CHECK4-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 01676// CHECK4-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11677// CHECK4-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11678// CHECK4-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41679// CHECK4-NEXT: store i32 0, ptr [[I]], align 41680// CHECK4-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41681// CHECK4-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]1682// CHECK4-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1683// CHECK4: omp.precond.then:1684// CHECK4-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 41685// CHECK4-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41686// CHECK4-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_COMB_UB]], align 41687// CHECK4-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41688// CHECK4-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41689// CHECK4-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41690// CHECK4-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 41691// 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)1692// CHECK4-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41693// CHECK4-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41694// CHECK4-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP7]], [[TMP8]]1695// CHECK4-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1696// CHECK4: cond.true:1697// CHECK4-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41698// CHECK4-NEXT: br label [[COND_END:%.*]]1699// CHECK4: cond.false:1700// CHECK4-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41701// CHECK4-NEXT: br label [[COND_END]]1702// CHECK4: cond.end:1703// CHECK4-NEXT: [[COND:%.*]] = phi i32 [ [[TMP9]], [[COND_TRUE]] ], [ [[TMP10]], [[COND_FALSE]] ]1704// CHECK4-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41705// CHECK4-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41706// CHECK4-NEXT: store i32 [[TMP11]], ptr [[DOTOMP_IV]], align 41707// CHECK4-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1708// CHECK4: omp.inner.for.cond:1709// CHECK4-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41710// CHECK4-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41711// CHECK4-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP12]], [[TMP13]]1712// CHECK4-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1713// CHECK4: omp.inner.for.body:1714// CHECK4-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41715// CHECK4-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41716// CHECK4-NEXT: [[TMP16:%.*]] = load i32, ptr [[N_ADDR]], align 41717// CHECK4-NEXT: store i32 [[TMP16]], ptr [[N_CASTED]], align 41718// CHECK4-NEXT: [[TMP17:%.*]] = load i32, ptr [[N_CASTED]], align 41719// CHECK4-NEXT: [[TMP18:%.*]] = load ptr, ptr [[G_ADDR]], align 41720// 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_l58.omp_outlined.omp_outlined, i32 [[TMP14]], i32 [[TMP15]], i32 [[TMP17]], ptr [[TMP0]], ptr [[TMP18]])1721// CHECK4-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1722// CHECK4: omp.inner.for.inc:1723// CHECK4-NEXT: [[TMP19:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41724// CHECK4-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41725// CHECK4-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP19]], [[TMP20]]1726// CHECK4-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 41727// CHECK4-NEXT: br label [[OMP_INNER_FOR_COND]]1728// CHECK4: omp.inner.for.end:1729// CHECK4-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1730// CHECK4: omp.loop.exit:1731// CHECK4-NEXT: [[TMP21:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41732// CHECK4-NEXT: [[TMP22:%.*]] = load i32, ptr [[TMP21]], align 41733// CHECK4-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP22]])1734// CHECK4-NEXT: br label [[OMP_PRECOND_END]]1735// CHECK4: omp.precond.end:1736// CHECK4-NEXT: ret void1737//1738//1739// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l58.omp_outlined.omp_outlined1740// 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]] {1741// CHECK4-NEXT: entry:1742// CHECK4-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41743// CHECK4-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41744// CHECK4-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i32, align 41745// CHECK4-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i32, align 41746// CHECK4-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41747// CHECK4-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41748// CHECK4-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 41749// CHECK4-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41750// CHECK4-NEXT: [[TMP:%.*]] = alloca i32, align 41751// CHECK4-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41752// CHECK4-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41753// CHECK4-NEXT: [[I:%.*]] = alloca i32, align 41754// CHECK4-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41755// CHECK4-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41756// CHECK4-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41757// CHECK4-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41758// CHECK4-NEXT: [[I3:%.*]] = alloca i32, align 41759// CHECK4-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41760// CHECK4-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41761// CHECK4-NEXT: store i32 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 41762// CHECK4-NEXT: store i32 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 41763// CHECK4-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41764// CHECK4-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41765// CHECK4-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 41766// CHECK4-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41767// CHECK4-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41768// CHECK4-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 41769// CHECK4-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41770// CHECK4-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 01771// CHECK4-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11772// CHECK4-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11773// CHECK4-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41774// CHECK4-NEXT: store i32 0, ptr [[I]], align 41775// CHECK4-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41776// CHECK4-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]1777// CHECK4-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1778// CHECK4: omp.precond.then:1779// CHECK4-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41780// CHECK4-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41781// CHECK4-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_UB]], align 41782// CHECK4-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTPREVIOUS_LB__ADDR]], align 41783// CHECK4-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTPREVIOUS_UB__ADDR]], align 41784// CHECK4-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_LB]], align 41785// CHECK4-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_UB]], align 41786// CHECK4-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41787// CHECK4-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41788// CHECK4-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41789// CHECK4-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41790// 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)1791// CHECK4-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41792// CHECK4-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41793// CHECK4-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]1794// CHECK4-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1795// CHECK4: cond.true:1796// CHECK4-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41797// CHECK4-NEXT: br label [[COND_END:%.*]]1798// CHECK4: cond.false:1799// CHECK4-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41800// CHECK4-NEXT: br label [[COND_END]]1801// CHECK4: cond.end:1802// CHECK4-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]1803// CHECK4-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41804// CHECK4-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41805// CHECK4-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 41806// CHECK4-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1807// CHECK4: omp.inner.for.cond:1808// CHECK4-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41809// CHECK4-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41810// CHECK4-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]1811// CHECK4-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1812// CHECK4: omp.inner.for.body:1813// CHECK4-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41814// CHECK4-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 11815// CHECK4-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1816// CHECK4-NEXT: store i32 [[ADD]], ptr [[I3]], align 41817// CHECK4-NEXT: [[TMP17:%.*]] = load ptr, ptr [[G_ADDR]], align 41818// CHECK4-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP17]], i32 01819// CHECK4-NEXT: [[TMP18:%.*]] = load i32, ptr [[ARRAYIDX]], align 41820// CHECK4-NEXT: [[TMP19:%.*]] = load i32, ptr [[I3]], align 41821// CHECK4-NEXT: [[ARRAYIDX6:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i32 0, i32 [[TMP19]]1822// CHECK4-NEXT: store i32 [[TMP18]], ptr [[ARRAYIDX6]], align 41823// CHECK4-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1824// CHECK4: omp.body.continue:1825// CHECK4-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1826// CHECK4: omp.inner.for.inc:1827// CHECK4-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41828// CHECK4-NEXT: [[ADD7:%.*]] = add nsw i32 [[TMP20]], 11829// CHECK4-NEXT: store i32 [[ADD7]], ptr [[DOTOMP_IV]], align 41830// CHECK4-NEXT: br label [[OMP_INNER_FOR_COND]]1831// CHECK4: omp.inner.for.end:1832// CHECK4-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1833// CHECK4: omp.loop.exit:1834// CHECK4-NEXT: [[TMP21:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41835// CHECK4-NEXT: [[TMP22:%.*]] = load i32, ptr [[TMP21]], align 41836// CHECK4-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP22]])1837// CHECK4-NEXT: br label [[OMP_PRECOND_END]]1838// CHECK4: omp.precond.end:1839// CHECK4-NEXT: ret void1840//1841//1842// CHECK10-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l511843// 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]+]] {1844// CHECK10-NEXT: entry:1845// CHECK10-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81846// CHECK10-NEXT: [[N_ADDR:%.*]] = alloca i64, align 81847// CHECK10-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81848// CHECK10-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 81849// CHECK10-NEXT: [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca i64, align 81850// CHECK10-NEXT: [[N_CASTED:%.*]] = alloca i64, align 81851// CHECK10-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB3:[0-9]+]])1852// CHECK10-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81853// CHECK10-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 81854// CHECK10-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81855// CHECK10-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 81856// CHECK10-NEXT: store i64 [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 81857// CHECK10-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 81858// CHECK10-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 41859// CHECK10-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR2]], align 41860// CHECK10-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB3]], i32 [[TMP0]], i32 [[TMP2]], i32 [[TMP3]])1861// CHECK10-NEXT: [[TMP4:%.*]] = load i32, ptr [[N_ADDR]], align 41862// CHECK10-NEXT: store i32 [[TMP4]], ptr [[N_CASTED]], align 41863// CHECK10-NEXT: [[TMP5:%.*]] = load i64, ptr [[N_CASTED]], align 81864// 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]])1865// CHECK10-NEXT: ret void1866//1867//1868// CHECK10-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined1869// CHECK10-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR0]] {1870// CHECK10-NEXT: entry:1871// CHECK10-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81872// CHECK10-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81873// CHECK10-NEXT: [[N_ADDR:%.*]] = alloca i64, align 81874// CHECK10-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81875// CHECK10-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41876// CHECK10-NEXT: [[TMP:%.*]] = alloca i32, align 41877// CHECK10-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41878// CHECK10-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41879// CHECK10-NEXT: [[I:%.*]] = alloca i32, align 41880// CHECK10-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41881// CHECK10-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41882// CHECK10-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41883// CHECK10-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41884// CHECK10-NEXT: [[I3:%.*]] = alloca i32, align 41885// CHECK10-NEXT: [[N_CASTED:%.*]] = alloca i64, align 81886// CHECK10-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81887// CHECK10-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81888// CHECK10-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 81889// CHECK10-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81890// CHECK10-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81891// CHECK10-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41892// CHECK10-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 41893// CHECK10-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41894// CHECK10-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 01895// CHECK10-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11896// CHECK10-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11897// CHECK10-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41898// CHECK10-NEXT: store i32 0, ptr [[I]], align 41899// CHECK10-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41900// CHECK10-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]1901// CHECK10-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1902// CHECK10: omp.precond.then:1903// CHECK10-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 41904// CHECK10-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41905// CHECK10-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_COMB_UB]], align 41906// CHECK10-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41907// CHECK10-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41908// CHECK10-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81909// CHECK10-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 41910// 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)1911// CHECK10-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41912// CHECK10-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41913// CHECK10-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP7]], [[TMP8]]1914// CHECK10-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1915// CHECK10: cond.true:1916// CHECK10-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41917// CHECK10-NEXT: br label [[COND_END:%.*]]1918// CHECK10: cond.false:1919// CHECK10-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41920// CHECK10-NEXT: br label [[COND_END]]1921// CHECK10: cond.end:1922// CHECK10-NEXT: [[COND:%.*]] = phi i32 [ [[TMP9]], [[COND_TRUE]] ], [ [[TMP10]], [[COND_FALSE]] ]1923// CHECK10-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41924// CHECK10-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41925// CHECK10-NEXT: store i32 [[TMP11]], ptr [[DOTOMP_IV]], align 41926// CHECK10-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1927// CHECK10: omp.inner.for.cond:1928// CHECK10-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41929// CHECK10-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41930// CHECK10-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP12]], [[TMP13]]1931// CHECK10-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1932// CHECK10: omp.inner.for.body:1933// CHECK10-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41934// CHECK10-NEXT: [[TMP15:%.*]] = zext i32 [[TMP14]] to i641935// CHECK10-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41936// CHECK10-NEXT: [[TMP17:%.*]] = zext i32 [[TMP16]] to i641937// CHECK10-NEXT: [[TMP18:%.*]] = load i32, ptr [[N_ADDR]], align 41938// CHECK10-NEXT: store i32 [[TMP18]], ptr [[N_CASTED]], align 41939// CHECK10-NEXT: [[TMP19:%.*]] = load i64, ptr [[N_CASTED]], align 81940// 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]])1941// CHECK10-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1942// CHECK10: omp.inner.for.inc:1943// CHECK10-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41944// CHECK10-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41945// CHECK10-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP20]], [[TMP21]]1946// CHECK10-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 41947// CHECK10-NEXT: br label [[OMP_INNER_FOR_COND]]1948// CHECK10: omp.inner.for.end:1949// CHECK10-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1950// CHECK10: omp.loop.exit:1951// CHECK10-NEXT: [[TMP22:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81952// CHECK10-NEXT: [[TMP23:%.*]] = load i32, ptr [[TMP22]], align 41953// CHECK10-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP23]])1954// CHECK10-NEXT: br label [[OMP_PRECOND_END]]1955// CHECK10: omp.precond.end:1956// CHECK10-NEXT: ret void1957//1958//1959// CHECK10-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined.omp_outlined1960// 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]] {1961// CHECK10-NEXT: entry:1962// CHECK10-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81963// CHECK10-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81964// CHECK10-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 81965// CHECK10-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 81966// CHECK10-NEXT: [[N_ADDR:%.*]] = alloca i64, align 81967// CHECK10-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81968// CHECK10-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41969// CHECK10-NEXT: [[TMP:%.*]] = alloca i32, align 41970// CHECK10-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41971// CHECK10-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41972// CHECK10-NEXT: [[I:%.*]] = alloca i32, align 41973// CHECK10-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41974// CHECK10-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41975// CHECK10-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41976// CHECK10-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41977// CHECK10-NEXT: [[I4:%.*]] = alloca i32, align 41978// CHECK10-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81979// CHECK10-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81980// CHECK10-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 81981// CHECK10-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 81982// CHECK10-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 81983// CHECK10-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81984// CHECK10-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81985// CHECK10-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41986// CHECK10-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 41987// CHECK10-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41988// CHECK10-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 01989// CHECK10-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11990// CHECK10-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11991// CHECK10-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41992// CHECK10-NEXT: store i32 0, ptr [[I]], align 41993// CHECK10-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41994// CHECK10-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]1995// CHECK10-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1996// CHECK10: omp.precond.then:1997// CHECK10-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41998// CHECK10-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41999// CHECK10-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_UB]], align 42000// CHECK10-NEXT: [[TMP5:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 82001// CHECK10-NEXT: [[CONV:%.*]] = trunc i64 [[TMP5]] to i322002// CHECK10-NEXT: [[TMP6:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 82003// CHECK10-NEXT: [[CONV3:%.*]] = trunc i64 [[TMP6]] to i322004// CHECK10-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 42005// CHECK10-NEXT: store i32 [[CONV3]], ptr [[DOTOMP_UB]], align 42006// CHECK10-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42007// CHECK10-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42008// CHECK10-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82009// CHECK10-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 42010// 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)2011// CHECK10-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42012// CHECK10-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42013// CHECK10-NEXT: [[CMP5:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]2014// CHECK10-NEXT: br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2015// CHECK10: cond.true:2016// CHECK10-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42017// CHECK10-NEXT: br label [[COND_END:%.*]]2018// CHECK10: cond.false:2019// CHECK10-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42020// CHECK10-NEXT: br label [[COND_END]]2021// CHECK10: cond.end:2022// CHECK10-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]2023// CHECK10-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 42024// CHECK10-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42025// CHECK10-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 42026// CHECK10-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2027// CHECK10: omp.inner.for.cond:2028// CHECK10-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42029// CHECK10-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42030// CHECK10-NEXT: [[CMP6:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]2031// CHECK10-NEXT: br i1 [[CMP6]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2032// CHECK10: omp.inner.for.body:2033// CHECK10-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42034// CHECK10-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 12035// CHECK10-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]2036// CHECK10-NEXT: store i32 [[ADD]], ptr [[I4]], align 42037// CHECK10-NEXT: [[TMP17:%.*]] = load i32, ptr [[I4]], align 42038// CHECK10-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP17]] to i642039// CHECK10-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]2040// CHECK10-NEXT: store i32 0, ptr [[ARRAYIDX]], align 42041// CHECK10-NEXT: [[TMP18:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82042// CHECK10-NEXT: [[TMP19:%.*]] = load i32, ptr [[TMP18]], align 42043// CHECK10-NEXT: [[TMP20:%.*]] = call i32 @__kmpc_cancel(ptr @[[GLOB3]], i32 [[TMP19]], i32 2)2044// CHECK10-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 02045// CHECK10-NEXT: br i1 [[TMP21]], label [[DOTCANCEL_EXIT:%.*]], label [[DOTCANCEL_CONTINUE:%.*]]2046// CHECK10: .cancel.exit:2047// CHECK10-NEXT: br label [[CANCEL_EXIT:%.*]]2048// CHECK10: .cancel.continue:2049// CHECK10-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]2050// CHECK10: omp.body.continue:2051// CHECK10-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2052// CHECK10: omp.inner.for.inc:2053// CHECK10-NEXT: [[TMP22:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42054// CHECK10-NEXT: [[ADD7:%.*]] = add nsw i32 [[TMP22]], 12055// CHECK10-NEXT: store i32 [[ADD7]], ptr [[DOTOMP_IV]], align 42056// CHECK10-NEXT: br label [[OMP_INNER_FOR_COND]]2057// CHECK10: omp.inner.for.end:2058// CHECK10-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2059// CHECK10: omp.loop.exit:2060// CHECK10-NEXT: [[TMP23:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82061// CHECK10-NEXT: [[TMP24:%.*]] = load i32, ptr [[TMP23]], align 42062// CHECK10-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP24]])2063// CHECK10-NEXT: br label [[OMP_PRECOND_END]]2064// CHECK10: cancel.exit:2065// CHECK10-NEXT: [[TMP25:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82066// CHECK10-NEXT: [[TMP26:%.*]] = load i32, ptr [[TMP25]], align 42067// CHECK10-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP26]])2068// CHECK10-NEXT: br label [[CANCEL_CONT:%.*]]2069// CHECK10: omp.precond.end:2070// CHECK10-NEXT: br label [[CANCEL_CONT]]2071// CHECK10: cancel.cont:2072// CHECK10-NEXT: ret void2073//2074//2075// CHECK10-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l582076// CHECK10-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], ptr noundef [[G:%.*]]) #[[ATTR0]] {2077// CHECK10-NEXT: entry:2078// CHECK10-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 82079// CHECK10-NEXT: [[N_ADDR:%.*]] = alloca i64, align 82080// CHECK10-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 82081// CHECK10-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 82082// CHECK10-NEXT: [[N_CASTED:%.*]] = alloca i64, align 82083// CHECK10-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 82084// CHECK10-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 82085// CHECK10-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 82086// CHECK10-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 82087// CHECK10-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 82088// CHECK10-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 42089// CHECK10-NEXT: store i32 [[TMP1]], ptr [[N_CASTED]], align 42090// CHECK10-NEXT: [[TMP2:%.*]] = load i64, ptr [[N_CASTED]], align 82091// CHECK10-NEXT: [[TMP3:%.*]] = load ptr, ptr [[G_ADDR]], align 82092// 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_l58.omp_outlined, i64 [[TMP2]], ptr [[TMP0]], ptr [[TMP3]])2093// CHECK10-NEXT: ret void2094//2095//2096// CHECK10-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l58.omp_outlined2097// 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]] {2098// CHECK10-NEXT: entry:2099// CHECK10-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82100// CHECK10-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82101// CHECK10-NEXT: [[N_ADDR:%.*]] = alloca i64, align 82102// CHECK10-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 82103// CHECK10-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 82104// CHECK10-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42105// CHECK10-NEXT: [[TMP:%.*]] = alloca i32, align 42106// CHECK10-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42107// CHECK10-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42108// CHECK10-NEXT: [[I:%.*]] = alloca i32, align 42109// CHECK10-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 42110// CHECK10-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 42111// CHECK10-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42112// CHECK10-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42113// CHECK10-NEXT: [[I3:%.*]] = alloca i32, align 42114// CHECK10-NEXT: [[N_CASTED:%.*]] = alloca i64, align 82115// CHECK10-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82116// CHECK10-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82117// CHECK10-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 82118// CHECK10-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 82119// CHECK10-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 82120// CHECK10-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 82121// CHECK10-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 42122// CHECK10-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 42123// CHECK10-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42124// CHECK10-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 02125// CHECK10-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12126// CHECK10-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12127// CHECK10-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42128// CHECK10-NEXT: store i32 0, ptr [[I]], align 42129// CHECK10-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42130// CHECK10-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]2131// CHECK10-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2132// CHECK10: omp.precond.then:2133// CHECK10-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 42134// CHECK10-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42135// CHECK10-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_COMB_UB]], align 42136// CHECK10-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42137// CHECK10-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42138// CHECK10-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82139// CHECK10-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 42140// 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)2141// CHECK10-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42142// CHECK10-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42143// CHECK10-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP7]], [[TMP8]]2144// CHECK10-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2145// CHECK10: cond.true:2146// CHECK10-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42147// CHECK10-NEXT: br label [[COND_END:%.*]]2148// CHECK10: cond.false:2149// CHECK10-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42150// CHECK10-NEXT: br label [[COND_END]]2151// CHECK10: cond.end:2152// CHECK10-NEXT: [[COND:%.*]] = phi i32 [ [[TMP9]], [[COND_TRUE]] ], [ [[TMP10]], [[COND_FALSE]] ]2153// CHECK10-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 42154// CHECK10-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 42155// CHECK10-NEXT: store i32 [[TMP11]], ptr [[DOTOMP_IV]], align 42156// CHECK10-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2157// CHECK10: omp.inner.for.cond:2158// CHECK10-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42159// CHECK10-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42160// CHECK10-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP12]], [[TMP13]]2161// CHECK10-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2162// CHECK10: omp.inner.for.body:2163// CHECK10-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 42164// CHECK10-NEXT: [[TMP15:%.*]] = zext i32 [[TMP14]] to i642165// CHECK10-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42166// CHECK10-NEXT: [[TMP17:%.*]] = zext i32 [[TMP16]] to i642167// CHECK10-NEXT: [[TMP18:%.*]] = load i32, ptr [[N_ADDR]], align 42168// CHECK10-NEXT: store i32 [[TMP18]], ptr [[N_CASTED]], align 42169// CHECK10-NEXT: [[TMP19:%.*]] = load i64, ptr [[N_CASTED]], align 82170// CHECK10-NEXT: [[TMP20:%.*]] = load ptr, ptr [[G_ADDR]], align 82171// 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_l58.omp_outlined.omp_outlined, i64 [[TMP15]], i64 [[TMP17]], i64 [[TMP19]], ptr [[TMP0]], ptr [[TMP20]])2172// CHECK10-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2173// CHECK10: omp.inner.for.inc:2174// CHECK10-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42175// CHECK10-NEXT: [[TMP22:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 42176// CHECK10-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP21]], [[TMP22]]2177// CHECK10-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 42178// CHECK10-NEXT: br label [[OMP_INNER_FOR_COND]]2179// CHECK10: omp.inner.for.end:2180// CHECK10-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2181// CHECK10: omp.loop.exit:2182// CHECK10-NEXT: [[TMP23:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82183// CHECK10-NEXT: [[TMP24:%.*]] = load i32, ptr [[TMP23]], align 42184// CHECK10-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP24]])2185// CHECK10-NEXT: br label [[OMP_PRECOND_END]]2186// CHECK10: omp.precond.end:2187// CHECK10-NEXT: ret void2188//2189//2190// CHECK10-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l58.omp_outlined.omp_outlined2191// 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]] {2192// CHECK10-NEXT: entry:2193// CHECK10-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82194// CHECK10-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82195// CHECK10-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 82196// CHECK10-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 82197// CHECK10-NEXT: [[N_ADDR:%.*]] = alloca i64, align 82198// CHECK10-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 82199// CHECK10-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 82200// CHECK10-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42201// CHECK10-NEXT: [[TMP:%.*]] = alloca i32, align 42202// CHECK10-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42203// CHECK10-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42204// CHECK10-NEXT: [[I:%.*]] = alloca i32, align 42205// CHECK10-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 42206// CHECK10-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 42207// CHECK10-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42208// CHECK10-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42209// CHECK10-NEXT: [[I4:%.*]] = alloca i32, align 42210// CHECK10-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82211// CHECK10-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82212// CHECK10-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 82213// CHECK10-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 82214// CHECK10-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 82215// CHECK10-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 82216// CHECK10-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 82217// CHECK10-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 82218// CHECK10-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 42219// CHECK10-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 42220// CHECK10-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42221// CHECK10-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 02222// CHECK10-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12223// CHECK10-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12224// CHECK10-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42225// CHECK10-NEXT: store i32 0, ptr [[I]], align 42226// CHECK10-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42227// CHECK10-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]2228// CHECK10-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2229// CHECK10: omp.precond.then:2230// CHECK10-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 42231// CHECK10-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42232// CHECK10-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_UB]], align 42233// CHECK10-NEXT: [[TMP5:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 82234// CHECK10-NEXT: [[CONV:%.*]] = trunc i64 [[TMP5]] to i322235// CHECK10-NEXT: [[TMP6:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 82236// CHECK10-NEXT: [[CONV3:%.*]] = trunc i64 [[TMP6]] to i322237// CHECK10-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 42238// CHECK10-NEXT: store i32 [[CONV3]], ptr [[DOTOMP_UB]], align 42239// CHECK10-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42240// CHECK10-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42241// CHECK10-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82242// CHECK10-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 42243// 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)2244// CHECK10-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42245// CHECK10-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42246// CHECK10-NEXT: [[CMP5:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]2247// CHECK10-NEXT: br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2248// CHECK10: cond.true:2249// CHECK10-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42250// CHECK10-NEXT: br label [[COND_END:%.*]]2251// CHECK10: cond.false:2252// CHECK10-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42253// CHECK10-NEXT: br label [[COND_END]]2254// CHECK10: cond.end:2255// CHECK10-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]2256// CHECK10-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 42257// CHECK10-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42258// CHECK10-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 42259// CHECK10-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2260// CHECK10: omp.inner.for.cond:2261// CHECK10-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42262// CHECK10-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42263// CHECK10-NEXT: [[CMP6:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]2264// CHECK10-NEXT: br i1 [[CMP6]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2265// CHECK10: omp.inner.for.body:2266// CHECK10-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42267// CHECK10-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 12268// CHECK10-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]2269// CHECK10-NEXT: store i32 [[ADD]], ptr [[I4]], align 42270// CHECK10-NEXT: [[TMP17:%.*]] = load ptr, ptr [[G_ADDR]], align 82271// CHECK10-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP17]], i64 02272// CHECK10-NEXT: [[TMP18:%.*]] = load i32, ptr [[ARRAYIDX]], align 42273// CHECK10-NEXT: [[TMP19:%.*]] = load i32, ptr [[I4]], align 42274// CHECK10-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP19]] to i642275// CHECK10-NEXT: [[ARRAYIDX7:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]2276// CHECK10-NEXT: store i32 [[TMP18]], ptr [[ARRAYIDX7]], align 42277// CHECK10-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]2278// CHECK10: omp.body.continue:2279// CHECK10-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2280// CHECK10: omp.inner.for.inc:2281// CHECK10-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42282// CHECK10-NEXT: [[ADD8:%.*]] = add nsw i32 [[TMP20]], 12283// CHECK10-NEXT: store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 42284// CHECK10-NEXT: br label [[OMP_INNER_FOR_COND]]2285// CHECK10: omp.inner.for.end:2286// CHECK10-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2287// CHECK10: omp.loop.exit:2288// CHECK10-NEXT: [[TMP21:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82289// CHECK10-NEXT: [[TMP22:%.*]] = load i32, ptr [[TMP21]], align 42290// CHECK10-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP22]])2291// CHECK10-NEXT: br label [[OMP_PRECOND_END]]2292// CHECK10: omp.precond.end:2293// CHECK10-NEXT: ret void2294//2295//2296// CHECK12-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l512297// 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]+]] {2298// CHECK12-NEXT: entry:2299// CHECK12-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42300// CHECK12-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42301// CHECK12-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42302// CHECK12-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 42303// CHECK12-NEXT: [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca i32, align 42304// CHECK12-NEXT: [[N_CASTED:%.*]] = alloca i32, align 42305// CHECK12-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB3:[0-9]+]])2306// CHECK12-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42307// CHECK12-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42308// CHECK12-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42309// CHECK12-NEXT: store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 42310// CHECK12-NEXT: store i32 [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 42311// CHECK12-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 42312// CHECK12-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 42313// CHECK12-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR2]], align 42314// CHECK12-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB3]], i32 [[TMP0]], i32 [[TMP2]], i32 [[TMP3]])2315// CHECK12-NEXT: [[TMP4:%.*]] = load i32, ptr [[N_ADDR]], align 42316// CHECK12-NEXT: store i32 [[TMP4]], ptr [[N_CASTED]], align 42317// CHECK12-NEXT: [[TMP5:%.*]] = load i32, ptr [[N_CASTED]], align 42318// 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]])2319// CHECK12-NEXT: ret void2320//2321//2322// CHECK12-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined2323// CHECK12-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR0]] {2324// CHECK12-NEXT: entry:2325// CHECK12-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42326// CHECK12-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42327// CHECK12-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42328// CHECK12-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42329// CHECK12-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42330// CHECK12-NEXT: [[TMP:%.*]] = alloca i32, align 42331// CHECK12-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42332// CHECK12-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42333// CHECK12-NEXT: [[I:%.*]] = alloca i32, align 42334// CHECK12-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 42335// CHECK12-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 42336// CHECK12-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42337// CHECK12-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42338// CHECK12-NEXT: [[I3:%.*]] = alloca i32, align 42339// CHECK12-NEXT: [[N_CASTED:%.*]] = alloca i32, align 42340// CHECK12-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42341// CHECK12-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42342// CHECK12-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42343// CHECK12-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42344// CHECK12-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42345// CHECK12-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 42346// CHECK12-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 42347// CHECK12-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42348// CHECK12-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 02349// CHECK12-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12350// CHECK12-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12351// CHECK12-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42352// CHECK12-NEXT: store i32 0, ptr [[I]], align 42353// CHECK12-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42354// CHECK12-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]2355// CHECK12-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2356// CHECK12: omp.precond.then:2357// CHECK12-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 42358// CHECK12-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42359// CHECK12-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_COMB_UB]], align 42360// CHECK12-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42361// CHECK12-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42362// CHECK12-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42363// CHECK12-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 42364// 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)2365// CHECK12-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42366// CHECK12-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42367// CHECK12-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP7]], [[TMP8]]2368// CHECK12-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2369// CHECK12: cond.true:2370// CHECK12-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42371// CHECK12-NEXT: br label [[COND_END:%.*]]2372// CHECK12: cond.false:2373// CHECK12-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42374// CHECK12-NEXT: br label [[COND_END]]2375// CHECK12: cond.end:2376// CHECK12-NEXT: [[COND:%.*]] = phi i32 [ [[TMP9]], [[COND_TRUE]] ], [ [[TMP10]], [[COND_FALSE]] ]2377// CHECK12-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 42378// CHECK12-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 42379// CHECK12-NEXT: store i32 [[TMP11]], ptr [[DOTOMP_IV]], align 42380// CHECK12-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2381// CHECK12: omp.inner.for.cond:2382// CHECK12-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42383// CHECK12-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42384// CHECK12-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP12]], [[TMP13]]2385// CHECK12-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2386// CHECK12: omp.inner.for.body:2387// CHECK12-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 42388// CHECK12-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42389// CHECK12-NEXT: [[TMP16:%.*]] = load i32, ptr [[N_ADDR]], align 42390// CHECK12-NEXT: store i32 [[TMP16]], ptr [[N_CASTED]], align 42391// CHECK12-NEXT: [[TMP17:%.*]] = load i32, ptr [[N_CASTED]], align 42392// 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]])2393// CHECK12-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2394// CHECK12: omp.inner.for.inc:2395// CHECK12-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42396// CHECK12-NEXT: [[TMP19:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 42397// CHECK12-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP18]], [[TMP19]]2398// CHECK12-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 42399// CHECK12-NEXT: br label [[OMP_INNER_FOR_COND]]2400// CHECK12: omp.inner.for.end:2401// CHECK12-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2402// CHECK12: omp.loop.exit:2403// CHECK12-NEXT: [[TMP20:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42404// CHECK12-NEXT: [[TMP21:%.*]] = load i32, ptr [[TMP20]], align 42405// CHECK12-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP21]])2406// CHECK12-NEXT: br label [[OMP_PRECOND_END]]2407// CHECK12: omp.precond.end:2408// CHECK12-NEXT: ret void2409//2410//2411// CHECK12-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l51.omp_outlined.omp_outlined2412// 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]] {2413// CHECK12-NEXT: entry:2414// CHECK12-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42415// CHECK12-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42416// CHECK12-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i32, align 42417// CHECK12-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i32, align 42418// CHECK12-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42419// CHECK12-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42420// CHECK12-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42421// CHECK12-NEXT: [[TMP:%.*]] = alloca i32, align 42422// CHECK12-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42423// CHECK12-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42424// CHECK12-NEXT: [[I:%.*]] = alloca i32, align 42425// CHECK12-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 42426// CHECK12-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 42427// CHECK12-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42428// CHECK12-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42429// CHECK12-NEXT: [[I3:%.*]] = alloca i32, align 42430// CHECK12-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42431// CHECK12-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42432// CHECK12-NEXT: store i32 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 42433// CHECK12-NEXT: store i32 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 42434// CHECK12-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42435// CHECK12-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42436// CHECK12-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42437// CHECK12-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 42438// CHECK12-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 42439// CHECK12-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42440// CHECK12-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 02441// CHECK12-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12442// CHECK12-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12443// CHECK12-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42444// CHECK12-NEXT: store i32 0, ptr [[I]], align 42445// CHECK12-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42446// CHECK12-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]2447// CHECK12-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2448// CHECK12: omp.precond.then:2449// CHECK12-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 42450// CHECK12-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42451// CHECK12-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_UB]], align 42452// CHECK12-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTPREVIOUS_LB__ADDR]], align 42453// CHECK12-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTPREVIOUS_UB__ADDR]], align 42454// CHECK12-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_LB]], align 42455// CHECK12-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_UB]], align 42456// CHECK12-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42457// CHECK12-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42458// CHECK12-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42459// CHECK12-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 42460// 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)2461// CHECK12-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42462// CHECK12-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42463// CHECK12-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]2464// CHECK12-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2465// CHECK12: cond.true:2466// CHECK12-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42467// CHECK12-NEXT: br label [[COND_END:%.*]]2468// CHECK12: cond.false:2469// CHECK12-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42470// CHECK12-NEXT: br label [[COND_END]]2471// CHECK12: cond.end:2472// CHECK12-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]2473// CHECK12-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 42474// CHECK12-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42475// CHECK12-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 42476// CHECK12-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2477// CHECK12: omp.inner.for.cond:2478// CHECK12-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42479// CHECK12-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42480// CHECK12-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]2481// CHECK12-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2482// CHECK12: omp.inner.for.body:2483// CHECK12-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42484// CHECK12-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 12485// CHECK12-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]2486// CHECK12-NEXT: store i32 [[ADD]], ptr [[I3]], align 42487// CHECK12-NEXT: [[TMP17:%.*]] = load i32, ptr [[I3]], align 42488// CHECK12-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i32 0, i32 [[TMP17]]2489// CHECK12-NEXT: store i32 0, ptr [[ARRAYIDX]], align 42490// CHECK12-NEXT: [[TMP18:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42491// CHECK12-NEXT: [[TMP19:%.*]] = load i32, ptr [[TMP18]], align 42492// CHECK12-NEXT: [[TMP20:%.*]] = call i32 @__kmpc_cancel(ptr @[[GLOB3]], i32 [[TMP19]], i32 2)2493// CHECK12-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 02494// CHECK12-NEXT: br i1 [[TMP21]], label [[DOTCANCEL_EXIT:%.*]], label [[DOTCANCEL_CONTINUE:%.*]]2495// CHECK12: .cancel.exit:2496// CHECK12-NEXT: br label [[CANCEL_EXIT:%.*]]2497// CHECK12: .cancel.continue:2498// CHECK12-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]2499// CHECK12: omp.body.continue:2500// CHECK12-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2501// CHECK12: omp.inner.for.inc:2502// CHECK12-NEXT: [[TMP22:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42503// CHECK12-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP22]], 12504// CHECK12-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 42505// CHECK12-NEXT: br label [[OMP_INNER_FOR_COND]]2506// CHECK12: omp.inner.for.end:2507// CHECK12-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2508// CHECK12: omp.loop.exit:2509// CHECK12-NEXT: [[TMP23:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42510// CHECK12-NEXT: [[TMP24:%.*]] = load i32, ptr [[TMP23]], align 42511// CHECK12-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP24]])2512// CHECK12-NEXT: br label [[OMP_PRECOND_END]]2513// CHECK12: cancel.exit:2514// CHECK12-NEXT: [[TMP25:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42515// CHECK12-NEXT: [[TMP26:%.*]] = load i32, ptr [[TMP25]], align 42516// CHECK12-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP26]])2517// CHECK12-NEXT: br label [[CANCEL_CONT:%.*]]2518// CHECK12: omp.precond.end:2519// CHECK12-NEXT: br label [[CANCEL_CONT]]2520// CHECK12: cancel.cont:2521// CHECK12-NEXT: ret void2522//2523//2524// CHECK12-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l582525// CHECK12-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]], ptr noundef [[G:%.*]]) #[[ATTR0]] {2526// CHECK12-NEXT: entry:2527// CHECK12-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42528// CHECK12-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42529// CHECK12-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42530// CHECK12-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 42531// CHECK12-NEXT: [[N_CASTED:%.*]] = alloca i32, align 42532// CHECK12-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42533// CHECK12-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42534// CHECK12-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42535// CHECK12-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 42536// CHECK12-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42537// CHECK12-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 42538// CHECK12-NEXT: store i32 [[TMP1]], ptr [[N_CASTED]], align 42539// CHECK12-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_CASTED]], align 42540// CHECK12-NEXT: [[TMP3:%.*]] = load ptr, ptr [[G_ADDR]], align 42541// 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_l58.omp_outlined, i32 [[TMP2]], ptr [[TMP0]], ptr [[TMP3]])2542// CHECK12-NEXT: ret void2543//2544//2545// CHECK12-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l58.omp_outlined2546// 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]] {2547// CHECK12-NEXT: entry:2548// CHECK12-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42549// CHECK12-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42550// CHECK12-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42551// CHECK12-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42552// CHECK12-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 42553// CHECK12-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42554// CHECK12-NEXT: [[TMP:%.*]] = alloca i32, align 42555// CHECK12-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42556// CHECK12-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42557// CHECK12-NEXT: [[I:%.*]] = alloca i32, align 42558// CHECK12-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 42559// CHECK12-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 42560// CHECK12-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42561// CHECK12-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42562// CHECK12-NEXT: [[I3:%.*]] = alloca i32, align 42563// CHECK12-NEXT: [[N_CASTED:%.*]] = alloca i32, align 42564// CHECK12-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42565// CHECK12-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42566// CHECK12-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42567// CHECK12-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42568// CHECK12-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 42569// CHECK12-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42570// CHECK12-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 42571// CHECK12-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 42572// CHECK12-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42573// CHECK12-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 02574// CHECK12-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12575// CHECK12-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12576// CHECK12-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42577// CHECK12-NEXT: store i32 0, ptr [[I]], align 42578// CHECK12-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42579// CHECK12-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]2580// CHECK12-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2581// CHECK12: omp.precond.then:2582// CHECK12-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 42583// CHECK12-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42584// CHECK12-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_COMB_UB]], align 42585// CHECK12-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42586// CHECK12-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42587// CHECK12-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42588// CHECK12-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 42589// 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)2590// CHECK12-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42591// CHECK12-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42592// CHECK12-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP7]], [[TMP8]]2593// CHECK12-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2594// CHECK12: cond.true:2595// CHECK12-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42596// CHECK12-NEXT: br label [[COND_END:%.*]]2597// CHECK12: cond.false:2598// CHECK12-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42599// CHECK12-NEXT: br label [[COND_END]]2600// CHECK12: cond.end:2601// CHECK12-NEXT: [[COND:%.*]] = phi i32 [ [[TMP9]], [[COND_TRUE]] ], [ [[TMP10]], [[COND_FALSE]] ]2602// CHECK12-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 42603// CHECK12-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 42604// CHECK12-NEXT: store i32 [[TMP11]], ptr [[DOTOMP_IV]], align 42605// CHECK12-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2606// CHECK12: omp.inner.for.cond:2607// CHECK12-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42608// CHECK12-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42609// CHECK12-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP12]], [[TMP13]]2610// CHECK12-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2611// CHECK12: omp.inner.for.body:2612// CHECK12-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 42613// CHECK12-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42614// CHECK12-NEXT: [[TMP16:%.*]] = load i32, ptr [[N_ADDR]], align 42615// CHECK12-NEXT: store i32 [[TMP16]], ptr [[N_CASTED]], align 42616// CHECK12-NEXT: [[TMP17:%.*]] = load i32, ptr [[N_CASTED]], align 42617// CHECK12-NEXT: [[TMP18:%.*]] = load ptr, ptr [[G_ADDR]], align 42618// 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_l58.omp_outlined.omp_outlined, i32 [[TMP14]], i32 [[TMP15]], i32 [[TMP17]], ptr [[TMP0]], ptr [[TMP18]])2619// CHECK12-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2620// CHECK12: omp.inner.for.inc:2621// CHECK12-NEXT: [[TMP19:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42622// CHECK12-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 42623// CHECK12-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP19]], [[TMP20]]2624// CHECK12-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 42625// CHECK12-NEXT: br label [[OMP_INNER_FOR_COND]]2626// CHECK12: omp.inner.for.end:2627// CHECK12-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2628// CHECK12: omp.loop.exit:2629// CHECK12-NEXT: [[TMP21:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42630// CHECK12-NEXT: [[TMP22:%.*]] = load i32, ptr [[TMP21]], align 42631// CHECK12-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP22]])2632// CHECK12-NEXT: br label [[OMP_PRECOND_END]]2633// CHECK12: omp.precond.end:2634// CHECK12-NEXT: ret void2635//2636//2637// CHECK12-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16target_teams_funPi_l58.omp_outlined.omp_outlined2638// 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]] {2639// CHECK12-NEXT: entry:2640// CHECK12-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42641// CHECK12-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42642// CHECK12-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i32, align 42643// CHECK12-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i32, align 42644// CHECK12-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42645// CHECK12-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42646// CHECK12-NEXT: [[G_ADDR:%.*]] = alloca ptr, align 42647// CHECK12-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42648// CHECK12-NEXT: [[TMP:%.*]] = alloca i32, align 42649// CHECK12-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42650// CHECK12-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42651// CHECK12-NEXT: [[I:%.*]] = alloca i32, align 42652// CHECK12-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 42653// CHECK12-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 42654// CHECK12-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42655// CHECK12-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42656// CHECK12-NEXT: [[I3:%.*]] = alloca i32, align 42657// CHECK12-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42658// CHECK12-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42659// CHECK12-NEXT: store i32 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 42660// CHECK12-NEXT: store i32 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 42661// CHECK12-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42662// CHECK12-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42663// CHECK12-NEXT: store ptr [[G]], ptr [[G_ADDR]], align 42664// CHECK12-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42665// CHECK12-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 42666// CHECK12-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 42667// CHECK12-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42668// CHECK12-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP2]], 02669// CHECK12-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12670// CHECK12-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12671// CHECK12-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42672// CHECK12-NEXT: store i32 0, ptr [[I]], align 42673// CHECK12-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42674// CHECK12-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP3]]2675// CHECK12-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2676// CHECK12: omp.precond.then:2677// CHECK12-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 42678// CHECK12-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42679// CHECK12-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_UB]], align 42680// CHECK12-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTPREVIOUS_LB__ADDR]], align 42681// CHECK12-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTPREVIOUS_UB__ADDR]], align 42682// CHECK12-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_LB]], align 42683// CHECK12-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_UB]], align 42684// CHECK12-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42685// CHECK12-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42686// CHECK12-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42687// CHECK12-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 42688// 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)2689// CHECK12-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42690// CHECK12-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42691// CHECK12-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]2692// CHECK12-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2693// CHECK12: cond.true:2694// CHECK12-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42695// CHECK12-NEXT: br label [[COND_END:%.*]]2696// CHECK12: cond.false:2697// CHECK12-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42698// CHECK12-NEXT: br label [[COND_END]]2699// CHECK12: cond.end:2700// CHECK12-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]2701// CHECK12-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 42702// CHECK12-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42703// CHECK12-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 42704// CHECK12-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2705// CHECK12: omp.inner.for.cond:2706// CHECK12-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42707// CHECK12-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42708// CHECK12-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]2709// CHECK12-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2710// CHECK12: omp.inner.for.body:2711// CHECK12-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42712// CHECK12-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 12713// CHECK12-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]2714// CHECK12-NEXT: store i32 [[ADD]], ptr [[I3]], align 42715// CHECK12-NEXT: [[TMP17:%.*]] = load ptr, ptr [[G_ADDR]], align 42716// CHECK12-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP17]], i32 02717// CHECK12-NEXT: [[TMP18:%.*]] = load i32, ptr [[ARRAYIDX]], align 42718// CHECK12-NEXT: [[TMP19:%.*]] = load i32, ptr [[I3]], align 42719// CHECK12-NEXT: [[ARRAYIDX6:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i32 0, i32 [[TMP19]]2720// CHECK12-NEXT: store i32 [[TMP18]], ptr [[ARRAYIDX6]], align 42721// CHECK12-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]2722// CHECK12: omp.body.continue:2723// CHECK12-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2724// CHECK12: omp.inner.for.inc:2725// CHECK12-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42726// CHECK12-NEXT: [[ADD7:%.*]] = add nsw i32 [[TMP20]], 12727// CHECK12-NEXT: store i32 [[ADD7]], ptr [[DOTOMP_IV]], align 42728// CHECK12-NEXT: br label [[OMP_INNER_FOR_COND]]2729// CHECK12: omp.inner.for.end:2730// CHECK12-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2731// CHECK12: omp.loop.exit:2732// CHECK12-NEXT: [[TMP21:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42733// CHECK12-NEXT: [[TMP22:%.*]] = load i32, ptr [[TMP21]], align 42734// CHECK12-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP22]])2735// CHECK12-NEXT: br label [[OMP_PRECOND_END]]2736// CHECK12: omp.precond.end:2737// CHECK12-NEXT: ret void2738//2739