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// expected-no-diagnostics3#ifndef HEADER4#define HEADER5 6// Test host codegen.7// RUN: %clang_cc1 -DCK1 -verify -Wno-vla -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK18// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s9// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK110// RUN: %clang_cc1 -DCK1 -verify -Wno-vla -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK311// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s12// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK313 14// RUN: %clang_cc1 -DCK1 -verify -Wno-vla -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"15// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s16// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"17// RUN: %clang_cc1 -DCK1 -verify -Wno-vla -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"18// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s19// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"20#ifdef CK121 22template <typename T, int X, long long Y>23struct SS{24  T a[X][Y];25 26  int foo(void) {27 28    #pragma omp target teams loop collapse(2)29    for(int i = 0; i < X; i++) {30      for(int j = 0; j < Y; j++) {31        a[i][j] = (T)0;32      }33    }34 35    // discard loop variables not needed here36 37 38    return a[0][0];39  }40};41 42int teams_template_struct(void) {43  SS<int, 123, 456> V;44  return V.foo();45 46}47#endif // CK148 49// Test host codegen.50// RUN: %clang_cc1 -DCK2 -verify -Wno-vla -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK951// RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s52// RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK953// RUN: %clang_cc1 -DCK2 -verify -Wno-vla -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK1154// RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s55// RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK1156 57// RUN: %clang_cc1 -DCK2 -verify -Wno-vla -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}}"58// RUN: %clang_cc1 -DCK2 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s59// RUN: %clang_cc1 -DCK2 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"60// RUN: %clang_cc1 -DCK2 -verify -Wno-vla -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}}"61// RUN: %clang_cc1 -DCK2 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s62// RUN: %clang_cc1 -DCK2 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"63#ifdef CK264 65template <typename T, int n, int m>66int tmain(T argc) {67  T a[n][m];68  #pragma omp target teams loop collapse(2)69  for(int i = 0; i < n; i++) {70    for(int j = 0; j < m; j++) {71      a[i][j] = (T)0;72    }73  }74  return 0;75}76 77int main (int argc, char **argv) {78  int n = 100;79  int m = 2;80  int a[n][m];81  #pragma omp target teams loop collapse(2)82  for(int i = 0; i < n; i++) {83    for(int j = 0; j < m; j++) {84      a[i][j] = 0;85    }86  }87  return tmain<int, 10, 2>(argc);88}89 90 91 92 93 94 95 96 97// discard loop variables not needed here98 99 100#endif // CK2101#endif // #ifndef HEADER102// CHECK1-LABEL: define {{[^@]+}}@_Z21teams_template_structv103// CHECK1-SAME: () #[[ATTR0:[0-9]+]] {104// CHECK1-NEXT:  entry:105// CHECK1-NEXT:    [[V:%.*]] = alloca [[STRUCT_SS:%.*]], align 4106// CHECK1-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_ZN2SSIiLi123ELx456EE3fooEv(ptr noundef nonnull align 4 dereferenceable(224352) [[V]])107// CHECK1-NEXT:    ret i32 [[CALL]]108//109//110// CHECK1-LABEL: define {{[^@]+}}@_ZN2SSIiLi123ELx456EE3fooEv111// CHECK1-SAME: (ptr noundef nonnull align 4 dereferenceable(224352) [[THIS:%.*]]) #[[ATTR0]] comdat {112// CHECK1-NEXT:  entry:113// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8114// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8115// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8116// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8117// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4118// CHECK1-NEXT:    [[_TMP2:%.*]] = alloca i32, align 4119// CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8120// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8121// CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8122// CHECK1-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[THIS1]], i32 0, i32 0123// CHECK1-NEXT:    [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0124// CHECK1-NEXT:    store ptr [[THIS1]], ptr [[TMP0]], align 8125// CHECK1-NEXT:    [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0126// CHECK1-NEXT:    store ptr [[A]], ptr [[TMP1]], align 8127// CHECK1-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0128// CHECK1-NEXT:    store ptr null, ptr [[TMP2]], align 8129// CHECK1-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0130// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0131// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0132// CHECK1-NEXT:    store i32 3, ptr [[TMP5]], align 4133// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1134// CHECK1-NEXT:    store i32 1, ptr [[TMP6]], align 4135// CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2136// CHECK1-NEXT:    store ptr [[TMP3]], ptr [[TMP7]], align 8137// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3138// CHECK1-NEXT:    store ptr [[TMP4]], ptr [[TMP8]], align 8139// CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4140// CHECK1-NEXT:    store ptr @.offload_sizes, ptr [[TMP9]], align 8141// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5142// CHECK1-NEXT:    store ptr @.offload_maptypes, ptr [[TMP10]], align 8143// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6144// CHECK1-NEXT:    store ptr null, ptr [[TMP11]], align 8145// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7146// CHECK1-NEXT:    store ptr null, ptr [[TMP12]], align 8147// CHECK1-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8148// CHECK1-NEXT:    store i64 56088, ptr [[TMP13]], align 8149// CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9150// CHECK1-NEXT:    store i64 0, ptr [[TMP14]], align 8151// CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10152// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4153// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11154// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4155// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12156// CHECK1-NEXT:    store i32 0, ptr [[TMP17]], align 4157// CHECK1-NEXT:    [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28.region_id, ptr [[KERNEL_ARGS]])158// CHECK1-NEXT:    [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0159// CHECK1-NEXT:    br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]160// CHECK1:       omp_offload.failed:161// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28(ptr [[THIS1]]) #[[ATTR2:[0-9]+]]162// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]163// CHECK1:       omp_offload.cont:164// CHECK1-NEXT:    [[A3:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 0165// CHECK1-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x [456 x i32]], ptr [[A3]], i64 0, i64 0166// CHECK1-NEXT:    [[ARRAYIDX4:%.*]] = getelementptr inbounds [456 x i32], ptr [[ARRAYIDX]], i64 0, i64 0167// CHECK1-NEXT:    [[TMP20:%.*]] = load i32, ptr [[ARRAYIDX4]], align 4168// CHECK1-NEXT:    ret i32 [[TMP20]]169//170//171// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28172// CHECK1-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1:[0-9]+]] {173// CHECK1-NEXT:  entry:174// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8175// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8176// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8177// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28.omp_outlined, ptr [[TMP0]])178// CHECK1-NEXT:    ret void179//180//181// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28.omp_outlined182// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {183// CHECK1-NEXT:  entry:184// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8185// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8186// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8187// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4188// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4189// CHECK1-NEXT:    [[_TMP1:%.*]] = alloca i32, align 4190// CHECK1-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4191// CHECK1-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4192// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4193// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4194// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4195// CHECK1-NEXT:    [[J:%.*]] = alloca i32, align 4196// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8197// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8198// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8199// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8200// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 4201// CHECK1-NEXT:    store i32 56087, ptr [[DOTOMP_COMB_UB]], align 4202// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4203// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4204// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8205// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4206// CHECK1-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)207// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4208// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 56087209// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]210// CHECK1:       cond.true:211// CHECK1-NEXT:    br label [[COND_END:%.*]]212// CHECK1:       cond.false:213// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4214// CHECK1-NEXT:    br label [[COND_END]]215// CHECK1:       cond.end:216// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 56087, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]217// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4218// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4219// CHECK1-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4220// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]221// CHECK1:       omp.inner.for.cond:222// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4223// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4224// CHECK1-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]225// CHECK1-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]226// CHECK1:       omp.inner.for.body:227// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4228// CHECK1-NEXT:    [[TMP9:%.*]] = zext i32 [[TMP8]] to i64229// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4230// CHECK1-NEXT:    [[TMP11:%.*]] = zext i32 [[TMP10]] to i64231// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28.omp_outlined.omp_outlined, i64 [[TMP9]], i64 [[TMP11]], ptr [[TMP0]])232// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]233// CHECK1:       omp.inner.for.inc:234// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4235// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4236// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP12]], [[TMP13]]237// CHECK1-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4238// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]239// CHECK1:       omp.inner.for.end:240// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]241// CHECK1:       omp.loop.exit:242// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])243// CHECK1-NEXT:    ret void244//245//246// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28.omp_outlined.omp_outlined247// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {248// CHECK1-NEXT:  entry:249// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8250// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8251// CHECK1-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8252// CHECK1-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8253// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8254// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4255// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4256// CHECK1-NEXT:    [[_TMP1:%.*]] = alloca i32, align 4257// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4258// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4259// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4260// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4261// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4262// CHECK1-NEXT:    [[J:%.*]] = alloca i32, align 4263// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8264// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8265// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8266// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8267// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8268// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8269// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4270// CHECK1-NEXT:    store i32 56087, ptr [[DOTOMP_UB]], align 4271// CHECK1-NEXT:    [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8272// CHECK1-NEXT:    [[CONV:%.*]] = trunc i64 [[TMP1]] to i32273// CHECK1-NEXT:    [[TMP2:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8274// CHECK1-NEXT:    [[CONV2:%.*]] = trunc i64 [[TMP2]] to i32275// CHECK1-NEXT:    store i32 [[CONV]], ptr [[DOTOMP_LB]], align 4276// CHECK1-NEXT:    store i32 [[CONV2]], ptr [[DOTOMP_UB]], align 4277// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4278// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4279// CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8280// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[TMP3]], align 4281// CHECK1-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP4]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)282// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4283// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP5]], 56087284// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]285// CHECK1:       cond.true:286// CHECK1-NEXT:    br label [[COND_END:%.*]]287// CHECK1:       cond.false:288// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4289// CHECK1-NEXT:    br label [[COND_END]]290// CHECK1:       cond.end:291// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 56087, [[COND_TRUE]] ], [ [[TMP6]], [[COND_FALSE]] ]292// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4293// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4294// CHECK1-NEXT:    store i32 [[TMP7]], ptr [[DOTOMP_IV]], align 4295// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]296// CHECK1:       omp.inner.for.cond:297// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4298// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4299// CHECK1-NEXT:    [[CMP3:%.*]] = icmp sle i32 [[TMP8]], [[TMP9]]300// CHECK1-NEXT:    br i1 [[CMP3]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]301// CHECK1:       omp.inner.for.body:302// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4303// CHECK1-NEXT:    [[DIV:%.*]] = sdiv i32 [[TMP10]], 456304// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[DIV]], 1305// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]306// CHECK1-NEXT:    store i32 [[ADD]], ptr [[I]], align 4307// CHECK1-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4308// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4309// CHECK1-NEXT:    [[DIV4:%.*]] = sdiv i32 [[TMP12]], 456310// CHECK1-NEXT:    [[MUL5:%.*]] = mul nsw i32 [[DIV4]], 456311// CHECK1-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP11]], [[MUL5]]312// CHECK1-NEXT:    [[MUL6:%.*]] = mul nsw i32 [[SUB]], 1313// CHECK1-NEXT:    [[ADD7:%.*]] = add nsw i32 0, [[MUL6]]314// CHECK1-NEXT:    store i32 [[ADD7]], ptr [[J]], align 4315// CHECK1-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP0]], i32 0, i32 0316// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[I]], align 4317// CHECK1-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP13]] to i64318// CHECK1-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x [456 x i32]], ptr [[A]], i64 0, i64 [[IDXPROM]]319// CHECK1-NEXT:    [[TMP14:%.*]] = load i32, ptr [[J]], align 4320// CHECK1-NEXT:    [[IDXPROM8:%.*]] = sext i32 [[TMP14]] to i64321// CHECK1-NEXT:    [[ARRAYIDX9:%.*]] = getelementptr inbounds [456 x i32], ptr [[ARRAYIDX]], i64 0, i64 [[IDXPROM8]]322// CHECK1-NEXT:    store i32 0, ptr [[ARRAYIDX9]], align 4323// CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]324// CHECK1:       omp.body.continue:325// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]326// CHECK1:       omp.inner.for.inc:327// CHECK1-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4328// CHECK1-NEXT:    [[ADD10:%.*]] = add nsw i32 [[TMP15]], 1329// CHECK1-NEXT:    store i32 [[ADD10]], ptr [[DOTOMP_IV]], align 4330// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]331// CHECK1:       omp.inner.for.end:332// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]333// CHECK1:       omp.loop.exit:334// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP4]])335// CHECK1-NEXT:    ret void336//337//338// CHECK3-LABEL: define {{[^@]+}}@_Z21teams_template_structv339// CHECK3-SAME: () #[[ATTR0:[0-9]+]] {340// CHECK3-NEXT:  entry:341// CHECK3-NEXT:    [[V:%.*]] = alloca [[STRUCT_SS:%.*]], align 4342// CHECK3-NEXT:    [[CALL:%.*]] = call noundef i32 @_ZN2SSIiLi123ELx456EE3fooEv(ptr noundef nonnull align 4 dereferenceable(224352) [[V]])343// CHECK3-NEXT:    ret i32 [[CALL]]344//345//346// CHECK3-LABEL: define {{[^@]+}}@_ZN2SSIiLi123ELx456EE3fooEv347// CHECK3-SAME: (ptr noundef nonnull align 4 dereferenceable(224352) [[THIS:%.*]]) #[[ATTR0]] comdat align 2 {348// CHECK3-NEXT:  entry:349// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4350// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4351// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4352// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4353// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 4354// CHECK3-NEXT:    [[_TMP2:%.*]] = alloca i32, align 4355// CHECK3-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8356// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4357// CHECK3-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4358// CHECK3-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[THIS1]], i32 0, i32 0359// CHECK3-NEXT:    [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0360// CHECK3-NEXT:    store ptr [[THIS1]], ptr [[TMP0]], align 4361// CHECK3-NEXT:    [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0362// CHECK3-NEXT:    store ptr [[A]], ptr [[TMP1]], align 4363// CHECK3-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0364// CHECK3-NEXT:    store ptr null, ptr [[TMP2]], align 4365// CHECK3-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0366// CHECK3-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0367// CHECK3-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0368// CHECK3-NEXT:    store i32 3, ptr [[TMP5]], align 4369// CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1370// CHECK3-NEXT:    store i32 1, ptr [[TMP6]], align 4371// CHECK3-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2372// CHECK3-NEXT:    store ptr [[TMP3]], ptr [[TMP7]], align 4373// CHECK3-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3374// CHECK3-NEXT:    store ptr [[TMP4]], ptr [[TMP8]], align 4375// CHECK3-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4376// CHECK3-NEXT:    store ptr @.offload_sizes, ptr [[TMP9]], align 4377// CHECK3-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5378// CHECK3-NEXT:    store ptr @.offload_maptypes, ptr [[TMP10]], align 4379// CHECK3-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6380// CHECK3-NEXT:    store ptr null, ptr [[TMP11]], align 4381// CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7382// CHECK3-NEXT:    store ptr null, ptr [[TMP12]], align 4383// CHECK3-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8384// CHECK3-NEXT:    store i64 56088, ptr [[TMP13]], align 8385// CHECK3-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9386// CHECK3-NEXT:    store i64 0, ptr [[TMP14]], align 8387// CHECK3-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10388// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4389// CHECK3-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11390// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4391// CHECK3-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12392// CHECK3-NEXT:    store i32 0, ptr [[TMP17]], align 4393// CHECK3-NEXT:    [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28.region_id, ptr [[KERNEL_ARGS]])394// CHECK3-NEXT:    [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0395// CHECK3-NEXT:    br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]396// CHECK3:       omp_offload.failed:397// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28(ptr [[THIS1]]) #[[ATTR2:[0-9]+]]398// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT]]399// CHECK3:       omp_offload.cont:400// CHECK3-NEXT:    [[A3:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 0401// CHECK3-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x [456 x i32]], ptr [[A3]], i32 0, i32 0402// CHECK3-NEXT:    [[ARRAYIDX4:%.*]] = getelementptr inbounds [456 x i32], ptr [[ARRAYIDX]], i32 0, i32 0403// CHECK3-NEXT:    [[TMP20:%.*]] = load i32, ptr [[ARRAYIDX4]], align 4404// CHECK3-NEXT:    ret i32 [[TMP20]]405//406//407// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28408// CHECK3-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1:[0-9]+]] {409// CHECK3-NEXT:  entry:410// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4411// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4412// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4413// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28.omp_outlined, ptr [[TMP0]])414// CHECK3-NEXT:    ret void415//416//417// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28.omp_outlined418// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {419// CHECK3-NEXT:  entry:420// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4421// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4422// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4423// CHECK3-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4424// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 4425// CHECK3-NEXT:    [[_TMP1:%.*]] = alloca i32, align 4426// CHECK3-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4427// CHECK3-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4428// CHECK3-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4429// CHECK3-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4430// CHECK3-NEXT:    [[I:%.*]] = alloca i32, align 4431// CHECK3-NEXT:    [[J:%.*]] = alloca i32, align 4432// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4433// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4434// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4435// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4436// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 4437// CHECK3-NEXT:    store i32 56087, ptr [[DOTOMP_COMB_UB]], align 4438// CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4439// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4440// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4441// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4442// CHECK3-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)443// CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4444// CHECK3-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 56087445// CHECK3-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]446// CHECK3:       cond.true:447// CHECK3-NEXT:    br label [[COND_END:%.*]]448// CHECK3:       cond.false:449// CHECK3-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4450// CHECK3-NEXT:    br label [[COND_END]]451// CHECK3:       cond.end:452// CHECK3-NEXT:    [[COND:%.*]] = phi i32 [ 56087, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]453// CHECK3-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4454// CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4455// CHECK3-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4456// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]457// CHECK3:       omp.inner.for.cond:458// CHECK3-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4459// CHECK3-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4460// CHECK3-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]461// CHECK3-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]462// CHECK3:       omp.inner.for.body:463// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4464// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4465// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28.omp_outlined.omp_outlined, i32 [[TMP8]], i32 [[TMP9]], ptr [[TMP0]])466// CHECK3-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]467// CHECK3:       omp.inner.for.inc:468// CHECK3-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4469// CHECK3-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4470// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP10]], [[TMP11]]471// CHECK3-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4472// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND]]473// CHECK3:       omp.inner.for.end:474// CHECK3-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]475// CHECK3:       omp.loop.exit:476// CHECK3-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])477// CHECK3-NEXT:    ret void478//479//480// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28.omp_outlined.omp_outlined481// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[DOTPREVIOUS_LB_:%.*]], i32 noundef [[DOTPREVIOUS_UB_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {482// CHECK3-NEXT:  entry:483// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4484// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4485// CHECK3-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i32, align 4486// CHECK3-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i32, align 4487// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4488// CHECK3-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4489// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 4490// CHECK3-NEXT:    [[_TMP1:%.*]] = alloca i32, align 4491// CHECK3-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4492// CHECK3-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4493// CHECK3-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4494// CHECK3-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4495// CHECK3-NEXT:    [[I:%.*]] = alloca i32, align 4496// CHECK3-NEXT:    [[J:%.*]] = alloca i32, align 4497// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4498// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4499// CHECK3-NEXT:    store i32 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 4500// CHECK3-NEXT:    store i32 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 4501// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4502// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4503// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4504// CHECK3-NEXT:    store i32 56087, ptr [[DOTOMP_UB]], align 4505// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTPREVIOUS_LB__ADDR]], align 4506// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTPREVIOUS_UB__ADDR]], align 4507// CHECK3-NEXT:    store i32 [[TMP1]], ptr [[DOTOMP_LB]], align 4508// CHECK3-NEXT:    store i32 [[TMP2]], ptr [[DOTOMP_UB]], align 4509// CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4510// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4511// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4512// CHECK3-NEXT:    [[TMP4:%.*]] = load i32, ptr [[TMP3]], align 4513// CHECK3-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP4]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)514// CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4515// CHECK3-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP5]], 56087516// CHECK3-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]517// CHECK3:       cond.true:518// CHECK3-NEXT:    br label [[COND_END:%.*]]519// CHECK3:       cond.false:520// CHECK3-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4521// CHECK3-NEXT:    br label [[COND_END]]522// CHECK3:       cond.end:523// CHECK3-NEXT:    [[COND:%.*]] = phi i32 [ 56087, [[COND_TRUE]] ], [ [[TMP6]], [[COND_FALSE]] ]524// CHECK3-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4525// CHECK3-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4526// CHECK3-NEXT:    store i32 [[TMP7]], ptr [[DOTOMP_IV]], align 4527// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]528// CHECK3:       omp.inner.for.cond:529// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4530// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4531// CHECK3-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP8]], [[TMP9]]532// CHECK3-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]533// CHECK3:       omp.inner.for.body:534// CHECK3-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4535// CHECK3-NEXT:    [[DIV:%.*]] = sdiv i32 [[TMP10]], 456536// CHECK3-NEXT:    [[MUL:%.*]] = mul nsw i32 [[DIV]], 1537// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]538// CHECK3-NEXT:    store i32 [[ADD]], ptr [[I]], align 4539// CHECK3-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4540// CHECK3-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4541// CHECK3-NEXT:    [[DIV3:%.*]] = sdiv i32 [[TMP12]], 456542// CHECK3-NEXT:    [[MUL4:%.*]] = mul nsw i32 [[DIV3]], 456543// CHECK3-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP11]], [[MUL4]]544// CHECK3-NEXT:    [[MUL5:%.*]] = mul nsw i32 [[SUB]], 1545// CHECK3-NEXT:    [[ADD6:%.*]] = add nsw i32 0, [[MUL5]]546// CHECK3-NEXT:    store i32 [[ADD6]], ptr [[J]], align 4547// CHECK3-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP0]], i32 0, i32 0548// CHECK3-NEXT:    [[TMP13:%.*]] = load i32, ptr [[I]], align 4549// CHECK3-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x [456 x i32]], ptr [[A]], i32 0, i32 [[TMP13]]550// CHECK3-NEXT:    [[TMP14:%.*]] = load i32, ptr [[J]], align 4551// CHECK3-NEXT:    [[ARRAYIDX7:%.*]] = getelementptr inbounds [456 x i32], ptr [[ARRAYIDX]], i32 0, i32 [[TMP14]]552// CHECK3-NEXT:    store i32 0, ptr [[ARRAYIDX7]], align 4553// CHECK3-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]554// CHECK3:       omp.body.continue:555// CHECK3-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]556// CHECK3:       omp.inner.for.inc:557// CHECK3-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4558// CHECK3-NEXT:    [[ADD8:%.*]] = add nsw i32 [[TMP15]], 1559// CHECK3-NEXT:    store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 4560// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND]]561// CHECK3:       omp.inner.for.end:562// CHECK3-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]563// CHECK3:       omp.loop.exit:564// CHECK3-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP4]])565// CHECK3-NEXT:    ret void566//567//568// CHECK9-LABEL: define {{[^@]+}}@main569// CHECK9-SAME: (i32 noundef signext [[ARGC:%.*]], ptr noundef [[ARGV:%.*]]) #[[ATTR0:[0-9]+]] {570// CHECK9-NEXT:  entry:571// CHECK9-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4572// CHECK9-NEXT:    [[ARGC_ADDR:%.*]] = alloca i32, align 4573// CHECK9-NEXT:    [[ARGV_ADDR:%.*]] = alloca ptr, align 8574// CHECK9-NEXT:    [[N:%.*]] = alloca i32, align 4575// CHECK9-NEXT:    [[M:%.*]] = alloca i32, align 4576// CHECK9-NEXT:    [[SAVED_STACK:%.*]] = alloca ptr, align 8577// CHECK9-NEXT:    [[__VLA_EXPR0:%.*]] = alloca i64, align 8578// CHECK9-NEXT:    [[__VLA_EXPR1:%.*]] = alloca i64, align 8579// CHECK9-NEXT:    [[N_CASTED:%.*]] = alloca i64, align 8580// CHECK9-NEXT:    [[M_CASTED:%.*]] = alloca i64, align 8581// CHECK9-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [5 x ptr], align 8582// CHECK9-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [5 x ptr], align 8583// CHECK9-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [5 x ptr], align 8584// CHECK9-NEXT:    [[DOTOFFLOAD_SIZES:%.*]] = alloca [5 x i64], align 8585// CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 4586// CHECK9-NEXT:    [[_TMP1:%.*]] = alloca i32, align 4587// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4588// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_2:%.*]] = alloca i32, align 4589// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_3:%.*]] = alloca i64, align 8590// CHECK9-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8591// CHECK9-NEXT:    store i32 0, ptr [[RETVAL]], align 4592// CHECK9-NEXT:    store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 4593// CHECK9-NEXT:    store ptr [[ARGV]], ptr [[ARGV_ADDR]], align 8594// CHECK9-NEXT:    store i32 100, ptr [[N]], align 4595// CHECK9-NEXT:    store i32 2, ptr [[M]], align 4596// CHECK9-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N]], align 4597// CHECK9-NEXT:    [[TMP1:%.*]] = zext i32 [[TMP0]] to i64598// CHECK9-NEXT:    [[TMP2:%.*]] = load i32, ptr [[M]], align 4599// CHECK9-NEXT:    [[TMP3:%.*]] = zext i32 [[TMP2]] to i64600// CHECK9-NEXT:    [[TMP4:%.*]] = call ptr @llvm.stacksave.p0()601// CHECK9-NEXT:    store ptr [[TMP4]], ptr [[SAVED_STACK]], align 8602// CHECK9-NEXT:    [[TMP5:%.*]] = mul nuw i64 [[TMP1]], [[TMP3]]603// CHECK9-NEXT:    [[VLA:%.*]] = alloca i32, i64 [[TMP5]], align 4604// CHECK9-NEXT:    store i64 [[TMP1]], ptr [[__VLA_EXPR0]], align 8605// CHECK9-NEXT:    store i64 [[TMP3]], ptr [[__VLA_EXPR1]], align 8606// CHECK9-NEXT:    [[TMP6:%.*]] = load i32, ptr [[N]], align 4607// CHECK9-NEXT:    store i32 [[TMP6]], ptr [[N_CASTED]], align 4608// CHECK9-NEXT:    [[TMP7:%.*]] = load i64, ptr [[N_CASTED]], align 8609// CHECK9-NEXT:    [[TMP8:%.*]] = load i32, ptr [[M]], align 4610// CHECK9-NEXT:    store i32 [[TMP8]], ptr [[M_CASTED]], align 4611// CHECK9-NEXT:    [[TMP9:%.*]] = load i64, ptr [[M_CASTED]], align 8612// CHECK9-NEXT:    [[TMP10:%.*]] = mul nuw i64 [[TMP1]], [[TMP3]]613// CHECK9-NEXT:    [[TMP11:%.*]] = mul nuw i64 [[TMP10]], 4614// CHECK9-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES]], ptr align 8 @.offload_sizes, i64 40, i1 false)615// CHECK9-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0616// CHECK9-NEXT:    store i64 [[TMP7]], ptr [[TMP12]], align 8617// CHECK9-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0618// CHECK9-NEXT:    store i64 [[TMP7]], ptr [[TMP13]], align 8619// CHECK9-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0620// CHECK9-NEXT:    store ptr null, ptr [[TMP14]], align 8621// CHECK9-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1622// CHECK9-NEXT:    store i64 [[TMP9]], ptr [[TMP15]], align 8623// CHECK9-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1624// CHECK9-NEXT:    store i64 [[TMP9]], ptr [[TMP16]], align 8625// CHECK9-NEXT:    [[TMP17:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1626// CHECK9-NEXT:    store ptr null, ptr [[TMP17]], align 8627// CHECK9-NEXT:    [[TMP18:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2628// CHECK9-NEXT:    store i64 [[TMP1]], ptr [[TMP18]], align 8629// CHECK9-NEXT:    [[TMP19:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2630// CHECK9-NEXT:    store i64 [[TMP1]], ptr [[TMP19]], align 8631// CHECK9-NEXT:    [[TMP20:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2632// CHECK9-NEXT:    store ptr null, ptr [[TMP20]], align 8633// CHECK9-NEXT:    [[TMP21:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3634// CHECK9-NEXT:    store i64 [[TMP3]], ptr [[TMP21]], align 8635// CHECK9-NEXT:    [[TMP22:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3636// CHECK9-NEXT:    store i64 [[TMP3]], ptr [[TMP22]], align 8637// CHECK9-NEXT:    [[TMP23:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3638// CHECK9-NEXT:    store ptr null, ptr [[TMP23]], align 8639// CHECK9-NEXT:    [[TMP24:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 4640// CHECK9-NEXT:    store ptr [[VLA]], ptr [[TMP24]], align 8641// CHECK9-NEXT:    [[TMP25:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 4642// CHECK9-NEXT:    store ptr [[VLA]], ptr [[TMP25]], align 8643// CHECK9-NEXT:    [[TMP26:%.*]] = getelementptr inbounds [5 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 4644// CHECK9-NEXT:    store i64 [[TMP11]], ptr [[TMP26]], align 8645// CHECK9-NEXT:    [[TMP27:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 4646// CHECK9-NEXT:    store ptr null, ptr [[TMP27]], align 8647// CHECK9-NEXT:    [[TMP28:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0648// CHECK9-NEXT:    [[TMP29:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0649// CHECK9-NEXT:    [[TMP30:%.*]] = getelementptr inbounds [5 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0650// CHECK9-NEXT:    [[TMP31:%.*]] = load i32, ptr [[N]], align 4651// CHECK9-NEXT:    store i32 [[TMP31]], ptr [[DOTCAPTURE_EXPR_]], align 4652// CHECK9-NEXT:    [[TMP32:%.*]] = load i32, ptr [[M]], align 4653// CHECK9-NEXT:    store i32 [[TMP32]], ptr [[DOTCAPTURE_EXPR_2]], align 4654// CHECK9-NEXT:    [[TMP33:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4655// CHECK9-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP33]], 0656// CHECK9-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1657// CHECK9-NEXT:    [[CONV:%.*]] = sext i32 [[DIV]] to i64658// CHECK9-NEXT:    [[TMP34:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 4659// CHECK9-NEXT:    [[SUB4:%.*]] = sub nsw i32 [[TMP34]], 0660// CHECK9-NEXT:    [[DIV5:%.*]] = sdiv i32 [[SUB4]], 1661// CHECK9-NEXT:    [[CONV6:%.*]] = sext i32 [[DIV5]] to i64662// CHECK9-NEXT:    [[MUL:%.*]] = mul nsw i64 [[CONV]], [[CONV6]]663// CHECK9-NEXT:    [[SUB7:%.*]] = sub nsw i64 [[MUL]], 1664// CHECK9-NEXT:    store i64 [[SUB7]], ptr [[DOTCAPTURE_EXPR_3]], align 8665// CHECK9-NEXT:    [[TMP35:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR_3]], align 8666// CHECK9-NEXT:    [[ADD:%.*]] = add nsw i64 [[TMP35]], 1667// CHECK9-NEXT:    [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0668// CHECK9-NEXT:    store i32 3, ptr [[TMP36]], align 4669// CHECK9-NEXT:    [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1670// CHECK9-NEXT:    store i32 5, ptr [[TMP37]], align 4671// CHECK9-NEXT:    [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2672// CHECK9-NEXT:    store ptr [[TMP28]], ptr [[TMP38]], align 8673// CHECK9-NEXT:    [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3674// CHECK9-NEXT:    store ptr [[TMP29]], ptr [[TMP39]], align 8675// CHECK9-NEXT:    [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4676// CHECK9-NEXT:    store ptr [[TMP30]], ptr [[TMP40]], align 8677// CHECK9-NEXT:    [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5678// CHECK9-NEXT:    store ptr @.offload_maptypes, ptr [[TMP41]], align 8679// CHECK9-NEXT:    [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6680// CHECK9-NEXT:    store ptr null, ptr [[TMP42]], align 8681// CHECK9-NEXT:    [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7682// CHECK9-NEXT:    store ptr null, ptr [[TMP43]], align 8683// CHECK9-NEXT:    [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8684// CHECK9-NEXT:    store i64 [[ADD]], ptr [[TMP44]], align 8685// CHECK9-NEXT:    [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9686// CHECK9-NEXT:    store i64 0, ptr [[TMP45]], align 8687// CHECK9-NEXT:    [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10688// CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP46]], align 4689// CHECK9-NEXT:    [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11690// CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP47]], align 4691// CHECK9-NEXT:    [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12692// CHECK9-NEXT:    store i32 0, ptr [[TMP48]], align 4693// CHECK9-NEXT:    [[TMP49:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l81.region_id, ptr [[KERNEL_ARGS]])694// CHECK9-NEXT:    [[TMP50:%.*]] = icmp ne i32 [[TMP49]], 0695// CHECK9-NEXT:    br i1 [[TMP50]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]696// CHECK9:       omp_offload.failed:697// CHECK9-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l81(i64 [[TMP7]], i64 [[TMP9]], i64 [[TMP1]], i64 [[TMP3]], ptr [[VLA]]) #[[ATTR3:[0-9]+]]698// CHECK9-NEXT:    br label [[OMP_OFFLOAD_CONT]]699// CHECK9:       omp_offload.cont:700// CHECK9-NEXT:    [[TMP51:%.*]] = load i32, ptr [[ARGC_ADDR]], align 4701// CHECK9-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_Z5tmainIiLi10ELi2EEiT_(i32 noundef signext [[TMP51]])702// CHECK9-NEXT:    store i32 [[CALL]], ptr [[RETVAL]], align 4703// CHECK9-NEXT:    [[TMP52:%.*]] = load ptr, ptr [[SAVED_STACK]], align 8704// CHECK9-NEXT:    call void @llvm.stackrestore.p0(ptr [[TMP52]])705// CHECK9-NEXT:    [[TMP53:%.*]] = load i32, ptr [[RETVAL]], align 4706// CHECK9-NEXT:    ret i32 [[TMP53]]707//708//709// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l81710// CHECK9-SAME: (i64 noundef [[N:%.*]], i64 noundef [[M:%.*]], i64 noundef [[VLA:%.*]], i64 noundef [[VLA1:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2:[0-9]+]] {711// CHECK9-NEXT:  entry:712// CHECK9-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 8713// CHECK9-NEXT:    [[M_ADDR:%.*]] = alloca i64, align 8714// CHECK9-NEXT:    [[VLA_ADDR:%.*]] = alloca i64, align 8715// CHECK9-NEXT:    [[VLA_ADDR2:%.*]] = alloca i64, align 8716// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8717// CHECK9-NEXT:    [[N_CASTED:%.*]] = alloca i64, align 8718// CHECK9-NEXT:    [[M_CASTED:%.*]] = alloca i64, align 8719// CHECK9-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 8720// CHECK9-NEXT:    store i64 [[M]], ptr [[M_ADDR]], align 8721// CHECK9-NEXT:    store i64 [[VLA]], ptr [[VLA_ADDR]], align 8722// CHECK9-NEXT:    store i64 [[VLA1]], ptr [[VLA_ADDR2]], align 8723// CHECK9-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8724// CHECK9-NEXT:    [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR]], align 8725// CHECK9-NEXT:    [[TMP1:%.*]] = load i64, ptr [[VLA_ADDR2]], align 8726// CHECK9-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 8727// CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[N_ADDR]], align 4728// CHECK9-NEXT:    store i32 [[TMP3]], ptr [[N_CASTED]], align 4729// CHECK9-NEXT:    [[TMP4:%.*]] = load i64, ptr [[N_CASTED]], align 8730// CHECK9-NEXT:    [[TMP5:%.*]] = load i32, ptr [[M_ADDR]], align 4731// CHECK9-NEXT:    store i32 [[TMP5]], ptr [[M_CASTED]], align 4732// CHECK9-NEXT:    [[TMP6:%.*]] = load i64, ptr [[M_CASTED]], align 8733// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 5, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l81.omp_outlined, i64 [[TMP4]], i64 [[TMP6]], i64 [[TMP0]], i64 [[TMP1]], ptr [[TMP2]])734// CHECK9-NEXT:    ret void735//736//737// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l81.omp_outlined738// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[N:%.*]], i64 noundef [[M:%.*]], i64 noundef [[VLA:%.*]], i64 noundef [[VLA1:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {739// CHECK9-NEXT:  entry:740// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8741// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8742// CHECK9-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 8743// CHECK9-NEXT:    [[M_ADDR:%.*]] = alloca i64, align 8744// CHECK9-NEXT:    [[VLA_ADDR:%.*]] = alloca i64, align 8745// CHECK9-NEXT:    [[VLA_ADDR2:%.*]] = alloca i64, align 8746// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8747// CHECK9-NEXT:    [[DOTOMP_IV:%.*]] = alloca i64, align 8748// CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 4749// CHECK9-NEXT:    [[_TMP3:%.*]] = alloca i32, align 4750// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4751// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_4:%.*]] = alloca i32, align 4752// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_5:%.*]] = alloca i64, align 8753// CHECK9-NEXT:    [[I:%.*]] = alloca i32, align 4754// CHECK9-NEXT:    [[J:%.*]] = alloca i32, align 4755// CHECK9-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i64, align 8756// CHECK9-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i64, align 8757// CHECK9-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i64, align 8758// CHECK9-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4759// CHECK9-NEXT:    [[I11:%.*]] = alloca i32, align 4760// CHECK9-NEXT:    [[J12:%.*]] = alloca i32, align 4761// CHECK9-NEXT:    [[N_CASTED:%.*]] = alloca i64, align 8762// CHECK9-NEXT:    [[M_CASTED:%.*]] = alloca i64, align 8763// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8764// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8765// CHECK9-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 8766// CHECK9-NEXT:    store i64 [[M]], ptr [[M_ADDR]], align 8767// CHECK9-NEXT:    store i64 [[VLA]], ptr [[VLA_ADDR]], align 8768// CHECK9-NEXT:    store i64 [[VLA1]], ptr [[VLA_ADDR2]], align 8769// CHECK9-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8770// CHECK9-NEXT:    [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR]], align 8771// CHECK9-NEXT:    [[TMP1:%.*]] = load i64, ptr [[VLA_ADDR2]], align 8772// CHECK9-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 8773// CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[N_ADDR]], align 4774// CHECK9-NEXT:    store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR_]], align 4775// CHECK9-NEXT:    [[TMP4:%.*]] = load i32, ptr [[M_ADDR]], align 4776// CHECK9-NEXT:    store i32 [[TMP4]], ptr [[DOTCAPTURE_EXPR_4]], align 4777// CHECK9-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4778// CHECK9-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP5]], 0779// CHECK9-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1780// CHECK9-NEXT:    [[CONV:%.*]] = sext i32 [[DIV]] to i64781// CHECK9-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 4782// CHECK9-NEXT:    [[SUB6:%.*]] = sub nsw i32 [[TMP6]], 0783// CHECK9-NEXT:    [[DIV7:%.*]] = sdiv i32 [[SUB6]], 1784// CHECK9-NEXT:    [[CONV8:%.*]] = sext i32 [[DIV7]] to i64785// CHECK9-NEXT:    [[MUL:%.*]] = mul nsw i64 [[CONV]], [[CONV8]]786// CHECK9-NEXT:    [[SUB9:%.*]] = sub nsw i64 [[MUL]], 1787// CHECK9-NEXT:    store i64 [[SUB9]], ptr [[DOTCAPTURE_EXPR_5]], align 8788// CHECK9-NEXT:    store i32 0, ptr [[I]], align 4789// CHECK9-NEXT:    store i32 0, ptr [[J]], align 4790// CHECK9-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4791// CHECK9-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP7]]792// CHECK9-NEXT:    br i1 [[CMP]], label [[LAND_LHS_TRUE:%.*]], label [[OMP_PRECOND_END:%.*]]793// CHECK9:       land.lhs.true:794// CHECK9-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 4795// CHECK9-NEXT:    [[CMP10:%.*]] = icmp slt i32 0, [[TMP8]]796// CHECK9-NEXT:    br i1 [[CMP10]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END]]797// CHECK9:       omp.precond.then:798// CHECK9-NEXT:    store i64 0, ptr [[DOTOMP_COMB_LB]], align 8799// CHECK9-NEXT:    [[TMP9:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR_5]], align 8800// CHECK9-NEXT:    store i64 [[TMP9]], ptr [[DOTOMP_COMB_UB]], align 8801// CHECK9-NEXT:    store i64 1, ptr [[DOTOMP_STRIDE]], align 8802// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4803// CHECK9-NEXT:    [[TMP10:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8804// CHECK9-NEXT:    [[TMP11:%.*]] = load i32, ptr [[TMP10]], align 4805// CHECK9-NEXT:    call void @__kmpc_for_static_init_8(ptr @[[GLOB1:[0-9]+]], i32 [[TMP11]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i64 1, i64 1)806// CHECK9-NEXT:    [[TMP12:%.*]] = load i64, ptr [[DOTOMP_COMB_UB]], align 8807// CHECK9-NEXT:    [[TMP13:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR_5]], align 8808// CHECK9-NEXT:    [[CMP13:%.*]] = icmp sgt i64 [[TMP12]], [[TMP13]]809// CHECK9-NEXT:    br i1 [[CMP13]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]810// CHECK9:       cond.true:811// CHECK9-NEXT:    [[TMP14:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR_5]], align 8812// CHECK9-NEXT:    br label [[COND_END:%.*]]813// CHECK9:       cond.false:814// CHECK9-NEXT:    [[TMP15:%.*]] = load i64, ptr [[DOTOMP_COMB_UB]], align 8815// CHECK9-NEXT:    br label [[COND_END]]816// CHECK9:       cond.end:817// CHECK9-NEXT:    [[COND:%.*]] = phi i64 [ [[TMP14]], [[COND_TRUE]] ], [ [[TMP15]], [[COND_FALSE]] ]818// CHECK9-NEXT:    store i64 [[COND]], ptr [[DOTOMP_COMB_UB]], align 8819// CHECK9-NEXT:    [[TMP16:%.*]] = load i64, ptr [[DOTOMP_COMB_LB]], align 8820// CHECK9-NEXT:    store i64 [[TMP16]], ptr [[DOTOMP_IV]], align 8821// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]822// CHECK9:       omp.inner.for.cond:823// CHECK9-NEXT:    [[TMP17:%.*]] = load i64, ptr [[DOTOMP_IV]], align 8824// CHECK9-NEXT:    [[TMP18:%.*]] = load i64, ptr [[DOTOMP_COMB_UB]], align 8825// CHECK9-NEXT:    [[CMP14:%.*]] = icmp sle i64 [[TMP17]], [[TMP18]]826// CHECK9-NEXT:    br i1 [[CMP14]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]827// CHECK9:       omp.inner.for.body:828// CHECK9-NEXT:    [[TMP19:%.*]] = load i64, ptr [[DOTOMP_COMB_LB]], align 8829// CHECK9-NEXT:    [[TMP20:%.*]] = load i64, ptr [[DOTOMP_COMB_UB]], align 8830// CHECK9-NEXT:    [[TMP21:%.*]] = load i32, ptr [[N_ADDR]], align 4831// CHECK9-NEXT:    store i32 [[TMP21]], ptr [[N_CASTED]], align 4832// CHECK9-NEXT:    [[TMP22:%.*]] = load i64, ptr [[N_CASTED]], align 8833// CHECK9-NEXT:    [[TMP23:%.*]] = load i32, ptr [[M_ADDR]], align 4834// CHECK9-NEXT:    store i32 [[TMP23]], ptr [[M_CASTED]], align 4835// CHECK9-NEXT:    [[TMP24:%.*]] = load i64, ptr [[M_CASTED]], align 8836// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 7, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l81.omp_outlined.omp_outlined, i64 [[TMP19]], i64 [[TMP20]], i64 [[TMP22]], i64 [[TMP24]], i64 [[TMP0]], i64 [[TMP1]], ptr [[TMP2]])837// CHECK9-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]838// CHECK9:       omp.inner.for.inc:839// CHECK9-NEXT:    [[TMP25:%.*]] = load i64, ptr [[DOTOMP_IV]], align 8840// CHECK9-NEXT:    [[TMP26:%.*]] = load i64, ptr [[DOTOMP_STRIDE]], align 8841// CHECK9-NEXT:    [[ADD:%.*]] = add nsw i64 [[TMP25]], [[TMP26]]842// CHECK9-NEXT:    store i64 [[ADD]], ptr [[DOTOMP_IV]], align 8843// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND]]844// CHECK9:       omp.inner.for.end:845// CHECK9-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]846// CHECK9:       omp.loop.exit:847// CHECK9-NEXT:    [[TMP27:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8848// CHECK9-NEXT:    [[TMP28:%.*]] = load i32, ptr [[TMP27]], align 4849// CHECK9-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP28]])850// CHECK9-NEXT:    br label [[OMP_PRECOND_END]]851// CHECK9:       omp.precond.end:852// CHECK9-NEXT:    ret void853//854//855// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l81.omp_outlined.omp_outlined856// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]], i64 noundef [[N:%.*]], i64 noundef [[M:%.*]], i64 noundef [[VLA:%.*]], i64 noundef [[VLA1:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {857// CHECK9-NEXT:  entry:858// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8859// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8860// CHECK9-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8861// CHECK9-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8862// CHECK9-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 8863// CHECK9-NEXT:    [[M_ADDR:%.*]] = alloca i64, align 8864// CHECK9-NEXT:    [[VLA_ADDR:%.*]] = alloca i64, align 8865// CHECK9-NEXT:    [[VLA_ADDR2:%.*]] = alloca i64, align 8866// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8867// CHECK9-NEXT:    [[DOTOMP_IV:%.*]] = alloca i64, align 8868// CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 4869// CHECK9-NEXT:    [[_TMP3:%.*]] = alloca i32, align 4870// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4871// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_4:%.*]] = alloca i32, align 4872// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_5:%.*]] = alloca i64, align 8873// CHECK9-NEXT:    [[I:%.*]] = alloca i32, align 4874// CHECK9-NEXT:    [[J:%.*]] = alloca i32, align 4875// CHECK9-NEXT:    [[DOTOMP_LB:%.*]] = alloca i64, align 8876// CHECK9-NEXT:    [[DOTOMP_UB:%.*]] = alloca i64, align 8877// CHECK9-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i64, align 8878// CHECK9-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4879// CHECK9-NEXT:    [[I11:%.*]] = alloca i32, align 4880// CHECK9-NEXT:    [[J12:%.*]] = alloca i32, align 4881// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8882// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8883// CHECK9-NEXT:    store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8884// CHECK9-NEXT:    store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8885// CHECK9-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 8886// CHECK9-NEXT:    store i64 [[M]], ptr [[M_ADDR]], align 8887// CHECK9-NEXT:    store i64 [[VLA]], ptr [[VLA_ADDR]], align 8888// CHECK9-NEXT:    store i64 [[VLA1]], ptr [[VLA_ADDR2]], align 8889// CHECK9-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8890// CHECK9-NEXT:    [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR]], align 8891// CHECK9-NEXT:    [[TMP1:%.*]] = load i64, ptr [[VLA_ADDR2]], align 8892// CHECK9-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 8893// CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[N_ADDR]], align 4894// CHECK9-NEXT:    store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR_]], align 4895// CHECK9-NEXT:    [[TMP4:%.*]] = load i32, ptr [[M_ADDR]], align 4896// CHECK9-NEXT:    store i32 [[TMP4]], ptr [[DOTCAPTURE_EXPR_4]], align 4897// CHECK9-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4898// CHECK9-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP5]], 0899// CHECK9-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1900// CHECK9-NEXT:    [[CONV:%.*]] = sext i32 [[DIV]] to i64901// CHECK9-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 4902// CHECK9-NEXT:    [[SUB6:%.*]] = sub nsw i32 [[TMP6]], 0903// CHECK9-NEXT:    [[DIV7:%.*]] = sdiv i32 [[SUB6]], 1904// CHECK9-NEXT:    [[CONV8:%.*]] = sext i32 [[DIV7]] to i64905// CHECK9-NEXT:    [[MUL:%.*]] = mul nsw i64 [[CONV]], [[CONV8]]906// CHECK9-NEXT:    [[SUB9:%.*]] = sub nsw i64 [[MUL]], 1907// CHECK9-NEXT:    store i64 [[SUB9]], ptr [[DOTCAPTURE_EXPR_5]], align 8908// CHECK9-NEXT:    store i32 0, ptr [[I]], align 4909// CHECK9-NEXT:    store i32 0, ptr [[J]], align 4910// CHECK9-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4911// CHECK9-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP7]]912// CHECK9-NEXT:    br i1 [[CMP]], label [[LAND_LHS_TRUE:%.*]], label [[OMP_PRECOND_END:%.*]]913// CHECK9:       land.lhs.true:914// CHECK9-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 4915// CHECK9-NEXT:    [[CMP10:%.*]] = icmp slt i32 0, [[TMP8]]916// CHECK9-NEXT:    br i1 [[CMP10]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END]]917// CHECK9:       omp.precond.then:918// CHECK9-NEXT:    store i64 0, ptr [[DOTOMP_LB]], align 8919// CHECK9-NEXT:    [[TMP9:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR_5]], align 8920// CHECK9-NEXT:    store i64 [[TMP9]], ptr [[DOTOMP_UB]], align 8921// CHECK9-NEXT:    [[TMP10:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8922// CHECK9-NEXT:    [[TMP11:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8923// CHECK9-NEXT:    store i64 [[TMP10]], ptr [[DOTOMP_LB]], align 8924// CHECK9-NEXT:    store i64 [[TMP11]], ptr [[DOTOMP_UB]], align 8925// CHECK9-NEXT:    store i64 1, ptr [[DOTOMP_STRIDE]], align 8926// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4927// CHECK9-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8928// CHECK9-NEXT:    [[TMP13:%.*]] = load i32, ptr [[TMP12]], align 4929// CHECK9-NEXT:    call void @__kmpc_for_static_init_8(ptr @[[GLOB2:[0-9]+]], i32 [[TMP13]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i64 1, i64 1)930// CHECK9-NEXT:    [[TMP14:%.*]] = load i64, ptr [[DOTOMP_UB]], align 8931// CHECK9-NEXT:    [[TMP15:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR_5]], align 8932// CHECK9-NEXT:    [[CMP13:%.*]] = icmp sgt i64 [[TMP14]], [[TMP15]]933// CHECK9-NEXT:    br i1 [[CMP13]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]934// CHECK9:       cond.true:935// CHECK9-NEXT:    [[TMP16:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR_5]], align 8936// CHECK9-NEXT:    br label [[COND_END:%.*]]937// CHECK9:       cond.false:938// CHECK9-NEXT:    [[TMP17:%.*]] = load i64, ptr [[DOTOMP_UB]], align 8939// CHECK9-NEXT:    br label [[COND_END]]940// CHECK9:       cond.end:941// CHECK9-NEXT:    [[COND:%.*]] = phi i64 [ [[TMP16]], [[COND_TRUE]] ], [ [[TMP17]], [[COND_FALSE]] ]942// CHECK9-NEXT:    store i64 [[COND]], ptr [[DOTOMP_UB]], align 8943// CHECK9-NEXT:    [[TMP18:%.*]] = load i64, ptr [[DOTOMP_LB]], align 8944// CHECK9-NEXT:    store i64 [[TMP18]], ptr [[DOTOMP_IV]], align 8945// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]946// CHECK9:       omp.inner.for.cond:947// CHECK9-NEXT:    [[TMP19:%.*]] = load i64, ptr [[DOTOMP_IV]], align 8948// CHECK9-NEXT:    [[TMP20:%.*]] = load i64, ptr [[DOTOMP_UB]], align 8949// CHECK9-NEXT:    [[CMP14:%.*]] = icmp sle i64 [[TMP19]], [[TMP20]]950// CHECK9-NEXT:    br i1 [[CMP14]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]951// CHECK9:       omp.inner.for.body:952// CHECK9-NEXT:    [[TMP21:%.*]] = load i64, ptr [[DOTOMP_IV]], align 8953// CHECK9-NEXT:    [[TMP22:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 4954// CHECK9-NEXT:    [[SUB15:%.*]] = sub nsw i32 [[TMP22]], 0955// CHECK9-NEXT:    [[DIV16:%.*]] = sdiv i32 [[SUB15]], 1956// CHECK9-NEXT:    [[MUL17:%.*]] = mul nsw i32 1, [[DIV16]]957// CHECK9-NEXT:    [[CONV18:%.*]] = sext i32 [[MUL17]] to i64958// CHECK9-NEXT:    [[DIV19:%.*]] = sdiv i64 [[TMP21]], [[CONV18]]959// CHECK9-NEXT:    [[MUL20:%.*]] = mul nsw i64 [[DIV19]], 1960// CHECK9-NEXT:    [[ADD:%.*]] = add nsw i64 0, [[MUL20]]961// CHECK9-NEXT:    [[CONV21:%.*]] = trunc i64 [[ADD]] to i32962// CHECK9-NEXT:    store i32 [[CONV21]], ptr [[I11]], align 4963// CHECK9-NEXT:    [[TMP23:%.*]] = load i64, ptr [[DOTOMP_IV]], align 8964// CHECK9-NEXT:    [[TMP24:%.*]] = load i64, ptr [[DOTOMP_IV]], align 8965// CHECK9-NEXT:    [[TMP25:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 4966// CHECK9-NEXT:    [[SUB22:%.*]] = sub nsw i32 [[TMP25]], 0967// CHECK9-NEXT:    [[DIV23:%.*]] = sdiv i32 [[SUB22]], 1968// CHECK9-NEXT:    [[MUL24:%.*]] = mul nsw i32 1, [[DIV23]]969// CHECK9-NEXT:    [[CONV25:%.*]] = sext i32 [[MUL24]] to i64970// CHECK9-NEXT:    [[DIV26:%.*]] = sdiv i64 [[TMP24]], [[CONV25]]971// CHECK9-NEXT:    [[TMP26:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 4972// CHECK9-NEXT:    [[SUB27:%.*]] = sub nsw i32 [[TMP26]], 0973// CHECK9-NEXT:    [[DIV28:%.*]] = sdiv i32 [[SUB27]], 1974// CHECK9-NEXT:    [[MUL29:%.*]] = mul nsw i32 1, [[DIV28]]975// CHECK9-NEXT:    [[CONV30:%.*]] = sext i32 [[MUL29]] to i64976// CHECK9-NEXT:    [[MUL31:%.*]] = mul nsw i64 [[DIV26]], [[CONV30]]977// CHECK9-NEXT:    [[SUB32:%.*]] = sub nsw i64 [[TMP23]], [[MUL31]]978// CHECK9-NEXT:    [[MUL33:%.*]] = mul nsw i64 [[SUB32]], 1979// CHECK9-NEXT:    [[ADD34:%.*]] = add nsw i64 0, [[MUL33]]980// CHECK9-NEXT:    [[CONV35:%.*]] = trunc i64 [[ADD34]] to i32981// CHECK9-NEXT:    store i32 [[CONV35]], ptr [[J12]], align 4982// CHECK9-NEXT:    [[TMP27:%.*]] = load i32, ptr [[I11]], align 4983// CHECK9-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP27]] to i64984// CHECK9-NEXT:    [[TMP28:%.*]] = mul nsw i64 [[IDXPROM]], [[TMP1]]985// CHECK9-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i64 [[TMP28]]986// CHECK9-NEXT:    [[TMP29:%.*]] = load i32, ptr [[J12]], align 4987// CHECK9-NEXT:    [[IDXPROM36:%.*]] = sext i32 [[TMP29]] to i64988// CHECK9-NEXT:    [[ARRAYIDX37:%.*]] = getelementptr inbounds i32, ptr [[ARRAYIDX]], i64 [[IDXPROM36]]989// CHECK9-NEXT:    store i32 0, ptr [[ARRAYIDX37]], align 4990// CHECK9-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]991// CHECK9:       omp.body.continue:992// CHECK9-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]993// CHECK9:       omp.inner.for.inc:994// CHECK9-NEXT:    [[TMP30:%.*]] = load i64, ptr [[DOTOMP_IV]], align 8995// CHECK9-NEXT:    [[ADD38:%.*]] = add nsw i64 [[TMP30]], 1996// CHECK9-NEXT:    store i64 [[ADD38]], ptr [[DOTOMP_IV]], align 8997// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND]]998// CHECK9:       omp.inner.for.end:999// CHECK9-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1000// CHECK9:       omp.loop.exit:1001// CHECK9-NEXT:    [[TMP31:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81002// CHECK9-NEXT:    [[TMP32:%.*]] = load i32, ptr [[TMP31]], align 41003// CHECK9-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP32]])1004// CHECK9-NEXT:    br label [[OMP_PRECOND_END]]1005// CHECK9:       omp.precond.end:1006// CHECK9-NEXT:    ret void1007//1008//1009// CHECK9-LABEL: define {{[^@]+}}@_Z5tmainIiLi10ELi2EEiT_1010// CHECK9-SAME: (i32 noundef signext [[ARGC:%.*]]) #[[ATTR5:[0-9]+]] comdat {1011// CHECK9-NEXT:  entry:1012// CHECK9-NEXT:    [[ARGC_ADDR:%.*]] = alloca i32, align 41013// CHECK9-NEXT:    [[A:%.*]] = alloca [10 x [2 x i32]], align 41014// CHECK9-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 81015// CHECK9-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 81016// CHECK9-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 81017// CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 41018// CHECK9-NEXT:    [[_TMP1:%.*]] = alloca i32, align 41019// CHECK9-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81020// CHECK9-NEXT:    store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 41021// CHECK9-NEXT:    [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01022// CHECK9-NEXT:    store ptr [[A]], ptr [[TMP0]], align 81023// CHECK9-NEXT:    [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01024// CHECK9-NEXT:    store ptr [[A]], ptr [[TMP1]], align 81025// CHECK9-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 01026// CHECK9-NEXT:    store ptr null, ptr [[TMP2]], align 81027// CHECK9-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01028// CHECK9-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01029// CHECK9-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01030// CHECK9-NEXT:    store i32 3, ptr [[TMP5]], align 41031// CHECK9-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11032// CHECK9-NEXT:    store i32 1, ptr [[TMP6]], align 41033// CHECK9-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21034// CHECK9-NEXT:    store ptr [[TMP3]], ptr [[TMP7]], align 81035// CHECK9-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31036// CHECK9-NEXT:    store ptr [[TMP4]], ptr [[TMP8]], align 81037// CHECK9-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41038// CHECK9-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP9]], align 81039// CHECK9-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51040// CHECK9-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP10]], align 81041// CHECK9-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61042// CHECK9-NEXT:    store ptr null, ptr [[TMP11]], align 81043// CHECK9-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71044// CHECK9-NEXT:    store ptr null, ptr [[TMP12]], align 81045// CHECK9-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81046// CHECK9-NEXT:    store i64 20, ptr [[TMP13]], align 81047// CHECK9-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91048// CHECK9-NEXT:    store i64 0, ptr [[TMP14]], align 81049// CHECK9-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101050// CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP15]], align 41051// CHECK9-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111052// CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP16]], align 41053// CHECK9-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121054// CHECK9-NEXT:    store i32 0, ptr [[TMP17]], align 41055// CHECK9-NEXT:    [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10ELi2EEiT__l68.region_id, ptr [[KERNEL_ARGS]])1056// CHECK9-NEXT:    [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 01057// CHECK9-NEXT:    br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1058// CHECK9:       omp_offload.failed:1059// CHECK9-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10ELi2EEiT__l68(ptr [[A]]) #[[ATTR3]]1060// CHECK9-NEXT:    br label [[OMP_OFFLOAD_CONT]]1061// CHECK9:       omp_offload.cont:1062// CHECK9-NEXT:    ret i32 01063//1064//1065// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10ELi2EEiT__l681066// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(80) [[A:%.*]]) #[[ATTR2]] {1067// CHECK9-NEXT:  entry:1068// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81069// CHECK9-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81070// CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81071// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10ELi2EEiT__l68.omp_outlined, ptr [[TMP0]])1072// CHECK9-NEXT:    ret void1073//1074//1075// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10ELi2EEiT__l68.omp_outlined1076// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(80) [[A:%.*]]) #[[ATTR2]] {1077// CHECK9-NEXT:  entry:1078// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81079// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81080// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81081// CHECK9-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41082// CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 41083// CHECK9-NEXT:    [[_TMP1:%.*]] = alloca i32, align 41084// CHECK9-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41085// CHECK9-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41086// CHECK9-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41087// CHECK9-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41088// CHECK9-NEXT:    [[I:%.*]] = alloca i32, align 41089// CHECK9-NEXT:    [[J:%.*]] = alloca i32, align 41090// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81091// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81092// CHECK9-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81093// CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81094// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 41095// CHECK9-NEXT:    store i32 19, ptr [[DOTOMP_COMB_UB]], align 41096// CHECK9-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41097// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41098// CHECK9-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81099// CHECK9-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41100// CHECK9-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1101// CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41102// CHECK9-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 191103// CHECK9-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1104// CHECK9:       cond.true:1105// CHECK9-NEXT:    br label [[COND_END:%.*]]1106// CHECK9:       cond.false:1107// CHECK9-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41108// CHECK9-NEXT:    br label [[COND_END]]1109// CHECK9:       cond.end:1110// CHECK9-NEXT:    [[COND:%.*]] = phi i32 [ 19, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]1111// CHECK9-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41112// CHECK9-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41113// CHECK9-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 41114// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1115// CHECK9:       omp.inner.for.cond:1116// CHECK9-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41117// CHECK9-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41118// CHECK9-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]1119// CHECK9-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1120// CHECK9:       omp.inner.for.body:1121// CHECK9-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41122// CHECK9-NEXT:    [[TMP9:%.*]] = zext i32 [[TMP8]] to i641123// CHECK9-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41124// CHECK9-NEXT:    [[TMP11:%.*]] = zext i32 [[TMP10]] to i641125// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10ELi2EEiT__l68.omp_outlined.omp_outlined, i64 [[TMP9]], i64 [[TMP11]], ptr [[TMP0]])1126// CHECK9-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1127// CHECK9:       omp.inner.for.inc:1128// CHECK9-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41129// CHECK9-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41130// CHECK9-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP12]], [[TMP13]]1131// CHECK9-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 41132// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND]]1133// CHECK9:       omp.inner.for.end:1134// CHECK9-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1135// CHECK9:       omp.loop.exit:1136// CHECK9-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])1137// CHECK9-NEXT:    ret void1138//1139//1140// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10ELi2EEiT__l68.omp_outlined.omp_outlined1141// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]], ptr noundef nonnull align 4 dereferenceable(80) [[A:%.*]]) #[[ATTR2]] {1142// CHECK9-NEXT:  entry:1143// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81144// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81145// CHECK9-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 81146// CHECK9-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 81147// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81148// CHECK9-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41149// CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 41150// CHECK9-NEXT:    [[_TMP1:%.*]] = alloca i32, align 41151// CHECK9-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 41152// CHECK9-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 41153// CHECK9-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41154// CHECK9-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41155// CHECK9-NEXT:    [[I:%.*]] = alloca i32, align 41156// CHECK9-NEXT:    [[J:%.*]] = alloca i32, align 41157// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81158// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81159// CHECK9-NEXT:    store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 81160// CHECK9-NEXT:    store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 81161// CHECK9-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81162// CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81163// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 41164// CHECK9-NEXT:    store i32 19, ptr [[DOTOMP_UB]], align 41165// CHECK9-NEXT:    [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 81166// CHECK9-NEXT:    [[CONV:%.*]] = trunc i64 [[TMP1]] to i321167// CHECK9-NEXT:    [[TMP2:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 81168// CHECK9-NEXT:    [[CONV2:%.*]] = trunc i64 [[TMP2]] to i321169// CHECK9-NEXT:    store i32 [[CONV]], ptr [[DOTOMP_LB]], align 41170// CHECK9-NEXT:    store i32 [[CONV2]], ptr [[DOTOMP_UB]], align 41171// CHECK9-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41172// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41173// CHECK9-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81174// CHECK9-NEXT:    [[TMP4:%.*]] = load i32, ptr [[TMP3]], align 41175// CHECK9-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB2]], i32 [[TMP4]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1176// CHECK9-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41177// CHECK9-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP5]], 191178// CHECK9-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1179// CHECK9:       cond.true:1180// CHECK9-NEXT:    br label [[COND_END:%.*]]1181// CHECK9:       cond.false:1182// CHECK9-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41183// CHECK9-NEXT:    br label [[COND_END]]1184// CHECK9:       cond.end:1185// CHECK9-NEXT:    [[COND:%.*]] = phi i32 [ 19, [[COND_TRUE]] ], [ [[TMP6]], [[COND_FALSE]] ]1186// CHECK9-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 41187// CHECK9-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41188// CHECK9-NEXT:    store i32 [[TMP7]], ptr [[DOTOMP_IV]], align 41189// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1190// CHECK9:       omp.inner.for.cond:1191// CHECK9-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41192// CHECK9-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41193// CHECK9-NEXT:    [[CMP3:%.*]] = icmp sle i32 [[TMP8]], [[TMP9]]1194// CHECK9-NEXT:    br i1 [[CMP3]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1195// CHECK9:       omp.inner.for.body:1196// CHECK9-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41197// CHECK9-NEXT:    [[DIV:%.*]] = sdiv i32 [[TMP10]], 21198// CHECK9-NEXT:    [[MUL:%.*]] = mul nsw i32 [[DIV]], 11199// CHECK9-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1200// CHECK9-NEXT:    store i32 [[ADD]], ptr [[I]], align 41201// CHECK9-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41202// CHECK9-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41203// CHECK9-NEXT:    [[DIV4:%.*]] = sdiv i32 [[TMP12]], 21204// CHECK9-NEXT:    [[MUL5:%.*]] = mul nsw i32 [[DIV4]], 21205// CHECK9-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP11]], [[MUL5]]1206// CHECK9-NEXT:    [[MUL6:%.*]] = mul nsw i32 [[SUB]], 11207// CHECK9-NEXT:    [[ADD7:%.*]] = add nsw i32 0, [[MUL6]]1208// CHECK9-NEXT:    store i32 [[ADD7]], ptr [[J]], align 41209// CHECK9-NEXT:    [[TMP13:%.*]] = load i32, ptr [[I]], align 41210// CHECK9-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP13]] to i641211// CHECK9-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [2 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]1212// CHECK9-NEXT:    [[TMP14:%.*]] = load i32, ptr [[J]], align 41213// CHECK9-NEXT:    [[IDXPROM8:%.*]] = sext i32 [[TMP14]] to i641214// CHECK9-NEXT:    [[ARRAYIDX9:%.*]] = getelementptr inbounds [2 x i32], ptr [[ARRAYIDX]], i64 0, i64 [[IDXPROM8]]1215// CHECK9-NEXT:    store i32 0, ptr [[ARRAYIDX9]], align 41216// CHECK9-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1217// CHECK9:       omp.body.continue:1218// CHECK9-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1219// CHECK9:       omp.inner.for.inc:1220// CHECK9-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41221// CHECK9-NEXT:    [[ADD10:%.*]] = add nsw i32 [[TMP15]], 11222// CHECK9-NEXT:    store i32 [[ADD10]], ptr [[DOTOMP_IV]], align 41223// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND]]1224// CHECK9:       omp.inner.for.end:1225// CHECK9-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1226// CHECK9:       omp.loop.exit:1227// CHECK9-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP4]])1228// CHECK9-NEXT:    ret void1229//1230//1231// CHECK11-LABEL: define {{[^@]+}}@main1232// CHECK11-SAME: (i32 noundef [[ARGC:%.*]], ptr noundef [[ARGV:%.*]]) #[[ATTR0:[0-9]+]] {1233// CHECK11-NEXT:  entry:1234// CHECK11-NEXT:    [[RETVAL:%.*]] = alloca i32, align 41235// CHECK11-NEXT:    [[ARGC_ADDR:%.*]] = alloca i32, align 41236// CHECK11-NEXT:    [[ARGV_ADDR:%.*]] = alloca ptr, align 41237// CHECK11-NEXT:    [[N:%.*]] = alloca i32, align 41238// CHECK11-NEXT:    [[M:%.*]] = alloca i32, align 41239// CHECK11-NEXT:    [[SAVED_STACK:%.*]] = alloca ptr, align 41240// CHECK11-NEXT:    [[__VLA_EXPR0:%.*]] = alloca i32, align 41241// CHECK11-NEXT:    [[__VLA_EXPR1:%.*]] = alloca i32, align 41242// CHECK11-NEXT:    [[N_CASTED:%.*]] = alloca i32, align 41243// CHECK11-NEXT:    [[M_CASTED:%.*]] = alloca i32, align 41244// CHECK11-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [5 x ptr], align 41245// CHECK11-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [5 x ptr], align 41246// CHECK11-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [5 x ptr], align 41247// CHECK11-NEXT:    [[DOTOFFLOAD_SIZES:%.*]] = alloca [5 x i64], align 41248// CHECK11-NEXT:    [[TMP:%.*]] = alloca i32, align 41249// CHECK11-NEXT:    [[_TMP1:%.*]] = alloca i32, align 41250// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41251// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_2:%.*]] = alloca i32, align 41252// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_3:%.*]] = alloca i64, align 81253// CHECK11-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81254// CHECK11-NEXT:    store i32 0, ptr [[RETVAL]], align 41255// CHECK11-NEXT:    store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 41256// CHECK11-NEXT:    store ptr [[ARGV]], ptr [[ARGV_ADDR]], align 41257// CHECK11-NEXT:    store i32 100, ptr [[N]], align 41258// CHECK11-NEXT:    store i32 2, ptr [[M]], align 41259// CHECK11-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N]], align 41260// CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[M]], align 41261// CHECK11-NEXT:    [[TMP2:%.*]] = call ptr @llvm.stacksave.p0()1262// CHECK11-NEXT:    store ptr [[TMP2]], ptr [[SAVED_STACK]], align 41263// CHECK11-NEXT:    [[TMP3:%.*]] = mul nuw i32 [[TMP0]], [[TMP1]]1264// CHECK11-NEXT:    [[VLA:%.*]] = alloca i32, i32 [[TMP3]], align 41265// CHECK11-NEXT:    store i32 [[TMP0]], ptr [[__VLA_EXPR0]], align 41266// CHECK11-NEXT:    store i32 [[TMP1]], ptr [[__VLA_EXPR1]], align 41267// CHECK11-NEXT:    [[TMP4:%.*]] = load i32, ptr [[N]], align 41268// CHECK11-NEXT:    store i32 [[TMP4]], ptr [[N_CASTED]], align 41269// CHECK11-NEXT:    [[TMP5:%.*]] = load i32, ptr [[N_CASTED]], align 41270// CHECK11-NEXT:    [[TMP6:%.*]] = load i32, ptr [[M]], align 41271// CHECK11-NEXT:    store i32 [[TMP6]], ptr [[M_CASTED]], align 41272// CHECK11-NEXT:    [[TMP7:%.*]] = load i32, ptr [[M_CASTED]], align 41273// CHECK11-NEXT:    [[TMP8:%.*]] = mul nuw i32 [[TMP0]], [[TMP1]]1274// CHECK11-NEXT:    [[TMP9:%.*]] = mul nuw i32 [[TMP8]], 41275// CHECK11-NEXT:    [[TMP10:%.*]] = sext i32 [[TMP9]] to i641276// CHECK11-NEXT:    call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[DOTOFFLOAD_SIZES]], ptr align 4 @.offload_sizes, i32 40, i1 false)1277// CHECK11-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01278// CHECK11-NEXT:    store i32 [[TMP5]], ptr [[TMP11]], align 41279// CHECK11-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01280// CHECK11-NEXT:    store i32 [[TMP5]], ptr [[TMP12]], align 41281// CHECK11-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01282// CHECK11-NEXT:    store ptr null, ptr [[TMP13]], align 41283// CHECK11-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11284// CHECK11-NEXT:    store i32 [[TMP7]], ptr [[TMP14]], align 41285// CHECK11-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11286// CHECK11-NEXT:    store i32 [[TMP7]], ptr [[TMP15]], align 41287// CHECK11-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 11288// CHECK11-NEXT:    store ptr null, ptr [[TMP16]], align 41289// CHECK11-NEXT:    [[TMP17:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21290// CHECK11-NEXT:    store i32 [[TMP0]], ptr [[TMP17]], align 41291// CHECK11-NEXT:    [[TMP18:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21292// CHECK11-NEXT:    store i32 [[TMP0]], ptr [[TMP18]], align 41293// CHECK11-NEXT:    [[TMP19:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 21294// CHECK11-NEXT:    store ptr null, ptr [[TMP19]], align 41295// CHECK11-NEXT:    [[TMP20:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 31296// CHECK11-NEXT:    store i32 [[TMP1]], ptr [[TMP20]], align 41297// CHECK11-NEXT:    [[TMP21:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 31298// CHECK11-NEXT:    store i32 [[TMP1]], ptr [[TMP21]], align 41299// CHECK11-NEXT:    [[TMP22:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 31300// CHECK11-NEXT:    store ptr null, ptr [[TMP22]], align 41301// CHECK11-NEXT:    [[TMP23:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 41302// CHECK11-NEXT:    store ptr [[VLA]], ptr [[TMP23]], align 41303// CHECK11-NEXT:    [[TMP24:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 41304// CHECK11-NEXT:    store ptr [[VLA]], ptr [[TMP24]], align 41305// CHECK11-NEXT:    [[TMP25:%.*]] = getelementptr inbounds [5 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 41306// CHECK11-NEXT:    store i64 [[TMP10]], ptr [[TMP25]], align 41307// CHECK11-NEXT:    [[TMP26:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 41308// CHECK11-NEXT:    store ptr null, ptr [[TMP26]], align 41309// CHECK11-NEXT:    [[TMP27:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01310// CHECK11-NEXT:    [[TMP28:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01311// CHECK11-NEXT:    [[TMP29:%.*]] = getelementptr inbounds [5 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 01312// CHECK11-NEXT:    [[TMP30:%.*]] = load i32, ptr [[N]], align 41313// CHECK11-NEXT:    store i32 [[TMP30]], ptr [[DOTCAPTURE_EXPR_]], align 41314// CHECK11-NEXT:    [[TMP31:%.*]] = load i32, ptr [[M]], align 41315// CHECK11-NEXT:    store i32 [[TMP31]], ptr [[DOTCAPTURE_EXPR_2]], align 41316// CHECK11-NEXT:    [[TMP32:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41317// CHECK11-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP32]], 01318// CHECK11-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 11319// CHECK11-NEXT:    [[CONV:%.*]] = sext i32 [[DIV]] to i641320// CHECK11-NEXT:    [[TMP33:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 41321// CHECK11-NEXT:    [[SUB4:%.*]] = sub nsw i32 [[TMP33]], 01322// CHECK11-NEXT:    [[DIV5:%.*]] = sdiv i32 [[SUB4]], 11323// CHECK11-NEXT:    [[CONV6:%.*]] = sext i32 [[DIV5]] to i641324// CHECK11-NEXT:    [[MUL:%.*]] = mul nsw i64 [[CONV]], [[CONV6]]1325// CHECK11-NEXT:    [[SUB7:%.*]] = sub nsw i64 [[MUL]], 11326// CHECK11-NEXT:    store i64 [[SUB7]], ptr [[DOTCAPTURE_EXPR_3]], align 81327// CHECK11-NEXT:    [[TMP34:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR_3]], align 81328// CHECK11-NEXT:    [[ADD:%.*]] = add nsw i64 [[TMP34]], 11329// CHECK11-NEXT:    [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01330// CHECK11-NEXT:    store i32 3, ptr [[TMP35]], align 41331// CHECK11-NEXT:    [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11332// CHECK11-NEXT:    store i32 5, ptr [[TMP36]], align 41333// CHECK11-NEXT:    [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21334// CHECK11-NEXT:    store ptr [[TMP27]], ptr [[TMP37]], align 41335// CHECK11-NEXT:    [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31336// CHECK11-NEXT:    store ptr [[TMP28]], ptr [[TMP38]], align 41337// CHECK11-NEXT:    [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41338// CHECK11-NEXT:    store ptr [[TMP29]], ptr [[TMP39]], align 41339// CHECK11-NEXT:    [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51340// CHECK11-NEXT:    store ptr @.offload_maptypes, ptr [[TMP40]], align 41341// CHECK11-NEXT:    [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61342// CHECK11-NEXT:    store ptr null, ptr [[TMP41]], align 41343// CHECK11-NEXT:    [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71344// CHECK11-NEXT:    store ptr null, ptr [[TMP42]], align 41345// CHECK11-NEXT:    [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81346// CHECK11-NEXT:    store i64 [[ADD]], ptr [[TMP43]], align 81347// CHECK11-NEXT:    [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91348// CHECK11-NEXT:    store i64 0, ptr [[TMP44]], align 81349// CHECK11-NEXT:    [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101350// CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP45]], align 41351// CHECK11-NEXT:    [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111352// CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP46]], align 41353// CHECK11-NEXT:    [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121354// CHECK11-NEXT:    store i32 0, ptr [[TMP47]], align 41355// CHECK11-NEXT:    [[TMP48:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l81.region_id, ptr [[KERNEL_ARGS]])1356// CHECK11-NEXT:    [[TMP49:%.*]] = icmp ne i32 [[TMP48]], 01357// CHECK11-NEXT:    br i1 [[TMP49]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1358// CHECK11:       omp_offload.failed:1359// CHECK11-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l81(i32 [[TMP5]], i32 [[TMP7]], i32 [[TMP0]], i32 [[TMP1]], ptr [[VLA]]) #[[ATTR3:[0-9]+]]1360// CHECK11-NEXT:    br label [[OMP_OFFLOAD_CONT]]1361// CHECK11:       omp_offload.cont:1362// CHECK11-NEXT:    [[TMP50:%.*]] = load i32, ptr [[ARGC_ADDR]], align 41363// CHECK11-NEXT:    [[CALL:%.*]] = call noundef i32 @_Z5tmainIiLi10ELi2EEiT_(i32 noundef [[TMP50]])1364// CHECK11-NEXT:    store i32 [[CALL]], ptr [[RETVAL]], align 41365// CHECK11-NEXT:    [[TMP51:%.*]] = load ptr, ptr [[SAVED_STACK]], align 41366// CHECK11-NEXT:    call void @llvm.stackrestore.p0(ptr [[TMP51]])1367// CHECK11-NEXT:    [[TMP52:%.*]] = load i32, ptr [[RETVAL]], align 41368// CHECK11-NEXT:    ret i32 [[TMP52]]1369//1370//1371// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l811372// CHECK11-SAME: (i32 noundef [[N:%.*]], i32 noundef [[M:%.*]], i32 noundef [[VLA:%.*]], i32 noundef [[VLA1:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2:[0-9]+]] {1373// CHECK11-NEXT:  entry:1374// CHECK11-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41375// CHECK11-NEXT:    [[M_ADDR:%.*]] = alloca i32, align 41376// CHECK11-NEXT:    [[VLA_ADDR:%.*]] = alloca i32, align 41377// CHECK11-NEXT:    [[VLA_ADDR2:%.*]] = alloca i32, align 41378// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 41379// CHECK11-NEXT:    [[N_CASTED:%.*]] = alloca i32, align 41380// CHECK11-NEXT:    [[M_CASTED:%.*]] = alloca i32, align 41381// CHECK11-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41382// CHECK11-NEXT:    store i32 [[M]], ptr [[M_ADDR]], align 41383// CHECK11-NEXT:    store i32 [[VLA]], ptr [[VLA_ADDR]], align 41384// CHECK11-NEXT:    store i32 [[VLA1]], ptr [[VLA_ADDR2]], align 41385// CHECK11-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 41386// CHECK11-NEXT:    [[TMP0:%.*]] = load i32, ptr [[VLA_ADDR]], align 41387// CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[VLA_ADDR2]], align 41388// CHECK11-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 41389// CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[N_ADDR]], align 41390// CHECK11-NEXT:    store i32 [[TMP3]], ptr [[N_CASTED]], align 41391// CHECK11-NEXT:    [[TMP4:%.*]] = load i32, ptr [[N_CASTED]], align 41392// CHECK11-NEXT:    [[TMP5:%.*]] = load i32, ptr [[M_ADDR]], align 41393// CHECK11-NEXT:    store i32 [[TMP5]], ptr [[M_CASTED]], align 41394// CHECK11-NEXT:    [[TMP6:%.*]] = load i32, ptr [[M_CASTED]], align 41395// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 5, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l81.omp_outlined, i32 [[TMP4]], i32 [[TMP6]], i32 [[TMP0]], i32 [[TMP1]], ptr [[TMP2]])1396// CHECK11-NEXT:    ret void1397//1398//1399// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l81.omp_outlined1400// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[N:%.*]], i32 noundef [[M:%.*]], i32 noundef [[VLA:%.*]], i32 noundef [[VLA1:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {1401// CHECK11-NEXT:  entry:1402// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41403// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41404// CHECK11-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41405// CHECK11-NEXT:    [[M_ADDR:%.*]] = alloca i32, align 41406// CHECK11-NEXT:    [[VLA_ADDR:%.*]] = alloca i32, align 41407// CHECK11-NEXT:    [[VLA_ADDR2:%.*]] = alloca i32, align 41408// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 41409// CHECK11-NEXT:    [[DOTOMP_IV:%.*]] = alloca i64, align 81410// CHECK11-NEXT:    [[TMP:%.*]] = alloca i32, align 41411// CHECK11-NEXT:    [[_TMP3:%.*]] = alloca i32, align 41412// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41413// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_4:%.*]] = alloca i32, align 41414// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_5:%.*]] = alloca i64, align 81415// CHECK11-NEXT:    [[I:%.*]] = alloca i32, align 41416// CHECK11-NEXT:    [[J:%.*]] = alloca i32, align 41417// CHECK11-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i64, align 81418// CHECK11-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i64, align 81419// CHECK11-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i64, align 81420// CHECK11-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41421// CHECK11-NEXT:    [[I11:%.*]] = alloca i32, align 41422// CHECK11-NEXT:    [[J12:%.*]] = alloca i32, align 41423// CHECK11-NEXT:    [[N_CASTED:%.*]] = alloca i32, align 41424// CHECK11-NEXT:    [[M_CASTED:%.*]] = alloca i32, align 41425// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41426// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41427// CHECK11-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41428// CHECK11-NEXT:    store i32 [[M]], ptr [[M_ADDR]], align 41429// CHECK11-NEXT:    store i32 [[VLA]], ptr [[VLA_ADDR]], align 41430// CHECK11-NEXT:    store i32 [[VLA1]], ptr [[VLA_ADDR2]], align 41431// CHECK11-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 41432// CHECK11-NEXT:    [[TMP0:%.*]] = load i32, ptr [[VLA_ADDR]], align 41433// CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[VLA_ADDR2]], align 41434// CHECK11-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 41435// CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[N_ADDR]], align 41436// CHECK11-NEXT:    store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR_]], align 41437// CHECK11-NEXT:    [[TMP4:%.*]] = load i32, ptr [[M_ADDR]], align 41438// CHECK11-NEXT:    store i32 [[TMP4]], ptr [[DOTCAPTURE_EXPR_4]], align 41439// CHECK11-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41440// CHECK11-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP5]], 01441// CHECK11-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 11442// CHECK11-NEXT:    [[CONV:%.*]] = sext i32 [[DIV]] to i641443// CHECK11-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 41444// CHECK11-NEXT:    [[SUB6:%.*]] = sub nsw i32 [[TMP6]], 01445// CHECK11-NEXT:    [[DIV7:%.*]] = sdiv i32 [[SUB6]], 11446// CHECK11-NEXT:    [[CONV8:%.*]] = sext i32 [[DIV7]] to i641447// CHECK11-NEXT:    [[MUL:%.*]] = mul nsw i64 [[CONV]], [[CONV8]]1448// CHECK11-NEXT:    [[SUB9:%.*]] = sub nsw i64 [[MUL]], 11449// CHECK11-NEXT:    store i64 [[SUB9]], ptr [[DOTCAPTURE_EXPR_5]], align 81450// CHECK11-NEXT:    store i32 0, ptr [[I]], align 41451// CHECK11-NEXT:    store i32 0, ptr [[J]], align 41452// CHECK11-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41453// CHECK11-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP7]]1454// CHECK11-NEXT:    br i1 [[CMP]], label [[LAND_LHS_TRUE:%.*]], label [[OMP_PRECOND_END:%.*]]1455// CHECK11:       land.lhs.true:1456// CHECK11-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 41457// CHECK11-NEXT:    [[CMP10:%.*]] = icmp slt i32 0, [[TMP8]]1458// CHECK11-NEXT:    br i1 [[CMP10]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END]]1459// CHECK11:       omp.precond.then:1460// CHECK11-NEXT:    store i64 0, ptr [[DOTOMP_COMB_LB]], align 81461// CHECK11-NEXT:    [[TMP9:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR_5]], align 81462// CHECK11-NEXT:    store i64 [[TMP9]], ptr [[DOTOMP_COMB_UB]], align 81463// CHECK11-NEXT:    store i64 1, ptr [[DOTOMP_STRIDE]], align 81464// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41465// CHECK11-NEXT:    [[TMP10:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41466// CHECK11-NEXT:    [[TMP11:%.*]] = load i32, ptr [[TMP10]], align 41467// CHECK11-NEXT:    call void @__kmpc_for_static_init_8(ptr @[[GLOB1:[0-9]+]], i32 [[TMP11]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i64 1, i64 1)1468// CHECK11-NEXT:    [[TMP12:%.*]] = load i64, ptr [[DOTOMP_COMB_UB]], align 81469// CHECK11-NEXT:    [[TMP13:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR_5]], align 81470// CHECK11-NEXT:    [[CMP13:%.*]] = icmp sgt i64 [[TMP12]], [[TMP13]]1471// CHECK11-NEXT:    br i1 [[CMP13]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1472// CHECK11:       cond.true:1473// CHECK11-NEXT:    [[TMP14:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR_5]], align 81474// CHECK11-NEXT:    br label [[COND_END:%.*]]1475// CHECK11:       cond.false:1476// CHECK11-NEXT:    [[TMP15:%.*]] = load i64, ptr [[DOTOMP_COMB_UB]], align 81477// CHECK11-NEXT:    br label [[COND_END]]1478// CHECK11:       cond.end:1479// CHECK11-NEXT:    [[COND:%.*]] = phi i64 [ [[TMP14]], [[COND_TRUE]] ], [ [[TMP15]], [[COND_FALSE]] ]1480// CHECK11-NEXT:    store i64 [[COND]], ptr [[DOTOMP_COMB_UB]], align 81481// CHECK11-NEXT:    [[TMP16:%.*]] = load i64, ptr [[DOTOMP_COMB_LB]], align 81482// CHECK11-NEXT:    store i64 [[TMP16]], ptr [[DOTOMP_IV]], align 81483// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1484// CHECK11:       omp.inner.for.cond:1485// CHECK11-NEXT:    [[TMP17:%.*]] = load i64, ptr [[DOTOMP_IV]], align 81486// CHECK11-NEXT:    [[TMP18:%.*]] = load i64, ptr [[DOTOMP_COMB_UB]], align 81487// CHECK11-NEXT:    [[CMP14:%.*]] = icmp sle i64 [[TMP17]], [[TMP18]]1488// CHECK11-NEXT:    br i1 [[CMP14]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1489// CHECK11:       omp.inner.for.body:1490// CHECK11-NEXT:    [[TMP19:%.*]] = load i64, ptr [[DOTOMP_COMB_LB]], align 81491// CHECK11-NEXT:    [[TMP20:%.*]] = trunc i64 [[TMP19]] to i321492// CHECK11-NEXT:    [[TMP21:%.*]] = load i64, ptr [[DOTOMP_COMB_UB]], align 81493// CHECK11-NEXT:    [[TMP22:%.*]] = trunc i64 [[TMP21]] to i321494// CHECK11-NEXT:    [[TMP23:%.*]] = load i32, ptr [[N_ADDR]], align 41495// CHECK11-NEXT:    store i32 [[TMP23]], ptr [[N_CASTED]], align 41496// CHECK11-NEXT:    [[TMP24:%.*]] = load i32, ptr [[N_CASTED]], align 41497// CHECK11-NEXT:    [[TMP25:%.*]] = load i32, ptr [[M_ADDR]], align 41498// CHECK11-NEXT:    store i32 [[TMP25]], ptr [[M_CASTED]], align 41499// CHECK11-NEXT:    [[TMP26:%.*]] = load i32, ptr [[M_CASTED]], align 41500// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 7, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l81.omp_outlined.omp_outlined, i32 [[TMP20]], i32 [[TMP22]], i32 [[TMP24]], i32 [[TMP26]], i32 [[TMP0]], i32 [[TMP1]], ptr [[TMP2]])1501// CHECK11-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1502// CHECK11:       omp.inner.for.inc:1503// CHECK11-NEXT:    [[TMP27:%.*]] = load i64, ptr [[DOTOMP_IV]], align 81504// CHECK11-NEXT:    [[TMP28:%.*]] = load i64, ptr [[DOTOMP_STRIDE]], align 81505// CHECK11-NEXT:    [[ADD:%.*]] = add nsw i64 [[TMP27]], [[TMP28]]1506// CHECK11-NEXT:    store i64 [[ADD]], ptr [[DOTOMP_IV]], align 81507// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND]]1508// CHECK11:       omp.inner.for.end:1509// CHECK11-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1510// CHECK11:       omp.loop.exit:1511// CHECK11-NEXT:    [[TMP29:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41512// CHECK11-NEXT:    [[TMP30:%.*]] = load i32, ptr [[TMP29]], align 41513// CHECK11-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP30]])1514// CHECK11-NEXT:    br label [[OMP_PRECOND_END]]1515// CHECK11:       omp.precond.end:1516// CHECK11-NEXT:    ret void1517//1518//1519// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l81.omp_outlined.omp_outlined1520// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[DOTPREVIOUS_LB_:%.*]], i32 noundef [[DOTPREVIOUS_UB_:%.*]], i32 noundef [[N:%.*]], i32 noundef [[M:%.*]], i32 noundef [[VLA:%.*]], i32 noundef [[VLA1:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {1521// CHECK11-NEXT:  entry:1522// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41523// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41524// CHECK11-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i32, align 41525// CHECK11-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i32, align 41526// CHECK11-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41527// CHECK11-NEXT:    [[M_ADDR:%.*]] = alloca i32, align 41528// CHECK11-NEXT:    [[VLA_ADDR:%.*]] = alloca i32, align 41529// CHECK11-NEXT:    [[VLA_ADDR2:%.*]] = alloca i32, align 41530// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 41531// CHECK11-NEXT:    [[DOTOMP_IV:%.*]] = alloca i64, align 81532// CHECK11-NEXT:    [[TMP:%.*]] = alloca i32, align 41533// CHECK11-NEXT:    [[_TMP3:%.*]] = alloca i32, align 41534// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41535// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_4:%.*]] = alloca i32, align 41536// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_5:%.*]] = alloca i64, align 81537// CHECK11-NEXT:    [[I:%.*]] = alloca i32, align 41538// CHECK11-NEXT:    [[J:%.*]] = alloca i32, align 41539// CHECK11-NEXT:    [[DOTOMP_LB:%.*]] = alloca i64, align 81540// CHECK11-NEXT:    [[DOTOMP_UB:%.*]] = alloca i64, align 81541// CHECK11-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i64, align 81542// CHECK11-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41543// CHECK11-NEXT:    [[I13:%.*]] = alloca i32, align 41544// CHECK11-NEXT:    [[J14:%.*]] = alloca i32, align 41545// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41546// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41547// CHECK11-NEXT:    store i32 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 41548// CHECK11-NEXT:    store i32 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 41549// CHECK11-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41550// CHECK11-NEXT:    store i32 [[M]], ptr [[M_ADDR]], align 41551// CHECK11-NEXT:    store i32 [[VLA]], ptr [[VLA_ADDR]], align 41552// CHECK11-NEXT:    store i32 [[VLA1]], ptr [[VLA_ADDR2]], align 41553// CHECK11-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 41554// CHECK11-NEXT:    [[TMP0:%.*]] = load i32, ptr [[VLA_ADDR]], align 41555// CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[VLA_ADDR2]], align 41556// CHECK11-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 41557// CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[N_ADDR]], align 41558// CHECK11-NEXT:    store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR_]], align 41559// CHECK11-NEXT:    [[TMP4:%.*]] = load i32, ptr [[M_ADDR]], align 41560// CHECK11-NEXT:    store i32 [[TMP4]], ptr [[DOTCAPTURE_EXPR_4]], align 41561// CHECK11-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41562// CHECK11-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP5]], 01563// CHECK11-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 11564// CHECK11-NEXT:    [[CONV:%.*]] = sext i32 [[DIV]] to i641565// CHECK11-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 41566// CHECK11-NEXT:    [[SUB6:%.*]] = sub nsw i32 [[TMP6]], 01567// CHECK11-NEXT:    [[DIV7:%.*]] = sdiv i32 [[SUB6]], 11568// CHECK11-NEXT:    [[CONV8:%.*]] = sext i32 [[DIV7]] to i641569// CHECK11-NEXT:    [[MUL:%.*]] = mul nsw i64 [[CONV]], [[CONV8]]1570// CHECK11-NEXT:    [[SUB9:%.*]] = sub nsw i64 [[MUL]], 11571// CHECK11-NEXT:    store i64 [[SUB9]], ptr [[DOTCAPTURE_EXPR_5]], align 81572// CHECK11-NEXT:    store i32 0, ptr [[I]], align 41573// CHECK11-NEXT:    store i32 0, ptr [[J]], align 41574// CHECK11-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41575// CHECK11-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP7]]1576// CHECK11-NEXT:    br i1 [[CMP]], label [[LAND_LHS_TRUE:%.*]], label [[OMP_PRECOND_END:%.*]]1577// CHECK11:       land.lhs.true:1578// CHECK11-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 41579// CHECK11-NEXT:    [[CMP10:%.*]] = icmp slt i32 0, [[TMP8]]1580// CHECK11-NEXT:    br i1 [[CMP10]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END]]1581// CHECK11:       omp.precond.then:1582// CHECK11-NEXT:    store i64 0, ptr [[DOTOMP_LB]], align 81583// CHECK11-NEXT:    [[TMP9:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR_5]], align 81584// CHECK11-NEXT:    store i64 [[TMP9]], ptr [[DOTOMP_UB]], align 81585// CHECK11-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTPREVIOUS_LB__ADDR]], align 41586// CHECK11-NEXT:    [[CONV11:%.*]] = zext i32 [[TMP10]] to i641587// CHECK11-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTPREVIOUS_UB__ADDR]], align 41588// CHECK11-NEXT:    [[CONV12:%.*]] = zext i32 [[TMP11]] to i641589// CHECK11-NEXT:    store i64 [[CONV11]], ptr [[DOTOMP_LB]], align 81590// CHECK11-NEXT:    store i64 [[CONV12]], ptr [[DOTOMP_UB]], align 81591// CHECK11-NEXT:    store i64 1, ptr [[DOTOMP_STRIDE]], align 81592// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41593// CHECK11-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41594// CHECK11-NEXT:    [[TMP13:%.*]] = load i32, ptr [[TMP12]], align 41595// CHECK11-NEXT:    call void @__kmpc_for_static_init_8(ptr @[[GLOB2:[0-9]+]], i32 [[TMP13]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i64 1, i64 1)1596// CHECK11-NEXT:    [[TMP14:%.*]] = load i64, ptr [[DOTOMP_UB]], align 81597// CHECK11-NEXT:    [[TMP15:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR_5]], align 81598// CHECK11-NEXT:    [[CMP15:%.*]] = icmp sgt i64 [[TMP14]], [[TMP15]]1599// CHECK11-NEXT:    br i1 [[CMP15]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1600// CHECK11:       cond.true:1601// CHECK11-NEXT:    [[TMP16:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR_5]], align 81602// CHECK11-NEXT:    br label [[COND_END:%.*]]1603// CHECK11:       cond.false:1604// CHECK11-NEXT:    [[TMP17:%.*]] = load i64, ptr [[DOTOMP_UB]], align 81605// CHECK11-NEXT:    br label [[COND_END]]1606// CHECK11:       cond.end:1607// CHECK11-NEXT:    [[COND:%.*]] = phi i64 [ [[TMP16]], [[COND_TRUE]] ], [ [[TMP17]], [[COND_FALSE]] ]1608// CHECK11-NEXT:    store i64 [[COND]], ptr [[DOTOMP_UB]], align 81609// CHECK11-NEXT:    [[TMP18:%.*]] = load i64, ptr [[DOTOMP_LB]], align 81610// CHECK11-NEXT:    store i64 [[TMP18]], ptr [[DOTOMP_IV]], align 81611// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1612// CHECK11:       omp.inner.for.cond:1613// CHECK11-NEXT:    [[TMP19:%.*]] = load i64, ptr [[DOTOMP_IV]], align 81614// CHECK11-NEXT:    [[TMP20:%.*]] = load i64, ptr [[DOTOMP_UB]], align 81615// CHECK11-NEXT:    [[CMP16:%.*]] = icmp sle i64 [[TMP19]], [[TMP20]]1616// CHECK11-NEXT:    br i1 [[CMP16]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1617// CHECK11:       omp.inner.for.body:1618// CHECK11-NEXT:    [[TMP21:%.*]] = load i64, ptr [[DOTOMP_IV]], align 81619// CHECK11-NEXT:    [[TMP22:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 41620// CHECK11-NEXT:    [[SUB17:%.*]] = sub nsw i32 [[TMP22]], 01621// CHECK11-NEXT:    [[DIV18:%.*]] = sdiv i32 [[SUB17]], 11622// CHECK11-NEXT:    [[MUL19:%.*]] = mul nsw i32 1, [[DIV18]]1623// CHECK11-NEXT:    [[CONV20:%.*]] = sext i32 [[MUL19]] to i641624// CHECK11-NEXT:    [[DIV21:%.*]] = sdiv i64 [[TMP21]], [[CONV20]]1625// CHECK11-NEXT:    [[MUL22:%.*]] = mul nsw i64 [[DIV21]], 11626// CHECK11-NEXT:    [[ADD:%.*]] = add nsw i64 0, [[MUL22]]1627// CHECK11-NEXT:    [[CONV23:%.*]] = trunc i64 [[ADD]] to i321628// CHECK11-NEXT:    store i32 [[CONV23]], ptr [[I13]], align 41629// CHECK11-NEXT:    [[TMP23:%.*]] = load i64, ptr [[DOTOMP_IV]], align 81630// CHECK11-NEXT:    [[TMP24:%.*]] = load i64, ptr [[DOTOMP_IV]], align 81631// CHECK11-NEXT:    [[TMP25:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 41632// CHECK11-NEXT:    [[SUB24:%.*]] = sub nsw i32 [[TMP25]], 01633// CHECK11-NEXT:    [[DIV25:%.*]] = sdiv i32 [[SUB24]], 11634// CHECK11-NEXT:    [[MUL26:%.*]] = mul nsw i32 1, [[DIV25]]1635// CHECK11-NEXT:    [[CONV27:%.*]] = sext i32 [[MUL26]] to i641636// CHECK11-NEXT:    [[DIV28:%.*]] = sdiv i64 [[TMP24]], [[CONV27]]1637// CHECK11-NEXT:    [[TMP26:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 41638// CHECK11-NEXT:    [[SUB29:%.*]] = sub nsw i32 [[TMP26]], 01639// CHECK11-NEXT:    [[DIV30:%.*]] = sdiv i32 [[SUB29]], 11640// CHECK11-NEXT:    [[MUL31:%.*]] = mul nsw i32 1, [[DIV30]]1641// CHECK11-NEXT:    [[CONV32:%.*]] = sext i32 [[MUL31]] to i641642// CHECK11-NEXT:    [[MUL33:%.*]] = mul nsw i64 [[DIV28]], [[CONV32]]1643// CHECK11-NEXT:    [[SUB34:%.*]] = sub nsw i64 [[TMP23]], [[MUL33]]1644// CHECK11-NEXT:    [[MUL35:%.*]] = mul nsw i64 [[SUB34]], 11645// CHECK11-NEXT:    [[ADD36:%.*]] = add nsw i64 0, [[MUL35]]1646// CHECK11-NEXT:    [[CONV37:%.*]] = trunc i64 [[ADD36]] to i321647// CHECK11-NEXT:    store i32 [[CONV37]], ptr [[J14]], align 41648// CHECK11-NEXT:    [[TMP27:%.*]] = load i32, ptr [[I13]], align 41649// CHECK11-NEXT:    [[TMP28:%.*]] = mul nsw i32 [[TMP27]], [[TMP1]]1650// CHECK11-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i32 [[TMP28]]1651// CHECK11-NEXT:    [[TMP29:%.*]] = load i32, ptr [[J14]], align 41652// CHECK11-NEXT:    [[ARRAYIDX38:%.*]] = getelementptr inbounds i32, ptr [[ARRAYIDX]], i32 [[TMP29]]1653// CHECK11-NEXT:    store i32 0, ptr [[ARRAYIDX38]], align 41654// CHECK11-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1655// CHECK11:       omp.body.continue:1656// CHECK11-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1657// CHECK11:       omp.inner.for.inc:1658// CHECK11-NEXT:    [[TMP30:%.*]] = load i64, ptr [[DOTOMP_IV]], align 81659// CHECK11-NEXT:    [[ADD39:%.*]] = add nsw i64 [[TMP30]], 11660// CHECK11-NEXT:    store i64 [[ADD39]], ptr [[DOTOMP_IV]], align 81661// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND]]1662// CHECK11:       omp.inner.for.end:1663// CHECK11-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1664// CHECK11:       omp.loop.exit:1665// CHECK11-NEXT:    [[TMP31:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41666// CHECK11-NEXT:    [[TMP32:%.*]] = load i32, ptr [[TMP31]], align 41667// CHECK11-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP32]])1668// CHECK11-NEXT:    br label [[OMP_PRECOND_END]]1669// CHECK11:       omp.precond.end:1670// CHECK11-NEXT:    ret void1671//1672//1673// CHECK11-LABEL: define {{[^@]+}}@_Z5tmainIiLi10ELi2EEiT_1674// CHECK11-SAME: (i32 noundef [[ARGC:%.*]]) #[[ATTR5:[0-9]+]] comdat {1675// CHECK11-NEXT:  entry:1676// CHECK11-NEXT:    [[ARGC_ADDR:%.*]] = alloca i32, align 41677// CHECK11-NEXT:    [[A:%.*]] = alloca [10 x [2 x i32]], align 41678// CHECK11-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 41679// CHECK11-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 41680// CHECK11-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 41681// CHECK11-NEXT:    [[TMP:%.*]] = alloca i32, align 41682// CHECK11-NEXT:    [[_TMP1:%.*]] = alloca i32, align 41683// CHECK11-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81684// CHECK11-NEXT:    store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 41685// CHECK11-NEXT:    [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01686// CHECK11-NEXT:    store ptr [[A]], ptr [[TMP0]], align 41687// CHECK11-NEXT:    [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01688// CHECK11-NEXT:    store ptr [[A]], ptr [[TMP1]], align 41689// CHECK11-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01690// CHECK11-NEXT:    store ptr null, ptr [[TMP2]], align 41691// CHECK11-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01692// CHECK11-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01693// CHECK11-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01694// CHECK11-NEXT:    store i32 3, ptr [[TMP5]], align 41695// CHECK11-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11696// CHECK11-NEXT:    store i32 1, ptr [[TMP6]], align 41697// CHECK11-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21698// CHECK11-NEXT:    store ptr [[TMP3]], ptr [[TMP7]], align 41699// CHECK11-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31700// CHECK11-NEXT:    store ptr [[TMP4]], ptr [[TMP8]], align 41701// CHECK11-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41702// CHECK11-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP9]], align 41703// CHECK11-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51704// CHECK11-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP10]], align 41705// CHECK11-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61706// CHECK11-NEXT:    store ptr null, ptr [[TMP11]], align 41707// CHECK11-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71708// CHECK11-NEXT:    store ptr null, ptr [[TMP12]], align 41709// CHECK11-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81710// CHECK11-NEXT:    store i64 20, ptr [[TMP13]], align 81711// CHECK11-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91712// CHECK11-NEXT:    store i64 0, ptr [[TMP14]], align 81713// CHECK11-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101714// CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP15]], align 41715// CHECK11-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111716// CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP16]], align 41717// CHECK11-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121718// CHECK11-NEXT:    store i32 0, ptr [[TMP17]], align 41719// CHECK11-NEXT:    [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB3]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10ELi2EEiT__l68.region_id, ptr [[KERNEL_ARGS]])1720// CHECK11-NEXT:    [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 01721// CHECK11-NEXT:    br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1722// CHECK11:       omp_offload.failed:1723// CHECK11-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10ELi2EEiT__l68(ptr [[A]]) #[[ATTR3]]1724// CHECK11-NEXT:    br label [[OMP_OFFLOAD_CONT]]1725// CHECK11:       omp_offload.cont:1726// CHECK11-NEXT:    ret i32 01727//1728//1729// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10ELi2EEiT__l681730// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(80) [[A:%.*]]) #[[ATTR2]] {1731// CHECK11-NEXT:  entry:1732// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 41733// CHECK11-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 41734// CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41735// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10ELi2EEiT__l68.omp_outlined, ptr [[TMP0]])1736// CHECK11-NEXT:    ret void1737//1738//1739// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10ELi2EEiT__l68.omp_outlined1740// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(80) [[A:%.*]]) #[[ATTR2]] {1741// CHECK11-NEXT:  entry:1742// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41743// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41744// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 41745// CHECK11-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41746// CHECK11-NEXT:    [[TMP:%.*]] = alloca i32, align 41747// CHECK11-NEXT:    [[_TMP1:%.*]] = alloca i32, align 41748// CHECK11-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41749// CHECK11-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41750// CHECK11-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41751// CHECK11-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41752// CHECK11-NEXT:    [[I:%.*]] = alloca i32, align 41753// CHECK11-NEXT:    [[J:%.*]] = alloca i32, align 41754// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41755// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41756// CHECK11-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 41757// CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41758// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 41759// CHECK11-NEXT:    store i32 19, ptr [[DOTOMP_COMB_UB]], align 41760// CHECK11-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41761// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41762// CHECK11-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41763// CHECK11-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41764// CHECK11-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1765// CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41766// CHECK11-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 191767// CHECK11-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1768// CHECK11:       cond.true:1769// CHECK11-NEXT:    br label [[COND_END:%.*]]1770// CHECK11:       cond.false:1771// CHECK11-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41772// CHECK11-NEXT:    br label [[COND_END]]1773// CHECK11:       cond.end:1774// CHECK11-NEXT:    [[COND:%.*]] = phi i32 [ 19, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]1775// CHECK11-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41776// CHECK11-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41777// CHECK11-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 41778// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1779// CHECK11:       omp.inner.for.cond:1780// CHECK11-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41781// CHECK11-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41782// CHECK11-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]1783// CHECK11-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1784// CHECK11:       omp.inner.for.body:1785// CHECK11-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41786// CHECK11-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41787// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10ELi2EEiT__l68.omp_outlined.omp_outlined, i32 [[TMP8]], i32 [[TMP9]], ptr [[TMP0]])1788// CHECK11-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1789// CHECK11:       omp.inner.for.inc:1790// CHECK11-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41791// CHECK11-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41792// CHECK11-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP10]], [[TMP11]]1793// CHECK11-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 41794// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND]]1795// CHECK11:       omp.inner.for.end:1796// CHECK11-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1797// CHECK11:       omp.loop.exit:1798// CHECK11-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])1799// CHECK11-NEXT:    ret void1800//1801//1802// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10ELi2EEiT__l68.omp_outlined.omp_outlined1803// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[DOTPREVIOUS_LB_:%.*]], i32 noundef [[DOTPREVIOUS_UB_:%.*]], ptr noundef nonnull align 4 dereferenceable(80) [[A:%.*]]) #[[ATTR2]] {1804// CHECK11-NEXT:  entry:1805// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41806// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41807// CHECK11-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i32, align 41808// CHECK11-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i32, align 41809// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 41810// CHECK11-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41811// CHECK11-NEXT:    [[TMP:%.*]] = alloca i32, align 41812// CHECK11-NEXT:    [[_TMP1:%.*]] = alloca i32, align 41813// CHECK11-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 41814// CHECK11-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 41815// CHECK11-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41816// CHECK11-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41817// CHECK11-NEXT:    [[I:%.*]] = alloca i32, align 41818// CHECK11-NEXT:    [[J:%.*]] = alloca i32, align 41819// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41820// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41821// CHECK11-NEXT:    store i32 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 41822// CHECK11-NEXT:    store i32 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 41823// CHECK11-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 41824// CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41825// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 41826// CHECK11-NEXT:    store i32 19, ptr [[DOTOMP_UB]], align 41827// CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTPREVIOUS_LB__ADDR]], align 41828// CHECK11-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTPREVIOUS_UB__ADDR]], align 41829// CHECK11-NEXT:    store i32 [[TMP1]], ptr [[DOTOMP_LB]], align 41830// CHECK11-NEXT:    store i32 [[TMP2]], ptr [[DOTOMP_UB]], align 41831// CHECK11-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41832// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41833// CHECK11-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41834// CHECK11-NEXT:    [[TMP4:%.*]] = load i32, ptr [[TMP3]], align 41835// CHECK11-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB2]], i32 [[TMP4]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1836// CHECK11-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41837// CHECK11-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP5]], 191838// CHECK11-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1839// CHECK11:       cond.true:1840// CHECK11-NEXT:    br label [[COND_END:%.*]]1841// CHECK11:       cond.false:1842// CHECK11-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41843// CHECK11-NEXT:    br label [[COND_END]]1844// CHECK11:       cond.end:1845// CHECK11-NEXT:    [[COND:%.*]] = phi i32 [ 19, [[COND_TRUE]] ], [ [[TMP6]], [[COND_FALSE]] ]1846// CHECK11-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 41847// CHECK11-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41848// CHECK11-NEXT:    store i32 [[TMP7]], ptr [[DOTOMP_IV]], align 41849// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1850// CHECK11:       omp.inner.for.cond:1851// CHECK11-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41852// CHECK11-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41853// CHECK11-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP8]], [[TMP9]]1854// CHECK11-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1855// CHECK11:       omp.inner.for.body:1856// CHECK11-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41857// CHECK11-NEXT:    [[DIV:%.*]] = sdiv i32 [[TMP10]], 21858// CHECK11-NEXT:    [[MUL:%.*]] = mul nsw i32 [[DIV]], 11859// CHECK11-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1860// CHECK11-NEXT:    store i32 [[ADD]], ptr [[I]], align 41861// CHECK11-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41862// CHECK11-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41863// CHECK11-NEXT:    [[DIV3:%.*]] = sdiv i32 [[TMP12]], 21864// CHECK11-NEXT:    [[MUL4:%.*]] = mul nsw i32 [[DIV3]], 21865// CHECK11-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP11]], [[MUL4]]1866// CHECK11-NEXT:    [[MUL5:%.*]] = mul nsw i32 [[SUB]], 11867// CHECK11-NEXT:    [[ADD6:%.*]] = add nsw i32 0, [[MUL5]]1868// CHECK11-NEXT:    store i32 [[ADD6]], ptr [[J]], align 41869// CHECK11-NEXT:    [[TMP13:%.*]] = load i32, ptr [[I]], align 41870// CHECK11-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [2 x i32]], ptr [[TMP0]], i32 0, i32 [[TMP13]]1871// CHECK11-NEXT:    [[TMP14:%.*]] = load i32, ptr [[J]], align 41872// CHECK11-NEXT:    [[ARRAYIDX7:%.*]] = getelementptr inbounds [2 x i32], ptr [[ARRAYIDX]], i32 0, i32 [[TMP14]]1873// CHECK11-NEXT:    store i32 0, ptr [[ARRAYIDX7]], align 41874// CHECK11-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1875// CHECK11:       omp.body.continue:1876// CHECK11-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1877// CHECK11:       omp.inner.for.inc:1878// CHECK11-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41879// CHECK11-NEXT:    [[ADD8:%.*]] = add nsw i32 [[TMP15]], 11880// CHECK11-NEXT:    store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 41881// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND]]1882// CHECK11:       omp.inner.for.end:1883// CHECK11-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1884// CHECK11:       omp.loop.exit:1885// CHECK11-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP4]])1886// CHECK11-NEXT:    ret void1887//1888