brintos / llvm-project-archived public Read only

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// REQUIRES: powerpc-registered-target3// REQUIRES: nvptx-registered-target4 5// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK16// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc7// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple nvptx64-nvidia-cuda -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK28// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple nvptx64-nvidia-cuda -fopenmp-targets=nvptx64-nvidia-cuda %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t9// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple nvptx64-nvidia-cuda -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - | FileCheck %s --check-prefix=CHECK310// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple nvptx64-nvidia-cuda -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - | FileCheck %s --check-prefix=CHECK311 12// expected-no-diagnostics13#ifndef HEADER14#define HEADER15 16template <typename T>17int foo(const T &t) {18  #pragma omp target parallel19  t();20  return 0;21}22 23struct S {24  int a = 15;25  int foo() {26    auto &&L = [&]() { return a; };27#pragma omp target28    L();29#pragma omp target parallel30    L();31    return a + ::foo(L);32  }33} s;34 35int main(int argc, char **argv) {36  int &b = argc;37  int &&c = 1;38  int *d = &argc;39  int a;40  auto &&L = [&]() { return argc + b + c + reinterpret_cast<long int>(d) + a; };41#pragma omp target firstprivate(argc) map(to : a)42  L();43#pragma omp target parallel44  L();45  return argc + s.foo();46}47 48#endif // HEADER49// CHECK1-LABEL: define {{[^@]+}}@main50// CHECK1-SAME: (i32 noundef signext [[ARGC:%.*]], ptr noundef [[ARGV:%.*]]) #[[ATTR0:[0-9]+]] {51// CHECK1-NEXT:  entry:52// CHECK1-NEXT:    [[RETVAL:%.*]] = alloca i32, align 453// CHECK1-NEXT:    [[ARGC_ADDR:%.*]] = alloca i32, align 454// CHECK1-NEXT:    [[ARGV_ADDR:%.*]] = alloca ptr, align 855// CHECK1-NEXT:    [[B:%.*]] = alloca ptr, align 856// CHECK1-NEXT:    [[C:%.*]] = alloca ptr, align 857// CHECK1-NEXT:    [[REF_TMP:%.*]] = alloca i32, align 458// CHECK1-NEXT:    [[D:%.*]] = alloca ptr, align 859// CHECK1-NEXT:    [[A:%.*]] = alloca i32, align 460// CHECK1-NEXT:    [[L:%.*]] = alloca ptr, align 861// CHECK1-NEXT:    [[REF_TMP1:%.*]] = alloca [[CLASS_ANON:%.*]], align 862// CHECK1-NEXT:    [[TMP:%.*]] = alloca ptr, align 863// CHECK1-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 864// CHECK1-NEXT:    [[_TMP3:%.*]] = alloca ptr, align 865// CHECK1-NEXT:    [[ARGC_CASTED:%.*]] = alloca i64, align 866// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [11 x ptr], align 867// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [11 x ptr], align 868// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [11 x ptr], align 869// CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 870// CHECK1-NEXT:    [[_TMP4:%.*]] = alloca ptr, align 871// CHECK1-NEXT:    [[_TMP5:%.*]] = alloca ptr, align 872// CHECK1-NEXT:    [[_TMP6:%.*]] = alloca ptr, align 873// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS7:%.*]] = alloca [11 x ptr], align 874// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS8:%.*]] = alloca [11 x ptr], align 875// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS9:%.*]] = alloca [11 x ptr], align 876// CHECK1-NEXT:    [[KERNEL_ARGS10:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 877// CHECK1-NEXT:    store i32 0, ptr [[RETVAL]], align 478// CHECK1-NEXT:    store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 479// CHECK1-NEXT:    store ptr [[ARGV]], ptr [[ARGV_ADDR]], align 880// CHECK1-NEXT:    store ptr [[ARGC_ADDR]], ptr [[B]], align 881// CHECK1-NEXT:    store i32 1, ptr [[REF_TMP]], align 482// CHECK1-NEXT:    store ptr [[REF_TMP]], ptr [[C]], align 883// CHECK1-NEXT:    store ptr [[ARGC_ADDR]], ptr [[D]], align 884// CHECK1-NEXT:    [[TMP0:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[REF_TMP1]], i32 0, i32 085// CHECK1-NEXT:    store ptr [[ARGC_ADDR]], ptr [[TMP0]], align 886// CHECK1-NEXT:    [[TMP1:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[REF_TMP1]], i32 0, i32 187// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[B]], align 888// CHECK1-NEXT:    store ptr [[TMP2]], ptr [[TMP1]], align 889// CHECK1-NEXT:    [[TMP3:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[REF_TMP1]], i32 0, i32 290// CHECK1-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[C]], align 891// CHECK1-NEXT:    store ptr [[TMP4]], ptr [[TMP3]], align 892// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[REF_TMP1]], i32 0, i32 393// CHECK1-NEXT:    store ptr [[D]], ptr [[TMP5]], align 894// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[REF_TMP1]], i32 0, i32 495// CHECK1-NEXT:    store ptr [[A]], ptr [[TMP6]], align 896// CHECK1-NEXT:    store ptr [[REF_TMP1]], ptr [[L]], align 897// CHECK1-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[B]], align 898// CHECK1-NEXT:    store ptr [[TMP7]], ptr [[TMP]], align 899// CHECK1-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[C]], align 8100// CHECK1-NEXT:    store ptr [[TMP8]], ptr [[_TMP2]], align 8101// CHECK1-NEXT:    [[TMP9:%.*]] = load ptr, ptr [[L]], align 8102// CHECK1-NEXT:    store ptr [[TMP9]], ptr [[_TMP3]], align 8103// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[ARGC_ADDR]], align 4104// CHECK1-NEXT:    store i32 [[TMP10]], ptr [[ARGC_CASTED]], align 4105// CHECK1-NEXT:    [[TMP11:%.*]] = load i64, ptr [[ARGC_CASTED]], align 8106// CHECK1-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[TMP]], align 8107// CHECK1-NEXT:    [[TMP13:%.*]] = load ptr, ptr [[_TMP2]], align 8108// CHECK1-NEXT:    [[TMP14:%.*]] = load ptr, ptr [[D]], align 8109// CHECK1-NEXT:    [[TMP15:%.*]] = load ptr, ptr [[_TMP3]], align 8110// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP15]], i32 0, i32 0111// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP15]], i32 0, i32 0112// CHECK1-NEXT:    [[TMP18:%.*]] = load ptr, ptr [[TMP17]], align 8113// CHECK1-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP15]], i32 0, i32 1114// CHECK1-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP15]], i32 0, i32 1115// CHECK1-NEXT:    [[TMP21:%.*]] = load ptr, ptr [[TMP20]], align 8116// CHECK1-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP15]], i32 0, i32 2117// CHECK1-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP15]], i32 0, i32 2118// CHECK1-NEXT:    [[TMP24:%.*]] = load ptr, ptr [[TMP23]], align 8119// CHECK1-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP15]], i32 0, i32 3120// CHECK1-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP15]], i32 0, i32 3121// CHECK1-NEXT:    [[TMP27:%.*]] = load ptr, ptr [[TMP26]], align 8122// CHECK1-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP15]], i32 0, i32 4123// CHECK1-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP15]], i32 0, i32 4124// CHECK1-NEXT:    [[TMP30:%.*]] = load ptr, ptr [[TMP29]], align 8125// CHECK1-NEXT:    [[TMP31:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0126// CHECK1-NEXT:    store i64 [[TMP11]], ptr [[TMP31]], align 8127// CHECK1-NEXT:    [[TMP32:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0128// CHECK1-NEXT:    store i64 [[TMP11]], ptr [[TMP32]], align 8129// CHECK1-NEXT:    [[TMP33:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0130// CHECK1-NEXT:    store ptr null, ptr [[TMP33]], align 8131// CHECK1-NEXT:    [[TMP34:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1132// CHECK1-NEXT:    store ptr [[TMP12]], ptr [[TMP34]], align 8133// CHECK1-NEXT:    [[TMP35:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1134// CHECK1-NEXT:    store ptr [[TMP12]], ptr [[TMP35]], align 8135// CHECK1-NEXT:    [[TMP36:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1136// CHECK1-NEXT:    store ptr null, ptr [[TMP36]], align 8137// CHECK1-NEXT:    [[TMP37:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2138// CHECK1-NEXT:    store ptr [[TMP13]], ptr [[TMP37]], align 8139// CHECK1-NEXT:    [[TMP38:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2140// CHECK1-NEXT:    store ptr [[TMP13]], ptr [[TMP38]], align 8141// CHECK1-NEXT:    [[TMP39:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2142// CHECK1-NEXT:    store ptr null, ptr [[TMP39]], align 8143// CHECK1-NEXT:    [[TMP40:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3144// CHECK1-NEXT:    store ptr [[TMP14]], ptr [[TMP40]], align 8145// CHECK1-NEXT:    [[TMP41:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3146// CHECK1-NEXT:    store ptr [[TMP14]], ptr [[TMP41]], align 8147// CHECK1-NEXT:    [[TMP42:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3148// CHECK1-NEXT:    store ptr null, ptr [[TMP42]], align 8149// CHECK1-NEXT:    [[TMP43:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 4150// CHECK1-NEXT:    store ptr [[A]], ptr [[TMP43]], align 8151// CHECK1-NEXT:    [[TMP44:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 4152// CHECK1-NEXT:    store ptr [[A]], ptr [[TMP44]], align 8153// CHECK1-NEXT:    [[TMP45:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 4154// CHECK1-NEXT:    store ptr null, ptr [[TMP45]], align 8155// CHECK1-NEXT:    [[TMP46:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 5156// CHECK1-NEXT:    store ptr [[TMP15]], ptr [[TMP46]], align 8157// CHECK1-NEXT:    [[TMP47:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 5158// CHECK1-NEXT:    store ptr [[TMP15]], ptr [[TMP47]], align 8159// CHECK1-NEXT:    [[TMP48:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 5160// CHECK1-NEXT:    store ptr null, ptr [[TMP48]], align 8161// CHECK1-NEXT:    [[TMP49:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 6162// CHECK1-NEXT:    store ptr [[TMP16]], ptr [[TMP49]], align 8163// CHECK1-NEXT:    [[TMP50:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 6164// CHECK1-NEXT:    store ptr [[TMP18]], ptr [[TMP50]], align 8165// CHECK1-NEXT:    [[TMP51:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 6166// CHECK1-NEXT:    store ptr null, ptr [[TMP51]], align 8167// CHECK1-NEXT:    [[TMP52:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 7168// CHECK1-NEXT:    store ptr [[TMP19]], ptr [[TMP52]], align 8169// CHECK1-NEXT:    [[TMP53:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 7170// CHECK1-NEXT:    store ptr [[TMP21]], ptr [[TMP53]], align 8171// CHECK1-NEXT:    [[TMP54:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 7172// CHECK1-NEXT:    store ptr null, ptr [[TMP54]], align 8173// CHECK1-NEXT:    [[TMP55:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 8174// CHECK1-NEXT:    store ptr [[TMP22]], ptr [[TMP55]], align 8175// CHECK1-NEXT:    [[TMP56:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 8176// CHECK1-NEXT:    store ptr [[TMP24]], ptr [[TMP56]], align 8177// CHECK1-NEXT:    [[TMP57:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 8178// CHECK1-NEXT:    store ptr null, ptr [[TMP57]], align 8179// CHECK1-NEXT:    [[TMP58:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 9180// CHECK1-NEXT:    store ptr [[TMP25]], ptr [[TMP58]], align 8181// CHECK1-NEXT:    [[TMP59:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 9182// CHECK1-NEXT:    store ptr [[TMP27]], ptr [[TMP59]], align 8183// CHECK1-NEXT:    [[TMP60:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 9184// CHECK1-NEXT:    store ptr null, ptr [[TMP60]], align 8185// CHECK1-NEXT:    [[TMP61:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 10186// CHECK1-NEXT:    store ptr [[TMP28]], ptr [[TMP61]], align 8187// CHECK1-NEXT:    [[TMP62:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 10188// CHECK1-NEXT:    store ptr [[TMP30]], ptr [[TMP62]], align 8189// CHECK1-NEXT:    [[TMP63:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 10190// CHECK1-NEXT:    store ptr null, ptr [[TMP63]], align 8191// CHECK1-NEXT:    [[TMP64:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0192// CHECK1-NEXT:    [[TMP65:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0193// CHECK1-NEXT:    [[TMP66:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0194// CHECK1-NEXT:    store i32 3, ptr [[TMP66]], align 4195// CHECK1-NEXT:    [[TMP67:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1196// CHECK1-NEXT:    store i32 11, ptr [[TMP67]], align 4197// CHECK1-NEXT:    [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2198// CHECK1-NEXT:    store ptr [[TMP64]], ptr [[TMP68]], align 8199// CHECK1-NEXT:    [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3200// CHECK1-NEXT:    store ptr [[TMP65]], ptr [[TMP69]], align 8201// CHECK1-NEXT:    [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4202// CHECK1-NEXT:    store ptr @.offload_sizes, ptr [[TMP70]], align 8203// CHECK1-NEXT:    [[TMP71:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5204// CHECK1-NEXT:    store ptr @.offload_maptypes, ptr [[TMP71]], align 8205// CHECK1-NEXT:    [[TMP72:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6206// CHECK1-NEXT:    store ptr null, ptr [[TMP72]], align 8207// CHECK1-NEXT:    [[TMP73:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7208// CHECK1-NEXT:    store ptr null, ptr [[TMP73]], align 8209// CHECK1-NEXT:    [[TMP74:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8210// CHECK1-NEXT:    store i64 0, ptr [[TMP74]], align 8211// CHECK1-NEXT:    [[TMP75:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9212// CHECK1-NEXT:    store i64 0, ptr [[TMP75]], align 8213// CHECK1-NEXT:    [[TMP76:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10214// CHECK1-NEXT:    store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP76]], align 4215// CHECK1-NEXT:    [[TMP77:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11216// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP77]], align 4217// CHECK1-NEXT:    [[TMP78:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12218// CHECK1-NEXT:    store i32 0, ptr [[TMP78]], align 4219// CHECK1-NEXT:    [[TMP79:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l41.region_id, ptr [[KERNEL_ARGS]])220// CHECK1-NEXT:    [[TMP80:%.*]] = icmp ne i32 [[TMP79]], 0221// CHECK1-NEXT:    br i1 [[TMP80]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]222// CHECK1:       omp_offload.failed:223// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l41(i64 [[TMP11]], ptr [[TMP12]], ptr [[TMP13]], ptr [[TMP14]], ptr [[A]], ptr [[TMP15]]) #[[ATTR4:[0-9]+]]224// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]225// CHECK1:       omp_offload.cont:226// CHECK1-NEXT:    [[TMP81:%.*]] = load ptr, ptr [[B]], align 8227// CHECK1-NEXT:    store ptr [[TMP81]], ptr [[_TMP4]], align 8228// CHECK1-NEXT:    [[TMP82:%.*]] = load ptr, ptr [[C]], align 8229// CHECK1-NEXT:    store ptr [[TMP82]], ptr [[_TMP5]], align 8230// CHECK1-NEXT:    [[TMP83:%.*]] = load ptr, ptr [[L]], align 8231// CHECK1-NEXT:    store ptr [[TMP83]], ptr [[_TMP6]], align 8232// CHECK1-NEXT:    [[TMP84:%.*]] = load ptr, ptr [[_TMP4]], align 8233// CHECK1-NEXT:    [[TMP85:%.*]] = load ptr, ptr [[_TMP5]], align 8234// CHECK1-NEXT:    [[TMP86:%.*]] = load ptr, ptr [[D]], align 8235// CHECK1-NEXT:    [[TMP87:%.*]] = load ptr, ptr [[_TMP6]], align 8236// CHECK1-NEXT:    [[TMP88:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP87]], i32 0, i32 0237// CHECK1-NEXT:    [[TMP89:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP87]], i32 0, i32 0238// CHECK1-NEXT:    [[TMP90:%.*]] = load ptr, ptr [[TMP89]], align 8239// CHECK1-NEXT:    [[TMP91:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP87]], i32 0, i32 1240// CHECK1-NEXT:    [[TMP92:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP87]], i32 0, i32 1241// CHECK1-NEXT:    [[TMP93:%.*]] = load ptr, ptr [[TMP92]], align 8242// CHECK1-NEXT:    [[TMP94:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP87]], i32 0, i32 2243// CHECK1-NEXT:    [[TMP95:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP87]], i32 0, i32 2244// CHECK1-NEXT:    [[TMP96:%.*]] = load ptr, ptr [[TMP95]], align 8245// CHECK1-NEXT:    [[TMP97:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP87]], i32 0, i32 3246// CHECK1-NEXT:    [[TMP98:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP87]], i32 0, i32 3247// CHECK1-NEXT:    [[TMP99:%.*]] = load ptr, ptr [[TMP98]], align 8248// CHECK1-NEXT:    [[TMP100:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP87]], i32 0, i32 4249// CHECK1-NEXT:    [[TMP101:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP87]], i32 0, i32 4250// CHECK1-NEXT:    [[TMP102:%.*]] = load ptr, ptr [[TMP101]], align 8251// CHECK1-NEXT:    [[TMP103:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 0252// CHECK1-NEXT:    store ptr [[ARGC_ADDR]], ptr [[TMP103]], align 8253// CHECK1-NEXT:    [[TMP104:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 0254// CHECK1-NEXT:    store ptr [[ARGC_ADDR]], ptr [[TMP104]], align 8255// CHECK1-NEXT:    [[TMP105:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 0256// CHECK1-NEXT:    store ptr null, ptr [[TMP105]], align 8257// CHECK1-NEXT:    [[TMP106:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 1258// CHECK1-NEXT:    store ptr [[TMP84]], ptr [[TMP106]], align 8259// CHECK1-NEXT:    [[TMP107:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 1260// CHECK1-NEXT:    store ptr [[TMP84]], ptr [[TMP107]], align 8261// CHECK1-NEXT:    [[TMP108:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 1262// CHECK1-NEXT:    store ptr null, ptr [[TMP108]], align 8263// CHECK1-NEXT:    [[TMP109:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 2264// CHECK1-NEXT:    store ptr [[TMP85]], ptr [[TMP109]], align 8265// CHECK1-NEXT:    [[TMP110:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 2266// CHECK1-NEXT:    store ptr [[TMP85]], ptr [[TMP110]], align 8267// CHECK1-NEXT:    [[TMP111:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 2268// CHECK1-NEXT:    store ptr null, ptr [[TMP111]], align 8269// CHECK1-NEXT:    [[TMP112:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 3270// CHECK1-NEXT:    store ptr [[TMP86]], ptr [[TMP112]], align 8271// CHECK1-NEXT:    [[TMP113:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 3272// CHECK1-NEXT:    store ptr [[TMP86]], ptr [[TMP113]], align 8273// CHECK1-NEXT:    [[TMP114:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 3274// CHECK1-NEXT:    store ptr null, ptr [[TMP114]], align 8275// CHECK1-NEXT:    [[TMP115:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 4276// CHECK1-NEXT:    store ptr [[A]], ptr [[TMP115]], align 8277// CHECK1-NEXT:    [[TMP116:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 4278// CHECK1-NEXT:    store ptr [[A]], ptr [[TMP116]], align 8279// CHECK1-NEXT:    [[TMP117:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 4280// CHECK1-NEXT:    store ptr null, ptr [[TMP117]], align 8281// CHECK1-NEXT:    [[TMP118:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 5282// CHECK1-NEXT:    store ptr [[TMP87]], ptr [[TMP118]], align 8283// CHECK1-NEXT:    [[TMP119:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 5284// CHECK1-NEXT:    store ptr [[TMP87]], ptr [[TMP119]], align 8285// CHECK1-NEXT:    [[TMP120:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 5286// CHECK1-NEXT:    store ptr null, ptr [[TMP120]], align 8287// CHECK1-NEXT:    [[TMP121:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 6288// CHECK1-NEXT:    store ptr [[TMP88]], ptr [[TMP121]], align 8289// CHECK1-NEXT:    [[TMP122:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 6290// CHECK1-NEXT:    store ptr [[TMP90]], ptr [[TMP122]], align 8291// CHECK1-NEXT:    [[TMP123:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 6292// CHECK1-NEXT:    store ptr null, ptr [[TMP123]], align 8293// CHECK1-NEXT:    [[TMP124:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 7294// CHECK1-NEXT:    store ptr [[TMP91]], ptr [[TMP124]], align 8295// CHECK1-NEXT:    [[TMP125:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 7296// CHECK1-NEXT:    store ptr [[TMP93]], ptr [[TMP125]], align 8297// CHECK1-NEXT:    [[TMP126:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 7298// CHECK1-NEXT:    store ptr null, ptr [[TMP126]], align 8299// CHECK1-NEXT:    [[TMP127:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 8300// CHECK1-NEXT:    store ptr [[TMP94]], ptr [[TMP127]], align 8301// CHECK1-NEXT:    [[TMP128:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 8302// CHECK1-NEXT:    store ptr [[TMP96]], ptr [[TMP128]], align 8303// CHECK1-NEXT:    [[TMP129:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 8304// CHECK1-NEXT:    store ptr null, ptr [[TMP129]], align 8305// CHECK1-NEXT:    [[TMP130:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 9306// CHECK1-NEXT:    store ptr [[TMP97]], ptr [[TMP130]], align 8307// CHECK1-NEXT:    [[TMP131:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 9308// CHECK1-NEXT:    store ptr [[TMP99]], ptr [[TMP131]], align 8309// CHECK1-NEXT:    [[TMP132:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 9310// CHECK1-NEXT:    store ptr null, ptr [[TMP132]], align 8311// CHECK1-NEXT:    [[TMP133:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 10312// CHECK1-NEXT:    store ptr [[TMP100]], ptr [[TMP133]], align 8313// CHECK1-NEXT:    [[TMP134:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 10314// CHECK1-NEXT:    store ptr [[TMP102]], ptr [[TMP134]], align 8315// CHECK1-NEXT:    [[TMP135:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 10316// CHECK1-NEXT:    store ptr null, ptr [[TMP135]], align 8317// CHECK1-NEXT:    [[TMP136:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 0318// CHECK1-NEXT:    [[TMP137:%.*]] = getelementptr inbounds [11 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 0319// CHECK1-NEXT:    [[TMP138:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 0320// CHECK1-NEXT:    store i32 3, ptr [[TMP138]], align 4321// CHECK1-NEXT:    [[TMP139:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 1322// CHECK1-NEXT:    store i32 11, ptr [[TMP139]], align 4323// CHECK1-NEXT:    [[TMP140:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 2324// CHECK1-NEXT:    store ptr [[TMP136]], ptr [[TMP140]], align 8325// CHECK1-NEXT:    [[TMP141:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 3326// CHECK1-NEXT:    store ptr [[TMP137]], ptr [[TMP141]], align 8327// CHECK1-NEXT:    [[TMP142:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 4328// CHECK1-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP142]], align 8329// CHECK1-NEXT:    [[TMP143:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 5330// CHECK1-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP143]], align 8331// CHECK1-NEXT:    [[TMP144:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 6332// CHECK1-NEXT:    store ptr null, ptr [[TMP144]], align 8333// CHECK1-NEXT:    [[TMP145:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 7334// CHECK1-NEXT:    store ptr null, ptr [[TMP145]], align 8335// CHECK1-NEXT:    [[TMP146:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 8336// CHECK1-NEXT:    store i64 0, ptr [[TMP146]], align 8337// CHECK1-NEXT:    [[TMP147:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 9338// CHECK1-NEXT:    store i64 0, ptr [[TMP147]], align 8339// CHECK1-NEXT:    [[TMP148:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 10340// CHECK1-NEXT:    store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP148]], align 4341// CHECK1-NEXT:    [[TMP149:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 11342// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP149]], align 4343// CHECK1-NEXT:    [[TMP150:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 12344// CHECK1-NEXT:    store i32 0, ptr [[TMP150]], align 4345// CHECK1-NEXT:    [[TMP151:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l43.region_id, ptr [[KERNEL_ARGS10]])346// CHECK1-NEXT:    [[TMP152:%.*]] = icmp ne i32 [[TMP151]], 0347// CHECK1-NEXT:    br i1 [[TMP152]], label [[OMP_OFFLOAD_FAILED11:%.*]], label [[OMP_OFFLOAD_CONT12:%.*]]348// CHECK1:       omp_offload.failed11:349// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l43(ptr [[ARGC_ADDR]], ptr [[TMP84]], ptr [[TMP85]], ptr [[TMP86]], ptr [[A]], ptr [[TMP87]]) #[[ATTR4]]350// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT12]]351// CHECK1:       omp_offload.cont12:352// CHECK1-NEXT:    [[TMP153:%.*]] = load i32, ptr [[ARGC_ADDR]], align 4353// CHECK1-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_ZN1S3fooEv(ptr noundef nonnull align 4 dereferenceable(4) @s)354// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP153]], [[CALL]]355// CHECK1-NEXT:    ret i32 [[ADD]]356//357//358// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l41359// CHECK1-SAME: (i64 noundef [[ARGC:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[C:%.*]], ptr noundef [[D:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 8 dereferenceable(40) [[L:%.*]]) #[[ATTR1:[0-9]+]] {360// CHECK1-NEXT:  entry:361// CHECK1-NEXT:    [[ARGC_ADDR:%.*]] = alloca i64, align 8362// CHECK1-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8363// CHECK1-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 8364// CHECK1-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 8365// CHECK1-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8366// CHECK1-NEXT:    [[L_ADDR:%.*]] = alloca ptr, align 8367// CHECK1-NEXT:    [[TMP:%.*]] = alloca ptr, align 8368// CHECK1-NEXT:    [[_TMP1:%.*]] = alloca ptr, align 8369// CHECK1-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 8370// CHECK1-NEXT:    [[L3:%.*]] = alloca [[CLASS_ANON:%.*]], align 8371// CHECK1-NEXT:    [[_TMP4:%.*]] = alloca ptr, align 8372// CHECK1-NEXT:    [[B5:%.*]] = alloca i32, align 4373// CHECK1-NEXT:    [[_TMP6:%.*]] = alloca ptr, align 8374// CHECK1-NEXT:    [[C7:%.*]] = alloca i32, align 4375// CHECK1-NEXT:    [[_TMP8:%.*]] = alloca ptr, align 8376// CHECK1-NEXT:    store i64 [[ARGC]], ptr [[ARGC_ADDR]], align 8377// CHECK1-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8378// CHECK1-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 8379// CHECK1-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 8380// CHECK1-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8381// CHECK1-NEXT:    store ptr [[L]], ptr [[L_ADDR]], align 8382// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 8383// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[C_ADDR]], align 8384// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 8385// CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[L_ADDR]], align 8386// CHECK1-NEXT:    store ptr [[TMP0]], ptr [[TMP]], align 8387// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[_TMP1]], align 8388// CHECK1-NEXT:    store ptr [[TMP3]], ptr [[_TMP2]], align 8389// CHECK1-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[_TMP2]], align 8390// CHECK1-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[L3]], ptr align 8 [[TMP4]], i64 40, i1 false)391// CHECK1-NEXT:    store ptr [[L3]], ptr [[_TMP4]], align 8392// CHECK1-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[TMP]], align 8393// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 4394// CHECK1-NEXT:    store i32 [[TMP6]], ptr [[B5]], align 4395// CHECK1-NEXT:    store ptr [[B5]], ptr [[_TMP6]], align 8396// CHECK1-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[_TMP1]], align 8397// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4398// CHECK1-NEXT:    store i32 [[TMP8]], ptr [[C7]], align 4399// CHECK1-NEXT:    store ptr [[C7]], ptr [[_TMP8]], align 8400// CHECK1-NEXT:    [[TMP9:%.*]] = load ptr, ptr [[_TMP4]], align 8401// CHECK1-NEXT:    [[CALL:%.*]] = call noundef i64 @"_ZZ4mainENK3$_0clEv"(ptr noundef nonnull align 8 dereferenceable(40) [[TMP9]])402// CHECK1-NEXT:    ret void403//404//405// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l43406// CHECK1-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[C:%.*]], ptr noundef [[D:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 8 dereferenceable(40) [[L:%.*]]) #[[ATTR1]] {407// CHECK1-NEXT:  entry:408// CHECK1-NEXT:    [[ARGC_ADDR:%.*]] = alloca ptr, align 8409// CHECK1-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8410// CHECK1-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 8411// CHECK1-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 8412// CHECK1-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8413// CHECK1-NEXT:    [[L_ADDR:%.*]] = alloca ptr, align 8414// CHECK1-NEXT:    [[TMP:%.*]] = alloca ptr, align 8415// CHECK1-NEXT:    [[_TMP1:%.*]] = alloca ptr, align 8416// CHECK1-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 8417// CHECK1-NEXT:    store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 8418// CHECK1-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8419// CHECK1-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 8420// CHECK1-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 8421// CHECK1-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8422// CHECK1-NEXT:    store ptr [[L]], ptr [[L_ADDR]], align 8423// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 8424// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 8425// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 8426// CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[A_ADDR]], align 8427// CHECK1-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[L_ADDR]], align 8428// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 8429// CHECK1-NEXT:    store ptr [[TMP2]], ptr [[_TMP1]], align 8430// CHECK1-NEXT:    store ptr [[TMP4]], ptr [[_TMP2]], align 8431// CHECK1-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[TMP]], align 8432// CHECK1-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[_TMP1]], align 8433// CHECK1-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[D_ADDR]], align 8434// CHECK1-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[_TMP2]], align 8435// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 6, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l43.omp_outlined, ptr [[TMP0]], ptr [[TMP5]], ptr [[TMP6]], ptr [[TMP7]], ptr [[TMP3]], ptr [[TMP8]])436// CHECK1-NEXT:    ret void437//438//439// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l43.omp_outlined440// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[C:%.*]], ptr noundef [[D:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 8 dereferenceable(40) [[L:%.*]]) #[[ATTR1]] {441// CHECK1-NEXT:  entry:442// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8443// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8444// CHECK1-NEXT:    [[ARGC_ADDR:%.*]] = alloca ptr, align 8445// CHECK1-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8446// CHECK1-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 8447// CHECK1-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 8448// CHECK1-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8449// CHECK1-NEXT:    [[L_ADDR:%.*]] = alloca ptr, align 8450// CHECK1-NEXT:    [[TMP:%.*]] = alloca ptr, align 8451// CHECK1-NEXT:    [[_TMP1:%.*]] = alloca ptr, align 8452// CHECK1-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 8453// CHECK1-NEXT:    [[L3:%.*]] = alloca [[CLASS_ANON:%.*]], align 8454// CHECK1-NEXT:    [[_TMP4:%.*]] = alloca ptr, align 8455// CHECK1-NEXT:    [[ARGC5:%.*]] = alloca i32, align 4456// CHECK1-NEXT:    [[B6:%.*]] = alloca i32, align 4457// CHECK1-NEXT:    [[_TMP7:%.*]] = alloca ptr, align 8458// CHECK1-NEXT:    [[C8:%.*]] = alloca i32, align 4459// CHECK1-NEXT:    [[_TMP9:%.*]] = alloca ptr, align 8460// CHECK1-NEXT:    [[A10:%.*]] = alloca i32, align 4461// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8462// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8463// CHECK1-NEXT:    store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 8464// CHECK1-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8465// CHECK1-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 8466// CHECK1-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 8467// CHECK1-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8468// CHECK1-NEXT:    store ptr [[L]], ptr [[L_ADDR]], align 8469// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 8470// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 8471// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 8472// CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[A_ADDR]], align 8473// CHECK1-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[L_ADDR]], align 8474// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 8475// CHECK1-NEXT:    store ptr [[TMP2]], ptr [[_TMP1]], align 8476// CHECK1-NEXT:    store ptr [[TMP4]], ptr [[_TMP2]], align 8477// CHECK1-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[_TMP2]], align 8478// CHECK1-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[L3]], ptr align 8 [[TMP5]], i64 40, i1 false)479// CHECK1-NEXT:    store ptr [[L3]], ptr [[_TMP4]], align 8480// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[TMP0]], align 4481// CHECK1-NEXT:    store i32 [[TMP6]], ptr [[ARGC5]], align 4482// CHECK1-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP]], align 8483// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4484// CHECK1-NEXT:    store i32 [[TMP8]], ptr [[B6]], align 4485// CHECK1-NEXT:    store ptr [[B6]], ptr [[_TMP7]], align 8486// CHECK1-NEXT:    [[TMP9:%.*]] = load ptr, ptr [[_TMP1]], align 8487// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[TMP9]], align 4488// CHECK1-NEXT:    store i32 [[TMP10]], ptr [[C8]], align 4489// CHECK1-NEXT:    store ptr [[C8]], ptr [[_TMP9]], align 8490// CHECK1-NEXT:    [[TMP11:%.*]] = load i32, ptr [[TMP3]], align 4491// CHECK1-NEXT:    store i32 [[TMP11]], ptr [[A10]], align 4492// CHECK1-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[_TMP4]], align 8493// CHECK1-NEXT:    [[CALL:%.*]] = call noundef i64 @"_ZZ4mainENK3$_0clEv"(ptr noundef nonnull align 8 dereferenceable(40) [[TMP12]])494// CHECK1-NEXT:    ret void495//496//497// CHECK1-LABEL: define {{[^@]+}}@_ZN1S3fooEv498// CHECK1-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) #[[ATTR3:[0-9]+]] comdat {499// CHECK1-NEXT:  entry:500// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8501// CHECK1-NEXT:    [[L:%.*]] = alloca ptr, align 8502// CHECK1-NEXT:    [[REF_TMP:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 8503// CHECK1-NEXT:    [[TMP:%.*]] = alloca ptr, align 8504// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 8505// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 8506// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 8507// CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8508// CHECK1-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 8509// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS3:%.*]] = alloca [3 x ptr], align 8510// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS4:%.*]] = alloca [3 x ptr], align 8511// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS5:%.*]] = alloca [3 x ptr], align 8512// CHECK1-NEXT:    [[KERNEL_ARGS6:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8513// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8514// CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8515// CHECK1-NEXT:    [[TMP0:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 0516// CHECK1-NEXT:    store ptr [[THIS1]], ptr [[TMP0]], align 8517// CHECK1-NEXT:    store ptr [[REF_TMP]], ptr [[L]], align 8518// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[L]], align 8519// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 8520// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[TMP]], align 8521// CHECK1-NEXT:    [[TMP3:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[TMP2]], i32 0, i32 0522// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[TMP2]], i32 0, i32 0523// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0524// CHECK1-NEXT:    store ptr [[THIS1]], ptr [[TMP5]], align 8525// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0526// CHECK1-NEXT:    store ptr [[THIS1]], ptr [[TMP6]], align 8527// CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0528// CHECK1-NEXT:    store ptr null, ptr [[TMP7]], align 8529// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1530// CHECK1-NEXT:    store ptr [[TMP2]], ptr [[TMP8]], align 8531// CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1532// CHECK1-NEXT:    store ptr [[TMP2]], ptr [[TMP9]], align 8533// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1534// CHECK1-NEXT:    store ptr null, ptr [[TMP10]], align 8535// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2536// CHECK1-NEXT:    store ptr [[TMP3]], ptr [[TMP11]], align 8537// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2538// CHECK1-NEXT:    store ptr [[TMP4]], ptr [[TMP12]], align 8539// CHECK1-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2540// CHECK1-NEXT:    store ptr null, ptr [[TMP13]], align 8541// CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0542// CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0543// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0544// CHECK1-NEXT:    store i32 3, ptr [[TMP16]], align 4545// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1546// CHECK1-NEXT:    store i32 3, ptr [[TMP17]], align 4547// CHECK1-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2548// CHECK1-NEXT:    store ptr [[TMP14]], ptr [[TMP18]], align 8549// CHECK1-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3550// CHECK1-NEXT:    store ptr [[TMP15]], ptr [[TMP19]], align 8551// CHECK1-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4552// CHECK1-NEXT:    store ptr @.offload_sizes.3, ptr [[TMP20]], align 8553// CHECK1-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5554// CHECK1-NEXT:    store ptr @.offload_maptypes.4, ptr [[TMP21]], align 8555// CHECK1-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6556// CHECK1-NEXT:    store ptr null, ptr [[TMP22]], align 8557// CHECK1-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7558// CHECK1-NEXT:    store ptr null, ptr [[TMP23]], align 8559// CHECK1-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8560// CHECK1-NEXT:    store i64 0, ptr [[TMP24]], align 8561// CHECK1-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9562// CHECK1-NEXT:    store i64 0, ptr [[TMP25]], align 8563// CHECK1-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10564// CHECK1-NEXT:    store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP26]], align 4565// CHECK1-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11566// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP27]], align 4567// CHECK1-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12568// CHECK1-NEXT:    store i32 0, ptr [[TMP28]], align 4569// CHECK1-NEXT:    [[TMP29:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l27.region_id, ptr [[KERNEL_ARGS]])570// CHECK1-NEXT:    [[TMP30:%.*]] = icmp ne i32 [[TMP29]], 0571// CHECK1-NEXT:    br i1 [[TMP30]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]572// CHECK1:       omp_offload.failed:573// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l27(ptr [[THIS1]], ptr [[TMP2]]) #[[ATTR4]]574// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]575// CHECK1:       omp_offload.cont:576// CHECK1-NEXT:    [[TMP31:%.*]] = load ptr, ptr [[L]], align 8577// CHECK1-NEXT:    store ptr [[TMP31]], ptr [[_TMP2]], align 8578// CHECK1-NEXT:    [[TMP32:%.*]] = load ptr, ptr [[_TMP2]], align 8579// CHECK1-NEXT:    [[TMP33:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[TMP32]], i32 0, i32 0580// CHECK1-NEXT:    [[TMP34:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[TMP32]], i32 0, i32 0581// CHECK1-NEXT:    [[TMP35:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0582// CHECK1-NEXT:    store ptr [[THIS1]], ptr [[TMP35]], align 8583// CHECK1-NEXT:    [[TMP36:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 0584// CHECK1-NEXT:    store ptr [[THIS1]], ptr [[TMP36]], align 8585// CHECK1-NEXT:    [[TMP37:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS5]], i64 0, i64 0586// CHECK1-NEXT:    store ptr null, ptr [[TMP37]], align 8587// CHECK1-NEXT:    [[TMP38:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 1588// CHECK1-NEXT:    store ptr [[TMP32]], ptr [[TMP38]], align 8589// CHECK1-NEXT:    [[TMP39:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 1590// CHECK1-NEXT:    store ptr [[TMP32]], ptr [[TMP39]], align 8591// CHECK1-NEXT:    [[TMP40:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS5]], i64 0, i64 1592// CHECK1-NEXT:    store ptr null, ptr [[TMP40]], align 8593// CHECK1-NEXT:    [[TMP41:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 2594// CHECK1-NEXT:    store ptr [[TMP33]], ptr [[TMP41]], align 8595// CHECK1-NEXT:    [[TMP42:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 2596// CHECK1-NEXT:    store ptr [[TMP34]], ptr [[TMP42]], align 8597// CHECK1-NEXT:    [[TMP43:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS5]], i64 0, i64 2598// CHECK1-NEXT:    store ptr null, ptr [[TMP43]], align 8599// CHECK1-NEXT:    [[TMP44:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0600// CHECK1-NEXT:    [[TMP45:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 0601// CHECK1-NEXT:    [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 0602// CHECK1-NEXT:    store i32 3, ptr [[TMP46]], align 4603// CHECK1-NEXT:    [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 1604// CHECK1-NEXT:    store i32 3, ptr [[TMP47]], align 4605// CHECK1-NEXT:    [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 2606// CHECK1-NEXT:    store ptr [[TMP44]], ptr [[TMP48]], align 8607// CHECK1-NEXT:    [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 3608// CHECK1-NEXT:    store ptr [[TMP45]], ptr [[TMP49]], align 8609// CHECK1-NEXT:    [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 4610// CHECK1-NEXT:    store ptr @.offload_sizes.5, ptr [[TMP50]], align 8611// CHECK1-NEXT:    [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 5612// CHECK1-NEXT:    store ptr @.offload_maptypes.6, ptr [[TMP51]], align 8613// CHECK1-NEXT:    [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 6614// CHECK1-NEXT:    store ptr null, ptr [[TMP52]], align 8615// CHECK1-NEXT:    [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 7616// CHECK1-NEXT:    store ptr null, ptr [[TMP53]], align 8617// CHECK1-NEXT:    [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 8618// CHECK1-NEXT:    store i64 0, ptr [[TMP54]], align 8619// CHECK1-NEXT:    [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 9620// CHECK1-NEXT:    store i64 0, ptr [[TMP55]], align 8621// CHECK1-NEXT:    [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 10622// CHECK1-NEXT:    store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP56]], align 4623// CHECK1-NEXT:    [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 11624// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP57]], align 4625// CHECK1-NEXT:    [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 12626// CHECK1-NEXT:    store i32 0, ptr [[TMP58]], align 4627// CHECK1-NEXT:    [[TMP59:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l29.region_id, ptr [[KERNEL_ARGS6]])628// CHECK1-NEXT:    [[TMP60:%.*]] = icmp ne i32 [[TMP59]], 0629// CHECK1-NEXT:    br i1 [[TMP60]], label [[OMP_OFFLOAD_FAILED7:%.*]], label [[OMP_OFFLOAD_CONT8:%.*]]630// CHECK1:       omp_offload.failed7:631// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l29(ptr [[THIS1]], ptr [[TMP32]]) #[[ATTR4]]632// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT8]]633// CHECK1:       omp_offload.cont8:634// CHECK1-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[THIS1]], i32 0, i32 0635// CHECK1-NEXT:    [[TMP61:%.*]] = load i32, ptr [[A]], align 4636// CHECK1-NEXT:    [[TMP62:%.*]] = load ptr, ptr [[L]], align 8637// CHECK1-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_Z3fooIZN1S3fooEvEUlvE_EiRKT_(ptr noundef nonnull align 8 dereferenceable(8) [[TMP62]])638// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP61]], [[CALL]]639// CHECK1-NEXT:    ret i32 [[ADD]]640//641//642// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l27643// CHECK1-SAME: (ptr noundef [[THIS:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[L:%.*]]) #[[ATTR1]] {644// CHECK1-NEXT:  entry:645// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8646// CHECK1-NEXT:    [[L_ADDR:%.*]] = alloca ptr, align 8647// CHECK1-NEXT:    [[TMP:%.*]] = alloca ptr, align 8648// CHECK1-NEXT:    [[L1:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 8649// CHECK1-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 8650// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8651// CHECK1-NEXT:    store ptr [[L]], ptr [[L_ADDR]], align 8652// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8653// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[L_ADDR]], align 8654// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 8655// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[TMP]], align 8656// CHECK1-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[L1]], ptr align 8 [[TMP2]], i64 8, i1 false)657// CHECK1-NEXT:    store ptr [[L1]], ptr [[_TMP2]], align 8658// CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[_TMP2]], align 8659// CHECK1-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_ZZN1S3fooEvENKUlvE_clEv(ptr noundef nonnull align 8 dereferenceable(8) [[TMP3]])660// CHECK1-NEXT:    ret void661//662//663// CHECK1-LABEL: define {{[^@]+}}@_ZZN1S3fooEvENKUlvE_clEv664// CHECK1-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[THIS:%.*]]) #[[ATTR3]] comdat {665// CHECK1-NEXT:  entry:666// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8667// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8668// CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8669// CHECK1-NEXT:    [[TMP0:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0:%.*]], ptr [[THIS1]], i32 0, i32 0670// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8671// CHECK1-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP1]], i32 0, i32 0672// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[A]], align 4673// CHECK1-NEXT:    ret i32 [[TMP2]]674//675//676// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l29677// CHECK1-SAME: (ptr noundef [[THIS:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[L:%.*]]) #[[ATTR1]] {678// CHECK1-NEXT:  entry:679// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8680// CHECK1-NEXT:    [[L_ADDR:%.*]] = alloca ptr, align 8681// CHECK1-NEXT:    [[TMP:%.*]] = alloca ptr, align 8682// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8683// CHECK1-NEXT:    store ptr [[L]], ptr [[L_ADDR]], align 8684// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8685// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[L_ADDR]], align 8686// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 8687// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[TMP]], align 8688// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l29.omp_outlined, ptr [[TMP0]], ptr [[TMP2]])689// CHECK1-NEXT:    ret void690//691//692// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l29.omp_outlined693// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[L:%.*]]) #[[ATTR1]] {694// CHECK1-NEXT:  entry:695// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8696// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8697// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8698// CHECK1-NEXT:    [[L_ADDR:%.*]] = alloca ptr, align 8699// CHECK1-NEXT:    [[TMP:%.*]] = alloca ptr, align 8700// CHECK1-NEXT:    [[L1:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 8701// CHECK1-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 8702// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8703// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8704// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8705// CHECK1-NEXT:    store ptr [[L]], ptr [[L_ADDR]], align 8706// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8707// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[L_ADDR]], align 8708// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 8709// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[TMP]], align 8710// CHECK1-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[L1]], ptr align 8 [[TMP2]], i64 8, i1 false)711// CHECK1-NEXT:    store ptr [[L1]], ptr [[_TMP2]], align 8712// CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[_TMP2]], align 8713// CHECK1-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_ZZN1S3fooEvENKUlvE_clEv(ptr noundef nonnull align 8 dereferenceable(8) [[TMP3]])714// CHECK1-NEXT:    ret void715//716//717// CHECK1-LABEL: define {{[^@]+}}@_Z3fooIZN1S3fooEvEUlvE_EiRKT_718// CHECK1-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[T:%.*]]) #[[ATTR3]] comdat {719// CHECK1-NEXT:  entry:720// CHECK1-NEXT:    [[T_ADDR:%.*]] = alloca ptr, align 8721// CHECK1-NEXT:    [[TMP:%.*]] = alloca ptr, align 8722// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x ptr], align 8723// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x ptr], align 8724// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [2 x ptr], align 8725// CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8726// CHECK1-NEXT:    store ptr [[T]], ptr [[T_ADDR]], align 8727// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[T_ADDR]], align 8728// CHECK1-NEXT:    store ptr [[TMP0]], ptr [[TMP]], align 8729// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[TMP]], align 8730// CHECK1-NEXT:    [[TMP2:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0:%.*]], ptr [[TMP1]], i32 0, i32 0731// CHECK1-NEXT:    [[TMP3:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[TMP1]], i32 0, i32 0732// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0733// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[TMP4]], align 8734// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0735// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[TMP5]], align 8736// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0737// CHECK1-NEXT:    store ptr null, ptr [[TMP6]], align 8738// CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1739// CHECK1-NEXT:    store ptr [[TMP2]], ptr [[TMP7]], align 8740// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1741// CHECK1-NEXT:    store ptr [[TMP3]], ptr [[TMP8]], align 8742// CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1743// CHECK1-NEXT:    store ptr null, ptr [[TMP9]], align 8744// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0745// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0746// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0747// CHECK1-NEXT:    store i32 3, ptr [[TMP12]], align 4748// CHECK1-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1749// CHECK1-NEXT:    store i32 2, ptr [[TMP13]], align 4750// CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2751// CHECK1-NEXT:    store ptr [[TMP10]], ptr [[TMP14]], align 8752// CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3753// CHECK1-NEXT:    store ptr [[TMP11]], ptr [[TMP15]], align 8754// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4755// CHECK1-NEXT:    store ptr @.offload_sizes.7, ptr [[TMP16]], align 8756// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5757// CHECK1-NEXT:    store ptr @.offload_maptypes.8, ptr [[TMP17]], align 8758// CHECK1-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6759// CHECK1-NEXT:    store ptr null, ptr [[TMP18]], align 8760// CHECK1-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7761// CHECK1-NEXT:    store ptr null, ptr [[TMP19]], align 8762// CHECK1-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8763// CHECK1-NEXT:    store i64 0, ptr [[TMP20]], align 8764// CHECK1-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9765// CHECK1-NEXT:    store i64 0, ptr [[TMP21]], align 8766// CHECK1-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10767// CHECK1-NEXT:    store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP22]], align 4768// CHECK1-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11769// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP23]], align 4770// CHECK1-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12771// CHECK1-NEXT:    store i32 0, ptr [[TMP24]], align 4772// CHECK1-NEXT:    [[TMP25:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooIZN1S3fooEvEUlvE_EiRKT__l18.region_id, ptr [[KERNEL_ARGS]])773// CHECK1-NEXT:    [[TMP26:%.*]] = icmp ne i32 [[TMP25]], 0774// CHECK1-NEXT:    br i1 [[TMP26]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]775// CHECK1:       omp_offload.failed:776// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooIZN1S3fooEvEUlvE_EiRKT__l18(ptr [[TMP1]]) #[[ATTR4]]777// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]778// CHECK1:       omp_offload.cont:779// CHECK1-NEXT:    ret i32 0780//781//782// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooIZN1S3fooEvEUlvE_EiRKT__l18783// CHECK1-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[T:%.*]]) #[[ATTR1]] {784// CHECK1-NEXT:  entry:785// CHECK1-NEXT:    [[T_ADDR:%.*]] = alloca ptr, align 8786// CHECK1-NEXT:    [[TMP:%.*]] = alloca ptr, align 8787// CHECK1-NEXT:    store ptr [[T]], ptr [[T_ADDR]], align 8788// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[T_ADDR]], align 8789// CHECK1-NEXT:    store ptr [[TMP0]], ptr [[TMP]], align 8790// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[TMP]], align 8791// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooIZN1S3fooEvEUlvE_EiRKT__l18.omp_outlined, ptr [[TMP1]])792// CHECK1-NEXT:    ret void793//794//795// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooIZN1S3fooEvEUlvE_EiRKT__l18.omp_outlined796// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[T:%.*]]) #[[ATTR1]] {797// CHECK1-NEXT:  entry:798// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8799// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8800// CHECK1-NEXT:    [[T_ADDR:%.*]] = alloca ptr, align 8801// CHECK1-NEXT:    [[TMP:%.*]] = alloca ptr, align 8802// CHECK1-NEXT:    [[T1:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 8803// CHECK1-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 8804// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8805// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8806// CHECK1-NEXT:    store ptr [[T]], ptr [[T_ADDR]], align 8807// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[T_ADDR]], align 8808// CHECK1-NEXT:    store ptr [[TMP0]], ptr [[TMP]], align 8809// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[TMP]], align 8810// CHECK1-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[T1]], ptr align 8 [[TMP1]], i64 8, i1 false)811// CHECK1-NEXT:    store ptr [[T1]], ptr [[_TMP2]], align 8812// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[_TMP2]], align 8813// CHECK1-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_ZZN1S3fooEvENKUlvE_clEv(ptr noundef nonnull align 8 dereferenceable(8) [[TMP2]])814// CHECK1-NEXT:    ret void815//816//817// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l27818// CHECK2-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[THIS:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[L:%.*]]) #[[ATTR0:[0-9]+]] {819// CHECK2-NEXT:  entry:820// CHECK2-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8821// CHECK2-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8822// CHECK2-NEXT:    [[L_ADDR:%.*]] = alloca ptr, align 8823// CHECK2-NEXT:    [[TMP:%.*]] = alloca ptr, align 8824// CHECK2-NEXT:    [[L1:%.*]] = alloca [[CLASS_ANON:%.*]], align 8825// CHECK2-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 8826// CHECK2-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8827// CHECK2-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8828// CHECK2-NEXT:    store ptr [[L]], ptr [[L_ADDR]], align 8829// CHECK2-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8830// CHECK2-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[L_ADDR]], align 8831// CHECK2-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 8832// CHECK2-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l27_kernel_environment, ptr [[DYN_PTR]])833// CHECK2-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -1834// CHECK2-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]835// CHECK2:       user_code.entry:836// CHECK2-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[TMP]], align 8837// CHECK2-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[L1]], ptr align 8 [[TMP3]], i64 8, i1 false)838// CHECK2-NEXT:    store ptr [[L1]], ptr [[_TMP2]], align 8839// CHECK2-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[_TMP2]], align 8840// CHECK2-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP4]], i32 0, i32 0841// CHECK2-NEXT:    store ptr [[TMP0]], ptr [[TMP5]], align 8842// CHECK2-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[_TMP2]], align 8843// CHECK2-NEXT:    [[CALL:%.*]] = call noundef i32 @_ZZN1S3fooEvENKUlvE_clEv(ptr noundef nonnull align 8 dereferenceable(8) [[TMP6]]) #[[ATTR7:[0-9]+]]844// CHECK2-NEXT:    call void @__kmpc_target_deinit()845// CHECK2-NEXT:    ret void846// CHECK2:       worker.exit:847// CHECK2-NEXT:    ret void848//849//850// CHECK2-LABEL: define {{[^@]+}}@_ZZN1S3fooEvENKUlvE_clEv851// CHECK2-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[THIS:%.*]]) #[[ATTR2:[0-9]+]] comdat align 2 {852// CHECK2-NEXT:  entry:853// CHECK2-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8854// CHECK2-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8855// CHECK2-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8856// CHECK2-NEXT:    [[TMP0:%.*]] = getelementptr inbounds nuw [[CLASS_ANON:%.*]], ptr [[THIS1]], i32 0, i32 0857// CHECK2-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8858// CHECK2-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP1]], i32 0, i32 0859// CHECK2-NEXT:    [[TMP2:%.*]] = load i32, ptr [[A]], align 4860// CHECK2-NEXT:    ret i32 [[TMP2]]861//862//863// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l29864// CHECK2-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[THIS:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[L:%.*]]) #[[ATTR3:[0-9]+]] {865// CHECK2-NEXT:  entry:866// CHECK2-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8867// CHECK2-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8868// CHECK2-NEXT:    [[L_ADDR:%.*]] = alloca ptr, align 8869// CHECK2-NEXT:    [[TMP:%.*]] = alloca ptr, align 8870// CHECK2-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [2 x ptr], align 8871// CHECK2-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8872// CHECK2-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8873// CHECK2-NEXT:    store ptr [[L]], ptr [[L_ADDR]], align 8874// CHECK2-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8875// CHECK2-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[L_ADDR]], align 8876// CHECK2-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 8877// CHECK2-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l29_kernel_environment, ptr [[DYN_PTR]])878// CHECK2-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -1879// CHECK2-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]880// CHECK2:       user_code.entry:881// CHECK2-NEXT:    [[TMP3:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])882// CHECK2-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[TMP]], align 8883// CHECK2-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 0884// CHECK2-NEXT:    store ptr [[TMP0]], ptr [[TMP5]], align 8885// CHECK2-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 1886// CHECK2-NEXT:    store ptr [[TMP4]], ptr [[TMP6]], align 8887// CHECK2-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP3]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l29_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i64 2)888// CHECK2-NEXT:    call void @__kmpc_target_deinit()889// CHECK2-NEXT:    ret void890// CHECK2:       worker.exit:891// CHECK2-NEXT:    ret void892//893//894// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l29_omp_outlined895// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[L:%.*]]) #[[ATTR4:[0-9]+]] {896// CHECK2-NEXT:  entry:897// CHECK2-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8898// CHECK2-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8899// CHECK2-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8900// CHECK2-NEXT:    [[L_ADDR:%.*]] = alloca ptr, align 8901// CHECK2-NEXT:    [[TMP:%.*]] = alloca ptr, align 8902// CHECK2-NEXT:    [[L1:%.*]] = alloca [[CLASS_ANON:%.*]], align 8903// CHECK2-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 8904// CHECK2-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8905// CHECK2-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8906// CHECK2-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8907// CHECK2-NEXT:    store ptr [[L]], ptr [[L_ADDR]], align 8908// CHECK2-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8909// CHECK2-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[L_ADDR]], align 8910// CHECK2-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 8911// CHECK2-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[TMP]], align 8912// CHECK2-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[L1]], ptr align 8 [[TMP2]], i64 8, i1 false)913// CHECK2-NEXT:    store ptr [[L1]], ptr [[_TMP2]], align 8914// CHECK2-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[_TMP2]], align 8915// CHECK2-NEXT:    [[TMP4:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP3]], i32 0, i32 0916// CHECK2-NEXT:    store ptr [[TMP0]], ptr [[TMP4]], align 8917// CHECK2-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[_TMP2]], align 8918// CHECK2-NEXT:    [[CALL:%.*]] = call noundef i32 @_ZZN1S3fooEvENKUlvE_clEv(ptr noundef nonnull align 8 dereferenceable(8) [[TMP5]]) #[[ATTR7]]919// CHECK2-NEXT:    ret void920//921//922// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l41923// CHECK2-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[ARGC:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[C:%.*]], ptr noundef [[D:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 8 dereferenceable(40) [[L:%.*]]) #[[ATTR0]] {924// CHECK2-NEXT:  entry:925// CHECK2-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8926// CHECK2-NEXT:    [[ARGC_ADDR:%.*]] = alloca i64, align 8927// CHECK2-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8928// CHECK2-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 8929// CHECK2-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 8930// CHECK2-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8931// CHECK2-NEXT:    [[L_ADDR:%.*]] = alloca ptr, align 8932// CHECK2-NEXT:    [[TMP:%.*]] = alloca ptr, align 8933// CHECK2-NEXT:    [[_TMP1:%.*]] = alloca ptr, align 8934// CHECK2-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 8935// CHECK2-NEXT:    [[L3:%.*]] = alloca [[CLASS_ANON_1:%.*]], align 8936// CHECK2-NEXT:    [[_TMP4:%.*]] = alloca ptr, align 8937// CHECK2-NEXT:    [[B5:%.*]] = alloca i32, align 4938// CHECK2-NEXT:    [[_TMP6:%.*]] = alloca ptr, align 8939// CHECK2-NEXT:    [[C7:%.*]] = alloca i32, align 4940// CHECK2-NEXT:    [[_TMP8:%.*]] = alloca ptr, align 8941// CHECK2-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8942// CHECK2-NEXT:    store i64 [[ARGC]], ptr [[ARGC_ADDR]], align 8943// CHECK2-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8944// CHECK2-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 8945// CHECK2-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 8946// CHECK2-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8947// CHECK2-NEXT:    store ptr [[L]], ptr [[L_ADDR]], align 8948// CHECK2-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 8949// CHECK2-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[C_ADDR]], align 8950// CHECK2-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 8951// CHECK2-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[L_ADDR]], align 8952// CHECK2-NEXT:    store ptr [[TMP0]], ptr [[TMP]], align 8953// CHECK2-NEXT:    store ptr [[TMP1]], ptr [[_TMP1]], align 8954// CHECK2-NEXT:    store ptr [[TMP3]], ptr [[_TMP2]], align 8955// CHECK2-NEXT:    [[TMP4:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l41_kernel_environment, ptr [[DYN_PTR]])956// CHECK2-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP4]], -1957// CHECK2-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]958// CHECK2:       user_code.entry:959// CHECK2-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[_TMP2]], align 8960// CHECK2-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[L3]], ptr align 8 [[TMP5]], i64 40, i1 false)961// CHECK2-NEXT:    store ptr [[L3]], ptr [[_TMP4]], align 8962// CHECK2-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[TMP]], align 8963// CHECK2-NEXT:    [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 4964// CHECK2-NEXT:    store i32 [[TMP7]], ptr [[B5]], align 4965// CHECK2-NEXT:    store ptr [[B5]], ptr [[_TMP6]], align 8966// CHECK2-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[_TMP1]], align 8967// CHECK2-NEXT:    [[TMP9:%.*]] = load i32, ptr [[TMP8]], align 4968// CHECK2-NEXT:    store i32 [[TMP9]], ptr [[C7]], align 4969// CHECK2-NEXT:    store ptr [[C7]], ptr [[_TMP8]], align 8970// CHECK2-NEXT:    [[TMP10:%.*]] = load ptr, ptr [[_TMP4]], align 8971// CHECK2-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[TMP10]], i32 0, i32 0972// CHECK2-NEXT:    store ptr [[ARGC_ADDR]], ptr [[TMP11]], align 8973// CHECK2-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[TMP10]], i32 0, i32 1974// CHECK2-NEXT:    [[TMP13:%.*]] = load ptr, ptr [[_TMP6]], align 8975// CHECK2-NEXT:    store ptr [[TMP13]], ptr [[TMP12]], align 8976// CHECK2-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[TMP10]], i32 0, i32 2977// CHECK2-NEXT:    [[TMP15:%.*]] = load ptr, ptr [[_TMP8]], align 8978// CHECK2-NEXT:    store ptr [[TMP15]], ptr [[TMP14]], align 8979// CHECK2-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[TMP10]], i32 0, i32 3980// CHECK2-NEXT:    store ptr [[D_ADDR]], ptr [[TMP16]], align 8981// CHECK2-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[TMP10]], i32 0, i32 4982// CHECK2-NEXT:    store ptr [[TMP2]], ptr [[TMP17]], align 8983// CHECK2-NEXT:    [[TMP18:%.*]] = load ptr, ptr [[_TMP4]], align 8984// CHECK2-NEXT:    [[CALL:%.*]] = call noundef i64 @"_ZZ4mainENK3$_0clEv"(ptr noundef nonnull align 8 dereferenceable(40) [[TMP18]]) #[[ATTR7]]985// CHECK2-NEXT:    call void @__kmpc_target_deinit()986// CHECK2-NEXT:    ret void987// CHECK2:       worker.exit:988// CHECK2-NEXT:    ret void989//990//991// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l43992// CHECK2-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[C:%.*]], ptr noundef [[D:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 8 dereferenceable(40) [[L:%.*]]) #[[ATTR3]] {993// CHECK2-NEXT:  entry:994// CHECK2-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8995// CHECK2-NEXT:    [[ARGC_ADDR:%.*]] = alloca ptr, align 8996// CHECK2-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8997// CHECK2-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 8998// CHECK2-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 8999// CHECK2-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81000// CHECK2-NEXT:    [[L_ADDR:%.*]] = alloca ptr, align 81001// CHECK2-NEXT:    [[TMP:%.*]] = alloca ptr, align 81002// CHECK2-NEXT:    [[_TMP1:%.*]] = alloca ptr, align 81003// CHECK2-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 81004// CHECK2-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [6 x ptr], align 81005// CHECK2-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81006// CHECK2-NEXT:    store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 81007// CHECK2-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 81008// CHECK2-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 81009// CHECK2-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 81010// CHECK2-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81011// CHECK2-NEXT:    store ptr [[L]], ptr [[L_ADDR]], align 81012// CHECK2-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 81013// CHECK2-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 81014// CHECK2-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 81015// CHECK2-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[A_ADDR]], align 81016// CHECK2-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[L_ADDR]], align 81017// CHECK2-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 81018// CHECK2-NEXT:    store ptr [[TMP2]], ptr [[_TMP1]], align 81019// CHECK2-NEXT:    store ptr [[TMP4]], ptr [[_TMP2]], align 81020// CHECK2-NEXT:    [[TMP5:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l43_kernel_environment, ptr [[DYN_PTR]])1021// CHECK2-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP5]], -11022// CHECK2-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]1023// CHECK2:       user_code.entry:1024// CHECK2-NEXT:    [[TMP6:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1025// CHECK2-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP]], align 81026// CHECK2-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[_TMP1]], align 81027// CHECK2-NEXT:    [[TMP9:%.*]] = load ptr, ptr [[D_ADDR]], align 81028// CHECK2-NEXT:    [[TMP10:%.*]] = load ptr, ptr [[_TMP2]], align 81029// CHECK2-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [6 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 01030// CHECK2-NEXT:    store ptr [[TMP0]], ptr [[TMP11]], align 81031// CHECK2-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [6 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 11032// CHECK2-NEXT:    store ptr [[TMP7]], ptr [[TMP12]], align 81033// CHECK2-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [6 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 21034// CHECK2-NEXT:    store ptr [[TMP8]], ptr [[TMP13]], align 81035// CHECK2-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [6 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 31036// CHECK2-NEXT:    store ptr [[TMP9]], ptr [[TMP14]], align 81037// CHECK2-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [6 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 41038// CHECK2-NEXT:    store ptr [[TMP3]], ptr [[TMP15]], align 81039// CHECK2-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [6 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 51040// CHECK2-NEXT:    store ptr [[TMP10]], ptr [[TMP16]], align 81041// CHECK2-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP6]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l43_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i64 6)1042// CHECK2-NEXT:    call void @__kmpc_target_deinit()1043// CHECK2-NEXT:    ret void1044// CHECK2:       worker.exit:1045// CHECK2-NEXT:    ret void1046//1047//1048// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l43_omp_outlined1049// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[C:%.*]], ptr noundef [[D:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 8 dereferenceable(40) [[L:%.*]]) #[[ATTR4]] {1050// CHECK2-NEXT:  entry:1051// CHECK2-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81052// CHECK2-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81053// CHECK2-NEXT:    [[ARGC_ADDR:%.*]] = alloca ptr, align 81054// CHECK2-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 81055// CHECK2-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 81056// CHECK2-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 81057// CHECK2-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81058// CHECK2-NEXT:    [[L_ADDR:%.*]] = alloca ptr, align 81059// CHECK2-NEXT:    [[TMP:%.*]] = alloca ptr, align 81060// CHECK2-NEXT:    [[_TMP1:%.*]] = alloca ptr, align 81061// CHECK2-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 81062// CHECK2-NEXT:    [[L3:%.*]] = alloca [[CLASS_ANON_1:%.*]], align 81063// CHECK2-NEXT:    [[_TMP4:%.*]] = alloca ptr, align 81064// CHECK2-NEXT:    [[ARGC5:%.*]] = alloca i32, align 41065// CHECK2-NEXT:    [[B6:%.*]] = alloca i32, align 41066// CHECK2-NEXT:    [[_TMP7:%.*]] = alloca ptr, align 81067// CHECK2-NEXT:    [[C8:%.*]] = alloca i32, align 41068// CHECK2-NEXT:    [[_TMP9:%.*]] = alloca ptr, align 81069// CHECK2-NEXT:    [[A10:%.*]] = alloca i32, align 41070// CHECK2-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81071// CHECK2-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81072// CHECK2-NEXT:    store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 81073// CHECK2-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 81074// CHECK2-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 81075// CHECK2-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 81076// CHECK2-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81077// CHECK2-NEXT:    store ptr [[L]], ptr [[L_ADDR]], align 81078// CHECK2-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 81079// CHECK2-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 81080// CHECK2-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 81081// CHECK2-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[A_ADDR]], align 81082// CHECK2-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[L_ADDR]], align 81083// CHECK2-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 81084// CHECK2-NEXT:    store ptr [[TMP2]], ptr [[_TMP1]], align 81085// CHECK2-NEXT:    store ptr [[TMP4]], ptr [[_TMP2]], align 81086// CHECK2-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[_TMP2]], align 81087// CHECK2-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[L3]], ptr align 8 [[TMP5]], i64 40, i1 false)1088// CHECK2-NEXT:    store ptr [[L3]], ptr [[_TMP4]], align 81089// CHECK2-NEXT:    [[TMP6:%.*]] = load i32, ptr [[TMP0]], align 41090// CHECK2-NEXT:    store i32 [[TMP6]], ptr [[ARGC5]], align 41091// CHECK2-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP]], align 81092// CHECK2-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41093// CHECK2-NEXT:    store i32 [[TMP8]], ptr [[B6]], align 41094// CHECK2-NEXT:    store ptr [[B6]], ptr [[_TMP7]], align 81095// CHECK2-NEXT:    [[TMP9:%.*]] = load ptr, ptr [[_TMP1]], align 81096// CHECK2-NEXT:    [[TMP10:%.*]] = load i32, ptr [[TMP9]], align 41097// CHECK2-NEXT:    store i32 [[TMP10]], ptr [[C8]], align 41098// CHECK2-NEXT:    store ptr [[C8]], ptr [[_TMP9]], align 81099// CHECK2-NEXT:    [[TMP11:%.*]] = load i32, ptr [[TMP3]], align 41100// CHECK2-NEXT:    store i32 [[TMP11]], ptr [[A10]], align 41101// CHECK2-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[_TMP4]], align 81102// CHECK2-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[TMP12]], i32 0, i32 01103// CHECK2-NEXT:    store ptr [[ARGC5]], ptr [[TMP13]], align 81104// CHECK2-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[TMP12]], i32 0, i32 11105// CHECK2-NEXT:    [[TMP15:%.*]] = load ptr, ptr [[_TMP7]], align 81106// CHECK2-NEXT:    store ptr [[TMP15]], ptr [[TMP14]], align 81107// CHECK2-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[TMP12]], i32 0, i32 21108// CHECK2-NEXT:    [[TMP17:%.*]] = load ptr, ptr [[_TMP9]], align 81109// CHECK2-NEXT:    store ptr [[TMP17]], ptr [[TMP16]], align 81110// CHECK2-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[TMP12]], i32 0, i32 31111// CHECK2-NEXT:    store ptr [[D_ADDR]], ptr [[TMP18]], align 81112// CHECK2-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[TMP12]], i32 0, i32 41113// CHECK2-NEXT:    store ptr [[A10]], ptr [[TMP19]], align 81114// CHECK2-NEXT:    [[TMP20:%.*]] = load ptr, ptr [[_TMP4]], align 81115// CHECK2-NEXT:    [[CALL:%.*]] = call noundef i64 @"_ZZ4mainENK3$_0clEv"(ptr noundef nonnull align 8 dereferenceable(40) [[TMP20]]) #[[ATTR7]]1116// CHECK2-NEXT:    ret void1117//1118//1119// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooIZN1S3fooEvEUlvE_EiRKT__l181120// CHECK2-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[T:%.*]]) #[[ATTR3]] {1121// CHECK2-NEXT:  entry:1122// CHECK2-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81123// CHECK2-NEXT:    [[T_ADDR:%.*]] = alloca ptr, align 81124// CHECK2-NEXT:    [[TMP:%.*]] = alloca ptr, align 81125// CHECK2-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 81126// CHECK2-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81127// CHECK2-NEXT:    store ptr [[T]], ptr [[T_ADDR]], align 81128// CHECK2-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[T_ADDR]], align 81129// CHECK2-NEXT:    store ptr [[TMP0]], ptr [[TMP]], align 81130// CHECK2-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooIZN1S3fooEvEUlvE_EiRKT__l18_kernel_environment, ptr [[DYN_PTR]])1131// CHECK2-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -11132// CHECK2-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]1133// CHECK2:       user_code.entry:1134// CHECK2-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1135// CHECK2-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[TMP]], align 81136// CHECK2-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 01137// CHECK2-NEXT:    store ptr [[TMP3]], ptr [[TMP4]], align 81138// CHECK2-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP2]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooIZN1S3fooEvEUlvE_EiRKT__l18_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i64 1)1139// CHECK2-NEXT:    call void @__kmpc_target_deinit()1140// CHECK2-NEXT:    ret void1141// CHECK2:       worker.exit:1142// CHECK2-NEXT:    ret void1143//1144//1145// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooIZN1S3fooEvEUlvE_EiRKT__l18_omp_outlined1146// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[T:%.*]]) #[[ATTR4]] {1147// CHECK2-NEXT:  entry:1148// CHECK2-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81149// CHECK2-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81150// CHECK2-NEXT:    [[T_ADDR:%.*]] = alloca ptr, align 81151// CHECK2-NEXT:    [[TMP:%.*]] = alloca ptr, align 81152// CHECK2-NEXT:    [[T1:%.*]] = alloca [[CLASS_ANON:%.*]], align 81153// CHECK2-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 81154// CHECK2-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81155// CHECK2-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81156// CHECK2-NEXT:    store ptr [[T]], ptr [[T_ADDR]], align 81157// CHECK2-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[T_ADDR]], align 81158// CHECK2-NEXT:    store ptr [[TMP0]], ptr [[TMP]], align 81159// CHECK2-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[TMP]], align 81160// CHECK2-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[T1]], ptr align 8 [[TMP1]], i64 8, i1 false)1161// CHECK2-NEXT:    store ptr [[T1]], ptr [[_TMP2]], align 81162// CHECK2-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[_TMP2]], align 81163// CHECK2-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[_TMP2]], align 81164// CHECK2-NEXT:    [[CALL:%.*]] = call noundef i32 @_ZZN1S3fooEvENKUlvE_clEv(ptr noundef nonnull align 8 dereferenceable(8) [[TMP3]]) #[[ATTR7]]1165// CHECK2-NEXT:    ret void1166//1167//1168// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l271169// CHECK3-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[THIS:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[L:%.*]]) #[[ATTR0:[0-9]+]] {1170// CHECK3-NEXT:  entry:1171// CHECK3-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81172// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 81173// CHECK3-NEXT:    [[L_ADDR:%.*]] = alloca ptr, align 81174// CHECK3-NEXT:    [[TMP:%.*]] = alloca ptr, align 81175// CHECK3-NEXT:    [[L1:%.*]] = alloca [[CLASS_ANON:%.*]], align 81176// CHECK3-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 81177// CHECK3-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81178// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 81179// CHECK3-NEXT:    store ptr [[L]], ptr [[L_ADDR]], align 81180// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81181// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[L_ADDR]], align 81182// CHECK3-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 81183// CHECK3-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l27_kernel_environment, ptr [[DYN_PTR]])1184// CHECK3-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -11185// CHECK3-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]1186// CHECK3:       user_code.entry:1187// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[TMP]], align 81188// CHECK3-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[L1]], ptr align 8 [[TMP3]], i64 8, i1 false)1189// CHECK3-NEXT:    store ptr [[L1]], ptr [[_TMP2]], align 81190// CHECK3-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[_TMP2]], align 81191// CHECK3-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP4]], i32 0, i32 01192// CHECK3-NEXT:    store ptr [[TMP0]], ptr [[TMP5]], align 81193// CHECK3-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[_TMP2]], align 81194// CHECK3-NEXT:    [[CALL:%.*]] = call noundef i32 @_ZZN1S3fooEvENKUlvE_clEv(ptr noundef nonnull align 8 dereferenceable(8) [[TMP6]]) #[[ATTR7:[0-9]+]]1195// CHECK3-NEXT:    call void @__kmpc_target_deinit()1196// CHECK3-NEXT:    ret void1197// CHECK3:       worker.exit:1198// CHECK3-NEXT:    ret void1199//1200//1201// CHECK3-LABEL: define {{[^@]+}}@_ZZN1S3fooEvENKUlvE_clEv1202// CHECK3-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[THIS:%.*]]) #[[ATTR2:[0-9]+]] comdat align 2 {1203// CHECK3-NEXT:  entry:1204// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 81205// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 81206// CHECK3-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81207// CHECK3-NEXT:    [[TMP0:%.*]] = getelementptr inbounds nuw [[CLASS_ANON:%.*]], ptr [[THIS1]], i32 0, i32 01208// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 81209// CHECK3-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP1]], i32 0, i32 01210// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[A]], align 41211// CHECK3-NEXT:    ret i32 [[TMP2]]1212//1213//1214// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l291215// CHECK3-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[THIS:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[L:%.*]]) #[[ATTR3:[0-9]+]] {1216// CHECK3-NEXT:  entry:1217// CHECK3-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81218// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 81219// CHECK3-NEXT:    [[L_ADDR:%.*]] = alloca ptr, align 81220// CHECK3-NEXT:    [[TMP:%.*]] = alloca ptr, align 81221// CHECK3-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [2 x ptr], align 81222// CHECK3-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81223// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 81224// CHECK3-NEXT:    store ptr [[L]], ptr [[L_ADDR]], align 81225// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81226// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[L_ADDR]], align 81227// CHECK3-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 81228// CHECK3-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l29_kernel_environment, ptr [[DYN_PTR]])1229// CHECK3-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -11230// CHECK3-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]1231// CHECK3:       user_code.entry:1232// CHECK3-NEXT:    [[TMP3:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])1233// CHECK3-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[TMP]], align 81234// CHECK3-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 01235// CHECK3-NEXT:    store ptr [[TMP0]], ptr [[TMP5]], align 81236// CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 11237// CHECK3-NEXT:    store ptr [[TMP4]], ptr [[TMP6]], align 81238// CHECK3-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP3]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l29_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i64 2)1239// CHECK3-NEXT:    call void @__kmpc_target_deinit()1240// CHECK3-NEXT:    ret void1241// CHECK3:       worker.exit:1242// CHECK3-NEXT:    ret void1243//1244//1245// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l29_omp_outlined1246// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[L:%.*]]) #[[ATTR4:[0-9]+]] {1247// CHECK3-NEXT:  entry:1248// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81249// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81250// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 81251// CHECK3-NEXT:    [[L_ADDR:%.*]] = alloca ptr, align 81252// CHECK3-NEXT:    [[TMP:%.*]] = alloca ptr, align 81253// CHECK3-NEXT:    [[L1:%.*]] = alloca [[CLASS_ANON:%.*]], align 81254// CHECK3-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 81255// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81256// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81257// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 81258// CHECK3-NEXT:    store ptr [[L]], ptr [[L_ADDR]], align 81259// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81260// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[L_ADDR]], align 81261// CHECK3-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 81262// CHECK3-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[TMP]], align 81263// CHECK3-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[L1]], ptr align 8 [[TMP2]], i64 8, i1 false)1264// CHECK3-NEXT:    store ptr [[L1]], ptr [[_TMP2]], align 81265// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[_TMP2]], align 81266// CHECK3-NEXT:    [[TMP4:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[TMP3]], i32 0, i32 01267// CHECK3-NEXT:    store ptr [[TMP0]], ptr [[TMP4]], align 81268// CHECK3-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[_TMP2]], align 81269// CHECK3-NEXT:    [[CALL:%.*]] = call noundef i32 @_ZZN1S3fooEvENKUlvE_clEv(ptr noundef nonnull align 8 dereferenceable(8) [[TMP5]]) #[[ATTR7]]1270// CHECK3-NEXT:    ret void1271//1272//1273// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l411274// CHECK3-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[ARGC:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[C:%.*]], ptr noundef [[D:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 8 dereferenceable(40) [[L:%.*]]) #[[ATTR0]] {1275// CHECK3-NEXT:  entry:1276// CHECK3-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81277// CHECK3-NEXT:    [[ARGC_ADDR:%.*]] = alloca i64, align 81278// CHECK3-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 81279// CHECK3-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 81280// CHECK3-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 81281// CHECK3-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81282// CHECK3-NEXT:    [[L_ADDR:%.*]] = alloca ptr, align 81283// CHECK3-NEXT:    [[TMP:%.*]] = alloca ptr, align 81284// CHECK3-NEXT:    [[_TMP1:%.*]] = alloca ptr, align 81285// CHECK3-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 81286// CHECK3-NEXT:    [[L3:%.*]] = alloca [[CLASS_ANON_1:%.*]], align 81287// CHECK3-NEXT:    [[_TMP4:%.*]] = alloca ptr, align 81288// CHECK3-NEXT:    [[B5:%.*]] = alloca i32, align 41289// CHECK3-NEXT:    [[_TMP6:%.*]] = alloca ptr, align 81290// CHECK3-NEXT:    [[C7:%.*]] = alloca i32, align 41291// CHECK3-NEXT:    [[_TMP8:%.*]] = alloca ptr, align 81292// CHECK3-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81293// CHECK3-NEXT:    store i64 [[ARGC]], ptr [[ARGC_ADDR]], align 81294// CHECK3-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 81295// CHECK3-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 81296// CHECK3-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 81297// CHECK3-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81298// CHECK3-NEXT:    store ptr [[L]], ptr [[L_ADDR]], align 81299// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 81300// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[C_ADDR]], align 81301// CHECK3-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[A_ADDR]], align 81302// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[L_ADDR]], align 81303// CHECK3-NEXT:    store ptr [[TMP0]], ptr [[TMP]], align 81304// CHECK3-NEXT:    store ptr [[TMP1]], ptr [[_TMP1]], align 81305// CHECK3-NEXT:    store ptr [[TMP3]], ptr [[_TMP2]], align 81306// CHECK3-NEXT:    [[TMP4:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l41_kernel_environment, ptr [[DYN_PTR]])1307// CHECK3-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP4]], -11308// CHECK3-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]1309// CHECK3:       user_code.entry:1310// CHECK3-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[_TMP2]], align 81311// CHECK3-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[L3]], ptr align 8 [[TMP5]], i64 40, i1 false)1312// CHECK3-NEXT:    store ptr [[L3]], ptr [[_TMP4]], align 81313// CHECK3-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[TMP]], align 81314// CHECK3-NEXT:    [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 41315// CHECK3-NEXT:    store i32 [[TMP7]], ptr [[B5]], align 41316// CHECK3-NEXT:    store ptr [[B5]], ptr [[_TMP6]], align 81317// CHECK3-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[_TMP1]], align 81318// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[TMP8]], align 41319// CHECK3-NEXT:    store i32 [[TMP9]], ptr [[C7]], align 41320// CHECK3-NEXT:    store ptr [[C7]], ptr [[_TMP8]], align 81321// CHECK3-NEXT:    [[TMP10:%.*]] = load ptr, ptr [[_TMP4]], align 81322// CHECK3-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[TMP10]], i32 0, i32 01323// CHECK3-NEXT:    store ptr [[ARGC_ADDR]], ptr [[TMP11]], align 81324// CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[TMP10]], i32 0, i32 11325// CHECK3-NEXT:    [[TMP13:%.*]] = load ptr, ptr [[_TMP6]], align 81326// CHECK3-NEXT:    store ptr [[TMP13]], ptr [[TMP12]], align 81327// CHECK3-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[TMP10]], i32 0, i32 21328// CHECK3-NEXT:    [[TMP15:%.*]] = load ptr, ptr [[_TMP8]], align 81329// CHECK3-NEXT:    store ptr [[TMP15]], ptr [[TMP14]], align 81330// CHECK3-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[TMP10]], i32 0, i32 31331// CHECK3-NEXT:    store ptr [[D_ADDR]], ptr [[TMP16]], align 81332// CHECK3-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[TMP10]], i32 0, i32 41333// CHECK3-NEXT:    store ptr [[TMP2]], ptr [[TMP17]], align 81334// CHECK3-NEXT:    [[TMP18:%.*]] = load ptr, ptr [[_TMP4]], align 81335// CHECK3-NEXT:    [[CALL:%.*]] = call noundef i64 @"_ZZ4mainENK3$_0clEv"(ptr noundef nonnull align 8 dereferenceable(40) [[TMP18]]) #[[ATTR7]]1336// CHECK3-NEXT:    call void @__kmpc_target_deinit()1337// CHECK3-NEXT:    ret void1338// CHECK3:       worker.exit:1339// CHECK3-NEXT:    ret void1340//1341//1342// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l431343// CHECK3-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[C:%.*]], ptr noundef [[D:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 8 dereferenceable(40) [[L:%.*]]) #[[ATTR3]] {1344// CHECK3-NEXT:  entry:1345// CHECK3-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81346// CHECK3-NEXT:    [[ARGC_ADDR:%.*]] = alloca ptr, align 81347// CHECK3-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 81348// CHECK3-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 81349// CHECK3-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 81350// CHECK3-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81351// CHECK3-NEXT:    [[L_ADDR:%.*]] = alloca ptr, align 81352// CHECK3-NEXT:    [[TMP:%.*]] = alloca ptr, align 81353// CHECK3-NEXT:    [[_TMP1:%.*]] = alloca ptr, align 81354// CHECK3-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 81355// CHECK3-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [6 x ptr], align 81356// CHECK3-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81357// CHECK3-NEXT:    store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 81358// CHECK3-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 81359// CHECK3-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 81360// CHECK3-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 81361// CHECK3-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81362// CHECK3-NEXT:    store ptr [[L]], ptr [[L_ADDR]], align 81363// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 81364// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 81365// CHECK3-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 81366// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[A_ADDR]], align 81367// CHECK3-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[L_ADDR]], align 81368// CHECK3-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 81369// CHECK3-NEXT:    store ptr [[TMP2]], ptr [[_TMP1]], align 81370// CHECK3-NEXT:    store ptr [[TMP4]], ptr [[_TMP2]], align 81371// CHECK3-NEXT:    [[TMP5:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l43_kernel_environment, ptr [[DYN_PTR]])1372// CHECK3-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP5]], -11373// CHECK3-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]1374// CHECK3:       user_code.entry:1375// CHECK3-NEXT:    [[TMP6:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1376// CHECK3-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP]], align 81377// CHECK3-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[_TMP1]], align 81378// CHECK3-NEXT:    [[TMP9:%.*]] = load ptr, ptr [[D_ADDR]], align 81379// CHECK3-NEXT:    [[TMP10:%.*]] = load ptr, ptr [[_TMP2]], align 81380// CHECK3-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [6 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 01381// CHECK3-NEXT:    store ptr [[TMP0]], ptr [[TMP11]], align 81382// CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [6 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 11383// CHECK3-NEXT:    store ptr [[TMP7]], ptr [[TMP12]], align 81384// CHECK3-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [6 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 21385// CHECK3-NEXT:    store ptr [[TMP8]], ptr [[TMP13]], align 81386// CHECK3-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [6 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 31387// CHECK3-NEXT:    store ptr [[TMP9]], ptr [[TMP14]], align 81388// CHECK3-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [6 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 41389// CHECK3-NEXT:    store ptr [[TMP3]], ptr [[TMP15]], align 81390// CHECK3-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [6 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 51391// CHECK3-NEXT:    store ptr [[TMP10]], ptr [[TMP16]], align 81392// CHECK3-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP6]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l43_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i64 6)1393// CHECK3-NEXT:    call void @__kmpc_target_deinit()1394// CHECK3-NEXT:    ret void1395// CHECK3:       worker.exit:1396// CHECK3-NEXT:    ret void1397//1398//1399// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l43_omp_outlined1400// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[C:%.*]], ptr noundef [[D:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 8 dereferenceable(40) [[L:%.*]]) #[[ATTR4]] {1401// CHECK3-NEXT:  entry:1402// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81403// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81404// CHECK3-NEXT:    [[ARGC_ADDR:%.*]] = alloca ptr, align 81405// CHECK3-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 81406// CHECK3-NEXT:    [[C_ADDR:%.*]] = alloca ptr, align 81407// CHECK3-NEXT:    [[D_ADDR:%.*]] = alloca ptr, align 81408// CHECK3-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81409// CHECK3-NEXT:    [[L_ADDR:%.*]] = alloca ptr, align 81410// CHECK3-NEXT:    [[TMP:%.*]] = alloca ptr, align 81411// CHECK3-NEXT:    [[_TMP1:%.*]] = alloca ptr, align 81412// CHECK3-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 81413// CHECK3-NEXT:    [[L3:%.*]] = alloca [[CLASS_ANON_1:%.*]], align 81414// CHECK3-NEXT:    [[_TMP4:%.*]] = alloca ptr, align 81415// CHECK3-NEXT:    [[ARGC5:%.*]] = alloca i32, align 41416// CHECK3-NEXT:    [[B6:%.*]] = alloca i32, align 41417// CHECK3-NEXT:    [[_TMP7:%.*]] = alloca ptr, align 81418// CHECK3-NEXT:    [[C8:%.*]] = alloca i32, align 41419// CHECK3-NEXT:    [[_TMP9:%.*]] = alloca ptr, align 81420// CHECK3-NEXT:    [[A10:%.*]] = alloca i32, align 41421// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81422// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81423// CHECK3-NEXT:    store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 81424// CHECK3-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 81425// CHECK3-NEXT:    store ptr [[C]], ptr [[C_ADDR]], align 81426// CHECK3-NEXT:    store ptr [[D]], ptr [[D_ADDR]], align 81427// CHECK3-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81428// CHECK3-NEXT:    store ptr [[L]], ptr [[L_ADDR]], align 81429// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 81430// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 81431// CHECK3-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 81432// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[A_ADDR]], align 81433// CHECK3-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[L_ADDR]], align 81434// CHECK3-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 81435// CHECK3-NEXT:    store ptr [[TMP2]], ptr [[_TMP1]], align 81436// CHECK3-NEXT:    store ptr [[TMP4]], ptr [[_TMP2]], align 81437// CHECK3-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[_TMP2]], align 81438// CHECK3-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[L3]], ptr align 8 [[TMP5]], i64 40, i1 false)1439// CHECK3-NEXT:    store ptr [[L3]], ptr [[_TMP4]], align 81440// CHECK3-NEXT:    [[TMP6:%.*]] = load i32, ptr [[TMP0]], align 41441// CHECK3-NEXT:    store i32 [[TMP6]], ptr [[ARGC5]], align 41442// CHECK3-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP]], align 81443// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 41444// CHECK3-NEXT:    store i32 [[TMP8]], ptr [[B6]], align 41445// CHECK3-NEXT:    store ptr [[B6]], ptr [[_TMP7]], align 81446// CHECK3-NEXT:    [[TMP9:%.*]] = load ptr, ptr [[_TMP1]], align 81447// CHECK3-NEXT:    [[TMP10:%.*]] = load i32, ptr [[TMP9]], align 41448// CHECK3-NEXT:    store i32 [[TMP10]], ptr [[C8]], align 41449// CHECK3-NEXT:    store ptr [[C8]], ptr [[_TMP9]], align 81450// CHECK3-NEXT:    [[TMP11:%.*]] = load i32, ptr [[TMP3]], align 41451// CHECK3-NEXT:    store i32 [[TMP11]], ptr [[A10]], align 41452// CHECK3-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[_TMP4]], align 81453// CHECK3-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[TMP12]], i32 0, i32 01454// CHECK3-NEXT:    store ptr [[ARGC5]], ptr [[TMP13]], align 81455// CHECK3-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[TMP12]], i32 0, i32 11456// CHECK3-NEXT:    [[TMP15:%.*]] = load ptr, ptr [[_TMP7]], align 81457// CHECK3-NEXT:    store ptr [[TMP15]], ptr [[TMP14]], align 81458// CHECK3-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[TMP12]], i32 0, i32 21459// CHECK3-NEXT:    [[TMP17:%.*]] = load ptr, ptr [[_TMP9]], align 81460// CHECK3-NEXT:    store ptr [[TMP17]], ptr [[TMP16]], align 81461// CHECK3-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[TMP12]], i32 0, i32 31462// CHECK3-NEXT:    store ptr [[D_ADDR]], ptr [[TMP18]], align 81463// CHECK3-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr [[TMP12]], i32 0, i32 41464// CHECK3-NEXT:    store ptr [[A10]], ptr [[TMP19]], align 81465// CHECK3-NEXT:    [[TMP20:%.*]] = load ptr, ptr [[_TMP4]], align 81466// CHECK3-NEXT:    [[CALL:%.*]] = call noundef i64 @"_ZZ4mainENK3$_0clEv"(ptr noundef nonnull align 8 dereferenceable(40) [[TMP20]]) #[[ATTR7]]1467// CHECK3-NEXT:    ret void1468//1469//1470// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooIZN1S3fooEvEUlvE_EiRKT__l181471// CHECK3-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[T:%.*]]) #[[ATTR3]] {1472// CHECK3-NEXT:  entry:1473// CHECK3-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81474// CHECK3-NEXT:    [[T_ADDR:%.*]] = alloca ptr, align 81475// CHECK3-NEXT:    [[TMP:%.*]] = alloca ptr, align 81476// CHECK3-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 81477// CHECK3-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81478// CHECK3-NEXT:    store ptr [[T]], ptr [[T_ADDR]], align 81479// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[T_ADDR]], align 81480// CHECK3-NEXT:    store ptr [[TMP0]], ptr [[TMP]], align 81481// CHECK3-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooIZN1S3fooEvEUlvE_EiRKT__l18_kernel_environment, ptr [[DYN_PTR]])1482// CHECK3-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -11483// CHECK3-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]1484// CHECK3:       user_code.entry:1485// CHECK3-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1486// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[TMP]], align 81487// CHECK3-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 01488// CHECK3-NEXT:    store ptr [[TMP3]], ptr [[TMP4]], align 81489// CHECK3-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP2]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooIZN1S3fooEvEUlvE_EiRKT__l18_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i64 1)1490// CHECK3-NEXT:    call void @__kmpc_target_deinit()1491// CHECK3-NEXT:    ret void1492// CHECK3:       worker.exit:1493// CHECK3-NEXT:    ret void1494//1495//1496// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooIZN1S3fooEvEUlvE_EiRKT__l18_omp_outlined1497// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[T:%.*]]) #[[ATTR4]] {1498// CHECK3-NEXT:  entry:1499// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81500// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81501// CHECK3-NEXT:    [[T_ADDR:%.*]] = alloca ptr, align 81502// CHECK3-NEXT:    [[TMP:%.*]] = alloca ptr, align 81503// CHECK3-NEXT:    [[T1:%.*]] = alloca [[CLASS_ANON:%.*]], align 81504// CHECK3-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 81505// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81506// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81507// CHECK3-NEXT:    store ptr [[T]], ptr [[T_ADDR]], align 81508// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[T_ADDR]], align 81509// CHECK3-NEXT:    store ptr [[TMP0]], ptr [[TMP]], align 81510// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[TMP]], align 81511// CHECK3-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[T1]], ptr align 8 [[TMP1]], i64 8, i1 false)1512// CHECK3-NEXT:    store ptr [[T1]], ptr [[_TMP2]], align 81513// CHECK3-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[_TMP2]], align 81514// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[_TMP2]], align 81515// CHECK3-NEXT:    [[CALL:%.*]] = call noundef i32 @_ZZN1S3fooEvENKUlvE_clEv(ptr noundef nonnull align 8 dereferenceable(8) [[TMP3]]) #[[ATTR7]]1516// CHECK3-NEXT:    ret void1517//1518