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=45 -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=45 -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=45 -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=45 -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=45 -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=45 -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=45 -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=45 -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=45 -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=45 -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=45 -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=45 -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=45 -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=45 -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=45 -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=45 -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=CHECK922// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -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=45 -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=45 -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=45 -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=CHECK1126 27// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -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=45 -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=45 -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=45 -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=45 -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=45 -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=45 -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=45 -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// Test host codegen.37// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK138// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s39// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK140// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK341// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s42// RUN: %clang_cc1 -fopenmp -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=CHECK343 44// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"45// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s46// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"47// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"48// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s49// RUN: %clang_cc1 -fopenmp-simd -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}}"50 51// Test target codegen - host bc file has to be created first.52// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc53// RUN: %clang_cc1 -verify -fopenmp -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=CHECK954// RUN: %clang_cc1 -fopenmp -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 %s55// RUN: %clang_cc1 -fopenmp -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=CHECK956// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc57// RUN: %clang_cc1 -verify -fopenmp -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=CHECK1158// RUN: %clang_cc1 -fopenmp -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 %s59// RUN: %clang_cc1 -fopenmp -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=CHECK1160 61// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc62// RUN: %clang_cc1 -verify -fopenmp-simd -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}}"63// RUN: %clang_cc1 -fopenmp-simd -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 %s64// RUN: %clang_cc1 -fopenmp-simd -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}}"65// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc66// RUN: %clang_cc1 -verify -fopenmp-simd -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}}"67// RUN: %clang_cc1 -fopenmp-simd -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 %s68// RUN: %clang_cc1 -fopenmp-simd -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}}"69 70// expected-no-diagnostics71#ifndef HEADER72#define HEADER73 74 75 76 77// We have 6 target regions78 79 80 81// Check target registration is registered as a Ctor.82 83 84template<typename tx>85tx ftemplate(int n) {86  tx a = 0;87 88  #pragma omp target teams ompx_dyn_cgroup_mem(tx(20))89  {90  }91 92  short b = 1;93  #pragma omp target teams num_teams(b) ompx_dyn_cgroup_mem(1024)94  {95    a += b;96  }97 98  return a;99}100 101static102int fstatic(int n) {103 104  #pragma omp target teams distribute parallel for simd num_teams(n) ompx_dyn_cgroup_mem(n*32)105  for (int i = 0; i < n ; ++i) {106  }107 108  #pragma omp target teams ompx_dyn_cgroup_mem(32+n) nowait109  {110  }111 112  return n+1;113}114 115struct S1 {116  double a;117 118  int r1(int n){119    int b = 1;120 121    #pragma omp target teams ompx_dyn_cgroup_mem(n-b)122    {123      this->a = (double)b + 1.5;124    }125 126    #pragma omp target ompx_dyn_cgroup_mem(1024)127    {128      this->a = 2.5;129    }130 131    return (int)a;132  }133};134 135int bar(int n){136  int a = 0;137 138  S1 S;139  a += S.r1(n);140 141  a += fstatic(n);142 143  a += ftemplate<int>(n);144 145  return a;146}147 148 149 150 151 152 153 154 155 156 157 158 159 160 161 162 163 164 165 166 167 168// Check that the offloading functions are emitted and that the parallel function169// is appropriately guarded.170 171 172 173 174 175 176#endif177 178// CHECK1-LABEL: define {{[^@]+}}@_Z3bari179// CHECK1-SAME: (i32 noundef signext [[N:%.*]]) #[[ATTR0:[0-9]+]] {180// CHECK1-NEXT:  entry:181// CHECK1-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4182// CHECK1-NEXT:    [[A:%.*]] = alloca i32, align 4183// CHECK1-NEXT:    [[S:%.*]] = alloca [[STRUCT_S1:%.*]], align 8184// CHECK1-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4185// CHECK1-NEXT:    store i32 0, ptr [[A]], align 4186// CHECK1-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4187// CHECK1-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_ZN2S12r1Ei(ptr noundef nonnull align 8 dereferenceable(8) [[S]], i32 noundef signext [[TMP0]])188// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A]], align 4189// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CALL]]190// CHECK1-NEXT:    store i32 [[ADD]], ptr [[A]], align 4191// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4192// CHECK1-NEXT:    [[CALL1:%.*]] = call noundef signext i32 @_ZL7fstatici(i32 noundef signext [[TMP2]])193// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A]], align 4194// CHECK1-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP3]], [[CALL1]]195// CHECK1-NEXT:    store i32 [[ADD2]], ptr [[A]], align 4196// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[N_ADDR]], align 4197// CHECK1-NEXT:    [[CALL3:%.*]] = call noundef signext i32 @_Z9ftemplateIiET_i(i32 noundef signext [[TMP4]])198// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[A]], align 4199// CHECK1-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP5]], [[CALL3]]200// CHECK1-NEXT:    store i32 [[ADD4]], ptr [[A]], align 4201// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[A]], align 4202// CHECK1-NEXT:    ret i32 [[TMP6]]203//204//205// CHECK1-LABEL: define {{[^@]+}}@_ZN2S12r1Ei206// CHECK1-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[THIS:%.*]], i32 noundef signext [[N:%.*]]) #[[ATTR0]] comdat {207// CHECK1-NEXT:  entry:208// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8209// CHECK1-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4210// CHECK1-NEXT:    [[B:%.*]] = alloca i32, align 4211// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4212// CHECK1-NEXT:    [[B_CASTED:%.*]] = alloca i64, align 8213// CHECK1-NEXT:    [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 8214// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 8215// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 8216// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 8217// CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8218// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS3:%.*]] = alloca [1 x ptr], align 8219// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS4:%.*]] = alloca [1 x ptr], align 8220// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS5:%.*]] = alloca [1 x ptr], align 8221// CHECK1-NEXT:    [[KERNEL_ARGS6:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8222// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8223// CHECK1-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4224// CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8225// CHECK1-NEXT:    store i32 1, ptr [[B]], align 4226// CHECK1-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4227// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[B]], align 4228// CHECK1-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP0]], [[TMP1]]229// CHECK1-NEXT:    store i32 [[SUB]], ptr [[DOTCAPTURE_EXPR_]], align 4230// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[B]], align 4231// CHECK1-NEXT:    store i32 [[TMP2]], ptr [[B_CASTED]], align 4232// CHECK1-NEXT:    [[TMP3:%.*]] = load i64, ptr [[B_CASTED]], align 8233// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4234// CHECK1-NEXT:    store i32 [[TMP4]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 4235// CHECK1-NEXT:    [[TMP5:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 8236// CHECK1-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[THIS1]], i32 0, i32 0237// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0238// CHECK1-NEXT:    store ptr [[THIS1]], ptr [[TMP6]], align 8239// CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0240// CHECK1-NEXT:    store ptr [[A]], ptr [[TMP7]], align 8241// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0242// CHECK1-NEXT:    store ptr null, ptr [[TMP8]], align 8243// CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1244// CHECK1-NEXT:    store i64 [[TMP3]], ptr [[TMP9]], align 8245// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1246// CHECK1-NEXT:    store i64 [[TMP3]], ptr [[TMP10]], align 8247// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1248// CHECK1-NEXT:    store ptr null, ptr [[TMP11]], align 8249// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2250// CHECK1-NEXT:    store i64 [[TMP5]], ptr [[TMP12]], align 8251// CHECK1-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2252// CHECK1-NEXT:    store i64 [[TMP5]], ptr [[TMP13]], align 8253// CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2254// CHECK1-NEXT:    store ptr null, ptr [[TMP14]], align 8255// CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0256// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0257// CHECK1-NEXT:    [[TMP17:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4258// CHECK1-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0259// CHECK1-NEXT:    store i32 3, ptr [[TMP18]], align 4260// CHECK1-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1261// CHECK1-NEXT:    store i32 3, ptr [[TMP19]], align 4262// CHECK1-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2263// CHECK1-NEXT:    store ptr [[TMP15]], ptr [[TMP20]], align 8264// CHECK1-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3265// CHECK1-NEXT:    store ptr [[TMP16]], ptr [[TMP21]], align 8266// CHECK1-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4267// CHECK1-NEXT:    store ptr @.offload_sizes, ptr [[TMP22]], align 8268// CHECK1-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5269// CHECK1-NEXT:    store ptr @.offload_maptypes, ptr [[TMP23]], align 8270// CHECK1-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6271// CHECK1-NEXT:    store ptr null, ptr [[TMP24]], align 8272// CHECK1-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7273// CHECK1-NEXT:    store ptr null, ptr [[TMP25]], align 8274// CHECK1-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8275// CHECK1-NEXT:    store i64 0, ptr [[TMP26]], align 8276// CHECK1-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9277// CHECK1-NEXT:    store i64 0, ptr [[TMP27]], align 8278// CHECK1-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10279// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP28]], align 4280// CHECK1-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11281// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP29]], align 4282// CHECK1-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12283// CHECK1-NEXT:    store i32 [[TMP17]], ptr [[TMP30]], align 4284// CHECK1-NEXT:    [[TMP31:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.region_id, ptr [[KERNEL_ARGS]])285// CHECK1-NEXT:    [[TMP32:%.*]] = icmp ne i32 [[TMP31]], 0286// CHECK1-NEXT:    br i1 [[TMP32]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]287// CHECK1:       omp_offload.failed:288// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121(ptr [[THIS1]], i64 [[TMP3]], i64 [[TMP5]]) #[[ATTR2:[0-9]+]]289// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]290// CHECK1:       omp_offload.cont:291// CHECK1-NEXT:    [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[THIS1]], i32 0, i32 0292// CHECK1-NEXT:    [[TMP33:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0293// CHECK1-NEXT:    store ptr [[THIS1]], ptr [[TMP33]], align 8294// CHECK1-NEXT:    [[TMP34:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 0295// CHECK1-NEXT:    store ptr [[A2]], ptr [[TMP34]], align 8296// CHECK1-NEXT:    [[TMP35:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS5]], i64 0, i64 0297// CHECK1-NEXT:    store ptr null, ptr [[TMP35]], align 8298// CHECK1-NEXT:    [[TMP36:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0299// CHECK1-NEXT:    [[TMP37:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 0300// CHECK1-NEXT:    [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 0301// CHECK1-NEXT:    store i32 3, ptr [[TMP38]], align 4302// CHECK1-NEXT:    [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 1303// CHECK1-NEXT:    store i32 1, ptr [[TMP39]], align 4304// CHECK1-NEXT:    [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 2305// CHECK1-NEXT:    store ptr [[TMP36]], ptr [[TMP40]], align 8306// CHECK1-NEXT:    [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 3307// CHECK1-NEXT:    store ptr [[TMP37]], ptr [[TMP41]], align 8308// CHECK1-NEXT:    [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 4309// CHECK1-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP42]], align 8310// CHECK1-NEXT:    [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 5311// CHECK1-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP43]], align 8312// CHECK1-NEXT:    [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 6313// CHECK1-NEXT:    store ptr null, ptr [[TMP44]], align 8314// CHECK1-NEXT:    [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 7315// CHECK1-NEXT:    store ptr null, ptr [[TMP45]], align 8316// CHECK1-NEXT:    [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 8317// CHECK1-NEXT:    store i64 0, ptr [[TMP46]], align 8318// CHECK1-NEXT:    [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 9319// CHECK1-NEXT:    store i64 0, ptr [[TMP47]], align 8320// CHECK1-NEXT:    [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 10321// CHECK1-NEXT:    store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP48]], align 4322// CHECK1-NEXT:    [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 11323// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP49]], align 4324// CHECK1-NEXT:    [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 12325// CHECK1-NEXT:    store i32 1024, ptr [[TMP50]], align 4326// CHECK1-NEXT:    [[TMP51:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126.region_id, ptr [[KERNEL_ARGS6]])327// CHECK1-NEXT:    [[TMP52:%.*]] = icmp ne i32 [[TMP51]], 0328// CHECK1-NEXT:    br i1 [[TMP52]], label [[OMP_OFFLOAD_FAILED7:%.*]], label [[OMP_OFFLOAD_CONT8:%.*]]329// CHECK1:       omp_offload.failed7:330// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126(ptr [[THIS1]]) #[[ATTR2]]331// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT8]]332// CHECK1:       omp_offload.cont8:333// CHECK1-NEXT:    [[A9:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[THIS1]], i32 0, i32 0334// CHECK1-NEXT:    [[TMP53:%.*]] = load double, ptr [[A9]], align 8335// CHECK1-NEXT:    [[CONV:%.*]] = fptosi double [[TMP53]] to i32336// CHECK1-NEXT:    ret i32 [[CONV]]337//338//339// CHECK1-LABEL: define {{[^@]+}}@_ZL7fstatici340// CHECK1-SAME: (i32 noundef signext [[N:%.*]]) #[[ATTR0]] {341// CHECK1-NEXT:  entry:342// CHECK1-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4343// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4344// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4345// CHECK1-NEXT:    [[N_CASTED:%.*]] = alloca i64, align 8346// CHECK1-NEXT:    [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 8347// CHECK1-NEXT:    [[DOTCAPTURE_EXPR__CASTED2:%.*]] = alloca i64, align 8348// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 8349// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 8350// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 8351// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4352// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_3:%.*]] = alloca i32, align 4353// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_4:%.*]] = alloca i32, align 4354// CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8355// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_6:%.*]] = alloca i32, align 4356// CHECK1-NEXT:    [[DOTCAPTURE_EXPR__CASTED8:%.*]] = alloca i64, align 8357// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS9:%.*]] = alloca [1 x ptr], align 8358// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS10:%.*]] = alloca [1 x ptr], align 8359// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS11:%.*]] = alloca [1 x ptr], align 8360// CHECK1-NEXT:    [[AGG_CAPTURED:%.*]] = alloca [[STRUCT_ANON:%.*]], align 4361// CHECK1-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])362// CHECK1-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4363// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 4364// CHECK1-NEXT:    store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 4365// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4366// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP2]], 32367// CHECK1-NEXT:    store i32 [[MUL]], ptr [[DOTCAPTURE_EXPR_1]], align 4368// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[N_ADDR]], align 4369// CHECK1-NEXT:    store i32 [[TMP3]], ptr [[N_CASTED]], align 4370// CHECK1-NEXT:    [[TMP4:%.*]] = load i64, ptr [[N_CASTED]], align 8371// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4372// CHECK1-NEXT:    store i32 [[TMP5]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 4373// CHECK1-NEXT:    [[TMP6:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 8374// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4375// CHECK1-NEXT:    store i32 [[TMP7]], ptr [[DOTCAPTURE_EXPR__CASTED2]], align 4376// CHECK1-NEXT:    [[TMP8:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED2]], align 8377// CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0378// CHECK1-NEXT:    store i64 [[TMP4]], ptr [[TMP9]], align 8379// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0380// CHECK1-NEXT:    store i64 [[TMP4]], ptr [[TMP10]], align 8381// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0382// CHECK1-NEXT:    store ptr null, ptr [[TMP11]], align 8383// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1384// CHECK1-NEXT:    store i64 [[TMP6]], ptr [[TMP12]], align 8385// CHECK1-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1386// CHECK1-NEXT:    store i64 [[TMP6]], ptr [[TMP13]], align 8387// CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1388// CHECK1-NEXT:    store ptr null, ptr [[TMP14]], align 8389// CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2390// CHECK1-NEXT:    store i64 [[TMP8]], ptr [[TMP15]], align 8391// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2392// CHECK1-NEXT:    store i64 [[TMP8]], ptr [[TMP16]], align 8393// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2394// CHECK1-NEXT:    store ptr null, ptr [[TMP17]], align 8395// CHECK1-NEXT:    [[TMP18:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0396// CHECK1-NEXT:    [[TMP19:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0397// CHECK1-NEXT:    [[TMP20:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4398// CHECK1-NEXT:    [[TMP21:%.*]] = load i32, ptr [[N_ADDR]], align 4399// CHECK1-NEXT:    store i32 [[TMP21]], ptr [[DOTCAPTURE_EXPR_3]], align 4400// CHECK1-NEXT:    [[TMP22:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_3]], align 4401// CHECK1-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP22]], 0402// CHECK1-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1403// CHECK1-NEXT:    [[SUB5:%.*]] = sub nsw i32 [[DIV]], 1404// CHECK1-NEXT:    store i32 [[SUB5]], ptr [[DOTCAPTURE_EXPR_4]], align 4405// CHECK1-NEXT:    [[TMP23:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 4406// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP23]], 1407// CHECK1-NEXT:    [[TMP24:%.*]] = zext i32 [[ADD]] to i64408// CHECK1-NEXT:    [[TMP25:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4409// CHECK1-NEXT:    [[TMP26:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP20]], 0410// CHECK1-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0411// CHECK1-NEXT:    store i32 3, ptr [[TMP27]], align 4412// CHECK1-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1413// CHECK1-NEXT:    store i32 3, ptr [[TMP28]], align 4414// CHECK1-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2415// CHECK1-NEXT:    store ptr [[TMP18]], ptr [[TMP29]], align 8416// CHECK1-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3417// CHECK1-NEXT:    store ptr [[TMP19]], ptr [[TMP30]], align 8418// CHECK1-NEXT:    [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4419// CHECK1-NEXT:    store ptr @.offload_sizes.3, ptr [[TMP31]], align 8420// CHECK1-NEXT:    [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5421// CHECK1-NEXT:    store ptr @.offload_maptypes.4, ptr [[TMP32]], align 8422// CHECK1-NEXT:    [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6423// CHECK1-NEXT:    store ptr null, ptr [[TMP33]], align 8424// CHECK1-NEXT:    [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7425// CHECK1-NEXT:    store ptr null, ptr [[TMP34]], align 8426// CHECK1-NEXT:    [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8427// CHECK1-NEXT:    store i64 [[TMP24]], ptr [[TMP35]], align 8428// CHECK1-NEXT:    [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9429// CHECK1-NEXT:    store i64 0, ptr [[TMP36]], align 8430// CHECK1-NEXT:    [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10431// CHECK1-NEXT:    store [3 x i32] [[TMP26]], ptr [[TMP37]], align 4432// CHECK1-NEXT:    [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11433// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP38]], align 4434// CHECK1-NEXT:    [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12435// CHECK1-NEXT:    store i32 [[TMP25]], ptr [[TMP39]], align 4436// CHECK1-NEXT:    [[TMP40:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[TMP20]], i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.region_id, ptr [[KERNEL_ARGS]])437// CHECK1-NEXT:    [[TMP41:%.*]] = icmp ne i32 [[TMP40]], 0438// CHECK1-NEXT:    br i1 [[TMP41]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]439// CHECK1:       omp_offload.failed:440// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104(i64 [[TMP4]], i64 [[TMP6]], i64 [[TMP8]]) #[[ATTR2]]441// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]442// CHECK1:       omp_offload.cont:443// CHECK1-NEXT:    [[TMP42:%.*]] = load i32, ptr [[N_ADDR]], align 4444// CHECK1-NEXT:    [[ADD7:%.*]] = add nsw i32 32, [[TMP42]]445// CHECK1-NEXT:    store i32 [[ADD7]], ptr [[DOTCAPTURE_EXPR_6]], align 4446// CHECK1-NEXT:    [[TMP43:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_6]], align 4447// CHECK1-NEXT:    store i32 [[TMP43]], ptr [[DOTCAPTURE_EXPR__CASTED8]], align 4448// CHECK1-NEXT:    [[TMP44:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED8]], align 8449// CHECK1-NEXT:    [[TMP45:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 0450// CHECK1-NEXT:    store i64 [[TMP44]], ptr [[TMP45]], align 8451// CHECK1-NEXT:    [[TMP46:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 0452// CHECK1-NEXT:    store i64 [[TMP44]], ptr [[TMP46]], align 8453// CHECK1-NEXT:    [[TMP47:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS11]], i64 0, i64 0454// CHECK1-NEXT:    store ptr null, ptr [[TMP47]], align 8455// CHECK1-NEXT:    [[TMP48:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 0456// CHECK1-NEXT:    [[TMP49:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 0457// CHECK1-NEXT:    [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT_ANON]], ptr [[AGG_CAPTURED]], i32 0, i32 0458// CHECK1-NEXT:    [[TMP51:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_6]], align 4459// CHECK1-NEXT:    store i32 [[TMP51]], ptr [[TMP50]], align 4460// CHECK1-NEXT:    [[TMP52:%.*]] = call ptr @__kmpc_omp_target_task_alloc(ptr @[[GLOB1]], i32 [[TMP0]], i32 1, i64 64, i64 4, ptr @.omp_task_entry., i64 -1)461// CHECK1-NEXT:    [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T_WITH_PRIVATES:%.*]], ptr [[TMP52]], i32 0, i32 0462// CHECK1-NEXT:    [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T:%.*]], ptr [[TMP53]], i32 0, i32 0463// CHECK1-NEXT:    [[TMP55:%.*]] = load ptr, ptr [[TMP54]], align 8464// CHECK1-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[TMP55]], ptr align 4 [[AGG_CAPTURED]], i64 4, i1 false)465// CHECK1-NEXT:    [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T_WITH_PRIVATES]], ptr [[TMP52]], i32 0, i32 1466// CHECK1-NEXT:    [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T:%.*]], ptr [[TMP56]], i32 0, i32 0467// CHECK1-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[TMP57]], ptr align 8 [[TMP48]], i64 8, i1 false)468// CHECK1-NEXT:    [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T]], ptr [[TMP56]], i32 0, i32 1469// CHECK1-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[TMP58]], ptr align 8 [[TMP49]], i64 8, i1 false)470// CHECK1-NEXT:    [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T]], ptr [[TMP56]], i32 0, i32 2471// CHECK1-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[TMP59]], ptr align 8 @.offload_sizes.5, i64 8, i1 false)472// CHECK1-NEXT:    [[TMP60:%.*]] = call i32 @__kmpc_omp_task(ptr @[[GLOB1]], i32 [[TMP0]], ptr [[TMP52]])473// CHECK1-NEXT:    [[TMP61:%.*]] = load i32, ptr [[N_ADDR]], align 4474// CHECK1-NEXT:    [[ADD12:%.*]] = add nsw i32 [[TMP61]], 1475// CHECK1-NEXT:    ret i32 [[ADD12]]476//477//478// CHECK1-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i479// CHECK1-SAME: (i32 noundef signext [[N:%.*]]) #[[ATTR0]] comdat {480// CHECK1-NEXT:  entry:481// CHECK1-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4482// CHECK1-NEXT:    [[A:%.*]] = alloca i32, align 4483// CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8484// CHECK1-NEXT:    [[B:%.*]] = alloca i16, align 2485// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i16, align 2486// CHECK1-NEXT:    [[A_CASTED:%.*]] = alloca i64, align 8487// CHECK1-NEXT:    [[B_CASTED:%.*]] = alloca i64, align 8488// CHECK1-NEXT:    [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 8489// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 8490// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 8491// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 8492// CHECK1-NEXT:    [[KERNEL_ARGS1:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8493// CHECK1-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4494// CHECK1-NEXT:    store i32 0, ptr [[A]], align 4495// CHECK1-NEXT:    [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0496// CHECK1-NEXT:    store i32 3, ptr [[TMP0]], align 4497// CHECK1-NEXT:    [[TMP1:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1498// CHECK1-NEXT:    store i32 0, ptr [[TMP1]], align 4499// CHECK1-NEXT:    [[TMP2:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2500// CHECK1-NEXT:    store ptr null, ptr [[TMP2]], align 8501// CHECK1-NEXT:    [[TMP3:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3502// CHECK1-NEXT:    store ptr null, ptr [[TMP3]], align 8503// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4504// CHECK1-NEXT:    store ptr null, ptr [[TMP4]], align 8505// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5506// CHECK1-NEXT:    store ptr null, ptr [[TMP5]], align 8507// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6508// CHECK1-NEXT:    store ptr null, ptr [[TMP6]], align 8509// CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7510// CHECK1-NEXT:    store ptr null, ptr [[TMP7]], align 8511// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8512// CHECK1-NEXT:    store i64 0, ptr [[TMP8]], align 8513// CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9514// CHECK1-NEXT:    store i64 0, ptr [[TMP9]], align 8515// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10516// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP10]], align 4517// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11518// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP11]], align 4519// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12520// CHECK1-NEXT:    store i32 20, ptr [[TMP12]], align 4521// CHECK1-NEXT:    [[TMP13:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l88.region_id, ptr [[KERNEL_ARGS]])522// CHECK1-NEXT:    [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 0523// CHECK1-NEXT:    br i1 [[TMP14]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]524// CHECK1:       omp_offload.failed:525// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l88() #[[ATTR2]]526// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]527// CHECK1:       omp_offload.cont:528// CHECK1-NEXT:    store i16 1, ptr [[B]], align 2529// CHECK1-NEXT:    [[TMP15:%.*]] = load i16, ptr [[B]], align 2530// CHECK1-NEXT:    store i16 [[TMP15]], ptr [[DOTCAPTURE_EXPR_]], align 2531// CHECK1-NEXT:    [[TMP16:%.*]] = load i32, ptr [[A]], align 4532// CHECK1-NEXT:    store i32 [[TMP16]], ptr [[A_CASTED]], align 4533// CHECK1-NEXT:    [[TMP17:%.*]] = load i64, ptr [[A_CASTED]], align 8534// CHECK1-NEXT:    [[TMP18:%.*]] = load i16, ptr [[B]], align 2535// CHECK1-NEXT:    store i16 [[TMP18]], ptr [[B_CASTED]], align 2536// CHECK1-NEXT:    [[TMP19:%.*]] = load i64, ptr [[B_CASTED]], align 8537// CHECK1-NEXT:    [[TMP20:%.*]] = load i16, ptr [[DOTCAPTURE_EXPR_]], align 2538// CHECK1-NEXT:    store i16 [[TMP20]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 2539// CHECK1-NEXT:    [[TMP21:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 8540// CHECK1-NEXT:    [[TMP22:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0541// CHECK1-NEXT:    store i64 [[TMP17]], ptr [[TMP22]], align 8542// CHECK1-NEXT:    [[TMP23:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0543// CHECK1-NEXT:    store i64 [[TMP17]], ptr [[TMP23]], align 8544// CHECK1-NEXT:    [[TMP24:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0545// CHECK1-NEXT:    store ptr null, ptr [[TMP24]], align 8546// CHECK1-NEXT:    [[TMP25:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1547// CHECK1-NEXT:    store i64 [[TMP19]], ptr [[TMP25]], align 8548// CHECK1-NEXT:    [[TMP26:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1549// CHECK1-NEXT:    store i64 [[TMP19]], ptr [[TMP26]], align 8550// CHECK1-NEXT:    [[TMP27:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1551// CHECK1-NEXT:    store ptr null, ptr [[TMP27]], align 8552// CHECK1-NEXT:    [[TMP28:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2553// CHECK1-NEXT:    store i64 [[TMP21]], ptr [[TMP28]], align 8554// CHECK1-NEXT:    [[TMP29:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2555// CHECK1-NEXT:    store i64 [[TMP21]], ptr [[TMP29]], align 8556// CHECK1-NEXT:    [[TMP30:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2557// CHECK1-NEXT:    store ptr null, ptr [[TMP30]], align 8558// CHECK1-NEXT:    [[TMP31:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0559// CHECK1-NEXT:    [[TMP32:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0560// CHECK1-NEXT:    [[TMP33:%.*]] = load i16, ptr [[DOTCAPTURE_EXPR_]], align 2561// CHECK1-NEXT:    [[TMP34:%.*]] = sext i16 [[TMP33]] to i32562// CHECK1-NEXT:    [[TMP35:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP34]], 0563// CHECK1-NEXT:    [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 0564// CHECK1-NEXT:    store i32 3, ptr [[TMP36]], align 4565// CHECK1-NEXT:    [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 1566// CHECK1-NEXT:    store i32 3, ptr [[TMP37]], align 4567// CHECK1-NEXT:    [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 2568// CHECK1-NEXT:    store ptr [[TMP31]], ptr [[TMP38]], align 8569// CHECK1-NEXT:    [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 3570// CHECK1-NEXT:    store ptr [[TMP32]], ptr [[TMP39]], align 8571// CHECK1-NEXT:    [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 4572// CHECK1-NEXT:    store ptr @.offload_sizes.7, ptr [[TMP40]], align 8573// CHECK1-NEXT:    [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 5574// CHECK1-NEXT:    store ptr @.offload_maptypes.8, ptr [[TMP41]], align 8575// CHECK1-NEXT:    [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 6576// CHECK1-NEXT:    store ptr null, ptr [[TMP42]], align 8577// CHECK1-NEXT:    [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 7578// CHECK1-NEXT:    store ptr null, ptr [[TMP43]], align 8579// CHECK1-NEXT:    [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 8580// CHECK1-NEXT:    store i64 0, ptr [[TMP44]], align 8581// CHECK1-NEXT:    [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 9582// CHECK1-NEXT:    store i64 0, ptr [[TMP45]], align 8583// CHECK1-NEXT:    [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 10584// CHECK1-NEXT:    store [3 x i32] [[TMP35]], ptr [[TMP46]], align 4585// CHECK1-NEXT:    [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 11586// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP47]], align 4587// CHECK1-NEXT:    [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 12588// CHECK1-NEXT:    store i32 1024, ptr [[TMP48]], align 4589// CHECK1-NEXT:    [[TMP49:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[TMP34]], i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93.region_id, ptr [[KERNEL_ARGS1]])590// CHECK1-NEXT:    [[TMP50:%.*]] = icmp ne i32 [[TMP49]], 0591// CHECK1-NEXT:    br i1 [[TMP50]], label [[OMP_OFFLOAD_FAILED2:%.*]], label [[OMP_OFFLOAD_CONT3:%.*]]592// CHECK1:       omp_offload.failed2:593// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93(i64 [[TMP17]], i64 [[TMP19]], i64 [[TMP21]]) #[[ATTR2]]594// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT3]]595// CHECK1:       omp_offload.cont3:596// CHECK1-NEXT:    [[TMP51:%.*]] = load i32, ptr [[A]], align 4597// CHECK1-NEXT:    ret i32 [[TMP51]]598//599//600// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121601// CHECK1-SAME: (ptr noundef [[THIS:%.*]], i64 noundef [[B:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1:[0-9]+]] {602// CHECK1-NEXT:  entry:603// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8604// CHECK1-NEXT:    [[B_ADDR:%.*]] = alloca i64, align 8605// CHECK1-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 8606// CHECK1-NEXT:    [[B_CASTED:%.*]] = alloca i64, align 8607// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8608// CHECK1-NEXT:    store i64 [[B]], ptr [[B_ADDR]], align 8609// CHECK1-NEXT:    store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 8610// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8611// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[B_ADDR]], align 4612// CHECK1-NEXT:    store i32 [[TMP1]], ptr [[B_CASTED]], align 4613// CHECK1-NEXT:    [[TMP2:%.*]] = load i64, ptr [[B_CASTED]], align 8614// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.omp_outlined, ptr [[TMP0]], i64 [[TMP2]])615// CHECK1-NEXT:    ret void616//617//618// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.omp_outlined619// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]], i64 noundef [[B:%.*]]) #[[ATTR1]] {620// CHECK1-NEXT:  entry:621// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8622// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8623// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8624// CHECK1-NEXT:    [[B_ADDR:%.*]] = alloca i64, align 8625// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8626// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8627// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8628// CHECK1-NEXT:    store i64 [[B]], ptr [[B_ADDR]], align 8629// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8630// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[B_ADDR]], align 4631// CHECK1-NEXT:    [[CONV:%.*]] = sitofp i32 [[TMP1]] to double632// CHECK1-NEXT:    [[ADD:%.*]] = fadd double [[CONV]], 1.500000e+00633// CHECK1-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 0634// CHECK1-NEXT:    store double [[ADD]], ptr [[A]], align 8635// CHECK1-NEXT:    ret void636//637//638// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126639// CHECK1-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1]] {640// CHECK1-NEXT:  entry:641// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8642// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8643// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8644// CHECK1-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 0645// CHECK1-NEXT:    store double 2.500000e+00, ptr [[A]], align 8646// CHECK1-NEXT:    ret void647//648//649// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104650// CHECK1-SAME: (i64 noundef [[N:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]], i64 noundef [[DOTCAPTURE_EXPR_1:%.*]]) #[[ATTR1]] {651// CHECK1-NEXT:  entry:652// CHECK1-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 8653// CHECK1-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 8654// CHECK1-NEXT:    [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca i64, align 8655// CHECK1-NEXT:    [[N_CASTED:%.*]] = alloca i64, align 8656// CHECK1-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])657// CHECK1-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 8658// CHECK1-NEXT:    store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 8659// CHECK1-NEXT:    store i64 [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 8660// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 4661// CHECK1-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 0)662// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4663// CHECK1-NEXT:    store i32 [[TMP2]], ptr [[N_CASTED]], align 4664// CHECK1-NEXT:    [[TMP3:%.*]] = load i64, ptr [[N_CASTED]], align 8665// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined, i64 [[TMP3]])666// CHECK1-NEXT:    ret void667//668//669// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined670// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[N:%.*]]) #[[ATTR1]] {671// CHECK1-NEXT:  entry:672// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8673// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8674// CHECK1-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 8675// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4676// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4677// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4678// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4679// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4680// CHECK1-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4681// CHECK1-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4682// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4683// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4684// CHECK1-NEXT:    [[I3:%.*]] = alloca i32, align 4685// CHECK1-NEXT:    [[N_CASTED:%.*]] = alloca i64, align 8686// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8687// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8688// CHECK1-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 8689// CHECK1-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4690// CHECK1-NEXT:    store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 4691// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4692// CHECK1-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP1]], 0693// CHECK1-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1694// CHECK1-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1695// CHECK1-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4696// CHECK1-NEXT:    store i32 0, ptr [[I]], align 4697// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4698// CHECK1-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP2]]699// CHECK1-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]700// CHECK1:       omp.precond.then:701// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 4702// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4703// CHECK1-NEXT:    store i32 [[TMP3]], ptr [[DOTOMP_COMB_UB]], align 4704// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4705// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4706// CHECK1-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8707// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 4708// CHECK1-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP5]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)709// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4710// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4711// CHECK1-NEXT:    [[CMP4:%.*]] = icmp sgt i32 [[TMP6]], [[TMP7]]712// CHECK1-NEXT:    br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]713// CHECK1:       cond.true:714// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4715// CHECK1-NEXT:    br label [[COND_END:%.*]]716// CHECK1:       cond.false:717// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4718// CHECK1-NEXT:    br label [[COND_END]]719// CHECK1:       cond.end:720// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP8]], [[COND_TRUE]] ], [ [[TMP9]], [[COND_FALSE]] ]721// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4722// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4723// CHECK1-NEXT:    store i32 [[TMP10]], ptr [[DOTOMP_IV]], align 4724// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]725// CHECK1:       omp.inner.for.cond:726// CHECK1-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11:![0-9]+]]727// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4, !llvm.access.group [[ACC_GRP11]]728// CHECK1-NEXT:    [[CMP5:%.*]] = icmp sle i32 [[TMP11]], [[TMP12]]729// CHECK1-NEXT:    br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]730// CHECK1:       omp.inner.for.body:731// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4, !llvm.access.group [[ACC_GRP11]]732// CHECK1-NEXT:    [[TMP14:%.*]] = zext i32 [[TMP13]] to i64733// CHECK1-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4, !llvm.access.group [[ACC_GRP11]]734// CHECK1-NEXT:    [[TMP16:%.*]] = zext i32 [[TMP15]] to i64735// CHECK1-NEXT:    [[TMP17:%.*]] = load i32, ptr [[N_ADDR]], align 4, !llvm.access.group [[ACC_GRP11]]736// CHECK1-NEXT:    store i32 [[TMP17]], ptr [[N_CASTED]], align 4, !llvm.access.group [[ACC_GRP11]]737// CHECK1-NEXT:    [[TMP18:%.*]] = load i64, ptr [[N_CASTED]], align 8, !llvm.access.group [[ACC_GRP11]]738// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined.omp_outlined, i64 [[TMP14]], i64 [[TMP16]], i64 [[TMP18]]), !llvm.access.group [[ACC_GRP11]]739// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]740// CHECK1:       omp.inner.for.inc:741// CHECK1-NEXT:    [[TMP19:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]742// CHECK1-NEXT:    [[TMP20:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4, !llvm.access.group [[ACC_GRP11]]743// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP19]], [[TMP20]]744// CHECK1-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]745// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP12:![0-9]+]]746// CHECK1:       omp.inner.for.end:747// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]748// CHECK1:       omp.loop.exit:749// CHECK1-NEXT:    [[TMP21:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8750// CHECK1-NEXT:    [[TMP22:%.*]] = load i32, ptr [[TMP21]], align 4751// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP22]])752// CHECK1-NEXT:    [[TMP23:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 4753// CHECK1-NEXT:    [[TMP24:%.*]] = icmp ne i32 [[TMP23]], 0754// CHECK1-NEXT:    br i1 [[TMP24]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]755// CHECK1:       .omp.final.then:756// CHECK1-NEXT:    [[TMP25:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4757// CHECK1-NEXT:    [[SUB6:%.*]] = sub nsw i32 [[TMP25]], 0758// CHECK1-NEXT:    [[DIV7:%.*]] = sdiv i32 [[SUB6]], 1759// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[DIV7]], 1760// CHECK1-NEXT:    [[ADD8:%.*]] = add nsw i32 0, [[MUL]]761// CHECK1-NEXT:    store i32 [[ADD8]], ptr [[I3]], align 4762// CHECK1-NEXT:    br label [[DOTOMP_FINAL_DONE]]763// CHECK1:       .omp.final.done:764// CHECK1-NEXT:    br label [[OMP_PRECOND_END]]765// CHECK1:       omp.precond.end:766// CHECK1-NEXT:    ret void767//768//769// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined.omp_outlined770// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]], i64 noundef [[N:%.*]]) #[[ATTR1]] {771// CHECK1-NEXT:  entry:772// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8773// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8774// CHECK1-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8775// CHECK1-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8776// CHECK1-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 8777// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4778// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4779// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4780// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4781// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4782// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4783// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4784// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4785// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4786// CHECK1-NEXT:    [[I4:%.*]] = alloca i32, align 4787// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8788// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8789// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8790// CHECK1-NEXT:    store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8791// CHECK1-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 8792// CHECK1-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4793// CHECK1-NEXT:    store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 4794// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4795// CHECK1-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP1]], 0796// CHECK1-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1797// CHECK1-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1798// CHECK1-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4799// CHECK1-NEXT:    store i32 0, ptr [[I]], align 4800// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4801// CHECK1-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP2]]802// CHECK1-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]803// CHECK1:       omp.precond.then:804// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4805// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4806// CHECK1-NEXT:    store i32 [[TMP3]], ptr [[DOTOMP_UB]], align 4807// CHECK1-NEXT:    [[TMP4:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8808// CHECK1-NEXT:    [[CONV:%.*]] = trunc i64 [[TMP4]] to i32809// CHECK1-NEXT:    [[TMP5:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8810// CHECK1-NEXT:    [[CONV3:%.*]] = trunc i64 [[TMP5]] to i32811// CHECK1-NEXT:    store i32 [[CONV]], ptr [[DOTOMP_LB]], align 4812// CHECK1-NEXT:    store i32 [[CONV3]], ptr [[DOTOMP_UB]], align 4813// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4814// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4815// CHECK1-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8816// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 4817// CHECK1-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB3:[0-9]+]], i32 [[TMP7]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)818// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4819// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4820// CHECK1-NEXT:    [[CMP5:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]821// CHECK1-NEXT:    br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]822// CHECK1:       cond.true:823// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4824// CHECK1-NEXT:    br label [[COND_END:%.*]]825// CHECK1:       cond.false:826// CHECK1-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4827// CHECK1-NEXT:    br label [[COND_END]]828// CHECK1:       cond.end:829// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]830// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4831// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4832// CHECK1-NEXT:    store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 4833// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]834// CHECK1:       omp.inner.for.cond:835// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP15:![0-9]+]]836// CHECK1-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP15]]837// CHECK1-NEXT:    [[CMP6:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]838// CHECK1-NEXT:    br i1 [[CMP6]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]839// CHECK1:       omp.inner.for.body:840// CHECK1-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP15]]841// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP15]], 1842// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]843// CHECK1-NEXT:    store i32 [[ADD]], ptr [[I4]], align 4, !llvm.access.group [[ACC_GRP15]]844// CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]845// CHECK1:       omp.body.continue:846// CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]847// CHECK1:       omp.inner.for.inc:848// CHECK1-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP15]]849// CHECK1-NEXT:    [[ADD7:%.*]] = add nsw i32 [[TMP16]], 1850// CHECK1-NEXT:    store i32 [[ADD7]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP15]]851// CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP16:![0-9]+]]852// CHECK1:       omp.inner.for.end:853// CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]854// CHECK1:       omp.loop.exit:855// CHECK1-NEXT:    [[TMP17:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8856// CHECK1-NEXT:    [[TMP18:%.*]] = load i32, ptr [[TMP17]], align 4857// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB3]], i32 [[TMP18]])858// CHECK1-NEXT:    [[TMP19:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 4859// CHECK1-NEXT:    [[TMP20:%.*]] = icmp ne i32 [[TMP19]], 0860// CHECK1-NEXT:    br i1 [[TMP20]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]861// CHECK1:       .omp.final.then:862// CHECK1-NEXT:    [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4863// CHECK1-NEXT:    [[SUB8:%.*]] = sub nsw i32 [[TMP21]], 0864// CHECK1-NEXT:    [[DIV9:%.*]] = sdiv i32 [[SUB8]], 1865// CHECK1-NEXT:    [[MUL10:%.*]] = mul nsw i32 [[DIV9]], 1866// CHECK1-NEXT:    [[ADD11:%.*]] = add nsw i32 0, [[MUL10]]867// CHECK1-NEXT:    store i32 [[ADD11]], ptr [[I4]], align 4868// CHECK1-NEXT:    br label [[DOTOMP_FINAL_DONE]]869// CHECK1:       .omp.final.done:870// CHECK1-NEXT:    br label [[OMP_PRECOND_END]]871// CHECK1:       omp.precond.end:872// CHECK1-NEXT:    ret void873//874//875// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108876// CHECK1-SAME: (i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1]] {877// CHECK1-NEXT:  entry:878// CHECK1-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 8879// CHECK1-NEXT:    store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 8880// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108.omp_outlined)881// CHECK1-NEXT:    ret void882//883//884// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108.omp_outlined885// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {886// CHECK1-NEXT:  entry:887// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8888// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8889// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8890// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8891// CHECK1-NEXT:    ret void892//893//894// CHECK1-LABEL: define {{[^@]+}}@.omp_task_privates_map.895// CHECK1-SAME: (ptr noalias noundef [[TMP0:%.*]], ptr noalias noundef [[TMP1:%.*]], ptr noalias noundef [[TMP2:%.*]], ptr noalias noundef [[TMP3:%.*]]) #[[ATTR3:[0-9]+]] {896// CHECK1-NEXT:  entry:897// CHECK1-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 8898// CHECK1-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 8899// CHECK1-NEXT:    [[DOTADDR2:%.*]] = alloca ptr, align 8900// CHECK1-NEXT:    [[DOTADDR3:%.*]] = alloca ptr, align 8901// CHECK1-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 8902// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 8903// CHECK1-NEXT:    store ptr [[TMP2]], ptr [[DOTADDR2]], align 8904// CHECK1-NEXT:    store ptr [[TMP3]], ptr [[DOTADDR3]], align 8905// CHECK1-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTADDR]], align 8906// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T:%.*]], ptr [[TMP4]], i32 0, i32 0907// CHECK1-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[DOTADDR1]], align 8908// CHECK1-NEXT:    store ptr [[TMP5]], ptr [[TMP6]], align 8909// CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T]], ptr [[TMP4]], i32 0, i32 1910// CHECK1-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[DOTADDR2]], align 8911// CHECK1-NEXT:    store ptr [[TMP7]], ptr [[TMP8]], align 8912// CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T]], ptr [[TMP4]], i32 0, i32 2913// CHECK1-NEXT:    [[TMP10:%.*]] = load ptr, ptr [[DOTADDR3]], align 8914// CHECK1-NEXT:    store ptr [[TMP9]], ptr [[TMP10]], align 8915// CHECK1-NEXT:    ret void916//917//918// CHECK1-LABEL: define {{[^@]+}}@.omp_task_entry.919// CHECK1-SAME: (i32 noundef signext [[TMP0:%.*]], ptr noalias noundef [[TMP1:%.*]]) #[[ATTR4:[0-9]+]] {920// CHECK1-NEXT:  entry:921// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR_I:%.*]] = alloca i32, align 4922// CHECK1-NEXT:    [[DOTPART_ID__ADDR_I:%.*]] = alloca ptr, align 8923// CHECK1-NEXT:    [[DOTPRIVATES__ADDR_I:%.*]] = alloca ptr, align 8924// CHECK1-NEXT:    [[DOTCOPY_FN__ADDR_I:%.*]] = alloca ptr, align 8925// CHECK1-NEXT:    [[DOTTASK_T__ADDR_I:%.*]] = alloca ptr, align 8926// CHECK1-NEXT:    [[__CONTEXT_ADDR_I:%.*]] = alloca ptr, align 8927// CHECK1-NEXT:    [[DOTFIRSTPRIV_PTR_ADDR_I:%.*]] = alloca ptr, align 8928// CHECK1-NEXT:    [[DOTFIRSTPRIV_PTR_ADDR1_I:%.*]] = alloca ptr, align 8929// CHECK1-NEXT:    [[DOTFIRSTPRIV_PTR_ADDR2_I:%.*]] = alloca ptr, align 8930// CHECK1-NEXT:    [[KERNEL_ARGS_I:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8931// CHECK1-NEXT:    [[DOTCAPTURE_EXPR__CASTED_I:%.*]] = alloca i64, align 8932// CHECK1-NEXT:    [[DOTADDR:%.*]] = alloca i32, align 4933// CHECK1-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 8934// CHECK1-NEXT:    store i32 [[TMP0]], ptr [[DOTADDR]], align 4935// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 8936// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTADDR]], align 4937// CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 8938// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T_WITH_PRIVATES:%.*]], ptr [[TMP3]], i32 0, i32 0939// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T:%.*]], ptr [[TMP4]], i32 0, i32 2940// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T]], ptr [[TMP4]], i32 0, i32 0941// CHECK1-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 8942// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T_WITH_PRIVATES]], ptr [[TMP3]], i32 0, i32 1943// CHECK1-NEXT:    call void @llvm.experimental.noalias.scope.decl(metadata [[META18:![0-9]+]])944// CHECK1-NEXT:    call void @llvm.experimental.noalias.scope.decl(metadata [[META21:![0-9]+]])945// CHECK1-NEXT:    call void @llvm.experimental.noalias.scope.decl(metadata [[META23:![0-9]+]])946// CHECK1-NEXT:    call void @llvm.experimental.noalias.scope.decl(metadata [[META25:![0-9]+]])947// CHECK1-NEXT:    store i32 [[TMP2]], ptr [[DOTGLOBAL_TID__ADDR_I]], align 4, !noalias [[META27:![0-9]+]]948// CHECK1-NEXT:    store ptr [[TMP5]], ptr [[DOTPART_ID__ADDR_I]], align 8, !noalias [[META27]]949// CHECK1-NEXT:    store ptr [[TMP8]], ptr [[DOTPRIVATES__ADDR_I]], align 8, !noalias [[META27]]950// CHECK1-NEXT:    store ptr @.omp_task_privates_map., ptr [[DOTCOPY_FN__ADDR_I]], align 8, !noalias [[META27]]951// CHECK1-NEXT:    store ptr [[TMP3]], ptr [[DOTTASK_T__ADDR_I]], align 8, !noalias [[META27]]952// CHECK1-NEXT:    store ptr [[TMP7]], ptr [[__CONTEXT_ADDR_I]], align 8, !noalias [[META27]]953// CHECK1-NEXT:    [[TMP9:%.*]] = load ptr, ptr [[__CONTEXT_ADDR_I]], align 8, !noalias [[META27]]954// CHECK1-NEXT:    [[TMP10:%.*]] = load ptr, ptr [[DOTCOPY_FN__ADDR_I]], align 8, !noalias [[META27]]955// CHECK1-NEXT:    [[TMP11:%.*]] = load ptr, ptr [[DOTPRIVATES__ADDR_I]], align 8, !noalias [[META27]]956// CHECK1-NEXT:    call void [[TMP10]](ptr [[TMP11]], ptr [[DOTFIRSTPRIV_PTR_ADDR_I]], ptr [[DOTFIRSTPRIV_PTR_ADDR1_I]], ptr [[DOTFIRSTPRIV_PTR_ADDR2_I]]) #[[ATTR2]]957// CHECK1-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[DOTFIRSTPRIV_PTR_ADDR_I]], align 8, !noalias [[META27]]958// CHECK1-NEXT:    [[TMP13:%.*]] = load ptr, ptr [[DOTFIRSTPRIV_PTR_ADDR1_I]], align 8, !noalias [[META27]]959// CHECK1-NEXT:    [[TMP14:%.*]] = load ptr, ptr [[DOTFIRSTPRIV_PTR_ADDR2_I]], align 8, !noalias [[META27]]960// CHECK1-NEXT:    [[TMP15:%.*]] = load i32, ptr [[TMP9]], align 4961// CHECK1-NEXT:    store i32 3, ptr [[KERNEL_ARGS_I]], align 4, !noalias [[META27]]962// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 1963// CHECK1-NEXT:    store i32 1, ptr [[TMP16]], align 4, !noalias [[META27]]964// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 2965// CHECK1-NEXT:    store ptr [[TMP12]], ptr [[TMP17]], align 8, !noalias [[META27]]966// CHECK1-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 3967// CHECK1-NEXT:    store ptr [[TMP13]], ptr [[TMP18]], align 8, !noalias [[META27]]968// CHECK1-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 4969// CHECK1-NEXT:    store ptr [[TMP14]], ptr [[TMP19]], align 8, !noalias [[META27]]970// CHECK1-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 5971// CHECK1-NEXT:    store ptr @.offload_maptypes.6, ptr [[TMP20]], align 8, !noalias [[META27]]972// CHECK1-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 6973// CHECK1-NEXT:    store ptr null, ptr [[TMP21]], align 8, !noalias [[META27]]974// CHECK1-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 7975// CHECK1-NEXT:    store ptr null, ptr [[TMP22]], align 8, !noalias [[META27]]976// CHECK1-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 8977// CHECK1-NEXT:    store i64 0, ptr [[TMP23]], align 8, !noalias [[META27]]978// CHECK1-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 9979// CHECK1-NEXT:    store i64 1, ptr [[TMP24]], align 8, !noalias [[META27]]980// CHECK1-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 10981// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP25]], align 4, !noalias [[META27]]982// CHECK1-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 11983// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP26]], align 4, !noalias [[META27]]984// CHECK1-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 12985// CHECK1-NEXT:    store i32 [[TMP15]], ptr [[TMP27]], align 4, !noalias [[META27]]986// CHECK1-NEXT:    [[TMP28:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108.region_id, ptr [[KERNEL_ARGS_I]])987// CHECK1-NEXT:    [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 0988// CHECK1-NEXT:    br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED_I:%.*]], label [[DOTOMP_OUTLINED__EXIT:%.*]]989// CHECK1:       omp_offload.failed.i:990// CHECK1-NEXT:    [[TMP30:%.*]] = load i32, ptr [[TMP9]], align 4991// CHECK1-NEXT:    store i32 [[TMP30]], ptr [[DOTCAPTURE_EXPR__CASTED_I]], align 4, !noalias [[META27]]992// CHECK1-NEXT:    [[TMP31:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED_I]], align 8, !noalias [[META27]]993// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108(i64 [[TMP31]]) #[[ATTR2]]994// CHECK1-NEXT:    br label [[DOTOMP_OUTLINED__EXIT]]995// CHECK1:       .omp_outlined..exit:996// CHECK1-NEXT:    ret i32 0997//998//999// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l881000// CHECK1-SAME: () #[[ATTR1]] {1001// CHECK1-NEXT:  entry:1002// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l88.omp_outlined)1003// CHECK1-NEXT:    ret void1004//1005//1006// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l88.omp_outlined1007// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {1008// CHECK1-NEXT:  entry:1009// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81010// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81011// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81012// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81013// CHECK1-NEXT:    ret void1014//1015//1016// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l931017// CHECK1-SAME: (i64 noundef [[A:%.*]], i64 noundef [[B:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1]] {1018// CHECK1-NEXT:  entry:1019// CHECK1-NEXT:    [[A_ADDR:%.*]] = alloca i64, align 81020// CHECK1-NEXT:    [[B_ADDR:%.*]] = alloca i64, align 81021// CHECK1-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 81022// CHECK1-NEXT:    [[A_CASTED:%.*]] = alloca i64, align 81023// CHECK1-NEXT:    [[B_CASTED:%.*]] = alloca i64, align 81024// CHECK1-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1025// CHECK1-NEXT:    store i64 [[A]], ptr [[A_ADDR]], align 81026// CHECK1-NEXT:    store i64 [[B]], ptr [[B_ADDR]], align 81027// CHECK1-NEXT:    store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 81028// CHECK1-NEXT:    [[TMP1:%.*]] = load i16, ptr [[DOTCAPTURE_EXPR__ADDR]], align 21029// CHECK1-NEXT:    [[TMP2:%.*]] = sext i16 [[TMP1]] to i321030// CHECK1-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 0)1031// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A_ADDR]], align 41032// CHECK1-NEXT:    store i32 [[TMP3]], ptr [[A_CASTED]], align 41033// CHECK1-NEXT:    [[TMP4:%.*]] = load i64, ptr [[A_CASTED]], align 81034// CHECK1-NEXT:    [[TMP5:%.*]] = load i16, ptr [[B_ADDR]], align 21035// CHECK1-NEXT:    store i16 [[TMP5]], ptr [[B_CASTED]], align 21036// CHECK1-NEXT:    [[TMP6:%.*]] = load i64, ptr [[B_CASTED]], align 81037// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93.omp_outlined, i64 [[TMP4]], i64 [[TMP6]])1038// CHECK1-NEXT:    ret void1039//1040//1041// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93.omp_outlined1042// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[A:%.*]], i64 noundef [[B:%.*]]) #[[ATTR1]] {1043// CHECK1-NEXT:  entry:1044// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81045// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81046// CHECK1-NEXT:    [[A_ADDR:%.*]] = alloca i64, align 81047// CHECK1-NEXT:    [[B_ADDR:%.*]] = alloca i64, align 81048// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81049// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81050// CHECK1-NEXT:    store i64 [[A]], ptr [[A_ADDR]], align 81051// CHECK1-NEXT:    store i64 [[B]], ptr [[B_ADDR]], align 81052// CHECK1-NEXT:    [[TMP0:%.*]] = load i16, ptr [[B_ADDR]], align 21053// CHECK1-NEXT:    [[CONV:%.*]] = sext i16 [[TMP0]] to i321054// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 41055// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CONV]]1056// CHECK1-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 41057// CHECK1-NEXT:    ret void1058//1059//1060// CHECK3-LABEL: define {{[^@]+}}@_Z3bari1061// CHECK3-SAME: (i32 noundef [[N:%.*]]) #[[ATTR0:[0-9]+]] {1062// CHECK3-NEXT:  entry:1063// CHECK3-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41064// CHECK3-NEXT:    [[A:%.*]] = alloca i32, align 41065// CHECK3-NEXT:    [[S:%.*]] = alloca [[STRUCT_S1:%.*]], align 41066// CHECK3-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41067// CHECK3-NEXT:    store i32 0, ptr [[A]], align 41068// CHECK3-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 41069// CHECK3-NEXT:    [[CALL:%.*]] = call noundef i32 @_ZN2S12r1Ei(ptr noundef nonnull align 4 dereferenceable(8) [[S]], i32 noundef [[TMP0]])1070// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A]], align 41071// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CALL]]1072// CHECK3-NEXT:    store i32 [[ADD]], ptr [[A]], align 41073// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 41074// CHECK3-NEXT:    [[CALL1:%.*]] = call noundef i32 @_ZL7fstatici(i32 noundef [[TMP2]])1075// CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A]], align 41076// CHECK3-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP3]], [[CALL1]]1077// CHECK3-NEXT:    store i32 [[ADD2]], ptr [[A]], align 41078// CHECK3-NEXT:    [[TMP4:%.*]] = load i32, ptr [[N_ADDR]], align 41079// CHECK3-NEXT:    [[CALL3:%.*]] = call noundef i32 @_Z9ftemplateIiET_i(i32 noundef [[TMP4]])1080// CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[A]], align 41081// CHECK3-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP5]], [[CALL3]]1082// CHECK3-NEXT:    store i32 [[ADD4]], ptr [[A]], align 41083// CHECK3-NEXT:    [[TMP6:%.*]] = load i32, ptr [[A]], align 41084// CHECK3-NEXT:    ret i32 [[TMP6]]1085//1086//1087// CHECK3-LABEL: define {{[^@]+}}@_ZN2S12r1Ei1088// CHECK3-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]], i32 noundef [[N:%.*]]) #[[ATTR0]] comdat align 2 {1089// CHECK3-NEXT:  entry:1090// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 41091// CHECK3-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41092// CHECK3-NEXT:    [[B:%.*]] = alloca i32, align 41093// CHECK3-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41094// CHECK3-NEXT:    [[B_CASTED:%.*]] = alloca i32, align 41095// CHECK3-NEXT:    [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i32, align 41096// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 41097// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 41098// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 41099// CHECK3-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81100// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS3:%.*]] = alloca [1 x ptr], align 41101// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS4:%.*]] = alloca [1 x ptr], align 41102// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS5:%.*]] = alloca [1 x ptr], align 41103// CHECK3-NEXT:    [[KERNEL_ARGS6:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 81104// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 41105// CHECK3-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41106// CHECK3-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41107// CHECK3-NEXT:    store i32 1, ptr [[B]], align 41108// CHECK3-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 41109// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[B]], align 41110// CHECK3-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP0]], [[TMP1]]1111// CHECK3-NEXT:    store i32 [[SUB]], ptr [[DOTCAPTURE_EXPR_]], align 41112// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[B]], align 41113// CHECK3-NEXT:    store i32 [[TMP2]], ptr [[B_CASTED]], align 41114// CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[B_CASTED]], align 41115// CHECK3-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41116// CHECK3-NEXT:    store i32 [[TMP4]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 41117// CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED]], align 41118// CHECK3-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[THIS1]], i32 0, i32 01119// CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01120// CHECK3-NEXT:    store ptr [[THIS1]], ptr [[TMP6]], align 41121// CHECK3-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01122// CHECK3-NEXT:    store ptr [[A]], ptr [[TMP7]], align 41123// CHECK3-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01124// CHECK3-NEXT:    store ptr null, ptr [[TMP8]], align 41125// CHECK3-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11126// CHECK3-NEXT:    store i32 [[TMP3]], ptr [[TMP9]], align 41127// CHECK3-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11128// CHECK3-NEXT:    store i32 [[TMP3]], ptr [[TMP10]], align 41129// CHECK3-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 11130// CHECK3-NEXT:    store ptr null, ptr [[TMP11]], align 41131// CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21132// CHECK3-NEXT:    store i32 [[TMP5]], ptr [[TMP12]], align 41133// CHECK3-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21134// CHECK3-NEXT:    store i32 [[TMP5]], ptr [[TMP13]], align 41135// CHECK3-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 21136// CHECK3-NEXT:    store ptr null, ptr [[TMP14]], align 41137// CHECK3-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01138// CHECK3-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01139// CHECK3-NEXT:    [[TMP17:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41140// CHECK3-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01141// CHECK3-NEXT:    store i32 3, ptr [[TMP18]], align 41142// CHECK3-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11143// CHECK3-NEXT:    store i32 3, ptr [[TMP19]], align 41144// CHECK3-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21145// CHECK3-NEXT:    store ptr [[TMP15]], ptr [[TMP20]], align 41146// CHECK3-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31147// CHECK3-NEXT:    store ptr [[TMP16]], ptr [[TMP21]], align 41148// CHECK3-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41149// CHECK3-NEXT:    store ptr @.offload_sizes, ptr [[TMP22]], align 41150// CHECK3-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51151// CHECK3-NEXT:    store ptr @.offload_maptypes, ptr [[TMP23]], align 41152// CHECK3-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61153// CHECK3-NEXT:    store ptr null, ptr [[TMP24]], align 41154// CHECK3-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71155// CHECK3-NEXT:    store ptr null, ptr [[TMP25]], align 41156// CHECK3-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81157// CHECK3-NEXT:    store i64 0, ptr [[TMP26]], align 81158// CHECK3-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91159// CHECK3-NEXT:    store i64 0, ptr [[TMP27]], align 81160// CHECK3-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101161// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP28]], align 41162// CHECK3-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111163// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP29]], align 41164// CHECK3-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121165// CHECK3-NEXT:    store i32 [[TMP17]], ptr [[TMP30]], align 41166// CHECK3-NEXT:    [[TMP31:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.region_id, ptr [[KERNEL_ARGS]])1167// CHECK3-NEXT:    [[TMP32:%.*]] = icmp ne i32 [[TMP31]], 01168// CHECK3-NEXT:    br i1 [[TMP32]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1169// CHECK3:       omp_offload.failed:1170// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121(ptr [[THIS1]], i32 [[TMP3]], i32 [[TMP5]]) #[[ATTR2:[0-9]+]]1171// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT]]1172// CHECK3:       omp_offload.cont:1173// CHECK3-NEXT:    [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[THIS1]], i32 0, i32 01174// CHECK3-NEXT:    [[TMP33:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 01175// CHECK3-NEXT:    store ptr [[THIS1]], ptr [[TMP33]], align 41176// CHECK3-NEXT:    [[TMP34:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 01177// CHECK3-NEXT:    store ptr [[A2]], ptr [[TMP34]], align 41178// CHECK3-NEXT:    [[TMP35:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS5]], i32 0, i32 01179// CHECK3-NEXT:    store ptr null, ptr [[TMP35]], align 41180// CHECK3-NEXT:    [[TMP36:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 01181// CHECK3-NEXT:    [[TMP37:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 01182// CHECK3-NEXT:    [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 01183// CHECK3-NEXT:    store i32 3, ptr [[TMP38]], align 41184// CHECK3-NEXT:    [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 11185// CHECK3-NEXT:    store i32 1, ptr [[TMP39]], align 41186// CHECK3-NEXT:    [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 21187// CHECK3-NEXT:    store ptr [[TMP36]], ptr [[TMP40]], align 41188// CHECK3-NEXT:    [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 31189// CHECK3-NEXT:    store ptr [[TMP37]], ptr [[TMP41]], align 41190// CHECK3-NEXT:    [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 41191// CHECK3-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP42]], align 41192// CHECK3-NEXT:    [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 51193// CHECK3-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP43]], align 41194// CHECK3-NEXT:    [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 61195// CHECK3-NEXT:    store ptr null, ptr [[TMP44]], align 41196// CHECK3-NEXT:    [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 71197// CHECK3-NEXT:    store ptr null, ptr [[TMP45]], align 41198// CHECK3-NEXT:    [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 81199// CHECK3-NEXT:    store i64 0, ptr [[TMP46]], align 81200// CHECK3-NEXT:    [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 91201// CHECK3-NEXT:    store i64 0, ptr [[TMP47]], align 81202// CHECK3-NEXT:    [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 101203// CHECK3-NEXT:    store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP48]], align 41204// CHECK3-NEXT:    [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 111205// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP49]], align 41206// CHECK3-NEXT:    [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 121207// CHECK3-NEXT:    store i32 1024, ptr [[TMP50]], align 41208// CHECK3-NEXT:    [[TMP51:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126.region_id, ptr [[KERNEL_ARGS6]])1209// CHECK3-NEXT:    [[TMP52:%.*]] = icmp ne i32 [[TMP51]], 01210// CHECK3-NEXT:    br i1 [[TMP52]], label [[OMP_OFFLOAD_FAILED7:%.*]], label [[OMP_OFFLOAD_CONT8:%.*]]1211// CHECK3:       omp_offload.failed7:1212// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126(ptr [[THIS1]]) #[[ATTR2]]1213// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT8]]1214// CHECK3:       omp_offload.cont8:1215// CHECK3-NEXT:    [[A9:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[THIS1]], i32 0, i32 01216// CHECK3-NEXT:    [[TMP53:%.*]] = load double, ptr [[A9]], align 41217// CHECK3-NEXT:    [[CONV:%.*]] = fptosi double [[TMP53]] to i321218// CHECK3-NEXT:    ret i32 [[CONV]]1219//1220//1221// CHECK3-LABEL: define {{[^@]+}}@_ZL7fstatici1222// CHECK3-SAME: (i32 noundef [[N:%.*]]) #[[ATTR0]] {1223// CHECK3-NEXT:  entry:1224// CHECK3-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41225// CHECK3-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41226// CHECK3-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41227// CHECK3-NEXT:    [[N_CASTED:%.*]] = alloca i32, align 41228// CHECK3-NEXT:    [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i32, align 41229// CHECK3-NEXT:    [[DOTCAPTURE_EXPR__CASTED2:%.*]] = alloca i32, align 41230// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 41231// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 41232// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 41233// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 41234// CHECK3-NEXT:    [[DOTCAPTURE_EXPR_3:%.*]] = alloca i32, align 41235// CHECK3-NEXT:    [[DOTCAPTURE_EXPR_4:%.*]] = alloca i32, align 41236// CHECK3-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81237// CHECK3-NEXT:    [[DOTCAPTURE_EXPR_6:%.*]] = alloca i32, align 41238// CHECK3-NEXT:    [[DOTCAPTURE_EXPR__CASTED8:%.*]] = alloca i32, align 41239// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS9:%.*]] = alloca [1 x ptr], align 41240// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS10:%.*]] = alloca [1 x ptr], align 41241// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS11:%.*]] = alloca [1 x ptr], align 41242// CHECK3-NEXT:    [[AGG_CAPTURED:%.*]] = alloca [[STRUCT_ANON:%.*]], align 41243// CHECK3-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1244// CHECK3-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41245// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41246// CHECK3-NEXT:    store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 41247// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 41248// CHECK3-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP2]], 321249// CHECK3-NEXT:    store i32 [[MUL]], ptr [[DOTCAPTURE_EXPR_1]], align 41250// CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[N_ADDR]], align 41251// CHECK3-NEXT:    store i32 [[TMP3]], ptr [[N_CASTED]], align 41252// CHECK3-NEXT:    [[TMP4:%.*]] = load i32, ptr [[N_CASTED]], align 41253// CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41254// CHECK3-NEXT:    store i32 [[TMP5]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 41255// CHECK3-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED]], align 41256// CHECK3-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41257// CHECK3-NEXT:    store i32 [[TMP7]], ptr [[DOTCAPTURE_EXPR__CASTED2]], align 41258// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED2]], align 41259// CHECK3-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01260// CHECK3-NEXT:    store i32 [[TMP4]], ptr [[TMP9]], align 41261// CHECK3-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01262// CHECK3-NEXT:    store i32 [[TMP4]], ptr [[TMP10]], align 41263// CHECK3-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01264// CHECK3-NEXT:    store ptr null, ptr [[TMP11]], align 41265// CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11266// CHECK3-NEXT:    store i32 [[TMP6]], ptr [[TMP12]], align 41267// CHECK3-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11268// CHECK3-NEXT:    store i32 [[TMP6]], ptr [[TMP13]], align 41269// CHECK3-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 11270// CHECK3-NEXT:    store ptr null, ptr [[TMP14]], align 41271// CHECK3-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21272// CHECK3-NEXT:    store i32 [[TMP8]], ptr [[TMP15]], align 41273// CHECK3-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21274// CHECK3-NEXT:    store i32 [[TMP8]], ptr [[TMP16]], align 41275// CHECK3-NEXT:    [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 21276// CHECK3-NEXT:    store ptr null, ptr [[TMP17]], align 41277// CHECK3-NEXT:    [[TMP18:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01278// CHECK3-NEXT:    [[TMP19:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01279// CHECK3-NEXT:    [[TMP20:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41280// CHECK3-NEXT:    [[TMP21:%.*]] = load i32, ptr [[N_ADDR]], align 41281// CHECK3-NEXT:    store i32 [[TMP21]], ptr [[DOTCAPTURE_EXPR_3]], align 41282// CHECK3-NEXT:    [[TMP22:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_3]], align 41283// CHECK3-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP22]], 01284// CHECK3-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 11285// CHECK3-NEXT:    [[SUB5:%.*]] = sub nsw i32 [[DIV]], 11286// CHECK3-NEXT:    store i32 [[SUB5]], ptr [[DOTCAPTURE_EXPR_4]], align 41287// CHECK3-NEXT:    [[TMP23:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 41288// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP23]], 11289// CHECK3-NEXT:    [[TMP24:%.*]] = zext i32 [[ADD]] to i641290// CHECK3-NEXT:    [[TMP25:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41291// CHECK3-NEXT:    [[TMP26:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP20]], 01292// CHECK3-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01293// CHECK3-NEXT:    store i32 3, ptr [[TMP27]], align 41294// CHECK3-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11295// CHECK3-NEXT:    store i32 3, ptr [[TMP28]], align 41296// CHECK3-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21297// CHECK3-NEXT:    store ptr [[TMP18]], ptr [[TMP29]], align 41298// CHECK3-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31299// CHECK3-NEXT:    store ptr [[TMP19]], ptr [[TMP30]], align 41300// CHECK3-NEXT:    [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41301// CHECK3-NEXT:    store ptr @.offload_sizes.3, ptr [[TMP31]], align 41302// CHECK3-NEXT:    [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51303// CHECK3-NEXT:    store ptr @.offload_maptypes.4, ptr [[TMP32]], align 41304// CHECK3-NEXT:    [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61305// CHECK3-NEXT:    store ptr null, ptr [[TMP33]], align 41306// CHECK3-NEXT:    [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71307// CHECK3-NEXT:    store ptr null, ptr [[TMP34]], align 41308// CHECK3-NEXT:    [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81309// CHECK3-NEXT:    store i64 [[TMP24]], ptr [[TMP35]], align 81310// CHECK3-NEXT:    [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91311// CHECK3-NEXT:    store i64 0, ptr [[TMP36]], align 81312// CHECK3-NEXT:    [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101313// CHECK3-NEXT:    store [3 x i32] [[TMP26]], ptr [[TMP37]], align 41314// CHECK3-NEXT:    [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111315// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP38]], align 41316// CHECK3-NEXT:    [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121317// CHECK3-NEXT:    store i32 [[TMP25]], ptr [[TMP39]], align 41318// CHECK3-NEXT:    [[TMP40:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[TMP20]], i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.region_id, ptr [[KERNEL_ARGS]])1319// CHECK3-NEXT:    [[TMP41:%.*]] = icmp ne i32 [[TMP40]], 01320// CHECK3-NEXT:    br i1 [[TMP41]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1321// CHECK3:       omp_offload.failed:1322// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104(i32 [[TMP4]], i32 [[TMP6]], i32 [[TMP8]]) #[[ATTR2]]1323// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT]]1324// CHECK3:       omp_offload.cont:1325// CHECK3-NEXT:    [[TMP42:%.*]] = load i32, ptr [[N_ADDR]], align 41326// CHECK3-NEXT:    [[ADD7:%.*]] = add nsw i32 32, [[TMP42]]1327// CHECK3-NEXT:    store i32 [[ADD7]], ptr [[DOTCAPTURE_EXPR_6]], align 41328// CHECK3-NEXT:    [[TMP43:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_6]], align 41329// CHECK3-NEXT:    store i32 [[TMP43]], ptr [[DOTCAPTURE_EXPR__CASTED8]], align 41330// CHECK3-NEXT:    [[TMP44:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED8]], align 41331// CHECK3-NEXT:    [[TMP45:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 01332// CHECK3-NEXT:    store i32 [[TMP44]], ptr [[TMP45]], align 41333// CHECK3-NEXT:    [[TMP46:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 01334// CHECK3-NEXT:    store i32 [[TMP44]], ptr [[TMP46]], align 41335// CHECK3-NEXT:    [[TMP47:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS11]], i32 0, i32 01336// CHECK3-NEXT:    store ptr null, ptr [[TMP47]], align 41337// CHECK3-NEXT:    [[TMP48:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 01338// CHECK3-NEXT:    [[TMP49:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 01339// CHECK3-NEXT:    [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT_ANON]], ptr [[AGG_CAPTURED]], i32 0, i32 01340// CHECK3-NEXT:    [[TMP51:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_6]], align 41341// CHECK3-NEXT:    store i32 [[TMP51]], ptr [[TMP50]], align 41342// CHECK3-NEXT:    [[TMP52:%.*]] = call ptr @__kmpc_omp_target_task_alloc(ptr @[[GLOB1]], i32 [[TMP0]], i32 1, i32 36, i32 4, ptr @.omp_task_entry., i64 -1)1343// CHECK3-NEXT:    [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T_WITH_PRIVATES:%.*]], ptr [[TMP52]], i32 0, i32 01344// CHECK3-NEXT:    [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T:%.*]], ptr [[TMP53]], i32 0, i32 01345// CHECK3-NEXT:    [[TMP55:%.*]] = load ptr, ptr [[TMP54]], align 41346// CHECK3-NEXT:    call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[TMP55]], ptr align 4 [[AGG_CAPTURED]], i32 4, i1 false)1347// CHECK3-NEXT:    [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T_WITH_PRIVATES]], ptr [[TMP52]], i32 0, i32 11348// CHECK3-NEXT:    [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T:%.*]], ptr [[TMP56]], i32 0, i32 01349// CHECK3-NEXT:    call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[TMP57]], ptr align 4 @.offload_sizes.5, i32 8, i1 false)1350// CHECK3-NEXT:    [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T]], ptr [[TMP56]], i32 0, i32 11351// CHECK3-NEXT:    call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[TMP58]], ptr align 4 [[TMP48]], i32 4, i1 false)1352// CHECK3-NEXT:    [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T]], ptr [[TMP56]], i32 0, i32 21353// CHECK3-NEXT:    call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[TMP59]], ptr align 4 [[TMP49]], i32 4, i1 false)1354// CHECK3-NEXT:    [[TMP60:%.*]] = call i32 @__kmpc_omp_task(ptr @[[GLOB1]], i32 [[TMP0]], ptr [[TMP52]])1355// CHECK3-NEXT:    [[TMP61:%.*]] = load i32, ptr [[N_ADDR]], align 41356// CHECK3-NEXT:    [[ADD12:%.*]] = add nsw i32 [[TMP61]], 11357// CHECK3-NEXT:    ret i32 [[ADD12]]1358//1359//1360// CHECK3-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i1361// CHECK3-SAME: (i32 noundef [[N:%.*]]) #[[ATTR0]] comdat {1362// CHECK3-NEXT:  entry:1363// CHECK3-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41364// CHECK3-NEXT:    [[A:%.*]] = alloca i32, align 41365// CHECK3-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81366// CHECK3-NEXT:    [[B:%.*]] = alloca i16, align 21367// CHECK3-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i16, align 21368// CHECK3-NEXT:    [[A_CASTED:%.*]] = alloca i32, align 41369// CHECK3-NEXT:    [[B_CASTED:%.*]] = alloca i32, align 41370// CHECK3-NEXT:    [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i32, align 41371// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 41372// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 41373// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 41374// CHECK3-NEXT:    [[KERNEL_ARGS1:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 81375// CHECK3-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41376// CHECK3-NEXT:    store i32 0, ptr [[A]], align 41377// CHECK3-NEXT:    [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01378// CHECK3-NEXT:    store i32 3, ptr [[TMP0]], align 41379// CHECK3-NEXT:    [[TMP1:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11380// CHECK3-NEXT:    store i32 0, ptr [[TMP1]], align 41381// CHECK3-NEXT:    [[TMP2:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21382// CHECK3-NEXT:    store ptr null, ptr [[TMP2]], align 41383// CHECK3-NEXT:    [[TMP3:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31384// CHECK3-NEXT:    store ptr null, ptr [[TMP3]], align 41385// CHECK3-NEXT:    [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41386// CHECK3-NEXT:    store ptr null, ptr [[TMP4]], align 41387// CHECK3-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51388// CHECK3-NEXT:    store ptr null, ptr [[TMP5]], align 41389// CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61390// CHECK3-NEXT:    store ptr null, ptr [[TMP6]], align 41391// CHECK3-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71392// CHECK3-NEXT:    store ptr null, ptr [[TMP7]], align 41393// CHECK3-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81394// CHECK3-NEXT:    store i64 0, ptr [[TMP8]], align 81395// CHECK3-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91396// CHECK3-NEXT:    store i64 0, ptr [[TMP9]], align 81397// CHECK3-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101398// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP10]], align 41399// CHECK3-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111400// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP11]], align 41401// CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121402// CHECK3-NEXT:    store i32 20, ptr [[TMP12]], align 41403// CHECK3-NEXT:    [[TMP13:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l88.region_id, ptr [[KERNEL_ARGS]])1404// CHECK3-NEXT:    [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 01405// CHECK3-NEXT:    br i1 [[TMP14]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1406// CHECK3:       omp_offload.failed:1407// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l88() #[[ATTR2]]1408// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT]]1409// CHECK3:       omp_offload.cont:1410// CHECK3-NEXT:    store i16 1, ptr [[B]], align 21411// CHECK3-NEXT:    [[TMP15:%.*]] = load i16, ptr [[B]], align 21412// CHECK3-NEXT:    store i16 [[TMP15]], ptr [[DOTCAPTURE_EXPR_]], align 21413// CHECK3-NEXT:    [[TMP16:%.*]] = load i32, ptr [[A]], align 41414// CHECK3-NEXT:    store i32 [[TMP16]], ptr [[A_CASTED]], align 41415// CHECK3-NEXT:    [[TMP17:%.*]] = load i32, ptr [[A_CASTED]], align 41416// CHECK3-NEXT:    [[TMP18:%.*]] = load i16, ptr [[B]], align 21417// CHECK3-NEXT:    store i16 [[TMP18]], ptr [[B_CASTED]], align 21418// CHECK3-NEXT:    [[TMP19:%.*]] = load i32, ptr [[B_CASTED]], align 41419// CHECK3-NEXT:    [[TMP20:%.*]] = load i16, ptr [[DOTCAPTURE_EXPR_]], align 21420// CHECK3-NEXT:    store i16 [[TMP20]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 21421// CHECK3-NEXT:    [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED]], align 41422// CHECK3-NEXT:    [[TMP22:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01423// CHECK3-NEXT:    store i32 [[TMP17]], ptr [[TMP22]], align 41424// CHECK3-NEXT:    [[TMP23:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01425// CHECK3-NEXT:    store i32 [[TMP17]], ptr [[TMP23]], align 41426// CHECK3-NEXT:    [[TMP24:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01427// CHECK3-NEXT:    store ptr null, ptr [[TMP24]], align 41428// CHECK3-NEXT:    [[TMP25:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11429// CHECK3-NEXT:    store i32 [[TMP19]], ptr [[TMP25]], align 41430// CHECK3-NEXT:    [[TMP26:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11431// CHECK3-NEXT:    store i32 [[TMP19]], ptr [[TMP26]], align 41432// CHECK3-NEXT:    [[TMP27:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 11433// CHECK3-NEXT:    store ptr null, ptr [[TMP27]], align 41434// CHECK3-NEXT:    [[TMP28:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21435// CHECK3-NEXT:    store i32 [[TMP21]], ptr [[TMP28]], align 41436// CHECK3-NEXT:    [[TMP29:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21437// CHECK3-NEXT:    store i32 [[TMP21]], ptr [[TMP29]], align 41438// CHECK3-NEXT:    [[TMP30:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 21439// CHECK3-NEXT:    store ptr null, ptr [[TMP30]], align 41440// CHECK3-NEXT:    [[TMP31:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01441// CHECK3-NEXT:    [[TMP32:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01442// CHECK3-NEXT:    [[TMP33:%.*]] = load i16, ptr [[DOTCAPTURE_EXPR_]], align 21443// CHECK3-NEXT:    [[TMP34:%.*]] = sext i16 [[TMP33]] to i321444// CHECK3-NEXT:    [[TMP35:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP34]], 01445// CHECK3-NEXT:    [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 01446// CHECK3-NEXT:    store i32 3, ptr [[TMP36]], align 41447// CHECK3-NEXT:    [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 11448// CHECK3-NEXT:    store i32 3, ptr [[TMP37]], align 41449// CHECK3-NEXT:    [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 21450// CHECK3-NEXT:    store ptr [[TMP31]], ptr [[TMP38]], align 41451// CHECK3-NEXT:    [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 31452// CHECK3-NEXT:    store ptr [[TMP32]], ptr [[TMP39]], align 41453// CHECK3-NEXT:    [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 41454// CHECK3-NEXT:    store ptr @.offload_sizes.7, ptr [[TMP40]], align 41455// CHECK3-NEXT:    [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 51456// CHECK3-NEXT:    store ptr @.offload_maptypes.8, ptr [[TMP41]], align 41457// CHECK3-NEXT:    [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 61458// CHECK3-NEXT:    store ptr null, ptr [[TMP42]], align 41459// CHECK3-NEXT:    [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 71460// CHECK3-NEXT:    store ptr null, ptr [[TMP43]], align 41461// CHECK3-NEXT:    [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 81462// CHECK3-NEXT:    store i64 0, ptr [[TMP44]], align 81463// CHECK3-NEXT:    [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 91464// CHECK3-NEXT:    store i64 0, ptr [[TMP45]], align 81465// CHECK3-NEXT:    [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 101466// CHECK3-NEXT:    store [3 x i32] [[TMP35]], ptr [[TMP46]], align 41467// CHECK3-NEXT:    [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 111468// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP47]], align 41469// CHECK3-NEXT:    [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 121470// CHECK3-NEXT:    store i32 1024, ptr [[TMP48]], align 41471// CHECK3-NEXT:    [[TMP49:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[TMP34]], i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93.region_id, ptr [[KERNEL_ARGS1]])1472// CHECK3-NEXT:    [[TMP50:%.*]] = icmp ne i32 [[TMP49]], 01473// CHECK3-NEXT:    br i1 [[TMP50]], label [[OMP_OFFLOAD_FAILED2:%.*]], label [[OMP_OFFLOAD_CONT3:%.*]]1474// CHECK3:       omp_offload.failed2:1475// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93(i32 [[TMP17]], i32 [[TMP19]], i32 [[TMP21]]) #[[ATTR2]]1476// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT3]]1477// CHECK3:       omp_offload.cont3:1478// CHECK3-NEXT:    [[TMP51:%.*]] = load i32, ptr [[A]], align 41479// CHECK3-NEXT:    ret i32 [[TMP51]]1480//1481//1482// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l1211483// CHECK3-SAME: (ptr noundef [[THIS:%.*]], i32 noundef [[B:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1:[0-9]+]] {1484// CHECK3-NEXT:  entry:1485// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 41486// CHECK3-NEXT:    [[B_ADDR:%.*]] = alloca i32, align 41487// CHECK3-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 41488// CHECK3-NEXT:    [[B_CASTED:%.*]] = alloca i32, align 41489// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 41490// CHECK3-NEXT:    store i32 [[B]], ptr [[B_ADDR]], align 41491// CHECK3-NEXT:    store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 41492// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41493// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[B_ADDR]], align 41494// CHECK3-NEXT:    store i32 [[TMP1]], ptr [[B_CASTED]], align 41495// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[B_CASTED]], align 41496// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.omp_outlined, ptr [[TMP0]], i32 [[TMP2]])1497// CHECK3-NEXT:    ret void1498//1499//1500// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.omp_outlined1501// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]], i32 noundef [[B:%.*]]) #[[ATTR1]] {1502// CHECK3-NEXT:  entry:1503// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41504// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41505// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 41506// CHECK3-NEXT:    [[B_ADDR:%.*]] = alloca i32, align 41507// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41508// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41509// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 41510// CHECK3-NEXT:    store i32 [[B]], ptr [[B_ADDR]], align 41511// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41512// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[B_ADDR]], align 41513// CHECK3-NEXT:    [[CONV:%.*]] = sitofp i32 [[TMP1]] to double1514// CHECK3-NEXT:    [[ADD:%.*]] = fadd double [[CONV]], 1.500000e+001515// CHECK3-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 01516// CHECK3-NEXT:    store double [[ADD]], ptr [[A]], align 41517// CHECK3-NEXT:    ret void1518//1519//1520// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l1261521// CHECK3-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1]] {1522// CHECK3-NEXT:  entry:1523// CHECK3-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 41524// CHECK3-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 41525// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41526// CHECK3-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 01527// CHECK3-NEXT:    store double 2.500000e+00, ptr [[A]], align 41528// CHECK3-NEXT:    ret void1529//1530//1531// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l1041532// CHECK3-SAME: (i32 noundef [[N:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]], i32 noundef [[DOTCAPTURE_EXPR_1:%.*]]) #[[ATTR1]] {1533// CHECK3-NEXT:  entry:1534// CHECK3-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41535// CHECK3-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 41536// CHECK3-NEXT:    [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca i32, align 41537// CHECK3-NEXT:    [[N_CASTED:%.*]] = alloca i32, align 41538// CHECK3-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1539// CHECK3-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41540// CHECK3-NEXT:    store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 41541// CHECK3-NEXT:    store i32 [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 41542// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 41543// CHECK3-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 0)1544// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 41545// CHECK3-NEXT:    store i32 [[TMP2]], ptr [[N_CASTED]], align 41546// CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[N_CASTED]], align 41547// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined, i32 [[TMP3]])1548// CHECK3-NEXT:    ret void1549//1550//1551// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined1552// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[N:%.*]]) #[[ATTR1]] {1553// CHECK3-NEXT:  entry:1554// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41555// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41556// CHECK3-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41557// CHECK3-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41558// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 41559// CHECK3-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41560// CHECK3-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41561// CHECK3-NEXT:    [[I:%.*]] = alloca i32, align 41562// CHECK3-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41563// CHECK3-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41564// CHECK3-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41565// CHECK3-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41566// CHECK3-NEXT:    [[I3:%.*]] = alloca i32, align 41567// CHECK3-NEXT:    [[N_CASTED:%.*]] = alloca i32, align 41568// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41569// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41570// CHECK3-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41571// CHECK3-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 41572// CHECK3-NEXT:    store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 41573// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41574// CHECK3-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP1]], 01575// CHECK3-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 11576// CHECK3-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11577// CHECK3-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41578// CHECK3-NEXT:    store i32 0, ptr [[I]], align 41579// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41580// CHECK3-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP2]]1581// CHECK3-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1582// CHECK3:       omp.precond.then:1583// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 41584// CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41585// CHECK3-NEXT:    store i32 [[TMP3]], ptr [[DOTOMP_COMB_UB]], align 41586// CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41587// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41588// CHECK3-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41589// CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 41590// CHECK3-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP5]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1591// CHECK3-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41592// CHECK3-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41593// CHECK3-NEXT:    [[CMP4:%.*]] = icmp sgt i32 [[TMP6]], [[TMP7]]1594// CHECK3-NEXT:    br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1595// CHECK3:       cond.true:1596// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41597// CHECK3-NEXT:    br label [[COND_END:%.*]]1598// CHECK3:       cond.false:1599// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41600// CHECK3-NEXT:    br label [[COND_END]]1601// CHECK3:       cond.end:1602// CHECK3-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP8]], [[COND_TRUE]] ], [ [[TMP9]], [[COND_FALSE]] ]1603// CHECK3-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41604// CHECK3-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41605// CHECK3-NEXT:    store i32 [[TMP10]], ptr [[DOTOMP_IV]], align 41606// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1607// CHECK3:       omp.inner.for.cond:1608// CHECK3-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12:![0-9]+]]1609// CHECK3-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4, !llvm.access.group [[ACC_GRP12]]1610// CHECK3-NEXT:    [[CMP5:%.*]] = icmp sle i32 [[TMP11]], [[TMP12]]1611// CHECK3-NEXT:    br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1612// CHECK3:       omp.inner.for.body:1613// CHECK3-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4, !llvm.access.group [[ACC_GRP12]]1614// CHECK3-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4, !llvm.access.group [[ACC_GRP12]]1615// CHECK3-NEXT:    [[TMP15:%.*]] = load i32, ptr [[N_ADDR]], align 4, !llvm.access.group [[ACC_GRP12]]1616// CHECK3-NEXT:    store i32 [[TMP15]], ptr [[N_CASTED]], align 4, !llvm.access.group [[ACC_GRP12]]1617// CHECK3-NEXT:    [[TMP16:%.*]] = load i32, ptr [[N_CASTED]], align 4, !llvm.access.group [[ACC_GRP12]]1618// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined.omp_outlined, i32 [[TMP13]], i32 [[TMP14]], i32 [[TMP16]]), !llvm.access.group [[ACC_GRP12]]1619// CHECK3-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1620// CHECK3:       omp.inner.for.inc:1621// CHECK3-NEXT:    [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]1622// CHECK3-NEXT:    [[TMP18:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4, !llvm.access.group [[ACC_GRP12]]1623// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP17]], [[TMP18]]1624// CHECK3-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]1625// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP13:![0-9]+]]1626// CHECK3:       omp.inner.for.end:1627// CHECK3-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1628// CHECK3:       omp.loop.exit:1629// CHECK3-NEXT:    [[TMP19:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41630// CHECK3-NEXT:    [[TMP20:%.*]] = load i32, ptr [[TMP19]], align 41631// CHECK3-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP20]])1632// CHECK3-NEXT:    [[TMP21:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 41633// CHECK3-NEXT:    [[TMP22:%.*]] = icmp ne i32 [[TMP21]], 01634// CHECK3-NEXT:    br i1 [[TMP22]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]1635// CHECK3:       .omp.final.then:1636// CHECK3-NEXT:    [[TMP23:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41637// CHECK3-NEXT:    [[SUB6:%.*]] = sub nsw i32 [[TMP23]], 01638// CHECK3-NEXT:    [[DIV7:%.*]] = sdiv i32 [[SUB6]], 11639// CHECK3-NEXT:    [[MUL:%.*]] = mul nsw i32 [[DIV7]], 11640// CHECK3-NEXT:    [[ADD8:%.*]] = add nsw i32 0, [[MUL]]1641// CHECK3-NEXT:    store i32 [[ADD8]], ptr [[I3]], align 41642// CHECK3-NEXT:    br label [[DOTOMP_FINAL_DONE]]1643// CHECK3:       .omp.final.done:1644// CHECK3-NEXT:    br label [[OMP_PRECOND_END]]1645// CHECK3:       omp.precond.end:1646// CHECK3-NEXT:    ret void1647//1648//1649// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined.omp_outlined1650// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[DOTPREVIOUS_LB_:%.*]], i32 noundef [[DOTPREVIOUS_UB_:%.*]], i32 noundef [[N:%.*]]) #[[ATTR1]] {1651// CHECK3-NEXT:  entry:1652// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41653// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41654// CHECK3-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i32, align 41655// CHECK3-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i32, align 41656// CHECK3-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41657// CHECK3-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41658// CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 41659// CHECK3-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41660// CHECK3-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41661// CHECK3-NEXT:    [[I:%.*]] = alloca i32, align 41662// CHECK3-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 41663// CHECK3-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 41664// CHECK3-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41665// CHECK3-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41666// CHECK3-NEXT:    [[I3:%.*]] = alloca i32, align 41667// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41668// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41669// CHECK3-NEXT:    store i32 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 41670// CHECK3-NEXT:    store i32 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 41671// CHECK3-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41672// CHECK3-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 41673// CHECK3-NEXT:    store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 41674// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41675// CHECK3-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP1]], 01676// CHECK3-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 11677// CHECK3-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11678// CHECK3-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41679// CHECK3-NEXT:    store i32 0, ptr [[I]], align 41680// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41681// CHECK3-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP2]]1682// CHECK3-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1683// CHECK3:       omp.precond.then:1684// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 41685// CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41686// CHECK3-NEXT:    store i32 [[TMP3]], ptr [[DOTOMP_UB]], align 41687// CHECK3-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTPREVIOUS_LB__ADDR]], align 41688// CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTPREVIOUS_UB__ADDR]], align 41689// CHECK3-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_LB]], align 41690// CHECK3-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_UB]], align 41691// CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41692// CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41693// CHECK3-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41694// CHECK3-NEXT:    [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 41695// CHECK3-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB3:[0-9]+]], i32 [[TMP7]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1696// CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41697// CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41698// CHECK3-NEXT:    [[CMP4:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]1699// CHECK3-NEXT:    br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1700// CHECK3:       cond.true:1701// CHECK3-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41702// CHECK3-NEXT:    br label [[COND_END:%.*]]1703// CHECK3:       cond.false:1704// CHECK3-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41705// CHECK3-NEXT:    br label [[COND_END]]1706// CHECK3:       cond.end:1707// CHECK3-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]1708// CHECK3-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 41709// CHECK3-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41710// CHECK3-NEXT:    store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 41711// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1712// CHECK3:       omp.inner.for.cond:1713// CHECK3-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16:![0-9]+]]1714// CHECK3-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP16]]1715// CHECK3-NEXT:    [[CMP5:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]1716// CHECK3-NEXT:    br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1717// CHECK3:       omp.inner.for.body:1718// CHECK3-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]1719// CHECK3-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP15]], 11720// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1721// CHECK3-NEXT:    store i32 [[ADD]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP16]]1722// CHECK3-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1723// CHECK3:       omp.body.continue:1724// CHECK3-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1725// CHECK3:       omp.inner.for.inc:1726// CHECK3-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]1727// CHECK3-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP16]], 11728// CHECK3-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]1729// CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP17:![0-9]+]]1730// CHECK3:       omp.inner.for.end:1731// CHECK3-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]1732// CHECK3:       omp.loop.exit:1733// CHECK3-NEXT:    [[TMP17:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41734// CHECK3-NEXT:    [[TMP18:%.*]] = load i32, ptr [[TMP17]], align 41735// CHECK3-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB3]], i32 [[TMP18]])1736// CHECK3-NEXT:    [[TMP19:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 41737// CHECK3-NEXT:    [[TMP20:%.*]] = icmp ne i32 [[TMP19]], 01738// CHECK3-NEXT:    br i1 [[TMP20]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]1739// CHECK3:       .omp.final.then:1740// CHECK3-NEXT:    [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41741// CHECK3-NEXT:    [[SUB7:%.*]] = sub nsw i32 [[TMP21]], 01742// CHECK3-NEXT:    [[DIV8:%.*]] = sdiv i32 [[SUB7]], 11743// CHECK3-NEXT:    [[MUL9:%.*]] = mul nsw i32 [[DIV8]], 11744// CHECK3-NEXT:    [[ADD10:%.*]] = add nsw i32 0, [[MUL9]]1745// CHECK3-NEXT:    store i32 [[ADD10]], ptr [[I3]], align 41746// CHECK3-NEXT:    br label [[DOTOMP_FINAL_DONE]]1747// CHECK3:       .omp.final.done:1748// CHECK3-NEXT:    br label [[OMP_PRECOND_END]]1749// CHECK3:       omp.precond.end:1750// CHECK3-NEXT:    ret void1751//1752//1753// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l1081754// CHECK3-SAME: (i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1]] {1755// CHECK3-NEXT:  entry:1756// CHECK3-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 41757// CHECK3-NEXT:    store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 41758// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108.omp_outlined)1759// CHECK3-NEXT:    ret void1760//1761//1762// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108.omp_outlined1763// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {1764// CHECK3-NEXT:  entry:1765// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41766// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41767// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41768// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41769// CHECK3-NEXT:    ret void1770//1771//1772// CHECK3-LABEL: define {{[^@]+}}@.omp_task_privates_map.1773// CHECK3-SAME: (ptr noalias noundef [[TMP0:%.*]], ptr noalias noundef [[TMP1:%.*]], ptr noalias noundef [[TMP2:%.*]], ptr noalias noundef [[TMP3:%.*]]) #[[ATTR3:[0-9]+]] {1774// CHECK3-NEXT:  entry:1775// CHECK3-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 41776// CHECK3-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 41777// CHECK3-NEXT:    [[DOTADDR2:%.*]] = alloca ptr, align 41778// CHECK3-NEXT:    [[DOTADDR3:%.*]] = alloca ptr, align 41779// CHECK3-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 41780// CHECK3-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 41781// CHECK3-NEXT:    store ptr [[TMP2]], ptr [[DOTADDR2]], align 41782// CHECK3-NEXT:    store ptr [[TMP3]], ptr [[DOTADDR3]], align 41783// CHECK3-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTADDR]], align 41784// CHECK3-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T:%.*]], ptr [[TMP4]], i32 0, i32 01785// CHECK3-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[DOTADDR3]], align 41786// CHECK3-NEXT:    store ptr [[TMP5]], ptr [[TMP6]], align 41787// CHECK3-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T]], ptr [[TMP4]], i32 0, i32 11788// CHECK3-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[DOTADDR1]], align 41789// CHECK3-NEXT:    store ptr [[TMP7]], ptr [[TMP8]], align 41790// CHECK3-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T]], ptr [[TMP4]], i32 0, i32 21791// CHECK3-NEXT:    [[TMP10:%.*]] = load ptr, ptr [[DOTADDR2]], align 41792// CHECK3-NEXT:    store ptr [[TMP9]], ptr [[TMP10]], align 41793// CHECK3-NEXT:    ret void1794//1795//1796// CHECK3-LABEL: define {{[^@]+}}@.omp_task_entry.1797// CHECK3-SAME: (i32 noundef [[TMP0:%.*]], ptr noalias noundef [[TMP1:%.*]]) #[[ATTR4:[0-9]+]] {1798// CHECK3-NEXT:  entry:1799// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR_I:%.*]] = alloca i32, align 41800// CHECK3-NEXT:    [[DOTPART_ID__ADDR_I:%.*]] = alloca ptr, align 41801// CHECK3-NEXT:    [[DOTPRIVATES__ADDR_I:%.*]] = alloca ptr, align 41802// CHECK3-NEXT:    [[DOTCOPY_FN__ADDR_I:%.*]] = alloca ptr, align 41803// CHECK3-NEXT:    [[DOTTASK_T__ADDR_I:%.*]] = alloca ptr, align 41804// CHECK3-NEXT:    [[__CONTEXT_ADDR_I:%.*]] = alloca ptr, align 41805// CHECK3-NEXT:    [[DOTFIRSTPRIV_PTR_ADDR_I:%.*]] = alloca ptr, align 41806// CHECK3-NEXT:    [[DOTFIRSTPRIV_PTR_ADDR1_I:%.*]] = alloca ptr, align 41807// CHECK3-NEXT:    [[DOTFIRSTPRIV_PTR_ADDR2_I:%.*]] = alloca ptr, align 41808// CHECK3-NEXT:    [[KERNEL_ARGS_I:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81809// CHECK3-NEXT:    [[DOTCAPTURE_EXPR__CASTED_I:%.*]] = alloca i32, align 41810// CHECK3-NEXT:    [[DOTADDR:%.*]] = alloca i32, align 41811// CHECK3-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 41812// CHECK3-NEXT:    store i32 [[TMP0]], ptr [[DOTADDR]], align 41813// CHECK3-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 41814// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTADDR]], align 41815// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 41816// CHECK3-NEXT:    [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T_WITH_PRIVATES:%.*]], ptr [[TMP3]], i32 0, i32 01817// CHECK3-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T:%.*]], ptr [[TMP4]], i32 0, i32 21818// CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T]], ptr [[TMP4]], i32 0, i32 01819// CHECK3-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 41820// CHECK3-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T_WITH_PRIVATES]], ptr [[TMP3]], i32 0, i32 11821// CHECK3-NEXT:    call void @llvm.experimental.noalias.scope.decl(metadata [[META19:![0-9]+]])1822// CHECK3-NEXT:    call void @llvm.experimental.noalias.scope.decl(metadata [[META22:![0-9]+]])1823// CHECK3-NEXT:    call void @llvm.experimental.noalias.scope.decl(metadata [[META24:![0-9]+]])1824// CHECK3-NEXT:    call void @llvm.experimental.noalias.scope.decl(metadata [[META26:![0-9]+]])1825// CHECK3-NEXT:    store i32 [[TMP2]], ptr [[DOTGLOBAL_TID__ADDR_I]], align 4, !noalias [[META28:![0-9]+]]1826// CHECK3-NEXT:    store ptr [[TMP5]], ptr [[DOTPART_ID__ADDR_I]], align 4, !noalias [[META28]]1827// CHECK3-NEXT:    store ptr [[TMP8]], ptr [[DOTPRIVATES__ADDR_I]], align 4, !noalias [[META28]]1828// CHECK3-NEXT:    store ptr @.omp_task_privates_map., ptr [[DOTCOPY_FN__ADDR_I]], align 4, !noalias [[META28]]1829// CHECK3-NEXT:    store ptr [[TMP3]], ptr [[DOTTASK_T__ADDR_I]], align 4, !noalias [[META28]]1830// CHECK3-NEXT:    store ptr [[TMP7]], ptr [[__CONTEXT_ADDR_I]], align 4, !noalias [[META28]]1831// CHECK3-NEXT:    [[TMP9:%.*]] = load ptr, ptr [[__CONTEXT_ADDR_I]], align 4, !noalias [[META28]]1832// CHECK3-NEXT:    [[TMP10:%.*]] = load ptr, ptr [[DOTCOPY_FN__ADDR_I]], align 4, !noalias [[META28]]1833// CHECK3-NEXT:    [[TMP11:%.*]] = load ptr, ptr [[DOTPRIVATES__ADDR_I]], align 4, !noalias [[META28]]1834// CHECK3-NEXT:    call void [[TMP10]](ptr [[TMP11]], ptr [[DOTFIRSTPRIV_PTR_ADDR_I]], ptr [[DOTFIRSTPRIV_PTR_ADDR1_I]], ptr [[DOTFIRSTPRIV_PTR_ADDR2_I]]) #[[ATTR2]]1835// CHECK3-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[DOTFIRSTPRIV_PTR_ADDR_I]], align 4, !noalias [[META28]]1836// CHECK3-NEXT:    [[TMP13:%.*]] = load ptr, ptr [[DOTFIRSTPRIV_PTR_ADDR1_I]], align 4, !noalias [[META28]]1837// CHECK3-NEXT:    [[TMP14:%.*]] = load ptr, ptr [[DOTFIRSTPRIV_PTR_ADDR2_I]], align 4, !noalias [[META28]]1838// CHECK3-NEXT:    [[TMP15:%.*]] = load i32, ptr [[TMP9]], align 41839// CHECK3-NEXT:    store i32 3, ptr [[KERNEL_ARGS_I]], align 4, !noalias [[META28]]1840// CHECK3-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 11841// CHECK3-NEXT:    store i32 1, ptr [[TMP16]], align 4, !noalias [[META28]]1842// CHECK3-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 21843// CHECK3-NEXT:    store ptr [[TMP12]], ptr [[TMP17]], align 4, !noalias [[META28]]1844// CHECK3-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 31845// CHECK3-NEXT:    store ptr [[TMP13]], ptr [[TMP18]], align 4, !noalias [[META28]]1846// CHECK3-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 41847// CHECK3-NEXT:    store ptr [[TMP14]], ptr [[TMP19]], align 4, !noalias [[META28]]1848// CHECK3-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 51849// CHECK3-NEXT:    store ptr @.offload_maptypes.6, ptr [[TMP20]], align 4, !noalias [[META28]]1850// CHECK3-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 61851// CHECK3-NEXT:    store ptr null, ptr [[TMP21]], align 4, !noalias [[META28]]1852// CHECK3-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 71853// CHECK3-NEXT:    store ptr null, ptr [[TMP22]], align 4, !noalias [[META28]]1854// CHECK3-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 81855// CHECK3-NEXT:    store i64 0, ptr [[TMP23]], align 8, !noalias [[META28]]1856// CHECK3-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 91857// CHECK3-NEXT:    store i64 1, ptr [[TMP24]], align 8, !noalias [[META28]]1858// CHECK3-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 101859// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP25]], align 4, !noalias [[META28]]1860// CHECK3-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 111861// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP26]], align 4, !noalias [[META28]]1862// CHECK3-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 121863// CHECK3-NEXT:    store i32 [[TMP15]], ptr [[TMP27]], align 4, !noalias [[META28]]1864// CHECK3-NEXT:    [[TMP28:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108.region_id, ptr [[KERNEL_ARGS_I]])1865// CHECK3-NEXT:    [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 01866// CHECK3-NEXT:    br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED_I:%.*]], label [[DOTOMP_OUTLINED__EXIT:%.*]]1867// CHECK3:       omp_offload.failed.i:1868// CHECK3-NEXT:    [[TMP30:%.*]] = load i32, ptr [[TMP9]], align 41869// CHECK3-NEXT:    store i32 [[TMP30]], ptr [[DOTCAPTURE_EXPR__CASTED_I]], align 4, !noalias [[META28]]1870// CHECK3-NEXT:    [[TMP31:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED_I]], align 4, !noalias [[META28]]1871// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108(i32 [[TMP31]]) #[[ATTR2]]1872// CHECK3-NEXT:    br label [[DOTOMP_OUTLINED__EXIT]]1873// CHECK3:       .omp_outlined..exit:1874// CHECK3-NEXT:    ret i32 01875//1876//1877// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l881878// CHECK3-SAME: () #[[ATTR1]] {1879// CHECK3-NEXT:  entry:1880// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l88.omp_outlined)1881// CHECK3-NEXT:    ret void1882//1883//1884// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l88.omp_outlined1885// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {1886// CHECK3-NEXT:  entry:1887// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41888// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41889// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41890// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41891// CHECK3-NEXT:    ret void1892//1893//1894// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l931895// CHECK3-SAME: (i32 noundef [[A:%.*]], i32 noundef [[B:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1]] {1896// CHECK3-NEXT:  entry:1897// CHECK3-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 41898// CHECK3-NEXT:    [[B_ADDR:%.*]] = alloca i32, align 41899// CHECK3-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 41900// CHECK3-NEXT:    [[A_CASTED:%.*]] = alloca i32, align 41901// CHECK3-NEXT:    [[B_CASTED:%.*]] = alloca i32, align 41902// CHECK3-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1903// CHECK3-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 41904// CHECK3-NEXT:    store i32 [[B]], ptr [[B_ADDR]], align 41905// CHECK3-NEXT:    store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 41906// CHECK3-NEXT:    [[TMP1:%.*]] = load i16, ptr [[DOTCAPTURE_EXPR__ADDR]], align 21907// CHECK3-NEXT:    [[TMP2:%.*]] = sext i16 [[TMP1]] to i321908// CHECK3-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 0)1909// CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A_ADDR]], align 41910// CHECK3-NEXT:    store i32 [[TMP3]], ptr [[A_CASTED]], align 41911// CHECK3-NEXT:    [[TMP4:%.*]] = load i32, ptr [[A_CASTED]], align 41912// CHECK3-NEXT:    [[TMP5:%.*]] = load i16, ptr [[B_ADDR]], align 21913// CHECK3-NEXT:    store i16 [[TMP5]], ptr [[B_CASTED]], align 21914// CHECK3-NEXT:    [[TMP6:%.*]] = load i32, ptr [[B_CASTED]], align 41915// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93.omp_outlined, i32 [[TMP4]], i32 [[TMP6]])1916// CHECK3-NEXT:    ret void1917//1918//1919// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93.omp_outlined1920// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[A:%.*]], i32 noundef [[B:%.*]]) #[[ATTR1]] {1921// CHECK3-NEXT:  entry:1922// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41923// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41924// CHECK3-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 41925// CHECK3-NEXT:    [[B_ADDR:%.*]] = alloca i32, align 41926// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41927// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41928// CHECK3-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 41929// CHECK3-NEXT:    store i32 [[B]], ptr [[B_ADDR]], align 41930// CHECK3-NEXT:    [[TMP0:%.*]] = load i16, ptr [[B_ADDR]], align 21931// CHECK3-NEXT:    [[CONV:%.*]] = sext i16 [[TMP0]] to i321932// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 41933// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CONV]]1934// CHECK3-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 41935// CHECK3-NEXT:    ret void1936//1937//1938// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l1041939// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[N:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]], i64 noundef [[DOTCAPTURE_EXPR_1:%.*]]) #[[ATTR0:[0-9]+]] {1940// CHECK9-NEXT:  entry:1941// CHECK9-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81942// CHECK9-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 81943// CHECK9-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 81944// CHECK9-NEXT:    [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca i64, align 81945// CHECK9-NEXT:    [[N_CASTED:%.*]] = alloca i64, align 81946// CHECK9-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB3:[0-9]+]])1947// CHECK9-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81948// CHECK9-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 81949// CHECK9-NEXT:    store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 81950// CHECK9-NEXT:    store i64 [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 81951// CHECK9-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 41952// CHECK9-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB3]], i32 [[TMP0]], i32 [[TMP1]], i32 0)1953// CHECK9-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 41954// CHECK9-NEXT:    store i32 [[TMP2]], ptr [[N_CASTED]], align 41955// CHECK9-NEXT:    [[TMP3:%.*]] = load i64, ptr [[N_CASTED]], align 81956// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined, i64 [[TMP3]])1957// CHECK9-NEXT:    ret void1958//1959//1960// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined1961// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[N:%.*]]) #[[ATTR0]] {1962// CHECK9-NEXT:  entry:1963// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81964// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81965// CHECK9-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 81966// CHECK9-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41967// CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 41968// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41969// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41970// CHECK9-NEXT:    [[I:%.*]] = alloca i32, align 41971// CHECK9-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41972// CHECK9-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41973// CHECK9-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41974// CHECK9-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41975// CHECK9-NEXT:    [[I3:%.*]] = alloca i32, align 41976// CHECK9-NEXT:    [[N_CASTED:%.*]] = alloca i64, align 81977// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81978// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81979// CHECK9-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 81980// CHECK9-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 41981// CHECK9-NEXT:    store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 41982// CHECK9-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41983// CHECK9-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP1]], 01984// CHECK9-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 11985// CHECK9-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11986// CHECK9-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41987// CHECK9-NEXT:    store i32 0, ptr [[I]], align 41988// CHECK9-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41989// CHECK9-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP2]]1990// CHECK9-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1991// CHECK9:       omp.precond.then:1992// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 41993// CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41994// CHECK9-NEXT:    store i32 [[TMP3]], ptr [[DOTOMP_COMB_UB]], align 41995// CHECK9-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 41996// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 41997// CHECK9-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81998// CHECK9-NEXT:    [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 41999// CHECK9-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP5]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2000// CHECK9-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42001// CHECK9-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42002// CHECK9-NEXT:    [[CMP4:%.*]] = icmp sgt i32 [[TMP6]], [[TMP7]]2003// CHECK9-NEXT:    br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2004// CHECK9:       cond.true:2005// CHECK9-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42006// CHECK9-NEXT:    br label [[COND_END:%.*]]2007// CHECK9:       cond.false:2008// CHECK9-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42009// CHECK9-NEXT:    br label [[COND_END]]2010// CHECK9:       cond.end:2011// CHECK9-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP8]], [[COND_TRUE]] ], [ [[TMP9]], [[COND_FALSE]] ]2012// CHECK9-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 42013// CHECK9-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 42014// CHECK9-NEXT:    store i32 [[TMP10]], ptr [[DOTOMP_IV]], align 42015// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]2016// CHECK9:       omp.inner.for.cond:2017// CHECK9-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP10:![0-9]+]]2018// CHECK9-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4, !llvm.access.group [[ACC_GRP10]]2019// CHECK9-NEXT:    [[CMP5:%.*]] = icmp sle i32 [[TMP11]], [[TMP12]]2020// CHECK9-NEXT:    br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2021// CHECK9:       omp.inner.for.body:2022// CHECK9-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4, !llvm.access.group [[ACC_GRP10]]2023// CHECK9-NEXT:    [[TMP14:%.*]] = zext i32 [[TMP13]] to i642024// CHECK9-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4, !llvm.access.group [[ACC_GRP10]]2025// CHECK9-NEXT:    [[TMP16:%.*]] = zext i32 [[TMP15]] to i642026// CHECK9-NEXT:    [[TMP17:%.*]] = load i32, ptr [[N_ADDR]], align 4, !llvm.access.group [[ACC_GRP10]]2027// CHECK9-NEXT:    store i32 [[TMP17]], ptr [[N_CASTED]], align 4, !llvm.access.group [[ACC_GRP10]]2028// CHECK9-NEXT:    [[TMP18:%.*]] = load i64, ptr [[N_CASTED]], align 8, !llvm.access.group [[ACC_GRP10]]2029// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined.omp_outlined, i64 [[TMP14]], i64 [[TMP16]], i64 [[TMP18]]), !llvm.access.group [[ACC_GRP10]]2030// CHECK9-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]2031// CHECK9:       omp.inner.for.inc:2032// CHECK9-NEXT:    [[TMP19:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP10]]2033// CHECK9-NEXT:    [[TMP20:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4, !llvm.access.group [[ACC_GRP10]]2034// CHECK9-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP19]], [[TMP20]]2035// CHECK9-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP10]]2036// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP11:![0-9]+]]2037// CHECK9:       omp.inner.for.end:2038// CHECK9-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]2039// CHECK9:       omp.loop.exit:2040// CHECK9-NEXT:    [[TMP21:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82041// CHECK9-NEXT:    [[TMP22:%.*]] = load i32, ptr [[TMP21]], align 42042// CHECK9-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP22]])2043// CHECK9-NEXT:    [[TMP23:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 42044// CHECK9-NEXT:    [[TMP24:%.*]] = icmp ne i32 [[TMP23]], 02045// CHECK9-NEXT:    br i1 [[TMP24]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]2046// CHECK9:       .omp.final.then:2047// CHECK9-NEXT:    [[TMP25:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42048// CHECK9-NEXT:    [[SUB6:%.*]] = sub nsw i32 [[TMP25]], 02049// CHECK9-NEXT:    [[DIV7:%.*]] = sdiv i32 [[SUB6]], 12050// CHECK9-NEXT:    [[MUL:%.*]] = mul nsw i32 [[DIV7]], 12051// CHECK9-NEXT:    [[ADD8:%.*]] = add nsw i32 0, [[MUL]]2052// CHECK9-NEXT:    store i32 [[ADD8]], ptr [[I3]], align 42053// CHECK9-NEXT:    br label [[DOTOMP_FINAL_DONE]]2054// CHECK9:       .omp.final.done:2055// CHECK9-NEXT:    br label [[OMP_PRECOND_END]]2056// CHECK9:       omp.precond.end:2057// CHECK9-NEXT:    ret void2058//2059//2060// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined.omp_outlined2061// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]], i64 noundef [[N:%.*]]) #[[ATTR0]] {2062// CHECK9-NEXT:  entry:2063// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82064// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82065// CHECK9-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 82066// CHECK9-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 82067// CHECK9-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 82068// CHECK9-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 42069// CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 42070// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42071// CHECK9-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42072// CHECK9-NEXT:    [[I:%.*]] = alloca i32, align 42073// CHECK9-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 42074// CHECK9-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 42075// CHECK9-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42076// CHECK9-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42077// CHECK9-NEXT:    [[I4:%.*]] = alloca i32, align 42078// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82079// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82080// CHECK9-NEXT:    store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 82081// CHECK9-NEXT:    store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 82082// CHECK9-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 82083// CHECK9-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 42084// CHECK9-NEXT:    store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 42085// CHECK9-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42086// CHECK9-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP1]], 02087// CHECK9-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 12088// CHECK9-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12089// CHECK9-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42090// CHECK9-NEXT:    store i32 0, ptr [[I]], align 42091// CHECK9-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42092// CHECK9-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP2]]2093// CHECK9-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2094// CHECK9:       omp.precond.then:2095// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 42096// CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42097// CHECK9-NEXT:    store i32 [[TMP3]], ptr [[DOTOMP_UB]], align 42098// CHECK9-NEXT:    [[TMP4:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 82099// CHECK9-NEXT:    [[CONV:%.*]] = trunc i64 [[TMP4]] to i322100// CHECK9-NEXT:    [[TMP5:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 82101// CHECK9-NEXT:    [[CONV3:%.*]] = trunc i64 [[TMP5]] to i322102// CHECK9-NEXT:    store i32 [[CONV]], ptr [[DOTOMP_LB]], align 42103// CHECK9-NEXT:    store i32 [[CONV3]], ptr [[DOTOMP_UB]], align 42104// CHECK9-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 42105// CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 42106// CHECK9-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82107// CHECK9-NEXT:    [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 42108// CHECK9-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP7]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2109// CHECK9-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42110// CHECK9-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42111// CHECK9-NEXT:    [[CMP5:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]2112// CHECK9-NEXT:    br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2113// CHECK9:       cond.true:2114// CHECK9-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42115// CHECK9-NEXT:    br label [[COND_END:%.*]]2116// CHECK9:       cond.false:2117// CHECK9-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42118// CHECK9-NEXT:    br label [[COND_END]]2119// CHECK9:       cond.end:2120// CHECK9-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]2121// CHECK9-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 42122// CHECK9-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42123// CHECK9-NEXT:    store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 42124// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]2125// CHECK9:       omp.inner.for.cond:2126// CHECK9-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP14:![0-9]+]]2127// CHECK9-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP14]]2128// CHECK9-NEXT:    [[CMP6:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]2129// CHECK9-NEXT:    br i1 [[CMP6]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2130// CHECK9:       omp.inner.for.body:2131// CHECK9-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP14]]2132// CHECK9-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP15]], 12133// CHECK9-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]2134// CHECK9-NEXT:    store i32 [[ADD]], ptr [[I4]], align 4, !llvm.access.group [[ACC_GRP14]]2135// CHECK9-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]2136// CHECK9:       omp.body.continue:2137// CHECK9-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]2138// CHECK9:       omp.inner.for.inc:2139// CHECK9-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP14]]2140// CHECK9-NEXT:    [[ADD7:%.*]] = add nsw i32 [[TMP16]], 12141// CHECK9-NEXT:    store i32 [[ADD7]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP14]]2142// CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP15:![0-9]+]]2143// CHECK9:       omp.inner.for.end:2144// CHECK9-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]2145// CHECK9:       omp.loop.exit:2146// CHECK9-NEXT:    [[TMP17:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82147// CHECK9-NEXT:    [[TMP18:%.*]] = load i32, ptr [[TMP17]], align 42148// CHECK9-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP18]])2149// CHECK9-NEXT:    [[TMP19:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 42150// CHECK9-NEXT:    [[TMP20:%.*]] = icmp ne i32 [[TMP19]], 02151// CHECK9-NEXT:    br i1 [[TMP20]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]2152// CHECK9:       .omp.final.then:2153// CHECK9-NEXT:    [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42154// CHECK9-NEXT:    [[SUB8:%.*]] = sub nsw i32 [[TMP21]], 02155// CHECK9-NEXT:    [[DIV9:%.*]] = sdiv i32 [[SUB8]], 12156// CHECK9-NEXT:    [[MUL10:%.*]] = mul nsw i32 [[DIV9]], 12157// CHECK9-NEXT:    [[ADD11:%.*]] = add nsw i32 0, [[MUL10]]2158// CHECK9-NEXT:    store i32 [[ADD11]], ptr [[I4]], align 42159// CHECK9-NEXT:    br label [[DOTOMP_FINAL_DONE]]2160// CHECK9:       .omp.final.done:2161// CHECK9-NEXT:    br label [[OMP_PRECOND_END]]2162// CHECK9:       omp.precond.end:2163// CHECK9-NEXT:    ret void2164//2165//2166// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l1082167// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0]] {2168// CHECK9-NEXT:  entry:2169// CHECK9-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 82170// CHECK9-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 82171// CHECK9-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 82172// CHECK9-NEXT:    store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 82173// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108.omp_outlined)2174// CHECK9-NEXT:    ret void2175//2176//2177// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108.omp_outlined2178// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {2179// CHECK9-NEXT:  entry:2180// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82181// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82182// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82183// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82184// CHECK9-NEXT:    ret void2185//2186//2187// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l1212188// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[THIS:%.*]], i64 noundef [[B:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0]] {2189// CHECK9-NEXT:  entry:2190// CHECK9-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 82191// CHECK9-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 82192// CHECK9-NEXT:    [[B_ADDR:%.*]] = alloca i64, align 82193// CHECK9-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 82194// CHECK9-NEXT:    [[B_CASTED:%.*]] = alloca i64, align 82195// CHECK9-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 82196// CHECK9-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 82197// CHECK9-NEXT:    store i64 [[B]], ptr [[B_ADDR]], align 82198// CHECK9-NEXT:    store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 82199// CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 82200// CHECK9-NEXT:    [[TMP1:%.*]] = load i32, ptr [[B_ADDR]], align 42201// CHECK9-NEXT:    store i32 [[TMP1]], ptr [[B_CASTED]], align 42202// CHECK9-NEXT:    [[TMP2:%.*]] = load i64, ptr [[B_CASTED]], align 82203// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.omp_outlined, ptr [[TMP0]], i64 [[TMP2]])2204// CHECK9-NEXT:    ret void2205//2206//2207// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.omp_outlined2208// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]], i64 noundef [[B:%.*]]) #[[ATTR0]] {2209// CHECK9-NEXT:  entry:2210// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82211// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82212// CHECK9-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 82213// CHECK9-NEXT:    [[B_ADDR:%.*]] = alloca i64, align 82214// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82215// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82216// CHECK9-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 82217// CHECK9-NEXT:    store i64 [[B]], ptr [[B_ADDR]], align 82218// CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 82219// CHECK9-NEXT:    [[TMP1:%.*]] = load i32, ptr [[B_ADDR]], align 42220// CHECK9-NEXT:    [[CONV:%.*]] = sitofp i32 [[TMP1]] to double2221// CHECK9-NEXT:    [[ADD:%.*]] = fadd double [[CONV]], 1.500000e+002222// CHECK9-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 02223// CHECK9-NEXT:    store double [[ADD]], ptr [[A]], align 82224// CHECK9-NEXT:    ret void2225//2226//2227// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l1262228// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR0]] {2229// CHECK9-NEXT:  entry:2230// CHECK9-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 82231// CHECK9-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 82232// CHECK9-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 82233// CHECK9-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 82234// CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 82235// CHECK9-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 02236// CHECK9-NEXT:    store double 2.500000e+00, ptr [[A]], align 82237// CHECK9-NEXT:    ret void2238//2239//2240// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l882241// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0]] {2242// CHECK9-NEXT:  entry:2243// CHECK9-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 82244// CHECK9-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 82245// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l88.omp_outlined)2246// CHECK9-NEXT:    ret void2247//2248//2249// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l88.omp_outlined2250// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {2251// CHECK9-NEXT:  entry:2252// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82253// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82254// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82255// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82256// CHECK9-NEXT:    ret void2257//2258//2259// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l932260// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[A:%.*]], i64 noundef [[B:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0]] {2261// CHECK9-NEXT:  entry:2262// CHECK9-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 82263// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca i64, align 82264// CHECK9-NEXT:    [[B_ADDR:%.*]] = alloca i64, align 82265// CHECK9-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 82266// CHECK9-NEXT:    [[A_CASTED:%.*]] = alloca i64, align 82267// CHECK9-NEXT:    [[B_CASTED:%.*]] = alloca i64, align 82268// CHECK9-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB3]])2269// CHECK9-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 82270// CHECK9-NEXT:    store i64 [[A]], ptr [[A_ADDR]], align 82271// CHECK9-NEXT:    store i64 [[B]], ptr [[B_ADDR]], align 82272// CHECK9-NEXT:    store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 82273// CHECK9-NEXT:    [[TMP1:%.*]] = load i16, ptr [[DOTCAPTURE_EXPR__ADDR]], align 22274// CHECK9-NEXT:    [[TMP2:%.*]] = sext i16 [[TMP1]] to i322275// CHECK9-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB3]], i32 [[TMP0]], i32 [[TMP2]], i32 0)2276// CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A_ADDR]], align 42277// CHECK9-NEXT:    store i32 [[TMP3]], ptr [[A_CASTED]], align 42278// CHECK9-NEXT:    [[TMP4:%.*]] = load i64, ptr [[A_CASTED]], align 82279// CHECK9-NEXT:    [[TMP5:%.*]] = load i16, ptr [[B_ADDR]], align 22280// CHECK9-NEXT:    store i16 [[TMP5]], ptr [[B_CASTED]], align 22281// CHECK9-NEXT:    [[TMP6:%.*]] = load i64, ptr [[B_CASTED]], align 82282// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93.omp_outlined, i64 [[TMP4]], i64 [[TMP6]])2283// CHECK9-NEXT:    ret void2284//2285//2286// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93.omp_outlined2287// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[A:%.*]], i64 noundef [[B:%.*]]) #[[ATTR0]] {2288// CHECK9-NEXT:  entry:2289// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82290// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82291// CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca i64, align 82292// CHECK9-NEXT:    [[B_ADDR:%.*]] = alloca i64, align 82293// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82294// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82295// CHECK9-NEXT:    store i64 [[A]], ptr [[A_ADDR]], align 82296// CHECK9-NEXT:    store i64 [[B]], ptr [[B_ADDR]], align 82297// CHECK9-NEXT:    [[TMP0:%.*]] = load i16, ptr [[B_ADDR]], align 22298// CHECK9-NEXT:    [[CONV:%.*]] = sext i16 [[TMP0]] to i322299// CHECK9-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 42300// CHECK9-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CONV]]2301// CHECK9-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 42302// CHECK9-NEXT:    ret void2303//2304//2305// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l1042306// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[N:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]], i32 noundef [[DOTCAPTURE_EXPR_1:%.*]]) #[[ATTR0:[0-9]+]] {2307// CHECK11-NEXT:  entry:2308// CHECK11-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42309// CHECK11-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 42310// CHECK11-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 42311// CHECK11-NEXT:    [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca i32, align 42312// CHECK11-NEXT:    [[N_CASTED:%.*]] = alloca i32, align 42313// CHECK11-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB3:[0-9]+]])2314// CHECK11-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42315// CHECK11-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 42316// CHECK11-NEXT:    store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 42317// CHECK11-NEXT:    store i32 [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 42318// CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 42319// CHECK11-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB3]], i32 [[TMP0]], i32 [[TMP1]], i32 0)2320// CHECK11-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 42321// CHECK11-NEXT:    store i32 [[TMP2]], ptr [[N_CASTED]], align 42322// CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[N_CASTED]], align 42323// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined, i32 [[TMP3]])2324// CHECK11-NEXT:    ret void2325//2326//2327// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined2328// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[N:%.*]]) #[[ATTR0]] {2329// CHECK11-NEXT:  entry:2330// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42331// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42332// CHECK11-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 42333// CHECK11-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 42334// CHECK11-NEXT:    [[TMP:%.*]] = alloca i32, align 42335// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42336// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42337// CHECK11-NEXT:    [[I:%.*]] = alloca i32, align 42338// CHECK11-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 42339// CHECK11-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 42340// CHECK11-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42341// CHECK11-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42342// CHECK11-NEXT:    [[I3:%.*]] = alloca i32, align 42343// CHECK11-NEXT:    [[N_CASTED:%.*]] = alloca i32, align 42344// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42345// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42346// CHECK11-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 42347// CHECK11-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 42348// CHECK11-NEXT:    store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 42349// CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42350// CHECK11-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP1]], 02351// CHECK11-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 12352// CHECK11-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12353// CHECK11-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42354// CHECK11-NEXT:    store i32 0, ptr [[I]], align 42355// CHECK11-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42356// CHECK11-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP2]]2357// CHECK11-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2358// CHECK11:       omp.precond.then:2359// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 42360// CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42361// CHECK11-NEXT:    store i32 [[TMP3]], ptr [[DOTOMP_COMB_UB]], align 42362// CHECK11-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 42363// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 42364// CHECK11-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42365// CHECK11-NEXT:    [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 42366// CHECK11-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP5]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2367// CHECK11-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42368// CHECK11-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42369// CHECK11-NEXT:    [[CMP4:%.*]] = icmp sgt i32 [[TMP6]], [[TMP7]]2370// CHECK11-NEXT:    br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2371// CHECK11:       cond.true:2372// CHECK11-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42373// CHECK11-NEXT:    br label [[COND_END:%.*]]2374// CHECK11:       cond.false:2375// CHECK11-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42376// CHECK11-NEXT:    br label [[COND_END]]2377// CHECK11:       cond.end:2378// CHECK11-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP8]], [[COND_TRUE]] ], [ [[TMP9]], [[COND_FALSE]] ]2379// CHECK11-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 42380// CHECK11-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 42381// CHECK11-NEXT:    store i32 [[TMP10]], ptr [[DOTOMP_IV]], align 42382// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]2383// CHECK11:       omp.inner.for.cond:2384// CHECK11-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11:![0-9]+]]2385// CHECK11-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4, !llvm.access.group [[ACC_GRP11]]2386// CHECK11-NEXT:    [[CMP5:%.*]] = icmp sle i32 [[TMP11]], [[TMP12]]2387// CHECK11-NEXT:    br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2388// CHECK11:       omp.inner.for.body:2389// CHECK11-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4, !llvm.access.group [[ACC_GRP11]]2390// CHECK11-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4, !llvm.access.group [[ACC_GRP11]]2391// CHECK11-NEXT:    [[TMP15:%.*]] = load i32, ptr [[N_ADDR]], align 4, !llvm.access.group [[ACC_GRP11]]2392// CHECK11-NEXT:    store i32 [[TMP15]], ptr [[N_CASTED]], align 4, !llvm.access.group [[ACC_GRP11]]2393// CHECK11-NEXT:    [[TMP16:%.*]] = load i32, ptr [[N_CASTED]], align 4, !llvm.access.group [[ACC_GRP11]]2394// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined.omp_outlined, i32 [[TMP13]], i32 [[TMP14]], i32 [[TMP16]]), !llvm.access.group [[ACC_GRP11]]2395// CHECK11-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]2396// CHECK11:       omp.inner.for.inc:2397// CHECK11-NEXT:    [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]2398// CHECK11-NEXT:    [[TMP18:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4, !llvm.access.group [[ACC_GRP11]]2399// CHECK11-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP17]], [[TMP18]]2400// CHECK11-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]2401// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP12:![0-9]+]]2402// CHECK11:       omp.inner.for.end:2403// CHECK11-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]2404// CHECK11:       omp.loop.exit:2405// CHECK11-NEXT:    [[TMP19:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42406// CHECK11-NEXT:    [[TMP20:%.*]] = load i32, ptr [[TMP19]], align 42407// CHECK11-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP20]])2408// CHECK11-NEXT:    [[TMP21:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 42409// CHECK11-NEXT:    [[TMP22:%.*]] = icmp ne i32 [[TMP21]], 02410// CHECK11-NEXT:    br i1 [[TMP22]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]2411// CHECK11:       .omp.final.then:2412// CHECK11-NEXT:    [[TMP23:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42413// CHECK11-NEXT:    [[SUB6:%.*]] = sub nsw i32 [[TMP23]], 02414// CHECK11-NEXT:    [[DIV7:%.*]] = sdiv i32 [[SUB6]], 12415// CHECK11-NEXT:    [[MUL:%.*]] = mul nsw i32 [[DIV7]], 12416// CHECK11-NEXT:    [[ADD8:%.*]] = add nsw i32 0, [[MUL]]2417// CHECK11-NEXT:    store i32 [[ADD8]], ptr [[I3]], align 42418// CHECK11-NEXT:    br label [[DOTOMP_FINAL_DONE]]2419// CHECK11:       .omp.final.done:2420// CHECK11-NEXT:    br label [[OMP_PRECOND_END]]2421// CHECK11:       omp.precond.end:2422// CHECK11-NEXT:    ret void2423//2424//2425// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined.omp_outlined2426// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[DOTPREVIOUS_LB_:%.*]], i32 noundef [[DOTPREVIOUS_UB_:%.*]], i32 noundef [[N:%.*]]) #[[ATTR0]] {2427// CHECK11-NEXT:  entry:2428// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42429// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42430// CHECK11-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i32, align 42431// CHECK11-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i32, align 42432// CHECK11-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 42433// CHECK11-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 42434// CHECK11-NEXT:    [[TMP:%.*]] = alloca i32, align 42435// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42436// CHECK11-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42437// CHECK11-NEXT:    [[I:%.*]] = alloca i32, align 42438// CHECK11-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 42439// CHECK11-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 42440// CHECK11-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42441// CHECK11-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42442// CHECK11-NEXT:    [[I3:%.*]] = alloca i32, align 42443// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42444// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42445// CHECK11-NEXT:    store i32 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 42446// CHECK11-NEXT:    store i32 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 42447// CHECK11-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 42448// CHECK11-NEXT:    [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 42449// CHECK11-NEXT:    store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 42450// CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42451// CHECK11-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP1]], 02452// CHECK11-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 12453// CHECK11-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12454// CHECK11-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42455// CHECK11-NEXT:    store i32 0, ptr [[I]], align 42456// CHECK11-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42457// CHECK11-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP2]]2458// CHECK11-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2459// CHECK11:       omp.precond.then:2460// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 42461// CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42462// CHECK11-NEXT:    store i32 [[TMP3]], ptr [[DOTOMP_UB]], align 42463// CHECK11-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTPREVIOUS_LB__ADDR]], align 42464// CHECK11-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTPREVIOUS_UB__ADDR]], align 42465// CHECK11-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_LB]], align 42466// CHECK11-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_UB]], align 42467// CHECK11-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 42468// CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 42469// CHECK11-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42470// CHECK11-NEXT:    [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 42471// CHECK11-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP7]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2472// CHECK11-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42473// CHECK11-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42474// CHECK11-NEXT:    [[CMP4:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]2475// CHECK11-NEXT:    br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2476// CHECK11:       cond.true:2477// CHECK11-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42478// CHECK11-NEXT:    br label [[COND_END:%.*]]2479// CHECK11:       cond.false:2480// CHECK11-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42481// CHECK11-NEXT:    br label [[COND_END]]2482// CHECK11:       cond.end:2483// CHECK11-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]2484// CHECK11-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 42485// CHECK11-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42486// CHECK11-NEXT:    store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 42487// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]2488// CHECK11:       omp.inner.for.cond:2489// CHECK11-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP15:![0-9]+]]2490// CHECK11-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP15]]2491// CHECK11-NEXT:    [[CMP5:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]2492// CHECK11-NEXT:    br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2493// CHECK11:       omp.inner.for.body:2494// CHECK11-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP15]]2495// CHECK11-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP15]], 12496// CHECK11-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]2497// CHECK11-NEXT:    store i32 [[ADD]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP15]]2498// CHECK11-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]2499// CHECK11:       omp.body.continue:2500// CHECK11-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]2501// CHECK11:       omp.inner.for.inc:2502// CHECK11-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP15]]2503// CHECK11-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP16]], 12504// CHECK11-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP15]]2505// CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP16:![0-9]+]]2506// CHECK11:       omp.inner.for.end:2507// CHECK11-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]2508// CHECK11:       omp.loop.exit:2509// CHECK11-NEXT:    [[TMP17:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42510// CHECK11-NEXT:    [[TMP18:%.*]] = load i32, ptr [[TMP17]], align 42511// CHECK11-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP18]])2512// CHECK11-NEXT:    [[TMP19:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 42513// CHECK11-NEXT:    [[TMP20:%.*]] = icmp ne i32 [[TMP19]], 02514// CHECK11-NEXT:    br i1 [[TMP20]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]2515// CHECK11:       .omp.final.then:2516// CHECK11-NEXT:    [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42517// CHECK11-NEXT:    [[SUB7:%.*]] = sub nsw i32 [[TMP21]], 02518// CHECK11-NEXT:    [[DIV8:%.*]] = sdiv i32 [[SUB7]], 12519// CHECK11-NEXT:    [[MUL9:%.*]] = mul nsw i32 [[DIV8]], 12520// CHECK11-NEXT:    [[ADD10:%.*]] = add nsw i32 0, [[MUL9]]2521// CHECK11-NEXT:    store i32 [[ADD10]], ptr [[I3]], align 42522// CHECK11-NEXT:    br label [[DOTOMP_FINAL_DONE]]2523// CHECK11:       .omp.final.done:2524// CHECK11-NEXT:    br label [[OMP_PRECOND_END]]2525// CHECK11:       omp.precond.end:2526// CHECK11-NEXT:    ret void2527//2528//2529// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l1082530// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0]] {2531// CHECK11-NEXT:  entry:2532// CHECK11-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42533// CHECK11-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 42534// CHECK11-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42535// CHECK11-NEXT:    store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 42536// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108.omp_outlined)2537// CHECK11-NEXT:    ret void2538//2539//2540// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108.omp_outlined2541// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {2542// CHECK11-NEXT:  entry:2543// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42544// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42545// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42546// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42547// CHECK11-NEXT:    ret void2548//2549//2550// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l1212551// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[THIS:%.*]], i32 noundef [[B:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0]] {2552// CHECK11-NEXT:  entry:2553// CHECK11-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42554// CHECK11-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 42555// CHECK11-NEXT:    [[B_ADDR:%.*]] = alloca i32, align 42556// CHECK11-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 42557// CHECK11-NEXT:    [[B_CASTED:%.*]] = alloca i32, align 42558// CHECK11-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42559// CHECK11-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 42560// CHECK11-NEXT:    store i32 [[B]], ptr [[B_ADDR]], align 42561// CHECK11-NEXT:    store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 42562// CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 42563// CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[B_ADDR]], align 42564// CHECK11-NEXT:    store i32 [[TMP1]], ptr [[B_CASTED]], align 42565// CHECK11-NEXT:    [[TMP2:%.*]] = load i32, ptr [[B_CASTED]], align 42566// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.omp_outlined, ptr [[TMP0]], i32 [[TMP2]])2567// CHECK11-NEXT:    ret void2568//2569//2570// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.omp_outlined2571// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]], i32 noundef [[B:%.*]]) #[[ATTR0]] {2572// CHECK11-NEXT:  entry:2573// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42574// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42575// CHECK11-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 42576// CHECK11-NEXT:    [[B_ADDR:%.*]] = alloca i32, align 42577// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42578// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42579// CHECK11-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 42580// CHECK11-NEXT:    store i32 [[B]], ptr [[B_ADDR]], align 42581// CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 42582// CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[B_ADDR]], align 42583// CHECK11-NEXT:    [[CONV:%.*]] = sitofp i32 [[TMP1]] to double2584// CHECK11-NEXT:    [[ADD:%.*]] = fadd double [[CONV]], 1.500000e+002585// CHECK11-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 02586// CHECK11-NEXT:    store double [[ADD]], ptr [[A]], align 42587// CHECK11-NEXT:    ret void2588//2589//2590// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l1262591// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR0]] {2592// CHECK11-NEXT:  entry:2593// CHECK11-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42594// CHECK11-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 42595// CHECK11-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42596// CHECK11-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 42597// CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 42598// CHECK11-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 02599// CHECK11-NEXT:    store double 2.500000e+00, ptr [[A]], align 42600// CHECK11-NEXT:    ret void2601//2602//2603// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l882604// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0]] {2605// CHECK11-NEXT:  entry:2606// CHECK11-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42607// CHECK11-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42608// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l88.omp_outlined)2609// CHECK11-NEXT:    ret void2610//2611//2612// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l88.omp_outlined2613// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {2614// CHECK11-NEXT:  entry:2615// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42616// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42617// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42618// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42619// CHECK11-NEXT:    ret void2620//2621//2622// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l932623// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[A:%.*]], i32 noundef [[B:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0]] {2624// CHECK11-NEXT:  entry:2625// CHECK11-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42626// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 42627// CHECK11-NEXT:    [[B_ADDR:%.*]] = alloca i32, align 42628// CHECK11-NEXT:    [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 42629// CHECK11-NEXT:    [[A_CASTED:%.*]] = alloca i32, align 42630// CHECK11-NEXT:    [[B_CASTED:%.*]] = alloca i32, align 42631// CHECK11-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB3]])2632// CHECK11-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42633// CHECK11-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 42634// CHECK11-NEXT:    store i32 [[B]], ptr [[B_ADDR]], align 42635// CHECK11-NEXT:    store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 42636// CHECK11-NEXT:    [[TMP1:%.*]] = load i16, ptr [[DOTCAPTURE_EXPR__ADDR]], align 22637// CHECK11-NEXT:    [[TMP2:%.*]] = sext i16 [[TMP1]] to i322638// CHECK11-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB3]], i32 [[TMP0]], i32 [[TMP2]], i32 0)2639// CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A_ADDR]], align 42640// CHECK11-NEXT:    store i32 [[TMP3]], ptr [[A_CASTED]], align 42641// CHECK11-NEXT:    [[TMP4:%.*]] = load i32, ptr [[A_CASTED]], align 42642// CHECK11-NEXT:    [[TMP5:%.*]] = load i16, ptr [[B_ADDR]], align 22643// CHECK11-NEXT:    store i16 [[TMP5]], ptr [[B_CASTED]], align 22644// CHECK11-NEXT:    [[TMP6:%.*]] = load i32, ptr [[B_CASTED]], align 42645// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93.omp_outlined, i32 [[TMP4]], i32 [[TMP6]])2646// CHECK11-NEXT:    ret void2647//2648//2649// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93.omp_outlined2650// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[A:%.*]], i32 noundef [[B:%.*]]) #[[ATTR0]] {2651// CHECK11-NEXT:  entry:2652// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42653// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42654// CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 42655// CHECK11-NEXT:    [[B_ADDR:%.*]] = alloca i32, align 42656// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42657// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42658// CHECK11-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 42659// CHECK11-NEXT:    store i32 [[B]], ptr [[B_ADDR]], align 42660// CHECK11-NEXT:    [[TMP0:%.*]] = load i16, ptr [[B_ADDR]], align 22661// CHECK11-NEXT:    [[CONV:%.*]] = sext i16 [[TMP0]] to i322662// CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 42663// CHECK11-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CONV]]2664// CHECK11-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 42665// CHECK11-NEXT:    ret void2666//2667