2435 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 29 #pragma omp target30 #pragma omp teams distribute num_teams(te), thread_limit(th)31 for(int i = 0; i < n; i++) {32 a[i] = 0;33 }34 35 #pragma omp target36 {{{37 #pragma omp teams distribute38 for(int i = 0; i < n; i++) {39 a[i] = 0;40 }41 }}}42 43 // outlined target regions44 45 46 47 48 return a[0];49}50 51#endif // CK152 53// Test host codegen.54// 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=CHECK955// RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s56// 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=CHECK957// 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=CHECK1158// RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s59// 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=CHECK1160 61// 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}}"62// 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 %s63// 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}}"64// 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}}"65// 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 %s66// 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}}"67#ifdef CK268 69int teams_local_arg(void) {70 int n = 100;71 int a[n];72 73 #pragma omp target74 #pragma omp teams distribute75 for(int i = 0; i < n; i++) {76 a[i] = 0;77 }78 79 // outlined target region80 81 82 return a[0];83}84#endif // CK285 86// Test host codegen.87// 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=CHECK1788// RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s89// 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=CHECK1790// 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=CHECK1991// RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s92// 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=CHECK1993 94// 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}}"95// 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 %s96// 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}}"97// 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}}"98// 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 %s99// 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}}"100#ifdef CK3101 102 103template <typename T, int X, long long Y>104struct SS{105 T a[X];106 float b;107 int foo(void) {108 109 #pragma omp target110 #pragma omp teams distribute111 for(int i = 0; i < X; i++) {112 a[i] = (T)0;113 }114 115 // outlined target region116 117 118 return a[0];119 }120};121 122int teams_template_struct(void) {123 SS<int, 123, 456> V;124 return V.foo();125 126}127#endif // CK3128 129// Test host codegen.130// 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=CHECK25131// RUN: %clang_cc1 -DCK4 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s132// 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=CHECK25133// 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=CHECK27134// RUN: %clang_cc1 -DCK4 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s135// 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=CHECK27136 137// 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}}"138// 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 %s139// 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}}"140// 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}}"141// 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 %s142// 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}}"143 144#ifdef CK4145 146template <typename T, int n>147int tmain(T argc) {148 T a[n];149 int te = n/128;150 int th = 128;151#pragma omp target152#pragma omp teams distribute num_teams(te) thread_limit(th)153 for(int i = 0; i < n; i++) {154 a[i] = (T)0;155 }156 return 0;157}158 159int main (int argc, char **argv) {160 int n = 100;161 int a[n];162#pragma omp target163#pragma omp teams distribute164 for(int i = 0; i < n; i++) {165 a[i] = 0;166 }167 return tmain<int, 10>(argc);168}169 170 171 172 173 174 175 176#endif // CK4177#endif178// CHECK1-LABEL: define {{[^@]+}}@_Z21teams_argument_globali179// CHECK1-SAME: (i32 noundef signext [[N:%.*]]) #[[ATTR0:[0-9]+]] {180// CHECK1-NEXT: entry:181// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i32, align 4182// CHECK1-NEXT: [[TE:%.*]] = alloca i32, align 4183// CHECK1-NEXT: [[TH:%.*]] = alloca i32, align 4184// CHECK1-NEXT: [[TE_CASTED:%.*]] = alloca i64, align 8185// CHECK1-NEXT: [[TH_CASTED:%.*]] = alloca i64, align 8186// CHECK1-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8187// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [4 x ptr], align 8188// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [4 x ptr], align 8189// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [4 x ptr], align 8190// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4191// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4192// CHECK1-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4193// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8194// CHECK1-NEXT: [[N_CASTED4:%.*]] = alloca i64, align 8195// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS5:%.*]] = alloca [2 x ptr], align 8196// CHECK1-NEXT: [[DOTOFFLOAD_PTRS6:%.*]] = alloca [2 x ptr], align 8197// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS7:%.*]] = alloca [2 x ptr], align 8198// CHECK1-NEXT: [[_TMP8:%.*]] = alloca i32, align 4199// CHECK1-NEXT: [[DOTCAPTURE_EXPR_9:%.*]] = alloca i32, align 4200// CHECK1-NEXT: [[DOTCAPTURE_EXPR_10:%.*]] = alloca i32, align 4201// CHECK1-NEXT: [[KERNEL_ARGS15:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8202// CHECK1-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 4203// CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4204// CHECK1-NEXT: [[DIV:%.*]] = sdiv i32 [[TMP0]], 128205// CHECK1-NEXT: store i32 [[DIV]], ptr [[TE]], align 4206// CHECK1-NEXT: store i32 128, ptr [[TH]], align 4207// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TE]], align 4208// CHECK1-NEXT: store i32 [[TMP1]], ptr [[TE_CASTED]], align 4209// CHECK1-NEXT: [[TMP2:%.*]] = load i64, ptr [[TE_CASTED]], align 8210// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[TH]], align 4211// CHECK1-NEXT: store i32 [[TMP3]], ptr [[TH_CASTED]], align 4212// CHECK1-NEXT: [[TMP4:%.*]] = load i64, ptr [[TH_CASTED]], align 8213// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[N_ADDR]], align 4214// CHECK1-NEXT: store i32 [[TMP5]], ptr [[N_CASTED]], align 4215// CHECK1-NEXT: [[TMP6:%.*]] = load i64, ptr [[N_CASTED]], align 8216// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0217// CHECK1-NEXT: store i64 [[TMP2]], ptr [[TMP7]], align 8218// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0219// CHECK1-NEXT: store i64 [[TMP2]], ptr [[TMP8]], align 8220// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0221// CHECK1-NEXT: store ptr null, ptr [[TMP9]], align 8222// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1223// CHECK1-NEXT: store i64 [[TMP4]], ptr [[TMP10]], align 8224// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1225// CHECK1-NEXT: store i64 [[TMP4]], ptr [[TMP11]], align 8226// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1227// CHECK1-NEXT: store ptr null, ptr [[TMP12]], align 8228// CHECK1-NEXT: [[TMP13:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2229// CHECK1-NEXT: store i64 [[TMP6]], ptr [[TMP13]], align 8230// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2231// CHECK1-NEXT: store i64 [[TMP6]], ptr [[TMP14]], align 8232// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2233// CHECK1-NEXT: store ptr null, ptr [[TMP15]], align 8234// CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3235// CHECK1-NEXT: store ptr @a, ptr [[TMP16]], align 8236// CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3237// CHECK1-NEXT: store ptr @a, ptr [[TMP17]], align 8238// CHECK1-NEXT: [[TMP18:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3239// CHECK1-NEXT: store ptr null, ptr [[TMP18]], align 8240// CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0241// CHECK1-NEXT: [[TMP20:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0242// CHECK1-NEXT: [[TMP21:%.*]] = load i32, ptr [[TE]], align 4243// CHECK1-NEXT: [[TMP22:%.*]] = load i32, ptr [[TH]], align 4244// CHECK1-NEXT: [[TMP23:%.*]] = load i32, ptr [[N_ADDR]], align 4245// CHECK1-NEXT: store i32 [[TMP23]], ptr [[DOTCAPTURE_EXPR_]], align 4246// CHECK1-NEXT: [[TMP24:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4247// CHECK1-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP24]], 0248// CHECK1-NEXT: [[DIV2:%.*]] = sdiv i32 [[SUB]], 1249// CHECK1-NEXT: [[SUB3:%.*]] = sub nsw i32 [[DIV2]], 1250// CHECK1-NEXT: store i32 [[SUB3]], ptr [[DOTCAPTURE_EXPR_1]], align 4251// CHECK1-NEXT: [[TMP25:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4252// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP25]], 1253// CHECK1-NEXT: [[TMP26:%.*]] = zext i32 [[ADD]] to i64254// CHECK1-NEXT: [[TMP27:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP21]], 0255// CHECK1-NEXT: [[TMP28:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP22]], 0256// CHECK1-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0257// CHECK1-NEXT: store i32 3, ptr [[TMP29]], align 4258// CHECK1-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1259// CHECK1-NEXT: store i32 4, ptr [[TMP30]], align 4260// CHECK1-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2261// CHECK1-NEXT: store ptr [[TMP19]], ptr [[TMP31]], align 8262// CHECK1-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3263// CHECK1-NEXT: store ptr [[TMP20]], ptr [[TMP32]], align 8264// CHECK1-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4265// CHECK1-NEXT: store ptr @.offload_sizes, ptr [[TMP33]], align 8266// CHECK1-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5267// CHECK1-NEXT: store ptr @.offload_maptypes, ptr [[TMP34]], align 8268// CHECK1-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6269// CHECK1-NEXT: store ptr null, ptr [[TMP35]], align 8270// CHECK1-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7271// CHECK1-NEXT: store ptr null, ptr [[TMP36]], align 8272// CHECK1-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8273// CHECK1-NEXT: store i64 [[TMP26]], ptr [[TMP37]], align 8274// CHECK1-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9275// CHECK1-NEXT: store i64 0, ptr [[TMP38]], align 8276// CHECK1-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10277// CHECK1-NEXT: store [3 x i32] [[TMP27]], ptr [[TMP39]], align 4278// CHECK1-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11279// CHECK1-NEXT: store [3 x i32] [[TMP28]], ptr [[TMP40]], align 4280// CHECK1-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12281// CHECK1-NEXT: store i32 0, ptr [[TMP41]], align 4282// 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_l29.region_id, ptr [[KERNEL_ARGS]])283// CHECK1-NEXT: [[TMP43:%.*]] = icmp ne i32 [[TMP42]], 0284// CHECK1-NEXT: br i1 [[TMP43]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]285// CHECK1: omp_offload.failed:286// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l29(i64 [[TMP2]], i64 [[TMP4]], i64 [[TMP6]], ptr @a) #[[ATTR2:[0-9]+]]287// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]288// CHECK1: omp_offload.cont:289// CHECK1-NEXT: [[TMP44:%.*]] = load i32, ptr [[N_ADDR]], align 4290// CHECK1-NEXT: store i32 [[TMP44]], ptr [[N_CASTED4]], align 4291// CHECK1-NEXT: [[TMP45:%.*]] = load i64, ptr [[N_CASTED4]], align 8292// CHECK1-NEXT: [[TMP46:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0293// CHECK1-NEXT: store i64 [[TMP45]], ptr [[TMP46]], align 8294// CHECK1-NEXT: [[TMP47:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS6]], i32 0, i32 0295// CHECK1-NEXT: store i64 [[TMP45]], ptr [[TMP47]], align 8296// CHECK1-NEXT: [[TMP48:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS7]], i64 0, i64 0297// CHECK1-NEXT: store ptr null, ptr [[TMP48]], align 8298// CHECK1-NEXT: [[TMP49:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 1299// CHECK1-NEXT: store ptr @a, ptr [[TMP49]], align 8300// CHECK1-NEXT: [[TMP50:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS6]], i32 0, i32 1301// CHECK1-NEXT: store ptr @a, ptr [[TMP50]], align 8302// CHECK1-NEXT: [[TMP51:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS7]], i64 0, i64 1303// CHECK1-NEXT: store ptr null, ptr [[TMP51]], align 8304// CHECK1-NEXT: [[TMP52:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0305// CHECK1-NEXT: [[TMP53:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS6]], i32 0, i32 0306// CHECK1-NEXT: [[TMP54:%.*]] = load i32, ptr [[N_ADDR]], align 4307// CHECK1-NEXT: store i32 [[TMP54]], ptr [[DOTCAPTURE_EXPR_9]], align 4308// CHECK1-NEXT: [[TMP55:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_9]], align 4309// CHECK1-NEXT: [[SUB11:%.*]] = sub nsw i32 [[TMP55]], 0310// CHECK1-NEXT: [[DIV12:%.*]] = sdiv i32 [[SUB11]], 1311// CHECK1-NEXT: [[SUB13:%.*]] = sub nsw i32 [[DIV12]], 1312// CHECK1-NEXT: store i32 [[SUB13]], ptr [[DOTCAPTURE_EXPR_10]], align 4313// CHECK1-NEXT: [[TMP56:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_10]], align 4314// CHECK1-NEXT: [[ADD14:%.*]] = add nsw i32 [[TMP56]], 1315// CHECK1-NEXT: [[TMP57:%.*]] = zext i32 [[ADD14]] to i64316// CHECK1-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 0317// CHECK1-NEXT: store i32 3, ptr [[TMP58]], align 4318// CHECK1-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 1319// CHECK1-NEXT: store i32 2, ptr [[TMP59]], align 4320// CHECK1-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 2321// CHECK1-NEXT: store ptr [[TMP52]], ptr [[TMP60]], align 8322// CHECK1-NEXT: [[TMP61:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 3323// CHECK1-NEXT: store ptr [[TMP53]], ptr [[TMP61]], align 8324// CHECK1-NEXT: [[TMP62:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 4325// CHECK1-NEXT: store ptr @.offload_sizes.1, ptr [[TMP62]], align 8326// CHECK1-NEXT: [[TMP63:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 5327// CHECK1-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP63]], align 8328// CHECK1-NEXT: [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 6329// CHECK1-NEXT: store ptr null, ptr [[TMP64]], align 8330// CHECK1-NEXT: [[TMP65:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 7331// CHECK1-NEXT: store ptr null, ptr [[TMP65]], align 8332// CHECK1-NEXT: [[TMP66:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 8333// CHECK1-NEXT: store i64 [[TMP57]], ptr [[TMP66]], align 8334// CHECK1-NEXT: [[TMP67:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 9335// CHECK1-NEXT: store i64 0, ptr [[TMP67]], align 8336// CHECK1-NEXT: [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 10337// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP68]], align 4338// CHECK1-NEXT: [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 11339// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP69]], align 4340// CHECK1-NEXT: [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 12341// CHECK1-NEXT: store i32 0, ptr [[TMP70]], align 4342// 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_l35.region_id, ptr [[KERNEL_ARGS15]])343// CHECK1-NEXT: [[TMP72:%.*]] = icmp ne i32 [[TMP71]], 0344// CHECK1-NEXT: br i1 [[TMP72]], label [[OMP_OFFLOAD_FAILED16:%.*]], label [[OMP_OFFLOAD_CONT17:%.*]]345// CHECK1: omp_offload.failed16:346// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l35(i64 [[TMP45]], ptr @a) #[[ATTR2]]347// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT17]]348// CHECK1: omp_offload.cont17:349// CHECK1-NEXT: [[TMP73:%.*]] = load i32, ptr @a, align 4350// CHECK1-NEXT: ret i32 [[TMP73]]351//352//353// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l29354// CHECK1-SAME: (i64 noundef [[TE:%.*]], i64 noundef [[TH:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(400) [[A:%.*]]) #[[ATTR1:[0-9]+]] {355// CHECK1-NEXT: entry:356// CHECK1-NEXT: [[TE_ADDR:%.*]] = alloca i64, align 8357// CHECK1-NEXT: [[TH_ADDR:%.*]] = alloca i64, align 8358// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8359// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8360// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB2]])361// CHECK1-NEXT: store i64 [[TE]], ptr [[TE_ADDR]], align 8362// CHECK1-NEXT: store i64 [[TH]], ptr [[TH_ADDR]], align 8363// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8364// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8365// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 8366// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[TE_ADDR]], align 4367// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[TH_ADDR]], align 4368// CHECK1-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB2]], i32 [[TMP0]], i32 [[TMP2]], i32 [[TMP3]])369// 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_l29.omp_outlined, ptr [[N_ADDR]], ptr [[TMP1]])370// CHECK1-NEXT: ret void371//372//373// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l29.omp_outlined374// 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]] {375// CHECK1-NEXT: entry:376// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8377// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8378// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca ptr, align 8379// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8380// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4381// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4382// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4383// CHECK1-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4384// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4385// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4386// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4387// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4388// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4389// CHECK1-NEXT: [[I3:%.*]] = alloca i32, align 4390// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8391// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8392// CHECK1-NEXT: store ptr [[N]], ptr [[N_ADDR]], align 8393// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8394// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 8395// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 8396// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP0]], align 4397// CHECK1-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4398// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4399// CHECK1-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0400// CHECK1-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1401// CHECK1-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1402// CHECK1-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4403// CHECK1-NEXT: store i32 0, ptr [[I]], align 4404// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4405// CHECK1-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]406// CHECK1-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]407// CHECK1: omp.precond.then:408// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4409// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4410// CHECK1-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_UB]], align 4411// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4412// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4413// CHECK1-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8414// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 4415// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP7]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)416// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4417// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4418// CHECK1-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]419// CHECK1-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]420// CHECK1: cond.true:421// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4422// CHECK1-NEXT: br label [[COND_END:%.*]]423// CHECK1: cond.false:424// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4425// CHECK1-NEXT: br label [[COND_END]]426// CHECK1: cond.end:427// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]428// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4429// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4430// CHECK1-NEXT: store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 4431// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]432// CHECK1: omp.inner.for.cond:433// CHECK1-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4434// CHECK1-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4435// CHECK1-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]436// CHECK1-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]437// CHECK1: omp.inner.for.body:438// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4439// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP15]], 1440// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]441// CHECK1-NEXT: store i32 [[ADD]], ptr [[I3]], align 4442// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[I3]], align 4443// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP16]] to i64444// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [100 x i32], ptr [[TMP1]], i64 0, i64 [[IDXPROM]]445// CHECK1-NEXT: store i32 0, ptr [[ARRAYIDX]], align 4446// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]447// CHECK1: omp.body.continue:448// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]449// CHECK1: omp.inner.for.inc:450// CHECK1-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4451// CHECK1-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP17]], 1452// CHECK1-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4453// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]454// CHECK1: omp.inner.for.end:455// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]456// CHECK1: omp.loop.exit:457// CHECK1-NEXT: [[TMP18:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8458// CHECK1-NEXT: [[TMP19:%.*]] = load i32, ptr [[TMP18]], align 4459// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP19]])460// CHECK1-NEXT: br label [[OMP_PRECOND_END]]461// CHECK1: omp.precond.end:462// CHECK1-NEXT: ret void463//464//465// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l35466// CHECK1-SAME: (i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(400) [[A:%.*]]) #[[ATTR1]] {467// CHECK1-NEXT: entry:468// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8469// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8470// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8471// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8472// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8473// 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_l35.omp_outlined, ptr [[N_ADDR]], ptr [[TMP0]])474// CHECK1-NEXT: ret void475//476//477// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l35.omp_outlined478// 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]] {479// CHECK1-NEXT: entry:480// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8481// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8482// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca ptr, align 8483// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8484// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4485// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4486// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4487// CHECK1-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4488// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4489// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4490// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4491// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4492// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4493// CHECK1-NEXT: [[I3:%.*]] = alloca i32, align 4494// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8495// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8496// CHECK1-NEXT: store ptr [[N]], ptr [[N_ADDR]], align 8497// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8498// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 8499// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 8500// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP0]], align 4501// CHECK1-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4502// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4503// CHECK1-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0504// CHECK1-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1505// CHECK1-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1506// CHECK1-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4507// CHECK1-NEXT: store i32 0, ptr [[I]], align 4508// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4509// CHECK1-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]510// CHECK1-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]511// CHECK1: omp.precond.then:512// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4513// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4514// CHECK1-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_UB]], align 4515// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4516// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4517// CHECK1-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8518// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 4519// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP7]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)520// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4521// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4522// CHECK1-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]523// CHECK1-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]524// CHECK1: cond.true:525// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4526// CHECK1-NEXT: br label [[COND_END:%.*]]527// CHECK1: cond.false:528// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4529// CHECK1-NEXT: br label [[COND_END]]530// CHECK1: cond.end:531// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]532// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4533// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4534// CHECK1-NEXT: store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 4535// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]536// CHECK1: omp.inner.for.cond:537// CHECK1-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4538// CHECK1-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4539// CHECK1-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]540// CHECK1-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]541// CHECK1: omp.inner.for.body:542// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4543// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP15]], 1544// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]545// CHECK1-NEXT: store i32 [[ADD]], ptr [[I3]], align 4546// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[I3]], align 4547// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP16]] to i64548// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [100 x i32], ptr [[TMP1]], i64 0, i64 [[IDXPROM]]549// CHECK1-NEXT: store i32 0, ptr [[ARRAYIDX]], align 4550// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]551// CHECK1: omp.body.continue:552// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]553// CHECK1: omp.inner.for.inc:554// CHECK1-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4555// CHECK1-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP17]], 1556// CHECK1-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4557// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]558// CHECK1: omp.inner.for.end:559// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]560// CHECK1: omp.loop.exit:561// CHECK1-NEXT: [[TMP18:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8562// CHECK1-NEXT: [[TMP19:%.*]] = load i32, ptr [[TMP18]], align 4563// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP19]])564// CHECK1-NEXT: br label [[OMP_PRECOND_END]]565// CHECK1: omp.precond.end:566// CHECK1-NEXT: ret void567//568//569// CHECK3-LABEL: define {{[^@]+}}@_Z21teams_argument_globali570// CHECK3-SAME: (i32 noundef [[N:%.*]]) #[[ATTR0:[0-9]+]] {571// CHECK3-NEXT: entry:572// CHECK3-NEXT: [[N_ADDR:%.*]] = alloca i32, align 4573// CHECK3-NEXT: [[TE:%.*]] = alloca i32, align 4574// CHECK3-NEXT: [[TH:%.*]] = alloca i32, align 4575// CHECK3-NEXT: [[TE_CASTED:%.*]] = alloca i32, align 4576// CHECK3-NEXT: [[TH_CASTED:%.*]] = alloca i32, align 4577// CHECK3-NEXT: [[N_CASTED:%.*]] = alloca i32, align 4578// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [4 x ptr], align 4579// CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [4 x ptr], align 4580// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [4 x ptr], align 4581// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4582// CHECK3-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4583// CHECK3-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4584// CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8585// CHECK3-NEXT: [[N_CASTED4:%.*]] = alloca i32, align 4586// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS5:%.*]] = alloca [2 x ptr], align 4587// CHECK3-NEXT: [[DOTOFFLOAD_PTRS6:%.*]] = alloca [2 x ptr], align 4588// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS7:%.*]] = alloca [2 x ptr], align 4589// CHECK3-NEXT: [[_TMP8:%.*]] = alloca i32, align 4590// CHECK3-NEXT: [[DOTCAPTURE_EXPR_9:%.*]] = alloca i32, align 4591// CHECK3-NEXT: [[DOTCAPTURE_EXPR_10:%.*]] = alloca i32, align 4592// CHECK3-NEXT: [[KERNEL_ARGS15:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8593// CHECK3-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 4594// CHECK3-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4595// CHECK3-NEXT: [[DIV:%.*]] = sdiv i32 [[TMP0]], 128596// CHECK3-NEXT: store i32 [[DIV]], ptr [[TE]], align 4597// CHECK3-NEXT: store i32 128, ptr [[TH]], align 4598// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[TE]], align 4599// CHECK3-NEXT: store i32 [[TMP1]], ptr [[TE_CASTED]], align 4600// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[TE_CASTED]], align 4601// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[TH]], align 4602// CHECK3-NEXT: store i32 [[TMP3]], ptr [[TH_CASTED]], align 4603// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[TH_CASTED]], align 4604// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[N_ADDR]], align 4605// CHECK3-NEXT: store i32 [[TMP5]], ptr [[N_CASTED]], align 4606// CHECK3-NEXT: [[TMP6:%.*]] = load i32, ptr [[N_CASTED]], align 4607// CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0608// CHECK3-NEXT: store i32 [[TMP2]], ptr [[TMP7]], align 4609// CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0610// CHECK3-NEXT: store i32 [[TMP2]], ptr [[TMP8]], align 4611// CHECK3-NEXT: [[TMP9:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0612// CHECK3-NEXT: store ptr null, ptr [[TMP9]], align 4613// CHECK3-NEXT: [[TMP10:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1614// CHECK3-NEXT: store i32 [[TMP4]], ptr [[TMP10]], align 4615// CHECK3-NEXT: [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1616// CHECK3-NEXT: store i32 [[TMP4]], ptr [[TMP11]], align 4617// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 1618// CHECK3-NEXT: store ptr null, ptr [[TMP12]], align 4619// CHECK3-NEXT: [[TMP13:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2620// CHECK3-NEXT: store i32 [[TMP6]], ptr [[TMP13]], align 4621// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2622// CHECK3-NEXT: store i32 [[TMP6]], ptr [[TMP14]], align 4623// CHECK3-NEXT: [[TMP15:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 2624// CHECK3-NEXT: store ptr null, ptr [[TMP15]], align 4625// CHECK3-NEXT: [[TMP16:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3626// CHECK3-NEXT: store ptr @a, ptr [[TMP16]], align 4627// CHECK3-NEXT: [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3628// CHECK3-NEXT: store ptr @a, ptr [[TMP17]], align 4629// CHECK3-NEXT: [[TMP18:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 3630// CHECK3-NEXT: store ptr null, ptr [[TMP18]], align 4631// CHECK3-NEXT: [[TMP19:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0632// CHECK3-NEXT: [[TMP20:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0633// CHECK3-NEXT: [[TMP21:%.*]] = load i32, ptr [[TE]], align 4634// CHECK3-NEXT: [[TMP22:%.*]] = load i32, ptr [[TH]], align 4635// CHECK3-NEXT: [[TMP23:%.*]] = load i32, ptr [[N_ADDR]], align 4636// CHECK3-NEXT: store i32 [[TMP23]], ptr [[DOTCAPTURE_EXPR_]], align 4637// CHECK3-NEXT: [[TMP24:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4638// CHECK3-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP24]], 0639// CHECK3-NEXT: [[DIV2:%.*]] = sdiv i32 [[SUB]], 1640// CHECK3-NEXT: [[SUB3:%.*]] = sub nsw i32 [[DIV2]], 1641// CHECK3-NEXT: store i32 [[SUB3]], ptr [[DOTCAPTURE_EXPR_1]], align 4642// CHECK3-NEXT: [[TMP25:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4643// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP25]], 1644// CHECK3-NEXT: [[TMP26:%.*]] = zext i32 [[ADD]] to i64645// CHECK3-NEXT: [[TMP27:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP21]], 0646// CHECK3-NEXT: [[TMP28:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP22]], 0647// CHECK3-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0648// CHECK3-NEXT: store i32 3, ptr [[TMP29]], align 4649// CHECK3-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1650// CHECK3-NEXT: store i32 4, ptr [[TMP30]], align 4651// CHECK3-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2652// CHECK3-NEXT: store ptr [[TMP19]], ptr [[TMP31]], align 4653// CHECK3-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3654// CHECK3-NEXT: store ptr [[TMP20]], ptr [[TMP32]], align 4655// CHECK3-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4656// CHECK3-NEXT: store ptr @.offload_sizes, ptr [[TMP33]], align 4657// CHECK3-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5658// CHECK3-NEXT: store ptr @.offload_maptypes, ptr [[TMP34]], align 4659// CHECK3-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6660// CHECK3-NEXT: store ptr null, ptr [[TMP35]], align 4661// CHECK3-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7662// CHECK3-NEXT: store ptr null, ptr [[TMP36]], align 4663// CHECK3-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8664// CHECK3-NEXT: store i64 [[TMP26]], ptr [[TMP37]], align 8665// CHECK3-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9666// CHECK3-NEXT: store i64 0, ptr [[TMP38]], align 8667// CHECK3-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10668// CHECK3-NEXT: store [3 x i32] [[TMP27]], ptr [[TMP39]], align 4669// CHECK3-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11670// CHECK3-NEXT: store [3 x i32] [[TMP28]], ptr [[TMP40]], align 4671// CHECK3-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12672// CHECK3-NEXT: store i32 0, ptr [[TMP41]], align 4673// 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_l29.region_id, ptr [[KERNEL_ARGS]])674// CHECK3-NEXT: [[TMP43:%.*]] = icmp ne i32 [[TMP42]], 0675// CHECK3-NEXT: br i1 [[TMP43]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]676// CHECK3: omp_offload.failed:677// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l29(i32 [[TMP2]], i32 [[TMP4]], i32 [[TMP6]], ptr @a) #[[ATTR2:[0-9]+]]678// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]]679// CHECK3: omp_offload.cont:680// CHECK3-NEXT: [[TMP44:%.*]] = load i32, ptr [[N_ADDR]], align 4681// CHECK3-NEXT: store i32 [[TMP44]], ptr [[N_CASTED4]], align 4682// CHECK3-NEXT: [[TMP45:%.*]] = load i32, ptr [[N_CASTED4]], align 4683// CHECK3-NEXT: [[TMP46:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0684// CHECK3-NEXT: store i32 [[TMP45]], ptr [[TMP46]], align 4685// CHECK3-NEXT: [[TMP47:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS6]], i32 0, i32 0686// CHECK3-NEXT: store i32 [[TMP45]], ptr [[TMP47]], align 4687// CHECK3-NEXT: [[TMP48:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS7]], i32 0, i32 0688// CHECK3-NEXT: store ptr null, ptr [[TMP48]], align 4689// CHECK3-NEXT: [[TMP49:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 1690// CHECK3-NEXT: store ptr @a, ptr [[TMP49]], align 4691// CHECK3-NEXT: [[TMP50:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS6]], i32 0, i32 1692// CHECK3-NEXT: store ptr @a, ptr [[TMP50]], align 4693// CHECK3-NEXT: [[TMP51:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS7]], i32 0, i32 1694// CHECK3-NEXT: store ptr null, ptr [[TMP51]], align 4695// CHECK3-NEXT: [[TMP52:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0696// CHECK3-NEXT: [[TMP53:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS6]], i32 0, i32 0697// CHECK3-NEXT: [[TMP54:%.*]] = load i32, ptr [[N_ADDR]], align 4698// CHECK3-NEXT: store i32 [[TMP54]], ptr [[DOTCAPTURE_EXPR_9]], align 4699// CHECK3-NEXT: [[TMP55:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_9]], align 4700// CHECK3-NEXT: [[SUB11:%.*]] = sub nsw i32 [[TMP55]], 0701// CHECK3-NEXT: [[DIV12:%.*]] = sdiv i32 [[SUB11]], 1702// CHECK3-NEXT: [[SUB13:%.*]] = sub nsw i32 [[DIV12]], 1703// CHECK3-NEXT: store i32 [[SUB13]], ptr [[DOTCAPTURE_EXPR_10]], align 4704// CHECK3-NEXT: [[TMP56:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_10]], align 4705// CHECK3-NEXT: [[ADD14:%.*]] = add nsw i32 [[TMP56]], 1706// CHECK3-NEXT: [[TMP57:%.*]] = zext i32 [[ADD14]] to i64707// CHECK3-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 0708// CHECK3-NEXT: store i32 3, ptr [[TMP58]], align 4709// CHECK3-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 1710// CHECK3-NEXT: store i32 2, ptr [[TMP59]], align 4711// CHECK3-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 2712// CHECK3-NEXT: store ptr [[TMP52]], ptr [[TMP60]], align 4713// CHECK3-NEXT: [[TMP61:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 3714// CHECK3-NEXT: store ptr [[TMP53]], ptr [[TMP61]], align 4715// CHECK3-NEXT: [[TMP62:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 4716// CHECK3-NEXT: store ptr @.offload_sizes.1, ptr [[TMP62]], align 4717// CHECK3-NEXT: [[TMP63:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 5718// CHECK3-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP63]], align 4719// CHECK3-NEXT: [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 6720// CHECK3-NEXT: store ptr null, ptr [[TMP64]], align 4721// CHECK3-NEXT: [[TMP65:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 7722// CHECK3-NEXT: store ptr null, ptr [[TMP65]], align 4723// CHECK3-NEXT: [[TMP66:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 8724// CHECK3-NEXT: store i64 [[TMP57]], ptr [[TMP66]], align 8725// CHECK3-NEXT: [[TMP67:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 9726// CHECK3-NEXT: store i64 0, ptr [[TMP67]], align 8727// CHECK3-NEXT: [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 10728// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP68]], align 4729// CHECK3-NEXT: [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 11730// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP69]], align 4731// CHECK3-NEXT: [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS15]], i32 0, i32 12732// CHECK3-NEXT: store i32 0, ptr [[TMP70]], align 4733// 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_l35.region_id, ptr [[KERNEL_ARGS15]])734// CHECK3-NEXT: [[TMP72:%.*]] = icmp ne i32 [[TMP71]], 0735// CHECK3-NEXT: br i1 [[TMP72]], label [[OMP_OFFLOAD_FAILED16:%.*]], label [[OMP_OFFLOAD_CONT17:%.*]]736// CHECK3: omp_offload.failed16:737// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l35(i32 [[TMP45]], ptr @a) #[[ATTR2]]738// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT17]]739// CHECK3: omp_offload.cont17:740// CHECK3-NEXT: [[TMP73:%.*]] = load i32, ptr @a, align 4741// CHECK3-NEXT: ret i32 [[TMP73]]742//743//744// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l29745// CHECK3-SAME: (i32 noundef [[TE:%.*]], i32 noundef [[TH:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(400) [[A:%.*]]) #[[ATTR1:[0-9]+]] {746// CHECK3-NEXT: entry:747// CHECK3-NEXT: [[TE_ADDR:%.*]] = alloca i32, align 4748// CHECK3-NEXT: [[TH_ADDR:%.*]] = alloca i32, align 4749// CHECK3-NEXT: [[N_ADDR:%.*]] = alloca i32, align 4750// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 4751// CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB2]])752// CHECK3-NEXT: store i32 [[TE]], ptr [[TE_ADDR]], align 4753// CHECK3-NEXT: store i32 [[TH]], ptr [[TH_ADDR]], align 4754// CHECK3-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 4755// CHECK3-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 4756// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 4757// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[TE_ADDR]], align 4758// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[TH_ADDR]], align 4759// CHECK3-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB2]], i32 [[TMP0]], i32 [[TMP2]], i32 [[TMP3]])760// 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_l29.omp_outlined, ptr [[N_ADDR]], ptr [[TMP1]])761// CHECK3-NEXT: ret void762//763//764// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l29.omp_outlined765// 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]] {766// CHECK3-NEXT: entry:767// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4768// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4769// CHECK3-NEXT: [[N_ADDR:%.*]] = alloca ptr, align 4770// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 4771// CHECK3-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4772// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4773// CHECK3-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4774// CHECK3-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4775// CHECK3-NEXT: [[I:%.*]] = alloca i32, align 4776// CHECK3-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4777// CHECK3-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4778// CHECK3-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4779// CHECK3-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4780// CHECK3-NEXT: [[I3:%.*]] = alloca i32, align 4781// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4782// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4783// CHECK3-NEXT: store ptr [[N]], ptr [[N_ADDR]], align 4784// CHECK3-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 4785// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 4786// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 4787// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP0]], align 4788// CHECK3-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4789// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4790// CHECK3-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0791// CHECK3-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1792// CHECK3-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1793// CHECK3-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4794// CHECK3-NEXT: store i32 0, ptr [[I]], align 4795// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4796// CHECK3-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]797// CHECK3-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]798// CHECK3: omp.precond.then:799// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4800// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4801// CHECK3-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_UB]], align 4802// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4803// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4804// CHECK3-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4805// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 4806// CHECK3-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP7]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)807// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4808// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4809// CHECK3-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]810// CHECK3-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]811// CHECK3: cond.true:812// CHECK3-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4813// CHECK3-NEXT: br label [[COND_END:%.*]]814// CHECK3: cond.false:815// CHECK3-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4816// CHECK3-NEXT: br label [[COND_END]]817// CHECK3: cond.end:818// CHECK3-NEXT: [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]819// CHECK3-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4820// CHECK3-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4821// CHECK3-NEXT: store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 4822// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]823// CHECK3: omp.inner.for.cond:824// CHECK3-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4825// CHECK3-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4826// CHECK3-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]827// CHECK3-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]828// CHECK3: omp.inner.for.body:829// CHECK3-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4830// CHECK3-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP15]], 1831// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]832// CHECK3-NEXT: store i32 [[ADD]], ptr [[I3]], align 4833// CHECK3-NEXT: [[TMP16:%.*]] = load i32, ptr [[I3]], align 4834// CHECK3-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [100 x i32], ptr [[TMP1]], i32 0, i32 [[TMP16]]835// CHECK3-NEXT: store i32 0, ptr [[ARRAYIDX]], align 4836// CHECK3-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]837// CHECK3: omp.body.continue:838// CHECK3-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]839// CHECK3: omp.inner.for.inc:840// CHECK3-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4841// CHECK3-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP17]], 1842// CHECK3-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4843// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND]]844// CHECK3: omp.inner.for.end:845// CHECK3-NEXT: br label [[OMP_LOOP_EXIT:%.*]]846// CHECK3: omp.loop.exit:847// CHECK3-NEXT: [[TMP18:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4848// CHECK3-NEXT: [[TMP19:%.*]] = load i32, ptr [[TMP18]], align 4849// CHECK3-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP19]])850// CHECK3-NEXT: br label [[OMP_PRECOND_END]]851// CHECK3: omp.precond.end:852// CHECK3-NEXT: ret void853//854//855// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l35856// CHECK3-SAME: (i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(400) [[A:%.*]]) #[[ATTR1]] {857// CHECK3-NEXT: entry:858// CHECK3-NEXT: [[N_ADDR:%.*]] = alloca i32, align 4859// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 4860// CHECK3-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 4861// CHECK3-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 4862// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 4863// 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_l35.omp_outlined, ptr [[N_ADDR]], ptr [[TMP0]])864// CHECK3-NEXT: ret void865//866//867// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z21teams_argument_globali_l35.omp_outlined868// 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]] {869// CHECK3-NEXT: entry:870// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4871// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4872// CHECK3-NEXT: [[N_ADDR:%.*]] = alloca ptr, align 4873// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 4874// CHECK3-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4875// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4876// CHECK3-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4877// CHECK3-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4878// CHECK3-NEXT: [[I:%.*]] = alloca i32, align 4879// CHECK3-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4880// CHECK3-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4881// CHECK3-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4882// CHECK3-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4883// CHECK3-NEXT: [[I3:%.*]] = alloca i32, align 4884// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4885// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4886// CHECK3-NEXT: store ptr [[N]], ptr [[N_ADDR]], align 4887// CHECK3-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 4888// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 4889// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 4890// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP0]], align 4891// CHECK3-NEXT: store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4892// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4893// CHECK3-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0894// CHECK3-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1895// CHECK3-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1896// CHECK3-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4897// CHECK3-NEXT: store i32 0, ptr [[I]], align 4898// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4899// CHECK3-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]900// CHECK3-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]901// CHECK3: omp.precond.then:902// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4903// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4904// CHECK3-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_UB]], align 4905// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4906// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4907// CHECK3-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4908// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 4909// CHECK3-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP7]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)910// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4911// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4912// CHECK3-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]913// CHECK3-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]914// CHECK3: cond.true:915// CHECK3-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4916// CHECK3-NEXT: br label [[COND_END:%.*]]917// CHECK3: cond.false:918// CHECK3-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4919// CHECK3-NEXT: br label [[COND_END]]920// CHECK3: cond.end:921// CHECK3-NEXT: [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]922// CHECK3-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4923// CHECK3-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4924// CHECK3-NEXT: store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 4925// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]926// CHECK3: omp.inner.for.cond:927// CHECK3-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4928// CHECK3-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4929// CHECK3-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]930// CHECK3-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]931// CHECK3: omp.inner.for.body:932// CHECK3-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4933// CHECK3-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP15]], 1934// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]935// CHECK3-NEXT: store i32 [[ADD]], ptr [[I3]], align 4936// CHECK3-NEXT: [[TMP16:%.*]] = load i32, ptr [[I3]], align 4937// CHECK3-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [100 x i32], ptr [[TMP1]], i32 0, i32 [[TMP16]]938// CHECK3-NEXT: store i32 0, ptr [[ARRAYIDX]], align 4939// CHECK3-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]940// CHECK3: omp.body.continue:941// CHECK3-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]942// CHECK3: omp.inner.for.inc:943// CHECK3-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4944// CHECK3-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP17]], 1945// CHECK3-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4946// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND]]947// CHECK3: omp.inner.for.end:948// CHECK3-NEXT: br label [[OMP_LOOP_EXIT:%.*]]949// CHECK3: omp.loop.exit:950// CHECK3-NEXT: [[TMP18:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4951// CHECK3-NEXT: [[TMP19:%.*]] = load i32, ptr [[TMP18]], align 4952// CHECK3-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP19]])953// CHECK3-NEXT: br label [[OMP_PRECOND_END]]954// CHECK3: omp.precond.end:955// CHECK3-NEXT: ret void956//957//958// CHECK9-LABEL: define {{[^@]+}}@_Z15teams_local_argv959// CHECK9-SAME: () #[[ATTR0:[0-9]+]] {960// CHECK9-NEXT: entry:961// CHECK9-NEXT: [[N:%.*]] = alloca i32, align 4962// CHECK9-NEXT: [[SAVED_STACK:%.*]] = alloca ptr, align 8963// CHECK9-NEXT: [[__VLA_EXPR0:%.*]] = alloca i64, align 8964// CHECK9-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8965// CHECK9-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 8966// CHECK9-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 8967// CHECK9-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 8968// CHECK9-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 8969// CHECK9-NEXT: [[TMP:%.*]] = alloca i32, align 4970// CHECK9-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4971// CHECK9-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4972// CHECK9-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8973// CHECK9-NEXT: store i32 100, ptr [[N]], align 4974// CHECK9-NEXT: [[TMP0:%.*]] = load i32, ptr [[N]], align 4975// CHECK9-NEXT: [[TMP1:%.*]] = zext i32 [[TMP0]] to i64976// CHECK9-NEXT: [[TMP2:%.*]] = call ptr @llvm.stacksave.p0()977// CHECK9-NEXT: store ptr [[TMP2]], ptr [[SAVED_STACK]], align 8978// CHECK9-NEXT: [[VLA:%.*]] = alloca i32, i64 [[TMP1]], align 4979// CHECK9-NEXT: store i64 [[TMP1]], ptr [[__VLA_EXPR0]], align 8980// CHECK9-NEXT: [[TMP3:%.*]] = load i32, ptr [[N]], align 4981// CHECK9-NEXT: store i32 [[TMP3]], ptr [[N_CASTED]], align 4982// CHECK9-NEXT: [[TMP4:%.*]] = load i64, ptr [[N_CASTED]], align 8983// CHECK9-NEXT: [[TMP5:%.*]] = mul nuw i64 [[TMP1]], 4984// CHECK9-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES]], ptr align 8 @.offload_sizes, i64 24, i1 false)985// CHECK9-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0986// CHECK9-NEXT: store i64 [[TMP4]], ptr [[TMP6]], align 8987// CHECK9-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0988// CHECK9-NEXT: store i64 [[TMP4]], ptr [[TMP7]], align 8989// CHECK9-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0990// CHECK9-NEXT: store ptr null, ptr [[TMP8]], align 8991// CHECK9-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1992// CHECK9-NEXT: store i64 [[TMP1]], ptr [[TMP9]], align 8993// CHECK9-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1994// CHECK9-NEXT: store i64 [[TMP1]], ptr [[TMP10]], align 8995// CHECK9-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1996// CHECK9-NEXT: store ptr null, ptr [[TMP11]], align 8997// CHECK9-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2998// CHECK9-NEXT: store ptr [[VLA]], ptr [[TMP12]], align 8999// CHECK9-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21000// CHECK9-NEXT: store ptr [[VLA]], ptr [[TMP13]], align 81001// CHECK9-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 21002// CHECK9-NEXT: store i64 [[TMP5]], ptr [[TMP14]], align 81003// CHECK9-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 21004// CHECK9-NEXT: store ptr null, ptr [[TMP15]], align 81005// CHECK9-NEXT: [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01006// CHECK9-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01007// CHECK9-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 01008// CHECK9-NEXT: [[TMP19:%.*]] = load i32, ptr [[N]], align 41009// CHECK9-NEXT: store i32 [[TMP19]], ptr [[DOTCAPTURE_EXPR_]], align 41010// CHECK9-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41011// CHECK9-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP20]], 01012// CHECK9-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11013// CHECK9-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11014// CHECK9-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41015// CHECK9-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41016// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP21]], 11017// CHECK9-NEXT: [[TMP22:%.*]] = zext i32 [[ADD]] to i641018// CHECK9-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01019// CHECK9-NEXT: store i32 3, ptr [[TMP23]], align 41020// CHECK9-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11021// CHECK9-NEXT: store i32 3, ptr [[TMP24]], align 41022// CHECK9-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21023// CHECK9-NEXT: store ptr [[TMP16]], ptr [[TMP25]], align 81024// CHECK9-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31025// CHECK9-NEXT: store ptr [[TMP17]], ptr [[TMP26]], align 81026// CHECK9-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41027// CHECK9-NEXT: store ptr [[TMP18]], ptr [[TMP27]], align 81028// CHECK9-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51029// CHECK9-NEXT: store ptr @.offload_maptypes, ptr [[TMP28]], align 81030// CHECK9-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61031// CHECK9-NEXT: store ptr null, ptr [[TMP29]], align 81032// CHECK9-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71033// CHECK9-NEXT: store ptr null, ptr [[TMP30]], align 81034// CHECK9-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81035// CHECK9-NEXT: store i64 [[TMP22]], ptr [[TMP31]], align 81036// CHECK9-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91037// CHECK9-NEXT: store i64 0, ptr [[TMP32]], align 81038// CHECK9-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101039// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP33]], align 41040// CHECK9-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111041// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP34]], align 41042// CHECK9-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121043// CHECK9-NEXT: store i32 0, ptr [[TMP35]], align 41044// 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_l73.region_id, ptr [[KERNEL_ARGS]])1045// CHECK9-NEXT: [[TMP37:%.*]] = icmp ne i32 [[TMP36]], 01046// CHECK9-NEXT: br i1 [[TMP37]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1047// CHECK9: omp_offload.failed:1048// CHECK9-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z15teams_local_argv_l73(i64 [[TMP4]], i64 [[TMP1]], ptr [[VLA]]) #[[ATTR3:[0-9]+]]1049// CHECK9-NEXT: br label [[OMP_OFFLOAD_CONT]]1050// CHECK9: omp_offload.cont:1051// CHECK9-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[VLA]], i64 01052// CHECK9-NEXT: [[TMP38:%.*]] = load i32, ptr [[ARRAYIDX]], align 41053// CHECK9-NEXT: [[TMP39:%.*]] = load ptr, ptr [[SAVED_STACK]], align 81054// CHECK9-NEXT: call void @llvm.stackrestore.p0(ptr [[TMP39]])1055// CHECK9-NEXT: ret i32 [[TMP38]]1056//1057//1058// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z15teams_local_argv_l731059// CHECK9-SAME: (i64 noundef [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2:[0-9]+]] {1060// CHECK9-NEXT: entry:1061// CHECK9-NEXT: [[N_ADDR:%.*]] = alloca i64, align 81062// CHECK9-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 81063// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81064// CHECK9-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 81065// CHECK9-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR]], align 81066// CHECK9-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81067// CHECK9-NEXT: [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR]], align 81068// CHECK9-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 81069// 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_l73.omp_outlined, ptr [[N_ADDR]], i64 [[TMP0]], ptr [[TMP1]])1070// CHECK9-NEXT: ret void1071//1072//1073// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z15teams_local_argv_l73.omp_outlined1074// 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]] {1075// CHECK9-NEXT: entry:1076// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81077// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81078// CHECK9-NEXT: [[N_ADDR:%.*]] = alloca ptr, align 81079// CHECK9-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 81080// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81081// CHECK9-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41082// CHECK9-NEXT: [[TMP:%.*]] = alloca i32, align 41083// CHECK9-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41084// CHECK9-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41085// CHECK9-NEXT: [[I:%.*]] = alloca i32, align 41086// CHECK9-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41087// CHECK9-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41088// CHECK9-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41089// CHECK9-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41090// CHECK9-NEXT: [[I3:%.*]] = alloca i32, align 41091// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81092// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81093// CHECK9-NEXT: store ptr [[N]], ptr [[N_ADDR]], align 81094// CHECK9-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR]], align 81095// CHECK9-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81096// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 81097// CHECK9-NEXT: [[TMP1:%.*]] = load i64, ptr [[VLA_ADDR]], align 81098// CHECK9-NEXT: [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 81099// CHECK9-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 41100// CHECK9-NEXT: store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR_]], align 41101// CHECK9-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41102// CHECK9-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP4]], 01103// CHECK9-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11104// CHECK9-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11105// CHECK9-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41106// CHECK9-NEXT: store i32 0, ptr [[I]], align 41107// CHECK9-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41108// CHECK9-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP5]]1109// CHECK9-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1110// CHECK9: omp.precond.then:1111// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41112// CHECK9-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41113// CHECK9-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_UB]], align 41114// CHECK9-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41115// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41116// CHECK9-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81117// CHECK9-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41118// CHECK9-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP8]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1119// CHECK9-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41120// CHECK9-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41121// CHECK9-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]1122// CHECK9-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1123// CHECK9: cond.true:1124// CHECK9-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41125// CHECK9-NEXT: br label [[COND_END:%.*]]1126// CHECK9: cond.false:1127// CHECK9-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41128// CHECK9-NEXT: br label [[COND_END]]1129// CHECK9: cond.end:1130// CHECK9-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]1131// CHECK9-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41132// CHECK9-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41133// CHECK9-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 41134// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1135// CHECK9: omp.inner.for.cond:1136// CHECK9-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41137// CHECK9-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41138// CHECK9-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]1139// CHECK9-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1140// CHECK9: omp.inner.for.body:1141// CHECK9-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41142// CHECK9-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 11143// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1144// CHECK9-NEXT: store i32 [[ADD]], ptr [[I3]], align 41145// CHECK9-NEXT: [[TMP17:%.*]] = load i32, ptr [[I3]], align 41146// CHECK9-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP17]] to i641147// CHECK9-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i64 [[IDXPROM]]1148// CHECK9-NEXT: store i32 0, ptr [[ARRAYIDX]], align 41149// CHECK9-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1150// CHECK9: omp.body.continue:1151// CHECK9-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1152// CHECK9: omp.inner.for.inc:1153// CHECK9-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41154// CHECK9-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP18]], 11155// CHECK9-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 41156// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND]]1157// CHECK9: omp.inner.for.end:1158// CHECK9-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1159// CHECK9: omp.loop.exit:1160// CHECK9-NEXT: [[TMP19:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81161// CHECK9-NEXT: [[TMP20:%.*]] = load i32, ptr [[TMP19]], align 41162// CHECK9-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP20]])1163// CHECK9-NEXT: br label [[OMP_PRECOND_END]]1164// CHECK9: omp.precond.end:1165// CHECK9-NEXT: ret void1166//1167//1168// CHECK11-LABEL: define {{[^@]+}}@_Z15teams_local_argv1169// CHECK11-SAME: () #[[ATTR0:[0-9]+]] {1170// CHECK11-NEXT: entry:1171// CHECK11-NEXT: [[N:%.*]] = alloca i32, align 41172// CHECK11-NEXT: [[SAVED_STACK:%.*]] = alloca ptr, align 41173// CHECK11-NEXT: [[__VLA_EXPR0:%.*]] = alloca i32, align 41174// CHECK11-NEXT: [[N_CASTED:%.*]] = alloca i32, align 41175// CHECK11-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 41176// CHECK11-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 41177// CHECK11-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 41178// CHECK11-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 41179// CHECK11-NEXT: [[TMP:%.*]] = alloca i32, align 41180// CHECK11-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41181// CHECK11-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41182// CHECK11-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81183// CHECK11-NEXT: store i32 100, ptr [[N]], align 41184// CHECK11-NEXT: [[TMP0:%.*]] = load i32, ptr [[N]], align 41185// CHECK11-NEXT: [[TMP1:%.*]] = call ptr @llvm.stacksave.p0()1186// CHECK11-NEXT: store ptr [[TMP1]], ptr [[SAVED_STACK]], align 41187// CHECK11-NEXT: [[VLA:%.*]] = alloca i32, i32 [[TMP0]], align 41188// CHECK11-NEXT: store i32 [[TMP0]], ptr [[__VLA_EXPR0]], align 41189// CHECK11-NEXT: [[TMP2:%.*]] = load i32, ptr [[N]], align 41190// CHECK11-NEXT: store i32 [[TMP2]], ptr [[N_CASTED]], align 41191// CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[N_CASTED]], align 41192// CHECK11-NEXT: [[TMP4:%.*]] = mul nuw i32 [[TMP0]], 41193// CHECK11-NEXT: [[TMP5:%.*]] = sext i32 [[TMP4]] to i641194// CHECK11-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[DOTOFFLOAD_SIZES]], ptr align 4 @.offload_sizes, i32 24, i1 false)1195// CHECK11-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01196// CHECK11-NEXT: store i32 [[TMP3]], ptr [[TMP6]], align 41197// CHECK11-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01198// CHECK11-NEXT: store i32 [[TMP3]], ptr [[TMP7]], align 41199// CHECK11-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01200// CHECK11-NEXT: store ptr null, ptr [[TMP8]], align 41201// CHECK11-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11202// CHECK11-NEXT: store i32 [[TMP0]], ptr [[TMP9]], align 41203// CHECK11-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11204// CHECK11-NEXT: store i32 [[TMP0]], ptr [[TMP10]], align 41205// CHECK11-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 11206// CHECK11-NEXT: store ptr null, ptr [[TMP11]], align 41207// CHECK11-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21208// CHECK11-NEXT: store ptr [[VLA]], ptr [[TMP12]], align 41209// CHECK11-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21210// CHECK11-NEXT: store ptr [[VLA]], ptr [[TMP13]], align 41211// CHECK11-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 21212// CHECK11-NEXT: store i64 [[TMP5]], ptr [[TMP14]], align 41213// CHECK11-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 21214// CHECK11-NEXT: store ptr null, ptr [[TMP15]], align 41215// CHECK11-NEXT: [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01216// CHECK11-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01217// CHECK11-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 01218// CHECK11-NEXT: [[TMP19:%.*]] = load i32, ptr [[N]], align 41219// CHECK11-NEXT: store i32 [[TMP19]], ptr [[DOTCAPTURE_EXPR_]], align 41220// CHECK11-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41221// CHECK11-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP20]], 01222// CHECK11-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11223// CHECK11-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11224// CHECK11-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41225// CHECK11-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41226// CHECK11-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP21]], 11227// CHECK11-NEXT: [[TMP22:%.*]] = zext i32 [[ADD]] to i641228// CHECK11-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01229// CHECK11-NEXT: store i32 3, ptr [[TMP23]], align 41230// CHECK11-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11231// CHECK11-NEXT: store i32 3, ptr [[TMP24]], align 41232// CHECK11-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21233// CHECK11-NEXT: store ptr [[TMP16]], ptr [[TMP25]], align 41234// CHECK11-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31235// CHECK11-NEXT: store ptr [[TMP17]], ptr [[TMP26]], align 41236// CHECK11-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41237// CHECK11-NEXT: store ptr [[TMP18]], ptr [[TMP27]], align 41238// CHECK11-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51239// CHECK11-NEXT: store ptr @.offload_maptypes, ptr [[TMP28]], align 41240// CHECK11-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61241// CHECK11-NEXT: store ptr null, ptr [[TMP29]], align 41242// CHECK11-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71243// CHECK11-NEXT: store ptr null, ptr [[TMP30]], align 41244// CHECK11-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81245// CHECK11-NEXT: store i64 [[TMP22]], ptr [[TMP31]], align 81246// CHECK11-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91247// CHECK11-NEXT: store i64 0, ptr [[TMP32]], align 81248// CHECK11-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101249// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP33]], align 41250// CHECK11-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111251// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP34]], align 41252// CHECK11-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121253// CHECK11-NEXT: store i32 0, ptr [[TMP35]], align 41254// 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_l73.region_id, ptr [[KERNEL_ARGS]])1255// CHECK11-NEXT: [[TMP37:%.*]] = icmp ne i32 [[TMP36]], 01256// CHECK11-NEXT: br i1 [[TMP37]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1257// CHECK11: omp_offload.failed:1258// CHECK11-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z15teams_local_argv_l73(i32 [[TMP3]], i32 [[TMP0]], ptr [[VLA]]) #[[ATTR3:[0-9]+]]1259// CHECK11-NEXT: br label [[OMP_OFFLOAD_CONT]]1260// CHECK11: omp_offload.cont:1261// CHECK11-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[VLA]], i32 01262// CHECK11-NEXT: [[TMP38:%.*]] = load i32, ptr [[ARRAYIDX]], align 41263// CHECK11-NEXT: [[TMP39:%.*]] = load ptr, ptr [[SAVED_STACK]], align 41264// CHECK11-NEXT: call void @llvm.stackrestore.p0(ptr [[TMP39]])1265// CHECK11-NEXT: ret i32 [[TMP38]]1266//1267//1268// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z15teams_local_argv_l731269// CHECK11-SAME: (i32 noundef [[N:%.*]], i32 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2:[0-9]+]] {1270// CHECK11-NEXT: entry:1271// CHECK11-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41272// CHECK11-NEXT: [[VLA_ADDR:%.*]] = alloca i32, align 41273// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41274// CHECK11-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41275// CHECK11-NEXT: store i32 [[VLA]], ptr [[VLA_ADDR]], align 41276// CHECK11-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41277// CHECK11-NEXT: [[TMP0:%.*]] = load i32, ptr [[VLA_ADDR]], align 41278// CHECK11-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 41279// 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_l73.omp_outlined, ptr [[N_ADDR]], i32 [[TMP0]], ptr [[TMP1]])1280// CHECK11-NEXT: ret void1281//1282//1283// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z15teams_local_argv_l73.omp_outlined1284// 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]] {1285// CHECK11-NEXT: entry:1286// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41287// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41288// CHECK11-NEXT: [[N_ADDR:%.*]] = alloca ptr, align 41289// CHECK11-NEXT: [[VLA_ADDR:%.*]] = alloca i32, align 41290// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41291// CHECK11-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41292// CHECK11-NEXT: [[TMP:%.*]] = alloca i32, align 41293// CHECK11-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41294// CHECK11-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41295// CHECK11-NEXT: [[I:%.*]] = alloca i32, align 41296// CHECK11-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41297// CHECK11-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41298// CHECK11-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41299// CHECK11-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41300// CHECK11-NEXT: [[I3:%.*]] = alloca i32, align 41301// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41302// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41303// CHECK11-NEXT: store ptr [[N]], ptr [[N_ADDR]], align 41304// CHECK11-NEXT: store i32 [[VLA]], ptr [[VLA_ADDR]], align 41305// CHECK11-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41306// CHECK11-NEXT: [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 41307// CHECK11-NEXT: [[TMP1:%.*]] = load i32, ptr [[VLA_ADDR]], align 41308// CHECK11-NEXT: [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 41309// CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 41310// CHECK11-NEXT: store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR_]], align 41311// CHECK11-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41312// CHECK11-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP4]], 01313// CHECK11-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11314// CHECK11-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11315// CHECK11-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41316// CHECK11-NEXT: store i32 0, ptr [[I]], align 41317// CHECK11-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41318// CHECK11-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP5]]1319// CHECK11-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1320// CHECK11: omp.precond.then:1321// CHECK11-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41322// CHECK11-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41323// CHECK11-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_UB]], align 41324// CHECK11-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41325// CHECK11-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41326// CHECK11-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41327// CHECK11-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41328// CHECK11-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP8]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1329// CHECK11-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41330// CHECK11-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41331// CHECK11-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]1332// CHECK11-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1333// CHECK11: cond.true:1334// CHECK11-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41335// CHECK11-NEXT: br label [[COND_END:%.*]]1336// CHECK11: cond.false:1337// CHECK11-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41338// CHECK11-NEXT: br label [[COND_END]]1339// CHECK11: cond.end:1340// CHECK11-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]1341// CHECK11-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41342// CHECK11-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41343// CHECK11-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 41344// CHECK11-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1345// CHECK11: omp.inner.for.cond:1346// CHECK11-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41347// CHECK11-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41348// CHECK11-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]1349// CHECK11-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1350// CHECK11: omp.inner.for.body:1351// CHECK11-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41352// CHECK11-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 11353// CHECK11-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1354// CHECK11-NEXT: store i32 [[ADD]], ptr [[I3]], align 41355// CHECK11-NEXT: [[TMP17:%.*]] = load i32, ptr [[I3]], align 41356// CHECK11-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i32 [[TMP17]]1357// CHECK11-NEXT: store i32 0, ptr [[ARRAYIDX]], align 41358// CHECK11-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1359// CHECK11: omp.body.continue:1360// CHECK11-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1361// CHECK11: omp.inner.for.inc:1362// CHECK11-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41363// CHECK11-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP18]], 11364// CHECK11-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 41365// CHECK11-NEXT: br label [[OMP_INNER_FOR_COND]]1366// CHECK11: omp.inner.for.end:1367// CHECK11-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1368// CHECK11: omp.loop.exit:1369// CHECK11-NEXT: [[TMP19:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41370// CHECK11-NEXT: [[TMP20:%.*]] = load i32, ptr [[TMP19]], align 41371// CHECK11-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP20]])1372// CHECK11-NEXT: br label [[OMP_PRECOND_END]]1373// CHECK11: omp.precond.end:1374// CHECK11-NEXT: ret void1375//1376//1377// CHECK17-LABEL: define {{[^@]+}}@_Z21teams_template_structv1378// CHECK17-SAME: () #[[ATTR0:[0-9]+]] {1379// CHECK17-NEXT: entry:1380// CHECK17-NEXT: [[V:%.*]] = alloca [[STRUCT_SS:%.*]], align 41381// CHECK17-NEXT: [[CALL:%.*]] = call noundef signext i32 @_ZN2SSIiLi123ELx456EE3fooEv(ptr noundef nonnull align 4 dereferenceable(496) [[V]])1382// CHECK17-NEXT: ret i32 [[CALL]]1383//1384//1385// CHECK17-LABEL: define {{[^@]+}}@_ZN2SSIiLi123ELx456EE3fooEv1386// CHECK17-SAME: (ptr noundef nonnull align 4 dereferenceable(496) [[THIS:%.*]]) #[[ATTR0]] comdat {1387// CHECK17-NEXT: entry:1388// CHECK17-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81389// CHECK17-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 81390// CHECK17-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 81391// CHECK17-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 81392// CHECK17-NEXT: [[TMP:%.*]] = alloca i32, align 41393// CHECK17-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81394// CHECK17-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81395// CHECK17-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81396// CHECK17-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[THIS1]], i32 0, i32 01397// CHECK17-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01398// CHECK17-NEXT: store ptr [[THIS1]], ptr [[TMP0]], align 81399// CHECK17-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01400// CHECK17-NEXT: store ptr [[A]], ptr [[TMP1]], align 81401// CHECK17-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 01402// CHECK17-NEXT: store ptr null, ptr [[TMP2]], align 81403// CHECK17-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01404// CHECK17-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01405// CHECK17-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01406// CHECK17-NEXT: store i32 3, ptr [[TMP5]], align 41407// CHECK17-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11408// CHECK17-NEXT: store i32 1, ptr [[TMP6]], align 41409// CHECK17-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21410// CHECK17-NEXT: store ptr [[TMP3]], ptr [[TMP7]], align 81411// CHECK17-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31412// CHECK17-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 81413// CHECK17-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41414// CHECK17-NEXT: store ptr @.offload_sizes, ptr [[TMP9]], align 81415// CHECK17-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51416// CHECK17-NEXT: store ptr @.offload_maptypes, ptr [[TMP10]], align 81417// CHECK17-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61418// CHECK17-NEXT: store ptr null, ptr [[TMP11]], align 81419// CHECK17-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71420// CHECK17-NEXT: store ptr null, ptr [[TMP12]], align 81421// CHECK17-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81422// CHECK17-NEXT: store i64 123, ptr [[TMP13]], align 81423// CHECK17-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91424// CHECK17-NEXT: store i64 0, ptr [[TMP14]], align 81425// CHECK17-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101426// CHECK17-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 41427// CHECK17-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111428// CHECK17-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 41429// CHECK17-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121430// CHECK17-NEXT: store i32 0, ptr [[TMP17]], align 41431// 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_l109.region_id, ptr [[KERNEL_ARGS]])1432// CHECK17-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 01433// CHECK17-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1434// CHECK17: omp_offload.failed:1435// CHECK17-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l109(ptr [[THIS1]]) #[[ATTR2:[0-9]+]]1436// CHECK17-NEXT: br label [[OMP_OFFLOAD_CONT]]1437// CHECK17: omp_offload.cont:1438// CHECK17-NEXT: [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 01439// CHECK17-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A2]], i64 0, i64 01440// CHECK17-NEXT: [[TMP20:%.*]] = load i32, ptr [[ARRAYIDX]], align 41441// CHECK17-NEXT: ret i32 [[TMP20]]1442//1443//1444// CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l1091445// CHECK17-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1:[0-9]+]] {1446// CHECK17-NEXT: entry:1447// CHECK17-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81448// CHECK17-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81449// CHECK17-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81450// CHECK17-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l109.omp_outlined, ptr [[TMP0]])1451// CHECK17-NEXT: ret void1452//1453//1454// CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l109.omp_outlined1455// CHECK17-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {1456// CHECK17-NEXT: entry:1457// CHECK17-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81458// CHECK17-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81459// CHECK17-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81460// CHECK17-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41461// CHECK17-NEXT: [[TMP:%.*]] = alloca i32, align 41462// CHECK17-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41463// CHECK17-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41464// CHECK17-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41465// CHECK17-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41466// CHECK17-NEXT: [[I:%.*]] = alloca i32, align 41467// CHECK17-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81468// CHECK17-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81469// CHECK17-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81470// CHECK17-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81471// CHECK17-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41472// CHECK17-NEXT: store i32 122, ptr [[DOTOMP_UB]], align 41473// CHECK17-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41474// CHECK17-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41475// CHECK17-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81476// CHECK17-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41477// CHECK17-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1478// CHECK17-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41479// CHECK17-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1221480// CHECK17-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1481// CHECK17: cond.true:1482// CHECK17-NEXT: br label [[COND_END:%.*]]1483// CHECK17: cond.false:1484// CHECK17-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41485// CHECK17-NEXT: br label [[COND_END]]1486// CHECK17: cond.end:1487// CHECK17-NEXT: [[COND:%.*]] = phi i32 [ 122, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]1488// CHECK17-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41489// CHECK17-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41490// CHECK17-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 41491// CHECK17-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1492// CHECK17: omp.inner.for.cond:1493// CHECK17-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41494// CHECK17-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41495// CHECK17-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]1496// CHECK17-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1497// CHECK17: omp.inner.for.body:1498// CHECK17-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41499// CHECK17-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 11500// CHECK17-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1501// CHECK17-NEXT: store i32 [[ADD]], ptr [[I]], align 41502// CHECK17-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP0]], i32 0, i32 01503// CHECK17-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 41504// CHECK17-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP9]] to i641505// CHECK17-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A]], i64 0, i64 [[IDXPROM]]1506// CHECK17-NEXT: store i32 0, ptr [[ARRAYIDX]], align 41507// CHECK17-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1508// CHECK17: omp.body.continue:1509// CHECK17-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1510// CHECK17: omp.inner.for.inc:1511// CHECK17-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41512// CHECK17-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP10]], 11513// CHECK17-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 41514// CHECK17-NEXT: br label [[OMP_INNER_FOR_COND]]1515// CHECK17: omp.inner.for.end:1516// CHECK17-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1517// CHECK17: omp.loop.exit:1518// CHECK17-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])1519// CHECK17-NEXT: ret void1520//1521//1522// CHECK19-LABEL: define {{[^@]+}}@_Z21teams_template_structv1523// CHECK19-SAME: () #[[ATTR0:[0-9]+]] {1524// CHECK19-NEXT: entry:1525// CHECK19-NEXT: [[V:%.*]] = alloca [[STRUCT_SS:%.*]], align 41526// CHECK19-NEXT: [[CALL:%.*]] = call noundef i32 @_ZN2SSIiLi123ELx456EE3fooEv(ptr noundef nonnull align 4 dereferenceable(496) [[V]])1527// CHECK19-NEXT: ret i32 [[CALL]]1528//1529//1530// CHECK19-LABEL: define {{[^@]+}}@_ZN2SSIiLi123ELx456EE3fooEv1531// CHECK19-SAME: (ptr noundef nonnull align 4 dereferenceable(496) [[THIS:%.*]]) #[[ATTR0]] comdat align 2 {1532// CHECK19-NEXT: entry:1533// CHECK19-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41534// CHECK19-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 41535// CHECK19-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 41536// CHECK19-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 41537// CHECK19-NEXT: [[TMP:%.*]] = alloca i32, align 41538// CHECK19-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81539// CHECK19-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41540// CHECK19-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41541// CHECK19-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[THIS1]], i32 0, i32 01542// CHECK19-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01543// CHECK19-NEXT: store ptr [[THIS1]], ptr [[TMP0]], align 41544// CHECK19-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01545// CHECK19-NEXT: store ptr [[A]], ptr [[TMP1]], align 41546// CHECK19-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01547// CHECK19-NEXT: store ptr null, ptr [[TMP2]], align 41548// CHECK19-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01549// CHECK19-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01550// CHECK19-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01551// CHECK19-NEXT: store i32 3, ptr [[TMP5]], align 41552// CHECK19-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11553// CHECK19-NEXT: store i32 1, ptr [[TMP6]], align 41554// CHECK19-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21555// CHECK19-NEXT: store ptr [[TMP3]], ptr [[TMP7]], align 41556// CHECK19-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31557// CHECK19-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 41558// CHECK19-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41559// CHECK19-NEXT: store ptr @.offload_sizes, ptr [[TMP9]], align 41560// CHECK19-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51561// CHECK19-NEXT: store ptr @.offload_maptypes, ptr [[TMP10]], align 41562// CHECK19-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61563// CHECK19-NEXT: store ptr null, ptr [[TMP11]], align 41564// CHECK19-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71565// CHECK19-NEXT: store ptr null, ptr [[TMP12]], align 41566// CHECK19-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81567// CHECK19-NEXT: store i64 123, ptr [[TMP13]], align 81568// CHECK19-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91569// CHECK19-NEXT: store i64 0, ptr [[TMP14]], align 81570// CHECK19-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101571// CHECK19-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 41572// CHECK19-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111573// CHECK19-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 41574// CHECK19-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121575// CHECK19-NEXT: store i32 0, ptr [[TMP17]], align 41576// 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_l109.region_id, ptr [[KERNEL_ARGS]])1577// CHECK19-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 01578// CHECK19-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1579// CHECK19: omp_offload.failed:1580// CHECK19-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l109(ptr [[THIS1]]) #[[ATTR2:[0-9]+]]1581// CHECK19-NEXT: br label [[OMP_OFFLOAD_CONT]]1582// CHECK19: omp_offload.cont:1583// CHECK19-NEXT: [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 01584// CHECK19-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A2]], i32 0, i32 01585// CHECK19-NEXT: [[TMP20:%.*]] = load i32, ptr [[ARRAYIDX]], align 41586// CHECK19-NEXT: ret i32 [[TMP20]]1587//1588//1589// CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l1091590// CHECK19-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1:[0-9]+]] {1591// CHECK19-NEXT: entry:1592// CHECK19-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41593// CHECK19-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41594// CHECK19-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41595// CHECK19-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l109.omp_outlined, ptr [[TMP0]])1596// CHECK19-NEXT: ret void1597//1598//1599// CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l109.omp_outlined1600// CHECK19-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {1601// CHECK19-NEXT: entry:1602// CHECK19-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41603// CHECK19-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41604// CHECK19-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41605// CHECK19-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41606// CHECK19-NEXT: [[TMP:%.*]] = alloca i32, align 41607// CHECK19-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41608// CHECK19-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41609// CHECK19-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41610// CHECK19-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41611// CHECK19-NEXT: [[I:%.*]] = alloca i32, align 41612// CHECK19-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41613// CHECK19-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41614// CHECK19-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41615// CHECK19-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41616// CHECK19-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41617// CHECK19-NEXT: store i32 122, ptr [[DOTOMP_UB]], align 41618// CHECK19-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41619// CHECK19-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41620// CHECK19-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41621// CHECK19-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41622// CHECK19-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1623// CHECK19-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41624// CHECK19-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1221625// CHECK19-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1626// CHECK19: cond.true:1627// CHECK19-NEXT: br label [[COND_END:%.*]]1628// CHECK19: cond.false:1629// CHECK19-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41630// CHECK19-NEXT: br label [[COND_END]]1631// CHECK19: cond.end:1632// CHECK19-NEXT: [[COND:%.*]] = phi i32 [ 122, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]1633// CHECK19-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41634// CHECK19-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41635// CHECK19-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 41636// CHECK19-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1637// CHECK19: omp.inner.for.cond:1638// CHECK19-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41639// CHECK19-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41640// CHECK19-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]1641// CHECK19-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1642// CHECK19: omp.inner.for.body:1643// CHECK19-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41644// CHECK19-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 11645// CHECK19-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1646// CHECK19-NEXT: store i32 [[ADD]], ptr [[I]], align 41647// CHECK19-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP0]], i32 0, i32 01648// CHECK19-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 41649// CHECK19-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [123 x i32], ptr [[A]], i32 0, i32 [[TMP9]]1650// CHECK19-NEXT: store i32 0, ptr [[ARRAYIDX]], align 41651// CHECK19-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1652// CHECK19: omp.body.continue:1653// CHECK19-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1654// CHECK19: omp.inner.for.inc:1655// CHECK19-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41656// CHECK19-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP10]], 11657// CHECK19-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 41658// CHECK19-NEXT: br label [[OMP_INNER_FOR_COND]]1659// CHECK19: omp.inner.for.end:1660// CHECK19-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1661// CHECK19: omp.loop.exit:1662// CHECK19-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])1663// CHECK19-NEXT: ret void1664//1665//1666// CHECK25-LABEL: define {{[^@]+}}@main1667// CHECK25-SAME: (i32 noundef signext [[ARGC:%.*]], ptr noundef [[ARGV:%.*]]) #[[ATTR0:[0-9]+]] {1668// CHECK25-NEXT: entry:1669// CHECK25-NEXT: [[RETVAL:%.*]] = alloca i32, align 41670// CHECK25-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 41671// CHECK25-NEXT: [[ARGV_ADDR:%.*]] = alloca ptr, align 81672// CHECK25-NEXT: [[N:%.*]] = alloca i32, align 41673// CHECK25-NEXT: [[SAVED_STACK:%.*]] = alloca ptr, align 81674// CHECK25-NEXT: [[__VLA_EXPR0:%.*]] = alloca i64, align 81675// CHECK25-NEXT: [[N_CASTED:%.*]] = alloca i64, align 81676// CHECK25-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 81677// CHECK25-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 81678// CHECK25-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 81679// CHECK25-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 81680// CHECK25-NEXT: [[TMP:%.*]] = alloca i32, align 41681// CHECK25-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41682// CHECK25-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41683// CHECK25-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81684// CHECK25-NEXT: store i32 0, ptr [[RETVAL]], align 41685// CHECK25-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 41686// CHECK25-NEXT: store ptr [[ARGV]], ptr [[ARGV_ADDR]], align 81687// CHECK25-NEXT: store i32 100, ptr [[N]], align 41688// CHECK25-NEXT: [[TMP0:%.*]] = load i32, ptr [[N]], align 41689// CHECK25-NEXT: [[TMP1:%.*]] = zext i32 [[TMP0]] to i641690// CHECK25-NEXT: [[TMP2:%.*]] = call ptr @llvm.stacksave.p0()1691// CHECK25-NEXT: store ptr [[TMP2]], ptr [[SAVED_STACK]], align 81692// CHECK25-NEXT: [[VLA:%.*]] = alloca i32, i64 [[TMP1]], align 41693// CHECK25-NEXT: store i64 [[TMP1]], ptr [[__VLA_EXPR0]], align 81694// CHECK25-NEXT: [[TMP3:%.*]] = load i32, ptr [[N]], align 41695// CHECK25-NEXT: store i32 [[TMP3]], ptr [[N_CASTED]], align 41696// CHECK25-NEXT: [[TMP4:%.*]] = load i64, ptr [[N_CASTED]], align 81697// CHECK25-NEXT: [[TMP5:%.*]] = mul nuw i64 [[TMP1]], 41698// CHECK25-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES]], ptr align 8 @.offload_sizes, i64 24, i1 false)1699// CHECK25-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01700// CHECK25-NEXT: store i64 [[TMP4]], ptr [[TMP6]], align 81701// CHECK25-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01702// CHECK25-NEXT: store i64 [[TMP4]], ptr [[TMP7]], align 81703// CHECK25-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 01704// CHECK25-NEXT: store ptr null, ptr [[TMP8]], align 81705// CHECK25-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11706// CHECK25-NEXT: store i64 [[TMP1]], ptr [[TMP9]], align 81707// CHECK25-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11708// CHECK25-NEXT: store i64 [[TMP1]], ptr [[TMP10]], align 81709// CHECK25-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 11710// CHECK25-NEXT: store ptr null, ptr [[TMP11]], align 81711// CHECK25-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21712// CHECK25-NEXT: store ptr [[VLA]], ptr [[TMP12]], align 81713// CHECK25-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21714// CHECK25-NEXT: store ptr [[VLA]], ptr [[TMP13]], align 81715// CHECK25-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 21716// CHECK25-NEXT: store i64 [[TMP5]], ptr [[TMP14]], align 81717// CHECK25-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 21718// CHECK25-NEXT: store ptr null, ptr [[TMP15]], align 81719// CHECK25-NEXT: [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01720// CHECK25-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01721// CHECK25-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 01722// CHECK25-NEXT: [[TMP19:%.*]] = load i32, ptr [[N]], align 41723// CHECK25-NEXT: store i32 [[TMP19]], ptr [[DOTCAPTURE_EXPR_]], align 41724// CHECK25-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41725// CHECK25-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP20]], 01726// CHECK25-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11727// CHECK25-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11728// CHECK25-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41729// CHECK25-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41730// CHECK25-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP21]], 11731// CHECK25-NEXT: [[TMP22:%.*]] = zext i32 [[ADD]] to i641732// CHECK25-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01733// CHECK25-NEXT: store i32 3, ptr [[TMP23]], align 41734// CHECK25-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11735// CHECK25-NEXT: store i32 3, ptr [[TMP24]], align 41736// CHECK25-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21737// CHECK25-NEXT: store ptr [[TMP16]], ptr [[TMP25]], align 81738// CHECK25-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31739// CHECK25-NEXT: store ptr [[TMP17]], ptr [[TMP26]], align 81740// CHECK25-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41741// CHECK25-NEXT: store ptr [[TMP18]], ptr [[TMP27]], align 81742// CHECK25-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51743// CHECK25-NEXT: store ptr @.offload_maptypes, ptr [[TMP28]], align 81744// CHECK25-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61745// CHECK25-NEXT: store ptr null, ptr [[TMP29]], align 81746// CHECK25-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71747// CHECK25-NEXT: store ptr null, ptr [[TMP30]], align 81748// CHECK25-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81749// CHECK25-NEXT: store i64 [[TMP22]], ptr [[TMP31]], align 81750// CHECK25-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91751// CHECK25-NEXT: store i64 0, ptr [[TMP32]], align 81752// CHECK25-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101753// CHECK25-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP33]], align 41754// CHECK25-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111755// CHECK25-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP34]], align 41756// CHECK25-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121757// CHECK25-NEXT: store i32 0, ptr [[TMP35]], align 41758// 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_l162.region_id, ptr [[KERNEL_ARGS]])1759// CHECK25-NEXT: [[TMP37:%.*]] = icmp ne i32 [[TMP36]], 01760// CHECK25-NEXT: br i1 [[TMP37]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1761// CHECK25: omp_offload.failed:1762// CHECK25-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l162(i64 [[TMP4]], i64 [[TMP1]], ptr [[VLA]]) #[[ATTR3:[0-9]+]]1763// CHECK25-NEXT: br label [[OMP_OFFLOAD_CONT]]1764// CHECK25: omp_offload.cont:1765// CHECK25-NEXT: [[TMP38:%.*]] = load i32, ptr [[ARGC_ADDR]], align 41766// CHECK25-NEXT: [[CALL:%.*]] = call noundef signext i32 @_Z5tmainIiLi10EEiT_(i32 noundef signext [[TMP38]])1767// CHECK25-NEXT: store i32 [[CALL]], ptr [[RETVAL]], align 41768// CHECK25-NEXT: [[TMP39:%.*]] = load ptr, ptr [[SAVED_STACK]], align 81769// CHECK25-NEXT: call void @llvm.stackrestore.p0(ptr [[TMP39]])1770// CHECK25-NEXT: [[TMP40:%.*]] = load i32, ptr [[RETVAL]], align 41771// CHECK25-NEXT: ret i32 [[TMP40]]1772//1773//1774// CHECK25-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l1621775// CHECK25-SAME: (i64 noundef [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2:[0-9]+]] {1776// CHECK25-NEXT: entry:1777// CHECK25-NEXT: [[N_ADDR:%.*]] = alloca i64, align 81778// CHECK25-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 81779// CHECK25-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81780// CHECK25-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 81781// CHECK25-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR]], align 81782// CHECK25-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81783// CHECK25-NEXT: [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR]], align 81784// CHECK25-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 81785// CHECK25-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l162.omp_outlined, ptr [[N_ADDR]], i64 [[TMP0]], ptr [[TMP1]])1786// CHECK25-NEXT: ret void1787//1788//1789// CHECK25-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l162.omp_outlined1790// 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]] {1791// CHECK25-NEXT: entry:1792// CHECK25-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81793// CHECK25-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81794// CHECK25-NEXT: [[N_ADDR:%.*]] = alloca ptr, align 81795// CHECK25-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 81796// CHECK25-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81797// CHECK25-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41798// CHECK25-NEXT: [[TMP:%.*]] = alloca i32, align 41799// CHECK25-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41800// CHECK25-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41801// CHECK25-NEXT: [[I:%.*]] = alloca i32, align 41802// CHECK25-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41803// CHECK25-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41804// CHECK25-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41805// CHECK25-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41806// CHECK25-NEXT: [[I3:%.*]] = alloca i32, align 41807// CHECK25-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81808// CHECK25-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81809// CHECK25-NEXT: store ptr [[N]], ptr [[N_ADDR]], align 81810// CHECK25-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR]], align 81811// CHECK25-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81812// CHECK25-NEXT: [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 81813// CHECK25-NEXT: [[TMP1:%.*]] = load i64, ptr [[VLA_ADDR]], align 81814// CHECK25-NEXT: [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 81815// CHECK25-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 41816// CHECK25-NEXT: store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR_]], align 41817// CHECK25-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41818// CHECK25-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP4]], 01819// CHECK25-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11820// CHECK25-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11821// CHECK25-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41822// CHECK25-NEXT: store i32 0, ptr [[I]], align 41823// CHECK25-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41824// CHECK25-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP5]]1825// CHECK25-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1826// CHECK25: omp.precond.then:1827// CHECK25-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41828// CHECK25-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41829// CHECK25-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_UB]], align 41830// CHECK25-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41831// CHECK25-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41832// CHECK25-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81833// CHECK25-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41834// CHECK25-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP8]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1835// CHECK25-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41836// CHECK25-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41837// CHECK25-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]1838// CHECK25-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1839// CHECK25: cond.true:1840// CHECK25-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41841// CHECK25-NEXT: br label [[COND_END:%.*]]1842// CHECK25: cond.false:1843// CHECK25-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41844// CHECK25-NEXT: br label [[COND_END]]1845// CHECK25: cond.end:1846// CHECK25-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]1847// CHECK25-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41848// CHECK25-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41849// CHECK25-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 41850// CHECK25-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1851// CHECK25: omp.inner.for.cond:1852// CHECK25-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41853// CHECK25-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41854// CHECK25-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]1855// CHECK25-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1856// CHECK25: omp.inner.for.body:1857// CHECK25-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41858// CHECK25-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 11859// CHECK25-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1860// CHECK25-NEXT: store i32 [[ADD]], ptr [[I3]], align 41861// CHECK25-NEXT: [[TMP17:%.*]] = load i32, ptr [[I3]], align 41862// CHECK25-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP17]] to i641863// CHECK25-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i64 [[IDXPROM]]1864// CHECK25-NEXT: store i32 0, ptr [[ARRAYIDX]], align 41865// CHECK25-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1866// CHECK25: omp.body.continue:1867// CHECK25-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1868// CHECK25: omp.inner.for.inc:1869// CHECK25-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 41870// CHECK25-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP18]], 11871// CHECK25-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 41872// CHECK25-NEXT: br label [[OMP_INNER_FOR_COND]]1873// CHECK25: omp.inner.for.end:1874// CHECK25-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1875// CHECK25: omp.loop.exit:1876// CHECK25-NEXT: [[TMP19:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81877// CHECK25-NEXT: [[TMP20:%.*]] = load i32, ptr [[TMP19]], align 41878// CHECK25-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP20]])1879// CHECK25-NEXT: br label [[OMP_PRECOND_END]]1880// CHECK25: omp.precond.end:1881// CHECK25-NEXT: ret void1882//1883//1884// CHECK25-LABEL: define {{[^@]+}}@_Z5tmainIiLi10EEiT_1885// CHECK25-SAME: (i32 noundef signext [[ARGC:%.*]]) #[[ATTR5:[0-9]+]] comdat {1886// CHECK25-NEXT: entry:1887// CHECK25-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 41888// CHECK25-NEXT: [[A:%.*]] = alloca [10 x i32], align 41889// CHECK25-NEXT: [[TE:%.*]] = alloca i32, align 41890// CHECK25-NEXT: [[TH:%.*]] = alloca i32, align 41891// CHECK25-NEXT: [[TE_CASTED:%.*]] = alloca i64, align 81892// CHECK25-NEXT: [[TH_CASTED:%.*]] = alloca i64, align 81893// CHECK25-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 81894// CHECK25-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 81895// CHECK25-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 81896// CHECK25-NEXT: [[TMP:%.*]] = alloca i32, align 41897// CHECK25-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81898// CHECK25-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 41899// CHECK25-NEXT: store i32 0, ptr [[TE]], align 41900// CHECK25-NEXT: store i32 128, ptr [[TH]], align 41901// CHECK25-NEXT: [[TMP0:%.*]] = load i32, ptr [[TE]], align 41902// CHECK25-NEXT: store i32 [[TMP0]], ptr [[TE_CASTED]], align 41903// CHECK25-NEXT: [[TMP1:%.*]] = load i64, ptr [[TE_CASTED]], align 81904// CHECK25-NEXT: [[TMP2:%.*]] = load i32, ptr [[TH]], align 41905// CHECK25-NEXT: store i32 [[TMP2]], ptr [[TH_CASTED]], align 41906// CHECK25-NEXT: [[TMP3:%.*]] = load i64, ptr [[TH_CASTED]], align 81907// CHECK25-NEXT: [[TMP4:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01908// CHECK25-NEXT: store i64 [[TMP1]], ptr [[TMP4]], align 81909// CHECK25-NEXT: [[TMP5:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01910// CHECK25-NEXT: store i64 [[TMP1]], ptr [[TMP5]], align 81911// CHECK25-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 01912// CHECK25-NEXT: store ptr null, ptr [[TMP6]], align 81913// CHECK25-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11914// CHECK25-NEXT: store i64 [[TMP3]], ptr [[TMP7]], align 81915// CHECK25-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11916// CHECK25-NEXT: store i64 [[TMP3]], ptr [[TMP8]], align 81917// CHECK25-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 11918// CHECK25-NEXT: store ptr null, ptr [[TMP9]], align 81919// CHECK25-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21920// CHECK25-NEXT: store ptr [[A]], ptr [[TMP10]], align 81921// CHECK25-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21922// CHECK25-NEXT: store ptr [[A]], ptr [[TMP11]], align 81923// CHECK25-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 21924// CHECK25-NEXT: store ptr null, ptr [[TMP12]], align 81925// CHECK25-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01926// CHECK25-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01927// CHECK25-NEXT: [[TMP15:%.*]] = load i32, ptr [[TE]], align 41928// CHECK25-NEXT: [[TMP16:%.*]] = load i32, ptr [[TH]], align 41929// CHECK25-NEXT: [[TMP17:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP15]], 01930// CHECK25-NEXT: [[TMP18:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP16]], 01931// CHECK25-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01932// CHECK25-NEXT: store i32 3, ptr [[TMP19]], align 41933// CHECK25-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11934// CHECK25-NEXT: store i32 3, ptr [[TMP20]], align 41935// CHECK25-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21936// CHECK25-NEXT: store ptr [[TMP13]], ptr [[TMP21]], align 81937// CHECK25-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31938// CHECK25-NEXT: store ptr [[TMP14]], ptr [[TMP22]], align 81939// CHECK25-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41940// CHECK25-NEXT: store ptr @.offload_sizes.1, ptr [[TMP23]], align 81941// CHECK25-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51942// CHECK25-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP24]], align 81943// CHECK25-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61944// CHECK25-NEXT: store ptr null, ptr [[TMP25]], align 81945// CHECK25-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71946// CHECK25-NEXT: store ptr null, ptr [[TMP26]], align 81947// CHECK25-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81948// CHECK25-NEXT: store i64 10, ptr [[TMP27]], align 81949// CHECK25-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91950// CHECK25-NEXT: store i64 0, ptr [[TMP28]], align 81951// CHECK25-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101952// CHECK25-NEXT: store [3 x i32] [[TMP17]], ptr [[TMP29]], align 41953// CHECK25-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111954// CHECK25-NEXT: store [3 x i32] [[TMP18]], ptr [[TMP30]], align 41955// CHECK25-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121956// CHECK25-NEXT: store i32 0, ptr [[TMP31]], align 41957// 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__l151.region_id, ptr [[KERNEL_ARGS]])1958// CHECK25-NEXT: [[TMP33:%.*]] = icmp ne i32 [[TMP32]], 01959// CHECK25-NEXT: br i1 [[TMP33]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1960// CHECK25: omp_offload.failed:1961// CHECK25-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l151(i64 [[TMP1]], i64 [[TMP3]], ptr [[A]]) #[[ATTR3]]1962// CHECK25-NEXT: br label [[OMP_OFFLOAD_CONT]]1963// CHECK25: omp_offload.cont:1964// CHECK25-NEXT: ret i32 01965//1966//1967// CHECK25-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l1511968// CHECK25-SAME: (i64 noundef [[TE:%.*]], i64 noundef [[TH:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {1969// CHECK25-NEXT: entry:1970// CHECK25-NEXT: [[TE_ADDR:%.*]] = alloca i64, align 81971// CHECK25-NEXT: [[TH_ADDR:%.*]] = alloca i64, align 81972// CHECK25-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81973// CHECK25-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB2]])1974// CHECK25-NEXT: store i64 [[TE]], ptr [[TE_ADDR]], align 81975// CHECK25-NEXT: store i64 [[TH]], ptr [[TH_ADDR]], align 81976// CHECK25-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 81977// CHECK25-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 81978// CHECK25-NEXT: [[TMP2:%.*]] = load i32, ptr [[TE_ADDR]], align 41979// CHECK25-NEXT: [[TMP3:%.*]] = load i32, ptr [[TH_ADDR]], align 41980// CHECK25-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB2]], i32 [[TMP0]], i32 [[TMP2]], i32 [[TMP3]])1981// CHECK25-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l151.omp_outlined, ptr [[TMP1]])1982// CHECK25-NEXT: ret void1983//1984//1985// CHECK25-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l151.omp_outlined1986// CHECK25-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {1987// CHECK25-NEXT: entry:1988// CHECK25-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81989// CHECK25-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81990// CHECK25-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 81991// CHECK25-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41992// CHECK25-NEXT: [[TMP:%.*]] = alloca i32, align 41993// CHECK25-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41994// CHECK25-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41995// CHECK25-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41996// CHECK25-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41997// CHECK25-NEXT: [[I:%.*]] = alloca i32, align 41998// CHECK25-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81999// CHECK25-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82000// CHECK25-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 82001// CHECK25-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 82002// CHECK25-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 42003// CHECK25-NEXT: store i32 9, ptr [[DOTOMP_UB]], align 42004// CHECK25-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42005// CHECK25-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42006// CHECK25-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82007// CHECK25-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 42008// CHECK25-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2009// CHECK25-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42010// CHECK25-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 92011// CHECK25-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2012// CHECK25: cond.true:2013// CHECK25-NEXT: br label [[COND_END:%.*]]2014// CHECK25: cond.false:2015// CHECK25-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42016// CHECK25-NEXT: br label [[COND_END]]2017// CHECK25: cond.end:2018// CHECK25-NEXT: [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]2019// CHECK25-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 42020// CHECK25-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42021// CHECK25-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 42022// CHECK25-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2023// CHECK25: omp.inner.for.cond:2024// CHECK25-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42025// CHECK25-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42026// CHECK25-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]2027// CHECK25-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2028// CHECK25: omp.inner.for.body:2029// CHECK25-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42030// CHECK25-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 12031// CHECK25-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]2032// CHECK25-NEXT: store i32 [[ADD]], ptr [[I]], align 42033// CHECK25-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 42034// CHECK25-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP9]] to i642035// CHECK25-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]2036// CHECK25-NEXT: store i32 0, ptr [[ARRAYIDX]], align 42037// CHECK25-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]2038// CHECK25: omp.body.continue:2039// CHECK25-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2040// CHECK25: omp.inner.for.inc:2041// CHECK25-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42042// CHECK25-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP10]], 12043// CHECK25-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 42044// CHECK25-NEXT: br label [[OMP_INNER_FOR_COND]]2045// CHECK25: omp.inner.for.end:2046// CHECK25-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2047// CHECK25: omp.loop.exit:2048// CHECK25-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])2049// CHECK25-NEXT: ret void2050//2051//2052// CHECK27-LABEL: define {{[^@]+}}@main2053// CHECK27-SAME: (i32 noundef [[ARGC:%.*]], ptr noundef [[ARGV:%.*]]) #[[ATTR0:[0-9]+]] {2054// CHECK27-NEXT: entry:2055// CHECK27-NEXT: [[RETVAL:%.*]] = alloca i32, align 42056// CHECK27-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 42057// CHECK27-NEXT: [[ARGV_ADDR:%.*]] = alloca ptr, align 42058// CHECK27-NEXT: [[N:%.*]] = alloca i32, align 42059// CHECK27-NEXT: [[SAVED_STACK:%.*]] = alloca ptr, align 42060// CHECK27-NEXT: [[__VLA_EXPR0:%.*]] = alloca i32, align 42061// CHECK27-NEXT: [[N_CASTED:%.*]] = alloca i32, align 42062// CHECK27-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 42063// CHECK27-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 42064// CHECK27-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 42065// CHECK27-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 42066// CHECK27-NEXT: [[TMP:%.*]] = alloca i32, align 42067// CHECK27-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42068// CHECK27-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42069// CHECK27-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 82070// CHECK27-NEXT: store i32 0, ptr [[RETVAL]], align 42071// CHECK27-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 42072// CHECK27-NEXT: store ptr [[ARGV]], ptr [[ARGV_ADDR]], align 42073// CHECK27-NEXT: store i32 100, ptr [[N]], align 42074// CHECK27-NEXT: [[TMP0:%.*]] = load i32, ptr [[N]], align 42075// CHECK27-NEXT: [[TMP1:%.*]] = call ptr @llvm.stacksave.p0()2076// CHECK27-NEXT: store ptr [[TMP1]], ptr [[SAVED_STACK]], align 42077// CHECK27-NEXT: [[VLA:%.*]] = alloca i32, i32 [[TMP0]], align 42078// CHECK27-NEXT: store i32 [[TMP0]], ptr [[__VLA_EXPR0]], align 42079// CHECK27-NEXT: [[TMP2:%.*]] = load i32, ptr [[N]], align 42080// CHECK27-NEXT: store i32 [[TMP2]], ptr [[N_CASTED]], align 42081// CHECK27-NEXT: [[TMP3:%.*]] = load i32, ptr [[N_CASTED]], align 42082// CHECK27-NEXT: [[TMP4:%.*]] = mul nuw i32 [[TMP0]], 42083// CHECK27-NEXT: [[TMP5:%.*]] = sext i32 [[TMP4]] to i642084// CHECK27-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[DOTOFFLOAD_SIZES]], ptr align 4 @.offload_sizes, i32 24, i1 false)2085// CHECK27-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 02086// CHECK27-NEXT: store i32 [[TMP3]], ptr [[TMP6]], align 42087// CHECK27-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 02088// CHECK27-NEXT: store i32 [[TMP3]], ptr [[TMP7]], align 42089// CHECK27-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 02090// CHECK27-NEXT: store ptr null, ptr [[TMP8]], align 42091// CHECK27-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 12092// CHECK27-NEXT: store i32 [[TMP0]], ptr [[TMP9]], align 42093// CHECK27-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 12094// CHECK27-NEXT: store i32 [[TMP0]], ptr [[TMP10]], align 42095// CHECK27-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 12096// CHECK27-NEXT: store ptr null, ptr [[TMP11]], align 42097// CHECK27-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 22098// CHECK27-NEXT: store ptr [[VLA]], ptr [[TMP12]], align 42099// CHECK27-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 22100// CHECK27-NEXT: store ptr [[VLA]], ptr [[TMP13]], align 42101// CHECK27-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 22102// CHECK27-NEXT: store i64 [[TMP5]], ptr [[TMP14]], align 42103// CHECK27-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 22104// CHECK27-NEXT: store ptr null, ptr [[TMP15]], align 42105// CHECK27-NEXT: [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 02106// CHECK27-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 02107// CHECK27-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 02108// CHECK27-NEXT: [[TMP19:%.*]] = load i32, ptr [[N]], align 42109// CHECK27-NEXT: store i32 [[TMP19]], ptr [[DOTCAPTURE_EXPR_]], align 42110// CHECK27-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42111// CHECK27-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP20]], 02112// CHECK27-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12113// CHECK27-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12114// CHECK27-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42115// CHECK27-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42116// CHECK27-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP21]], 12117// CHECK27-NEXT: [[TMP22:%.*]] = zext i32 [[ADD]] to i642118// CHECK27-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 02119// CHECK27-NEXT: store i32 3, ptr [[TMP23]], align 42120// CHECK27-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12121// CHECK27-NEXT: store i32 3, ptr [[TMP24]], align 42122// CHECK27-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 22123// CHECK27-NEXT: store ptr [[TMP16]], ptr [[TMP25]], align 42124// CHECK27-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 32125// CHECK27-NEXT: store ptr [[TMP17]], ptr [[TMP26]], align 42126// CHECK27-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 42127// CHECK27-NEXT: store ptr [[TMP18]], ptr [[TMP27]], align 42128// CHECK27-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 52129// CHECK27-NEXT: store ptr @.offload_maptypes, ptr [[TMP28]], align 42130// CHECK27-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 62131// CHECK27-NEXT: store ptr null, ptr [[TMP29]], align 42132// CHECK27-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 72133// CHECK27-NEXT: store ptr null, ptr [[TMP30]], align 42134// CHECK27-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 82135// CHECK27-NEXT: store i64 [[TMP22]], ptr [[TMP31]], align 82136// CHECK27-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 92137// CHECK27-NEXT: store i64 0, ptr [[TMP32]], align 82138// CHECK27-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 102139// CHECK27-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP33]], align 42140// CHECK27-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 112141// CHECK27-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP34]], align 42142// CHECK27-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 122143// CHECK27-NEXT: store i32 0, ptr [[TMP35]], align 42144// 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_l162.region_id, ptr [[KERNEL_ARGS]])2145// CHECK27-NEXT: [[TMP37:%.*]] = icmp ne i32 [[TMP36]], 02146// CHECK27-NEXT: br i1 [[TMP37]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]2147// CHECK27: omp_offload.failed:2148// CHECK27-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l162(i32 [[TMP3]], i32 [[TMP0]], ptr [[VLA]]) #[[ATTR3:[0-9]+]]2149// CHECK27-NEXT: br label [[OMP_OFFLOAD_CONT]]2150// CHECK27: omp_offload.cont:2151// CHECK27-NEXT: [[TMP38:%.*]] = load i32, ptr [[ARGC_ADDR]], align 42152// CHECK27-NEXT: [[CALL:%.*]] = call noundef i32 @_Z5tmainIiLi10EEiT_(i32 noundef [[TMP38]])2153// CHECK27-NEXT: store i32 [[CALL]], ptr [[RETVAL]], align 42154// CHECK27-NEXT: [[TMP39:%.*]] = load ptr, ptr [[SAVED_STACK]], align 42155// CHECK27-NEXT: call void @llvm.stackrestore.p0(ptr [[TMP39]])2156// CHECK27-NEXT: [[TMP40:%.*]] = load i32, ptr [[RETVAL]], align 42157// CHECK27-NEXT: ret i32 [[TMP40]]2158//2159//2160// CHECK27-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l1622161// CHECK27-SAME: (i32 noundef [[N:%.*]], i32 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2:[0-9]+]] {2162// CHECK27-NEXT: entry:2163// CHECK27-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42164// CHECK27-NEXT: [[VLA_ADDR:%.*]] = alloca i32, align 42165// CHECK27-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42166// CHECK27-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42167// CHECK27-NEXT: store i32 [[VLA]], ptr [[VLA_ADDR]], align 42168// CHECK27-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42169// CHECK27-NEXT: [[TMP0:%.*]] = load i32, ptr [[VLA_ADDR]], align 42170// CHECK27-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 42171// CHECK27-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l162.omp_outlined, ptr [[N_ADDR]], i32 [[TMP0]], ptr [[TMP1]])2172// CHECK27-NEXT: ret void2173//2174//2175// CHECK27-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l162.omp_outlined2176// 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]] {2177// CHECK27-NEXT: entry:2178// CHECK27-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42179// CHECK27-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42180// CHECK27-NEXT: [[N_ADDR:%.*]] = alloca ptr, align 42181// CHECK27-NEXT: [[VLA_ADDR:%.*]] = alloca i32, align 42182// CHECK27-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42183// CHECK27-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42184// CHECK27-NEXT: [[TMP:%.*]] = alloca i32, align 42185// CHECK27-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42186// CHECK27-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42187// CHECK27-NEXT: [[I:%.*]] = alloca i32, align 42188// CHECK27-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 42189// CHECK27-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 42190// CHECK27-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42191// CHECK27-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42192// CHECK27-NEXT: [[I3:%.*]] = alloca i32, align 42193// CHECK27-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42194// CHECK27-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42195// CHECK27-NEXT: store ptr [[N]], ptr [[N_ADDR]], align 42196// CHECK27-NEXT: store i32 [[VLA]], ptr [[VLA_ADDR]], align 42197// CHECK27-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42198// CHECK27-NEXT: [[TMP0:%.*]] = load ptr, ptr [[N_ADDR]], align 42199// CHECK27-NEXT: [[TMP1:%.*]] = load i32, ptr [[VLA_ADDR]], align 42200// CHECK27-NEXT: [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 42201// CHECK27-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 42202// CHECK27-NEXT: store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR_]], align 42203// CHECK27-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42204// CHECK27-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP4]], 02205// CHECK27-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12206// CHECK27-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12207// CHECK27-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42208// CHECK27-NEXT: store i32 0, ptr [[I]], align 42209// CHECK27-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42210// CHECK27-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP5]]2211// CHECK27-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2212// CHECK27: omp.precond.then:2213// CHECK27-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 42214// CHECK27-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42215// CHECK27-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_UB]], align 42216// CHECK27-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42217// CHECK27-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42218// CHECK27-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42219// CHECK27-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 42220// CHECK27-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP8]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2221// CHECK27-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42222// CHECK27-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42223// CHECK27-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP9]], [[TMP10]]2224// CHECK27-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2225// CHECK27: cond.true:2226// CHECK27-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42227// CHECK27-NEXT: br label [[COND_END:%.*]]2228// CHECK27: cond.false:2229// CHECK27-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42230// CHECK27-NEXT: br label [[COND_END]]2231// CHECK27: cond.end:2232// CHECK27-NEXT: [[COND:%.*]] = phi i32 [ [[TMP11]], [[COND_TRUE]] ], [ [[TMP12]], [[COND_FALSE]] ]2233// CHECK27-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 42234// CHECK27-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42235// CHECK27-NEXT: store i32 [[TMP13]], ptr [[DOTOMP_IV]], align 42236// CHECK27-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2237// CHECK27: omp.inner.for.cond:2238// CHECK27-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42239// CHECK27-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42240// CHECK27-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP14]], [[TMP15]]2241// CHECK27-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2242// CHECK27: omp.inner.for.body:2243// CHECK27-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42244// CHECK27-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP16]], 12245// CHECK27-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]2246// CHECK27-NEXT: store i32 [[ADD]], ptr [[I3]], align 42247// CHECK27-NEXT: [[TMP17:%.*]] = load i32, ptr [[I3]], align 42248// CHECK27-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i32 [[TMP17]]2249// CHECK27-NEXT: store i32 0, ptr [[ARRAYIDX]], align 42250// CHECK27-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]2251// CHECK27: omp.body.continue:2252// CHECK27-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2253// CHECK27: omp.inner.for.inc:2254// CHECK27-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42255// CHECK27-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP18]], 12256// CHECK27-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 42257// CHECK27-NEXT: br label [[OMP_INNER_FOR_COND]]2258// CHECK27: omp.inner.for.end:2259// CHECK27-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2260// CHECK27: omp.loop.exit:2261// CHECK27-NEXT: [[TMP19:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42262// CHECK27-NEXT: [[TMP20:%.*]] = load i32, ptr [[TMP19]], align 42263// CHECK27-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP20]])2264// CHECK27-NEXT: br label [[OMP_PRECOND_END]]2265// CHECK27: omp.precond.end:2266// CHECK27-NEXT: ret void2267//2268//2269// CHECK27-LABEL: define {{[^@]+}}@_Z5tmainIiLi10EEiT_2270// CHECK27-SAME: (i32 noundef [[ARGC:%.*]]) #[[ATTR5:[0-9]+]] comdat {2271// CHECK27-NEXT: entry:2272// CHECK27-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 42273// CHECK27-NEXT: [[A:%.*]] = alloca [10 x i32], align 42274// CHECK27-NEXT: [[TE:%.*]] = alloca i32, align 42275// CHECK27-NEXT: [[TH:%.*]] = alloca i32, align 42276// CHECK27-NEXT: [[TE_CASTED:%.*]] = alloca i32, align 42277// CHECK27-NEXT: [[TH_CASTED:%.*]] = alloca i32, align 42278// CHECK27-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 42279// CHECK27-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 42280// CHECK27-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 42281// CHECK27-NEXT: [[TMP:%.*]] = alloca i32, align 42282// CHECK27-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 82283// CHECK27-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 42284// CHECK27-NEXT: store i32 0, ptr [[TE]], align 42285// CHECK27-NEXT: store i32 128, ptr [[TH]], align 42286// CHECK27-NEXT: [[TMP0:%.*]] = load i32, ptr [[TE]], align 42287// CHECK27-NEXT: store i32 [[TMP0]], ptr [[TE_CASTED]], align 42288// CHECK27-NEXT: [[TMP1:%.*]] = load i32, ptr [[TE_CASTED]], align 42289// CHECK27-NEXT: [[TMP2:%.*]] = load i32, ptr [[TH]], align 42290// CHECK27-NEXT: store i32 [[TMP2]], ptr [[TH_CASTED]], align 42291// CHECK27-NEXT: [[TMP3:%.*]] = load i32, ptr [[TH_CASTED]], align 42292// CHECK27-NEXT: [[TMP4:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 02293// CHECK27-NEXT: store i32 [[TMP1]], ptr [[TMP4]], align 42294// CHECK27-NEXT: [[TMP5:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 02295// CHECK27-NEXT: store i32 [[TMP1]], ptr [[TMP5]], align 42296// CHECK27-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 02297// CHECK27-NEXT: store ptr null, ptr [[TMP6]], align 42298// CHECK27-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 12299// CHECK27-NEXT: store i32 [[TMP3]], ptr [[TMP7]], align 42300// CHECK27-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 12301// CHECK27-NEXT: store i32 [[TMP3]], ptr [[TMP8]], align 42302// CHECK27-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 12303// CHECK27-NEXT: store ptr null, ptr [[TMP9]], align 42304// CHECK27-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 22305// CHECK27-NEXT: store ptr [[A]], ptr [[TMP10]], align 42306// CHECK27-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 22307// CHECK27-NEXT: store ptr [[A]], ptr [[TMP11]], align 42308// CHECK27-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 22309// CHECK27-NEXT: store ptr null, ptr [[TMP12]], align 42310// CHECK27-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 02311// CHECK27-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 02312// CHECK27-NEXT: [[TMP15:%.*]] = load i32, ptr [[TE]], align 42313// CHECK27-NEXT: [[TMP16:%.*]] = load i32, ptr [[TH]], align 42314// CHECK27-NEXT: [[TMP17:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP15]], 02315// CHECK27-NEXT: [[TMP18:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP16]], 02316// CHECK27-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 02317// CHECK27-NEXT: store i32 3, ptr [[TMP19]], align 42318// CHECK27-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12319// CHECK27-NEXT: store i32 3, ptr [[TMP20]], align 42320// CHECK27-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 22321// CHECK27-NEXT: store ptr [[TMP13]], ptr [[TMP21]], align 42322// CHECK27-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 32323// CHECK27-NEXT: store ptr [[TMP14]], ptr [[TMP22]], align 42324// CHECK27-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 42325// CHECK27-NEXT: store ptr @.offload_sizes.1, ptr [[TMP23]], align 42326// CHECK27-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 52327// CHECK27-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP24]], align 42328// CHECK27-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 62329// CHECK27-NEXT: store ptr null, ptr [[TMP25]], align 42330// CHECK27-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 72331// CHECK27-NEXT: store ptr null, ptr [[TMP26]], align 42332// CHECK27-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 82333// CHECK27-NEXT: store i64 10, ptr [[TMP27]], align 82334// CHECK27-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 92335// CHECK27-NEXT: store i64 0, ptr [[TMP28]], align 82336// CHECK27-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 102337// CHECK27-NEXT: store [3 x i32] [[TMP17]], ptr [[TMP29]], align 42338// CHECK27-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 112339// CHECK27-NEXT: store [3 x i32] [[TMP18]], ptr [[TMP30]], align 42340// CHECK27-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 122341// CHECK27-NEXT: store i32 0, ptr [[TMP31]], align 42342// 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__l151.region_id, ptr [[KERNEL_ARGS]])2343// CHECK27-NEXT: [[TMP33:%.*]] = icmp ne i32 [[TMP32]], 02344// CHECK27-NEXT: br i1 [[TMP33]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]2345// CHECK27: omp_offload.failed:2346// CHECK27-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l151(i32 [[TMP1]], i32 [[TMP3]], ptr [[A]]) #[[ATTR3]]2347// CHECK27-NEXT: br label [[OMP_OFFLOAD_CONT]]2348// CHECK27: omp_offload.cont:2349// CHECK27-NEXT: ret i32 02350//2351//2352// CHECK27-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l1512353// CHECK27-SAME: (i32 noundef [[TE:%.*]], i32 noundef [[TH:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {2354// CHECK27-NEXT: entry:2355// CHECK27-NEXT: [[TE_ADDR:%.*]] = alloca i32, align 42356// CHECK27-NEXT: [[TH_ADDR:%.*]] = alloca i32, align 42357// CHECK27-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42358// CHECK27-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB2]])2359// CHECK27-NEXT: store i32 [[TE]], ptr [[TE_ADDR]], align 42360// CHECK27-NEXT: store i32 [[TH]], ptr [[TH_ADDR]], align 42361// CHECK27-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42362// CHECK27-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 42363// CHECK27-NEXT: [[TMP2:%.*]] = load i32, ptr [[TE_ADDR]], align 42364// CHECK27-NEXT: [[TMP3:%.*]] = load i32, ptr [[TH_ADDR]], align 42365// CHECK27-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB2]], i32 [[TMP0]], i32 [[TMP2]], i32 [[TMP3]])2366// CHECK27-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l151.omp_outlined, ptr [[TMP1]])2367// CHECK27-NEXT: ret void2368//2369//2370// CHECK27-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiLi10EEiT__l151.omp_outlined2371// CHECK27-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR2]] {2372// CHECK27-NEXT: entry:2373// CHECK27-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42374// CHECK27-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42375// CHECK27-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 42376// CHECK27-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42377// CHECK27-NEXT: [[TMP:%.*]] = alloca i32, align 42378// CHECK27-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 42379// CHECK27-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 42380// CHECK27-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42381// CHECK27-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42382// CHECK27-NEXT: [[I:%.*]] = alloca i32, align 42383// CHECK27-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42384// CHECK27-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42385// CHECK27-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 42386// CHECK27-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 42387// CHECK27-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 42388// CHECK27-NEXT: store i32 9, ptr [[DOTOMP_UB]], align 42389// CHECK27-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42390// CHECK27-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42391// CHECK27-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42392// CHECK27-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 42393// CHECK27-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2394// CHECK27-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42395// CHECK27-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 92396// CHECK27-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2397// CHECK27: cond.true:2398// CHECK27-NEXT: br label [[COND_END:%.*]]2399// CHECK27: cond.false:2400// CHECK27-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42401// CHECK27-NEXT: br label [[COND_END]]2402// CHECK27: cond.end:2403// CHECK27-NEXT: [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]2404// CHECK27-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 42405// CHECK27-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42406// CHECK27-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 42407// CHECK27-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2408// CHECK27: omp.inner.for.cond:2409// CHECK27-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42410// CHECK27-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42411// CHECK27-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]2412// CHECK27-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2413// CHECK27: omp.inner.for.body:2414// CHECK27-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42415// CHECK27-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 12416// CHECK27-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]2417// CHECK27-NEXT: store i32 [[ADD]], ptr [[I]], align 42418// CHECK27-NEXT: [[TMP9:%.*]] = load i32, ptr [[I]], align 42419// CHECK27-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP9]]2420// CHECK27-NEXT: store i32 0, ptr [[ARRAYIDX]], align 42421// CHECK27-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]2422// CHECK27: omp.body.continue:2423// CHECK27-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2424// CHECK27: omp.inner.for.inc:2425// CHECK27-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 42426// CHECK27-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP10]], 12427// CHECK27-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 42428// CHECK27-NEXT: br label [[OMP_INNER_FOR_COND]]2429// CHECK27: omp.inner.for.end:2430// CHECK27-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2431// CHECK27: omp.loop.exit:2432// CHECK27-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])2433// CHECK27-NEXT: ret void2434//2435