1605 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]+" --prefix-filecheck-ir-name _2// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s3 4// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s5// expected-no-diagnostics6 7struct ST {8 int *a;9};10typedef int arr[10];11typedef ST STarr[10];12struct SA {13 const int da[5] = { 0 };14 ST g[10];15 STarr &rg = g;16 int i;17 int &j = i;18 int *k = &j;19 int *&z = k;20 int aa[10];21 arr &raa = aa;22 void func(int arg) {23#pragma omp target has_device_addr(k)24 {k++;}25#pragma omp target has_device_addr(z)26 {z++;}27#pragma omp target has_device_addr(aa)28 {aa[0]=1;}29#pragma omp target has_device_addr(raa)30 {raa[0] = 10;}31#pragma omp target has_device_addr(g)32 {g[0].a= &i;}33#pragma omp target has_device_addr(da)34 {int a = da[1];}35 return;36 }37};38 39struct SB {40 unsigned A;41 unsigned B;42 float Arr[100];43 float *Ptr;44 float *foo() {45 return &Arr[0];46 }47};48 49struct SC {50 unsigned A : 2;51 unsigned B : 3;52 unsigned C;53 unsigned D;54 float Arr[100];55 SB S;56 SB ArrS[100];57 SB *PtrS;58 SB *&RPtrS;59 float *Ptr;60 61 SC(SB *&_RPtrS) : RPtrS(_RPtrS) {}62};63 64union SD {65 unsigned A;66 float B;67};68 69struct S1;70extern S1 a;71class S2 {72 mutable int a;73public:74 S2():a(0) { }75 S2(S2 &s2):a(s2.a) { }76 static float S2s;77 static const float S2sc;78};79const float S2::S2sc = 0;80const S2 b;81const S2 ba[5];82class S3 {83 int a;84public:85 S3():a(0) { }86 S3(S3 &s3):a(s3.a) { }87};88const S3 c;89const S3 ca[5];90extern const int f;91class S4 {92 int a;93 S4();94 S4(const S4 &s4);95public:96 S4(int v):a(v) { }97};98class S5 {99 int a;100 S5():a(0) {}101 S5(const S5 &s5):a(s5.a) { }102public:103 S5(int v):a(v) { }104};105 106S3 h;107#pragma omp threadprivate(h)108 109typedef struct {110 int a;111} S6;112 113template <typename T>114T tmain(T argc) {115 const T da[5] = { 0 };116 S6 h[10];117 auto &rh = h;118 T i;119 T &j = i;120 T *k = &j;121 T *&z = k;122 T aa[10];123#pragma omp target has_device_addr(k)124 {k++;}125#pragma omp target has_device_addr(z)126 {z++;}127#pragma omp target has_device_addr(aa)128 {T a = aa[0];}129#pragma omp target has_device_addr(h)130 {int a = h[0].a;}131 return 0;132}133 134 135int main(int argc, char **argv) {136 const int da[5] = { 0 };137 S6 h[10];138 auto &rh = h;139 int i;140 int &j = i;141 int *k = &j;142 int *&z = k;143 int aa[10];144 auto &raa = aa;145#pragma omp target has_device_addr(k)146 {k++;}147#pragma omp target has_device_addr(z)148 {z++;}149#pragma omp target has_device_addr(aa)150 {aa[0]=1;}151#pragma omp target has_device_addr(raa)152 {int a = raa[0];}153#pragma omp target has_device_addr(h)154 {int a = h[1].a;}155#pragma omp target has_device_addr(da[1:3])156 {int a = da[1];}157 return tmain<int>(argc) + *tmain<int *>(&argc);158}159 160struct SomeKernel {161 int targetDev;162 float devPtr;163 SomeKernel();164 ~SomeKernel();165 166 template<unsigned int nRHS>167 void apply() {168 #pragma omp target has_device_addr(devPtr) device(targetDev)169 {170 devPtr++;171 targetDev++;172 }173 }174};175 176void use_template() {177 SomeKernel aKern;178 aKern.apply<32>();179}180// CHECK-LABEL: define {{[^@]+}}@__cxx_global_var_init181// CHECK-SAME: () #[[ATTR0:[0-9]+]] {182// CHECK-NEXT: entry:183// CHECK-NEXT: call void @_ZN2S2C1Ev(ptr noundef nonnull align 4 dereferenceable(4) @_ZL1b)184// CHECK-NEXT: ret void185//186//187// CHECK-LABEL: define {{[^@]+}}@_ZN2S2C1Ev188// CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1:[0-9]+]] comdat {189// CHECK-NEXT: entry:190// CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8191// CHECK-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8192// CHECK-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8193// CHECK-NEXT: call void @_ZN2S2C2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]])194// CHECK-NEXT: ret void195//196//197// CHECK-LABEL: define {{[^@]+}}@_ZN2S2C2Ev198// CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {199// CHECK-NEXT: entry:200// CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8201// CHECK-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8202// CHECK-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8203// CHECK-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[CLASS_S2:%.*]], ptr [[THIS1]], i32 0, i32 0204// CHECK-NEXT: store i32 0, ptr [[A]], align 4205// CHECK-NEXT: ret void206//207//208// CHECK-LABEL: define {{[^@]+}}@__cxx_global_var_init.1209// CHECK-SAME: () #[[ATTR0]] {210// CHECK-NEXT: entry:211// CHECK-NEXT: br label [[ARRAYCTOR_LOOP:%.*]]212// CHECK: arrayctor.loop:213// CHECK-NEXT: [[ARRAYCTOR_CUR:%.*]] = phi ptr [ @_ZL2ba, [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ]214// CHECK-NEXT: call void @_ZN2S2C1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]])215// CHECK-NEXT: [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[CLASS_S2:%.*]], ptr [[ARRAYCTOR_CUR]], i64 1216// CHECK-NEXT: [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], getelementptr inbounds ([[CLASS_S2]], ptr @_ZL2ba, i64 5)217// CHECK-NEXT: br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]]218// CHECK: arrayctor.cont:219// CHECK-NEXT: ret void220//221//222// CHECK-LABEL: define {{[^@]+}}@__cxx_global_var_init.2223// CHECK-SAME: () #[[ATTR0]] {224// CHECK-NEXT: entry:225// CHECK-NEXT: call void @_ZN2S3C1Ev(ptr noundef nonnull align 4 dereferenceable(4) @_ZL1c)226// CHECK-NEXT: ret void227//228//229// CHECK-LABEL: define {{[^@]+}}@_ZN2S3C1Ev230// CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {231// CHECK-NEXT: entry:232// CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8233// CHECK-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8234// CHECK-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8235// CHECK-NEXT: call void @_ZN2S3C2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]])236// CHECK-NEXT: ret void237//238//239// CHECK-LABEL: define {{[^@]+}}@_ZN2S3C2Ev240// CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {241// CHECK-NEXT: entry:242// CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8243// CHECK-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8244// CHECK-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8245// CHECK-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[CLASS_S3:%.*]], ptr [[THIS1]], i32 0, i32 0246// CHECK-NEXT: store i32 0, ptr [[A]], align 4247// CHECK-NEXT: ret void248//249//250// CHECK-LABEL: define {{[^@]+}}@__cxx_global_var_init.3251// CHECK-SAME: () #[[ATTR0]] {252// CHECK-NEXT: entry:253// CHECK-NEXT: br label [[ARRAYCTOR_LOOP:%.*]]254// CHECK: arrayctor.loop:255// CHECK-NEXT: [[ARRAYCTOR_CUR:%.*]] = phi ptr [ @_ZL2ca, [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ]256// CHECK-NEXT: call void @_ZN2S3C1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]])257// CHECK-NEXT: [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[CLASS_S3:%.*]], ptr [[ARRAYCTOR_CUR]], i64 1258// CHECK-NEXT: [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], getelementptr inbounds ([[CLASS_S3]], ptr @_ZL2ca, i64 5)259// CHECK-NEXT: br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]]260// CHECK: arrayctor.cont:261// CHECK-NEXT: ret void262//263//264// CHECK-LABEL: define {{[^@]+}}@__cxx_global_var_init.4265// CHECK-SAME: () #[[ATTR0]] {266// CHECK-NEXT: entry:267// CHECK-NEXT: call void @_ZN2S3C1Ev(ptr noundef nonnull align 4 dereferenceable(4) @h)268// CHECK-NEXT: ret void269//270//271// CHECK-LABEL: define {{[^@]+}}@main272// CHECK-SAME: (i32 noundef signext [[ARGC:%.*]], ptr noundef [[ARGV:%.*]]) #[[ATTR2:[0-9]+]] {273// CHECK-NEXT: entry:274// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4275// CHECK-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 4276// CHECK-NEXT: [[ARGV_ADDR:%.*]] = alloca ptr, align 8277// CHECK-NEXT: [[DA:%.*]] = alloca [5 x i32], align 4278// CHECK-NEXT: [[H:%.*]] = alloca [10 x %struct.S6], align 4279// CHECK-NEXT: [[RH:%.*]] = alloca ptr, align 8280// CHECK-NEXT: [[I:%.*]] = alloca i32, align 4281// CHECK-NEXT: [[J:%.*]] = alloca ptr, align 8282// CHECK-NEXT: [[K:%.*]] = alloca ptr, align 8283// CHECK-NEXT: [[Z:%.*]] = alloca ptr, align 8284// CHECK-NEXT: [[AA:%.*]] = alloca [10 x i32], align 4285// CHECK-NEXT: [[RAA:%.*]] = alloca ptr, align 8286// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8287// CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8288// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8289// CHECK-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8290// CHECK-NEXT: [[TMP:%.*]] = alloca ptr, align 8291// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS1:%.*]] = alloca [1 x ptr], align 8292// CHECK-NEXT: [[DOTOFFLOAD_PTRS2:%.*]] = alloca [1 x ptr], align 8293// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS3:%.*]] = alloca [1 x ptr], align 8294// CHECK-NEXT: [[KERNEL_ARGS4:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8295// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS7:%.*]] = alloca [1 x ptr], align 8296// CHECK-NEXT: [[DOTOFFLOAD_PTRS8:%.*]] = alloca [1 x ptr], align 8297// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS9:%.*]] = alloca [1 x ptr], align 8298// CHECK-NEXT: [[KERNEL_ARGS10:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8299// CHECK-NEXT: [[_TMP13:%.*]] = alloca ptr, align 8300// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS14:%.*]] = alloca [1 x ptr], align 8301// CHECK-NEXT: [[DOTOFFLOAD_PTRS15:%.*]] = alloca [1 x ptr], align 8302// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS16:%.*]] = alloca [1 x ptr], align 8303// CHECK-NEXT: [[KERNEL_ARGS17:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8304// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS20:%.*]] = alloca [1 x ptr], align 8305// CHECK-NEXT: [[DOTOFFLOAD_PTRS21:%.*]] = alloca [1 x ptr], align 8306// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS22:%.*]] = alloca [1 x ptr], align 8307// CHECK-NEXT: [[KERNEL_ARGS23:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8308// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS26:%.*]] = alloca [1 x ptr], align 8309// CHECK-NEXT: [[DOTOFFLOAD_PTRS27:%.*]] = alloca [1 x ptr], align 8310// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS28:%.*]] = alloca [1 x ptr], align 8311// CHECK-NEXT: [[KERNEL_ARGS29:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8312// CHECK-NEXT: store i32 0, ptr [[RETVAL]], align 4313// CHECK-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 4314// CHECK-NEXT: store ptr [[ARGV]], ptr [[ARGV_ADDR]], align 8315// CHECK-NEXT: call void @llvm.memset.p0.i64(ptr align 4 [[DA]], i8 0, i64 20, i1 false)316// CHECK-NEXT: store ptr [[H]], ptr [[RH]], align 8317// CHECK-NEXT: store ptr [[I]], ptr [[J]], align 8318// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[J]], align 8319// CHECK-NEXT: store ptr [[TMP0]], ptr [[K]], align 8320// CHECK-NEXT: store ptr [[K]], ptr [[Z]], align 8321// CHECK-NEXT: store ptr [[AA]], ptr [[RAA]], align 8322// CHECK-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0323// CHECK-NEXT: store ptr [[K]], ptr [[TMP1]], align 8324// CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0325// CHECK-NEXT: store ptr [[K]], ptr [[TMP2]], align 8326// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0327// CHECK-NEXT: store ptr null, ptr [[TMP3]], align 8328// CHECK-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0329// CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0330// CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0331// CHECK-NEXT: store i32 3, ptr [[TMP6]], align 4332// CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1333// CHECK-NEXT: store i32 1, ptr [[TMP7]], align 4334// CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2335// CHECK-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 8336// CHECK-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3337// CHECK-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 8338// CHECK-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4339// CHECK-NEXT: store ptr @.offload_sizes, ptr [[TMP10]], align 8340// CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5341// CHECK-NEXT: store ptr @.offload_maptypes, ptr [[TMP11]], align 8342// CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6343// CHECK-NEXT: store ptr null, ptr [[TMP12]], align 8344// CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7345// CHECK-NEXT: store ptr null, ptr [[TMP13]], align 8346// CHECK-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8347// CHECK-NEXT: store i64 0, ptr [[TMP14]], align 8348// CHECK-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9349// CHECK-NEXT: store i64 0, ptr [[TMP15]], align 8350// CHECK-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10351// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP16]], align 4352// CHECK-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11353// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4354// CHECK-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12355// CHECK-NEXT: store i32 0, ptr [[TMP18]], align 4356// CHECK-NEXT: [[TMP19:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l145.region_id, ptr [[KERNEL_ARGS]])357// CHECK-NEXT: [[TMP20:%.*]] = icmp ne i32 [[TMP19]], 0358// CHECK-NEXT: br i1 [[TMP20]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]359// CHECK: omp_offload.failed:360// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l145(ptr [[K]]) #[[ATTR5:[0-9]+]]361// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT]]362// CHECK: omp_offload.cont:363// CHECK-NEXT: [[TMP21:%.*]] = load ptr, ptr [[Z]], align 8364// CHECK-NEXT: store ptr [[TMP21]], ptr [[TMP]], align 8365// CHECK-NEXT: [[TMP22:%.*]] = load ptr, ptr [[TMP]], align 8366// CHECK-NEXT: [[TMP23:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0367// CHECK-NEXT: store ptr [[TMP22]], ptr [[TMP23]], align 8368// CHECK-NEXT: [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 0369// CHECK-NEXT: store ptr [[TMP22]], ptr [[TMP24]], align 8370// CHECK-NEXT: [[TMP25:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS3]], i64 0, i64 0371// CHECK-NEXT: store ptr null, ptr [[TMP25]], align 8372// CHECK-NEXT: [[TMP26:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0373// CHECK-NEXT: [[TMP27:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 0374// CHECK-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 0375// CHECK-NEXT: store i32 3, ptr [[TMP28]], align 4376// CHECK-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 1377// CHECK-NEXT: store i32 1, ptr [[TMP29]], align 4378// CHECK-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 2379// CHECK-NEXT: store ptr [[TMP26]], ptr [[TMP30]], align 8380// CHECK-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 3381// CHECK-NEXT: store ptr [[TMP27]], ptr [[TMP31]], align 8382// CHECK-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 4383// CHECK-NEXT: store ptr @.offload_sizes.5, ptr [[TMP32]], align 8384// CHECK-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 5385// CHECK-NEXT: store ptr @.offload_maptypes.6, ptr [[TMP33]], align 8386// CHECK-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 6387// CHECK-NEXT: store ptr null, ptr [[TMP34]], align 8388// CHECK-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 7389// CHECK-NEXT: store ptr null, ptr [[TMP35]], align 8390// CHECK-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 8391// CHECK-NEXT: store i64 0, ptr [[TMP36]], align 8392// CHECK-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 9393// CHECK-NEXT: store i64 0, ptr [[TMP37]], align 8394// CHECK-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 10395// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP38]], align 4396// CHECK-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 11397// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP39]], align 4398// CHECK-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 12399// CHECK-NEXT: store i32 0, ptr [[TMP40]], align 4400// CHECK-NEXT: [[TMP41:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l147.region_id, ptr [[KERNEL_ARGS4]])401// CHECK-NEXT: [[TMP42:%.*]] = icmp ne i32 [[TMP41]], 0402// CHECK-NEXT: br i1 [[TMP42]], label [[OMP_OFFLOAD_FAILED5:%.*]], label [[OMP_OFFLOAD_CONT6:%.*]]403// CHECK: omp_offload.failed5:404// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l147(ptr [[TMP22]]) #[[ATTR5]]405// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT6]]406// CHECK: omp_offload.cont6:407// CHECK-NEXT: [[TMP43:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 0408// CHECK-NEXT: store ptr [[AA]], ptr [[TMP43]], align 8409// CHECK-NEXT: [[TMP44:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 0410// CHECK-NEXT: store ptr [[AA]], ptr [[TMP44]], align 8411// CHECK-NEXT: [[TMP45:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 0412// CHECK-NEXT: store ptr null, ptr [[TMP45]], align 8413// CHECK-NEXT: [[TMP46:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 0414// CHECK-NEXT: [[TMP47:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 0415// CHECK-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 0416// CHECK-NEXT: store i32 3, ptr [[TMP48]], align 4417// CHECK-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 1418// CHECK-NEXT: store i32 1, ptr [[TMP49]], align 4419// CHECK-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 2420// CHECK-NEXT: store ptr [[TMP46]], ptr [[TMP50]], align 8421// CHECK-NEXT: [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 3422// CHECK-NEXT: store ptr [[TMP47]], ptr [[TMP51]], align 8423// CHECK-NEXT: [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 4424// CHECK-NEXT: store ptr @.offload_sizes.7, ptr [[TMP52]], align 8425// CHECK-NEXT: [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 5426// CHECK-NEXT: store ptr @.offload_maptypes.8, ptr [[TMP53]], align 8427// CHECK-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 6428// CHECK-NEXT: store ptr null, ptr [[TMP54]], align 8429// CHECK-NEXT: [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 7430// CHECK-NEXT: store ptr null, ptr [[TMP55]], align 8431// CHECK-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 8432// CHECK-NEXT: store i64 0, ptr [[TMP56]], align 8433// CHECK-NEXT: [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 9434// CHECK-NEXT: store i64 0, ptr [[TMP57]], align 8435// CHECK-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 10436// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP58]], align 4437// CHECK-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 11438// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP59]], align 4439// CHECK-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 12440// CHECK-NEXT: store i32 0, ptr [[TMP60]], align 4441// CHECK-NEXT: [[TMP61:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l149.region_id, ptr [[KERNEL_ARGS10]])442// CHECK-NEXT: [[TMP62:%.*]] = icmp ne i32 [[TMP61]], 0443// CHECK-NEXT: br i1 [[TMP62]], label [[OMP_OFFLOAD_FAILED11:%.*]], label [[OMP_OFFLOAD_CONT12:%.*]]444// CHECK: omp_offload.failed11:445// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l149(ptr [[AA]]) #[[ATTR5]]446// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT12]]447// CHECK: omp_offload.cont12:448// CHECK-NEXT: [[TMP63:%.*]] = load ptr, ptr [[RAA]], align 8449// CHECK-NEXT: store ptr [[TMP63]], ptr [[_TMP13]], align 8450// CHECK-NEXT: [[TMP64:%.*]] = load ptr, ptr [[_TMP13]], align 8451// CHECK-NEXT: [[TMP65:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS14]], i32 0, i32 0452// CHECK-NEXT: store ptr [[TMP64]], ptr [[TMP65]], align 8453// CHECK-NEXT: [[TMP66:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS15]], i32 0, i32 0454// CHECK-NEXT: store ptr [[TMP64]], ptr [[TMP66]], align 8455// CHECK-NEXT: [[TMP67:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS16]], i64 0, i64 0456// CHECK-NEXT: store ptr null, ptr [[TMP67]], align 8457// CHECK-NEXT: [[TMP68:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS14]], i32 0, i32 0458// CHECK-NEXT: [[TMP69:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS15]], i32 0, i32 0459// CHECK-NEXT: [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 0460// CHECK-NEXT: store i32 3, ptr [[TMP70]], align 4461// CHECK-NEXT: [[TMP71:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 1462// CHECK-NEXT: store i32 1, ptr [[TMP71]], align 4463// CHECK-NEXT: [[TMP72:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 2464// CHECK-NEXT: store ptr [[TMP68]], ptr [[TMP72]], align 8465// CHECK-NEXT: [[TMP73:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 3466// CHECK-NEXT: store ptr [[TMP69]], ptr [[TMP73]], align 8467// CHECK-NEXT: [[TMP74:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 4468// CHECK-NEXT: store ptr @.offload_sizes.9, ptr [[TMP74]], align 8469// CHECK-NEXT: [[TMP75:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 5470// CHECK-NEXT: store ptr @.offload_maptypes.10, ptr [[TMP75]], align 8471// CHECK-NEXT: [[TMP76:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 6472// CHECK-NEXT: store ptr null, ptr [[TMP76]], align 8473// CHECK-NEXT: [[TMP77:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 7474// CHECK-NEXT: store ptr null, ptr [[TMP77]], align 8475// CHECK-NEXT: [[TMP78:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 8476// CHECK-NEXT: store i64 0, ptr [[TMP78]], align 8477// CHECK-NEXT: [[TMP79:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 9478// CHECK-NEXT: store i64 0, ptr [[TMP79]], align 8479// CHECK-NEXT: [[TMP80:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 10480// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP80]], align 4481// CHECK-NEXT: [[TMP81:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 11482// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP81]], align 4483// CHECK-NEXT: [[TMP82:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 12484// CHECK-NEXT: store i32 0, ptr [[TMP82]], align 4485// CHECK-NEXT: [[TMP83:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l151.region_id, ptr [[KERNEL_ARGS17]])486// CHECK-NEXT: [[TMP84:%.*]] = icmp ne i32 [[TMP83]], 0487// CHECK-NEXT: br i1 [[TMP84]], label [[OMP_OFFLOAD_FAILED18:%.*]], label [[OMP_OFFLOAD_CONT19:%.*]]488// CHECK: omp_offload.failed18:489// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l151(ptr [[TMP64]]) #[[ATTR5]]490// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT19]]491// CHECK: omp_offload.cont19:492// CHECK-NEXT: [[TMP85:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS20]], i32 0, i32 0493// CHECK-NEXT: store ptr [[H]], ptr [[TMP85]], align 8494// CHECK-NEXT: [[TMP86:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS21]], i32 0, i32 0495// CHECK-NEXT: store ptr [[H]], ptr [[TMP86]], align 8496// CHECK-NEXT: [[TMP87:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS22]], i64 0, i64 0497// CHECK-NEXT: store ptr null, ptr [[TMP87]], align 8498// CHECK-NEXT: [[TMP88:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS20]], i32 0, i32 0499// CHECK-NEXT: [[TMP89:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS21]], i32 0, i32 0500// CHECK-NEXT: [[TMP90:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 0501// CHECK-NEXT: store i32 3, ptr [[TMP90]], align 4502// CHECK-NEXT: [[TMP91:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 1503// CHECK-NEXT: store i32 1, ptr [[TMP91]], align 4504// CHECK-NEXT: [[TMP92:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 2505// CHECK-NEXT: store ptr [[TMP88]], ptr [[TMP92]], align 8506// CHECK-NEXT: [[TMP93:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 3507// CHECK-NEXT: store ptr [[TMP89]], ptr [[TMP93]], align 8508// CHECK-NEXT: [[TMP94:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 4509// CHECK-NEXT: store ptr @.offload_sizes.11, ptr [[TMP94]], align 8510// CHECK-NEXT: [[TMP95:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 5511// CHECK-NEXT: store ptr @.offload_maptypes.12, ptr [[TMP95]], align 8512// CHECK-NEXT: [[TMP96:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 6513// CHECK-NEXT: store ptr null, ptr [[TMP96]], align 8514// CHECK-NEXT: [[TMP97:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 7515// CHECK-NEXT: store ptr null, ptr [[TMP97]], align 8516// CHECK-NEXT: [[TMP98:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 8517// CHECK-NEXT: store i64 0, ptr [[TMP98]], align 8518// CHECK-NEXT: [[TMP99:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 9519// CHECK-NEXT: store i64 0, ptr [[TMP99]], align 8520// CHECK-NEXT: [[TMP100:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 10521// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP100]], align 4522// CHECK-NEXT: [[TMP101:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 11523// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP101]], align 4524// CHECK-NEXT: [[TMP102:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 12525// CHECK-NEXT: store i32 0, ptr [[TMP102]], align 4526// CHECK-NEXT: [[TMP103:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l153.region_id, ptr [[KERNEL_ARGS23]])527// CHECK-NEXT: [[TMP104:%.*]] = icmp ne i32 [[TMP103]], 0528// CHECK-NEXT: br i1 [[TMP104]], label [[OMP_OFFLOAD_FAILED24:%.*]], label [[OMP_OFFLOAD_CONT25:%.*]]529// CHECK: omp_offload.failed24:530// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l153(ptr [[H]]) #[[ATTR5]]531// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT25]]532// CHECK: omp_offload.cont25:533// CHECK-NEXT: [[TMP105:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS26]], i32 0, i32 0534// CHECK-NEXT: store ptr [[DA]], ptr [[TMP105]], align 8535// CHECK-NEXT: [[TMP106:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS27]], i32 0, i32 0536// CHECK-NEXT: store ptr [[DA]], ptr [[TMP106]], align 8537// CHECK-NEXT: [[TMP107:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS28]], i64 0, i64 0538// CHECK-NEXT: store ptr null, ptr [[TMP107]], align 8539// CHECK-NEXT: [[TMP108:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS26]], i32 0, i32 0540// CHECK-NEXT: [[TMP109:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS27]], i32 0, i32 0541// CHECK-NEXT: [[TMP110:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 0542// CHECK-NEXT: store i32 3, ptr [[TMP110]], align 4543// CHECK-NEXT: [[TMP111:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 1544// CHECK-NEXT: store i32 1, ptr [[TMP111]], align 4545// CHECK-NEXT: [[TMP112:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 2546// CHECK-NEXT: store ptr [[TMP108]], ptr [[TMP112]], align 8547// CHECK-NEXT: [[TMP113:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 3548// CHECK-NEXT: store ptr [[TMP109]], ptr [[TMP113]], align 8549// CHECK-NEXT: [[TMP114:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 4550// CHECK-NEXT: store ptr @.offload_sizes.13, ptr [[TMP114]], align 8551// CHECK-NEXT: [[TMP115:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 5552// CHECK-NEXT: store ptr @.offload_maptypes.14, ptr [[TMP115]], align 8553// CHECK-NEXT: [[TMP116:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 6554// CHECK-NEXT: store ptr null, ptr [[TMP116]], align 8555// CHECK-NEXT: [[TMP117:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 7556// CHECK-NEXT: store ptr null, ptr [[TMP117]], align 8557// CHECK-NEXT: [[TMP118:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 8558// CHECK-NEXT: store i64 0, ptr [[TMP118]], align 8559// CHECK-NEXT: [[TMP119:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 9560// CHECK-NEXT: store i64 0, ptr [[TMP119]], align 8561// CHECK-NEXT: [[TMP120:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 10562// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP120]], align 4563// CHECK-NEXT: [[TMP121:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 11564// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP121]], align 4565// CHECK-NEXT: [[TMP122:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 12566// CHECK-NEXT: store i32 0, ptr [[TMP122]], align 4567// CHECK-NEXT: [[TMP123:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l155.region_id, ptr [[KERNEL_ARGS29]])568// CHECK-NEXT: [[TMP124:%.*]] = icmp ne i32 [[TMP123]], 0569// CHECK-NEXT: br i1 [[TMP124]], label [[OMP_OFFLOAD_FAILED30:%.*]], label [[OMP_OFFLOAD_CONT31:%.*]]570// CHECK: omp_offload.failed30:571// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l155(ptr [[DA]]) #[[ATTR5]]572// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT31]]573// CHECK: omp_offload.cont31:574// CHECK-NEXT: [[TMP125:%.*]] = load i32, ptr [[ARGC_ADDR]], align 4575// CHECK-NEXT: [[CALL:%.*]] = call noundef signext i32 @_Z5tmainIiET_S0_(i32 noundef signext [[TMP125]])576// CHECK-NEXT: [[CALL32:%.*]] = call noundef ptr @_Z5tmainIPiET_S1_(ptr noundef [[ARGC_ADDR]])577// CHECK-NEXT: [[TMP126:%.*]] = load i32, ptr [[CALL32]], align 4578// CHECK-NEXT: [[ADD:%.*]] = add nsw i32 [[CALL]], [[TMP126]]579// CHECK-NEXT: ret i32 [[ADD]]580//581//582// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l145583// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[K:%.*]]) #[[ATTR4:[0-9]+]] {584// CHECK-NEXT: entry:585// CHECK-NEXT: [[K_ADDR:%.*]] = alloca ptr, align 8586// CHECK-NEXT: store ptr [[K]], ptr [[K_ADDR]], align 8587// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[K_ADDR]], align 8588// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8589// CHECK-NEXT: [[INCDEC_PTR:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP1]], i32 1590// CHECK-NEXT: store ptr [[INCDEC_PTR]], ptr [[TMP0]], align 8591// CHECK-NEXT: ret void592//593//594// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l147595// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[Z:%.*]]) #[[ATTR4]] {596// CHECK-NEXT: entry:597// CHECK-NEXT: [[Z_ADDR:%.*]] = alloca ptr, align 8598// CHECK-NEXT: [[TMP:%.*]] = alloca ptr, align 8599// CHECK-NEXT: store ptr [[Z]], ptr [[Z_ADDR]], align 8600// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[Z_ADDR]], align 8601// CHECK-NEXT: store ptr [[TMP0]], ptr [[TMP]], align 8602// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP]], align 8603// CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[TMP1]], align 8604// CHECK-NEXT: [[INCDEC_PTR:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP2]], i32 1605// CHECK-NEXT: store ptr [[INCDEC_PTR]], ptr [[TMP1]], align 8606// CHECK-NEXT: ret void607//608//609// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l149610// CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[AA:%.*]]) #[[ATTR4]] {611// CHECK-NEXT: entry:612// CHECK-NEXT: [[AA_ADDR:%.*]] = alloca ptr, align 8613// CHECK-NEXT: store ptr [[AA]], ptr [[AA_ADDR]], align 8614// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 8615// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 0616// CHECK-NEXT: store i32 1, ptr [[ARRAYIDX]], align 4617// CHECK-NEXT: ret void618//619//620// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l151621// CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[RAA:%.*]]) #[[ATTR4]] {622// CHECK-NEXT: entry:623// CHECK-NEXT: [[RAA_ADDR:%.*]] = alloca ptr, align 8624// CHECK-NEXT: [[TMP:%.*]] = alloca ptr, align 8625// CHECK-NEXT: [[A:%.*]] = alloca i32, align 4626// CHECK-NEXT: store ptr [[RAA]], ptr [[RAA_ADDR]], align 8627// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[RAA_ADDR]], align 8628// CHECK-NEXT: store ptr [[TMP0]], ptr [[TMP]], align 8629// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP]], align 8630// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP1]], i64 0, i64 0631// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[ARRAYIDX]], align 4632// CHECK-NEXT: store i32 [[TMP2]], ptr [[A]], align 4633// CHECK-NEXT: ret void634//635//636// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l153637// CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[H:%.*]]) #[[ATTR4]] {638// CHECK-NEXT: entry:639// CHECK-NEXT: [[H_ADDR:%.*]] = alloca ptr, align 8640// CHECK-NEXT: [[A:%.*]] = alloca i32, align 4641// CHECK-NEXT: store ptr [[H]], ptr [[H_ADDR]], align 8642// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[H_ADDR]], align 8643// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x %struct.S6], ptr [[TMP0]], i64 0, i64 1644// CHECK-NEXT: [[A1:%.*]] = getelementptr inbounds nuw [[STRUCT_S6:%.*]], ptr [[ARRAYIDX]], i32 0, i32 0645// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[A1]], align 4646// CHECK-NEXT: store i32 [[TMP1]], ptr [[A]], align 4647// CHECK-NEXT: ret void648//649//650// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l155651// CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(20) [[DA:%.*]]) #[[ATTR4]] {652// CHECK-NEXT: entry:653// CHECK-NEXT: [[DA_ADDR:%.*]] = alloca ptr, align 8654// CHECK-NEXT: [[A:%.*]] = alloca i32, align 4655// CHECK-NEXT: store ptr [[DA]], ptr [[DA_ADDR]], align 8656// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DA_ADDR]], align 8657// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [5 x i32], ptr [[TMP0]], i64 0, i64 1658// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[ARRAYIDX]], align 4659// CHECK-NEXT: store i32 [[TMP1]], ptr [[A]], align 4660// CHECK-NEXT: ret void661//662//663// CHECK-LABEL: define {{[^@]+}}@_Z5tmainIiET_S0_664// CHECK-SAME: (i32 noundef signext [[ARGC:%.*]]) #[[ATTR1]] comdat {665// CHECK-NEXT: entry:666// CHECK-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 4667// CHECK-NEXT: [[DA:%.*]] = alloca [5 x i32], align 4668// CHECK-NEXT: [[H:%.*]] = alloca [10 x %struct.S6], align 4669// CHECK-NEXT: [[RH:%.*]] = alloca ptr, align 8670// CHECK-NEXT: [[I:%.*]] = alloca i32, align 4671// CHECK-NEXT: [[J:%.*]] = alloca ptr, align 8672// CHECK-NEXT: [[K:%.*]] = alloca ptr, align 8673// CHECK-NEXT: [[Z:%.*]] = alloca ptr, align 8674// CHECK-NEXT: [[AA:%.*]] = alloca [10 x i32], align 4675// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8676// CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8677// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8678// CHECK-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8679// CHECK-NEXT: [[TMP:%.*]] = alloca ptr, align 8680// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS1:%.*]] = alloca [1 x ptr], align 8681// CHECK-NEXT: [[DOTOFFLOAD_PTRS2:%.*]] = alloca [1 x ptr], align 8682// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS3:%.*]] = alloca [1 x ptr], align 8683// CHECK-NEXT: [[KERNEL_ARGS4:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8684// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS7:%.*]] = alloca [1 x ptr], align 8685// CHECK-NEXT: [[DOTOFFLOAD_PTRS8:%.*]] = alloca [1 x ptr], align 8686// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS9:%.*]] = alloca [1 x ptr], align 8687// CHECK-NEXT: [[KERNEL_ARGS10:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8688// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS13:%.*]] = alloca [1 x ptr], align 8689// CHECK-NEXT: [[DOTOFFLOAD_PTRS14:%.*]] = alloca [1 x ptr], align 8690// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS15:%.*]] = alloca [1 x ptr], align 8691// CHECK-NEXT: [[KERNEL_ARGS16:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8692// CHECK-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 4693// CHECK-NEXT: call void @llvm.memset.p0.i64(ptr align 4 [[DA]], i8 0, i64 20, i1 false)694// CHECK-NEXT: store ptr [[H]], ptr [[RH]], align 8695// CHECK-NEXT: store ptr [[I]], ptr [[J]], align 8696// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[J]], align 8697// CHECK-NEXT: store ptr [[TMP0]], ptr [[K]], align 8698// CHECK-NEXT: store ptr [[K]], ptr [[Z]], align 8699// CHECK-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0700// CHECK-NEXT: store ptr [[K]], ptr [[TMP1]], align 8701// CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0702// CHECK-NEXT: store ptr [[K]], ptr [[TMP2]], align 8703// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0704// CHECK-NEXT: store ptr null, ptr [[TMP3]], align 8705// CHECK-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0706// CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0707// CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0708// CHECK-NEXT: store i32 3, ptr [[TMP6]], align 4709// CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1710// CHECK-NEXT: store i32 1, ptr [[TMP7]], align 4711// CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2712// CHECK-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 8713// CHECK-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3714// CHECK-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 8715// CHECK-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4716// CHECK-NEXT: store ptr @.offload_sizes.15, ptr [[TMP10]], align 8717// CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5718// CHECK-NEXT: store ptr @.offload_maptypes.16, ptr [[TMP11]], align 8719// CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6720// CHECK-NEXT: store ptr null, ptr [[TMP12]], align 8721// CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7722// CHECK-NEXT: store ptr null, ptr [[TMP13]], align 8723// CHECK-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8724// CHECK-NEXT: store i64 0, ptr [[TMP14]], align 8725// CHECK-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9726// CHECK-NEXT: store i64 0, ptr [[TMP15]], align 8727// CHECK-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10728// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP16]], align 4729// CHECK-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11730// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4731// CHECK-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12732// CHECK-NEXT: store i32 0, ptr [[TMP18]], align 4733// CHECK-NEXT: [[TMP19:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l123.region_id, ptr [[KERNEL_ARGS]])734// CHECK-NEXT: [[TMP20:%.*]] = icmp ne i32 [[TMP19]], 0735// CHECK-NEXT: br i1 [[TMP20]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]736// CHECK: omp_offload.failed:737// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l123(ptr [[K]]) #[[ATTR5]]738// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT]]739// CHECK: omp_offload.cont:740// CHECK-NEXT: [[TMP21:%.*]] = load ptr, ptr [[Z]], align 8741// CHECK-NEXT: store ptr [[TMP21]], ptr [[TMP]], align 8742// CHECK-NEXT: [[TMP22:%.*]] = load ptr, ptr [[TMP]], align 8743// CHECK-NEXT: [[TMP23:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0744// CHECK-NEXT: store ptr [[TMP22]], ptr [[TMP23]], align 8745// CHECK-NEXT: [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 0746// CHECK-NEXT: store ptr [[TMP22]], ptr [[TMP24]], align 8747// CHECK-NEXT: [[TMP25:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS3]], i64 0, i64 0748// CHECK-NEXT: store ptr null, ptr [[TMP25]], align 8749// CHECK-NEXT: [[TMP26:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0750// CHECK-NEXT: [[TMP27:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 0751// CHECK-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 0752// CHECK-NEXT: store i32 3, ptr [[TMP28]], align 4753// CHECK-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 1754// CHECK-NEXT: store i32 1, ptr [[TMP29]], align 4755// CHECK-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 2756// CHECK-NEXT: store ptr [[TMP26]], ptr [[TMP30]], align 8757// CHECK-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 3758// CHECK-NEXT: store ptr [[TMP27]], ptr [[TMP31]], align 8759// CHECK-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 4760// CHECK-NEXT: store ptr @.offload_sizes.17, ptr [[TMP32]], align 8761// CHECK-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 5762// CHECK-NEXT: store ptr @.offload_maptypes.18, ptr [[TMP33]], align 8763// CHECK-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 6764// CHECK-NEXT: store ptr null, ptr [[TMP34]], align 8765// CHECK-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 7766// CHECK-NEXT: store ptr null, ptr [[TMP35]], align 8767// CHECK-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 8768// CHECK-NEXT: store i64 0, ptr [[TMP36]], align 8769// CHECK-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 9770// CHECK-NEXT: store i64 0, ptr [[TMP37]], align 8771// CHECK-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 10772// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP38]], align 4773// CHECK-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 11774// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP39]], align 4775// CHECK-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 12776// CHECK-NEXT: store i32 0, ptr [[TMP40]], align 4777// CHECK-NEXT: [[TMP41:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l125.region_id, ptr [[KERNEL_ARGS4]])778// CHECK-NEXT: [[TMP42:%.*]] = icmp ne i32 [[TMP41]], 0779// CHECK-NEXT: br i1 [[TMP42]], label [[OMP_OFFLOAD_FAILED5:%.*]], label [[OMP_OFFLOAD_CONT6:%.*]]780// CHECK: omp_offload.failed5:781// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l125(ptr [[TMP22]]) #[[ATTR5]]782// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT6]]783// CHECK: omp_offload.cont6:784// CHECK-NEXT: [[TMP43:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 0785// CHECK-NEXT: store ptr [[AA]], ptr [[TMP43]], align 8786// CHECK-NEXT: [[TMP44:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 0787// CHECK-NEXT: store ptr [[AA]], ptr [[TMP44]], align 8788// CHECK-NEXT: [[TMP45:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 0789// CHECK-NEXT: store ptr null, ptr [[TMP45]], align 8790// CHECK-NEXT: [[TMP46:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 0791// CHECK-NEXT: [[TMP47:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 0792// CHECK-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 0793// CHECK-NEXT: store i32 3, ptr [[TMP48]], align 4794// CHECK-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 1795// CHECK-NEXT: store i32 1, ptr [[TMP49]], align 4796// CHECK-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 2797// CHECK-NEXT: store ptr [[TMP46]], ptr [[TMP50]], align 8798// CHECK-NEXT: [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 3799// CHECK-NEXT: store ptr [[TMP47]], ptr [[TMP51]], align 8800// CHECK-NEXT: [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 4801// CHECK-NEXT: store ptr @.offload_sizes.19, ptr [[TMP52]], align 8802// CHECK-NEXT: [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 5803// CHECK-NEXT: store ptr @.offload_maptypes.20, ptr [[TMP53]], align 8804// CHECK-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 6805// CHECK-NEXT: store ptr null, ptr [[TMP54]], align 8806// CHECK-NEXT: [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 7807// CHECK-NEXT: store ptr null, ptr [[TMP55]], align 8808// CHECK-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 8809// CHECK-NEXT: store i64 0, ptr [[TMP56]], align 8810// CHECK-NEXT: [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 9811// CHECK-NEXT: store i64 0, ptr [[TMP57]], align 8812// CHECK-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 10813// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP58]], align 4814// CHECK-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 11815// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP59]], align 4816// CHECK-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 12817// CHECK-NEXT: store i32 0, ptr [[TMP60]], align 4818// CHECK-NEXT: [[TMP61:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l127.region_id, ptr [[KERNEL_ARGS10]])819// CHECK-NEXT: [[TMP62:%.*]] = icmp ne i32 [[TMP61]], 0820// CHECK-NEXT: br i1 [[TMP62]], label [[OMP_OFFLOAD_FAILED11:%.*]], label [[OMP_OFFLOAD_CONT12:%.*]]821// CHECK: omp_offload.failed11:822// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l127(ptr [[AA]]) #[[ATTR5]]823// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT12]]824// CHECK: omp_offload.cont12:825// CHECK-NEXT: [[TMP63:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS13]], i32 0, i32 0826// CHECK-NEXT: store ptr [[H]], ptr [[TMP63]], align 8827// CHECK-NEXT: [[TMP64:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS14]], i32 0, i32 0828// CHECK-NEXT: store ptr [[H]], ptr [[TMP64]], align 8829// CHECK-NEXT: [[TMP65:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS15]], i64 0, i64 0830// CHECK-NEXT: store ptr null, ptr [[TMP65]], align 8831// CHECK-NEXT: [[TMP66:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS13]], i32 0, i32 0832// CHECK-NEXT: [[TMP67:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS14]], i32 0, i32 0833// CHECK-NEXT: [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 0834// CHECK-NEXT: store i32 3, ptr [[TMP68]], align 4835// CHECK-NEXT: [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 1836// CHECK-NEXT: store i32 1, ptr [[TMP69]], align 4837// CHECK-NEXT: [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 2838// CHECK-NEXT: store ptr [[TMP66]], ptr [[TMP70]], align 8839// CHECK-NEXT: [[TMP71:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 3840// CHECK-NEXT: store ptr [[TMP67]], ptr [[TMP71]], align 8841// CHECK-NEXT: [[TMP72:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 4842// CHECK-NEXT: store ptr @.offload_sizes.21, ptr [[TMP72]], align 8843// CHECK-NEXT: [[TMP73:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 5844// CHECK-NEXT: store ptr @.offload_maptypes.22, ptr [[TMP73]], align 8845// CHECK-NEXT: [[TMP74:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 6846// CHECK-NEXT: store ptr null, ptr [[TMP74]], align 8847// CHECK-NEXT: [[TMP75:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 7848// CHECK-NEXT: store ptr null, ptr [[TMP75]], align 8849// CHECK-NEXT: [[TMP76:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 8850// CHECK-NEXT: store i64 0, ptr [[TMP76]], align 8851// CHECK-NEXT: [[TMP77:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 9852// CHECK-NEXT: store i64 0, ptr [[TMP77]], align 8853// CHECK-NEXT: [[TMP78:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 10854// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP78]], align 4855// CHECK-NEXT: [[TMP79:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 11856// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP79]], align 4857// CHECK-NEXT: [[TMP80:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 12858// CHECK-NEXT: store i32 0, ptr [[TMP80]], align 4859// CHECK-NEXT: [[TMP81:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l129.region_id, ptr [[KERNEL_ARGS16]])860// CHECK-NEXT: [[TMP82:%.*]] = icmp ne i32 [[TMP81]], 0861// CHECK-NEXT: br i1 [[TMP82]], label [[OMP_OFFLOAD_FAILED17:%.*]], label [[OMP_OFFLOAD_CONT18:%.*]]862// CHECK: omp_offload.failed17:863// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l129(ptr [[H]]) #[[ATTR5]]864// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT18]]865// CHECK: omp_offload.cont18:866// CHECK-NEXT: ret i32 0867//868//869// CHECK-LABEL: define {{[^@]+}}@_Z5tmainIPiET_S1_870// CHECK-SAME: (ptr noundef [[ARGC:%.*]]) #[[ATTR1]] comdat {871// CHECK-NEXT: entry:872// CHECK-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 8873// CHECK-NEXT: [[DA:%.*]] = alloca [5 x ptr], align 8874// CHECK-NEXT: [[H:%.*]] = alloca [10 x %struct.S6], align 4875// CHECK-NEXT: [[RH:%.*]] = alloca ptr, align 8876// CHECK-NEXT: [[I:%.*]] = alloca ptr, align 8877// CHECK-NEXT: [[J:%.*]] = alloca ptr, align 8878// CHECK-NEXT: [[K:%.*]] = alloca ptr, align 8879// CHECK-NEXT: [[Z:%.*]] = alloca ptr, align 8880// CHECK-NEXT: [[AA:%.*]] = alloca [10 x ptr], align 8881// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8882// CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8883// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8884// CHECK-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8885// CHECK-NEXT: [[TMP:%.*]] = alloca ptr, align 8886// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS1:%.*]] = alloca [1 x ptr], align 8887// CHECK-NEXT: [[DOTOFFLOAD_PTRS2:%.*]] = alloca [1 x ptr], align 8888// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS3:%.*]] = alloca [1 x ptr], align 8889// CHECK-NEXT: [[KERNEL_ARGS4:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8890// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS7:%.*]] = alloca [1 x ptr], align 8891// CHECK-NEXT: [[DOTOFFLOAD_PTRS8:%.*]] = alloca [1 x ptr], align 8892// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS9:%.*]] = alloca [1 x ptr], align 8893// CHECK-NEXT: [[KERNEL_ARGS10:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8894// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS13:%.*]] = alloca [1 x ptr], align 8895// CHECK-NEXT: [[DOTOFFLOAD_PTRS14:%.*]] = alloca [1 x ptr], align 8896// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS15:%.*]] = alloca [1 x ptr], align 8897// CHECK-NEXT: [[KERNEL_ARGS16:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8898// CHECK-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 8899// CHECK-NEXT: call void @llvm.memset.p0.i64(ptr align 8 [[DA]], i8 0, i64 40, i1 false)900// CHECK-NEXT: store ptr [[H]], ptr [[RH]], align 8901// CHECK-NEXT: store ptr [[I]], ptr [[J]], align 8902// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[J]], align 8903// CHECK-NEXT: store ptr [[TMP0]], ptr [[K]], align 8904// CHECK-NEXT: store ptr [[K]], ptr [[Z]], align 8905// CHECK-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0906// CHECK-NEXT: store ptr [[K]], ptr [[TMP1]], align 8907// CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0908// CHECK-NEXT: store ptr [[K]], ptr [[TMP2]], align 8909// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0910// CHECK-NEXT: store ptr null, ptr [[TMP3]], align 8911// CHECK-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0912// CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0913// CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0914// CHECK-NEXT: store i32 3, ptr [[TMP6]], align 4915// CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1916// CHECK-NEXT: store i32 1, ptr [[TMP7]], align 4917// CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2918// CHECK-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 8919// CHECK-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3920// CHECK-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 8921// CHECK-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4922// CHECK-NEXT: store ptr @.offload_sizes.23, ptr [[TMP10]], align 8923// CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5924// CHECK-NEXT: store ptr @.offload_maptypes.24, ptr [[TMP11]], align 8925// CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6926// CHECK-NEXT: store ptr null, ptr [[TMP12]], align 8927// CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7928// CHECK-NEXT: store ptr null, ptr [[TMP13]], align 8929// CHECK-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8930// CHECK-NEXT: store i64 0, ptr [[TMP14]], align 8931// CHECK-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9932// CHECK-NEXT: store i64 0, ptr [[TMP15]], align 8933// CHECK-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10934// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP16]], align 4935// CHECK-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11936// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4937// CHECK-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12938// CHECK-NEXT: store i32 0, ptr [[TMP18]], align 4939// CHECK-NEXT: [[TMP19:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l123.region_id, ptr [[KERNEL_ARGS]])940// CHECK-NEXT: [[TMP20:%.*]] = icmp ne i32 [[TMP19]], 0941// CHECK-NEXT: br i1 [[TMP20]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]942// CHECK: omp_offload.failed:943// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l123(ptr [[K]]) #[[ATTR5]]944// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT]]945// CHECK: omp_offload.cont:946// CHECK-NEXT: [[TMP21:%.*]] = load ptr, ptr [[Z]], align 8947// CHECK-NEXT: store ptr [[TMP21]], ptr [[TMP]], align 8948// CHECK-NEXT: [[TMP22:%.*]] = load ptr, ptr [[TMP]], align 8949// CHECK-NEXT: [[TMP23:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0950// CHECK-NEXT: store ptr [[TMP22]], ptr [[TMP23]], align 8951// CHECK-NEXT: [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 0952// CHECK-NEXT: store ptr [[TMP22]], ptr [[TMP24]], align 8953// CHECK-NEXT: [[TMP25:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS3]], i64 0, i64 0954// CHECK-NEXT: store ptr null, ptr [[TMP25]], align 8955// CHECK-NEXT: [[TMP26:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0956// CHECK-NEXT: [[TMP27:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 0957// CHECK-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 0958// CHECK-NEXT: store i32 3, ptr [[TMP28]], align 4959// CHECK-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 1960// CHECK-NEXT: store i32 1, ptr [[TMP29]], align 4961// CHECK-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 2962// CHECK-NEXT: store ptr [[TMP26]], ptr [[TMP30]], align 8963// CHECK-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 3964// CHECK-NEXT: store ptr [[TMP27]], ptr [[TMP31]], align 8965// CHECK-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 4966// CHECK-NEXT: store ptr @.offload_sizes.25, ptr [[TMP32]], align 8967// CHECK-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 5968// CHECK-NEXT: store ptr @.offload_maptypes.26, ptr [[TMP33]], align 8969// CHECK-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 6970// CHECK-NEXT: store ptr null, ptr [[TMP34]], align 8971// CHECK-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 7972// CHECK-NEXT: store ptr null, ptr [[TMP35]], align 8973// CHECK-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 8974// CHECK-NEXT: store i64 0, ptr [[TMP36]], align 8975// CHECK-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 9976// CHECK-NEXT: store i64 0, ptr [[TMP37]], align 8977// CHECK-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 10978// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP38]], align 4979// CHECK-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 11980// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP39]], align 4981// CHECK-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 12982// CHECK-NEXT: store i32 0, ptr [[TMP40]], align 4983// CHECK-NEXT: [[TMP41:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l125.region_id, ptr [[KERNEL_ARGS4]])984// CHECK-NEXT: [[TMP42:%.*]] = icmp ne i32 [[TMP41]], 0985// CHECK-NEXT: br i1 [[TMP42]], label [[OMP_OFFLOAD_FAILED5:%.*]], label [[OMP_OFFLOAD_CONT6:%.*]]986// CHECK: omp_offload.failed5:987// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l125(ptr [[TMP22]]) #[[ATTR5]]988// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT6]]989// CHECK: omp_offload.cont6:990// CHECK-NEXT: [[TMP43:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 0991// CHECK-NEXT: store ptr [[AA]], ptr [[TMP43]], align 8992// CHECK-NEXT: [[TMP44:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 0993// CHECK-NEXT: store ptr [[AA]], ptr [[TMP44]], align 8994// CHECK-NEXT: [[TMP45:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 0995// CHECK-NEXT: store ptr null, ptr [[TMP45]], align 8996// CHECK-NEXT: [[TMP46:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 0997// CHECK-NEXT: [[TMP47:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 0998// CHECK-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 0999// CHECK-NEXT: store i32 3, ptr [[TMP48]], align 41000// CHECK-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 11001// CHECK-NEXT: store i32 1, ptr [[TMP49]], align 41002// CHECK-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 21003// CHECK-NEXT: store ptr [[TMP46]], ptr [[TMP50]], align 81004// CHECK-NEXT: [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 31005// CHECK-NEXT: store ptr [[TMP47]], ptr [[TMP51]], align 81006// CHECK-NEXT: [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 41007// CHECK-NEXT: store ptr @.offload_sizes.27, ptr [[TMP52]], align 81008// CHECK-NEXT: [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 51009// CHECK-NEXT: store ptr @.offload_maptypes.28, ptr [[TMP53]], align 81010// CHECK-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 61011// CHECK-NEXT: store ptr null, ptr [[TMP54]], align 81012// CHECK-NEXT: [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 71013// CHECK-NEXT: store ptr null, ptr [[TMP55]], align 81014// CHECK-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 81015// CHECK-NEXT: store i64 0, ptr [[TMP56]], align 81016// CHECK-NEXT: [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 91017// CHECK-NEXT: store i64 0, ptr [[TMP57]], align 81018// CHECK-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 101019// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP58]], align 41020// CHECK-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 111021// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP59]], align 41022// CHECK-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 121023// CHECK-NEXT: store i32 0, ptr [[TMP60]], align 41024// CHECK-NEXT: [[TMP61:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l127.region_id, ptr [[KERNEL_ARGS10]])1025// CHECK-NEXT: [[TMP62:%.*]] = icmp ne i32 [[TMP61]], 01026// CHECK-NEXT: br i1 [[TMP62]], label [[OMP_OFFLOAD_FAILED11:%.*]], label [[OMP_OFFLOAD_CONT12:%.*]]1027// CHECK: omp_offload.failed11:1028// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l127(ptr [[AA]]) #[[ATTR5]]1029// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT12]]1030// CHECK: omp_offload.cont12:1031// CHECK-NEXT: [[TMP63:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS13]], i32 0, i32 01032// CHECK-NEXT: store ptr [[H]], ptr [[TMP63]], align 81033// CHECK-NEXT: [[TMP64:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS14]], i32 0, i32 01034// CHECK-NEXT: store ptr [[H]], ptr [[TMP64]], align 81035// CHECK-NEXT: [[TMP65:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS15]], i64 0, i64 01036// CHECK-NEXT: store ptr null, ptr [[TMP65]], align 81037// CHECK-NEXT: [[TMP66:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS13]], i32 0, i32 01038// CHECK-NEXT: [[TMP67:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS14]], i32 0, i32 01039// CHECK-NEXT: [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 01040// CHECK-NEXT: store i32 3, ptr [[TMP68]], align 41041// CHECK-NEXT: [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 11042// CHECK-NEXT: store i32 1, ptr [[TMP69]], align 41043// CHECK-NEXT: [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 21044// CHECK-NEXT: store ptr [[TMP66]], ptr [[TMP70]], align 81045// CHECK-NEXT: [[TMP71:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 31046// CHECK-NEXT: store ptr [[TMP67]], ptr [[TMP71]], align 81047// CHECK-NEXT: [[TMP72:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 41048// CHECK-NEXT: store ptr @.offload_sizes.29, ptr [[TMP72]], align 81049// CHECK-NEXT: [[TMP73:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 51050// CHECK-NEXT: store ptr @.offload_maptypes.30, ptr [[TMP73]], align 81051// CHECK-NEXT: [[TMP74:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 61052// CHECK-NEXT: store ptr null, ptr [[TMP74]], align 81053// CHECK-NEXT: [[TMP75:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 71054// CHECK-NEXT: store ptr null, ptr [[TMP75]], align 81055// CHECK-NEXT: [[TMP76:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 81056// CHECK-NEXT: store i64 0, ptr [[TMP76]], align 81057// CHECK-NEXT: [[TMP77:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 91058// CHECK-NEXT: store i64 0, ptr [[TMP77]], align 81059// CHECK-NEXT: [[TMP78:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 101060// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP78]], align 41061// CHECK-NEXT: [[TMP79:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 111062// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP79]], align 41063// CHECK-NEXT: [[TMP80:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 121064// CHECK-NEXT: store i32 0, ptr [[TMP80]], align 41065// CHECK-NEXT: [[TMP81:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l129.region_id, ptr [[KERNEL_ARGS16]])1066// CHECK-NEXT: [[TMP82:%.*]] = icmp ne i32 [[TMP81]], 01067// CHECK-NEXT: br i1 [[TMP82]], label [[OMP_OFFLOAD_FAILED17:%.*]], label [[OMP_OFFLOAD_CONT18:%.*]]1068// CHECK: omp_offload.failed17:1069// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l129(ptr [[H]]) #[[ATTR5]]1070// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT18]]1071// CHECK: omp_offload.cont18:1072// CHECK-NEXT: ret ptr null1073//1074//1075// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l1231076// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[K:%.*]]) #[[ATTR4]] {1077// CHECK-NEXT: entry:1078// CHECK-NEXT: [[K_ADDR:%.*]] = alloca ptr, align 81079// CHECK-NEXT: store ptr [[K]], ptr [[K_ADDR]], align 81080// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[K_ADDR]], align 81081// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 81082// CHECK-NEXT: [[INCDEC_PTR:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP1]], i32 11083// CHECK-NEXT: store ptr [[INCDEC_PTR]], ptr [[TMP0]], align 81084// CHECK-NEXT: ret void1085//1086//1087// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l1251088// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[Z:%.*]]) #[[ATTR4]] {1089// CHECK-NEXT: entry:1090// CHECK-NEXT: [[Z_ADDR:%.*]] = alloca ptr, align 81091// CHECK-NEXT: [[TMP:%.*]] = alloca ptr, align 81092// CHECK-NEXT: store ptr [[Z]], ptr [[Z_ADDR]], align 81093// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[Z_ADDR]], align 81094// CHECK-NEXT: store ptr [[TMP0]], ptr [[TMP]], align 81095// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP]], align 81096// CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[TMP1]], align 81097// CHECK-NEXT: [[INCDEC_PTR:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP2]], i32 11098// CHECK-NEXT: store ptr [[INCDEC_PTR]], ptr [[TMP1]], align 81099// CHECK-NEXT: ret void1100//1101//1102// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l1271103// CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[AA:%.*]]) #[[ATTR4]] {1104// CHECK-NEXT: entry:1105// CHECK-NEXT: [[AA_ADDR:%.*]] = alloca ptr, align 81106// CHECK-NEXT: [[A:%.*]] = alloca i32, align 41107// CHECK-NEXT: store ptr [[AA]], ptr [[AA_ADDR]], align 81108// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 81109// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 01110// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[ARRAYIDX]], align 41111// CHECK-NEXT: store i32 [[TMP1]], ptr [[A]], align 41112// CHECK-NEXT: ret void1113//1114//1115// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l1291116// CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[H:%.*]]) #[[ATTR4]] {1117// CHECK-NEXT: entry:1118// CHECK-NEXT: [[H_ADDR:%.*]] = alloca ptr, align 81119// CHECK-NEXT: [[A:%.*]] = alloca i32, align 41120// CHECK-NEXT: store ptr [[H]], ptr [[H_ADDR]], align 81121// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[H_ADDR]], align 81122// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x %struct.S6], ptr [[TMP0]], i64 0, i64 01123// CHECK-NEXT: [[A1:%.*]] = getelementptr inbounds nuw [[STRUCT_S6:%.*]], ptr [[ARRAYIDX]], i32 0, i32 01124// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[A1]], align 41125// CHECK-NEXT: store i32 [[TMP1]], ptr [[A]], align 41126// CHECK-NEXT: ret void1127//1128//1129// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l1231130// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[K:%.*]]) #[[ATTR4]] {1131// CHECK-NEXT: entry:1132// CHECK-NEXT: [[K_ADDR:%.*]] = alloca ptr, align 81133// CHECK-NEXT: store ptr [[K]], ptr [[K_ADDR]], align 81134// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[K_ADDR]], align 81135// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 81136// CHECK-NEXT: [[INCDEC_PTR:%.*]] = getelementptr inbounds nuw ptr, ptr [[TMP1]], i32 11137// CHECK-NEXT: store ptr [[INCDEC_PTR]], ptr [[TMP0]], align 81138// CHECK-NEXT: ret void1139//1140//1141// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l1251142// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[Z:%.*]]) #[[ATTR4]] {1143// CHECK-NEXT: entry:1144// CHECK-NEXT: [[Z_ADDR:%.*]] = alloca ptr, align 81145// CHECK-NEXT: [[TMP:%.*]] = alloca ptr, align 81146// CHECK-NEXT: store ptr [[Z]], ptr [[Z_ADDR]], align 81147// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[Z_ADDR]], align 81148// CHECK-NEXT: store ptr [[TMP0]], ptr [[TMP]], align 81149// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP]], align 81150// CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[TMP1]], align 81151// CHECK-NEXT: [[INCDEC_PTR:%.*]] = getelementptr inbounds nuw ptr, ptr [[TMP2]], i32 11152// CHECK-NEXT: store ptr [[INCDEC_PTR]], ptr [[TMP1]], align 81153// CHECK-NEXT: ret void1154//1155//1156// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l1271157// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(80) [[AA:%.*]]) #[[ATTR4]] {1158// CHECK-NEXT: entry:1159// CHECK-NEXT: [[AA_ADDR:%.*]] = alloca ptr, align 81160// CHECK-NEXT: [[A:%.*]] = alloca ptr, align 81161// CHECK-NEXT: store ptr [[AA]], ptr [[AA_ADDR]], align 81162// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 81163// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x ptr], ptr [[TMP0]], i64 0, i64 01164// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[ARRAYIDX]], align 81165// CHECK-NEXT: store ptr [[TMP1]], ptr [[A]], align 81166// CHECK-NEXT: ret void1167//1168//1169// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l1291170// CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[H:%.*]]) #[[ATTR4]] {1171// CHECK-NEXT: entry:1172// CHECK-NEXT: [[H_ADDR:%.*]] = alloca ptr, align 81173// CHECK-NEXT: [[A:%.*]] = alloca i32, align 41174// CHECK-NEXT: store ptr [[H]], ptr [[H_ADDR]], align 81175// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[H_ADDR]], align 81176// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x %struct.S6], ptr [[TMP0]], i64 0, i64 01177// CHECK-NEXT: [[A1:%.*]] = getelementptr inbounds nuw [[STRUCT_S6:%.*]], ptr [[ARRAYIDX]], i32 0, i32 01178// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[A1]], align 41179// CHECK-NEXT: store i32 [[TMP1]], ptr [[A]], align 41180// CHECK-NEXT: ret void1181//1182//1183// CHECK-LABEL: define {{[^@]+}}@_Z12use_templatev1184// CHECK-SAME: () #[[ATTR1]] {1185// CHECK-NEXT: entry:1186// CHECK-NEXT: [[AKERN:%.*]] = alloca [[STRUCT_SOMEKERNEL:%.*]], align 41187// CHECK-NEXT: call void @_ZN10SomeKernelC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AKERN]])1188// CHECK-NEXT: call void @_ZN10SomeKernel5applyILj32EEEvv(ptr noundef nonnull align 4 dereferenceable(8) [[AKERN]])1189// CHECK-NEXT: call void @_ZN10SomeKernelD1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AKERN]]) #[[ATTR5]]1190// CHECK-NEXT: ret void1191//1192//1193// CHECK-LABEL: define {{[^@]+}}@_ZN10SomeKernel5applyILj32EEEvv1194// CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]]) #[[ATTR1]] comdat {1195// CHECK-NEXT: entry:1196// CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81197// CHECK-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41198// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 81199// CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 81200// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 81201// CHECK-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 81202// CHECK-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81203// CHECK-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81204// CHECK-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81205// CHECK-NEXT: [[TARGETDEV:%.*]] = getelementptr inbounds nuw [[STRUCT_SOMEKERNEL:%.*]], ptr [[THIS1]], i32 0, i32 01206// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[TARGETDEV]], align 41207// CHECK-NEXT: store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 41208// CHECK-NEXT: [[DEVPTR:%.*]] = getelementptr inbounds nuw [[STRUCT_SOMEKERNEL]], ptr [[THIS1]], i32 0, i32 11209// CHECK-NEXT: [[TARGETDEV2:%.*]] = getelementptr inbounds nuw [[STRUCT_SOMEKERNEL]], ptr [[THIS1]], i32 0, i32 01210// CHECK-NEXT: [[TMP1:%.*]] = getelementptr float, ptr [[DEVPTR]], i32 11211// CHECK-NEXT: [[TMP2:%.*]] = ptrtoint ptr [[TMP1]] to i641212// CHECK-NEXT: [[TMP3:%.*]] = ptrtoint ptr [[TARGETDEV2]] to i641213// CHECK-NEXT: [[TMP4:%.*]] = sub i64 [[TMP2]], [[TMP3]]1214// CHECK-NEXT: [[TMP5:%.*]] = sdiv exact i64 [[TMP4]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)1215// CHECK-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES]], ptr align 8 @.offload_sizes.31, i64 24, i1 false)1216// CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01217// CHECK-NEXT: store ptr [[THIS1]], ptr [[TMP6]], align 81218// CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01219// CHECK-NEXT: store ptr [[TARGETDEV2]], ptr [[TMP7]], align 81220// CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 01221// CHECK-NEXT: store i64 [[TMP5]], ptr [[TMP8]], align 81222// CHECK-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 01223// CHECK-NEXT: store ptr null, ptr [[TMP9]], align 81224// CHECK-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11225// CHECK-NEXT: store ptr [[THIS1]], ptr [[TMP10]], align 81226// CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11227// CHECK-NEXT: store ptr [[DEVPTR]], ptr [[TMP11]], align 81228// CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 11229// CHECK-NEXT: store ptr null, ptr [[TMP12]], align 81230// CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21231// CHECK-NEXT: store ptr [[THIS1]], ptr [[TMP13]], align 81232// CHECK-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21233// CHECK-NEXT: store ptr [[TARGETDEV2]], ptr [[TMP14]], align 81234// CHECK-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 21235// CHECK-NEXT: store ptr null, ptr [[TMP15]], align 81236// CHECK-NEXT: [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01237// CHECK-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01238// CHECK-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 01239// CHECK-NEXT: [[TMP19:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41240// CHECK-NEXT: [[TMP20:%.*]] = sext i32 [[TMP19]] to i641241// CHECK-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01242// CHECK-NEXT: store i32 3, ptr [[TMP21]], align 41243// CHECK-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11244// CHECK-NEXT: store i32 3, ptr [[TMP22]], align 41245// CHECK-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21246// CHECK-NEXT: store ptr [[TMP16]], ptr [[TMP23]], align 81247// CHECK-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31248// CHECK-NEXT: store ptr [[TMP17]], ptr [[TMP24]], align 81249// CHECK-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41250// CHECK-NEXT: store ptr [[TMP18]], ptr [[TMP25]], align 81251// CHECK-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51252// CHECK-NEXT: store ptr @.offload_maptypes.32, ptr [[TMP26]], align 81253// CHECK-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61254// CHECK-NEXT: store ptr null, ptr [[TMP27]], align 81255// CHECK-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71256// CHECK-NEXT: store ptr null, ptr [[TMP28]], align 81257// CHECK-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81258// CHECK-NEXT: store i64 0, ptr [[TMP29]], align 81259// CHECK-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91260// CHECK-NEXT: store i64 0, ptr [[TMP30]], align 81261// CHECK-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101262// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP31]], align 41263// CHECK-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111264// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP32]], align 41265// CHECK-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121266// CHECK-NEXT: store i32 0, ptr [[TMP33]], align 41267// CHECK-NEXT: [[TMP34:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 [[TMP20]], i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN10SomeKernel5applyILj32EEEvv_l168.region_id, ptr [[KERNEL_ARGS]])1268// CHECK-NEXT: [[TMP35:%.*]] = icmp ne i32 [[TMP34]], 01269// CHECK-NEXT: br i1 [[TMP35]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1270// CHECK: omp_offload.failed:1271// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN10SomeKernel5applyILj32EEEvv_l168(ptr [[THIS1]]) #[[ATTR5]]1272// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT]]1273// CHECK: omp_offload.cont:1274// CHECK-NEXT: ret void1275//1276//1277// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN10SomeKernel5applyILj32EEEvv_l1681278// CHECK-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR4]] {1279// CHECK-NEXT: entry:1280// CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81281// CHECK-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81282// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81283// CHECK-NEXT: [[DEVPTR:%.*]] = getelementptr inbounds nuw [[STRUCT_SOMEKERNEL:%.*]], ptr [[TMP0]], i32 0, i32 11284// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr [[DEVPTR]], align 41285// CHECK-NEXT: [[INC:%.*]] = fadd float [[TMP1]], 1.000000e+001286// CHECK-NEXT: store float [[INC]], ptr [[DEVPTR]], align 41287// CHECK-NEXT: [[TARGETDEV:%.*]] = getelementptr inbounds nuw [[STRUCT_SOMEKERNEL]], ptr [[TMP0]], i32 0, i32 01288// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[TARGETDEV]], align 41289// CHECK-NEXT: [[INC1:%.*]] = add nsw i32 [[TMP2]], 11290// CHECK-NEXT: store i32 [[INC1]], ptr [[TARGETDEV]], align 41291// CHECK-NEXT: ret void1292//1293//1294// CHECK-LABEL: define {{[^@]+}}@_GLOBAL__sub_I_target_has_device_addr_codegen.cpp1295// CHECK-SAME: () #[[ATTR0]] {1296// CHECK-NEXT: entry:1297// CHECK-NEXT: call void @__cxx_global_var_init()1298// CHECK-NEXT: call void @__cxx_global_var_init.1()1299// CHECK-NEXT: call void @__cxx_global_var_init.2()1300// CHECK-NEXT: call void @__cxx_global_var_init.3()1301// CHECK-NEXT: ret void1302//1303//1304// CHECK-LABEL: define {{[^@]+}}@__tls_init1305// CHECK-SAME: () #[[ATTR0]] {1306// CHECK-NEXT: entry:1307// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr @__tls_guard, align 11308// CHECK-NEXT: [[GUARD_UNINITIALIZED:%.*]] = icmp eq i8 [[TMP0]], 01309// CHECK-NEXT: br i1 [[GUARD_UNINITIALIZED]], label [[INIT:%.*]], label [[EXIT:%.*]], !prof [[PROF18:![0-9]+]]1310// CHECK: init:1311// CHECK-NEXT: store i8 1, ptr @__tls_guard, align 11312// CHECK-NEXT: call void @__cxx_global_var_init.4()1313// CHECK-NEXT: br label [[EXIT]]1314// CHECK: exit:1315// CHECK-NEXT: ret void1316//1317//1318// CHECK-LABEL: define {{[^@]+}}@_ZTW1h1319// CHECK-SAME: () #[[ATTR0]] comdat {1320// CHECK-NEXT: call void @_ZTH1h()1321// CHECK-NEXT: [[TMP1:%.*]] = call align 4 ptr @llvm.threadlocal.address.p0(ptr align 4 @h)1322// CHECK-NEXT: ret ptr [[TMP1]]1323//1324//1325// SIMD-ONLY0-LABEL: define {{[^@]+}}@__cxx_global_var_init1326// SIMD-ONLY0-SAME: () #[[ATTR0:[0-9]+]] {1327// SIMD-ONLY0-NEXT: entry:1328// SIMD-ONLY0-NEXT: call void @_ZN2S2C1Ev(ptr noundef nonnull align 4 dereferenceable(4) @_ZL1b)1329// SIMD-ONLY0-NEXT: ret void1330//1331//1332// SIMD-ONLY0-LABEL: define {{[^@]+}}@_ZN2S2C1Ev1333// SIMD-ONLY0-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1:[0-9]+]] comdat {1334// SIMD-ONLY0-NEXT: entry:1335// SIMD-ONLY0-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81336// SIMD-ONLY0-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81337// SIMD-ONLY0-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81338// SIMD-ONLY0-NEXT: call void @_ZN2S2C2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]])1339// SIMD-ONLY0-NEXT: ret void1340//1341//1342// SIMD-ONLY0-LABEL: define {{[^@]+}}@__cxx_global_var_init.11343// SIMD-ONLY0-SAME: () #[[ATTR0]] {1344// SIMD-ONLY0-NEXT: entry:1345// SIMD-ONLY0-NEXT: br label [[ARRAYCTOR_LOOP:%.*]]1346// SIMD-ONLY0: arrayctor.loop:1347// SIMD-ONLY0-NEXT: [[ARRAYCTOR_CUR:%.*]] = phi ptr [ @_ZL2ba, [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ]1348// SIMD-ONLY0-NEXT: call void @_ZN2S2C1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]])1349// SIMD-ONLY0-NEXT: [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[CLASS_S2:%.*]], ptr [[ARRAYCTOR_CUR]], i64 11350// SIMD-ONLY0-NEXT: [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], getelementptr inbounds ([[CLASS_S2]], ptr @_ZL2ba, i64 5)1351// SIMD-ONLY0-NEXT: br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]]1352// SIMD-ONLY0: arrayctor.cont:1353// SIMD-ONLY0-NEXT: ret void1354//1355//1356// SIMD-ONLY0-LABEL: define {{[^@]+}}@__cxx_global_var_init.21357// SIMD-ONLY0-SAME: () #[[ATTR0]] {1358// SIMD-ONLY0-NEXT: entry:1359// SIMD-ONLY0-NEXT: call void @_ZN2S3C1Ev(ptr noundef nonnull align 4 dereferenceable(4) @_ZL1c)1360// SIMD-ONLY0-NEXT: ret void1361//1362//1363// SIMD-ONLY0-LABEL: define {{[^@]+}}@_ZN2S3C1Ev1364// SIMD-ONLY0-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {1365// SIMD-ONLY0-NEXT: entry:1366// SIMD-ONLY0-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81367// SIMD-ONLY0-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81368// SIMD-ONLY0-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81369// SIMD-ONLY0-NEXT: call void @_ZN2S3C2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]])1370// SIMD-ONLY0-NEXT: ret void1371//1372//1373// SIMD-ONLY0-LABEL: define {{[^@]+}}@__cxx_global_var_init.31374// SIMD-ONLY0-SAME: () #[[ATTR0]] {1375// SIMD-ONLY0-NEXT: entry:1376// SIMD-ONLY0-NEXT: br label [[ARRAYCTOR_LOOP:%.*]]1377// SIMD-ONLY0: arrayctor.loop:1378// SIMD-ONLY0-NEXT: [[ARRAYCTOR_CUR:%.*]] = phi ptr [ @_ZL2ca, [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ]1379// SIMD-ONLY0-NEXT: call void @_ZN2S3C1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]])1380// SIMD-ONLY0-NEXT: [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[CLASS_S3:%.*]], ptr [[ARRAYCTOR_CUR]], i64 11381// SIMD-ONLY0-NEXT: [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], getelementptr inbounds ([[CLASS_S3]], ptr @_ZL2ca, i64 5)1382// SIMD-ONLY0-NEXT: br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]]1383// SIMD-ONLY0: arrayctor.cont:1384// SIMD-ONLY0-NEXT: ret void1385//1386//1387// SIMD-ONLY0-LABEL: define {{[^@]+}}@__cxx_global_var_init.41388// SIMD-ONLY0-SAME: () #[[ATTR0]] {1389// SIMD-ONLY0-NEXT: entry:1390// SIMD-ONLY0-NEXT: call void @_ZN2S3C1Ev(ptr noundef nonnull align 4 dereferenceable(4) @h)1391// SIMD-ONLY0-NEXT: ret void1392//1393//1394// SIMD-ONLY0-LABEL: define {{[^@]+}}@main1395// SIMD-ONLY0-SAME: (i32 noundef signext [[ARGC:%.*]], ptr noundef [[ARGV:%.*]]) #[[ATTR2:[0-9]+]] {1396// SIMD-ONLY0-NEXT: entry:1397// SIMD-ONLY0-NEXT: [[RETVAL:%.*]] = alloca i32, align 41398// SIMD-ONLY0-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 41399// SIMD-ONLY0-NEXT: [[ARGV_ADDR:%.*]] = alloca ptr, align 81400// SIMD-ONLY0-NEXT: [[DA:%.*]] = alloca [5 x i32], align 41401// SIMD-ONLY0-NEXT: [[H:%.*]] = alloca [10 x %struct.S6], align 41402// SIMD-ONLY0-NEXT: [[RH:%.*]] = alloca ptr, align 81403// SIMD-ONLY0-NEXT: [[I:%.*]] = alloca i32, align 41404// SIMD-ONLY0-NEXT: [[J:%.*]] = alloca ptr, align 81405// SIMD-ONLY0-NEXT: [[K:%.*]] = alloca ptr, align 81406// SIMD-ONLY0-NEXT: [[Z:%.*]] = alloca ptr, align 81407// SIMD-ONLY0-NEXT: [[AA:%.*]] = alloca [10 x i32], align 41408// SIMD-ONLY0-NEXT: [[RAA:%.*]] = alloca ptr, align 81409// SIMD-ONLY0-NEXT: [[TMP:%.*]] = alloca ptr, align 81410// SIMD-ONLY0-NEXT: [[_TMP2:%.*]] = alloca ptr, align 81411// SIMD-ONLY0-NEXT: [[A:%.*]] = alloca i32, align 41412// SIMD-ONLY0-NEXT: [[A4:%.*]] = alloca i32, align 41413// SIMD-ONLY0-NEXT: [[A7:%.*]] = alloca i32, align 41414// SIMD-ONLY0-NEXT: store i32 0, ptr [[RETVAL]], align 41415// SIMD-ONLY0-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 41416// SIMD-ONLY0-NEXT: store ptr [[ARGV]], ptr [[ARGV_ADDR]], align 81417// SIMD-ONLY0-NEXT: call void @llvm.memset.p0.i64(ptr align 4 [[DA]], i8 0, i64 20, i1 false)1418// SIMD-ONLY0-NEXT: store ptr [[H]], ptr [[RH]], align 81419// SIMD-ONLY0-NEXT: store ptr [[I]], ptr [[J]], align 81420// SIMD-ONLY0-NEXT: [[TMP0:%.*]] = load ptr, ptr [[J]], align 81421// SIMD-ONLY0-NEXT: store ptr [[TMP0]], ptr [[K]], align 81422// SIMD-ONLY0-NEXT: store ptr [[K]], ptr [[Z]], align 81423// SIMD-ONLY0-NEXT: store ptr [[AA]], ptr [[RAA]], align 81424// SIMD-ONLY0-NEXT: [[TMP1:%.*]] = load ptr, ptr [[K]], align 81425// SIMD-ONLY0-NEXT: [[INCDEC_PTR:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP1]], i32 11426// SIMD-ONLY0-NEXT: store ptr [[INCDEC_PTR]], ptr [[K]], align 81427// SIMD-ONLY0-NEXT: [[TMP2:%.*]] = load ptr, ptr [[Z]], align 81428// SIMD-ONLY0-NEXT: store ptr [[TMP2]], ptr [[TMP]], align 81429// SIMD-ONLY0-NEXT: [[TMP3:%.*]] = load ptr, ptr [[Z]], align 81430// SIMD-ONLY0-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP]], align 81431// SIMD-ONLY0-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 81432// SIMD-ONLY0-NEXT: [[INCDEC_PTR1:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP5]], i32 11433// SIMD-ONLY0-NEXT: store ptr [[INCDEC_PTR1]], ptr [[TMP4]], align 81434// SIMD-ONLY0-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[AA]], i64 0, i64 01435// SIMD-ONLY0-NEXT: store i32 1, ptr [[ARRAYIDX]], align 41436// SIMD-ONLY0-NEXT: [[TMP6:%.*]] = load ptr, ptr [[RAA]], align 81437// SIMD-ONLY0-NEXT: store ptr [[TMP6]], ptr [[_TMP2]], align 81438// SIMD-ONLY0-NEXT: [[TMP7:%.*]] = load ptr, ptr [[RAA]], align 81439// SIMD-ONLY0-NEXT: [[TMP8:%.*]] = load ptr, ptr [[_TMP2]], align 81440// SIMD-ONLY0-NEXT: [[ARRAYIDX3:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP8]], i64 0, i64 01441// SIMD-ONLY0-NEXT: [[TMP9:%.*]] = load i32, ptr [[ARRAYIDX3]], align 41442// SIMD-ONLY0-NEXT: store i32 [[TMP9]], ptr [[A]], align 41443// SIMD-ONLY0-NEXT: [[ARRAYIDX5:%.*]] = getelementptr inbounds [10 x %struct.S6], ptr [[H]], i64 0, i64 11444// SIMD-ONLY0-NEXT: [[A6:%.*]] = getelementptr inbounds nuw [[STRUCT_S6:%.*]], ptr [[ARRAYIDX5]], i32 0, i32 01445// SIMD-ONLY0-NEXT: [[TMP10:%.*]] = load i32, ptr [[A6]], align 41446// SIMD-ONLY0-NEXT: store i32 [[TMP10]], ptr [[A4]], align 41447// SIMD-ONLY0-NEXT: [[ARRAYIDX8:%.*]] = getelementptr inbounds [5 x i32], ptr [[DA]], i64 0, i64 11448// SIMD-ONLY0-NEXT: [[TMP11:%.*]] = load i32, ptr [[ARRAYIDX8]], align 41449// SIMD-ONLY0-NEXT: store i32 [[TMP11]], ptr [[A7]], align 41450// SIMD-ONLY0-NEXT: [[TMP12:%.*]] = load i32, ptr [[ARGC_ADDR]], align 41451// SIMD-ONLY0-NEXT: [[CALL:%.*]] = call noundef signext i32 @_Z5tmainIiET_S0_(i32 noundef signext [[TMP12]])1452// SIMD-ONLY0-NEXT: [[CALL9:%.*]] = call noundef ptr @_Z5tmainIPiET_S1_(ptr noundef [[ARGC_ADDR]])1453// SIMD-ONLY0-NEXT: [[TMP13:%.*]] = load i32, ptr [[CALL9]], align 41454// SIMD-ONLY0-NEXT: [[ADD:%.*]] = add nsw i32 [[CALL]], [[TMP13]]1455// SIMD-ONLY0-NEXT: ret i32 [[ADD]]1456//1457//1458// SIMD-ONLY0-LABEL: define {{[^@]+}}@_Z5tmainIiET_S0_1459// SIMD-ONLY0-SAME: (i32 noundef signext [[ARGC:%.*]]) #[[ATTR1]] comdat {1460// SIMD-ONLY0-NEXT: entry:1461// SIMD-ONLY0-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 41462// SIMD-ONLY0-NEXT: [[DA:%.*]] = alloca [5 x i32], align 41463// SIMD-ONLY0-NEXT: [[H:%.*]] = alloca [10 x %struct.S6], align 41464// SIMD-ONLY0-NEXT: [[RH:%.*]] = alloca ptr, align 81465// SIMD-ONLY0-NEXT: [[I:%.*]] = alloca i32, align 41466// SIMD-ONLY0-NEXT: [[J:%.*]] = alloca ptr, align 81467// SIMD-ONLY0-NEXT: [[K:%.*]] = alloca ptr, align 81468// SIMD-ONLY0-NEXT: [[Z:%.*]] = alloca ptr, align 81469// SIMD-ONLY0-NEXT: [[AA:%.*]] = alloca [10 x i32], align 41470// SIMD-ONLY0-NEXT: [[TMP:%.*]] = alloca ptr, align 81471// SIMD-ONLY0-NEXT: [[A:%.*]] = alloca i32, align 41472// SIMD-ONLY0-NEXT: [[A2:%.*]] = alloca i32, align 41473// SIMD-ONLY0-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 41474// SIMD-ONLY0-NEXT: call void @llvm.memset.p0.i64(ptr align 4 [[DA]], i8 0, i64 20, i1 false)1475// SIMD-ONLY0-NEXT: store ptr [[H]], ptr [[RH]], align 81476// SIMD-ONLY0-NEXT: store ptr [[I]], ptr [[J]], align 81477// SIMD-ONLY0-NEXT: [[TMP0:%.*]] = load ptr, ptr [[J]], align 81478// SIMD-ONLY0-NEXT: store ptr [[TMP0]], ptr [[K]], align 81479// SIMD-ONLY0-NEXT: store ptr [[K]], ptr [[Z]], align 81480// SIMD-ONLY0-NEXT: [[TMP1:%.*]] = load ptr, ptr [[K]], align 81481// SIMD-ONLY0-NEXT: [[INCDEC_PTR:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP1]], i32 11482// SIMD-ONLY0-NEXT: store ptr [[INCDEC_PTR]], ptr [[K]], align 81483// SIMD-ONLY0-NEXT: [[TMP2:%.*]] = load ptr, ptr [[Z]], align 81484// SIMD-ONLY0-NEXT: store ptr [[TMP2]], ptr [[TMP]], align 81485// SIMD-ONLY0-NEXT: [[TMP3:%.*]] = load ptr, ptr [[Z]], align 81486// SIMD-ONLY0-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP]], align 81487// SIMD-ONLY0-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 81488// SIMD-ONLY0-NEXT: [[INCDEC_PTR1:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP5]], i32 11489// SIMD-ONLY0-NEXT: store ptr [[INCDEC_PTR1]], ptr [[TMP4]], align 81490// SIMD-ONLY0-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[AA]], i64 0, i64 01491// SIMD-ONLY0-NEXT: [[TMP6:%.*]] = load i32, ptr [[ARRAYIDX]], align 41492// SIMD-ONLY0-NEXT: store i32 [[TMP6]], ptr [[A]], align 41493// SIMD-ONLY0-NEXT: [[ARRAYIDX3:%.*]] = getelementptr inbounds [10 x %struct.S6], ptr [[H]], i64 0, i64 01494// SIMD-ONLY0-NEXT: [[A4:%.*]] = getelementptr inbounds nuw [[STRUCT_S6:%.*]], ptr [[ARRAYIDX3]], i32 0, i32 01495// SIMD-ONLY0-NEXT: [[TMP7:%.*]] = load i32, ptr [[A4]], align 41496// SIMD-ONLY0-NEXT: store i32 [[TMP7]], ptr [[A2]], align 41497// SIMD-ONLY0-NEXT: ret i32 01498//1499//1500// SIMD-ONLY0-LABEL: define {{[^@]+}}@_Z5tmainIPiET_S1_1501// SIMD-ONLY0-SAME: (ptr noundef [[ARGC:%.*]]) #[[ATTR1]] comdat {1502// SIMD-ONLY0-NEXT: entry:1503// SIMD-ONLY0-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 81504// SIMD-ONLY0-NEXT: [[DA:%.*]] = alloca [5 x ptr], align 81505// SIMD-ONLY0-NEXT: [[H:%.*]] = alloca [10 x %struct.S6], align 41506// SIMD-ONLY0-NEXT: [[RH:%.*]] = alloca ptr, align 81507// SIMD-ONLY0-NEXT: [[I:%.*]] = alloca ptr, align 81508// SIMD-ONLY0-NEXT: [[J:%.*]] = alloca ptr, align 81509// SIMD-ONLY0-NEXT: [[K:%.*]] = alloca ptr, align 81510// SIMD-ONLY0-NEXT: [[Z:%.*]] = alloca ptr, align 81511// SIMD-ONLY0-NEXT: [[AA:%.*]] = alloca [10 x ptr], align 81512// SIMD-ONLY0-NEXT: [[TMP:%.*]] = alloca ptr, align 81513// SIMD-ONLY0-NEXT: [[A:%.*]] = alloca ptr, align 81514// SIMD-ONLY0-NEXT: [[A2:%.*]] = alloca i32, align 41515// SIMD-ONLY0-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 81516// SIMD-ONLY0-NEXT: call void @llvm.memset.p0.i64(ptr align 8 [[DA]], i8 0, i64 40, i1 false)1517// SIMD-ONLY0-NEXT: store ptr [[H]], ptr [[RH]], align 81518// SIMD-ONLY0-NEXT: store ptr [[I]], ptr [[J]], align 81519// SIMD-ONLY0-NEXT: [[TMP0:%.*]] = load ptr, ptr [[J]], align 81520// SIMD-ONLY0-NEXT: store ptr [[TMP0]], ptr [[K]], align 81521// SIMD-ONLY0-NEXT: store ptr [[K]], ptr [[Z]], align 81522// SIMD-ONLY0-NEXT: [[TMP1:%.*]] = load ptr, ptr [[K]], align 81523// SIMD-ONLY0-NEXT: [[INCDEC_PTR:%.*]] = getelementptr inbounds nuw ptr, ptr [[TMP1]], i32 11524// SIMD-ONLY0-NEXT: store ptr [[INCDEC_PTR]], ptr [[K]], align 81525// SIMD-ONLY0-NEXT: [[TMP2:%.*]] = load ptr, ptr [[Z]], align 81526// SIMD-ONLY0-NEXT: store ptr [[TMP2]], ptr [[TMP]], align 81527// SIMD-ONLY0-NEXT: [[TMP3:%.*]] = load ptr, ptr [[Z]], align 81528// SIMD-ONLY0-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP]], align 81529// SIMD-ONLY0-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 81530// SIMD-ONLY0-NEXT: [[INCDEC_PTR1:%.*]] = getelementptr inbounds nuw ptr, ptr [[TMP5]], i32 11531// SIMD-ONLY0-NEXT: store ptr [[INCDEC_PTR1]], ptr [[TMP4]], align 81532// SIMD-ONLY0-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x ptr], ptr [[AA]], i64 0, i64 01533// SIMD-ONLY0-NEXT: [[TMP6:%.*]] = load ptr, ptr [[ARRAYIDX]], align 81534// SIMD-ONLY0-NEXT: store ptr [[TMP6]], ptr [[A]], align 81535// SIMD-ONLY0-NEXT: [[ARRAYIDX3:%.*]] = getelementptr inbounds [10 x %struct.S6], ptr [[H]], i64 0, i64 01536// SIMD-ONLY0-NEXT: [[A4:%.*]] = getelementptr inbounds nuw [[STRUCT_S6:%.*]], ptr [[ARRAYIDX3]], i32 0, i32 01537// SIMD-ONLY0-NEXT: [[TMP7:%.*]] = load i32, ptr [[A4]], align 41538// SIMD-ONLY0-NEXT: store i32 [[TMP7]], ptr [[A2]], align 41539// SIMD-ONLY0-NEXT: ret ptr null1540//1541//1542// SIMD-ONLY0-LABEL: define {{[^@]+}}@_Z12use_templatev1543// SIMD-ONLY0-SAME: () #[[ATTR1]] {1544// SIMD-ONLY0-NEXT: entry:1545// SIMD-ONLY0-NEXT: [[AKERN:%.*]] = alloca [[STRUCT_SOMEKERNEL:%.*]], align 41546// SIMD-ONLY0-NEXT: call void @_ZN10SomeKernelC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AKERN]])1547// SIMD-ONLY0-NEXT: call void @_ZN10SomeKernel5applyILj32EEEvv(ptr noundef nonnull align 4 dereferenceable(8) [[AKERN]])1548// SIMD-ONLY0-NEXT: call void @_ZN10SomeKernelD1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AKERN]]) #[[ATTR6:[0-9]+]]1549// SIMD-ONLY0-NEXT: ret void1550//1551//1552// SIMD-ONLY0-LABEL: define {{[^@]+}}@_ZN10SomeKernel5applyILj32EEEvv1553// SIMD-ONLY0-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]]) #[[ATTR1]] comdat {1554// SIMD-ONLY0-NEXT: entry:1555// SIMD-ONLY0-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81556// SIMD-ONLY0-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41557// SIMD-ONLY0-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81558// SIMD-ONLY0-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81559// SIMD-ONLY0-NEXT: [[TARGETDEV:%.*]] = getelementptr inbounds nuw [[STRUCT_SOMEKERNEL:%.*]], ptr [[THIS1]], i32 0, i32 01560// SIMD-ONLY0-NEXT: [[TMP0:%.*]] = load i32, ptr [[TARGETDEV]], align 41561// SIMD-ONLY0-NEXT: store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 41562// SIMD-ONLY0-NEXT: [[DEVPTR:%.*]] = getelementptr inbounds nuw [[STRUCT_SOMEKERNEL]], ptr [[THIS1]], i32 0, i32 11563// SIMD-ONLY0-NEXT: [[TMP1:%.*]] = load float, ptr [[DEVPTR]], align 41564// SIMD-ONLY0-NEXT: [[INC:%.*]] = fadd float [[TMP1]], 1.000000e+001565// SIMD-ONLY0-NEXT: store float [[INC]], ptr [[DEVPTR]], align 41566// SIMD-ONLY0-NEXT: [[TARGETDEV2:%.*]] = getelementptr inbounds nuw [[STRUCT_SOMEKERNEL]], ptr [[THIS1]], i32 0, i32 01567// SIMD-ONLY0-NEXT: [[TMP2:%.*]] = load i32, ptr [[TARGETDEV2]], align 41568// SIMD-ONLY0-NEXT: [[INC3:%.*]] = add nsw i32 [[TMP2]], 11569// SIMD-ONLY0-NEXT: store i32 [[INC3]], ptr [[TARGETDEV2]], align 41570// SIMD-ONLY0-NEXT: ret void1571//1572//1573// SIMD-ONLY0-LABEL: define {{[^@]+}}@_ZN2S2C2Ev1574// SIMD-ONLY0-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {1575// SIMD-ONLY0-NEXT: entry:1576// SIMD-ONLY0-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81577// SIMD-ONLY0-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81578// SIMD-ONLY0-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81579// SIMD-ONLY0-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[CLASS_S2:%.*]], ptr [[THIS1]], i32 0, i32 01580// SIMD-ONLY0-NEXT: store i32 0, ptr [[A]], align 41581// SIMD-ONLY0-NEXT: ret void1582//1583//1584// SIMD-ONLY0-LABEL: define {{[^@]+}}@_ZN2S3C2Ev1585// SIMD-ONLY0-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {1586// SIMD-ONLY0-NEXT: entry:1587// SIMD-ONLY0-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81588// SIMD-ONLY0-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81589// SIMD-ONLY0-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81590// SIMD-ONLY0-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[CLASS_S3:%.*]], ptr [[THIS1]], i32 0, i32 01591// SIMD-ONLY0-NEXT: store i32 0, ptr [[A]], align 41592// SIMD-ONLY0-NEXT: ret void1593//1594//1595// SIMD-ONLY0-LABEL: define {{[^@]+}}@_GLOBAL__sub_I_target_has_device_addr_codegen.cpp1596// SIMD-ONLY0-SAME: () #[[ATTR0]] {1597// SIMD-ONLY0-NEXT: entry:1598// SIMD-ONLY0-NEXT: call void @__cxx_global_var_init()1599// SIMD-ONLY0-NEXT: call void @__cxx_global_var_init.1()1600// SIMD-ONLY0-NEXT: call void @__cxx_global_var_init.2()1601// SIMD-ONLY0-NEXT: call void @__cxx_global_var_init.3()1602// SIMD-ONLY0-NEXT: call void @__cxx_global_var_init.4()1603// SIMD-ONLY0-NEXT: ret void1604//1605