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// expected-no-diagnostics3#ifndef HEADER4#define HEADER5// Test host codegen.6// RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK17// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s8// RUN: %clang_cc1 -DCK1 -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=CHECK19// RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK310// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s11// RUN: %clang_cc1 -DCK1 -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=CHECK312 13// RUN: %clang_cc1 -DCK1 -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}}"14// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s15// RUN: %clang_cc1 -DCK1 -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}}"16// RUN: %clang_cc1 -DCK1 -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}}"17// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s18// RUN: %clang_cc1 -DCK1 -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}}"19#ifdef CK120 21int Gbla;22long long Gblb;23int &Gblc = Gbla;24 25int teams_argument_global_local(int a){26  int comp = 1;27 28  int la = 23;29  float lc = 25.0;30 31  #pragma omp target32  #pragma omp teams33  {34    ++comp;35  }36 37  #pragma omp target38  {{{39    #pragma omp teams40    {41      ++comp;42    }43  }}}44 45 46  #pragma omp target47  #pragma omp teams num_teams(la)48  {49    ++comp;50  }51 52 53  #pragma omp target54  #pragma omp teams thread_limit(la)55  {56    ++comp;57  }58 59 60 61 62  #pragma omp target63  #pragma omp teams num_teams(Gbla+a) thread_limit(Gblb+(long long)lc)64  {65    ++comp;66  }67 68 69 70 71  #pragma omp target72  #pragma omp teams num_teams(Gblc+1) thread_limit(Gblc+2)73  {74    comp += Gblc;75  }76 77  return comp;78}79 80#endif // CK181 82// Test host codegen.83// RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK984// RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s85// RUN: %clang_cc1 -DCK2 -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=CHECK986// RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK1187// RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s88// RUN: %clang_cc1 -DCK2 -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=CHECK1189 90// RUN: %clang_cc1 -DCK2 -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}}"91// RUN: %clang_cc1 -DCK2 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s92// RUN: %clang_cc1 -DCK2 -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}}"93// RUN: %clang_cc1 -DCK2 -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}}"94// RUN: %clang_cc1 -DCK2 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s95// RUN: %clang_cc1 -DCK2 -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}}"96#ifdef CK297 98template <typename T>99struct SS{100  T a;101  float b;102};103 104SS<int> Gbla;105SS<long long> Gblb;106 107int teams_template_arg(void) {108  int comp = 1;109 110  SS<int> la;111  SS<long long> lb;112 113 114 115 116  #pragma omp target117  #pragma omp teams num_teams(Gbla.a) thread_limit((long long)la.b)118  {119    ++comp;120  }121 122 123 124 125  #pragma omp target126  #pragma omp teams num_teams((long long)lb.b) thread_limit(Gblb.a)127  {128    ++comp;129  }130  return comp;131}132#endif // CK2133 134// Test host codegen.135// RUN: %clang_cc1 -DCK3 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK17136// RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s137// RUN: %clang_cc1 -DCK3 -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=CHECK17138// RUN: %clang_cc1 -DCK3 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK19139// RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s140// RUN: %clang_cc1 -DCK3 -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=CHECK19141 142// RUN: %clang_cc1 -DCK3 -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}}"143// RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s144// RUN: %clang_cc1 -DCK3 -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}}"145// RUN: %clang_cc1 -DCK3 -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}}"146// RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s147// RUN: %clang_cc1 -DCK3 -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}}"148#ifdef CK3149 150 151template <typename T, int X, long long Y>152struct SS{153  T a;154  float b;155 156  int foo(void) {157    int comp = 1;158 159 160 161    #pragma omp target162    #pragma omp teams num_teams(a) thread_limit(X)163    {164      ++comp;165    }166 167 168 169    #pragma omp target170    #pragma omp teams num_teams(Y) thread_limit((int)b+X)171    {172      ++comp;173    }174    return comp;175  }176};177 178int teams_template_struct(void) {179  SS<int, 123, 456> V;180  return V.foo();181 182}183#endif // CK3184 185// Test target codegen - host bc file has to be created first.186// RUN: %clang_cc1 -DCK4 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc187// RUN: %clang_cc1 -DCK4 -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=CHECK25188// RUN: %clang_cc1 -DCK4 -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 %s189// RUN: %clang_cc1 -DCK4 -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=CHECK25190// RUN: %clang_cc1 -DCK4 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc191// RUN: %clang_cc1 -DCK4 -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=CHECK27192// RUN: %clang_cc1 -DCK4 -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 %s193// RUN: %clang_cc1 -DCK4 -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=CHECK27194 195// RUN: %clang_cc1 -DCK4 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc196// RUN: %clang_cc1 -DCK4 -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}}"197// RUN: %clang_cc1 -DCK4 -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 %s198// RUN: %clang_cc1 -DCK4 -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}}"199// RUN: %clang_cc1 -DCK4 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc200// RUN: %clang_cc1 -DCK4 -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}}"201// RUN: %clang_cc1 -DCK4 -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 %s202// RUN: %clang_cc1 -DCK4 -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}}"203 204#ifdef CK4205 206 207template <typename T>208int tmain(T argc) {209#pragma omp target210#pragma omp teams211  argc = 0;212  return 0;213}214 215int main (int argc, char **argv) {216#pragma omp target217#pragma omp teams218  argc = 0;219  return tmain(argv);220}221 222 223 224 225#endif // CK4226 227// Test target codegen - host bc file has to be created first.228// RUN: %clang_cc1 -DCK5 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc229// RUN: %clang_cc1 -DCK5 -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=CHECK33230// RUN: %clang_cc1 -DCK5 -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 %s231// RUN: %clang_cc1 -DCK5 -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=CHECK33232// RUN: %clang_cc1 -DCK5 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc233// RUN: %clang_cc1 -DCK5 -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=CHECK35234// RUN: %clang_cc1 -DCK5 -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 %s235// RUN: %clang_cc1 -DCK5 -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=CHECK35236 237// RUN: %clang_cc1 -DCK5 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc238// RUN: %clang_cc1 -DCK5 -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}}"239// RUN: %clang_cc1 -DCK5 -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 %s240// RUN: %clang_cc1 -DCK5 -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}}"241// RUN: %clang_cc1 -DCK5 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc242// RUN: %clang_cc1 -DCK5 -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}}"243// RUN: %clang_cc1 -DCK5 -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 %s244// RUN: %clang_cc1 -DCK5 -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}}"245 246// expected-no-diagnostics247#ifdef CK5248 249 250template <typename T>251int tmain(T argc) {252  int a = 10;253  int b = 5;254#pragma omp target255#pragma omp teams num_teams(a) thread_limit(b)256  {257  argc = 0;258  }259  return 0;260}261 262int main (int argc, char **argv) {263  int a = 20;264  int b = 5;265#pragma omp target266#pragma omp teams num_teams(a) thread_limit(b)267  {268  argc = 0;269  }270  return tmain(argv);271}272 273 274 275#endif // CK5276 277// Test host codegen.278// RUN: %clang_cc1 -DCK6 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK41279// RUN: %clang_cc1 -DCK6 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s280// RUN: %clang_cc1 -DCK6 -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=CHECK41281// RUN: %clang_cc1 -DCK6 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK43282// RUN: %clang_cc1 -DCK6 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s283// RUN: %clang_cc1 -DCK6 -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=CHECK43284 285// RUN: %clang_cc1 -DCK6 -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}}"286// RUN: %clang_cc1 -DCK6 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s287// RUN: %clang_cc1 -DCK6 -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}}"288// RUN: %clang_cc1 -DCK6 -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}}"289// RUN: %clang_cc1 -DCK6 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s290// RUN: %clang_cc1 -DCK6 -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}}"291#ifdef CK6292 293void foo() {294#pragma omp teams295  ;296}297 298#endif // CK6299 300#endif301// CHECK1-LABEL: define {{[^@]+}}@_Z27teams_argument_global_locali302// CHECK1-SAME: (i32 noundef signext [[A:%.*]]) #[[ATTR0:[0-9]+]] {303// CHECK1-NEXT:  entry:304// CHECK1-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 4305// CHECK1-NEXT:    [[COMP:%.*]] = alloca i32, align 4306// CHECK1-NEXT:    [[LA:%.*]] = alloca i32, align 4307// CHECK1-NEXT:    [[LC:%.*]] = alloca float, align 4308// CHECK1-NEXT:    [[COMP_CASTED:%.*]] = alloca i64, align 8309// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8310// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8311// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8312// CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8313// CHECK1-NEXT:    [[COMP_CASTED1:%.*]] = alloca i64, align 8314// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS2:%.*]] = alloca [1 x ptr], align 8315// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS3:%.*]] = alloca [1 x ptr], align 8316// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS4:%.*]] = alloca [1 x ptr], align 8317// CHECK1-NEXT:    [[KERNEL_ARGS5:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8318// CHECK1-NEXT:    [[LA_CASTED:%.*]] = alloca i64, align 8319// CHECK1-NEXT:    [[COMP_CASTED8:%.*]] = alloca i64, align 8320// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS9:%.*]] = alloca [2 x ptr], align 8321// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS10:%.*]] = alloca [2 x ptr], align 8322// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS11:%.*]] = alloca [2 x ptr], align 8323// CHECK1-NEXT:    [[KERNEL_ARGS12:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8324// CHECK1-NEXT:    [[LA_CASTED15:%.*]] = alloca i64, align 8325// CHECK1-NEXT:    [[COMP_CASTED16:%.*]] = alloca i64, align 8326// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS17:%.*]] = alloca [2 x ptr], align 8327// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS18:%.*]] = alloca [2 x ptr], align 8328// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS19:%.*]] = alloca [2 x ptr], align 8329// CHECK1-NEXT:    [[KERNEL_ARGS20:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8330// CHECK1-NEXT:    [[GBLA_CASTED:%.*]] = alloca i64, align 8331// CHECK1-NEXT:    [[A_CASTED:%.*]] = alloca i64, align 8332// CHECK1-NEXT:    [[GBLB_CASTED:%.*]] = alloca i64, align 8333// CHECK1-NEXT:    [[LC_CASTED:%.*]] = alloca i64, align 8334// CHECK1-NEXT:    [[COMP_CASTED23:%.*]] = alloca i64, align 8335// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS24:%.*]] = alloca [5 x ptr], align 8336// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS25:%.*]] = alloca [5 x ptr], align 8337// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS26:%.*]] = alloca [5 x ptr], align 8338// CHECK1-NEXT:    [[KERNEL_ARGS28:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8339// CHECK1-NEXT:    [[TMP:%.*]] = alloca ptr, align 8340// CHECK1-NEXT:    [[GBLC_CASTED:%.*]] = alloca i64, align 8341// CHECK1-NEXT:    [[COMP_CASTED31:%.*]] = alloca i64, align 8342// CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS32:%.*]] = alloca [2 x ptr], align 8343// CHECK1-NEXT:    [[DOTOFFLOAD_PTRS33:%.*]] = alloca [2 x ptr], align 8344// CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS34:%.*]] = alloca [2 x ptr], align 8345// CHECK1-NEXT:    [[_TMP35:%.*]] = alloca ptr, align 8346// CHECK1-NEXT:    [[_TMP37:%.*]] = alloca ptr, align 8347// CHECK1-NEXT:    [[KERNEL_ARGS39:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8348// CHECK1-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 4349// CHECK1-NEXT:    store i32 1, ptr [[COMP]], align 4350// CHECK1-NEXT:    store i32 23, ptr [[LA]], align 4351// CHECK1-NEXT:    store float 2.500000e+01, ptr [[LC]], align 4352// CHECK1-NEXT:    [[TMP0:%.*]] = load i32, ptr [[COMP]], align 4353// CHECK1-NEXT:    store i32 [[TMP0]], ptr [[COMP_CASTED]], align 4354// CHECK1-NEXT:    [[TMP1:%.*]] = load i64, ptr [[COMP_CASTED]], align 8355// CHECK1-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0356// CHECK1-NEXT:    store i64 [[TMP1]], ptr [[TMP2]], align 8357// CHECK1-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0358// CHECK1-NEXT:    store i64 [[TMP1]], ptr [[TMP3]], align 8359// CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0360// CHECK1-NEXT:    store ptr null, ptr [[TMP4]], align 8361// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0362// CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0363// CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0364// CHECK1-NEXT:    store i32 3, ptr [[TMP7]], align 4365// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1366// CHECK1-NEXT:    store i32 1, ptr [[TMP8]], align 4367// CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2368// CHECK1-NEXT:    store ptr [[TMP5]], ptr [[TMP9]], align 8369// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3370// CHECK1-NEXT:    store ptr [[TMP6]], ptr [[TMP10]], align 8371// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4372// CHECK1-NEXT:    store ptr @.offload_sizes, ptr [[TMP11]], align 8373// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5374// CHECK1-NEXT:    store ptr @.offload_maptypes, ptr [[TMP12]], align 8375// CHECK1-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6376// CHECK1-NEXT:    store ptr null, ptr [[TMP13]], align 8377// CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7378// CHECK1-NEXT:    store ptr null, ptr [[TMP14]], align 8379// CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8380// CHECK1-NEXT:    store i64 0, ptr [[TMP15]], align 8381// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9382// CHECK1-NEXT:    store i64 0, ptr [[TMP16]], align 8383// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10384// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4385// CHECK1-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11386// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP18]], align 4387// CHECK1-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12388// CHECK1-NEXT:    store i32 0, ptr [[TMP19]], align 4389// CHECK1-NEXT:    [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31.region_id, ptr [[KERNEL_ARGS]])390// CHECK1-NEXT:    [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0391// CHECK1-NEXT:    br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]392// CHECK1:       omp_offload.failed:393// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31(i64 [[TMP1]]) #[[ATTR2:[0-9]+]]394// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]395// CHECK1:       omp_offload.cont:396// CHECK1-NEXT:    [[TMP22:%.*]] = load i32, ptr [[COMP]], align 4397// CHECK1-NEXT:    store i32 [[TMP22]], ptr [[COMP_CASTED1]], align 4398// CHECK1-NEXT:    [[TMP23:%.*]] = load i64, ptr [[COMP_CASTED1]], align 8399// CHECK1-NEXT:    [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0400// CHECK1-NEXT:    store i64 [[TMP23]], ptr [[TMP24]], align 8401// CHECK1-NEXT:    [[TMP25:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 0402// CHECK1-NEXT:    store i64 [[TMP23]], ptr [[TMP25]], align 8403// CHECK1-NEXT:    [[TMP26:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i64 0, i64 0404// CHECK1-NEXT:    store ptr null, ptr [[TMP26]], align 8405// CHECK1-NEXT:    [[TMP27:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0406// CHECK1-NEXT:    [[TMP28:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 0407// CHECK1-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 0408// CHECK1-NEXT:    store i32 3, ptr [[TMP29]], align 4409// CHECK1-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 1410// CHECK1-NEXT:    store i32 1, ptr [[TMP30]], align 4411// CHECK1-NEXT:    [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 2412// CHECK1-NEXT:    store ptr [[TMP27]], ptr [[TMP31]], align 8413// CHECK1-NEXT:    [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 3414// CHECK1-NEXT:    store ptr [[TMP28]], ptr [[TMP32]], align 8415// CHECK1-NEXT:    [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 4416// CHECK1-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP33]], align 8417// CHECK1-NEXT:    [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 5418// CHECK1-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP34]], align 8419// CHECK1-NEXT:    [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 6420// CHECK1-NEXT:    store ptr null, ptr [[TMP35]], align 8421// CHECK1-NEXT:    [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 7422// CHECK1-NEXT:    store ptr null, ptr [[TMP36]], align 8423// CHECK1-NEXT:    [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 8424// CHECK1-NEXT:    store i64 0, ptr [[TMP37]], align 8425// CHECK1-NEXT:    [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 9426// CHECK1-NEXT:    store i64 0, ptr [[TMP38]], align 8427// CHECK1-NEXT:    [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 10428// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP39]], align 4429// CHECK1-NEXT:    [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 11430// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP40]], align 4431// CHECK1-NEXT:    [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 12432// CHECK1-NEXT:    store i32 0, ptr [[TMP41]], align 4433// CHECK1-NEXT:    [[TMP42:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37.region_id, ptr [[KERNEL_ARGS5]])434// CHECK1-NEXT:    [[TMP43:%.*]] = icmp ne i32 [[TMP42]], 0435// CHECK1-NEXT:    br i1 [[TMP43]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]]436// CHECK1:       omp_offload.failed6:437// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37(i64 [[TMP23]]) #[[ATTR2]]438// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT7]]439// CHECK1:       omp_offload.cont7:440// CHECK1-NEXT:    [[TMP44:%.*]] = load i32, ptr [[LA]], align 4441// CHECK1-NEXT:    store i32 [[TMP44]], ptr [[LA_CASTED]], align 4442// CHECK1-NEXT:    [[TMP45:%.*]] = load i64, ptr [[LA_CASTED]], align 8443// CHECK1-NEXT:    [[TMP46:%.*]] = load i32, ptr [[COMP]], align 4444// CHECK1-NEXT:    store i32 [[TMP46]], ptr [[COMP_CASTED8]], align 4445// CHECK1-NEXT:    [[TMP47:%.*]] = load i64, ptr [[COMP_CASTED8]], align 8446// CHECK1-NEXT:    [[TMP48:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 0447// CHECK1-NEXT:    store i64 [[TMP45]], ptr [[TMP48]], align 8448// CHECK1-NEXT:    [[TMP49:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 0449// CHECK1-NEXT:    store i64 [[TMP45]], ptr [[TMP49]], align 8450// CHECK1-NEXT:    [[TMP50:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS11]], i64 0, i64 0451// CHECK1-NEXT:    store ptr null, ptr [[TMP50]], align 8452// CHECK1-NEXT:    [[TMP51:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 1453// CHECK1-NEXT:    store i64 [[TMP47]], ptr [[TMP51]], align 8454// CHECK1-NEXT:    [[TMP52:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 1455// CHECK1-NEXT:    store i64 [[TMP47]], ptr [[TMP52]], align 8456// CHECK1-NEXT:    [[TMP53:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS11]], i64 0, i64 1457// CHECK1-NEXT:    store ptr null, ptr [[TMP53]], align 8458// CHECK1-NEXT:    [[TMP54:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 0459// CHECK1-NEXT:    [[TMP55:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 0460// CHECK1-NEXT:    [[TMP56:%.*]] = load i32, ptr [[LA]], align 4461// CHECK1-NEXT:    [[TMP57:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP56]], 0462// CHECK1-NEXT:    [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 0463// CHECK1-NEXT:    store i32 3, ptr [[TMP58]], align 4464// CHECK1-NEXT:    [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 1465// CHECK1-NEXT:    store i32 2, ptr [[TMP59]], align 4466// CHECK1-NEXT:    [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 2467// CHECK1-NEXT:    store ptr [[TMP54]], ptr [[TMP60]], align 8468// CHECK1-NEXT:    [[TMP61:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 3469// CHECK1-NEXT:    store ptr [[TMP55]], ptr [[TMP61]], align 8470// CHECK1-NEXT:    [[TMP62:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 4471// CHECK1-NEXT:    store ptr @.offload_sizes.3, ptr [[TMP62]], align 8472// CHECK1-NEXT:    [[TMP63:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 5473// CHECK1-NEXT:    store ptr @.offload_maptypes.4, ptr [[TMP63]], align 8474// CHECK1-NEXT:    [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 6475// CHECK1-NEXT:    store ptr null, ptr [[TMP64]], align 8476// CHECK1-NEXT:    [[TMP65:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 7477// CHECK1-NEXT:    store ptr null, ptr [[TMP65]], align 8478// CHECK1-NEXT:    [[TMP66:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 8479// CHECK1-NEXT:    store i64 0, ptr [[TMP66]], align 8480// CHECK1-NEXT:    [[TMP67:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 9481// CHECK1-NEXT:    store i64 0, ptr [[TMP67]], align 8482// CHECK1-NEXT:    [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 10483// CHECK1-NEXT:    store [3 x i32] [[TMP57]], ptr [[TMP68]], align 4484// CHECK1-NEXT:    [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 11485// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP69]], align 4486// CHECK1-NEXT:    [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 12487// CHECK1-NEXT:    store i32 0, ptr [[TMP70]], align 4488// CHECK1-NEXT:    [[TMP71:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[TMP56]], i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46.region_id, ptr [[KERNEL_ARGS12]])489// CHECK1-NEXT:    [[TMP72:%.*]] = icmp ne i32 [[TMP71]], 0490// CHECK1-NEXT:    br i1 [[TMP72]], label [[OMP_OFFLOAD_FAILED13:%.*]], label [[OMP_OFFLOAD_CONT14:%.*]]491// CHECK1:       omp_offload.failed13:492// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46(i64 [[TMP45]], i64 [[TMP47]]) #[[ATTR2]]493// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT14]]494// CHECK1:       omp_offload.cont14:495// CHECK1-NEXT:    [[TMP73:%.*]] = load i32, ptr [[LA]], align 4496// CHECK1-NEXT:    store i32 [[TMP73]], ptr [[LA_CASTED15]], align 4497// CHECK1-NEXT:    [[TMP74:%.*]] = load i64, ptr [[LA_CASTED15]], align 8498// CHECK1-NEXT:    [[TMP75:%.*]] = load i32, ptr [[COMP]], align 4499// CHECK1-NEXT:    store i32 [[TMP75]], ptr [[COMP_CASTED16]], align 4500// CHECK1-NEXT:    [[TMP76:%.*]] = load i64, ptr [[COMP_CASTED16]], align 8501// CHECK1-NEXT:    [[TMP77:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS17]], i32 0, i32 0502// CHECK1-NEXT:    store i64 [[TMP74]], ptr [[TMP77]], align 8503// CHECK1-NEXT:    [[TMP78:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS18]], i32 0, i32 0504// CHECK1-NEXT:    store i64 [[TMP74]], ptr [[TMP78]], align 8505// CHECK1-NEXT:    [[TMP79:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS19]], i64 0, i64 0506// CHECK1-NEXT:    store ptr null, ptr [[TMP79]], align 8507// CHECK1-NEXT:    [[TMP80:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS17]], i32 0, i32 1508// CHECK1-NEXT:    store i64 [[TMP76]], ptr [[TMP80]], align 8509// CHECK1-NEXT:    [[TMP81:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS18]], i32 0, i32 1510// CHECK1-NEXT:    store i64 [[TMP76]], ptr [[TMP81]], align 8511// CHECK1-NEXT:    [[TMP82:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS19]], i64 0, i64 1512// CHECK1-NEXT:    store ptr null, ptr [[TMP82]], align 8513// CHECK1-NEXT:    [[TMP83:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS17]], i32 0, i32 0514// CHECK1-NEXT:    [[TMP84:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS18]], i32 0, i32 0515// CHECK1-NEXT:    [[TMP85:%.*]] = load i32, ptr [[LA]], align 4516// CHECK1-NEXT:    [[TMP86:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP85]], 0517// CHECK1-NEXT:    [[TMP87:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 0518// CHECK1-NEXT:    store i32 3, ptr [[TMP87]], align 4519// CHECK1-NEXT:    [[TMP88:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 1520// CHECK1-NEXT:    store i32 2, ptr [[TMP88]], align 4521// CHECK1-NEXT:    [[TMP89:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 2522// CHECK1-NEXT:    store ptr [[TMP83]], ptr [[TMP89]], align 8523// CHECK1-NEXT:    [[TMP90:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 3524// CHECK1-NEXT:    store ptr [[TMP84]], ptr [[TMP90]], align 8525// CHECK1-NEXT:    [[TMP91:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 4526// CHECK1-NEXT:    store ptr @.offload_sizes.5, ptr [[TMP91]], align 8527// CHECK1-NEXT:    [[TMP92:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 5528// CHECK1-NEXT:    store ptr @.offload_maptypes.6, ptr [[TMP92]], align 8529// CHECK1-NEXT:    [[TMP93:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 6530// CHECK1-NEXT:    store ptr null, ptr [[TMP93]], align 8531// CHECK1-NEXT:    [[TMP94:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 7532// CHECK1-NEXT:    store ptr null, ptr [[TMP94]], align 8533// CHECK1-NEXT:    [[TMP95:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 8534// CHECK1-NEXT:    store i64 0, ptr [[TMP95]], align 8535// CHECK1-NEXT:    [[TMP96:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 9536// CHECK1-NEXT:    store i64 0, ptr [[TMP96]], align 8537// CHECK1-NEXT:    [[TMP97:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 10538// CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP97]], align 4539// CHECK1-NEXT:    [[TMP98:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 11540// CHECK1-NEXT:    store [3 x i32] [[TMP86]], ptr [[TMP98]], align 4541// CHECK1-NEXT:    [[TMP99:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 12542// CHECK1-NEXT:    store i32 0, ptr [[TMP99]], align 4543// CHECK1-NEXT:    [[TMP100:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 [[TMP85]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53.region_id, ptr [[KERNEL_ARGS20]])544// CHECK1-NEXT:    [[TMP101:%.*]] = icmp ne i32 [[TMP100]], 0545// CHECK1-NEXT:    br i1 [[TMP101]], label [[OMP_OFFLOAD_FAILED21:%.*]], label [[OMP_OFFLOAD_CONT22:%.*]]546// CHECK1:       omp_offload.failed21:547// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53(i64 [[TMP74]], i64 [[TMP76]]) #[[ATTR2]]548// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT22]]549// CHECK1:       omp_offload.cont22:550// CHECK1-NEXT:    [[TMP102:%.*]] = load i32, ptr @Gbla, align 4551// CHECK1-NEXT:    store i32 [[TMP102]], ptr [[GBLA_CASTED]], align 4552// CHECK1-NEXT:    [[TMP103:%.*]] = load i64, ptr [[GBLA_CASTED]], align 8553// CHECK1-NEXT:    [[TMP104:%.*]] = load i32, ptr [[A_ADDR]], align 4554// CHECK1-NEXT:    store i32 [[TMP104]], ptr [[A_CASTED]], align 4555// CHECK1-NEXT:    [[TMP105:%.*]] = load i64, ptr [[A_CASTED]], align 8556// CHECK1-NEXT:    [[TMP106:%.*]] = load i64, ptr @Gblb, align 8557// CHECK1-NEXT:    store i64 [[TMP106]], ptr [[GBLB_CASTED]], align 8558// CHECK1-NEXT:    [[TMP107:%.*]] = load i64, ptr [[GBLB_CASTED]], align 8559// CHECK1-NEXT:    [[TMP108:%.*]] = load float, ptr [[LC]], align 4560// CHECK1-NEXT:    store float [[TMP108]], ptr [[LC_CASTED]], align 4561// CHECK1-NEXT:    [[TMP109:%.*]] = load i64, ptr [[LC_CASTED]], align 8562// CHECK1-NEXT:    [[TMP110:%.*]] = load i32, ptr [[COMP]], align 4563// CHECK1-NEXT:    store i32 [[TMP110]], ptr [[COMP_CASTED23]], align 4564// CHECK1-NEXT:    [[TMP111:%.*]] = load i64, ptr [[COMP_CASTED23]], align 8565// CHECK1-NEXT:    [[TMP112:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 0566// CHECK1-NEXT:    store i64 [[TMP103]], ptr [[TMP112]], align 8567// CHECK1-NEXT:    [[TMP113:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 0568// CHECK1-NEXT:    store i64 [[TMP103]], ptr [[TMP113]], align 8569// CHECK1-NEXT:    [[TMP114:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS26]], i64 0, i64 0570// CHECK1-NEXT:    store ptr null, ptr [[TMP114]], align 8571// CHECK1-NEXT:    [[TMP115:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 1572// CHECK1-NEXT:    store i64 [[TMP105]], ptr [[TMP115]], align 8573// CHECK1-NEXT:    [[TMP116:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 1574// CHECK1-NEXT:    store i64 [[TMP105]], ptr [[TMP116]], align 8575// CHECK1-NEXT:    [[TMP117:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS26]], i64 0, i64 1576// CHECK1-NEXT:    store ptr null, ptr [[TMP117]], align 8577// CHECK1-NEXT:    [[TMP118:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 2578// CHECK1-NEXT:    store i64 [[TMP107]], ptr [[TMP118]], align 8579// CHECK1-NEXT:    [[TMP119:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 2580// CHECK1-NEXT:    store i64 [[TMP107]], ptr [[TMP119]], align 8581// CHECK1-NEXT:    [[TMP120:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS26]], i64 0, i64 2582// CHECK1-NEXT:    store ptr null, ptr [[TMP120]], align 8583// CHECK1-NEXT:    [[TMP121:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 3584// CHECK1-NEXT:    store i64 [[TMP109]], ptr [[TMP121]], align 8585// CHECK1-NEXT:    [[TMP122:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 3586// CHECK1-NEXT:    store i64 [[TMP109]], ptr [[TMP122]], align 8587// CHECK1-NEXT:    [[TMP123:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS26]], i64 0, i64 3588// CHECK1-NEXT:    store ptr null, ptr [[TMP123]], align 8589// CHECK1-NEXT:    [[TMP124:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 4590// CHECK1-NEXT:    store i64 [[TMP111]], ptr [[TMP124]], align 8591// CHECK1-NEXT:    [[TMP125:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 4592// CHECK1-NEXT:    store i64 [[TMP111]], ptr [[TMP125]], align 8593// CHECK1-NEXT:    [[TMP126:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS26]], i64 0, i64 4594// CHECK1-NEXT:    store ptr null, ptr [[TMP126]], align 8595// CHECK1-NEXT:    [[TMP127:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 0596// CHECK1-NEXT:    [[TMP128:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 0597// CHECK1-NEXT:    [[TMP129:%.*]] = load i32, ptr @Gbla, align 4598// CHECK1-NEXT:    [[TMP130:%.*]] = load i32, ptr [[A_ADDR]], align 4599// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP129]], [[TMP130]]600// CHECK1-NEXT:    [[TMP131:%.*]] = load i64, ptr @Gblb, align 8601// CHECK1-NEXT:    [[TMP132:%.*]] = load float, ptr [[LC]], align 4602// CHECK1-NEXT:    [[CONV:%.*]] = fptosi float [[TMP132]] to i64603// CHECK1-NEXT:    [[ADD27:%.*]] = add nsw i64 [[TMP131]], [[CONV]]604// CHECK1-NEXT:    [[TMP133:%.*]] = trunc i64 [[ADD27]] to i32605// CHECK1-NEXT:    [[TMP134:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[ADD]], 0606// CHECK1-NEXT:    [[TMP135:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP133]], 0607// CHECK1-NEXT:    [[TMP136:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 0608// CHECK1-NEXT:    store i32 3, ptr [[TMP136]], align 4609// CHECK1-NEXT:    [[TMP137:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 1610// CHECK1-NEXT:    store i32 5, ptr [[TMP137]], align 4611// CHECK1-NEXT:    [[TMP138:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 2612// CHECK1-NEXT:    store ptr [[TMP127]], ptr [[TMP138]], align 8613// CHECK1-NEXT:    [[TMP139:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 3614// CHECK1-NEXT:    store ptr [[TMP128]], ptr [[TMP139]], align 8615// CHECK1-NEXT:    [[TMP140:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 4616// CHECK1-NEXT:    store ptr @.offload_sizes.7, ptr [[TMP140]], align 8617// CHECK1-NEXT:    [[TMP141:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 5618// CHECK1-NEXT:    store ptr @.offload_maptypes.8, ptr [[TMP141]], align 8619// CHECK1-NEXT:    [[TMP142:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 6620// CHECK1-NEXT:    store ptr null, ptr [[TMP142]], align 8621// CHECK1-NEXT:    [[TMP143:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 7622// CHECK1-NEXT:    store ptr null, ptr [[TMP143]], align 8623// CHECK1-NEXT:    [[TMP144:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 8624// CHECK1-NEXT:    store i64 0, ptr [[TMP144]], align 8625// CHECK1-NEXT:    [[TMP145:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 9626// CHECK1-NEXT:    store i64 0, ptr [[TMP145]], align 8627// CHECK1-NEXT:    [[TMP146:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 10628// CHECK1-NEXT:    store [3 x i32] [[TMP134]], ptr [[TMP146]], align 4629// CHECK1-NEXT:    [[TMP147:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 11630// CHECK1-NEXT:    store [3 x i32] [[TMP135]], ptr [[TMP147]], align 4631// CHECK1-NEXT:    [[TMP148:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 12632// CHECK1-NEXT:    store i32 0, ptr [[TMP148]], align 4633// CHECK1-NEXT:    [[TMP149:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[ADD]], i32 [[TMP133]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62.region_id, ptr [[KERNEL_ARGS28]])634// CHECK1-NEXT:    [[TMP150:%.*]] = icmp ne i32 [[TMP149]], 0635// CHECK1-NEXT:    br i1 [[TMP150]], label [[OMP_OFFLOAD_FAILED29:%.*]], label [[OMP_OFFLOAD_CONT30:%.*]]636// CHECK1:       omp_offload.failed29:637// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62(i64 [[TMP103]], i64 [[TMP105]], i64 [[TMP107]], i64 [[TMP109]], i64 [[TMP111]]) #[[ATTR2]]638// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT30]]639// CHECK1:       omp_offload.cont30:640// CHECK1-NEXT:    [[TMP151:%.*]] = load ptr, ptr @Gblc, align 8641// CHECK1-NEXT:    store ptr [[TMP151]], ptr [[TMP]], align 8642// CHECK1-NEXT:    [[TMP152:%.*]] = load i32, ptr @Gbla, align 4643// CHECK1-NEXT:    store i32 [[TMP152]], ptr [[GBLC_CASTED]], align 4644// CHECK1-NEXT:    [[TMP153:%.*]] = load i64, ptr [[GBLC_CASTED]], align 8645// CHECK1-NEXT:    [[TMP154:%.*]] = load i32, ptr [[COMP]], align 4646// CHECK1-NEXT:    store i32 [[TMP154]], ptr [[COMP_CASTED31]], align 4647// CHECK1-NEXT:    [[TMP155:%.*]] = load i64, ptr [[COMP_CASTED31]], align 8648// CHECK1-NEXT:    [[TMP156:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS32]], i32 0, i32 0649// CHECK1-NEXT:    store i64 [[TMP153]], ptr [[TMP156]], align 8650// CHECK1-NEXT:    [[TMP157:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS33]], i32 0, i32 0651// CHECK1-NEXT:    store i64 [[TMP153]], ptr [[TMP157]], align 8652// CHECK1-NEXT:    [[TMP158:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS34]], i64 0, i64 0653// CHECK1-NEXT:    store ptr null, ptr [[TMP158]], align 8654// CHECK1-NEXT:    [[TMP159:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS32]], i32 0, i32 1655// CHECK1-NEXT:    store i64 [[TMP155]], ptr [[TMP159]], align 8656// CHECK1-NEXT:    [[TMP160:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS33]], i32 0, i32 1657// CHECK1-NEXT:    store i64 [[TMP155]], ptr [[TMP160]], align 8658// CHECK1-NEXT:    [[TMP161:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS34]], i64 0, i64 1659// CHECK1-NEXT:    store ptr null, ptr [[TMP161]], align 8660// CHECK1-NEXT:    [[TMP162:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS32]], i32 0, i32 0661// CHECK1-NEXT:    [[TMP163:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS33]], i32 0, i32 0662// CHECK1-NEXT:    [[TMP164:%.*]] = load ptr, ptr @Gblc, align 8663// CHECK1-NEXT:    store ptr [[TMP164]], ptr [[_TMP35]], align 8664// CHECK1-NEXT:    [[TMP165:%.*]] = load i32, ptr @Gbla, align 4665// CHECK1-NEXT:    [[ADD36:%.*]] = add nsw i32 [[TMP165]], 1666// CHECK1-NEXT:    [[TMP166:%.*]] = load ptr, ptr @Gblc, align 8667// CHECK1-NEXT:    store ptr [[TMP166]], ptr [[_TMP37]], align 8668// CHECK1-NEXT:    [[TMP167:%.*]] = load i32, ptr @Gbla, align 4669// CHECK1-NEXT:    [[ADD38:%.*]] = add nsw i32 [[TMP167]], 2670// CHECK1-NEXT:    [[TMP168:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[ADD36]], 0671// CHECK1-NEXT:    [[TMP169:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[ADD38]], 0672// CHECK1-NEXT:    [[TMP170:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 0673// CHECK1-NEXT:    store i32 3, ptr [[TMP170]], align 4674// CHECK1-NEXT:    [[TMP171:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 1675// CHECK1-NEXT:    store i32 2, ptr [[TMP171]], align 4676// CHECK1-NEXT:    [[TMP172:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 2677// CHECK1-NEXT:    store ptr [[TMP162]], ptr [[TMP172]], align 8678// CHECK1-NEXT:    [[TMP173:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 3679// CHECK1-NEXT:    store ptr [[TMP163]], ptr [[TMP173]], align 8680// CHECK1-NEXT:    [[TMP174:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 4681// CHECK1-NEXT:    store ptr @.offload_sizes.9, ptr [[TMP174]], align 8682// CHECK1-NEXT:    [[TMP175:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 5683// CHECK1-NEXT:    store ptr @.offload_maptypes.10, ptr [[TMP175]], align 8684// CHECK1-NEXT:    [[TMP176:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 6685// CHECK1-NEXT:    store ptr null, ptr [[TMP176]], align 8686// CHECK1-NEXT:    [[TMP177:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 7687// CHECK1-NEXT:    store ptr null, ptr [[TMP177]], align 8688// CHECK1-NEXT:    [[TMP178:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 8689// CHECK1-NEXT:    store i64 0, ptr [[TMP178]], align 8690// CHECK1-NEXT:    [[TMP179:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 9691// CHECK1-NEXT:    store i64 0, ptr [[TMP179]], align 8692// CHECK1-NEXT:    [[TMP180:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 10693// CHECK1-NEXT:    store [3 x i32] [[TMP168]], ptr [[TMP180]], align 4694// CHECK1-NEXT:    [[TMP181:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 11695// CHECK1-NEXT:    store [3 x i32] [[TMP169]], ptr [[TMP181]], align 4696// CHECK1-NEXT:    [[TMP182:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 12697// CHECK1-NEXT:    store i32 0, ptr [[TMP182]], align 4698// CHECK1-NEXT:    [[TMP183:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[ADD36]], i32 [[ADD38]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71.region_id, ptr [[KERNEL_ARGS39]])699// CHECK1-NEXT:    [[TMP184:%.*]] = icmp ne i32 [[TMP183]], 0700// CHECK1-NEXT:    br i1 [[TMP184]], label [[OMP_OFFLOAD_FAILED40:%.*]], label [[OMP_OFFLOAD_CONT41:%.*]]701// CHECK1:       omp_offload.failed40:702// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71(i64 [[TMP153]], i64 [[TMP155]]) #[[ATTR2]]703// CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT41]]704// CHECK1:       omp_offload.cont41:705// CHECK1-NEXT:    [[TMP185:%.*]] = load i32, ptr [[COMP]], align 4706// CHECK1-NEXT:    ret i32 [[TMP185]]707//708//709// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31710// CHECK1-SAME: (i64 noundef [[COMP:%.*]]) #[[ATTR1:[0-9]+]] {711// CHECK1-NEXT:  entry:712// CHECK1-NEXT:    [[COMP_ADDR:%.*]] = alloca i64, align 8713// CHECK1-NEXT:    store i64 [[COMP]], ptr [[COMP_ADDR]], align 8714// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31.omp_outlined, ptr [[COMP_ADDR]])715// CHECK1-NEXT:    ret void716//717//718// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31.omp_outlined719// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] {720// CHECK1-NEXT:  entry:721// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8722// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8723// CHECK1-NEXT:    [[COMP_ADDR:%.*]] = alloca ptr, align 8724// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8725// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8726// CHECK1-NEXT:    store ptr [[COMP]], ptr [[COMP_ADDR]], align 8727// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 8728// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4729// CHECK1-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP1]], 1730// CHECK1-NEXT:    store i32 [[INC]], ptr [[TMP0]], align 4731// CHECK1-NEXT:    ret void732//733//734// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37735// CHECK1-SAME: (i64 noundef [[COMP:%.*]]) #[[ATTR1]] {736// CHECK1-NEXT:  entry:737// CHECK1-NEXT:    [[COMP_ADDR:%.*]] = alloca i64, align 8738// CHECK1-NEXT:    store i64 [[COMP]], ptr [[COMP_ADDR]], align 8739// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37.omp_outlined, ptr [[COMP_ADDR]])740// CHECK1-NEXT:    ret void741//742//743// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37.omp_outlined744// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] {745// CHECK1-NEXT:  entry:746// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8747// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8748// CHECK1-NEXT:    [[COMP_ADDR:%.*]] = alloca ptr, align 8749// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8750// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8751// CHECK1-NEXT:    store ptr [[COMP]], ptr [[COMP_ADDR]], align 8752// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 8753// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4754// CHECK1-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP1]], 1755// CHECK1-NEXT:    store i32 [[INC]], ptr [[TMP0]], align 4756// CHECK1-NEXT:    ret void757//758//759// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46760// CHECK1-SAME: (i64 noundef [[LA:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1]] {761// CHECK1-NEXT:  entry:762// CHECK1-NEXT:    [[LA_ADDR:%.*]] = alloca i64, align 8763// CHECK1-NEXT:    [[COMP_ADDR:%.*]] = alloca i64, align 8764// CHECK1-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])765// CHECK1-NEXT:    store i64 [[LA]], ptr [[LA_ADDR]], align 8766// CHECK1-NEXT:    store i64 [[COMP]], ptr [[COMP_ADDR]], align 8767// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[LA_ADDR]], align 4768// CHECK1-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 0)769// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46.omp_outlined, ptr [[COMP_ADDR]])770// CHECK1-NEXT:    ret void771//772//773// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46.omp_outlined774// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] {775// CHECK1-NEXT:  entry:776// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8777// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8778// CHECK1-NEXT:    [[COMP_ADDR:%.*]] = alloca ptr, align 8779// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8780// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8781// CHECK1-NEXT:    store ptr [[COMP]], ptr [[COMP_ADDR]], align 8782// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 8783// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4784// CHECK1-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP1]], 1785// CHECK1-NEXT:    store i32 [[INC]], ptr [[TMP0]], align 4786// CHECK1-NEXT:    ret void787//788//789// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53790// CHECK1-SAME: (i64 noundef [[LA:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1]] {791// CHECK1-NEXT:  entry:792// CHECK1-NEXT:    [[LA_ADDR:%.*]] = alloca i64, align 8793// CHECK1-NEXT:    [[COMP_ADDR:%.*]] = alloca i64, align 8794// CHECK1-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])795// CHECK1-NEXT:    store i64 [[LA]], ptr [[LA_ADDR]], align 8796// CHECK1-NEXT:    store i64 [[COMP]], ptr [[COMP_ADDR]], align 8797// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[LA_ADDR]], align 4798// CHECK1-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 0, i32 [[TMP1]])799// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53.omp_outlined, ptr [[COMP_ADDR]])800// CHECK1-NEXT:    ret void801//802//803// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53.omp_outlined804// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] {805// CHECK1-NEXT:  entry:806// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8807// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8808// CHECK1-NEXT:    [[COMP_ADDR:%.*]] = alloca ptr, align 8809// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8810// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8811// CHECK1-NEXT:    store ptr [[COMP]], ptr [[COMP_ADDR]], align 8812// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 8813// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4814// CHECK1-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP1]], 1815// CHECK1-NEXT:    store i32 [[INC]], ptr [[TMP0]], align 4816// CHECK1-NEXT:    ret void817//818//819// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62820// CHECK1-SAME: (i64 noundef [[GBLA:%.*]], i64 noundef [[A:%.*]], i64 noundef [[GBLB:%.*]], i64 noundef [[LC:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1]] {821// CHECK1-NEXT:  entry:822// CHECK1-NEXT:    [[GBLA_ADDR:%.*]] = alloca i64, align 8823// CHECK1-NEXT:    [[A_ADDR:%.*]] = alloca i64, align 8824// CHECK1-NEXT:    [[GBLB_ADDR:%.*]] = alloca i64, align 8825// CHECK1-NEXT:    [[LC_ADDR:%.*]] = alloca i64, align 8826// CHECK1-NEXT:    [[COMP_ADDR:%.*]] = alloca i64, align 8827// CHECK1-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])828// CHECK1-NEXT:    store i64 [[GBLA]], ptr [[GBLA_ADDR]], align 8829// CHECK1-NEXT:    store i64 [[A]], ptr [[A_ADDR]], align 8830// CHECK1-NEXT:    store i64 [[GBLB]], ptr [[GBLB_ADDR]], align 8831// CHECK1-NEXT:    store i64 [[LC]], ptr [[LC_ADDR]], align 8832// CHECK1-NEXT:    store i64 [[COMP]], ptr [[COMP_ADDR]], align 8833// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[GBLA_ADDR]], align 4834// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[A_ADDR]], align 4835// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], [[TMP2]]836// CHECK1-NEXT:    [[TMP3:%.*]] = load i64, ptr [[GBLB_ADDR]], align 8837// CHECK1-NEXT:    [[TMP4:%.*]] = load float, ptr [[LC_ADDR]], align 4838// CHECK1-NEXT:    [[CONV:%.*]] = fptosi float [[TMP4]] to i64839// CHECK1-NEXT:    [[ADD1:%.*]] = add nsw i64 [[TMP3]], [[CONV]]840// CHECK1-NEXT:    [[TMP5:%.*]] = trunc i64 [[ADD1]] to i32841// CHECK1-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[ADD]], i32 [[TMP5]])842// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62.omp_outlined, ptr [[COMP_ADDR]])843// CHECK1-NEXT:    ret void844//845//846// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62.omp_outlined847// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] {848// CHECK1-NEXT:  entry:849// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8850// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8851// CHECK1-NEXT:    [[COMP_ADDR:%.*]] = alloca ptr, align 8852// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8853// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8854// CHECK1-NEXT:    store ptr [[COMP]], ptr [[COMP_ADDR]], align 8855// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 8856// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4857// CHECK1-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP1]], 1858// CHECK1-NEXT:    store i32 [[INC]], ptr [[TMP0]], align 4859// CHECK1-NEXT:    ret void860//861//862// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71863// CHECK1-SAME: (i64 noundef [[GBLC:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1]] {864// CHECK1-NEXT:  entry:865// CHECK1-NEXT:    [[GBLC_ADDR:%.*]] = alloca i64, align 8866// CHECK1-NEXT:    [[COMP_ADDR:%.*]] = alloca i64, align 8867// CHECK1-NEXT:    [[TMP:%.*]] = alloca ptr, align 8868// CHECK1-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])869// CHECK1-NEXT:    store i64 [[GBLC]], ptr [[GBLC_ADDR]], align 8870// CHECK1-NEXT:    store i64 [[COMP]], ptr [[COMP_ADDR]], align 8871// CHECK1-NEXT:    store ptr [[GBLC_ADDR]], ptr [[TMP]], align 8872// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr @Gbla, align 4873// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], 1874// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr @Gbla, align 4875// CHECK1-NEXT:    [[ADD1:%.*]] = add nsw i32 [[TMP2]], 2876// CHECK1-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[ADD]], i32 [[ADD1]])877// CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[TMP]], align 8878// CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71.omp_outlined, ptr [[COMP_ADDR]], ptr [[TMP3]])879// CHECK1-NEXT:    ret void880//881//882// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71.omp_outlined883// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[GBLC:%.*]]) #[[ATTR1]] {884// CHECK1-NEXT:  entry:885// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8886// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8887// CHECK1-NEXT:    [[COMP_ADDR:%.*]] = alloca ptr, align 8888// CHECK1-NEXT:    [[GBLC_ADDR:%.*]] = alloca ptr, align 8889// CHECK1-NEXT:    [[TMP:%.*]] = alloca ptr, align 8890// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8891// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8892// CHECK1-NEXT:    store ptr [[COMP]], ptr [[COMP_ADDR]], align 8893// CHECK1-NEXT:    store ptr [[GBLC]], ptr [[GBLC_ADDR]], align 8894// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 8895// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[GBLC_ADDR]], align 8896// CHECK1-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 8897// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr @Gbla, align 4898// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 4899// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP3]], [[TMP2]]900// CHECK1-NEXT:    store i32 [[ADD]], ptr [[TMP0]], align 4901// CHECK1-NEXT:    ret void902//903//904// CHECK3-LABEL: define {{[^@]+}}@_Z27teams_argument_global_locali905// CHECK3-SAME: (i32 noundef [[A:%.*]]) #[[ATTR0:[0-9]+]] {906// CHECK3-NEXT:  entry:907// CHECK3-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 4908// CHECK3-NEXT:    [[COMP:%.*]] = alloca i32, align 4909// CHECK3-NEXT:    [[LA:%.*]] = alloca i32, align 4910// CHECK3-NEXT:    [[LC:%.*]] = alloca float, align 4911// CHECK3-NEXT:    [[COMP_CASTED:%.*]] = alloca i32, align 4912// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4913// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4914// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4915// CHECK3-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8916// CHECK3-NEXT:    [[COMP_CASTED1:%.*]] = alloca i32, align 4917// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS2:%.*]] = alloca [1 x ptr], align 4918// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS3:%.*]] = alloca [1 x ptr], align 4919// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS4:%.*]] = alloca [1 x ptr], align 4920// CHECK3-NEXT:    [[KERNEL_ARGS5:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8921// CHECK3-NEXT:    [[LA_CASTED:%.*]] = alloca i32, align 4922// CHECK3-NEXT:    [[COMP_CASTED8:%.*]] = alloca i32, align 4923// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS9:%.*]] = alloca [2 x ptr], align 4924// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS10:%.*]] = alloca [2 x ptr], align 4925// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS11:%.*]] = alloca [2 x ptr], align 4926// CHECK3-NEXT:    [[KERNEL_ARGS12:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8927// CHECK3-NEXT:    [[LA_CASTED15:%.*]] = alloca i32, align 4928// CHECK3-NEXT:    [[COMP_CASTED16:%.*]] = alloca i32, align 4929// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS17:%.*]] = alloca [2 x ptr], align 4930// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS18:%.*]] = alloca [2 x ptr], align 4931// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS19:%.*]] = alloca [2 x ptr], align 4932// CHECK3-NEXT:    [[KERNEL_ARGS20:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8933// CHECK3-NEXT:    [[GBLA_CASTED:%.*]] = alloca i32, align 4934// CHECK3-NEXT:    [[A_CASTED:%.*]] = alloca i32, align 4935// CHECK3-NEXT:    [[LC_CASTED:%.*]] = alloca i32, align 4936// CHECK3-NEXT:    [[COMP_CASTED23:%.*]] = alloca i32, align 4937// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS24:%.*]] = alloca [5 x ptr], align 4938// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS25:%.*]] = alloca [5 x ptr], align 4939// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS26:%.*]] = alloca [5 x ptr], align 4940// CHECK3-NEXT:    [[KERNEL_ARGS28:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8941// CHECK3-NEXT:    [[TMP:%.*]] = alloca ptr, align 4942// CHECK3-NEXT:    [[GBLC_CASTED:%.*]] = alloca i32, align 4943// CHECK3-NEXT:    [[COMP_CASTED31:%.*]] = alloca i32, align 4944// CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS32:%.*]] = alloca [2 x ptr], align 4945// CHECK3-NEXT:    [[DOTOFFLOAD_PTRS33:%.*]] = alloca [2 x ptr], align 4946// CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS34:%.*]] = alloca [2 x ptr], align 4947// CHECK3-NEXT:    [[_TMP35:%.*]] = alloca ptr, align 4948// CHECK3-NEXT:    [[_TMP37:%.*]] = alloca ptr, align 4949// CHECK3-NEXT:    [[KERNEL_ARGS39:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8950// CHECK3-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 4951// CHECK3-NEXT:    store i32 1, ptr [[COMP]], align 4952// CHECK3-NEXT:    store i32 23, ptr [[LA]], align 4953// CHECK3-NEXT:    store float 2.500000e+01, ptr [[LC]], align 4954// CHECK3-NEXT:    [[TMP0:%.*]] = load i32, ptr [[COMP]], align 4955// CHECK3-NEXT:    store i32 [[TMP0]], ptr [[COMP_CASTED]], align 4956// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[COMP_CASTED]], align 4957// CHECK3-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0958// CHECK3-NEXT:    store i32 [[TMP1]], ptr [[TMP2]], align 4959// CHECK3-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0960// CHECK3-NEXT:    store i32 [[TMP1]], ptr [[TMP3]], align 4961// CHECK3-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0962// CHECK3-NEXT:    store ptr null, ptr [[TMP4]], align 4963// CHECK3-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0964// CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0965// CHECK3-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0966// CHECK3-NEXT:    store i32 3, ptr [[TMP7]], align 4967// CHECK3-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1968// CHECK3-NEXT:    store i32 1, ptr [[TMP8]], align 4969// CHECK3-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2970// CHECK3-NEXT:    store ptr [[TMP5]], ptr [[TMP9]], align 4971// CHECK3-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3972// CHECK3-NEXT:    store ptr [[TMP6]], ptr [[TMP10]], align 4973// CHECK3-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4974// CHECK3-NEXT:    store ptr @.offload_sizes, ptr [[TMP11]], align 4975// CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5976// CHECK3-NEXT:    store ptr @.offload_maptypes, ptr [[TMP12]], align 4977// CHECK3-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6978// CHECK3-NEXT:    store ptr null, ptr [[TMP13]], align 4979// CHECK3-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7980// CHECK3-NEXT:    store ptr null, ptr [[TMP14]], align 4981// CHECK3-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8982// CHECK3-NEXT:    store i64 0, ptr [[TMP15]], align 8983// CHECK3-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9984// CHECK3-NEXT:    store i64 0, ptr [[TMP16]], align 8985// CHECK3-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10986// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4987// CHECK3-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11988// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP18]], align 4989// CHECK3-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12990// CHECK3-NEXT:    store i32 0, ptr [[TMP19]], align 4991// CHECK3-NEXT:    [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31.region_id, ptr [[KERNEL_ARGS]])992// CHECK3-NEXT:    [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0993// CHECK3-NEXT:    br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]994// CHECK3:       omp_offload.failed:995// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31(i32 [[TMP1]]) #[[ATTR2:[0-9]+]]996// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT]]997// CHECK3:       omp_offload.cont:998// CHECK3-NEXT:    [[TMP22:%.*]] = load i32, ptr [[COMP]], align 4999// CHECK3-NEXT:    store i32 [[TMP22]], ptr [[COMP_CASTED1]], align 41000// CHECK3-NEXT:    [[TMP23:%.*]] = load i32, ptr [[COMP_CASTED1]], align 41001// CHECK3-NEXT:    [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 01002// CHECK3-NEXT:    store i32 [[TMP23]], ptr [[TMP24]], align 41003// CHECK3-NEXT:    [[TMP25:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 01004// CHECK3-NEXT:    store i32 [[TMP23]], ptr [[TMP25]], align 41005// CHECK3-NEXT:    [[TMP26:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i32 0, i32 01006// CHECK3-NEXT:    store ptr null, ptr [[TMP26]], align 41007// CHECK3-NEXT:    [[TMP27:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 01008// CHECK3-NEXT:    [[TMP28:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 01009// CHECK3-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 01010// CHECK3-NEXT:    store i32 3, ptr [[TMP29]], align 41011// CHECK3-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 11012// CHECK3-NEXT:    store i32 1, ptr [[TMP30]], align 41013// CHECK3-NEXT:    [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 21014// CHECK3-NEXT:    store ptr [[TMP27]], ptr [[TMP31]], align 41015// CHECK3-NEXT:    [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 31016// CHECK3-NEXT:    store ptr [[TMP28]], ptr [[TMP32]], align 41017// CHECK3-NEXT:    [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 41018// CHECK3-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP33]], align 41019// CHECK3-NEXT:    [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 51020// CHECK3-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP34]], align 41021// CHECK3-NEXT:    [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 61022// CHECK3-NEXT:    store ptr null, ptr [[TMP35]], align 41023// CHECK3-NEXT:    [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 71024// CHECK3-NEXT:    store ptr null, ptr [[TMP36]], align 41025// CHECK3-NEXT:    [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 81026// CHECK3-NEXT:    store i64 0, ptr [[TMP37]], align 81027// CHECK3-NEXT:    [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 91028// CHECK3-NEXT:    store i64 0, ptr [[TMP38]], align 81029// CHECK3-NEXT:    [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 101030// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP39]], align 41031// CHECK3-NEXT:    [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 111032// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP40]], align 41033// CHECK3-NEXT:    [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 121034// CHECK3-NEXT:    store i32 0, ptr [[TMP41]], align 41035// CHECK3-NEXT:    [[TMP42:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37.region_id, ptr [[KERNEL_ARGS5]])1036// CHECK3-NEXT:    [[TMP43:%.*]] = icmp ne i32 [[TMP42]], 01037// CHECK3-NEXT:    br i1 [[TMP43]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]]1038// CHECK3:       omp_offload.failed6:1039// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37(i32 [[TMP23]]) #[[ATTR2]]1040// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT7]]1041// CHECK3:       omp_offload.cont7:1042// CHECK3-NEXT:    [[TMP44:%.*]] = load i32, ptr [[LA]], align 41043// CHECK3-NEXT:    store i32 [[TMP44]], ptr [[LA_CASTED]], align 41044// CHECK3-NEXT:    [[TMP45:%.*]] = load i32, ptr [[LA_CASTED]], align 41045// CHECK3-NEXT:    [[TMP46:%.*]] = load i32, ptr [[COMP]], align 41046// CHECK3-NEXT:    store i32 [[TMP46]], ptr [[COMP_CASTED8]], align 41047// CHECK3-NEXT:    [[TMP47:%.*]] = load i32, ptr [[COMP_CASTED8]], align 41048// CHECK3-NEXT:    [[TMP48:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 01049// CHECK3-NEXT:    store i32 [[TMP45]], ptr [[TMP48]], align 41050// CHECK3-NEXT:    [[TMP49:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 01051// CHECK3-NEXT:    store i32 [[TMP45]], ptr [[TMP49]], align 41052// CHECK3-NEXT:    [[TMP50:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS11]], i32 0, i32 01053// CHECK3-NEXT:    store ptr null, ptr [[TMP50]], align 41054// CHECK3-NEXT:    [[TMP51:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 11055// CHECK3-NEXT:    store i32 [[TMP47]], ptr [[TMP51]], align 41056// CHECK3-NEXT:    [[TMP52:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 11057// CHECK3-NEXT:    store i32 [[TMP47]], ptr [[TMP52]], align 41058// CHECK3-NEXT:    [[TMP53:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS11]], i32 0, i32 11059// CHECK3-NEXT:    store ptr null, ptr [[TMP53]], align 41060// CHECK3-NEXT:    [[TMP54:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 01061// CHECK3-NEXT:    [[TMP55:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 01062// CHECK3-NEXT:    [[TMP56:%.*]] = load i32, ptr [[LA]], align 41063// CHECK3-NEXT:    [[TMP57:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP56]], 01064// CHECK3-NEXT:    [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 01065// CHECK3-NEXT:    store i32 3, ptr [[TMP58]], align 41066// CHECK3-NEXT:    [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 11067// CHECK3-NEXT:    store i32 2, ptr [[TMP59]], align 41068// CHECK3-NEXT:    [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 21069// CHECK3-NEXT:    store ptr [[TMP54]], ptr [[TMP60]], align 41070// CHECK3-NEXT:    [[TMP61:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 31071// CHECK3-NEXT:    store ptr [[TMP55]], ptr [[TMP61]], align 41072// CHECK3-NEXT:    [[TMP62:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 41073// CHECK3-NEXT:    store ptr @.offload_sizes.3, ptr [[TMP62]], align 41074// CHECK3-NEXT:    [[TMP63:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 51075// CHECK3-NEXT:    store ptr @.offload_maptypes.4, ptr [[TMP63]], align 41076// CHECK3-NEXT:    [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 61077// CHECK3-NEXT:    store ptr null, ptr [[TMP64]], align 41078// CHECK3-NEXT:    [[TMP65:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 71079// CHECK3-NEXT:    store ptr null, ptr [[TMP65]], align 41080// CHECK3-NEXT:    [[TMP66:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 81081// CHECK3-NEXT:    store i64 0, ptr [[TMP66]], align 81082// CHECK3-NEXT:    [[TMP67:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 91083// CHECK3-NEXT:    store i64 0, ptr [[TMP67]], align 81084// CHECK3-NEXT:    [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 101085// CHECK3-NEXT:    store [3 x i32] [[TMP57]], ptr [[TMP68]], align 41086// CHECK3-NEXT:    [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 111087// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP69]], align 41088// CHECK3-NEXT:    [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 121089// CHECK3-NEXT:    store i32 0, ptr [[TMP70]], align 41090// CHECK3-NEXT:    [[TMP71:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[TMP56]], i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46.region_id, ptr [[KERNEL_ARGS12]])1091// CHECK3-NEXT:    [[TMP72:%.*]] = icmp ne i32 [[TMP71]], 01092// CHECK3-NEXT:    br i1 [[TMP72]], label [[OMP_OFFLOAD_FAILED13:%.*]], label [[OMP_OFFLOAD_CONT14:%.*]]1093// CHECK3:       omp_offload.failed13:1094// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46(i32 [[TMP45]], i32 [[TMP47]]) #[[ATTR2]]1095// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT14]]1096// CHECK3:       omp_offload.cont14:1097// CHECK3-NEXT:    [[TMP73:%.*]] = load i32, ptr [[LA]], align 41098// CHECK3-NEXT:    store i32 [[TMP73]], ptr [[LA_CASTED15]], align 41099// CHECK3-NEXT:    [[TMP74:%.*]] = load i32, ptr [[LA_CASTED15]], align 41100// CHECK3-NEXT:    [[TMP75:%.*]] = load i32, ptr [[COMP]], align 41101// CHECK3-NEXT:    store i32 [[TMP75]], ptr [[COMP_CASTED16]], align 41102// CHECK3-NEXT:    [[TMP76:%.*]] = load i32, ptr [[COMP_CASTED16]], align 41103// CHECK3-NEXT:    [[TMP77:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS17]], i32 0, i32 01104// CHECK3-NEXT:    store i32 [[TMP74]], ptr [[TMP77]], align 41105// CHECK3-NEXT:    [[TMP78:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS18]], i32 0, i32 01106// CHECK3-NEXT:    store i32 [[TMP74]], ptr [[TMP78]], align 41107// CHECK3-NEXT:    [[TMP79:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS19]], i32 0, i32 01108// CHECK3-NEXT:    store ptr null, ptr [[TMP79]], align 41109// CHECK3-NEXT:    [[TMP80:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS17]], i32 0, i32 11110// CHECK3-NEXT:    store i32 [[TMP76]], ptr [[TMP80]], align 41111// CHECK3-NEXT:    [[TMP81:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS18]], i32 0, i32 11112// CHECK3-NEXT:    store i32 [[TMP76]], ptr [[TMP81]], align 41113// CHECK3-NEXT:    [[TMP82:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS19]], i32 0, i32 11114// CHECK3-NEXT:    store ptr null, ptr [[TMP82]], align 41115// CHECK3-NEXT:    [[TMP83:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS17]], i32 0, i32 01116// CHECK3-NEXT:    [[TMP84:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS18]], i32 0, i32 01117// CHECK3-NEXT:    [[TMP85:%.*]] = load i32, ptr [[LA]], align 41118// CHECK3-NEXT:    [[TMP86:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP85]], 01119// CHECK3-NEXT:    [[TMP87:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 01120// CHECK3-NEXT:    store i32 3, ptr [[TMP87]], align 41121// CHECK3-NEXT:    [[TMP88:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 11122// CHECK3-NEXT:    store i32 2, ptr [[TMP88]], align 41123// CHECK3-NEXT:    [[TMP89:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 21124// CHECK3-NEXT:    store ptr [[TMP83]], ptr [[TMP89]], align 41125// CHECK3-NEXT:    [[TMP90:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 31126// CHECK3-NEXT:    store ptr [[TMP84]], ptr [[TMP90]], align 41127// CHECK3-NEXT:    [[TMP91:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 41128// CHECK3-NEXT:    store ptr @.offload_sizes.5, ptr [[TMP91]], align 41129// CHECK3-NEXT:    [[TMP92:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 51130// CHECK3-NEXT:    store ptr @.offload_maptypes.6, ptr [[TMP92]], align 41131// CHECK3-NEXT:    [[TMP93:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 61132// CHECK3-NEXT:    store ptr null, ptr [[TMP93]], align 41133// CHECK3-NEXT:    [[TMP94:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 71134// CHECK3-NEXT:    store ptr null, ptr [[TMP94]], align 41135// CHECK3-NEXT:    [[TMP95:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 81136// CHECK3-NEXT:    store i64 0, ptr [[TMP95]], align 81137// CHECK3-NEXT:    [[TMP96:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 91138// CHECK3-NEXT:    store i64 0, ptr [[TMP96]], align 81139// CHECK3-NEXT:    [[TMP97:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 101140// CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP97]], align 41141// CHECK3-NEXT:    [[TMP98:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 111142// CHECK3-NEXT:    store [3 x i32] [[TMP86]], ptr [[TMP98]], align 41143// CHECK3-NEXT:    [[TMP99:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 121144// CHECK3-NEXT:    store i32 0, ptr [[TMP99]], align 41145// CHECK3-NEXT:    [[TMP100:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 [[TMP85]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53.region_id, ptr [[KERNEL_ARGS20]])1146// CHECK3-NEXT:    [[TMP101:%.*]] = icmp ne i32 [[TMP100]], 01147// CHECK3-NEXT:    br i1 [[TMP101]], label [[OMP_OFFLOAD_FAILED21:%.*]], label [[OMP_OFFLOAD_CONT22:%.*]]1148// CHECK3:       omp_offload.failed21:1149// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53(i32 [[TMP74]], i32 [[TMP76]]) #[[ATTR2]]1150// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT22]]1151// CHECK3:       omp_offload.cont22:1152// CHECK3-NEXT:    [[TMP102:%.*]] = load i32, ptr @Gbla, align 41153// CHECK3-NEXT:    store i32 [[TMP102]], ptr [[GBLA_CASTED]], align 41154// CHECK3-NEXT:    [[TMP103:%.*]] = load i32, ptr [[GBLA_CASTED]], align 41155// CHECK3-NEXT:    [[TMP104:%.*]] = load i32, ptr [[A_ADDR]], align 41156// CHECK3-NEXT:    store i32 [[TMP104]], ptr [[A_CASTED]], align 41157// CHECK3-NEXT:    [[TMP105:%.*]] = load i32, ptr [[A_CASTED]], align 41158// CHECK3-NEXT:    [[TMP106:%.*]] = load float, ptr [[LC]], align 41159// CHECK3-NEXT:    store float [[TMP106]], ptr [[LC_CASTED]], align 41160// CHECK3-NEXT:    [[TMP107:%.*]] = load i32, ptr [[LC_CASTED]], align 41161// CHECK3-NEXT:    [[TMP108:%.*]] = load i32, ptr [[COMP]], align 41162// CHECK3-NEXT:    store i32 [[TMP108]], ptr [[COMP_CASTED23]], align 41163// CHECK3-NEXT:    [[TMP109:%.*]] = load i32, ptr [[COMP_CASTED23]], align 41164// CHECK3-NEXT:    [[TMP110:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 01165// CHECK3-NEXT:    store i32 [[TMP103]], ptr [[TMP110]], align 41166// CHECK3-NEXT:    [[TMP111:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 01167// CHECK3-NEXT:    store i32 [[TMP103]], ptr [[TMP111]], align 41168// CHECK3-NEXT:    [[TMP112:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS26]], i32 0, i32 01169// CHECK3-NEXT:    store ptr null, ptr [[TMP112]], align 41170// CHECK3-NEXT:    [[TMP113:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 11171// CHECK3-NEXT:    store i32 [[TMP105]], ptr [[TMP113]], align 41172// CHECK3-NEXT:    [[TMP114:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 11173// CHECK3-NEXT:    store i32 [[TMP105]], ptr [[TMP114]], align 41174// CHECK3-NEXT:    [[TMP115:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS26]], i32 0, i32 11175// CHECK3-NEXT:    store ptr null, ptr [[TMP115]], align 41176// CHECK3-NEXT:    [[TMP116:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 21177// CHECK3-NEXT:    store ptr @Gblb, ptr [[TMP116]], align 41178// CHECK3-NEXT:    [[TMP117:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 21179// CHECK3-NEXT:    store ptr @Gblb, ptr [[TMP117]], align 41180// CHECK3-NEXT:    [[TMP118:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS26]], i32 0, i32 21181// CHECK3-NEXT:    store ptr null, ptr [[TMP118]], align 41182// CHECK3-NEXT:    [[TMP119:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 31183// CHECK3-NEXT:    store i32 [[TMP107]], ptr [[TMP119]], align 41184// CHECK3-NEXT:    [[TMP120:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 31185// CHECK3-NEXT:    store i32 [[TMP107]], ptr [[TMP120]], align 41186// CHECK3-NEXT:    [[TMP121:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS26]], i32 0, i32 31187// CHECK3-NEXT:    store ptr null, ptr [[TMP121]], align 41188// CHECK3-NEXT:    [[TMP122:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 41189// CHECK3-NEXT:    store i32 [[TMP109]], ptr [[TMP122]], align 41190// CHECK3-NEXT:    [[TMP123:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 41191// CHECK3-NEXT:    store i32 [[TMP109]], ptr [[TMP123]], align 41192// CHECK3-NEXT:    [[TMP124:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS26]], i32 0, i32 41193// CHECK3-NEXT:    store ptr null, ptr [[TMP124]], align 41194// CHECK3-NEXT:    [[TMP125:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 01195// CHECK3-NEXT:    [[TMP126:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 01196// CHECK3-NEXT:    [[TMP127:%.*]] = load i32, ptr @Gbla, align 41197// CHECK3-NEXT:    [[TMP128:%.*]] = load i32, ptr [[A_ADDR]], align 41198// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP127]], [[TMP128]]1199// CHECK3-NEXT:    [[TMP129:%.*]] = load i64, ptr @Gblb, align 81200// CHECK3-NEXT:    [[TMP130:%.*]] = load float, ptr [[LC]], align 41201// CHECK3-NEXT:    [[CONV:%.*]] = fptosi float [[TMP130]] to i641202// CHECK3-NEXT:    [[ADD27:%.*]] = add nsw i64 [[TMP129]], [[CONV]]1203// CHECK3-NEXT:    [[TMP131:%.*]] = trunc i64 [[ADD27]] to i321204// CHECK3-NEXT:    [[TMP132:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[ADD]], 01205// CHECK3-NEXT:    [[TMP133:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP131]], 01206// CHECK3-NEXT:    [[TMP134:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 01207// CHECK3-NEXT:    store i32 3, ptr [[TMP134]], align 41208// CHECK3-NEXT:    [[TMP135:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 11209// CHECK3-NEXT:    store i32 5, ptr [[TMP135]], align 41210// CHECK3-NEXT:    [[TMP136:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 21211// CHECK3-NEXT:    store ptr [[TMP125]], ptr [[TMP136]], align 41212// CHECK3-NEXT:    [[TMP137:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 31213// CHECK3-NEXT:    store ptr [[TMP126]], ptr [[TMP137]], align 41214// CHECK3-NEXT:    [[TMP138:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 41215// CHECK3-NEXT:    store ptr @.offload_sizes.7, ptr [[TMP138]], align 41216// CHECK3-NEXT:    [[TMP139:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 51217// CHECK3-NEXT:    store ptr @.offload_maptypes.8, ptr [[TMP139]], align 41218// CHECK3-NEXT:    [[TMP140:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 61219// CHECK3-NEXT:    store ptr null, ptr [[TMP140]], align 41220// CHECK3-NEXT:    [[TMP141:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 71221// CHECK3-NEXT:    store ptr null, ptr [[TMP141]], align 41222// CHECK3-NEXT:    [[TMP142:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 81223// CHECK3-NEXT:    store i64 0, ptr [[TMP142]], align 81224// CHECK3-NEXT:    [[TMP143:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 91225// CHECK3-NEXT:    store i64 0, ptr [[TMP143]], align 81226// CHECK3-NEXT:    [[TMP144:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 101227// CHECK3-NEXT:    store [3 x i32] [[TMP132]], ptr [[TMP144]], align 41228// CHECK3-NEXT:    [[TMP145:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 111229// CHECK3-NEXT:    store [3 x i32] [[TMP133]], ptr [[TMP145]], align 41230// CHECK3-NEXT:    [[TMP146:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 121231// CHECK3-NEXT:    store i32 0, ptr [[TMP146]], align 41232// CHECK3-NEXT:    [[TMP147:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[ADD]], i32 [[TMP131]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62.region_id, ptr [[KERNEL_ARGS28]])1233// CHECK3-NEXT:    [[TMP148:%.*]] = icmp ne i32 [[TMP147]], 01234// CHECK3-NEXT:    br i1 [[TMP148]], label [[OMP_OFFLOAD_FAILED29:%.*]], label [[OMP_OFFLOAD_CONT30:%.*]]1235// CHECK3:       omp_offload.failed29:1236// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62(i32 [[TMP103]], i32 [[TMP105]], ptr @Gblb, i32 [[TMP107]], i32 [[TMP109]]) #[[ATTR2]]1237// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT30]]1238// CHECK3:       omp_offload.cont30:1239// CHECK3-NEXT:    [[TMP149:%.*]] = load ptr, ptr @Gblc, align 41240// CHECK3-NEXT:    store ptr [[TMP149]], ptr [[TMP]], align 41241// CHECK3-NEXT:    [[TMP150:%.*]] = load i32, ptr @Gbla, align 41242// CHECK3-NEXT:    store i32 [[TMP150]], ptr [[GBLC_CASTED]], align 41243// CHECK3-NEXT:    [[TMP151:%.*]] = load i32, ptr [[GBLC_CASTED]], align 41244// CHECK3-NEXT:    [[TMP152:%.*]] = load i32, ptr [[COMP]], align 41245// CHECK3-NEXT:    store i32 [[TMP152]], ptr [[COMP_CASTED31]], align 41246// CHECK3-NEXT:    [[TMP153:%.*]] = load i32, ptr [[COMP_CASTED31]], align 41247// CHECK3-NEXT:    [[TMP154:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS32]], i32 0, i32 01248// CHECK3-NEXT:    store i32 [[TMP151]], ptr [[TMP154]], align 41249// CHECK3-NEXT:    [[TMP155:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS33]], i32 0, i32 01250// CHECK3-NEXT:    store i32 [[TMP151]], ptr [[TMP155]], align 41251// CHECK3-NEXT:    [[TMP156:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS34]], i32 0, i32 01252// CHECK3-NEXT:    store ptr null, ptr [[TMP156]], align 41253// CHECK3-NEXT:    [[TMP157:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS32]], i32 0, i32 11254// CHECK3-NEXT:    store i32 [[TMP153]], ptr [[TMP157]], align 41255// CHECK3-NEXT:    [[TMP158:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS33]], i32 0, i32 11256// CHECK3-NEXT:    store i32 [[TMP153]], ptr [[TMP158]], align 41257// CHECK3-NEXT:    [[TMP159:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS34]], i32 0, i32 11258// CHECK3-NEXT:    store ptr null, ptr [[TMP159]], align 41259// CHECK3-NEXT:    [[TMP160:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS32]], i32 0, i32 01260// CHECK3-NEXT:    [[TMP161:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS33]], i32 0, i32 01261// CHECK3-NEXT:    [[TMP162:%.*]] = load ptr, ptr @Gblc, align 41262// CHECK3-NEXT:    store ptr [[TMP162]], ptr [[_TMP35]], align 41263// CHECK3-NEXT:    [[TMP163:%.*]] = load i32, ptr @Gbla, align 41264// CHECK3-NEXT:    [[ADD36:%.*]] = add nsw i32 [[TMP163]], 11265// CHECK3-NEXT:    [[TMP164:%.*]] = load ptr, ptr @Gblc, align 41266// CHECK3-NEXT:    store ptr [[TMP164]], ptr [[_TMP37]], align 41267// CHECK3-NEXT:    [[TMP165:%.*]] = load i32, ptr @Gbla, align 41268// CHECK3-NEXT:    [[ADD38:%.*]] = add nsw i32 [[TMP165]], 21269// CHECK3-NEXT:    [[TMP166:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[ADD36]], 01270// CHECK3-NEXT:    [[TMP167:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[ADD38]], 01271// CHECK3-NEXT:    [[TMP168:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 01272// CHECK3-NEXT:    store i32 3, ptr [[TMP168]], align 41273// CHECK3-NEXT:    [[TMP169:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 11274// CHECK3-NEXT:    store i32 2, ptr [[TMP169]], align 41275// CHECK3-NEXT:    [[TMP170:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 21276// CHECK3-NEXT:    store ptr [[TMP160]], ptr [[TMP170]], align 41277// CHECK3-NEXT:    [[TMP171:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 31278// CHECK3-NEXT:    store ptr [[TMP161]], ptr [[TMP171]], align 41279// CHECK3-NEXT:    [[TMP172:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 41280// CHECK3-NEXT:    store ptr @.offload_sizes.9, ptr [[TMP172]], align 41281// CHECK3-NEXT:    [[TMP173:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 51282// CHECK3-NEXT:    store ptr @.offload_maptypes.10, ptr [[TMP173]], align 41283// CHECK3-NEXT:    [[TMP174:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 61284// CHECK3-NEXT:    store ptr null, ptr [[TMP174]], align 41285// CHECK3-NEXT:    [[TMP175:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 71286// CHECK3-NEXT:    store ptr null, ptr [[TMP175]], align 41287// CHECK3-NEXT:    [[TMP176:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 81288// CHECK3-NEXT:    store i64 0, ptr [[TMP176]], align 81289// CHECK3-NEXT:    [[TMP177:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 91290// CHECK3-NEXT:    store i64 0, ptr [[TMP177]], align 81291// CHECK3-NEXT:    [[TMP178:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 101292// CHECK3-NEXT:    store [3 x i32] [[TMP166]], ptr [[TMP178]], align 41293// CHECK3-NEXT:    [[TMP179:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 111294// CHECK3-NEXT:    store [3 x i32] [[TMP167]], ptr [[TMP179]], align 41295// CHECK3-NEXT:    [[TMP180:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 121296// CHECK3-NEXT:    store i32 0, ptr [[TMP180]], align 41297// CHECK3-NEXT:    [[TMP181:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[ADD36]], i32 [[ADD38]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71.region_id, ptr [[KERNEL_ARGS39]])1298// CHECK3-NEXT:    [[TMP182:%.*]] = icmp ne i32 [[TMP181]], 01299// CHECK3-NEXT:    br i1 [[TMP182]], label [[OMP_OFFLOAD_FAILED40:%.*]], label [[OMP_OFFLOAD_CONT41:%.*]]1300// CHECK3:       omp_offload.failed40:1301// CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71(i32 [[TMP151]], i32 [[TMP153]]) #[[ATTR2]]1302// CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT41]]1303// CHECK3:       omp_offload.cont41:1304// CHECK3-NEXT:    [[TMP183:%.*]] = load i32, ptr [[COMP]], align 41305// CHECK3-NEXT:    ret i32 [[TMP183]]1306//1307//1308// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l311309// CHECK3-SAME: (i32 noundef [[COMP:%.*]]) #[[ATTR1:[0-9]+]] {1310// CHECK3-NEXT:  entry:1311// CHECK3-NEXT:    [[COMP_ADDR:%.*]] = alloca i32, align 41312// CHECK3-NEXT:    store i32 [[COMP]], ptr [[COMP_ADDR]], align 41313// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31.omp_outlined, ptr [[COMP_ADDR]])1314// CHECK3-NEXT:    ret void1315//1316//1317// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31.omp_outlined1318// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] {1319// CHECK3-NEXT:  entry:1320// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41321// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41322// CHECK3-NEXT:    [[COMP_ADDR:%.*]] = alloca ptr, align 41323// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41324// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41325// CHECK3-NEXT:    store ptr [[COMP]], ptr [[COMP_ADDR]], align 41326// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 41327// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 41328// CHECK3-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP1]], 11329// CHECK3-NEXT:    store i32 [[INC]], ptr [[TMP0]], align 41330// CHECK3-NEXT:    ret void1331//1332//1333// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l371334// CHECK3-SAME: (i32 noundef [[COMP:%.*]]) #[[ATTR1]] {1335// CHECK3-NEXT:  entry:1336// CHECK3-NEXT:    [[COMP_ADDR:%.*]] = alloca i32, align 41337// CHECK3-NEXT:    store i32 [[COMP]], ptr [[COMP_ADDR]], align 41338// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37.omp_outlined, ptr [[COMP_ADDR]])1339// CHECK3-NEXT:    ret void1340//1341//1342// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37.omp_outlined1343// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] {1344// CHECK3-NEXT:  entry:1345// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41346// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41347// CHECK3-NEXT:    [[COMP_ADDR:%.*]] = alloca ptr, align 41348// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41349// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41350// CHECK3-NEXT:    store ptr [[COMP]], ptr [[COMP_ADDR]], align 41351// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 41352// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 41353// CHECK3-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP1]], 11354// CHECK3-NEXT:    store i32 [[INC]], ptr [[TMP0]], align 41355// CHECK3-NEXT:    ret void1356//1357//1358// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l461359// CHECK3-SAME: (i32 noundef [[LA:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1]] {1360// CHECK3-NEXT:  entry:1361// CHECK3-NEXT:    [[LA_ADDR:%.*]] = alloca i32, align 41362// CHECK3-NEXT:    [[COMP_ADDR:%.*]] = alloca i32, align 41363// CHECK3-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1364// CHECK3-NEXT:    store i32 [[LA]], ptr [[LA_ADDR]], align 41365// CHECK3-NEXT:    store i32 [[COMP]], ptr [[COMP_ADDR]], align 41366// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[LA_ADDR]], align 41367// CHECK3-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 0)1368// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46.omp_outlined, ptr [[COMP_ADDR]])1369// CHECK3-NEXT:    ret void1370//1371//1372// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46.omp_outlined1373// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] {1374// CHECK3-NEXT:  entry:1375// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41376// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41377// CHECK3-NEXT:    [[COMP_ADDR:%.*]] = alloca ptr, align 41378// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41379// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41380// CHECK3-NEXT:    store ptr [[COMP]], ptr [[COMP_ADDR]], align 41381// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 41382// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 41383// CHECK3-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP1]], 11384// CHECK3-NEXT:    store i32 [[INC]], ptr [[TMP0]], align 41385// CHECK3-NEXT:    ret void1386//1387//1388// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l531389// CHECK3-SAME: (i32 noundef [[LA:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1]] {1390// CHECK3-NEXT:  entry:1391// CHECK3-NEXT:    [[LA_ADDR:%.*]] = alloca i32, align 41392// CHECK3-NEXT:    [[COMP_ADDR:%.*]] = alloca i32, align 41393// CHECK3-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1394// CHECK3-NEXT:    store i32 [[LA]], ptr [[LA_ADDR]], align 41395// CHECK3-NEXT:    store i32 [[COMP]], ptr [[COMP_ADDR]], align 41396// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[LA_ADDR]], align 41397// CHECK3-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 0, i32 [[TMP1]])1398// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53.omp_outlined, ptr [[COMP_ADDR]])1399// CHECK3-NEXT:    ret void1400//1401//1402// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53.omp_outlined1403// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] {1404// CHECK3-NEXT:  entry:1405// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41406// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41407// CHECK3-NEXT:    [[COMP_ADDR:%.*]] = alloca ptr, align 41408// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41409// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41410// CHECK3-NEXT:    store ptr [[COMP]], ptr [[COMP_ADDR]], align 41411// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 41412// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 41413// CHECK3-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP1]], 11414// CHECK3-NEXT:    store i32 [[INC]], ptr [[TMP0]], align 41415// CHECK3-NEXT:    ret void1416//1417//1418// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l621419// CHECK3-SAME: (i32 noundef [[GBLA:%.*]], i32 noundef [[A:%.*]], ptr noundef nonnull align 4 dereferenceable(8) [[GBLB:%.*]], i32 noundef [[LC:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1]] {1420// CHECK3-NEXT:  entry:1421// CHECK3-NEXT:    [[GBLA_ADDR:%.*]] = alloca i32, align 41422// CHECK3-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 41423// CHECK3-NEXT:    [[GBLB_ADDR:%.*]] = alloca ptr, align 41424// CHECK3-NEXT:    [[LC_ADDR:%.*]] = alloca i32, align 41425// CHECK3-NEXT:    [[COMP_ADDR:%.*]] = alloca i32, align 41426// CHECK3-NEXT:    [[GBLB1:%.*]] = alloca i64, align 81427// CHECK3-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1428// CHECK3-NEXT:    store i32 [[GBLA]], ptr [[GBLA_ADDR]], align 41429// CHECK3-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 41430// CHECK3-NEXT:    store ptr [[GBLB]], ptr [[GBLB_ADDR]], align 41431// CHECK3-NEXT:    store i32 [[LC]], ptr [[LC_ADDR]], align 41432// CHECK3-NEXT:    store i32 [[COMP]], ptr [[COMP_ADDR]], align 41433// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[GBLB_ADDR]], align 41434// CHECK3-NEXT:    [[TMP2:%.*]] = load i64, ptr [[TMP1]], align 81435// CHECK3-NEXT:    store i64 [[TMP2]], ptr [[GBLB1]], align 81436// CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[GBLA_ADDR]], align 41437// CHECK3-NEXT:    [[TMP4:%.*]] = load i32, ptr [[A_ADDR]], align 41438// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP3]], [[TMP4]]1439// CHECK3-NEXT:    [[TMP5:%.*]] = load i64, ptr [[GBLB1]], align 81440// CHECK3-NEXT:    [[TMP6:%.*]] = load float, ptr [[LC_ADDR]], align 41441// CHECK3-NEXT:    [[CONV:%.*]] = fptosi float [[TMP6]] to i641442// CHECK3-NEXT:    [[ADD2:%.*]] = add nsw i64 [[TMP5]], [[CONV]]1443// CHECK3-NEXT:    [[TMP7:%.*]] = trunc i64 [[ADD2]] to i321444// CHECK3-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[ADD]], i32 [[TMP7]])1445// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62.omp_outlined, ptr [[COMP_ADDR]])1446// CHECK3-NEXT:    ret void1447//1448//1449// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62.omp_outlined1450// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] {1451// CHECK3-NEXT:  entry:1452// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41453// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41454// CHECK3-NEXT:    [[COMP_ADDR:%.*]] = alloca ptr, align 41455// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41456// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41457// CHECK3-NEXT:    store ptr [[COMP]], ptr [[COMP_ADDR]], align 41458// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 41459// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 41460// CHECK3-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP1]], 11461// CHECK3-NEXT:    store i32 [[INC]], ptr [[TMP0]], align 41462// CHECK3-NEXT:    ret void1463//1464//1465// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l711466// CHECK3-SAME: (i32 noundef [[GBLC:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1]] {1467// CHECK3-NEXT:  entry:1468// CHECK3-NEXT:    [[GBLC_ADDR:%.*]] = alloca i32, align 41469// CHECK3-NEXT:    [[COMP_ADDR:%.*]] = alloca i32, align 41470// CHECK3-NEXT:    [[TMP:%.*]] = alloca ptr, align 41471// CHECK3-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1472// CHECK3-NEXT:    store i32 [[GBLC]], ptr [[GBLC_ADDR]], align 41473// CHECK3-NEXT:    store i32 [[COMP]], ptr [[COMP_ADDR]], align 41474// CHECK3-NEXT:    store ptr [[GBLC_ADDR]], ptr [[TMP]], align 41475// CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr @Gbla, align 41476// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], 11477// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr @Gbla, align 41478// CHECK3-NEXT:    [[ADD1:%.*]] = add nsw i32 [[TMP2]], 21479// CHECK3-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[ADD]], i32 [[ADD1]])1480// CHECK3-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[TMP]], align 41481// CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71.omp_outlined, ptr [[COMP_ADDR]], ptr [[TMP3]])1482// CHECK3-NEXT:    ret void1483//1484//1485// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71.omp_outlined1486// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[GBLC:%.*]]) #[[ATTR1]] {1487// CHECK3-NEXT:  entry:1488// CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41489// CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41490// CHECK3-NEXT:    [[COMP_ADDR:%.*]] = alloca ptr, align 41491// CHECK3-NEXT:    [[GBLC_ADDR:%.*]] = alloca ptr, align 41492// CHECK3-NEXT:    [[TMP:%.*]] = alloca ptr, align 41493// CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41494// CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41495// CHECK3-NEXT:    store ptr [[COMP]], ptr [[COMP_ADDR]], align 41496// CHECK3-NEXT:    store ptr [[GBLC]], ptr [[GBLC_ADDR]], align 41497// CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 41498// CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[GBLC_ADDR]], align 41499// CHECK3-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 41500// CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr @Gbla, align 41501// CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 41502// CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP3]], [[TMP2]]1503// CHECK3-NEXT:    store i32 [[ADD]], ptr [[TMP0]], align 41504// CHECK3-NEXT:    ret void1505//1506//1507// CHECK9-LABEL: define {{[^@]+}}@_Z18teams_template_argv1508// CHECK9-SAME: () #[[ATTR0:[0-9]+]] {1509// CHECK9-NEXT:  entry:1510// CHECK9-NEXT:    [[COMP:%.*]] = alloca i32, align 41511// CHECK9-NEXT:    [[LA:%.*]] = alloca [[STRUCT_SS:%.*]], align 41512// CHECK9-NEXT:    [[LB:%.*]] = alloca [[STRUCT_SS_0:%.*]], align 81513// CHECK9-NEXT:    [[COMP_CASTED:%.*]] = alloca i64, align 81514// CHECK9-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 81515// CHECK9-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 81516// CHECK9-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 81517// CHECK9-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81518// CHECK9-NEXT:    [[COMP_CASTED1:%.*]] = alloca i64, align 81519// CHECK9-NEXT:    [[DOTOFFLOAD_BASEPTRS2:%.*]] = alloca [3 x ptr], align 81520// CHECK9-NEXT:    [[DOTOFFLOAD_PTRS3:%.*]] = alloca [3 x ptr], align 81521// CHECK9-NEXT:    [[DOTOFFLOAD_MAPPERS4:%.*]] = alloca [3 x ptr], align 81522// CHECK9-NEXT:    [[KERNEL_ARGS7:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 81523// CHECK9-NEXT:    store i32 1, ptr [[COMP]], align 41524// CHECK9-NEXT:    [[TMP0:%.*]] = load i32, ptr [[COMP]], align 41525// CHECK9-NEXT:    store i32 [[TMP0]], ptr [[COMP_CASTED]], align 41526// CHECK9-NEXT:    [[TMP1:%.*]] = load i64, ptr [[COMP_CASTED]], align 81527// CHECK9-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01528// CHECK9-NEXT:    store ptr @Gbla, ptr [[TMP2]], align 81529// CHECK9-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01530// CHECK9-NEXT:    store ptr @Gbla, ptr [[TMP3]], align 81531// CHECK9-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 01532// CHECK9-NEXT:    store ptr null, ptr [[TMP4]], align 81533// CHECK9-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11534// CHECK9-NEXT:    store ptr [[LA]], ptr [[TMP5]], align 81535// CHECK9-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11536// CHECK9-NEXT:    store ptr [[LA]], ptr [[TMP6]], align 81537// CHECK9-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 11538// CHECK9-NEXT:    store ptr null, ptr [[TMP7]], align 81539// CHECK9-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21540// CHECK9-NEXT:    store i64 [[TMP1]], ptr [[TMP8]], align 81541// CHECK9-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21542// CHECK9-NEXT:    store i64 [[TMP1]], ptr [[TMP9]], align 81543// CHECK9-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 21544// CHECK9-NEXT:    store ptr null, ptr [[TMP10]], align 81545// CHECK9-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01546// CHECK9-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01547// CHECK9-NEXT:    [[TMP13:%.*]] = load i32, ptr @Gbla, align 41548// CHECK9-NEXT:    [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[LA]], i32 0, i32 11549// CHECK9-NEXT:    [[TMP14:%.*]] = load float, ptr [[B]], align 41550// CHECK9-NEXT:    [[CONV:%.*]] = fptosi float [[TMP14]] to i641551// CHECK9-NEXT:    [[TMP15:%.*]] = trunc i64 [[CONV]] to i321552// CHECK9-NEXT:    [[TMP16:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP13]], 01553// CHECK9-NEXT:    [[TMP17:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP15]], 01554// CHECK9-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01555// CHECK9-NEXT:    store i32 3, ptr [[TMP18]], align 41556// CHECK9-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11557// CHECK9-NEXT:    store i32 3, ptr [[TMP19]], align 41558// CHECK9-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21559// CHECK9-NEXT:    store ptr [[TMP11]], ptr [[TMP20]], align 81560// CHECK9-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31561// CHECK9-NEXT:    store ptr [[TMP12]], ptr [[TMP21]], align 81562// CHECK9-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41563// CHECK9-NEXT:    store ptr @.offload_sizes, ptr [[TMP22]], align 81564// CHECK9-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51565// CHECK9-NEXT:    store ptr @.offload_maptypes, ptr [[TMP23]], align 81566// CHECK9-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61567// CHECK9-NEXT:    store ptr null, ptr [[TMP24]], align 81568// CHECK9-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71569// CHECK9-NEXT:    store ptr null, ptr [[TMP25]], align 81570// CHECK9-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81571// CHECK9-NEXT:    store i64 0, ptr [[TMP26]], align 81572// CHECK9-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91573// CHECK9-NEXT:    store i64 0, ptr [[TMP27]], align 81574// CHECK9-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101575// CHECK9-NEXT:    store [3 x i32] [[TMP16]], ptr [[TMP28]], align 41576// CHECK9-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111577// CHECK9-NEXT:    store [3 x i32] [[TMP17]], ptr [[TMP29]], align 41578// CHECK9-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121579// CHECK9-NEXT:    store i32 0, ptr [[TMP30]], align 41580// CHECK9-NEXT:    [[TMP31:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 [[TMP13]], i32 [[TMP15]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116.region_id, ptr [[KERNEL_ARGS]])1581// CHECK9-NEXT:    [[TMP32:%.*]] = icmp ne i32 [[TMP31]], 01582// CHECK9-NEXT:    br i1 [[TMP32]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1583// CHECK9:       omp_offload.failed:1584// CHECK9-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116(ptr @Gbla, ptr [[LA]], i64 [[TMP1]]) #[[ATTR2:[0-9]+]]1585// CHECK9-NEXT:    br label [[OMP_OFFLOAD_CONT]]1586// CHECK9:       omp_offload.cont:1587// CHECK9-NEXT:    [[TMP33:%.*]] = load i32, ptr [[COMP]], align 41588// CHECK9-NEXT:    store i32 [[TMP33]], ptr [[COMP_CASTED1]], align 41589// CHECK9-NEXT:    [[TMP34:%.*]] = load i64, ptr [[COMP_CASTED1]], align 81590// CHECK9-NEXT:    [[TMP35:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 01591// CHECK9-NEXT:    store ptr [[LB]], ptr [[TMP35]], align 81592// CHECK9-NEXT:    [[TMP36:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 01593// CHECK9-NEXT:    store ptr [[LB]], ptr [[TMP36]], align 81594// CHECK9-NEXT:    [[TMP37:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i64 0, i64 01595// CHECK9-NEXT:    store ptr null, ptr [[TMP37]], align 81596// CHECK9-NEXT:    [[TMP38:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 11597// CHECK9-NEXT:    store ptr @Gblb, ptr [[TMP38]], align 81598// CHECK9-NEXT:    [[TMP39:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 11599// CHECK9-NEXT:    store ptr @Gblb, ptr [[TMP39]], align 81600// CHECK9-NEXT:    [[TMP40:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i64 0, i64 11601// CHECK9-NEXT:    store ptr null, ptr [[TMP40]], align 81602// CHECK9-NEXT:    [[TMP41:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 21603// CHECK9-NEXT:    store i64 [[TMP34]], ptr [[TMP41]], align 81604// CHECK9-NEXT:    [[TMP42:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 21605// CHECK9-NEXT:    store i64 [[TMP34]], ptr [[TMP42]], align 81606// CHECK9-NEXT:    [[TMP43:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i64 0, i64 21607// CHECK9-NEXT:    store ptr null, ptr [[TMP43]], align 81608// CHECK9-NEXT:    [[TMP44:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 01609// CHECK9-NEXT:    [[TMP45:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 01610// CHECK9-NEXT:    [[B5:%.*]] = getelementptr inbounds nuw [[STRUCT_SS_0]], ptr [[LB]], i32 0, i32 11611// CHECK9-NEXT:    [[TMP46:%.*]] = load float, ptr [[B5]], align 81612// CHECK9-NEXT:    [[CONV6:%.*]] = fptosi float [[TMP46]] to i641613// CHECK9-NEXT:    [[TMP47:%.*]] = trunc i64 [[CONV6]] to i321614// CHECK9-NEXT:    [[TMP48:%.*]] = load i64, ptr @Gblb, align 81615// CHECK9-NEXT:    [[TMP49:%.*]] = trunc i64 [[TMP48]] to i321616// CHECK9-NEXT:    [[TMP50:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP47]], 01617// CHECK9-NEXT:    [[TMP51:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP49]], 01618// CHECK9-NEXT:    [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 01619// CHECK9-NEXT:    store i32 3, ptr [[TMP52]], align 41620// CHECK9-NEXT:    [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 11621// CHECK9-NEXT:    store i32 3, ptr [[TMP53]], align 41622// CHECK9-NEXT:    [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 21623// CHECK9-NEXT:    store ptr [[TMP44]], ptr [[TMP54]], align 81624// CHECK9-NEXT:    [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 31625// CHECK9-NEXT:    store ptr [[TMP45]], ptr [[TMP55]], align 81626// CHECK9-NEXT:    [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 41627// CHECK9-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP56]], align 81628// CHECK9-NEXT:    [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 51629// CHECK9-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP57]], align 81630// CHECK9-NEXT:    [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 61631// CHECK9-NEXT:    store ptr null, ptr [[TMP58]], align 81632// CHECK9-NEXT:    [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 71633// CHECK9-NEXT:    store ptr null, ptr [[TMP59]], align 81634// CHECK9-NEXT:    [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 81635// CHECK9-NEXT:    store i64 0, ptr [[TMP60]], align 81636// CHECK9-NEXT:    [[TMP61:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 91637// CHECK9-NEXT:    store i64 0, ptr [[TMP61]], align 81638// CHECK9-NEXT:    [[TMP62:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 101639// CHECK9-NEXT:    store [3 x i32] [[TMP50]], ptr [[TMP62]], align 41640// CHECK9-NEXT:    [[TMP63:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 111641// CHECK9-NEXT:    store [3 x i32] [[TMP51]], ptr [[TMP63]], align 41642// CHECK9-NEXT:    [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 121643// CHECK9-NEXT:    store i32 0, ptr [[TMP64]], align 41644// CHECK9-NEXT:    [[TMP65:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[TMP47]], i32 [[TMP49]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125.region_id, ptr [[KERNEL_ARGS7]])1645// CHECK9-NEXT:    [[TMP66:%.*]] = icmp ne i32 [[TMP65]], 01646// CHECK9-NEXT:    br i1 [[TMP66]], label [[OMP_OFFLOAD_FAILED8:%.*]], label [[OMP_OFFLOAD_CONT9:%.*]]1647// CHECK9:       omp_offload.failed8:1648// CHECK9-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125(ptr [[LB]], ptr @Gblb, i64 [[TMP34]]) #[[ATTR2]]1649// CHECK9-NEXT:    br label [[OMP_OFFLOAD_CONT9]]1650// CHECK9:       omp_offload.cont9:1651// CHECK9-NEXT:    [[TMP67:%.*]] = load i32, ptr [[COMP]], align 41652// CHECK9-NEXT:    ret i32 [[TMP67]]1653//1654//1655// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l1161656// CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[GBLA:%.*]], ptr noundef nonnull align 4 dereferenceable(8) [[LA:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1:[0-9]+]] {1657// CHECK9-NEXT:  entry:1658// CHECK9-NEXT:    [[GBLA_ADDR:%.*]] = alloca ptr, align 81659// CHECK9-NEXT:    [[LA_ADDR:%.*]] = alloca ptr, align 81660// CHECK9-NEXT:    [[COMP_ADDR:%.*]] = alloca i64, align 81661// CHECK9-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1662// CHECK9-NEXT:    store ptr [[GBLA]], ptr [[GBLA_ADDR]], align 81663// CHECK9-NEXT:    store ptr [[LA]], ptr [[LA_ADDR]], align 81664// CHECK9-NEXT:    store i64 [[COMP]], ptr [[COMP_ADDR]], align 81665// CHECK9-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[GBLA_ADDR]], align 81666// CHECK9-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[LA_ADDR]], align 81667// CHECK9-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP1]], i32 0, i32 01668// CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A]], align 41669// CHECK9-NEXT:    [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[TMP2]], i32 0, i32 11670// CHECK9-NEXT:    [[TMP4:%.*]] = load float, ptr [[B]], align 41671// CHECK9-NEXT:    [[CONV:%.*]] = fptosi float [[TMP4]] to i641672// CHECK9-NEXT:    [[TMP5:%.*]] = trunc i64 [[CONV]] to i321673// CHECK9-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP3]], i32 [[TMP5]])1674// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116.omp_outlined, ptr [[COMP_ADDR]])1675// CHECK9-NEXT:    ret void1676//1677//1678// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116.omp_outlined1679// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] {1680// CHECK9-NEXT:  entry:1681// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81682// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81683// CHECK9-NEXT:    [[COMP_ADDR:%.*]] = alloca ptr, align 81684// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81685// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81686// CHECK9-NEXT:    store ptr [[COMP]], ptr [[COMP_ADDR]], align 81687// CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 81688// CHECK9-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 41689// CHECK9-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP1]], 11690// CHECK9-NEXT:    store i32 [[INC]], ptr [[TMP0]], align 41691// CHECK9-NEXT:    ret void1692//1693//1694// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l1251695// CHECK9-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[LB:%.*]], ptr noundef nonnull align 8 dereferenceable(16) [[GBLB:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1]] {1696// CHECK9-NEXT:  entry:1697// CHECK9-NEXT:    [[LB_ADDR:%.*]] = alloca ptr, align 81698// CHECK9-NEXT:    [[GBLB_ADDR:%.*]] = alloca ptr, align 81699// CHECK9-NEXT:    [[COMP_ADDR:%.*]] = alloca i64, align 81700// CHECK9-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1701// CHECK9-NEXT:    store ptr [[LB]], ptr [[LB_ADDR]], align 81702// CHECK9-NEXT:    store ptr [[GBLB]], ptr [[GBLB_ADDR]], align 81703// CHECK9-NEXT:    store i64 [[COMP]], ptr [[COMP_ADDR]], align 81704// CHECK9-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[LB_ADDR]], align 81705// CHECK9-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[GBLB_ADDR]], align 81706// CHECK9-NEXT:    [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS_0:%.*]], ptr [[TMP1]], i32 0, i32 11707// CHECK9-NEXT:    [[TMP3:%.*]] = load float, ptr [[B]], align 81708// CHECK9-NEXT:    [[CONV:%.*]] = fptosi float [[TMP3]] to i641709// CHECK9-NEXT:    [[TMP4:%.*]] = trunc i64 [[CONV]] to i321710// CHECK9-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS_0]], ptr [[TMP2]], i32 0, i32 01711// CHECK9-NEXT:    [[TMP5:%.*]] = load i64, ptr [[A]], align 81712// CHECK9-NEXT:    [[TMP6:%.*]] = trunc i64 [[TMP5]] to i321713// CHECK9-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP4]], i32 [[TMP6]])1714// CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125.omp_outlined, ptr [[COMP_ADDR]])1715// CHECK9-NEXT:    ret void1716//1717//1718// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125.omp_outlined1719// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] {1720// CHECK9-NEXT:  entry:1721// CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81722// CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81723// CHECK9-NEXT:    [[COMP_ADDR:%.*]] = alloca ptr, align 81724// CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81725// CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81726// CHECK9-NEXT:    store ptr [[COMP]], ptr [[COMP_ADDR]], align 81727// CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 81728// CHECK9-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 41729// CHECK9-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP1]], 11730// CHECK9-NEXT:    store i32 [[INC]], ptr [[TMP0]], align 41731// CHECK9-NEXT:    ret void1732//1733//1734// CHECK11-LABEL: define {{[^@]+}}@_Z18teams_template_argv1735// CHECK11-SAME: () #[[ATTR0:[0-9]+]] {1736// CHECK11-NEXT:  entry:1737// CHECK11-NEXT:    [[COMP:%.*]] = alloca i32, align 41738// CHECK11-NEXT:    [[LA:%.*]] = alloca [[STRUCT_SS:%.*]], align 41739// CHECK11-NEXT:    [[LB:%.*]] = alloca [[STRUCT_SS_0:%.*]], align 41740// CHECK11-NEXT:    [[COMP_CASTED:%.*]] = alloca i32, align 41741// CHECK11-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 41742// CHECK11-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 41743// CHECK11-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 41744// CHECK11-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81745// CHECK11-NEXT:    [[COMP_CASTED1:%.*]] = alloca i32, align 41746// CHECK11-NEXT:    [[DOTOFFLOAD_BASEPTRS2:%.*]] = alloca [3 x ptr], align 41747// CHECK11-NEXT:    [[DOTOFFLOAD_PTRS3:%.*]] = alloca [3 x ptr], align 41748// CHECK11-NEXT:    [[DOTOFFLOAD_MAPPERS4:%.*]] = alloca [3 x ptr], align 41749// CHECK11-NEXT:    [[KERNEL_ARGS7:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 81750// CHECK11-NEXT:    store i32 1, ptr [[COMP]], align 41751// CHECK11-NEXT:    [[TMP0:%.*]] = load i32, ptr [[COMP]], align 41752// CHECK11-NEXT:    store i32 [[TMP0]], ptr [[COMP_CASTED]], align 41753// CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[COMP_CASTED]], align 41754// CHECK11-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01755// CHECK11-NEXT:    store ptr @Gbla, ptr [[TMP2]], align 41756// CHECK11-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01757// CHECK11-NEXT:    store ptr @Gbla, ptr [[TMP3]], align 41758// CHECK11-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01759// CHECK11-NEXT:    store ptr null, ptr [[TMP4]], align 41760// CHECK11-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11761// CHECK11-NEXT:    store ptr [[LA]], ptr [[TMP5]], align 41762// CHECK11-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11763// CHECK11-NEXT:    store ptr [[LA]], ptr [[TMP6]], align 41764// CHECK11-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 11765// CHECK11-NEXT:    store ptr null, ptr [[TMP7]], align 41766// CHECK11-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21767// CHECK11-NEXT:    store i32 [[TMP1]], ptr [[TMP8]], align 41768// CHECK11-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21769// CHECK11-NEXT:    store i32 [[TMP1]], ptr [[TMP9]], align 41770// CHECK11-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 21771// CHECK11-NEXT:    store ptr null, ptr [[TMP10]], align 41772// CHECK11-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01773// CHECK11-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01774// CHECK11-NEXT:    [[TMP13:%.*]] = load i32, ptr @Gbla, align 41775// CHECK11-NEXT:    [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[LA]], i32 0, i32 11776// CHECK11-NEXT:    [[TMP14:%.*]] = load float, ptr [[B]], align 41777// CHECK11-NEXT:    [[CONV:%.*]] = fptosi float [[TMP14]] to i641778// CHECK11-NEXT:    [[TMP15:%.*]] = trunc i64 [[CONV]] to i321779// CHECK11-NEXT:    [[TMP16:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP13]], 01780// CHECK11-NEXT:    [[TMP17:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP15]], 01781// CHECK11-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01782// CHECK11-NEXT:    store i32 3, ptr [[TMP18]], align 41783// CHECK11-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11784// CHECK11-NEXT:    store i32 3, ptr [[TMP19]], align 41785// CHECK11-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21786// CHECK11-NEXT:    store ptr [[TMP11]], ptr [[TMP20]], align 41787// CHECK11-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31788// CHECK11-NEXT:    store ptr [[TMP12]], ptr [[TMP21]], align 41789// CHECK11-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41790// CHECK11-NEXT:    store ptr @.offload_sizes, ptr [[TMP22]], align 41791// CHECK11-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51792// CHECK11-NEXT:    store ptr @.offload_maptypes, ptr [[TMP23]], align 41793// CHECK11-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61794// CHECK11-NEXT:    store ptr null, ptr [[TMP24]], align 41795// CHECK11-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71796// CHECK11-NEXT:    store ptr null, ptr [[TMP25]], align 41797// CHECK11-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81798// CHECK11-NEXT:    store i64 0, ptr [[TMP26]], align 81799// CHECK11-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91800// CHECK11-NEXT:    store i64 0, ptr [[TMP27]], align 81801// CHECK11-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101802// CHECK11-NEXT:    store [3 x i32] [[TMP16]], ptr [[TMP28]], align 41803// CHECK11-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111804// CHECK11-NEXT:    store [3 x i32] [[TMP17]], ptr [[TMP29]], align 41805// CHECK11-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121806// CHECK11-NEXT:    store i32 0, ptr [[TMP30]], align 41807// CHECK11-NEXT:    [[TMP31:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 [[TMP13]], i32 [[TMP15]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116.region_id, ptr [[KERNEL_ARGS]])1808// CHECK11-NEXT:    [[TMP32:%.*]] = icmp ne i32 [[TMP31]], 01809// CHECK11-NEXT:    br i1 [[TMP32]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1810// CHECK11:       omp_offload.failed:1811// CHECK11-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116(ptr @Gbla, ptr [[LA]], i32 [[TMP1]]) #[[ATTR2:[0-9]+]]1812// CHECK11-NEXT:    br label [[OMP_OFFLOAD_CONT]]1813// CHECK11:       omp_offload.cont:1814// CHECK11-NEXT:    [[TMP33:%.*]] = load i32, ptr [[COMP]], align 41815// CHECK11-NEXT:    store i32 [[TMP33]], ptr [[COMP_CASTED1]], align 41816// CHECK11-NEXT:    [[TMP34:%.*]] = load i32, ptr [[COMP_CASTED1]], align 41817// CHECK11-NEXT:    [[TMP35:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 01818// CHECK11-NEXT:    store ptr [[LB]], ptr [[TMP35]], align 41819// CHECK11-NEXT:    [[TMP36:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 01820// CHECK11-NEXT:    store ptr [[LB]], ptr [[TMP36]], align 41821// CHECK11-NEXT:    [[TMP37:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i32 0, i32 01822// CHECK11-NEXT:    store ptr null, ptr [[TMP37]], align 41823// CHECK11-NEXT:    [[TMP38:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 11824// CHECK11-NEXT:    store ptr @Gblb, ptr [[TMP38]], align 41825// CHECK11-NEXT:    [[TMP39:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 11826// CHECK11-NEXT:    store ptr @Gblb, ptr [[TMP39]], align 41827// CHECK11-NEXT:    [[TMP40:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i32 0, i32 11828// CHECK11-NEXT:    store ptr null, ptr [[TMP40]], align 41829// CHECK11-NEXT:    [[TMP41:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 21830// CHECK11-NEXT:    store i32 [[TMP34]], ptr [[TMP41]], align 41831// CHECK11-NEXT:    [[TMP42:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 21832// CHECK11-NEXT:    store i32 [[TMP34]], ptr [[TMP42]], align 41833// CHECK11-NEXT:    [[TMP43:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i32 0, i32 21834// CHECK11-NEXT:    store ptr null, ptr [[TMP43]], align 41835// CHECK11-NEXT:    [[TMP44:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 01836// CHECK11-NEXT:    [[TMP45:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 01837// CHECK11-NEXT:    [[B5:%.*]] = getelementptr inbounds nuw [[STRUCT_SS_0]], ptr [[LB]], i32 0, i32 11838// CHECK11-NEXT:    [[TMP46:%.*]] = load float, ptr [[B5]], align 41839// CHECK11-NEXT:    [[CONV6:%.*]] = fptosi float [[TMP46]] to i641840// CHECK11-NEXT:    [[TMP47:%.*]] = trunc i64 [[CONV6]] to i321841// CHECK11-NEXT:    [[TMP48:%.*]] = load i64, ptr @Gblb, align 41842// CHECK11-NEXT:    [[TMP49:%.*]] = trunc i64 [[TMP48]] to i321843// CHECK11-NEXT:    [[TMP50:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP47]], 01844// CHECK11-NEXT:    [[TMP51:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP49]], 01845// CHECK11-NEXT:    [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 01846// CHECK11-NEXT:    store i32 3, ptr [[TMP52]], align 41847// CHECK11-NEXT:    [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 11848// CHECK11-NEXT:    store i32 3, ptr [[TMP53]], align 41849// CHECK11-NEXT:    [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 21850// CHECK11-NEXT:    store ptr [[TMP44]], ptr [[TMP54]], align 41851// CHECK11-NEXT:    [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 31852// CHECK11-NEXT:    store ptr [[TMP45]], ptr [[TMP55]], align 41853// CHECK11-NEXT:    [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 41854// CHECK11-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP56]], align 41855// CHECK11-NEXT:    [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 51856// CHECK11-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP57]], align 41857// CHECK11-NEXT:    [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 61858// CHECK11-NEXT:    store ptr null, ptr [[TMP58]], align 41859// CHECK11-NEXT:    [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 71860// CHECK11-NEXT:    store ptr null, ptr [[TMP59]], align 41861// CHECK11-NEXT:    [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 81862// CHECK11-NEXT:    store i64 0, ptr [[TMP60]], align 81863// CHECK11-NEXT:    [[TMP61:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 91864// CHECK11-NEXT:    store i64 0, ptr [[TMP61]], align 81865// CHECK11-NEXT:    [[TMP62:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 101866// CHECK11-NEXT:    store [3 x i32] [[TMP50]], ptr [[TMP62]], align 41867// CHECK11-NEXT:    [[TMP63:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 111868// CHECK11-NEXT:    store [3 x i32] [[TMP51]], ptr [[TMP63]], align 41869// CHECK11-NEXT:    [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 121870// CHECK11-NEXT:    store i32 0, ptr [[TMP64]], align 41871// CHECK11-NEXT:    [[TMP65:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[TMP47]], i32 [[TMP49]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125.region_id, ptr [[KERNEL_ARGS7]])1872// CHECK11-NEXT:    [[TMP66:%.*]] = icmp ne i32 [[TMP65]], 01873// CHECK11-NEXT:    br i1 [[TMP66]], label [[OMP_OFFLOAD_FAILED8:%.*]], label [[OMP_OFFLOAD_CONT9:%.*]]1874// CHECK11:       omp_offload.failed8:1875// CHECK11-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125(ptr [[LB]], ptr @Gblb, i32 [[TMP34]]) #[[ATTR2]]1876// CHECK11-NEXT:    br label [[OMP_OFFLOAD_CONT9]]1877// CHECK11:       omp_offload.cont9:1878// CHECK11-NEXT:    [[TMP67:%.*]] = load i32, ptr [[COMP]], align 41879// CHECK11-NEXT:    ret i32 [[TMP67]]1880//1881//1882// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l1161883// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[GBLA:%.*]], ptr noundef nonnull align 4 dereferenceable(8) [[LA:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1:[0-9]+]] {1884// CHECK11-NEXT:  entry:1885// CHECK11-NEXT:    [[GBLA_ADDR:%.*]] = alloca ptr, align 41886// CHECK11-NEXT:    [[LA_ADDR:%.*]] = alloca ptr, align 41887// CHECK11-NEXT:    [[COMP_ADDR:%.*]] = alloca i32, align 41888// CHECK11-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1889// CHECK11-NEXT:    store ptr [[GBLA]], ptr [[GBLA_ADDR]], align 41890// CHECK11-NEXT:    store ptr [[LA]], ptr [[LA_ADDR]], align 41891// CHECK11-NEXT:    store i32 [[COMP]], ptr [[COMP_ADDR]], align 41892// CHECK11-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[GBLA_ADDR]], align 41893// CHECK11-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[LA_ADDR]], align 41894// CHECK11-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP1]], i32 0, i32 01895// CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A]], align 41896// CHECK11-NEXT:    [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[TMP2]], i32 0, i32 11897// CHECK11-NEXT:    [[TMP4:%.*]] = load float, ptr [[B]], align 41898// CHECK11-NEXT:    [[CONV:%.*]] = fptosi float [[TMP4]] to i641899// CHECK11-NEXT:    [[TMP5:%.*]] = trunc i64 [[CONV]] to i321900// CHECK11-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP3]], i32 [[TMP5]])1901// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116.omp_outlined, ptr [[COMP_ADDR]])1902// CHECK11-NEXT:    ret void1903//1904//1905// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116.omp_outlined1906// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] {1907// CHECK11-NEXT:  entry:1908// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41909// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41910// CHECK11-NEXT:    [[COMP_ADDR:%.*]] = alloca ptr, align 41911// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41912// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41913// CHECK11-NEXT:    store ptr [[COMP]], ptr [[COMP_ADDR]], align 41914// CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 41915// CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 41916// CHECK11-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP1]], 11917// CHECK11-NEXT:    store i32 [[INC]], ptr [[TMP0]], align 41918// CHECK11-NEXT:    ret void1919//1920//1921// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l1251922// CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(12) [[LB:%.*]], ptr noundef nonnull align 4 dereferenceable(12) [[GBLB:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1]] {1923// CHECK11-NEXT:  entry:1924// CHECK11-NEXT:    [[LB_ADDR:%.*]] = alloca ptr, align 41925// CHECK11-NEXT:    [[GBLB_ADDR:%.*]] = alloca ptr, align 41926// CHECK11-NEXT:    [[COMP_ADDR:%.*]] = alloca i32, align 41927// CHECK11-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1928// CHECK11-NEXT:    store ptr [[LB]], ptr [[LB_ADDR]], align 41929// CHECK11-NEXT:    store ptr [[GBLB]], ptr [[GBLB_ADDR]], align 41930// CHECK11-NEXT:    store i32 [[COMP]], ptr [[COMP_ADDR]], align 41931// CHECK11-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[LB_ADDR]], align 41932// CHECK11-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[GBLB_ADDR]], align 41933// CHECK11-NEXT:    [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS_0:%.*]], ptr [[TMP1]], i32 0, i32 11934// CHECK11-NEXT:    [[TMP3:%.*]] = load float, ptr [[B]], align 41935// CHECK11-NEXT:    [[CONV:%.*]] = fptosi float [[TMP3]] to i641936// CHECK11-NEXT:    [[TMP4:%.*]] = trunc i64 [[CONV]] to i321937// CHECK11-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS_0]], ptr [[TMP2]], i32 0, i32 01938// CHECK11-NEXT:    [[TMP5:%.*]] = load i64, ptr [[A]], align 41939// CHECK11-NEXT:    [[TMP6:%.*]] = trunc i64 [[TMP5]] to i321940// CHECK11-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP4]], i32 [[TMP6]])1941// CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125.omp_outlined, ptr [[COMP_ADDR]])1942// CHECK11-NEXT:    ret void1943//1944//1945// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125.omp_outlined1946// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] {1947// CHECK11-NEXT:  entry:1948// CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41949// CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41950// CHECK11-NEXT:    [[COMP_ADDR:%.*]] = alloca ptr, align 41951// CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41952// CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41953// CHECK11-NEXT:    store ptr [[COMP]], ptr [[COMP_ADDR]], align 41954// CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 41955// CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 41956// CHECK11-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP1]], 11957// CHECK11-NEXT:    store i32 [[INC]], ptr [[TMP0]], align 41958// CHECK11-NEXT:    ret void1959//1960//1961// CHECK17-LABEL: define {{[^@]+}}@_Z21teams_template_structv1962// CHECK17-SAME: () #[[ATTR0:[0-9]+]] {1963// CHECK17-NEXT:  entry:1964// CHECK17-NEXT:    [[V:%.*]] = alloca [[STRUCT_SS:%.*]], align 41965// CHECK17-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_ZN2SSIiLi123ELx456EE3fooEv(ptr noundef nonnull align 4 dereferenceable(8) [[V]])1966// CHECK17-NEXT:    ret i32 [[CALL]]1967//1968//1969// CHECK17-LABEL: define {{[^@]+}}@_ZN2SSIiLi123ELx456EE3fooEv1970// CHECK17-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]]) #[[ATTR0]] comdat {1971// CHECK17-NEXT:  entry:1972// CHECK17-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 81973// CHECK17-NEXT:    [[COMP:%.*]] = alloca i32, align 41974// CHECK17-NEXT:    [[COMP_CASTED:%.*]] = alloca i64, align 81975// CHECK17-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x ptr], align 81976// CHECK17-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x ptr], align 81977// CHECK17-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [2 x ptr], align 81978// CHECK17-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81979// CHECK17-NEXT:    [[COMP_CASTED3:%.*]] = alloca i64, align 81980// CHECK17-NEXT:    [[DOTOFFLOAD_BASEPTRS4:%.*]] = alloca [2 x ptr], align 81981// CHECK17-NEXT:    [[DOTOFFLOAD_PTRS5:%.*]] = alloca [2 x ptr], align 81982// CHECK17-NEXT:    [[DOTOFFLOAD_MAPPERS6:%.*]] = alloca [2 x ptr], align 81983// CHECK17-NEXT:    [[KERNEL_ARGS8:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 81984// CHECK17-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 81985// CHECK17-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81986// CHECK17-NEXT:    store i32 1, ptr [[COMP]], align 41987// CHECK17-NEXT:    [[TMP0:%.*]] = load i32, ptr [[COMP]], align 41988// CHECK17-NEXT:    store i32 [[TMP0]], ptr [[COMP_CASTED]], align 41989// CHECK17-NEXT:    [[TMP1:%.*]] = load i64, ptr [[COMP_CASTED]], align 81990// CHECK17-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[THIS1]], i32 0, i32 01991// CHECK17-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01992// CHECK17-NEXT:    store ptr [[THIS1]], ptr [[TMP2]], align 81993// CHECK17-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01994// CHECK17-NEXT:    store ptr [[A]], ptr [[TMP3]], align 81995// CHECK17-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 01996// CHECK17-NEXT:    store ptr null, ptr [[TMP4]], align 81997// CHECK17-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11998// CHECK17-NEXT:    store i64 [[TMP1]], ptr [[TMP5]], align 81999// CHECK17-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 12000// CHECK17-NEXT:    store i64 [[TMP1]], ptr [[TMP6]], align 82001// CHECK17-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 12002// CHECK17-NEXT:    store ptr null, ptr [[TMP7]], align 82003// CHECK17-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 02004// CHECK17-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 02005// CHECK17-NEXT:    [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 02006// CHECK17-NEXT:    [[TMP10:%.*]] = load i32, ptr [[A2]], align 42007// CHECK17-NEXT:    [[TMP11:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP10]], 02008// CHECK17-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 02009// CHECK17-NEXT:    store i32 3, ptr [[TMP12]], align 42010// CHECK17-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12011// CHECK17-NEXT:    store i32 2, ptr [[TMP13]], align 42012// CHECK17-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 22013// CHECK17-NEXT:    store ptr [[TMP8]], ptr [[TMP14]], align 82014// CHECK17-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 32015// CHECK17-NEXT:    store ptr [[TMP9]], ptr [[TMP15]], align 82016// CHECK17-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 42017// CHECK17-NEXT:    store ptr @.offload_sizes, ptr [[TMP16]], align 82018// CHECK17-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 52019// CHECK17-NEXT:    store ptr @.offload_maptypes, ptr [[TMP17]], align 82020// CHECK17-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 62021// CHECK17-NEXT:    store ptr null, ptr [[TMP18]], align 82022// CHECK17-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 72023// CHECK17-NEXT:    store ptr null, ptr [[TMP19]], align 82024// CHECK17-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 82025// CHECK17-NEXT:    store i64 0, ptr [[TMP20]], align 82026// CHECK17-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 92027// CHECK17-NEXT:    store i64 0, ptr [[TMP21]], align 82028// CHECK17-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 102029// CHECK17-NEXT:    store [3 x i32] [[TMP11]], ptr [[TMP22]], align 42030// CHECK17-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 112031// CHECK17-NEXT:    store [3 x i32] [i32 123, i32 0, i32 0], ptr [[TMP23]], align 42032// CHECK17-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 122033// CHECK17-NEXT:    store i32 0, ptr [[TMP24]], align 42034// CHECK17-NEXT:    [[TMP25:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 [[TMP10]], i32 123, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161.region_id, ptr [[KERNEL_ARGS]])2035// CHECK17-NEXT:    [[TMP26:%.*]] = icmp ne i32 [[TMP25]], 02036// CHECK17-NEXT:    br i1 [[TMP26]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]2037// CHECK17:       omp_offload.failed:2038// CHECK17-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161(ptr [[THIS1]], i64 [[TMP1]]) #[[ATTR2:[0-9]+]]2039// CHECK17-NEXT:    br label [[OMP_OFFLOAD_CONT]]2040// CHECK17:       omp_offload.cont:2041// CHECK17-NEXT:    [[TMP27:%.*]] = load i32, ptr [[COMP]], align 42042// CHECK17-NEXT:    store i32 [[TMP27]], ptr [[COMP_CASTED3]], align 42043// CHECK17-NEXT:    [[TMP28:%.*]] = load i64, ptr [[COMP_CASTED3]], align 82044// CHECK17-NEXT:    [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 12045// CHECK17-NEXT:    [[TMP29:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 02046// CHECK17-NEXT:    store ptr [[THIS1]], ptr [[TMP29]], align 82047// CHECK17-NEXT:    [[TMP30:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 02048// CHECK17-NEXT:    store ptr [[B]], ptr [[TMP30]], align 82049// CHECK17-NEXT:    [[TMP31:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i64 0, i64 02050// CHECK17-NEXT:    store ptr null, ptr [[TMP31]], align 82051// CHECK17-NEXT:    [[TMP32:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 12052// CHECK17-NEXT:    store i64 [[TMP28]], ptr [[TMP32]], align 82053// CHECK17-NEXT:    [[TMP33:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 12054// CHECK17-NEXT:    store i64 [[TMP28]], ptr [[TMP33]], align 82055// CHECK17-NEXT:    [[TMP34:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i64 0, i64 12056// CHECK17-NEXT:    store ptr null, ptr [[TMP34]], align 82057// CHECK17-NEXT:    [[TMP35:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 02058// CHECK17-NEXT:    [[TMP36:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 02059// CHECK17-NEXT:    [[B7:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 12060// CHECK17-NEXT:    [[TMP37:%.*]] = load float, ptr [[B7]], align 42061// CHECK17-NEXT:    [[CONV:%.*]] = fptosi float [[TMP37]] to i322062// CHECK17-NEXT:    [[ADD:%.*]] = add nsw i32 [[CONV]], 1232063// CHECK17-NEXT:    [[TMP38:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[ADD]], 02064// CHECK17-NEXT:    [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 02065// CHECK17-NEXT:    store i32 3, ptr [[TMP39]], align 42066// CHECK17-NEXT:    [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 12067// CHECK17-NEXT:    store i32 2, ptr [[TMP40]], align 42068// CHECK17-NEXT:    [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 22069// CHECK17-NEXT:    store ptr [[TMP35]], ptr [[TMP41]], align 82070// CHECK17-NEXT:    [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 32071// CHECK17-NEXT:    store ptr [[TMP36]], ptr [[TMP42]], align 82072// CHECK17-NEXT:    [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 42073// CHECK17-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP43]], align 82074// CHECK17-NEXT:    [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 52075// CHECK17-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP44]], align 82076// CHECK17-NEXT:    [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 62077// CHECK17-NEXT:    store ptr null, ptr [[TMP45]], align 82078// CHECK17-NEXT:    [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 72079// CHECK17-NEXT:    store ptr null, ptr [[TMP46]], align 82080// CHECK17-NEXT:    [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 82081// CHECK17-NEXT:    store i64 0, ptr [[TMP47]], align 82082// CHECK17-NEXT:    [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 92083// CHECK17-NEXT:    store i64 0, ptr [[TMP48]], align 82084// CHECK17-NEXT:    [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 102085// CHECK17-NEXT:    store [3 x i32] [i32 456, i32 0, i32 0], ptr [[TMP49]], align 42086// CHECK17-NEXT:    [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 112087// CHECK17-NEXT:    store [3 x i32] [[TMP38]], ptr [[TMP50]], align 42088// CHECK17-NEXT:    [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 122089// CHECK17-NEXT:    store i32 0, ptr [[TMP51]], align 42090// CHECK17-NEXT:    [[TMP52:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 456, i32 [[ADD]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169.region_id, ptr [[KERNEL_ARGS8]])2091// CHECK17-NEXT:    [[TMP53:%.*]] = icmp ne i32 [[TMP52]], 02092// CHECK17-NEXT:    br i1 [[TMP53]], label [[OMP_OFFLOAD_FAILED9:%.*]], label [[OMP_OFFLOAD_CONT10:%.*]]2093// CHECK17:       omp_offload.failed9:2094// CHECK17-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169(ptr [[THIS1]], i64 [[TMP28]]) #[[ATTR2]]2095// CHECK17-NEXT:    br label [[OMP_OFFLOAD_CONT10]]2096// CHECK17:       omp_offload.cont10:2097// CHECK17-NEXT:    [[TMP54:%.*]] = load i32, ptr [[COMP]], align 42098// CHECK17-NEXT:    ret i32 [[TMP54]]2099//2100//2101// CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l1612102// CHECK17-SAME: (ptr noundef [[THIS:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1:[0-9]+]] {2103// CHECK17-NEXT:  entry:2104// CHECK17-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 82105// CHECK17-NEXT:    [[COMP_ADDR:%.*]] = alloca i64, align 82106// CHECK17-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])2107// CHECK17-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 82108// CHECK17-NEXT:    store i64 [[COMP]], ptr [[COMP_ADDR]], align 82109// CHECK17-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 82110// CHECK17-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP1]], i32 0, i32 02111// CHECK17-NEXT:    [[TMP2:%.*]] = load i32, ptr [[A]], align 42112// CHECK17-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 123)2113// CHECK17-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161.omp_outlined, ptr [[COMP_ADDR]])2114// CHECK17-NEXT:    ret void2115//2116//2117// CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161.omp_outlined2118// CHECK17-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] {2119// CHECK17-NEXT:  entry:2120// CHECK17-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82121// CHECK17-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82122// CHECK17-NEXT:    [[COMP_ADDR:%.*]] = alloca ptr, align 82123// CHECK17-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82124// CHECK17-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82125// CHECK17-NEXT:    store ptr [[COMP]], ptr [[COMP_ADDR]], align 82126// CHECK17-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 82127// CHECK17-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 42128// CHECK17-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP1]], 12129// CHECK17-NEXT:    store i32 [[INC]], ptr [[TMP0]], align 42130// CHECK17-NEXT:    ret void2131//2132//2133// CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l1692134// CHECK17-SAME: (ptr noundef [[THIS:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1]] {2135// CHECK17-NEXT:  entry:2136// CHECK17-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 82137// CHECK17-NEXT:    [[COMP_ADDR:%.*]] = alloca i64, align 82138// CHECK17-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])2139// CHECK17-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 82140// CHECK17-NEXT:    store i64 [[COMP]], ptr [[COMP_ADDR]], align 82141// CHECK17-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 82142// CHECK17-NEXT:    [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP1]], i32 0, i32 12143// CHECK17-NEXT:    [[TMP2:%.*]] = load float, ptr [[B]], align 42144// CHECK17-NEXT:    [[CONV:%.*]] = fptosi float [[TMP2]] to i322145// CHECK17-NEXT:    [[ADD:%.*]] = add nsw i32 [[CONV]], 1232146// CHECK17-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 456, i32 [[ADD]])2147// CHECK17-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169.omp_outlined, ptr [[COMP_ADDR]])2148// CHECK17-NEXT:    ret void2149//2150//2151// CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169.omp_outlined2152// CHECK17-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] {2153// CHECK17-NEXT:  entry:2154// CHECK17-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82155// CHECK17-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82156// CHECK17-NEXT:    [[COMP_ADDR:%.*]] = alloca ptr, align 82157// CHECK17-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82158// CHECK17-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82159// CHECK17-NEXT:    store ptr [[COMP]], ptr [[COMP_ADDR]], align 82160// CHECK17-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 82161// CHECK17-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 42162// CHECK17-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP1]], 12163// CHECK17-NEXT:    store i32 [[INC]], ptr [[TMP0]], align 42164// CHECK17-NEXT:    ret void2165//2166//2167// CHECK19-LABEL: define {{[^@]+}}@_Z21teams_template_structv2168// CHECK19-SAME: () #[[ATTR0:[0-9]+]] {2169// CHECK19-NEXT:  entry:2170// CHECK19-NEXT:    [[V:%.*]] = alloca [[STRUCT_SS:%.*]], align 42171// CHECK19-NEXT:    [[CALL:%.*]] = call noundef i32 @_ZN2SSIiLi123ELx456EE3fooEv(ptr noundef nonnull align 4 dereferenceable(8) [[V]])2172// CHECK19-NEXT:    ret i32 [[CALL]]2173//2174//2175// CHECK19-LABEL: define {{[^@]+}}@_ZN2SSIiLi123ELx456EE3fooEv2176// CHECK19-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]]) #[[ATTR0]] comdat align 2 {2177// CHECK19-NEXT:  entry:2178// CHECK19-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 42179// CHECK19-NEXT:    [[COMP:%.*]] = alloca i32, align 42180// CHECK19-NEXT:    [[COMP_CASTED:%.*]] = alloca i32, align 42181// CHECK19-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x ptr], align 42182// CHECK19-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x ptr], align 42183// CHECK19-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [2 x ptr], align 42184// CHECK19-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 82185// CHECK19-NEXT:    [[COMP_CASTED3:%.*]] = alloca i32, align 42186// CHECK19-NEXT:    [[DOTOFFLOAD_BASEPTRS4:%.*]] = alloca [2 x ptr], align 42187// CHECK19-NEXT:    [[DOTOFFLOAD_PTRS5:%.*]] = alloca [2 x ptr], align 42188// CHECK19-NEXT:    [[DOTOFFLOAD_MAPPERS6:%.*]] = alloca [2 x ptr], align 42189// CHECK19-NEXT:    [[KERNEL_ARGS8:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 82190// CHECK19-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 42191// CHECK19-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 42192// CHECK19-NEXT:    store i32 1, ptr [[COMP]], align 42193// CHECK19-NEXT:    [[TMP0:%.*]] = load i32, ptr [[COMP]], align 42194// CHECK19-NEXT:    store i32 [[TMP0]], ptr [[COMP_CASTED]], align 42195// CHECK19-NEXT:    [[TMP1:%.*]] = load i32, ptr [[COMP_CASTED]], align 42196// CHECK19-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[THIS1]], i32 0, i32 02197// CHECK19-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 02198// CHECK19-NEXT:    store ptr [[THIS1]], ptr [[TMP2]], align 42199// CHECK19-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 02200// CHECK19-NEXT:    store ptr [[A]], ptr [[TMP3]], align 42201// CHECK19-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 02202// CHECK19-NEXT:    store ptr null, ptr [[TMP4]], align 42203// CHECK19-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 12204// CHECK19-NEXT:    store i32 [[TMP1]], ptr [[TMP5]], align 42205// CHECK19-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 12206// CHECK19-NEXT:    store i32 [[TMP1]], ptr [[TMP6]], align 42207// CHECK19-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 12208// CHECK19-NEXT:    store ptr null, ptr [[TMP7]], align 42209// CHECK19-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 02210// CHECK19-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 02211// CHECK19-NEXT:    [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 02212// CHECK19-NEXT:    [[TMP10:%.*]] = load i32, ptr [[A2]], align 42213// CHECK19-NEXT:    [[TMP11:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP10]], 02214// CHECK19-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 02215// CHECK19-NEXT:    store i32 3, ptr [[TMP12]], align 42216// CHECK19-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12217// CHECK19-NEXT:    store i32 2, ptr [[TMP13]], align 42218// CHECK19-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 22219// CHECK19-NEXT:    store ptr [[TMP8]], ptr [[TMP14]], align 42220// CHECK19-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 32221// CHECK19-NEXT:    store ptr [[TMP9]], ptr [[TMP15]], align 42222// CHECK19-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 42223// CHECK19-NEXT:    store ptr @.offload_sizes, ptr [[TMP16]], align 42224// CHECK19-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 52225// CHECK19-NEXT:    store ptr @.offload_maptypes, ptr [[TMP17]], align 42226// CHECK19-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 62227// CHECK19-NEXT:    store ptr null, ptr [[TMP18]], align 42228// CHECK19-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 72229// CHECK19-NEXT:    store ptr null, ptr [[TMP19]], align 42230// CHECK19-NEXT:    [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 82231// CHECK19-NEXT:    store i64 0, ptr [[TMP20]], align 82232// CHECK19-NEXT:    [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 92233// CHECK19-NEXT:    store i64 0, ptr [[TMP21]], align 82234// CHECK19-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 102235// CHECK19-NEXT:    store [3 x i32] [[TMP11]], ptr [[TMP22]], align 42236// CHECK19-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 112237// CHECK19-NEXT:    store [3 x i32] [i32 123, i32 0, i32 0], ptr [[TMP23]], align 42238// CHECK19-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 122239// CHECK19-NEXT:    store i32 0, ptr [[TMP24]], align 42240// CHECK19-NEXT:    [[TMP25:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 [[TMP10]], i32 123, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161.region_id, ptr [[KERNEL_ARGS]])2241// CHECK19-NEXT:    [[TMP26:%.*]] = icmp ne i32 [[TMP25]], 02242// CHECK19-NEXT:    br i1 [[TMP26]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]2243// CHECK19:       omp_offload.failed:2244// CHECK19-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161(ptr [[THIS1]], i32 [[TMP1]]) #[[ATTR2:[0-9]+]]2245// CHECK19-NEXT:    br label [[OMP_OFFLOAD_CONT]]2246// CHECK19:       omp_offload.cont:2247// CHECK19-NEXT:    [[TMP27:%.*]] = load i32, ptr [[COMP]], align 42248// CHECK19-NEXT:    store i32 [[TMP27]], ptr [[COMP_CASTED3]], align 42249// CHECK19-NEXT:    [[TMP28:%.*]] = load i32, ptr [[COMP_CASTED3]], align 42250// CHECK19-NEXT:    [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 12251// CHECK19-NEXT:    [[TMP29:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 02252// CHECK19-NEXT:    store ptr [[THIS1]], ptr [[TMP29]], align 42253// CHECK19-NEXT:    [[TMP30:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 02254// CHECK19-NEXT:    store ptr [[B]], ptr [[TMP30]], align 42255// CHECK19-NEXT:    [[TMP31:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i32 0, i32 02256// CHECK19-NEXT:    store ptr null, ptr [[TMP31]], align 42257// CHECK19-NEXT:    [[TMP32:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 12258// CHECK19-NEXT:    store i32 [[TMP28]], ptr [[TMP32]], align 42259// CHECK19-NEXT:    [[TMP33:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 12260// CHECK19-NEXT:    store i32 [[TMP28]], ptr [[TMP33]], align 42261// CHECK19-NEXT:    [[TMP34:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i32 0, i32 12262// CHECK19-NEXT:    store ptr null, ptr [[TMP34]], align 42263// CHECK19-NEXT:    [[TMP35:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 02264// CHECK19-NEXT:    [[TMP36:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 02265// CHECK19-NEXT:    [[B7:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 12266// CHECK19-NEXT:    [[TMP37:%.*]] = load float, ptr [[B7]], align 42267// CHECK19-NEXT:    [[CONV:%.*]] = fptosi float [[TMP37]] to i322268// CHECK19-NEXT:    [[ADD:%.*]] = add nsw i32 [[CONV]], 1232269// CHECK19-NEXT:    [[TMP38:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[ADD]], 02270// CHECK19-NEXT:    [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 02271// CHECK19-NEXT:    store i32 3, ptr [[TMP39]], align 42272// CHECK19-NEXT:    [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 12273// CHECK19-NEXT:    store i32 2, ptr [[TMP40]], align 42274// CHECK19-NEXT:    [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 22275// CHECK19-NEXT:    store ptr [[TMP35]], ptr [[TMP41]], align 42276// CHECK19-NEXT:    [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 32277// CHECK19-NEXT:    store ptr [[TMP36]], ptr [[TMP42]], align 42278// CHECK19-NEXT:    [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 42279// CHECK19-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP43]], align 42280// CHECK19-NEXT:    [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 52281// CHECK19-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP44]], align 42282// CHECK19-NEXT:    [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 62283// CHECK19-NEXT:    store ptr null, ptr [[TMP45]], align 42284// CHECK19-NEXT:    [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 72285// CHECK19-NEXT:    store ptr null, ptr [[TMP46]], align 42286// CHECK19-NEXT:    [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 82287// CHECK19-NEXT:    store i64 0, ptr [[TMP47]], align 82288// CHECK19-NEXT:    [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 92289// CHECK19-NEXT:    store i64 0, ptr [[TMP48]], align 82290// CHECK19-NEXT:    [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 102291// CHECK19-NEXT:    store [3 x i32] [i32 456, i32 0, i32 0], ptr [[TMP49]], align 42292// CHECK19-NEXT:    [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 112293// CHECK19-NEXT:    store [3 x i32] [[TMP38]], ptr [[TMP50]], align 42294// CHECK19-NEXT:    [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 122295// CHECK19-NEXT:    store i32 0, ptr [[TMP51]], align 42296// CHECK19-NEXT:    [[TMP52:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 456, i32 [[ADD]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169.region_id, ptr [[KERNEL_ARGS8]])2297// CHECK19-NEXT:    [[TMP53:%.*]] = icmp ne i32 [[TMP52]], 02298// CHECK19-NEXT:    br i1 [[TMP53]], label [[OMP_OFFLOAD_FAILED9:%.*]], label [[OMP_OFFLOAD_CONT10:%.*]]2299// CHECK19:       omp_offload.failed9:2300// CHECK19-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169(ptr [[THIS1]], i32 [[TMP28]]) #[[ATTR2]]2301// CHECK19-NEXT:    br label [[OMP_OFFLOAD_CONT10]]2302// CHECK19:       omp_offload.cont10:2303// CHECK19-NEXT:    [[TMP54:%.*]] = load i32, ptr [[COMP]], align 42304// CHECK19-NEXT:    ret i32 [[TMP54]]2305//2306//2307// CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l1612308// CHECK19-SAME: (ptr noundef [[THIS:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1:[0-9]+]] {2309// CHECK19-NEXT:  entry:2310// CHECK19-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 42311// CHECK19-NEXT:    [[COMP_ADDR:%.*]] = alloca i32, align 42312// CHECK19-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])2313// CHECK19-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 42314// CHECK19-NEXT:    store i32 [[COMP]], ptr [[COMP_ADDR]], align 42315// CHECK19-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 42316// CHECK19-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP1]], i32 0, i32 02317// CHECK19-NEXT:    [[TMP2:%.*]] = load i32, ptr [[A]], align 42318// CHECK19-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 123)2319// CHECK19-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161.omp_outlined, ptr [[COMP_ADDR]])2320// CHECK19-NEXT:    ret void2321//2322//2323// CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161.omp_outlined2324// CHECK19-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] {2325// CHECK19-NEXT:  entry:2326// CHECK19-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42327// CHECK19-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42328// CHECK19-NEXT:    [[COMP_ADDR:%.*]] = alloca ptr, align 42329// CHECK19-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42330// CHECK19-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42331// CHECK19-NEXT:    store ptr [[COMP]], ptr [[COMP_ADDR]], align 42332// CHECK19-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 42333// CHECK19-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 42334// CHECK19-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP1]], 12335// CHECK19-NEXT:    store i32 [[INC]], ptr [[TMP0]], align 42336// CHECK19-NEXT:    ret void2337//2338//2339// CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l1692340// CHECK19-SAME: (ptr noundef [[THIS:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1]] {2341// CHECK19-NEXT:  entry:2342// CHECK19-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 42343// CHECK19-NEXT:    [[COMP_ADDR:%.*]] = alloca i32, align 42344// CHECK19-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])2345// CHECK19-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 42346// CHECK19-NEXT:    store i32 [[COMP]], ptr [[COMP_ADDR]], align 42347// CHECK19-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 42348// CHECK19-NEXT:    [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP1]], i32 0, i32 12349// CHECK19-NEXT:    [[TMP2:%.*]] = load float, ptr [[B]], align 42350// CHECK19-NEXT:    [[CONV:%.*]] = fptosi float [[TMP2]] to i322351// CHECK19-NEXT:    [[ADD:%.*]] = add nsw i32 [[CONV]], 1232352// CHECK19-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 456, i32 [[ADD]])2353// CHECK19-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169.omp_outlined, ptr [[COMP_ADDR]])2354// CHECK19-NEXT:    ret void2355//2356//2357// CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169.omp_outlined2358// CHECK19-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] {2359// CHECK19-NEXT:  entry:2360// CHECK19-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42361// CHECK19-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42362// CHECK19-NEXT:    [[COMP_ADDR:%.*]] = alloca ptr, align 42363// CHECK19-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42364// CHECK19-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42365// CHECK19-NEXT:    store ptr [[COMP]], ptr [[COMP_ADDR]], align 42366// CHECK19-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 42367// CHECK19-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 42368// CHECK19-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP1]], 12369// CHECK19-NEXT:    store i32 [[INC]], ptr [[TMP0]], align 42370// CHECK19-NEXT:    ret void2371//2372//2373// CHECK25-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l2162374// CHECK25-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[ARGC:%.*]]) #[[ATTR0:[0-9]+]] {2375// CHECK25-NEXT:  entry:2376// CHECK25-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 82377// CHECK25-NEXT:    [[ARGC_ADDR:%.*]] = alloca i64, align 82378// CHECK25-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 82379// CHECK25-NEXT:    store i64 [[ARGC]], ptr [[ARGC_ADDR]], align 82380// CHECK25-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1:[0-9]+]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l216.omp_outlined, ptr [[ARGC_ADDR]])2381// CHECK25-NEXT:    ret void2382//2383//2384// CHECK25-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l216.omp_outlined2385// CHECK25-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR0]] {2386// CHECK25-NEXT:  entry:2387// CHECK25-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82388// CHECK25-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82389// CHECK25-NEXT:    [[ARGC_ADDR:%.*]] = alloca ptr, align 82390// CHECK25-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82391// CHECK25-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82392// CHECK25-NEXT:    store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 82393// CHECK25-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 82394// CHECK25-NEXT:    store i32 0, ptr [[TMP0]], align 42395// CHECK25-NEXT:    ret void2396//2397//2398// CHECK25-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l2092399// CHECK25-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[ARGC:%.*]]) #[[ATTR0]] {2400// CHECK25-NEXT:  entry:2401// CHECK25-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 82402// CHECK25-NEXT:    [[ARGC_ADDR:%.*]] = alloca ptr, align 82403// CHECK25-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 82404// CHECK25-NEXT:    store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 82405// CHECK25-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l209.omp_outlined, ptr [[ARGC_ADDR]])2406// CHECK25-NEXT:    ret void2407//2408//2409// CHECK25-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l209.omp_outlined2410// CHECK25-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[ARGC:%.*]]) #[[ATTR0]] {2411// CHECK25-NEXT:  entry:2412// CHECK25-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82413// CHECK25-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82414// CHECK25-NEXT:    [[ARGC_ADDR:%.*]] = alloca ptr, align 82415// CHECK25-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82416// CHECK25-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82417// CHECK25-NEXT:    store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 82418// CHECK25-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 82419// CHECK25-NEXT:    store ptr null, ptr [[TMP0]], align 82420// CHECK25-NEXT:    ret void2421//2422//2423// CHECK27-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l2162424// CHECK27-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[ARGC:%.*]]) #[[ATTR0:[0-9]+]] {2425// CHECK27-NEXT:  entry:2426// CHECK27-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42427// CHECK27-NEXT:    [[ARGC_ADDR:%.*]] = alloca i32, align 42428// CHECK27-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42429// CHECK27-NEXT:    store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 42430// CHECK27-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1:[0-9]+]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l216.omp_outlined, ptr [[ARGC_ADDR]])2431// CHECK27-NEXT:    ret void2432//2433//2434// CHECK27-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l216.omp_outlined2435// CHECK27-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR0]] {2436// CHECK27-NEXT:  entry:2437// CHECK27-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42438// CHECK27-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42439// CHECK27-NEXT:    [[ARGC_ADDR:%.*]] = alloca ptr, align 42440// CHECK27-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42441// CHECK27-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42442// CHECK27-NEXT:    store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 42443// CHECK27-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 42444// CHECK27-NEXT:    store i32 0, ptr [[TMP0]], align 42445// CHECK27-NEXT:    ret void2446//2447//2448// CHECK27-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l2092449// CHECK27-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[ARGC:%.*]]) #[[ATTR0]] {2450// CHECK27-NEXT:  entry:2451// CHECK27-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42452// CHECK27-NEXT:    [[ARGC_ADDR:%.*]] = alloca ptr, align 42453// CHECK27-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42454// CHECK27-NEXT:    store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 42455// CHECK27-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l209.omp_outlined, ptr [[ARGC_ADDR]])2456// CHECK27-NEXT:    ret void2457//2458//2459// CHECK27-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l209.omp_outlined2460// CHECK27-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR0]] {2461// CHECK27-NEXT:  entry:2462// CHECK27-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42463// CHECK27-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42464// CHECK27-NEXT:    [[ARGC_ADDR:%.*]] = alloca ptr, align 42465// CHECK27-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42466// CHECK27-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42467// CHECK27-NEXT:    store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 42468// CHECK27-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 42469// CHECK27-NEXT:    store ptr null, ptr [[TMP0]], align 42470// CHECK27-NEXT:    ret void2471//2472//2473// CHECK33-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l2652474// CHECK33-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[A:%.*]], i64 noundef [[B:%.*]], i64 noundef [[ARGC:%.*]]) #[[ATTR0:[0-9]+]] {2475// CHECK33-NEXT:  entry:2476// CHECK33-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 82477// CHECK33-NEXT:    [[A_ADDR:%.*]] = alloca i64, align 82478// CHECK33-NEXT:    [[B_ADDR:%.*]] = alloca i64, align 82479// CHECK33-NEXT:    [[ARGC_ADDR:%.*]] = alloca i64, align 82480// CHECK33-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])2481// CHECK33-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 82482// CHECK33-NEXT:    store i64 [[A]], ptr [[A_ADDR]], align 82483// CHECK33-NEXT:    store i64 [[B]], ptr [[B_ADDR]], align 82484// CHECK33-NEXT:    store i64 [[ARGC]], ptr [[ARGC_ADDR]], align 82485// CHECK33-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 42486// CHECK33-NEXT:    [[TMP2:%.*]] = load i32, ptr [[B_ADDR]], align 42487// CHECK33-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]])2488// CHECK33-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l265.omp_outlined, ptr [[ARGC_ADDR]])2489// CHECK33-NEXT:    ret void2490//2491//2492// CHECK33-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l265.omp_outlined2493// CHECK33-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR0]] {2494// CHECK33-NEXT:  entry:2495// CHECK33-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82496// CHECK33-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82497// CHECK33-NEXT:    [[ARGC_ADDR:%.*]] = alloca ptr, align 82498// CHECK33-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82499// CHECK33-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82500// CHECK33-NEXT:    store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 82501// CHECK33-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 82502// CHECK33-NEXT:    store i32 0, ptr [[TMP0]], align 42503// CHECK33-NEXT:    ret void2504//2505//2506// CHECK33-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l2542507// CHECK33-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[A:%.*]], i64 noundef [[B:%.*]], ptr noundef [[ARGC:%.*]]) #[[ATTR0]] {2508// CHECK33-NEXT:  entry:2509// CHECK33-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 82510// CHECK33-NEXT:    [[A_ADDR:%.*]] = alloca i64, align 82511// CHECK33-NEXT:    [[B_ADDR:%.*]] = alloca i64, align 82512// CHECK33-NEXT:    [[ARGC_ADDR:%.*]] = alloca ptr, align 82513// CHECK33-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])2514// CHECK33-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 82515// CHECK33-NEXT:    store i64 [[A]], ptr [[A_ADDR]], align 82516// CHECK33-NEXT:    store i64 [[B]], ptr [[B_ADDR]], align 82517// CHECK33-NEXT:    store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 82518// CHECK33-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 42519// CHECK33-NEXT:    [[TMP2:%.*]] = load i32, ptr [[B_ADDR]], align 42520// CHECK33-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]])2521// CHECK33-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l254.omp_outlined, ptr [[ARGC_ADDR]])2522// CHECK33-NEXT:    ret void2523//2524//2525// CHECK33-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l254.omp_outlined2526// CHECK33-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[ARGC:%.*]]) #[[ATTR0]] {2527// CHECK33-NEXT:  entry:2528// CHECK33-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82529// CHECK33-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82530// CHECK33-NEXT:    [[ARGC_ADDR:%.*]] = alloca ptr, align 82531// CHECK33-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82532// CHECK33-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82533// CHECK33-NEXT:    store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 82534// CHECK33-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 82535// CHECK33-NEXT:    store ptr null, ptr [[TMP0]], align 82536// CHECK33-NEXT:    ret void2537//2538//2539// CHECK35-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l2652540// CHECK35-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[A:%.*]], i32 noundef [[B:%.*]], i32 noundef [[ARGC:%.*]]) #[[ATTR0:[0-9]+]] {2541// CHECK35-NEXT:  entry:2542// CHECK35-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42543// CHECK35-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 42544// CHECK35-NEXT:    [[B_ADDR:%.*]] = alloca i32, align 42545// CHECK35-NEXT:    [[ARGC_ADDR:%.*]] = alloca i32, align 42546// CHECK35-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])2547// CHECK35-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42548// CHECK35-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 42549// CHECK35-NEXT:    store i32 [[B]], ptr [[B_ADDR]], align 42550// CHECK35-NEXT:    store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 42551// CHECK35-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 42552// CHECK35-NEXT:    [[TMP2:%.*]] = load i32, ptr [[B_ADDR]], align 42553// CHECK35-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]])2554// CHECK35-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l265.omp_outlined, ptr [[ARGC_ADDR]])2555// CHECK35-NEXT:    ret void2556//2557//2558// CHECK35-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l265.omp_outlined2559// CHECK35-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR0]] {2560// CHECK35-NEXT:  entry:2561// CHECK35-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42562// CHECK35-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42563// CHECK35-NEXT:    [[ARGC_ADDR:%.*]] = alloca ptr, align 42564// CHECK35-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42565// CHECK35-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42566// CHECK35-NEXT:    store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 42567// CHECK35-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 42568// CHECK35-NEXT:    store i32 0, ptr [[TMP0]], align 42569// CHECK35-NEXT:    ret void2570//2571//2572// CHECK35-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l2542573// CHECK35-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[A:%.*]], i32 noundef [[B:%.*]], ptr noundef [[ARGC:%.*]]) #[[ATTR0]] {2574// CHECK35-NEXT:  entry:2575// CHECK35-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42576// CHECK35-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 42577// CHECK35-NEXT:    [[B_ADDR:%.*]] = alloca i32, align 42578// CHECK35-NEXT:    [[ARGC_ADDR:%.*]] = alloca ptr, align 42579// CHECK35-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])2580// CHECK35-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42581// CHECK35-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 42582// CHECK35-NEXT:    store i32 [[B]], ptr [[B_ADDR]], align 42583// CHECK35-NEXT:    store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 42584// CHECK35-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 42585// CHECK35-NEXT:    [[TMP2:%.*]] = load i32, ptr [[B_ADDR]], align 42586// CHECK35-NEXT:    call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]])2587// CHECK35-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l254.omp_outlined, ptr [[ARGC_ADDR]])2588// CHECK35-NEXT:    ret void2589//2590//2591// CHECK35-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l254.omp_outlined2592// CHECK35-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR0]] {2593// CHECK35-NEXT:  entry:2594// CHECK35-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42595// CHECK35-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42596// CHECK35-NEXT:    [[ARGC_ADDR:%.*]] = alloca ptr, align 42597// CHECK35-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42598// CHECK35-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42599// CHECK35-NEXT:    store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 42600// CHECK35-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 42601// CHECK35-NEXT:    store ptr null, ptr [[TMP0]], align 42602// CHECK35-NEXT:    ret void2603//2604//2605// CHECK41-LABEL: define {{[^@]+}}@_Z3foov2606// CHECK41-SAME: () #[[ATTR0:[0-9]+]] {2607// CHECK41-NEXT:  entry:2608// CHECK41-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1:[0-9]+]], i32 0, ptr @_Z3foov.omp_outlined)2609// CHECK41-NEXT:    ret void2610//2611//2612// CHECK41-LABEL: define {{[^@]+}}@_Z3foov.omp_outlined2613// CHECK41-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] {2614// CHECK41-NEXT:  entry:2615// CHECK41-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82616// CHECK41-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82617// CHECK41-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82618// CHECK41-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82619// CHECK41-NEXT:    ret void2620//2621//2622// CHECK43-LABEL: define {{[^@]+}}@_Z3foov2623// CHECK43-SAME: () #[[ATTR0:[0-9]+]] {2624// CHECK43-NEXT:  entry:2625// CHECK43-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1:[0-9]+]], i32 0, ptr @_Z3foov.omp_outlined)2626// CHECK43-NEXT:    ret void2627//2628//2629// CHECK43-LABEL: define {{[^@]+}}@_Z3foov.omp_outlined2630// CHECK43-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] {2631// CHECK43-NEXT:  entry:2632// CHECK43-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42633// CHECK43-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42634// CHECK43-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42635// CHECK43-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42636// CHECK43-NEXT:    ret void2637//2638