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// Test host codegen.3// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=60 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK14// RUN: %clang_cc1 -fopenmp -fopenmp-version=60 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s5// RUN: %clang_cc1 -fopenmp -fopenmp-version=60 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK16// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=60 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK37// RUN: %clang_cc1 -fopenmp -fopenmp-version=60 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s8// RUN: %clang_cc1 -fopenmp -fopenmp-version=60 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK39 10// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=60 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"11// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=60 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s12// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=60 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"13// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=60 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"14// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=60 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s15// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=60 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"16 17// Test target codegen - host bc file has to be created first.18// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=60 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc19// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=60 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK920// RUN: %clang_cc1 -fopenmp -fopenmp-version=60 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s21// RUN: %clang_cc1 -fopenmp -fopenmp-version=60 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK1022// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=60 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc23// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=60 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix=CHECK1124// RUN: %clang_cc1 -fopenmp -fopenmp-version=60 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s25// RUN: %clang_cc1 -fopenmp -fopenmp-version=60 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK1226 27// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=60 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc28// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=60 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"29// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=60 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s30// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=60 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"31// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=60 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc32// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=60 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"33// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=60 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s34// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=60 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"35 36// expected-no-diagnostics37 38#ifndef HEADER39#define HEADER40 41template<typename tx>42tx ftemplate(int n) {43  tx a = 0;44 45  #pragma omp parallel num_threads(strict: tx(20)) severity(fatal) message("msg")46  {47  }48 49  short b = 1;50  #pragma omp parallel num_threads(strict: b) severity(warning) message("msg")51  {52    a += b;53  }54 55  return a;56}57 58static59int fstatic(int n) {60 61  #pragma omp target parallel num_threads(strict: n) message("msg")62  {63  }64 65  #pragma omp target parallel num_threads(strict: 32+n) severity(warning)66  {67  }68 69  return n+1;70}71 72struct S1 {73  double a;74 75  int r1(int n){76    int b = 1;77 78    #pragma omp parallel num_threads(strict: n-b) severity(warning) message("msg")79    {80      this->a = (double)b + 1.5;81    }82 83    #pragma omp parallel num_threads(strict: 1024) severity(fatal)84    {85      this->a = 2.5;86    }87 88    return (int)a;89  }90};91 92int bar(int n){93  int a = 0;94 95  #pragma omp target96  {97    S1 S;98    a += S.r1(n);99  }100 101  a += fstatic(n);102 103  #pragma omp target104  a += ftemplate<int>(n);105 106  return a;107}108 109#endif110// CHECK1-LABEL: define {{[^@]+}}@_Z3bari111// CHECK1-SAME: (i32 noundef signext [[N:%.*]]) #[[ATTR0:[0-9]+]] {112// CHECK1-NEXT:  entry:113// CHECK1-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4114// CHECK1-NEXT:    [[A:%.*]] = alloca i32, align 4115// CHECK1-NEXT:    [[A_CASTED:%.*]] = alloca i64, align 8116// CHECK1-NEXT:    [[N_CASTED:%.*]] = alloca i64, align 8117// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x ptr], align 8118// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x ptr], align 8119// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [2 x ptr], align 8120// CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8121// CHECK1-NEXT:    [[A_CASTED1:%.*]] = alloca i64, align 8122// CHECK1-NEXT:    [[N_CASTED2:%.*]] = alloca i64, align 8123// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS3:%.*]] = alloca [2 x ptr], align 8124// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS4:%.*]] = alloca [2 x ptr], align 8125// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS5:%.*]] = alloca [2 x ptr], align 8126// CHECK1-NEXT:    [[KERNEL_ARGS6:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8127// CHECK1-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4128// CHECK1-NEXT:    store i32 0, ptr [[A]], align 4129// CHECK1-NEXT:    [[TMP0:%.*]] = load i32, ptr [[A]], align 4130// CHECK1-NEXT:    store i32 [[TMP0]], ptr [[A_CASTED]], align 4131// CHECK1-NEXT:    [[TMP1:%.*]] = load i64, ptr [[A_CASTED]], align 8132// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4133// CHECK1-NEXT:    store i32 [[TMP2]], ptr [[N_CASTED]], align 4134// CHECK1-NEXT:    [[TMP3:%.*]] = load i64, ptr [[N_CASTED]], align 8135// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0136// CHECK1-NEXT:    store i64 [[TMP1]], ptr [[TMP4]], align 8137// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0138// CHECK1-NEXT:    store i64 [[TMP1]], ptr [[TMP5]], align 8139// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0140// CHECK1-NEXT:    store ptr null, ptr [[TMP6]], align 8141// CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1142// CHECK1-NEXT:    store i64 [[TMP3]], ptr [[TMP7]], align 8143// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1144// CHECK1-NEXT:    store i64 [[TMP3]], ptr [[TMP8]], align 8145// CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1146// CHECK1-NEXT:    store ptr null, ptr [[TMP9]], align 8147// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0148// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0149// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0150// CHECK1-NEXT:    store i32 3, ptr [[TMP12]], align 4151// CHECK1-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1152// CHECK1-NEXT:    store i32 2, ptr [[TMP13]], align 4153// CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2154// CHECK1-NEXT:    store ptr [[TMP10]], ptr [[TMP14]], align 8155// CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3156// CHECK1-NEXT:    store ptr [[TMP11]], ptr [[TMP15]], align 8157// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4158// CHECK1-NEXT:    store ptr @.offload_sizes, ptr [[TMP16]], align 8159// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5160// CHECK1-NEXT:    store ptr @.offload_maptypes, ptr [[TMP17]], align 8161// CHECK1-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6162// CHECK1-NEXT:    store ptr null, ptr [[TMP18]], align 8163// CHECK1-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7164// CHECK1-NEXT:    store ptr null, ptr [[TMP19]], align 8165// CHECK1-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8166// CHECK1-NEXT:    store i64 0, ptr [[TMP20]], align 8167// CHECK1-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9168// CHECK1-NEXT:    store i64 0, ptr [[TMP21]], align 8169// CHECK1-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10170// CHECK1-NEXT:    store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP22]], align 4171// CHECK1-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11172// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP23]], align 4173// CHECK1-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12174// CHECK1-NEXT:    store i32 0, ptr [[TMP24]], align 4175// CHECK1-NEXT:    [[TMP25:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3bari_l95.region_id, ptr [[KERNEL_ARGS]])176// CHECK1-NEXT:    [[TMP26:%.*]] = icmp ne i32 [[TMP25]], 0177// CHECK1-NEXT:    br i1 [[TMP26]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]178// CHECK1:       omp_offload.failed:179// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3bari_l95(i64 [[TMP1]], i64 [[TMP3]]) #[[ATTR2:[0-9]+]]180// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]181// CHECK1:       omp_offload.cont:182// CHECK1-NEXT:    [[TMP27:%.*]] = load i32, ptr [[N_ADDR]], align 4183// CHECK1-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_ZL7fstatici(i32 noundef signext [[TMP27]])184// CHECK1-NEXT:    [[TMP28:%.*]] = load i32, ptr [[A]], align 4185// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP28]], [[CALL]]186// CHECK1-NEXT:    store i32 [[ADD]], ptr [[A]], align 4187// CHECK1-NEXT:    [[TMP29:%.*]] = load i32, ptr [[A]], align 4188// CHECK1-NEXT:    store i32 [[TMP29]], ptr [[A_CASTED1]], align 4189// CHECK1-NEXT:    [[TMP30:%.*]] = load i64, ptr [[A_CASTED1]], align 8190// CHECK1-NEXT:    [[TMP31:%.*]] = load i32, ptr [[N_ADDR]], align 4191// CHECK1-NEXT:    store i32 [[TMP31]], ptr [[N_CASTED2]], align 4192// CHECK1-NEXT:    [[TMP32:%.*]] = load i64, ptr [[N_CASTED2]], align 8193// CHECK1-NEXT:    [[TMP33:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0194// CHECK1-NEXT:    store i64 [[TMP30]], ptr [[TMP33]], align 8195// CHECK1-NEXT:    [[TMP34:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 0196// CHECK1-NEXT:    store i64 [[TMP30]], ptr [[TMP34]], align 8197// CHECK1-NEXT:    [[TMP35:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS5]], i64 0, i64 0198// CHECK1-NEXT:    store ptr null, ptr [[TMP35]], align 8199// CHECK1-NEXT:    [[TMP36:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 1200// CHECK1-NEXT:    store i64 [[TMP32]], ptr [[TMP36]], align 8201// CHECK1-NEXT:    [[TMP37:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 1202// CHECK1-NEXT:    store i64 [[TMP32]], ptr [[TMP37]], align 8203// CHECK1-NEXT:    [[TMP38:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS5]], i64 0, i64 1204// CHECK1-NEXT:    store ptr null, ptr [[TMP38]], align 8205// CHECK1-NEXT:    [[TMP39:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0206// CHECK1-NEXT:    [[TMP40:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 0207// CHECK1-NEXT:    [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 0208// CHECK1-NEXT:    store i32 3, ptr [[TMP41]], align 4209// CHECK1-NEXT:    [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 1210// CHECK1-NEXT:    store i32 2, ptr [[TMP42]], align 4211// CHECK1-NEXT:    [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 2212// CHECK1-NEXT:    store ptr [[TMP39]], ptr [[TMP43]], align 8213// CHECK1-NEXT:    [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 3214// CHECK1-NEXT:    store ptr [[TMP40]], ptr [[TMP44]], align 8215// CHECK1-NEXT:    [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 4216// CHECK1-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP45]], align 8217// CHECK1-NEXT:    [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 5218// CHECK1-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP46]], align 8219// CHECK1-NEXT:    [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 6220// CHECK1-NEXT:    store ptr null, ptr [[TMP47]], align 8221// CHECK1-NEXT:    [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 7222// CHECK1-NEXT:    store ptr null, ptr [[TMP48]], align 8223// CHECK1-NEXT:    [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 8224// CHECK1-NEXT:    store i64 0, ptr [[TMP49]], align 8225// CHECK1-NEXT:    [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 9226// CHECK1-NEXT:    store i64 0, ptr [[TMP50]], align 8227// CHECK1-NEXT:    [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 10228// CHECK1-NEXT:    store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP51]], align 4229// CHECK1-NEXT:    [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 11230// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP52]], align 4231// CHECK1-NEXT:    [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 12232// CHECK1-NEXT:    store i32 0, ptr [[TMP53]], align 4233// CHECK1-NEXT:    [[TMP54:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3bari_l103.region_id, ptr [[KERNEL_ARGS6]])234// CHECK1-NEXT:    [[TMP55:%.*]] = icmp ne i32 [[TMP54]], 0235// CHECK1-NEXT:    br i1 [[TMP55]], label [[OMP_OFFLOAD_FAILED7:%.*]], label [[OMP_OFFLOAD_CONT8:%.*]]236// CHECK1:       omp_offload.failed7:237// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3bari_l103(i64 [[TMP30]], i64 [[TMP32]]) #[[ATTR2]]238// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT8]]239// CHECK1:       omp_offload.cont8:240// CHECK1-NEXT:    [[TMP56:%.*]] = load i32, ptr [[A]], align 4241// CHECK1-NEXT:    ret i32 [[TMP56]]242//243//244// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3bari_l95245// CHECK1-SAME: (i64 noundef [[A:%.*]], i64 noundef [[N:%.*]]) #[[ATTR1:[0-9]+]] {246// CHECK1-NEXT:  entry:247// CHECK1-NEXT:    [[A_ADDR:%.*]] = alloca i64, align 8248// CHECK1-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 8249// CHECK1-NEXT:    [[S:%.*]] = alloca [[STRUCT_S1:%.*]], align 8250// CHECK1-NEXT:    store i64 [[A]], ptr [[A_ADDR]], align 8251// CHECK1-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 8252// CHECK1-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4253// CHECK1-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_ZN2S12r1Ei(ptr noundef nonnull align 8 dereferenceable(8) [[S]], i32 noundef signext [[TMP0]])254// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 4255// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CALL]]256// CHECK1-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 4257// CHECK1-NEXT:    ret void258//259//260// CHECK1-LABEL: define {{[^@]+}}@_ZN2S12r1Ei261// CHECK1-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[THIS:%.*]], i32 noundef signext [[N:%.*]]) #[[ATTR0]] comdat {262// CHECK1-NEXT:  entry:263// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8264// CHECK1-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4265// CHECK1-NEXT:    [[B:%.*]] = alloca i32, align 4266// CHECK1-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])267// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8268// CHECK1-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4269// CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8270// CHECK1-NEXT:    store i32 1, ptr [[B]], align 4271// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 4272// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[B]], align 4273// CHECK1-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP1]], [[TMP2]]274// CHECK1-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[SUB]], i32 1, ptr @.str)275// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 2, ptr @_ZN2S12r1Ei.omp_outlined, ptr [[THIS1]], ptr [[B]])276// CHECK1-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 1024, i32 2, ptr null)277// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 1, ptr @_ZN2S12r1Ei.omp_outlined.3, ptr [[THIS1]])278// CHECK1-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[THIS1]], i32 0, i32 0279// CHECK1-NEXT:    [[TMP3:%.*]] = load double, ptr [[A]], align 8280// CHECK1-NEXT:    [[CONV:%.*]] = fptosi double [[TMP3]] to i32281// CHECK1-NEXT:    ret i32 [[CONV]]282//283//284// CHECK1-LABEL: define {{[^@]+}}@_ZL7fstatici285// CHECK1-SAME: (i32 noundef signext [[N:%.*]]) #[[ATTR0]] {286// CHECK1-NEXT:  entry:287// CHECK1-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4288// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4289// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca ptr, align 8290// CHECK1-NEXT:    [[TMP:%.*]] = alloca ptr, align 8291// CHECK1-NEXT:    [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 8292// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x ptr], align 8293// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x ptr], align 8294// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [2 x ptr], align 8295// CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8296// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_2:%.*]] = alloca i32, align 4297// CHECK1-NEXT:    [[DOTCAPTURE_EXPR__CASTED3:%.*]] = alloca i64, align 8298// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS4:%.*]] = alloca [1 x ptr], align 8299// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS5:%.*]] = alloca [1 x ptr], align 8300// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS6:%.*]] = alloca [1 x ptr], align 8301// CHECK1-NEXT:    [[KERNEL_ARGS7:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8302// CHECK1-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4303// CHECK1-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4304// CHECK1-NEXT:    store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 4305// CHECK1-NEXT:    store ptr @.str, ptr [[DOTCAPTURE_EXPR_1]], align 8306// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTCAPTURE_EXPR_1]], align 8, !nonnull [[META11:![0-9]+]]307// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 8308// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4309// CHECK1-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 4310// CHECK1-NEXT:    [[TMP3:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 8311// CHECK1-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[TMP]], align 8, !nonnull [[META11]]312// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0313// CHECK1-NEXT:    store i64 [[TMP3]], ptr [[TMP5]], align 8314// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0315// CHECK1-NEXT:    store i64 [[TMP3]], ptr [[TMP6]], align 8316// CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0317// CHECK1-NEXT:    store ptr null, ptr [[TMP7]], align 8318// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1319// CHECK1-NEXT:    store ptr [[TMP4]], ptr [[TMP8]], align 8320// CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1321// CHECK1-NEXT:    store ptr [[TMP4]], ptr [[TMP9]], align 8322// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1323// CHECK1-NEXT:    store ptr null, ptr [[TMP10]], align 8324// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0325// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0326// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4327// CHECK1-NEXT:    [[TMP14:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP13]], 0328// CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0329// CHECK1-NEXT:    store i32 3, ptr [[TMP15]], align 4330// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1331// CHECK1-NEXT:    store i32 2, ptr [[TMP16]], align 4332// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2333// CHECK1-NEXT:    store ptr [[TMP11]], ptr [[TMP17]], align 8334// CHECK1-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3335// CHECK1-NEXT:    store ptr [[TMP12]], ptr [[TMP18]], align 8336// CHECK1-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4337// CHECK1-NEXT:    store ptr @.offload_sizes.4, ptr [[TMP19]], align 8338// CHECK1-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5339// CHECK1-NEXT:    store ptr @.offload_maptypes.5, ptr [[TMP20]], align 8340// CHECK1-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6341// CHECK1-NEXT:    store ptr null, ptr [[TMP21]], align 8342// CHECK1-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7343// CHECK1-NEXT:    store ptr null, ptr [[TMP22]], align 8344// CHECK1-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8345// CHECK1-NEXT:    store i64 0, ptr [[TMP23]], align 8346// CHECK1-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9347// CHECK1-NEXT:    store i64 0, ptr [[TMP24]], align 8348// CHECK1-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10349// CHECK1-NEXT:    store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP25]], align 4350// CHECK1-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11351// CHECK1-NEXT:    store [3 x i32] [[TMP14]], ptr [[TMP26]], align 4352// CHECK1-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12353// CHECK1-NEXT:    store i32 0, ptr [[TMP27]], align 4354// CHECK1-NEXT:    [[TMP28:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 1, i32 [[TMP13]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l61.region_id, ptr [[KERNEL_ARGS]])355// CHECK1-NEXT:    [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 0356// CHECK1-NEXT:    br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]357// CHECK1:       omp_offload.failed:358// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l61(i64 [[TMP3]], ptr [[TMP4]]) #[[ATTR2]]359// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]360// CHECK1:       omp_offload.cont:361// CHECK1-NEXT:    [[TMP30:%.*]] = load i32, ptr [[N_ADDR]], align 4362// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 32, [[TMP30]]363// CHECK1-NEXT:    store i32 [[ADD]], ptr [[DOTCAPTURE_EXPR_2]], align 4364// CHECK1-NEXT:    [[TMP31:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 4365// CHECK1-NEXT:    store i32 [[TMP31]], ptr [[DOTCAPTURE_EXPR__CASTED3]], align 4366// CHECK1-NEXT:    [[TMP32:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED3]], align 8367// CHECK1-NEXT:    [[TMP33:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 0368// CHECK1-NEXT:    store i64 [[TMP32]], ptr [[TMP33]], align 8369// CHECK1-NEXT:    [[TMP34:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 0370// CHECK1-NEXT:    store i64 [[TMP32]], ptr [[TMP34]], align 8371// CHECK1-NEXT:    [[TMP35:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i64 0, i64 0372// CHECK1-NEXT:    store ptr null, ptr [[TMP35]], align 8373// CHECK1-NEXT:    [[TMP36:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 0374// CHECK1-NEXT:    [[TMP37:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 0375// CHECK1-NEXT:    [[TMP38:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 4376// CHECK1-NEXT:    [[TMP39:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP38]], 0377// CHECK1-NEXT:    [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 0378// CHECK1-NEXT:    store i32 3, ptr [[TMP40]], align 4379// CHECK1-NEXT:    [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 1380// CHECK1-NEXT:    store i32 1, ptr [[TMP41]], align 4381// CHECK1-NEXT:    [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 2382// CHECK1-NEXT:    store ptr [[TMP36]], ptr [[TMP42]], align 8383// CHECK1-NEXT:    [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 3384// CHECK1-NEXT:    store ptr [[TMP37]], ptr [[TMP43]], align 8385// CHECK1-NEXT:    [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 4386// CHECK1-NEXT:    store ptr @.offload_sizes.6, ptr [[TMP44]], align 8387// CHECK1-NEXT:    [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 5388// CHECK1-NEXT:    store ptr @.offload_maptypes.7, ptr [[TMP45]], align 8389// CHECK1-NEXT:    [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 6390// CHECK1-NEXT:    store ptr null, ptr [[TMP46]], align 8391// CHECK1-NEXT:    [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 7392// CHECK1-NEXT:    store ptr null, ptr [[TMP47]], align 8393// CHECK1-NEXT:    [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 8394// CHECK1-NEXT:    store i64 0, ptr [[TMP48]], align 8395// CHECK1-NEXT:    [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 9396// CHECK1-NEXT:    store i64 0, ptr [[TMP49]], align 8397// CHECK1-NEXT:    [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 10398// CHECK1-NEXT:    store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP50]], align 4399// CHECK1-NEXT:    [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 11400// CHECK1-NEXT:    store [3 x i32] [[TMP39]], ptr [[TMP51]], align 4401// CHECK1-NEXT:    [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 12402// CHECK1-NEXT:    store i32 0, ptr [[TMP52]], align 4403// CHECK1-NEXT:    [[TMP53:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 1, i32 [[TMP38]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l65.region_id, ptr [[KERNEL_ARGS7]])404// CHECK1-NEXT:    [[TMP54:%.*]] = icmp ne i32 [[TMP53]], 0405// CHECK1-NEXT:    br i1 [[TMP54]], label [[OMP_OFFLOAD_FAILED8:%.*]], label [[OMP_OFFLOAD_CONT9:%.*]]406// CHECK1:       omp_offload.failed8:407// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l65(i64 [[TMP32]]) #[[ATTR2]]408// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT9]]409// CHECK1:       omp_offload.cont9:410// CHECK1-NEXT:    [[TMP55:%.*]] = load i32, ptr [[N_ADDR]], align 4411// CHECK1-NEXT:    [[ADD10:%.*]] = add nsw i32 [[TMP55]], 1412// CHECK1-NEXT:    ret i32 [[ADD10]]413//414//415// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3bari_l103416// CHECK1-SAME: (i64 noundef [[A:%.*]], i64 noundef [[N:%.*]]) #[[ATTR1]] {417// CHECK1-NEXT:  entry:418// CHECK1-NEXT:    [[A_ADDR:%.*]] = alloca i64, align 8419// CHECK1-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 8420// CHECK1-NEXT:    store i64 [[A]], ptr [[A_ADDR]], align 8421// CHECK1-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 8422// CHECK1-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4423// CHECK1-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_Z9ftemplateIiET_i(i32 noundef signext [[TMP0]])424// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 4425// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CALL]]426// CHECK1-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 4427// CHECK1-NEXT:    ret void428//429//430// CHECK1-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i431// CHECK1-SAME: (i32 noundef signext [[N:%.*]]) #[[ATTR0]] comdat {432// CHECK1-NEXT:  entry:433// CHECK1-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4434// CHECK1-NEXT:    [[A:%.*]] = alloca i32, align 4435// CHECK1-NEXT:    [[B:%.*]] = alloca i16, align 2436// CHECK1-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])437// CHECK1-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4438// CHECK1-NEXT:    store i32 0, ptr [[A]], align 4439// CHECK1-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 20, i32 2, ptr @.str)440// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @_Z9ftemplateIiET_i.omp_outlined)441// CHECK1-NEXT:    store i16 1, ptr [[B]], align 2442// CHECK1-NEXT:    [[TMP1:%.*]] = load i16, ptr [[B]], align 2443// CHECK1-NEXT:    [[TMP2:%.*]] = sext i16 [[TMP1]] to i32444// CHECK1-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 1, ptr @.str)445// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 2, ptr @_Z9ftemplateIiET_i.omp_outlined.8, ptr [[A]], ptr [[B]])446// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A]], align 4447// CHECK1-NEXT:    ret i32 [[TMP3]]448//449//450// CHECK1-LABEL: define {{[^@]+}}@_ZN2S12r1Ei.omp_outlined451// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]]) #[[ATTR1]] {452// CHECK1-NEXT:  entry:453// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8454// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8455// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8456// CHECK1-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8457// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8458// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8459// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8460// CHECK1-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8461// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8462// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 8, !nonnull [[META11]], !align [[META12:![0-9]+]]463// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4464// CHECK1-NEXT:    [[CONV:%.*]] = sitofp i32 [[TMP2]] to double465// CHECK1-NEXT:    [[ADD:%.*]] = fadd double [[CONV]], 1.500000e+00466// CHECK1-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 0467// CHECK1-NEXT:    store double [[ADD]], ptr [[A]], align 8468// CHECK1-NEXT:    ret void469//470//471// CHECK1-LABEL: define {{[^@]+}}@_ZN2S12r1Ei.omp_outlined.3472// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {473// CHECK1-NEXT:  entry:474// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8475// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8476// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8477// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8478// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8479// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8480// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8481// CHECK1-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 0482// CHECK1-NEXT:    store double 2.500000e+00, ptr [[A]], align 8483// CHECK1-NEXT:    ret void484//485//486// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l61487// CHECK1-SAME: (i64 noundef [[DOTCAPTURE_EXPR_:%.*]], ptr noundef nonnull align 1 dereferenceable(4) [[DOTCAPTURE_EXPR_1:%.*]]) #[[ATTR1]] {488// CHECK1-NEXT:  entry:489// CHECK1-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 8490// CHECK1-NEXT:    [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca ptr, align 8491// CHECK1-NEXT:    [[TMP:%.*]] = alloca ptr, align 8492// CHECK1-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])493// CHECK1-NEXT:    store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 8494// CHECK1-NEXT:    store ptr [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 8495// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTCAPTURE_EXPR__ADDR2]], align 8, !nonnull [[META11]]496// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 8497// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 4498// CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[TMP]], align 8, !nonnull [[META11]]499// CHECK1-NEXT:    [[ARRAYDECAY:%.*]] = getelementptr inbounds [4 x i8], ptr [[TMP3]], i64 0, i64 0500// CHECK1-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 2, ptr [[ARRAYDECAY]])501// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l61.omp_outlined)502// CHECK1-NEXT:    ret void503//504//505// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l61.omp_outlined506// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {507// CHECK1-NEXT:  entry:508// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8509// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8510// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8511// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8512// CHECK1-NEXT:    ret void513//514//515// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l65516// CHECK1-SAME: (i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1]] {517// CHECK1-NEXT:  entry:518// CHECK1-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 8519// CHECK1-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])520// CHECK1-NEXT:    store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 8521// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 4522// CHECK1-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 1, ptr null)523// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l65.omp_outlined)524// CHECK1-NEXT:    ret void525//526//527// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l65.omp_outlined528// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {529// CHECK1-NEXT:  entry:530// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8531// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8532// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8533// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8534// CHECK1-NEXT:    ret void535//536//537// CHECK1-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i.omp_outlined538// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {539// CHECK1-NEXT:  entry:540// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8541// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8542// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8543// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8544// CHECK1-NEXT:    ret void545//546//547// CHECK1-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i.omp_outlined.8548// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[B:%.*]]) #[[ATTR1]] {549// CHECK1-NEXT:  entry:550// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8551// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8552// CHECK1-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8553// CHECK1-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8554// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8555// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8556// CHECK1-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8557// CHECK1-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8558// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8, !nonnull [[META11]], !align [[META12]]559// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 8, !nonnull [[META11]], !align [[META15:![0-9]+]]560// CHECK1-NEXT:    [[TMP2:%.*]] = load i16, ptr [[TMP1]], align 2561// CHECK1-NEXT:    [[CONV:%.*]] = sext i16 [[TMP2]] to i32562// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 4563// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP3]], [[CONV]]564// CHECK1-NEXT:    store i32 [[ADD]], ptr [[TMP0]], align 4565// CHECK1-NEXT:    ret void566//567//568// CHECK3-LABEL: define {{[^@]+}}@_Z3bari569// CHECK3-SAME: (i32 noundef [[N:%.*]]) #[[ATTR0:[0-9]+]] {570// CHECK3-NEXT:  entry:571// CHECK3-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4572// CHECK3-NEXT:    [[A:%.*]] = alloca i32, align 4573// CHECK3-NEXT:    [[A_CASTED:%.*]] = alloca i32, align 4574// CHECK3-NEXT:    [[N_CASTED:%.*]] = alloca i32, align 4575// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x ptr], align 4576// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x ptr], align 4577// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [2 x ptr], align 4578// CHECK3-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8579// CHECK3-NEXT:    [[A_CASTED1:%.*]] = alloca i32, align 4580// CHECK3-NEXT:    [[N_CASTED2:%.*]] = alloca i32, align 4581// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS3:%.*]] = alloca [2 x ptr], align 4582// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS4:%.*]] = alloca [2 x ptr], align 4583// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS5:%.*]] = alloca [2 x ptr], align 4584// CHECK3-NEXT:    [[KERNEL_ARGS6:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8585// CHECK3-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4586// CHECK3-NEXT:    store i32 0, ptr [[A]], align 4587// CHECK3-NEXT:    [[TMP0:%.*]] = load i32, ptr [[A]], align 4588// CHECK3-NEXT:    store i32 [[TMP0]], ptr [[A_CASTED]], align 4589// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_CASTED]], align 4590// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4591// CHECK3-NEXT:    store i32 [[TMP2]], ptr [[N_CASTED]], align 4592// CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[N_CASTED]], align 4593// CHECK3-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0594// CHECK3-NEXT:    store i32 [[TMP1]], ptr [[TMP4]], align 4595// CHECK3-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0596// CHECK3-NEXT:    store i32 [[TMP1]], ptr [[TMP5]], align 4597// CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0598// CHECK3-NEXT:    store ptr null, ptr [[TMP6]], align 4599// CHECK3-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1600// CHECK3-NEXT:    store i32 [[TMP3]], ptr [[TMP7]], align 4601// CHECK3-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1602// CHECK3-NEXT:    store i32 [[TMP3]], ptr [[TMP8]], align 4603// CHECK3-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 1604// CHECK3-NEXT:    store ptr null, ptr [[TMP9]], align 4605// CHECK3-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0606// CHECK3-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0607// CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0608// CHECK3-NEXT:    store i32 3, ptr [[TMP12]], align 4609// CHECK3-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1610// CHECK3-NEXT:    store i32 2, ptr [[TMP13]], align 4611// CHECK3-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2612// CHECK3-NEXT:    store ptr [[TMP10]], ptr [[TMP14]], align 4613// CHECK3-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3614// CHECK3-NEXT:    store ptr [[TMP11]], ptr [[TMP15]], align 4615// CHECK3-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4616// CHECK3-NEXT:    store ptr @.offload_sizes, ptr [[TMP16]], align 4617// CHECK3-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5618// CHECK3-NEXT:    store ptr @.offload_maptypes, ptr [[TMP17]], align 4619// CHECK3-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6620// CHECK3-NEXT:    store ptr null, ptr [[TMP18]], align 4621// CHECK3-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7622// CHECK3-NEXT:    store ptr null, ptr [[TMP19]], align 4623// CHECK3-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8624// CHECK3-NEXT:    store i64 0, ptr [[TMP20]], align 8625// CHECK3-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9626// CHECK3-NEXT:    store i64 0, ptr [[TMP21]], align 8627// CHECK3-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10628// CHECK3-NEXT:    store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP22]], align 4629// CHECK3-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11630// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP23]], align 4631// CHECK3-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12632// CHECK3-NEXT:    store i32 0, ptr [[TMP24]], align 4633// CHECK3-NEXT:    [[TMP25:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3bari_l95.region_id, ptr [[KERNEL_ARGS]])634// CHECK3-NEXT:    [[TMP26:%.*]] = icmp ne i32 [[TMP25]], 0635// CHECK3-NEXT:    br i1 [[TMP26]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]636// CHECK3:       omp_offload.failed:637// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3bari_l95(i32 [[TMP1]], i32 [[TMP3]]) #[[ATTR2:[0-9]+]]638// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT]]639// CHECK3:       omp_offload.cont:640// CHECK3-NEXT:    [[TMP27:%.*]] = load i32, ptr [[N_ADDR]], align 4641// CHECK3-NEXT:    [[CALL:%.*]] = call noundef i32 @_ZL7fstatici(i32 noundef [[TMP27]])642// CHECK3-NEXT:    [[TMP28:%.*]] = load i32, ptr [[A]], align 4643// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP28]], [[CALL]]644// CHECK3-NEXT:    store i32 [[ADD]], ptr [[A]], align 4645// CHECK3-NEXT:    [[TMP29:%.*]] = load i32, ptr [[A]], align 4646// CHECK3-NEXT:    store i32 [[TMP29]], ptr [[A_CASTED1]], align 4647// CHECK3-NEXT:    [[TMP30:%.*]] = load i32, ptr [[A_CASTED1]], align 4648// CHECK3-NEXT:    [[TMP31:%.*]] = load i32, ptr [[N_ADDR]], align 4649// CHECK3-NEXT:    store i32 [[TMP31]], ptr [[N_CASTED2]], align 4650// CHECK3-NEXT:    [[TMP32:%.*]] = load i32, ptr [[N_CASTED2]], align 4651// CHECK3-NEXT:    [[TMP33:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0652// CHECK3-NEXT:    store i32 [[TMP30]], ptr [[TMP33]], align 4653// CHECK3-NEXT:    [[TMP34:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 0654// CHECK3-NEXT:    store i32 [[TMP30]], ptr [[TMP34]], align 4655// CHECK3-NEXT:    [[TMP35:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS5]], i32 0, i32 0656// CHECK3-NEXT:    store ptr null, ptr [[TMP35]], align 4657// CHECK3-NEXT:    [[TMP36:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 1658// CHECK3-NEXT:    store i32 [[TMP32]], ptr [[TMP36]], align 4659// CHECK3-NEXT:    [[TMP37:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 1660// CHECK3-NEXT:    store i32 [[TMP32]], ptr [[TMP37]], align 4661// CHECK3-NEXT:    [[TMP38:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS5]], i32 0, i32 1662// CHECK3-NEXT:    store ptr null, ptr [[TMP38]], align 4663// CHECK3-NEXT:    [[TMP39:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0664// CHECK3-NEXT:    [[TMP40:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 0665// CHECK3-NEXT:    [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 0666// CHECK3-NEXT:    store i32 3, ptr [[TMP41]], align 4667// CHECK3-NEXT:    [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 1668// CHECK3-NEXT:    store i32 2, ptr [[TMP42]], align 4669// CHECK3-NEXT:    [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 2670// CHECK3-NEXT:    store ptr [[TMP39]], ptr [[TMP43]], align 4671// CHECK3-NEXT:    [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 3672// CHECK3-NEXT:    store ptr [[TMP40]], ptr [[TMP44]], align 4673// CHECK3-NEXT:    [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 4674// CHECK3-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP45]], align 4675// CHECK3-NEXT:    [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 5676// CHECK3-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP46]], align 4677// CHECK3-NEXT:    [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 6678// CHECK3-NEXT:    store ptr null, ptr [[TMP47]], align 4679// CHECK3-NEXT:    [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 7680// CHECK3-NEXT:    store ptr null, ptr [[TMP48]], align 4681// CHECK3-NEXT:    [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 8682// CHECK3-NEXT:    store i64 0, ptr [[TMP49]], align 8683// CHECK3-NEXT:    [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 9684// CHECK3-NEXT:    store i64 0, ptr [[TMP50]], align 8685// CHECK3-NEXT:    [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 10686// CHECK3-NEXT:    store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP51]], align 4687// CHECK3-NEXT:    [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 11688// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP52]], align 4689// CHECK3-NEXT:    [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 12690// CHECK3-NEXT:    store i32 0, ptr [[TMP53]], align 4691// CHECK3-NEXT:    [[TMP54:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3bari_l103.region_id, ptr [[KERNEL_ARGS6]])692// CHECK3-NEXT:    [[TMP55:%.*]] = icmp ne i32 [[TMP54]], 0693// CHECK3-NEXT:    br i1 [[TMP55]], label [[OMP_OFFLOAD_FAILED7:%.*]], label [[OMP_OFFLOAD_CONT8:%.*]]694// CHECK3:       omp_offload.failed7:695// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3bari_l103(i32 [[TMP30]], i32 [[TMP32]]) #[[ATTR2]]696// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT8]]697// CHECK3:       omp_offload.cont8:698// CHECK3-NEXT:    [[TMP56:%.*]] = load i32, ptr [[A]], align 4699// CHECK3-NEXT:    ret i32 [[TMP56]]700//701//702// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3bari_l95703// CHECK3-SAME: (i32 noundef [[A:%.*]], i32 noundef [[N:%.*]]) #[[ATTR1:[0-9]+]] {704// CHECK3-NEXT:  entry:705// CHECK3-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 4706// CHECK3-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4707// CHECK3-NEXT:    [[S:%.*]] = alloca [[STRUCT_S1:%.*]], align 4708// CHECK3-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 4709// CHECK3-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4710// CHECK3-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4711// CHECK3-NEXT:    [[CALL:%.*]] = call noundef i32 @_ZN2S12r1Ei(ptr noundef nonnull align 4 dereferenceable(8) [[S]], i32 noundef [[TMP0]])712// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 4713// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CALL]]714// CHECK3-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 4715// CHECK3-NEXT:    ret void716//717//718// CHECK3-LABEL: define {{[^@]+}}@_ZN2S12r1Ei719// CHECK3-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]], i32 noundef [[N:%.*]]) #[[ATTR0]] comdat align 2 {720// CHECK3-NEXT:  entry:721// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4722// CHECK3-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4723// CHECK3-NEXT:    [[B:%.*]] = alloca i32, align 4724// CHECK3-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])725// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4726// CHECK3-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4727// CHECK3-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4728// CHECK3-NEXT:    store i32 1, ptr [[B]], align 4729// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 4730// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[B]], align 4731// CHECK3-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP1]], [[TMP2]]732// CHECK3-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[SUB]], i32 1, ptr @.str)733// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 2, ptr @_ZN2S12r1Ei.omp_outlined, ptr [[THIS1]], ptr [[B]])734// CHECK3-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 1024, i32 2, ptr null)735// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 1, ptr @_ZN2S12r1Ei.omp_outlined.3, ptr [[THIS1]])736// CHECK3-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[THIS1]], i32 0, i32 0737// CHECK3-NEXT:    [[TMP3:%.*]] = load double, ptr [[A]], align 4738// CHECK3-NEXT:    [[CONV:%.*]] = fptosi double [[TMP3]] to i32739// CHECK3-NEXT:    ret i32 [[CONV]]740//741//742// CHECK3-LABEL: define {{[^@]+}}@_ZL7fstatici743// CHECK3-SAME: (i32 noundef [[N:%.*]]) #[[ATTR0]] {744// CHECK3-NEXT:  entry:745// CHECK3-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4746// CHECK3-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4747// CHECK3-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca ptr, align 4748// CHECK3-NEXT:    [[TMP:%.*]] = alloca ptr, align 4749// CHECK3-NEXT:    [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i32, align 4750// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x ptr], align 4751// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x ptr], align 4752// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [2 x ptr], align 4753// CHECK3-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8754// CHECK3-NEXT:    [[DOTCAPTURE_EXPR_2:%.*]] = alloca i32, align 4755// CHECK3-NEXT:    [[DOTCAPTURE_EXPR__CASTED3:%.*]] = alloca i32, align 4756// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS4:%.*]] = alloca [1 x ptr], align 4757// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS5:%.*]] = alloca [1 x ptr], align 4758// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS6:%.*]] = alloca [1 x ptr], align 4759// CHECK3-NEXT:    [[KERNEL_ARGS7:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8760// CHECK3-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4761// CHECK3-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4762// CHECK3-NEXT:    store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 4763// CHECK3-NEXT:    store ptr @.str, ptr [[DOTCAPTURE_EXPR_1]], align 4764// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTCAPTURE_EXPR_1]], align 4, !nonnull [[META12:![0-9]+]]765// CHECK3-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 4766// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4767// CHECK3-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 4768// CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED]], align 4769// CHECK3-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[TMP]], align 4, !nonnull [[META12]]770// CHECK3-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0771// CHECK3-NEXT:    store i32 [[TMP3]], ptr [[TMP5]], align 4772// CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0773// CHECK3-NEXT:    store i32 [[TMP3]], ptr [[TMP6]], align 4774// CHECK3-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0775// CHECK3-NEXT:    store ptr null, ptr [[TMP7]], align 4776// CHECK3-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1777// CHECK3-NEXT:    store ptr [[TMP4]], ptr [[TMP8]], align 4778// CHECK3-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1779// CHECK3-NEXT:    store ptr [[TMP4]], ptr [[TMP9]], align 4780// CHECK3-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 1781// CHECK3-NEXT:    store ptr null, ptr [[TMP10]], align 4782// CHECK3-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0783// CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0784// CHECK3-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4785// CHECK3-NEXT:    [[TMP14:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP13]], 0786// CHECK3-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0787// CHECK3-NEXT:    store i32 3, ptr [[TMP15]], align 4788// CHECK3-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1789// CHECK3-NEXT:    store i32 2, ptr [[TMP16]], align 4790// CHECK3-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2791// CHECK3-NEXT:    store ptr [[TMP11]], ptr [[TMP17]], align 4792// CHECK3-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3793// CHECK3-NEXT:    store ptr [[TMP12]], ptr [[TMP18]], align 4794// CHECK3-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4795// CHECK3-NEXT:    store ptr @.offload_sizes.4, ptr [[TMP19]], align 4796// CHECK3-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5797// CHECK3-NEXT:    store ptr @.offload_maptypes.5, ptr [[TMP20]], align 4798// CHECK3-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6799// CHECK3-NEXT:    store ptr null, ptr [[TMP21]], align 4800// CHECK3-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7801// CHECK3-NEXT:    store ptr null, ptr [[TMP22]], align 4802// CHECK3-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8803// CHECK3-NEXT:    store i64 0, ptr [[TMP23]], align 8804// CHECK3-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9805// CHECK3-NEXT:    store i64 0, ptr [[TMP24]], align 8806// CHECK3-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10807// CHECK3-NEXT:    store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP25]], align 4808// CHECK3-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11809// CHECK3-NEXT:    store [3 x i32] [[TMP14]], ptr [[TMP26]], align 4810// CHECK3-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12811// CHECK3-NEXT:    store i32 0, ptr [[TMP27]], align 4812// CHECK3-NEXT:    [[TMP28:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 1, i32 [[TMP13]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l61.region_id, ptr [[KERNEL_ARGS]])813// CHECK3-NEXT:    [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 0814// CHECK3-NEXT:    br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]815// CHECK3:       omp_offload.failed:816// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l61(i32 [[TMP3]], ptr [[TMP4]]) #[[ATTR2]]817// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT]]818// CHECK3:       omp_offload.cont:819// CHECK3-NEXT:    [[TMP30:%.*]] = load i32, ptr [[N_ADDR]], align 4820// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 32, [[TMP30]]821// CHECK3-NEXT:    store i32 [[ADD]], ptr [[DOTCAPTURE_EXPR_2]], align 4822// CHECK3-NEXT:    [[TMP31:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 4823// CHECK3-NEXT:    store i32 [[TMP31]], ptr [[DOTCAPTURE_EXPR__CASTED3]], align 4824// CHECK3-NEXT:    [[TMP32:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED3]], align 4825// CHECK3-NEXT:    [[TMP33:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 0826// CHECK3-NEXT:    store i32 [[TMP32]], ptr [[TMP33]], align 4827// CHECK3-NEXT:    [[TMP34:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 0828// CHECK3-NEXT:    store i32 [[TMP32]], ptr [[TMP34]], align 4829// CHECK3-NEXT:    [[TMP35:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i32 0, i32 0830// CHECK3-NEXT:    store ptr null, ptr [[TMP35]], align 4831// CHECK3-NEXT:    [[TMP36:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 0832// CHECK3-NEXT:    [[TMP37:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 0833// CHECK3-NEXT:    [[TMP38:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 4834// CHECK3-NEXT:    [[TMP39:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP38]], 0835// CHECK3-NEXT:    [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 0836// CHECK3-NEXT:    store i32 3, ptr [[TMP40]], align 4837// CHECK3-NEXT:    [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 1838// CHECK3-NEXT:    store i32 1, ptr [[TMP41]], align 4839// CHECK3-NEXT:    [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 2840// CHECK3-NEXT:    store ptr [[TMP36]], ptr [[TMP42]], align 4841// CHECK3-NEXT:    [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 3842// CHECK3-NEXT:    store ptr [[TMP37]], ptr [[TMP43]], align 4843// CHECK3-NEXT:    [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 4844// CHECK3-NEXT:    store ptr @.offload_sizes.6, ptr [[TMP44]], align 4845// CHECK3-NEXT:    [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 5846// CHECK3-NEXT:    store ptr @.offload_maptypes.7, ptr [[TMP45]], align 4847// CHECK3-NEXT:    [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 6848// CHECK3-NEXT:    store ptr null, ptr [[TMP46]], align 4849// CHECK3-NEXT:    [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 7850// CHECK3-NEXT:    store ptr null, ptr [[TMP47]], align 4851// CHECK3-NEXT:    [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 8852// CHECK3-NEXT:    store i64 0, ptr [[TMP48]], align 8853// CHECK3-NEXT:    [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 9854// CHECK3-NEXT:    store i64 0, ptr [[TMP49]], align 8855// CHECK3-NEXT:    [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 10856// CHECK3-NEXT:    store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP50]], align 4857// CHECK3-NEXT:    [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 11858// CHECK3-NEXT:    store [3 x i32] [[TMP39]], ptr [[TMP51]], align 4859// CHECK3-NEXT:    [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 12860// CHECK3-NEXT:    store i32 0, ptr [[TMP52]], align 4861// CHECK3-NEXT:    [[TMP53:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 1, i32 [[TMP38]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l65.region_id, ptr [[KERNEL_ARGS7]])862// CHECK3-NEXT:    [[TMP54:%.*]] = icmp ne i32 [[TMP53]], 0863// CHECK3-NEXT:    br i1 [[TMP54]], label [[OMP_OFFLOAD_FAILED8:%.*]], label [[OMP_OFFLOAD_CONT9:%.*]]864// CHECK3:       omp_offload.failed8:865// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l65(i32 [[TMP32]]) #[[ATTR2]]866// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT9]]867// CHECK3:       omp_offload.cont9:868// CHECK3-NEXT:    [[TMP55:%.*]] = load i32, ptr [[N_ADDR]], align 4869// CHECK3-NEXT:    [[ADD10:%.*]] = add nsw i32 [[TMP55]], 1870// CHECK3-NEXT:    ret i32 [[ADD10]]871//872//873// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3bari_l103874// CHECK3-SAME: (i32 noundef [[A:%.*]], i32 noundef [[N:%.*]]) #[[ATTR1]] {875// CHECK3-NEXT:  entry:876// CHECK3-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 4877// CHECK3-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4878// CHECK3-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 4879// CHECK3-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4880// CHECK3-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4881// CHECK3-NEXT:    [[CALL:%.*]] = call noundef i32 @_Z9ftemplateIiET_i(i32 noundef [[TMP0]])882// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 4883// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CALL]]884// CHECK3-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 4885// CHECK3-NEXT:    ret void886//887//888// CHECK3-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i889// CHECK3-SAME: (i32 noundef [[N:%.*]]) #[[ATTR0]] comdat {890// CHECK3-NEXT:  entry:891// CHECK3-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4892// CHECK3-NEXT:    [[A:%.*]] = alloca i32, align 4893// CHECK3-NEXT:    [[B:%.*]] = alloca i16, align 2894// CHECK3-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])895// CHECK3-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4896// CHECK3-NEXT:    store i32 0, ptr [[A]], align 4897// CHECK3-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 20, i32 2, ptr @.str)898// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @_Z9ftemplateIiET_i.omp_outlined)899// CHECK3-NEXT:    store i16 1, ptr [[B]], align 2900// CHECK3-NEXT:    [[TMP1:%.*]] = load i16, ptr [[B]], align 2901// CHECK3-NEXT:    [[TMP2:%.*]] = sext i16 [[TMP1]] to i32902// CHECK3-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 1, ptr @.str)903// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 2, ptr @_Z9ftemplateIiET_i.omp_outlined.8, ptr [[A]], ptr [[B]])904// CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A]], align 4905// CHECK3-NEXT:    ret i32 [[TMP3]]906//907//908// CHECK3-LABEL: define {{[^@]+}}@_ZN2S12r1Ei.omp_outlined909// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]]) #[[ATTR1]] {910// CHECK3-NEXT:  entry:911// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4912// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4913// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4914// CHECK3-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4915// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4916// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4917// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4918// CHECK3-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4919// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4920// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 4, !nonnull [[META12]], !align [[META13:![0-9]+]]921// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4922// CHECK3-NEXT:    [[CONV:%.*]] = sitofp i32 [[TMP2]] to double923// CHECK3-NEXT:    [[ADD:%.*]] = fadd double [[CONV]], 1.500000e+00924// CHECK3-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 0925// CHECK3-NEXT:    store double [[ADD]], ptr [[A]], align 4926// CHECK3-NEXT:    ret void927//928//929// CHECK3-LABEL: define {{[^@]+}}@_ZN2S12r1Ei.omp_outlined.3930// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {931// CHECK3-NEXT:  entry:932// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4933// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4934// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4935// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4936// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4937// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4938// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4939// CHECK3-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 0940// CHECK3-NEXT:    store double 2.500000e+00, ptr [[A]], align 4941// CHECK3-NEXT:    ret void942//943//944// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l61945// CHECK3-SAME: (i32 noundef [[DOTCAPTURE_EXPR_:%.*]], ptr noundef nonnull align 1 dereferenceable(4) [[DOTCAPTURE_EXPR_1:%.*]]) #[[ATTR1]] {946// CHECK3-NEXT:  entry:947// CHECK3-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 4948// CHECK3-NEXT:    [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca ptr, align 4949// CHECK3-NEXT:    [[TMP:%.*]] = alloca ptr, align 4950// CHECK3-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])951// CHECK3-NEXT:    store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 4952// CHECK3-NEXT:    store ptr [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 4953// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTCAPTURE_EXPR__ADDR2]], align 4, !nonnull [[META12]]954// CHECK3-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 4955// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 4956// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[TMP]], align 4, !nonnull [[META12]]957// CHECK3-NEXT:    [[ARRAYDECAY:%.*]] = getelementptr inbounds [4 x i8], ptr [[TMP3]], i32 0, i32 0958// CHECK3-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 2, ptr [[ARRAYDECAY]])959// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l61.omp_outlined)960// CHECK3-NEXT:    ret void961//962//963// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l61.omp_outlined964// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {965// CHECK3-NEXT:  entry:966// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4967// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4968// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4969// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4970// CHECK3-NEXT:    ret void971//972//973// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l65974// CHECK3-SAME: (i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1]] {975// CHECK3-NEXT:  entry:976// CHECK3-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 4977// CHECK3-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])978// CHECK3-NEXT:    store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 4979// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 4980// CHECK3-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 1, ptr null)981// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l65.omp_outlined)982// CHECK3-NEXT:    ret void983//984//985// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l65.omp_outlined986// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {987// CHECK3-NEXT:  entry:988// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4989// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4990// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4991// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4992// CHECK3-NEXT:    ret void993//994//995// CHECK3-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i.omp_outlined996// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {997// CHECK3-NEXT:  entry:998// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4999// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41000// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41001// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41002// CHECK3-NEXT:    ret void1003//1004//1005// CHECK3-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i.omp_outlined.81006// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[B:%.*]]) #[[ATTR1]] {1007// CHECK3-NEXT:  entry:1008// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41009// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41010// CHECK3-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 41011// CHECK3-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 41012// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41013// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41014// CHECK3-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 41015// CHECK3-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 41016// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 4, !nonnull [[META12]], !align [[META13]]1017// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 4, !nonnull [[META12]], !align [[META16:![0-9]+]]1018// CHECK3-NEXT:    [[TMP2:%.*]] = load i16, ptr [[TMP1]], align 21019// CHECK3-NEXT:    [[CONV:%.*]] = sext i16 [[TMP2]] to i321020// CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 41021// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP3]], [[CONV]]1022// CHECK3-NEXT:    store i32 [[ADD]], ptr [[TMP0]], align 41023// CHECK3-NEXT:    ret void1024//1025//1026// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l611027// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]], ptr noundef nonnull align 1 dereferenceable(4) [[DOTCAPTURE_EXPR_1:%.*]]) #[[ATTR0:[0-9]+]] {1028// CHECK9-NEXT:  entry:1029// CHECK9-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81030// CHECK9-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 81031// CHECK9-NEXT:    [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca ptr, align 81032// CHECK9-NEXT:    [[TMP:%.*]] = alloca ptr, align 81033// CHECK9-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])1034// CHECK9-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81035// CHECK9-NEXT:    store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 81036// CHECK9-NEXT:    store ptr [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 81037// CHECK9-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTCAPTURE_EXPR__ADDR2]], align 8, !nonnull [[META12:![0-9]+]]1038// CHECK9-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 81039// CHECK9-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 41040// CHECK9-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[TMP]], align 8, !nonnull [[META12]]1041// CHECK9-NEXT:    [[ARRAYDECAY:%.*]] = getelementptr inbounds [4 x i8], ptr [[TMP3]], i64 0, i64 01042// CHECK9-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 2, ptr [[ARRAYDECAY]])1043// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l61.omp_outlined)1044// CHECK9-NEXT:    ret void1045//1046//1047// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l61.omp_outlined1048// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {1049// CHECK9-NEXT:  entry:1050// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81051// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81052// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81053// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81054// CHECK9-NEXT:    ret void1055//1056//1057// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l651058// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0]] {1059// CHECK9-NEXT:  entry:1060// CHECK9-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81061// CHECK9-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 81062// CHECK9-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1063// CHECK9-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81064// CHECK9-NEXT:    store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 81065// CHECK9-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 41066// CHECK9-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 1, ptr null)1067// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l65.omp_outlined)1068// CHECK9-NEXT:    ret void1069//1070//1071// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l65.omp_outlined1072// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {1073// CHECK9-NEXT:  entry:1074// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81075// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81076// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81077// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81078// CHECK9-NEXT:    ret void1079//1080//1081// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3bari_l951082// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[A:%.*]], i64 noundef [[N:%.*]]) #[[ATTR0]] {1083// CHECK9-NEXT:  entry:1084// CHECK9-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81085// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca i64, align 81086// CHECK9-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 81087// CHECK9-NEXT:    [[S:%.*]] = alloca [[STRUCT_S1:%.*]], align 81088// CHECK9-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81089// CHECK9-NEXT:    store i64 [[A]], ptr [[A_ADDR]], align 81090// CHECK9-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 81091// CHECK9-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 41092// CHECK9-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_ZN2S12r1Ei(ptr noundef nonnull align 8 dereferenceable(8) [[S]], i32 noundef signext [[TMP0]])1093// CHECK9-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 41094// CHECK9-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CALL]]1095// CHECK9-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 41096// CHECK9-NEXT:    ret void1097//1098//1099// CHECK9-LABEL: define {{[^@]+}}@_ZN2S12r1Ei1100// CHECK9-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[THIS:%.*]], i32 noundef signext [[N:%.*]]) #[[ATTR2:[0-9]+]] comdat {1101// CHECK9-NEXT:  entry:1102// CHECK9-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 81103// CHECK9-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41104// CHECK9-NEXT:    [[B:%.*]] = alloca i32, align 41105// CHECK9-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1106// CHECK9-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 81107// CHECK9-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41108// CHECK9-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81109// CHECK9-NEXT:    store i32 1, ptr [[B]], align 41110// CHECK9-NEXT:    [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41111// CHECK9-NEXT:    [[TMP2:%.*]] = load i32, ptr [[B]], align 41112// CHECK9-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP1]], [[TMP2]]1113// CHECK9-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[SUB]], i32 1, ptr @.str)1114// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 2, ptr @_ZN2S12r1Ei.omp_outlined, ptr [[THIS1]], ptr [[B]])1115// CHECK9-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 1024, i32 2, ptr null)1116// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 1, ptr @_ZN2S12r1Ei.omp_outlined.1, ptr [[THIS1]])1117// CHECK9-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[THIS1]], i32 0, i32 01118// CHECK9-NEXT:    [[TMP3:%.*]] = load double, ptr [[A]], align 81119// CHECK9-NEXT:    [[CONV:%.*]] = fptosi double [[TMP3]] to i321120// CHECK9-NEXT:    ret i32 [[CONV]]1121//1122//1123// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3bari_l1031124// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[A:%.*]], i64 noundef [[N:%.*]]) #[[ATTR0]] {1125// CHECK9-NEXT:  entry:1126// CHECK9-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81127// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca i64, align 81128// CHECK9-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 81129// CHECK9-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81130// CHECK9-NEXT:    store i64 [[A]], ptr [[A_ADDR]], align 81131// CHECK9-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 81132// CHECK9-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 41133// CHECK9-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_Z9ftemplateIiET_i(i32 noundef signext [[TMP0]])1134// CHECK9-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 41135// CHECK9-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CALL]]1136// CHECK9-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 41137// CHECK9-NEXT:    ret void1138//1139//1140// CHECK9-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i1141// CHECK9-SAME: (i32 noundef signext [[N:%.*]]) #[[ATTR2]] comdat {1142// CHECK9-NEXT:  entry:1143// CHECK9-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41144// CHECK9-NEXT:    [[A:%.*]] = alloca i32, align 41145// CHECK9-NEXT:    [[B:%.*]] = alloca i16, align 21146// CHECK9-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1147// CHECK9-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41148// CHECK9-NEXT:    store i32 0, ptr [[A]], align 41149// CHECK9-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 20, i32 2, ptr @.str)1150// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @_Z9ftemplateIiET_i.omp_outlined)1151// CHECK9-NEXT:    store i16 1, ptr [[B]], align 21152// CHECK9-NEXT:    [[TMP1:%.*]] = load i16, ptr [[B]], align 21153// CHECK9-NEXT:    [[TMP2:%.*]] = sext i16 [[TMP1]] to i321154// CHECK9-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 1, ptr @.str)1155// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 2, ptr @_Z9ftemplateIiET_i.omp_outlined.2, ptr [[A]], ptr [[B]])1156// CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A]], align 41157// CHECK9-NEXT:    ret i32 [[TMP3]]1158//1159//1160// CHECK9-LABEL: define {{[^@]+}}@_ZN2S12r1Ei.omp_outlined1161// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]]) #[[ATTR0]] {1162// CHECK9-NEXT:  entry:1163// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81164// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81165// CHECK9-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 81166// CHECK9-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 81167// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81168// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81169// CHECK9-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 81170// CHECK9-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 81171// CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81172// CHECK9-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 8, !nonnull [[META12]], !align [[META15:![0-9]+]]1173// CHECK9-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41174// CHECK9-NEXT:    [[CONV:%.*]] = sitofp i32 [[TMP2]] to double1175// CHECK9-NEXT:    [[ADD:%.*]] = fadd double [[CONV]], 1.500000e+001176// CHECK9-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 01177// CHECK9-NEXT:    store double [[ADD]], ptr [[A]], align 81178// CHECK9-NEXT:    ret void1179//1180//1181// CHECK9-LABEL: define {{[^@]+}}@_ZN2S12r1Ei.omp_outlined.11182// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR0]] {1183// CHECK9-NEXT:  entry:1184// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81185// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81186// CHECK9-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 81187// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81188// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81189// CHECK9-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 81190// CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81191// CHECK9-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 01192// CHECK9-NEXT:    store double 2.500000e+00, ptr [[A]], align 81193// CHECK9-NEXT:    ret void1194//1195//1196// CHECK9-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i.omp_outlined1197// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {1198// CHECK9-NEXT:  entry:1199// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81200// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81201// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81202// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81203// CHECK9-NEXT:    ret void1204//1205//1206// CHECK9-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i.omp_outlined.21207// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[B:%.*]]) #[[ATTR0]] {1208// CHECK9-NEXT:  entry:1209// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81210// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81211// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81212// CHECK9-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 81213// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81214// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81215// CHECK9-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81216// CHECK9-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 81217// CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8, !nonnull [[META12]], !align [[META15]]1218// CHECK9-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 8, !nonnull [[META12]], !align [[META16:![0-9]+]]1219// CHECK9-NEXT:    [[TMP2:%.*]] = load i16, ptr [[TMP1]], align 21220// CHECK9-NEXT:    [[CONV:%.*]] = sext i16 [[TMP2]] to i321221// CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 41222// CHECK9-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP3]], [[CONV]]1223// CHECK9-NEXT:    store i32 [[ADD]], ptr [[TMP0]], align 41224// CHECK9-NEXT:    ret void1225//1226//1227// CHECK10-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3bari_l951228// CHECK10-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[A:%.*]], i64 noundef [[N:%.*]]) #[[ATTR0:[0-9]+]] {1229// CHECK10-NEXT:  entry:1230// CHECK10-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81231// CHECK10-NEXT:    [[A_ADDR:%.*]] = alloca i64, align 81232// CHECK10-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 81233// CHECK10-NEXT:    [[S:%.*]] = alloca [[STRUCT_S1:%.*]], align 81234// CHECK10-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81235// CHECK10-NEXT:    store i64 [[A]], ptr [[A_ADDR]], align 81236// CHECK10-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 81237// CHECK10-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 41238// CHECK10-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_ZN2S12r1Ei(ptr noundef nonnull align 8 dereferenceable(8) [[S]], i32 noundef signext [[TMP0]])1239// CHECK10-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 41240// CHECK10-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CALL]]1241// CHECK10-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 41242// CHECK10-NEXT:    ret void1243//1244//1245// CHECK10-LABEL: define {{[^@]+}}@_ZN2S12r1Ei1246// CHECK10-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[THIS:%.*]], i32 noundef signext [[N:%.*]]) #[[ATTR1:[0-9]+]] comdat {1247// CHECK10-NEXT:  entry:1248// CHECK10-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 81249// CHECK10-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41250// CHECK10-NEXT:    [[B:%.*]] = alloca i32, align 41251// CHECK10-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])1252// CHECK10-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 81253// CHECK10-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41254// CHECK10-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81255// CHECK10-NEXT:    store i32 1, ptr [[B]], align 41256// CHECK10-NEXT:    [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41257// CHECK10-NEXT:    [[TMP2:%.*]] = load i32, ptr [[B]], align 41258// CHECK10-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP1]], [[TMP2]]1259// CHECK10-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[SUB]], i32 1, ptr @.str)1260// CHECK10-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 2, ptr @_ZN2S12r1Ei.omp_outlined, ptr [[THIS1]], ptr [[B]])1261// CHECK10-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 1024, i32 2, ptr null)1262// CHECK10-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 1, ptr @_ZN2S12r1Ei.omp_outlined.1, ptr [[THIS1]])1263// CHECK10-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[THIS1]], i32 0, i32 01264// CHECK10-NEXT:    [[TMP3:%.*]] = load double, ptr [[A]], align 81265// CHECK10-NEXT:    [[CONV:%.*]] = fptosi double [[TMP3]] to i321266// CHECK10-NEXT:    ret i32 [[CONV]]1267//1268//1269// CHECK10-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3bari_l1031270// CHECK10-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[A:%.*]], i64 noundef [[N:%.*]]) #[[ATTR0]] {1271// CHECK10-NEXT:  entry:1272// CHECK10-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81273// CHECK10-NEXT:    [[A_ADDR:%.*]] = alloca i64, align 81274// CHECK10-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 81275// CHECK10-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81276// CHECK10-NEXT:    store i64 [[A]], ptr [[A_ADDR]], align 81277// CHECK10-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 81278// CHECK10-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 41279// CHECK10-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_Z9ftemplateIiET_i(i32 noundef signext [[TMP0]])1280// CHECK10-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 41281// CHECK10-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CALL]]1282// CHECK10-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 41283// CHECK10-NEXT:    ret void1284//1285//1286// CHECK10-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i1287// CHECK10-SAME: (i32 noundef signext [[N:%.*]]) #[[ATTR1]] comdat {1288// CHECK10-NEXT:  entry:1289// CHECK10-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41290// CHECK10-NEXT:    [[A:%.*]] = alloca i32, align 41291// CHECK10-NEXT:    [[B:%.*]] = alloca i16, align 21292// CHECK10-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1293// CHECK10-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41294// CHECK10-NEXT:    store i32 0, ptr [[A]], align 41295// CHECK10-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 20, i32 2, ptr @.str)1296// CHECK10-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @_Z9ftemplateIiET_i.omp_outlined)1297// CHECK10-NEXT:    store i16 1, ptr [[B]], align 21298// CHECK10-NEXT:    [[TMP1:%.*]] = load i16, ptr [[B]], align 21299// CHECK10-NEXT:    [[TMP2:%.*]] = sext i16 [[TMP1]] to i321300// CHECK10-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 1, ptr @.str)1301// CHECK10-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 2, ptr @_Z9ftemplateIiET_i.omp_outlined.2, ptr [[A]], ptr [[B]])1302// CHECK10-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A]], align 41303// CHECK10-NEXT:    ret i32 [[TMP3]]1304//1305//1306// CHECK10-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l611307// CHECK10-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]], ptr noundef nonnull align 1 dereferenceable(4) [[DOTCAPTURE_EXPR_1:%.*]]) #[[ATTR0]] {1308// CHECK10-NEXT:  entry:1309// CHECK10-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81310// CHECK10-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 81311// CHECK10-NEXT:    [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca ptr, align 81312// CHECK10-NEXT:    [[TMP:%.*]] = alloca ptr, align 81313// CHECK10-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1314// CHECK10-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81315// CHECK10-NEXT:    store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 81316// CHECK10-NEXT:    store ptr [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 81317// CHECK10-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTCAPTURE_EXPR__ADDR2]], align 8, !nonnull [[META12:![0-9]+]]1318// CHECK10-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 81319// CHECK10-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 41320// CHECK10-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[TMP]], align 8, !nonnull [[META12]]1321// CHECK10-NEXT:    [[ARRAYDECAY:%.*]] = getelementptr inbounds [4 x i8], ptr [[TMP3]], i64 0, i64 01322// CHECK10-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 2, ptr [[ARRAYDECAY]])1323// CHECK10-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l61.omp_outlined)1324// CHECK10-NEXT:    ret void1325//1326//1327// CHECK10-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l61.omp_outlined1328// CHECK10-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {1329// CHECK10-NEXT:  entry:1330// CHECK10-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81331// CHECK10-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81332// CHECK10-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81333// CHECK10-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81334// CHECK10-NEXT:    ret void1335//1336//1337// CHECK10-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l651338// CHECK10-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0]] {1339// CHECK10-NEXT:  entry:1340// CHECK10-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81341// CHECK10-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 81342// CHECK10-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1343// CHECK10-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81344// CHECK10-NEXT:    store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 81345// CHECK10-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 41346// CHECK10-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 1, ptr null)1347// CHECK10-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l65.omp_outlined)1348// CHECK10-NEXT:    ret void1349//1350//1351// CHECK10-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l65.omp_outlined1352// CHECK10-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {1353// CHECK10-NEXT:  entry:1354// CHECK10-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81355// CHECK10-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81356// CHECK10-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81357// CHECK10-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81358// CHECK10-NEXT:    ret void1359//1360//1361// CHECK10-LABEL: define {{[^@]+}}@_ZN2S12r1Ei.omp_outlined1362// CHECK10-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]]) #[[ATTR0]] {1363// CHECK10-NEXT:  entry:1364// CHECK10-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81365// CHECK10-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81366// CHECK10-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 81367// CHECK10-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 81368// CHECK10-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81369// CHECK10-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81370// CHECK10-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 81371// CHECK10-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 81372// CHECK10-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81373// CHECK10-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 8, !nonnull [[META12]], !align [[META15:![0-9]+]]1374// CHECK10-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41375// CHECK10-NEXT:    [[CONV:%.*]] = sitofp i32 [[TMP2]] to double1376// CHECK10-NEXT:    [[ADD:%.*]] = fadd double [[CONV]], 1.500000e+001377// CHECK10-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 01378// CHECK10-NEXT:    store double [[ADD]], ptr [[A]], align 81379// CHECK10-NEXT:    ret void1380//1381//1382// CHECK10-LABEL: define {{[^@]+}}@_ZN2S12r1Ei.omp_outlined.11383// CHECK10-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR0]] {1384// CHECK10-NEXT:  entry:1385// CHECK10-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81386// CHECK10-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81387// CHECK10-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 81388// CHECK10-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81389// CHECK10-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81390// CHECK10-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 81391// CHECK10-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81392// CHECK10-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 01393// CHECK10-NEXT:    store double 2.500000e+00, ptr [[A]], align 81394// CHECK10-NEXT:    ret void1395//1396//1397// CHECK10-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i.omp_outlined1398// CHECK10-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {1399// CHECK10-NEXT:  entry:1400// CHECK10-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81401// CHECK10-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81402// CHECK10-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81403// CHECK10-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81404// CHECK10-NEXT:    ret void1405//1406//1407// CHECK10-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i.omp_outlined.21408// CHECK10-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[B:%.*]]) #[[ATTR0]] {1409// CHECK10-NEXT:  entry:1410// CHECK10-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81411// CHECK10-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81412// CHECK10-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 81413// CHECK10-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 81414// CHECK10-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81415// CHECK10-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81416// CHECK10-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 81417// CHECK10-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 81418// CHECK10-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8, !nonnull [[META12]], !align [[META15]]1419// CHECK10-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 8, !nonnull [[META12]], !align [[META16:![0-9]+]]1420// CHECK10-NEXT:    [[TMP2:%.*]] = load i16, ptr [[TMP1]], align 21421// CHECK10-NEXT:    [[CONV:%.*]] = sext i16 [[TMP2]] to i321422// CHECK10-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 41423// CHECK10-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP3]], [[CONV]]1424// CHECK10-NEXT:    store i32 [[ADD]], ptr [[TMP0]], align 41425// CHECK10-NEXT:    ret void1426//1427//1428// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l611429// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]], ptr noundef nonnull align 1 dereferenceable(4) [[DOTCAPTURE_EXPR_1:%.*]]) #[[ATTR0:[0-9]+]] {1430// CHECK11-NEXT:  entry:1431// CHECK11-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41432// CHECK11-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 41433// CHECK11-NEXT:    [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca ptr, align 41434// CHECK11-NEXT:    [[TMP:%.*]] = alloca ptr, align 41435// CHECK11-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])1436// CHECK11-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41437// CHECK11-NEXT:    store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 41438// CHECK11-NEXT:    store ptr [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 41439// CHECK11-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTCAPTURE_EXPR__ADDR2]], align 4, !nonnull [[META13:![0-9]+]]1440// CHECK11-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 41441// CHECK11-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 41442// CHECK11-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[TMP]], align 4, !nonnull [[META13]]1443// CHECK11-NEXT:    [[ARRAYDECAY:%.*]] = getelementptr inbounds [4 x i8], ptr [[TMP3]], i32 0, i32 01444// CHECK11-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 2, ptr [[ARRAYDECAY]])1445// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l61.omp_outlined)1446// CHECK11-NEXT:    ret void1447//1448//1449// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l61.omp_outlined1450// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {1451// CHECK11-NEXT:  entry:1452// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41453// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41454// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41455// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41456// CHECK11-NEXT:    ret void1457//1458//1459// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l651460// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0]] {1461// CHECK11-NEXT:  entry:1462// CHECK11-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41463// CHECK11-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 41464// CHECK11-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1465// CHECK11-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41466// CHECK11-NEXT:    store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 41467// CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 41468// CHECK11-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 1, ptr null)1469// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l65.omp_outlined)1470// CHECK11-NEXT:    ret void1471//1472//1473// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l65.omp_outlined1474// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {1475// CHECK11-NEXT:  entry:1476// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41477// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41478// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41479// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41480// CHECK11-NEXT:    ret void1481//1482//1483// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3bari_l951484// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[A:%.*]], i32 noundef [[N:%.*]]) #[[ATTR0]] {1485// CHECK11-NEXT:  entry:1486// CHECK11-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41487// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 41488// CHECK11-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41489// CHECK11-NEXT:    [[S:%.*]] = alloca [[STRUCT_S1:%.*]], align 41490// CHECK11-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41491// CHECK11-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 41492// CHECK11-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41493// CHECK11-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 41494// CHECK11-NEXT:    [[CALL:%.*]] = call noundef i32 @_ZN2S12r1Ei(ptr noundef nonnull align 4 dereferenceable(8) [[S]], i32 noundef [[TMP0]])1495// CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 41496// CHECK11-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CALL]]1497// CHECK11-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 41498// CHECK11-NEXT:    ret void1499//1500//1501// CHECK11-LABEL: define {{[^@]+}}@_ZN2S12r1Ei1502// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]], i32 noundef [[N:%.*]]) #[[ATTR2:[0-9]+]] comdat align 2 {1503// CHECK11-NEXT:  entry:1504// CHECK11-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 41505// CHECK11-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41506// CHECK11-NEXT:    [[B:%.*]] = alloca i32, align 41507// CHECK11-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1508// CHECK11-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 41509// CHECK11-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41510// CHECK11-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41511// CHECK11-NEXT:    store i32 1, ptr [[B]], align 41512// CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41513// CHECK11-NEXT:    [[TMP2:%.*]] = load i32, ptr [[B]], align 41514// CHECK11-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP1]], [[TMP2]]1515// CHECK11-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[SUB]], i32 1, ptr @.str)1516// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 2, ptr @_ZN2S12r1Ei.omp_outlined, ptr [[THIS1]], ptr [[B]])1517// CHECK11-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 1024, i32 2, ptr null)1518// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 1, ptr @_ZN2S12r1Ei.omp_outlined.1, ptr [[THIS1]])1519// CHECK11-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[THIS1]], i32 0, i32 01520// CHECK11-NEXT:    [[TMP3:%.*]] = load double, ptr [[A]], align 41521// CHECK11-NEXT:    [[CONV:%.*]] = fptosi double [[TMP3]] to i321522// CHECK11-NEXT:    ret i32 [[CONV]]1523//1524//1525// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3bari_l1031526// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[A:%.*]], i32 noundef [[N:%.*]]) #[[ATTR0]] {1527// CHECK11-NEXT:  entry:1528// CHECK11-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41529// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 41530// CHECK11-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41531// CHECK11-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41532// CHECK11-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 41533// CHECK11-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41534// CHECK11-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 41535// CHECK11-NEXT:    [[CALL:%.*]] = call noundef i32 @_Z9ftemplateIiET_i(i32 noundef [[TMP0]])1536// CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 41537// CHECK11-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CALL]]1538// CHECK11-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 41539// CHECK11-NEXT:    ret void1540//1541//1542// CHECK11-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i1543// CHECK11-SAME: (i32 noundef [[N:%.*]]) #[[ATTR2]] comdat {1544// CHECK11-NEXT:  entry:1545// CHECK11-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41546// CHECK11-NEXT:    [[A:%.*]] = alloca i32, align 41547// CHECK11-NEXT:    [[B:%.*]] = alloca i16, align 21548// CHECK11-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1549// CHECK11-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41550// CHECK11-NEXT:    store i32 0, ptr [[A]], align 41551// CHECK11-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 20, i32 2, ptr @.str)1552// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @_Z9ftemplateIiET_i.omp_outlined)1553// CHECK11-NEXT:    store i16 1, ptr [[B]], align 21554// CHECK11-NEXT:    [[TMP1:%.*]] = load i16, ptr [[B]], align 21555// CHECK11-NEXT:    [[TMP2:%.*]] = sext i16 [[TMP1]] to i321556// CHECK11-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 1, ptr @.str)1557// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 2, ptr @_Z9ftemplateIiET_i.omp_outlined.2, ptr [[A]], ptr [[B]])1558// CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A]], align 41559// CHECK11-NEXT:    ret i32 [[TMP3]]1560//1561//1562// CHECK11-LABEL: define {{[^@]+}}@_ZN2S12r1Ei.omp_outlined1563// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]]) #[[ATTR0]] {1564// CHECK11-NEXT:  entry:1565// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41566// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41567// CHECK11-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 41568// CHECK11-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 41569// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41570// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41571// CHECK11-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 41572// CHECK11-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 41573// CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41574// CHECK11-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 4, !nonnull [[META13]], !align [[META16:![0-9]+]]1575// CHECK11-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41576// CHECK11-NEXT:    [[CONV:%.*]] = sitofp i32 [[TMP2]] to double1577// CHECK11-NEXT:    [[ADD:%.*]] = fadd double [[CONV]], 1.500000e+001578// CHECK11-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 01579// CHECK11-NEXT:    store double [[ADD]], ptr [[A]], align 41580// CHECK11-NEXT:    ret void1581//1582//1583// CHECK11-LABEL: define {{[^@]+}}@_ZN2S12r1Ei.omp_outlined.11584// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR0]] {1585// CHECK11-NEXT:  entry:1586// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41587// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41588// CHECK11-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 41589// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41590// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41591// CHECK11-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 41592// CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41593// CHECK11-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 01594// CHECK11-NEXT:    store double 2.500000e+00, ptr [[A]], align 41595// CHECK11-NEXT:    ret void1596//1597//1598// CHECK11-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i.omp_outlined1599// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {1600// CHECK11-NEXT:  entry:1601// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41602// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41603// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41604// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41605// CHECK11-NEXT:    ret void1606//1607//1608// CHECK11-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i.omp_outlined.21609// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[B:%.*]]) #[[ATTR0]] {1610// CHECK11-NEXT:  entry:1611// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41612// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41613// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 41614// CHECK11-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 41615// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41616// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41617// CHECK11-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 41618// CHECK11-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 41619// CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 4, !nonnull [[META13]], !align [[META16]]1620// CHECK11-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 4, !nonnull [[META13]], !align [[META17:![0-9]+]]1621// CHECK11-NEXT:    [[TMP2:%.*]] = load i16, ptr [[TMP1]], align 21622// CHECK11-NEXT:    [[CONV:%.*]] = sext i16 [[TMP2]] to i321623// CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 41624// CHECK11-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP3]], [[CONV]]1625// CHECK11-NEXT:    store i32 [[ADD]], ptr [[TMP0]], align 41626// CHECK11-NEXT:    ret void1627//1628//1629// CHECK12-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3bari_l951630// CHECK12-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[A:%.*]], i32 noundef [[N:%.*]]) #[[ATTR0:[0-9]+]] {1631// CHECK12-NEXT:  entry:1632// CHECK12-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41633// CHECK12-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 41634// CHECK12-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41635// CHECK12-NEXT:    [[S:%.*]] = alloca [[STRUCT_S1:%.*]], align 41636// CHECK12-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41637// CHECK12-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 41638// CHECK12-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41639// CHECK12-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 41640// CHECK12-NEXT:    [[CALL:%.*]] = call noundef i32 @_ZN2S12r1Ei(ptr noundef nonnull align 4 dereferenceable(8) [[S]], i32 noundef [[TMP0]])1641// CHECK12-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 41642// CHECK12-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CALL]]1643// CHECK12-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 41644// CHECK12-NEXT:    ret void1645//1646//1647// CHECK12-LABEL: define {{[^@]+}}@_ZN2S12r1Ei1648// CHECK12-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]], i32 noundef [[N:%.*]]) #[[ATTR1:[0-9]+]] comdat align 2 {1649// CHECK12-NEXT:  entry:1650// CHECK12-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 41651// CHECK12-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41652// CHECK12-NEXT:    [[B:%.*]] = alloca i32, align 41653// CHECK12-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])1654// CHECK12-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 41655// CHECK12-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41656// CHECK12-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41657// CHECK12-NEXT:    store i32 1, ptr [[B]], align 41658// CHECK12-NEXT:    [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41659// CHECK12-NEXT:    [[TMP2:%.*]] = load i32, ptr [[B]], align 41660// CHECK12-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP1]], [[TMP2]]1661// CHECK12-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[SUB]], i32 1, ptr @.str)1662// CHECK12-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 2, ptr @_ZN2S12r1Ei.omp_outlined, ptr [[THIS1]], ptr [[B]])1663// CHECK12-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 1024, i32 2, ptr null)1664// CHECK12-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 1, ptr @_ZN2S12r1Ei.omp_outlined.1, ptr [[THIS1]])1665// CHECK12-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[THIS1]], i32 0, i32 01666// CHECK12-NEXT:    [[TMP3:%.*]] = load double, ptr [[A]], align 41667// CHECK12-NEXT:    [[CONV:%.*]] = fptosi double [[TMP3]] to i321668// CHECK12-NEXT:    ret i32 [[CONV]]1669//1670//1671// CHECK12-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3bari_l1031672// CHECK12-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[A:%.*]], i32 noundef [[N:%.*]]) #[[ATTR0]] {1673// CHECK12-NEXT:  entry:1674// CHECK12-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41675// CHECK12-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 41676// CHECK12-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41677// CHECK12-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41678// CHECK12-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 41679// CHECK12-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41680// CHECK12-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 41681// CHECK12-NEXT:    [[CALL:%.*]] = call noundef i32 @_Z9ftemplateIiET_i(i32 noundef [[TMP0]])1682// CHECK12-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 41683// CHECK12-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CALL]]1684// CHECK12-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 41685// CHECK12-NEXT:    ret void1686//1687//1688// CHECK12-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i1689// CHECK12-SAME: (i32 noundef [[N:%.*]]) #[[ATTR1]] comdat {1690// CHECK12-NEXT:  entry:1691// CHECK12-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41692// CHECK12-NEXT:    [[A:%.*]] = alloca i32, align 41693// CHECK12-NEXT:    [[B:%.*]] = alloca i16, align 21694// CHECK12-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1695// CHECK12-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41696// CHECK12-NEXT:    store i32 0, ptr [[A]], align 41697// CHECK12-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 20, i32 2, ptr @.str)1698// CHECK12-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @_Z9ftemplateIiET_i.omp_outlined)1699// CHECK12-NEXT:    store i16 1, ptr [[B]], align 21700// CHECK12-NEXT:    [[TMP1:%.*]] = load i16, ptr [[B]], align 21701// CHECK12-NEXT:    [[TMP2:%.*]] = sext i16 [[TMP1]] to i321702// CHECK12-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 1, ptr @.str)1703// CHECK12-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 2, ptr @_Z9ftemplateIiET_i.omp_outlined.2, ptr [[A]], ptr [[B]])1704// CHECK12-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A]], align 41705// CHECK12-NEXT:    ret i32 [[TMP3]]1706//1707//1708// CHECK12-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l611709// CHECK12-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]], ptr noundef nonnull align 1 dereferenceable(4) [[DOTCAPTURE_EXPR_1:%.*]]) #[[ATTR0]] {1710// CHECK12-NEXT:  entry:1711// CHECK12-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41712// CHECK12-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 41713// CHECK12-NEXT:    [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca ptr, align 41714// CHECK12-NEXT:    [[TMP:%.*]] = alloca ptr, align 41715// CHECK12-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1716// CHECK12-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41717// CHECK12-NEXT:    store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 41718// CHECK12-NEXT:    store ptr [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 41719// CHECK12-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTCAPTURE_EXPR__ADDR2]], align 4, !nonnull [[META13:![0-9]+]]1720// CHECK12-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 41721// CHECK12-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 41722// CHECK12-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[TMP]], align 4, !nonnull [[META13]]1723// CHECK12-NEXT:    [[ARRAYDECAY:%.*]] = getelementptr inbounds [4 x i8], ptr [[TMP3]], i32 0, i32 01724// CHECK12-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 2, ptr [[ARRAYDECAY]])1725// CHECK12-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l61.omp_outlined)1726// CHECK12-NEXT:    ret void1727//1728//1729// CHECK12-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l61.omp_outlined1730// CHECK12-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {1731// CHECK12-NEXT:  entry:1732// CHECK12-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41733// CHECK12-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41734// CHECK12-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41735// CHECK12-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41736// CHECK12-NEXT:    ret void1737//1738//1739// CHECK12-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l651740// CHECK12-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0]] {1741// CHECK12-NEXT:  entry:1742// CHECK12-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41743// CHECK12-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 41744// CHECK12-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1745// CHECK12-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41746// CHECK12-NEXT:    store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 41747// CHECK12-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 41748// CHECK12-NEXT:    call void @__kmpc_push_num_threads_strict(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 1, ptr null)1749// CHECK12-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l65.omp_outlined)1750// CHECK12-NEXT:    ret void1751//1752//1753// CHECK12-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l65.omp_outlined1754// CHECK12-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {1755// CHECK12-NEXT:  entry:1756// CHECK12-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41757// CHECK12-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41758// CHECK12-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41759// CHECK12-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41760// CHECK12-NEXT:    ret void1761//1762//1763// CHECK12-LABEL: define {{[^@]+}}@_ZN2S12r1Ei.omp_outlined1764// CHECK12-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]]) #[[ATTR0]] {1765// CHECK12-NEXT:  entry:1766// CHECK12-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41767// CHECK12-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41768// CHECK12-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 41769// CHECK12-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 41770// CHECK12-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41771// CHECK12-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41772// CHECK12-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 41773// CHECK12-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 41774// CHECK12-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41775// CHECK12-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 4, !nonnull [[META13]], !align [[META16:![0-9]+]]1776// CHECK12-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 41777// CHECK12-NEXT:    [[CONV:%.*]] = sitofp i32 [[TMP2]] to double1778// CHECK12-NEXT:    [[ADD:%.*]] = fadd double [[CONV]], 1.500000e+001779// CHECK12-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 01780// CHECK12-NEXT:    store double [[ADD]], ptr [[A]], align 41781// CHECK12-NEXT:    ret void1782//1783//1784// CHECK12-LABEL: define {{[^@]+}}@_ZN2S12r1Ei.omp_outlined.11785// CHECK12-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR0]] {1786// CHECK12-NEXT:  entry:1787// CHECK12-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41788// CHECK12-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41789// CHECK12-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 41790// CHECK12-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41791// CHECK12-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41792// CHECK12-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 41793// CHECK12-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41794// CHECK12-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 01795// CHECK12-NEXT:    store double 2.500000e+00, ptr [[A]], align 41796// CHECK12-NEXT:    ret void1797//1798//1799// CHECK12-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i.omp_outlined1800// CHECK12-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {1801// CHECK12-NEXT:  entry:1802// CHECK12-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41803// CHECK12-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41804// CHECK12-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41805// CHECK12-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41806// CHECK12-NEXT:    ret void1807//1808//1809// CHECK12-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i.omp_outlined.21810// CHECK12-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[B:%.*]]) #[[ATTR0]] {1811// CHECK12-NEXT:  entry:1812// CHECK12-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41813// CHECK12-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41814// CHECK12-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 41815// CHECK12-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 41816// CHECK12-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41817// CHECK12-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41818// CHECK12-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 41819// CHECK12-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 41820// CHECK12-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 4, !nonnull [[META13]], !align [[META16]]1821// CHECK12-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 4, !nonnull [[META13]], !align [[META17:![0-9]+]]1822// CHECK12-NEXT:    [[TMP2:%.*]] = load i16, ptr [[TMP1]], align 21823// CHECK12-NEXT:    [[CONV:%.*]] = sext i16 [[TMP2]] to i321824// CHECK12-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 41825// CHECK12-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP3]], [[CONV]]1826// CHECK12-NEXT:    store i32 [[ADD]], ptr [[TMP0]], align 41827// CHECK12-NEXT:    ret void1828//1829