1713 lines · cpp
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _2// RUN: %clang_cc1 -DLAMBDA -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK13// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s4// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK15// RUN: %clang_cc1 -DLAMBDA -verify -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK36// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s7// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK38 9// RUN: %clang_cc1 -DLAMBDA -verify -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"10// RUN: %clang_cc1 -DLAMBDA -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s11// RUN: %clang_cc1 -DLAMBDA -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"12// RUN: %clang_cc1 -DLAMBDA -verify -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"13// RUN: %clang_cc1 -DLAMBDA -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s14// RUN: %clang_cc1 -DLAMBDA -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"15 16// RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK917// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s18// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK919// RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK1120// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s21// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK1122 23// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"24// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s25// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"26// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"27// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s28// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"29// expected-no-diagnostics30#ifndef HEADER31#define HEADER32template <class T>33struct S {34 T f;35 S(T a) : f(a) {}36 S() : f() {}37 operator T() { return T(); }38 ~S() {}39};40 41volatile int g __attribute__((aligned(128))) = 1212;42 43struct SS {44 int a;45 int b : 4;46 int &c;47 SS(int &d) : a(0), b(0), c(d) {48#pragma omp target49#pragma omp teams private(a, b, c)50#ifdef LAMBDA51 [&]() {52 ++this->a, --b, (this)->c /= 1;53 }();54#else55 ++this->a, --b, c /= 1;56#endif57 }58};59 60template<typename T>61struct SST {62 T a;63 SST() : a(T()) {64#pragma omp target65#pragma omp teams private(a)66#ifdef LAMBDA67 [&]() {68 [&]() {69 ++this->a;70 }();71 }();72#else73 ++(this)->a;74#endif75 }76};77 78template <typename T>79T tmain() {80 S<T> test;81 SST<T> sst;82 T t_var __attribute__((aligned(128))) = T();83 T vec[] __attribute__((aligned(128))) = {1, 2};84 S<T> s_arr[] __attribute__((aligned(128))) = {1, 2};85 S<T> var __attribute__((aligned(128))) (3);86#pragma omp target87#pragma omp teams private(t_var, vec, s_arr, var)88 {89 vec[0] = t_var;90 s_arr[0] = var;91 }92 return T();93}94 95int main() {96 static int sivar;97 SS ss(sivar);98#ifdef LAMBDA99 100 101 // lambda and target region in main102 103 // target region in struct constructor104 105 // offloading function in struct constructor106 107 // outlined teams region in struct constructor108 // call void [[INNER_LAMBDA_CONSTR:@.+]](ptr109 110 // inner lambda in struct constructor111 // define{{.*}} void [[INNER_LAMBDA_CONSTR]](ptr112 113 114 // ret115 116 [&]() {117#pragma omp target118#pragma omp teams private(g, sivar)119 {120 121 g = 1;122 sivar = 2;123 [&]() {124 g = 2;125 sivar = 4;126 }();127 }128 }();129 return 0;130#else131 S<float> test;132 int t_var = 0;133 int vec[] = {1, 2};134 S<float> s_arr[] = {1, 2};135 S<float> var(3);136#pragma omp target137#pragma omp teams private(t_var, vec, s_arr, var, sivar)138 {139 vec[0] = t_var;140 s_arr[0] = var;141 sivar = 3;142 }143 return tmain<int>();144#endif145}146 147 148// target region in main function149 150 151// template tmain152 153// target in SS constructor154 155 156// target in tmain template157 158 159// SST constructor160 161// target in SST constructor162 163 164#endif165 166// CHECK1-LABEL: define {{[^@]+}}@main167// CHECK1-SAME: () #[[ATTR0:[0-9]+]] {168// CHECK1-NEXT: entry:169// CHECK1-NEXT: [[RETVAL:%.*]] = alloca i32, align 4170// CHECK1-NEXT: [[SS:%.*]] = alloca [[STRUCT_SS:%.*]], align 8171// CHECK1-NEXT: [[REF_TMP:%.*]] = alloca [[CLASS_ANON:%.*]], align 1172// CHECK1-NEXT: store i32 0, ptr [[RETVAL]], align 4173// CHECK1-NEXT: call void @_ZN2SSC1ERi(ptr noundef nonnull align 8 dereferenceable(16) [[SS]], ptr noundef nonnull align 4 dereferenceable(4) @_ZZ4mainE5sivar)174// CHECK1-NEXT: call void @"_ZZ4mainENK3$_0clEv"(ptr noundef nonnull align 1 dereferenceable(1) [[REF_TMP]])175// CHECK1-NEXT: ret i32 0176//177//178// CHECK1-LABEL: define {{[^@]+}}@_ZN2SSC1ERi179// CHECK1-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[THIS:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[D:%.*]]) unnamed_addr #[[ATTR1:[0-9]+]] comdat {180// CHECK1-NEXT: entry:181// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8182// CHECK1-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 8183// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8184// CHECK1-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 8185// CHECK1-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8186// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[D_ADDR]], align 8187// CHECK1-NEXT: call void @_ZN2SSC2ERi(ptr noundef nonnull align 8 dereferenceable(16) [[THIS1]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP0]])188// CHECK1-NEXT: ret void189//190//191// CHECK1-LABEL: define {{[^@]+}}@_ZN2SSC2ERi192// CHECK1-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[THIS:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[D:%.*]]) unnamed_addr #[[ATTR1]] comdat {193// CHECK1-NEXT: entry:194// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8195// CHECK1-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 8196// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8197// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8198// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8199// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8200// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8201// CHECK1-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 8202// CHECK1-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8203// CHECK1-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[THIS1]], i32 0, i32 0204// CHECK1-NEXT: store i32 0, ptr [[A]], align 8205// CHECK1-NEXT: [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 1206// CHECK1-NEXT: [[BF_LOAD:%.*]] = load i8, ptr [[B]], align 4207// CHECK1-NEXT: [[BF_CLEAR:%.*]] = and i8 [[BF_LOAD]], -16208// CHECK1-NEXT: [[BF_SET:%.*]] = or i8 [[BF_CLEAR]], 0209// CHECK1-NEXT: store i8 [[BF_SET]], ptr [[B]], align 4210// CHECK1-NEXT: [[C:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 2211// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[D_ADDR]], align 8212// CHECK1-NEXT: store ptr [[TMP0]], ptr [[C]], align 8213// CHECK1-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0214// CHECK1-NEXT: store ptr [[THIS1]], ptr [[TMP1]], align 8215// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0216// CHECK1-NEXT: store ptr [[THIS1]], ptr [[TMP2]], align 8217// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0218// CHECK1-NEXT: store ptr null, ptr [[TMP3]], align 8219// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0220// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0221// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0222// CHECK1-NEXT: store i32 3, ptr [[TMP6]], align 4223// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1224// CHECK1-NEXT: store i32 1, ptr [[TMP7]], align 4225// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2226// CHECK1-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 8227// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3228// CHECK1-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 8229// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4230// CHECK1-NEXT: store ptr @.offload_sizes, ptr [[TMP10]], align 8231// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5232// CHECK1-NEXT: store ptr @.offload_maptypes, ptr [[TMP11]], align 8233// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6234// CHECK1-NEXT: store ptr null, ptr [[TMP12]], align 8235// CHECK1-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7236// CHECK1-NEXT: store ptr null, ptr [[TMP13]], align 8237// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8238// CHECK1-NEXT: store i64 0, ptr [[TMP14]], align 8239// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9240// CHECK1-NEXT: store i64 0, ptr [[TMP15]], align 8241// CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10242// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4243// CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11244// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4245// CHECK1-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12246// CHECK1-NEXT: store i32 0, ptr [[TMP18]], align 4247// CHECK1-NEXT: [[TMP19:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSC1ERi_l48.region_id, ptr [[KERNEL_ARGS]])248// CHECK1-NEXT: [[TMP20:%.*]] = icmp ne i32 [[TMP19]], 0249// CHECK1-NEXT: br i1 [[TMP20]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]250// CHECK1: omp_offload.failed:251// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSC1ERi_l48(ptr [[THIS1]]) #[[ATTR3:[0-9]+]]252// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]253// CHECK1: omp_offload.cont:254// CHECK1-NEXT: ret void255//256//257// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSC1ERi_l48258// CHECK1-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR2:[0-9]+]] {259// CHECK1-NEXT: entry:260// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8261// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8262// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8263// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSC1ERi_l48.omp_outlined, ptr [[TMP0]])264// CHECK1-NEXT: ret void265//266//267// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSC1ERi_l48.omp_outlined268// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR2]] {269// CHECK1-NEXT: entry:270// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8271// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8272// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8273// CHECK1-NEXT: [[A:%.*]] = alloca i32, align 4274// CHECK1-NEXT: [[TMP:%.*]] = alloca ptr, align 8275// CHECK1-NEXT: [[B:%.*]] = alloca i32, align 4276// CHECK1-NEXT: [[C:%.*]] = alloca i32, align 4277// CHECK1-NEXT: [[_TMP1:%.*]] = alloca ptr, align 8278// CHECK1-NEXT: [[REF_TMP:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 8279// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8280// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8281// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8282// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8283// CHECK1-NEXT: store ptr [[A]], ptr [[TMP]], align 8284// CHECK1-NEXT: store ptr [[C]], ptr [[_TMP1]], align 8285// CHECK1-NEXT: [[TMP1:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 0286// CHECK1-NEXT: store ptr [[TMP0]], ptr [[TMP1]], align 8287// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 1288// CHECK1-NEXT: [[TMP3:%.*]] = load ptr, ptr [[TMP]], align 8289// CHECK1-NEXT: store ptr [[TMP3]], ptr [[TMP2]], align 8290// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 2291// CHECK1-NEXT: store ptr [[B]], ptr [[TMP4]], align 8292// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 3293// CHECK1-NEXT: [[TMP6:%.*]] = load ptr, ptr [[_TMP1]], align 8294// CHECK1-NEXT: store ptr [[TMP6]], ptr [[TMP5]], align 8295// CHECK1-NEXT: call void @_ZZN2SSC1ERiENKUlvE_clEv(ptr noundef nonnull align 8 dereferenceable(32) [[REF_TMP]])296// CHECK1-NEXT: ret void297//298//299// CHECK1-LABEL: define {{[^@]+}}@_ZZN2SSC1ERiENKUlvE_clEv300// CHECK1-SAME: (ptr noundef nonnull align 8 dereferenceable(32) [[THIS:%.*]]) #[[ATTR1]] {301// CHECK1-NEXT: entry:302// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8303// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8304// CHECK1-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8305// CHECK1-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0:%.*]], ptr [[THIS1]], i32 0, i32 0306// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8307// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[THIS1]], i32 0, i32 1308// CHECK1-NEXT: [[TMP3:%.*]] = load ptr, ptr [[TMP2]], align 8309// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[TMP3]], align 4310// CHECK1-NEXT: [[INC:%.*]] = add nsw i32 [[TMP4]], 1311// CHECK1-NEXT: store i32 [[INC]], ptr [[TMP3]], align 4312// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[THIS1]], i32 0, i32 2313// CHECK1-NEXT: [[TMP6:%.*]] = load ptr, ptr [[TMP5]], align 8314// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 4315// CHECK1-NEXT: [[DEC:%.*]] = add nsw i32 [[TMP7]], -1316// CHECK1-NEXT: store i32 [[DEC]], ptr [[TMP6]], align 4317// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[THIS1]], i32 0, i32 3318// CHECK1-NEXT: [[TMP9:%.*]] = load ptr, ptr [[TMP8]], align 8319// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[TMP9]], align 4320// CHECK1-NEXT: [[DIV:%.*]] = sdiv i32 [[TMP10]], 1321// CHECK1-NEXT: store i32 [[DIV]], ptr [[TMP9]], align 4322// CHECK1-NEXT: ret void323//324//325// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l117326// CHECK1-SAME: () #[[ATTR2]] {327// CHECK1-NEXT: entry:328// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l117.omp_outlined)329// CHECK1-NEXT: ret void330//331//332// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l117.omp_outlined333// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR2]] {334// CHECK1-NEXT: entry:335// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8336// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8337// CHECK1-NEXT: [[G:%.*]] = alloca i32, align 128338// CHECK1-NEXT: [[SIVAR:%.*]] = alloca i32, align 4339// CHECK1-NEXT: [[REF_TMP:%.*]] = alloca [[CLASS_ANON_1:%.*]], align 8340// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8341// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8342// CHECK1-NEXT: store i32 1, ptr [[G]], align 128343// CHECK1-NEXT: store i32 2, ptr [[SIVAR]], align 4344// CHECK1-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[REF_TMP]], i32 0, i32 0345// CHECK1-NEXT: store ptr [[G]], ptr [[TMP0]], align 8346// CHECK1-NEXT: [[TMP1:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[REF_TMP]], i32 0, i32 1347// CHECK1-NEXT: store ptr [[SIVAR]], ptr [[TMP1]], align 8348// CHECK1-NEXT: call void @"_ZZZ4mainENK3$_0clEvENKUlvE_clEv"(ptr noundef nonnull align 8 dereferenceable(16) [[REF_TMP]])349// CHECK1-NEXT: ret void350//351//352// CHECK3-LABEL: define {{[^@]+}}@main353// CHECK3-SAME: () #[[ATTR0:[0-9]+]] {354// CHECK3-NEXT: entry:355// CHECK3-NEXT: [[RETVAL:%.*]] = alloca i32, align 4356// CHECK3-NEXT: [[SS:%.*]] = alloca [[STRUCT_SS:%.*]], align 4357// CHECK3-NEXT: [[REF_TMP:%.*]] = alloca [[CLASS_ANON:%.*]], align 1358// CHECK3-NEXT: store i32 0, ptr [[RETVAL]], align 4359// CHECK3-NEXT: call void @_ZN2SSC1ERi(ptr noundef nonnull align 4 dereferenceable(12) [[SS]], ptr noundef nonnull align 4 dereferenceable(4) @_ZZ4mainE5sivar)360// CHECK3-NEXT: call void @"_ZZ4mainENK3$_0clEv"(ptr noundef nonnull align 1 dereferenceable(1) [[REF_TMP]])361// CHECK3-NEXT: ret i32 0362//363//364// CHECK3-LABEL: define {{[^@]+}}@_ZN2SSC1ERi365// CHECK3-SAME: (ptr noundef nonnull align 4 dereferenceable(12) [[THIS:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[D:%.*]]) unnamed_addr #[[ATTR1:[0-9]+]] comdat align 2 {366// CHECK3-NEXT: entry:367// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4368// CHECK3-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 4369// CHECK3-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4370// CHECK3-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 4371// CHECK3-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4372// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[D_ADDR]], align 4373// CHECK3-NEXT: call void @_ZN2SSC2ERi(ptr noundef nonnull align 4 dereferenceable(12) [[THIS1]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP0]])374// CHECK3-NEXT: ret void375//376//377// CHECK3-LABEL: define {{[^@]+}}@_ZN2SSC2ERi378// CHECK3-SAME: (ptr noundef nonnull align 4 dereferenceable(12) [[THIS:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[D:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {379// CHECK3-NEXT: entry:380// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4381// CHECK3-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 4382// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4383// CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4384// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4385// CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8386// CHECK3-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4387// CHECK3-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 4388// CHECK3-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4389// CHECK3-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[THIS1]], i32 0, i32 0390// CHECK3-NEXT: store i32 0, ptr [[A]], align 4391// CHECK3-NEXT: [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 1392// CHECK3-NEXT: [[BF_LOAD:%.*]] = load i8, ptr [[B]], align 4393// CHECK3-NEXT: [[BF_CLEAR:%.*]] = and i8 [[BF_LOAD]], -16394// CHECK3-NEXT: [[BF_SET:%.*]] = or i8 [[BF_CLEAR]], 0395// CHECK3-NEXT: store i8 [[BF_SET]], ptr [[B]], align 4396// CHECK3-NEXT: [[C:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 2397// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[D_ADDR]], align 4398// CHECK3-NEXT: store ptr [[TMP0]], ptr [[C]], align 4399// CHECK3-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0400// CHECK3-NEXT: store ptr [[THIS1]], ptr [[TMP1]], align 4401// CHECK3-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0402// CHECK3-NEXT: store ptr [[THIS1]], ptr [[TMP2]], align 4403// CHECK3-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0404// CHECK3-NEXT: store ptr null, ptr [[TMP3]], align 4405// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0406// CHECK3-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0407// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0408// CHECK3-NEXT: store i32 3, ptr [[TMP6]], align 4409// CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1410// CHECK3-NEXT: store i32 1, ptr [[TMP7]], align 4411// CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2412// CHECK3-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 4413// CHECK3-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3414// CHECK3-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 4415// CHECK3-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4416// CHECK3-NEXT: store ptr @.offload_sizes, ptr [[TMP10]], align 4417// CHECK3-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5418// CHECK3-NEXT: store ptr @.offload_maptypes, ptr [[TMP11]], align 4419// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6420// CHECK3-NEXT: store ptr null, ptr [[TMP12]], align 4421// CHECK3-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7422// CHECK3-NEXT: store ptr null, ptr [[TMP13]], align 4423// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8424// CHECK3-NEXT: store i64 0, ptr [[TMP14]], align 8425// CHECK3-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9426// CHECK3-NEXT: store i64 0, ptr [[TMP15]], align 8427// CHECK3-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10428// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4429// CHECK3-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11430// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4431// CHECK3-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12432// CHECK3-NEXT: store i32 0, ptr [[TMP18]], align 4433// CHECK3-NEXT: [[TMP19:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSC1ERi_l48.region_id, ptr [[KERNEL_ARGS]])434// CHECK3-NEXT: [[TMP20:%.*]] = icmp ne i32 [[TMP19]], 0435// CHECK3-NEXT: br i1 [[TMP20]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]436// CHECK3: omp_offload.failed:437// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSC1ERi_l48(ptr [[THIS1]]) #[[ATTR3:[0-9]+]]438// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]]439// CHECK3: omp_offload.cont:440// CHECK3-NEXT: ret void441//442//443// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSC1ERi_l48444// CHECK3-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR2:[0-9]+]] {445// CHECK3-NEXT: entry:446// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4447// CHECK3-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4448// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4449// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSC1ERi_l48.omp_outlined, ptr [[TMP0]])450// CHECK3-NEXT: ret void451//452//453// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSC1ERi_l48.omp_outlined454// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR2]] {455// CHECK3-NEXT: entry:456// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4457// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4458// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4459// CHECK3-NEXT: [[A:%.*]] = alloca i32, align 4460// CHECK3-NEXT: [[TMP:%.*]] = alloca ptr, align 4461// CHECK3-NEXT: [[B:%.*]] = alloca i32, align 4462// CHECK3-NEXT: [[C:%.*]] = alloca i32, align 4463// CHECK3-NEXT: [[_TMP1:%.*]] = alloca ptr, align 4464// CHECK3-NEXT: [[REF_TMP:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 4465// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4466// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4467// CHECK3-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4468// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4469// CHECK3-NEXT: store ptr [[A]], ptr [[TMP]], align 4470// CHECK3-NEXT: store ptr [[C]], ptr [[_TMP1]], align 4471// CHECK3-NEXT: [[TMP1:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 0472// CHECK3-NEXT: store ptr [[TMP0]], ptr [[TMP1]], align 4473// CHECK3-NEXT: [[TMP2:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 1474// CHECK3-NEXT: [[TMP3:%.*]] = load ptr, ptr [[TMP]], align 4475// CHECK3-NEXT: store ptr [[TMP3]], ptr [[TMP2]], align 4476// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 2477// CHECK3-NEXT: store ptr [[B]], ptr [[TMP4]], align 4478// CHECK3-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 3479// CHECK3-NEXT: [[TMP6:%.*]] = load ptr, ptr [[_TMP1]], align 4480// CHECK3-NEXT: store ptr [[TMP6]], ptr [[TMP5]], align 4481// CHECK3-NEXT: call void @_ZZN2SSC1ERiENKUlvE_clEv(ptr noundef nonnull align 4 dereferenceable(16) [[REF_TMP]])482// CHECK3-NEXT: ret void483//484//485// CHECK3-LABEL: define {{[^@]+}}@_ZZN2SSC1ERiENKUlvE_clEv486// CHECK3-SAME: (ptr noundef nonnull align 4 dereferenceable(16) [[THIS:%.*]]) #[[ATTR1]] align 2 {487// CHECK3-NEXT: entry:488// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4489// CHECK3-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4490// CHECK3-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4491// CHECK3-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0:%.*]], ptr [[THIS1]], i32 0, i32 0492// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 4493// CHECK3-NEXT: [[TMP2:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[THIS1]], i32 0, i32 1494// CHECK3-NEXT: [[TMP3:%.*]] = load ptr, ptr [[TMP2]], align 4495// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[TMP3]], align 4496// CHECK3-NEXT: [[INC:%.*]] = add nsw i32 [[TMP4]], 1497// CHECK3-NEXT: store i32 [[INC]], ptr [[TMP3]], align 4498// CHECK3-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[THIS1]], i32 0, i32 2499// CHECK3-NEXT: [[TMP6:%.*]] = load ptr, ptr [[TMP5]], align 4500// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 4501// CHECK3-NEXT: [[DEC:%.*]] = add nsw i32 [[TMP7]], -1502// CHECK3-NEXT: store i32 [[DEC]], ptr [[TMP6]], align 4503// CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[THIS1]], i32 0, i32 3504// CHECK3-NEXT: [[TMP9:%.*]] = load ptr, ptr [[TMP8]], align 4505// CHECK3-NEXT: [[TMP10:%.*]] = load i32, ptr [[TMP9]], align 4506// CHECK3-NEXT: [[DIV:%.*]] = sdiv i32 [[TMP10]], 1507// CHECK3-NEXT: store i32 [[DIV]], ptr [[TMP9]], align 4508// CHECK3-NEXT: ret void509//510//511// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l117512// CHECK3-SAME: () #[[ATTR2]] {513// CHECK3-NEXT: entry:514// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l117.omp_outlined)515// CHECK3-NEXT: ret void516//517//518// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l117.omp_outlined519// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR2]] {520// CHECK3-NEXT: entry:521// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4522// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4523// CHECK3-NEXT: [[G:%.*]] = alloca i32, align 128524// CHECK3-NEXT: [[SIVAR:%.*]] = alloca i32, align 4525// CHECK3-NEXT: [[REF_TMP:%.*]] = alloca [[CLASS_ANON_1:%.*]], align 4526// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4527// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4528// CHECK3-NEXT: store i32 1, ptr [[G]], align 128529// CHECK3-NEXT: store i32 2, ptr [[SIVAR]], align 4530// CHECK3-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[REF_TMP]], i32 0, i32 0531// CHECK3-NEXT: store ptr [[G]], ptr [[TMP0]], align 4532// CHECK3-NEXT: [[TMP1:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[REF_TMP]], i32 0, i32 1533// CHECK3-NEXT: store ptr [[SIVAR]], ptr [[TMP1]], align 4534// CHECK3-NEXT: call void @"_ZZZ4mainENK3$_0clEvENKUlvE_clEv"(ptr noundef nonnull align 4 dereferenceable(8) [[REF_TMP]])535// CHECK3-NEXT: ret void536//537//538// CHECK9-LABEL: define {{[^@]+}}@main539// CHECK9-SAME: () #[[ATTR0:[0-9]+]] {540// CHECK9-NEXT: entry:541// CHECK9-NEXT: [[RETVAL:%.*]] = alloca i32, align 4542// CHECK9-NEXT: [[SS:%.*]] = alloca [[STRUCT_SS:%.*]], align 8543// CHECK9-NEXT: [[TEST:%.*]] = alloca [[STRUCT_S:%.*]], align 4544// CHECK9-NEXT: [[T_VAR:%.*]] = alloca i32, align 4545// CHECK9-NEXT: [[VEC:%.*]] = alloca [2 x i32], align 4546// CHECK9-NEXT: [[S_ARR:%.*]] = alloca [2 x %struct.S], align 4547// CHECK9-NEXT: [[VAR:%.*]] = alloca [[STRUCT_S]], align 4548// CHECK9-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8549// CHECK9-NEXT: store i32 0, ptr [[RETVAL]], align 4550// CHECK9-NEXT: call void @_ZN2SSC1ERi(ptr noundef nonnull align 8 dereferenceable(16) [[SS]], ptr noundef nonnull align 4 dereferenceable(4) @_ZZ4mainE5sivar)551// CHECK9-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[TEST]])552// CHECK9-NEXT: store i32 0, ptr [[T_VAR]], align 4553// CHECK9-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[VEC]], ptr align 4 @__const.main.vec, i64 8, i1 false)554// CHECK9-NEXT: call void @_ZN1SIfEC1Ef(ptr noundef nonnull align 4 dereferenceable(4) [[S_ARR]], float noundef 1.000000e+00)555// CHECK9-NEXT: [[ARRAYINIT_ELEMENT:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[S_ARR]], i64 1556// CHECK9-NEXT: call void @_ZN1SIfEC1Ef(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYINIT_ELEMENT]], float noundef 2.000000e+00)557// CHECK9-NEXT: call void @_ZN1SIfEC1Ef(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]], float noundef 3.000000e+00)558// CHECK9-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0559// CHECK9-NEXT: store i32 3, ptr [[TMP0]], align 4560// CHECK9-NEXT: [[TMP1:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1561// CHECK9-NEXT: store i32 0, ptr [[TMP1]], align 4562// CHECK9-NEXT: [[TMP2:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2563// CHECK9-NEXT: store ptr null, ptr [[TMP2]], align 8564// CHECK9-NEXT: [[TMP3:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3565// CHECK9-NEXT: store ptr null, ptr [[TMP3]], align 8566// CHECK9-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4567// CHECK9-NEXT: store ptr null, ptr [[TMP4]], align 8568// CHECK9-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5569// CHECK9-NEXT: store ptr null, ptr [[TMP5]], align 8570// CHECK9-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6571// CHECK9-NEXT: store ptr null, ptr [[TMP6]], align 8572// CHECK9-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7573// CHECK9-NEXT: store ptr null, ptr [[TMP7]], align 8574// CHECK9-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8575// CHECK9-NEXT: store i64 0, ptr [[TMP8]], align 8576// CHECK9-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9577// CHECK9-NEXT: store i64 0, ptr [[TMP9]], align 8578// CHECK9-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10579// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP10]], align 4580// CHECK9-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11581// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP11]], align 4582// CHECK9-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12583// CHECK9-NEXT: store i32 0, ptr [[TMP12]], align 4584// CHECK9-NEXT: [[TMP13:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l136.region_id, ptr [[KERNEL_ARGS]])585// CHECK9-NEXT: [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 0586// CHECK9-NEXT: br i1 [[TMP14]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]587// CHECK9: omp_offload.failed:588// CHECK9-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l136() #[[ATTR4:[0-9]+]]589// CHECK9-NEXT: br label [[OMP_OFFLOAD_CONT]]590// CHECK9: omp_offload.cont:591// CHECK9-NEXT: [[CALL:%.*]] = call noundef signext i32 @_Z5tmainIiET_v()592// CHECK9-NEXT: store i32 [[CALL]], ptr [[RETVAL]], align 4593// CHECK9-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]]) #[[ATTR4]]594// CHECK9-NEXT: [[ARRAY_BEGIN:%.*]] = getelementptr inbounds [2 x %struct.S], ptr [[S_ARR]], i32 0, i32 0595// CHECK9-NEXT: [[TMP15:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAY_BEGIN]], i64 2596// CHECK9-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]597// CHECK9: arraydestroy.body:598// CHECK9-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP15]], [[OMP_OFFLOAD_CONT]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]599// CHECK9-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i64 -1600// CHECK9-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR4]]601// CHECK9-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN]]602// CHECK9-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE1:%.*]], label [[ARRAYDESTROY_BODY]]603// CHECK9: arraydestroy.done1:604// CHECK9-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[TEST]]) #[[ATTR4]]605// CHECK9-NEXT: [[TMP16:%.*]] = load i32, ptr [[RETVAL]], align 4606// CHECK9-NEXT: ret i32 [[TMP16]]607//608//609// CHECK9-LABEL: define {{[^@]+}}@_ZN2SSC1ERi610// CHECK9-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[THIS:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[D:%.*]]) unnamed_addr #[[ATTR1:[0-9]+]] comdat {611// CHECK9-NEXT: entry:612// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8613// CHECK9-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 8614// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8615// CHECK9-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 8616// CHECK9-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8617// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[D_ADDR]], align 8618// CHECK9-NEXT: call void @_ZN2SSC2ERi(ptr noundef nonnull align 8 dereferenceable(16) [[THIS1]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP0]])619// CHECK9-NEXT: ret void620//621//622// CHECK9-LABEL: define {{[^@]+}}@_ZN1SIfEC1Ev623// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {624// CHECK9-NEXT: entry:625// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8626// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8627// CHECK9-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8628// CHECK9-NEXT: call void @_ZN1SIfEC2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]])629// CHECK9-NEXT: ret void630//631//632// CHECK9-LABEL: define {{[^@]+}}@_ZN1SIfEC1Ef633// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]], float noundef [[A:%.*]]) unnamed_addr #[[ATTR1]] comdat {634// CHECK9-NEXT: entry:635// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8636// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca float, align 4637// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8638// CHECK9-NEXT: store float [[A]], ptr [[A_ADDR]], align 4639// CHECK9-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8640// CHECK9-NEXT: [[TMP0:%.*]] = load float, ptr [[A_ADDR]], align 4641// CHECK9-NEXT: call void @_ZN1SIfEC2Ef(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]], float noundef [[TMP0]])642// CHECK9-NEXT: ret void643//644//645// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l136646// CHECK9-SAME: () #[[ATTR3:[0-9]+]] {647// CHECK9-NEXT: entry:648// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l136.omp_outlined)649// CHECK9-NEXT: ret void650//651//652// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l136.omp_outlined653// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR3]] {654// CHECK9-NEXT: entry:655// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8656// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8657// CHECK9-NEXT: [[T_VAR:%.*]] = alloca i32, align 4658// CHECK9-NEXT: [[VEC:%.*]] = alloca [2 x i32], align 4659// CHECK9-NEXT: [[S_ARR:%.*]] = alloca [2 x %struct.S], align 4660// CHECK9-NEXT: [[VAR:%.*]] = alloca [[STRUCT_S:%.*]], align 4661// CHECK9-NEXT: [[SIVAR:%.*]] = alloca i32, align 4662// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8663// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8664// CHECK9-NEXT: [[ARRAY_BEGIN:%.*]] = getelementptr inbounds [2 x %struct.S], ptr [[S_ARR]], i32 0, i32 0665// CHECK9-NEXT: [[ARRAYCTOR_END:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAY_BEGIN]], i64 2666// CHECK9-NEXT: br label [[ARRAYCTOR_LOOP:%.*]]667// CHECK9: arrayctor.loop:668// CHECK9-NEXT: [[ARRAYCTOR_CUR:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ]669// CHECK9-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]])670// CHECK9-NEXT: [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYCTOR_CUR]], i64 1671// CHECK9-NEXT: [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], [[ARRAYCTOR_END]]672// CHECK9-NEXT: br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]]673// CHECK9: arrayctor.cont:674// CHECK9-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]])675// CHECK9-NEXT: [[TMP0:%.*]] = load i32, ptr [[T_VAR]], align 4676// CHECK9-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [2 x i32], ptr [[VEC]], i64 0, i64 0677// CHECK9-NEXT: store i32 [[TMP0]], ptr [[ARRAYIDX]], align 4678// CHECK9-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [2 x %struct.S], ptr [[S_ARR]], i64 0, i64 0679// CHECK9-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[ARRAYIDX1]], ptr align 4 [[VAR]], i64 4, i1 false)680// CHECK9-NEXT: store i32 3, ptr [[SIVAR]], align 4681// CHECK9-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]]) #[[ATTR4]]682// CHECK9-NEXT: [[ARRAY_BEGIN2:%.*]] = getelementptr inbounds [2 x %struct.S], ptr [[S_ARR]], i32 0, i32 0683// CHECK9-NEXT: [[TMP1:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAY_BEGIN2]], i64 2684// CHECK9-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]685// CHECK9: arraydestroy.body:686// CHECK9-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP1]], [[ARRAYCTOR_CONT]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]687// CHECK9-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i64 -1688// CHECK9-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR4]]689// CHECK9-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN2]]690// CHECK9-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE3:%.*]], label [[ARRAYDESTROY_BODY]]691// CHECK9: arraydestroy.done3:692// CHECK9-NEXT: ret void693//694//695// CHECK9-LABEL: define {{[^@]+}}@_ZN1SIfED1Ev696// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {697// CHECK9-NEXT: entry:698// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8699// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8700// CHECK9-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8701// CHECK9-NEXT: call void @_ZN1SIfED2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]]) #[[ATTR4]]702// CHECK9-NEXT: ret void703//704//705// CHECK9-LABEL: define {{[^@]+}}@_Z5tmainIiET_v706// CHECK9-SAME: () #[[ATTR1]] comdat {707// CHECK9-NEXT: entry:708// CHECK9-NEXT: [[RETVAL:%.*]] = alloca i32, align 4709// CHECK9-NEXT: [[TEST:%.*]] = alloca [[STRUCT_S_0:%.*]], align 4710// CHECK9-NEXT: [[SST:%.*]] = alloca [[STRUCT_SST:%.*]], align 4711// CHECK9-NEXT: [[T_VAR:%.*]] = alloca i32, align 128712// CHECK9-NEXT: [[VEC:%.*]] = alloca [2 x i32], align 128713// CHECK9-NEXT: [[S_ARR:%.*]] = alloca [2 x %struct.S.0], align 128714// CHECK9-NEXT: [[VAR:%.*]] = alloca [[STRUCT_S_0]], align 128715// CHECK9-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8716// CHECK9-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[TEST]])717// CHECK9-NEXT: call void @_ZN3SSTIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[SST]])718// CHECK9-NEXT: store i32 0, ptr [[T_VAR]], align 128719// CHECK9-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 128 [[VEC]], ptr align 128 @__const._Z5tmainIiET_v.vec, i64 8, i1 false)720// CHECK9-NEXT: call void @_ZN1SIiEC1Ei(ptr noundef nonnull align 4 dereferenceable(4) [[S_ARR]], i32 noundef signext 1)721// CHECK9-NEXT: [[ARRAYINIT_ELEMENT:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[S_ARR]], i64 1722// CHECK9-NEXT: call void @_ZN1SIiEC1Ei(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYINIT_ELEMENT]], i32 noundef signext 2)723// CHECK9-NEXT: call void @_ZN1SIiEC1Ei(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]], i32 noundef signext 3)724// CHECK9-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0725// CHECK9-NEXT: store i32 3, ptr [[TMP0]], align 4726// CHECK9-NEXT: [[TMP1:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1727// CHECK9-NEXT: store i32 0, ptr [[TMP1]], align 4728// CHECK9-NEXT: [[TMP2:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2729// CHECK9-NEXT: store ptr null, ptr [[TMP2]], align 8730// CHECK9-NEXT: [[TMP3:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3731// CHECK9-NEXT: store ptr null, ptr [[TMP3]], align 8732// CHECK9-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4733// CHECK9-NEXT: store ptr null, ptr [[TMP4]], align 8734// CHECK9-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5735// CHECK9-NEXT: store ptr null, ptr [[TMP5]], align 8736// CHECK9-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6737// CHECK9-NEXT: store ptr null, ptr [[TMP6]], align 8738// CHECK9-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7739// CHECK9-NEXT: store ptr null, ptr [[TMP7]], align 8740// CHECK9-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8741// CHECK9-NEXT: store i64 0, ptr [[TMP8]], align 8742// CHECK9-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9743// CHECK9-NEXT: store i64 0, ptr [[TMP9]], align 8744// CHECK9-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10745// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP10]], align 4746// CHECK9-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11747// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP11]], align 4748// CHECK9-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12749// CHECK9-NEXT: store i32 0, ptr [[TMP12]], align 4750// CHECK9-NEXT: [[TMP13:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l86.region_id, ptr [[KERNEL_ARGS]])751// CHECK9-NEXT: [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 0752// CHECK9-NEXT: br i1 [[TMP14]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]753// CHECK9: omp_offload.failed:754// CHECK9-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l86() #[[ATTR4]]755// CHECK9-NEXT: br label [[OMP_OFFLOAD_CONT]]756// CHECK9: omp_offload.cont:757// CHECK9-NEXT: store i32 0, ptr [[RETVAL]], align 4758// CHECK9-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]]) #[[ATTR4]]759// CHECK9-NEXT: [[ARRAY_BEGIN:%.*]] = getelementptr inbounds [2 x %struct.S.0], ptr [[S_ARR]], i32 0, i32 0760// CHECK9-NEXT: [[TMP15:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAY_BEGIN]], i64 2761// CHECK9-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]762// CHECK9: arraydestroy.body:763// CHECK9-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP15]], [[OMP_OFFLOAD_CONT]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]764// CHECK9-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i64 -1765// CHECK9-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR4]]766// CHECK9-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN]]767// CHECK9-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE1:%.*]], label [[ARRAYDESTROY_BODY]]768// CHECK9: arraydestroy.done1:769// CHECK9-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[TEST]]) #[[ATTR4]]770// CHECK9-NEXT: [[TMP16:%.*]] = load i32, ptr [[RETVAL]], align 4771// CHECK9-NEXT: ret i32 [[TMP16]]772//773//774// CHECK9-LABEL: define {{[^@]+}}@_ZN2SSC2ERi775// CHECK9-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[THIS:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[D:%.*]]) unnamed_addr #[[ATTR1]] comdat {776// CHECK9-NEXT: entry:777// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8778// CHECK9-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 8779// CHECK9-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8780// CHECK9-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8781// CHECK9-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8782// CHECK9-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8783// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8784// CHECK9-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 8785// CHECK9-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8786// CHECK9-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[THIS1]], i32 0, i32 0787// CHECK9-NEXT: store i32 0, ptr [[A]], align 8788// CHECK9-NEXT: [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 1789// CHECK9-NEXT: [[BF_LOAD:%.*]] = load i8, ptr [[B]], align 4790// CHECK9-NEXT: [[BF_CLEAR:%.*]] = and i8 [[BF_LOAD]], -16791// CHECK9-NEXT: [[BF_SET:%.*]] = or i8 [[BF_CLEAR]], 0792// CHECK9-NEXT: store i8 [[BF_SET]], ptr [[B]], align 4793// CHECK9-NEXT: [[C:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 2794// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[D_ADDR]], align 8795// CHECK9-NEXT: store ptr [[TMP0]], ptr [[C]], align 8796// CHECK9-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0797// CHECK9-NEXT: store ptr [[THIS1]], ptr [[TMP1]], align 8798// CHECK9-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0799// CHECK9-NEXT: store ptr [[THIS1]], ptr [[TMP2]], align 8800// CHECK9-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0801// CHECK9-NEXT: store ptr null, ptr [[TMP3]], align 8802// CHECK9-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0803// CHECK9-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0804// CHECK9-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0805// CHECK9-NEXT: store i32 3, ptr [[TMP6]], align 4806// CHECK9-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1807// CHECK9-NEXT: store i32 1, ptr [[TMP7]], align 4808// CHECK9-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2809// CHECK9-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 8810// CHECK9-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3811// CHECK9-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 8812// CHECK9-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4813// CHECK9-NEXT: store ptr @.offload_sizes, ptr [[TMP10]], align 8814// CHECK9-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5815// CHECK9-NEXT: store ptr @.offload_maptypes, ptr [[TMP11]], align 8816// CHECK9-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6817// CHECK9-NEXT: store ptr null, ptr [[TMP12]], align 8818// CHECK9-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7819// CHECK9-NEXT: store ptr null, ptr [[TMP13]], align 8820// CHECK9-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8821// CHECK9-NEXT: store i64 0, ptr [[TMP14]], align 8822// CHECK9-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9823// CHECK9-NEXT: store i64 0, ptr [[TMP15]], align 8824// CHECK9-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10825// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4826// CHECK9-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11827// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4828// CHECK9-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12829// CHECK9-NEXT: store i32 0, ptr [[TMP18]], align 4830// CHECK9-NEXT: [[TMP19:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSC1ERi_l48.region_id, ptr [[KERNEL_ARGS]])831// CHECK9-NEXT: [[TMP20:%.*]] = icmp ne i32 [[TMP19]], 0832// CHECK9-NEXT: br i1 [[TMP20]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]833// CHECK9: omp_offload.failed:834// CHECK9-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSC1ERi_l48(ptr [[THIS1]]) #[[ATTR4]]835// CHECK9-NEXT: br label [[OMP_OFFLOAD_CONT]]836// CHECK9: omp_offload.cont:837// CHECK9-NEXT: ret void838//839//840// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSC1ERi_l48841// CHECK9-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR3]] {842// CHECK9-NEXT: entry:843// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8844// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8845// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8846// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSC1ERi_l48.omp_outlined, ptr [[TMP0]])847// CHECK9-NEXT: ret void848//849//850// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSC1ERi_l48.omp_outlined851// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR3]] {852// CHECK9-NEXT: entry:853// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8854// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8855// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8856// CHECK9-NEXT: [[A:%.*]] = alloca i32, align 4857// CHECK9-NEXT: [[TMP:%.*]] = alloca ptr, align 8858// CHECK9-NEXT: [[B:%.*]] = alloca i32, align 4859// CHECK9-NEXT: [[C:%.*]] = alloca i32, align 4860// CHECK9-NEXT: [[_TMP1:%.*]] = alloca ptr, align 8861// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8862// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8863// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8864// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8865// CHECK9-NEXT: store ptr [[A]], ptr [[TMP]], align 8866// CHECK9-NEXT: store ptr [[C]], ptr [[_TMP1]], align 8867// CHECK9-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP]], align 8868// CHECK9-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4869// CHECK9-NEXT: [[INC:%.*]] = add nsw i32 [[TMP2]], 1870// CHECK9-NEXT: store i32 [[INC]], ptr [[TMP1]], align 4871// CHECK9-NEXT: [[TMP3:%.*]] = load i32, ptr [[B]], align 4872// CHECK9-NEXT: [[DEC:%.*]] = add nsw i32 [[TMP3]], -1873// CHECK9-NEXT: store i32 [[DEC]], ptr [[B]], align 4874// CHECK9-NEXT: [[TMP4:%.*]] = load ptr, ptr [[_TMP1]], align 8875// CHECK9-NEXT: [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 4876// CHECK9-NEXT: [[DIV:%.*]] = sdiv i32 [[TMP5]], 1877// CHECK9-NEXT: store i32 [[DIV]], ptr [[TMP4]], align 4878// CHECK9-NEXT: ret void879//880//881// CHECK9-LABEL: define {{[^@]+}}@_ZN1SIfEC2Ev882// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {883// CHECK9-NEXT: entry:884// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8885// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8886// CHECK9-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8887// CHECK9-NEXT: [[F:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[THIS1]], i32 0, i32 0888// CHECK9-NEXT: store float 0.000000e+00, ptr [[F]], align 4889// CHECK9-NEXT: ret void890//891//892// CHECK9-LABEL: define {{[^@]+}}@_ZN1SIfEC2Ef893// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]], float noundef [[A:%.*]]) unnamed_addr #[[ATTR1]] comdat {894// CHECK9-NEXT: entry:895// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8896// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca float, align 4897// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8898// CHECK9-NEXT: store float [[A]], ptr [[A_ADDR]], align 4899// CHECK9-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8900// CHECK9-NEXT: [[F:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[THIS1]], i32 0, i32 0901// CHECK9-NEXT: [[TMP0:%.*]] = load float, ptr [[A_ADDR]], align 4902// CHECK9-NEXT: store float [[TMP0]], ptr [[F]], align 4903// CHECK9-NEXT: ret void904//905//906// CHECK9-LABEL: define {{[^@]+}}@_ZN1SIfED2Ev907// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {908// CHECK9-NEXT: entry:909// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8910// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8911// CHECK9-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8912// CHECK9-NEXT: ret void913//914//915// CHECK9-LABEL: define {{[^@]+}}@_ZN1SIiEC1Ev916// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {917// CHECK9-NEXT: entry:918// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8919// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8920// CHECK9-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8921// CHECK9-NEXT: call void @_ZN1SIiEC2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]])922// CHECK9-NEXT: ret void923//924//925// CHECK9-LABEL: define {{[^@]+}}@_ZN3SSTIiEC1Ev926// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {927// CHECK9-NEXT: entry:928// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8929// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8930// CHECK9-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8931// CHECK9-NEXT: call void @_ZN3SSTIiEC2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]])932// CHECK9-NEXT: ret void933//934//935// CHECK9-LABEL: define {{[^@]+}}@_ZN1SIiEC1Ei936// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]], i32 noundef signext [[A:%.*]]) unnamed_addr #[[ATTR1]] comdat {937// CHECK9-NEXT: entry:938// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8939// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4940// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8941// CHECK9-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4942// CHECK9-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8943// CHECK9-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 4944// CHECK9-NEXT: call void @_ZN1SIiEC2Ei(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]], i32 noundef signext [[TMP0]])945// CHECK9-NEXT: ret void946//947//948// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l86949// CHECK9-SAME: () #[[ATTR3]] {950// CHECK9-NEXT: entry:951// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l86.omp_outlined)952// CHECK9-NEXT: ret void953//954//955// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l86.omp_outlined956// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR3]] {957// CHECK9-NEXT: entry:958// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8959// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8960// CHECK9-NEXT: [[T_VAR:%.*]] = alloca i32, align 128961// CHECK9-NEXT: [[VEC:%.*]] = alloca [2 x i32], align 128962// CHECK9-NEXT: [[S_ARR:%.*]] = alloca [2 x %struct.S.0], align 128963// CHECK9-NEXT: [[VAR:%.*]] = alloca [[STRUCT_S_0:%.*]], align 128964// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8965// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8966// CHECK9-NEXT: [[ARRAY_BEGIN:%.*]] = getelementptr inbounds [2 x %struct.S.0], ptr [[S_ARR]], i32 0, i32 0967// CHECK9-NEXT: [[ARRAYCTOR_END:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAY_BEGIN]], i64 2968// CHECK9-NEXT: br label [[ARRAYCTOR_LOOP:%.*]]969// CHECK9: arrayctor.loop:970// CHECK9-NEXT: [[ARRAYCTOR_CUR:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ]971// CHECK9-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]])972// CHECK9-NEXT: [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYCTOR_CUR]], i64 1973// CHECK9-NEXT: [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], [[ARRAYCTOR_END]]974// CHECK9-NEXT: br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]]975// CHECK9: arrayctor.cont:976// CHECK9-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]])977// CHECK9-NEXT: [[TMP0:%.*]] = load i32, ptr [[T_VAR]], align 128978// CHECK9-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [2 x i32], ptr [[VEC]], i64 0, i64 0979// CHECK9-NEXT: store i32 [[TMP0]], ptr [[ARRAYIDX]], align 128980// CHECK9-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [2 x %struct.S.0], ptr [[S_ARR]], i64 0, i64 0981// CHECK9-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 128 [[ARRAYIDX1]], ptr align 128 [[VAR]], i64 4, i1 false)982// CHECK9-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]]) #[[ATTR4]]983// CHECK9-NEXT: [[ARRAY_BEGIN2:%.*]] = getelementptr inbounds [2 x %struct.S.0], ptr [[S_ARR]], i32 0, i32 0984// CHECK9-NEXT: [[TMP1:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAY_BEGIN2]], i64 2985// CHECK9-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]986// CHECK9: arraydestroy.body:987// CHECK9-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP1]], [[ARRAYCTOR_CONT]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]988// CHECK9-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i64 -1989// CHECK9-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR4]]990// CHECK9-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN2]]991// CHECK9-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE3:%.*]], label [[ARRAYDESTROY_BODY]]992// CHECK9: arraydestroy.done3:993// CHECK9-NEXT: ret void994//995//996// CHECK9-LABEL: define {{[^@]+}}@_ZN1SIiED1Ev997// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {998// CHECK9-NEXT: entry:999// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81000// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81001// CHECK9-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81002// CHECK9-NEXT: call void @_ZN1SIiED2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]]) #[[ATTR4]]1003// CHECK9-NEXT: ret void1004//1005//1006// CHECK9-LABEL: define {{[^@]+}}@_ZN1SIiEC2Ev1007// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {1008// CHECK9-NEXT: entry:1009// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81010// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81011// CHECK9-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81012// CHECK9-NEXT: [[F:%.*]] = getelementptr inbounds nuw [[STRUCT_S_0:%.*]], ptr [[THIS1]], i32 0, i32 01013// CHECK9-NEXT: store i32 0, ptr [[F]], align 41014// CHECK9-NEXT: ret void1015//1016//1017// CHECK9-LABEL: define {{[^@]+}}@_ZN3SSTIiEC2Ev1018// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {1019// CHECK9-NEXT: entry:1020// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81021// CHECK9-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 81022// CHECK9-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 81023// CHECK9-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 81024// CHECK9-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81025// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81026// CHECK9-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81027// CHECK9-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SST:%.*]], ptr [[THIS1]], i32 0, i32 01028// CHECK9-NEXT: store i32 0, ptr [[A]], align 41029// CHECK9-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01030// CHECK9-NEXT: store ptr [[THIS1]], ptr [[TMP0]], align 81031// CHECK9-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01032// CHECK9-NEXT: store ptr [[THIS1]], ptr [[TMP1]], align 81033// CHECK9-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 01034// CHECK9-NEXT: store ptr null, ptr [[TMP2]], align 81035// CHECK9-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01036// CHECK9-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01037// CHECK9-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01038// CHECK9-NEXT: store i32 3, ptr [[TMP5]], align 41039// CHECK9-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11040// CHECK9-NEXT: store i32 1, ptr [[TMP6]], align 41041// CHECK9-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21042// CHECK9-NEXT: store ptr [[TMP3]], ptr [[TMP7]], align 81043// CHECK9-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31044// CHECK9-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 81045// CHECK9-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41046// CHECK9-NEXT: store ptr @.offload_sizes.1, ptr [[TMP9]], align 81047// CHECK9-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51048// CHECK9-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP10]], align 81049// CHECK9-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61050// CHECK9-NEXT: store ptr null, ptr [[TMP11]], align 81051// CHECK9-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71052// CHECK9-NEXT: store ptr null, ptr [[TMP12]], align 81053// CHECK9-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81054// CHECK9-NEXT: store i64 0, ptr [[TMP13]], align 81055// CHECK9-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91056// CHECK9-NEXT: store i64 0, ptr [[TMP14]], align 81057// CHECK9-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101058// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 41059// CHECK9-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111060// CHECK9-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 41061// CHECK9-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121062// CHECK9-NEXT: store i32 0, ptr [[TMP17]], align 41063// CHECK9-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN3SSTIiEC1Ev_l64.region_id, ptr [[KERNEL_ARGS]])1064// CHECK9-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 01065// CHECK9-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1066// CHECK9: omp_offload.failed:1067// CHECK9-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN3SSTIiEC1Ev_l64(ptr [[THIS1]]) #[[ATTR4]]1068// CHECK9-NEXT: br label [[OMP_OFFLOAD_CONT]]1069// CHECK9: omp_offload.cont:1070// CHECK9-NEXT: ret void1071//1072//1073// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN3SSTIiEC1Ev_l641074// CHECK9-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR3]] {1075// CHECK9-NEXT: entry:1076// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81077// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81078// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81079// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN3SSTIiEC1Ev_l64.omp_outlined, ptr [[TMP0]])1080// CHECK9-NEXT: ret void1081//1082//1083// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN3SSTIiEC1Ev_l64.omp_outlined1084// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR3]] {1085// CHECK9-NEXT: entry:1086// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81087// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81088// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81089// CHECK9-NEXT: [[A:%.*]] = alloca i32, align 41090// CHECK9-NEXT: [[TMP:%.*]] = alloca ptr, align 81091// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81092// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81093// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81094// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81095// CHECK9-NEXT: store ptr [[A]], ptr [[TMP]], align 81096// CHECK9-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP]], align 81097// CHECK9-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41098// CHECK9-NEXT: [[INC:%.*]] = add nsw i32 [[TMP2]], 11099// CHECK9-NEXT: store i32 [[INC]], ptr [[TMP1]], align 41100// CHECK9-NEXT: ret void1101//1102//1103// CHECK9-LABEL: define {{[^@]+}}@_ZN1SIiEC2Ei1104// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]], i32 noundef signext [[A:%.*]]) unnamed_addr #[[ATTR1]] comdat {1105// CHECK9-NEXT: entry:1106// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81107// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca i32, align 41108// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81109// CHECK9-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 41110// CHECK9-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81111// CHECK9-NEXT: [[F:%.*]] = getelementptr inbounds nuw [[STRUCT_S_0:%.*]], ptr [[THIS1]], i32 0, i32 01112// CHECK9-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 41113// CHECK9-NEXT: store i32 [[TMP0]], ptr [[F]], align 41114// CHECK9-NEXT: ret void1115//1116//1117// CHECK9-LABEL: define {{[^@]+}}@_ZN1SIiED2Ev1118// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {1119// CHECK9-NEXT: entry:1120// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81121// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81122// CHECK9-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81123// CHECK9-NEXT: ret void1124//1125//1126// CHECK11-LABEL: define {{[^@]+}}@main1127// CHECK11-SAME: () #[[ATTR0:[0-9]+]] {1128// CHECK11-NEXT: entry:1129// CHECK11-NEXT: [[RETVAL:%.*]] = alloca i32, align 41130// CHECK11-NEXT: [[SS:%.*]] = alloca [[STRUCT_SS:%.*]], align 41131// CHECK11-NEXT: [[TEST:%.*]] = alloca [[STRUCT_S:%.*]], align 41132// CHECK11-NEXT: [[T_VAR:%.*]] = alloca i32, align 41133// CHECK11-NEXT: [[VEC:%.*]] = alloca [2 x i32], align 41134// CHECK11-NEXT: [[S_ARR:%.*]] = alloca [2 x %struct.S], align 41135// CHECK11-NEXT: [[VAR:%.*]] = alloca [[STRUCT_S]], align 41136// CHECK11-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81137// CHECK11-NEXT: store i32 0, ptr [[RETVAL]], align 41138// CHECK11-NEXT: call void @_ZN2SSC1ERi(ptr noundef nonnull align 4 dereferenceable(12) [[SS]], ptr noundef nonnull align 4 dereferenceable(4) @_ZZ4mainE5sivar)1139// CHECK11-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[TEST]])1140// CHECK11-NEXT: store i32 0, ptr [[T_VAR]], align 41141// CHECK11-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[VEC]], ptr align 4 @__const.main.vec, i32 8, i1 false)1142// CHECK11-NEXT: call void @_ZN1SIfEC1Ef(ptr noundef nonnull align 4 dereferenceable(4) [[S_ARR]], float noundef 1.000000e+00)1143// CHECK11-NEXT: [[ARRAYINIT_ELEMENT:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[S_ARR]], i32 11144// CHECK11-NEXT: call void @_ZN1SIfEC1Ef(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYINIT_ELEMENT]], float noundef 2.000000e+00)1145// CHECK11-NEXT: call void @_ZN1SIfEC1Ef(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]], float noundef 3.000000e+00)1146// CHECK11-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01147// CHECK11-NEXT: store i32 3, ptr [[TMP0]], align 41148// CHECK11-NEXT: [[TMP1:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11149// CHECK11-NEXT: store i32 0, ptr [[TMP1]], align 41150// CHECK11-NEXT: [[TMP2:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21151// CHECK11-NEXT: store ptr null, ptr [[TMP2]], align 41152// CHECK11-NEXT: [[TMP3:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31153// CHECK11-NEXT: store ptr null, ptr [[TMP3]], align 41154// CHECK11-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41155// CHECK11-NEXT: store ptr null, ptr [[TMP4]], align 41156// CHECK11-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51157// CHECK11-NEXT: store ptr null, ptr [[TMP5]], align 41158// CHECK11-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61159// CHECK11-NEXT: store ptr null, ptr [[TMP6]], align 41160// CHECK11-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71161// CHECK11-NEXT: store ptr null, ptr [[TMP7]], align 41162// CHECK11-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81163// CHECK11-NEXT: store i64 0, ptr [[TMP8]], align 81164// CHECK11-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91165// CHECK11-NEXT: store i64 0, ptr [[TMP9]], align 81166// CHECK11-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101167// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP10]], align 41168// CHECK11-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111169// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP11]], align 41170// CHECK11-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121171// CHECK11-NEXT: store i32 0, ptr [[TMP12]], align 41172// CHECK11-NEXT: [[TMP13:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l136.region_id, ptr [[KERNEL_ARGS]])1173// CHECK11-NEXT: [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 01174// CHECK11-NEXT: br i1 [[TMP14]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1175// CHECK11: omp_offload.failed:1176// CHECK11-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l136() #[[ATTR4:[0-9]+]]1177// CHECK11-NEXT: br label [[OMP_OFFLOAD_CONT]]1178// CHECK11: omp_offload.cont:1179// CHECK11-NEXT: [[CALL:%.*]] = call noundef i32 @_Z5tmainIiET_v()1180// CHECK11-NEXT: store i32 [[CALL]], ptr [[RETVAL]], align 41181// CHECK11-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]]) #[[ATTR4]]1182// CHECK11-NEXT: [[ARRAY_BEGIN:%.*]] = getelementptr inbounds [2 x %struct.S], ptr [[S_ARR]], i32 0, i32 01183// CHECK11-NEXT: [[TMP15:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAY_BEGIN]], i32 21184// CHECK11-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]1185// CHECK11: arraydestroy.body:1186// CHECK11-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP15]], [[OMP_OFFLOAD_CONT]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]1187// CHECK11-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i32 -11188// CHECK11-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR4]]1189// CHECK11-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN]]1190// CHECK11-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE1:%.*]], label [[ARRAYDESTROY_BODY]]1191// CHECK11: arraydestroy.done1:1192// CHECK11-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[TEST]]) #[[ATTR4]]1193// CHECK11-NEXT: [[TMP16:%.*]] = load i32, ptr [[RETVAL]], align 41194// CHECK11-NEXT: ret i32 [[TMP16]]1195//1196//1197// CHECK11-LABEL: define {{[^@]+}}@_ZN2SSC1ERi1198// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(12) [[THIS:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[D:%.*]]) unnamed_addr #[[ATTR1:[0-9]+]] comdat align 2 {1199// CHECK11-NEXT: entry:1200// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41201// CHECK11-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 41202// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41203// CHECK11-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 41204// CHECK11-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41205// CHECK11-NEXT: [[TMP0:%.*]] = load ptr, ptr [[D_ADDR]], align 41206// CHECK11-NEXT: call void @_ZN2SSC2ERi(ptr noundef nonnull align 4 dereferenceable(12) [[THIS1]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP0]])1207// CHECK11-NEXT: ret void1208//1209//1210// CHECK11-LABEL: define {{[^@]+}}@_ZN1SIfEC1Ev1211// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {1212// CHECK11-NEXT: entry:1213// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41214// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41215// CHECK11-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41216// CHECK11-NEXT: call void @_ZN1SIfEC2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]])1217// CHECK11-NEXT: ret void1218//1219//1220// CHECK11-LABEL: define {{[^@]+}}@_ZN1SIfEC1Ef1221// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]], float noundef [[A:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {1222// CHECK11-NEXT: entry:1223// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41224// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca float, align 41225// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41226// CHECK11-NEXT: store float [[A]], ptr [[A_ADDR]], align 41227// CHECK11-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41228// CHECK11-NEXT: [[TMP0:%.*]] = load float, ptr [[A_ADDR]], align 41229// CHECK11-NEXT: call void @_ZN1SIfEC2Ef(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]], float noundef [[TMP0]])1230// CHECK11-NEXT: ret void1231//1232//1233// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l1361234// CHECK11-SAME: () #[[ATTR3:[0-9]+]] {1235// CHECK11-NEXT: entry:1236// CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l136.omp_outlined)1237// CHECK11-NEXT: ret void1238//1239//1240// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l136.omp_outlined1241// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR3]] {1242// CHECK11-NEXT: entry:1243// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41244// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41245// CHECK11-NEXT: [[T_VAR:%.*]] = alloca i32, align 41246// CHECK11-NEXT: [[VEC:%.*]] = alloca [2 x i32], align 41247// CHECK11-NEXT: [[S_ARR:%.*]] = alloca [2 x %struct.S], align 41248// CHECK11-NEXT: [[VAR:%.*]] = alloca [[STRUCT_S:%.*]], align 41249// CHECK11-NEXT: [[SIVAR:%.*]] = alloca i32, align 41250// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41251// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41252// CHECK11-NEXT: [[ARRAY_BEGIN:%.*]] = getelementptr inbounds [2 x %struct.S], ptr [[S_ARR]], i32 0, i32 01253// CHECK11-NEXT: [[ARRAYCTOR_END:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAY_BEGIN]], i32 21254// CHECK11-NEXT: br label [[ARRAYCTOR_LOOP:%.*]]1255// CHECK11: arrayctor.loop:1256// CHECK11-NEXT: [[ARRAYCTOR_CUR:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ]1257// CHECK11-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]])1258// CHECK11-NEXT: [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYCTOR_CUR]], i32 11259// CHECK11-NEXT: [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], [[ARRAYCTOR_END]]1260// CHECK11-NEXT: br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]]1261// CHECK11: arrayctor.cont:1262// CHECK11-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]])1263// CHECK11-NEXT: [[TMP0:%.*]] = load i32, ptr [[T_VAR]], align 41264// CHECK11-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [2 x i32], ptr [[VEC]], i32 0, i32 01265// CHECK11-NEXT: store i32 [[TMP0]], ptr [[ARRAYIDX]], align 41266// CHECK11-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [2 x %struct.S], ptr [[S_ARR]], i32 0, i32 01267// CHECK11-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[ARRAYIDX1]], ptr align 4 [[VAR]], i32 4, i1 false)1268// CHECK11-NEXT: store i32 3, ptr [[SIVAR]], align 41269// CHECK11-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]]) #[[ATTR4]]1270// CHECK11-NEXT: [[ARRAY_BEGIN2:%.*]] = getelementptr inbounds [2 x %struct.S], ptr [[S_ARR]], i32 0, i32 01271// CHECK11-NEXT: [[TMP1:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAY_BEGIN2]], i32 21272// CHECK11-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]1273// CHECK11: arraydestroy.body:1274// CHECK11-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP1]], [[ARRAYCTOR_CONT]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]1275// CHECK11-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i32 -11276// CHECK11-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR4]]1277// CHECK11-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN2]]1278// CHECK11-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE3:%.*]], label [[ARRAYDESTROY_BODY]]1279// CHECK11: arraydestroy.done3:1280// CHECK11-NEXT: ret void1281//1282//1283// CHECK11-LABEL: define {{[^@]+}}@_ZN1SIfED1Ev1284// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {1285// CHECK11-NEXT: entry:1286// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41287// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41288// CHECK11-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41289// CHECK11-NEXT: call void @_ZN1SIfED2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]]) #[[ATTR4]]1290// CHECK11-NEXT: ret void1291//1292//1293// CHECK11-LABEL: define {{[^@]+}}@_Z5tmainIiET_v1294// CHECK11-SAME: () #[[ATTR1]] comdat {1295// CHECK11-NEXT: entry:1296// CHECK11-NEXT: [[RETVAL:%.*]] = alloca i32, align 41297// CHECK11-NEXT: [[TEST:%.*]] = alloca [[STRUCT_S_0:%.*]], align 41298// CHECK11-NEXT: [[SST:%.*]] = alloca [[STRUCT_SST:%.*]], align 41299// CHECK11-NEXT: [[T_VAR:%.*]] = alloca i32, align 1281300// CHECK11-NEXT: [[VEC:%.*]] = alloca [2 x i32], align 1281301// CHECK11-NEXT: [[S_ARR:%.*]] = alloca [2 x %struct.S.0], align 1281302// CHECK11-NEXT: [[VAR:%.*]] = alloca [[STRUCT_S_0]], align 1281303// CHECK11-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81304// CHECK11-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[TEST]])1305// CHECK11-NEXT: call void @_ZN3SSTIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[SST]])1306// CHECK11-NEXT: store i32 0, ptr [[T_VAR]], align 1281307// CHECK11-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 128 [[VEC]], ptr align 128 @__const._Z5tmainIiET_v.vec, i32 8, i1 false)1308// CHECK11-NEXT: call void @_ZN1SIiEC1Ei(ptr noundef nonnull align 4 dereferenceable(4) [[S_ARR]], i32 noundef 1)1309// CHECK11-NEXT: [[ARRAYINIT_ELEMENT:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[S_ARR]], i32 11310// CHECK11-NEXT: call void @_ZN1SIiEC1Ei(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYINIT_ELEMENT]], i32 noundef 2)1311// CHECK11-NEXT: call void @_ZN1SIiEC1Ei(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]], i32 noundef 3)1312// CHECK11-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01313// CHECK11-NEXT: store i32 3, ptr [[TMP0]], align 41314// CHECK11-NEXT: [[TMP1:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11315// CHECK11-NEXT: store i32 0, ptr [[TMP1]], align 41316// CHECK11-NEXT: [[TMP2:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21317// CHECK11-NEXT: store ptr null, ptr [[TMP2]], align 41318// CHECK11-NEXT: [[TMP3:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31319// CHECK11-NEXT: store ptr null, ptr [[TMP3]], align 41320// CHECK11-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41321// CHECK11-NEXT: store ptr null, ptr [[TMP4]], align 41322// CHECK11-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51323// CHECK11-NEXT: store ptr null, ptr [[TMP5]], align 41324// CHECK11-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61325// CHECK11-NEXT: store ptr null, ptr [[TMP6]], align 41326// CHECK11-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71327// CHECK11-NEXT: store ptr null, ptr [[TMP7]], align 41328// CHECK11-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81329// CHECK11-NEXT: store i64 0, ptr [[TMP8]], align 81330// CHECK11-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91331// CHECK11-NEXT: store i64 0, ptr [[TMP9]], align 81332// CHECK11-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101333// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP10]], align 41334// CHECK11-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111335// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP11]], align 41336// CHECK11-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121337// CHECK11-NEXT: store i32 0, ptr [[TMP12]], align 41338// CHECK11-NEXT: [[TMP13:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l86.region_id, ptr [[KERNEL_ARGS]])1339// CHECK11-NEXT: [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 01340// CHECK11-NEXT: br i1 [[TMP14]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1341// CHECK11: omp_offload.failed:1342// CHECK11-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l86() #[[ATTR4]]1343// CHECK11-NEXT: br label [[OMP_OFFLOAD_CONT]]1344// CHECK11: omp_offload.cont:1345// CHECK11-NEXT: store i32 0, ptr [[RETVAL]], align 41346// CHECK11-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]]) #[[ATTR4]]1347// CHECK11-NEXT: [[ARRAY_BEGIN:%.*]] = getelementptr inbounds [2 x %struct.S.0], ptr [[S_ARR]], i32 0, i32 01348// CHECK11-NEXT: [[TMP15:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAY_BEGIN]], i32 21349// CHECK11-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]1350// CHECK11: arraydestroy.body:1351// CHECK11-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP15]], [[OMP_OFFLOAD_CONT]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]1352// CHECK11-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i32 -11353// CHECK11-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR4]]1354// CHECK11-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN]]1355// CHECK11-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE1:%.*]], label [[ARRAYDESTROY_BODY]]1356// CHECK11: arraydestroy.done1:1357// CHECK11-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[TEST]]) #[[ATTR4]]1358// CHECK11-NEXT: [[TMP16:%.*]] = load i32, ptr [[RETVAL]], align 41359// CHECK11-NEXT: ret i32 [[TMP16]]1360//1361//1362// CHECK11-LABEL: define {{[^@]+}}@_ZN2SSC2ERi1363// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(12) [[THIS:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[D:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {1364// CHECK11-NEXT: entry:1365// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41366// CHECK11-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 41367// CHECK11-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 41368// CHECK11-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 41369// CHECK11-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 41370// CHECK11-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81371// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41372// CHECK11-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 41373// CHECK11-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41374// CHECK11-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[THIS1]], i32 0, i32 01375// CHECK11-NEXT: store i32 0, ptr [[A]], align 41376// CHECK11-NEXT: [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 11377// CHECK11-NEXT: [[BF_LOAD:%.*]] = load i8, ptr [[B]], align 41378// CHECK11-NEXT: [[BF_CLEAR:%.*]] = and i8 [[BF_LOAD]], -161379// CHECK11-NEXT: [[BF_SET:%.*]] = or i8 [[BF_CLEAR]], 01380// CHECK11-NEXT: store i8 [[BF_SET]], ptr [[B]], align 41381// CHECK11-NEXT: [[C:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 21382// CHECK11-NEXT: [[TMP0:%.*]] = load ptr, ptr [[D_ADDR]], align 41383// CHECK11-NEXT: store ptr [[TMP0]], ptr [[C]], align 41384// CHECK11-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01385// CHECK11-NEXT: store ptr [[THIS1]], ptr [[TMP1]], align 41386// CHECK11-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01387// CHECK11-NEXT: store ptr [[THIS1]], ptr [[TMP2]], align 41388// CHECK11-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01389// CHECK11-NEXT: store ptr null, ptr [[TMP3]], align 41390// CHECK11-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01391// CHECK11-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01392// CHECK11-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01393// CHECK11-NEXT: store i32 3, ptr [[TMP6]], align 41394// CHECK11-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11395// CHECK11-NEXT: store i32 1, ptr [[TMP7]], align 41396// CHECK11-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21397// CHECK11-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 41398// CHECK11-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31399// CHECK11-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 41400// CHECK11-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41401// CHECK11-NEXT: store ptr @.offload_sizes, ptr [[TMP10]], align 41402// CHECK11-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51403// CHECK11-NEXT: store ptr @.offload_maptypes, ptr [[TMP11]], align 41404// CHECK11-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61405// CHECK11-NEXT: store ptr null, ptr [[TMP12]], align 41406// CHECK11-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71407// CHECK11-NEXT: store ptr null, ptr [[TMP13]], align 41408// CHECK11-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81409// CHECK11-NEXT: store i64 0, ptr [[TMP14]], align 81410// CHECK11-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91411// CHECK11-NEXT: store i64 0, ptr [[TMP15]], align 81412// CHECK11-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101413// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 41414// CHECK11-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111415// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 41416// CHECK11-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121417// CHECK11-NEXT: store i32 0, ptr [[TMP18]], align 41418// CHECK11-NEXT: [[TMP19:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSC1ERi_l48.region_id, ptr [[KERNEL_ARGS]])1419// CHECK11-NEXT: [[TMP20:%.*]] = icmp ne i32 [[TMP19]], 01420// CHECK11-NEXT: br i1 [[TMP20]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1421// CHECK11: omp_offload.failed:1422// CHECK11-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSC1ERi_l48(ptr [[THIS1]]) #[[ATTR4]]1423// CHECK11-NEXT: br label [[OMP_OFFLOAD_CONT]]1424// CHECK11: omp_offload.cont:1425// CHECK11-NEXT: ret void1426//1427//1428// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSC1ERi_l481429// CHECK11-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR3]] {1430// CHECK11-NEXT: entry:1431// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41432// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41433// CHECK11-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41434// CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSC1ERi_l48.omp_outlined, ptr [[TMP0]])1435// CHECK11-NEXT: ret void1436//1437//1438// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSC1ERi_l48.omp_outlined1439// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR3]] {1440// CHECK11-NEXT: entry:1441// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41442// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41443// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41444// CHECK11-NEXT: [[A:%.*]] = alloca i32, align 41445// CHECK11-NEXT: [[TMP:%.*]] = alloca ptr, align 41446// CHECK11-NEXT: [[B:%.*]] = alloca i32, align 41447// CHECK11-NEXT: [[C:%.*]] = alloca i32, align 41448// CHECK11-NEXT: [[_TMP1:%.*]] = alloca ptr, align 41449// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41450// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41451// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41452// CHECK11-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41453// CHECK11-NEXT: store ptr [[A]], ptr [[TMP]], align 41454// CHECK11-NEXT: store ptr [[C]], ptr [[_TMP1]], align 41455// CHECK11-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP]], align 41456// CHECK11-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41457// CHECK11-NEXT: [[INC:%.*]] = add nsw i32 [[TMP2]], 11458// CHECK11-NEXT: store i32 [[INC]], ptr [[TMP1]], align 41459// CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[B]], align 41460// CHECK11-NEXT: [[DEC:%.*]] = add nsw i32 [[TMP3]], -11461// CHECK11-NEXT: store i32 [[DEC]], ptr [[B]], align 41462// CHECK11-NEXT: [[TMP4:%.*]] = load ptr, ptr [[_TMP1]], align 41463// CHECK11-NEXT: [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 41464// CHECK11-NEXT: [[DIV:%.*]] = sdiv i32 [[TMP5]], 11465// CHECK11-NEXT: store i32 [[DIV]], ptr [[TMP4]], align 41466// CHECK11-NEXT: ret void1467//1468//1469// CHECK11-LABEL: define {{[^@]+}}@_ZN1SIfEC2Ev1470// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {1471// CHECK11-NEXT: entry:1472// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41473// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41474// CHECK11-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41475// CHECK11-NEXT: [[F:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[THIS1]], i32 0, i32 01476// CHECK11-NEXT: store float 0.000000e+00, ptr [[F]], align 41477// CHECK11-NEXT: ret void1478//1479//1480// CHECK11-LABEL: define {{[^@]+}}@_ZN1SIfEC2Ef1481// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]], float noundef [[A:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {1482// CHECK11-NEXT: entry:1483// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41484// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca float, align 41485// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41486// CHECK11-NEXT: store float [[A]], ptr [[A_ADDR]], align 41487// CHECK11-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41488// CHECK11-NEXT: [[F:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[THIS1]], i32 0, i32 01489// CHECK11-NEXT: [[TMP0:%.*]] = load float, ptr [[A_ADDR]], align 41490// CHECK11-NEXT: store float [[TMP0]], ptr [[F]], align 41491// CHECK11-NEXT: ret void1492//1493//1494// CHECK11-LABEL: define {{[^@]+}}@_ZN1SIfED2Ev1495// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {1496// CHECK11-NEXT: entry:1497// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41498// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41499// CHECK11-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41500// CHECK11-NEXT: ret void1501//1502//1503// CHECK11-LABEL: define {{[^@]+}}@_ZN1SIiEC1Ev1504// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {1505// CHECK11-NEXT: entry:1506// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41507// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41508// CHECK11-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41509// CHECK11-NEXT: call void @_ZN1SIiEC2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]])1510// CHECK11-NEXT: ret void1511//1512//1513// CHECK11-LABEL: define {{[^@]+}}@_ZN3SSTIiEC1Ev1514// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {1515// CHECK11-NEXT: entry:1516// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41517// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41518// CHECK11-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41519// CHECK11-NEXT: call void @_ZN3SSTIiEC2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]])1520// CHECK11-NEXT: ret void1521//1522//1523// CHECK11-LABEL: define {{[^@]+}}@_ZN1SIiEC1Ei1524// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]], i32 noundef [[A:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {1525// CHECK11-NEXT: entry:1526// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41527// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca i32, align 41528// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41529// CHECK11-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 41530// CHECK11-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41531// CHECK11-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 41532// CHECK11-NEXT: call void @_ZN1SIiEC2Ei(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]], i32 noundef [[TMP0]])1533// CHECK11-NEXT: ret void1534//1535//1536// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l861537// CHECK11-SAME: () #[[ATTR3]] {1538// CHECK11-NEXT: entry:1539// CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l86.omp_outlined)1540// CHECK11-NEXT: ret void1541//1542//1543// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l86.omp_outlined1544// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR3]] {1545// CHECK11-NEXT: entry:1546// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41547// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41548// CHECK11-NEXT: [[T_VAR:%.*]] = alloca i32, align 1281549// CHECK11-NEXT: [[VEC:%.*]] = alloca [2 x i32], align 1281550// CHECK11-NEXT: [[S_ARR:%.*]] = alloca [2 x %struct.S.0], align 1281551// CHECK11-NEXT: [[VAR:%.*]] = alloca [[STRUCT_S_0:%.*]], align 1281552// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41553// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41554// CHECK11-NEXT: [[ARRAY_BEGIN:%.*]] = getelementptr inbounds [2 x %struct.S.0], ptr [[S_ARR]], i32 0, i32 01555// CHECK11-NEXT: [[ARRAYCTOR_END:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAY_BEGIN]], i32 21556// CHECK11-NEXT: br label [[ARRAYCTOR_LOOP:%.*]]1557// CHECK11: arrayctor.loop:1558// CHECK11-NEXT: [[ARRAYCTOR_CUR:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ]1559// CHECK11-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]])1560// CHECK11-NEXT: [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYCTOR_CUR]], i32 11561// CHECK11-NEXT: [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], [[ARRAYCTOR_END]]1562// CHECK11-NEXT: br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]]1563// CHECK11: arrayctor.cont:1564// CHECK11-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]])1565// CHECK11-NEXT: [[TMP0:%.*]] = load i32, ptr [[T_VAR]], align 1281566// CHECK11-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [2 x i32], ptr [[VEC]], i32 0, i32 01567// CHECK11-NEXT: store i32 [[TMP0]], ptr [[ARRAYIDX]], align 1281568// CHECK11-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [2 x %struct.S.0], ptr [[S_ARR]], i32 0, i32 01569// CHECK11-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 128 [[ARRAYIDX1]], ptr align 128 [[VAR]], i32 4, i1 false)1570// CHECK11-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]]) #[[ATTR4]]1571// CHECK11-NEXT: [[ARRAY_BEGIN2:%.*]] = getelementptr inbounds [2 x %struct.S.0], ptr [[S_ARR]], i32 0, i32 01572// CHECK11-NEXT: [[TMP1:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAY_BEGIN2]], i32 21573// CHECK11-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]1574// CHECK11: arraydestroy.body:1575// CHECK11-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP1]], [[ARRAYCTOR_CONT]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]1576// CHECK11-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i32 -11577// CHECK11-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR4]]1578// CHECK11-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN2]]1579// CHECK11-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE3:%.*]], label [[ARRAYDESTROY_BODY]]1580// CHECK11: arraydestroy.done3:1581// CHECK11-NEXT: ret void1582//1583//1584// CHECK11-LABEL: define {{[^@]+}}@_ZN1SIiED1Ev1585// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {1586// CHECK11-NEXT: entry:1587// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41588// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41589// CHECK11-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41590// CHECK11-NEXT: call void @_ZN1SIiED2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]]) #[[ATTR4]]1591// CHECK11-NEXT: ret void1592//1593//1594// CHECK11-LABEL: define {{[^@]+}}@_ZN1SIiEC2Ev1595// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {1596// CHECK11-NEXT: entry:1597// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41598// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41599// CHECK11-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41600// CHECK11-NEXT: [[F:%.*]] = getelementptr inbounds nuw [[STRUCT_S_0:%.*]], ptr [[THIS1]], i32 0, i32 01601// CHECK11-NEXT: store i32 0, ptr [[F]], align 41602// CHECK11-NEXT: ret void1603//1604//1605// CHECK11-LABEL: define {{[^@]+}}@_ZN3SSTIiEC2Ev1606// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {1607// CHECK11-NEXT: entry:1608// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41609// CHECK11-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 41610// CHECK11-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 41611// CHECK11-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 41612// CHECK11-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81613// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41614// CHECK11-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41615// CHECK11-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SST:%.*]], ptr [[THIS1]], i32 0, i32 01616// CHECK11-NEXT: store i32 0, ptr [[A]], align 41617// CHECK11-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01618// CHECK11-NEXT: store ptr [[THIS1]], ptr [[TMP0]], align 41619// CHECK11-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01620// CHECK11-NEXT: store ptr [[THIS1]], ptr [[TMP1]], align 41621// CHECK11-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01622// CHECK11-NEXT: store ptr null, ptr [[TMP2]], align 41623// CHECK11-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01624// CHECK11-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01625// CHECK11-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01626// CHECK11-NEXT: store i32 3, ptr [[TMP5]], align 41627// CHECK11-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11628// CHECK11-NEXT: store i32 1, ptr [[TMP6]], align 41629// CHECK11-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21630// CHECK11-NEXT: store ptr [[TMP3]], ptr [[TMP7]], align 41631// CHECK11-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31632// CHECK11-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 41633// CHECK11-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41634// CHECK11-NEXT: store ptr @.offload_sizes.1, ptr [[TMP9]], align 41635// CHECK11-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51636// CHECK11-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP10]], align 41637// CHECK11-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61638// CHECK11-NEXT: store ptr null, ptr [[TMP11]], align 41639// CHECK11-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71640// CHECK11-NEXT: store ptr null, ptr [[TMP12]], align 41641// CHECK11-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81642// CHECK11-NEXT: store i64 0, ptr [[TMP13]], align 81643// CHECK11-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91644// CHECK11-NEXT: store i64 0, ptr [[TMP14]], align 81645// CHECK11-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101646// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP15]], align 41647// CHECK11-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111648// CHECK11-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 41649// CHECK11-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121650// CHECK11-NEXT: store i32 0, ptr [[TMP17]], align 41651// CHECK11-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN3SSTIiEC1Ev_l64.region_id, ptr [[KERNEL_ARGS]])1652// CHECK11-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 01653// CHECK11-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1654// CHECK11: omp_offload.failed:1655// CHECK11-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN3SSTIiEC1Ev_l64(ptr [[THIS1]]) #[[ATTR4]]1656// CHECK11-NEXT: br label [[OMP_OFFLOAD_CONT]]1657// CHECK11: omp_offload.cont:1658// CHECK11-NEXT: ret void1659//1660//1661// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN3SSTIiEC1Ev_l641662// CHECK11-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR3]] {1663// CHECK11-NEXT: entry:1664// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41665// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41666// CHECK11-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41667// CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN3SSTIiEC1Ev_l64.omp_outlined, ptr [[TMP0]])1668// CHECK11-NEXT: ret void1669//1670//1671// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN3SSTIiEC1Ev_l64.omp_outlined1672// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR3]] {1673// CHECK11-NEXT: entry:1674// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41675// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41676// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41677// CHECK11-NEXT: [[A:%.*]] = alloca i32, align 41678// CHECK11-NEXT: [[TMP:%.*]] = alloca ptr, align 41679// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41680// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41681// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41682// CHECK11-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41683// CHECK11-NEXT: store ptr [[A]], ptr [[TMP]], align 41684// CHECK11-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP]], align 41685// CHECK11-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41686// CHECK11-NEXT: [[INC:%.*]] = add nsw i32 [[TMP2]], 11687// CHECK11-NEXT: store i32 [[INC]], ptr [[TMP1]], align 41688// CHECK11-NEXT: ret void1689//1690//1691// CHECK11-LABEL: define {{[^@]+}}@_ZN1SIiEC2Ei1692// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]], i32 noundef [[A:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {1693// CHECK11-NEXT: entry:1694// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41695// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca i32, align 41696// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41697// CHECK11-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 41698// CHECK11-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41699// CHECK11-NEXT: [[F:%.*]] = getelementptr inbounds nuw [[STRUCT_S_0:%.*]], ptr [[THIS1]], i32 0, i32 01700// CHECK11-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 41701// CHECK11-NEXT: store i32 [[TMP0]], ptr [[F]], align 41702// CHECK11-NEXT: ret void1703//1704//1705// CHECK11-LABEL: define {{[^@]+}}@_ZN1SIiED2Ev1706// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {1707// CHECK11-NEXT: entry:1708// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41709// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41710// CHECK11-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41711// CHECK11-NEXT: ret void1712//1713