3023 lines · cpp
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _2// expected-no-diagnostics3#ifndef HEADER4#define HEADER5 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 target teams distribute29 for(int i = 0; i < X; i++) {30 a[i] = (T)0;31 }32 #pragma omp target teams distribute dist_schedule(static)33 for(int i = 0; i < X; i++) {34 a[i] = (T)0;35 }36 #pragma omp target teams distribute dist_schedule(static, X/2)37 for(int i = 0; i < X; i++) {38 a[i] = (T)0;39 }40 41 42 43 44 45 46 return a[0];47 }48};49 50int teams_template_struct(void) {51 SS<int, 123, 456> V;52 return V.foo();53 54}55#endif // CK156 57// Test host codegen.58// 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=CHECK959// RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s60// 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=CHECK961// 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=CHECK1162// RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s63// 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=CHECK1164 65// 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}}"66// 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 %s67// 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}}"68// 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}}"69// 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 %s70// 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}}"71#ifdef CK272 73template <typename T, int n>74int tmain(T argc) {75 T a[n];76#pragma omp target teams distribute77 for(int i = 0; i < n; i++) {78 a[i] = (T)0;79 }80#pragma omp target teams distribute dist_schedule(static)81 for(int i = 0; i < n; i++) {82 a[i] = (T)0;83 }84#pragma omp target teams distribute dist_schedule(static, n)85 for(int i = 0; i < n; i++) {86 a[i] = (T)0;87 }88 return 0;89}90 91int main (int argc, char **argv) {92 int n = 100;93 int a[n];94#pragma omp target teams distribute95 for(int i = 0; i < n; i++) {96 a[i] = 0;97 }98#pragma omp target teams distribute dist_schedule(static)99 for(int i = 0; i < n; i++) {100 a[i] = 0;101 }102#pragma omp target teams distribute dist_schedule(static, n)103 for(int i = 0; i < n; i++) {104 a[i] = 0;105 }106 return tmain<int, 10>(argc);107}108 109 110 111 112 113 114 115 116 117 118 119 120 121 122 123#endif // CK2124#endif // #ifndef HEADER125// CHECK1-LABEL: define {{[^@]+}}@_Z21teams_template_structv126// CHECK1-SAME: () #[[ATTR0:[0-9]+]] {127// CHECK1-NEXT: entry:128// CHECK1-NEXT: [[V:%.*]] = alloca [[STRUCT_SS:%.*]], align 4129// CHECK1-NEXT: [[CALL:%.*]] = call noundef signext i32 @_ZN2SSIiLi123ELx456EE3fooEv(ptr noundef nonnull align 4 dereferenceable(496) [[V]])130// CHECK1-NEXT: ret i32 [[CALL]]131//132//133// CHECK1-LABEL: define {{[^@]+}}@_ZN2SSIiLi123ELx456EE3fooEv134// CHECK1-SAME: (ptr noundef nonnull align 4 dereferenceable(496) [[THIS:%.*]]) #[[ATTR0]] comdat {135// CHECK1-NEXT: entry:136// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8137// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8138// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8139// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8140// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4141// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8142// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS3:%.*]] = alloca [1 x ptr], align 8143// CHECK1-NEXT: [[DOTOFFLOAD_PTRS4:%.*]] = alloca [1 x ptr], align 8144// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS5:%.*]] = alloca [1 x ptr], align 8145// CHECK1-NEXT: [[_TMP6:%.*]] = alloca i32, align 4146// CHECK1-NEXT: [[KERNEL_ARGS7:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8147// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS11:%.*]] = alloca [1 x ptr], align 8148// CHECK1-NEXT: [[DOTOFFLOAD_PTRS12:%.*]] = alloca [1 x ptr], align 8149// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS13:%.*]] = alloca [1 x ptr], align 8150// CHECK1-NEXT: [[_TMP14:%.*]] = alloca i32, align 4151// CHECK1-NEXT: [[KERNEL_ARGS15:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8152// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8153// CHECK1-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8154// CHECK1-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[THIS1]], i32 0, i32 0155// CHECK1-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0156// CHECK1-NEXT: store ptr [[THIS1]], ptr [[TMP0]], align 8157// CHECK1-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0158// CHECK1-NEXT: store ptr [[A]], ptr [[TMP1]], align 8159// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0160// CHECK1-NEXT: store ptr null, ptr [[TMP2]], align 8161// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0162// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0163// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0164// CHECK1-NEXT: store i32 3, ptr [[TMP5]], align 4165// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1166// CHECK1-NEXT: store i32 1, ptr [[TMP6]], align 4167// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2168// CHECK1-NEXT: store ptr [[TMP3]], ptr [[TMP7]], align 8169// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3170// CHECK1-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 8171// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4172// CHECK1-NEXT: store ptr @.offload_sizes, ptr [[TMP9]], align 8173// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5174// CHECK1-NEXT: store ptr @.offload_maptypes, ptr [[TMP10]], align 8175// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6176// CHECK1-NEXT: store ptr null, ptr [[TMP11]], align 8177// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7178// CHECK1-NEXT: store ptr null, ptr [[TMP12]], align 8179// CHECK1-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8180// CHECK1-NEXT: store i64 123, ptr [[TMP13]], align 8181// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9182// CHECK1-NEXT: store i64 0, ptr [[TMP14]], align 8183// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10184// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4185// CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11186// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4187// CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12188// CHECK1-NEXT: store i32 0, ptr [[TMP17]], align 4189// 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]])190// CHECK1-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0191// CHECK1-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]192// CHECK1: omp_offload.failed:193// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28(ptr [[THIS1]]) #[[ATTR2:[0-9]+]]194// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]195// CHECK1: omp_offload.cont:196// CHECK1-NEXT: [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 0197// CHECK1-NEXT: [[TMP20:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0198// CHECK1-NEXT: store ptr [[THIS1]], ptr [[TMP20]], align 8199// CHECK1-NEXT: [[TMP21:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 0200// CHECK1-NEXT: store ptr [[A2]], ptr [[TMP21]], align 8201// CHECK1-NEXT: [[TMP22:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS5]], i64 0, i64 0202// CHECK1-NEXT: store ptr null, ptr [[TMP22]], align 8203// CHECK1-NEXT: [[TMP23:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0204// CHECK1-NEXT: [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 0205// CHECK1-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 0206// CHECK1-NEXT: store i32 3, ptr [[TMP25]], align 4207// CHECK1-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 1208// CHECK1-NEXT: store i32 1, ptr [[TMP26]], align 4209// CHECK1-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 2210// CHECK1-NEXT: store ptr [[TMP23]], ptr [[TMP27]], align 8211// CHECK1-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 3212// CHECK1-NEXT: store ptr [[TMP24]], ptr [[TMP28]], align 8213// CHECK1-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 4214// CHECK1-NEXT: store ptr @.offload_sizes.1, ptr [[TMP29]], align 8215// CHECK1-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 5216// CHECK1-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP30]], align 8217// CHECK1-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 6218// CHECK1-NEXT: store ptr null, ptr [[TMP31]], align 8219// CHECK1-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 7220// CHECK1-NEXT: store ptr null, ptr [[TMP32]], align 8221// CHECK1-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 8222// CHECK1-NEXT: store i64 123, ptr [[TMP33]], align 8223// CHECK1-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 9224// CHECK1-NEXT: store i64 0, ptr [[TMP34]], align 8225// CHECK1-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 10226// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP35]], align 4227// CHECK1-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 11228// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP36]], align 4229// CHECK1-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 12230// CHECK1-NEXT: store i32 0, ptr [[TMP37]], align 4231// 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_l32.region_id, ptr [[KERNEL_ARGS7]])232// CHECK1-NEXT: [[TMP39:%.*]] = icmp ne i32 [[TMP38]], 0233// CHECK1-NEXT: br i1 [[TMP39]], label [[OMP_OFFLOAD_FAILED8:%.*]], label [[OMP_OFFLOAD_CONT9:%.*]]234// CHECK1: omp_offload.failed8:235// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l32(ptr [[THIS1]]) #[[ATTR2]]236// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT9]]237// CHECK1: omp_offload.cont9:238// CHECK1-NEXT: [[A10:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 0239// CHECK1-NEXT: [[TMP40:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS11]], i32 0, i32 0240// CHECK1-NEXT: store ptr [[THIS1]], ptr [[TMP40]], align 8241// CHECK1-NEXT: [[TMP41:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS12]], i32 0, i32 0242// CHECK1-NEXT: store ptr [[A10]], ptr [[TMP41]], align 8243// CHECK1-NEXT: [[TMP42:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS13]], i64 0, i64 0244// CHECK1-NEXT: store ptr null, ptr [[TMP42]], align 8245// CHECK1-NEXT: [[TMP43:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS11]], i32 0, i32 0246// CHECK1-NEXT: [[TMP44:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS12]], i32 0, i32 0247// CHECK1-NEXT: [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 0248// CHECK1-NEXT: store i32 3, ptr [[TMP45]], align 4249// CHECK1-NEXT: [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 1250// CHECK1-NEXT: store i32 1, ptr [[TMP46]], align 4251// CHECK1-NEXT: [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 2252// CHECK1-NEXT: store ptr [[TMP43]], ptr [[TMP47]], align 8253// CHECK1-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 3254// CHECK1-NEXT: store ptr [[TMP44]], ptr [[TMP48]], align 8255// CHECK1-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 4256// CHECK1-NEXT: store ptr @.offload_sizes.3, ptr [[TMP49]], align 8257// CHECK1-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 5258// CHECK1-NEXT: store ptr @.offload_maptypes.4, ptr [[TMP50]], align 8259// CHECK1-NEXT: [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 6260// CHECK1-NEXT: store ptr null, ptr [[TMP51]], align 8261// CHECK1-NEXT: [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 7262// CHECK1-NEXT: store ptr null, ptr [[TMP52]], align 8263// CHECK1-NEXT: [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 8264// CHECK1-NEXT: store i64 123, ptr [[TMP53]], align 8265// CHECK1-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 9266// CHECK1-NEXT: store i64 0, ptr [[TMP54]], align 8267// CHECK1-NEXT: [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 10268// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP55]], align 4269// CHECK1-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 11270// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP56]], align 4271// CHECK1-NEXT: [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 12272// CHECK1-NEXT: store i32 0, ptr [[TMP57]], align 4273// 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_l36.region_id, ptr [[KERNEL_ARGS15]])274// CHECK1-NEXT: [[TMP59:%.*]] = icmp ne i32 [[TMP58]], 0275// CHECK1-NEXT: br i1 [[TMP59]], label [[OMP_OFFLOAD_FAILED16:%.*]], label [[OMP_OFFLOAD_CONT17:%.*]]276// CHECK1: omp_offload.failed16:277// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l36(ptr [[THIS1]]) #[[ATTR2]]278// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT17]]279// CHECK1: omp_offload.cont17:280// CHECK1-NEXT: [[A18:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 0281// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A18]], i64 0, i64 0282// CHECK1-NEXT: [[TMP60:%.*]] = load i32, ptr [[ARRAYIDX]], align 4283// CHECK1-NEXT: ret i32 [[TMP60]]284//285//286// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28287// CHECK1-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1:[0-9]+]] {288// CHECK1-NEXT: entry:289// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8290// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8291// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8292// 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]])293// CHECK1-NEXT: ret void294//295//296// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28.omp_outlined297// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {298// CHECK1-NEXT: entry:299// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8300// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8301// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8302// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4303// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4304// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4305// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4306// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4307// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4308// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4309// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8310// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8311// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8312// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8313// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4314// CHECK1-NEXT: store i32 122, ptr [[DOTOMP_UB]], align 4315// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4316// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4317// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8318// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4319// 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)320// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4321// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 122322// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]323// CHECK1: cond.true:324// CHECK1-NEXT: br label [[COND_END:%.*]]325// CHECK1: cond.false:326// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4327// CHECK1-NEXT: br label [[COND_END]]328// CHECK1: cond.end:329// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 122, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]330// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4331// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4332// CHECK1-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4333// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]334// CHECK1: omp.inner.for.cond:335// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4336// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4337// CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]338// CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]339// CHECK1: omp.inner.for.body:340// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4341// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1342// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]343// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4344// CHECK1-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP0]], i32 0, i32 0345// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4346// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP9]] to i64347// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A]], i64 0, i64 [[IDXPROM]]348// CHECK1-NEXT: store i32 0, ptr [[ARRAYIDX]], align 4349// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]350// CHECK1: omp.body.continue:351// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]352// CHECK1: omp.inner.for.inc:353// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4354// CHECK1-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP10]], 1355// CHECK1-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4356// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]357// CHECK1: omp.inner.for.end:358// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]359// CHECK1: omp.loop.exit:360// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])361// CHECK1-NEXT: ret void362//363//364// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l32365// CHECK1-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1]] {366// CHECK1-NEXT: entry:367// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8368// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8369// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8370// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l32.omp_outlined, ptr [[TMP0]])371// CHECK1-NEXT: ret void372//373//374// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l32.omp_outlined375// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {376// CHECK1-NEXT: entry:377// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8378// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8379// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8380// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4381// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4382// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4383// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4384// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4385// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4386// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4387// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8388// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8389// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8390// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8391// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4392// CHECK1-NEXT: store i32 122, ptr [[DOTOMP_UB]], align 4393// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4394// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4395// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8396// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4397// 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)398// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4399// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 122400// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]401// CHECK1: cond.true:402// CHECK1-NEXT: br label [[COND_END:%.*]]403// CHECK1: cond.false:404// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4405// CHECK1-NEXT: br label [[COND_END]]406// CHECK1: cond.end:407// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 122, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]408// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4409// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4410// CHECK1-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4411// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]412// CHECK1: omp.inner.for.cond:413// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4414// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4415// CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]416// CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]417// CHECK1: omp.inner.for.body:418// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4419// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1420// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]421// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4422// CHECK1-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP0]], i32 0, i32 0423// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4424// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP9]] to i64425// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A]], i64 0, i64 [[IDXPROM]]426// CHECK1-NEXT: store i32 0, ptr [[ARRAYIDX]], align 4427// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]428// CHECK1: omp.body.continue:429// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]430// CHECK1: omp.inner.for.inc:431// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4432// CHECK1-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP10]], 1433// CHECK1-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4434// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]435// CHECK1: omp.inner.for.end:436// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]437// CHECK1: omp.loop.exit:438// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])439// CHECK1-NEXT: ret void440//441//442// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l36443// CHECK1-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1]] {444// CHECK1-NEXT: entry:445// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8446// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8447// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8448// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l36.omp_outlined, ptr [[TMP0]])449// CHECK1-NEXT: ret void450//451//452// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l36.omp_outlined453// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {454// CHECK1-NEXT: entry:455// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8456// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8457// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8458// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4459// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4460// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4461// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4462// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4463// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4464// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4465// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8466// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8467// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8468// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8469// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4470// CHECK1-NEXT: store i32 122, ptr [[DOTOMP_UB]], align 4471// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4472// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4473// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8474// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4475// 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)476// CHECK1-NEXT: br label [[OMP_DISPATCH_COND:%.*]]477// CHECK1: omp.dispatch.cond:478// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4479// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 122480// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]481// CHECK1: cond.true:482// CHECK1-NEXT: br label [[COND_END:%.*]]483// CHECK1: cond.false:484// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4485// CHECK1-NEXT: br label [[COND_END]]486// CHECK1: cond.end:487// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 122, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]488// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4489// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4490// CHECK1-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4491// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4492// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4493// CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]494// CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]495// CHECK1: omp.dispatch.body:496// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]497// CHECK1: omp.inner.for.cond:498// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP8:![0-9]+]]499// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP8]]500// CHECK1-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP8]], [[TMP9]]501// CHECK1-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]502// CHECK1: omp.inner.for.body:503// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP8]]504// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP10]], 1505// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]506// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP8]]507// CHECK1-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP0]], i32 0, i32 0508// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP8]]509// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP11]] to i64510// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A]], i64 0, i64 [[IDXPROM]]511// CHECK1-NEXT: store i32 0, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP8]]512// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]513// CHECK1: omp.body.continue:514// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]515// CHECK1: omp.inner.for.inc:516// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP8]]517// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP12]], 1518// CHECK1-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP8]]519// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP9:![0-9]+]]520// CHECK1: omp.inner.for.end:521// CHECK1-NEXT: br label [[OMP_DISPATCH_INC:%.*]]522// CHECK1: omp.dispatch.inc:523// CHECK1-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4524// CHECK1-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4525// CHECK1-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP13]], [[TMP14]]526// CHECK1-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_LB]], align 4527// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4528// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4529// CHECK1-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP15]], [[TMP16]]530// CHECK1-NEXT: store i32 [[ADD5]], ptr [[DOTOMP_UB]], align 4531// CHECK1-NEXT: br label [[OMP_DISPATCH_COND]]532// CHECK1: omp.dispatch.end:533// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])534// CHECK1-NEXT: ret void535//536//537// CHECK3-LABEL: define {{[^@]+}}@_Z21teams_template_structv538// CHECK3-SAME: () #[[ATTR0:[0-9]+]] {539// CHECK3-NEXT: entry:540// CHECK3-NEXT: [[V:%.*]] = alloca [[STRUCT_SS:%.*]], align 4541// CHECK3-NEXT: [[CALL:%.*]] = call noundef i32 @_ZN2SSIiLi123ELx456EE3fooEv(ptr noundef nonnull align 4 dereferenceable(496) [[V]])542// CHECK3-NEXT: ret i32 [[CALL]]543//544//545// CHECK3-LABEL: define {{[^@]+}}@_ZN2SSIiLi123ELx456EE3fooEv546// CHECK3-SAME: (ptr noundef nonnull align 4 dereferenceable(496) [[THIS:%.*]]) #[[ATTR0]] comdat align 2 {547// CHECK3-NEXT: entry:548// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4549// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4550// CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4551// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4552// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4553// CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8554// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS3:%.*]] = alloca [1 x ptr], align 4555// CHECK3-NEXT: [[DOTOFFLOAD_PTRS4:%.*]] = alloca [1 x ptr], align 4556// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS5:%.*]] = alloca [1 x ptr], align 4557// CHECK3-NEXT: [[_TMP6:%.*]] = alloca i32, align 4558// CHECK3-NEXT: [[KERNEL_ARGS7:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8559// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS11:%.*]] = alloca [1 x ptr], align 4560// CHECK3-NEXT: [[DOTOFFLOAD_PTRS12:%.*]] = alloca [1 x ptr], align 4561// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS13:%.*]] = alloca [1 x ptr], align 4562// CHECK3-NEXT: [[_TMP14:%.*]] = alloca i32, align 4563// CHECK3-NEXT: [[KERNEL_ARGS15:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8564// CHECK3-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4565// CHECK3-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4566// CHECK3-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[THIS1]], i32 0, i32 0567// CHECK3-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0568// CHECK3-NEXT: store ptr [[THIS1]], ptr [[TMP0]], align 4569// CHECK3-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0570// CHECK3-NEXT: store ptr [[A]], ptr [[TMP1]], align 4571// CHECK3-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0572// CHECK3-NEXT: store ptr null, ptr [[TMP2]], align 4573// CHECK3-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0574// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0575// CHECK3-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0576// CHECK3-NEXT: store i32 3, ptr [[TMP5]], align 4577// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1578// CHECK3-NEXT: store i32 1, ptr [[TMP6]], align 4579// CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2580// CHECK3-NEXT: store ptr [[TMP3]], ptr [[TMP7]], align 4581// CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3582// CHECK3-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 4583// CHECK3-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4584// CHECK3-NEXT: store ptr @.offload_sizes, ptr [[TMP9]], align 4585// CHECK3-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5586// CHECK3-NEXT: store ptr @.offload_maptypes, ptr [[TMP10]], align 4587// CHECK3-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6588// CHECK3-NEXT: store ptr null, ptr [[TMP11]], align 4589// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7590// CHECK3-NEXT: store ptr null, ptr [[TMP12]], align 4591// CHECK3-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8592// CHECK3-NEXT: store i64 123, ptr [[TMP13]], align 8593// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9594// CHECK3-NEXT: store i64 0, ptr [[TMP14]], align 8595// CHECK3-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10596// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 4597// CHECK3-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11598// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4599// CHECK3-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12600// CHECK3-NEXT: store i32 0, ptr [[TMP17]], align 4601// 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]])602// CHECK3-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0603// CHECK3-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]604// CHECK3: omp_offload.failed:605// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28(ptr [[THIS1]]) #[[ATTR2:[0-9]+]]606// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]]607// CHECK3: omp_offload.cont:608// CHECK3-NEXT: [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 0609// CHECK3-NEXT: [[TMP20:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0610// CHECK3-NEXT: store ptr [[THIS1]], ptr [[TMP20]], align 4611// CHECK3-NEXT: [[TMP21:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 0612// CHECK3-NEXT: store ptr [[A2]], ptr [[TMP21]], align 4613// CHECK3-NEXT: [[TMP22:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS5]], i32 0, i32 0614// CHECK3-NEXT: store ptr null, ptr [[TMP22]], align 4615// CHECK3-NEXT: [[TMP23:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0616// CHECK3-NEXT: [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 0617// CHECK3-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 0618// CHECK3-NEXT: store i32 3, ptr [[TMP25]], align 4619// CHECK3-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 1620// CHECK3-NEXT: store i32 1, ptr [[TMP26]], align 4621// CHECK3-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 2622// CHECK3-NEXT: store ptr [[TMP23]], ptr [[TMP27]], align 4623// CHECK3-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 3624// CHECK3-NEXT: store ptr [[TMP24]], ptr [[TMP28]], align 4625// CHECK3-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 4626// CHECK3-NEXT: store ptr @.offload_sizes.1, ptr [[TMP29]], align 4627// CHECK3-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 5628// CHECK3-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP30]], align 4629// CHECK3-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 6630// CHECK3-NEXT: store ptr null, ptr [[TMP31]], align 4631// CHECK3-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 7632// CHECK3-NEXT: store ptr null, ptr [[TMP32]], align 4633// CHECK3-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 8634// CHECK3-NEXT: store i64 123, ptr [[TMP33]], align 8635// CHECK3-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 9636// CHECK3-NEXT: store i64 0, ptr [[TMP34]], align 8637// CHECK3-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 10638// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP35]], align 4639// CHECK3-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 11640// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP36]], align 4641// CHECK3-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 12642// CHECK3-NEXT: store i32 0, ptr [[TMP37]], align 4643// 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_l32.region_id, ptr [[KERNEL_ARGS7]])644// CHECK3-NEXT: [[TMP39:%.*]] = icmp ne i32 [[TMP38]], 0645// CHECK3-NEXT: br i1 [[TMP39]], label [[OMP_OFFLOAD_FAILED8:%.*]], label [[OMP_OFFLOAD_CONT9:%.*]]646// CHECK3: omp_offload.failed8:647// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l32(ptr [[THIS1]]) #[[ATTR2]]648// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT9]]649// CHECK3: omp_offload.cont9:650// CHECK3-NEXT: [[A10:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 0651// CHECK3-NEXT: [[TMP40:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS11]], i32 0, i32 0652// CHECK3-NEXT: store ptr [[THIS1]], ptr [[TMP40]], align 4653// CHECK3-NEXT: [[TMP41:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS12]], i32 0, i32 0654// CHECK3-NEXT: store ptr [[A10]], ptr [[TMP41]], align 4655// CHECK3-NEXT: [[TMP42:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS13]], i32 0, i32 0656// CHECK3-NEXT: store ptr null, ptr [[TMP42]], align 4657// CHECK3-NEXT: [[TMP43:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS11]], i32 0, i32 0658// CHECK3-NEXT: [[TMP44:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS12]], i32 0, i32 0659// CHECK3-NEXT: [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 0660// CHECK3-NEXT: store i32 3, ptr [[TMP45]], align 4661// CHECK3-NEXT: [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 1662// CHECK3-NEXT: store i32 1, ptr [[TMP46]], align 4663// CHECK3-NEXT: [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 2664// CHECK3-NEXT: store ptr [[TMP43]], ptr [[TMP47]], align 4665// CHECK3-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 3666// CHECK3-NEXT: store ptr [[TMP44]], ptr [[TMP48]], align 4667// CHECK3-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 4668// CHECK3-NEXT: store ptr @.offload_sizes.3, ptr [[TMP49]], align 4669// CHECK3-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 5670// CHECK3-NEXT: store ptr @.offload_maptypes.4, ptr [[TMP50]], align 4671// CHECK3-NEXT: [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 6672// CHECK3-NEXT: store ptr null, ptr [[TMP51]], align 4673// CHECK3-NEXT: [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 7674// CHECK3-NEXT: store ptr null, ptr [[TMP52]], align 4675// CHECK3-NEXT: [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 8676// CHECK3-NEXT: store i64 123, ptr [[TMP53]], align 8677// CHECK3-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 9678// CHECK3-NEXT: store i64 0, ptr [[TMP54]], align 8679// CHECK3-NEXT: [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 10680// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP55]], align 4681// CHECK3-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 11682// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP56]], align 4683// CHECK3-NEXT: [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 12684// CHECK3-NEXT: store i32 0, ptr [[TMP57]], align 4685// 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_l36.region_id, ptr [[KERNEL_ARGS15]])686// CHECK3-NEXT: [[TMP59:%.*]] = icmp ne i32 [[TMP58]], 0687// CHECK3-NEXT: br i1 [[TMP59]], label [[OMP_OFFLOAD_FAILED16:%.*]], label [[OMP_OFFLOAD_CONT17:%.*]]688// CHECK3: omp_offload.failed16:689// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l36(ptr [[THIS1]]) #[[ATTR2]]690// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT17]]691// CHECK3: omp_offload.cont17:692// CHECK3-NEXT: [[A18:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 0693// CHECK3-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A18]], i32 0, i32 0694// CHECK3-NEXT: [[TMP60:%.*]] = load i32, ptr [[ARRAYIDX]], align 4695// CHECK3-NEXT: ret i32 [[TMP60]]696//697//698// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28699// CHECK3-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1:[0-9]+]] {700// CHECK3-NEXT: entry:701// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4702// CHECK3-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4703// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4704// 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]])705// CHECK3-NEXT: ret void706//707//708// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l28.omp_outlined709// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {710// CHECK3-NEXT: entry:711// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4712// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4713// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4714// CHECK3-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4715// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4716// CHECK3-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4717// CHECK3-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4718// CHECK3-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4719// CHECK3-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4720// CHECK3-NEXT: [[I:%.*]] = alloca i32, align 4721// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4722// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4723// CHECK3-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4724// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4725// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4726// CHECK3-NEXT: store i32 122, ptr [[DOTOMP_UB]], align 4727// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4728// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4729// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4730// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4731// 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)732// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4733// CHECK3-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 122734// CHECK3-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]735// CHECK3: cond.true:736// CHECK3-NEXT: br label [[COND_END:%.*]]737// CHECK3: cond.false:738// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4739// CHECK3-NEXT: br label [[COND_END]]740// CHECK3: cond.end:741// CHECK3-NEXT: [[COND:%.*]] = phi i32 [ 122, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]742// CHECK3-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4743// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4744// CHECK3-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4745// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]746// CHECK3: omp.inner.for.cond:747// CHECK3-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4748// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4749// CHECK3-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]750// CHECK3-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]751// CHECK3: omp.inner.for.body:752// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4753// CHECK3-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1754// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]755// CHECK3-NEXT: store i32 [[ADD]], ptr [[I]], align 4756// CHECK3-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP0]], i32 0, i32 0757// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4758// CHECK3-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A]], i32 0, i32 [[TMP9]]759// CHECK3-NEXT: store i32 0, ptr [[ARRAYIDX]], align 4760// CHECK3-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]761// CHECK3: omp.body.continue:762// CHECK3-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]763// CHECK3: omp.inner.for.inc:764// CHECK3-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4765// CHECK3-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP10]], 1766// CHECK3-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4767// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND]]768// CHECK3: omp.inner.for.end:769// CHECK3-NEXT: br label [[OMP_LOOP_EXIT:%.*]]770// CHECK3: omp.loop.exit:771// CHECK3-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])772// CHECK3-NEXT: ret void773//774//775// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l32776// CHECK3-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1]] {777// CHECK3-NEXT: entry:778// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4779// CHECK3-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4780// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4781// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l32.omp_outlined, ptr [[TMP0]])782// CHECK3-NEXT: ret void783//784//785// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l32.omp_outlined786// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {787// CHECK3-NEXT: entry:788// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4789// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4790// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4791// CHECK3-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4792// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4793// CHECK3-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4794// CHECK3-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4795// CHECK3-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4796// CHECK3-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4797// CHECK3-NEXT: [[I:%.*]] = alloca i32, align 4798// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4799// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4800// CHECK3-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4801// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4802// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4803// CHECK3-NEXT: store i32 122, ptr [[DOTOMP_UB]], align 4804// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4805// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4806// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4807// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4808// 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)809// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4810// CHECK3-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 122811// CHECK3-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]812// CHECK3: cond.true:813// CHECK3-NEXT: br label [[COND_END:%.*]]814// CHECK3: cond.false:815// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4816// CHECK3-NEXT: br label [[COND_END]]817// CHECK3: cond.end:818// CHECK3-NEXT: [[COND:%.*]] = phi i32 [ 122, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]819// CHECK3-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4820// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4821// CHECK3-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4822// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]823// CHECK3: omp.inner.for.cond:824// CHECK3-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4825// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4826// CHECK3-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]827// CHECK3-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]828// CHECK3: omp.inner.for.body:829// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4830// CHECK3-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1831// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]832// CHECK3-NEXT: store i32 [[ADD]], ptr [[I]], align 4833// CHECK3-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP0]], i32 0, i32 0834// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 4835// CHECK3-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A]], i32 0, i32 [[TMP9]]836// CHECK3-NEXT: store i32 0, ptr [[ARRAYIDX]], align 4837// CHECK3-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]838// CHECK3: omp.body.continue:839// CHECK3-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]840// CHECK3: omp.inner.for.inc:841// CHECK3-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4842// CHECK3-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP10]], 1843// CHECK3-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4844// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND]]845// CHECK3: omp.inner.for.end:846// CHECK3-NEXT: br label [[OMP_LOOP_EXIT:%.*]]847// CHECK3: omp.loop.exit:848// CHECK3-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])849// CHECK3-NEXT: ret void850//851//852// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l36853// CHECK3-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1]] {854// CHECK3-NEXT: entry:855// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4856// CHECK3-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4857// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4858// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l36.omp_outlined, ptr [[TMP0]])859// CHECK3-NEXT: ret void860//861//862// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l36.omp_outlined863// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {864// CHECK3-NEXT: entry:865// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4866// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4867// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4868// CHECK3-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4869// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4870// CHECK3-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4871// CHECK3-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4872// CHECK3-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4873// CHECK3-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4874// CHECK3-NEXT: [[I:%.*]] = alloca i32, align 4875// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4876// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4877// CHECK3-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4878// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4879// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4880// CHECK3-NEXT: store i32 122, ptr [[DOTOMP_UB]], align 4881// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4882// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4883// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4884// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4885// 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)886// CHECK3-NEXT: br label [[OMP_DISPATCH_COND:%.*]]887// CHECK3: omp.dispatch.cond:888// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4889// CHECK3-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 122890// CHECK3-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]891// CHECK3: cond.true:892// CHECK3-NEXT: br label [[COND_END:%.*]]893// CHECK3: cond.false:894// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4895// CHECK3-NEXT: br label [[COND_END]]896// CHECK3: cond.end:897// CHECK3-NEXT: [[COND:%.*]] = phi i32 [ 122, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]898// CHECK3-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4899// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4900// CHECK3-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4901// CHECK3-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4902// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4903// CHECK3-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]904// CHECK3-NEXT: br i1 [[CMP1]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]905// CHECK3: omp.dispatch.body:906// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]907// CHECK3: omp.inner.for.cond:908// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP9:![0-9]+]]909// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP9]]910// CHECK3-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP8]], [[TMP9]]911// CHECK3-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]912// CHECK3: omp.inner.for.body:913// CHECK3-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP9]]914// CHECK3-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP10]], 1915// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]916// CHECK3-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP9]]917// CHECK3-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP0]], i32 0, i32 0918// CHECK3-NEXT: [[TMP11:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP9]]919// CHECK3-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A]], i32 0, i32 [[TMP11]]920// CHECK3-NEXT: store i32 0, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP9]]921// CHECK3-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]922// CHECK3: omp.body.continue:923// CHECK3-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]924// CHECK3: omp.inner.for.inc:925// CHECK3-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP9]]926// CHECK3-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP12]], 1927// CHECK3-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP9]]928// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP10:![0-9]+]]929// CHECK3: omp.inner.for.end:930// CHECK3-NEXT: br label [[OMP_DISPATCH_INC:%.*]]931// CHECK3: omp.dispatch.inc:932// CHECK3-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4933// CHECK3-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4934// CHECK3-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP13]], [[TMP14]]935// CHECK3-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_LB]], align 4936// CHECK3-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4937// CHECK3-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4938// CHECK3-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP15]], [[TMP16]]939// CHECK3-NEXT: store i32 [[ADD5]], ptr [[DOTOMP_UB]], align 4940// CHECK3-NEXT: br label [[OMP_DISPATCH_COND]]941// CHECK3: omp.dispatch.end:942// CHECK3-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])943// CHECK3-NEXT: ret void944//945//946// CHECK9-LABEL: define {{[^@]+}}@main947// CHECK9-SAME: (i32 noundef signext [[ARGC:%.*]], ptr noundef [[ARGV:%.*]]) #[[ATTR0:[0-9]+]] {948// CHECK9-NEXT: entry:949// CHECK9-NEXT: [[RETVAL:%.*]] = alloca i32, align 4950// CHECK9-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 4951// CHECK9-NEXT: [[ARGV_ADDR:%.*]] = alloca ptr, align 8952// CHECK9-NEXT: [[N:%.*]] = alloca i32, align 4953// CHECK9-NEXT: [[SAVED_STACK:%.*]] = alloca ptr, align 8954// CHECK9-NEXT: [[__VLA_EXPR0:%.*]] = alloca i64, align 8955// CHECK9-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8956// CHECK9-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 8957// CHECK9-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 8958// CHECK9-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 8959// CHECK9-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 8960// CHECK9-NEXT: [[TMP:%.*]] = alloca i32, align 4961// CHECK9-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4962// CHECK9-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4963// CHECK9-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8964// CHECK9-NEXT: [[N_CASTED3:%.*]] = alloca i64, align 8965// CHECK9-NEXT: [[DOTOFFLOAD_BASEPTRS4:%.*]] = alloca [3 x ptr], align 8966// CHECK9-NEXT: [[DOTOFFLOAD_PTRS5:%.*]] = alloca [3 x ptr], align 8967// CHECK9-NEXT: [[DOTOFFLOAD_MAPPERS6:%.*]] = alloca [3 x ptr], align 8968// CHECK9-NEXT: [[DOTOFFLOAD_SIZES7:%.*]] = alloca [3 x i64], align 8969// CHECK9-NEXT: [[_TMP8:%.*]] = alloca i32, align 4970// CHECK9-NEXT: [[DOTCAPTURE_EXPR_9:%.*]] = alloca i32, align 4971// CHECK9-NEXT: [[DOTCAPTURE_EXPR_10:%.*]] = alloca i32, align 4972// CHECK9-NEXT: [[KERNEL_ARGS15:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8973// CHECK9-NEXT: [[DOTCAPTURE_EXPR_18:%.*]] = alloca i32, align 4974// CHECK9-NEXT: [[N_CASTED19:%.*]] = alloca i64, align 8975// CHECK9-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 8976// CHECK9-NEXT: [[DOTOFFLOAD_BASEPTRS20:%.*]] = alloca [4 x ptr], align 8977// CHECK9-NEXT: [[DOTOFFLOAD_PTRS21:%.*]] = alloca [4 x ptr], align 8978// CHECK9-NEXT: [[DOTOFFLOAD_MAPPERS22:%.*]] = alloca [4 x ptr], align 8979// CHECK9-NEXT: [[DOTOFFLOAD_SIZES23:%.*]] = alloca [4 x i64], align 8980// CHECK9-NEXT: [[_TMP24:%.*]] = alloca i32, align 4981// CHECK9-NEXT: [[DOTCAPTURE_EXPR_25:%.*]] = alloca i32, align 4982// CHECK9-NEXT: [[DOTCAPTURE_EXPR_26:%.*]] = alloca i32, align 4983// CHECK9-NEXT: [[KERNEL_ARGS31:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8984// CHECK9-NEXT: store i32 0, ptr [[RETVAL]], align 4985// CHECK9-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 4986// CHECK9-NEXT: store ptr [[ARGV]], ptr [[ARGV_ADDR]], align 8987// CHECK9-NEXT: store i32 100, ptr [[N]], align 4988// CHECK9-NEXT: [[TMP0:%.*]] = load i32, ptr [[N]], align 4989// CHECK9-NEXT: [[TMP1:%.*]] = zext i32 [[TMP0]] to i64990// CHECK9-NEXT: [[TMP2:%.*]] = call ptr @llvm.stacksave.p0()991// CHECK9-NEXT: store ptr [[TMP2]], ptr [[SAVED_STACK]], align 8992// CHECK9-NEXT: [[VLA:%.*]] = alloca i32, i64 [[TMP1]], align 4993// CHECK9-NEXT: store i64 [[TMP1]], ptr [[__VLA_EXPR0]], align 8994// CHECK9-NEXT: [[TMP3:%.*]] = load i32, ptr [[N]], align 4995// CHECK9-NEXT: store i32 [[TMP3]], ptr [[N_CASTED]], align 4996// CHECK9-NEXT: [[TMP4:%.*]] = load i64, ptr [[N_CASTED]], align 8997// CHECK9-NEXT: [[TMP5:%.*]] = mul nuw i64 [[TMP1]], 4998// CHECK9-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES]], ptr align 8 @.offload_sizes, i64 24, i1 false)999// CHECK9-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01000// CHECK9-NEXT: store i64 [[TMP4]], ptr [[TMP6]], align 81001// CHECK9-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01002// CHECK9-NEXT: store i64 [[TMP4]], ptr [[TMP7]], align 81003// CHECK9-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 01004// CHECK9-NEXT: store ptr null, ptr [[TMP8]], align 81005// CHECK9-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11006// CHECK9-NEXT: store i64 [[TMP1]], ptr [[TMP9]], align 81007// CHECK9-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11008// CHECK9-NEXT: store i64 [[TMP1]], ptr [[TMP10]], align 81009// CHECK9-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 11010// CHECK9-NEXT: store ptr null, ptr [[TMP11]], align 81011// CHECK9-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21012// CHECK9-NEXT: store ptr [[VLA]], ptr [[TMP12]], align 81013// CHECK9-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21014// CHECK9-NEXT: store ptr [[VLA]], ptr [[TMP13]], align 81015// CHECK9-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 21016// CHECK9-NEXT: store i64 [[TMP5]], ptr [[TMP14]], align 81017// CHECK9-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 21018// CHECK9-NEXT: store ptr null, ptr [[TMP15]], align 81019// CHECK9-NEXT: [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01020// CHECK9-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01021// CHECK9-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 01022// CHECK9-NEXT: [[TMP19:%.*]] = load i32, ptr [[N]], align 41023// CHECK9-NEXT: store i32 [[TMP19]], ptr [[DOTCAPTURE_EXPR_]], align 41024// CHECK9-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41025// CHECK9-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP20]], 01026// CHECK9-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11027// CHECK9-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11028// CHECK9-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41029// CHECK9-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41030// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP21]], 11031// CHECK9-NEXT: [[TMP22:%.*]] = zext i32 [[ADD]] to i641032// CHECK9-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01033// CHECK9-NEXT: store i32 3, ptr [[TMP23]], align 41034// CHECK9-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11035// CHECK9-NEXT: store i32 3, ptr [[TMP24]], align 41036// CHECK9-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21037// CHECK9-NEXT: store ptr [[TMP16]], ptr [[TMP25]], align 81038// CHECK9-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31039// CHECK9-NEXT: store ptr [[TMP17]], ptr [[TMP26]], align 81040// CHECK9-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41041// CHECK9-NEXT: store ptr [[TMP18]], ptr [[TMP27]], align 81042// CHECK9-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51043// CHECK9-NEXT: store ptr @.offload_maptypes, ptr [[TMP28]], align 81044// CHECK9-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61045// CHECK9-NEXT: store ptr null, ptr [[TMP29]], align 81046// CHECK9-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71047// CHECK9-NEXT: store ptr null, ptr [[TMP30]], align 81048// CHECK9-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81049// CHECK9-NEXT: store i64 [[TMP22]], ptr [[TMP31]], align 81050// CHECK9-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91051// CHECK9-NEXT: store i64 0, ptr [[TMP32]], align 81052// CHECK9-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101053// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP33]], align 41054// CHECK9-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111055// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP34]], align 41056// CHECK9-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121057// CHECK9-NEXT: store i32 0, ptr [[TMP35]], align 41058// 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_l94.region_id, ptr [[KERNEL_ARGS]])1059// CHECK9-NEXT: [[TMP37:%.*]] = icmp ne i32 [[TMP36]], 01060// CHECK9-NEXT: br i1 [[TMP37]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1061// CHECK9: omp_offload.failed:1062// CHECK9-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l94(i64 [[TMP4]], i64 [[TMP1]], ptr [[VLA]]) #[[ATTR3:[0-9]+]]1063// CHECK9-NEXT: br label [[OMP_OFFLOAD_CONT]]1064// CHECK9: omp_offload.cont:1065// CHECK9-NEXT: [[TMP38:%.*]] = load i32, ptr [[N]], align 41066// CHECK9-NEXT: store i32 [[TMP38]], ptr [[N_CASTED3]], align 41067// CHECK9-NEXT: [[TMP39:%.*]] = load i64, ptr [[N_CASTED3]], align 81068// CHECK9-NEXT: [[TMP40:%.*]] = mul nuw i64 [[TMP1]], 41069// CHECK9-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES7]], ptr align 8 @.offload_sizes.1, i64 24, i1 false)1070// CHECK9-NEXT: [[TMP41:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 01071// CHECK9-NEXT: store i64 [[TMP39]], ptr [[TMP41]], align 81072// CHECK9-NEXT: [[TMP42:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 01073// CHECK9-NEXT: store i64 [[TMP39]], ptr [[TMP42]], align 81074// CHECK9-NEXT: [[TMP43:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i64 0, i64 01075// CHECK9-NEXT: store ptr null, ptr [[TMP43]], align 81076// CHECK9-NEXT: [[TMP44:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 11077// CHECK9-NEXT: store i64 [[TMP1]], ptr [[TMP44]], align 81078// CHECK9-NEXT: [[TMP45:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 11079// CHECK9-NEXT: store i64 [[TMP1]], ptr [[TMP45]], align 81080// CHECK9-NEXT: [[TMP46:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i64 0, i64 11081// CHECK9-NEXT: store ptr null, ptr [[TMP46]], align 81082// CHECK9-NEXT: [[TMP47:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 21083// CHECK9-NEXT: store ptr [[VLA]], ptr [[TMP47]], align 81084// CHECK9-NEXT: [[TMP48:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 21085// CHECK9-NEXT: store ptr [[VLA]], ptr [[TMP48]], align 81086// CHECK9-NEXT: [[TMP49:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES7]], i32 0, i32 21087// CHECK9-NEXT: store i64 [[TMP40]], ptr [[TMP49]], align 81088// CHECK9-NEXT: [[TMP50:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i64 0, i64 21089// CHECK9-NEXT: store ptr null, ptr [[TMP50]], align 81090// CHECK9-NEXT: [[TMP51:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 01091// CHECK9-NEXT: [[TMP52:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 01092// CHECK9-NEXT: [[TMP53:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES7]], i32 0, i32 01093// CHECK9-NEXT: [[TMP54:%.*]] = load i32, ptr [[N]], align 41094// CHECK9-NEXT: store i32 [[TMP54]], ptr [[DOTCAPTURE_EXPR_9]], align 41095// CHECK9-NEXT: [[TMP55:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_9]], align 41096// CHECK9-NEXT: [[SUB11:%.*]] = sub nsw i32 [[TMP55]], 01097// CHECK9-NEXT: [[DIV12:%.*]] = sdiv i32 [[SUB11]], 11098// CHECK9-NEXT: [[SUB13:%.*]] = sub nsw i32 [[DIV12]], 11099// CHECK9-NEXT: store i32 [[SUB13]], ptr [[DOTCAPTURE_EXPR_10]], align 41100// CHECK9-NEXT: [[TMP56:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_10]], align 41101// CHECK9-NEXT: [[ADD14:%.*]] = add nsw i32 [[TMP56]], 11102// CHECK9-NEXT: [[TMP57:%.*]] = zext i32 [[ADD14]] to i641103// CHECK9-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 01104// CHECK9-NEXT: store i32 3, ptr [[TMP58]], align 41105// CHECK9-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 11106// CHECK9-NEXT: store i32 3, ptr [[TMP59]], align 41107// CHECK9-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 21108// CHECK9-NEXT: store ptr [[TMP51]], ptr [[TMP60]], align 81109// CHECK9-NEXT: [[TMP61:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 31110// CHECK9-NEXT: store ptr [[TMP52]], ptr [[TMP61]], align 81111// CHECK9-NEXT: [[TMP62:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 41112// CHECK9-NEXT: store ptr [[TMP53]], ptr [[TMP62]], align 81113// CHECK9-NEXT: [[TMP63:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 51114// CHECK9-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP63]], align 81115// CHECK9-NEXT: [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 61116// CHECK9-NEXT: store ptr null, ptr [[TMP64]], align 81117// CHECK9-NEXT: [[TMP65:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 71118// CHECK9-NEXT: store ptr null, ptr [[TMP65]], align 81119// CHECK9-NEXT: [[TMP66:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 81120// CHECK9-NEXT: store i64 [[TMP57]], ptr [[TMP66]], align 81121// CHECK9-NEXT: [[TMP67:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 91122// CHECK9-NEXT: store i64 0, ptr [[TMP67]], align 81123// CHECK9-NEXT: [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 101124// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP68]], align 41125// CHECK9-NEXT: [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 111126// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP69]], align 41127// CHECK9-NEXT: [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 121128// CHECK9-NEXT: store i32 0, ptr [[TMP70]], align 41129// 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_l98.region_id, ptr [[KERNEL_ARGS15]])1130// CHECK9-NEXT: [[TMP72:%.*]] = icmp ne i32 [[TMP71]], 01131// CHECK9-NEXT: br i1 [[TMP72]], label [[OMP_OFFLOAD_FAILED16:%.*]], label [[OMP_OFFLOAD_CONT17:%.*]]1132// CHECK9: omp_offload.failed16:1133// CHECK9-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l98(i64 [[TMP39]], i64 [[TMP1]], ptr [[VLA]]) #[[ATTR3]]1134// CHECK9-NEXT: br label [[OMP_OFFLOAD_CONT17]]1135// CHECK9: omp_offload.cont17:1136// CHECK9-NEXT: [[TMP73:%.*]] = load i32, ptr [[N]], align 41137// CHECK9-NEXT: store i32 [[TMP73]], ptr [[DOTCAPTURE_EXPR_18]], align 41138// CHECK9-NEXT: [[TMP74:%.*]] = load i32, ptr [[N]], align 41139// CHECK9-NEXT: store i32 [[TMP74]], ptr [[N_CASTED19]], align 41140// CHECK9-NEXT: [[TMP75:%.*]] = load i64, ptr [[N_CASTED19]], align 81141// CHECK9-NEXT: [[TMP76:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_18]], align 41142// CHECK9-NEXT: store i32 [[TMP76]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 41143// CHECK9-NEXT: [[TMP77:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 81144// CHECK9-NEXT: [[TMP78:%.*]] = mul nuw i64 [[TMP1]], 41145// CHECK9-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES23]], ptr align 8 @.offload_sizes.3, i64 32, i1 false)1146// CHECK9-NEXT: [[TMP79:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS20]], i32 0, i32 01147// CHECK9-NEXT: store i64 [[TMP75]], ptr [[TMP79]], align 81148// CHECK9-NEXT: [[TMP80:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS21]], i32 0, i32 01149// CHECK9-NEXT: store i64 [[TMP75]], ptr [[TMP80]], align 81150// CHECK9-NEXT: [[TMP81:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS22]], i64 0, i64 01151// CHECK9-NEXT: store ptr null, ptr [[TMP81]], align 81152// CHECK9-NEXT: [[TMP82:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS20]], i32 0, i32 11153// CHECK9-NEXT: store i64 [[TMP1]], ptr [[TMP82]], align 81154// CHECK9-NEXT: [[TMP83:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS21]], i32 0, i32 11155// CHECK9-NEXT: store i64 [[TMP1]], ptr [[TMP83]], align 81156// CHECK9-NEXT: [[TMP84:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS22]], i64 0, i64 11157// CHECK9-NEXT: store ptr null, ptr [[TMP84]], align 81158// CHECK9-NEXT: [[TMP85:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS20]], i32 0, i32 21159// CHECK9-NEXT: store ptr [[VLA]], ptr [[TMP85]], align 81160// CHECK9-NEXT: [[TMP86:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS21]], i32 0, i32 21161// CHECK9-NEXT: store ptr [[VLA]], ptr [[TMP86]], align 81162// CHECK9-NEXT: [[TMP87:%.*]] = getelementptr inbounds [4 x i64], ptr [[DOTOFFLOAD_SIZES23]], i32 0, i32 21163// CHECK9-NEXT: store i64 [[TMP78]], ptr [[TMP87]], align 81164// CHECK9-NEXT: [[TMP88:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS22]], i64 0, i64 21165// CHECK9-NEXT: store ptr null, ptr [[TMP88]], align 81166// CHECK9-NEXT: [[TMP89:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS20]], i32 0, i32 31167// CHECK9-NEXT: store i64 [[TMP77]], ptr [[TMP89]], align 81168// CHECK9-NEXT: [[TMP90:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS21]], i32 0, i32 31169// CHECK9-NEXT: store i64 [[TMP77]], ptr [[TMP90]], align 81170// CHECK9-NEXT: [[TMP91:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS22]], i64 0, i64 31171// CHECK9-NEXT: store ptr null, ptr [[TMP91]], align 81172// CHECK9-NEXT: [[TMP92:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS20]], i32 0, i32 01173// CHECK9-NEXT: [[TMP93:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS21]], i32 0, i32 01174// CHECK9-NEXT: [[TMP94:%.*]] = getelementptr inbounds [4 x i64], ptr [[DOTOFFLOAD_SIZES23]], i32 0, i32 01175// CHECK9-NEXT: [[TMP95:%.*]] = load i32, ptr [[N]], align 41176// CHECK9-NEXT: store i32 [[TMP95]], ptr [[DOTCAPTURE_EXPR_25]], align 41177// CHECK9-NEXT: [[TMP96:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_25]], align 41178// CHECK9-NEXT: [[SUB27:%.*]] = sub nsw i32 [[TMP96]], 01179// CHECK9-NEXT: [[DIV28:%.*]] = sdiv i32 [[SUB27]], 11180// CHECK9-NEXT: [[SUB29:%.*]] = sub nsw i32 [[DIV28]], 11181// CHECK9-NEXT: store i32 [[SUB29]], ptr [[DOTCAPTURE_EXPR_26]], align 41182// CHECK9-NEXT: [[TMP97:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_26]], align 41183// CHECK9-NEXT: [[ADD30:%.*]] = add nsw i32 [[TMP97]], 11184// CHECK9-NEXT: [[TMP98:%.*]] = zext i32 [[ADD30]] to i641185// CHECK9-NEXT: [[TMP99:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 01186// CHECK9-NEXT: store i32 3, ptr [[TMP99]], align 41187// CHECK9-NEXT: [[TMP100:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 11188// CHECK9-NEXT: store i32 4, ptr [[TMP100]], align 41189// CHECK9-NEXT: [[TMP101:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 21190// CHECK9-NEXT: store ptr [[TMP92]], ptr [[TMP101]], align 81191// CHECK9-NEXT: [[TMP102:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 31192// CHECK9-NEXT: store ptr [[TMP93]], ptr [[TMP102]], align 81193// CHECK9-NEXT: [[TMP103:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 41194// CHECK9-NEXT: store ptr [[TMP94]], ptr [[TMP103]], align 81195// CHECK9-NEXT: [[TMP104:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 51196// CHECK9-NEXT: store ptr @.offload_maptypes.4, ptr [[TMP104]], align 81197// CHECK9-NEXT: [[TMP105:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 61198// CHECK9-NEXT: store ptr null, ptr [[TMP105]], align 81199// CHECK9-NEXT: [[TMP106:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 71200// CHECK9-NEXT: store ptr null, ptr [[TMP106]], align 81201// CHECK9-NEXT: [[TMP107:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 81202// CHECK9-NEXT: store i64 [[TMP98]], ptr [[TMP107]], align 81203// CHECK9-NEXT: [[TMP108:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 91204// CHECK9-NEXT: store i64 0, ptr [[TMP108]], align 81205// CHECK9-NEXT: [[TMP109:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 101206// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP109]], align 41207// CHECK9-NEXT: [[TMP110:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 111208// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP110]], align 41209// CHECK9-NEXT: [[TMP111:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 121210// CHECK9-NEXT: store i32 0, ptr [[TMP111]], align 41211// CHECK9-NEXT: [[TMP112:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l102.region_id, ptr [[KERNEL_ARGS31]])1212// CHECK9-NEXT: [[TMP113:%.*]] = icmp ne i32 [[TMP112]], 01213// CHECK9-NEXT: br i1 [[TMP113]], label [[OMP_OFFLOAD_FAILED32:%.*]], label [[OMP_OFFLOAD_CONT33:%.*]]1214// CHECK9: omp_offload.failed32:1215// CHECK9-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l102(i64 [[TMP75]], i64 [[TMP1]], ptr [[VLA]], i64 [[TMP77]]) #[[ATTR3]]1216// CHECK9-NEXT: br label [[OMP_OFFLOAD_CONT33]]1217// CHECK9: omp_offload.cont33:1218// CHECK9-NEXT: [[TMP114:%.*]] = load i32, ptr [[ARGC_ADDR]], align 41219// CHECK9-NEXT: [[CALL:%.*]] = call noundef signext i32 @_Z5tmainIiLi10EEiT_(i32 noundef signext [[TMP114]])1220// CHECK9-NEXT: store i32 [[CALL]], ptr [[RETVAL]], align 41221// CHECK9-NEXT: [[TMP115:%.*]] = load ptr, ptr [[SAVED_STACK]], align 81222// CHECK9-NEXT: call void @llvm.stackrestore.p0(ptr [[TMP115]])1223// CHECK9-NEXT: [[TMP116:%.*]] = load i32, ptr [[RETVAL]], align 41224// CHECK9-NEXT: ret i32 [[TMP116]]1225//1226//1227// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l941228// CHECK9-SAME: (i64 noundef [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2:[0-9]+]] {1229// CHECK9-NEXT: entry:1230// CHECK9-NEXT: [[N_ADDR:%.*]] = alloca i64, align 81231// CHECK9-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 81232// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81233// CHECK9-NEXT: [[N_CASTED:%.*]] = alloca i64, align 81234// CHECK9-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 81235// CHECK9-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR]], align 81236// CHECK9-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81237// CHECK9-NEXT: [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR]], align 81238// CHECK9-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 81239// CHECK9-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 41240// CHECK9-NEXT: store i32 [[TMP2]], ptr [[N_CASTED]], align 41241// CHECK9-NEXT: [[TMP3:%.*]] = load i64, ptr [[N_CASTED]], align 81242// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l94.omp_outlined, i64 [[TMP3]], i64 [[TMP0]], ptr [[TMP1]])1243// CHECK9-NEXT: ret void1244//1245//1246// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l94.omp_outlined1247// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {1248// CHECK9-NEXT: entry:1249// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81250// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81251// CHECK9-NEXT: [[N_ADDR:%.*]] = alloca i64, align 81252// CHECK9-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 81253// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81254// CHECK9-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41255// CHECK9-NEXT: [[TMP:%.*]] = alloca i32, align 41256// CHECK9-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41257// CHECK9-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41258// CHECK9-NEXT: [[I:%.*]] = alloca i32, align 41259// CHECK9-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41260// CHECK9-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41261// CHECK9-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41262// CHECK9-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41263// CHECK9-NEXT: [[I3:%.*]] = alloca i32, align 41264// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81265// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81266// CHECK9-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 81267// CHECK9-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR]], align 81268// CHECK9-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81269// CHECK9-NEXT: [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR]], align 81270// CHECK9-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 81271// CHECK9-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 41272// CHECK9-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 41273// CHECK9-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41274// CHECK9-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP3]], 01275// CHECK9-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11276// CHECK9-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11277// CHECK9-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41278// CHECK9-NEXT: store i32 0, ptr [[I]], align 41279// CHECK9-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41280// CHECK9-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]1281// CHECK9-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1282// CHECK9: omp.precond.then:1283// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41284// CHECK9-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41285// CHECK9-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_UB]], align 41286// CHECK9-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41287// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41288// CHECK9-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81289// CHECK9-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 41290// CHECK9-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP7]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1291// CHECK9-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41292// CHECK9-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41293// CHECK9-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]1294// CHECK9-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1295// CHECK9: cond.true:1296// CHECK9-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41297// CHECK9-NEXT: br label [[COND_END:%.*]]1298// CHECK9: cond.false:1299// CHECK9-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41300// CHECK9-NEXT: br label [[COND_END]]1301// CHECK9: cond.end:1302// CHECK9-NEXT: [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]1303// CHECK9-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41304// CHECK9-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41305// CHECK9-NEXT: store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 41306// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1307// CHECK9: omp.inner.for.cond:1308// CHECK9-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41309// CHECK9-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41310// CHECK9-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]1311// CHECK9-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1312// CHECK9: omp.inner.for.body:1313// CHECK9-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41314// CHECK9-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP15]], 11315// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1316// CHECK9-NEXT: store i32 [[ADD]], ptr [[I3]], align 41317// CHECK9-NEXT: [[TMP16:%.*]] = load i32, ptr [[I3]], align 41318// CHECK9-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP16]] to i641319// CHECK9-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 [[IDXPROM]]1320// CHECK9-NEXT: store i32 0, ptr [[ARRAYIDX]], align 41321// CHECK9-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1322// CHECK9: omp.body.continue:1323// CHECK9-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1324// CHECK9: omp.inner.for.inc:1325// CHECK9-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41326// CHECK9-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP17]], 11327// CHECK9-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 41328// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND]]1329// CHECK9: omp.inner.for.end:1330// CHECK9-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1331// CHECK9: omp.loop.exit:1332// CHECK9-NEXT: [[TMP18:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81333// CHECK9-NEXT: [[TMP19:%.*]] = load i32, ptr [[TMP18]], align 41334// CHECK9-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP19]])1335// CHECK9-NEXT: br label [[OMP_PRECOND_END]]1336// CHECK9: omp.precond.end:1337// CHECK9-NEXT: ret void1338//1339//1340// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l981341// CHECK9-SAME: (i64 noundef [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {1342// CHECK9-NEXT: entry:1343// CHECK9-NEXT: [[N_ADDR:%.*]] = alloca i64, align 81344// CHECK9-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 81345// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81346// CHECK9-NEXT: [[N_CASTED:%.*]] = alloca i64, align 81347// CHECK9-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 81348// CHECK9-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR]], align 81349// CHECK9-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81350// CHECK9-NEXT: [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR]], align 81351// CHECK9-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 81352// CHECK9-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 41353// CHECK9-NEXT: store i32 [[TMP2]], ptr [[N_CASTED]], align 41354// CHECK9-NEXT: [[TMP3:%.*]] = load i64, ptr [[N_CASTED]], align 81355// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l98.omp_outlined, i64 [[TMP3]], i64 [[TMP0]], ptr [[TMP1]])1356// CHECK9-NEXT: ret void1357//1358//1359// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l98.omp_outlined1360// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {1361// CHECK9-NEXT: entry:1362// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81363// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81364// CHECK9-NEXT: [[N_ADDR:%.*]] = alloca i64, align 81365// CHECK9-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 81366// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81367// CHECK9-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41368// CHECK9-NEXT: [[TMP:%.*]] = alloca i32, align 41369// CHECK9-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41370// CHECK9-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41371// CHECK9-NEXT: [[I:%.*]] = alloca i32, align 41372// CHECK9-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41373// CHECK9-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41374// CHECK9-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41375// CHECK9-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41376// CHECK9-NEXT: [[I3:%.*]] = alloca i32, align 41377// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81378// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81379// CHECK9-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 81380// CHECK9-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR]], align 81381// CHECK9-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81382// CHECK9-NEXT: [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR]], align 81383// CHECK9-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 81384// CHECK9-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 41385// CHECK9-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 41386// CHECK9-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41387// CHECK9-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP3]], 01388// CHECK9-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11389// CHECK9-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11390// CHECK9-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41391// CHECK9-NEXT: store i32 0, ptr [[I]], align 41392// CHECK9-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41393// CHECK9-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]1394// CHECK9-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1395// CHECK9: omp.precond.then:1396// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41397// CHECK9-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41398// CHECK9-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_UB]], align 41399// CHECK9-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41400// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41401// CHECK9-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81402// CHECK9-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 41403// CHECK9-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP7]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1404// CHECK9-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41405// CHECK9-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41406// CHECK9-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]1407// CHECK9-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1408// CHECK9: cond.true:1409// CHECK9-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41410// CHECK9-NEXT: br label [[COND_END:%.*]]1411// CHECK9: cond.false:1412// CHECK9-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41413// CHECK9-NEXT: br label [[COND_END]]1414// CHECK9: cond.end:1415// CHECK9-NEXT: [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]1416// CHECK9-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41417// CHECK9-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41418// CHECK9-NEXT: store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 41419// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1420// CHECK9: omp.inner.for.cond:1421// CHECK9-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41422// CHECK9-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41423// CHECK9-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]1424// CHECK9-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1425// CHECK9: omp.inner.for.body:1426// CHECK9-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41427// CHECK9-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP15]], 11428// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1429// CHECK9-NEXT: store i32 [[ADD]], ptr [[I3]], align 41430// CHECK9-NEXT: [[TMP16:%.*]] = load i32, ptr [[I3]], align 41431// CHECK9-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP16]] to i641432// CHECK9-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 [[IDXPROM]]1433// CHECK9-NEXT: store i32 0, ptr [[ARRAYIDX]], align 41434// CHECK9-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1435// CHECK9: omp.body.continue:1436// CHECK9-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1437// CHECK9: omp.inner.for.inc:1438// CHECK9-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41439// CHECK9-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP17]], 11440// CHECK9-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 41441// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND]]1442// CHECK9: omp.inner.for.end:1443// CHECK9-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1444// CHECK9: omp.loop.exit:1445// CHECK9-NEXT: [[TMP18:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81446// CHECK9-NEXT: [[TMP19:%.*]] = load i32, ptr [[TMP18]], align 41447// CHECK9-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP19]])1448// CHECK9-NEXT: br label [[OMP_PRECOND_END]]1449// CHECK9: omp.precond.end:1450// CHECK9-NEXT: ret void1451//1452//1453// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l1021454// CHECK9-SAME: (i64 noundef [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR2]] {1455// CHECK9-NEXT: entry:1456// CHECK9-NEXT: [[N_ADDR:%.*]] = alloca i64, align 81457// CHECK9-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 81458// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81459// CHECK9-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 81460// CHECK9-NEXT: [[N_CASTED:%.*]] = alloca i64, align 81461// CHECK9-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 81462// CHECK9-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 81463// CHECK9-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR]], align 81464// CHECK9-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81465// CHECK9-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 81466// CHECK9-NEXT: [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR]], align 81467// CHECK9-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 81468// CHECK9-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 41469// CHECK9-NEXT: store i32 [[TMP2]], ptr [[N_CASTED]], align 41470// CHECK9-NEXT: [[TMP3:%.*]] = load i64, ptr [[N_CASTED]], align 81471// CHECK9-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 41472// CHECK9-NEXT: store i32 [[TMP4]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 41473// CHECK9-NEXT: [[TMP5:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 81474// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 4, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l102.omp_outlined, i64 [[TMP3]], i64 [[TMP0]], ptr [[TMP1]], i64 [[TMP5]])1475// CHECK9-NEXT: ret void1476//1477//1478// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l102.omp_outlined1479// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR2]] {1480// CHECK9-NEXT: entry:1481// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81482// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81483// CHECK9-NEXT: [[N_ADDR:%.*]] = alloca i64, align 81484// CHECK9-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 81485// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81486// CHECK9-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 81487// CHECK9-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41488// CHECK9-NEXT: [[TMP:%.*]] = alloca i32, align 41489// CHECK9-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41490// CHECK9-NEXT: [[DOTCAPTURE_EXPR_2:%.*]] = alloca i32, align 41491// CHECK9-NEXT: [[I:%.*]] = alloca i32, align 41492// CHECK9-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41493// CHECK9-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41494// CHECK9-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41495// CHECK9-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41496// CHECK9-NEXT: [[I4:%.*]] = alloca i32, align 41497// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81498// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81499// CHECK9-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 81500// CHECK9-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR]], align 81501// CHECK9-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81502// CHECK9-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 81503// CHECK9-NEXT: [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR]], align 81504// CHECK9-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 81505// CHECK9-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 41506// CHECK9-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_1]], align 41507// CHECK9-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41508// CHECK9-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP3]], 01509// CHECK9-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11510// CHECK9-NEXT: [[SUB3:%.*]] = sub nsw i32 [[DIV]], 11511// CHECK9-NEXT: store i32 [[SUB3]], ptr [[DOTCAPTURE_EXPR_2]], align 41512// CHECK9-NEXT: store i32 0, ptr [[I]], align 41513// CHECK9-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41514// CHECK9-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]1515// CHECK9-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1516// CHECK9: omp.precond.then:1517// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41518// CHECK9-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 41519// CHECK9-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_UB]], align 41520// CHECK9-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41521// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41522// CHECK9-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 41523// CHECK9-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81524// CHECK9-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41525// CHECK9-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP8]], i32 91, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 [[TMP6]])1526// CHECK9-NEXT: br label [[OMP_DISPATCH_COND:%.*]]1527// CHECK9: omp.dispatch.cond:1528// CHECK9-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41529// CHECK9-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 41530// CHECK9-NEXT: [[CMP5:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]1531// CHECK9-NEXT: br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1532// CHECK9: cond.true:1533// CHECK9-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 41534// CHECK9-NEXT: br label [[COND_END:%.*]]1535// CHECK9: cond.false:1536// CHECK9-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41537// CHECK9-NEXT: br label [[COND_END]]1538// CHECK9: cond.end:1539// CHECK9-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]1540// CHECK9-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41541// CHECK9-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41542// CHECK9-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 41543// CHECK9-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41544// CHECK9-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41545// CHECK9-NEXT: [[CMP6:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]1546// CHECK9-NEXT: br i1 [[CMP6]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]1547// CHECK9: omp.dispatch.body:1548// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1549// CHECK9: omp.inner.for.cond:1550// CHECK9-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11:![0-9]+]]1551// CHECK9-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP11]]1552// CHECK9-NEXT: [[CMP7:%.*]] = icmp sle i32 [[TMP16]], [[TMP17]]1553// CHECK9-NEXT: br i1 [[CMP7]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1554// CHECK9: omp.inner.for.body:1555// CHECK9-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]1556// CHECK9-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP18]], 11557// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1558// CHECK9-NEXT: store i32 [[ADD]], ptr [[I4]], align 4, !llvm.access.group [[ACC_GRP11]]1559// CHECK9-NEXT: [[TMP19:%.*]] = load i32, ptr [[I4]], align 4, !llvm.access.group [[ACC_GRP11]]1560// CHECK9-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP19]] to i641561// CHECK9-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 [[IDXPROM]]1562// CHECK9-NEXT: store i32 0, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP11]]1563// CHECK9-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1564// CHECK9: omp.body.continue:1565// CHECK9-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1566// CHECK9: omp.inner.for.inc:1567// CHECK9-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]1568// CHECK9-NEXT: [[ADD8:%.*]] = add nsw i32 [[TMP20]], 11569// CHECK9-NEXT: store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]1570// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP12:![0-9]+]]1571// CHECK9: omp.inner.for.end:1572// CHECK9-NEXT: br label [[OMP_DISPATCH_INC:%.*]]1573// CHECK9: omp.dispatch.inc:1574// CHECK9-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41575// CHECK9-NEXT: [[TMP22:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41576// CHECK9-NEXT: [[ADD9:%.*]] = add nsw i32 [[TMP21]], [[TMP22]]1577// CHECK9-NEXT: store i32 [[ADD9]], ptr [[DOTOMP_LB]], align 41578// CHECK9-NEXT: [[TMP23:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41579// CHECK9-NEXT: [[TMP24:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41580// CHECK9-NEXT: [[ADD10:%.*]] = add nsw i32 [[TMP23]], [[TMP24]]1581// CHECK9-NEXT: store i32 [[ADD10]], ptr [[DOTOMP_UB]], align 41582// CHECK9-NEXT: br label [[OMP_DISPATCH_COND]]1583// CHECK9: omp.dispatch.end:1584// CHECK9-NEXT: [[TMP25:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81585// CHECK9-NEXT: [[TMP26:%.*]] = load i32, ptr [[TMP25]], align 41586// CHECK9-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP26]])1587// CHECK9-NEXT: br label [[OMP_PRECOND_END]]1588// CHECK9: omp.precond.end:1589// CHECK9-NEXT: ret void1590//1591//1592// CHECK9-LABEL: define {{[^@]+}}@_Z5tmainIiLi10EEiT_1593// CHECK9-SAME: (i32 noundef signext [[ARGC:%.*]]) #[[ATTR5:[0-9]+]] comdat {1594// CHECK9-NEXT: entry:1595// CHECK9-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 41596// CHECK9-NEXT: [[A:%.*]] = alloca [10 x i32], align 41597// CHECK9-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 81598// CHECK9-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 81599// CHECK9-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 81600// CHECK9-NEXT: [[TMP:%.*]] = alloca i32, align 41601// CHECK9-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81602// CHECK9-NEXT: [[DOTOFFLOAD_BASEPTRS1:%.*]] = alloca [1 x ptr], align 81603// CHECK9-NEXT: [[DOTOFFLOAD_PTRS2:%.*]] = alloca [1 x ptr], align 81604// CHECK9-NEXT: [[DOTOFFLOAD_MAPPERS3:%.*]] = alloca [1 x ptr], align 81605// CHECK9-NEXT: [[_TMP4:%.*]] = alloca i32, align 41606// CHECK9-NEXT: [[KERNEL_ARGS5:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 81607// CHECK9-NEXT: [[DOTOFFLOAD_BASEPTRS8:%.*]] = alloca [1 x ptr], align 81608// CHECK9-NEXT: [[DOTOFFLOAD_PTRS9:%.*]] = alloca [1 x ptr], align 81609// CHECK9-NEXT: [[DOTOFFLOAD_MAPPERS10:%.*]] = alloca [1 x ptr], align 81610// CHECK9-NEXT: [[_TMP11:%.*]] = alloca i32, align 41611// CHECK9-NEXT: [[KERNEL_ARGS12:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 81612// CHECK9-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 41613// CHECK9-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01614// CHECK9-NEXT: store ptr [[A]], ptr [[TMP0]], align 81615// CHECK9-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01616// CHECK9-NEXT: store ptr [[A]], ptr [[TMP1]], align 81617// CHECK9-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 01618// CHECK9-NEXT: store ptr null, ptr [[TMP2]], align 81619// CHECK9-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01620// CHECK9-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01621// CHECK9-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01622// CHECK9-NEXT: store i32 3, ptr [[TMP5]], align 41623// CHECK9-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11624// CHECK9-NEXT: store i32 1, ptr [[TMP6]], align 41625// CHECK9-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21626// CHECK9-NEXT: store ptr [[TMP3]], ptr [[TMP7]], align 81627// CHECK9-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31628// CHECK9-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 81629// CHECK9-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41630// CHECK9-NEXT: store ptr @.offload_sizes.5, ptr [[TMP9]], align 81631// CHECK9-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51632// CHECK9-NEXT: store ptr @.offload_maptypes.6, ptr [[TMP10]], align 81633// CHECK9-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61634// CHECK9-NEXT: store ptr null, ptr [[TMP11]], align 81635// CHECK9-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71636// CHECK9-NEXT: store ptr null, ptr [[TMP12]], align 81637// CHECK9-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81638// CHECK9-NEXT: store i64 10, ptr [[TMP13]], align 81639// CHECK9-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91640// CHECK9-NEXT: store i64 0, ptr [[TMP14]], align 81641// CHECK9-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101642// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 41643// CHECK9-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111644// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 41645// CHECK9-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121646// CHECK9-NEXT: store i32 0, ptr [[TMP17]], align 41647// 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__l76.region_id, ptr [[KERNEL_ARGS]])1648// CHECK9-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 01649// CHECK9-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1650// CHECK9: omp_offload.failed:1651// CHECK9-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l76(ptr [[A]]) #[[ATTR3]]1652// CHECK9-NEXT: br label [[OMP_OFFLOAD_CONT]]1653// CHECK9: omp_offload.cont:1654// CHECK9-NEXT: [[TMP20:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 01655// CHECK9-NEXT: store ptr [[A]], ptr [[TMP20]], align 81656// CHECK9-NEXT: [[TMP21:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 01657// CHECK9-NEXT: store ptr [[A]], ptr [[TMP21]], align 81658// CHECK9-NEXT: [[TMP22:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS3]], i64 0, i64 01659// CHECK9-NEXT: store ptr null, ptr [[TMP22]], align 81660// CHECK9-NEXT: [[TMP23:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 01661// CHECK9-NEXT: [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 01662// CHECK9-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 01663// CHECK9-NEXT: store i32 3, ptr [[TMP25]], align 41664// CHECK9-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 11665// CHECK9-NEXT: store i32 1, ptr [[TMP26]], align 41666// CHECK9-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 21667// CHECK9-NEXT: store ptr [[TMP23]], ptr [[TMP27]], align 81668// CHECK9-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 31669// CHECK9-NEXT: store ptr [[TMP24]], ptr [[TMP28]], align 81670// CHECK9-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 41671// CHECK9-NEXT: store ptr @.offload_sizes.7, ptr [[TMP29]], align 81672// CHECK9-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 51673// CHECK9-NEXT: store ptr @.offload_maptypes.8, ptr [[TMP30]], align 81674// CHECK9-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 61675// CHECK9-NEXT: store ptr null, ptr [[TMP31]], align 81676// CHECK9-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 71677// CHECK9-NEXT: store ptr null, ptr [[TMP32]], align 81678// CHECK9-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 81679// CHECK9-NEXT: store i64 10, ptr [[TMP33]], align 81680// CHECK9-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 91681// CHECK9-NEXT: store i64 0, ptr [[TMP34]], align 81682// CHECK9-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 101683// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP35]], align 41684// CHECK9-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 111685// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP36]], align 41686// CHECK9-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 121687// CHECK9-NEXT: store i32 0, ptr [[TMP37]], align 41688// 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__l80.region_id, ptr [[KERNEL_ARGS5]])1689// CHECK9-NEXT: [[TMP39:%.*]] = icmp ne i32 [[TMP38]], 01690// CHECK9-NEXT: br i1 [[TMP39]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]]1691// CHECK9: omp_offload.failed6:1692// CHECK9-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l80(ptr [[A]]) #[[ATTR3]]1693// CHECK9-NEXT: br label [[OMP_OFFLOAD_CONT7]]1694// CHECK9: omp_offload.cont7:1695// CHECK9-NEXT: [[TMP40:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 01696// CHECK9-NEXT: store ptr [[A]], ptr [[TMP40]], align 81697// CHECK9-NEXT: [[TMP41:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 01698// CHECK9-NEXT: store ptr [[A]], ptr [[TMP41]], align 81699// CHECK9-NEXT: [[TMP42:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS10]], i64 0, i64 01700// CHECK9-NEXT: store ptr null, ptr [[TMP42]], align 81701// CHECK9-NEXT: [[TMP43:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 01702// CHECK9-NEXT: [[TMP44:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 01703// CHECK9-NEXT: [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 01704// CHECK9-NEXT: store i32 3, ptr [[TMP45]], align 41705// CHECK9-NEXT: [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 11706// CHECK9-NEXT: store i32 1, ptr [[TMP46]], align 41707// CHECK9-NEXT: [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 21708// CHECK9-NEXT: store ptr [[TMP43]], ptr [[TMP47]], align 81709// CHECK9-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 31710// CHECK9-NEXT: store ptr [[TMP44]], ptr [[TMP48]], align 81711// CHECK9-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 41712// CHECK9-NEXT: store ptr @.offload_sizes.9, ptr [[TMP49]], align 81713// CHECK9-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 51714// CHECK9-NEXT: store ptr @.offload_maptypes.10, ptr [[TMP50]], align 81715// CHECK9-NEXT: [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 61716// CHECK9-NEXT: store ptr null, ptr [[TMP51]], align 81717// CHECK9-NEXT: [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 71718// CHECK9-NEXT: store ptr null, ptr [[TMP52]], align 81719// CHECK9-NEXT: [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 81720// CHECK9-NEXT: store i64 10, ptr [[TMP53]], align 81721// CHECK9-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 91722// CHECK9-NEXT: store i64 0, ptr [[TMP54]], align 81723// CHECK9-NEXT: [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 101724// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP55]], align 41725// CHECK9-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 111726// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP56]], align 41727// CHECK9-NEXT: [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 121728// CHECK9-NEXT: store i32 0, ptr [[TMP57]], align 41729// 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__l84.region_id, ptr [[KERNEL_ARGS12]])1730// CHECK9-NEXT: [[TMP59:%.*]] = icmp ne i32 [[TMP58]], 01731// CHECK9-NEXT: br i1 [[TMP59]], label [[OMP_OFFLOAD_FAILED13:%.*]], label [[OMP_OFFLOAD_CONT14:%.*]]1732// CHECK9: omp_offload.failed13:1733// CHECK9-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l84(ptr [[A]]) #[[ATTR3]]1734// CHECK9-NEXT: br label [[OMP_OFFLOAD_CONT14]]1735// CHECK9: omp_offload.cont14:1736// CHECK9-NEXT: ret i32 01737//1738//1739// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l761740// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {1741// CHECK9-NEXT: entry:1742// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81743// CHECK9-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81744// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81745// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l76.omp_outlined, ptr [[TMP0]])1746// CHECK9-NEXT: ret void1747//1748//1749// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l76.omp_outlined1750// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {1751// CHECK9-NEXT: entry:1752// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81753// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81754// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81755// CHECK9-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41756// CHECK9-NEXT: [[TMP:%.*]] = alloca i32, align 41757// CHECK9-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41758// CHECK9-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41759// CHECK9-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41760// CHECK9-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41761// CHECK9-NEXT: [[I:%.*]] = alloca i32, align 41762// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81763// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81764// CHECK9-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81765// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81766// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41767// CHECK9-NEXT: store i32 9, ptr [[DOTOMP_UB]], align 41768// CHECK9-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41769// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41770// CHECK9-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81771// CHECK9-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41772// 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)1773// CHECK9-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41774// CHECK9-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 91775// CHECK9-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1776// CHECK9: cond.true:1777// CHECK9-NEXT: br label [[COND_END:%.*]]1778// CHECK9: cond.false:1779// CHECK9-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41780// CHECK9-NEXT: br label [[COND_END]]1781// CHECK9: cond.end:1782// CHECK9-NEXT: [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]1783// CHECK9-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41784// CHECK9-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41785// CHECK9-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 41786// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1787// CHECK9: omp.inner.for.cond:1788// CHECK9-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41789// CHECK9-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41790// CHECK9-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]1791// CHECK9-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1792// CHECK9: omp.inner.for.body:1793// CHECK9-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41794// CHECK9-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 11795// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1796// CHECK9-NEXT: store i32 [[ADD]], ptr [[I]], align 41797// CHECK9-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 41798// CHECK9-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP9]] to i641799// CHECK9-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]1800// CHECK9-NEXT: store i32 0, ptr [[ARRAYIDX]], align 41801// CHECK9-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1802// CHECK9: omp.body.continue:1803// CHECK9-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1804// CHECK9: omp.inner.for.inc:1805// CHECK9-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41806// CHECK9-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP10]], 11807// CHECK9-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 41808// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND]]1809// CHECK9: omp.inner.for.end:1810// CHECK9-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1811// CHECK9: omp.loop.exit:1812// CHECK9-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])1813// CHECK9-NEXT: ret void1814//1815//1816// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l801817// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {1818// CHECK9-NEXT: entry:1819// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81820// CHECK9-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81821// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81822// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l80.omp_outlined, ptr [[TMP0]])1823// CHECK9-NEXT: ret void1824//1825//1826// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l80.omp_outlined1827// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {1828// CHECK9-NEXT: entry:1829// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81830// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81831// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81832// CHECK9-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41833// CHECK9-NEXT: [[TMP:%.*]] = alloca i32, align 41834// CHECK9-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41835// CHECK9-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41836// CHECK9-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41837// CHECK9-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41838// CHECK9-NEXT: [[I:%.*]] = alloca i32, align 41839// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81840// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81841// CHECK9-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81842// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81843// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41844// CHECK9-NEXT: store i32 9, ptr [[DOTOMP_UB]], align 41845// CHECK9-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41846// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41847// CHECK9-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81848// CHECK9-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41849// 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)1850// CHECK9-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41851// CHECK9-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 91852// CHECK9-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1853// CHECK9: cond.true:1854// CHECK9-NEXT: br label [[COND_END:%.*]]1855// CHECK9: cond.false:1856// CHECK9-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41857// CHECK9-NEXT: br label [[COND_END]]1858// CHECK9: cond.end:1859// CHECK9-NEXT: [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]1860// CHECK9-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41861// CHECK9-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41862// CHECK9-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 41863// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1864// CHECK9: omp.inner.for.cond:1865// CHECK9-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41866// CHECK9-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41867// CHECK9-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]1868// CHECK9-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1869// CHECK9: omp.inner.for.body:1870// CHECK9-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41871// CHECK9-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 11872// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1873// CHECK9-NEXT: store i32 [[ADD]], ptr [[I]], align 41874// CHECK9-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 41875// CHECK9-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP9]] to i641876// CHECK9-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]1877// CHECK9-NEXT: store i32 0, ptr [[ARRAYIDX]], align 41878// CHECK9-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1879// CHECK9: omp.body.continue:1880// CHECK9-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1881// CHECK9: omp.inner.for.inc:1882// CHECK9-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41883// CHECK9-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP10]], 11884// CHECK9-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 41885// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND]]1886// CHECK9: omp.inner.for.end:1887// CHECK9-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1888// CHECK9: omp.loop.exit:1889// CHECK9-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])1890// CHECK9-NEXT: ret void1891//1892//1893// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l841894// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {1895// CHECK9-NEXT: entry:1896// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81897// CHECK9-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81898// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81899// 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]])1900// CHECK9-NEXT: ret void1901//1902//1903// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l84.omp_outlined1904// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {1905// CHECK9-NEXT: entry:1906// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81907// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81908// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81909// CHECK9-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41910// CHECK9-NEXT: [[TMP:%.*]] = alloca i32, align 41911// CHECK9-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41912// CHECK9-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41913// CHECK9-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41914// CHECK9-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41915// CHECK9-NEXT: [[I:%.*]] = alloca i32, align 41916// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81917// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81918// CHECK9-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81919// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 81920// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41921// CHECK9-NEXT: store i32 9, ptr [[DOTOMP_UB]], align 41922// CHECK9-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41923// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41924// CHECK9-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81925// CHECK9-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41926// 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)1927// CHECK9-NEXT: br label [[OMP_DISPATCH_COND:%.*]]1928// CHECK9: omp.dispatch.cond:1929// CHECK9-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41930// CHECK9-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 91931// CHECK9-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1932// CHECK9: cond.true:1933// CHECK9-NEXT: br label [[COND_END:%.*]]1934// CHECK9: cond.false:1935// CHECK9-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41936// CHECK9-NEXT: br label [[COND_END]]1937// CHECK9: cond.end:1938// CHECK9-NEXT: [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]1939// CHECK9-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41940// CHECK9-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41941// CHECK9-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 41942// CHECK9-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41943// CHECK9-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41944// CHECK9-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]1945// CHECK9-NEXT: br i1 [[CMP1]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]1946// CHECK9: omp.dispatch.body:1947// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1948// CHECK9: omp.inner.for.cond:1949// CHECK9-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP14:![0-9]+]]1950// CHECK9-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP14]]1951// CHECK9-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP8]], [[TMP9]]1952// CHECK9-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1953// CHECK9: omp.inner.for.body:1954// CHECK9-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP14]]1955// CHECK9-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP10]], 11956// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1957// CHECK9-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP14]]1958// CHECK9-NEXT: [[TMP11:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP14]]1959// CHECK9-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP11]] to i641960// CHECK9-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]1961// CHECK9-NEXT: store i32 0, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP14]]1962// CHECK9-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1963// CHECK9: omp.body.continue:1964// CHECK9-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1965// CHECK9: omp.inner.for.inc:1966// CHECK9-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP14]]1967// CHECK9-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP12]], 11968// CHECK9-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP14]]1969// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP15:![0-9]+]]1970// CHECK9: omp.inner.for.end:1971// CHECK9-NEXT: br label [[OMP_DISPATCH_INC:%.*]]1972// CHECK9: omp.dispatch.inc:1973// CHECK9-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41974// CHECK9-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41975// CHECK9-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP13]], [[TMP14]]1976// CHECK9-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_LB]], align 41977// CHECK9-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41978// CHECK9-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 41979// CHECK9-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP15]], [[TMP16]]1980// CHECK9-NEXT: store i32 [[ADD5]], ptr [[DOTOMP_UB]], align 41981// CHECK9-NEXT: br label [[OMP_DISPATCH_COND]]1982// CHECK9: omp.dispatch.end:1983// CHECK9-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])1984// CHECK9-NEXT: ret void1985//1986//1987// CHECK11-LABEL: define {{[^@]+}}@main1988// CHECK11-SAME: (i32 noundef [[ARGC:%.*]], ptr noundef [[ARGV:%.*]]) #[[ATTR0:[0-9]+]] {1989// CHECK11-NEXT: entry:1990// CHECK11-NEXT: [[RETVAL:%.*]] = alloca i32, align 41991// CHECK11-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 41992// CHECK11-NEXT: [[ARGV_ADDR:%.*]] = alloca ptr, align 41993// CHECK11-NEXT: [[N:%.*]] = alloca i32, align 41994// CHECK11-NEXT: [[SAVED_STACK:%.*]] = alloca ptr, align 41995// CHECK11-NEXT: [[__VLA_EXPR0:%.*]] = alloca i32, align 41996// CHECK11-NEXT: [[N_CASTED:%.*]] = alloca i32, align 41997// CHECK11-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 41998// CHECK11-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 41999// CHECK11-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 42000// CHECK11-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 42001// CHECK11-NEXT: [[TMP:%.*]] = alloca i32, align 42002// CHECK11-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42003// CHECK11-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42004// CHECK11-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 82005// CHECK11-NEXT: [[N_CASTED3:%.*]] = alloca i32, align 42006// CHECK11-NEXT: [[DOTOFFLOAD_BASEPTRS4:%.*]] = alloca [3 x ptr], align 42007// CHECK11-NEXT: [[DOTOFFLOAD_PTRS5:%.*]] = alloca [3 x ptr], align 42008// CHECK11-NEXT: [[DOTOFFLOAD_MAPPERS6:%.*]] = alloca [3 x ptr], align 42009// CHECK11-NEXT: [[DOTOFFLOAD_SIZES7:%.*]] = alloca [3 x i64], align 42010// CHECK11-NEXT: [[_TMP8:%.*]] = alloca i32, align 42011// CHECK11-NEXT: [[DOTCAPTURE_EXPR_9:%.*]] = alloca i32, align 42012// CHECK11-NEXT: [[DOTCAPTURE_EXPR_10:%.*]] = alloca i32, align 42013// CHECK11-NEXT: [[KERNEL_ARGS15:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 82014// CHECK11-NEXT: [[DOTCAPTURE_EXPR_18:%.*]] = alloca i32, align 42015// CHECK11-NEXT: [[N_CASTED19:%.*]] = alloca i32, align 42016// CHECK11-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i32, align 42017// CHECK11-NEXT: [[DOTOFFLOAD_BASEPTRS20:%.*]] = alloca [4 x ptr], align 42018// CHECK11-NEXT: [[DOTOFFLOAD_PTRS21:%.*]] = alloca [4 x ptr], align 42019// CHECK11-NEXT: [[DOTOFFLOAD_MAPPERS22:%.*]] = alloca [4 x ptr], align 42020// CHECK11-NEXT: [[DOTOFFLOAD_SIZES23:%.*]] = alloca [4 x i64], align 42021// CHECK11-NEXT: [[_TMP24:%.*]] = alloca i32, align 42022// CHECK11-NEXT: [[DOTCAPTURE_EXPR_25:%.*]] = alloca i32, align 42023// CHECK11-NEXT: [[DOTCAPTURE_EXPR_26:%.*]] = alloca i32, align 42024// CHECK11-NEXT: [[KERNEL_ARGS31:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 82025// CHECK11-NEXT: store i32 0, ptr [[RETVAL]], align 42026// CHECK11-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 42027// CHECK11-NEXT: store ptr [[ARGV]], ptr [[ARGV_ADDR]], align 42028// CHECK11-NEXT: store i32 100, ptr [[N]], align 42029// CHECK11-NEXT: [[TMP0:%.*]] = load i32, ptr [[N]], align 42030// CHECK11-NEXT: [[TMP1:%.*]] = call ptr @llvm.stacksave.p0()2031// CHECK11-NEXT: store ptr [[TMP1]], ptr [[SAVED_STACK]], align 42032// CHECK11-NEXT: [[VLA:%.*]] = alloca i32, i32 [[TMP0]], align 42033// CHECK11-NEXT: store i32 [[TMP0]], ptr [[__VLA_EXPR0]], align 42034// CHECK11-NEXT: [[TMP2:%.*]] = load i32, ptr [[N]], align 42035// CHECK11-NEXT: store i32 [[TMP2]], ptr [[N_CASTED]], align 42036// CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[N_CASTED]], align 42037// CHECK11-NEXT: [[TMP4:%.*]] = mul nuw i32 [[TMP0]], 42038// CHECK11-NEXT: [[TMP5:%.*]] = sext i32 [[TMP4]] to i642039// CHECK11-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[DOTOFFLOAD_SIZES]], ptr align 4 @.offload_sizes, i32 24, i1 false)2040// CHECK11-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 02041// CHECK11-NEXT: store i32 [[TMP3]], ptr [[TMP6]], align 42042// CHECK11-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 02043// CHECK11-NEXT: store i32 [[TMP3]], ptr [[TMP7]], align 42044// CHECK11-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 02045// CHECK11-NEXT: store ptr null, ptr [[TMP8]], align 42046// CHECK11-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 12047// CHECK11-NEXT: store i32 [[TMP0]], ptr [[TMP9]], align 42048// CHECK11-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 12049// CHECK11-NEXT: store i32 [[TMP0]], ptr [[TMP10]], align 42050// CHECK11-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 12051// CHECK11-NEXT: store ptr null, ptr [[TMP11]], align 42052// CHECK11-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 22053// CHECK11-NEXT: store ptr [[VLA]], ptr [[TMP12]], align 42054// CHECK11-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 22055// CHECK11-NEXT: store ptr [[VLA]], ptr [[TMP13]], align 42056// CHECK11-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 22057// CHECK11-NEXT: store i64 [[TMP5]], ptr [[TMP14]], align 42058// CHECK11-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 22059// CHECK11-NEXT: store ptr null, ptr [[TMP15]], align 42060// CHECK11-NEXT: [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 02061// CHECK11-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 02062// CHECK11-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 02063// CHECK11-NEXT: [[TMP19:%.*]] = load i32, ptr [[N]], align 42064// CHECK11-NEXT: store i32 [[TMP19]], ptr [[DOTCAPTURE_EXPR_]], align 42065// CHECK11-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42066// CHECK11-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP20]], 02067// CHECK11-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12068// CHECK11-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12069// CHECK11-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42070// CHECK11-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42071// CHECK11-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP21]], 12072// CHECK11-NEXT: [[TMP22:%.*]] = zext i32 [[ADD]] to i642073// CHECK11-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 02074// CHECK11-NEXT: store i32 3, ptr [[TMP23]], align 42075// CHECK11-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12076// CHECK11-NEXT: store i32 3, ptr [[TMP24]], align 42077// CHECK11-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 22078// CHECK11-NEXT: store ptr [[TMP16]], ptr [[TMP25]], align 42079// CHECK11-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 32080// CHECK11-NEXT: store ptr [[TMP17]], ptr [[TMP26]], align 42081// CHECK11-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 42082// CHECK11-NEXT: store ptr [[TMP18]], ptr [[TMP27]], align 42083// CHECK11-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 52084// CHECK11-NEXT: store ptr @.offload_maptypes, ptr [[TMP28]], align 42085// CHECK11-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 62086// CHECK11-NEXT: store ptr null, ptr [[TMP29]], align 42087// CHECK11-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 72088// CHECK11-NEXT: store ptr null, ptr [[TMP30]], align 42089// CHECK11-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 82090// CHECK11-NEXT: store i64 [[TMP22]], ptr [[TMP31]], align 82091// CHECK11-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 92092// CHECK11-NEXT: store i64 0, ptr [[TMP32]], align 82093// CHECK11-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 102094// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP33]], align 42095// CHECK11-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 112096// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP34]], align 42097// CHECK11-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 122098// CHECK11-NEXT: store i32 0, ptr [[TMP35]], align 42099// 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_l94.region_id, ptr [[KERNEL_ARGS]])2100// CHECK11-NEXT: [[TMP37:%.*]] = icmp ne i32 [[TMP36]], 02101// CHECK11-NEXT: br i1 [[TMP37]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]2102// CHECK11: omp_offload.failed:2103// CHECK11-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l94(i32 [[TMP3]], i32 [[TMP0]], ptr [[VLA]]) #[[ATTR3:[0-9]+]]2104// CHECK11-NEXT: br label [[OMP_OFFLOAD_CONT]]2105// CHECK11: omp_offload.cont:2106// CHECK11-NEXT: [[TMP38:%.*]] = load i32, ptr [[N]], align 42107// CHECK11-NEXT: store i32 [[TMP38]], ptr [[N_CASTED3]], align 42108// CHECK11-NEXT: [[TMP39:%.*]] = load i32, ptr [[N_CASTED3]], align 42109// CHECK11-NEXT: [[TMP40:%.*]] = mul nuw i32 [[TMP0]], 42110// CHECK11-NEXT: [[TMP41:%.*]] = sext i32 [[TMP40]] to i642111// CHECK11-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[DOTOFFLOAD_SIZES7]], ptr align 4 @.offload_sizes.1, i32 24, i1 false)2112// CHECK11-NEXT: [[TMP42:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 02113// CHECK11-NEXT: store i32 [[TMP39]], ptr [[TMP42]], align 42114// CHECK11-NEXT: [[TMP43:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 02115// CHECK11-NEXT: store i32 [[TMP39]], ptr [[TMP43]], align 42116// CHECK11-NEXT: [[TMP44:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i32 0, i32 02117// CHECK11-NEXT: store ptr null, ptr [[TMP44]], align 42118// CHECK11-NEXT: [[TMP45:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 12119// CHECK11-NEXT: store i32 [[TMP0]], ptr [[TMP45]], align 42120// CHECK11-NEXT: [[TMP46:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 12121// CHECK11-NEXT: store i32 [[TMP0]], ptr [[TMP46]], align 42122// CHECK11-NEXT: [[TMP47:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i32 0, i32 12123// CHECK11-NEXT: store ptr null, ptr [[TMP47]], align 42124// CHECK11-NEXT: [[TMP48:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 22125// CHECK11-NEXT: store ptr [[VLA]], ptr [[TMP48]], align 42126// CHECK11-NEXT: [[TMP49:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 22127// CHECK11-NEXT: store ptr [[VLA]], ptr [[TMP49]], align 42128// CHECK11-NEXT: [[TMP50:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES7]], i32 0, i32 22129// CHECK11-NEXT: store i64 [[TMP41]], ptr [[TMP50]], align 42130// CHECK11-NEXT: [[TMP51:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i32 0, i32 22131// CHECK11-NEXT: store ptr null, ptr [[TMP51]], align 42132// CHECK11-NEXT: [[TMP52:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 02133// CHECK11-NEXT: [[TMP53:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 02134// CHECK11-NEXT: [[TMP54:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES7]], i32 0, i32 02135// CHECK11-NEXT: [[TMP55:%.*]] = load i32, ptr [[N]], align 42136// CHECK11-NEXT: store i32 [[TMP55]], ptr [[DOTCAPTURE_EXPR_9]], align 42137// CHECK11-NEXT: [[TMP56:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_9]], align 42138// CHECK11-NEXT: [[SUB11:%.*]] = sub nsw i32 [[TMP56]], 02139// CHECK11-NEXT: [[DIV12:%.*]] = sdiv i32 [[SUB11]], 12140// CHECK11-NEXT: [[SUB13:%.*]] = sub nsw i32 [[DIV12]], 12141// CHECK11-NEXT: store i32 [[SUB13]], ptr [[DOTCAPTURE_EXPR_10]], align 42142// CHECK11-NEXT: [[TMP57:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_10]], align 42143// CHECK11-NEXT: [[ADD14:%.*]] = add nsw i32 [[TMP57]], 12144// CHECK11-NEXT: [[TMP58:%.*]] = zext i32 [[ADD14]] to i642145// CHECK11-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 02146// CHECK11-NEXT: store i32 3, ptr [[TMP59]], align 42147// CHECK11-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 12148// CHECK11-NEXT: store i32 3, ptr [[TMP60]], align 42149// CHECK11-NEXT: [[TMP61:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 22150// CHECK11-NEXT: store ptr [[TMP52]], ptr [[TMP61]], align 42151// CHECK11-NEXT: [[TMP62:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 32152// CHECK11-NEXT: store ptr [[TMP53]], ptr [[TMP62]], align 42153// CHECK11-NEXT: [[TMP63:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 42154// CHECK11-NEXT: store ptr [[TMP54]], ptr [[TMP63]], align 42155// CHECK11-NEXT: [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 52156// CHECK11-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP64]], align 42157// CHECK11-NEXT: [[TMP65:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 62158// CHECK11-NEXT: store ptr null, ptr [[TMP65]], align 42159// CHECK11-NEXT: [[TMP66:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 72160// CHECK11-NEXT: store ptr null, ptr [[TMP66]], align 42161// CHECK11-NEXT: [[TMP67:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 82162// CHECK11-NEXT: store i64 [[TMP58]], ptr [[TMP67]], align 82163// CHECK11-NEXT: [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 92164// CHECK11-NEXT: store i64 0, ptr [[TMP68]], align 82165// CHECK11-NEXT: [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 102166// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP69]], align 42167// CHECK11-NEXT: [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 112168// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP70]], align 42169// CHECK11-NEXT: [[TMP71:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 122170// CHECK11-NEXT: store i32 0, ptr [[TMP71]], align 42171// 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_l98.region_id, ptr [[KERNEL_ARGS15]])2172// CHECK11-NEXT: [[TMP73:%.*]] = icmp ne i32 [[TMP72]], 02173// CHECK11-NEXT: br i1 [[TMP73]], label [[OMP_OFFLOAD_FAILED16:%.*]], label [[OMP_OFFLOAD_CONT17:%.*]]2174// CHECK11: omp_offload.failed16:2175// CHECK11-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l98(i32 [[TMP39]], i32 [[TMP0]], ptr [[VLA]]) #[[ATTR3]]2176// CHECK11-NEXT: br label [[OMP_OFFLOAD_CONT17]]2177// CHECK11: omp_offload.cont17:2178// CHECK11-NEXT: [[TMP74:%.*]] = load i32, ptr [[N]], align 42179// CHECK11-NEXT: store i32 [[TMP74]], ptr [[DOTCAPTURE_EXPR_18]], align 42180// CHECK11-NEXT: [[TMP75:%.*]] = load i32, ptr [[N]], align 42181// CHECK11-NEXT: store i32 [[TMP75]], ptr [[N_CASTED19]], align 42182// CHECK11-NEXT: [[TMP76:%.*]] = load i32, ptr [[N_CASTED19]], align 42183// CHECK11-NEXT: [[TMP77:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_18]], align 42184// CHECK11-NEXT: store i32 [[TMP77]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 42185// CHECK11-NEXT: [[TMP78:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED]], align 42186// CHECK11-NEXT: [[TMP79:%.*]] = mul nuw i32 [[TMP0]], 42187// CHECK11-NEXT: [[TMP80:%.*]] = sext i32 [[TMP79]] to i642188// CHECK11-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[DOTOFFLOAD_SIZES23]], ptr align 4 @.offload_sizes.3, i32 32, i1 false)2189// CHECK11-NEXT: [[TMP81:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS20]], i32 0, i32 02190// CHECK11-NEXT: store i32 [[TMP76]], ptr [[TMP81]], align 42191// CHECK11-NEXT: [[TMP82:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS21]], i32 0, i32 02192// CHECK11-NEXT: store i32 [[TMP76]], ptr [[TMP82]], align 42193// CHECK11-NEXT: [[TMP83:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS22]], i32 0, i32 02194// CHECK11-NEXT: store ptr null, ptr [[TMP83]], align 42195// CHECK11-NEXT: [[TMP84:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS20]], i32 0, i32 12196// CHECK11-NEXT: store i32 [[TMP0]], ptr [[TMP84]], align 42197// CHECK11-NEXT: [[TMP85:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS21]], i32 0, i32 12198// CHECK11-NEXT: store i32 [[TMP0]], ptr [[TMP85]], align 42199// CHECK11-NEXT: [[TMP86:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS22]], i32 0, i32 12200// CHECK11-NEXT: store ptr null, ptr [[TMP86]], align 42201// CHECK11-NEXT: [[TMP87:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS20]], i32 0, i32 22202// CHECK11-NEXT: store ptr [[VLA]], ptr [[TMP87]], align 42203// CHECK11-NEXT: [[TMP88:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS21]], i32 0, i32 22204// CHECK11-NEXT: store ptr [[VLA]], ptr [[TMP88]], align 42205// CHECK11-NEXT: [[TMP89:%.*]] = getelementptr inbounds [4 x i64], ptr [[DOTOFFLOAD_SIZES23]], i32 0, i32 22206// CHECK11-NEXT: store i64 [[TMP80]], ptr [[TMP89]], align 42207// CHECK11-NEXT: [[TMP90:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS22]], i32 0, i32 22208// CHECK11-NEXT: store ptr null, ptr [[TMP90]], align 42209// CHECK11-NEXT: [[TMP91:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS20]], i32 0, i32 32210// CHECK11-NEXT: store i32 [[TMP78]], ptr [[TMP91]], align 42211// CHECK11-NEXT: [[TMP92:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS21]], i32 0, i32 32212// CHECK11-NEXT: store i32 [[TMP78]], ptr [[TMP92]], align 42213// CHECK11-NEXT: [[TMP93:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS22]], i32 0, i32 32214// CHECK11-NEXT: store ptr null, ptr [[TMP93]], align 42215// CHECK11-NEXT: [[TMP94:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS20]], i32 0, i32 02216// CHECK11-NEXT: [[TMP95:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS21]], i32 0, i32 02217// CHECK11-NEXT: [[TMP96:%.*]] = getelementptr inbounds [4 x i64], ptr [[DOTOFFLOAD_SIZES23]], i32 0, i32 02218// CHECK11-NEXT: [[TMP97:%.*]] = load i32, ptr [[N]], align 42219// CHECK11-NEXT: store i32 [[TMP97]], ptr [[DOTCAPTURE_EXPR_25]], align 42220// CHECK11-NEXT: [[TMP98:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_25]], align 42221// CHECK11-NEXT: [[SUB27:%.*]] = sub nsw i32 [[TMP98]], 02222// CHECK11-NEXT: [[DIV28:%.*]] = sdiv i32 [[SUB27]], 12223// CHECK11-NEXT: [[SUB29:%.*]] = sub nsw i32 [[DIV28]], 12224// CHECK11-NEXT: store i32 [[SUB29]], ptr [[DOTCAPTURE_EXPR_26]], align 42225// CHECK11-NEXT: [[TMP99:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_26]], align 42226// CHECK11-NEXT: [[ADD30:%.*]] = add nsw i32 [[TMP99]], 12227// CHECK11-NEXT: [[TMP100:%.*]] = zext i32 [[ADD30]] to i642228// CHECK11-NEXT: [[TMP101:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 02229// CHECK11-NEXT: store i32 3, ptr [[TMP101]], align 42230// CHECK11-NEXT: [[TMP102:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 12231// CHECK11-NEXT: store i32 4, ptr [[TMP102]], align 42232// CHECK11-NEXT: [[TMP103:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 22233// CHECK11-NEXT: store ptr [[TMP94]], ptr [[TMP103]], align 42234// CHECK11-NEXT: [[TMP104:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 32235// CHECK11-NEXT: store ptr [[TMP95]], ptr [[TMP104]], align 42236// CHECK11-NEXT: [[TMP105:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 42237// CHECK11-NEXT: store ptr [[TMP96]], ptr [[TMP105]], align 42238// CHECK11-NEXT: [[TMP106:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 52239// CHECK11-NEXT: store ptr @.offload_maptypes.4, ptr [[TMP106]], align 42240// CHECK11-NEXT: [[TMP107:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 62241// CHECK11-NEXT: store ptr null, ptr [[TMP107]], align 42242// CHECK11-NEXT: [[TMP108:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 72243// CHECK11-NEXT: store ptr null, ptr [[TMP108]], align 42244// CHECK11-NEXT: [[TMP109:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 82245// CHECK11-NEXT: store i64 [[TMP100]], ptr [[TMP109]], align 82246// CHECK11-NEXT: [[TMP110:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 92247// CHECK11-NEXT: store i64 0, ptr [[TMP110]], align 82248// CHECK11-NEXT: [[TMP111:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 102249// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP111]], align 42250// CHECK11-NEXT: [[TMP112:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 112251// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP112]], align 42252// CHECK11-NEXT: [[TMP113:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS31]], i32 0, i32 122253// CHECK11-NEXT: store i32 0, ptr [[TMP113]], align 42254// CHECK11-NEXT: [[TMP114:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l102.region_id, ptr [[KERNEL_ARGS31]])2255// CHECK11-NEXT: [[TMP115:%.*]] = icmp ne i32 [[TMP114]], 02256// CHECK11-NEXT: br i1 [[TMP115]], label [[OMP_OFFLOAD_FAILED32:%.*]], label [[OMP_OFFLOAD_CONT33:%.*]]2257// CHECK11: omp_offload.failed32:2258// CHECK11-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l102(i32 [[TMP76]], i32 [[TMP0]], ptr [[VLA]], i32 [[TMP78]]) #[[ATTR3]]2259// CHECK11-NEXT: br label [[OMP_OFFLOAD_CONT33]]2260// CHECK11: omp_offload.cont33:2261// CHECK11-NEXT: [[TMP116:%.*]] = load i32, ptr [[ARGC_ADDR]], align 42262// CHECK11-NEXT: [[CALL:%.*]] = call noundef i32 @_Z5tmainIiLi10EEiT_(i32 noundef [[TMP116]])2263// CHECK11-NEXT: store i32 [[CALL]], ptr [[RETVAL]], align 42264// CHECK11-NEXT: [[TMP117:%.*]] = load ptr, ptr [[SAVED_STACK]], align 42265// CHECK11-NEXT: call void @llvm.stackrestore.p0(ptr [[TMP117]])2266// CHECK11-NEXT: [[TMP118:%.*]] = load i32, ptr [[RETVAL]], align 42267// CHECK11-NEXT: ret i32 [[TMP118]]2268//2269//2270// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l942271// CHECK11-SAME: (i32 noundef [[N:%.*]], i32 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2:[0-9]+]] {2272// CHECK11-NEXT: entry:2273// CHECK11-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42274// CHECK11-NEXT: [[VLA_ADDR:%.*]] = alloca i32, align 42275// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42276// CHECK11-NEXT: [[N_CASTED:%.*]] = alloca i32, align 42277// CHECK11-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42278// CHECK11-NEXT: store i32 [[VLA]], ptr [[VLA_ADDR]], align 42279// CHECK11-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42280// CHECK11-NEXT: [[TMP0:%.*]] = load i32, ptr [[VLA_ADDR]], align 42281// CHECK11-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 42282// CHECK11-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 42283// CHECK11-NEXT: store i32 [[TMP2]], ptr [[N_CASTED]], align 42284// CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[N_CASTED]], align 42285// CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l94.omp_outlined, i32 [[TMP3]], i32 [[TMP0]], ptr [[TMP1]])2286// CHECK11-NEXT: ret void2287//2288//2289// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l94.omp_outlined2290// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[N:%.*]], i32 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {2291// CHECK11-NEXT: entry:2292// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42293// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42294// CHECK11-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42295// CHECK11-NEXT: [[VLA_ADDR:%.*]] = alloca i32, align 42296// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42297// CHECK11-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42298// CHECK11-NEXT: [[TMP:%.*]] = alloca i32, align 42299// CHECK11-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42300// CHECK11-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42301// CHECK11-NEXT: [[I:%.*]] = alloca i32, align 42302// CHECK11-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 42303// CHECK11-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 42304// CHECK11-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42305// CHECK11-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42306// CHECK11-NEXT: [[I3:%.*]] = alloca i32, align 42307// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42308// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42309// CHECK11-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42310// CHECK11-NEXT: store i32 [[VLA]], ptr [[VLA_ADDR]], align 42311// CHECK11-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42312// CHECK11-NEXT: [[TMP0:%.*]] = load i32, ptr [[VLA_ADDR]], align 42313// CHECK11-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 42314// CHECK11-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 42315// CHECK11-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 42316// CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42317// CHECK11-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP3]], 02318// CHECK11-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12319// CHECK11-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12320// CHECK11-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42321// CHECK11-NEXT: store i32 0, ptr [[I]], align 42322// CHECK11-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42323// CHECK11-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]2324// CHECK11-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2325// CHECK11: omp.precond.then:2326// CHECK11-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 42327// CHECK11-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42328// CHECK11-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_UB]], align 42329// CHECK11-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42330// CHECK11-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42331// CHECK11-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42332// CHECK11-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 42333// CHECK11-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP7]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2334// CHECK11-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42335// CHECK11-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42336// CHECK11-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]2337// CHECK11-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2338// CHECK11: cond.true:2339// CHECK11-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42340// CHECK11-NEXT: br label [[COND_END:%.*]]2341// CHECK11: cond.false:2342// CHECK11-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42343// CHECK11-NEXT: br label [[COND_END]]2344// CHECK11: cond.end:2345// CHECK11-NEXT: [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]2346// CHECK11-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 42347// CHECK11-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42348// CHECK11-NEXT: store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 42349// CHECK11-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2350// CHECK11: omp.inner.for.cond:2351// CHECK11-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42352// CHECK11-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42353// CHECK11-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]2354// CHECK11-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2355// CHECK11: omp.inner.for.body:2356// CHECK11-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42357// CHECK11-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP15]], 12358// CHECK11-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]2359// CHECK11-NEXT: store i32 [[ADD]], ptr [[I3]], align 42360// CHECK11-NEXT: [[TMP16:%.*]] = load i32, ptr [[I3]], align 42361// CHECK11-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i32 [[TMP16]]2362// CHECK11-NEXT: store i32 0, ptr [[ARRAYIDX]], align 42363// CHECK11-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]2364// CHECK11: omp.body.continue:2365// CHECK11-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2366// CHECK11: omp.inner.for.inc:2367// CHECK11-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42368// CHECK11-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP17]], 12369// CHECK11-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 42370// CHECK11-NEXT: br label [[OMP_INNER_FOR_COND]]2371// CHECK11: omp.inner.for.end:2372// CHECK11-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2373// CHECK11: omp.loop.exit:2374// CHECK11-NEXT: [[TMP18:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42375// CHECK11-NEXT: [[TMP19:%.*]] = load i32, ptr [[TMP18]], align 42376// CHECK11-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP19]])2377// CHECK11-NEXT: br label [[OMP_PRECOND_END]]2378// CHECK11: omp.precond.end:2379// CHECK11-NEXT: ret void2380//2381//2382// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l982383// CHECK11-SAME: (i32 noundef [[N:%.*]], i32 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {2384// CHECK11-NEXT: entry:2385// CHECK11-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42386// CHECK11-NEXT: [[VLA_ADDR:%.*]] = alloca i32, align 42387// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42388// CHECK11-NEXT: [[N_CASTED:%.*]] = alloca i32, align 42389// CHECK11-NEXT: store i32 [[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 i32, ptr [[VLA_ADDR]], align 42393// CHECK11-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 42394// CHECK11-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 42395// CHECK11-NEXT: store i32 [[TMP2]], ptr [[N_CASTED]], align 42396// CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[N_CASTED]], align 42397// CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l98.omp_outlined, i32 [[TMP3]], i32 [[TMP0]], ptr [[TMP1]])2398// CHECK11-NEXT: ret void2399//2400//2401// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l98.omp_outlined2402// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[N:%.*]], i32 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {2403// CHECK11-NEXT: entry:2404// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42405// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42406// CHECK11-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42407// CHECK11-NEXT: [[VLA_ADDR:%.*]] = alloca i32, align 42408// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42409// CHECK11-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42410// CHECK11-NEXT: [[TMP:%.*]] = alloca i32, align 42411// CHECK11-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42412// CHECK11-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42413// CHECK11-NEXT: [[I:%.*]] = alloca i32, align 42414// CHECK11-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 42415// CHECK11-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 42416// CHECK11-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42417// CHECK11-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42418// CHECK11-NEXT: [[I3:%.*]] = alloca i32, align 42419// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42420// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42421// CHECK11-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42422// CHECK11-NEXT: store i32 [[VLA]], ptr [[VLA_ADDR]], align 42423// CHECK11-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42424// CHECK11-NEXT: [[TMP0:%.*]] = load i32, ptr [[VLA_ADDR]], align 42425// CHECK11-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 42426// CHECK11-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 42427// CHECK11-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 42428// CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42429// CHECK11-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP3]], 02430// CHECK11-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12431// CHECK11-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12432// CHECK11-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42433// CHECK11-NEXT: store i32 0, ptr [[I]], align 42434// CHECK11-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42435// CHECK11-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]2436// CHECK11-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2437// CHECK11: omp.precond.then:2438// CHECK11-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 42439// CHECK11-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42440// CHECK11-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_UB]], align 42441// CHECK11-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42442// CHECK11-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42443// CHECK11-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42444// CHECK11-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 42445// CHECK11-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP7]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2446// CHECK11-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42447// CHECK11-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42448// CHECK11-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]2449// CHECK11-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2450// CHECK11: cond.true:2451// CHECK11-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42452// CHECK11-NEXT: br label [[COND_END:%.*]]2453// CHECK11: cond.false:2454// CHECK11-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42455// CHECK11-NEXT: br label [[COND_END]]2456// CHECK11: cond.end:2457// CHECK11-NEXT: [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]2458// CHECK11-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 42459// CHECK11-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42460// CHECK11-NEXT: store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 42461// CHECK11-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2462// CHECK11: omp.inner.for.cond:2463// CHECK11-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42464// CHECK11-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42465// CHECK11-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]2466// CHECK11-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2467// CHECK11: omp.inner.for.body:2468// CHECK11-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42469// CHECK11-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP15]], 12470// CHECK11-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]2471// CHECK11-NEXT: store i32 [[ADD]], ptr [[I3]], align 42472// CHECK11-NEXT: [[TMP16:%.*]] = load i32, ptr [[I3]], align 42473// CHECK11-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i32 [[TMP16]]2474// CHECK11-NEXT: store i32 0, ptr [[ARRAYIDX]], align 42475// CHECK11-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]2476// CHECK11: omp.body.continue:2477// CHECK11-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2478// CHECK11: omp.inner.for.inc:2479// CHECK11-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42480// CHECK11-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP17]], 12481// CHECK11-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 42482// CHECK11-NEXT: br label [[OMP_INNER_FOR_COND]]2483// CHECK11: omp.inner.for.end:2484// CHECK11-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2485// CHECK11: omp.loop.exit:2486// CHECK11-NEXT: [[TMP18:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42487// CHECK11-NEXT: [[TMP19:%.*]] = load i32, ptr [[TMP18]], align 42488// CHECK11-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP19]])2489// CHECK11-NEXT: br label [[OMP_PRECOND_END]]2490// CHECK11: omp.precond.end:2491// CHECK11-NEXT: ret void2492//2493//2494// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l1022495// CHECK11-SAME: (i32 noundef [[N:%.*]], i32 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR2]] {2496// CHECK11-NEXT: entry:2497// CHECK11-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42498// CHECK11-NEXT: [[VLA_ADDR:%.*]] = alloca i32, align 42499// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42500// CHECK11-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 42501// CHECK11-NEXT: [[N_CASTED:%.*]] = alloca i32, align 42502// CHECK11-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i32, align 42503// CHECK11-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42504// CHECK11-NEXT: store i32 [[VLA]], ptr [[VLA_ADDR]], align 42505// CHECK11-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42506// CHECK11-NEXT: store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 42507// CHECK11-NEXT: [[TMP0:%.*]] = load i32, ptr [[VLA_ADDR]], align 42508// CHECK11-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 42509// CHECK11-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 42510// CHECK11-NEXT: store i32 [[TMP2]], ptr [[N_CASTED]], align 42511// CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[N_CASTED]], align 42512// CHECK11-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 42513// CHECK11-NEXT: store i32 [[TMP4]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 42514// CHECK11-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED]], align 42515// CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 4, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l102.omp_outlined, i32 [[TMP3]], i32 [[TMP0]], ptr [[TMP1]], i32 [[TMP5]])2516// CHECK11-NEXT: ret void2517//2518//2519// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l102.omp_outlined2520// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[N:%.*]], i32 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR2]] {2521// CHECK11-NEXT: entry:2522// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42523// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42524// CHECK11-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42525// CHECK11-NEXT: [[VLA_ADDR:%.*]] = alloca i32, align 42526// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42527// CHECK11-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 42528// CHECK11-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42529// CHECK11-NEXT: [[TMP:%.*]] = alloca i32, align 42530// CHECK11-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42531// CHECK11-NEXT: [[DOTCAPTURE_EXPR_2:%.*]] = alloca i32, align 42532// CHECK11-NEXT: [[I:%.*]] = alloca i32, align 42533// CHECK11-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 42534// CHECK11-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 42535// CHECK11-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42536// CHECK11-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42537// CHECK11-NEXT: [[I4:%.*]] = alloca i32, align 42538// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42539// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42540// CHECK11-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42541// CHECK11-NEXT: store i32 [[VLA]], ptr [[VLA_ADDR]], align 42542// CHECK11-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42543// CHECK11-NEXT: store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 42544// CHECK11-NEXT: [[TMP0:%.*]] = load i32, ptr [[VLA_ADDR]], align 42545// CHECK11-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 42546// CHECK11-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 42547// CHECK11-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_1]], align 42548// CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42549// CHECK11-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP3]], 02550// CHECK11-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12551// CHECK11-NEXT: [[SUB3:%.*]] = sub nsw i32 [[DIV]], 12552// CHECK11-NEXT: store i32 [[SUB3]], ptr [[DOTCAPTURE_EXPR_2]], align 42553// CHECK11-NEXT: store i32 0, ptr [[I]], align 42554// CHECK11-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42555// CHECK11-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]2556// CHECK11-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2557// CHECK11: omp.precond.then:2558// CHECK11-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 42559// CHECK11-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 42560// CHECK11-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_UB]], align 42561// CHECK11-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42562// CHECK11-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42563// CHECK11-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 42564// CHECK11-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42565// CHECK11-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 42566// CHECK11-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP8]], i32 91, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 [[TMP6]])2567// CHECK11-NEXT: br label [[OMP_DISPATCH_COND:%.*]]2568// CHECK11: omp.dispatch.cond:2569// CHECK11-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42570// CHECK11-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 42571// CHECK11-NEXT: [[CMP5:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]2572// CHECK11-NEXT: br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2573// CHECK11: cond.true:2574// CHECK11-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 42575// CHECK11-NEXT: br label [[COND_END:%.*]]2576// CHECK11: cond.false:2577// CHECK11-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42578// CHECK11-NEXT: br label [[COND_END]]2579// CHECK11: cond.end:2580// CHECK11-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]2581// CHECK11-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 42582// CHECK11-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42583// CHECK11-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 42584// CHECK11-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42585// CHECK11-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42586// CHECK11-NEXT: [[CMP6:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]2587// CHECK11-NEXT: br i1 [[CMP6]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]2588// CHECK11: omp.dispatch.body:2589// CHECK11-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2590// CHECK11: omp.inner.for.cond:2591// CHECK11-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12:![0-9]+]]2592// CHECK11-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP12]]2593// CHECK11-NEXT: [[CMP7:%.*]] = icmp sle i32 [[TMP16]], [[TMP17]]2594// CHECK11-NEXT: br i1 [[CMP7]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2595// CHECK11: omp.inner.for.body:2596// CHECK11-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]2597// CHECK11-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP18]], 12598// CHECK11-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]2599// CHECK11-NEXT: store i32 [[ADD]], ptr [[I4]], align 4, !llvm.access.group [[ACC_GRP12]]2600// CHECK11-NEXT: [[TMP19:%.*]] = load i32, ptr [[I4]], align 4, !llvm.access.group [[ACC_GRP12]]2601// CHECK11-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i32 [[TMP19]]2602// CHECK11-NEXT: store i32 0, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP12]]2603// CHECK11-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]2604// CHECK11: omp.body.continue:2605// CHECK11-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2606// CHECK11: omp.inner.for.inc:2607// CHECK11-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]2608// CHECK11-NEXT: [[ADD8:%.*]] = add nsw i32 [[TMP20]], 12609// CHECK11-NEXT: store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]2610// CHECK11-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP13:![0-9]+]]2611// CHECK11: omp.inner.for.end:2612// CHECK11-NEXT: br label [[OMP_DISPATCH_INC:%.*]]2613// CHECK11: omp.dispatch.inc:2614// CHECK11-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42615// CHECK11-NEXT: [[TMP22:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 42616// CHECK11-NEXT: [[ADD9:%.*]] = add nsw i32 [[TMP21]], [[TMP22]]2617// CHECK11-NEXT: store i32 [[ADD9]], ptr [[DOTOMP_LB]], align 42618// CHECK11-NEXT: [[TMP23:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42619// CHECK11-NEXT: [[TMP24:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 42620// CHECK11-NEXT: [[ADD10:%.*]] = add nsw i32 [[TMP23]], [[TMP24]]2621// CHECK11-NEXT: store i32 [[ADD10]], ptr [[DOTOMP_UB]], align 42622// CHECK11-NEXT: br label [[OMP_DISPATCH_COND]]2623// CHECK11: omp.dispatch.end:2624// CHECK11-NEXT: [[TMP25:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42625// CHECK11-NEXT: [[TMP26:%.*]] = load i32, ptr [[TMP25]], align 42626// CHECK11-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP26]])2627// CHECK11-NEXT: br label [[OMP_PRECOND_END]]2628// CHECK11: omp.precond.end:2629// CHECK11-NEXT: ret void2630//2631//2632// CHECK11-LABEL: define {{[^@]+}}@_Z5tmainIiLi10EEiT_2633// CHECK11-SAME: (i32 noundef [[ARGC:%.*]]) #[[ATTR5:[0-9]+]] comdat {2634// CHECK11-NEXT: entry:2635// CHECK11-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 42636// CHECK11-NEXT: [[A:%.*]] = alloca [10 x i32], align 42637// CHECK11-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 42638// CHECK11-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 42639// CHECK11-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 42640// CHECK11-NEXT: [[TMP:%.*]] = alloca i32, align 42641// CHECK11-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 82642// CHECK11-NEXT: [[DOTOFFLOAD_BASEPTRS1:%.*]] = alloca [1 x ptr], align 42643// CHECK11-NEXT: [[DOTOFFLOAD_PTRS2:%.*]] = alloca [1 x ptr], align 42644// CHECK11-NEXT: [[DOTOFFLOAD_MAPPERS3:%.*]] = alloca [1 x ptr], align 42645// CHECK11-NEXT: [[_TMP4:%.*]] = alloca i32, align 42646// CHECK11-NEXT: [[KERNEL_ARGS5:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 82647// CHECK11-NEXT: [[DOTOFFLOAD_BASEPTRS8:%.*]] = alloca [1 x ptr], align 42648// CHECK11-NEXT: [[DOTOFFLOAD_PTRS9:%.*]] = alloca [1 x ptr], align 42649// CHECK11-NEXT: [[DOTOFFLOAD_MAPPERS10:%.*]] = alloca [1 x ptr], align 42650// CHECK11-NEXT: [[_TMP11:%.*]] = alloca i32, align 42651// CHECK11-NEXT: [[KERNEL_ARGS12:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 82652// CHECK11-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 42653// CHECK11-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 02654// CHECK11-NEXT: store ptr [[A]], ptr [[TMP0]], align 42655// CHECK11-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 02656// CHECK11-NEXT: store ptr [[A]], ptr [[TMP1]], align 42657// CHECK11-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 02658// CHECK11-NEXT: store ptr null, ptr [[TMP2]], align 42659// CHECK11-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 02660// CHECK11-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 02661// CHECK11-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 02662// CHECK11-NEXT: store i32 3, ptr [[TMP5]], align 42663// CHECK11-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12664// CHECK11-NEXT: store i32 1, ptr [[TMP6]], align 42665// CHECK11-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 22666// CHECK11-NEXT: store ptr [[TMP3]], ptr [[TMP7]], align 42667// CHECK11-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 32668// CHECK11-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 42669// CHECK11-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 42670// CHECK11-NEXT: store ptr @.offload_sizes.5, ptr [[TMP9]], align 42671// CHECK11-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 52672// CHECK11-NEXT: store ptr @.offload_maptypes.6, ptr [[TMP10]], align 42673// CHECK11-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 62674// CHECK11-NEXT: store ptr null, ptr [[TMP11]], align 42675// CHECK11-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 72676// CHECK11-NEXT: store ptr null, ptr [[TMP12]], align 42677// CHECK11-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 82678// CHECK11-NEXT: store i64 10, ptr [[TMP13]], align 82679// CHECK11-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 92680// CHECK11-NEXT: store i64 0, ptr [[TMP14]], align 82681// CHECK11-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 102682// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 42683// CHECK11-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 112684// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 42685// CHECK11-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 122686// CHECK11-NEXT: store i32 0, ptr [[TMP17]], align 42687// 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__l76.region_id, ptr [[KERNEL_ARGS]])2688// CHECK11-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 02689// CHECK11-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]2690// CHECK11: omp_offload.failed:2691// CHECK11-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l76(ptr [[A]]) #[[ATTR3]]2692// CHECK11-NEXT: br label [[OMP_OFFLOAD_CONT]]2693// CHECK11: omp_offload.cont:2694// CHECK11-NEXT: [[TMP20:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 02695// CHECK11-NEXT: store ptr [[A]], ptr [[TMP20]], align 42696// CHECK11-NEXT: [[TMP21:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 02697// CHECK11-NEXT: store ptr [[A]], ptr [[TMP21]], align 42698// CHECK11-NEXT: [[TMP22:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS3]], i32 0, i32 02699// CHECK11-NEXT: store ptr null, ptr [[TMP22]], align 42700// CHECK11-NEXT: [[TMP23:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 02701// CHECK11-NEXT: [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 02702// CHECK11-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 02703// CHECK11-NEXT: store i32 3, ptr [[TMP25]], align 42704// CHECK11-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 12705// CHECK11-NEXT: store i32 1, ptr [[TMP26]], align 42706// CHECK11-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 22707// CHECK11-NEXT: store ptr [[TMP23]], ptr [[TMP27]], align 42708// CHECK11-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 32709// CHECK11-NEXT: store ptr [[TMP24]], ptr [[TMP28]], align 42710// CHECK11-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 42711// CHECK11-NEXT: store ptr @.offload_sizes.7, ptr [[TMP29]], align 42712// CHECK11-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 52713// CHECK11-NEXT: store ptr @.offload_maptypes.8, ptr [[TMP30]], align 42714// CHECK11-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 62715// CHECK11-NEXT: store ptr null, ptr [[TMP31]], align 42716// CHECK11-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 72717// CHECK11-NEXT: store ptr null, ptr [[TMP32]], align 42718// CHECK11-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 82719// CHECK11-NEXT: store i64 10, ptr [[TMP33]], align 82720// CHECK11-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 92721// CHECK11-NEXT: store i64 0, ptr [[TMP34]], align 82722// CHECK11-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 102723// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP35]], align 42724// CHECK11-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 112725// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP36]], align 42726// CHECK11-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 122727// CHECK11-NEXT: store i32 0, ptr [[TMP37]], align 42728// 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__l80.region_id, ptr [[KERNEL_ARGS5]])2729// CHECK11-NEXT: [[TMP39:%.*]] = icmp ne i32 [[TMP38]], 02730// CHECK11-NEXT: br i1 [[TMP39]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]]2731// CHECK11: omp_offload.failed6:2732// CHECK11-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l80(ptr [[A]]) #[[ATTR3]]2733// CHECK11-NEXT: br label [[OMP_OFFLOAD_CONT7]]2734// CHECK11: omp_offload.cont7:2735// CHECK11-NEXT: [[TMP40:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 02736// CHECK11-NEXT: store ptr [[A]], ptr [[TMP40]], align 42737// CHECK11-NEXT: [[TMP41:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 02738// CHECK11-NEXT: store ptr [[A]], ptr [[TMP41]], align 42739// CHECK11-NEXT: [[TMP42:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS10]], i32 0, i32 02740// CHECK11-NEXT: store ptr null, ptr [[TMP42]], align 42741// CHECK11-NEXT: [[TMP43:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 02742// CHECK11-NEXT: [[TMP44:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS9]], i32 0, i32 02743// CHECK11-NEXT: [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 02744// CHECK11-NEXT: store i32 3, ptr [[TMP45]], align 42745// CHECK11-NEXT: [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 12746// CHECK11-NEXT: store i32 1, ptr [[TMP46]], align 42747// CHECK11-NEXT: [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 22748// CHECK11-NEXT: store ptr [[TMP43]], ptr [[TMP47]], align 42749// CHECK11-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 32750// CHECK11-NEXT: store ptr [[TMP44]], ptr [[TMP48]], align 42751// CHECK11-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 42752// CHECK11-NEXT: store ptr @.offload_sizes.9, ptr [[TMP49]], align 42753// CHECK11-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 52754// CHECK11-NEXT: store ptr @.offload_maptypes.10, ptr [[TMP50]], align 42755// CHECK11-NEXT: [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 62756// CHECK11-NEXT: store ptr null, ptr [[TMP51]], align 42757// CHECK11-NEXT: [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 72758// CHECK11-NEXT: store ptr null, ptr [[TMP52]], align 42759// CHECK11-NEXT: [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 82760// CHECK11-NEXT: store i64 10, ptr [[TMP53]], align 82761// CHECK11-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 92762// CHECK11-NEXT: store i64 0, ptr [[TMP54]], align 82763// CHECK11-NEXT: [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 102764// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP55]], align 42765// CHECK11-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 112766// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP56]], align 42767// CHECK11-NEXT: [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 122768// CHECK11-NEXT: store i32 0, ptr [[TMP57]], align 42769// 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__l84.region_id, ptr [[KERNEL_ARGS12]])2770// CHECK11-NEXT: [[TMP59:%.*]] = icmp ne i32 [[TMP58]], 02771// CHECK11-NEXT: br i1 [[TMP59]], label [[OMP_OFFLOAD_FAILED13:%.*]], label [[OMP_OFFLOAD_CONT14:%.*]]2772// CHECK11: omp_offload.failed13:2773// CHECK11-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l84(ptr [[A]]) #[[ATTR3]]2774// CHECK11-NEXT: br label [[OMP_OFFLOAD_CONT14]]2775// CHECK11: omp_offload.cont14:2776// CHECK11-NEXT: ret i32 02777//2778//2779// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l762780// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {2781// CHECK11-NEXT: entry:2782// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42783// CHECK11-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42784// CHECK11-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42785// CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l76.omp_outlined, ptr [[TMP0]])2786// CHECK11-NEXT: ret void2787//2788//2789// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l76.omp_outlined2790// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {2791// CHECK11-NEXT: entry:2792// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42793// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42794// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42795// CHECK11-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42796// CHECK11-NEXT: [[TMP:%.*]] = alloca i32, align 42797// CHECK11-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 42798// CHECK11-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 42799// CHECK11-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42800// CHECK11-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42801// CHECK11-NEXT: [[I:%.*]] = alloca i32, align 42802// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42803// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42804// CHECK11-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42805// CHECK11-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42806// CHECK11-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 42807// CHECK11-NEXT: store i32 9, ptr [[DOTOMP_UB]], align 42808// CHECK11-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42809// CHECK11-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42810// CHECK11-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42811// CHECK11-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 42812// 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)2813// CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42814// CHECK11-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 92815// CHECK11-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2816// CHECK11: cond.true:2817// CHECK11-NEXT: br label [[COND_END:%.*]]2818// CHECK11: cond.false:2819// CHECK11-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42820// CHECK11-NEXT: br label [[COND_END]]2821// CHECK11: cond.end:2822// CHECK11-NEXT: [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]2823// CHECK11-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 42824// CHECK11-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42825// CHECK11-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 42826// CHECK11-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2827// CHECK11: omp.inner.for.cond:2828// CHECK11-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42829// CHECK11-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42830// CHECK11-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]2831// CHECK11-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2832// CHECK11: omp.inner.for.body:2833// CHECK11-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42834// CHECK11-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 12835// CHECK11-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]2836// CHECK11-NEXT: store i32 [[ADD]], ptr [[I]], align 42837// CHECK11-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 42838// CHECK11-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP9]]2839// CHECK11-NEXT: store i32 0, ptr [[ARRAYIDX]], align 42840// CHECK11-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]2841// CHECK11: omp.body.continue:2842// CHECK11-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2843// CHECK11: omp.inner.for.inc:2844// CHECK11-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42845// CHECK11-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP10]], 12846// CHECK11-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 42847// CHECK11-NEXT: br label [[OMP_INNER_FOR_COND]]2848// CHECK11: omp.inner.for.end:2849// CHECK11-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2850// CHECK11: omp.loop.exit:2851// CHECK11-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])2852// CHECK11-NEXT: ret void2853//2854//2855// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l802856// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {2857// CHECK11-NEXT: entry:2858// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42859// CHECK11-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42860// CHECK11-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42861// CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l80.omp_outlined, ptr [[TMP0]])2862// CHECK11-NEXT: ret void2863//2864//2865// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l80.omp_outlined2866// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {2867// CHECK11-NEXT: entry:2868// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42869// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42870// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42871// CHECK11-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42872// CHECK11-NEXT: [[TMP:%.*]] = alloca i32, align 42873// CHECK11-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 42874// CHECK11-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 42875// CHECK11-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42876// CHECK11-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42877// CHECK11-NEXT: [[I:%.*]] = alloca i32, align 42878// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42879// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42880// CHECK11-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42881// CHECK11-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42882// CHECK11-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 42883// CHECK11-NEXT: store i32 9, ptr [[DOTOMP_UB]], align 42884// CHECK11-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42885// CHECK11-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42886// CHECK11-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42887// CHECK11-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 42888// 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)2889// CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42890// CHECK11-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 92891// CHECK11-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2892// CHECK11: cond.true:2893// CHECK11-NEXT: br label [[COND_END:%.*]]2894// CHECK11: cond.false:2895// CHECK11-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42896// CHECK11-NEXT: br label [[COND_END]]2897// CHECK11: cond.end:2898// CHECK11-NEXT: [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]2899// CHECK11-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 42900// CHECK11-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42901// CHECK11-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 42902// CHECK11-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2903// CHECK11: omp.inner.for.cond:2904// CHECK11-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42905// CHECK11-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42906// CHECK11-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]2907// CHECK11-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2908// CHECK11: omp.inner.for.body:2909// CHECK11-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42910// CHECK11-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 12911// CHECK11-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]2912// CHECK11-NEXT: store i32 [[ADD]], ptr [[I]], align 42913// CHECK11-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 42914// CHECK11-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP9]]2915// CHECK11-NEXT: store i32 0, ptr [[ARRAYIDX]], align 42916// CHECK11-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]2917// CHECK11: omp.body.continue:2918// CHECK11-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2919// CHECK11: omp.inner.for.inc:2920// CHECK11-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42921// CHECK11-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP10]], 12922// CHECK11-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 42923// CHECK11-NEXT: br label [[OMP_INNER_FOR_COND]]2924// CHECK11: omp.inner.for.end:2925// CHECK11-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2926// CHECK11: omp.loop.exit:2927// CHECK11-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])2928// CHECK11-NEXT: ret void2929//2930//2931// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l842932// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {2933// CHECK11-NEXT: entry:2934// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42935// CHECK11-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42936// CHECK11-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42937// 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]])2938// CHECK11-NEXT: ret void2939//2940//2941// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l84.omp_outlined2942// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {2943// CHECK11-NEXT: entry:2944// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42945// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42946// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42947// CHECK11-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42948// CHECK11-NEXT: [[TMP:%.*]] = alloca i32, align 42949// CHECK11-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 42950// CHECK11-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 42951// CHECK11-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42952// CHECK11-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42953// CHECK11-NEXT: [[I:%.*]] = alloca i32, align 42954// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42955// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42956// CHECK11-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42957// CHECK11-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42958// CHECK11-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 42959// CHECK11-NEXT: store i32 9, ptr [[DOTOMP_UB]], align 42960// CHECK11-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42961// CHECK11-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42962// CHECK11-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42963// CHECK11-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 42964// 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)2965// CHECK11-NEXT: br label [[OMP_DISPATCH_COND:%.*]]2966// CHECK11: omp.dispatch.cond:2967// CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42968// CHECK11-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 92969// CHECK11-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2970// CHECK11: cond.true:2971// CHECK11-NEXT: br label [[COND_END:%.*]]2972// CHECK11: cond.false:2973// CHECK11-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42974// CHECK11-NEXT: br label [[COND_END]]2975// CHECK11: cond.end:2976// CHECK11-NEXT: [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]2977// CHECK11-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 42978// CHECK11-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42979// CHECK11-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 42980// CHECK11-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42981// CHECK11-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42982// CHECK11-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]2983// CHECK11-NEXT: br i1 [[CMP1]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]2984// CHECK11: omp.dispatch.body:2985// CHECK11-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2986// CHECK11: omp.inner.for.cond:2987// CHECK11-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP15:![0-9]+]]2988// CHECK11-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP15]]2989// CHECK11-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP8]], [[TMP9]]2990// CHECK11-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2991// CHECK11: omp.inner.for.body:2992// CHECK11-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP15]]2993// CHECK11-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP10]], 12994// CHECK11-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]2995// CHECK11-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP15]]2996// CHECK11-NEXT: [[TMP11:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP15]]2997// CHECK11-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP11]]2998// CHECK11-NEXT: store i32 0, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP15]]2999// CHECK11-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]3000// CHECK11: omp.body.continue:3001// CHECK11-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]3002// CHECK11: omp.inner.for.inc:3003// CHECK11-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP15]]3004// CHECK11-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP12]], 13005// CHECK11-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP15]]3006// CHECK11-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP16:![0-9]+]]3007// CHECK11: omp.inner.for.end:3008// CHECK11-NEXT: br label [[OMP_DISPATCH_INC:%.*]]3009// CHECK11: omp.dispatch.inc:3010// CHECK11-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 43011// CHECK11-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 43012// CHECK11-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP13]], [[TMP14]]3013// CHECK11-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_LB]], align 43014// CHECK11-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 43015// CHECK11-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 43016// CHECK11-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP15]], [[TMP16]]3017// CHECK11-NEXT: store i32 [[ADD5]], ptr [[DOTOMP_UB]], align 43018// CHECK11-NEXT: br label [[OMP_DISPATCH_COND]]3019// CHECK11: omp.dispatch.end:3020// CHECK11-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])3021// CHECK11-NEXT: ret void3022//3023