brintos / llvm-project-archived public Read only

1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _2// Test host codegen.3// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK14// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s5// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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=CHECK16// RUN: %clang_cc1 -verify -fopenmp-version=45 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK37// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s8// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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=CHECK39// Test host codegen.10// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK111// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s12// RUN: %clang_cc1 -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=CHECK113// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK314// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s15// RUN: %clang_cc1 -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=CHECK316 17// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"18// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s19// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -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}}"20// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"21// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s22// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -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}}"23 24// RUN: %clang_cc1 -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}}"25// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s26// RUN: %clang_cc1 -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}}"27// RUN: %clang_cc1 -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}}"28// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s29// RUN: %clang_cc1 -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}}"30 31// Test target codegen - host bc file has to be created first. (no significant differences with host version of target region)32// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc33// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -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=CHECK1734// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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 %s35// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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=CHECK1736// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc37// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -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=CHECK1938// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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 %s39// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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=CHECK1940 41// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc42// RUN: %clang_cc1 -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}}"43// RUN: %clang_cc1 -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 %s44// RUN: %clang_cc1 -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}}"45// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc46// RUN: %clang_cc1 -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}}"47// RUN: %clang_cc1 -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 %s48// RUN: %clang_cc1 -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}}"49 50// expected-no-diagnostics51#ifndef HEADER52#define HEADER53 54 55void without_schedule_clause(float *a, float *b, float *c, float *d) {56  #pragma omp target57  #pragma omp teams58  #pragma omp distribute59  for (int i = 33; i < 32000000; i += 7) {60    a[i] = b[i] * c[i] * d[i];61  }62}63 64// ... loop body ...65 66 67void static_not_chunked(float *a, float *b, float *c, float *d) {68  #pragma omp target69  #pragma omp teams70  #pragma omp distribute dist_schedule(static)71  for (int i = 32000000; i > 33; i += -7) {72        a[i] = b[i] * c[i] * d[i];73  }74}75 76// ... loop body ...77 78 79void static_chunked(float *a, float *b, float *c, float *d) {80  #pragma omp target81  #pragma omp teams82#pragma omp distribute dist_schedule(static, 5)83  for (unsigned i = 131071; i <= 2147483647; i += 127) {84    a[i] = b[i] * c[i] * d[i];85  }86}87 88// ... loop body ...89 90void test_precond() {91  char a = 0;92  #pragma omp target93  #pragma omp teams94  #pragma omp distribute95  for(char i = a; i < 10; ++i);96}97 98// a is passed as a parameter to the outlined functions99// ..many loads of %0..100 101// no templates for now, as these require special handling in target regions and/or declare target102 103 104template <typename T>105T ftemplate() {106  short aa = 0;107 108#pragma omp target109#pragma omp teams110#pragma omp distribute dist_schedule(static, aa)111  for (int i = 0; i < 100; i++) {112  }113  return T();114}115 116int fint(void) { return ftemplate<int>(); }117 118#endif119// CHECK1-LABEL: define {{[^@]+}}@_Z23without_schedule_clausePfS_S_S_120// CHECK1-SAME: (ptr noundef [[A:%.*]], ptr noundef [[B:%.*]], ptr noundef [[C:%.*]], ptr noundef [[D:%.*]]) #[[ATTR0:[0-9]+]] {121// CHECK1-NEXT:  entry:122// CHECK1-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8123// CHECK1-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8124// CHECK1-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 8125// CHECK1-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 8126// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [4 x ptr], align 8127// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [4 x ptr], align 8128// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [4 x ptr], align 8129// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4130// CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8131// CHECK1-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8132// CHECK1-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8133// CHECK1-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 8134// CHECK1-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 8135// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8136// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 8137// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 8138// CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 8139// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0140// CHECK1-NEXT:    store ptr [[TMP0]], ptr [[TMP4]], align 8141// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0142// CHECK1-NEXT:    store ptr [[TMP0]], ptr [[TMP5]], align 8143// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0144// CHECK1-NEXT:    store ptr null, ptr [[TMP6]], align 8145// CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1146// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[TMP7]], align 8147// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1148// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[TMP8]], align 8149// CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1150// CHECK1-NEXT:    store ptr null, ptr [[TMP9]], align 8151// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2152// CHECK1-NEXT:    store ptr [[TMP2]], ptr [[TMP10]], align 8153// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2154// CHECK1-NEXT:    store ptr [[TMP2]], ptr [[TMP11]], align 8155// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2156// CHECK1-NEXT:    store ptr null, ptr [[TMP12]], align 8157// CHECK1-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3158// CHECK1-NEXT:    store ptr [[TMP3]], ptr [[TMP13]], align 8159// CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3160// CHECK1-NEXT:    store ptr [[TMP3]], ptr [[TMP14]], align 8161// CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3162// CHECK1-NEXT:    store ptr null, ptr [[TMP15]], align 8163// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0164// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0165// CHECK1-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0166// CHECK1-NEXT:    store i32 3, ptr [[TMP18]], align 4167// CHECK1-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1168// CHECK1-NEXT:    store i32 4, ptr [[TMP19]], align 4169// CHECK1-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2170// CHECK1-NEXT:    store ptr [[TMP16]], ptr [[TMP20]], align 8171// CHECK1-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3172// CHECK1-NEXT:    store ptr [[TMP17]], ptr [[TMP21]], align 8173// CHECK1-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4174// CHECK1-NEXT:    store ptr @.offload_sizes, ptr [[TMP22]], align 8175// CHECK1-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5176// CHECK1-NEXT:    store ptr @.offload_maptypes, ptr [[TMP23]], align 8177// CHECK1-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6178// CHECK1-NEXT:    store ptr null, ptr [[TMP24]], align 8179// CHECK1-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7180// CHECK1-NEXT:    store ptr null, ptr [[TMP25]], align 8181// CHECK1-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8182// CHECK1-NEXT:    store i64 4571424, ptr [[TMP26]], align 8183// CHECK1-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9184// CHECK1-NEXT:    store i64 0, ptr [[TMP27]], align 8185// CHECK1-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10186// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP28]], align 4187// CHECK1-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11188// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP29]], align 4189// CHECK1-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12190// CHECK1-NEXT:    store i32 0, ptr [[TMP30]], align 4191// CHECK1-NEXT:    [[TMP31:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z23without_schedule_clausePfS_S_S__l56.region_id, ptr [[KERNEL_ARGS]])192// CHECK1-NEXT:    [[TMP32:%.*]] = icmp ne i32 [[TMP31]], 0193// CHECK1-NEXT:    br i1 [[TMP32]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]194// CHECK1:       omp_offload.failed:195// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z23without_schedule_clausePfS_S_S__l56(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], ptr [[TMP3]]) #[[ATTR2:[0-9]+]]196// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]197// CHECK1:       omp_offload.cont:198// CHECK1-NEXT:    ret void199//200//201// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z23without_schedule_clausePfS_S_S__l56202// CHECK1-SAME: (ptr noundef [[A:%.*]], ptr noundef [[B:%.*]], ptr noundef [[C:%.*]], ptr noundef [[D:%.*]]) #[[ATTR1:[0-9]+]] {203// CHECK1-NEXT:  entry:204// CHECK1-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8205// CHECK1-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8206// CHECK1-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 8207// CHECK1-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 8208// CHECK1-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8209// CHECK1-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8210// CHECK1-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 8211// CHECK1-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 8212// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 4, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z23without_schedule_clausePfS_S_S__l56.omp_outlined, ptr [[A_ADDR]], ptr [[B_ADDR]], ptr [[C_ADDR]], ptr [[D_ADDR]])213// CHECK1-NEXT:    ret void214//215//216// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z23without_schedule_clausePfS_S_S__l56.omp_outlined217// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[A:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[B:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[C:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[D:%.*]]) #[[ATTR1]] {218// CHECK1-NEXT:  entry:219// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8220// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8221// CHECK1-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8222// CHECK1-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8223// CHECK1-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 8224// CHECK1-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 8225// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4226// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4227// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4228// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4229// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4230// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4231// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4232// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8233// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8234// CHECK1-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8235// CHECK1-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8236// CHECK1-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 8237// CHECK1-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 8238// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8239// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 8240// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 8241// CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 8242// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4243// CHECK1-NEXT:    store i32 4571423, ptr [[DOTOMP_UB]], align 4244// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4245// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4246// CHECK1-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8247// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 4248// CHECK1-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP5]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)249// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4250// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP6]], 4571423251// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]252// CHECK1:       cond.true:253// CHECK1-NEXT:    br label [[COND_END:%.*]]254// CHECK1:       cond.false:255// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4256// CHECK1-NEXT:    br label [[COND_END]]257// CHECK1:       cond.end:258// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 4571423, [[COND_TRUE]] ], [ [[TMP7]], [[COND_FALSE]] ]259// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4260// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4261// CHECK1-NEXT:    store i32 [[TMP8]], ptr [[DOTOMP_IV]], align 4262// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]263// CHECK1:       omp.inner.for.cond:264// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4265// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4266// CHECK1-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP9]], [[TMP10]]267// CHECK1-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]268// CHECK1:       omp.inner.for.body:269// CHECK1-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4270// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP11]], 7271// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 33, [[MUL]]272// CHECK1-NEXT:    store i32 [[ADD]], ptr [[I]], align 4273// CHECK1-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[TMP1]], align 8274// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[I]], align 4275// CHECK1-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP13]] to i64276// CHECK1-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds float, ptr [[TMP12]], i64 [[IDXPROM]]277// CHECK1-NEXT:    [[TMP14:%.*]] = load float, ptr [[ARRAYIDX]], align 4278// CHECK1-NEXT:    [[TMP15:%.*]] = load ptr, ptr [[TMP2]], align 8279// CHECK1-NEXT:    [[TMP16:%.*]] = load i32, ptr [[I]], align 4280// CHECK1-NEXT:    [[IDXPROM2:%.*]] = sext i32 [[TMP16]] to i64281// CHECK1-NEXT:    [[ARRAYIDX3:%.*]] = getelementptr inbounds float, ptr [[TMP15]], i64 [[IDXPROM2]]282// CHECK1-NEXT:    [[TMP17:%.*]] = load float, ptr [[ARRAYIDX3]], align 4283// CHECK1-NEXT:    [[MUL4:%.*]] = fmul float [[TMP14]], [[TMP17]]284// CHECK1-NEXT:    [[TMP18:%.*]] = load ptr, ptr [[TMP3]], align 8285// CHECK1-NEXT:    [[TMP19:%.*]] = load i32, ptr [[I]], align 4286// CHECK1-NEXT:    [[IDXPROM5:%.*]] = sext i32 [[TMP19]] to i64287// CHECK1-NEXT:    [[ARRAYIDX6:%.*]] = getelementptr inbounds float, ptr [[TMP18]], i64 [[IDXPROM5]]288// CHECK1-NEXT:    [[TMP20:%.*]] = load float, ptr [[ARRAYIDX6]], align 4289// CHECK1-NEXT:    [[MUL7:%.*]] = fmul float [[MUL4]], [[TMP20]]290// CHECK1-NEXT:    [[TMP21:%.*]] = load ptr, ptr [[TMP0]], align 8291// CHECK1-NEXT:    [[TMP22:%.*]] = load i32, ptr [[I]], align 4292// CHECK1-NEXT:    [[IDXPROM8:%.*]] = sext i32 [[TMP22]] to i64293// CHECK1-NEXT:    [[ARRAYIDX9:%.*]] = getelementptr inbounds float, ptr [[TMP21]], i64 [[IDXPROM8]]294// CHECK1-NEXT:    store float [[MUL7]], ptr [[ARRAYIDX9]], align 4295// CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]296// CHECK1:       omp.body.continue:297// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]298// CHECK1:       omp.inner.for.inc:299// CHECK1-NEXT:    [[TMP23:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4300// CHECK1-NEXT:    [[ADD10:%.*]] = add nsw i32 [[TMP23]], 1301// CHECK1-NEXT:    store i32 [[ADD10]], ptr [[DOTOMP_IV]], align 4302// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]303// CHECK1:       omp.inner.for.end:304// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]305// CHECK1:       omp.loop.exit:306// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP5]])307// CHECK1-NEXT:    ret void308//309//310// CHECK1-LABEL: define {{[^@]+}}@_Z18static_not_chunkedPfS_S_S_311// CHECK1-SAME: (ptr noundef [[A:%.*]], ptr noundef [[B:%.*]], ptr noundef [[C:%.*]], ptr noundef [[D:%.*]]) #[[ATTR0]] {312// CHECK1-NEXT:  entry:313// CHECK1-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8314// CHECK1-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8315// CHECK1-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 8316// CHECK1-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 8317// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [4 x ptr], align 8318// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [4 x ptr], align 8319// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [4 x ptr], align 8320// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4321// CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8322// CHECK1-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8323// CHECK1-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8324// CHECK1-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 8325// CHECK1-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 8326// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8327// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 8328// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 8329// CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 8330// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0331// CHECK1-NEXT:    store ptr [[TMP0]], ptr [[TMP4]], align 8332// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0333// CHECK1-NEXT:    store ptr [[TMP0]], ptr [[TMP5]], align 8334// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0335// CHECK1-NEXT:    store ptr null, ptr [[TMP6]], align 8336// CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1337// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[TMP7]], align 8338// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1339// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[TMP8]], align 8340// CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1341// CHECK1-NEXT:    store ptr null, ptr [[TMP9]], align 8342// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2343// CHECK1-NEXT:    store ptr [[TMP2]], ptr [[TMP10]], align 8344// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2345// CHECK1-NEXT:    store ptr [[TMP2]], ptr [[TMP11]], align 8346// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2347// CHECK1-NEXT:    store ptr null, ptr [[TMP12]], align 8348// CHECK1-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3349// CHECK1-NEXT:    store ptr [[TMP3]], ptr [[TMP13]], align 8350// CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3351// CHECK1-NEXT:    store ptr [[TMP3]], ptr [[TMP14]], align 8352// CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3353// CHECK1-NEXT:    store ptr null, ptr [[TMP15]], align 8354// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0355// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0356// CHECK1-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0357// CHECK1-NEXT:    store i32 3, ptr [[TMP18]], align 4358// CHECK1-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1359// CHECK1-NEXT:    store i32 4, ptr [[TMP19]], align 4360// CHECK1-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2361// CHECK1-NEXT:    store ptr [[TMP16]], ptr [[TMP20]], align 8362// CHECK1-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3363// CHECK1-NEXT:    store ptr [[TMP17]], ptr [[TMP21]], align 8364// CHECK1-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4365// CHECK1-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP22]], align 8366// CHECK1-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5367// CHECK1-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP23]], align 8368// CHECK1-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6369// CHECK1-NEXT:    store ptr null, ptr [[TMP24]], align 8370// CHECK1-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7371// CHECK1-NEXT:    store ptr null, ptr [[TMP25]], align 8372// CHECK1-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8373// CHECK1-NEXT:    store i64 4571424, ptr [[TMP26]], align 8374// CHECK1-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9375// CHECK1-NEXT:    store i64 0, ptr [[TMP27]], align 8376// CHECK1-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10377// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP28]], align 4378// CHECK1-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11379// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP29]], align 4380// CHECK1-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12381// CHECK1-NEXT:    store i32 0, ptr [[TMP30]], align 4382// CHECK1-NEXT:    [[TMP31:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18static_not_chunkedPfS_S_S__l68.region_id, ptr [[KERNEL_ARGS]])383// CHECK1-NEXT:    [[TMP32:%.*]] = icmp ne i32 [[TMP31]], 0384// CHECK1-NEXT:    br i1 [[TMP32]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]385// CHECK1:       omp_offload.failed:386// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18static_not_chunkedPfS_S_S__l68(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], ptr [[TMP3]]) #[[ATTR2]]387// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]388// CHECK1:       omp_offload.cont:389// CHECK1-NEXT:    ret void390//391//392// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18static_not_chunkedPfS_S_S__l68393// CHECK1-SAME: (ptr noundef [[A:%.*]], ptr noundef [[B:%.*]], ptr noundef [[C:%.*]], ptr noundef [[D:%.*]]) #[[ATTR1]] {394// CHECK1-NEXT:  entry:395// CHECK1-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8396// CHECK1-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8397// CHECK1-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 8398// CHECK1-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 8399// CHECK1-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8400// CHECK1-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8401// CHECK1-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 8402// CHECK1-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 8403// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 4, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18static_not_chunkedPfS_S_S__l68.omp_outlined, ptr [[A_ADDR]], ptr [[B_ADDR]], ptr [[C_ADDR]], ptr [[D_ADDR]])404// CHECK1-NEXT:    ret void405//406//407// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18static_not_chunkedPfS_S_S__l68.omp_outlined408// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[A:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[B:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[C:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[D:%.*]]) #[[ATTR1]] {409// CHECK1-NEXT:  entry:410// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8411// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8412// CHECK1-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8413// CHECK1-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8414// CHECK1-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 8415// CHECK1-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 8416// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4417// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4418// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4419// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4420// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4421// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4422// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4423// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8424// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8425// CHECK1-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8426// CHECK1-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8427// CHECK1-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 8428// CHECK1-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 8429// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8430// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 8431// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 8432// CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 8433// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4434// CHECK1-NEXT:    store i32 4571423, ptr [[DOTOMP_UB]], align 4435// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4436// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4437// CHECK1-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8438// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 4439// CHECK1-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP5]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)440// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4441// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP6]], 4571423442// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]443// CHECK1:       cond.true:444// CHECK1-NEXT:    br label [[COND_END:%.*]]445// CHECK1:       cond.false:446// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4447// CHECK1-NEXT:    br label [[COND_END]]448// CHECK1:       cond.end:449// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 4571423, [[COND_TRUE]] ], [ [[TMP7]], [[COND_FALSE]] ]450// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4451// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4452// CHECK1-NEXT:    store i32 [[TMP8]], ptr [[DOTOMP_IV]], align 4453// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]454// CHECK1:       omp.inner.for.cond:455// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4456// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4457// CHECK1-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP9]], [[TMP10]]458// CHECK1-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]459// CHECK1:       omp.inner.for.body:460// CHECK1-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4461// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP11]], 7462// CHECK1-NEXT:    [[SUB:%.*]] = sub nsw i32 32000000, [[MUL]]463// CHECK1-NEXT:    store i32 [[SUB]], ptr [[I]], align 4464// CHECK1-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[TMP1]], align 8465// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[I]], align 4466// CHECK1-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP13]] to i64467// CHECK1-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds float, ptr [[TMP12]], i64 [[IDXPROM]]468// CHECK1-NEXT:    [[TMP14:%.*]] = load float, ptr [[ARRAYIDX]], align 4469// CHECK1-NEXT:    [[TMP15:%.*]] = load ptr, ptr [[TMP2]], align 8470// CHECK1-NEXT:    [[TMP16:%.*]] = load i32, ptr [[I]], align 4471// CHECK1-NEXT:    [[IDXPROM2:%.*]] = sext i32 [[TMP16]] to i64472// CHECK1-NEXT:    [[ARRAYIDX3:%.*]] = getelementptr inbounds float, ptr [[TMP15]], i64 [[IDXPROM2]]473// CHECK1-NEXT:    [[TMP17:%.*]] = load float, ptr [[ARRAYIDX3]], align 4474// CHECK1-NEXT:    [[MUL4:%.*]] = fmul float [[TMP14]], [[TMP17]]475// CHECK1-NEXT:    [[TMP18:%.*]] = load ptr, ptr [[TMP3]], align 8476// CHECK1-NEXT:    [[TMP19:%.*]] = load i32, ptr [[I]], align 4477// CHECK1-NEXT:    [[IDXPROM5:%.*]] = sext i32 [[TMP19]] to i64478// CHECK1-NEXT:    [[ARRAYIDX6:%.*]] = getelementptr inbounds float, ptr [[TMP18]], i64 [[IDXPROM5]]479// CHECK1-NEXT:    [[TMP20:%.*]] = load float, ptr [[ARRAYIDX6]], align 4480// CHECK1-NEXT:    [[MUL7:%.*]] = fmul float [[MUL4]], [[TMP20]]481// CHECK1-NEXT:    [[TMP21:%.*]] = load ptr, ptr [[TMP0]], align 8482// CHECK1-NEXT:    [[TMP22:%.*]] = load i32, ptr [[I]], align 4483// CHECK1-NEXT:    [[IDXPROM8:%.*]] = sext i32 [[TMP22]] to i64484// CHECK1-NEXT:    [[ARRAYIDX9:%.*]] = getelementptr inbounds float, ptr [[TMP21]], i64 [[IDXPROM8]]485// CHECK1-NEXT:    store float [[MUL7]], ptr [[ARRAYIDX9]], align 4486// CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]487// CHECK1:       omp.body.continue:488// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]489// CHECK1:       omp.inner.for.inc:490// CHECK1-NEXT:    [[TMP23:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4491// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP23]], 1492// CHECK1-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4493// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]494// CHECK1:       omp.inner.for.end:495// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]496// CHECK1:       omp.loop.exit:497// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP5]])498// CHECK1-NEXT:    ret void499//500//501// CHECK1-LABEL: define {{[^@]+}}@_Z14static_chunkedPfS_S_S_502// CHECK1-SAME: (ptr noundef [[A:%.*]], ptr noundef [[B:%.*]], ptr noundef [[C:%.*]], ptr noundef [[D:%.*]]) #[[ATTR0]] {503// CHECK1-NEXT:  entry:504// CHECK1-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8505// CHECK1-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8506// CHECK1-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 8507// CHECK1-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 8508// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [4 x ptr], align 8509// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [4 x ptr], align 8510// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [4 x ptr], align 8511// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4512// CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8513// CHECK1-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8514// CHECK1-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8515// CHECK1-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 8516// CHECK1-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 8517// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8518// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 8519// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 8520// CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 8521// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0522// CHECK1-NEXT:    store ptr [[TMP0]], ptr [[TMP4]], align 8523// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0524// CHECK1-NEXT:    store ptr [[TMP0]], ptr [[TMP5]], align 8525// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0526// CHECK1-NEXT:    store ptr null, ptr [[TMP6]], align 8527// CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1528// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[TMP7]], align 8529// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1530// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[TMP8]], align 8531// CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1532// CHECK1-NEXT:    store ptr null, ptr [[TMP9]], align 8533// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2534// CHECK1-NEXT:    store ptr [[TMP2]], ptr [[TMP10]], align 8535// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2536// CHECK1-NEXT:    store ptr [[TMP2]], ptr [[TMP11]], align 8537// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2538// CHECK1-NEXT:    store ptr null, ptr [[TMP12]], align 8539// CHECK1-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3540// CHECK1-NEXT:    store ptr [[TMP3]], ptr [[TMP13]], align 8541// CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3542// CHECK1-NEXT:    store ptr [[TMP3]], ptr [[TMP14]], align 8543// CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3544// CHECK1-NEXT:    store ptr null, ptr [[TMP15]], align 8545// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0546// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0547// CHECK1-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0548// CHECK1-NEXT:    store i32 3, ptr [[TMP18]], align 4549// CHECK1-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1550// CHECK1-NEXT:    store i32 4, ptr [[TMP19]], align 4551// CHECK1-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2552// CHECK1-NEXT:    store ptr [[TMP16]], ptr [[TMP20]], align 8553// CHECK1-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3554// CHECK1-NEXT:    store ptr [[TMP17]], ptr [[TMP21]], align 8555// CHECK1-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4556// CHECK1-NEXT:    store ptr @.offload_sizes.3, ptr [[TMP22]], align 8557// CHECK1-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5558// CHECK1-NEXT:    store ptr @.offload_maptypes.4, ptr [[TMP23]], align 8559// CHECK1-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6560// CHECK1-NEXT:    store ptr null, ptr [[TMP24]], align 8561// CHECK1-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7562// CHECK1-NEXT:    store ptr null, ptr [[TMP25]], align 8563// CHECK1-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8564// CHECK1-NEXT:    store i64 16908289, ptr [[TMP26]], align 8565// CHECK1-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9566// CHECK1-NEXT:    store i64 0, ptr [[TMP27]], align 8567// CHECK1-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10568// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP28]], align 4569// CHECK1-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11570// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP29]], align 4571// CHECK1-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12572// CHECK1-NEXT:    store i32 0, ptr [[TMP30]], align 4573// CHECK1-NEXT:    [[TMP31:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z14static_chunkedPfS_S_S__l80.region_id, ptr [[KERNEL_ARGS]])574// CHECK1-NEXT:    [[TMP32:%.*]] = icmp ne i32 [[TMP31]], 0575// CHECK1-NEXT:    br i1 [[TMP32]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]576// CHECK1:       omp_offload.failed:577// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z14static_chunkedPfS_S_S__l80(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], ptr [[TMP3]]) #[[ATTR2]]578// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]579// CHECK1:       omp_offload.cont:580// CHECK1-NEXT:    ret void581//582//583// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z14static_chunkedPfS_S_S__l80584// CHECK1-SAME: (ptr noundef [[A:%.*]], ptr noundef [[B:%.*]], ptr noundef [[C:%.*]], ptr noundef [[D:%.*]]) #[[ATTR1]] {585// CHECK1-NEXT:  entry:586// CHECK1-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8587// CHECK1-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8588// CHECK1-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 8589// CHECK1-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 8590// CHECK1-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8591// CHECK1-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8592// CHECK1-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 8593// CHECK1-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 8594// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 4, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z14static_chunkedPfS_S_S__l80.omp_outlined, ptr [[A_ADDR]], ptr [[B_ADDR]], ptr [[C_ADDR]], ptr [[D_ADDR]])595// CHECK1-NEXT:    ret void596//597//598// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z14static_chunkedPfS_S_S__l80.omp_outlined599// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[A:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[B:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[C:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[D:%.*]]) #[[ATTR1]] {600// CHECK1-NEXT:  entry:601// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8602// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8603// CHECK1-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8604// CHECK1-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8605// CHECK1-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 8606// CHECK1-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 8607// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4608// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4609// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4610// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4611// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4612// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4613// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4614// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8615// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8616// CHECK1-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8617// CHECK1-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8618// CHECK1-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 8619// CHECK1-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 8620// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8621// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 8622// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 8623// CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 8624// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4625// CHECK1-NEXT:    store i32 16908288, ptr [[DOTOMP_UB]], align 4626// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4627// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4628// CHECK1-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8629// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 4630// CHECK1-NEXT:    call void @__kmpc_for_static_init_4u(ptr @[[GLOB1]], i32 [[TMP5]], i32 91, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 5)631// CHECK1-NEXT:    br label [[OMP_DISPATCH_COND:%.*]]632// CHECK1:       omp.dispatch.cond:633// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4634// CHECK1-NEXT:    [[CMP:%.*]] = icmp ugt i32 [[TMP6]], 16908288635// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]636// CHECK1:       cond.true:637// CHECK1-NEXT:    br label [[COND_END:%.*]]638// CHECK1:       cond.false:639// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4640// CHECK1-NEXT:    br label [[COND_END]]641// CHECK1:       cond.end:642// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 16908288, [[COND_TRUE]] ], [ [[TMP7]], [[COND_FALSE]] ]643// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4644// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4645// CHECK1-NEXT:    store i32 [[TMP8]], ptr [[DOTOMP_IV]], align 4646// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4647// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4648// CHECK1-NEXT:    [[CMP1:%.*]] = icmp ule i32 [[TMP9]], [[TMP10]]649// CHECK1-NEXT:    br i1 [[CMP1]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]650// CHECK1:       omp.dispatch.body:651// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]652// CHECK1:       omp.inner.for.cond:653// CHECK1-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP10:![0-9]+]]654// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP10]]655// CHECK1-NEXT:    [[CMP2:%.*]] = icmp ule i32 [[TMP11]], [[TMP12]]656// CHECK1-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]657// CHECK1:       omp.inner.for.body:658// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP10]]659// CHECK1-NEXT:    [[MUL:%.*]] = mul i32 [[TMP13]], 127660// CHECK1-NEXT:    [[ADD:%.*]] = add i32 131071, [[MUL]]661// CHECK1-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP10]]662// CHECK1-NEXT:    [[TMP14:%.*]] = load ptr, ptr [[TMP1]], align 8, !llvm.access.group [[ACC_GRP10]]663// CHECK1-NEXT:    [[TMP15:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP10]]664// CHECK1-NEXT:    [[IDXPROM:%.*]] = zext i32 [[TMP15]] to i64665// CHECK1-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds nuw float, ptr [[TMP14]], i64 [[IDXPROM]]666// CHECK1-NEXT:    [[TMP16:%.*]] = load float, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP10]]667// CHECK1-NEXT:    [[TMP17:%.*]] = load ptr, ptr [[TMP2]], align 8, !llvm.access.group [[ACC_GRP10]]668// CHECK1-NEXT:    [[TMP18:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP10]]669// CHECK1-NEXT:    [[IDXPROM3:%.*]] = zext i32 [[TMP18]] to i64670// CHECK1-NEXT:    [[ARRAYIDX4:%.*]] = getelementptr inbounds nuw float, ptr [[TMP17]], i64 [[IDXPROM3]]671// CHECK1-NEXT:    [[TMP19:%.*]] = load float, ptr [[ARRAYIDX4]], align 4, !llvm.access.group [[ACC_GRP10]]672// CHECK1-NEXT:    [[MUL5:%.*]] = fmul float [[TMP16]], [[TMP19]]673// CHECK1-NEXT:    [[TMP20:%.*]] = load ptr, ptr [[TMP3]], align 8, !llvm.access.group [[ACC_GRP10]]674// CHECK1-NEXT:    [[TMP21:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP10]]675// CHECK1-NEXT:    [[IDXPROM6:%.*]] = zext i32 [[TMP21]] to i64676// CHECK1-NEXT:    [[ARRAYIDX7:%.*]] = getelementptr inbounds nuw float, ptr [[TMP20]], i64 [[IDXPROM6]]677// CHECK1-NEXT:    [[TMP22:%.*]] = load float, ptr [[ARRAYIDX7]], align 4, !llvm.access.group [[ACC_GRP10]]678// CHECK1-NEXT:    [[MUL8:%.*]] = fmul float [[MUL5]], [[TMP22]]679// CHECK1-NEXT:    [[TMP23:%.*]] = load ptr, ptr [[TMP0]], align 8, !llvm.access.group [[ACC_GRP10]]680// CHECK1-NEXT:    [[TMP24:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP10]]681// CHECK1-NEXT:    [[IDXPROM9:%.*]] = zext i32 [[TMP24]] to i64682// CHECK1-NEXT:    [[ARRAYIDX10:%.*]] = getelementptr inbounds nuw float, ptr [[TMP23]], i64 [[IDXPROM9]]683// CHECK1-NEXT:    store float [[MUL8]], ptr [[ARRAYIDX10]], align 4, !llvm.access.group [[ACC_GRP10]]684// CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]685// CHECK1:       omp.body.continue:686// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]687// CHECK1:       omp.inner.for.inc:688// CHECK1-NEXT:    [[TMP25:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP10]]689// CHECK1-NEXT:    [[ADD11:%.*]] = add i32 [[TMP25]], 1690// CHECK1-NEXT:    store i32 [[ADD11]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP10]]691// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP11:![0-9]+]]692// CHECK1:       omp.inner.for.end:693// CHECK1-NEXT:    br label [[OMP_DISPATCH_INC:%.*]]694// CHECK1:       omp.dispatch.inc:695// CHECK1-NEXT:    [[TMP26:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4696// CHECK1-NEXT:    [[TMP27:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4697// CHECK1-NEXT:    [[ADD12:%.*]] = add i32 [[TMP26]], [[TMP27]]698// CHECK1-NEXT:    store i32 [[ADD12]], ptr [[DOTOMP_LB]], align 4699// CHECK1-NEXT:    [[TMP28:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4700// CHECK1-NEXT:    [[TMP29:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4701// CHECK1-NEXT:    [[ADD13:%.*]] = add i32 [[TMP28]], [[TMP29]]702// CHECK1-NEXT:    store i32 [[ADD13]], ptr [[DOTOMP_UB]], align 4703// CHECK1-NEXT:    br label [[OMP_DISPATCH_COND]]704// CHECK1:       omp.dispatch.end:705// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP5]])706// CHECK1-NEXT:    ret void707//708//709// CHECK1-LABEL: define {{[^@]+}}@_Z12test_precondv710// CHECK1-SAME: () #[[ATTR0]] {711// CHECK1-NEXT:  entry:712// CHECK1-NEXT:    [[A:%.*]] = alloca i8, align 1713// CHECK1-NEXT:    [[A_CASTED:%.*]] = alloca i64, align 8714// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8715// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8716// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8717// CHECK1-NEXT:    [[TMP:%.*]] = alloca i8, align 1718// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i8, align 1719// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4720// CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8721// CHECK1-NEXT:    store i8 0, ptr [[A]], align 1722// CHECK1-NEXT:    [[TMP0:%.*]] = load i8, ptr [[A]], align 1723// CHECK1-NEXT:    store i8 [[TMP0]], ptr [[A_CASTED]], align 1724// CHECK1-NEXT:    [[TMP1:%.*]] = load i64, ptr [[A_CASTED]], align 8725// CHECK1-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0726// CHECK1-NEXT:    store i64 [[TMP1]], ptr [[TMP2]], align 8727// CHECK1-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0728// CHECK1-NEXT:    store i64 [[TMP1]], ptr [[TMP3]], align 8729// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0730// CHECK1-NEXT:    store ptr null, ptr [[TMP4]], align 8731// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0732// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0733// CHECK1-NEXT:    [[TMP7:%.*]] = load i8, ptr [[A]], align 1734// CHECK1-NEXT:    store i8 [[TMP7]], ptr [[DOTCAPTURE_EXPR_]], align 1735// CHECK1-NEXT:    [[TMP8:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 1736// CHECK1-NEXT:    [[CONV:%.*]] = sext i8 [[TMP8]] to i32737// CHECK1-NEXT:    [[SUB:%.*]] = sub i32 10, [[CONV]]738// CHECK1-NEXT:    [[SUB2:%.*]] = sub i32 [[SUB]], 1739// CHECK1-NEXT:    [[ADD:%.*]] = add i32 [[SUB2]], 1740// CHECK1-NEXT:    [[DIV:%.*]] = udiv i32 [[ADD]], 1741// CHECK1-NEXT:    [[SUB3:%.*]] = sub nsw i32 [[DIV]], 1742// CHECK1-NEXT:    store i32 [[SUB3]], ptr [[DOTCAPTURE_EXPR_1]], align 4743// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4744// CHECK1-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP9]], 1745// CHECK1-NEXT:    [[TMP10:%.*]] = zext i32 [[ADD4]] to i64746// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0747// CHECK1-NEXT:    store i32 3, ptr [[TMP11]], align 4748// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1749// CHECK1-NEXT:    store i32 1, ptr [[TMP12]], align 4750// CHECK1-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2751// CHECK1-NEXT:    store ptr [[TMP5]], ptr [[TMP13]], align 8752// CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3753// CHECK1-NEXT:    store ptr [[TMP6]], ptr [[TMP14]], align 8754// CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4755// CHECK1-NEXT:    store ptr @.offload_sizes.5, ptr [[TMP15]], align 8756// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5757// CHECK1-NEXT:    store ptr @.offload_maptypes.6, ptr [[TMP16]], align 8758// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6759// CHECK1-NEXT:    store ptr null, ptr [[TMP17]], align 8760// CHECK1-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7761// CHECK1-NEXT:    store ptr null, ptr [[TMP18]], align 8762// CHECK1-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8763// CHECK1-NEXT:    store i64 [[TMP10]], ptr [[TMP19]], align 8764// CHECK1-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9765// CHECK1-NEXT:    store i64 0, ptr [[TMP20]], align 8766// CHECK1-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10767// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP21]], align 4768// CHECK1-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11769// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP22]], align 4770// CHECK1-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12771// CHECK1-NEXT:    store i32 0, ptr [[TMP23]], align 4772// CHECK1-NEXT:    [[TMP24:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z12test_precondv_l92.region_id, ptr [[KERNEL_ARGS]])773// CHECK1-NEXT:    [[TMP25:%.*]] = icmp ne i32 [[TMP24]], 0774// CHECK1-NEXT:    br i1 [[TMP25]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]775// CHECK1:       omp_offload.failed:776// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z12test_precondv_l92(i64 [[TMP1]]) #[[ATTR2]]777// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]778// CHECK1:       omp_offload.cont:779// CHECK1-NEXT:    ret void780//781//782// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z12test_precondv_l92783// CHECK1-SAME: (i64 noundef [[A:%.*]]) #[[ATTR1]] {784// CHECK1-NEXT:  entry:785// CHECK1-NEXT:    [[A_ADDR:%.*]] = alloca i64, align 8786// CHECK1-NEXT:    store i64 [[A]], ptr [[A_ADDR]], align 8787// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z12test_precondv_l92.omp_outlined, ptr [[A_ADDR]])788// CHECK1-NEXT:    ret void789//790//791// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z12test_precondv_l92.omp_outlined792// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 1 dereferenceable(1) [[A:%.*]]) #[[ATTR1]] {793// CHECK1-NEXT:  entry:794// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8795// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8796// CHECK1-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8797// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4798// CHECK1-NEXT:    [[TMP:%.*]] = alloca i8, align 1799// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i8, align 1800// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4801// CHECK1-NEXT:    [[I:%.*]] = alloca i8, align 1802// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4803// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4804// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4805// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4806// CHECK1-NEXT:    [[I5:%.*]] = alloca i8, align 1807// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8808// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8809// CHECK1-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8810// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8811// CHECK1-NEXT:    [[TMP1:%.*]] = load i8, ptr [[TMP0]], align 1812// CHECK1-NEXT:    store i8 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 1813// CHECK1-NEXT:    [[TMP2:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 1814// CHECK1-NEXT:    [[CONV:%.*]] = sext i8 [[TMP2]] to i32815// CHECK1-NEXT:    [[SUB:%.*]] = sub i32 10, [[CONV]]816// CHECK1-NEXT:    [[SUB2:%.*]] = sub i32 [[SUB]], 1817// CHECK1-NEXT:    [[ADD:%.*]] = add i32 [[SUB2]], 1818// CHECK1-NEXT:    [[DIV:%.*]] = udiv i32 [[ADD]], 1819// CHECK1-NEXT:    [[SUB3:%.*]] = sub nsw i32 [[DIV]], 1820// CHECK1-NEXT:    store i32 [[SUB3]], ptr [[DOTCAPTURE_EXPR_1]], align 4821// CHECK1-NEXT:    [[TMP3:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 1822// CHECK1-NEXT:    store i8 [[TMP3]], ptr [[I]], align 1823// CHECK1-NEXT:    [[TMP4:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 1824// CHECK1-NEXT:    [[CONV4:%.*]] = sext i8 [[TMP4]] to i32825// CHECK1-NEXT:    [[CMP:%.*]] = icmp slt i32 [[CONV4]], 10826// CHECK1-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]827// CHECK1:       omp.precond.then:828// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4829// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4830// CHECK1-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_UB]], align 4831// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4832// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4833// CHECK1-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8834// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 4835// CHECK1-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP7]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)836// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4837// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4838// CHECK1-NEXT:    [[CMP6:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]839// CHECK1-NEXT:    br i1 [[CMP6]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]840// CHECK1:       cond.true:841// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4842// CHECK1-NEXT:    br label [[COND_END:%.*]]843// CHECK1:       cond.false:844// CHECK1-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4845// CHECK1-NEXT:    br label [[COND_END]]846// CHECK1:       cond.end:847// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]848// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4849// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4850// CHECK1-NEXT:    store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 4851// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]852// CHECK1:       omp.inner.for.cond:853// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4854// CHECK1-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4855// CHECK1-NEXT:    [[CMP7:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]856// CHECK1-NEXT:    br i1 [[CMP7]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]857// CHECK1:       omp.inner.for.body:858// CHECK1-NEXT:    [[TMP15:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 1859// CHECK1-NEXT:    [[CONV8:%.*]] = sext i8 [[TMP15]] to i32860// CHECK1-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4861// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP16]], 1862// CHECK1-NEXT:    [[ADD9:%.*]] = add nsw i32 [[CONV8]], [[MUL]]863// CHECK1-NEXT:    [[CONV10:%.*]] = trunc i32 [[ADD9]] to i8864// CHECK1-NEXT:    store i8 [[CONV10]], ptr [[I5]], align 1865// CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]866// CHECK1:       omp.body.continue:867// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]868// CHECK1:       omp.inner.for.inc:869// CHECK1-NEXT:    [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4870// CHECK1-NEXT:    [[ADD11:%.*]] = add nsw i32 [[TMP17]], 1871// CHECK1-NEXT:    store i32 [[ADD11]], ptr [[DOTOMP_IV]], align 4872// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]873// CHECK1:       omp.inner.for.end:874// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]875// CHECK1:       omp.loop.exit:876// CHECK1-NEXT:    [[TMP18:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8877// CHECK1-NEXT:    [[TMP19:%.*]] = load i32, ptr [[TMP18]], align 4878// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP19]])879// CHECK1-NEXT:    br label [[OMP_PRECOND_END]]880// CHECK1:       omp.precond.end:881// CHECK1-NEXT:    ret void882//883//884// CHECK1-LABEL: define {{[^@]+}}@_Z4fintv885// CHECK1-SAME: () #[[ATTR0]] {886// CHECK1-NEXT:  entry:887// CHECK1-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_Z9ftemplateIiET_v()888// CHECK1-NEXT:    ret i32 [[CALL]]889//890//891// CHECK1-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_v892// CHECK1-SAME: () #[[ATTR0]] comdat {893// CHECK1-NEXT:  entry:894// CHECK1-NEXT:    [[AA:%.*]] = alloca i16, align 2895// CHECK1-NEXT:    [[AA_CASTED:%.*]] = alloca i64, align 8896// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8897// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8898// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8899// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4900// CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8901// CHECK1-NEXT:    store i16 0, ptr [[AA]], align 2902// CHECK1-NEXT:    [[TMP0:%.*]] = load i16, ptr [[AA]], align 2903// CHECK1-NEXT:    store i16 [[TMP0]], ptr [[AA_CASTED]], align 2904// CHECK1-NEXT:    [[TMP1:%.*]] = load i64, ptr [[AA_CASTED]], align 8905// CHECK1-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0906// CHECK1-NEXT:    store i64 [[TMP1]], ptr [[TMP2]], align 8907// CHECK1-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0908// CHECK1-NEXT:    store i64 [[TMP1]], ptr [[TMP3]], align 8909// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0910// CHECK1-NEXT:    store ptr null, ptr [[TMP4]], align 8911// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0912// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0913// CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0914// CHECK1-NEXT:    store i32 3, ptr [[TMP7]], align 4915// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1916// CHECK1-NEXT:    store i32 1, ptr [[TMP8]], align 4917// CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2918// CHECK1-NEXT:    store ptr [[TMP5]], ptr [[TMP9]], align 8919// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3920// CHECK1-NEXT:    store ptr [[TMP6]], ptr [[TMP10]], align 8921// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4922// CHECK1-NEXT:    store ptr @.offload_sizes.7, ptr [[TMP11]], align 8923// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5924// CHECK1-NEXT:    store ptr @.offload_maptypes.8, ptr [[TMP12]], align 8925// CHECK1-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6926// CHECK1-NEXT:    store ptr null, ptr [[TMP13]], align 8927// CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7928// CHECK1-NEXT:    store ptr null, ptr [[TMP14]], align 8929// CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8930// CHECK1-NEXT:    store i64 100, ptr [[TMP15]], align 8931// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9932// CHECK1-NEXT:    store i64 0, ptr [[TMP16]], align 8933// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10934// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4935// CHECK1-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11936// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP18]], align 4937// CHECK1-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12938// CHECK1-NEXT:    store i32 0, ptr [[TMP19]], align 4939// CHECK1-NEXT:    [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_v_l108.region_id, ptr [[KERNEL_ARGS]])940// CHECK1-NEXT:    [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0941// CHECK1-NEXT:    br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]942// CHECK1:       omp_offload.failed:943// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_v_l108(i64 [[TMP1]]) #[[ATTR2]]944// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]945// CHECK1:       omp_offload.cont:946// CHECK1-NEXT:    ret i32 0947//948//949// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_v_l108950// CHECK1-SAME: (i64 noundef [[AA:%.*]]) #[[ATTR1]] {951// CHECK1-NEXT:  entry:952// CHECK1-NEXT:    [[AA_ADDR:%.*]] = alloca i64, align 8953// CHECK1-NEXT:    store i64 [[AA]], ptr [[AA_ADDR]], align 8954// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_v_l108.omp_outlined, ptr [[AA_ADDR]])955// CHECK1-NEXT:    ret void956//957//958// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_v_l108.omp_outlined959// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[AA:%.*]]) #[[ATTR1]] {960// CHECK1-NEXT:  entry:961// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8962// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8963// CHECK1-NEXT:    [[AA_ADDR:%.*]] = alloca ptr, align 8964// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4965// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4966// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4967// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4968// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4969// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4970// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4971// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8972// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8973// CHECK1-NEXT:    store ptr [[AA]], ptr [[AA_ADDR]], align 8974// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 8975// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4976// CHECK1-NEXT:    store i32 99, ptr [[DOTOMP_UB]], align 4977// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4978// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4979// CHECK1-NEXT:    [[TMP1:%.*]] = load i16, ptr [[TMP0]], align 2980// CHECK1-NEXT:    [[CONV:%.*]] = sext i16 [[TMP1]] to i32981// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8982// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4983// CHECK1-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP3]], i32 91, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 [[CONV]])984// CHECK1-NEXT:    br label [[OMP_DISPATCH_COND:%.*]]985// CHECK1:       omp.dispatch.cond:986// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4987// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 99988// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]989// CHECK1:       cond.true:990// CHECK1-NEXT:    br label [[COND_END:%.*]]991// CHECK1:       cond.false:992// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4993// CHECK1-NEXT:    br label [[COND_END]]994// CHECK1:       cond.end:995// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]996// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4997// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4998// CHECK1-NEXT:    store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 4999// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41000// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41001// CHECK1-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]1002// CHECK1-NEXT:    br i1 [[CMP1]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]1003// CHECK1:       omp.dispatch.body:1004// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1005// CHECK1:       omp.inner.for.cond:1006// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP13:![0-9]+]]1007// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP13]]1008// CHECK1-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP9]], [[TMP10]]1009// CHECK1-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1010// CHECK1:       omp.inner.for.body:1011// CHECK1-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP13]]1012// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP11]], 11013// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1014// CHECK1-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP13]]1015// CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1016// CHECK1:       omp.body.continue:1017// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1018// CHECK1:       omp.inner.for.inc:1019// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP13]]1020// CHECK1-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP12]], 11021// CHECK1-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP13]]1022// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP14:![0-9]+]]1023// CHECK1:       omp.inner.for.end:1024// CHECK1-NEXT:    br label [[OMP_DISPATCH_INC:%.*]]1025// CHECK1:       omp.dispatch.inc:1026// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41027// CHECK1-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41028// CHECK1-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP13]], [[TMP14]]1029// CHECK1-NEXT:    store i32 [[ADD4]], ptr [[DOTOMP_LB]], align 41030// CHECK1-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41031// CHECK1-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41032// CHECK1-NEXT:    [[ADD5:%.*]] = add nsw i32 [[TMP15]], [[TMP16]]1033// CHECK1-NEXT:    store i32 [[ADD5]], ptr [[DOTOMP_UB]], align 41034// CHECK1-NEXT:    br label [[OMP_DISPATCH_COND]]1035// CHECK1:       omp.dispatch.end:1036// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP3]])1037// CHECK1-NEXT:    ret void1038//1039//1040// CHECK3-LABEL: define {{[^@]+}}@_Z23without_schedule_clausePfS_S_S_1041// CHECK3-SAME: (ptr noundef [[A:%.*]], ptr noundef [[B:%.*]], ptr noundef [[C:%.*]], ptr noundef [[D:%.*]]) #[[ATTR0:[0-9]+]] {1042// CHECK3-NEXT:  entry:1043// CHECK3-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 41044// CHECK3-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 41045// CHECK3-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 41046// CHECK3-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 41047// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [4 x ptr], align 41048// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [4 x ptr], align 41049// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [4 x ptr], align 41050// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 41051// CHECK3-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81052// CHECK3-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 41053// CHECK3-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 41054// CHECK3-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 41055// CHECK3-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 41056// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41057// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 41058// CHECK3-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 41059// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 41060// CHECK3-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01061// CHECK3-NEXT:    store ptr [[TMP0]], ptr [[TMP4]], align 41062// CHECK3-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01063// CHECK3-NEXT:    store ptr [[TMP0]], ptr [[TMP5]], align 41064// CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01065// CHECK3-NEXT:    store ptr null, ptr [[TMP6]], align 41066// CHECK3-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11067// CHECK3-NEXT:    store ptr [[TMP1]], ptr [[TMP7]], align 41068// CHECK3-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11069// CHECK3-NEXT:    store ptr [[TMP1]], ptr [[TMP8]], align 41070// CHECK3-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 11071// CHECK3-NEXT:    store ptr null, ptr [[TMP9]], align 41072// CHECK3-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21073// CHECK3-NEXT:    store ptr [[TMP2]], ptr [[TMP10]], align 41074// CHECK3-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21075// CHECK3-NEXT:    store ptr [[TMP2]], ptr [[TMP11]], align 41076// CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 21077// CHECK3-NEXT:    store ptr null, ptr [[TMP12]], align 41078// CHECK3-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 31079// CHECK3-NEXT:    store ptr [[TMP3]], ptr [[TMP13]], align 41080// CHECK3-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 31081// CHECK3-NEXT:    store ptr [[TMP3]], ptr [[TMP14]], align 41082// CHECK3-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 31083// CHECK3-NEXT:    store ptr null, ptr [[TMP15]], align 41084// CHECK3-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01085// CHECK3-NEXT:    [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01086// CHECK3-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01087// CHECK3-NEXT:    store i32 3, ptr [[TMP18]], align 41088// CHECK3-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11089// CHECK3-NEXT:    store i32 4, ptr [[TMP19]], align 41090// CHECK3-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21091// CHECK3-NEXT:    store ptr [[TMP16]], ptr [[TMP20]], align 41092// CHECK3-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31093// CHECK3-NEXT:    store ptr [[TMP17]], ptr [[TMP21]], align 41094// CHECK3-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41095// CHECK3-NEXT:    store ptr @.offload_sizes, ptr [[TMP22]], align 41096// CHECK3-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51097// CHECK3-NEXT:    store ptr @.offload_maptypes, ptr [[TMP23]], align 41098// CHECK3-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61099// CHECK3-NEXT:    store ptr null, ptr [[TMP24]], align 41100// CHECK3-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71101// CHECK3-NEXT:    store ptr null, ptr [[TMP25]], align 41102// CHECK3-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81103// CHECK3-NEXT:    store i64 4571424, ptr [[TMP26]], align 81104// CHECK3-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91105// CHECK3-NEXT:    store i64 0, ptr [[TMP27]], align 81106// CHECK3-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101107// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP28]], align 41108// CHECK3-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111109// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP29]], align 41110// CHECK3-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121111// CHECK3-NEXT:    store i32 0, ptr [[TMP30]], align 41112// CHECK3-NEXT:    [[TMP31:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z23without_schedule_clausePfS_S_S__l56.region_id, ptr [[KERNEL_ARGS]])1113// CHECK3-NEXT:    [[TMP32:%.*]] = icmp ne i32 [[TMP31]], 01114// CHECK3-NEXT:    br i1 [[TMP32]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1115// CHECK3:       omp_offload.failed:1116// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z23without_schedule_clausePfS_S_S__l56(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], ptr [[TMP3]]) #[[ATTR2:[0-9]+]]1117// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT]]1118// CHECK3:       omp_offload.cont:1119// CHECK3-NEXT:    ret void1120//1121//1122// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z23without_schedule_clausePfS_S_S__l561123// CHECK3-SAME: (ptr noundef [[A:%.*]], ptr noundef [[B:%.*]], ptr noundef [[C:%.*]], ptr noundef [[D:%.*]]) #[[ATTR1:[0-9]+]] {1124// CHECK3-NEXT:  entry:1125// CHECK3-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 41126// CHECK3-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 41127// CHECK3-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 41128// CHECK3-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 41129// CHECK3-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 41130// CHECK3-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 41131// CHECK3-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 41132// CHECK3-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 41133// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 4, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z23without_schedule_clausePfS_S_S__l56.omp_outlined, ptr [[A_ADDR]], ptr [[B_ADDR]], ptr [[C_ADDR]], ptr [[D_ADDR]])1134// CHECK3-NEXT:    ret void1135//1136//1137// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z23without_schedule_clausePfS_S_S__l56.omp_outlined1138// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[C:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[D:%.*]]) #[[ATTR1]] {1139// CHECK3-NEXT:  entry:1140// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41141// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41142// CHECK3-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 41143// CHECK3-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 41144// CHECK3-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 41145// CHECK3-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 41146// CHECK3-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41147// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 41148// CHECK3-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 41149// CHECK3-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 41150// CHECK3-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41151// CHECK3-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41152// CHECK3-NEXT:    [[I:%.*]] = alloca i32, align 41153// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41154// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41155// CHECK3-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 41156// CHECK3-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 41157// CHECK3-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 41158// CHECK3-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 41159// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41160// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 41161// CHECK3-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 41162// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 41163// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 41164// CHECK3-NEXT:    store i32 4571423, ptr [[DOTOMP_UB]], align 41165// CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41166// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41167// CHECK3-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41168// CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 41169// CHECK3-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP5]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1170// CHECK3-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41171// CHECK3-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP6]], 45714231172// CHECK3-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1173// CHECK3:       cond.true:1174// CHECK3-NEXT:    br label [[COND_END:%.*]]1175// CHECK3:       cond.false:1176// CHECK3-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41177// CHECK3-NEXT:    br label [[COND_END]]1178// CHECK3:       cond.end:1179// CHECK3-NEXT:    [[COND:%.*]] = phi i32 [ 4571423, [[COND_TRUE]] ], [ [[TMP7]], [[COND_FALSE]] ]1180// CHECK3-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 41181// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41182// CHECK3-NEXT:    store i32 [[TMP8]], ptr [[DOTOMP_IV]], align 41183// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1184// CHECK3:       omp.inner.for.cond:1185// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41186// CHECK3-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41187// CHECK3-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP9]], [[TMP10]]1188// CHECK3-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1189// CHECK3:       omp.inner.for.body:1190// CHECK3-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41191// CHECK3-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP11]], 71192// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 33, [[MUL]]1193// CHECK3-NEXT:    store i32 [[ADD]], ptr [[I]], align 41194// CHECK3-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[TMP1]], align 41195// CHECK3-NEXT:    [[TMP13:%.*]] = load i32, ptr [[I]], align 41196// CHECK3-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds float, ptr [[TMP12]], i32 [[TMP13]]1197// CHECK3-NEXT:    [[TMP14:%.*]] = load float, ptr [[ARRAYIDX]], align 41198// CHECK3-NEXT:    [[TMP15:%.*]] = load ptr, ptr [[TMP2]], align 41199// CHECK3-NEXT:    [[TMP16:%.*]] = load i32, ptr [[I]], align 41200// CHECK3-NEXT:    [[ARRAYIDX2:%.*]] = getelementptr inbounds float, ptr [[TMP15]], i32 [[TMP16]]1201// CHECK3-NEXT:    [[TMP17:%.*]] = load float, ptr [[ARRAYIDX2]], align 41202// CHECK3-NEXT:    [[MUL3:%.*]] = fmul float [[TMP14]], [[TMP17]]1203// CHECK3-NEXT:    [[TMP18:%.*]] = load ptr, ptr [[TMP3]], align 41204// CHECK3-NEXT:    [[TMP19:%.*]] = load i32, ptr [[I]], align 41205// CHECK3-NEXT:    [[ARRAYIDX4:%.*]] = getelementptr inbounds float, ptr [[TMP18]], i32 [[TMP19]]1206// CHECK3-NEXT:    [[TMP20:%.*]] = load float, ptr [[ARRAYIDX4]], align 41207// CHECK3-NEXT:    [[MUL5:%.*]] = fmul float [[MUL3]], [[TMP20]]1208// CHECK3-NEXT:    [[TMP21:%.*]] = load ptr, ptr [[TMP0]], align 41209// CHECK3-NEXT:    [[TMP22:%.*]] = load i32, ptr [[I]], align 41210// CHECK3-NEXT:    [[ARRAYIDX6:%.*]] = getelementptr inbounds float, ptr [[TMP21]], i32 [[TMP22]]1211// CHECK3-NEXT:    store float [[MUL5]], ptr [[ARRAYIDX6]], align 41212// CHECK3-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1213// CHECK3:       omp.body.continue:1214// CHECK3-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1215// CHECK3:       omp.inner.for.inc:1216// CHECK3-NEXT:    [[TMP23:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41217// CHECK3-NEXT:    [[ADD7:%.*]] = add nsw i32 [[TMP23]], 11218// CHECK3-NEXT:    store i32 [[ADD7]], ptr [[DOTOMP_IV]], align 41219// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND]]1220// CHECK3:       omp.inner.for.end:1221// CHECK3-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1222// CHECK3:       omp.loop.exit:1223// CHECK3-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP5]])1224// CHECK3-NEXT:    ret void1225//1226//1227// CHECK3-LABEL: define {{[^@]+}}@_Z18static_not_chunkedPfS_S_S_1228// CHECK3-SAME: (ptr noundef [[A:%.*]], ptr noundef [[B:%.*]], ptr noundef [[C:%.*]], ptr noundef [[D:%.*]]) #[[ATTR0]] {1229// CHECK3-NEXT:  entry:1230// CHECK3-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 41231// CHECK3-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 41232// CHECK3-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 41233// CHECK3-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 41234// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [4 x ptr], align 41235// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [4 x ptr], align 41236// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [4 x ptr], align 41237// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 41238// CHECK3-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81239// CHECK3-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 41240// CHECK3-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 41241// CHECK3-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 41242// CHECK3-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 41243// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41244// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 41245// CHECK3-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 41246// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 41247// CHECK3-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01248// CHECK3-NEXT:    store ptr [[TMP0]], ptr [[TMP4]], align 41249// CHECK3-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01250// CHECK3-NEXT:    store ptr [[TMP0]], ptr [[TMP5]], align 41251// CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01252// CHECK3-NEXT:    store ptr null, ptr [[TMP6]], align 41253// CHECK3-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11254// CHECK3-NEXT:    store ptr [[TMP1]], ptr [[TMP7]], align 41255// CHECK3-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11256// CHECK3-NEXT:    store ptr [[TMP1]], ptr [[TMP8]], align 41257// CHECK3-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 11258// CHECK3-NEXT:    store ptr null, ptr [[TMP9]], align 41259// CHECK3-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21260// CHECK3-NEXT:    store ptr [[TMP2]], ptr [[TMP10]], align 41261// CHECK3-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21262// CHECK3-NEXT:    store ptr [[TMP2]], ptr [[TMP11]], align 41263// CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 21264// CHECK3-NEXT:    store ptr null, ptr [[TMP12]], align 41265// CHECK3-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 31266// CHECK3-NEXT:    store ptr [[TMP3]], ptr [[TMP13]], align 41267// CHECK3-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 31268// CHECK3-NEXT:    store ptr [[TMP3]], ptr [[TMP14]], align 41269// CHECK3-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 31270// CHECK3-NEXT:    store ptr null, ptr [[TMP15]], align 41271// CHECK3-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01272// CHECK3-NEXT:    [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01273// CHECK3-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01274// CHECK3-NEXT:    store i32 3, ptr [[TMP18]], align 41275// CHECK3-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11276// CHECK3-NEXT:    store i32 4, ptr [[TMP19]], align 41277// CHECK3-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21278// CHECK3-NEXT:    store ptr [[TMP16]], ptr [[TMP20]], align 41279// CHECK3-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31280// CHECK3-NEXT:    store ptr [[TMP17]], ptr [[TMP21]], align 41281// CHECK3-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41282// CHECK3-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP22]], align 41283// CHECK3-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51284// CHECK3-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP23]], align 41285// CHECK3-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61286// CHECK3-NEXT:    store ptr null, ptr [[TMP24]], align 41287// CHECK3-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71288// CHECK3-NEXT:    store ptr null, ptr [[TMP25]], align 41289// CHECK3-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81290// CHECK3-NEXT:    store i64 4571424, ptr [[TMP26]], align 81291// CHECK3-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91292// CHECK3-NEXT:    store i64 0, ptr [[TMP27]], align 81293// CHECK3-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101294// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP28]], align 41295// CHECK3-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111296// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP29]], align 41297// CHECK3-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121298// CHECK3-NEXT:    store i32 0, ptr [[TMP30]], align 41299// CHECK3-NEXT:    [[TMP31:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18static_not_chunkedPfS_S_S__l68.region_id, ptr [[KERNEL_ARGS]])1300// CHECK3-NEXT:    [[TMP32:%.*]] = icmp ne i32 [[TMP31]], 01301// CHECK3-NEXT:    br i1 [[TMP32]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1302// CHECK3:       omp_offload.failed:1303// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18static_not_chunkedPfS_S_S__l68(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], ptr [[TMP3]]) #[[ATTR2]]1304// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT]]1305// CHECK3:       omp_offload.cont:1306// CHECK3-NEXT:    ret void1307//1308//1309// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18static_not_chunkedPfS_S_S__l681310// CHECK3-SAME: (ptr noundef [[A:%.*]], ptr noundef [[B:%.*]], ptr noundef [[C:%.*]], ptr noundef [[D:%.*]]) #[[ATTR1]] {1311// CHECK3-NEXT:  entry:1312// CHECK3-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 41313// CHECK3-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 41314// CHECK3-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 41315// CHECK3-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 41316// CHECK3-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 41317// CHECK3-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 41318// CHECK3-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 41319// CHECK3-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 41320// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 4, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18static_not_chunkedPfS_S_S__l68.omp_outlined, ptr [[A_ADDR]], ptr [[B_ADDR]], ptr [[C_ADDR]], ptr [[D_ADDR]])1321// CHECK3-NEXT:    ret void1322//1323//1324// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18static_not_chunkedPfS_S_S__l68.omp_outlined1325// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[C:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[D:%.*]]) #[[ATTR1]] {1326// CHECK3-NEXT:  entry:1327// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41328// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41329// CHECK3-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 41330// CHECK3-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 41331// CHECK3-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 41332// CHECK3-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 41333// CHECK3-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41334// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 41335// CHECK3-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 41336// CHECK3-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 41337// CHECK3-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41338// CHECK3-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41339// CHECK3-NEXT:    [[I:%.*]] = alloca i32, align 41340// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41341// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41342// CHECK3-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 41343// CHECK3-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 41344// CHECK3-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 41345// CHECK3-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 41346// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41347// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 41348// CHECK3-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 41349// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 41350// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 41351// CHECK3-NEXT:    store i32 4571423, ptr [[DOTOMP_UB]], align 41352// CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41353// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41354// CHECK3-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41355// CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 41356// CHECK3-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP5]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1357// CHECK3-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41358// CHECK3-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP6]], 45714231359// CHECK3-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1360// CHECK3:       cond.true:1361// CHECK3-NEXT:    br label [[COND_END:%.*]]1362// CHECK3:       cond.false:1363// CHECK3-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41364// CHECK3-NEXT:    br label [[COND_END]]1365// CHECK3:       cond.end:1366// CHECK3-NEXT:    [[COND:%.*]] = phi i32 [ 4571423, [[COND_TRUE]] ], [ [[TMP7]], [[COND_FALSE]] ]1367// CHECK3-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 41368// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41369// CHECK3-NEXT:    store i32 [[TMP8]], ptr [[DOTOMP_IV]], align 41370// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1371// CHECK3:       omp.inner.for.cond:1372// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41373// CHECK3-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41374// CHECK3-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP9]], [[TMP10]]1375// CHECK3-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1376// CHECK3:       omp.inner.for.body:1377// CHECK3-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41378// CHECK3-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP11]], 71379// CHECK3-NEXT:    [[SUB:%.*]] = sub nsw i32 32000000, [[MUL]]1380// CHECK3-NEXT:    store i32 [[SUB]], ptr [[I]], align 41381// CHECK3-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[TMP1]], align 41382// CHECK3-NEXT:    [[TMP13:%.*]] = load i32, ptr [[I]], align 41383// CHECK3-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds float, ptr [[TMP12]], i32 [[TMP13]]1384// CHECK3-NEXT:    [[TMP14:%.*]] = load float, ptr [[ARRAYIDX]], align 41385// CHECK3-NEXT:    [[TMP15:%.*]] = load ptr, ptr [[TMP2]], align 41386// CHECK3-NEXT:    [[TMP16:%.*]] = load i32, ptr [[I]], align 41387// CHECK3-NEXT:    [[ARRAYIDX2:%.*]] = getelementptr inbounds float, ptr [[TMP15]], i32 [[TMP16]]1388// CHECK3-NEXT:    [[TMP17:%.*]] = load float, ptr [[ARRAYIDX2]], align 41389// CHECK3-NEXT:    [[MUL3:%.*]] = fmul float [[TMP14]], [[TMP17]]1390// CHECK3-NEXT:    [[TMP18:%.*]] = load ptr, ptr [[TMP3]], align 41391// CHECK3-NEXT:    [[TMP19:%.*]] = load i32, ptr [[I]], align 41392// CHECK3-NEXT:    [[ARRAYIDX4:%.*]] = getelementptr inbounds float, ptr [[TMP18]], i32 [[TMP19]]1393// CHECK3-NEXT:    [[TMP20:%.*]] = load float, ptr [[ARRAYIDX4]], align 41394// CHECK3-NEXT:    [[MUL5:%.*]] = fmul float [[MUL3]], [[TMP20]]1395// CHECK3-NEXT:    [[TMP21:%.*]] = load ptr, ptr [[TMP0]], align 41396// CHECK3-NEXT:    [[TMP22:%.*]] = load i32, ptr [[I]], align 41397// CHECK3-NEXT:    [[ARRAYIDX6:%.*]] = getelementptr inbounds float, ptr [[TMP21]], i32 [[TMP22]]1398// CHECK3-NEXT:    store float [[MUL5]], ptr [[ARRAYIDX6]], align 41399// CHECK3-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1400// CHECK3:       omp.body.continue:1401// CHECK3-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1402// CHECK3:       omp.inner.for.inc:1403// CHECK3-NEXT:    [[TMP23:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41404// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP23]], 11405// CHECK3-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 41406// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND]]1407// CHECK3:       omp.inner.for.end:1408// CHECK3-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1409// CHECK3:       omp.loop.exit:1410// CHECK3-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP5]])1411// CHECK3-NEXT:    ret void1412//1413//1414// CHECK3-LABEL: define {{[^@]+}}@_Z14static_chunkedPfS_S_S_1415// CHECK3-SAME: (ptr noundef [[A:%.*]], ptr noundef [[B:%.*]], ptr noundef [[C:%.*]], ptr noundef [[D:%.*]]) #[[ATTR0]] {1416// CHECK3-NEXT:  entry:1417// CHECK3-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 41418// CHECK3-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 41419// CHECK3-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 41420// CHECK3-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 41421// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [4 x ptr], align 41422// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [4 x ptr], align 41423// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [4 x ptr], align 41424// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 41425// CHECK3-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81426// CHECK3-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 41427// CHECK3-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 41428// CHECK3-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 41429// CHECK3-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 41430// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41431// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 41432// CHECK3-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 41433// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 41434// CHECK3-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01435// CHECK3-NEXT:    store ptr [[TMP0]], ptr [[TMP4]], align 41436// CHECK3-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01437// CHECK3-NEXT:    store ptr [[TMP0]], ptr [[TMP5]], align 41438// CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01439// CHECK3-NEXT:    store ptr null, ptr [[TMP6]], align 41440// CHECK3-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11441// CHECK3-NEXT:    store ptr [[TMP1]], ptr [[TMP7]], align 41442// CHECK3-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11443// CHECK3-NEXT:    store ptr [[TMP1]], ptr [[TMP8]], align 41444// CHECK3-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 11445// CHECK3-NEXT:    store ptr null, ptr [[TMP9]], align 41446// CHECK3-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21447// CHECK3-NEXT:    store ptr [[TMP2]], ptr [[TMP10]], align 41448// CHECK3-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21449// CHECK3-NEXT:    store ptr [[TMP2]], ptr [[TMP11]], align 41450// CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 21451// CHECK3-NEXT:    store ptr null, ptr [[TMP12]], align 41452// CHECK3-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 31453// CHECK3-NEXT:    store ptr [[TMP3]], ptr [[TMP13]], align 41454// CHECK3-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 31455// CHECK3-NEXT:    store ptr [[TMP3]], ptr [[TMP14]], align 41456// CHECK3-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 31457// CHECK3-NEXT:    store ptr null, ptr [[TMP15]], align 41458// CHECK3-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01459// CHECK3-NEXT:    [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01460// CHECK3-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01461// CHECK3-NEXT:    store i32 3, ptr [[TMP18]], align 41462// CHECK3-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11463// CHECK3-NEXT:    store i32 4, ptr [[TMP19]], align 41464// CHECK3-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21465// CHECK3-NEXT:    store ptr [[TMP16]], ptr [[TMP20]], align 41466// CHECK3-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31467// CHECK3-NEXT:    store ptr [[TMP17]], ptr [[TMP21]], align 41468// CHECK3-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41469// CHECK3-NEXT:    store ptr @.offload_sizes.3, ptr [[TMP22]], align 41470// CHECK3-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51471// CHECK3-NEXT:    store ptr @.offload_maptypes.4, ptr [[TMP23]], align 41472// CHECK3-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61473// CHECK3-NEXT:    store ptr null, ptr [[TMP24]], align 41474// CHECK3-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71475// CHECK3-NEXT:    store ptr null, ptr [[TMP25]], align 41476// CHECK3-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81477// CHECK3-NEXT:    store i64 16908289, ptr [[TMP26]], align 81478// CHECK3-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91479// CHECK3-NEXT:    store i64 0, ptr [[TMP27]], align 81480// CHECK3-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101481// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP28]], align 41482// CHECK3-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111483// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP29]], align 41484// CHECK3-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121485// CHECK3-NEXT:    store i32 0, ptr [[TMP30]], align 41486// CHECK3-NEXT:    [[TMP31:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z14static_chunkedPfS_S_S__l80.region_id, ptr [[KERNEL_ARGS]])1487// CHECK3-NEXT:    [[TMP32:%.*]] = icmp ne i32 [[TMP31]], 01488// CHECK3-NEXT:    br i1 [[TMP32]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1489// CHECK3:       omp_offload.failed:1490// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z14static_chunkedPfS_S_S__l80(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], ptr [[TMP3]]) #[[ATTR2]]1491// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT]]1492// CHECK3:       omp_offload.cont:1493// CHECK3-NEXT:    ret void1494//1495//1496// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z14static_chunkedPfS_S_S__l801497// CHECK3-SAME: (ptr noundef [[A:%.*]], ptr noundef [[B:%.*]], ptr noundef [[C:%.*]], ptr noundef [[D:%.*]]) #[[ATTR1]] {1498// CHECK3-NEXT:  entry:1499// CHECK3-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 41500// CHECK3-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 41501// CHECK3-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 41502// CHECK3-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 41503// CHECK3-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 41504// CHECK3-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 41505// CHECK3-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 41506// CHECK3-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 41507// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 4, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z14static_chunkedPfS_S_S__l80.omp_outlined, ptr [[A_ADDR]], ptr [[B_ADDR]], ptr [[C_ADDR]], ptr [[D_ADDR]])1508// CHECK3-NEXT:    ret void1509//1510//1511// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z14static_chunkedPfS_S_S__l80.omp_outlined1512// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[C:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[D:%.*]]) #[[ATTR1]] {1513// CHECK3-NEXT:  entry:1514// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41515// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41516// CHECK3-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 41517// CHECK3-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 41518// CHECK3-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 41519// CHECK3-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 41520// CHECK3-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41521// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 41522// CHECK3-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 41523// CHECK3-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 41524// CHECK3-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41525// CHECK3-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41526// CHECK3-NEXT:    [[I:%.*]] = alloca i32, align 41527// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41528// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41529// CHECK3-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 41530// CHECK3-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 41531// CHECK3-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 41532// CHECK3-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 41533// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41534// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 41535// CHECK3-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 41536// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 41537// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 41538// CHECK3-NEXT:    store i32 16908288, ptr [[DOTOMP_UB]], align 41539// CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41540// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41541// CHECK3-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41542// CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 41543// CHECK3-NEXT:    call void @__kmpc_for_static_init_4u(ptr @[[GLOB1]], i32 [[TMP5]], i32 91, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 5)1544// CHECK3-NEXT:    br label [[OMP_DISPATCH_COND:%.*]]1545// CHECK3:       omp.dispatch.cond:1546// CHECK3-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41547// CHECK3-NEXT:    [[CMP:%.*]] = icmp ugt i32 [[TMP6]], 169082881548// CHECK3-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1549// CHECK3:       cond.true:1550// CHECK3-NEXT:    br label [[COND_END:%.*]]1551// CHECK3:       cond.false:1552// CHECK3-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41553// CHECK3-NEXT:    br label [[COND_END]]1554// CHECK3:       cond.end:1555// CHECK3-NEXT:    [[COND:%.*]] = phi i32 [ 16908288, [[COND_TRUE]] ], [ [[TMP7]], [[COND_FALSE]] ]1556// CHECK3-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 41557// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41558// CHECK3-NEXT:    store i32 [[TMP8]], ptr [[DOTOMP_IV]], align 41559// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41560// CHECK3-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41561// CHECK3-NEXT:    [[CMP1:%.*]] = icmp ule i32 [[TMP9]], [[TMP10]]1562// CHECK3-NEXT:    br i1 [[CMP1]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]1563// CHECK3:       omp.dispatch.body:1564// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1565// CHECK3:       omp.inner.for.cond:1566// CHECK3-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11:![0-9]+]]1567// CHECK3-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP11]]1568// CHECK3-NEXT:    [[CMP2:%.*]] = icmp ule i32 [[TMP11]], [[TMP12]]1569// CHECK3-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1570// CHECK3:       omp.inner.for.body:1571// CHECK3-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]1572// CHECK3-NEXT:    [[MUL:%.*]] = mul i32 [[TMP13]], 1271573// CHECK3-NEXT:    [[ADD:%.*]] = add i32 131071, [[MUL]]1574// CHECK3-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP11]]1575// CHECK3-NEXT:    [[TMP14:%.*]] = load ptr, ptr [[TMP1]], align 4, !llvm.access.group [[ACC_GRP11]]1576// CHECK3-NEXT:    [[TMP15:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP11]]1577// CHECK3-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds nuw float, ptr [[TMP14]], i32 [[TMP15]]1578// CHECK3-NEXT:    [[TMP16:%.*]] = load float, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP11]]1579// CHECK3-NEXT:    [[TMP17:%.*]] = load ptr, ptr [[TMP2]], align 4, !llvm.access.group [[ACC_GRP11]]1580// CHECK3-NEXT:    [[TMP18:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP11]]1581// CHECK3-NEXT:    [[ARRAYIDX3:%.*]] = getelementptr inbounds nuw float, ptr [[TMP17]], i32 [[TMP18]]1582// CHECK3-NEXT:    [[TMP19:%.*]] = load float, ptr [[ARRAYIDX3]], align 4, !llvm.access.group [[ACC_GRP11]]1583// CHECK3-NEXT:    [[MUL4:%.*]] = fmul float [[TMP16]], [[TMP19]]1584// CHECK3-NEXT:    [[TMP20:%.*]] = load ptr, ptr [[TMP3]], align 4, !llvm.access.group [[ACC_GRP11]]1585// CHECK3-NEXT:    [[TMP21:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP11]]1586// CHECK3-NEXT:    [[ARRAYIDX5:%.*]] = getelementptr inbounds nuw float, ptr [[TMP20]], i32 [[TMP21]]1587// CHECK3-NEXT:    [[TMP22:%.*]] = load float, ptr [[ARRAYIDX5]], align 4, !llvm.access.group [[ACC_GRP11]]1588// CHECK3-NEXT:    [[MUL6:%.*]] = fmul float [[MUL4]], [[TMP22]]1589// CHECK3-NEXT:    [[TMP23:%.*]] = load ptr, ptr [[TMP0]], align 4, !llvm.access.group [[ACC_GRP11]]1590// CHECK3-NEXT:    [[TMP24:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP11]]1591// CHECK3-NEXT:    [[ARRAYIDX7:%.*]] = getelementptr inbounds nuw float, ptr [[TMP23]], i32 [[TMP24]]1592// CHECK3-NEXT:    store float [[MUL6]], ptr [[ARRAYIDX7]], align 4, !llvm.access.group [[ACC_GRP11]]1593// CHECK3-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1594// CHECK3:       omp.body.continue:1595// CHECK3-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1596// CHECK3:       omp.inner.for.inc:1597// CHECK3-NEXT:    [[TMP25:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]1598// CHECK3-NEXT:    [[ADD8:%.*]] = add i32 [[TMP25]], 11599// CHECK3-NEXT:    store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]1600// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP12:![0-9]+]]1601// CHECK3:       omp.inner.for.end:1602// CHECK3-NEXT:    br label [[OMP_DISPATCH_INC:%.*]]1603// CHECK3:       omp.dispatch.inc:1604// CHECK3-NEXT:    [[TMP26:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41605// CHECK3-NEXT:    [[TMP27:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41606// CHECK3-NEXT:    [[ADD9:%.*]] = add i32 [[TMP26]], [[TMP27]]1607// CHECK3-NEXT:    store i32 [[ADD9]], ptr [[DOTOMP_LB]], align 41608// CHECK3-NEXT:    [[TMP28:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41609// CHECK3-NEXT:    [[TMP29:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41610// CHECK3-NEXT:    [[ADD10:%.*]] = add i32 [[TMP28]], [[TMP29]]1611// CHECK3-NEXT:    store i32 [[ADD10]], ptr [[DOTOMP_UB]], align 41612// CHECK3-NEXT:    br label [[OMP_DISPATCH_COND]]1613// CHECK3:       omp.dispatch.end:1614// CHECK3-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP5]])1615// CHECK3-NEXT:    ret void1616//1617//1618// CHECK3-LABEL: define {{[^@]+}}@_Z12test_precondv1619// CHECK3-SAME: () #[[ATTR0]] {1620// CHECK3-NEXT:  entry:1621// CHECK3-NEXT:    [[A:%.*]] = alloca i8, align 11622// CHECK3-NEXT:    [[A_CASTED:%.*]] = alloca i32, align 41623// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 41624// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 41625// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 41626// CHECK3-NEXT:    [[TMP:%.*]] = alloca i8, align 11627// CHECK3-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i8, align 11628// CHECK3-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41629// CHECK3-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81630// CHECK3-NEXT:    store i8 0, ptr [[A]], align 11631// CHECK3-NEXT:    [[TMP0:%.*]] = load i8, ptr [[A]], align 11632// CHECK3-NEXT:    store i8 [[TMP0]], ptr [[A_CASTED]], align 11633// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_CASTED]], align 41634// CHECK3-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01635// CHECK3-NEXT:    store i32 [[TMP1]], ptr [[TMP2]], align 41636// CHECK3-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01637// CHECK3-NEXT:    store i32 [[TMP1]], ptr [[TMP3]], align 41638// CHECK3-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01639// CHECK3-NEXT:    store ptr null, ptr [[TMP4]], align 41640// CHECK3-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01641// CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01642// CHECK3-NEXT:    [[TMP7:%.*]] = load i8, ptr [[A]], align 11643// CHECK3-NEXT:    store i8 [[TMP7]], ptr [[DOTCAPTURE_EXPR_]], align 11644// CHECK3-NEXT:    [[TMP8:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 11645// CHECK3-NEXT:    [[CONV:%.*]] = sext i8 [[TMP8]] to i321646// CHECK3-NEXT:    [[SUB:%.*]] = sub i32 10, [[CONV]]1647// CHECK3-NEXT:    [[SUB2:%.*]] = sub i32 [[SUB]], 11648// CHECK3-NEXT:    [[ADD:%.*]] = add i32 [[SUB2]], 11649// CHECK3-NEXT:    [[DIV:%.*]] = udiv i32 [[ADD]], 11650// CHECK3-NEXT:    [[SUB3:%.*]] = sub nsw i32 [[DIV]], 11651// CHECK3-NEXT:    store i32 [[SUB3]], ptr [[DOTCAPTURE_EXPR_1]], align 41652// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41653// CHECK3-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP9]], 11654// CHECK3-NEXT:    [[TMP10:%.*]] = zext i32 [[ADD4]] to i641655// CHECK3-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01656// CHECK3-NEXT:    store i32 3, ptr [[TMP11]], align 41657// CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11658// CHECK3-NEXT:    store i32 1, ptr [[TMP12]], align 41659// CHECK3-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21660// CHECK3-NEXT:    store ptr [[TMP5]], ptr [[TMP13]], align 41661// CHECK3-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31662// CHECK3-NEXT:    store ptr [[TMP6]], ptr [[TMP14]], align 41663// CHECK3-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41664// CHECK3-NEXT:    store ptr @.offload_sizes.5, ptr [[TMP15]], align 41665// CHECK3-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51666// CHECK3-NEXT:    store ptr @.offload_maptypes.6, ptr [[TMP16]], align 41667// CHECK3-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61668// CHECK3-NEXT:    store ptr null, ptr [[TMP17]], align 41669// CHECK3-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71670// CHECK3-NEXT:    store ptr null, ptr [[TMP18]], align 41671// CHECK3-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81672// CHECK3-NEXT:    store i64 [[TMP10]], ptr [[TMP19]], align 81673// CHECK3-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91674// CHECK3-NEXT:    store i64 0, ptr [[TMP20]], align 81675// CHECK3-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101676// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP21]], align 41677// CHECK3-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111678// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP22]], align 41679// CHECK3-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121680// CHECK3-NEXT:    store i32 0, ptr [[TMP23]], align 41681// CHECK3-NEXT:    [[TMP24:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z12test_precondv_l92.region_id, ptr [[KERNEL_ARGS]])1682// CHECK3-NEXT:    [[TMP25:%.*]] = icmp ne i32 [[TMP24]], 01683// CHECK3-NEXT:    br i1 [[TMP25]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1684// CHECK3:       omp_offload.failed:1685// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z12test_precondv_l92(i32 [[TMP1]]) #[[ATTR2]]1686// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT]]1687// CHECK3:       omp_offload.cont:1688// CHECK3-NEXT:    ret void1689//1690//1691// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z12test_precondv_l921692// CHECK3-SAME: (i32 noundef [[A:%.*]]) #[[ATTR1]] {1693// CHECK3-NEXT:  entry:1694// CHECK3-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 41695// CHECK3-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 41696// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z12test_precondv_l92.omp_outlined, ptr [[A_ADDR]])1697// CHECK3-NEXT:    ret void1698//1699//1700// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z12test_precondv_l92.omp_outlined1701// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 1 dereferenceable(1) [[A:%.*]]) #[[ATTR1]] {1702// CHECK3-NEXT:  entry:1703// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41704// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41705// CHECK3-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 41706// CHECK3-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41707// CHECK3-NEXT:    [[TMP:%.*]] = alloca i8, align 11708// CHECK3-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i8, align 11709// CHECK3-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41710// CHECK3-NEXT:    [[I:%.*]] = alloca i8, align 11711// CHECK3-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 41712// CHECK3-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 41713// CHECK3-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41714// CHECK3-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41715// CHECK3-NEXT:    [[I5:%.*]] = alloca i8, align 11716// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41717// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41718// CHECK3-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 41719// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41720// CHECK3-NEXT:    [[TMP1:%.*]] = load i8, ptr [[TMP0]], align 11721// CHECK3-NEXT:    store i8 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 11722// CHECK3-NEXT:    [[TMP2:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 11723// CHECK3-NEXT:    [[CONV:%.*]] = sext i8 [[TMP2]] to i321724// CHECK3-NEXT:    [[SUB:%.*]] = sub i32 10, [[CONV]]1725// CHECK3-NEXT:    [[SUB2:%.*]] = sub i32 [[SUB]], 11726// CHECK3-NEXT:    [[ADD:%.*]] = add i32 [[SUB2]], 11727// CHECK3-NEXT:    [[DIV:%.*]] = udiv i32 [[ADD]], 11728// CHECK3-NEXT:    [[SUB3:%.*]] = sub nsw i32 [[DIV]], 11729// CHECK3-NEXT:    store i32 [[SUB3]], ptr [[DOTCAPTURE_EXPR_1]], align 41730// CHECK3-NEXT:    [[TMP3:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 11731// CHECK3-NEXT:    store i8 [[TMP3]], ptr [[I]], align 11732// CHECK3-NEXT:    [[TMP4:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 11733// CHECK3-NEXT:    [[CONV4:%.*]] = sext i8 [[TMP4]] to i321734// CHECK3-NEXT:    [[CMP:%.*]] = icmp slt i32 [[CONV4]], 101735// CHECK3-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1736// CHECK3:       omp.precond.then:1737// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 41738// CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41739// CHECK3-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_UB]], align 41740// CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41741// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41742// CHECK3-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41743// CHECK3-NEXT:    [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 41744// CHECK3-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP7]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1745// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41746// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41747// CHECK3-NEXT:    [[CMP6:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]1748// CHECK3-NEXT:    br i1 [[CMP6]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1749// CHECK3:       cond.true:1750// CHECK3-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41751// CHECK3-NEXT:    br label [[COND_END:%.*]]1752// CHECK3:       cond.false:1753// CHECK3-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41754// CHECK3-NEXT:    br label [[COND_END]]1755// CHECK3:       cond.end:1756// CHECK3-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]1757// CHECK3-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 41758// CHECK3-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41759// CHECK3-NEXT:    store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 41760// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1761// CHECK3:       omp.inner.for.cond:1762// CHECK3-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41763// CHECK3-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41764// CHECK3-NEXT:    [[CMP7:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]1765// CHECK3-NEXT:    br i1 [[CMP7]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1766// CHECK3:       omp.inner.for.body:1767// CHECK3-NEXT:    [[TMP15:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 11768// CHECK3-NEXT:    [[CONV8:%.*]] = sext i8 [[TMP15]] to i321769// CHECK3-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41770// CHECK3-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP16]], 11771// CHECK3-NEXT:    [[ADD9:%.*]] = add nsw i32 [[CONV8]], [[MUL]]1772// CHECK3-NEXT:    [[CONV10:%.*]] = trunc i32 [[ADD9]] to i81773// CHECK3-NEXT:    store i8 [[CONV10]], ptr [[I5]], align 11774// CHECK3-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1775// CHECK3:       omp.body.continue:1776// CHECK3-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1777// CHECK3:       omp.inner.for.inc:1778// CHECK3-NEXT:    [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41779// CHECK3-NEXT:    [[ADD11:%.*]] = add nsw i32 [[TMP17]], 11780// CHECK3-NEXT:    store i32 [[ADD11]], ptr [[DOTOMP_IV]], align 41781// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND]]1782// CHECK3:       omp.inner.for.end:1783// CHECK3-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1784// CHECK3:       omp.loop.exit:1785// CHECK3-NEXT:    [[TMP18:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41786// CHECK3-NEXT:    [[TMP19:%.*]] = load i32, ptr [[TMP18]], align 41787// CHECK3-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP19]])1788// CHECK3-NEXT:    br label [[OMP_PRECOND_END]]1789// CHECK3:       omp.precond.end:1790// CHECK3-NEXT:    ret void1791//1792//1793// CHECK3-LABEL: define {{[^@]+}}@_Z4fintv1794// CHECK3-SAME: () #[[ATTR0]] {1795// CHECK3-NEXT:  entry:1796// CHECK3-NEXT:    [[CALL:%.*]] = call noundef i32 @_Z9ftemplateIiET_v()1797// CHECK3-NEXT:    ret i32 [[CALL]]1798//1799//1800// CHECK3-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_v1801// CHECK3-SAME: () #[[ATTR0]] comdat {1802// CHECK3-NEXT:  entry:1803// CHECK3-NEXT:    [[AA:%.*]] = alloca i16, align 21804// CHECK3-NEXT:    [[AA_CASTED:%.*]] = alloca i32, align 41805// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 41806// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 41807// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 41808// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 41809// CHECK3-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81810// CHECK3-NEXT:    store i16 0, ptr [[AA]], align 21811// CHECK3-NEXT:    [[TMP0:%.*]] = load i16, ptr [[AA]], align 21812// CHECK3-NEXT:    store i16 [[TMP0]], ptr [[AA_CASTED]], align 21813// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[AA_CASTED]], align 41814// CHECK3-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01815// CHECK3-NEXT:    store i32 [[TMP1]], ptr [[TMP2]], align 41816// CHECK3-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01817// CHECK3-NEXT:    store i32 [[TMP1]], ptr [[TMP3]], align 41818// CHECK3-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01819// CHECK3-NEXT:    store ptr null, ptr [[TMP4]], align 41820// CHECK3-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01821// CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01822// CHECK3-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01823// CHECK3-NEXT:    store i32 3, ptr [[TMP7]], align 41824// CHECK3-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11825// CHECK3-NEXT:    store i32 1, ptr [[TMP8]], align 41826// CHECK3-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21827// CHECK3-NEXT:    store ptr [[TMP5]], ptr [[TMP9]], align 41828// CHECK3-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31829// CHECK3-NEXT:    store ptr [[TMP6]], ptr [[TMP10]], align 41830// CHECK3-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41831// CHECK3-NEXT:    store ptr @.offload_sizes.7, ptr [[TMP11]], align 41832// CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51833// CHECK3-NEXT:    store ptr @.offload_maptypes.8, ptr [[TMP12]], align 41834// CHECK3-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61835// CHECK3-NEXT:    store ptr null, ptr [[TMP13]], align 41836// CHECK3-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71837// CHECK3-NEXT:    store ptr null, ptr [[TMP14]], align 41838// CHECK3-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81839// CHECK3-NEXT:    store i64 100, ptr [[TMP15]], align 81840// CHECK3-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91841// CHECK3-NEXT:    store i64 0, ptr [[TMP16]], align 81842// CHECK3-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101843// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP17]], align 41844// CHECK3-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111845// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP18]], align 41846// CHECK3-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121847// CHECK3-NEXT:    store i32 0, ptr [[TMP19]], align 41848// CHECK3-NEXT:    [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_v_l108.region_id, ptr [[KERNEL_ARGS]])1849// CHECK3-NEXT:    [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 01850// CHECK3-NEXT:    br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1851// CHECK3:       omp_offload.failed:1852// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_v_l108(i32 [[TMP1]]) #[[ATTR2]]1853// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT]]1854// CHECK3:       omp_offload.cont:1855// CHECK3-NEXT:    ret i32 01856//1857//1858// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_v_l1081859// CHECK3-SAME: (i32 noundef [[AA:%.*]]) #[[ATTR1]] {1860// CHECK3-NEXT:  entry:1861// CHECK3-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 41862// CHECK3-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 41863// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_v_l108.omp_outlined, ptr [[AA_ADDR]])1864// CHECK3-NEXT:    ret void1865//1866//1867// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_v_l108.omp_outlined1868// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[AA:%.*]]) #[[ATTR1]] {1869// CHECK3-NEXT:  entry:1870// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41871// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41872// CHECK3-NEXT:    [[AA_ADDR:%.*]] = alloca ptr, align 41873// CHECK3-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41874// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 41875// CHECK3-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 41876// CHECK3-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 41877// CHECK3-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41878// CHECK3-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41879// CHECK3-NEXT:    [[I:%.*]] = alloca i32, align 41880// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41881// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41882// CHECK3-NEXT:    store ptr [[AA]], ptr [[AA_ADDR]], align 41883// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 41884// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 41885// CHECK3-NEXT:    store i32 99, ptr [[DOTOMP_UB]], align 41886// CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41887// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41888// CHECK3-NEXT:    [[TMP1:%.*]] = load i16, ptr [[TMP0]], align 21889// CHECK3-NEXT:    [[CONV:%.*]] = sext i16 [[TMP1]] to i321890// CHECK3-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41891// CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 41892// CHECK3-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP3]], i32 91, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 [[CONV]])1893// CHECK3-NEXT:    br label [[OMP_DISPATCH_COND:%.*]]1894// CHECK3:       omp.dispatch.cond:1895// CHECK3-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41896// CHECK3-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 991897// CHECK3-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1898// CHECK3:       cond.true:1899// CHECK3-NEXT:    br label [[COND_END:%.*]]1900// CHECK3:       cond.false:1901// CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41902// CHECK3-NEXT:    br label [[COND_END]]1903// CHECK3:       cond.end:1904// CHECK3-NEXT:    [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]1905// CHECK3-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 41906// CHECK3-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41907// CHECK3-NEXT:    store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 41908// CHECK3-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41909// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41910// CHECK3-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]1911// CHECK3-NEXT:    br i1 [[CMP1]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]1912// CHECK3:       omp.dispatch.body:1913// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1914// CHECK3:       omp.inner.for.cond:1915// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP14:![0-9]+]]1916// CHECK3-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP14]]1917// CHECK3-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP9]], [[TMP10]]1918// CHECK3-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1919// CHECK3:       omp.inner.for.body:1920// CHECK3-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP14]]1921// CHECK3-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP11]], 11922// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1923// CHECK3-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP14]]1924// CHECK3-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1925// CHECK3:       omp.body.continue:1926// CHECK3-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1927// CHECK3:       omp.inner.for.inc:1928// CHECK3-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP14]]1929// CHECK3-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP12]], 11930// CHECK3-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP14]]1931// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP15:![0-9]+]]1932// CHECK3:       omp.inner.for.end:1933// CHECK3-NEXT:    br label [[OMP_DISPATCH_INC:%.*]]1934// CHECK3:       omp.dispatch.inc:1935// CHECK3-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41936// CHECK3-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41937// CHECK3-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP13]], [[TMP14]]1938// CHECK3-NEXT:    store i32 [[ADD4]], ptr [[DOTOMP_LB]], align 41939// CHECK3-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41940// CHECK3-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41941// CHECK3-NEXT:    [[ADD5:%.*]] = add nsw i32 [[TMP15]], [[TMP16]]1942// CHECK3-NEXT:    store i32 [[ADD5]], ptr [[DOTOMP_UB]], align 41943// CHECK3-NEXT:    br label [[OMP_DISPATCH_COND]]1944// CHECK3:       omp.dispatch.end:1945// CHECK3-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP3]])1946// CHECK3-NEXT:    ret void1947//1948//1949// CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z23without_schedule_clausePfS_S_S__l561950// CHECK17-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[A:%.*]], ptr noundef [[B:%.*]], ptr noundef [[C:%.*]], ptr noundef [[D:%.*]]) #[[ATTR0:[0-9]+]] {1951// CHECK17-NEXT:  entry:1952// CHECK17-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81953// CHECK17-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81954// CHECK17-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 81955// CHECK17-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 81956// CHECK17-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 81957// CHECK17-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81958// CHECK17-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81959// CHECK17-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 81960// CHECK17-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 81961// CHECK17-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 81962// CHECK17-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2:[0-9]+]], i32 4, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z23without_schedule_clausePfS_S_S__l56.omp_outlined, ptr [[A_ADDR]], ptr [[B_ADDR]], ptr [[C_ADDR]], ptr [[D_ADDR]])1963// CHECK17-NEXT:    ret void1964//1965//1966// CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z23without_schedule_clausePfS_S_S__l56.omp_outlined1967// CHECK17-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[A:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[B:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[C:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[D:%.*]]) #[[ATTR0]] {1968// CHECK17-NEXT:  entry:1969// CHECK17-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81970// CHECK17-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81971// CHECK17-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81972// CHECK17-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 81973// CHECK17-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 81974// CHECK17-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 81975// CHECK17-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41976// CHECK17-NEXT:    [[TMP:%.*]] = alloca i32, align 41977// CHECK17-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 41978// CHECK17-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 41979// CHECK17-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41980// CHECK17-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41981// CHECK17-NEXT:    [[I:%.*]] = alloca i32, align 41982// CHECK17-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81983// CHECK17-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81984// CHECK17-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81985// CHECK17-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 81986// CHECK17-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 81987// CHECK17-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 81988// CHECK17-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81989// CHECK17-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 81990// CHECK17-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 81991// CHECK17-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 81992// CHECK17-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 41993// CHECK17-NEXT:    store i32 4571423, ptr [[DOTOMP_UB]], align 41994// CHECK17-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41995// CHECK17-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41996// CHECK17-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81997// CHECK17-NEXT:    [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 41998// CHECK17-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP5]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1999// CHECK17-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42000// CHECK17-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP6]], 45714232001// CHECK17-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2002// CHECK17:       cond.true:2003// CHECK17-NEXT:    br label [[COND_END:%.*]]2004// CHECK17:       cond.false:2005// CHECK17-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42006// CHECK17-NEXT:    br label [[COND_END]]2007// CHECK17:       cond.end:2008// CHECK17-NEXT:    [[COND:%.*]] = phi i32 [ 4571423, [[COND_TRUE]] ], [ [[TMP7]], [[COND_FALSE]] ]2009// CHECK17-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 42010// CHECK17-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42011// CHECK17-NEXT:    store i32 [[TMP8]], ptr [[DOTOMP_IV]], align 42012// CHECK17-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]2013// CHECK17:       omp.inner.for.cond:2014// CHECK17-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42015// CHECK17-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42016// CHECK17-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP9]], [[TMP10]]2017// CHECK17-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2018// CHECK17:       omp.inner.for.body:2019// CHECK17-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42020// CHECK17-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP11]], 72021// CHECK17-NEXT:    [[ADD:%.*]] = add nsw i32 33, [[MUL]]2022// CHECK17-NEXT:    store i32 [[ADD]], ptr [[I]], align 42023// CHECK17-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[TMP1]], align 82024// CHECK17-NEXT:    [[TMP13:%.*]] = load i32, ptr [[I]], align 42025// CHECK17-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP13]] to i642026// CHECK17-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds float, ptr [[TMP12]], i64 [[IDXPROM]]2027// CHECK17-NEXT:    [[TMP14:%.*]] = load float, ptr [[ARRAYIDX]], align 42028// CHECK17-NEXT:    [[TMP15:%.*]] = load ptr, ptr [[TMP2]], align 82029// CHECK17-NEXT:    [[TMP16:%.*]] = load i32, ptr [[I]], align 42030// CHECK17-NEXT:    [[IDXPROM2:%.*]] = sext i32 [[TMP16]] to i642031// CHECK17-NEXT:    [[ARRAYIDX3:%.*]] = getelementptr inbounds float, ptr [[TMP15]], i64 [[IDXPROM2]]2032// CHECK17-NEXT:    [[TMP17:%.*]] = load float, ptr [[ARRAYIDX3]], align 42033// CHECK17-NEXT:    [[MUL4:%.*]] = fmul float [[TMP14]], [[TMP17]]2034// CHECK17-NEXT:    [[TMP18:%.*]] = load ptr, ptr [[TMP3]], align 82035// CHECK17-NEXT:    [[TMP19:%.*]] = load i32, ptr [[I]], align 42036// CHECK17-NEXT:    [[IDXPROM5:%.*]] = sext i32 [[TMP19]] to i642037// CHECK17-NEXT:    [[ARRAYIDX6:%.*]] = getelementptr inbounds float, ptr [[TMP18]], i64 [[IDXPROM5]]2038// CHECK17-NEXT:    [[TMP20:%.*]] = load float, ptr [[ARRAYIDX6]], align 42039// CHECK17-NEXT:    [[MUL7:%.*]] = fmul float [[MUL4]], [[TMP20]]2040// CHECK17-NEXT:    [[TMP21:%.*]] = load ptr, ptr [[TMP0]], align 82041// CHECK17-NEXT:    [[TMP22:%.*]] = load i32, ptr [[I]], align 42042// CHECK17-NEXT:    [[IDXPROM8:%.*]] = sext i32 [[TMP22]] to i642043// CHECK17-NEXT:    [[ARRAYIDX9:%.*]] = getelementptr inbounds float, ptr [[TMP21]], i64 [[IDXPROM8]]2044// CHECK17-NEXT:    store float [[MUL7]], ptr [[ARRAYIDX9]], align 42045// CHECK17-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]2046// CHECK17:       omp.body.continue:2047// CHECK17-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]2048// CHECK17:       omp.inner.for.inc:2049// CHECK17-NEXT:    [[TMP23:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42050// CHECK17-NEXT:    [[ADD10:%.*]] = add nsw i32 [[TMP23]], 12051// CHECK17-NEXT:    store i32 [[ADD10]], ptr [[DOTOMP_IV]], align 42052// CHECK17-NEXT:    br label [[OMP_INNER_FOR_COND]]2053// CHECK17:       omp.inner.for.end:2054// CHECK17-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]2055// CHECK17:       omp.loop.exit:2056// CHECK17-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP5]])2057// CHECK17-NEXT:    ret void2058//2059//2060// CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18static_not_chunkedPfS_S_S__l682061// CHECK17-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[A:%.*]], ptr noundef [[B:%.*]], ptr noundef [[C:%.*]], ptr noundef [[D:%.*]]) #[[ATTR0]] {2062// CHECK17-NEXT:  entry:2063// CHECK17-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 82064// CHECK17-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 82065// CHECK17-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 82066// CHECK17-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 82067// CHECK17-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 82068// CHECK17-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 82069// CHECK17-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 82070// CHECK17-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 82071// CHECK17-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 82072// CHECK17-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 82073// CHECK17-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 4, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18static_not_chunkedPfS_S_S__l68.omp_outlined, ptr [[A_ADDR]], ptr [[B_ADDR]], ptr [[C_ADDR]], ptr [[D_ADDR]])2074// CHECK17-NEXT:    ret void2075//2076//2077// CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18static_not_chunkedPfS_S_S__l68.omp_outlined2078// CHECK17-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[A:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[B:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[C:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[D:%.*]]) #[[ATTR0]] {2079// CHECK17-NEXT:  entry:2080// CHECK17-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82081// CHECK17-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82082// CHECK17-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 82083// CHECK17-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 82084// CHECK17-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 82085// CHECK17-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 82086// CHECK17-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 42087// CHECK17-NEXT:    [[TMP:%.*]] = alloca i32, align 42088// CHECK17-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 42089// CHECK17-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 42090// CHECK17-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42091// CHECK17-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42092// CHECK17-NEXT:    [[I:%.*]] = alloca i32, align 42093// CHECK17-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82094// CHECK17-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82095// CHECK17-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 82096// CHECK17-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 82097// CHECK17-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 82098// CHECK17-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 82099// CHECK17-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 82100// CHECK17-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 82101// CHECK17-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 82102// CHECK17-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 82103// CHECK17-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 42104// CHECK17-NEXT:    store i32 4571423, ptr [[DOTOMP_UB]], align 42105// CHECK17-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 42106// CHECK17-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 42107// CHECK17-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82108// CHECK17-NEXT:    [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 42109// CHECK17-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP5]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2110// CHECK17-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42111// CHECK17-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP6]], 45714232112// CHECK17-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2113// CHECK17:       cond.true:2114// CHECK17-NEXT:    br label [[COND_END:%.*]]2115// CHECK17:       cond.false:2116// CHECK17-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42117// CHECK17-NEXT:    br label [[COND_END]]2118// CHECK17:       cond.end:2119// CHECK17-NEXT:    [[COND:%.*]] = phi i32 [ 4571423, [[COND_TRUE]] ], [ [[TMP7]], [[COND_FALSE]] ]2120// CHECK17-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 42121// CHECK17-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42122// CHECK17-NEXT:    store i32 [[TMP8]], ptr [[DOTOMP_IV]], align 42123// CHECK17-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]2124// CHECK17:       omp.inner.for.cond:2125// CHECK17-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42126// CHECK17-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42127// CHECK17-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP9]], [[TMP10]]2128// CHECK17-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2129// CHECK17:       omp.inner.for.body:2130// CHECK17-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42131// CHECK17-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP11]], 72132// CHECK17-NEXT:    [[SUB:%.*]] = sub nsw i32 32000000, [[MUL]]2133// CHECK17-NEXT:    store i32 [[SUB]], ptr [[I]], align 42134// CHECK17-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[TMP1]], align 82135// CHECK17-NEXT:    [[TMP13:%.*]] = load i32, ptr [[I]], align 42136// CHECK17-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP13]] to i642137// CHECK17-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds float, ptr [[TMP12]], i64 [[IDXPROM]]2138// CHECK17-NEXT:    [[TMP14:%.*]] = load float, ptr [[ARRAYIDX]], align 42139// CHECK17-NEXT:    [[TMP15:%.*]] = load ptr, ptr [[TMP2]], align 82140// CHECK17-NEXT:    [[TMP16:%.*]] = load i32, ptr [[I]], align 42141// CHECK17-NEXT:    [[IDXPROM2:%.*]] = sext i32 [[TMP16]] to i642142// CHECK17-NEXT:    [[ARRAYIDX3:%.*]] = getelementptr inbounds float, ptr [[TMP15]], i64 [[IDXPROM2]]2143// CHECK17-NEXT:    [[TMP17:%.*]] = load float, ptr [[ARRAYIDX3]], align 42144// CHECK17-NEXT:    [[MUL4:%.*]] = fmul float [[TMP14]], [[TMP17]]2145// CHECK17-NEXT:    [[TMP18:%.*]] = load ptr, ptr [[TMP3]], align 82146// CHECK17-NEXT:    [[TMP19:%.*]] = load i32, ptr [[I]], align 42147// CHECK17-NEXT:    [[IDXPROM5:%.*]] = sext i32 [[TMP19]] to i642148// CHECK17-NEXT:    [[ARRAYIDX6:%.*]] = getelementptr inbounds float, ptr [[TMP18]], i64 [[IDXPROM5]]2149// CHECK17-NEXT:    [[TMP20:%.*]] = load float, ptr [[ARRAYIDX6]], align 42150// CHECK17-NEXT:    [[MUL7:%.*]] = fmul float [[MUL4]], [[TMP20]]2151// CHECK17-NEXT:    [[TMP21:%.*]] = load ptr, ptr [[TMP0]], align 82152// CHECK17-NEXT:    [[TMP22:%.*]] = load i32, ptr [[I]], align 42153// CHECK17-NEXT:    [[IDXPROM8:%.*]] = sext i32 [[TMP22]] to i642154// CHECK17-NEXT:    [[ARRAYIDX9:%.*]] = getelementptr inbounds float, ptr [[TMP21]], i64 [[IDXPROM8]]2155// CHECK17-NEXT:    store float [[MUL7]], ptr [[ARRAYIDX9]], align 42156// CHECK17-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]2157// CHECK17:       omp.body.continue:2158// CHECK17-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]2159// CHECK17:       omp.inner.for.inc:2160// CHECK17-NEXT:    [[TMP23:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42161// CHECK17-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP23]], 12162// CHECK17-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 42163// CHECK17-NEXT:    br label [[OMP_INNER_FOR_COND]]2164// CHECK17:       omp.inner.for.end:2165// CHECK17-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]2166// CHECK17:       omp.loop.exit:2167// CHECK17-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP5]])2168// CHECK17-NEXT:    ret void2169//2170//2171// CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z14static_chunkedPfS_S_S__l802172// CHECK17-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[A:%.*]], ptr noundef [[B:%.*]], ptr noundef [[C:%.*]], ptr noundef [[D:%.*]]) #[[ATTR0]] {2173// CHECK17-NEXT:  entry:2174// CHECK17-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 82175// CHECK17-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 82176// CHECK17-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 82177// CHECK17-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 82178// CHECK17-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 82179// CHECK17-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 82180// CHECK17-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 82181// CHECK17-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 82182// CHECK17-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 82183// CHECK17-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 82184// CHECK17-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 4, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z14static_chunkedPfS_S_S__l80.omp_outlined, ptr [[A_ADDR]], ptr [[B_ADDR]], ptr [[C_ADDR]], ptr [[D_ADDR]])2185// CHECK17-NEXT:    ret void2186//2187//2188// CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z14static_chunkedPfS_S_S__l80.omp_outlined2189// CHECK17-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[A:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[B:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[C:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[D:%.*]]) #[[ATTR0]] {2190// CHECK17-NEXT:  entry:2191// CHECK17-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82192// CHECK17-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82193// CHECK17-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 82194// CHECK17-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 82195// CHECK17-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 82196// CHECK17-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 82197// CHECK17-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 42198// CHECK17-NEXT:    [[TMP:%.*]] = alloca i32, align 42199// CHECK17-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 42200// CHECK17-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 42201// CHECK17-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42202// CHECK17-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42203// CHECK17-NEXT:    [[I:%.*]] = alloca i32, align 42204// CHECK17-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82205// CHECK17-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82206// CHECK17-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 82207// CHECK17-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 82208// CHECK17-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 82209// CHECK17-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 82210// CHECK17-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 82211// CHECK17-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 82212// CHECK17-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 82213// CHECK17-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 82214// CHECK17-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 42215// CHECK17-NEXT:    store i32 16908288, ptr [[DOTOMP_UB]], align 42216// CHECK17-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 42217// CHECK17-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 42218// CHECK17-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82219// CHECK17-NEXT:    [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 42220// CHECK17-NEXT:    call void @__kmpc_for_static_init_4u(ptr @[[GLOB1]], i32 [[TMP5]], i32 91, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 5)2221// CHECK17-NEXT:    br label [[OMP_DISPATCH_COND:%.*]]2222// CHECK17:       omp.dispatch.cond:2223// CHECK17-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42224// CHECK17-NEXT:    [[CMP:%.*]] = icmp ugt i32 [[TMP6]], 169082882225// CHECK17-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2226// CHECK17:       cond.true:2227// CHECK17-NEXT:    br label [[COND_END:%.*]]2228// CHECK17:       cond.false:2229// CHECK17-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42230// CHECK17-NEXT:    br label [[COND_END]]2231// CHECK17:       cond.end:2232// CHECK17-NEXT:    [[COND:%.*]] = phi i32 [ 16908288, [[COND_TRUE]] ], [ [[TMP7]], [[COND_FALSE]] ]2233// CHECK17-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 42234// CHECK17-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42235// CHECK17-NEXT:    store i32 [[TMP8]], ptr [[DOTOMP_IV]], align 42236// CHECK17-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42237// CHECK17-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42238// CHECK17-NEXT:    [[CMP1:%.*]] = icmp ule i32 [[TMP9]], [[TMP10]]2239// CHECK17-NEXT:    br i1 [[CMP1]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]2240// CHECK17:       omp.dispatch.body:2241// CHECK17-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]2242// CHECK17:       omp.inner.for.cond:2243// CHECK17-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11:![0-9]+]]2244// CHECK17-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP11]]2245// CHECK17-NEXT:    [[CMP2:%.*]] = icmp ule i32 [[TMP11]], [[TMP12]]2246// CHECK17-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2247// CHECK17:       omp.inner.for.body:2248// CHECK17-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]2249// CHECK17-NEXT:    [[MUL:%.*]] = mul i32 [[TMP13]], 1272250// CHECK17-NEXT:    [[ADD:%.*]] = add i32 131071, [[MUL]]2251// CHECK17-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP11]]2252// CHECK17-NEXT:    [[TMP14:%.*]] = load ptr, ptr [[TMP1]], align 8, !llvm.access.group [[ACC_GRP11]]2253// CHECK17-NEXT:    [[TMP15:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP11]]2254// CHECK17-NEXT:    [[IDXPROM:%.*]] = zext i32 [[TMP15]] to i642255// CHECK17-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds nuw float, ptr [[TMP14]], i64 [[IDXPROM]]2256// CHECK17-NEXT:    [[TMP16:%.*]] = load float, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP11]]2257// CHECK17-NEXT:    [[TMP17:%.*]] = load ptr, ptr [[TMP2]], align 8, !llvm.access.group [[ACC_GRP11]]2258// CHECK17-NEXT:    [[TMP18:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP11]]2259// CHECK17-NEXT:    [[IDXPROM3:%.*]] = zext i32 [[TMP18]] to i642260// CHECK17-NEXT:    [[ARRAYIDX4:%.*]] = getelementptr inbounds nuw float, ptr [[TMP17]], i64 [[IDXPROM3]]2261// CHECK17-NEXT:    [[TMP19:%.*]] = load float, ptr [[ARRAYIDX4]], align 4, !llvm.access.group [[ACC_GRP11]]2262// CHECK17-NEXT:    [[MUL5:%.*]] = fmul float [[TMP16]], [[TMP19]]2263// CHECK17-NEXT:    [[TMP20:%.*]] = load ptr, ptr [[TMP3]], align 8, !llvm.access.group [[ACC_GRP11]]2264// CHECK17-NEXT:    [[TMP21:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP11]]2265// CHECK17-NEXT:    [[IDXPROM6:%.*]] = zext i32 [[TMP21]] to i642266// CHECK17-NEXT:    [[ARRAYIDX7:%.*]] = getelementptr inbounds nuw float, ptr [[TMP20]], i64 [[IDXPROM6]]2267// CHECK17-NEXT:    [[TMP22:%.*]] = load float, ptr [[ARRAYIDX7]], align 4, !llvm.access.group [[ACC_GRP11]]2268// CHECK17-NEXT:    [[MUL8:%.*]] = fmul float [[MUL5]], [[TMP22]]2269// CHECK17-NEXT:    [[TMP23:%.*]] = load ptr, ptr [[TMP0]], align 8, !llvm.access.group [[ACC_GRP11]]2270// CHECK17-NEXT:    [[TMP24:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP11]]2271// CHECK17-NEXT:    [[IDXPROM9:%.*]] = zext i32 [[TMP24]] to i642272// CHECK17-NEXT:    [[ARRAYIDX10:%.*]] = getelementptr inbounds nuw float, ptr [[TMP23]], i64 [[IDXPROM9]]2273// CHECK17-NEXT:    store float [[MUL8]], ptr [[ARRAYIDX10]], align 4, !llvm.access.group [[ACC_GRP11]]2274// CHECK17-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]2275// CHECK17:       omp.body.continue:2276// CHECK17-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]2277// CHECK17:       omp.inner.for.inc:2278// CHECK17-NEXT:    [[TMP25:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]2279// CHECK17-NEXT:    [[ADD11:%.*]] = add i32 [[TMP25]], 12280// CHECK17-NEXT:    store i32 [[ADD11]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]2281// CHECK17-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP12:![0-9]+]]2282// CHECK17:       omp.inner.for.end:2283// CHECK17-NEXT:    br label [[OMP_DISPATCH_INC:%.*]]2284// CHECK17:       omp.dispatch.inc:2285// CHECK17-NEXT:    [[TMP26:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42286// CHECK17-NEXT:    [[TMP27:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 42287// CHECK17-NEXT:    [[ADD12:%.*]] = add i32 [[TMP26]], [[TMP27]]2288// CHECK17-NEXT:    store i32 [[ADD12]], ptr [[DOTOMP_LB]], align 42289// CHECK17-NEXT:    [[TMP28:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42290// CHECK17-NEXT:    [[TMP29:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 42291// CHECK17-NEXT:    [[ADD13:%.*]] = add i32 [[TMP28]], [[TMP29]]2292// CHECK17-NEXT:    store i32 [[ADD13]], ptr [[DOTOMP_UB]], align 42293// CHECK17-NEXT:    br label [[OMP_DISPATCH_COND]]2294// CHECK17:       omp.dispatch.end:2295// CHECK17-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP5]])2296// CHECK17-NEXT:    ret void2297//2298//2299// CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z12test_precondv_l922300// CHECK17-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[A:%.*]]) #[[ATTR0]] {2301// CHECK17-NEXT:  entry:2302// CHECK17-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 82303// CHECK17-NEXT:    [[A_ADDR:%.*]] = alloca i64, align 82304// CHECK17-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 82305// CHECK17-NEXT:    store i64 [[A]], ptr [[A_ADDR]], align 82306// CHECK17-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z12test_precondv_l92.omp_outlined, ptr [[A_ADDR]])2307// CHECK17-NEXT:    ret void2308//2309//2310// CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z12test_precondv_l92.omp_outlined2311// CHECK17-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 1 dereferenceable(1) [[A:%.*]]) #[[ATTR0]] {2312// CHECK17-NEXT:  entry:2313// CHECK17-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82314// CHECK17-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82315// CHECK17-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 82316// CHECK17-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 42317// CHECK17-NEXT:    [[TMP:%.*]] = alloca i8, align 12318// CHECK17-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i8, align 12319// CHECK17-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42320// CHECK17-NEXT:    [[I:%.*]] = alloca i8, align 12321// CHECK17-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 42322// CHECK17-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 42323// CHECK17-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42324// CHECK17-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42325// CHECK17-NEXT:    [[I5:%.*]] = alloca i8, align 12326// CHECK17-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82327// CHECK17-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82328// CHECK17-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 82329// CHECK17-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 82330// CHECK17-NEXT:    [[TMP1:%.*]] = load i8, ptr [[TMP0]], align 12331// CHECK17-NEXT:    store i8 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 12332// CHECK17-NEXT:    [[TMP2:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 12333// CHECK17-NEXT:    [[CONV:%.*]] = sext i8 [[TMP2]] to i322334// CHECK17-NEXT:    [[SUB:%.*]] = sub i32 10, [[CONV]]2335// CHECK17-NEXT:    [[SUB2:%.*]] = sub i32 [[SUB]], 12336// CHECK17-NEXT:    [[ADD:%.*]] = add i32 [[SUB2]], 12337// CHECK17-NEXT:    [[DIV:%.*]] = udiv i32 [[ADD]], 12338// CHECK17-NEXT:    [[SUB3:%.*]] = sub nsw i32 [[DIV]], 12339// CHECK17-NEXT:    store i32 [[SUB3]], ptr [[DOTCAPTURE_EXPR_1]], align 42340// CHECK17-NEXT:    [[TMP3:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 12341// CHECK17-NEXT:    store i8 [[TMP3]], ptr [[I]], align 12342// CHECK17-NEXT:    [[TMP4:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 12343// CHECK17-NEXT:    [[CONV4:%.*]] = sext i8 [[TMP4]] to i322344// CHECK17-NEXT:    [[CMP:%.*]] = icmp slt i32 [[CONV4]], 102345// CHECK17-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2346// CHECK17:       omp.precond.then:2347// CHECK17-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 42348// CHECK17-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42349// CHECK17-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_UB]], align 42350// CHECK17-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 42351// CHECK17-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 42352// CHECK17-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82353// CHECK17-NEXT:    [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 42354// CHECK17-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP7]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2355// CHECK17-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42356// CHECK17-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42357// CHECK17-NEXT:    [[CMP6:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]2358// CHECK17-NEXT:    br i1 [[CMP6]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2359// CHECK17:       cond.true:2360// CHECK17-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42361// CHECK17-NEXT:    br label [[COND_END:%.*]]2362// CHECK17:       cond.false:2363// CHECK17-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42364// CHECK17-NEXT:    br label [[COND_END]]2365// CHECK17:       cond.end:2366// CHECK17-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]2367// CHECK17-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 42368// CHECK17-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42369// CHECK17-NEXT:    store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 42370// CHECK17-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]2371// CHECK17:       omp.inner.for.cond:2372// CHECK17-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42373// CHECK17-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42374// CHECK17-NEXT:    [[CMP7:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]2375// CHECK17-NEXT:    br i1 [[CMP7]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2376// CHECK17:       omp.inner.for.body:2377// CHECK17-NEXT:    [[TMP15:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 12378// CHECK17-NEXT:    [[CONV8:%.*]] = sext i8 [[TMP15]] to i322379// CHECK17-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42380// CHECK17-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP16]], 12381// CHECK17-NEXT:    [[ADD9:%.*]] = add nsw i32 [[CONV8]], [[MUL]]2382// CHECK17-NEXT:    [[CONV10:%.*]] = trunc i32 [[ADD9]] to i82383// CHECK17-NEXT:    store i8 [[CONV10]], ptr [[I5]], align 12384// CHECK17-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]2385// CHECK17:       omp.body.continue:2386// CHECK17-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]2387// CHECK17:       omp.inner.for.inc:2388// CHECK17-NEXT:    [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42389// CHECK17-NEXT:    [[ADD11:%.*]] = add nsw i32 [[TMP17]], 12390// CHECK17-NEXT:    store i32 [[ADD11]], ptr [[DOTOMP_IV]], align 42391// CHECK17-NEXT:    br label [[OMP_INNER_FOR_COND]]2392// CHECK17:       omp.inner.for.end:2393// CHECK17-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]2394// CHECK17:       omp.loop.exit:2395// CHECK17-NEXT:    [[TMP18:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82396// CHECK17-NEXT:    [[TMP19:%.*]] = load i32, ptr [[TMP18]], align 42397// CHECK17-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP19]])2398// CHECK17-NEXT:    br label [[OMP_PRECOND_END]]2399// CHECK17:       omp.precond.end:2400// CHECK17-NEXT:    ret void2401//2402//2403// CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_v_l1082404// CHECK17-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[AA:%.*]]) #[[ATTR0]] {2405// CHECK17-NEXT:  entry:2406// CHECK17-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 82407// CHECK17-NEXT:    [[AA_ADDR:%.*]] = alloca i64, align 82408// CHECK17-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 82409// CHECK17-NEXT:    store i64 [[AA]], ptr [[AA_ADDR]], align 82410// CHECK17-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_v_l108.omp_outlined, ptr [[AA_ADDR]])2411// CHECK17-NEXT:    ret void2412//2413//2414// CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_v_l108.omp_outlined2415// CHECK17-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[AA:%.*]]) #[[ATTR0]] {2416// CHECK17-NEXT:  entry:2417// CHECK17-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82418// CHECK17-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82419// CHECK17-NEXT:    [[AA_ADDR:%.*]] = alloca ptr, align 82420// CHECK17-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 42421// CHECK17-NEXT:    [[TMP:%.*]] = alloca i32, align 42422// CHECK17-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 42423// CHECK17-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 42424// CHECK17-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42425// CHECK17-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42426// CHECK17-NEXT:    [[I:%.*]] = alloca i32, align 42427// CHECK17-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82428// CHECK17-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82429// CHECK17-NEXT:    store ptr [[AA]], ptr [[AA_ADDR]], align 82430// CHECK17-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 82431// CHECK17-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 42432// CHECK17-NEXT:    store i32 99, ptr [[DOTOMP_UB]], align 42433// CHECK17-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 42434// CHECK17-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 42435// CHECK17-NEXT:    [[TMP1:%.*]] = load i16, ptr [[TMP0]], align 22436// CHECK17-NEXT:    [[CONV:%.*]] = sext i16 [[TMP1]] to i322437// CHECK17-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82438// CHECK17-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 42439// CHECK17-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP3]], i32 91, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 [[CONV]])2440// CHECK17-NEXT:    br label [[OMP_DISPATCH_COND:%.*]]2441// CHECK17:       omp.dispatch.cond:2442// CHECK17-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42443// CHECK17-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 992444// CHECK17-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2445// CHECK17:       cond.true:2446// CHECK17-NEXT:    br label [[COND_END:%.*]]2447// CHECK17:       cond.false:2448// CHECK17-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42449// CHECK17-NEXT:    br label [[COND_END]]2450// CHECK17:       cond.end:2451// CHECK17-NEXT:    [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]2452// CHECK17-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 42453// CHECK17-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42454// CHECK17-NEXT:    store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 42455// CHECK17-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42456// CHECK17-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42457// CHECK17-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]2458// CHECK17-NEXT:    br i1 [[CMP1]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]2459// CHECK17:       omp.dispatch.body:2460// CHECK17-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]2461// CHECK17:       omp.inner.for.cond:2462// CHECK17-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP14:![0-9]+]]2463// CHECK17-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP14]]2464// CHECK17-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP9]], [[TMP10]]2465// CHECK17-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2466// CHECK17:       omp.inner.for.body:2467// CHECK17-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP14]]2468// CHECK17-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP11]], 12469// CHECK17-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]2470// CHECK17-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP14]]2471// CHECK17-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]2472// CHECK17:       omp.body.continue:2473// CHECK17-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]2474// CHECK17:       omp.inner.for.inc:2475// CHECK17-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP14]]2476// CHECK17-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP12]], 12477// CHECK17-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP14]]2478// CHECK17-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP15:![0-9]+]]2479// CHECK17:       omp.inner.for.end:2480// CHECK17-NEXT:    br label [[OMP_DISPATCH_INC:%.*]]2481// CHECK17:       omp.dispatch.inc:2482// CHECK17-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42483// CHECK17-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 42484// CHECK17-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP13]], [[TMP14]]2485// CHECK17-NEXT:    store i32 [[ADD4]], ptr [[DOTOMP_LB]], align 42486// CHECK17-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42487// CHECK17-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 42488// CHECK17-NEXT:    [[ADD5:%.*]] = add nsw i32 [[TMP15]], [[TMP16]]2489// CHECK17-NEXT:    store i32 [[ADD5]], ptr [[DOTOMP_UB]], align 42490// CHECK17-NEXT:    br label [[OMP_DISPATCH_COND]]2491// CHECK17:       omp.dispatch.end:2492// CHECK17-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP3]])2493// CHECK17-NEXT:    ret void2494//2495//2496// CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z23without_schedule_clausePfS_S_S__l562497// CHECK19-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[A:%.*]], ptr noundef [[B:%.*]], ptr noundef [[C:%.*]], ptr noundef [[D:%.*]]) #[[ATTR0:[0-9]+]] {2498// CHECK19-NEXT:  entry:2499// CHECK19-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42500// CHECK19-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 42501// CHECK19-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 42502// CHECK19-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 42503// CHECK19-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 42504// CHECK19-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42505// CHECK19-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 42506// CHECK19-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 42507// CHECK19-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 42508// CHECK19-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 42509// CHECK19-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2:[0-9]+]], i32 4, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z23without_schedule_clausePfS_S_S__l56.omp_outlined, ptr [[A_ADDR]], ptr [[B_ADDR]], ptr [[C_ADDR]], ptr [[D_ADDR]])2510// CHECK19-NEXT:    ret void2511//2512//2513// CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z23without_schedule_clausePfS_S_S__l56.omp_outlined2514// CHECK19-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[C:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[D:%.*]]) #[[ATTR0]] {2515// CHECK19-NEXT:  entry:2516// CHECK19-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42517// CHECK19-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42518// CHECK19-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 42519// CHECK19-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 42520// CHECK19-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 42521// CHECK19-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 42522// CHECK19-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 42523// CHECK19-NEXT:    [[TMP:%.*]] = alloca i32, align 42524// CHECK19-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 42525// CHECK19-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 42526// CHECK19-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42527// CHECK19-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42528// CHECK19-NEXT:    [[I:%.*]] = alloca i32, align 42529// CHECK19-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42530// CHECK19-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42531// CHECK19-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 42532// CHECK19-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 42533// CHECK19-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 42534// CHECK19-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 42535// CHECK19-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42536// CHECK19-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 42537// CHECK19-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 42538// CHECK19-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 42539// CHECK19-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 42540// CHECK19-NEXT:    store i32 4571423, ptr [[DOTOMP_UB]], align 42541// CHECK19-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 42542// CHECK19-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 42543// CHECK19-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42544// CHECK19-NEXT:    [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 42545// CHECK19-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP5]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2546// CHECK19-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42547// CHECK19-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP6]], 45714232548// CHECK19-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2549// CHECK19:       cond.true:2550// CHECK19-NEXT:    br label [[COND_END:%.*]]2551// CHECK19:       cond.false:2552// CHECK19-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42553// CHECK19-NEXT:    br label [[COND_END]]2554// CHECK19:       cond.end:2555// CHECK19-NEXT:    [[COND:%.*]] = phi i32 [ 4571423, [[COND_TRUE]] ], [ [[TMP7]], [[COND_FALSE]] ]2556// CHECK19-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 42557// CHECK19-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42558// CHECK19-NEXT:    store i32 [[TMP8]], ptr [[DOTOMP_IV]], align 42559// CHECK19-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]2560// CHECK19:       omp.inner.for.cond:2561// CHECK19-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42562// CHECK19-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42563// CHECK19-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP9]], [[TMP10]]2564// CHECK19-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2565// CHECK19:       omp.inner.for.body:2566// CHECK19-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42567// CHECK19-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP11]], 72568// CHECK19-NEXT:    [[ADD:%.*]] = add nsw i32 33, [[MUL]]2569// CHECK19-NEXT:    store i32 [[ADD]], ptr [[I]], align 42570// CHECK19-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[TMP1]], align 42571// CHECK19-NEXT:    [[TMP13:%.*]] = load i32, ptr [[I]], align 42572// CHECK19-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds float, ptr [[TMP12]], i32 [[TMP13]]2573// CHECK19-NEXT:    [[TMP14:%.*]] = load float, ptr [[ARRAYIDX]], align 42574// CHECK19-NEXT:    [[TMP15:%.*]] = load ptr, ptr [[TMP2]], align 42575// CHECK19-NEXT:    [[TMP16:%.*]] = load i32, ptr [[I]], align 42576// CHECK19-NEXT:    [[ARRAYIDX2:%.*]] = getelementptr inbounds float, ptr [[TMP15]], i32 [[TMP16]]2577// CHECK19-NEXT:    [[TMP17:%.*]] = load float, ptr [[ARRAYIDX2]], align 42578// CHECK19-NEXT:    [[MUL3:%.*]] = fmul float [[TMP14]], [[TMP17]]2579// CHECK19-NEXT:    [[TMP18:%.*]] = load ptr, ptr [[TMP3]], align 42580// CHECK19-NEXT:    [[TMP19:%.*]] = load i32, ptr [[I]], align 42581// CHECK19-NEXT:    [[ARRAYIDX4:%.*]] = getelementptr inbounds float, ptr [[TMP18]], i32 [[TMP19]]2582// CHECK19-NEXT:    [[TMP20:%.*]] = load float, ptr [[ARRAYIDX4]], align 42583// CHECK19-NEXT:    [[MUL5:%.*]] = fmul float [[MUL3]], [[TMP20]]2584// CHECK19-NEXT:    [[TMP21:%.*]] = load ptr, ptr [[TMP0]], align 42585// CHECK19-NEXT:    [[TMP22:%.*]] = load i32, ptr [[I]], align 42586// CHECK19-NEXT:    [[ARRAYIDX6:%.*]] = getelementptr inbounds float, ptr [[TMP21]], i32 [[TMP22]]2587// CHECK19-NEXT:    store float [[MUL5]], ptr [[ARRAYIDX6]], align 42588// CHECK19-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]2589// CHECK19:       omp.body.continue:2590// CHECK19-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]2591// CHECK19:       omp.inner.for.inc:2592// CHECK19-NEXT:    [[TMP23:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42593// CHECK19-NEXT:    [[ADD7:%.*]] = add nsw i32 [[TMP23]], 12594// CHECK19-NEXT:    store i32 [[ADD7]], ptr [[DOTOMP_IV]], align 42595// CHECK19-NEXT:    br label [[OMP_INNER_FOR_COND]]2596// CHECK19:       omp.inner.for.end:2597// CHECK19-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]2598// CHECK19:       omp.loop.exit:2599// CHECK19-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP5]])2600// CHECK19-NEXT:    ret void2601//2602//2603// CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18static_not_chunkedPfS_S_S__l682604// CHECK19-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[A:%.*]], ptr noundef [[B:%.*]], ptr noundef [[C:%.*]], ptr noundef [[D:%.*]]) #[[ATTR0]] {2605// CHECK19-NEXT:  entry:2606// CHECK19-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42607// CHECK19-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 42608// CHECK19-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 42609// CHECK19-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 42610// CHECK19-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 42611// CHECK19-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42612// CHECK19-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 42613// CHECK19-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 42614// CHECK19-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 42615// CHECK19-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 42616// CHECK19-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 4, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18static_not_chunkedPfS_S_S__l68.omp_outlined, ptr [[A_ADDR]], ptr [[B_ADDR]], ptr [[C_ADDR]], ptr [[D_ADDR]])2617// CHECK19-NEXT:    ret void2618//2619//2620// CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18static_not_chunkedPfS_S_S__l68.omp_outlined2621// CHECK19-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[C:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[D:%.*]]) #[[ATTR0]] {2622// CHECK19-NEXT:  entry:2623// CHECK19-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42624// CHECK19-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42625// CHECK19-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 42626// CHECK19-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 42627// CHECK19-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 42628// CHECK19-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 42629// CHECK19-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 42630// CHECK19-NEXT:    [[TMP:%.*]] = alloca i32, align 42631// CHECK19-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 42632// CHECK19-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 42633// CHECK19-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42634// CHECK19-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42635// CHECK19-NEXT:    [[I:%.*]] = alloca i32, align 42636// CHECK19-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42637// CHECK19-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42638// CHECK19-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 42639// CHECK19-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 42640// CHECK19-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 42641// CHECK19-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 42642// CHECK19-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42643// CHECK19-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 42644// CHECK19-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 42645// CHECK19-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 42646// CHECK19-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 42647// CHECK19-NEXT:    store i32 4571423, ptr [[DOTOMP_UB]], align 42648// CHECK19-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 42649// CHECK19-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 42650// CHECK19-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42651// CHECK19-NEXT:    [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 42652// CHECK19-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP5]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2653// CHECK19-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42654// CHECK19-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP6]], 45714232655// CHECK19-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2656// CHECK19:       cond.true:2657// CHECK19-NEXT:    br label [[COND_END:%.*]]2658// CHECK19:       cond.false:2659// CHECK19-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42660// CHECK19-NEXT:    br label [[COND_END]]2661// CHECK19:       cond.end:2662// CHECK19-NEXT:    [[COND:%.*]] = phi i32 [ 4571423, [[COND_TRUE]] ], [ [[TMP7]], [[COND_FALSE]] ]2663// CHECK19-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 42664// CHECK19-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42665// CHECK19-NEXT:    store i32 [[TMP8]], ptr [[DOTOMP_IV]], align 42666// CHECK19-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]2667// CHECK19:       omp.inner.for.cond:2668// CHECK19-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42669// CHECK19-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42670// CHECK19-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP9]], [[TMP10]]2671// CHECK19-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2672// CHECK19:       omp.inner.for.body:2673// CHECK19-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42674// CHECK19-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP11]], 72675// CHECK19-NEXT:    [[SUB:%.*]] = sub nsw i32 32000000, [[MUL]]2676// CHECK19-NEXT:    store i32 [[SUB]], ptr [[I]], align 42677// CHECK19-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[TMP1]], align 42678// CHECK19-NEXT:    [[TMP13:%.*]] = load i32, ptr [[I]], align 42679// CHECK19-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds float, ptr [[TMP12]], i32 [[TMP13]]2680// CHECK19-NEXT:    [[TMP14:%.*]] = load float, ptr [[ARRAYIDX]], align 42681// CHECK19-NEXT:    [[TMP15:%.*]] = load ptr, ptr [[TMP2]], align 42682// CHECK19-NEXT:    [[TMP16:%.*]] = load i32, ptr [[I]], align 42683// CHECK19-NEXT:    [[ARRAYIDX2:%.*]] = getelementptr inbounds float, ptr [[TMP15]], i32 [[TMP16]]2684// CHECK19-NEXT:    [[TMP17:%.*]] = load float, ptr [[ARRAYIDX2]], align 42685// CHECK19-NEXT:    [[MUL3:%.*]] = fmul float [[TMP14]], [[TMP17]]2686// CHECK19-NEXT:    [[TMP18:%.*]] = load ptr, ptr [[TMP3]], align 42687// CHECK19-NEXT:    [[TMP19:%.*]] = load i32, ptr [[I]], align 42688// CHECK19-NEXT:    [[ARRAYIDX4:%.*]] = getelementptr inbounds float, ptr [[TMP18]], i32 [[TMP19]]2689// CHECK19-NEXT:    [[TMP20:%.*]] = load float, ptr [[ARRAYIDX4]], align 42690// CHECK19-NEXT:    [[MUL5:%.*]] = fmul float [[MUL3]], [[TMP20]]2691// CHECK19-NEXT:    [[TMP21:%.*]] = load ptr, ptr [[TMP0]], align 42692// CHECK19-NEXT:    [[TMP22:%.*]] = load i32, ptr [[I]], align 42693// CHECK19-NEXT:    [[ARRAYIDX6:%.*]] = getelementptr inbounds float, ptr [[TMP21]], i32 [[TMP22]]2694// CHECK19-NEXT:    store float [[MUL5]], ptr [[ARRAYIDX6]], align 42695// CHECK19-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]2696// CHECK19:       omp.body.continue:2697// CHECK19-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]2698// CHECK19:       omp.inner.for.inc:2699// CHECK19-NEXT:    [[TMP23:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42700// CHECK19-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP23]], 12701// CHECK19-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 42702// CHECK19-NEXT:    br label [[OMP_INNER_FOR_COND]]2703// CHECK19:       omp.inner.for.end:2704// CHECK19-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]2705// CHECK19:       omp.loop.exit:2706// CHECK19-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP5]])2707// CHECK19-NEXT:    ret void2708//2709//2710// CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z14static_chunkedPfS_S_S__l802711// CHECK19-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[A:%.*]], ptr noundef [[B:%.*]], ptr noundef [[C:%.*]], ptr noundef [[D:%.*]]) #[[ATTR0]] {2712// CHECK19-NEXT:  entry:2713// CHECK19-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42714// CHECK19-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 42715// CHECK19-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 42716// CHECK19-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 42717// CHECK19-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 42718// CHECK19-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42719// CHECK19-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 42720// CHECK19-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 42721// CHECK19-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 42722// CHECK19-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 42723// CHECK19-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 4, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z14static_chunkedPfS_S_S__l80.omp_outlined, ptr [[A_ADDR]], ptr [[B_ADDR]], ptr [[C_ADDR]], ptr [[D_ADDR]])2724// CHECK19-NEXT:    ret void2725//2726//2727// CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z14static_chunkedPfS_S_S__l80.omp_outlined2728// CHECK19-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[C:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[D:%.*]]) #[[ATTR0]] {2729// CHECK19-NEXT:  entry:2730// CHECK19-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42731// CHECK19-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42732// CHECK19-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 42733// CHECK19-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 42734// CHECK19-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 42735// CHECK19-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 42736// CHECK19-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 42737// CHECK19-NEXT:    [[TMP:%.*]] = alloca i32, align 42738// CHECK19-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 42739// CHECK19-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 42740// CHECK19-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42741// CHECK19-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42742// CHECK19-NEXT:    [[I:%.*]] = alloca i32, align 42743// CHECK19-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42744// CHECK19-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42745// CHECK19-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 42746// CHECK19-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 42747// CHECK19-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 42748// CHECK19-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 42749// CHECK19-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42750// CHECK19-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 42751// CHECK19-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 42752// CHECK19-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 42753// CHECK19-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 42754// CHECK19-NEXT:    store i32 16908288, ptr [[DOTOMP_UB]], align 42755// CHECK19-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 42756// CHECK19-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 42757// CHECK19-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42758// CHECK19-NEXT:    [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 42759// CHECK19-NEXT:    call void @__kmpc_for_static_init_4u(ptr @[[GLOB1]], i32 [[TMP5]], i32 91, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 5)2760// CHECK19-NEXT:    br label [[OMP_DISPATCH_COND:%.*]]2761// CHECK19:       omp.dispatch.cond:2762// CHECK19-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42763// CHECK19-NEXT:    [[CMP:%.*]] = icmp ugt i32 [[TMP6]], 169082882764// CHECK19-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2765// CHECK19:       cond.true:2766// CHECK19-NEXT:    br label [[COND_END:%.*]]2767// CHECK19:       cond.false:2768// CHECK19-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42769// CHECK19-NEXT:    br label [[COND_END]]2770// CHECK19:       cond.end:2771// CHECK19-NEXT:    [[COND:%.*]] = phi i32 [ 16908288, [[COND_TRUE]] ], [ [[TMP7]], [[COND_FALSE]] ]2772// CHECK19-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 42773// CHECK19-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42774// CHECK19-NEXT:    store i32 [[TMP8]], ptr [[DOTOMP_IV]], align 42775// CHECK19-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42776// CHECK19-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42777// CHECK19-NEXT:    [[CMP1:%.*]] = icmp ule i32 [[TMP9]], [[TMP10]]2778// CHECK19-NEXT:    br i1 [[CMP1]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]2779// CHECK19:       omp.dispatch.body:2780// CHECK19-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]2781// CHECK19:       omp.inner.for.cond:2782// CHECK19-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12:![0-9]+]]2783// CHECK19-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP12]]2784// CHECK19-NEXT:    [[CMP2:%.*]] = icmp ule i32 [[TMP11]], [[TMP12]]2785// CHECK19-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2786// CHECK19:       omp.inner.for.body:2787// CHECK19-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]2788// CHECK19-NEXT:    [[MUL:%.*]] = mul i32 [[TMP13]], 1272789// CHECK19-NEXT:    [[ADD:%.*]] = add i32 131071, [[MUL]]2790// CHECK19-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP12]]2791// CHECK19-NEXT:    [[TMP14:%.*]] = load ptr, ptr [[TMP1]], align 4, !llvm.access.group [[ACC_GRP12]]2792// CHECK19-NEXT:    [[TMP15:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP12]]2793// CHECK19-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds nuw float, ptr [[TMP14]], i32 [[TMP15]]2794// CHECK19-NEXT:    [[TMP16:%.*]] = load float, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP12]]2795// CHECK19-NEXT:    [[TMP17:%.*]] = load ptr, ptr [[TMP2]], align 4, !llvm.access.group [[ACC_GRP12]]2796// CHECK19-NEXT:    [[TMP18:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP12]]2797// CHECK19-NEXT:    [[ARRAYIDX3:%.*]] = getelementptr inbounds nuw float, ptr [[TMP17]], i32 [[TMP18]]2798// CHECK19-NEXT:    [[TMP19:%.*]] = load float, ptr [[ARRAYIDX3]], align 4, !llvm.access.group [[ACC_GRP12]]2799// CHECK19-NEXT:    [[MUL4:%.*]] = fmul float [[TMP16]], [[TMP19]]2800// CHECK19-NEXT:    [[TMP20:%.*]] = load ptr, ptr [[TMP3]], align 4, !llvm.access.group [[ACC_GRP12]]2801// CHECK19-NEXT:    [[TMP21:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP12]]2802// CHECK19-NEXT:    [[ARRAYIDX5:%.*]] = getelementptr inbounds nuw float, ptr [[TMP20]], i32 [[TMP21]]2803// CHECK19-NEXT:    [[TMP22:%.*]] = load float, ptr [[ARRAYIDX5]], align 4, !llvm.access.group [[ACC_GRP12]]2804// CHECK19-NEXT:    [[MUL6:%.*]] = fmul float [[MUL4]], [[TMP22]]2805// CHECK19-NEXT:    [[TMP23:%.*]] = load ptr, ptr [[TMP0]], align 4, !llvm.access.group [[ACC_GRP12]]2806// CHECK19-NEXT:    [[TMP24:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP12]]2807// CHECK19-NEXT:    [[ARRAYIDX7:%.*]] = getelementptr inbounds nuw float, ptr [[TMP23]], i32 [[TMP24]]2808// CHECK19-NEXT:    store float [[MUL6]], ptr [[ARRAYIDX7]], align 4, !llvm.access.group [[ACC_GRP12]]2809// CHECK19-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]2810// CHECK19:       omp.body.continue:2811// CHECK19-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]2812// CHECK19:       omp.inner.for.inc:2813// CHECK19-NEXT:    [[TMP25:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]2814// CHECK19-NEXT:    [[ADD8:%.*]] = add i32 [[TMP25]], 12815// CHECK19-NEXT:    store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]2816// CHECK19-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP13:![0-9]+]]2817// CHECK19:       omp.inner.for.end:2818// CHECK19-NEXT:    br label [[OMP_DISPATCH_INC:%.*]]2819// CHECK19:       omp.dispatch.inc:2820// CHECK19-NEXT:    [[TMP26:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42821// CHECK19-NEXT:    [[TMP27:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 42822// CHECK19-NEXT:    [[ADD9:%.*]] = add i32 [[TMP26]], [[TMP27]]2823// CHECK19-NEXT:    store i32 [[ADD9]], ptr [[DOTOMP_LB]], align 42824// CHECK19-NEXT:    [[TMP28:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42825// CHECK19-NEXT:    [[TMP29:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 42826// CHECK19-NEXT:    [[ADD10:%.*]] = add i32 [[TMP28]], [[TMP29]]2827// CHECK19-NEXT:    store i32 [[ADD10]], ptr [[DOTOMP_UB]], align 42828// CHECK19-NEXT:    br label [[OMP_DISPATCH_COND]]2829// CHECK19:       omp.dispatch.end:2830// CHECK19-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP5]])2831// CHECK19-NEXT:    ret void2832//2833//2834// CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z12test_precondv_l922835// CHECK19-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[A:%.*]]) #[[ATTR0]] {2836// CHECK19-NEXT:  entry:2837// CHECK19-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42838// CHECK19-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 42839// CHECK19-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42840// CHECK19-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 42841// CHECK19-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z12test_precondv_l92.omp_outlined, ptr [[A_ADDR]])2842// CHECK19-NEXT:    ret void2843//2844//2845// CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z12test_precondv_l92.omp_outlined2846// CHECK19-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 1 dereferenceable(1) [[A:%.*]]) #[[ATTR0]] {2847// CHECK19-NEXT:  entry:2848// CHECK19-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42849// CHECK19-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42850// CHECK19-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 42851// CHECK19-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 42852// CHECK19-NEXT:    [[TMP:%.*]] = alloca i8, align 12853// CHECK19-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i8, align 12854// CHECK19-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42855// CHECK19-NEXT:    [[I:%.*]] = alloca i8, align 12856// CHECK19-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 42857// CHECK19-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 42858// CHECK19-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42859// CHECK19-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42860// CHECK19-NEXT:    [[I5:%.*]] = alloca i8, align 12861// CHECK19-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42862// CHECK19-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42863// CHECK19-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 42864// CHECK19-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42865// CHECK19-NEXT:    [[TMP1:%.*]] = load i8, ptr [[TMP0]], align 12866// CHECK19-NEXT:    store i8 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 12867// CHECK19-NEXT:    [[TMP2:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 12868// CHECK19-NEXT:    [[CONV:%.*]] = sext i8 [[TMP2]] to i322869// CHECK19-NEXT:    [[SUB:%.*]] = sub i32 10, [[CONV]]2870// CHECK19-NEXT:    [[SUB2:%.*]] = sub i32 [[SUB]], 12871// CHECK19-NEXT:    [[ADD:%.*]] = add i32 [[SUB2]], 12872// CHECK19-NEXT:    [[DIV:%.*]] = udiv i32 [[ADD]], 12873// CHECK19-NEXT:    [[SUB3:%.*]] = sub nsw i32 [[DIV]], 12874// CHECK19-NEXT:    store i32 [[SUB3]], ptr [[DOTCAPTURE_EXPR_1]], align 42875// CHECK19-NEXT:    [[TMP3:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 12876// CHECK19-NEXT:    store i8 [[TMP3]], ptr [[I]], align 12877// CHECK19-NEXT:    [[TMP4:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 12878// CHECK19-NEXT:    [[CONV4:%.*]] = sext i8 [[TMP4]] to i322879// CHECK19-NEXT:    [[CMP:%.*]] = icmp slt i32 [[CONV4]], 102880// CHECK19-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2881// CHECK19:       omp.precond.then:2882// CHECK19-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 42883// CHECK19-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42884// CHECK19-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_UB]], align 42885// CHECK19-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 42886// CHECK19-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 42887// CHECK19-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42888// CHECK19-NEXT:    [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 42889// CHECK19-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP7]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2890// CHECK19-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42891// CHECK19-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42892// CHECK19-NEXT:    [[CMP6:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]2893// CHECK19-NEXT:    br i1 [[CMP6]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2894// CHECK19:       cond.true:2895// CHECK19-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42896// CHECK19-NEXT:    br label [[COND_END:%.*]]2897// CHECK19:       cond.false:2898// CHECK19-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42899// CHECK19-NEXT:    br label [[COND_END]]2900// CHECK19:       cond.end:2901// CHECK19-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]2902// CHECK19-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 42903// CHECK19-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42904// CHECK19-NEXT:    store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 42905// CHECK19-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]2906// CHECK19:       omp.inner.for.cond:2907// CHECK19-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42908// CHECK19-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42909// CHECK19-NEXT:    [[CMP7:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]2910// CHECK19-NEXT:    br i1 [[CMP7]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2911// CHECK19:       omp.inner.for.body:2912// CHECK19-NEXT:    [[TMP15:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 12913// CHECK19-NEXT:    [[CONV8:%.*]] = sext i8 [[TMP15]] to i322914// CHECK19-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42915// CHECK19-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP16]], 12916// CHECK19-NEXT:    [[ADD9:%.*]] = add nsw i32 [[CONV8]], [[MUL]]2917// CHECK19-NEXT:    [[CONV10:%.*]] = trunc i32 [[ADD9]] to i82918// CHECK19-NEXT:    store i8 [[CONV10]], ptr [[I5]], align 12919// CHECK19-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]2920// CHECK19:       omp.body.continue:2921// CHECK19-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]2922// CHECK19:       omp.inner.for.inc:2923// CHECK19-NEXT:    [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42924// CHECK19-NEXT:    [[ADD11:%.*]] = add nsw i32 [[TMP17]], 12925// CHECK19-NEXT:    store i32 [[ADD11]], ptr [[DOTOMP_IV]], align 42926// CHECK19-NEXT:    br label [[OMP_INNER_FOR_COND]]2927// CHECK19:       omp.inner.for.end:2928// CHECK19-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]2929// CHECK19:       omp.loop.exit:2930// CHECK19-NEXT:    [[TMP18:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42931// CHECK19-NEXT:    [[TMP19:%.*]] = load i32, ptr [[TMP18]], align 42932// CHECK19-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP19]])2933// CHECK19-NEXT:    br label [[OMP_PRECOND_END]]2934// CHECK19:       omp.precond.end:2935// CHECK19-NEXT:    ret void2936//2937//2938// CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_v_l1082939// CHECK19-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[AA:%.*]]) #[[ATTR0]] {2940// CHECK19-NEXT:  entry:2941// CHECK19-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42942// CHECK19-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 42943// CHECK19-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42944// CHECK19-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 42945// CHECK19-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_v_l108.omp_outlined, ptr [[AA_ADDR]])2946// CHECK19-NEXT:    ret void2947//2948//2949// CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_v_l108.omp_outlined2950// CHECK19-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[AA:%.*]]) #[[ATTR0]] {2951// CHECK19-NEXT:  entry:2952// CHECK19-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42953// CHECK19-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42954// CHECK19-NEXT:    [[AA_ADDR:%.*]] = alloca ptr, align 42955// CHECK19-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 42956// CHECK19-NEXT:    [[TMP:%.*]] = alloca i32, align 42957// CHECK19-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 42958// CHECK19-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 42959// CHECK19-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42960// CHECK19-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42961// CHECK19-NEXT:    [[I:%.*]] = alloca i32, align 42962// CHECK19-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42963// CHECK19-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42964// CHECK19-NEXT:    store ptr [[AA]], ptr [[AA_ADDR]], align 42965// CHECK19-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 42966// CHECK19-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 42967// CHECK19-NEXT:    store i32 99, ptr [[DOTOMP_UB]], align 42968// CHECK19-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 42969// CHECK19-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 42970// CHECK19-NEXT:    [[TMP1:%.*]] = load i16, ptr [[TMP0]], align 22971// CHECK19-NEXT:    [[CONV:%.*]] = sext i16 [[TMP1]] to i322972// CHECK19-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42973// CHECK19-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 42974// CHECK19-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP3]], i32 91, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 [[CONV]])2975// CHECK19-NEXT:    br label [[OMP_DISPATCH_COND:%.*]]2976// CHECK19:       omp.dispatch.cond:2977// CHECK19-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42978// CHECK19-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 992979// CHECK19-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2980// CHECK19:       cond.true:2981// CHECK19-NEXT:    br label [[COND_END:%.*]]2982// CHECK19:       cond.false:2983// CHECK19-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42984// CHECK19-NEXT:    br label [[COND_END]]2985// CHECK19:       cond.end:2986// CHECK19-NEXT:    [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]2987// CHECK19-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 42988// CHECK19-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42989// CHECK19-NEXT:    store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 42990// CHECK19-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42991// CHECK19-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42992// CHECK19-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]2993// CHECK19-NEXT:    br i1 [[CMP1]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]2994// CHECK19:       omp.dispatch.body:2995// CHECK19-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]2996// CHECK19:       omp.inner.for.cond:2997// CHECK19-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP15:![0-9]+]]2998// CHECK19-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP15]]2999// CHECK19-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP9]], [[TMP10]]3000// CHECK19-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]3001// CHECK19:       omp.inner.for.body:3002// CHECK19-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP15]]3003// CHECK19-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP11]], 13004// CHECK19-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]3005// CHECK19-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP15]]3006// CHECK19-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]3007// CHECK19:       omp.body.continue:3008// CHECK19-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]3009// CHECK19:       omp.inner.for.inc:3010// CHECK19-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP15]]3011// CHECK19-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP12]], 13012// CHECK19-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP15]]3013// CHECK19-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP16:![0-9]+]]3014// CHECK19:       omp.inner.for.end:3015// CHECK19-NEXT:    br label [[OMP_DISPATCH_INC:%.*]]3016// CHECK19:       omp.dispatch.inc:3017// CHECK19-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 43018// CHECK19-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 43019// CHECK19-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP13]], [[TMP14]]3020// CHECK19-NEXT:    store i32 [[ADD4]], ptr [[DOTOMP_LB]], align 43021// CHECK19-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 43022// CHECK19-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 43023// CHECK19-NEXT:    [[ADD5:%.*]] = add nsw i32 [[TMP15]], [[TMP16]]3024// CHECK19-NEXT:    store i32 [[ADD5]], ptr [[DOTOMP_UB]], align 43025// CHECK19-NEXT:    br label [[OMP_DISPATCH_COND]]3026// CHECK19:       omp.dispatch.end:3027// CHECK19-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP3]])3028// CHECK19-NEXT:    ret void3029//3030