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];25  float b;26  int foo(void) {27 28    #pragma omp target29    #pragma omp teams distribute30    for(int i = 0; i < X; i++) {31      a[i] = (T)0;32    }33    #pragma omp target34    #pragma omp teams distribute dist_schedule(static)35    for(int i = 0; i < X; i++) {36      a[i] = (T)0;37    }38    #pragma omp target39    #pragma omp teams distribute dist_schedule(static, X/2)40    for(int i = 0; i < X; i++) {41      a[i] = (T)0;42    }43 44 45 46 47 48 49    return a[0];50  }51};52 53int teams_template_struct(void) {54  SS<int, 123, 456> V;55  return V.foo();56 57}58#endif // CK159 60// Test host codegen.61// 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=CHECK962// RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s63// 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=CHECK964// 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=CHECK1165// RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s66// 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=CHECK1167 68// 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}}"69// 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 %s70// 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}}"71// 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}}"72// 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 %s73// 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}}"74#ifdef CK275 76template <typename T, int n>77int tmain(T argc) {78  T a[n];79#pragma omp target80#pragma omp teams distribute81  for(int i = 0; i < n; i++) {82    a[i] = (T)0;83  }84#pragma omp target85#pragma omp teams distribute dist_schedule(static)86  for(int i = 0; i < n; i++) {87    a[i] = (T)0;88  }89#pragma omp target90#pragma omp teams distribute dist_schedule(static, n)91  for(int i = 0; i < n; i++) {92    a[i] = (T)0;93  }94  return 0;95}96 97int main (int argc, char **argv) {98  int n = 100;99  int a[n];100#pragma omp target101#pragma omp teams distribute102  for(int i = 0; i < n; i++) {103    a[i] = 0;104  }105#pragma omp target106#pragma omp teams distribute dist_schedule(static)107  for(int i = 0; i < n; i++) {108    a[i] = 0;109  }110#pragma omp target111#pragma omp teams distribute dist_schedule(static, n)112  for(int i = 0; i < n; i++) {113    a[i] = 0;114  }115  return tmain<int, 10>(argc);116}117 118 119 120 121 122 123 124 125 126 127 128 129 130 131 132#endif // CK2133#endif // #ifndef HEADER134// CHECK1-LABEL: define {{[^@]+}}@_Z21teams_template_structv135// CHECK1-SAME: () #[[ATTR0:[0-9]+]] {136// CHECK1-NEXT:  entry:137// CHECK1-NEXT:    [[V:%.*]] = alloca [[STRUCT_SS:%.*]], align 4138// CHECK1-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_ZN2SSIiLi123ELx456EE3fooEv(ptr noundef nonnull align 4 dereferenceable(496) [[V]])139// CHECK1-NEXT:    ret i32 [[CALL]]140//141//142// CHECK1-LABEL: define {{[^@]+}}@_ZN2SSIiLi123ELx456EE3fooEv143// CHECK1-SAME: (ptr noundef nonnull align 4 dereferenceable(496) [[THIS:%.*]]) #[[ATTR0]] comdat {144// CHECK1-NEXT:  entry:145// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8146// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8147// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8148// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8149// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4150// CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8151// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS3:%.*]] = alloca [1 x ptr], align 8152// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS4:%.*]] = alloca [1 x ptr], align 8153// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS5:%.*]] = alloca [1 x ptr], align 8154// CHECK1-NEXT:    [[_TMP6:%.*]] = alloca i32, align 4155// CHECK1-NEXT:    [[KERNEL_ARGS7:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8156// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS11:%.*]] = alloca [1 x ptr], align 8157// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS12:%.*]] = alloca [1 x ptr], align 8158// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS13:%.*]] = alloca [1 x ptr], align 8159// CHECK1-NEXT:    [[_TMP14:%.*]] = alloca i32, align 4160// CHECK1-NEXT:    [[KERNEL_ARGS15:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8161// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8162// CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8163// CHECK1-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[THIS1]], i32 0, i32 0164// CHECK1-NEXT:    [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0165// CHECK1-NEXT:    store ptr [[THIS1]], ptr [[TMP0]], align 8166// CHECK1-NEXT:    [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0167// CHECK1-NEXT:    store ptr [[A]], ptr [[TMP1]], align 8168// CHECK1-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0169// CHECK1-NEXT:    store ptr null, ptr [[TMP2]], align 8170// CHECK1-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0171// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0172// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0173// CHECK1-NEXT:    store i32 3, ptr [[TMP5]], align 4174// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1175// CHECK1-NEXT:    store i32 1, ptr [[TMP6]], align 4176// CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2177// CHECK1-NEXT:    store ptr [[TMP3]], ptr [[TMP7]], align 8178// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3179// CHECK1-NEXT:    store ptr [[TMP4]], ptr [[TMP8]], align 8180// CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4181// CHECK1-NEXT:    store ptr @.offload_sizes, ptr [[TMP9]], align 8182// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5183// CHECK1-NEXT:    store ptr @.offload_maptypes, ptr [[TMP10]], align 8184// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6185// CHECK1-NEXT:    store ptr null, ptr [[TMP11]], align 8186// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7187// CHECK1-NEXT:    store ptr null, ptr [[TMP12]], align 8188// CHECK1-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8189// CHECK1-NEXT:    store i64 123, ptr [[TMP13]], align 8190// CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9191// CHECK1-NEXT:    store i64 0, ptr [[TMP14]], align 8192// CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10193// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4194// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11195// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4196// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12197// CHECK1-NEXT:    store i32 0, ptr [[TMP17]], align 4198// CHECK1-NEXT:    [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28.region_id, ptr [[KERNEL_ARGS]])199// CHECK1-NEXT:    [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0200// CHECK1-NEXT:    br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]201// CHECK1:       omp_offload.failed:202// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28(ptr [[THIS1]]) #[[ATTR2:[0-9]+]]203// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]204// CHECK1:       omp_offload.cont:205// CHECK1-NEXT:    [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 0206// CHECK1-NEXT:    [[TMP20:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0207// CHECK1-NEXT:    store ptr [[THIS1]], ptr [[TMP20]], align 8208// CHECK1-NEXT:    [[TMP21:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 0209// CHECK1-NEXT:    store ptr [[A2]], ptr [[TMP21]], align 8210// CHECK1-NEXT:    [[TMP22:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS5]], i64 0, i64 0211// CHECK1-NEXT:    store ptr null, ptr [[TMP22]], align 8212// CHECK1-NEXT:    [[TMP23:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0213// CHECK1-NEXT:    [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 0214// CHECK1-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 0215// CHECK1-NEXT:    store i32 3, ptr [[TMP25]], align 4216// CHECK1-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 1217// CHECK1-NEXT:    store i32 1, ptr [[TMP26]], align 4218// CHECK1-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 2219// CHECK1-NEXT:    store ptr [[TMP23]], ptr [[TMP27]], align 8220// CHECK1-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 3221// CHECK1-NEXT:    store ptr [[TMP24]], ptr [[TMP28]], align 8222// CHECK1-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 4223// CHECK1-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP29]], align 8224// CHECK1-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 5225// CHECK1-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP30]], align 8226// CHECK1-NEXT:    [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 6227// CHECK1-NEXT:    store ptr null, ptr [[TMP31]], align 8228// CHECK1-NEXT:    [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 7229// CHECK1-NEXT:    store ptr null, ptr [[TMP32]], align 8230// CHECK1-NEXT:    [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 8231// CHECK1-NEXT:    store i64 123, ptr [[TMP33]], align 8232// CHECK1-NEXT:    [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 9233// CHECK1-NEXT:    store i64 0, ptr [[TMP34]], align 8234// CHECK1-NEXT:    [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 10235// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP35]], align 4236// CHECK1-NEXT:    [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 11237// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP36]], align 4238// CHECK1-NEXT:    [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 12239// CHECK1-NEXT:    store i32 0, ptr [[TMP37]], align 4240// CHECK1-NEXT:    [[TMP38:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l33.region_id, ptr [[KERNEL_ARGS7]])241// CHECK1-NEXT:    [[TMP39:%.*]] = icmp ne i32 [[TMP38]], 0242// CHECK1-NEXT:    br i1 [[TMP39]], label [[OMP_OFFLOAD_FAILED8:%.*]], label [[OMP_OFFLOAD_CONT9:%.*]]243// CHECK1:       omp_offload.failed8:244// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l33(ptr [[THIS1]]) #[[ATTR2]]245// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT9]]246// CHECK1:       omp_offload.cont9:247// CHECK1-NEXT:    [[A10:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 0248// CHECK1-NEXT:    [[TMP40:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS11]], i32 0, i32 0249// CHECK1-NEXT:    store ptr [[THIS1]], ptr [[TMP40]], align 8250// CHECK1-NEXT:    [[TMP41:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS12]], i32 0, i32 0251// CHECK1-NEXT:    store ptr [[A10]], ptr [[TMP41]], align 8252// CHECK1-NEXT:    [[TMP42:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS13]], i64 0, i64 0253// CHECK1-NEXT:    store ptr null, ptr [[TMP42]], align 8254// CHECK1-NEXT:    [[TMP43:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS11]], i32 0, i32 0255// CHECK1-NEXT:    [[TMP44:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS12]], i32 0, i32 0256// CHECK1-NEXT:    [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 0257// CHECK1-NEXT:    store i32 3, ptr [[TMP45]], align 4258// CHECK1-NEXT:    [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 1259// CHECK1-NEXT:    store i32 1, ptr [[TMP46]], align 4260// CHECK1-NEXT:    [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 2261// CHECK1-NEXT:    store ptr [[TMP43]], ptr [[TMP47]], align 8262// CHECK1-NEXT:    [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 3263// CHECK1-NEXT:    store ptr [[TMP44]], ptr [[TMP48]], align 8264// CHECK1-NEXT:    [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 4265// CHECK1-NEXT:    store ptr @.offload_sizes.3, ptr [[TMP49]], align 8266// CHECK1-NEXT:    [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 5267// CHECK1-NEXT:    store ptr @.offload_maptypes.4, ptr [[TMP50]], align 8268// CHECK1-NEXT:    [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 6269// CHECK1-NEXT:    store ptr null, ptr [[TMP51]], align 8270// CHECK1-NEXT:    [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 7271// CHECK1-NEXT:    store ptr null, ptr [[TMP52]], align 8272// CHECK1-NEXT:    [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 8273// CHECK1-NEXT:    store i64 123, ptr [[TMP53]], align 8274// CHECK1-NEXT:    [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 9275// CHECK1-NEXT:    store i64 0, ptr [[TMP54]], align 8276// CHECK1-NEXT:    [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 10277// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP55]], align 4278// CHECK1-NEXT:    [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 11279// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP56]], align 4280// CHECK1-NEXT:    [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 12281// CHECK1-NEXT:    store i32 0, ptr [[TMP57]], align 4282// CHECK1-NEXT:    [[TMP58:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l38.region_id, ptr [[KERNEL_ARGS15]])283// CHECK1-NEXT:    [[TMP59:%.*]] = icmp ne i32 [[TMP58]], 0284// CHECK1-NEXT:    br i1 [[TMP59]], label [[OMP_OFFLOAD_FAILED16:%.*]], label [[OMP_OFFLOAD_CONT17:%.*]]285// CHECK1:       omp_offload.failed16:286// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l38(ptr [[THIS1]]) #[[ATTR2]]287// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT17]]288// CHECK1:       omp_offload.cont17:289// CHECK1-NEXT:    [[A18:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 0290// CHECK1-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A18]], i64 0, i64 0291// CHECK1-NEXT:    [[TMP60:%.*]] = load i32, ptr [[ARRAYIDX]], align 4292// CHECK1-NEXT:    ret i32 [[TMP60]]293//294//295// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28296// CHECK1-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1:[0-9]+]] {297// CHECK1-NEXT:  entry:298// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8299// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8300// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8301// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28.omp_outlined, ptr [[TMP0]])302// CHECK1-NEXT:    ret void303//304//305// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28.omp_outlined306// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {307// CHECK1-NEXT:  entry:308// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8309// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8310// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8311// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4312// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4313// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4314// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4315// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4316// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4317// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4318// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8319// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8320// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8321// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8322// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4323// CHECK1-NEXT:    store i32 122, ptr [[DOTOMP_UB]], align 4324// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4325// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4326// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8327// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4328// CHECK1-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)329// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4330// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 122331// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]332// CHECK1:       cond.true:333// CHECK1-NEXT:    br label [[COND_END:%.*]]334// CHECK1:       cond.false:335// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4336// CHECK1-NEXT:    br label [[COND_END]]337// CHECK1:       cond.end:338// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 122, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]339// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4340// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4341// CHECK1-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4342// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]343// CHECK1:       omp.inner.for.cond:344// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4345// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4346// CHECK1-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]347// CHECK1-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]348// CHECK1:       omp.inner.for.body:349// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4350// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1351// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]352// CHECK1-NEXT:    store i32 [[ADD]], ptr [[I]], align 4353// CHECK1-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP0]], i32 0, i32 0354// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[I]], align 4355// CHECK1-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP9]] to i64356// CHECK1-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A]], i64 0, i64 [[IDXPROM]]357// CHECK1-NEXT:    store i32 0, ptr [[ARRAYIDX]], align 4358// CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]359// CHECK1:       omp.body.continue:360// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]361// CHECK1:       omp.inner.for.inc:362// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4363// CHECK1-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP10]], 1364// CHECK1-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4365// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]366// CHECK1:       omp.inner.for.end:367// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]368// CHECK1:       omp.loop.exit:369// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])370// CHECK1-NEXT:    ret void371//372//373// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l33374// CHECK1-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1]] {375// CHECK1-NEXT:  entry:376// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8377// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8378// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8379// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l33.omp_outlined, ptr [[TMP0]])380// CHECK1-NEXT:    ret void381//382//383// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l33.omp_outlined384// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {385// CHECK1-NEXT:  entry:386// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8387// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8388// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8389// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4390// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4391// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4392// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4393// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4394// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4395// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4396// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8397// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8398// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8399// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8400// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4401// CHECK1-NEXT:    store i32 122, ptr [[DOTOMP_UB]], align 4402// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4403// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4404// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8405// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4406// CHECK1-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)407// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4408// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 122409// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]410// CHECK1:       cond.true:411// CHECK1-NEXT:    br label [[COND_END:%.*]]412// CHECK1:       cond.false:413// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4414// CHECK1-NEXT:    br label [[COND_END]]415// CHECK1:       cond.end:416// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 122, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]417// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4418// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4419// CHECK1-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4420// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]421// CHECK1:       omp.inner.for.cond:422// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4423// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4424// CHECK1-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]425// CHECK1-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]426// CHECK1:       omp.inner.for.body:427// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4428// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1429// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]430// CHECK1-NEXT:    store i32 [[ADD]], ptr [[I]], align 4431// CHECK1-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP0]], i32 0, i32 0432// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[I]], align 4433// CHECK1-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP9]] to i64434// CHECK1-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A]], i64 0, i64 [[IDXPROM]]435// CHECK1-NEXT:    store i32 0, ptr [[ARRAYIDX]], align 4436// CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]437// CHECK1:       omp.body.continue:438// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]439// CHECK1:       omp.inner.for.inc:440// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4441// CHECK1-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP10]], 1442// CHECK1-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4443// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]444// CHECK1:       omp.inner.for.end:445// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]446// CHECK1:       omp.loop.exit:447// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])448// CHECK1-NEXT:    ret void449//450//451// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l38452// CHECK1-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1]] {453// CHECK1-NEXT:  entry:454// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8455// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8456// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8457// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l38.omp_outlined, ptr [[TMP0]])458// CHECK1-NEXT:    ret void459//460//461// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l38.omp_outlined462// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {463// CHECK1-NEXT:  entry:464// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8465// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8466// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8467// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4468// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4469// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4470// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4471// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4472// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4473// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4474// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8475// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8476// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8477// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8478// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4479// CHECK1-NEXT:    store i32 122, ptr [[DOTOMP_UB]], align 4480// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4481// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4482// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8483// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4484// CHECK1-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 91, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 61)485// CHECK1-NEXT:    br label [[OMP_DISPATCH_COND:%.*]]486// CHECK1:       omp.dispatch.cond:487// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4488// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 122489// CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]490// CHECK1:       cond.true:491// CHECK1-NEXT:    br label [[COND_END:%.*]]492// CHECK1:       cond.false:493// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4494// CHECK1-NEXT:    br label [[COND_END]]495// CHECK1:       cond.end:496// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 122, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]497// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4498// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4499// CHECK1-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4500// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4501// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4502// CHECK1-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]503// CHECK1-NEXT:    br i1 [[CMP1]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]504// CHECK1:       omp.dispatch.body:505// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]506// CHECK1:       omp.inner.for.cond:507// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP8:![0-9]+]]508// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP8]]509// CHECK1-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP8]], [[TMP9]]510// CHECK1-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]511// CHECK1:       omp.inner.for.body:512// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP8]]513// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP10]], 1514// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]515// CHECK1-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP8]]516// CHECK1-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP0]], i32 0, i32 0517// CHECK1-NEXT:    [[TMP11:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP8]]518// CHECK1-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP11]] to i64519// CHECK1-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A]], i64 0, i64 [[IDXPROM]]520// CHECK1-NEXT:    store i32 0, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP8]]521// CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]522// CHECK1:       omp.body.continue:523// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]524// CHECK1:       omp.inner.for.inc:525// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP8]]526// CHECK1-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP12]], 1527// CHECK1-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP8]]528// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP9:![0-9]+]]529// CHECK1:       omp.inner.for.end:530// CHECK1-NEXT:    br label [[OMP_DISPATCH_INC:%.*]]531// CHECK1:       omp.dispatch.inc:532// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4533// CHECK1-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4534// CHECK1-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP13]], [[TMP14]]535// CHECK1-NEXT:    store i32 [[ADD4]], ptr [[DOTOMP_LB]], align 4536// CHECK1-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4537// CHECK1-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4538// CHECK1-NEXT:    [[ADD5:%.*]] = add nsw i32 [[TMP15]], [[TMP16]]539// CHECK1-NEXT:    store i32 [[ADD5]], ptr [[DOTOMP_UB]], align 4540// CHECK1-NEXT:    br label [[OMP_DISPATCH_COND]]541// CHECK1:       omp.dispatch.end:542// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])543// CHECK1-NEXT:    ret void544//545//546// CHECK3-LABEL: define {{[^@]+}}@_Z21teams_template_structv547// CHECK3-SAME: () #[[ATTR0:[0-9]+]] {548// CHECK3-NEXT:  entry:549// CHECK3-NEXT:    [[V:%.*]] = alloca [[STRUCT_SS:%.*]], align 4550// CHECK3-NEXT:    [[CALL:%.*]] = call noundef i32 @_ZN2SSIiLi123ELx456EE3fooEv(ptr noundef nonnull align 4 dereferenceable(496) [[V]])551// CHECK3-NEXT:    ret i32 [[CALL]]552//553//554// CHECK3-LABEL: define {{[^@]+}}@_ZN2SSIiLi123ELx456EE3fooEv555// CHECK3-SAME: (ptr noundef nonnull align 4 dereferenceable(496) [[THIS:%.*]]) #[[ATTR0]] comdat align 2 {556// CHECK3-NEXT:  entry:557// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4558// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4559// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4560// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4561// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 4562// CHECK3-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8563// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS3:%.*]] = alloca [1 x ptr], align 4564// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS4:%.*]] = alloca [1 x ptr], align 4565// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS5:%.*]] = alloca [1 x ptr], align 4566// CHECK3-NEXT:    [[_TMP6:%.*]] = alloca i32, align 4567// CHECK3-NEXT:    [[KERNEL_ARGS7:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8568// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS11:%.*]] = alloca [1 x ptr], align 4569// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS12:%.*]] = alloca [1 x ptr], align 4570// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS13:%.*]] = alloca [1 x ptr], align 4571// CHECK3-NEXT:    [[_TMP14:%.*]] = alloca i32, align 4572// CHECK3-NEXT:    [[KERNEL_ARGS15:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8573// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4574// CHECK3-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4575// CHECK3-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[THIS1]], i32 0, i32 0576// CHECK3-NEXT:    [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0577// CHECK3-NEXT:    store ptr [[THIS1]], ptr [[TMP0]], align 4578// CHECK3-NEXT:    [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0579// CHECK3-NEXT:    store ptr [[A]], ptr [[TMP1]], align 4580// CHECK3-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0581// CHECK3-NEXT:    store ptr null, ptr [[TMP2]], align 4582// CHECK3-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0583// CHECK3-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0584// CHECK3-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0585// CHECK3-NEXT:    store i32 3, ptr [[TMP5]], align 4586// CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1587// CHECK3-NEXT:    store i32 1, ptr [[TMP6]], align 4588// CHECK3-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2589// CHECK3-NEXT:    store ptr [[TMP3]], ptr [[TMP7]], align 4590// CHECK3-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3591// CHECK3-NEXT:    store ptr [[TMP4]], ptr [[TMP8]], align 4592// CHECK3-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4593// CHECK3-NEXT:    store ptr @.offload_sizes, ptr [[TMP9]], align 4594// CHECK3-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5595// CHECK3-NEXT:    store ptr @.offload_maptypes, ptr [[TMP10]], align 4596// CHECK3-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6597// CHECK3-NEXT:    store ptr null, ptr [[TMP11]], align 4598// CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7599// CHECK3-NEXT:    store ptr null, ptr [[TMP12]], align 4600// CHECK3-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8601// CHECK3-NEXT:    store i64 123, ptr [[TMP13]], align 8602// CHECK3-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9603// CHECK3-NEXT:    store i64 0, ptr [[TMP14]], align 8604// CHECK3-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10605// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4606// CHECK3-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11607// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4608// CHECK3-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12609// CHECK3-NEXT:    store i32 0, ptr [[TMP17]], align 4610// CHECK3-NEXT:    [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28.region_id, ptr [[KERNEL_ARGS]])611// CHECK3-NEXT:    [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0612// CHECK3-NEXT:    br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]613// CHECK3:       omp_offload.failed:614// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28(ptr [[THIS1]]) #[[ATTR2:[0-9]+]]615// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT]]616// CHECK3:       omp_offload.cont:617// CHECK3-NEXT:    [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 0618// CHECK3-NEXT:    [[TMP20:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0619// CHECK3-NEXT:    store ptr [[THIS1]], ptr [[TMP20]], align 4620// CHECK3-NEXT:    [[TMP21:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 0621// CHECK3-NEXT:    store ptr [[A2]], ptr [[TMP21]], align 4622// CHECK3-NEXT:    [[TMP22:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS5]], i32 0, i32 0623// CHECK3-NEXT:    store ptr null, ptr [[TMP22]], align 4624// CHECK3-NEXT:    [[TMP23:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0625// CHECK3-NEXT:    [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 0626// CHECK3-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 0627// CHECK3-NEXT:    store i32 3, ptr [[TMP25]], align 4628// CHECK3-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 1629// CHECK3-NEXT:    store i32 1, ptr [[TMP26]], align 4630// CHECK3-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 2631// CHECK3-NEXT:    store ptr [[TMP23]], ptr [[TMP27]], align 4632// CHECK3-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 3633// CHECK3-NEXT:    store ptr [[TMP24]], ptr [[TMP28]], align 4634// CHECK3-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 4635// CHECK3-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP29]], align 4636// CHECK3-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 5637// CHECK3-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP30]], align 4638// CHECK3-NEXT:    [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 6639// CHECK3-NEXT:    store ptr null, ptr [[TMP31]], align 4640// CHECK3-NEXT:    [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 7641// CHECK3-NEXT:    store ptr null, ptr [[TMP32]], align 4642// CHECK3-NEXT:    [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 8643// CHECK3-NEXT:    store i64 123, ptr [[TMP33]], align 8644// CHECK3-NEXT:    [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 9645// CHECK3-NEXT:    store i64 0, ptr [[TMP34]], align 8646// CHECK3-NEXT:    [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 10647// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP35]], align 4648// CHECK3-NEXT:    [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 11649// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP36]], align 4650// CHECK3-NEXT:    [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 12651// CHECK3-NEXT:    store i32 0, ptr [[TMP37]], align 4652// CHECK3-NEXT:    [[TMP38:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l33.region_id, ptr [[KERNEL_ARGS7]])653// CHECK3-NEXT:    [[TMP39:%.*]] = icmp ne i32 [[TMP38]], 0654// CHECK3-NEXT:    br i1 [[TMP39]], label [[OMP_OFFLOAD_FAILED8:%.*]], label [[OMP_OFFLOAD_CONT9:%.*]]655// CHECK3:       omp_offload.failed8:656// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l33(ptr [[THIS1]]) #[[ATTR2]]657// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT9]]658// CHECK3:       omp_offload.cont9:659// CHECK3-NEXT:    [[A10:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 0660// CHECK3-NEXT:    [[TMP40:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS11]], i32 0, i32 0661// CHECK3-NEXT:    store ptr [[THIS1]], ptr [[TMP40]], align 4662// CHECK3-NEXT:    [[TMP41:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS12]], i32 0, i32 0663// CHECK3-NEXT:    store ptr [[A10]], ptr [[TMP41]], align 4664// CHECK3-NEXT:    [[TMP42:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS13]], i32 0, i32 0665// CHECK3-NEXT:    store ptr null, ptr [[TMP42]], align 4666// CHECK3-NEXT:    [[TMP43:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS11]], i32 0, i32 0667// CHECK3-NEXT:    [[TMP44:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS12]], i32 0, i32 0668// CHECK3-NEXT:    [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 0669// CHECK3-NEXT:    store i32 3, ptr [[TMP45]], align 4670// CHECK3-NEXT:    [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 1671// CHECK3-NEXT:    store i32 1, ptr [[TMP46]], align 4672// CHECK3-NEXT:    [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 2673// CHECK3-NEXT:    store ptr [[TMP43]], ptr [[TMP47]], align 4674// CHECK3-NEXT:    [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 3675// CHECK3-NEXT:    store ptr [[TMP44]], ptr [[TMP48]], align 4676// CHECK3-NEXT:    [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 4677// CHECK3-NEXT:    store ptr @.offload_sizes.3, ptr [[TMP49]], align 4678// CHECK3-NEXT:    [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 5679// CHECK3-NEXT:    store ptr @.offload_maptypes.4, ptr [[TMP50]], align 4680// CHECK3-NEXT:    [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 6681// CHECK3-NEXT:    store ptr null, ptr [[TMP51]], align 4682// CHECK3-NEXT:    [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 7683// CHECK3-NEXT:    store ptr null, ptr [[TMP52]], align 4684// CHECK3-NEXT:    [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 8685// CHECK3-NEXT:    store i64 123, ptr [[TMP53]], align 8686// CHECK3-NEXT:    [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 9687// CHECK3-NEXT:    store i64 0, ptr [[TMP54]], align 8688// CHECK3-NEXT:    [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 10689// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP55]], align 4690// CHECK3-NEXT:    [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 11691// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP56]], align 4692// CHECK3-NEXT:    [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 12693// CHECK3-NEXT:    store i32 0, ptr [[TMP57]], align 4694// CHECK3-NEXT:    [[TMP58:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l38.region_id, ptr [[KERNEL_ARGS15]])695// CHECK3-NEXT:    [[TMP59:%.*]] = icmp ne i32 [[TMP58]], 0696// CHECK3-NEXT:    br i1 [[TMP59]], label [[OMP_OFFLOAD_FAILED16:%.*]], label [[OMP_OFFLOAD_CONT17:%.*]]697// CHECK3:       omp_offload.failed16:698// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l38(ptr [[THIS1]]) #[[ATTR2]]699// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT17]]700// CHECK3:       omp_offload.cont17:701// CHECK3-NEXT:    [[A18:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 0702// CHECK3-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A18]], i32 0, i32 0703// CHECK3-NEXT:    [[TMP60:%.*]] = load i32, ptr [[ARRAYIDX]], align 4704// CHECK3-NEXT:    ret i32 [[TMP60]]705//706//707// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28708// CHECK3-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1:[0-9]+]] {709// CHECK3-NEXT:  entry:710// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4711// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4712// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4713// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28.omp_outlined, ptr [[TMP0]])714// CHECK3-NEXT:    ret void715//716//717// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28.omp_outlined718// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {719// CHECK3-NEXT:  entry:720// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4721// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4722// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4723// CHECK3-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4724// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 4725// CHECK3-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4726// CHECK3-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4727// CHECK3-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4728// CHECK3-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4729// CHECK3-NEXT:    [[I:%.*]] = alloca i32, align 4730// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4731// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4732// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4733// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4734// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4735// CHECK3-NEXT:    store i32 122, ptr [[DOTOMP_UB]], align 4736// CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4737// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4738// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4739// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4740// CHECK3-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)741// CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4742// CHECK3-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 122743// CHECK3-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]744// CHECK3:       cond.true:745// CHECK3-NEXT:    br label [[COND_END:%.*]]746// CHECK3:       cond.false:747// CHECK3-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4748// CHECK3-NEXT:    br label [[COND_END]]749// CHECK3:       cond.end:750// CHECK3-NEXT:    [[COND:%.*]] = phi i32 [ 122, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]751// CHECK3-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4752// CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4753// CHECK3-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4754// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]755// CHECK3:       omp.inner.for.cond:756// CHECK3-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4757// CHECK3-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4758// CHECK3-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]759// CHECK3-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]760// CHECK3:       omp.inner.for.body:761// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4762// CHECK3-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1763// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]764// CHECK3-NEXT:    store i32 [[ADD]], ptr [[I]], align 4765// CHECK3-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP0]], i32 0, i32 0766// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[I]], align 4767// CHECK3-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A]], i32 0, i32 [[TMP9]]768// CHECK3-NEXT:    store i32 0, ptr [[ARRAYIDX]], align 4769// CHECK3-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]770// CHECK3:       omp.body.continue:771// CHECK3-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]772// CHECK3:       omp.inner.for.inc:773// CHECK3-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4774// CHECK3-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP10]], 1775// CHECK3-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4776// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND]]777// CHECK3:       omp.inner.for.end:778// CHECK3-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]779// CHECK3:       omp.loop.exit:780// CHECK3-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])781// CHECK3-NEXT:    ret void782//783//784// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l33785// CHECK3-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1]] {786// CHECK3-NEXT:  entry:787// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4788// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4789// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4790// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l33.omp_outlined, ptr [[TMP0]])791// CHECK3-NEXT:    ret void792//793//794// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l33.omp_outlined795// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {796// CHECK3-NEXT:  entry:797// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4798// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4799// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4800// CHECK3-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4801// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 4802// CHECK3-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4803// CHECK3-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4804// CHECK3-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4805// CHECK3-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4806// CHECK3-NEXT:    [[I:%.*]] = alloca i32, align 4807// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4808// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4809// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4810// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4811// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4812// CHECK3-NEXT:    store i32 122, ptr [[DOTOMP_UB]], align 4813// CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4814// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4815// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4816// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4817// CHECK3-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)818// CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4819// CHECK3-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 122820// CHECK3-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]821// CHECK3:       cond.true:822// CHECK3-NEXT:    br label [[COND_END:%.*]]823// CHECK3:       cond.false:824// CHECK3-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4825// CHECK3-NEXT:    br label [[COND_END]]826// CHECK3:       cond.end:827// CHECK3-NEXT:    [[COND:%.*]] = phi i32 [ 122, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]828// CHECK3-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4829// CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4830// CHECK3-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4831// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]832// CHECK3:       omp.inner.for.cond:833// CHECK3-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4834// CHECK3-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4835// CHECK3-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]836// CHECK3-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]837// CHECK3:       omp.inner.for.body:838// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4839// CHECK3-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1840// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]841// CHECK3-NEXT:    store i32 [[ADD]], ptr [[I]], align 4842// CHECK3-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP0]], i32 0, i32 0843// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[I]], align 4844// CHECK3-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A]], i32 0, i32 [[TMP9]]845// CHECK3-NEXT:    store i32 0, ptr [[ARRAYIDX]], align 4846// CHECK3-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]847// CHECK3:       omp.body.continue:848// CHECK3-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]849// CHECK3:       omp.inner.for.inc:850// CHECK3-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4851// CHECK3-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP10]], 1852// CHECK3-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4853// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND]]854// CHECK3:       omp.inner.for.end:855// CHECK3-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]856// CHECK3:       omp.loop.exit:857// CHECK3-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])858// CHECK3-NEXT:    ret void859//860//861// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l38862// CHECK3-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1]] {863// CHECK3-NEXT:  entry:864// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4865// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4866// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4867// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l38.omp_outlined, ptr [[TMP0]])868// CHECK3-NEXT:    ret void869//870//871// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l38.omp_outlined872// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {873// CHECK3-NEXT:  entry:874// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4875// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4876// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4877// CHECK3-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4878// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 4879// CHECK3-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4880// CHECK3-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4881// CHECK3-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4882// CHECK3-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4883// CHECK3-NEXT:    [[I:%.*]] = alloca i32, align 4884// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4885// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4886// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4887// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4888// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4889// CHECK3-NEXT:    store i32 122, ptr [[DOTOMP_UB]], align 4890// CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4891// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4892// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4893// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4894// CHECK3-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 91, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 61)895// CHECK3-NEXT:    br label [[OMP_DISPATCH_COND:%.*]]896// CHECK3:       omp.dispatch.cond:897// CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4898// CHECK3-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 122899// CHECK3-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]900// CHECK3:       cond.true:901// CHECK3-NEXT:    br label [[COND_END:%.*]]902// CHECK3:       cond.false:903// CHECK3-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4904// CHECK3-NEXT:    br label [[COND_END]]905// CHECK3:       cond.end:906// CHECK3-NEXT:    [[COND:%.*]] = phi i32 [ 122, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]907// CHECK3-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4908// CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4909// CHECK3-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4910// CHECK3-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4911// CHECK3-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4912// CHECK3-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]913// CHECK3-NEXT:    br i1 [[CMP1]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]914// CHECK3:       omp.dispatch.body:915// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]916// CHECK3:       omp.inner.for.cond:917// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP9:![0-9]+]]918// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP9]]919// CHECK3-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP8]], [[TMP9]]920// CHECK3-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]921// CHECK3:       omp.inner.for.body:922// CHECK3-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP9]]923// CHECK3-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP10]], 1924// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]925// CHECK3-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP9]]926// CHECK3-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP0]], i32 0, i32 0927// CHECK3-NEXT:    [[TMP11:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP9]]928// CHECK3-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A]], i32 0, i32 [[TMP11]]929// CHECK3-NEXT:    store i32 0, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP9]]930// CHECK3-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]931// CHECK3:       omp.body.continue:932// CHECK3-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]933// CHECK3:       omp.inner.for.inc:934// CHECK3-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP9]]935// CHECK3-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP12]], 1936// CHECK3-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP9]]937// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP10:![0-9]+]]938// CHECK3:       omp.inner.for.end:939// CHECK3-NEXT:    br label [[OMP_DISPATCH_INC:%.*]]940// CHECK3:       omp.dispatch.inc:941// CHECK3-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4942// CHECK3-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4943// CHECK3-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP13]], [[TMP14]]944// CHECK3-NEXT:    store i32 [[ADD4]], ptr [[DOTOMP_LB]], align 4945// CHECK3-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4946// CHECK3-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4947// CHECK3-NEXT:    [[ADD5:%.*]] = add nsw i32 [[TMP15]], [[TMP16]]948// CHECK3-NEXT:    store i32 [[ADD5]], ptr [[DOTOMP_UB]], align 4949// CHECK3-NEXT:    br label [[OMP_DISPATCH_COND]]950// CHECK3:       omp.dispatch.end:951// CHECK3-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])952// CHECK3-NEXT:    ret void953//954//955// CHECK9-LABEL: define {{[^@]+}}@main956// CHECK9-SAME: (i32 noundef signext [[ARGC:%.*]], ptr noundef [[ARGV:%.*]]) #[[ATTR0:[0-9]+]] {957// CHECK9-NEXT:  entry:958// CHECK9-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4959// CHECK9-NEXT:    [[ARGC_ADDR:%.*]] = alloca i32, align 4960// CHECK9-NEXT:    [[ARGV_ADDR:%.*]] = alloca ptr, align 8961// CHECK9-NEXT:    [[N:%.*]] = alloca i32, align 4962// CHECK9-NEXT:    [[SAVED_STACK:%.*]] = alloca ptr, align 8963// CHECK9-NEXT:    [[__VLA_EXPR0:%.*]] = alloca i64, align 8964// CHECK9-NEXT:    [[N_CASTED:%.*]] = alloca i64, align 8965// CHECK9-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 8966// CHECK9-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 8967// CHECK9-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 8968// CHECK9-NEXT:    [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 8969// CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 4970// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4971// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4972// CHECK9-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8973// CHECK9-NEXT:    [[N_CASTED3:%.*]] = alloca i64, align 8974// CHECK9-NEXT:    [[DOTOFFLOAD_BASEPTRS4:%.*]] = alloca [3 x ptr], align 8975// CHECK9-NEXT:    [[DOTOFFLOAD_PTRS5:%.*]] = alloca [3 x ptr], align 8976// CHECK9-NEXT:    [[DOTOFFLOAD_MAPPERS6:%.*]] = alloca [3 x ptr], align 8977// CHECK9-NEXT:    [[DOTOFFLOAD_SIZES7:%.*]] = alloca [3 x i64], align 8978// CHECK9-NEXT:    [[_TMP8:%.*]] = alloca i32, align 4979// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_9:%.*]] = alloca i32, align 4980// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_10:%.*]] = alloca i32, align 4981// CHECK9-NEXT:    [[KERNEL_ARGS15:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8982// CHECK9-NEXT:    [[N_CASTED18:%.*]] = alloca i64, align 8983// CHECK9-NEXT:    [[DOTOFFLOAD_BASEPTRS19:%.*]] = alloca [3 x ptr], align 8984// CHECK9-NEXT:    [[DOTOFFLOAD_PTRS20:%.*]] = alloca [3 x ptr], align 8985// CHECK9-NEXT:    [[DOTOFFLOAD_MAPPERS21:%.*]] = alloca [3 x ptr], align 8986// CHECK9-NEXT:    [[DOTOFFLOAD_SIZES22:%.*]] = alloca [3 x i64], align 8987// CHECK9-NEXT:    [[_TMP23:%.*]] = alloca i32, align 4988// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_24:%.*]] = alloca i32, align 4989// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_25:%.*]] = alloca i32, align 4990// CHECK9-NEXT:    [[KERNEL_ARGS30:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8991// CHECK9-NEXT:    store i32 0, ptr [[RETVAL]], align 4992// CHECK9-NEXT:    store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 4993// CHECK9-NEXT:    store ptr [[ARGV]], ptr [[ARGV_ADDR]], align 8994// CHECK9-NEXT:    store i32 100, ptr [[N]], align 4995// CHECK9-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N]], align 4996// CHECK9-NEXT:    [[TMP1:%.*]] = zext i32 [[TMP0]] to i64997// CHECK9-NEXT:    [[TMP2:%.*]] = call ptr @llvm.stacksave.p0()998// CHECK9-NEXT:    store ptr [[TMP2]], ptr [[SAVED_STACK]], align 8999// CHECK9-NEXT:    [[VLA:%.*]] = alloca i32, i64 [[TMP1]], align 41000// CHECK9-NEXT:    store i64 [[TMP1]], ptr [[__VLA_EXPR0]], align 81001// CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[N]], align 41002// CHECK9-NEXT:    store i32 [[TMP3]], ptr [[N_CASTED]], align 41003// CHECK9-NEXT:    [[TMP4:%.*]] = load i64, ptr [[N_CASTED]], align 81004// CHECK9-NEXT:    [[TMP5:%.*]] = mul nuw i64 [[TMP1]], 41005// CHECK9-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES]], ptr align 8 @.offload_sizes, i64 24, i1 false)1006// CHECK9-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01007// CHECK9-NEXT:    store i64 [[TMP4]], ptr [[TMP6]], align 81008// CHECK9-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01009// CHECK9-NEXT:    store i64 [[TMP4]], ptr [[TMP7]], align 81010// CHECK9-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 01011// CHECK9-NEXT:    store ptr null, ptr [[TMP8]], align 81012// CHECK9-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11013// CHECK9-NEXT:    store i64 [[TMP1]], ptr [[TMP9]], align 81014// CHECK9-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11015// CHECK9-NEXT:    store i64 [[TMP1]], ptr [[TMP10]], align 81016// CHECK9-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 11017// CHECK9-NEXT:    store ptr null, ptr [[TMP11]], align 81018// CHECK9-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21019// CHECK9-NEXT:    store ptr [[VLA]], ptr [[TMP12]], align 81020// CHECK9-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21021// CHECK9-NEXT:    store ptr [[VLA]], ptr [[TMP13]], align 81022// CHECK9-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 21023// CHECK9-NEXT:    store i64 [[TMP5]], ptr [[TMP14]], align 81024// CHECK9-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 21025// CHECK9-NEXT:    store ptr null, ptr [[TMP15]], align 81026// CHECK9-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01027// CHECK9-NEXT:    [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01028// CHECK9-NEXT:    [[TMP18:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 01029// CHECK9-NEXT:    [[TMP19:%.*]] = load i32, ptr [[N]], align 41030// CHECK9-NEXT:    store i32 [[TMP19]], ptr [[DOTCAPTURE_EXPR_]], align 41031// CHECK9-NEXT:    [[TMP20:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41032// CHECK9-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP20]], 01033// CHECK9-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 11034// CHECK9-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11035// CHECK9-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41036// CHECK9-NEXT:    [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41037// CHECK9-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP21]], 11038// CHECK9-NEXT:    [[TMP22:%.*]] = zext i32 [[ADD]] to i641039// CHECK9-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01040// CHECK9-NEXT:    store i32 3, ptr [[TMP23]], align 41041// CHECK9-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11042// CHECK9-NEXT:    store i32 3, ptr [[TMP24]], align 41043// CHECK9-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21044// CHECK9-NEXT:    store ptr [[TMP16]], ptr [[TMP25]], align 81045// CHECK9-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31046// CHECK9-NEXT:    store ptr [[TMP17]], ptr [[TMP26]], align 81047// CHECK9-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41048// CHECK9-NEXT:    store ptr [[TMP18]], ptr [[TMP27]], align 81049// CHECK9-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51050// CHECK9-NEXT:    store ptr @.offload_maptypes, ptr [[TMP28]], align 81051// CHECK9-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61052// CHECK9-NEXT:    store ptr null, ptr [[TMP29]], align 81053// CHECK9-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71054// CHECK9-NEXT:    store ptr null, ptr [[TMP30]], align 81055// CHECK9-NEXT:    [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81056// CHECK9-NEXT:    store i64 [[TMP22]], ptr [[TMP31]], align 81057// CHECK9-NEXT:    [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91058// CHECK9-NEXT:    store i64 0, ptr [[TMP32]], align 81059// CHECK9-NEXT:    [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101060// CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP33]], align 41061// CHECK9-NEXT:    [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111062// CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP34]], align 41063// CHECK9-NEXT:    [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121064// CHECK9-NEXT:    store i32 0, ptr [[TMP35]], align 41065// CHECK9-NEXT:    [[TMP36:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l100.region_id, ptr [[KERNEL_ARGS]])1066// CHECK9-NEXT:    [[TMP37:%.*]] = icmp ne i32 [[TMP36]], 01067// CHECK9-NEXT:    br i1 [[TMP37]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1068// CHECK9:       omp_offload.failed:1069// CHECK9-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l100(i64 [[TMP4]], i64 [[TMP1]], ptr [[VLA]]) #[[ATTR3:[0-9]+]]1070// CHECK9-NEXT:    br label [[OMP_OFFLOAD_CONT]]1071// CHECK9:       omp_offload.cont:1072// CHECK9-NEXT:    [[TMP38:%.*]] = load i32, ptr [[N]], align 41073// CHECK9-NEXT:    store i32 [[TMP38]], ptr [[N_CASTED3]], align 41074// CHECK9-NEXT:    [[TMP39:%.*]] = load i64, ptr [[N_CASTED3]], align 81075// CHECK9-NEXT:    [[TMP40:%.*]] = mul nuw i64 [[TMP1]], 41076// CHECK9-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES7]], ptr align 8 @.offload_sizes.1, i64 24, i1 false)1077// CHECK9-NEXT:    [[TMP41:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 01078// CHECK9-NEXT:    store i64 [[TMP39]], ptr [[TMP41]], align 81079// CHECK9-NEXT:    [[TMP42:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 01080// CHECK9-NEXT:    store i64 [[TMP39]], ptr [[TMP42]], align 81081// CHECK9-NEXT:    [[TMP43:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i64 0, i64 01082// CHECK9-NEXT:    store ptr null, ptr [[TMP43]], align 81083// CHECK9-NEXT:    [[TMP44:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 11084// CHECK9-NEXT:    store i64 [[TMP1]], ptr [[TMP44]], align 81085// CHECK9-NEXT:    [[TMP45:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 11086// CHECK9-NEXT:    store i64 [[TMP1]], ptr [[TMP45]], align 81087// CHECK9-NEXT:    [[TMP46:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i64 0, i64 11088// CHECK9-NEXT:    store ptr null, ptr [[TMP46]], align 81089// CHECK9-NEXT:    [[TMP47:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 21090// CHECK9-NEXT:    store ptr [[VLA]], ptr [[TMP47]], align 81091// CHECK9-NEXT:    [[TMP48:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 21092// CHECK9-NEXT:    store ptr [[VLA]], ptr [[TMP48]], align 81093// CHECK9-NEXT:    [[TMP49:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES7]], i32 0, i32 21094// CHECK9-NEXT:    store i64 [[TMP40]], ptr [[TMP49]], align 81095// CHECK9-NEXT:    [[TMP50:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i64 0, i64 21096// CHECK9-NEXT:    store ptr null, ptr [[TMP50]], align 81097// CHECK9-NEXT:    [[TMP51:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 01098// CHECK9-NEXT:    [[TMP52:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 01099// CHECK9-NEXT:    [[TMP53:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES7]], i32 0, i32 01100// CHECK9-NEXT:    [[TMP54:%.*]] = load i32, ptr [[N]], align 41101// CHECK9-NEXT:    store i32 [[TMP54]], ptr [[DOTCAPTURE_EXPR_9]], align 41102// CHECK9-NEXT:    [[TMP55:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_9]], align 41103// CHECK9-NEXT:    [[SUB11:%.*]] = sub nsw i32 [[TMP55]], 01104// CHECK9-NEXT:    [[DIV12:%.*]] = sdiv i32 [[SUB11]], 11105// CHECK9-NEXT:    [[SUB13:%.*]] = sub nsw i32 [[DIV12]], 11106// CHECK9-NEXT:    store i32 [[SUB13]], ptr [[DOTCAPTURE_EXPR_10]], align 41107// CHECK9-NEXT:    [[TMP56:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_10]], align 41108// CHECK9-NEXT:    [[ADD14:%.*]] = add nsw i32 [[TMP56]], 11109// CHECK9-NEXT:    [[TMP57:%.*]] = zext i32 [[ADD14]] to i641110// CHECK9-NEXT:    [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 01111// CHECK9-NEXT:    store i32 3, ptr [[TMP58]], align 41112// CHECK9-NEXT:    [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 11113// CHECK9-NEXT:    store i32 3, ptr [[TMP59]], align 41114// CHECK9-NEXT:    [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 21115// CHECK9-NEXT:    store ptr [[TMP51]], ptr [[TMP60]], align 81116// CHECK9-NEXT:    [[TMP61:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 31117// CHECK9-NEXT:    store ptr [[TMP52]], ptr [[TMP61]], align 81118// CHECK9-NEXT:    [[TMP62:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 41119// CHECK9-NEXT:    store ptr [[TMP53]], ptr [[TMP62]], align 81120// CHECK9-NEXT:    [[TMP63:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 51121// CHECK9-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP63]], align 81122// CHECK9-NEXT:    [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 61123// CHECK9-NEXT:    store ptr null, ptr [[TMP64]], align 81124// CHECK9-NEXT:    [[TMP65:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 71125// CHECK9-NEXT:    store ptr null, ptr [[TMP65]], align 81126// CHECK9-NEXT:    [[TMP66:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 81127// CHECK9-NEXT:    store i64 [[TMP57]], ptr [[TMP66]], align 81128// CHECK9-NEXT:    [[TMP67:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 91129// CHECK9-NEXT:    store i64 0, ptr [[TMP67]], align 81130// CHECK9-NEXT:    [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 101131// CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP68]], align 41132// CHECK9-NEXT:    [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 111133// CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP69]], align 41134// CHECK9-NEXT:    [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 121135// CHECK9-NEXT:    store i32 0, ptr [[TMP70]], align 41136// CHECK9-NEXT:    [[TMP71:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l105.region_id, ptr [[KERNEL_ARGS15]])1137// CHECK9-NEXT:    [[TMP72:%.*]] = icmp ne i32 [[TMP71]], 01138// CHECK9-NEXT:    br i1 [[TMP72]], label [[OMP_OFFLOAD_FAILED16:%.*]], label [[OMP_OFFLOAD_CONT17:%.*]]1139// CHECK9:       omp_offload.failed16:1140// CHECK9-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l105(i64 [[TMP39]], i64 [[TMP1]], ptr [[VLA]]) #[[ATTR3]]1141// CHECK9-NEXT:    br label [[OMP_OFFLOAD_CONT17]]1142// CHECK9:       omp_offload.cont17:1143// CHECK9-NEXT:    [[TMP73:%.*]] = load i32, ptr [[N]], align 41144// CHECK9-NEXT:    store i32 [[TMP73]], ptr [[N_CASTED18]], align 41145// CHECK9-NEXT:    [[TMP74:%.*]] = load i64, ptr [[N_CASTED18]], align 81146// CHECK9-NEXT:    [[TMP75:%.*]] = mul nuw i64 [[TMP1]], 41147// CHECK9-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES22]], ptr align 8 @.offload_sizes.3, i64 24, i1 false)1148// CHECK9-NEXT:    [[TMP76:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS19]], i32 0, i32 01149// CHECK9-NEXT:    store i64 [[TMP74]], ptr [[TMP76]], align 81150// CHECK9-NEXT:    [[TMP77:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS20]], i32 0, i32 01151// CHECK9-NEXT:    store i64 [[TMP74]], ptr [[TMP77]], align 81152// CHECK9-NEXT:    [[TMP78:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS21]], i64 0, i64 01153// CHECK9-NEXT:    store ptr null, ptr [[TMP78]], align 81154// CHECK9-NEXT:    [[TMP79:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS19]], i32 0, i32 11155// CHECK9-NEXT:    store i64 [[TMP1]], ptr [[TMP79]], align 81156// CHECK9-NEXT:    [[TMP80:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS20]], i32 0, i32 11157// CHECK9-NEXT:    store i64 [[TMP1]], ptr [[TMP80]], align 81158// CHECK9-NEXT:    [[TMP81:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS21]], i64 0, i64 11159// CHECK9-NEXT:    store ptr null, ptr [[TMP81]], align 81160// CHECK9-NEXT:    [[TMP82:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS19]], i32 0, i32 21161// CHECK9-NEXT:    store ptr [[VLA]], ptr [[TMP82]], align 81162// CHECK9-NEXT:    [[TMP83:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS20]], i32 0, i32 21163// CHECK9-NEXT:    store ptr [[VLA]], ptr [[TMP83]], align 81164// CHECK9-NEXT:    [[TMP84:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES22]], i32 0, i32 21165// CHECK9-NEXT:    store i64 [[TMP75]], ptr [[TMP84]], align 81166// CHECK9-NEXT:    [[TMP85:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS21]], i64 0, i64 21167// CHECK9-NEXT:    store ptr null, ptr [[TMP85]], align 81168// CHECK9-NEXT:    [[TMP86:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS19]], i32 0, i32 01169// CHECK9-NEXT:    [[TMP87:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS20]], i32 0, i32 01170// CHECK9-NEXT:    [[TMP88:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES22]], i32 0, i32 01171// CHECK9-NEXT:    [[TMP89:%.*]] = load i32, ptr [[N]], align 41172// CHECK9-NEXT:    store i32 [[TMP89]], ptr [[DOTCAPTURE_EXPR_24]], align 41173// CHECK9-NEXT:    [[TMP90:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_24]], align 41174// CHECK9-NEXT:    [[SUB26:%.*]] = sub nsw i32 [[TMP90]], 01175// CHECK9-NEXT:    [[DIV27:%.*]] = sdiv i32 [[SUB26]], 11176// CHECK9-NEXT:    [[SUB28:%.*]] = sub nsw i32 [[DIV27]], 11177// CHECK9-NEXT:    store i32 [[SUB28]], ptr [[DOTCAPTURE_EXPR_25]], align 41178// CHECK9-NEXT:    [[TMP91:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_25]], align 41179// CHECK9-NEXT:    [[ADD29:%.*]] = add nsw i32 [[TMP91]], 11180// CHECK9-NEXT:    [[TMP92:%.*]] = zext i32 [[ADD29]] to i641181// CHECK9-NEXT:    [[TMP93:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 01182// CHECK9-NEXT:    store i32 3, ptr [[TMP93]], align 41183// CHECK9-NEXT:    [[TMP94:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 11184// CHECK9-NEXT:    store i32 3, ptr [[TMP94]], align 41185// CHECK9-NEXT:    [[TMP95:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 21186// CHECK9-NEXT:    store ptr [[TMP86]], ptr [[TMP95]], align 81187// CHECK9-NEXT:    [[TMP96:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 31188// CHECK9-NEXT:    store ptr [[TMP87]], ptr [[TMP96]], align 81189// CHECK9-NEXT:    [[TMP97:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 41190// CHECK9-NEXT:    store ptr [[TMP88]], ptr [[TMP97]], align 81191// CHECK9-NEXT:    [[TMP98:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 51192// CHECK9-NEXT:    store ptr @.offload_maptypes.4, ptr [[TMP98]], align 81193// CHECK9-NEXT:    [[TMP99:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 61194// CHECK9-NEXT:    store ptr null, ptr [[TMP99]], align 81195// CHECK9-NEXT:    [[TMP100:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 71196// CHECK9-NEXT:    store ptr null, ptr [[TMP100]], align 81197// CHECK9-NEXT:    [[TMP101:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 81198// CHECK9-NEXT:    store i64 [[TMP92]], ptr [[TMP101]], align 81199// CHECK9-NEXT:    [[TMP102:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 91200// CHECK9-NEXT:    store i64 0, ptr [[TMP102]], align 81201// CHECK9-NEXT:    [[TMP103:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 101202// CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP103]], align 41203// CHECK9-NEXT:    [[TMP104:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 111204// CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP104]], align 41205// CHECK9-NEXT:    [[TMP105:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 121206// CHECK9-NEXT:    store i32 0, ptr [[TMP105]], align 41207// CHECK9-NEXT:    [[TMP106:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l110.region_id, ptr [[KERNEL_ARGS30]])1208// CHECK9-NEXT:    [[TMP107:%.*]] = icmp ne i32 [[TMP106]], 01209// CHECK9-NEXT:    br i1 [[TMP107]], label [[OMP_OFFLOAD_FAILED31:%.*]], label [[OMP_OFFLOAD_CONT32:%.*]]1210// CHECK9:       omp_offload.failed31:1211// CHECK9-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l110(i64 [[TMP74]], i64 [[TMP1]], ptr [[VLA]]) #[[ATTR3]]1212// CHECK9-NEXT:    br label [[OMP_OFFLOAD_CONT32]]1213// CHECK9:       omp_offload.cont32:1214// CHECK9-NEXT:    [[TMP108:%.*]] = load i32, ptr [[ARGC_ADDR]], align 41215// CHECK9-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_Z5tmainIiLi10EEiT_(i32 noundef signext [[TMP108]])1216// CHECK9-NEXT:    store i32 [[CALL]], ptr [[RETVAL]], align 41217// CHECK9-NEXT:    [[TMP109:%.*]] = load ptr, ptr [[SAVED_STACK]], align 81218// CHECK9-NEXT:    call void @llvm.stackrestore.p0(ptr [[TMP109]])1219// CHECK9-NEXT:    [[TMP110:%.*]] = load i32, ptr [[RETVAL]], align 41220// CHECK9-NEXT:    ret i32 [[TMP110]]1221//1222//1223// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l1001224// CHECK9-SAME: (i64 noundef [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2:[0-9]+]] {1225// CHECK9-NEXT:  entry:1226// CHECK9-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 81227// CHECK9-NEXT:    [[VLA_ADDR:%.*]] = alloca i64, align 81228// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81229// CHECK9-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 81230// CHECK9-NEXT:    store i64 [[VLA]], ptr [[VLA_ADDR]], align 81231// CHECK9-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81232// CHECK9-NEXT:    [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR]], align 81233// CHECK9-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 81234// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l100.omp_outlined, ptr [[N_ADDR]], i64 [[TMP0]], ptr [[TMP1]])1235// CHECK9-NEXT:    ret void1236//1237//1238// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l100.omp_outlined1239// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {1240// CHECK9-NEXT:  entry:1241// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81242// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81243// CHECK9-NEXT:    [[N_ADDR:%.*]] = alloca ptr, align 81244// CHECK9-NEXT:    [[VLA_ADDR:%.*]] = alloca i64, align 81245// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81246// CHECK9-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41247// CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 41248// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41249// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41250// CHECK9-NEXT:    [[I:%.*]] = alloca i32, align 41251// CHECK9-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 41252// CHECK9-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 41253// CHECK9-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41254// CHECK9-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41255// CHECK9-NEXT:    [[I3:%.*]] = alloca i32, align 41256// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81257// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81258// CHECK9-NEXT:    store ptr [[N]], ptr [[N_ADDR]], align 81259// CHECK9-NEXT:    store i64 [[VLA]], ptr [[VLA_ADDR]], align 81260// CHECK9-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81261// CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 81262// CHECK9-NEXT:    [[TMP1:%.*]] = load i64, ptr [[VLA_ADDR]], align 81263// CHECK9-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 81264// CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 41265// CHECK9-NEXT:    store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR_]], align 41266// CHECK9-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41267// CHECK9-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP4]], 01268// CHECK9-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 11269// CHECK9-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11270// CHECK9-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41271// CHECK9-NEXT:    store i32 0, ptr [[I]], align 41272// CHECK9-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41273// CHECK9-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP5]]1274// CHECK9-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1275// CHECK9:       omp.precond.then:1276// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 41277// CHECK9-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41278// CHECK9-NEXT:    store i32 [[TMP6]], ptr [[DOTOMP_UB]], align 41279// CHECK9-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41280// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41281// CHECK9-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81282// CHECK9-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41283// CHECK9-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP8]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1284// CHECK9-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41285// CHECK9-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41286// CHECK9-NEXT:    [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]1287// CHECK9-NEXT:    br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1288// CHECK9:       cond.true:1289// CHECK9-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41290// CHECK9-NEXT:    br label [[COND_END:%.*]]1291// CHECK9:       cond.false:1292// CHECK9-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41293// CHECK9-NEXT:    br label [[COND_END]]1294// CHECK9:       cond.end:1295// CHECK9-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]1296// CHECK9-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 41297// CHECK9-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41298// CHECK9-NEXT:    store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 41299// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1300// CHECK9:       omp.inner.for.cond:1301// CHECK9-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41302// CHECK9-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41303// CHECK9-NEXT:    [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]1304// CHECK9-NEXT:    br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1305// CHECK9:       omp.inner.for.body:1306// CHECK9-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41307// CHECK9-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP16]], 11308// CHECK9-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1309// CHECK9-NEXT:    store i32 [[ADD]], ptr [[I3]], align 41310// CHECK9-NEXT:    [[TMP17:%.*]] = load i32, ptr [[I3]], align 41311// CHECK9-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP17]] to i641312// CHECK9-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i64 [[IDXPROM]]1313// CHECK9-NEXT:    store i32 0, ptr [[ARRAYIDX]], align 41314// CHECK9-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1315// CHECK9:       omp.body.continue:1316// CHECK9-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1317// CHECK9:       omp.inner.for.inc:1318// CHECK9-NEXT:    [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41319// CHECK9-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP18]], 11320// CHECK9-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 41321// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND]]1322// CHECK9:       omp.inner.for.end:1323// CHECK9-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1324// CHECK9:       omp.loop.exit:1325// CHECK9-NEXT:    [[TMP19:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81326// CHECK9-NEXT:    [[TMP20:%.*]] = load i32, ptr [[TMP19]], align 41327// CHECK9-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP20]])1328// CHECK9-NEXT:    br label [[OMP_PRECOND_END]]1329// CHECK9:       omp.precond.end:1330// CHECK9-NEXT:    ret void1331//1332//1333// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l1051334// CHECK9-SAME: (i64 noundef [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {1335// CHECK9-NEXT:  entry:1336// CHECK9-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 81337// CHECK9-NEXT:    [[VLA_ADDR:%.*]] = alloca i64, align 81338// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81339// CHECK9-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 81340// CHECK9-NEXT:    store i64 [[VLA]], ptr [[VLA_ADDR]], align 81341// CHECK9-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81342// CHECK9-NEXT:    [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR]], align 81343// CHECK9-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 81344// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l105.omp_outlined, ptr [[N_ADDR]], i64 [[TMP0]], ptr [[TMP1]])1345// CHECK9-NEXT:    ret void1346//1347//1348// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l105.omp_outlined1349// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {1350// CHECK9-NEXT:  entry:1351// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81352// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81353// CHECK9-NEXT:    [[N_ADDR:%.*]] = alloca ptr, align 81354// CHECK9-NEXT:    [[VLA_ADDR:%.*]] = alloca i64, align 81355// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81356// CHECK9-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41357// CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 41358// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41359// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41360// CHECK9-NEXT:    [[I:%.*]] = alloca i32, align 41361// CHECK9-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 41362// CHECK9-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 41363// CHECK9-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41364// CHECK9-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41365// CHECK9-NEXT:    [[I3:%.*]] = alloca i32, align 41366// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81367// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81368// CHECK9-NEXT:    store ptr [[N]], ptr [[N_ADDR]], align 81369// CHECK9-NEXT:    store i64 [[VLA]], ptr [[VLA_ADDR]], align 81370// CHECK9-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81371// CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 81372// CHECK9-NEXT:    [[TMP1:%.*]] = load i64, ptr [[VLA_ADDR]], align 81373// CHECK9-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 81374// CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 41375// CHECK9-NEXT:    store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR_]], align 41376// CHECK9-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41377// CHECK9-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP4]], 01378// CHECK9-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 11379// CHECK9-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11380// CHECK9-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41381// CHECK9-NEXT:    store i32 0, ptr [[I]], align 41382// CHECK9-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41383// CHECK9-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP5]]1384// CHECK9-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1385// CHECK9:       omp.precond.then:1386// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 41387// CHECK9-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41388// CHECK9-NEXT:    store i32 [[TMP6]], ptr [[DOTOMP_UB]], align 41389// CHECK9-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41390// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41391// CHECK9-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81392// CHECK9-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41393// CHECK9-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP8]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1394// CHECK9-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41395// CHECK9-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41396// CHECK9-NEXT:    [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]1397// CHECK9-NEXT:    br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1398// CHECK9:       cond.true:1399// CHECK9-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41400// CHECK9-NEXT:    br label [[COND_END:%.*]]1401// CHECK9:       cond.false:1402// CHECK9-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41403// CHECK9-NEXT:    br label [[COND_END]]1404// CHECK9:       cond.end:1405// CHECK9-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]1406// CHECK9-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 41407// CHECK9-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41408// CHECK9-NEXT:    store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 41409// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1410// CHECK9:       omp.inner.for.cond:1411// CHECK9-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41412// CHECK9-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41413// CHECK9-NEXT:    [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]1414// CHECK9-NEXT:    br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1415// CHECK9:       omp.inner.for.body:1416// CHECK9-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41417// CHECK9-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP16]], 11418// CHECK9-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1419// CHECK9-NEXT:    store i32 [[ADD]], ptr [[I3]], align 41420// CHECK9-NEXT:    [[TMP17:%.*]] = load i32, ptr [[I3]], align 41421// CHECK9-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP17]] to i641422// CHECK9-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i64 [[IDXPROM]]1423// CHECK9-NEXT:    store i32 0, ptr [[ARRAYIDX]], align 41424// CHECK9-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1425// CHECK9:       omp.body.continue:1426// CHECK9-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1427// CHECK9:       omp.inner.for.inc:1428// CHECK9-NEXT:    [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41429// CHECK9-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP18]], 11430// CHECK9-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 41431// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND]]1432// CHECK9:       omp.inner.for.end:1433// CHECK9-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1434// CHECK9:       omp.loop.exit:1435// CHECK9-NEXT:    [[TMP19:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81436// CHECK9-NEXT:    [[TMP20:%.*]] = load i32, ptr [[TMP19]], align 41437// CHECK9-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP20]])1438// CHECK9-NEXT:    br label [[OMP_PRECOND_END]]1439// CHECK9:       omp.precond.end:1440// CHECK9-NEXT:    ret void1441//1442//1443// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l1101444// CHECK9-SAME: (i64 noundef [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {1445// CHECK9-NEXT:  entry:1446// CHECK9-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 81447// CHECK9-NEXT:    [[VLA_ADDR:%.*]] = alloca i64, align 81448// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81449// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41450// CHECK9-NEXT:    [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 81451// CHECK9-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 81452// CHECK9-NEXT:    store i64 [[VLA]], ptr [[VLA_ADDR]], align 81453// CHECK9-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81454// CHECK9-NEXT:    [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR]], align 81455// CHECK9-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 81456// CHECK9-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 41457// CHECK9-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 41458// CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41459// CHECK9-NEXT:    store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 41460// CHECK9-NEXT:    [[TMP4:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 81461// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 4, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l110.omp_outlined, ptr [[N_ADDR]], i64 [[TMP0]], ptr [[TMP1]], i64 [[TMP4]])1462// CHECK9-NEXT:    ret void1463//1464//1465// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l110.omp_outlined1466// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR2]] {1467// CHECK9-NEXT:  entry:1468// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81469// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81470// CHECK9-NEXT:    [[N_ADDR:%.*]] = alloca ptr, align 81471// CHECK9-NEXT:    [[VLA_ADDR:%.*]] = alloca i64, align 81472// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81473// CHECK9-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 81474// CHECK9-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41475// CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 41476// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41477// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_2:%.*]] = alloca i32, align 41478// CHECK9-NEXT:    [[I:%.*]] = alloca i32, align 41479// CHECK9-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 41480// CHECK9-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 41481// CHECK9-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41482// CHECK9-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41483// CHECK9-NEXT:    [[I4:%.*]] = alloca i32, align 41484// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81485// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81486// CHECK9-NEXT:    store ptr [[N]], ptr [[N_ADDR]], align 81487// CHECK9-NEXT:    store i64 [[VLA]], ptr [[VLA_ADDR]], align 81488// CHECK9-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81489// CHECK9-NEXT:    store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 81490// CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 81491// CHECK9-NEXT:    [[TMP1:%.*]] = load i64, ptr [[VLA_ADDR]], align 81492// CHECK9-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 81493// CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 41494// CHECK9-NEXT:    store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR_1]], align 41495// CHECK9-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41496// CHECK9-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP4]], 01497// CHECK9-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 11498// CHECK9-NEXT:    [[SUB3:%.*]] = sub nsw i32 [[DIV]], 11499// CHECK9-NEXT:    store i32 [[SUB3]], ptr [[DOTCAPTURE_EXPR_2]], align 41500// CHECK9-NEXT:    store i32 0, ptr [[I]], align 41501// CHECK9-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41502// CHECK9-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP5]]1503// CHECK9-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1504// CHECK9:       omp.precond.then:1505// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 41506// CHECK9-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 41507// CHECK9-NEXT:    store i32 [[TMP6]], ptr [[DOTOMP_UB]], align 41508// CHECK9-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41509// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41510// CHECK9-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 41511// CHECK9-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81512// CHECK9-NEXT:    [[TMP9:%.*]] = load i32, ptr [[TMP8]], align 41513// CHECK9-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP9]], i32 91, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 [[TMP7]])1514// CHECK9-NEXT:    br label [[OMP_DISPATCH_COND:%.*]]1515// CHECK9:       omp.dispatch.cond:1516// CHECK9-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41517// CHECK9-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 41518// CHECK9-NEXT:    [[CMP5:%.*]] = icmp sgt i32 [[TMP10]], [[TMP11]]1519// CHECK9-NEXT:    br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1520// CHECK9:       cond.true:1521// CHECK9-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 41522// CHECK9-NEXT:    br label [[COND_END:%.*]]1523// CHECK9:       cond.false:1524// CHECK9-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41525// CHECK9-NEXT:    br label [[COND_END]]1526// CHECK9:       cond.end:1527// CHECK9-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP12]], [[COND_TRUE]] ], [ [[TMP13]], [[COND_FALSE]] ]1528// CHECK9-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 41529// CHECK9-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41530// CHECK9-NEXT:    store i32 [[TMP14]], ptr [[DOTOMP_IV]], align 41531// CHECK9-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41532// CHECK9-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41533// CHECK9-NEXT:    [[CMP6:%.*]] = icmp sle i32 [[TMP15]], [[TMP16]]1534// CHECK9-NEXT:    br i1 [[CMP6]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]1535// CHECK9:       omp.dispatch.body:1536// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1537// CHECK9:       omp.inner.for.cond:1538// CHECK9-NEXT:    [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11:![0-9]+]]1539// CHECK9-NEXT:    [[TMP18:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP11]]1540// CHECK9-NEXT:    [[CMP7:%.*]] = icmp sle i32 [[TMP17]], [[TMP18]]1541// CHECK9-NEXT:    br i1 [[CMP7]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1542// CHECK9:       omp.inner.for.body:1543// CHECK9-NEXT:    [[TMP19:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]1544// CHECK9-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP19]], 11545// CHECK9-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1546// CHECK9-NEXT:    store i32 [[ADD]], ptr [[I4]], align 4, !llvm.access.group [[ACC_GRP11]]1547// CHECK9-NEXT:    [[TMP20:%.*]] = load i32, ptr [[I4]], align 4, !llvm.access.group [[ACC_GRP11]]1548// CHECK9-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP20]] to i641549// CHECK9-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i64 [[IDXPROM]]1550// CHECK9-NEXT:    store i32 0, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP11]]1551// CHECK9-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1552// CHECK9:       omp.body.continue:1553// CHECK9-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1554// CHECK9:       omp.inner.for.inc:1555// CHECK9-NEXT:    [[TMP21:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]1556// CHECK9-NEXT:    [[ADD8:%.*]] = add nsw i32 [[TMP21]], 11557// CHECK9-NEXT:    store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]1558// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP12:![0-9]+]]1559// CHECK9:       omp.inner.for.end:1560// CHECK9-NEXT:    br label [[OMP_DISPATCH_INC:%.*]]1561// CHECK9:       omp.dispatch.inc:1562// CHECK9-NEXT:    [[TMP22:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41563// CHECK9-NEXT:    [[TMP23:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41564// CHECK9-NEXT:    [[ADD9:%.*]] = add nsw i32 [[TMP22]], [[TMP23]]1565// CHECK9-NEXT:    store i32 [[ADD9]], ptr [[DOTOMP_LB]], align 41566// CHECK9-NEXT:    [[TMP24:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41567// CHECK9-NEXT:    [[TMP25:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41568// CHECK9-NEXT:    [[ADD10:%.*]] = add nsw i32 [[TMP24]], [[TMP25]]1569// CHECK9-NEXT:    store i32 [[ADD10]], ptr [[DOTOMP_UB]], align 41570// CHECK9-NEXT:    br label [[OMP_DISPATCH_COND]]1571// CHECK9:       omp.dispatch.end:1572// CHECK9-NEXT:    [[TMP26:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81573// CHECK9-NEXT:    [[TMP27:%.*]] = load i32, ptr [[TMP26]], align 41574// CHECK9-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP27]])1575// CHECK9-NEXT:    br label [[OMP_PRECOND_END]]1576// CHECK9:       omp.precond.end:1577// CHECK9-NEXT:    ret void1578//1579//1580// CHECK9-LABEL: define {{[^@]+}}@_Z5tmainIiLi10EEiT_1581// CHECK9-SAME: (i32 noundef signext [[ARGC:%.*]]) #[[ATTR5:[0-9]+]] comdat {1582// CHECK9-NEXT:  entry:1583// CHECK9-NEXT:    [[ARGC_ADDR:%.*]] = alloca i32, align 41584// CHECK9-NEXT:    [[A:%.*]] = alloca [10 x i32], align 41585// CHECK9-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 81586// CHECK9-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 81587// CHECK9-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 81588// CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 41589// CHECK9-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81590// CHECK9-NEXT:    [[DOTOFFLOAD_BASEPTRS1:%.*]] = alloca [1 x ptr], align 81591// CHECK9-NEXT:    [[DOTOFFLOAD_PTRS2:%.*]] = alloca [1 x ptr], align 81592// CHECK9-NEXT:    [[DOTOFFLOAD_MAPPERS3:%.*]] = alloca [1 x ptr], align 81593// CHECK9-NEXT:    [[_TMP4:%.*]] = alloca i32, align 41594// CHECK9-NEXT:    [[KERNEL_ARGS5:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 81595// CHECK9-NEXT:    [[DOTOFFLOAD_BASEPTRS8:%.*]] = alloca [1 x ptr], align 81596// CHECK9-NEXT:    [[DOTOFFLOAD_PTRS9:%.*]] = alloca [1 x ptr], align 81597// CHECK9-NEXT:    [[DOTOFFLOAD_MAPPERS10:%.*]] = alloca [1 x ptr], align 81598// CHECK9-NEXT:    [[_TMP11:%.*]] = alloca i32, align 41599// CHECK9-NEXT:    [[KERNEL_ARGS12:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 81600// CHECK9-NEXT:    store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 41601// CHECK9-NEXT:    [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01602// CHECK9-NEXT:    store ptr [[A]], ptr [[TMP0]], align 81603// CHECK9-NEXT:    [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01604// CHECK9-NEXT:    store ptr [[A]], ptr [[TMP1]], align 81605// CHECK9-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 01606// CHECK9-NEXT:    store ptr null, ptr [[TMP2]], align 81607// CHECK9-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01608// CHECK9-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01609// CHECK9-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01610// CHECK9-NEXT:    store i32 3, ptr [[TMP5]], align 41611// CHECK9-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11612// CHECK9-NEXT:    store i32 1, ptr [[TMP6]], align 41613// CHECK9-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21614// CHECK9-NEXT:    store ptr [[TMP3]], ptr [[TMP7]], align 81615// CHECK9-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31616// CHECK9-NEXT:    store ptr [[TMP4]], ptr [[TMP8]], align 81617// CHECK9-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41618// CHECK9-NEXT:    store ptr @.offload_sizes.5, ptr [[TMP9]], align 81619// CHECK9-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51620// CHECK9-NEXT:    store ptr @.offload_maptypes.6, ptr [[TMP10]], align 81621// CHECK9-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61622// CHECK9-NEXT:    store ptr null, ptr [[TMP11]], align 81623// CHECK9-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71624// CHECK9-NEXT:    store ptr null, ptr [[TMP12]], align 81625// CHECK9-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81626// CHECK9-NEXT:    store i64 10, ptr [[TMP13]], align 81627// CHECK9-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91628// CHECK9-NEXT:    store i64 0, ptr [[TMP14]], align 81629// CHECK9-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101630// CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP15]], align 41631// CHECK9-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111632// CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP16]], align 41633// CHECK9-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121634// CHECK9-NEXT:    store i32 0, ptr [[TMP17]], align 41635// CHECK9-NEXT:    [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l79.region_id, ptr [[KERNEL_ARGS]])1636// CHECK9-NEXT:    [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 01637// CHECK9-NEXT:    br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1638// CHECK9:       omp_offload.failed:1639// CHECK9-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l79(ptr [[A]]) #[[ATTR3]]1640// CHECK9-NEXT:    br label [[OMP_OFFLOAD_CONT]]1641// CHECK9:       omp_offload.cont:1642// CHECK9-NEXT:    [[TMP20:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 01643// CHECK9-NEXT:    store ptr [[A]], ptr [[TMP20]], align 81644// CHECK9-NEXT:    [[TMP21:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 01645// CHECK9-NEXT:    store ptr [[A]], ptr [[TMP21]], align 81646// CHECK9-NEXT:    [[TMP22:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS3]], i64 0, i64 01647// CHECK9-NEXT:    store ptr null, ptr [[TMP22]], align 81648// CHECK9-NEXT:    [[TMP23:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 01649// CHECK9-NEXT:    [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 01650// CHECK9-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 01651// CHECK9-NEXT:    store i32 3, ptr [[TMP25]], align 41652// CHECK9-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 11653// CHECK9-NEXT:    store i32 1, ptr [[TMP26]], align 41654// CHECK9-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 21655// CHECK9-NEXT:    store ptr [[TMP23]], ptr [[TMP27]], align 81656// CHECK9-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 31657// CHECK9-NEXT:    store ptr [[TMP24]], ptr [[TMP28]], align 81658// CHECK9-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 41659// CHECK9-NEXT:    store ptr @.offload_sizes.7, ptr [[TMP29]], align 81660// CHECK9-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 51661// CHECK9-NEXT:    store ptr @.offload_maptypes.8, ptr [[TMP30]], align 81662// CHECK9-NEXT:    [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 61663// CHECK9-NEXT:    store ptr null, ptr [[TMP31]], align 81664// CHECK9-NEXT:    [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 71665// CHECK9-NEXT:    store ptr null, ptr [[TMP32]], align 81666// CHECK9-NEXT:    [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 81667// CHECK9-NEXT:    store i64 10, ptr [[TMP33]], align 81668// CHECK9-NEXT:    [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 91669// CHECK9-NEXT:    store i64 0, ptr [[TMP34]], align 81670// CHECK9-NEXT:    [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 101671// CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP35]], align 41672// CHECK9-NEXT:    [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 111673// CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP36]], align 41674// CHECK9-NEXT:    [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 121675// CHECK9-NEXT:    store i32 0, ptr [[TMP37]], align 41676// CHECK9-NEXT:    [[TMP38:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l84.region_id, ptr [[KERNEL_ARGS5]])1677// CHECK9-NEXT:    [[TMP39:%.*]] = icmp ne i32 [[TMP38]], 01678// CHECK9-NEXT:    br i1 [[TMP39]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]]1679// CHECK9:       omp_offload.failed6:1680// CHECK9-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l84(ptr [[A]]) #[[ATTR3]]1681// CHECK9-NEXT:    br label [[OMP_OFFLOAD_CONT7]]1682// CHECK9:       omp_offload.cont7:1683// CHECK9-NEXT:    [[TMP40:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 01684// CHECK9-NEXT:    store ptr [[A]], ptr [[TMP40]], align 81685// CHECK9-NEXT:    [[TMP41:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 01686// CHECK9-NEXT:    store ptr [[A]], ptr [[TMP41]], align 81687// CHECK9-NEXT:    [[TMP42:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS10]], i64 0, i64 01688// CHECK9-NEXT:    store ptr null, ptr [[TMP42]], align 81689// CHECK9-NEXT:    [[TMP43:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 01690// CHECK9-NEXT:    [[TMP44:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 01691// CHECK9-NEXT:    [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 01692// CHECK9-NEXT:    store i32 3, ptr [[TMP45]], align 41693// CHECK9-NEXT:    [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 11694// CHECK9-NEXT:    store i32 1, ptr [[TMP46]], align 41695// CHECK9-NEXT:    [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 21696// CHECK9-NEXT:    store ptr [[TMP43]], ptr [[TMP47]], align 81697// CHECK9-NEXT:    [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 31698// CHECK9-NEXT:    store ptr [[TMP44]], ptr [[TMP48]], align 81699// CHECK9-NEXT:    [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 41700// CHECK9-NEXT:    store ptr @.offload_sizes.9, ptr [[TMP49]], align 81701// CHECK9-NEXT:    [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 51702// CHECK9-NEXT:    store ptr @.offload_maptypes.10, ptr [[TMP50]], align 81703// CHECK9-NEXT:    [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 61704// CHECK9-NEXT:    store ptr null, ptr [[TMP51]], align 81705// CHECK9-NEXT:    [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 71706// CHECK9-NEXT:    store ptr null, ptr [[TMP52]], align 81707// CHECK9-NEXT:    [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 81708// CHECK9-NEXT:    store i64 10, ptr [[TMP53]], align 81709// CHECK9-NEXT:    [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 91710// CHECK9-NEXT:    store i64 0, ptr [[TMP54]], align 81711// CHECK9-NEXT:    [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 101712// CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP55]], align 41713// CHECK9-NEXT:    [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 111714// CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP56]], align 41715// CHECK9-NEXT:    [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 121716// CHECK9-NEXT:    store i32 0, ptr [[TMP57]], align 41717// CHECK9-NEXT:    [[TMP58:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l89.region_id, ptr [[KERNEL_ARGS12]])1718// CHECK9-NEXT:    [[TMP59:%.*]] = icmp ne i32 [[TMP58]], 01719// CHECK9-NEXT:    br i1 [[TMP59]], label [[OMP_OFFLOAD_FAILED13:%.*]], label [[OMP_OFFLOAD_CONT14:%.*]]1720// CHECK9:       omp_offload.failed13:1721// CHECK9-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l89(ptr [[A]]) #[[ATTR3]]1722// CHECK9-NEXT:    br label [[OMP_OFFLOAD_CONT14]]1723// CHECK9:       omp_offload.cont14:1724// CHECK9-NEXT:    ret i32 01725//1726//1727// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l791728// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {1729// CHECK9-NEXT:  entry:1730// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81731// CHECK9-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81732// CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81733// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l79.omp_outlined, ptr [[TMP0]])1734// CHECK9-NEXT:    ret void1735//1736//1737// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l79.omp_outlined1738// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {1739// CHECK9-NEXT:  entry:1740// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81741// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81742// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81743// CHECK9-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41744// CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 41745// CHECK9-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 41746// CHECK9-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 41747// CHECK9-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41748// CHECK9-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41749// CHECK9-NEXT:    [[I:%.*]] = alloca i32, align 41750// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81751// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81752// CHECK9-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81753// CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81754// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 41755// CHECK9-NEXT:    store i32 9, ptr [[DOTOMP_UB]], align 41756// CHECK9-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41757// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41758// CHECK9-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81759// CHECK9-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41760// CHECK9-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1761// CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41762// CHECK9-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 91763// CHECK9-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1764// CHECK9:       cond.true:1765// CHECK9-NEXT:    br label [[COND_END:%.*]]1766// CHECK9:       cond.false:1767// CHECK9-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41768// CHECK9-NEXT:    br label [[COND_END]]1769// CHECK9:       cond.end:1770// CHECK9-NEXT:    [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]1771// CHECK9-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 41772// CHECK9-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41773// CHECK9-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 41774// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1775// CHECK9:       omp.inner.for.cond:1776// CHECK9-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41777// CHECK9-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41778// CHECK9-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]1779// CHECK9-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1780// CHECK9:       omp.inner.for.body:1781// CHECK9-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41782// CHECK9-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP8]], 11783// CHECK9-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1784// CHECK9-NEXT:    store i32 [[ADD]], ptr [[I]], align 41785// CHECK9-NEXT:    [[TMP9:%.*]] = load i32, ptr [[I]], align 41786// CHECK9-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP9]] to i641787// CHECK9-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]1788// CHECK9-NEXT:    store i32 0, ptr [[ARRAYIDX]], align 41789// CHECK9-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1790// CHECK9:       omp.body.continue:1791// CHECK9-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1792// CHECK9:       omp.inner.for.inc:1793// CHECK9-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41794// CHECK9-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP10]], 11795// CHECK9-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 41796// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND]]1797// CHECK9:       omp.inner.for.end:1798// CHECK9-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1799// CHECK9:       omp.loop.exit:1800// CHECK9-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])1801// CHECK9-NEXT:    ret void1802//1803//1804// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l841805// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {1806// CHECK9-NEXT:  entry:1807// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81808// CHECK9-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81809// CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81810// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l84.omp_outlined, ptr [[TMP0]])1811// CHECK9-NEXT:    ret void1812//1813//1814// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l84.omp_outlined1815// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {1816// CHECK9-NEXT:  entry:1817// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81818// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81819// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81820// CHECK9-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41821// CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 41822// CHECK9-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 41823// CHECK9-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 41824// CHECK9-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41825// CHECK9-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41826// CHECK9-NEXT:    [[I:%.*]] = alloca i32, align 41827// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81828// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81829// CHECK9-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81830// CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81831// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 41832// CHECK9-NEXT:    store i32 9, ptr [[DOTOMP_UB]], align 41833// CHECK9-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41834// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41835// CHECK9-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81836// CHECK9-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41837// CHECK9-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1838// CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41839// CHECK9-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 91840// CHECK9-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1841// CHECK9:       cond.true:1842// CHECK9-NEXT:    br label [[COND_END:%.*]]1843// CHECK9:       cond.false:1844// CHECK9-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41845// CHECK9-NEXT:    br label [[COND_END]]1846// CHECK9:       cond.end:1847// CHECK9-NEXT:    [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]1848// CHECK9-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 41849// CHECK9-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41850// CHECK9-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 41851// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1852// CHECK9:       omp.inner.for.cond:1853// CHECK9-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41854// CHECK9-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41855// CHECK9-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]1856// CHECK9-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1857// CHECK9:       omp.inner.for.body:1858// CHECK9-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41859// CHECK9-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP8]], 11860// CHECK9-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1861// CHECK9-NEXT:    store i32 [[ADD]], ptr [[I]], align 41862// CHECK9-NEXT:    [[TMP9:%.*]] = load i32, ptr [[I]], align 41863// CHECK9-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP9]] to i641864// CHECK9-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]1865// CHECK9-NEXT:    store i32 0, ptr [[ARRAYIDX]], align 41866// CHECK9-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1867// CHECK9:       omp.body.continue:1868// CHECK9-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1869// CHECK9:       omp.inner.for.inc:1870// CHECK9-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41871// CHECK9-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP10]], 11872// CHECK9-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 41873// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND]]1874// CHECK9:       omp.inner.for.end:1875// CHECK9-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1876// CHECK9:       omp.loop.exit:1877// CHECK9-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])1878// CHECK9-NEXT:    ret void1879//1880//1881// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l891882// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {1883// CHECK9-NEXT:  entry:1884// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81885// CHECK9-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81886// CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81887// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l89.omp_outlined, ptr [[TMP0]])1888// CHECK9-NEXT:    ret void1889//1890//1891// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l89.omp_outlined1892// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {1893// CHECK9-NEXT:  entry:1894// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81895// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81896// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81897// CHECK9-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41898// CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 41899// CHECK9-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 41900// CHECK9-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 41901// CHECK9-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41902// CHECK9-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41903// CHECK9-NEXT:    [[I:%.*]] = alloca i32, align 41904// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81905// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81906// CHECK9-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81907// CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81908// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 41909// CHECK9-NEXT:    store i32 9, ptr [[DOTOMP_UB]], align 41910// CHECK9-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41911// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41912// CHECK9-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81913// CHECK9-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41914// CHECK9-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 91, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 10)1915// CHECK9-NEXT:    br label [[OMP_DISPATCH_COND:%.*]]1916// CHECK9:       omp.dispatch.cond:1917// CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41918// CHECK9-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 91919// CHECK9-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1920// CHECK9:       cond.true:1921// CHECK9-NEXT:    br label [[COND_END:%.*]]1922// CHECK9:       cond.false:1923// CHECK9-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41924// CHECK9-NEXT:    br label [[COND_END]]1925// CHECK9:       cond.end:1926// CHECK9-NEXT:    [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]1927// CHECK9-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 41928// CHECK9-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41929// CHECK9-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 41930// CHECK9-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41931// CHECK9-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41932// CHECK9-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]1933// CHECK9-NEXT:    br i1 [[CMP1]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]1934// CHECK9:       omp.dispatch.body:1935// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1936// CHECK9:       omp.inner.for.cond:1937// CHECK9-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP14:![0-9]+]]1938// CHECK9-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP14]]1939// CHECK9-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP8]], [[TMP9]]1940// CHECK9-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1941// CHECK9:       omp.inner.for.body:1942// CHECK9-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP14]]1943// CHECK9-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP10]], 11944// CHECK9-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1945// CHECK9-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP14]]1946// CHECK9-NEXT:    [[TMP11:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP14]]1947// CHECK9-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP11]] to i641948// CHECK9-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]1949// CHECK9-NEXT:    store i32 0, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP14]]1950// CHECK9-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1951// CHECK9:       omp.body.continue:1952// CHECK9-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1953// CHECK9:       omp.inner.for.inc:1954// CHECK9-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP14]]1955// CHECK9-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP12]], 11956// CHECK9-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP14]]1957// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP15:![0-9]+]]1958// CHECK9:       omp.inner.for.end:1959// CHECK9-NEXT:    br label [[OMP_DISPATCH_INC:%.*]]1960// CHECK9:       omp.dispatch.inc:1961// CHECK9-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41962// CHECK9-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41963// CHECK9-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP13]], [[TMP14]]1964// CHECK9-NEXT:    store i32 [[ADD4]], ptr [[DOTOMP_LB]], align 41965// CHECK9-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41966// CHECK9-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41967// CHECK9-NEXT:    [[ADD5:%.*]] = add nsw i32 [[TMP15]], [[TMP16]]1968// CHECK9-NEXT:    store i32 [[ADD5]], ptr [[DOTOMP_UB]], align 41969// CHECK9-NEXT:    br label [[OMP_DISPATCH_COND]]1970// CHECK9:       omp.dispatch.end:1971// CHECK9-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])1972// CHECK9-NEXT:    ret void1973//1974//1975// CHECK11-LABEL: define {{[^@]+}}@main1976// CHECK11-SAME: (i32 noundef [[ARGC:%.*]], ptr noundef [[ARGV:%.*]]) #[[ATTR0:[0-9]+]] {1977// CHECK11-NEXT:  entry:1978// CHECK11-NEXT:    [[RETVAL:%.*]] = alloca i32, align 41979// CHECK11-NEXT:    [[ARGC_ADDR:%.*]] = alloca i32, align 41980// CHECK11-NEXT:    [[ARGV_ADDR:%.*]] = alloca ptr, align 41981// CHECK11-NEXT:    [[N:%.*]] = alloca i32, align 41982// CHECK11-NEXT:    [[SAVED_STACK:%.*]] = alloca ptr, align 41983// CHECK11-NEXT:    [[__VLA_EXPR0:%.*]] = alloca i32, align 41984// CHECK11-NEXT:    [[N_CASTED:%.*]] = alloca i32, align 41985// CHECK11-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 41986// CHECK11-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 41987// CHECK11-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 41988// CHECK11-NEXT:    [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 41989// CHECK11-NEXT:    [[TMP:%.*]] = alloca i32, align 41990// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41991// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41992// CHECK11-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81993// CHECK11-NEXT:    [[N_CASTED3:%.*]] = alloca i32, align 41994// CHECK11-NEXT:    [[DOTOFFLOAD_BASEPTRS4:%.*]] = alloca [3 x ptr], align 41995// CHECK11-NEXT:    [[DOTOFFLOAD_PTRS5:%.*]] = alloca [3 x ptr], align 41996// CHECK11-NEXT:    [[DOTOFFLOAD_MAPPERS6:%.*]] = alloca [3 x ptr], align 41997// CHECK11-NEXT:    [[DOTOFFLOAD_SIZES7:%.*]] = alloca [3 x i64], align 41998// CHECK11-NEXT:    [[_TMP8:%.*]] = alloca i32, align 41999// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_9:%.*]] = alloca i32, align 42000// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_10:%.*]] = alloca i32, align 42001// CHECK11-NEXT:    [[KERNEL_ARGS15:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 82002// CHECK11-NEXT:    [[N_CASTED18:%.*]] = alloca i32, align 42003// CHECK11-NEXT:    [[DOTOFFLOAD_BASEPTRS19:%.*]] = alloca [3 x ptr], align 42004// CHECK11-NEXT:    [[DOTOFFLOAD_PTRS20:%.*]] = alloca [3 x ptr], align 42005// CHECK11-NEXT:    [[DOTOFFLOAD_MAPPERS21:%.*]] = alloca [3 x ptr], align 42006// CHECK11-NEXT:    [[DOTOFFLOAD_SIZES22:%.*]] = alloca [3 x i64], align 42007// CHECK11-NEXT:    [[_TMP23:%.*]] = alloca i32, align 42008// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_24:%.*]] = alloca i32, align 42009// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_25:%.*]] = alloca i32, align 42010// CHECK11-NEXT:    [[KERNEL_ARGS30:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 82011// CHECK11-NEXT:    store i32 0, ptr [[RETVAL]], align 42012// CHECK11-NEXT:    store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 42013// CHECK11-NEXT:    store ptr [[ARGV]], ptr [[ARGV_ADDR]], align 42014// CHECK11-NEXT:    store i32 100, ptr [[N]], align 42015// CHECK11-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N]], align 42016// CHECK11-NEXT:    [[TMP1:%.*]] = call ptr @llvm.stacksave.p0()2017// CHECK11-NEXT:    store ptr [[TMP1]], ptr [[SAVED_STACK]], align 42018// CHECK11-NEXT:    [[VLA:%.*]] = alloca i32, i32 [[TMP0]], align 42019// CHECK11-NEXT:    store i32 [[TMP0]], ptr [[__VLA_EXPR0]], align 42020// CHECK11-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N]], align 42021// CHECK11-NEXT:    store i32 [[TMP2]], ptr [[N_CASTED]], align 42022// CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[N_CASTED]], align 42023// CHECK11-NEXT:    [[TMP4:%.*]] = mul nuw i32 [[TMP0]], 42024// CHECK11-NEXT:    [[TMP5:%.*]] = sext i32 [[TMP4]] to i642025// CHECK11-NEXT:    call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[DOTOFFLOAD_SIZES]], ptr align 4 @.offload_sizes, i32 24, i1 false)2026// CHECK11-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 02027// CHECK11-NEXT:    store i32 [[TMP3]], ptr [[TMP6]], align 42028// CHECK11-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 02029// CHECK11-NEXT:    store i32 [[TMP3]], ptr [[TMP7]], align 42030// CHECK11-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 02031// CHECK11-NEXT:    store ptr null, ptr [[TMP8]], align 42032// CHECK11-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 12033// CHECK11-NEXT:    store i32 [[TMP0]], ptr [[TMP9]], align 42034// CHECK11-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 12035// CHECK11-NEXT:    store i32 [[TMP0]], ptr [[TMP10]], align 42036// CHECK11-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 12037// CHECK11-NEXT:    store ptr null, ptr [[TMP11]], align 42038// CHECK11-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 22039// CHECK11-NEXT:    store ptr [[VLA]], ptr [[TMP12]], align 42040// CHECK11-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 22041// CHECK11-NEXT:    store ptr [[VLA]], ptr [[TMP13]], align 42042// CHECK11-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 22043// CHECK11-NEXT:    store i64 [[TMP5]], ptr [[TMP14]], align 42044// CHECK11-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 22045// CHECK11-NEXT:    store ptr null, ptr [[TMP15]], align 42046// CHECK11-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 02047// CHECK11-NEXT:    [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 02048// CHECK11-NEXT:    [[TMP18:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 02049// CHECK11-NEXT:    [[TMP19:%.*]] = load i32, ptr [[N]], align 42050// CHECK11-NEXT:    store i32 [[TMP19]], ptr [[DOTCAPTURE_EXPR_]], align 42051// CHECK11-NEXT:    [[TMP20:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42052// CHECK11-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP20]], 02053// CHECK11-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 12054// CHECK11-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12055// CHECK11-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42056// CHECK11-NEXT:    [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42057// CHECK11-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP21]], 12058// CHECK11-NEXT:    [[TMP22:%.*]] = zext i32 [[ADD]] to i642059// CHECK11-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 02060// CHECK11-NEXT:    store i32 3, ptr [[TMP23]], align 42061// CHECK11-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12062// CHECK11-NEXT:    store i32 3, ptr [[TMP24]], align 42063// CHECK11-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 22064// CHECK11-NEXT:    store ptr [[TMP16]], ptr [[TMP25]], align 42065// CHECK11-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 32066// CHECK11-NEXT:    store ptr [[TMP17]], ptr [[TMP26]], align 42067// CHECK11-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 42068// CHECK11-NEXT:    store ptr [[TMP18]], ptr [[TMP27]], align 42069// CHECK11-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 52070// CHECK11-NEXT:    store ptr @.offload_maptypes, ptr [[TMP28]], align 42071// CHECK11-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 62072// CHECK11-NEXT:    store ptr null, ptr [[TMP29]], align 42073// CHECK11-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 72074// CHECK11-NEXT:    store ptr null, ptr [[TMP30]], align 42075// CHECK11-NEXT:    [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 82076// CHECK11-NEXT:    store i64 [[TMP22]], ptr [[TMP31]], align 82077// CHECK11-NEXT:    [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 92078// CHECK11-NEXT:    store i64 0, ptr [[TMP32]], align 82079// CHECK11-NEXT:    [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 102080// CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP33]], align 42081// CHECK11-NEXT:    [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 112082// CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP34]], align 42083// CHECK11-NEXT:    [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 122084// CHECK11-NEXT:    store i32 0, ptr [[TMP35]], align 42085// CHECK11-NEXT:    [[TMP36:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l100.region_id, ptr [[KERNEL_ARGS]])2086// CHECK11-NEXT:    [[TMP37:%.*]] = icmp ne i32 [[TMP36]], 02087// CHECK11-NEXT:    br i1 [[TMP37]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]2088// CHECK11:       omp_offload.failed:2089// CHECK11-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l100(i32 [[TMP3]], i32 [[TMP0]], ptr [[VLA]]) #[[ATTR3:[0-9]+]]2090// CHECK11-NEXT:    br label [[OMP_OFFLOAD_CONT]]2091// CHECK11:       omp_offload.cont:2092// CHECK11-NEXT:    [[TMP38:%.*]] = load i32, ptr [[N]], align 42093// CHECK11-NEXT:    store i32 [[TMP38]], ptr [[N_CASTED3]], align 42094// CHECK11-NEXT:    [[TMP39:%.*]] = load i32, ptr [[N_CASTED3]], align 42095// CHECK11-NEXT:    [[TMP40:%.*]] = mul nuw i32 [[TMP0]], 42096// CHECK11-NEXT:    [[TMP41:%.*]] = sext i32 [[TMP40]] to i642097// CHECK11-NEXT:    call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[DOTOFFLOAD_SIZES7]], ptr align 4 @.offload_sizes.1, i32 24, i1 false)2098// CHECK11-NEXT:    [[TMP42:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 02099// CHECK11-NEXT:    store i32 [[TMP39]], ptr [[TMP42]], align 42100// CHECK11-NEXT:    [[TMP43:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 02101// CHECK11-NEXT:    store i32 [[TMP39]], ptr [[TMP43]], align 42102// CHECK11-NEXT:    [[TMP44:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i32 0, i32 02103// CHECK11-NEXT:    store ptr null, ptr [[TMP44]], align 42104// CHECK11-NEXT:    [[TMP45:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 12105// CHECK11-NEXT:    store i32 [[TMP0]], ptr [[TMP45]], align 42106// CHECK11-NEXT:    [[TMP46:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 12107// CHECK11-NEXT:    store i32 [[TMP0]], ptr [[TMP46]], align 42108// CHECK11-NEXT:    [[TMP47:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i32 0, i32 12109// CHECK11-NEXT:    store ptr null, ptr [[TMP47]], align 42110// CHECK11-NEXT:    [[TMP48:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 22111// CHECK11-NEXT:    store ptr [[VLA]], ptr [[TMP48]], align 42112// CHECK11-NEXT:    [[TMP49:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 22113// CHECK11-NEXT:    store ptr [[VLA]], ptr [[TMP49]], align 42114// CHECK11-NEXT:    [[TMP50:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES7]], i32 0, i32 22115// CHECK11-NEXT:    store i64 [[TMP41]], ptr [[TMP50]], align 42116// CHECK11-NEXT:    [[TMP51:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i32 0, i32 22117// CHECK11-NEXT:    store ptr null, ptr [[TMP51]], align 42118// CHECK11-NEXT:    [[TMP52:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 02119// CHECK11-NEXT:    [[TMP53:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 02120// CHECK11-NEXT:    [[TMP54:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES7]], i32 0, i32 02121// CHECK11-NEXT:    [[TMP55:%.*]] = load i32, ptr [[N]], align 42122// CHECK11-NEXT:    store i32 [[TMP55]], ptr [[DOTCAPTURE_EXPR_9]], align 42123// CHECK11-NEXT:    [[TMP56:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_9]], align 42124// CHECK11-NEXT:    [[SUB11:%.*]] = sub nsw i32 [[TMP56]], 02125// CHECK11-NEXT:    [[DIV12:%.*]] = sdiv i32 [[SUB11]], 12126// CHECK11-NEXT:    [[SUB13:%.*]] = sub nsw i32 [[DIV12]], 12127// CHECK11-NEXT:    store i32 [[SUB13]], ptr [[DOTCAPTURE_EXPR_10]], align 42128// CHECK11-NEXT:    [[TMP57:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_10]], align 42129// CHECK11-NEXT:    [[ADD14:%.*]] = add nsw i32 [[TMP57]], 12130// CHECK11-NEXT:    [[TMP58:%.*]] = zext i32 [[ADD14]] to i642131// CHECK11-NEXT:    [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 02132// CHECK11-NEXT:    store i32 3, ptr [[TMP59]], align 42133// CHECK11-NEXT:    [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 12134// CHECK11-NEXT:    store i32 3, ptr [[TMP60]], align 42135// CHECK11-NEXT:    [[TMP61:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 22136// CHECK11-NEXT:    store ptr [[TMP52]], ptr [[TMP61]], align 42137// CHECK11-NEXT:    [[TMP62:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 32138// CHECK11-NEXT:    store ptr [[TMP53]], ptr [[TMP62]], align 42139// CHECK11-NEXT:    [[TMP63:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 42140// CHECK11-NEXT:    store ptr [[TMP54]], ptr [[TMP63]], align 42141// CHECK11-NEXT:    [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 52142// CHECK11-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP64]], align 42143// CHECK11-NEXT:    [[TMP65:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 62144// CHECK11-NEXT:    store ptr null, ptr [[TMP65]], align 42145// CHECK11-NEXT:    [[TMP66:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 72146// CHECK11-NEXT:    store ptr null, ptr [[TMP66]], align 42147// CHECK11-NEXT:    [[TMP67:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 82148// CHECK11-NEXT:    store i64 [[TMP58]], ptr [[TMP67]], align 82149// CHECK11-NEXT:    [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 92150// CHECK11-NEXT:    store i64 0, ptr [[TMP68]], align 82151// CHECK11-NEXT:    [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 102152// CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP69]], align 42153// CHECK11-NEXT:    [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 112154// CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP70]], align 42155// CHECK11-NEXT:    [[TMP71:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 122156// CHECK11-NEXT:    store i32 0, ptr [[TMP71]], align 42157// CHECK11-NEXT:    [[TMP72:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l105.region_id, ptr [[KERNEL_ARGS15]])2158// CHECK11-NEXT:    [[TMP73:%.*]] = icmp ne i32 [[TMP72]], 02159// CHECK11-NEXT:    br i1 [[TMP73]], label [[OMP_OFFLOAD_FAILED16:%.*]], label [[OMP_OFFLOAD_CONT17:%.*]]2160// CHECK11:       omp_offload.failed16:2161// CHECK11-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l105(i32 [[TMP39]], i32 [[TMP0]], ptr [[VLA]]) #[[ATTR3]]2162// CHECK11-NEXT:    br label [[OMP_OFFLOAD_CONT17]]2163// CHECK11:       omp_offload.cont17:2164// CHECK11-NEXT:    [[TMP74:%.*]] = load i32, ptr [[N]], align 42165// CHECK11-NEXT:    store i32 [[TMP74]], ptr [[N_CASTED18]], align 42166// CHECK11-NEXT:    [[TMP75:%.*]] = load i32, ptr [[N_CASTED18]], align 42167// CHECK11-NEXT:    [[TMP76:%.*]] = mul nuw i32 [[TMP0]], 42168// CHECK11-NEXT:    [[TMP77:%.*]] = sext i32 [[TMP76]] to i642169// CHECK11-NEXT:    call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[DOTOFFLOAD_SIZES22]], ptr align 4 @.offload_sizes.3, i32 24, i1 false)2170// CHECK11-NEXT:    [[TMP78:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS19]], i32 0, i32 02171// CHECK11-NEXT:    store i32 [[TMP75]], ptr [[TMP78]], align 42172// CHECK11-NEXT:    [[TMP79:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS20]], i32 0, i32 02173// CHECK11-NEXT:    store i32 [[TMP75]], ptr [[TMP79]], align 42174// CHECK11-NEXT:    [[TMP80:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS21]], i32 0, i32 02175// CHECK11-NEXT:    store ptr null, ptr [[TMP80]], align 42176// CHECK11-NEXT:    [[TMP81:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS19]], i32 0, i32 12177// CHECK11-NEXT:    store i32 [[TMP0]], ptr [[TMP81]], align 42178// CHECK11-NEXT:    [[TMP82:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS20]], i32 0, i32 12179// CHECK11-NEXT:    store i32 [[TMP0]], ptr [[TMP82]], align 42180// CHECK11-NEXT:    [[TMP83:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS21]], i32 0, i32 12181// CHECK11-NEXT:    store ptr null, ptr [[TMP83]], align 42182// CHECK11-NEXT:    [[TMP84:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS19]], i32 0, i32 22183// CHECK11-NEXT:    store ptr [[VLA]], ptr [[TMP84]], align 42184// CHECK11-NEXT:    [[TMP85:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS20]], i32 0, i32 22185// CHECK11-NEXT:    store ptr [[VLA]], ptr [[TMP85]], align 42186// CHECK11-NEXT:    [[TMP86:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES22]], i32 0, i32 22187// CHECK11-NEXT:    store i64 [[TMP77]], ptr [[TMP86]], align 42188// CHECK11-NEXT:    [[TMP87:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS21]], i32 0, i32 22189// CHECK11-NEXT:    store ptr null, ptr [[TMP87]], align 42190// CHECK11-NEXT:    [[TMP88:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS19]], i32 0, i32 02191// CHECK11-NEXT:    [[TMP89:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS20]], i32 0, i32 02192// CHECK11-NEXT:    [[TMP90:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES22]], i32 0, i32 02193// CHECK11-NEXT:    [[TMP91:%.*]] = load i32, ptr [[N]], align 42194// CHECK11-NEXT:    store i32 [[TMP91]], ptr [[DOTCAPTURE_EXPR_24]], align 42195// CHECK11-NEXT:    [[TMP92:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_24]], align 42196// CHECK11-NEXT:    [[SUB26:%.*]] = sub nsw i32 [[TMP92]], 02197// CHECK11-NEXT:    [[DIV27:%.*]] = sdiv i32 [[SUB26]], 12198// CHECK11-NEXT:    [[SUB28:%.*]] = sub nsw i32 [[DIV27]], 12199// CHECK11-NEXT:    store i32 [[SUB28]], ptr [[DOTCAPTURE_EXPR_25]], align 42200// CHECK11-NEXT:    [[TMP93:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_25]], align 42201// CHECK11-NEXT:    [[ADD29:%.*]] = add nsw i32 [[TMP93]], 12202// CHECK11-NEXT:    [[TMP94:%.*]] = zext i32 [[ADD29]] to i642203// CHECK11-NEXT:    [[TMP95:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 02204// CHECK11-NEXT:    store i32 3, ptr [[TMP95]], align 42205// CHECK11-NEXT:    [[TMP96:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 12206// CHECK11-NEXT:    store i32 3, ptr [[TMP96]], align 42207// CHECK11-NEXT:    [[TMP97:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 22208// CHECK11-NEXT:    store ptr [[TMP88]], ptr [[TMP97]], align 42209// CHECK11-NEXT:    [[TMP98:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 32210// CHECK11-NEXT:    store ptr [[TMP89]], ptr [[TMP98]], align 42211// CHECK11-NEXT:    [[TMP99:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 42212// CHECK11-NEXT:    store ptr [[TMP90]], ptr [[TMP99]], align 42213// CHECK11-NEXT:    [[TMP100:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 52214// CHECK11-NEXT:    store ptr @.offload_maptypes.4, ptr [[TMP100]], align 42215// CHECK11-NEXT:    [[TMP101:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 62216// CHECK11-NEXT:    store ptr null, ptr [[TMP101]], align 42217// CHECK11-NEXT:    [[TMP102:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 72218// CHECK11-NEXT:    store ptr null, ptr [[TMP102]], align 42219// CHECK11-NEXT:    [[TMP103:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 82220// CHECK11-NEXT:    store i64 [[TMP94]], ptr [[TMP103]], align 82221// CHECK11-NEXT:    [[TMP104:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 92222// CHECK11-NEXT:    store i64 0, ptr [[TMP104]], align 82223// CHECK11-NEXT:    [[TMP105:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 102224// CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP105]], align 42225// CHECK11-NEXT:    [[TMP106:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 112226// CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP106]], align 42227// CHECK11-NEXT:    [[TMP107:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS30]], i32 0, i32 122228// CHECK11-NEXT:    store i32 0, ptr [[TMP107]], align 42229// CHECK11-NEXT:    [[TMP108:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l110.region_id, ptr [[KERNEL_ARGS30]])2230// CHECK11-NEXT:    [[TMP109:%.*]] = icmp ne i32 [[TMP108]], 02231// CHECK11-NEXT:    br i1 [[TMP109]], label [[OMP_OFFLOAD_FAILED31:%.*]], label [[OMP_OFFLOAD_CONT32:%.*]]2232// CHECK11:       omp_offload.failed31:2233// CHECK11-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l110(i32 [[TMP75]], i32 [[TMP0]], ptr [[VLA]]) #[[ATTR3]]2234// CHECK11-NEXT:    br label [[OMP_OFFLOAD_CONT32]]2235// CHECK11:       omp_offload.cont32:2236// CHECK11-NEXT:    [[TMP110:%.*]] = load i32, ptr [[ARGC_ADDR]], align 42237// CHECK11-NEXT:    [[CALL:%.*]] = call noundef i32 @_Z5tmainIiLi10EEiT_(i32 noundef [[TMP110]])2238// CHECK11-NEXT:    store i32 [[CALL]], ptr [[RETVAL]], align 42239// CHECK11-NEXT:    [[TMP111:%.*]] = load ptr, ptr [[SAVED_STACK]], align 42240// CHECK11-NEXT:    call void @llvm.stackrestore.p0(ptr [[TMP111]])2241// CHECK11-NEXT:    [[TMP112:%.*]] = load i32, ptr [[RETVAL]], align 42242// CHECK11-NEXT:    ret i32 [[TMP112]]2243//2244//2245// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l1002246// CHECK11-SAME: (i32 noundef [[N:%.*]], i32 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2:[0-9]+]] {2247// CHECK11-NEXT:  entry:2248// CHECK11-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 42249// CHECK11-NEXT:    [[VLA_ADDR:%.*]] = alloca i32, align 42250// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 42251// CHECK11-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 42252// CHECK11-NEXT:    store i32 [[VLA]], ptr [[VLA_ADDR]], align 42253// CHECK11-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 42254// CHECK11-NEXT:    [[TMP0:%.*]] = load i32, ptr [[VLA_ADDR]], align 42255// CHECK11-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 42256// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l100.omp_outlined, ptr [[N_ADDR]], i32 [[TMP0]], ptr [[TMP1]])2257// CHECK11-NEXT:    ret void2258//2259//2260// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l100.omp_outlined2261// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]], i32 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {2262// CHECK11-NEXT:  entry:2263// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42264// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42265// CHECK11-NEXT:    [[N_ADDR:%.*]] = alloca ptr, align 42266// CHECK11-NEXT:    [[VLA_ADDR:%.*]] = alloca i32, align 42267// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 42268// CHECK11-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 42269// CHECK11-NEXT:    [[TMP:%.*]] = alloca i32, align 42270// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42271// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42272// CHECK11-NEXT:    [[I:%.*]] = alloca i32, align 42273// CHECK11-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 42274// CHECK11-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 42275// CHECK11-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42276// CHECK11-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42277// CHECK11-NEXT:    [[I3:%.*]] = alloca i32, align 42278// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42279// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42280// CHECK11-NEXT:    store ptr [[N]], ptr [[N_ADDR]], align 42281// CHECK11-NEXT:    store i32 [[VLA]], ptr [[VLA_ADDR]], align 42282// CHECK11-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 42283// CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 42284// CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[VLA_ADDR]], align 42285// CHECK11-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 42286// CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 42287// CHECK11-NEXT:    store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR_]], align 42288// CHECK11-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42289// CHECK11-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP4]], 02290// CHECK11-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 12291// CHECK11-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12292// CHECK11-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42293// CHECK11-NEXT:    store i32 0, ptr [[I]], align 42294// CHECK11-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42295// CHECK11-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP5]]2296// CHECK11-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2297// CHECK11:       omp.precond.then:2298// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 42299// CHECK11-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42300// CHECK11-NEXT:    store i32 [[TMP6]], ptr [[DOTOMP_UB]], align 42301// CHECK11-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 42302// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 42303// CHECK11-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42304// CHECK11-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 42305// CHECK11-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP8]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2306// CHECK11-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42307// CHECK11-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42308// CHECK11-NEXT:    [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]2309// CHECK11-NEXT:    br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2310// CHECK11:       cond.true:2311// CHECK11-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42312// CHECK11-NEXT:    br label [[COND_END:%.*]]2313// CHECK11:       cond.false:2314// CHECK11-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42315// CHECK11-NEXT:    br label [[COND_END]]2316// CHECK11:       cond.end:2317// CHECK11-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]2318// CHECK11-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 42319// CHECK11-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42320// CHECK11-NEXT:    store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 42321// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]2322// CHECK11:       omp.inner.for.cond:2323// CHECK11-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42324// CHECK11-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42325// CHECK11-NEXT:    [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]2326// CHECK11-NEXT:    br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2327// CHECK11:       omp.inner.for.body:2328// CHECK11-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42329// CHECK11-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP16]], 12330// CHECK11-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]2331// CHECK11-NEXT:    store i32 [[ADD]], ptr [[I3]], align 42332// CHECK11-NEXT:    [[TMP17:%.*]] = load i32, ptr [[I3]], align 42333// CHECK11-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i32 [[TMP17]]2334// CHECK11-NEXT:    store i32 0, ptr [[ARRAYIDX]], align 42335// CHECK11-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]2336// CHECK11:       omp.body.continue:2337// CHECK11-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]2338// CHECK11:       omp.inner.for.inc:2339// CHECK11-NEXT:    [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42340// CHECK11-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP18]], 12341// CHECK11-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 42342// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND]]2343// CHECK11:       omp.inner.for.end:2344// CHECK11-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]2345// CHECK11:       omp.loop.exit:2346// CHECK11-NEXT:    [[TMP19:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42347// CHECK11-NEXT:    [[TMP20:%.*]] = load i32, ptr [[TMP19]], align 42348// CHECK11-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP20]])2349// CHECK11-NEXT:    br label [[OMP_PRECOND_END]]2350// CHECK11:       omp.precond.end:2351// CHECK11-NEXT:    ret void2352//2353//2354// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l1052355// CHECK11-SAME: (i32 noundef [[N:%.*]], i32 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {2356// CHECK11-NEXT:  entry:2357// CHECK11-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 42358// CHECK11-NEXT:    [[VLA_ADDR:%.*]] = alloca i32, align 42359// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 42360// CHECK11-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 42361// CHECK11-NEXT:    store i32 [[VLA]], ptr [[VLA_ADDR]], align 42362// CHECK11-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 42363// CHECK11-NEXT:    [[TMP0:%.*]] = load i32, ptr [[VLA_ADDR]], align 42364// CHECK11-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 42365// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l105.omp_outlined, ptr [[N_ADDR]], i32 [[TMP0]], ptr [[TMP1]])2366// CHECK11-NEXT:    ret void2367//2368//2369// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l105.omp_outlined2370// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]], i32 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {2371// CHECK11-NEXT:  entry:2372// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42373// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42374// CHECK11-NEXT:    [[N_ADDR:%.*]] = alloca ptr, align 42375// CHECK11-NEXT:    [[VLA_ADDR:%.*]] = alloca i32, align 42376// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 42377// CHECK11-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 42378// CHECK11-NEXT:    [[TMP:%.*]] = alloca i32, align 42379// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42380// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42381// CHECK11-NEXT:    [[I:%.*]] = alloca i32, align 42382// CHECK11-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 42383// CHECK11-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 42384// CHECK11-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42385// CHECK11-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42386// CHECK11-NEXT:    [[I3:%.*]] = alloca i32, align 42387// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42388// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42389// CHECK11-NEXT:    store ptr [[N]], ptr [[N_ADDR]], align 42390// CHECK11-NEXT:    store i32 [[VLA]], ptr [[VLA_ADDR]], align 42391// CHECK11-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 42392// CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 42393// CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[VLA_ADDR]], align 42394// CHECK11-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 42395// CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 42396// CHECK11-NEXT:    store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR_]], align 42397// CHECK11-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42398// CHECK11-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP4]], 02399// CHECK11-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 12400// CHECK11-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12401// CHECK11-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42402// CHECK11-NEXT:    store i32 0, ptr [[I]], align 42403// CHECK11-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42404// CHECK11-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP5]]2405// CHECK11-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2406// CHECK11:       omp.precond.then:2407// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 42408// CHECK11-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42409// CHECK11-NEXT:    store i32 [[TMP6]], ptr [[DOTOMP_UB]], align 42410// CHECK11-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 42411// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 42412// CHECK11-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42413// CHECK11-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 42414// CHECK11-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP8]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2415// CHECK11-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42416// CHECK11-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42417// CHECK11-NEXT:    [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]2418// CHECK11-NEXT:    br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2419// CHECK11:       cond.true:2420// CHECK11-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42421// CHECK11-NEXT:    br label [[COND_END:%.*]]2422// CHECK11:       cond.false:2423// CHECK11-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42424// CHECK11-NEXT:    br label [[COND_END]]2425// CHECK11:       cond.end:2426// CHECK11-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]2427// CHECK11-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 42428// CHECK11-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42429// CHECK11-NEXT:    store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 42430// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]2431// CHECK11:       omp.inner.for.cond:2432// CHECK11-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42433// CHECK11-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42434// CHECK11-NEXT:    [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]2435// CHECK11-NEXT:    br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2436// CHECK11:       omp.inner.for.body:2437// CHECK11-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42438// CHECK11-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP16]], 12439// CHECK11-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]2440// CHECK11-NEXT:    store i32 [[ADD]], ptr [[I3]], align 42441// CHECK11-NEXT:    [[TMP17:%.*]] = load i32, ptr [[I3]], align 42442// CHECK11-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i32 [[TMP17]]2443// CHECK11-NEXT:    store i32 0, ptr [[ARRAYIDX]], align 42444// CHECK11-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]2445// CHECK11:       omp.body.continue:2446// CHECK11-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]2447// CHECK11:       omp.inner.for.inc:2448// CHECK11-NEXT:    [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42449// CHECK11-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP18]], 12450// CHECK11-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 42451// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND]]2452// CHECK11:       omp.inner.for.end:2453// CHECK11-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]2454// CHECK11:       omp.loop.exit:2455// CHECK11-NEXT:    [[TMP19:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42456// CHECK11-NEXT:    [[TMP20:%.*]] = load i32, ptr [[TMP19]], align 42457// CHECK11-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP20]])2458// CHECK11-NEXT:    br label [[OMP_PRECOND_END]]2459// CHECK11:       omp.precond.end:2460// CHECK11-NEXT:    ret void2461//2462//2463// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l1102464// CHECK11-SAME: (i32 noundef [[N:%.*]], i32 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {2465// CHECK11-NEXT:  entry:2466// CHECK11-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 42467// CHECK11-NEXT:    [[VLA_ADDR:%.*]] = alloca i32, align 42468// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 42469// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42470// CHECK11-NEXT:    [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i32, align 42471// CHECK11-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 42472// CHECK11-NEXT:    store i32 [[VLA]], ptr [[VLA_ADDR]], align 42473// CHECK11-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 42474// CHECK11-NEXT:    [[TMP0:%.*]] = load i32, ptr [[VLA_ADDR]], align 42475// CHECK11-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 42476// CHECK11-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 42477// CHECK11-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 42478// CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42479// CHECK11-NEXT:    store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 42480// CHECK11-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED]], align 42481// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 4, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l110.omp_outlined, ptr [[N_ADDR]], i32 [[TMP0]], ptr [[TMP1]], i32 [[TMP4]])2482// CHECK11-NEXT:    ret void2483//2484//2485// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l110.omp_outlined2486// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]], i32 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR2]] {2487// CHECK11-NEXT:  entry:2488// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42489// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42490// CHECK11-NEXT:    [[N_ADDR:%.*]] = alloca ptr, align 42491// CHECK11-NEXT:    [[VLA_ADDR:%.*]] = alloca i32, align 42492// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 42493// CHECK11-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 42494// CHECK11-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 42495// CHECK11-NEXT:    [[TMP:%.*]] = alloca i32, align 42496// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42497// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_2:%.*]] = alloca i32, align 42498// CHECK11-NEXT:    [[I:%.*]] = alloca i32, align 42499// CHECK11-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 42500// CHECK11-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 42501// CHECK11-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42502// CHECK11-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42503// CHECK11-NEXT:    [[I4:%.*]] = alloca i32, align 42504// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42505// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42506// CHECK11-NEXT:    store ptr [[N]], ptr [[N_ADDR]], align 42507// CHECK11-NEXT:    store i32 [[VLA]], ptr [[VLA_ADDR]], align 42508// CHECK11-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 42509// CHECK11-NEXT:    store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 42510// CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 42511// CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[VLA_ADDR]], align 42512// CHECK11-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 42513// CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 42514// CHECK11-NEXT:    store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR_1]], align 42515// CHECK11-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42516// CHECK11-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP4]], 02517// CHECK11-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 12518// CHECK11-NEXT:    [[SUB3:%.*]] = sub nsw i32 [[DIV]], 12519// CHECK11-NEXT:    store i32 [[SUB3]], ptr [[DOTCAPTURE_EXPR_2]], align 42520// CHECK11-NEXT:    store i32 0, ptr [[I]], align 42521// CHECK11-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42522// CHECK11-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP5]]2523// CHECK11-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2524// CHECK11:       omp.precond.then:2525// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 42526// CHECK11-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 42527// CHECK11-NEXT:    store i32 [[TMP6]], ptr [[DOTOMP_UB]], align 42528// CHECK11-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 42529// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 42530// CHECK11-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 42531// CHECK11-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42532// CHECK11-NEXT:    [[TMP9:%.*]] = load i32, ptr [[TMP8]], align 42533// CHECK11-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP9]], i32 91, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 [[TMP7]])2534// CHECK11-NEXT:    br label [[OMP_DISPATCH_COND:%.*]]2535// CHECK11:       omp.dispatch.cond:2536// CHECK11-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42537// CHECK11-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 42538// CHECK11-NEXT:    [[CMP5:%.*]] = icmp sgt i32 [[TMP10]], [[TMP11]]2539// CHECK11-NEXT:    br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2540// CHECK11:       cond.true:2541// CHECK11-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 42542// CHECK11-NEXT:    br label [[COND_END:%.*]]2543// CHECK11:       cond.false:2544// CHECK11-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42545// CHECK11-NEXT:    br label [[COND_END]]2546// CHECK11:       cond.end:2547// CHECK11-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP12]], [[COND_TRUE]] ], [ [[TMP13]], [[COND_FALSE]] ]2548// CHECK11-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 42549// CHECK11-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42550// CHECK11-NEXT:    store i32 [[TMP14]], ptr [[DOTOMP_IV]], align 42551// CHECK11-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42552// CHECK11-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42553// CHECK11-NEXT:    [[CMP6:%.*]] = icmp sle i32 [[TMP15]], [[TMP16]]2554// CHECK11-NEXT:    br i1 [[CMP6]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]2555// CHECK11:       omp.dispatch.body:2556// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]2557// CHECK11:       omp.inner.for.cond:2558// CHECK11-NEXT:    [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12:![0-9]+]]2559// CHECK11-NEXT:    [[TMP18:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP12]]2560// CHECK11-NEXT:    [[CMP7:%.*]] = icmp sle i32 [[TMP17]], [[TMP18]]2561// CHECK11-NEXT:    br i1 [[CMP7]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2562// CHECK11:       omp.inner.for.body:2563// CHECK11-NEXT:    [[TMP19:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]2564// CHECK11-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP19]], 12565// CHECK11-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]2566// CHECK11-NEXT:    store i32 [[ADD]], ptr [[I4]], align 4, !llvm.access.group [[ACC_GRP12]]2567// CHECK11-NEXT:    [[TMP20:%.*]] = load i32, ptr [[I4]], align 4, !llvm.access.group [[ACC_GRP12]]2568// CHECK11-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i32 [[TMP20]]2569// CHECK11-NEXT:    store i32 0, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP12]]2570// CHECK11-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]2571// CHECK11:       omp.body.continue:2572// CHECK11-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]2573// CHECK11:       omp.inner.for.inc:2574// CHECK11-NEXT:    [[TMP21:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]2575// CHECK11-NEXT:    [[ADD8:%.*]] = add nsw i32 [[TMP21]], 12576// CHECK11-NEXT:    store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]2577// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP13:![0-9]+]]2578// CHECK11:       omp.inner.for.end:2579// CHECK11-NEXT:    br label [[OMP_DISPATCH_INC:%.*]]2580// CHECK11:       omp.dispatch.inc:2581// CHECK11-NEXT:    [[TMP22:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42582// CHECK11-NEXT:    [[TMP23:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 42583// CHECK11-NEXT:    [[ADD9:%.*]] = add nsw i32 [[TMP22]], [[TMP23]]2584// CHECK11-NEXT:    store i32 [[ADD9]], ptr [[DOTOMP_LB]], align 42585// CHECK11-NEXT:    [[TMP24:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42586// CHECK11-NEXT:    [[TMP25:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 42587// CHECK11-NEXT:    [[ADD10:%.*]] = add nsw i32 [[TMP24]], [[TMP25]]2588// CHECK11-NEXT:    store i32 [[ADD10]], ptr [[DOTOMP_UB]], align 42589// CHECK11-NEXT:    br label [[OMP_DISPATCH_COND]]2590// CHECK11:       omp.dispatch.end:2591// CHECK11-NEXT:    [[TMP26:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42592// CHECK11-NEXT:    [[TMP27:%.*]] = load i32, ptr [[TMP26]], align 42593// CHECK11-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP27]])2594// CHECK11-NEXT:    br label [[OMP_PRECOND_END]]2595// CHECK11:       omp.precond.end:2596// CHECK11-NEXT:    ret void2597//2598//2599// CHECK11-LABEL: define {{[^@]+}}@_Z5tmainIiLi10EEiT_2600// CHECK11-SAME: (i32 noundef [[ARGC:%.*]]) #[[ATTR5:[0-9]+]] comdat {2601// CHECK11-NEXT:  entry:2602// CHECK11-NEXT:    [[ARGC_ADDR:%.*]] = alloca i32, align 42603// CHECK11-NEXT:    [[A:%.*]] = alloca [10 x i32], align 42604// CHECK11-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 42605// CHECK11-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 42606// CHECK11-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 42607// CHECK11-NEXT:    [[TMP:%.*]] = alloca i32, align 42608// CHECK11-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 82609// CHECK11-NEXT:    [[DOTOFFLOAD_BASEPTRS1:%.*]] = alloca [1 x ptr], align 42610// CHECK11-NEXT:    [[DOTOFFLOAD_PTRS2:%.*]] = alloca [1 x ptr], align 42611// CHECK11-NEXT:    [[DOTOFFLOAD_MAPPERS3:%.*]] = alloca [1 x ptr], align 42612// CHECK11-NEXT:    [[_TMP4:%.*]] = alloca i32, align 42613// CHECK11-NEXT:    [[KERNEL_ARGS5:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 82614// CHECK11-NEXT:    [[DOTOFFLOAD_BASEPTRS8:%.*]] = alloca [1 x ptr], align 42615// CHECK11-NEXT:    [[DOTOFFLOAD_PTRS9:%.*]] = alloca [1 x ptr], align 42616// CHECK11-NEXT:    [[DOTOFFLOAD_MAPPERS10:%.*]] = alloca [1 x ptr], align 42617// CHECK11-NEXT:    [[_TMP11:%.*]] = alloca i32, align 42618// CHECK11-NEXT:    [[KERNEL_ARGS12:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 82619// CHECK11-NEXT:    store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 42620// CHECK11-NEXT:    [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 02621// CHECK11-NEXT:    store ptr [[A]], ptr [[TMP0]], align 42622// CHECK11-NEXT:    [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 02623// CHECK11-NEXT:    store ptr [[A]], ptr [[TMP1]], align 42624// CHECK11-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 02625// CHECK11-NEXT:    store ptr null, ptr [[TMP2]], align 42626// CHECK11-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 02627// CHECK11-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 02628// CHECK11-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 02629// CHECK11-NEXT:    store i32 3, ptr [[TMP5]], align 42630// CHECK11-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12631// CHECK11-NEXT:    store i32 1, ptr [[TMP6]], align 42632// CHECK11-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 22633// CHECK11-NEXT:    store ptr [[TMP3]], ptr [[TMP7]], align 42634// CHECK11-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 32635// CHECK11-NEXT:    store ptr [[TMP4]], ptr [[TMP8]], align 42636// CHECK11-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 42637// CHECK11-NEXT:    store ptr @.offload_sizes.5, ptr [[TMP9]], align 42638// CHECK11-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 52639// CHECK11-NEXT:    store ptr @.offload_maptypes.6, ptr [[TMP10]], align 42640// CHECK11-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 62641// CHECK11-NEXT:    store ptr null, ptr [[TMP11]], align 42642// CHECK11-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 72643// CHECK11-NEXT:    store ptr null, ptr [[TMP12]], align 42644// CHECK11-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 82645// CHECK11-NEXT:    store i64 10, ptr [[TMP13]], align 82646// CHECK11-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 92647// CHECK11-NEXT:    store i64 0, ptr [[TMP14]], align 82648// CHECK11-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 102649// CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP15]], align 42650// CHECK11-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 112651// CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP16]], align 42652// CHECK11-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 122653// CHECK11-NEXT:    store i32 0, ptr [[TMP17]], align 42654// CHECK11-NEXT:    [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l79.region_id, ptr [[KERNEL_ARGS]])2655// CHECK11-NEXT:    [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 02656// CHECK11-NEXT:    br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]2657// CHECK11:       omp_offload.failed:2658// CHECK11-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l79(ptr [[A]]) #[[ATTR3]]2659// CHECK11-NEXT:    br label [[OMP_OFFLOAD_CONT]]2660// CHECK11:       omp_offload.cont:2661// CHECK11-NEXT:    [[TMP20:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 02662// CHECK11-NEXT:    store ptr [[A]], ptr [[TMP20]], align 42663// CHECK11-NEXT:    [[TMP21:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 02664// CHECK11-NEXT:    store ptr [[A]], ptr [[TMP21]], align 42665// CHECK11-NEXT:    [[TMP22:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS3]], i32 0, i32 02666// CHECK11-NEXT:    store ptr null, ptr [[TMP22]], align 42667// CHECK11-NEXT:    [[TMP23:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 02668// CHECK11-NEXT:    [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 02669// CHECK11-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 02670// CHECK11-NEXT:    store i32 3, ptr [[TMP25]], align 42671// CHECK11-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 12672// CHECK11-NEXT:    store i32 1, ptr [[TMP26]], align 42673// CHECK11-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 22674// CHECK11-NEXT:    store ptr [[TMP23]], ptr [[TMP27]], align 42675// CHECK11-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 32676// CHECK11-NEXT:    store ptr [[TMP24]], ptr [[TMP28]], align 42677// CHECK11-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 42678// CHECK11-NEXT:    store ptr @.offload_sizes.7, ptr [[TMP29]], align 42679// CHECK11-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 52680// CHECK11-NEXT:    store ptr @.offload_maptypes.8, ptr [[TMP30]], align 42681// CHECK11-NEXT:    [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 62682// CHECK11-NEXT:    store ptr null, ptr [[TMP31]], align 42683// CHECK11-NEXT:    [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 72684// CHECK11-NEXT:    store ptr null, ptr [[TMP32]], align 42685// CHECK11-NEXT:    [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 82686// CHECK11-NEXT:    store i64 10, ptr [[TMP33]], align 82687// CHECK11-NEXT:    [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 92688// CHECK11-NEXT:    store i64 0, ptr [[TMP34]], align 82689// CHECK11-NEXT:    [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 102690// CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP35]], align 42691// CHECK11-NEXT:    [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 112692// CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP36]], align 42693// CHECK11-NEXT:    [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 122694// CHECK11-NEXT:    store i32 0, ptr [[TMP37]], align 42695// CHECK11-NEXT:    [[TMP38:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l84.region_id, ptr [[KERNEL_ARGS5]])2696// CHECK11-NEXT:    [[TMP39:%.*]] = icmp ne i32 [[TMP38]], 02697// CHECK11-NEXT:    br i1 [[TMP39]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]]2698// CHECK11:       omp_offload.failed6:2699// CHECK11-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l84(ptr [[A]]) #[[ATTR3]]2700// CHECK11-NEXT:    br label [[OMP_OFFLOAD_CONT7]]2701// CHECK11:       omp_offload.cont7:2702// CHECK11-NEXT:    [[TMP40:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 02703// CHECK11-NEXT:    store ptr [[A]], ptr [[TMP40]], align 42704// CHECK11-NEXT:    [[TMP41:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 02705// CHECK11-NEXT:    store ptr [[A]], ptr [[TMP41]], align 42706// CHECK11-NEXT:    [[TMP42:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS10]], i32 0, i32 02707// CHECK11-NEXT:    store ptr null, ptr [[TMP42]], align 42708// CHECK11-NEXT:    [[TMP43:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 02709// CHECK11-NEXT:    [[TMP44:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 02710// CHECK11-NEXT:    [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 02711// CHECK11-NEXT:    store i32 3, ptr [[TMP45]], align 42712// CHECK11-NEXT:    [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 12713// CHECK11-NEXT:    store i32 1, ptr [[TMP46]], align 42714// CHECK11-NEXT:    [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 22715// CHECK11-NEXT:    store ptr [[TMP43]], ptr [[TMP47]], align 42716// CHECK11-NEXT:    [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 32717// CHECK11-NEXT:    store ptr [[TMP44]], ptr [[TMP48]], align 42718// CHECK11-NEXT:    [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 42719// CHECK11-NEXT:    store ptr @.offload_sizes.9, ptr [[TMP49]], align 42720// CHECK11-NEXT:    [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 52721// CHECK11-NEXT:    store ptr @.offload_maptypes.10, ptr [[TMP50]], align 42722// CHECK11-NEXT:    [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 62723// CHECK11-NEXT:    store ptr null, ptr [[TMP51]], align 42724// CHECK11-NEXT:    [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 72725// CHECK11-NEXT:    store ptr null, ptr [[TMP52]], align 42726// CHECK11-NEXT:    [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 82727// CHECK11-NEXT:    store i64 10, ptr [[TMP53]], align 82728// CHECK11-NEXT:    [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 92729// CHECK11-NEXT:    store i64 0, ptr [[TMP54]], align 82730// CHECK11-NEXT:    [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 102731// CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP55]], align 42732// CHECK11-NEXT:    [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 112733// CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP56]], align 42734// CHECK11-NEXT:    [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 122735// CHECK11-NEXT:    store i32 0, ptr [[TMP57]], align 42736// CHECK11-NEXT:    [[TMP58:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l89.region_id, ptr [[KERNEL_ARGS12]])2737// CHECK11-NEXT:    [[TMP59:%.*]] = icmp ne i32 [[TMP58]], 02738// CHECK11-NEXT:    br i1 [[TMP59]], label [[OMP_OFFLOAD_FAILED13:%.*]], label [[OMP_OFFLOAD_CONT14:%.*]]2739// CHECK11:       omp_offload.failed13:2740// CHECK11-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l89(ptr [[A]]) #[[ATTR3]]2741// CHECK11-NEXT:    br label [[OMP_OFFLOAD_CONT14]]2742// CHECK11:       omp_offload.cont14:2743// CHECK11-NEXT:    ret i32 02744//2745//2746// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l792747// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {2748// CHECK11-NEXT:  entry:2749// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 42750// CHECK11-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 42751// CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42752// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l79.omp_outlined, ptr [[TMP0]])2753// CHECK11-NEXT:    ret void2754//2755//2756// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l79.omp_outlined2757// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {2758// CHECK11-NEXT:  entry:2759// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42760// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42761// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 42762// CHECK11-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 42763// CHECK11-NEXT:    [[TMP:%.*]] = alloca i32, align 42764// CHECK11-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 42765// CHECK11-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 42766// CHECK11-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42767// CHECK11-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42768// CHECK11-NEXT:    [[I:%.*]] = alloca i32, align 42769// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42770// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42771// CHECK11-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 42772// CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42773// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 42774// CHECK11-NEXT:    store i32 9, ptr [[DOTOMP_UB]], align 42775// CHECK11-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 42776// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 42777// CHECK11-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42778// CHECK11-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 42779// CHECK11-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2780// CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42781// CHECK11-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 92782// CHECK11-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2783// CHECK11:       cond.true:2784// CHECK11-NEXT:    br label [[COND_END:%.*]]2785// CHECK11:       cond.false:2786// CHECK11-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42787// CHECK11-NEXT:    br label [[COND_END]]2788// CHECK11:       cond.end:2789// CHECK11-NEXT:    [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]2790// CHECK11-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 42791// CHECK11-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42792// CHECK11-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 42793// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]2794// CHECK11:       omp.inner.for.cond:2795// CHECK11-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42796// CHECK11-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42797// CHECK11-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]2798// CHECK11-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2799// CHECK11:       omp.inner.for.body:2800// CHECK11-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42801// CHECK11-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP8]], 12802// CHECK11-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]2803// CHECK11-NEXT:    store i32 [[ADD]], ptr [[I]], align 42804// CHECK11-NEXT:    [[TMP9:%.*]] = load i32, ptr [[I]], align 42805// CHECK11-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP9]]2806// CHECK11-NEXT:    store i32 0, ptr [[ARRAYIDX]], align 42807// CHECK11-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]2808// CHECK11:       omp.body.continue:2809// CHECK11-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]2810// CHECK11:       omp.inner.for.inc:2811// CHECK11-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42812// CHECK11-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP10]], 12813// CHECK11-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 42814// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND]]2815// CHECK11:       omp.inner.for.end:2816// CHECK11-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]2817// CHECK11:       omp.loop.exit:2818// CHECK11-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])2819// CHECK11-NEXT:    ret void2820//2821//2822// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l842823// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {2824// CHECK11-NEXT:  entry:2825// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 42826// CHECK11-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 42827// CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42828// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l84.omp_outlined, ptr [[TMP0]])2829// CHECK11-NEXT:    ret void2830//2831//2832// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l84.omp_outlined2833// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {2834// CHECK11-NEXT:  entry:2835// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42836// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42837// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 42838// CHECK11-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 42839// CHECK11-NEXT:    [[TMP:%.*]] = alloca i32, align 42840// CHECK11-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 42841// CHECK11-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 42842// CHECK11-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42843// CHECK11-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42844// CHECK11-NEXT:    [[I:%.*]] = alloca i32, align 42845// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42846// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42847// CHECK11-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 42848// CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42849// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 42850// CHECK11-NEXT:    store i32 9, ptr [[DOTOMP_UB]], align 42851// CHECK11-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 42852// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 42853// CHECK11-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42854// CHECK11-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 42855// CHECK11-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2856// CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42857// CHECK11-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 92858// CHECK11-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2859// CHECK11:       cond.true:2860// CHECK11-NEXT:    br label [[COND_END:%.*]]2861// CHECK11:       cond.false:2862// CHECK11-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42863// CHECK11-NEXT:    br label [[COND_END]]2864// CHECK11:       cond.end:2865// CHECK11-NEXT:    [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]2866// CHECK11-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 42867// CHECK11-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42868// CHECK11-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 42869// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]2870// CHECK11:       omp.inner.for.cond:2871// CHECK11-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42872// CHECK11-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42873// CHECK11-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]2874// CHECK11-NEXT:    br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2875// CHECK11:       omp.inner.for.body:2876// CHECK11-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42877// CHECK11-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP8]], 12878// CHECK11-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]2879// CHECK11-NEXT:    store i32 [[ADD]], ptr [[I]], align 42880// CHECK11-NEXT:    [[TMP9:%.*]] = load i32, ptr [[I]], align 42881// CHECK11-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP9]]2882// CHECK11-NEXT:    store i32 0, ptr [[ARRAYIDX]], align 42883// CHECK11-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]2884// CHECK11:       omp.body.continue:2885// CHECK11-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]2886// CHECK11:       omp.inner.for.inc:2887// CHECK11-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42888// CHECK11-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP10]], 12889// CHECK11-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 42890// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND]]2891// CHECK11:       omp.inner.for.end:2892// CHECK11-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]2893// CHECK11:       omp.loop.exit:2894// CHECK11-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])2895// CHECK11-NEXT:    ret void2896//2897//2898// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l892899// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {2900// CHECK11-NEXT:  entry:2901// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 42902// CHECK11-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 42903// CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42904// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l89.omp_outlined, ptr [[TMP0]])2905// CHECK11-NEXT:    ret void2906//2907//2908// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l89.omp_outlined2909// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {2910// CHECK11-NEXT:  entry:2911// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42912// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42913// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 42914// CHECK11-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 42915// CHECK11-NEXT:    [[TMP:%.*]] = alloca i32, align 42916// CHECK11-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 42917// CHECK11-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 42918// CHECK11-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42919// CHECK11-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42920// CHECK11-NEXT:    [[I:%.*]] = alloca i32, align 42921// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42922// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42923// CHECK11-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 42924// CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42925// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 42926// CHECK11-NEXT:    store i32 9, ptr [[DOTOMP_UB]], align 42927// CHECK11-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 42928// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 42929// CHECK11-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42930// CHECK11-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 42931// CHECK11-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 91, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 10)2932// CHECK11-NEXT:    br label [[OMP_DISPATCH_COND:%.*]]2933// CHECK11:       omp.dispatch.cond:2934// CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42935// CHECK11-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 92936// CHECK11-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2937// CHECK11:       cond.true:2938// CHECK11-NEXT:    br label [[COND_END:%.*]]2939// CHECK11:       cond.false:2940// CHECK11-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42941// CHECK11-NEXT:    br label [[COND_END]]2942// CHECK11:       cond.end:2943// CHECK11-NEXT:    [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]2944// CHECK11-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 42945// CHECK11-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42946// CHECK11-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 42947// CHECK11-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42948// CHECK11-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42949// CHECK11-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]2950// CHECK11-NEXT:    br i1 [[CMP1]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]2951// CHECK11:       omp.dispatch.body:2952// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]2953// CHECK11:       omp.inner.for.cond:2954// CHECK11-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP15:![0-9]+]]2955// CHECK11-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP15]]2956// CHECK11-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP8]], [[TMP9]]2957// CHECK11-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2958// CHECK11:       omp.inner.for.body:2959// CHECK11-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP15]]2960// CHECK11-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP10]], 12961// CHECK11-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]2962// CHECK11-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP15]]2963// CHECK11-NEXT:    [[TMP11:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP15]]2964// CHECK11-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP11]]2965// CHECK11-NEXT:    store i32 0, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP15]]2966// CHECK11-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]2967// CHECK11:       omp.body.continue:2968// CHECK11-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]2969// CHECK11:       omp.inner.for.inc:2970// CHECK11-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP15]]2971// CHECK11-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP12]], 12972// CHECK11-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP15]]2973// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP16:![0-9]+]]2974// CHECK11:       omp.inner.for.end:2975// CHECK11-NEXT:    br label [[OMP_DISPATCH_INC:%.*]]2976// CHECK11:       omp.dispatch.inc:2977// CHECK11-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42978// CHECK11-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 42979// CHECK11-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP13]], [[TMP14]]2980// CHECK11-NEXT:    store i32 [[ADD4]], ptr [[DOTOMP_LB]], align 42981// CHECK11-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42982// CHECK11-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 42983// CHECK11-NEXT:    [[ADD5:%.*]] = add nsw i32 [[TMP15]], [[TMP16]]2984// CHECK11-NEXT:    store i32 [[ADD5]], ptr [[DOTOMP_UB]], align 42985// CHECK11-NEXT:    br label [[OMP_DISPATCH_COND]]2986// CHECK11:       omp.dispatch.end:2987// CHECK11-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])2988// CHECK11-NEXT:    ret void2989//2990