2434 lines · cpp
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _2// expected-no-diagnostics3#ifndef HEADER4#define HEADER5// Test host codegen.6// RUN: %clang_cc1 -DCK1 -verify -Wno-vla -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK17// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s8// 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=CHECK19// 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=CHECK310// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s11// 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=CHECK312 13// 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}}"14// 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 %s15// 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}}"16// 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}}"17// 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 %s18// 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}}"19#ifdef CK120 21int a[100];22 23int teams_argument_global(int n){24 int te = n / 128;25 int th = 128;26 // discard n_addr27 28 #pragma omp target29 #pragma omp teams loop num_teams(te), thread_limit(th)30 for(int i = 0; i < n; i++) {31 a[i] = 0;32 }33 34 #pragma omp target35 {{{36 #pragma omp teams loop37 for(int i = 0; i < n; i++) {38 a[i] = 0;39 }40 }}}41 42 // outlined target regions43 44 45 46 47 return a[0];48}49 50#endif // CK151 52// Test host codegen.53// 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=CHECK954// RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s55// 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=CHECK956// 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=CHECK1157// RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s58// 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=CHECK1159 60// 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}}"61// 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 %s62// 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}}"63// 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}}"64// 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 %s65// 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}}"66#ifdef CK267 68int teams_local_arg(void) {69 int n = 100;70 int a[n];71 72 #pragma omp target73 #pragma omp teams loop74 for(int i = 0; i < n; i++) {75 a[i] = 0;76 }77 78 // outlined target region79 80 81 return a[0];82}83#endif // CK284 85// Test host codegen.86// RUN: %clang_cc1 -DCK3 -verify -Wno-vla -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK1787// RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s88// RUN: %clang_cc1 -DCK3 -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=CHECK1789// RUN: %clang_cc1 -DCK3 -verify -Wno-vla -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK1990// RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s91// RUN: %clang_cc1 -DCK3 -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=CHECK1992 93// RUN: %clang_cc1 -DCK3 -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}}"94// RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s95// RUN: %clang_cc1 -DCK3 -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}}"96// RUN: %clang_cc1 -DCK3 -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}}"97// RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s98// RUN: %clang_cc1 -DCK3 -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}}"99#ifdef CK3100 101 102template <typename T, int X, long long Y>103struct SS{104 T a[X];105 float b;106 int foo(void) {107 108 #pragma omp target109 #pragma omp teams loop110 for(int i = 0; i < X; i++) {111 a[i] = (T)0;112 }113 114 // outlined target region115 116 117 return a[0];118 }119};120 121int teams_template_struct(void) {122 SS<int, 123, 456> V;123 return V.foo();124 125}126#endif // CK3127 128// Test host codegen.129// RUN: %clang_cc1 -DCK4 -verify -Wno-vla -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK25130// RUN: %clang_cc1 -DCK4 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s131// RUN: %clang_cc1 -DCK4 -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=CHECK25132// RUN: %clang_cc1 -DCK4 -verify -Wno-vla -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK27133// RUN: %clang_cc1 -DCK4 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s134// RUN: %clang_cc1 -DCK4 -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=CHECK27135 136// RUN: %clang_cc1 -DCK4 -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}}"137// RUN: %clang_cc1 -DCK4 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s138// RUN: %clang_cc1 -DCK4 -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}}"139// RUN: %clang_cc1 -DCK4 -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}}"140// RUN: %clang_cc1 -DCK4 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s141// RUN: %clang_cc1 -DCK4 -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}}"142 143#ifdef CK4144 145template <typename T, int n>146int tmain(T argc) {147 T a[n];148 int te = n/128;149 int th = 128;150#pragma omp target151#pragma omp teams loop num_teams(te) thread_limit(th)152 for(int i = 0; i < n; i++) {153 a[i] = (T)0;154 }155 return 0;156}157 158int main (int argc, char **argv) {159 int n = 100;160 int a[n];161#pragma omp target162#pragma omp teams loop163 for(int i = 0; i < n; i++) {164 a[i] = 0;165 }166 return tmain<int, 10>(argc);167}168 169 170 171 172 173 174 175#endif // CK4176#endif177// CHECK1-LABEL: define {{[^@]+}}@_Z21teams_argument_globali178// CHECK1-SAME: (i32 noundef signext [[N:%.*]]) #[[ATTR0:[0-9]+]] {179// CHECK1-NEXT: entry:180// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i32, align 4181// CHECK1-NEXT: [[TE:%.*]] = alloca i32, align 4182// CHECK1-NEXT: [[TH:%.*]] = alloca i32, align 4183// CHECK1-NEXT: [[TE_CASTED:%.*]] = alloca i64, align 8184// CHECK1-NEXT: [[TH_CASTED:%.*]] = alloca i64, align 8185// CHECK1-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8186// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [4 x ptr], align 8187// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [4 x ptr], align 8188// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [4 x ptr], align 8189// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4190// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4191// CHECK1-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4192// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8193// CHECK1-NEXT: [[N_CASTED4:%.*]] = alloca i64, align 8194// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS5:%.*]] = alloca [2 x ptr], align 8195// CHECK1-NEXT: [[DOTOFFLOAD_PTRS6:%.*]] = alloca [2 x ptr], align 8196// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS7:%.*]] = alloca [2 x ptr], align 8197// CHECK1-NEXT: [[_TMP8:%.*]] = alloca i32, align 4198// CHECK1-NEXT: [[DOTCAPTURE_EXPR_9:%.*]] = alloca i32, align 4199// CHECK1-NEXT: [[DOTCAPTURE_EXPR_10:%.*]] = alloca i32, align 4200// CHECK1-NEXT: [[KERNEL_ARGS15:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8201// CHECK1-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 4202// CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4203// CHECK1-NEXT: [[DIV:%.*]] = sdiv i32 [[TMP0]], 128204// CHECK1-NEXT: store i32 [[DIV]], ptr [[TE]], align 4205// CHECK1-NEXT: store i32 128, ptr [[TH]], align 4206// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TE]], align 4207// CHECK1-NEXT: store i32 [[TMP1]], ptr [[TE_CASTED]], align 4208// CHECK1-NEXT: [[TMP2:%.*]] = load i64, ptr [[TE_CASTED]], align 8209// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[TH]], align 4210// CHECK1-NEXT: store i32 [[TMP3]], ptr [[TH_CASTED]], align 4211// CHECK1-NEXT: [[TMP4:%.*]] = load i64, ptr [[TH_CASTED]], align 8212// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[N_ADDR]], align 4213// CHECK1-NEXT: store i32 [[TMP5]], ptr [[N_CASTED]], align 4214// CHECK1-NEXT: [[TMP6:%.*]] = load i64, ptr [[N_CASTED]], align 8215// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0216// CHECK1-NEXT: store i64 [[TMP2]], ptr [[TMP7]], align 8217// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0218// CHECK1-NEXT: store i64 [[TMP2]], ptr [[TMP8]], align 8219// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0220// CHECK1-NEXT: store ptr null, ptr [[TMP9]], align 8221// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1222// CHECK1-NEXT: store i64 [[TMP4]], ptr [[TMP10]], align 8223// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1224// CHECK1-NEXT: store i64 [[TMP4]], ptr [[TMP11]], align 8225// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1226// CHECK1-NEXT: store ptr null, ptr [[TMP12]], align 8227// CHECK1-NEXT: [[TMP13:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2228// CHECK1-NEXT: store i64 [[TMP6]], ptr [[TMP13]], align 8229// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2230// CHECK1-NEXT: store i64 [[TMP6]], ptr [[TMP14]], align 8231// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2232// CHECK1-NEXT: store ptr null, ptr [[TMP15]], align 8233// CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3234// CHECK1-NEXT: store ptr @a, ptr [[TMP16]], align 8235// CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3236// CHECK1-NEXT: store ptr @a, ptr [[TMP17]], align 8237// CHECK1-NEXT: [[TMP18:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3238// CHECK1-NEXT: store ptr null, ptr [[TMP18]], align 8239// CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0240// CHECK1-NEXT: [[TMP20:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0241// CHECK1-NEXT: [[TMP21:%.*]] = load i32, ptr [[TE]], align 4242// CHECK1-NEXT: [[TMP22:%.*]] = load i32, ptr [[TH]], align 4243// CHECK1-NEXT: [[TMP23:%.*]] = load i32, ptr [[N_ADDR]], align 4244// CHECK1-NEXT: store i32 [[TMP23]], ptr [[DOTCAPTURE_EXPR_]], align 4245// CHECK1-NEXT: [[TMP24:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4246// CHECK1-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP24]], 0247// CHECK1-NEXT: [[DIV2:%.*]] = sdiv i32 [[SUB]], 1248// CHECK1-NEXT: [[SUB3:%.*]] = sub nsw i32 [[DIV2]], 1249// CHECK1-NEXT: store i32 [[SUB3]], ptr [[DOTCAPTURE_EXPR_1]], align 4250// CHECK1-NEXT: [[TMP25:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4251// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP25]], 1252// CHECK1-NEXT: [[TMP26:%.*]] = zext i32 [[ADD]] to i64253// CHECK1-NEXT: [[TMP27:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP21]], 0254// CHECK1-NEXT: [[TMP28:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP22]], 0255// CHECK1-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0256// CHECK1-NEXT: store i32 3, ptr [[TMP29]], align 4257// CHECK1-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1258// CHECK1-NEXT: store i32 4, ptr [[TMP30]], align 4259// CHECK1-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2260// CHECK1-NEXT: store ptr [[TMP19]], ptr [[TMP31]], align 8261// CHECK1-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3262// CHECK1-NEXT: store ptr [[TMP20]], ptr [[TMP32]], align 8263// CHECK1-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4264// CHECK1-NEXT: store ptr @.offload_sizes, ptr [[TMP33]], align 8265// CHECK1-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5266// CHECK1-NEXT: store ptr @.offload_maptypes, ptr [[TMP34]], align 8267// CHECK1-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6268// CHECK1-NEXT: store ptr null, ptr [[TMP35]], align 8269// CHECK1-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7270// CHECK1-NEXT: store ptr null, ptr [[TMP36]], align 8271// CHECK1-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8272// CHECK1-NEXT: store i64 [[TMP26]], ptr [[TMP37]], align 8273// CHECK1-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9274// CHECK1-NEXT: store i64 0, ptr [[TMP38]], align 8275// CHECK1-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10276// CHECK1-NEXT: store [3 x i32] [[TMP27]], ptr [[TMP39]], align 4277// CHECK1-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11278// CHECK1-NEXT: store [3 x i32] [[TMP28]], ptr [[TMP40]], align 4279// CHECK1-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12280// CHECK1-NEXT: store i32 0, ptr [[TMP41]], align 4281// CHECK1-NEXT: [[TMP42:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2:[0-9]+]], i64 -1, i32 [[TMP21]], i32 [[TMP22]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l28.region_id, ptr [[KERNEL_ARGS]])282// CHECK1-NEXT: [[TMP43:%.*]] = icmp ne i32 [[TMP42]], 0283// CHECK1-NEXT: br i1 [[TMP43]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]284// CHECK1: omp_offload.failed:285// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l28(i64 [[TMP2]], i64 [[TMP4]], i64 [[TMP6]], ptr @a) #[[ATTR2:[0-9]+]]286// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]287// CHECK1: omp_offload.cont:288// CHECK1-NEXT: [[TMP44:%.*]] = load i32, ptr [[N_ADDR]], align 4289// CHECK1-NEXT: store i32 [[TMP44]], ptr [[N_CASTED4]], align 4290// CHECK1-NEXT: [[TMP45:%.*]] = load i64, ptr [[N_CASTED4]], align 8291// CHECK1-NEXT: [[TMP46:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0292// CHECK1-NEXT: store i64 [[TMP45]], ptr [[TMP46]], align 8293// CHECK1-NEXT: [[TMP47:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS6]], i32 0, i32 0294// CHECK1-NEXT: store i64 [[TMP45]], ptr [[TMP47]], align 8295// CHECK1-NEXT: [[TMP48:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS7]], i64 0, i64 0296// CHECK1-NEXT: store ptr null, ptr [[TMP48]], align 8297// CHECK1-NEXT: [[TMP49:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 1298// CHECK1-NEXT: store ptr @a, ptr [[TMP49]], align 8299// CHECK1-NEXT: [[TMP50:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS6]], i32 0, i32 1300// CHECK1-NEXT: store ptr @a, ptr [[TMP50]], align 8301// CHECK1-NEXT: [[TMP51:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS7]], i64 0, i64 1302// CHECK1-NEXT: store ptr null, ptr [[TMP51]], align 8303// CHECK1-NEXT: [[TMP52:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0304// CHECK1-NEXT: [[TMP53:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS6]], i32 0, i32 0305// CHECK1-NEXT: [[TMP54:%.*]] = load i32, ptr [[N_ADDR]], align 4306// CHECK1-NEXT: store i32 [[TMP54]], ptr [[DOTCAPTURE_EXPR_9]], align 4307// CHECK1-NEXT: [[TMP55:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_9]], align 4308// CHECK1-NEXT: [[SUB11:%.*]] = sub nsw i32 [[TMP55]], 0309// CHECK1-NEXT: [[DIV12:%.*]] = sdiv i32 [[SUB11]], 1310// CHECK1-NEXT: [[SUB13:%.*]] = sub nsw i32 [[DIV12]], 1311// CHECK1-NEXT: store i32 [[SUB13]], ptr [[DOTCAPTURE_EXPR_10]], align 4312// CHECK1-NEXT: [[TMP56:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_10]], align 4313// CHECK1-NEXT: [[ADD14:%.*]] = add nsw i32 [[TMP56]], 1314// CHECK1-NEXT: [[TMP57:%.*]] = zext i32 [[ADD14]] to i64315// CHECK1-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 0316// CHECK1-NEXT: store i32 3, ptr [[TMP58]], align 4317// CHECK1-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 1318// CHECK1-NEXT: store i32 2, ptr [[TMP59]], align 4319// CHECK1-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 2320// CHECK1-NEXT: store ptr [[TMP52]], ptr [[TMP60]], align 8321// CHECK1-NEXT: [[TMP61:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 3322// CHECK1-NEXT: store ptr [[TMP53]], ptr [[TMP61]], align 8323// CHECK1-NEXT: [[TMP62:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 4324// CHECK1-NEXT: store ptr @.offload_sizes.1, ptr [[TMP62]], align 8325// CHECK1-NEXT: [[TMP63:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 5326// CHECK1-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP63]], align 8327// CHECK1-NEXT: [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 6328// CHECK1-NEXT: store ptr null, ptr [[TMP64]], align 8329// CHECK1-NEXT: [[TMP65:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 7330// CHECK1-NEXT: store ptr null, ptr [[TMP65]], align 8331// CHECK1-NEXT: [[TMP66:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 8332// CHECK1-NEXT: store i64 [[TMP57]], ptr [[TMP66]], align 8333// CHECK1-NEXT: [[TMP67:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 9334// CHECK1-NEXT: store i64 0, ptr [[TMP67]], align 8335// CHECK1-NEXT: [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 10336// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP68]], align 4337// CHECK1-NEXT: [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 11338// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP69]], align 4339// CHECK1-NEXT: [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 12340// CHECK1-NEXT: store i32 0, ptr [[TMP70]], align 4341// CHECK1-NEXT: [[TMP71:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l34.region_id, ptr [[KERNEL_ARGS15]])342// CHECK1-NEXT: [[TMP72:%.*]] = icmp ne i32 [[TMP71]], 0343// CHECK1-NEXT: br i1 [[TMP72]], label [[OMP_OFFLOAD_FAILED16:%.*]], label [[OMP_OFFLOAD_CONT17:%.*]]344// CHECK1: omp_offload.failed16:345// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l34(i64 [[TMP45]], ptr @a) #[[ATTR2]]346// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT17]]347// CHECK1: omp_offload.cont17:348// CHECK1-NEXT: [[TMP73:%.*]] = load i32, ptr @a, align 4349// CHECK1-NEXT: ret i32 [[TMP73]]350//351//352// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l28353// CHECK1-SAME: (i64 noundef [[TE:%.*]], i64 noundef [[TH:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(400) [[A:%.*]]) #[[ATTR1:[0-9]+]] {354// CHECK1-NEXT: entry:355// CHECK1-NEXT: [[TE_ADDR:%.*]] = alloca i64, align 8356// CHECK1-NEXT: [[TH_ADDR:%.*]] = alloca i64, align 8357// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8358// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8359// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB2]])360// CHECK1-NEXT: store i64 [[TE]], ptr [[TE_ADDR]], align 8361// CHECK1-NEXT: store i64 [[TH]], ptr [[TH_ADDR]], align 8362// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8363// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8364// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 8365// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[TE_ADDR]], align 4366// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[TH_ADDR]], align 4367// CHECK1-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB2]], i32 [[TMP0]], i32 [[TMP2]], i32 [[TMP3]])368// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l28.omp_outlined, ptr [[N_ADDR]], ptr [[TMP1]])369// CHECK1-NEXT: ret void370//371//372// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l28.omp_outlined373// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(400) [[A:%.*]]) #[[ATTR1]] {374// CHECK1-NEXT: entry:375// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8376// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8377// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca ptr, align 8378// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8379// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4380// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4381// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4382// CHECK1-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4383// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4384// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4385// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4386// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4387// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4388// CHECK1-NEXT: [[I3:%.*]] = alloca i32, align 4389// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8390// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8391// CHECK1-NEXT: store ptr [[N]], ptr [[N_ADDR]], align 8392// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8393// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 8394// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 8395// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP0]], align 4396// CHECK1-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4397// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4398// CHECK1-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0399// CHECK1-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1400// CHECK1-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1401// CHECK1-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4402// CHECK1-NEXT: store i32 0, ptr [[I]], align 4403// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4404// CHECK1-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]405// CHECK1-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]406// CHECK1: omp.precond.then:407// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4408// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4409// CHECK1-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_COMB_UB]], align 4410// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4411// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4412// CHECK1-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8413// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 4414// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP7]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)415// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4416// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4417// CHECK1-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]418// CHECK1-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]419// CHECK1: cond.true:420// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4421// CHECK1-NEXT: br label [[COND_END:%.*]]422// CHECK1: cond.false:423// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4424// CHECK1-NEXT: br label [[COND_END]]425// CHECK1: cond.end:426// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]427// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4428// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4429// CHECK1-NEXT: store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 4430// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]431// CHECK1: omp.inner.for.cond:432// CHECK1-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4433// CHECK1-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4434// CHECK1-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]435// CHECK1-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]436// CHECK1: omp.inner.for.body:437// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4438// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP15]], 1439// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]440// CHECK1-NEXT: store i32 [[ADD]], ptr [[I3]], align 4441// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[I3]], align 4442// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP16]] to i64443// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [100 x i32], ptr [[TMP1]], i64 0, i64 [[IDXPROM]]444// CHECK1-NEXT: store i32 0, ptr [[ARRAYIDX]], align 4445// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]446// CHECK1: omp.body.continue:447// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]448// CHECK1: omp.inner.for.inc:449// CHECK1-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4450// CHECK1-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP17]], 1451// CHECK1-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4452// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]453// CHECK1: omp.inner.for.end:454// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]455// CHECK1: omp.loop.exit:456// CHECK1-NEXT: [[TMP18:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8457// CHECK1-NEXT: [[TMP19:%.*]] = load i32, ptr [[TMP18]], align 4458// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP19]])459// CHECK1-NEXT: br label [[OMP_PRECOND_END]]460// CHECK1: omp.precond.end:461// CHECK1-NEXT: ret void462//463//464// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l34465// CHECK1-SAME: (i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(400) [[A:%.*]]) #[[ATTR1]] {466// CHECK1-NEXT: entry:467// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8468// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8469// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8470// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8471// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8472// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l34.omp_outlined, ptr [[N_ADDR]], ptr [[TMP0]])473// CHECK1-NEXT: ret void474//475//476// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l34.omp_outlined477// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(400) [[A:%.*]]) #[[ATTR1]] {478// CHECK1-NEXT: entry:479// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8480// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8481// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca ptr, align 8482// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8483// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4484// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4485// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4486// CHECK1-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4487// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4488// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4489// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4490// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4491// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4492// CHECK1-NEXT: [[I3:%.*]] = alloca i32, align 4493// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8494// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8495// CHECK1-NEXT: store ptr [[N]], ptr [[N_ADDR]], align 8496// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8497// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 8498// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 8499// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP0]], align 4500// CHECK1-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4501// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4502// CHECK1-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0503// CHECK1-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1504// CHECK1-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1505// CHECK1-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4506// CHECK1-NEXT: store i32 0, ptr [[I]], align 4507// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4508// CHECK1-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]509// CHECK1-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]510// CHECK1: omp.precond.then:511// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4512// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4513// CHECK1-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_COMB_UB]], align 4514// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4515// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4516// CHECK1-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8517// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 4518// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP7]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)519// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4520// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4521// CHECK1-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]522// CHECK1-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]523// CHECK1: cond.true:524// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4525// CHECK1-NEXT: br label [[COND_END:%.*]]526// CHECK1: cond.false:527// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4528// CHECK1-NEXT: br label [[COND_END]]529// CHECK1: cond.end:530// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]531// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4532// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4533// CHECK1-NEXT: store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 4534// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]535// CHECK1: omp.inner.for.cond:536// CHECK1-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4537// CHECK1-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4538// CHECK1-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]539// CHECK1-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]540// CHECK1: omp.inner.for.body:541// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4542// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP15]], 1543// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]544// CHECK1-NEXT: store i32 [[ADD]], ptr [[I3]], align 4545// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[I3]], align 4546// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP16]] to i64547// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [100 x i32], ptr [[TMP1]], i64 0, i64 [[IDXPROM]]548// CHECK1-NEXT: store i32 0, ptr [[ARRAYIDX]], align 4549// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]550// CHECK1: omp.body.continue:551// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]552// CHECK1: omp.inner.for.inc:553// CHECK1-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4554// CHECK1-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP17]], 1555// CHECK1-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4556// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]557// CHECK1: omp.inner.for.end:558// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]559// CHECK1: omp.loop.exit:560// CHECK1-NEXT: [[TMP18:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8561// CHECK1-NEXT: [[TMP19:%.*]] = load i32, ptr [[TMP18]], align 4562// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP19]])563// CHECK1-NEXT: br label [[OMP_PRECOND_END]]564// CHECK1: omp.precond.end:565// CHECK1-NEXT: ret void566//567//568// CHECK3-LABEL: define {{[^@]+}}@_Z21teams_argument_globali569// CHECK3-SAME: (i32 noundef [[N:%.*]]) #[[ATTR0:[0-9]+]] {570// CHECK3-NEXT: entry:571// CHECK3-NEXT: [[N_ADDR:%.*]] = alloca i32, align 4572// CHECK3-NEXT: [[TE:%.*]] = alloca i32, align 4573// CHECK3-NEXT: [[TH:%.*]] = alloca i32, align 4574// CHECK3-NEXT: [[TE_CASTED:%.*]] = alloca i32, align 4575// CHECK3-NEXT: [[TH_CASTED:%.*]] = alloca i32, align 4576// CHECK3-NEXT: [[N_CASTED:%.*]] = alloca i32, align 4577// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [4 x ptr], align 4578// CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [4 x ptr], align 4579// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [4 x ptr], align 4580// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4581// CHECK3-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4582// CHECK3-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4583// CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8584// CHECK3-NEXT: [[N_CASTED4:%.*]] = alloca i32, align 4585// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS5:%.*]] = alloca [2 x ptr], align 4586// CHECK3-NEXT: [[DOTOFFLOAD_PTRS6:%.*]] = alloca [2 x ptr], align 4587// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS7:%.*]] = alloca [2 x ptr], align 4588// CHECK3-NEXT: [[_TMP8:%.*]] = alloca i32, align 4589// CHECK3-NEXT: [[DOTCAPTURE_EXPR_9:%.*]] = alloca i32, align 4590// CHECK3-NEXT: [[DOTCAPTURE_EXPR_10:%.*]] = alloca i32, align 4591// CHECK3-NEXT: [[KERNEL_ARGS15:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8592// CHECK3-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 4593// CHECK3-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4594// CHECK3-NEXT: [[DIV:%.*]] = sdiv i32 [[TMP0]], 128595// CHECK3-NEXT: store i32 [[DIV]], ptr [[TE]], align 4596// CHECK3-NEXT: store i32 128, ptr [[TH]], align 4597// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[TE]], align 4598// CHECK3-NEXT: store i32 [[TMP1]], ptr [[TE_CASTED]], align 4599// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[TE_CASTED]], align 4600// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[TH]], align 4601// CHECK3-NEXT: store i32 [[TMP3]], ptr [[TH_CASTED]], align 4602// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[TH_CASTED]], align 4603// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[N_ADDR]], align 4604// CHECK3-NEXT: store i32 [[TMP5]], ptr [[N_CASTED]], align 4605// CHECK3-NEXT: [[TMP6:%.*]] = load i32, ptr [[N_CASTED]], align 4606// CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0607// CHECK3-NEXT: store i32 [[TMP2]], ptr [[TMP7]], align 4608// CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0609// CHECK3-NEXT: store i32 [[TMP2]], ptr [[TMP8]], align 4610// CHECK3-NEXT: [[TMP9:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0611// CHECK3-NEXT: store ptr null, ptr [[TMP9]], align 4612// CHECK3-NEXT: [[TMP10:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1613// CHECK3-NEXT: store i32 [[TMP4]], ptr [[TMP10]], align 4614// CHECK3-NEXT: [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1615// CHECK3-NEXT: store i32 [[TMP4]], ptr [[TMP11]], align 4616// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 1617// CHECK3-NEXT: store ptr null, ptr [[TMP12]], align 4618// CHECK3-NEXT: [[TMP13:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2619// CHECK3-NEXT: store i32 [[TMP6]], ptr [[TMP13]], align 4620// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2621// CHECK3-NEXT: store i32 [[TMP6]], ptr [[TMP14]], align 4622// CHECK3-NEXT: [[TMP15:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 2623// CHECK3-NEXT: store ptr null, ptr [[TMP15]], align 4624// CHECK3-NEXT: [[TMP16:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3625// CHECK3-NEXT: store ptr @a, ptr [[TMP16]], align 4626// CHECK3-NEXT: [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3627// CHECK3-NEXT: store ptr @a, ptr [[TMP17]], align 4628// CHECK3-NEXT: [[TMP18:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 3629// CHECK3-NEXT: store ptr null, ptr [[TMP18]], align 4630// CHECK3-NEXT: [[TMP19:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0631// CHECK3-NEXT: [[TMP20:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0632// CHECK3-NEXT: [[TMP21:%.*]] = load i32, ptr [[TE]], align 4633// CHECK3-NEXT: [[TMP22:%.*]] = load i32, ptr [[TH]], align 4634// CHECK3-NEXT: [[TMP23:%.*]] = load i32, ptr [[N_ADDR]], align 4635// CHECK3-NEXT: store i32 [[TMP23]], ptr [[DOTCAPTURE_EXPR_]], align 4636// CHECK3-NEXT: [[TMP24:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4637// CHECK3-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP24]], 0638// CHECK3-NEXT: [[DIV2:%.*]] = sdiv i32 [[SUB]], 1639// CHECK3-NEXT: [[SUB3:%.*]] = sub nsw i32 [[DIV2]], 1640// CHECK3-NEXT: store i32 [[SUB3]], ptr [[DOTCAPTURE_EXPR_1]], align 4641// CHECK3-NEXT: [[TMP25:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4642// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP25]], 1643// CHECK3-NEXT: [[TMP26:%.*]] = zext i32 [[ADD]] to i64644// CHECK3-NEXT: [[TMP27:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP21]], 0645// CHECK3-NEXT: [[TMP28:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP22]], 0646// CHECK3-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0647// CHECK3-NEXT: store i32 3, ptr [[TMP29]], align 4648// CHECK3-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1649// CHECK3-NEXT: store i32 4, ptr [[TMP30]], align 4650// CHECK3-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2651// CHECK3-NEXT: store ptr [[TMP19]], ptr [[TMP31]], align 4652// CHECK3-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3653// CHECK3-NEXT: store ptr [[TMP20]], ptr [[TMP32]], align 4654// CHECK3-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4655// CHECK3-NEXT: store ptr @.offload_sizes, ptr [[TMP33]], align 4656// CHECK3-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5657// CHECK3-NEXT: store ptr @.offload_maptypes, ptr [[TMP34]], align 4658// CHECK3-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6659// CHECK3-NEXT: store ptr null, ptr [[TMP35]], align 4660// CHECK3-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7661// CHECK3-NEXT: store ptr null, ptr [[TMP36]], align 4662// CHECK3-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8663// CHECK3-NEXT: store i64 [[TMP26]], ptr [[TMP37]], align 8664// CHECK3-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9665// CHECK3-NEXT: store i64 0, ptr [[TMP38]], align 8666// CHECK3-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10667// CHECK3-NEXT: store [3 x i32] [[TMP27]], ptr [[TMP39]], align 4668// CHECK3-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11669// CHECK3-NEXT: store [3 x i32] [[TMP28]], ptr [[TMP40]], align 4670// CHECK3-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12671// CHECK3-NEXT: store i32 0, ptr [[TMP41]], align 4672// CHECK3-NEXT: [[TMP42:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2:[0-9]+]], i64 -1, i32 [[TMP21]], i32 [[TMP22]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l28.region_id, ptr [[KERNEL_ARGS]])673// CHECK3-NEXT: [[TMP43:%.*]] = icmp ne i32 [[TMP42]], 0674// CHECK3-NEXT: br i1 [[TMP43]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]675// CHECK3: omp_offload.failed:676// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l28(i32 [[TMP2]], i32 [[TMP4]], i32 [[TMP6]], ptr @a) #[[ATTR2:[0-9]+]]677// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]]678// CHECK3: omp_offload.cont:679// CHECK3-NEXT: [[TMP44:%.*]] = load i32, ptr [[N_ADDR]], align 4680// CHECK3-NEXT: store i32 [[TMP44]], ptr [[N_CASTED4]], align 4681// CHECK3-NEXT: [[TMP45:%.*]] = load i32, ptr [[N_CASTED4]], align 4682// CHECK3-NEXT: [[TMP46:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0683// CHECK3-NEXT: store i32 [[TMP45]], ptr [[TMP46]], align 4684// CHECK3-NEXT: [[TMP47:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS6]], i32 0, i32 0685// CHECK3-NEXT: store i32 [[TMP45]], ptr [[TMP47]], align 4686// CHECK3-NEXT: [[TMP48:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS7]], i32 0, i32 0687// CHECK3-NEXT: store ptr null, ptr [[TMP48]], align 4688// CHECK3-NEXT: [[TMP49:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 1689// CHECK3-NEXT: store ptr @a, ptr [[TMP49]], align 4690// CHECK3-NEXT: [[TMP50:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS6]], i32 0, i32 1691// CHECK3-NEXT: store ptr @a, ptr [[TMP50]], align 4692// CHECK3-NEXT: [[TMP51:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS7]], i32 0, i32 1693// CHECK3-NEXT: store ptr null, ptr [[TMP51]], align 4694// CHECK3-NEXT: [[TMP52:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0695// CHECK3-NEXT: [[TMP53:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS6]], i32 0, i32 0696// CHECK3-NEXT: [[TMP54:%.*]] = load i32, ptr [[N_ADDR]], align 4697// CHECK3-NEXT: store i32 [[TMP54]], ptr [[DOTCAPTURE_EXPR_9]], align 4698// CHECK3-NEXT: [[TMP55:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_9]], align 4699// CHECK3-NEXT: [[SUB11:%.*]] = sub nsw i32 [[TMP55]], 0700// CHECK3-NEXT: [[DIV12:%.*]] = sdiv i32 [[SUB11]], 1701// CHECK3-NEXT: [[SUB13:%.*]] = sub nsw i32 [[DIV12]], 1702// CHECK3-NEXT: store i32 [[SUB13]], ptr [[DOTCAPTURE_EXPR_10]], align 4703// CHECK3-NEXT: [[TMP56:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_10]], align 4704// CHECK3-NEXT: [[ADD14:%.*]] = add nsw i32 [[TMP56]], 1705// CHECK3-NEXT: [[TMP57:%.*]] = zext i32 [[ADD14]] to i64706// CHECK3-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 0707// CHECK3-NEXT: store i32 3, ptr [[TMP58]], align 4708// CHECK3-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 1709// CHECK3-NEXT: store i32 2, ptr [[TMP59]], align 4710// CHECK3-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 2711// CHECK3-NEXT: store ptr [[TMP52]], ptr [[TMP60]], align 4712// CHECK3-NEXT: [[TMP61:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 3713// CHECK3-NEXT: store ptr [[TMP53]], ptr [[TMP61]], align 4714// CHECK3-NEXT: [[TMP62:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 4715// CHECK3-NEXT: store ptr @.offload_sizes.1, ptr [[TMP62]], align 4716// CHECK3-NEXT: [[TMP63:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 5717// CHECK3-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP63]], align 4718// CHECK3-NEXT: [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 6719// CHECK3-NEXT: store ptr null, ptr [[TMP64]], align 4720// CHECK3-NEXT: [[TMP65:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 7721// CHECK3-NEXT: store ptr null, ptr [[TMP65]], align 4722// CHECK3-NEXT: [[TMP66:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 8723// CHECK3-NEXT: store i64 [[TMP57]], ptr [[TMP66]], align 8724// CHECK3-NEXT: [[TMP67:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 9725// CHECK3-NEXT: store i64 0, ptr [[TMP67]], align 8726// CHECK3-NEXT: [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 10727// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP68]], align 4728// CHECK3-NEXT: [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 11729// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP69]], align 4730// CHECK3-NEXT: [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 12731// CHECK3-NEXT: store i32 0, ptr [[TMP70]], align 4732// CHECK3-NEXT: [[TMP71:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l34.region_id, ptr [[KERNEL_ARGS15]])733// CHECK3-NEXT: [[TMP72:%.*]] = icmp ne i32 [[TMP71]], 0734// CHECK3-NEXT: br i1 [[TMP72]], label [[OMP_OFFLOAD_FAILED16:%.*]], label [[OMP_OFFLOAD_CONT17:%.*]]735// CHECK3: omp_offload.failed16:736// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l34(i32 [[TMP45]], ptr @a) #[[ATTR2]]737// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT17]]738// CHECK3: omp_offload.cont17:739// CHECK3-NEXT: [[TMP73:%.*]] = load i32, ptr @a, align 4740// CHECK3-NEXT: ret i32 [[TMP73]]741//742//743// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l28744// CHECK3-SAME: (i32 noundef [[TE:%.*]], i32 noundef [[TH:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(400) [[A:%.*]]) #[[ATTR1:[0-9]+]] {745// CHECK3-NEXT: entry:746// CHECK3-NEXT: [[TE_ADDR:%.*]] = alloca i32, align 4747// CHECK3-NEXT: [[TH_ADDR:%.*]] = alloca i32, align 4748// CHECK3-NEXT: [[N_ADDR:%.*]] = alloca i32, align 4749// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 4750// CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB2]])751// CHECK3-NEXT: store i32 [[TE]], ptr [[TE_ADDR]], align 4752// CHECK3-NEXT: store i32 [[TH]], ptr [[TH_ADDR]], align 4753// CHECK3-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 4754// CHECK3-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 4755// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 4756// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[TE_ADDR]], align 4757// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[TH_ADDR]], align 4758// CHECK3-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB2]], i32 [[TMP0]], i32 [[TMP2]], i32 [[TMP3]])759// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l28.omp_outlined, ptr [[N_ADDR]], ptr [[TMP1]])760// CHECK3-NEXT: ret void761//762//763// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l28.omp_outlined764// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(400) [[A:%.*]]) #[[ATTR1]] {765// CHECK3-NEXT: entry:766// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4767// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4768// CHECK3-NEXT: [[N_ADDR:%.*]] = alloca ptr, align 4769// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 4770// CHECK3-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4771// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4772// CHECK3-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4773// CHECK3-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4774// CHECK3-NEXT: [[I:%.*]] = alloca i32, align 4775// CHECK3-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4776// CHECK3-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4777// CHECK3-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4778// CHECK3-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4779// CHECK3-NEXT: [[I3:%.*]] = alloca i32, align 4780// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4781// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4782// CHECK3-NEXT: store ptr [[N]], ptr [[N_ADDR]], align 4783// CHECK3-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 4784// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 4785// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 4786// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP0]], align 4787// CHECK3-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4788// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4789// CHECK3-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0790// CHECK3-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1791// CHECK3-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1792// CHECK3-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4793// CHECK3-NEXT: store i32 0, ptr [[I]], align 4794// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4795// CHECK3-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]796// CHECK3-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]797// CHECK3: omp.precond.then:798// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4799// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4800// CHECK3-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_COMB_UB]], align 4801// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4802// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4803// CHECK3-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4804// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 4805// CHECK3-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP7]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)806// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4807// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4808// CHECK3-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]809// CHECK3-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]810// CHECK3: cond.true:811// CHECK3-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4812// CHECK3-NEXT: br label [[COND_END:%.*]]813// CHECK3: cond.false:814// CHECK3-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4815// CHECK3-NEXT: br label [[COND_END]]816// CHECK3: cond.end:817// CHECK3-NEXT: [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]818// CHECK3-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4819// CHECK3-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4820// CHECK3-NEXT: store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 4821// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]822// CHECK3: omp.inner.for.cond:823// CHECK3-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4824// CHECK3-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4825// CHECK3-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]826// CHECK3-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]827// CHECK3: omp.inner.for.body:828// CHECK3-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4829// CHECK3-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP15]], 1830// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]831// CHECK3-NEXT: store i32 [[ADD]], ptr [[I3]], align 4832// CHECK3-NEXT: [[TMP16:%.*]] = load i32, ptr [[I3]], align 4833// CHECK3-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [100 x i32], ptr [[TMP1]], i32 0, i32 [[TMP16]]834// CHECK3-NEXT: store i32 0, ptr [[ARRAYIDX]], align 4835// CHECK3-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]836// CHECK3: omp.body.continue:837// CHECK3-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]838// CHECK3: omp.inner.for.inc:839// CHECK3-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4840// CHECK3-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP17]], 1841// CHECK3-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4842// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND]]843// CHECK3: omp.inner.for.end:844// CHECK3-NEXT: br label [[OMP_LOOP_EXIT:%.*]]845// CHECK3: omp.loop.exit:846// CHECK3-NEXT: [[TMP18:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4847// CHECK3-NEXT: [[TMP19:%.*]] = load i32, ptr [[TMP18]], align 4848// CHECK3-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP19]])849// CHECK3-NEXT: br label [[OMP_PRECOND_END]]850// CHECK3: omp.precond.end:851// CHECK3-NEXT: ret void852//853//854// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l34855// CHECK3-SAME: (i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(400) [[A:%.*]]) #[[ATTR1]] {856// CHECK3-NEXT: entry:857// CHECK3-NEXT: [[N_ADDR:%.*]] = alloca i32, align 4858// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 4859// CHECK3-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 4860// CHECK3-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 4861// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 4862// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l34.omp_outlined, ptr [[N_ADDR]], ptr [[TMP0]])863// CHECK3-NEXT: ret void864//865//866// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l34.omp_outlined867// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(400) [[A:%.*]]) #[[ATTR1]] {868// CHECK3-NEXT: entry:869// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4870// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4871// CHECK3-NEXT: [[N_ADDR:%.*]] = alloca ptr, align 4872// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 4873// CHECK3-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4874// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4875// CHECK3-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4876// CHECK3-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4877// CHECK3-NEXT: [[I:%.*]] = alloca i32, align 4878// CHECK3-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4879// CHECK3-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4880// CHECK3-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4881// CHECK3-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4882// CHECK3-NEXT: [[I3:%.*]] = alloca i32, align 4883// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4884// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4885// CHECK3-NEXT: store ptr [[N]], ptr [[N_ADDR]], align 4886// CHECK3-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 4887// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 4888// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 4889// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP0]], align 4890// CHECK3-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4891// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4892// CHECK3-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0893// CHECK3-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1894// CHECK3-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1895// CHECK3-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4896// CHECK3-NEXT: store i32 0, ptr [[I]], align 4897// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4898// CHECK3-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]899// CHECK3-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]900// CHECK3: omp.precond.then:901// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4902// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4903// CHECK3-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_COMB_UB]], align 4904// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4905// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4906// CHECK3-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4907// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 4908// CHECK3-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP7]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)909// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4910// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4911// CHECK3-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]912// CHECK3-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]913// CHECK3: cond.true:914// CHECK3-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4915// CHECK3-NEXT: br label [[COND_END:%.*]]916// CHECK3: cond.false:917// CHECK3-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4918// CHECK3-NEXT: br label [[COND_END]]919// CHECK3: cond.end:920// CHECK3-NEXT: [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]921// CHECK3-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4922// CHECK3-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4923// CHECK3-NEXT: store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 4924// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]925// CHECK3: omp.inner.for.cond:926// CHECK3-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4927// CHECK3-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4928// CHECK3-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]929// CHECK3-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]930// CHECK3: omp.inner.for.body:931// CHECK3-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4932// CHECK3-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP15]], 1933// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]934// CHECK3-NEXT: store i32 [[ADD]], ptr [[I3]], align 4935// CHECK3-NEXT: [[TMP16:%.*]] = load i32, ptr [[I3]], align 4936// CHECK3-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [100 x i32], ptr [[TMP1]], i32 0, i32 [[TMP16]]937// CHECK3-NEXT: store i32 0, ptr [[ARRAYIDX]], align 4938// CHECK3-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]939// CHECK3: omp.body.continue:940// CHECK3-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]941// CHECK3: omp.inner.for.inc:942// CHECK3-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4943// CHECK3-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP17]], 1944// CHECK3-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4945// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND]]946// CHECK3: omp.inner.for.end:947// CHECK3-NEXT: br label [[OMP_LOOP_EXIT:%.*]]948// CHECK3: omp.loop.exit:949// CHECK3-NEXT: [[TMP18:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4950// CHECK3-NEXT: [[TMP19:%.*]] = load i32, ptr [[TMP18]], align 4951// CHECK3-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP19]])952// CHECK3-NEXT: br label [[OMP_PRECOND_END]]953// CHECK3: omp.precond.end:954// CHECK3-NEXT: ret void955//956//957// CHECK9-LABEL: define {{[^@]+}}@_Z15teams_local_argv958// CHECK9-SAME: () #[[ATTR0:[0-9]+]] {959// CHECK9-NEXT: entry:960// CHECK9-NEXT: [[N:%.*]] = alloca i32, align 4961// CHECK9-NEXT: [[SAVED_STACK:%.*]] = alloca ptr, align 8962// CHECK9-NEXT: [[__VLA_EXPR0:%.*]] = alloca i64, align 8963// CHECK9-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8964// CHECK9-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 8965// CHECK9-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 8966// CHECK9-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 8967// CHECK9-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 8968// CHECK9-NEXT: [[TMP:%.*]] = alloca i32, align 4969// CHECK9-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4970// CHECK9-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4971// CHECK9-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8972// CHECK9-NEXT: store i32 100, ptr [[N]], align 4973// CHECK9-NEXT: [[TMP0:%.*]] = load i32, ptr [[N]], align 4974// CHECK9-NEXT: [[TMP1:%.*]] = zext i32 [[TMP0]] to i64975// CHECK9-NEXT: [[TMP2:%.*]] = call ptr @llvm.stacksave.p0()976// CHECK9-NEXT: store ptr [[TMP2]], ptr [[SAVED_STACK]], align 8977// CHECK9-NEXT: [[VLA:%.*]] = alloca i32, i64 [[TMP1]], align 4978// CHECK9-NEXT: store i64 [[TMP1]], ptr [[__VLA_EXPR0]], align 8979// CHECK9-NEXT: [[TMP3:%.*]] = load i32, ptr [[N]], align 4980// CHECK9-NEXT: store i32 [[TMP3]], ptr [[N_CASTED]], align 4981// CHECK9-NEXT: [[TMP4:%.*]] = load i64, ptr [[N_CASTED]], align 8982// CHECK9-NEXT: [[TMP5:%.*]] = mul nuw i64 [[TMP1]], 4983// CHECK9-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES]], ptr align 8 @.offload_sizes, i64 24, i1 false)984// CHECK9-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0985// CHECK9-NEXT: store i64 [[TMP4]], ptr [[TMP6]], align 8986// CHECK9-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0987// CHECK9-NEXT: store i64 [[TMP4]], ptr [[TMP7]], align 8988// CHECK9-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0989// CHECK9-NEXT: store ptr null, ptr [[TMP8]], align 8990// CHECK9-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1991// CHECK9-NEXT: store i64 [[TMP1]], ptr [[TMP9]], align 8992// CHECK9-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1993// CHECK9-NEXT: store i64 [[TMP1]], ptr [[TMP10]], align 8994// CHECK9-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1995// CHECK9-NEXT: store ptr null, ptr [[TMP11]], align 8996// CHECK9-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2997// CHECK9-NEXT: store ptr [[VLA]], ptr [[TMP12]], align 8998// CHECK9-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2999// CHECK9-NEXT: store ptr [[VLA]], ptr [[TMP13]], align 81000// CHECK9-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 21001// CHECK9-NEXT: store i64 [[TMP5]], ptr [[TMP14]], align 81002// CHECK9-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 21003// CHECK9-NEXT: store ptr null, ptr [[TMP15]], align 81004// CHECK9-NEXT: [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01005// CHECK9-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01006// CHECK9-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 01007// CHECK9-NEXT: [[TMP19:%.*]] = load i32, ptr [[N]], align 41008// CHECK9-NEXT: store i32 [[TMP19]], ptr [[DOTCAPTURE_EXPR_]], align 41009// CHECK9-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41010// CHECK9-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP20]], 01011// CHECK9-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11012// CHECK9-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11013// CHECK9-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41014// CHECK9-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41015// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP21]], 11016// CHECK9-NEXT: [[TMP22:%.*]] = zext i32 [[ADD]] to i641017// CHECK9-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01018// CHECK9-NEXT: store i32 3, ptr [[TMP23]], align 41019// CHECK9-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11020// CHECK9-NEXT: store i32 3, ptr [[TMP24]], align 41021// CHECK9-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21022// CHECK9-NEXT: store ptr [[TMP16]], ptr [[TMP25]], align 81023// CHECK9-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31024// CHECK9-NEXT: store ptr [[TMP17]], ptr [[TMP26]], align 81025// CHECK9-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41026// CHECK9-NEXT: store ptr [[TMP18]], ptr [[TMP27]], align 81027// CHECK9-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51028// CHECK9-NEXT: store ptr @.offload_maptypes, ptr [[TMP28]], align 81029// CHECK9-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61030// CHECK9-NEXT: store ptr null, ptr [[TMP29]], align 81031// CHECK9-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71032// CHECK9-NEXT: store ptr null, ptr [[TMP30]], align 81033// CHECK9-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81034// CHECK9-NEXT: store i64 [[TMP22]], ptr [[TMP31]], align 81035// CHECK9-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91036// CHECK9-NEXT: store i64 0, ptr [[TMP32]], align 81037// CHECK9-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101038// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP33]], align 41039// CHECK9-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111040// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP34]], align 41041// CHECK9-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121042// CHECK9-NEXT: store i32 0, ptr [[TMP35]], align 41043// 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]+}}__Z15teams_local_argv_l72.region_id, ptr [[KERNEL_ARGS]])1044// CHECK9-NEXT: [[TMP37:%.*]] = icmp ne i32 [[TMP36]], 01045// CHECK9-NEXT: br i1 [[TMP37]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1046// CHECK9: omp_offload.failed:1047// CHECK9-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z15teams_local_argv_l72(i64 [[TMP4]], i64 [[TMP1]], ptr [[VLA]]) #[[ATTR3:[0-9]+]]1048// CHECK9-NEXT: br label [[OMP_OFFLOAD_CONT]]1049// CHECK9: omp_offload.cont:1050// CHECK9-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[VLA]], i64 01051// CHECK9-NEXT: [[TMP38:%.*]] = load i32, ptr [[ARRAYIDX]], align 41052// CHECK9-NEXT: [[TMP39:%.*]] = load ptr, ptr [[SAVED_STACK]], align 81053// CHECK9-NEXT: call void @llvm.stackrestore.p0(ptr [[TMP39]])1054// CHECK9-NEXT: ret i32 [[TMP38]]1055//1056//1057// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z15teams_local_argv_l721058// CHECK9-SAME: (i64 noundef [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2:[0-9]+]] {1059// CHECK9-NEXT: entry:1060// CHECK9-NEXT: [[N_ADDR:%.*]] = alloca i64, align 81061// CHECK9-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 81062// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81063// CHECK9-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 81064// CHECK9-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR]], align 81065// CHECK9-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81066// CHECK9-NEXT: [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR]], align 81067// CHECK9-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 81068// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z15teams_local_argv_l72.omp_outlined, ptr [[N_ADDR]], i64 [[TMP0]], ptr [[TMP1]])1069// CHECK9-NEXT: ret void1070//1071//1072// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z15teams_local_argv_l72.omp_outlined1073// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {1074// CHECK9-NEXT: entry:1075// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81076// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81077// CHECK9-NEXT: [[N_ADDR:%.*]] = alloca ptr, align 81078// CHECK9-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 81079// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81080// CHECK9-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41081// CHECK9-NEXT: [[TMP:%.*]] = alloca i32, align 41082// CHECK9-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41083// CHECK9-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41084// CHECK9-NEXT: [[I:%.*]] = alloca i32, align 41085// CHECK9-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41086// CHECK9-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41087// CHECK9-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41088// CHECK9-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41089// CHECK9-NEXT: [[I3:%.*]] = alloca i32, align 41090// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81091// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81092// CHECK9-NEXT: store ptr [[N]], ptr [[N_ADDR]], align 81093// CHECK9-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR]], align 81094// CHECK9-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81095// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 81096// CHECK9-NEXT: [[TMP1:%.*]] = load i64, ptr [[VLA_ADDR]], align 81097// CHECK9-NEXT: [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 81098// CHECK9-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 41099// CHECK9-NEXT: store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR_]], align 41100// CHECK9-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41101// CHECK9-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP4]], 01102// CHECK9-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11103// CHECK9-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11104// CHECK9-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41105// CHECK9-NEXT: store i32 0, ptr [[I]], align 41106// CHECK9-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41107// CHECK9-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP5]]1108// CHECK9-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1109// CHECK9: omp.precond.then:1110// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 41111// CHECK9-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41112// CHECK9-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_COMB_UB]], align 41113// CHECK9-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41114// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41115// CHECK9-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81116// CHECK9-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41117// CHECK9-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP8]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1118// CHECK9-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41119// CHECK9-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41120// CHECK9-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]1121// CHECK9-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1122// CHECK9: cond.true:1123// CHECK9-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41124// CHECK9-NEXT: br label [[COND_END:%.*]]1125// CHECK9: cond.false:1126// CHECK9-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41127// CHECK9-NEXT: br label [[COND_END]]1128// CHECK9: cond.end:1129// CHECK9-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]1130// CHECK9-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41131// CHECK9-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41132// CHECK9-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 41133// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1134// CHECK9: omp.inner.for.cond:1135// CHECK9-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41136// CHECK9-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41137// CHECK9-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]1138// CHECK9-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1139// CHECK9: omp.inner.for.body:1140// CHECK9-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41141// CHECK9-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 11142// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1143// CHECK9-NEXT: store i32 [[ADD]], ptr [[I3]], align 41144// CHECK9-NEXT: [[TMP17:%.*]] = load i32, ptr [[I3]], align 41145// CHECK9-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP17]] to i641146// CHECK9-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i64 [[IDXPROM]]1147// CHECK9-NEXT: store i32 0, ptr [[ARRAYIDX]], align 41148// CHECK9-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1149// CHECK9: omp.body.continue:1150// CHECK9-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1151// CHECK9: omp.inner.for.inc:1152// CHECK9-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41153// CHECK9-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP18]], 11154// CHECK9-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 41155// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND]]1156// CHECK9: omp.inner.for.end:1157// CHECK9-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1158// CHECK9: omp.loop.exit:1159// CHECK9-NEXT: [[TMP19:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81160// CHECK9-NEXT: [[TMP20:%.*]] = load i32, ptr [[TMP19]], align 41161// CHECK9-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP20]])1162// CHECK9-NEXT: br label [[OMP_PRECOND_END]]1163// CHECK9: omp.precond.end:1164// CHECK9-NEXT: ret void1165//1166//1167// CHECK11-LABEL: define {{[^@]+}}@_Z15teams_local_argv1168// CHECK11-SAME: () #[[ATTR0:[0-9]+]] {1169// CHECK11-NEXT: entry:1170// CHECK11-NEXT: [[N:%.*]] = alloca i32, align 41171// CHECK11-NEXT: [[SAVED_STACK:%.*]] = alloca ptr, align 41172// CHECK11-NEXT: [[__VLA_EXPR0:%.*]] = alloca i32, align 41173// CHECK11-NEXT: [[N_CASTED:%.*]] = alloca i32, align 41174// CHECK11-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 41175// CHECK11-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 41176// CHECK11-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 41177// CHECK11-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 41178// CHECK11-NEXT: [[TMP:%.*]] = alloca i32, align 41179// CHECK11-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41180// CHECK11-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41181// CHECK11-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81182// CHECK11-NEXT: store i32 100, ptr [[N]], align 41183// CHECK11-NEXT: [[TMP0:%.*]] = load i32, ptr [[N]], align 41184// CHECK11-NEXT: [[TMP1:%.*]] = call ptr @llvm.stacksave.p0()1185// CHECK11-NEXT: store ptr [[TMP1]], ptr [[SAVED_STACK]], align 41186// CHECK11-NEXT: [[VLA:%.*]] = alloca i32, i32 [[TMP0]], align 41187// CHECK11-NEXT: store i32 [[TMP0]], ptr [[__VLA_EXPR0]], align 41188// CHECK11-NEXT: [[TMP2:%.*]] = load i32, ptr [[N]], align 41189// CHECK11-NEXT: store i32 [[TMP2]], ptr [[N_CASTED]], align 41190// CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[N_CASTED]], align 41191// CHECK11-NEXT: [[TMP4:%.*]] = mul nuw i32 [[TMP0]], 41192// CHECK11-NEXT: [[TMP5:%.*]] = sext i32 [[TMP4]] to i641193// CHECK11-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[DOTOFFLOAD_SIZES]], ptr align 4 @.offload_sizes, i32 24, i1 false)1194// CHECK11-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01195// CHECK11-NEXT: store i32 [[TMP3]], ptr [[TMP6]], align 41196// CHECK11-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01197// CHECK11-NEXT: store i32 [[TMP3]], ptr [[TMP7]], align 41198// CHECK11-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01199// CHECK11-NEXT: store ptr null, ptr [[TMP8]], align 41200// CHECK11-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11201// CHECK11-NEXT: store i32 [[TMP0]], ptr [[TMP9]], align 41202// CHECK11-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11203// CHECK11-NEXT: store i32 [[TMP0]], ptr [[TMP10]], align 41204// CHECK11-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 11205// CHECK11-NEXT: store ptr null, ptr [[TMP11]], align 41206// CHECK11-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21207// CHECK11-NEXT: store ptr [[VLA]], ptr [[TMP12]], align 41208// CHECK11-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21209// CHECK11-NEXT: store ptr [[VLA]], ptr [[TMP13]], align 41210// CHECK11-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 21211// CHECK11-NEXT: store i64 [[TMP5]], ptr [[TMP14]], align 41212// CHECK11-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 21213// CHECK11-NEXT: store ptr null, ptr [[TMP15]], align 41214// CHECK11-NEXT: [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01215// CHECK11-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01216// CHECK11-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 01217// CHECK11-NEXT: [[TMP19:%.*]] = load i32, ptr [[N]], align 41218// CHECK11-NEXT: store i32 [[TMP19]], ptr [[DOTCAPTURE_EXPR_]], align 41219// CHECK11-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41220// CHECK11-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP20]], 01221// CHECK11-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11222// CHECK11-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11223// CHECK11-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41224// CHECK11-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41225// CHECK11-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP21]], 11226// CHECK11-NEXT: [[TMP22:%.*]] = zext i32 [[ADD]] to i641227// CHECK11-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01228// CHECK11-NEXT: store i32 3, ptr [[TMP23]], align 41229// CHECK11-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11230// CHECK11-NEXT: store i32 3, ptr [[TMP24]], align 41231// CHECK11-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21232// CHECK11-NEXT: store ptr [[TMP16]], ptr [[TMP25]], align 41233// CHECK11-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31234// CHECK11-NEXT: store ptr [[TMP17]], ptr [[TMP26]], align 41235// CHECK11-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41236// CHECK11-NEXT: store ptr [[TMP18]], ptr [[TMP27]], align 41237// CHECK11-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51238// CHECK11-NEXT: store ptr @.offload_maptypes, ptr [[TMP28]], align 41239// CHECK11-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61240// CHECK11-NEXT: store ptr null, ptr [[TMP29]], align 41241// CHECK11-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71242// CHECK11-NEXT: store ptr null, ptr [[TMP30]], align 41243// CHECK11-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81244// CHECK11-NEXT: store i64 [[TMP22]], ptr [[TMP31]], align 81245// CHECK11-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91246// CHECK11-NEXT: store i64 0, ptr [[TMP32]], align 81247// CHECK11-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101248// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP33]], align 41249// CHECK11-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111250// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP34]], align 41251// CHECK11-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121252// CHECK11-NEXT: store i32 0, ptr [[TMP35]], align 41253// 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]+}}__Z15teams_local_argv_l72.region_id, ptr [[KERNEL_ARGS]])1254// CHECK11-NEXT: [[TMP37:%.*]] = icmp ne i32 [[TMP36]], 01255// CHECK11-NEXT: br i1 [[TMP37]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1256// CHECK11: omp_offload.failed:1257// CHECK11-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z15teams_local_argv_l72(i32 [[TMP3]], i32 [[TMP0]], ptr [[VLA]]) #[[ATTR3:[0-9]+]]1258// CHECK11-NEXT: br label [[OMP_OFFLOAD_CONT]]1259// CHECK11: omp_offload.cont:1260// CHECK11-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[VLA]], i32 01261// CHECK11-NEXT: [[TMP38:%.*]] = load i32, ptr [[ARRAYIDX]], align 41262// CHECK11-NEXT: [[TMP39:%.*]] = load ptr, ptr [[SAVED_STACK]], align 41263// CHECK11-NEXT: call void @llvm.stackrestore.p0(ptr [[TMP39]])1264// CHECK11-NEXT: ret i32 [[TMP38]]1265//1266//1267// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z15teams_local_argv_l721268// CHECK11-SAME: (i32 noundef [[N:%.*]], i32 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2:[0-9]+]] {1269// CHECK11-NEXT: entry:1270// CHECK11-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41271// CHECK11-NEXT: [[VLA_ADDR:%.*]] = alloca i32, align 41272// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41273// CHECK11-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41274// CHECK11-NEXT: store i32 [[VLA]], ptr [[VLA_ADDR]], align 41275// CHECK11-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41276// CHECK11-NEXT: [[TMP0:%.*]] = load i32, ptr [[VLA_ADDR]], align 41277// CHECK11-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 41278// CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z15teams_local_argv_l72.omp_outlined, ptr [[N_ADDR]], i32 [[TMP0]], ptr [[TMP1]])1279// CHECK11-NEXT: ret void1280//1281//1282// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z15teams_local_argv_l72.omp_outlined1283// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]], i32 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {1284// CHECK11-NEXT: entry:1285// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41286// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41287// CHECK11-NEXT: [[N_ADDR:%.*]] = alloca ptr, align 41288// CHECK11-NEXT: [[VLA_ADDR:%.*]] = alloca i32, align 41289// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41290// CHECK11-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41291// CHECK11-NEXT: [[TMP:%.*]] = alloca i32, align 41292// CHECK11-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41293// CHECK11-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41294// CHECK11-NEXT: [[I:%.*]] = alloca i32, align 41295// CHECK11-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41296// CHECK11-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41297// CHECK11-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41298// CHECK11-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41299// CHECK11-NEXT: [[I3:%.*]] = alloca i32, align 41300// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41301// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41302// CHECK11-NEXT: store ptr [[N]], ptr [[N_ADDR]], align 41303// CHECK11-NEXT: store i32 [[VLA]], ptr [[VLA_ADDR]], align 41304// CHECK11-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41305// CHECK11-NEXT: [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 41306// CHECK11-NEXT: [[TMP1:%.*]] = load i32, ptr [[VLA_ADDR]], align 41307// CHECK11-NEXT: [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 41308// CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 41309// CHECK11-NEXT: store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR_]], align 41310// CHECK11-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41311// CHECK11-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP4]], 01312// CHECK11-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11313// CHECK11-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11314// CHECK11-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41315// CHECK11-NEXT: store i32 0, ptr [[I]], align 41316// CHECK11-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41317// CHECK11-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP5]]1318// CHECK11-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1319// CHECK11: omp.precond.then:1320// CHECK11-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 41321// CHECK11-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41322// CHECK11-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_COMB_UB]], align 41323// CHECK11-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41324// CHECK11-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41325// CHECK11-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41326// CHECK11-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41327// CHECK11-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP8]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1328// CHECK11-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41329// CHECK11-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41330// CHECK11-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]1331// CHECK11-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1332// CHECK11: cond.true:1333// CHECK11-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41334// CHECK11-NEXT: br label [[COND_END:%.*]]1335// CHECK11: cond.false:1336// CHECK11-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41337// CHECK11-NEXT: br label [[COND_END]]1338// CHECK11: cond.end:1339// CHECK11-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]1340// CHECK11-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41341// CHECK11-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41342// CHECK11-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 41343// CHECK11-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1344// CHECK11: omp.inner.for.cond:1345// CHECK11-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41346// CHECK11-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41347// CHECK11-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]1348// CHECK11-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1349// CHECK11: omp.inner.for.body:1350// CHECK11-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41351// CHECK11-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 11352// CHECK11-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1353// CHECK11-NEXT: store i32 [[ADD]], ptr [[I3]], align 41354// CHECK11-NEXT: [[TMP17:%.*]] = load i32, ptr [[I3]], align 41355// CHECK11-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i32 [[TMP17]]1356// CHECK11-NEXT: store i32 0, ptr [[ARRAYIDX]], align 41357// CHECK11-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1358// CHECK11: omp.body.continue:1359// CHECK11-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1360// CHECK11: omp.inner.for.inc:1361// CHECK11-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41362// CHECK11-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP18]], 11363// CHECK11-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 41364// CHECK11-NEXT: br label [[OMP_INNER_FOR_COND]]1365// CHECK11: omp.inner.for.end:1366// CHECK11-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1367// CHECK11: omp.loop.exit:1368// CHECK11-NEXT: [[TMP19:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41369// CHECK11-NEXT: [[TMP20:%.*]] = load i32, ptr [[TMP19]], align 41370// CHECK11-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP20]])1371// CHECK11-NEXT: br label [[OMP_PRECOND_END]]1372// CHECK11: omp.precond.end:1373// CHECK11-NEXT: ret void1374//1375//1376// CHECK17-LABEL: define {{[^@]+}}@_Z21teams_template_structv1377// CHECK17-SAME: () #[[ATTR0:[0-9]+]] {1378// CHECK17-NEXT: entry:1379// CHECK17-NEXT: [[V:%.*]] = alloca [[STRUCT_SS:%.*]], align 41380// CHECK17-NEXT: [[CALL:%.*]] = call noundef signext i32 @_ZN2SSIiLi123ELx456EE3fooEv(ptr noundef nonnull align 4 dereferenceable(496) [[V]])1381// CHECK17-NEXT: ret i32 [[CALL]]1382//1383//1384// CHECK17-LABEL: define {{[^@]+}}@_ZN2SSIiLi123ELx456EE3fooEv1385// CHECK17-SAME: (ptr noundef nonnull align 4 dereferenceable(496) [[THIS:%.*]]) #[[ATTR0]] comdat {1386// CHECK17-NEXT: entry:1387// CHECK17-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81388// CHECK17-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 81389// CHECK17-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 81390// CHECK17-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 81391// CHECK17-NEXT: [[TMP:%.*]] = alloca i32, align 41392// CHECK17-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81393// CHECK17-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81394// CHECK17-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81395// CHECK17-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[THIS1]], i32 0, i32 01396// CHECK17-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01397// CHECK17-NEXT: store ptr [[THIS1]], ptr [[TMP0]], align 81398// CHECK17-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01399// CHECK17-NEXT: store ptr [[A]], ptr [[TMP1]], align 81400// CHECK17-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 01401// CHECK17-NEXT: store ptr null, ptr [[TMP2]], align 81402// CHECK17-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01403// CHECK17-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01404// CHECK17-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01405// CHECK17-NEXT: store i32 3, ptr [[TMP5]], align 41406// CHECK17-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11407// CHECK17-NEXT: store i32 1, ptr [[TMP6]], align 41408// CHECK17-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21409// CHECK17-NEXT: store ptr [[TMP3]], ptr [[TMP7]], align 81410// CHECK17-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31411// CHECK17-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 81412// CHECK17-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41413// CHECK17-NEXT: store ptr @.offload_sizes, ptr [[TMP9]], align 81414// CHECK17-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51415// CHECK17-NEXT: store ptr @.offload_maptypes, ptr [[TMP10]], align 81416// CHECK17-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61417// CHECK17-NEXT: store ptr null, ptr [[TMP11]], align 81418// CHECK17-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71419// CHECK17-NEXT: store ptr null, ptr [[TMP12]], align 81420// CHECK17-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81421// CHECK17-NEXT: store i64 123, ptr [[TMP13]], align 81422// CHECK17-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91423// CHECK17-NEXT: store i64 0, ptr [[TMP14]], align 81424// CHECK17-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101425// CHECK17-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 41426// CHECK17-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111427// CHECK17-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 41428// CHECK17-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121429// CHECK17-NEXT: store i32 0, ptr [[TMP17]], align 41430// CHECK17-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_l108.region_id, ptr [[KERNEL_ARGS]])1431// CHECK17-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 01432// CHECK17-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1433// CHECK17: omp_offload.failed:1434// CHECK17-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l108(ptr [[THIS1]]) #[[ATTR2:[0-9]+]]1435// CHECK17-NEXT: br label [[OMP_OFFLOAD_CONT]]1436// CHECK17: omp_offload.cont:1437// CHECK17-NEXT: [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 01438// CHECK17-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A2]], i64 0, i64 01439// CHECK17-NEXT: [[TMP20:%.*]] = load i32, ptr [[ARRAYIDX]], align 41440// CHECK17-NEXT: ret i32 [[TMP20]]1441//1442//1443// CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l1081444// CHECK17-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1:[0-9]+]] {1445// CHECK17-NEXT: entry:1446// CHECK17-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81447// CHECK17-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81448// CHECK17-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81449// CHECK17-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l108.omp_outlined, ptr [[TMP0]])1450// CHECK17-NEXT: ret void1451//1452//1453// CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l108.omp_outlined1454// CHECK17-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {1455// CHECK17-NEXT: entry:1456// CHECK17-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81457// CHECK17-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81458// CHECK17-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81459// CHECK17-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41460// CHECK17-NEXT: [[TMP:%.*]] = alloca i32, align 41461// CHECK17-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41462// CHECK17-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41463// CHECK17-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41464// CHECK17-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41465// CHECK17-NEXT: [[I:%.*]] = alloca i32, align 41466// CHECK17-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81467// CHECK17-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81468// CHECK17-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81469// CHECK17-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81470// CHECK17-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 41471// CHECK17-NEXT: store i32 122, ptr [[DOTOMP_COMB_UB]], align 41472// CHECK17-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41473// CHECK17-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41474// CHECK17-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81475// CHECK17-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41476// CHECK17-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1477// CHECK17-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41478// CHECK17-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1221479// CHECK17-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1480// CHECK17: cond.true:1481// CHECK17-NEXT: br label [[COND_END:%.*]]1482// CHECK17: cond.false:1483// CHECK17-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41484// CHECK17-NEXT: br label [[COND_END]]1485// CHECK17: cond.end:1486// CHECK17-NEXT: [[COND:%.*]] = phi i32 [ 122, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]1487// CHECK17-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41488// CHECK17-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41489// CHECK17-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 41490// CHECK17-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1491// CHECK17: omp.inner.for.cond:1492// CHECK17-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41493// CHECK17-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41494// CHECK17-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]1495// CHECK17-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1496// CHECK17: omp.inner.for.body:1497// CHECK17-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41498// CHECK17-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 11499// CHECK17-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1500// CHECK17-NEXT: store i32 [[ADD]], ptr [[I]], align 41501// CHECK17-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP0]], i32 0, i32 01502// CHECK17-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 41503// CHECK17-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP9]] to i641504// CHECK17-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A]], i64 0, i64 [[IDXPROM]]1505// CHECK17-NEXT: store i32 0, ptr [[ARRAYIDX]], align 41506// CHECK17-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1507// CHECK17: omp.body.continue:1508// CHECK17-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1509// CHECK17: omp.inner.for.inc:1510// CHECK17-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41511// CHECK17-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP10]], 11512// CHECK17-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 41513// CHECK17-NEXT: br label [[OMP_INNER_FOR_COND]]1514// CHECK17: omp.inner.for.end:1515// CHECK17-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1516// CHECK17: omp.loop.exit:1517// CHECK17-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])1518// CHECK17-NEXT: ret void1519//1520//1521// CHECK19-LABEL: define {{[^@]+}}@_Z21teams_template_structv1522// CHECK19-SAME: () #[[ATTR0:[0-9]+]] {1523// CHECK19-NEXT: entry:1524// CHECK19-NEXT: [[V:%.*]] = alloca [[STRUCT_SS:%.*]], align 41525// CHECK19-NEXT: [[CALL:%.*]] = call noundef i32 @_ZN2SSIiLi123ELx456EE3fooEv(ptr noundef nonnull align 4 dereferenceable(496) [[V]])1526// CHECK19-NEXT: ret i32 [[CALL]]1527//1528//1529// CHECK19-LABEL: define {{[^@]+}}@_ZN2SSIiLi123ELx456EE3fooEv1530// CHECK19-SAME: (ptr noundef nonnull align 4 dereferenceable(496) [[THIS:%.*]]) #[[ATTR0]] comdat align 2 {1531// CHECK19-NEXT: entry:1532// CHECK19-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41533// CHECK19-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 41534// CHECK19-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 41535// CHECK19-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 41536// CHECK19-NEXT: [[TMP:%.*]] = alloca i32, align 41537// CHECK19-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81538// CHECK19-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41539// CHECK19-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41540// CHECK19-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[THIS1]], i32 0, i32 01541// CHECK19-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01542// CHECK19-NEXT: store ptr [[THIS1]], ptr [[TMP0]], align 41543// CHECK19-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01544// CHECK19-NEXT: store ptr [[A]], ptr [[TMP1]], align 41545// CHECK19-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01546// CHECK19-NEXT: store ptr null, ptr [[TMP2]], align 41547// CHECK19-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01548// CHECK19-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01549// CHECK19-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01550// CHECK19-NEXT: store i32 3, ptr [[TMP5]], align 41551// CHECK19-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11552// CHECK19-NEXT: store i32 1, ptr [[TMP6]], align 41553// CHECK19-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21554// CHECK19-NEXT: store ptr [[TMP3]], ptr [[TMP7]], align 41555// CHECK19-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31556// CHECK19-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 41557// CHECK19-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41558// CHECK19-NEXT: store ptr @.offload_sizes, ptr [[TMP9]], align 41559// CHECK19-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51560// CHECK19-NEXT: store ptr @.offload_maptypes, ptr [[TMP10]], align 41561// CHECK19-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61562// CHECK19-NEXT: store ptr null, ptr [[TMP11]], align 41563// CHECK19-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71564// CHECK19-NEXT: store ptr null, ptr [[TMP12]], align 41565// CHECK19-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81566// CHECK19-NEXT: store i64 123, ptr [[TMP13]], align 81567// CHECK19-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91568// CHECK19-NEXT: store i64 0, ptr [[TMP14]], align 81569// CHECK19-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101570// CHECK19-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 41571// CHECK19-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111572// CHECK19-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 41573// CHECK19-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121574// CHECK19-NEXT: store i32 0, ptr [[TMP17]], align 41575// CHECK19-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_l108.region_id, ptr [[KERNEL_ARGS]])1576// CHECK19-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 01577// CHECK19-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1578// CHECK19: omp_offload.failed:1579// CHECK19-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l108(ptr [[THIS1]]) #[[ATTR2:[0-9]+]]1580// CHECK19-NEXT: br label [[OMP_OFFLOAD_CONT]]1581// CHECK19: omp_offload.cont:1582// CHECK19-NEXT: [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 01583// CHECK19-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A2]], i32 0, i32 01584// CHECK19-NEXT: [[TMP20:%.*]] = load i32, ptr [[ARRAYIDX]], align 41585// CHECK19-NEXT: ret i32 [[TMP20]]1586//1587//1588// CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l1081589// CHECK19-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1:[0-9]+]] {1590// CHECK19-NEXT: entry:1591// CHECK19-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41592// CHECK19-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41593// CHECK19-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41594// CHECK19-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l108.omp_outlined, ptr [[TMP0]])1595// CHECK19-NEXT: ret void1596//1597//1598// CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l108.omp_outlined1599// CHECK19-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {1600// CHECK19-NEXT: entry:1601// CHECK19-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41602// CHECK19-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41603// CHECK19-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41604// CHECK19-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41605// CHECK19-NEXT: [[TMP:%.*]] = alloca i32, align 41606// CHECK19-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41607// CHECK19-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41608// CHECK19-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41609// CHECK19-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41610// CHECK19-NEXT: [[I:%.*]] = alloca i32, align 41611// CHECK19-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41612// CHECK19-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41613// CHECK19-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41614// CHECK19-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41615// CHECK19-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 41616// CHECK19-NEXT: store i32 122, ptr [[DOTOMP_COMB_UB]], align 41617// CHECK19-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41618// CHECK19-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41619// CHECK19-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41620// CHECK19-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41621// CHECK19-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1622// CHECK19-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41623// CHECK19-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1221624// CHECK19-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1625// CHECK19: cond.true:1626// CHECK19-NEXT: br label [[COND_END:%.*]]1627// CHECK19: cond.false:1628// CHECK19-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41629// CHECK19-NEXT: br label [[COND_END]]1630// CHECK19: cond.end:1631// CHECK19-NEXT: [[COND:%.*]] = phi i32 [ 122, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]1632// CHECK19-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41633// CHECK19-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41634// CHECK19-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 41635// CHECK19-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1636// CHECK19: omp.inner.for.cond:1637// CHECK19-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41638// CHECK19-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41639// CHECK19-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]1640// CHECK19-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1641// CHECK19: omp.inner.for.body:1642// CHECK19-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41643// CHECK19-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 11644// CHECK19-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1645// CHECK19-NEXT: store i32 [[ADD]], ptr [[I]], align 41646// CHECK19-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP0]], i32 0, i32 01647// CHECK19-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 41648// CHECK19-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A]], i32 0, i32 [[TMP9]]1649// CHECK19-NEXT: store i32 0, ptr [[ARRAYIDX]], align 41650// CHECK19-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1651// CHECK19: omp.body.continue:1652// CHECK19-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1653// CHECK19: omp.inner.for.inc:1654// CHECK19-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41655// CHECK19-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP10]], 11656// CHECK19-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 41657// CHECK19-NEXT: br label [[OMP_INNER_FOR_COND]]1658// CHECK19: omp.inner.for.end:1659// CHECK19-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1660// CHECK19: omp.loop.exit:1661// CHECK19-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])1662// CHECK19-NEXT: ret void1663//1664//1665// CHECK25-LABEL: define {{[^@]+}}@main1666// CHECK25-SAME: (i32 noundef signext [[ARGC:%.*]], ptr noundef [[ARGV:%.*]]) #[[ATTR0:[0-9]+]] {1667// CHECK25-NEXT: entry:1668// CHECK25-NEXT: [[RETVAL:%.*]] = alloca i32, align 41669// CHECK25-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 41670// CHECK25-NEXT: [[ARGV_ADDR:%.*]] = alloca ptr, align 81671// CHECK25-NEXT: [[N:%.*]] = alloca i32, align 41672// CHECK25-NEXT: [[SAVED_STACK:%.*]] = alloca ptr, align 81673// CHECK25-NEXT: [[__VLA_EXPR0:%.*]] = alloca i64, align 81674// CHECK25-NEXT: [[N_CASTED:%.*]] = alloca i64, align 81675// CHECK25-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 81676// CHECK25-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 81677// CHECK25-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 81678// CHECK25-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 81679// CHECK25-NEXT: [[TMP:%.*]] = alloca i32, align 41680// CHECK25-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41681// CHECK25-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41682// CHECK25-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81683// CHECK25-NEXT: store i32 0, ptr [[RETVAL]], align 41684// CHECK25-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 41685// CHECK25-NEXT: store ptr [[ARGV]], ptr [[ARGV_ADDR]], align 81686// CHECK25-NEXT: store i32 100, ptr [[N]], align 41687// CHECK25-NEXT: [[TMP0:%.*]] = load i32, ptr [[N]], align 41688// CHECK25-NEXT: [[TMP1:%.*]] = zext i32 [[TMP0]] to i641689// CHECK25-NEXT: [[TMP2:%.*]] = call ptr @llvm.stacksave.p0()1690// CHECK25-NEXT: store ptr [[TMP2]], ptr [[SAVED_STACK]], align 81691// CHECK25-NEXT: [[VLA:%.*]] = alloca i32, i64 [[TMP1]], align 41692// CHECK25-NEXT: store i64 [[TMP1]], ptr [[__VLA_EXPR0]], align 81693// CHECK25-NEXT: [[TMP3:%.*]] = load i32, ptr [[N]], align 41694// CHECK25-NEXT: store i32 [[TMP3]], ptr [[N_CASTED]], align 41695// CHECK25-NEXT: [[TMP4:%.*]] = load i64, ptr [[N_CASTED]], align 81696// CHECK25-NEXT: [[TMP5:%.*]] = mul nuw i64 [[TMP1]], 41697// CHECK25-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES]], ptr align 8 @.offload_sizes, i64 24, i1 false)1698// CHECK25-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01699// CHECK25-NEXT: store i64 [[TMP4]], ptr [[TMP6]], align 81700// CHECK25-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01701// CHECK25-NEXT: store i64 [[TMP4]], ptr [[TMP7]], align 81702// CHECK25-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 01703// CHECK25-NEXT: store ptr null, ptr [[TMP8]], align 81704// CHECK25-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11705// CHECK25-NEXT: store i64 [[TMP1]], ptr [[TMP9]], align 81706// CHECK25-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11707// CHECK25-NEXT: store i64 [[TMP1]], ptr [[TMP10]], align 81708// CHECK25-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 11709// CHECK25-NEXT: store ptr null, ptr [[TMP11]], align 81710// CHECK25-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21711// CHECK25-NEXT: store ptr [[VLA]], ptr [[TMP12]], align 81712// CHECK25-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21713// CHECK25-NEXT: store ptr [[VLA]], ptr [[TMP13]], align 81714// CHECK25-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 21715// CHECK25-NEXT: store i64 [[TMP5]], ptr [[TMP14]], align 81716// CHECK25-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 21717// CHECK25-NEXT: store ptr null, ptr [[TMP15]], align 81718// CHECK25-NEXT: [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01719// CHECK25-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01720// CHECK25-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 01721// CHECK25-NEXT: [[TMP19:%.*]] = load i32, ptr [[N]], align 41722// CHECK25-NEXT: store i32 [[TMP19]], ptr [[DOTCAPTURE_EXPR_]], align 41723// CHECK25-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41724// CHECK25-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP20]], 01725// CHECK25-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11726// CHECK25-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11727// CHECK25-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41728// CHECK25-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41729// CHECK25-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP21]], 11730// CHECK25-NEXT: [[TMP22:%.*]] = zext i32 [[ADD]] to i641731// CHECK25-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01732// CHECK25-NEXT: store i32 3, ptr [[TMP23]], align 41733// CHECK25-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11734// CHECK25-NEXT: store i32 3, ptr [[TMP24]], align 41735// CHECK25-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21736// CHECK25-NEXT: store ptr [[TMP16]], ptr [[TMP25]], align 81737// CHECK25-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31738// CHECK25-NEXT: store ptr [[TMP17]], ptr [[TMP26]], align 81739// CHECK25-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41740// CHECK25-NEXT: store ptr [[TMP18]], ptr [[TMP27]], align 81741// CHECK25-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51742// CHECK25-NEXT: store ptr @.offload_maptypes, ptr [[TMP28]], align 81743// CHECK25-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61744// CHECK25-NEXT: store ptr null, ptr [[TMP29]], align 81745// CHECK25-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71746// CHECK25-NEXT: store ptr null, ptr [[TMP30]], align 81747// CHECK25-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81748// CHECK25-NEXT: store i64 [[TMP22]], ptr [[TMP31]], align 81749// CHECK25-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91750// CHECK25-NEXT: store i64 0, ptr [[TMP32]], align 81751// CHECK25-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101752// CHECK25-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP33]], align 41753// CHECK25-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111754// CHECK25-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP34]], align 41755// CHECK25-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121756// CHECK25-NEXT: store i32 0, ptr [[TMP35]], align 41757// CHECK25-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_l161.region_id, ptr [[KERNEL_ARGS]])1758// CHECK25-NEXT: [[TMP37:%.*]] = icmp ne i32 [[TMP36]], 01759// CHECK25-NEXT: br i1 [[TMP37]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1760// CHECK25: omp_offload.failed:1761// CHECK25-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l161(i64 [[TMP4]], i64 [[TMP1]], ptr [[VLA]]) #[[ATTR3:[0-9]+]]1762// CHECK25-NEXT: br label [[OMP_OFFLOAD_CONT]]1763// CHECK25: omp_offload.cont:1764// CHECK25-NEXT: [[TMP38:%.*]] = load i32, ptr [[ARGC_ADDR]], align 41765// CHECK25-NEXT: [[CALL:%.*]] = call noundef signext i32 @_Z5tmainIiLi10EEiT_(i32 noundef signext [[TMP38]])1766// CHECK25-NEXT: store i32 [[CALL]], ptr [[RETVAL]], align 41767// CHECK25-NEXT: [[TMP39:%.*]] = load ptr, ptr [[SAVED_STACK]], align 81768// CHECK25-NEXT: call void @llvm.stackrestore.p0(ptr [[TMP39]])1769// CHECK25-NEXT: [[TMP40:%.*]] = load i32, ptr [[RETVAL]], align 41770// CHECK25-NEXT: ret i32 [[TMP40]]1771//1772//1773// CHECK25-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l1611774// CHECK25-SAME: (i64 noundef [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2:[0-9]+]] {1775// CHECK25-NEXT: entry:1776// CHECK25-NEXT: [[N_ADDR:%.*]] = alloca i64, align 81777// CHECK25-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 81778// CHECK25-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81779// CHECK25-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 81780// CHECK25-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR]], align 81781// CHECK25-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81782// CHECK25-NEXT: [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR]], align 81783// CHECK25-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 81784// CHECK25-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l161.omp_outlined, ptr [[N_ADDR]], i64 [[TMP0]], ptr [[TMP1]])1785// CHECK25-NEXT: ret void1786//1787//1788// CHECK25-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l161.omp_outlined1789// CHECK25-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {1790// CHECK25-NEXT: entry:1791// CHECK25-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81792// CHECK25-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81793// CHECK25-NEXT: [[N_ADDR:%.*]] = alloca ptr, align 81794// CHECK25-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 81795// CHECK25-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81796// CHECK25-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41797// CHECK25-NEXT: [[TMP:%.*]] = alloca i32, align 41798// CHECK25-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41799// CHECK25-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41800// CHECK25-NEXT: [[I:%.*]] = alloca i32, align 41801// CHECK25-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41802// CHECK25-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41803// CHECK25-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41804// CHECK25-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41805// CHECK25-NEXT: [[I3:%.*]] = alloca i32, align 41806// CHECK25-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81807// CHECK25-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81808// CHECK25-NEXT: store ptr [[N]], ptr [[N_ADDR]], align 81809// CHECK25-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR]], align 81810// CHECK25-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81811// CHECK25-NEXT: [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 81812// CHECK25-NEXT: [[TMP1:%.*]] = load i64, ptr [[VLA_ADDR]], align 81813// CHECK25-NEXT: [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 81814// CHECK25-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 41815// CHECK25-NEXT: store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR_]], align 41816// CHECK25-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41817// CHECK25-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP4]], 01818// CHECK25-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11819// CHECK25-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11820// CHECK25-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41821// CHECK25-NEXT: store i32 0, ptr [[I]], align 41822// CHECK25-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41823// CHECK25-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP5]]1824// CHECK25-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1825// CHECK25: omp.precond.then:1826// CHECK25-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 41827// CHECK25-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41828// CHECK25-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_COMB_UB]], align 41829// CHECK25-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41830// CHECK25-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41831// CHECK25-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81832// CHECK25-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41833// CHECK25-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP8]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1834// CHECK25-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41835// CHECK25-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41836// CHECK25-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]1837// CHECK25-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1838// CHECK25: cond.true:1839// CHECK25-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41840// CHECK25-NEXT: br label [[COND_END:%.*]]1841// CHECK25: cond.false:1842// CHECK25-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41843// CHECK25-NEXT: br label [[COND_END]]1844// CHECK25: cond.end:1845// CHECK25-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]1846// CHECK25-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41847// CHECK25-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41848// CHECK25-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 41849// CHECK25-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1850// CHECK25: omp.inner.for.cond:1851// CHECK25-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41852// CHECK25-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41853// CHECK25-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]1854// CHECK25-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1855// CHECK25: omp.inner.for.body:1856// CHECK25-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41857// CHECK25-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 11858// CHECK25-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1859// CHECK25-NEXT: store i32 [[ADD]], ptr [[I3]], align 41860// CHECK25-NEXT: [[TMP17:%.*]] = load i32, ptr [[I3]], align 41861// CHECK25-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP17]] to i641862// CHECK25-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i64 [[IDXPROM]]1863// CHECK25-NEXT: store i32 0, ptr [[ARRAYIDX]], align 41864// CHECK25-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1865// CHECK25: omp.body.continue:1866// CHECK25-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1867// CHECK25: omp.inner.for.inc:1868// CHECK25-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41869// CHECK25-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP18]], 11870// CHECK25-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 41871// CHECK25-NEXT: br label [[OMP_INNER_FOR_COND]]1872// CHECK25: omp.inner.for.end:1873// CHECK25-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1874// CHECK25: omp.loop.exit:1875// CHECK25-NEXT: [[TMP19:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81876// CHECK25-NEXT: [[TMP20:%.*]] = load i32, ptr [[TMP19]], align 41877// CHECK25-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP20]])1878// CHECK25-NEXT: br label [[OMP_PRECOND_END]]1879// CHECK25: omp.precond.end:1880// CHECK25-NEXT: ret void1881//1882//1883// CHECK25-LABEL: define {{[^@]+}}@_Z5tmainIiLi10EEiT_1884// CHECK25-SAME: (i32 noundef signext [[ARGC:%.*]]) #[[ATTR5:[0-9]+]] comdat {1885// CHECK25-NEXT: entry:1886// CHECK25-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 41887// CHECK25-NEXT: [[A:%.*]] = alloca [10 x i32], align 41888// CHECK25-NEXT: [[TE:%.*]] = alloca i32, align 41889// CHECK25-NEXT: [[TH:%.*]] = alloca i32, align 41890// CHECK25-NEXT: [[TE_CASTED:%.*]] = alloca i64, align 81891// CHECK25-NEXT: [[TH_CASTED:%.*]] = alloca i64, align 81892// CHECK25-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 81893// CHECK25-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 81894// CHECK25-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 81895// CHECK25-NEXT: [[TMP:%.*]] = alloca i32, align 41896// CHECK25-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81897// CHECK25-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 41898// CHECK25-NEXT: store i32 0, ptr [[TE]], align 41899// CHECK25-NEXT: store i32 128, ptr [[TH]], align 41900// CHECK25-NEXT: [[TMP0:%.*]] = load i32, ptr [[TE]], align 41901// CHECK25-NEXT: store i32 [[TMP0]], ptr [[TE_CASTED]], align 41902// CHECK25-NEXT: [[TMP1:%.*]] = load i64, ptr [[TE_CASTED]], align 81903// CHECK25-NEXT: [[TMP2:%.*]] = load i32, ptr [[TH]], align 41904// CHECK25-NEXT: store i32 [[TMP2]], ptr [[TH_CASTED]], align 41905// CHECK25-NEXT: [[TMP3:%.*]] = load i64, ptr [[TH_CASTED]], align 81906// CHECK25-NEXT: [[TMP4:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01907// CHECK25-NEXT: store i64 [[TMP1]], ptr [[TMP4]], align 81908// CHECK25-NEXT: [[TMP5:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01909// CHECK25-NEXT: store i64 [[TMP1]], ptr [[TMP5]], align 81910// CHECK25-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 01911// CHECK25-NEXT: store ptr null, ptr [[TMP6]], align 81912// CHECK25-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11913// CHECK25-NEXT: store i64 [[TMP3]], ptr [[TMP7]], align 81914// CHECK25-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11915// CHECK25-NEXT: store i64 [[TMP3]], ptr [[TMP8]], align 81916// CHECK25-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 11917// CHECK25-NEXT: store ptr null, ptr [[TMP9]], align 81918// CHECK25-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21919// CHECK25-NEXT: store ptr [[A]], ptr [[TMP10]], align 81920// CHECK25-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21921// CHECK25-NEXT: store ptr [[A]], ptr [[TMP11]], align 81922// CHECK25-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 21923// CHECK25-NEXT: store ptr null, ptr [[TMP12]], align 81924// CHECK25-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01925// CHECK25-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01926// CHECK25-NEXT: [[TMP15:%.*]] = load i32, ptr [[TE]], align 41927// CHECK25-NEXT: [[TMP16:%.*]] = load i32, ptr [[TH]], align 41928// CHECK25-NEXT: [[TMP17:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP15]], 01929// CHECK25-NEXT: [[TMP18:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP16]], 01930// CHECK25-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01931// CHECK25-NEXT: store i32 3, ptr [[TMP19]], align 41932// CHECK25-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11933// CHECK25-NEXT: store i32 3, ptr [[TMP20]], align 41934// CHECK25-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21935// CHECK25-NEXT: store ptr [[TMP13]], ptr [[TMP21]], align 81936// CHECK25-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31937// CHECK25-NEXT: store ptr [[TMP14]], ptr [[TMP22]], align 81938// CHECK25-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41939// CHECK25-NEXT: store ptr @.offload_sizes.1, ptr [[TMP23]], align 81940// CHECK25-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51941// CHECK25-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP24]], align 81942// CHECK25-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61943// CHECK25-NEXT: store ptr null, ptr [[TMP25]], align 81944// CHECK25-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71945// CHECK25-NEXT: store ptr null, ptr [[TMP26]], align 81946// CHECK25-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81947// CHECK25-NEXT: store i64 10, ptr [[TMP27]], align 81948// CHECK25-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91949// CHECK25-NEXT: store i64 0, ptr [[TMP28]], align 81950// CHECK25-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101951// CHECK25-NEXT: store [3 x i32] [[TMP17]], ptr [[TMP29]], align 41952// CHECK25-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111953// CHECK25-NEXT: store [3 x i32] [[TMP18]], ptr [[TMP30]], align 41954// CHECK25-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121955// CHECK25-NEXT: store i32 0, ptr [[TMP31]], align 41956// CHECK25-NEXT: [[TMP32:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 [[TMP15]], i32 [[TMP16]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l150.region_id, ptr [[KERNEL_ARGS]])1957// CHECK25-NEXT: [[TMP33:%.*]] = icmp ne i32 [[TMP32]], 01958// CHECK25-NEXT: br i1 [[TMP33]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1959// CHECK25: omp_offload.failed:1960// CHECK25-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l150(i64 [[TMP1]], i64 [[TMP3]], ptr [[A]]) #[[ATTR3]]1961// CHECK25-NEXT: br label [[OMP_OFFLOAD_CONT]]1962// CHECK25: omp_offload.cont:1963// CHECK25-NEXT: ret i32 01964//1965//1966// CHECK25-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l1501967// CHECK25-SAME: (i64 noundef [[TE:%.*]], i64 noundef [[TH:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {1968// CHECK25-NEXT: entry:1969// CHECK25-NEXT: [[TE_ADDR:%.*]] = alloca i64, align 81970// CHECK25-NEXT: [[TH_ADDR:%.*]] = alloca i64, align 81971// CHECK25-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81972// CHECK25-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB2]])1973// CHECK25-NEXT: store i64 [[TE]], ptr [[TE_ADDR]], align 81974// CHECK25-NEXT: store i64 [[TH]], ptr [[TH_ADDR]], align 81975// CHECK25-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81976// CHECK25-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 81977// CHECK25-NEXT: [[TMP2:%.*]] = load i32, ptr [[TE_ADDR]], align 41978// CHECK25-NEXT: [[TMP3:%.*]] = load i32, ptr [[TH_ADDR]], align 41979// CHECK25-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB2]], i32 [[TMP0]], i32 [[TMP2]], i32 [[TMP3]])1980// CHECK25-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l150.omp_outlined, ptr [[TMP1]])1981// CHECK25-NEXT: ret void1982//1983//1984// CHECK25-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l150.omp_outlined1985// CHECK25-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {1986// CHECK25-NEXT: entry:1987// CHECK25-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81988// CHECK25-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81989// CHECK25-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81990// CHECK25-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41991// CHECK25-NEXT: [[TMP:%.*]] = alloca i32, align 41992// CHECK25-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41993// CHECK25-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41994// CHECK25-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41995// CHECK25-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41996// CHECK25-NEXT: [[I:%.*]] = alloca i32, align 41997// CHECK25-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81998// CHECK25-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81999// CHECK25-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 82000// CHECK25-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 82001// CHECK25-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 42002// CHECK25-NEXT: store i32 9, ptr [[DOTOMP_COMB_UB]], align 42003// CHECK25-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42004// CHECK25-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42005// CHECK25-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82006// CHECK25-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 42007// CHECK25-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2008// CHECK25-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42009// CHECK25-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 92010// CHECK25-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2011// CHECK25: cond.true:2012// CHECK25-NEXT: br label [[COND_END:%.*]]2013// CHECK25: cond.false:2014// CHECK25-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42015// CHECK25-NEXT: br label [[COND_END]]2016// CHECK25: cond.end:2017// CHECK25-NEXT: [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]2018// CHECK25-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 42019// CHECK25-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 42020// CHECK25-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 42021// CHECK25-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2022// CHECK25: omp.inner.for.cond:2023// CHECK25-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42024// CHECK25-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42025// CHECK25-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]2026// CHECK25-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2027// CHECK25: omp.inner.for.body:2028// CHECK25-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42029// CHECK25-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 12030// CHECK25-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]2031// CHECK25-NEXT: store i32 [[ADD]], ptr [[I]], align 42032// CHECK25-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 42033// CHECK25-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP9]] to i642034// CHECK25-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]2035// CHECK25-NEXT: store i32 0, ptr [[ARRAYIDX]], align 42036// CHECK25-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]2037// CHECK25: omp.body.continue:2038// CHECK25-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2039// CHECK25: omp.inner.for.inc:2040// CHECK25-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42041// CHECK25-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP10]], 12042// CHECK25-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 42043// CHECK25-NEXT: br label [[OMP_INNER_FOR_COND]]2044// CHECK25: omp.inner.for.end:2045// CHECK25-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2046// CHECK25: omp.loop.exit:2047// CHECK25-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])2048// CHECK25-NEXT: ret void2049//2050//2051// CHECK27-LABEL: define {{[^@]+}}@main2052// CHECK27-SAME: (i32 noundef [[ARGC:%.*]], ptr noundef [[ARGV:%.*]]) #[[ATTR0:[0-9]+]] {2053// CHECK27-NEXT: entry:2054// CHECK27-NEXT: [[RETVAL:%.*]] = alloca i32, align 42055// CHECK27-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 42056// CHECK27-NEXT: [[ARGV_ADDR:%.*]] = alloca ptr, align 42057// CHECK27-NEXT: [[N:%.*]] = alloca i32, align 42058// CHECK27-NEXT: [[SAVED_STACK:%.*]] = alloca ptr, align 42059// CHECK27-NEXT: [[__VLA_EXPR0:%.*]] = alloca i32, align 42060// CHECK27-NEXT: [[N_CASTED:%.*]] = alloca i32, align 42061// CHECK27-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 42062// CHECK27-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 42063// CHECK27-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 42064// CHECK27-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 42065// CHECK27-NEXT: [[TMP:%.*]] = alloca i32, align 42066// CHECK27-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42067// CHECK27-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42068// CHECK27-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 82069// CHECK27-NEXT: store i32 0, ptr [[RETVAL]], align 42070// CHECK27-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 42071// CHECK27-NEXT: store ptr [[ARGV]], ptr [[ARGV_ADDR]], align 42072// CHECK27-NEXT: store i32 100, ptr [[N]], align 42073// CHECK27-NEXT: [[TMP0:%.*]] = load i32, ptr [[N]], align 42074// CHECK27-NEXT: [[TMP1:%.*]] = call ptr @llvm.stacksave.p0()2075// CHECK27-NEXT: store ptr [[TMP1]], ptr [[SAVED_STACK]], align 42076// CHECK27-NEXT: [[VLA:%.*]] = alloca i32, i32 [[TMP0]], align 42077// CHECK27-NEXT: store i32 [[TMP0]], ptr [[__VLA_EXPR0]], align 42078// CHECK27-NEXT: [[TMP2:%.*]] = load i32, ptr [[N]], align 42079// CHECK27-NEXT: store i32 [[TMP2]], ptr [[N_CASTED]], align 42080// CHECK27-NEXT: [[TMP3:%.*]] = load i32, ptr [[N_CASTED]], align 42081// CHECK27-NEXT: [[TMP4:%.*]] = mul nuw i32 [[TMP0]], 42082// CHECK27-NEXT: [[TMP5:%.*]] = sext i32 [[TMP4]] to i642083// CHECK27-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[DOTOFFLOAD_SIZES]], ptr align 4 @.offload_sizes, i32 24, i1 false)2084// CHECK27-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 02085// CHECK27-NEXT: store i32 [[TMP3]], ptr [[TMP6]], align 42086// CHECK27-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 02087// CHECK27-NEXT: store i32 [[TMP3]], ptr [[TMP7]], align 42088// CHECK27-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 02089// CHECK27-NEXT: store ptr null, ptr [[TMP8]], align 42090// CHECK27-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 12091// CHECK27-NEXT: store i32 [[TMP0]], ptr [[TMP9]], align 42092// CHECK27-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 12093// CHECK27-NEXT: store i32 [[TMP0]], ptr [[TMP10]], align 42094// CHECK27-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 12095// CHECK27-NEXT: store ptr null, ptr [[TMP11]], align 42096// CHECK27-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 22097// CHECK27-NEXT: store ptr [[VLA]], ptr [[TMP12]], align 42098// CHECK27-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 22099// CHECK27-NEXT: store ptr [[VLA]], ptr [[TMP13]], align 42100// CHECK27-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 22101// CHECK27-NEXT: store i64 [[TMP5]], ptr [[TMP14]], align 42102// CHECK27-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 22103// CHECK27-NEXT: store ptr null, ptr [[TMP15]], align 42104// CHECK27-NEXT: [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 02105// CHECK27-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 02106// CHECK27-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 02107// CHECK27-NEXT: [[TMP19:%.*]] = load i32, ptr [[N]], align 42108// CHECK27-NEXT: store i32 [[TMP19]], ptr [[DOTCAPTURE_EXPR_]], align 42109// CHECK27-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42110// CHECK27-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP20]], 02111// CHECK27-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12112// CHECK27-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12113// CHECK27-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42114// CHECK27-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42115// CHECK27-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP21]], 12116// CHECK27-NEXT: [[TMP22:%.*]] = zext i32 [[ADD]] to i642117// CHECK27-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 02118// CHECK27-NEXT: store i32 3, ptr [[TMP23]], align 42119// CHECK27-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12120// CHECK27-NEXT: store i32 3, ptr [[TMP24]], align 42121// CHECK27-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 22122// CHECK27-NEXT: store ptr [[TMP16]], ptr [[TMP25]], align 42123// CHECK27-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 32124// CHECK27-NEXT: store ptr [[TMP17]], ptr [[TMP26]], align 42125// CHECK27-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 42126// CHECK27-NEXT: store ptr [[TMP18]], ptr [[TMP27]], align 42127// CHECK27-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 52128// CHECK27-NEXT: store ptr @.offload_maptypes, ptr [[TMP28]], align 42129// CHECK27-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 62130// CHECK27-NEXT: store ptr null, ptr [[TMP29]], align 42131// CHECK27-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 72132// CHECK27-NEXT: store ptr null, ptr [[TMP30]], align 42133// CHECK27-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 82134// CHECK27-NEXT: store i64 [[TMP22]], ptr [[TMP31]], align 82135// CHECK27-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 92136// CHECK27-NEXT: store i64 0, ptr [[TMP32]], align 82137// CHECK27-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 102138// CHECK27-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP33]], align 42139// CHECK27-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 112140// CHECK27-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP34]], align 42141// CHECK27-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 122142// CHECK27-NEXT: store i32 0, ptr [[TMP35]], align 42143// CHECK27-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_l161.region_id, ptr [[KERNEL_ARGS]])2144// CHECK27-NEXT: [[TMP37:%.*]] = icmp ne i32 [[TMP36]], 02145// CHECK27-NEXT: br i1 [[TMP37]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]2146// CHECK27: omp_offload.failed:2147// CHECK27-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l161(i32 [[TMP3]], i32 [[TMP0]], ptr [[VLA]]) #[[ATTR3:[0-9]+]]2148// CHECK27-NEXT: br label [[OMP_OFFLOAD_CONT]]2149// CHECK27: omp_offload.cont:2150// CHECK27-NEXT: [[TMP38:%.*]] = load i32, ptr [[ARGC_ADDR]], align 42151// CHECK27-NEXT: [[CALL:%.*]] = call noundef i32 @_Z5tmainIiLi10EEiT_(i32 noundef [[TMP38]])2152// CHECK27-NEXT: store i32 [[CALL]], ptr [[RETVAL]], align 42153// CHECK27-NEXT: [[TMP39:%.*]] = load ptr, ptr [[SAVED_STACK]], align 42154// CHECK27-NEXT: call void @llvm.stackrestore.p0(ptr [[TMP39]])2155// CHECK27-NEXT: [[TMP40:%.*]] = load i32, ptr [[RETVAL]], align 42156// CHECK27-NEXT: ret i32 [[TMP40]]2157//2158//2159// CHECK27-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l1612160// CHECK27-SAME: (i32 noundef [[N:%.*]], i32 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2:[0-9]+]] {2161// CHECK27-NEXT: entry:2162// CHECK27-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42163// CHECK27-NEXT: [[VLA_ADDR:%.*]] = alloca i32, align 42164// CHECK27-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42165// CHECK27-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42166// CHECK27-NEXT: store i32 [[VLA]], ptr [[VLA_ADDR]], align 42167// CHECK27-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42168// CHECK27-NEXT: [[TMP0:%.*]] = load i32, ptr [[VLA_ADDR]], align 42169// CHECK27-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 42170// CHECK27-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l161.omp_outlined, ptr [[N_ADDR]], i32 [[TMP0]], ptr [[TMP1]])2171// CHECK27-NEXT: ret void2172//2173//2174// CHECK27-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l161.omp_outlined2175// CHECK27-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]], i32 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {2176// CHECK27-NEXT: entry:2177// CHECK27-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42178// CHECK27-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42179// CHECK27-NEXT: [[N_ADDR:%.*]] = alloca ptr, align 42180// CHECK27-NEXT: [[VLA_ADDR:%.*]] = alloca i32, align 42181// CHECK27-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42182// CHECK27-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42183// CHECK27-NEXT: [[TMP:%.*]] = alloca i32, align 42184// CHECK27-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42185// CHECK27-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42186// CHECK27-NEXT: [[I:%.*]] = alloca i32, align 42187// CHECK27-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 42188// CHECK27-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 42189// CHECK27-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42190// CHECK27-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42191// CHECK27-NEXT: [[I3:%.*]] = alloca i32, align 42192// CHECK27-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42193// CHECK27-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42194// CHECK27-NEXT: store ptr [[N]], ptr [[N_ADDR]], align 42195// CHECK27-NEXT: store i32 [[VLA]], ptr [[VLA_ADDR]], align 42196// CHECK27-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42197// CHECK27-NEXT: [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 42198// CHECK27-NEXT: [[TMP1:%.*]] = load i32, ptr [[VLA_ADDR]], align 42199// CHECK27-NEXT: [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 42200// CHECK27-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 42201// CHECK27-NEXT: store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR_]], align 42202// CHECK27-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42203// CHECK27-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP4]], 02204// CHECK27-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12205// CHECK27-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12206// CHECK27-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42207// CHECK27-NEXT: store i32 0, ptr [[I]], align 42208// CHECK27-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42209// CHECK27-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP5]]2210// CHECK27-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2211// CHECK27: omp.precond.then:2212// CHECK27-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 42213// CHECK27-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42214// CHECK27-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_COMB_UB]], align 42215// CHECK27-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42216// CHECK27-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42217// CHECK27-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42218// CHECK27-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 42219// CHECK27-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP8]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2220// CHECK27-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42221// CHECK27-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42222// CHECK27-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]2223// CHECK27-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2224// CHECK27: cond.true:2225// CHECK27-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42226// CHECK27-NEXT: br label [[COND_END:%.*]]2227// CHECK27: cond.false:2228// CHECK27-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42229// CHECK27-NEXT: br label [[COND_END]]2230// CHECK27: cond.end:2231// CHECK27-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]2232// CHECK27-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 42233// CHECK27-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 42234// CHECK27-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 42235// CHECK27-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2236// CHECK27: omp.inner.for.cond:2237// CHECK27-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42238// CHECK27-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42239// CHECK27-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]2240// CHECK27-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2241// CHECK27: omp.inner.for.body:2242// CHECK27-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42243// CHECK27-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 12244// CHECK27-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]2245// CHECK27-NEXT: store i32 [[ADD]], ptr [[I3]], align 42246// CHECK27-NEXT: [[TMP17:%.*]] = load i32, ptr [[I3]], align 42247// CHECK27-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i32 [[TMP17]]2248// CHECK27-NEXT: store i32 0, ptr [[ARRAYIDX]], align 42249// CHECK27-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]2250// CHECK27: omp.body.continue:2251// CHECK27-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2252// CHECK27: omp.inner.for.inc:2253// CHECK27-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42254// CHECK27-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP18]], 12255// CHECK27-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 42256// CHECK27-NEXT: br label [[OMP_INNER_FOR_COND]]2257// CHECK27: omp.inner.for.end:2258// CHECK27-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2259// CHECK27: omp.loop.exit:2260// CHECK27-NEXT: [[TMP19:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42261// CHECK27-NEXT: [[TMP20:%.*]] = load i32, ptr [[TMP19]], align 42262// CHECK27-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP20]])2263// CHECK27-NEXT: br label [[OMP_PRECOND_END]]2264// CHECK27: omp.precond.end:2265// CHECK27-NEXT: ret void2266//2267//2268// CHECK27-LABEL: define {{[^@]+}}@_Z5tmainIiLi10EEiT_2269// CHECK27-SAME: (i32 noundef [[ARGC:%.*]]) #[[ATTR5:[0-9]+]] comdat {2270// CHECK27-NEXT: entry:2271// CHECK27-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 42272// CHECK27-NEXT: [[A:%.*]] = alloca [10 x i32], align 42273// CHECK27-NEXT: [[TE:%.*]] = alloca i32, align 42274// CHECK27-NEXT: [[TH:%.*]] = alloca i32, align 42275// CHECK27-NEXT: [[TE_CASTED:%.*]] = alloca i32, align 42276// CHECK27-NEXT: [[TH_CASTED:%.*]] = alloca i32, align 42277// CHECK27-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 42278// CHECK27-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 42279// CHECK27-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 42280// CHECK27-NEXT: [[TMP:%.*]] = alloca i32, align 42281// CHECK27-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 82282// CHECK27-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 42283// CHECK27-NEXT: store i32 0, ptr [[TE]], align 42284// CHECK27-NEXT: store i32 128, ptr [[TH]], align 42285// CHECK27-NEXT: [[TMP0:%.*]] = load i32, ptr [[TE]], align 42286// CHECK27-NEXT: store i32 [[TMP0]], ptr [[TE_CASTED]], align 42287// CHECK27-NEXT: [[TMP1:%.*]] = load i32, ptr [[TE_CASTED]], align 42288// CHECK27-NEXT: [[TMP2:%.*]] = load i32, ptr [[TH]], align 42289// CHECK27-NEXT: store i32 [[TMP2]], ptr [[TH_CASTED]], align 42290// CHECK27-NEXT: [[TMP3:%.*]] = load i32, ptr [[TH_CASTED]], align 42291// CHECK27-NEXT: [[TMP4:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 02292// CHECK27-NEXT: store i32 [[TMP1]], ptr [[TMP4]], align 42293// CHECK27-NEXT: [[TMP5:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 02294// CHECK27-NEXT: store i32 [[TMP1]], ptr [[TMP5]], align 42295// CHECK27-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 02296// CHECK27-NEXT: store ptr null, ptr [[TMP6]], align 42297// CHECK27-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 12298// CHECK27-NEXT: store i32 [[TMP3]], ptr [[TMP7]], align 42299// CHECK27-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 12300// CHECK27-NEXT: store i32 [[TMP3]], ptr [[TMP8]], align 42301// CHECK27-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 12302// CHECK27-NEXT: store ptr null, ptr [[TMP9]], align 42303// CHECK27-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 22304// CHECK27-NEXT: store ptr [[A]], ptr [[TMP10]], align 42305// CHECK27-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 22306// CHECK27-NEXT: store ptr [[A]], ptr [[TMP11]], align 42307// CHECK27-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 22308// CHECK27-NEXT: store ptr null, ptr [[TMP12]], align 42309// CHECK27-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 02310// CHECK27-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 02311// CHECK27-NEXT: [[TMP15:%.*]] = load i32, ptr [[TE]], align 42312// CHECK27-NEXT: [[TMP16:%.*]] = load i32, ptr [[TH]], align 42313// CHECK27-NEXT: [[TMP17:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP15]], 02314// CHECK27-NEXT: [[TMP18:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP16]], 02315// CHECK27-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 02316// CHECK27-NEXT: store i32 3, ptr [[TMP19]], align 42317// CHECK27-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12318// CHECK27-NEXT: store i32 3, ptr [[TMP20]], align 42319// CHECK27-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 22320// CHECK27-NEXT: store ptr [[TMP13]], ptr [[TMP21]], align 42321// CHECK27-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 32322// CHECK27-NEXT: store ptr [[TMP14]], ptr [[TMP22]], align 42323// CHECK27-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 42324// CHECK27-NEXT: store ptr @.offload_sizes.1, ptr [[TMP23]], align 42325// CHECK27-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 52326// CHECK27-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP24]], align 42327// CHECK27-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 62328// CHECK27-NEXT: store ptr null, ptr [[TMP25]], align 42329// CHECK27-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 72330// CHECK27-NEXT: store ptr null, ptr [[TMP26]], align 42331// CHECK27-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 82332// CHECK27-NEXT: store i64 10, ptr [[TMP27]], align 82333// CHECK27-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 92334// CHECK27-NEXT: store i64 0, ptr [[TMP28]], align 82335// CHECK27-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 102336// CHECK27-NEXT: store [3 x i32] [[TMP17]], ptr [[TMP29]], align 42337// CHECK27-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 112338// CHECK27-NEXT: store [3 x i32] [[TMP18]], ptr [[TMP30]], align 42339// CHECK27-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 122340// CHECK27-NEXT: store i32 0, ptr [[TMP31]], align 42341// CHECK27-NEXT: [[TMP32:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 [[TMP15]], i32 [[TMP16]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l150.region_id, ptr [[KERNEL_ARGS]])2342// CHECK27-NEXT: [[TMP33:%.*]] = icmp ne i32 [[TMP32]], 02343// CHECK27-NEXT: br i1 [[TMP33]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]2344// CHECK27: omp_offload.failed:2345// CHECK27-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l150(i32 [[TMP1]], i32 [[TMP3]], ptr [[A]]) #[[ATTR3]]2346// CHECK27-NEXT: br label [[OMP_OFFLOAD_CONT]]2347// CHECK27: omp_offload.cont:2348// CHECK27-NEXT: ret i32 02349//2350//2351// CHECK27-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l1502352// CHECK27-SAME: (i32 noundef [[TE:%.*]], i32 noundef [[TH:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {2353// CHECK27-NEXT: entry:2354// CHECK27-NEXT: [[TE_ADDR:%.*]] = alloca i32, align 42355// CHECK27-NEXT: [[TH_ADDR:%.*]] = alloca i32, align 42356// CHECK27-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42357// CHECK27-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB2]])2358// CHECK27-NEXT: store i32 [[TE]], ptr [[TE_ADDR]], align 42359// CHECK27-NEXT: store i32 [[TH]], ptr [[TH_ADDR]], align 42360// CHECK27-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42361// CHECK27-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 42362// CHECK27-NEXT: [[TMP2:%.*]] = load i32, ptr [[TE_ADDR]], align 42363// CHECK27-NEXT: [[TMP3:%.*]] = load i32, ptr [[TH_ADDR]], align 42364// CHECK27-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB2]], i32 [[TMP0]], i32 [[TMP2]], i32 [[TMP3]])2365// CHECK27-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l150.omp_outlined, ptr [[TMP1]])2366// CHECK27-NEXT: ret void2367//2368//2369// CHECK27-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l150.omp_outlined2370// CHECK27-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {2371// CHECK27-NEXT: entry:2372// CHECK27-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42373// CHECK27-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42374// CHECK27-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42375// CHECK27-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42376// CHECK27-NEXT: [[TMP:%.*]] = alloca i32, align 42377// CHECK27-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 42378// CHECK27-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 42379// CHECK27-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42380// CHECK27-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42381// CHECK27-NEXT: [[I:%.*]] = alloca i32, align 42382// CHECK27-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42383// CHECK27-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42384// CHECK27-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42385// CHECK27-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42386// CHECK27-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 42387// CHECK27-NEXT: store i32 9, ptr [[DOTOMP_COMB_UB]], align 42388// CHECK27-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42389// CHECK27-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42390// CHECK27-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42391// CHECK27-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 42392// CHECK27-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2393// CHECK27-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42394// CHECK27-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 92395// CHECK27-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2396// CHECK27: cond.true:2397// CHECK27-NEXT: br label [[COND_END:%.*]]2398// CHECK27: cond.false:2399// CHECK27-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42400// CHECK27-NEXT: br label [[COND_END]]2401// CHECK27: cond.end:2402// CHECK27-NEXT: [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]2403// CHECK27-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 42404// CHECK27-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 42405// CHECK27-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 42406// CHECK27-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2407// CHECK27: omp.inner.for.cond:2408// CHECK27-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42409// CHECK27-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42410// CHECK27-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]2411// CHECK27-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2412// CHECK27: omp.inner.for.body:2413// CHECK27-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42414// CHECK27-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 12415// CHECK27-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]2416// CHECK27-NEXT: store i32 [[ADD]], ptr [[I]], align 42417// CHECK27-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 42418// CHECK27-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP9]]2419// CHECK27-NEXT: store i32 0, ptr [[ARRAYIDX]], align 42420// CHECK27-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]2421// CHECK27: omp.body.continue:2422// CHECK27-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2423// CHECK27: omp.inner.for.inc:2424// CHECK27-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42425// CHECK27-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP10]], 12426// CHECK27-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 42427// CHECK27-NEXT: br label [[OMP_INNER_FOR_COND]]2428// CHECK27: omp.inner.for.end:2429// CHECK27-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2430// CHECK27: omp.loop.exit:2431// CHECK27-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])2432// CHECK27-NEXT: ret void2433//2434