clang/test/OpenMP/distribute_private_codegen.cpp

   1 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _
   2 // RUN: %clang_cc1 -DLAMBDA -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK1
   3 // RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
   4 // RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK1
   5 // RUN: %clang_cc1 -DLAMBDA -verify -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK3
   6 // RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
   7 // RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK3
   8
   9 // RUN: %clang_cc1 -DLAMBDA -verify -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
  10 // RUN: %clang_cc1 -DLAMBDA -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
  11 // RUN: %clang_cc1 -DLAMBDA -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
  12 // RUN: %clang_cc1 -DLAMBDA -verify -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
  13 // RUN: %clang_cc1 -DLAMBDA -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
  14 // RUN: %clang_cc1 -DLAMBDA -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
  15
  16 // RUN: %clang_cc1  -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK9
  17 // RUN: %clang_cc1  -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
  18 // RUN: %clang_cc1  -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK9
  19 // RUN: %clang_cc1  -verify -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK11
  20 // RUN: %clang_cc1  -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
  21 // RUN: %clang_cc1  -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK11
  22
  23 // RUN: %clang_cc1  -verify -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
  24 // RUN: %clang_cc1  -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
  25 // RUN: %clang_cc1  -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
  26 // RUN: %clang_cc1  -verify -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
  27 // RUN: %clang_cc1  -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
  28 // RUN: %clang_cc1  -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"
  29 // expected-no-diagnostics
  30 #ifndef HEADER
  31 #define HEADER
  32
  33 template <class T>
  34 struct S {
  35   T f;
  36   S(T a) : f(a) {}
  37   S() : f() {}
  38   operator T() { return T(); }
  39   ~S() {}
  40 };
  41
  42 template <typename T>
  43 T tmain() {
  44   S<T> test;
  45   T t_var = T();
  46   T vec[] = {1, 2};
  47   S<T> s_arr[] = {1, 2};
  48   S<T> &var = test;
  49   #pragma omp target
  50   #pragma omp teams
  51 #pragma omp distribute private(t_var, vec, s_arr, s_arr, var, var)
  52   for (int i = 0; i < 2; ++i) {
  53     vec[i] = t_var;
  54     s_arr[i] = var;
  55   }
  56   return T();
  57 }
  58
  59 int main() {
  60   static int svar;
  61   volatile double g;
  62   volatile double &g1 = g;
  63
  64   #ifdef LAMBDA
  65   [&]() {
  66     static float sfvar;
  67
  68     #pragma omp target
  69     #pragma omp teams
  70 #pragma omp distribute private(g, g1, svar, sfvar)
  71     for (int i = 0; i < 2; ++i) {
  72       g = 1;
  73       g1 = 1;
  74       svar = 3;
  75       sfvar = 4.0;
  76       [&]() {
  77         g = 2;
  78         g1 = 2;
  79         svar = 4;
  80         sfvar = 8.0;
  81
  82       }();
  83     }
  84   }();
  85   return 0;
  86   #else
  87   S<float> test;
  88   int t_var = 0;
  89   int vec[] = {1, 2};
  90   S<float> s_arr[] = {1, 2};
  91   S<float> &var = test;
  92
  93   #pragma omp target
  94   #pragma omp teams
  95 #pragma omp distribute private(t_var, vec, s_arr, s_arr, var, var, svar)
  96   for (int i = 0; i < 2; ++i) {
  97     vec[i] = t_var;
  98     s_arr[i] = var;
  99   }
 100   int i;
 101
 102   #pragma omp target
 103   #pragma omp teams
 104 #pragma omp distribute private(i)
 105   for (i = 0; i < 2; ++i) {
 106     ;
 107   }
 108   return tmain<int>();
 109   #endif
 110 }
 111
 112
 113
 114
 115
 116 #endif
 117 // CHECK1-LABEL: define {{[^@]+}}@main
 118 // CHECK1-SAME: () #[[ATTR0:[0-9]+]] {
 119 // CHECK1-NEXT:  entry:
 120 // CHECK1-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4
 121 // CHECK1-NEXT:    [[G:%.*]] = alloca double, align 8
 122 // CHECK1-NEXT:    [[G1:%.*]] = alloca ptr, align 8
 123 // CHECK1-NEXT:    [[REF_TMP:%.*]] = alloca [[CLASS_ANON:%.*]], align 1
 124 // CHECK1-NEXT:    store i32 0, ptr [[RETVAL]], align 4
 125 // CHECK1-NEXT:    store ptr [[G]], ptr [[G1]], align 8
 126 // CHECK1-NEXT:    call void @"_ZZ4mainENK3$_0clEv"(ptr noundef nonnull align 1 dereferenceable(1) [[REF_TMP]])
 127 // CHECK1-NEXT:    ret i32 0
 128 //
 129 //
 130 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68
 131 // CHECK1-SAME: () #[[ATTR2:[0-9]+]] {
 132 // CHECK1-NEXT:  entry:
 133 // CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2:[0-9]+]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68.omp_outlined)
 134 // CHECK1-NEXT:    ret void
 135 //
 136 //
 137 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68.omp_outlined
 138 // CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR2]] {
 139 // CHECK1-NEXT:  entry:
 140 // CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
 141 // CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
 142 // CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
 143 // CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4
 144 // CHECK1-NEXT:    [[_TMP1:%.*]] = alloca ptr, align 8
 145 // CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4
 146 // CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4
 147 // CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
 148 // CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
 149 // CHECK1-NEXT:    [[G:%.*]] = alloca double, align 8
 150 // CHECK1-NEXT:    [[G1:%.*]] = alloca double, align 8
 151 // CHECK1-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 8
 152 // CHECK1-NEXT:    [[SVAR:%.*]] = alloca i32, align 4
 153 // CHECK1-NEXT:    [[SFVAR:%.*]] = alloca float, align 4
 154 // CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4
 155 // CHECK1-NEXT:    [[REF_TMP:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 8
 156 // CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
 157 // CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
 158 // CHECK1-NEXT:    store ptr undef, ptr [[_TMP1]], align 8
 159 // CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4
 160 // CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 4
 161 // CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4
 162 // CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4
 163 // CHECK1-NEXT:    store ptr [[G1]], ptr [[_TMP2]], align 8
 164 // CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
 165 // CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
 166 // CHECK1-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP1]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
 167 // CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
 168 // CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 1
 169 // CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
 170 // CHECK1:       cond.true:
 171 // CHECK1-NEXT:    br label [[COND_END:%.*]]
 172 // CHECK1:       cond.false:
 173 // CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
 174 // CHECK1-NEXT:    br label [[COND_END]]
 175 // CHECK1:       cond.end:
 176 // CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]
 177 // CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4
 178 // CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4
 179 // CHECK1-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4
 180 // CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
 181 // CHECK1:       omp.inner.for.cond:
 182 // CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
 183 // CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
 184 // CHECK1-NEXT:    [[CMP3:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]
 185 // CHECK1-NEXT:    br i1 [[CMP3]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
 186 // CHECK1:       omp.inner.for.body:
 187 // CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
 188 // CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
 189 // CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
 190 // CHECK1-NEXT:    store i32 [[ADD]], ptr [[I]], align 4
 191 // CHECK1-NEXT:    store double 1.000000e+00, ptr [[G]], align 8
 192 // CHECK1-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[_TMP2]], align 8
 193 // CHECK1-NEXT:    store volatile double 1.000000e+00, ptr [[TMP8]], align 8
 194 // CHECK1-NEXT:    store i32 3, ptr [[SVAR]], align 4
 195 // CHECK1-NEXT:    store float 4.000000e+00, ptr [[SFVAR]], align 4
 196 // CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 0
 197 // CHECK1-NEXT:    store ptr [[G]], ptr [[TMP9]], align 8
 198 // CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 1
 199 // CHECK1-NEXT:    [[TMP11:%.*]] = load ptr, ptr [[_TMP2]], align 8
 200 // CHECK1-NEXT:    store ptr [[TMP11]], ptr [[TMP10]], align 8
 201 // CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 2
 202 // CHECK1-NEXT:    store ptr [[SVAR]], ptr [[TMP12]], align 8
 203 // CHECK1-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 3
 204 // CHECK1-NEXT:    store ptr [[SFVAR]], ptr [[TMP13]], align 8
 205 // CHECK1-NEXT:    call void @"_ZZZ4mainENK3$_0clEvENKUlvE_clEv"(ptr noundef nonnull align 8 dereferenceable(32) [[REF_TMP]])
 206 // CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
 207 // CHECK1:       omp.body.continue:
 208 // CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
 209 // CHECK1:       omp.inner.for.inc:
 210 // CHECK1-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
 211 // CHECK1-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP14]], 1
 212 // CHECK1-NEXT:    store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4
 213 // CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]
 214 // CHECK1:       omp.inner.for.end:
 215 // CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
 216 // CHECK1:       omp.loop.exit:
 217 // CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])
 218 // CHECK1-NEXT:    ret void
 219 //
 220 //
 221 // CHECK1-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg
 222 // CHECK1-SAME: () #[[ATTR4:[0-9]+]] {
 223 // CHECK1-NEXT:  entry:
 224 // CHECK1-NEXT:    call void @__tgt_register_requires(i64 1)
 225 // CHECK1-NEXT:    ret void
 226 //
 227 //
 228 // CHECK3-LABEL: define {{[^@]+}}@main
 229 // CHECK3-SAME: () #[[ATTR0:[0-9]+]] {
 230 // CHECK3-NEXT:  entry:
 231 // CHECK3-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4
 232 // CHECK3-NEXT:    [[G:%.*]] = alloca double, align 8
 233 // CHECK3-NEXT:    [[G1:%.*]] = alloca ptr, align 4
 234 // CHECK3-NEXT:    [[REF_TMP:%.*]] = alloca [[CLASS_ANON:%.*]], align 1
 235 // CHECK3-NEXT:    store i32 0, ptr [[RETVAL]], align 4
 236 // CHECK3-NEXT:    store ptr [[G]], ptr [[G1]], align 4
 237 // CHECK3-NEXT:    call void @"_ZZ4mainENK3$_0clEv"(ptr noundef nonnull align 1 dereferenceable(1) [[REF_TMP]])
 238 // CHECK3-NEXT:    ret i32 0
 239 //
 240 //
 241 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68
 242 // CHECK3-SAME: () #[[ATTR2:[0-9]+]] {
 243 // CHECK3-NEXT:  entry:
 244 // CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2:[0-9]+]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68.omp_outlined)
 245 // CHECK3-NEXT:    ret void
 246 //
 247 //
 248 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68.omp_outlined
 249 // CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR2]] {
 250 // CHECK3-NEXT:  entry:
 251 // CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
 252 // CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
 253 // CHECK3-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
 254 // CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 4
 255 // CHECK3-NEXT:    [[_TMP1:%.*]] = alloca ptr, align 4
 256 // CHECK3-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4
 257 // CHECK3-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4
 258 // CHECK3-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
 259 // CHECK3-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
 260 // CHECK3-NEXT:    [[G:%.*]] = alloca double, align 8
 261 // CHECK3-NEXT:    [[G1:%.*]] = alloca double, align 8
 262 // CHECK3-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 4
 263 // CHECK3-NEXT:    [[SVAR:%.*]] = alloca i32, align 4
 264 // CHECK3-NEXT:    [[SFVAR:%.*]] = alloca float, align 4
 265 // CHECK3-NEXT:    [[I:%.*]] = alloca i32, align 4
 266 // CHECK3-NEXT:    [[REF_TMP:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 4
 267 // CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
 268 // CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
 269 // CHECK3-NEXT:    store ptr undef, ptr [[_TMP1]], align 4
 270 // CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4
 271 // CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 4
 272 // CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4
 273 // CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4
 274 // CHECK3-NEXT:    store ptr [[G1]], ptr [[_TMP2]], align 4
 275 // CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
 276 // CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
 277 // CHECK3-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP1]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
 278 // CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
 279 // CHECK3-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 1
 280 // CHECK3-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
 281 // CHECK3:       cond.true:
 282 // CHECK3-NEXT:    br label [[COND_END:%.*]]
 283 // CHECK3:       cond.false:
 284 // CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
 285 // CHECK3-NEXT:    br label [[COND_END]]
 286 // CHECK3:       cond.end:
 287 // CHECK3-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]
 288 // CHECK3-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4
 289 // CHECK3-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4
 290 // CHECK3-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4
 291 // CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
 292 // CHECK3:       omp.inner.for.cond:
 293 // CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
 294 // CHECK3-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
 295 // CHECK3-NEXT:    [[CMP3:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]
 296 // CHECK3-NEXT:    br i1 [[CMP3]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
 297 // CHECK3:       omp.inner.for.body:
 298 // CHECK3-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
 299 // CHECK3-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
 300 // CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
 301 // CHECK3-NEXT:    store i32 [[ADD]], ptr [[I]], align 4
 302 // CHECK3-NEXT:    store double 1.000000e+00, ptr [[G]], align 8
 303 // CHECK3-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[_TMP2]], align 4
 304 // CHECK3-NEXT:    store volatile double 1.000000e+00, ptr [[TMP8]], align 4
 305 // CHECK3-NEXT:    store i32 3, ptr [[SVAR]], align 4
 306 // CHECK3-NEXT:    store float 4.000000e+00, ptr [[SFVAR]], align 4
 307 // CHECK3-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 0
 308 // CHECK3-NEXT:    store ptr [[G]], ptr [[TMP9]], align 4
 309 // CHECK3-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 1
 310 // CHECK3-NEXT:    [[TMP11:%.*]] = load ptr, ptr [[_TMP2]], align 4
 311 // CHECK3-NEXT:    store ptr [[TMP11]], ptr [[TMP10]], align 4
 312 // CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 2
 313 // CHECK3-NEXT:    store ptr [[SVAR]], ptr [[TMP12]], align 4
 314 // CHECK3-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 3
 315 // CHECK3-NEXT:    store ptr [[SFVAR]], ptr [[TMP13]], align 4
 316 // CHECK3-NEXT:    call void @"_ZZZ4mainENK3$_0clEvENKUlvE_clEv"(ptr noundef nonnull align 4 dereferenceable(16) [[REF_TMP]])
 317 // CHECK3-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
 318 // CHECK3:       omp.body.continue:
 319 // CHECK3-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
 320 // CHECK3:       omp.inner.for.inc:
 321 // CHECK3-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
 322 // CHECK3-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP14]], 1
 323 // CHECK3-NEXT:    store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4
 324 // CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND]]
 325 // CHECK3:       omp.inner.for.end:
 326 // CHECK3-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
 327 // CHECK3:       omp.loop.exit:
 328 // CHECK3-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])
 329 // CHECK3-NEXT:    ret void
 330 //
 331 //
 332 // CHECK3-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg
 333 // CHECK3-SAME: () #[[ATTR4:[0-9]+]] {
 334 // CHECK3-NEXT:  entry:
 335 // CHECK3-NEXT:    call void @__tgt_register_requires(i64 1)
 336 // CHECK3-NEXT:    ret void
 337 //
 338 //
 339 // CHECK9-LABEL: define {{[^@]+}}@main
 340 // CHECK9-SAME: () #[[ATTR0:[0-9]+]] {
 341 // CHECK9-NEXT:  entry:
 342 // CHECK9-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4
 343 // CHECK9-NEXT:    [[G:%.*]] = alloca double, align 8
 344 // CHECK9-NEXT:    [[G1:%.*]] = alloca ptr, align 8
 345 // CHECK9-NEXT:    [[TEST:%.*]] = alloca [[STRUCT_S:%.*]], align 4
 346 // CHECK9-NEXT:    [[T_VAR:%.*]] = alloca i32, align 4
 347 // CHECK9-NEXT:    [[VEC:%.*]] = alloca [2 x i32], align 4
 348 // CHECK9-NEXT:    [[S_ARR:%.*]] = alloca [2 x %struct.S], align 4
 349 // CHECK9-NEXT:    [[VAR:%.*]] = alloca ptr, align 8
 350 // CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 4
 351 // CHECK9-NEXT:    [[_TMP1:%.*]] = alloca ptr, align 8
 352 // CHECK9-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8
 353 // CHECK9-NEXT:    [[I:%.*]] = alloca i32, align 4
 354 // CHECK9-NEXT:    [[_TMP2:%.*]] = alloca i32, align 4
 355 // CHECK9-NEXT:    [[KERNEL_ARGS3:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8
 356 // CHECK9-NEXT:    store i32 0, ptr [[RETVAL]], align 4
 357 // CHECK9-NEXT:    store ptr [[G]], ptr [[G1]], align 8
 358 // CHECK9-NEXT:    call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[TEST]])
 359 // CHECK9-NEXT:    store i32 0, ptr [[T_VAR]], align 4
 360 // CHECK9-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[VEC]], ptr align 4 @__const.main.vec, i64 8, i1 false)
 361 // CHECK9-NEXT:    [[ARRAYINIT_BEGIN:%.*]] = getelementptr inbounds [2 x %struct.S], ptr [[S_ARR]], i64 0, i64 0
 362 // CHECK9-NEXT:    call void @_ZN1SIfEC1Ef(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYINIT_BEGIN]], float noundef 1.000000e+00)
 363 // CHECK9-NEXT:    [[ARRAYINIT_ELEMENT:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYINIT_BEGIN]], i64 1
 364 // CHECK9-NEXT:    call void @_ZN1SIfEC1Ef(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYINIT_ELEMENT]], float noundef 2.000000e+00)
 365 // CHECK9-NEXT:    store ptr [[TEST]], ptr [[VAR]], align 8
 366 // CHECK9-NEXT:    store ptr undef, ptr [[_TMP1]], align 8
 367 // CHECK9-NEXT:    [[TMP0:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0
 368 // CHECK9-NEXT:    store i32 2, ptr [[TMP0]], align 4
 369 // CHECK9-NEXT:    [[TMP1:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1
 370 // CHECK9-NEXT:    store i32 0, ptr [[TMP1]], align 4
 371 // CHECK9-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2
 372 // CHECK9-NEXT:    store ptr null, ptr [[TMP2]], align 8
 373 // CHECK9-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3
 374 // CHECK9-NEXT:    store ptr null, ptr [[TMP3]], align 8
 375 // CHECK9-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4
 376 // CHECK9-NEXT:    store ptr null, ptr [[TMP4]], align 8
 377 // CHECK9-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5
 378 // CHECK9-NEXT:    store ptr null, ptr [[TMP5]], align 8
 379 // CHECK9-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6
 380 // CHECK9-NEXT:    store ptr null, ptr [[TMP6]], align 8
 381 // CHECK9-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7
 382 // CHECK9-NEXT:    store ptr null, ptr [[TMP7]], align 8
 383 // CHECK9-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8
 384 // CHECK9-NEXT:    store i64 2, ptr [[TMP8]], align 8
 385 // CHECK9-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9
 386 // CHECK9-NEXT:    store i64 0, ptr [[TMP9]], align 8
 387 // CHECK9-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10
 388 // CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP10]], align 4
 389 // CHECK9-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11
 390 // CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP11]], align 4
 391 // CHECK9-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12
 392 // CHECK9-NEXT:    store i32 0, ptr [[TMP12]], align 4
 393 // CHECK9-NEXT:    [[TMP13:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l93.region_id, ptr [[KERNEL_ARGS]])
 394 // CHECK9-NEXT:    [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 0
 395 // CHECK9-NEXT:    br i1 [[TMP14]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
 396 // CHECK9:       omp_offload.failed:
 397 // CHECK9-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l93() #[[ATTR4:[0-9]+]]
 398 // CHECK9-NEXT:    br label [[OMP_OFFLOAD_CONT]]
 399 // CHECK9:       omp_offload.cont:
 400 // CHECK9-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 0
 401 // CHECK9-NEXT:    store i32 2, ptr [[TMP15]], align 4
 402 // CHECK9-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 1
 403 // CHECK9-NEXT:    store i32 0, ptr [[TMP16]], align 4
 404 // CHECK9-NEXT:    [[TMP17:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 2
 405 // CHECK9-NEXT:    store ptr null, ptr [[TMP17]], align 8
 406 // CHECK9-NEXT:    [[TMP18:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 3
 407 // CHECK9-NEXT:    store ptr null, ptr [[TMP18]], align 8
 408 // CHECK9-NEXT:    [[TMP19:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 4
 409 // CHECK9-NEXT:    store ptr null, ptr [[TMP19]], align 8
 410 // CHECK9-NEXT:    [[TMP20:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 5
 411 // CHECK9-NEXT:    store ptr null, ptr [[TMP20]], align 8
 412 // CHECK9-NEXT:    [[TMP21:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 6
 413 // CHECK9-NEXT:    store ptr null, ptr [[TMP21]], align 8
 414 // CHECK9-NEXT:    [[TMP22:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 7
 415 // CHECK9-NEXT:    store ptr null, ptr [[TMP22]], align 8
 416 // CHECK9-NEXT:    [[TMP23:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 8
 417 // CHECK9-NEXT:    store i64 2, ptr [[TMP23]], align 8
 418 // CHECK9-NEXT:    [[TMP24:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 9
 419 // CHECK9-NEXT:    store i64 0, ptr [[TMP24]], align 8
 420 // CHECK9-NEXT:    [[TMP25:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 10
 421 // CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP25]], align 4
 422 // CHECK9-NEXT:    [[TMP26:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 11
 423 // CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP26]], align 4
 424 // CHECK9-NEXT:    [[TMP27:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 12
 425 // CHECK9-NEXT:    store i32 0, ptr [[TMP27]], align 4
 426 // CHECK9-NEXT:    [[TMP28:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l102.region_id, ptr [[KERNEL_ARGS3]])
 427 // CHECK9-NEXT:    [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 0
 428 // CHECK9-NEXT:    br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED4:%.*]], label [[OMP_OFFLOAD_CONT5:%.*]]
 429 // CHECK9:       omp_offload.failed4:
 430 // CHECK9-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l102() #[[ATTR4]]
 431 // CHECK9-NEXT:    br label [[OMP_OFFLOAD_CONT5]]
 432 // CHECK9:       omp_offload.cont5:
 433 // CHECK9-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_Z5tmainIiET_v()
 434 // CHECK9-NEXT:    store i32 [[CALL]], ptr [[RETVAL]], align 4
 435 // CHECK9-NEXT:    [[ARRAY_BEGIN:%.*]] = getelementptr inbounds [2 x %struct.S], ptr [[S_ARR]], i32 0, i32 0
 436 // CHECK9-NEXT:    [[TMP30:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAY_BEGIN]], i64 2
 437 // CHECK9-NEXT:    br label [[ARRAYDESTROY_BODY:%.*]]
 438 // CHECK9:       arraydestroy.body:
 439 // CHECK9-NEXT:    [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP30]], [[OMP_OFFLOAD_CONT5]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
 440 // CHECK9-NEXT:    [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i64 -1
 441 // CHECK9-NEXT:    call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR4]]
 442 // CHECK9-NEXT:    [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN]]
 443 // CHECK9-NEXT:    br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE6:%.*]], label [[ARRAYDESTROY_BODY]]
 444 // CHECK9:       arraydestroy.done6:
 445 // CHECK9-NEXT:    call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[TEST]]) #[[ATTR4]]
 446 // CHECK9-NEXT:    [[TMP31:%.*]] = load i32, ptr [[RETVAL]], align 4
 447 // CHECK9-NEXT:    ret i32 [[TMP31]]
 448 //
 449 //
 450 // CHECK9-LABEL: define {{[^@]+}}@_ZN1SIfEC1Ev
 451 // CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1:[0-9]+]] comdat {
 452 // CHECK9-NEXT:  entry:
 453 // CHECK9-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 454 // CHECK9-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 455 // CHECK9-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 456 // CHECK9-NEXT:    call void @_ZN1SIfEC2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]])
 457 // CHECK9-NEXT:    ret void
 458 //
 459 //
 460 // CHECK9-LABEL: define {{[^@]+}}@_ZN1SIfEC1Ef
 461 // CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]], float noundef [[A:%.*]]) unnamed_addr #[[ATTR1]] comdat {
 462 // CHECK9-NEXT:  entry:
 463 // CHECK9-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 464 // CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca float, align 4
 465 // CHECK9-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 466 // CHECK9-NEXT:    store float [[A]], ptr [[A_ADDR]], align 4
 467 // CHECK9-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 468 // CHECK9-NEXT:    [[TMP0:%.*]] = load float, ptr [[A_ADDR]], align 4
 469 // CHECK9-NEXT:    call void @_ZN1SIfEC2Ef(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]], float noundef [[TMP0]])
 470 // CHECK9-NEXT:    ret void
 471 //
 472 //
 473 // CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l93
 474 // CHECK9-SAME: () #[[ATTR3:[0-9]+]] {
 475 // CHECK9-NEXT:  entry:
 476 // CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l93.omp_outlined)
 477 // CHECK9-NEXT:    ret void
 478 //
 479 //
 480 // CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l93.omp_outlined
 481 // CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR3]] {
 482 // CHECK9-NEXT:  entry:
 483 // CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
 484 // CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
 485 // CHECK9-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
 486 // CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 4
 487 // CHECK9-NEXT:    [[_TMP1:%.*]] = alloca ptr, align 8
 488 // CHECK9-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4
 489 // CHECK9-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4
 490 // CHECK9-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
 491 // CHECK9-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
 492 // CHECK9-NEXT:    [[T_VAR:%.*]] = alloca i32, align 4
 493 // CHECK9-NEXT:    [[VEC:%.*]] = alloca [2 x i32], align 4
 494 // CHECK9-NEXT:    [[S_ARR:%.*]] = alloca [2 x %struct.S], align 4
 495 // CHECK9-NEXT:    [[VAR:%.*]] = alloca [[STRUCT_S:%.*]], align 4
 496 // CHECK9-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 8
 497 // CHECK9-NEXT:    [[SVAR:%.*]] = alloca i32, align 4
 498 // CHECK9-NEXT:    [[I:%.*]] = alloca i32, align 4
 499 // CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
 500 // CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
 501 // CHECK9-NEXT:    store ptr undef, ptr [[_TMP1]], align 8
 502 // CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4
 503 // CHECK9-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 4
 504 // CHECK9-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4
 505 // CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4
 506 // CHECK9-NEXT:    [[ARRAY_BEGIN:%.*]] = getelementptr inbounds [2 x %struct.S], ptr [[S_ARR]], i32 0, i32 0
 507 // CHECK9-NEXT:    [[ARRAYCTOR_END:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAY_BEGIN]], i64 2
 508 // CHECK9-NEXT:    br label [[ARRAYCTOR_LOOP:%.*]]
 509 // CHECK9:       arrayctor.loop:
 510 // CHECK9-NEXT:    [[ARRAYCTOR_CUR:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ]
 511 // CHECK9-NEXT:    call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]])
 512 // CHECK9-NEXT:    [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYCTOR_CUR]], i64 1
 513 // CHECK9-NEXT:    [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], [[ARRAYCTOR_END]]
 514 // CHECK9-NEXT:    br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]]
 515 // CHECK9:       arrayctor.cont:
 516 // CHECK9-NEXT:    call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]])
 517 // CHECK9-NEXT:    store ptr [[VAR]], ptr [[_TMP2]], align 8
 518 // CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
 519 // CHECK9-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
 520 // CHECK9-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP1]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
 521 // CHECK9-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
 522 // CHECK9-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 1
 523 // CHECK9-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
 524 // CHECK9:       cond.true:
 525 // CHECK9-NEXT:    br label [[COND_END:%.*]]
 526 // CHECK9:       cond.false:
 527 // CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
 528 // CHECK9-NEXT:    br label [[COND_END]]
 529 // CHECK9:       cond.end:
 530 // CHECK9-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]
 531 // CHECK9-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4
 532 // CHECK9-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4
 533 // CHECK9-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4
 534 // CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
 535 // CHECK9:       omp.inner.for.cond:
 536 // CHECK9-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
 537 // CHECK9-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
 538 // CHECK9-NEXT:    [[CMP3:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]
 539 // CHECK9-NEXT:    br i1 [[CMP3]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_COND_CLEANUP:%.*]]
 540 // CHECK9:       omp.inner.for.cond.cleanup:
 541 // CHECK9-NEXT:    br label [[OMP_INNER_FOR_END:%.*]]
 542 // CHECK9:       omp.inner.for.body:
 543 // CHECK9-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
 544 // CHECK9-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
 545 // CHECK9-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
 546 // CHECK9-NEXT:    store i32 [[ADD]], ptr [[I]], align 4
 547 // CHECK9-NEXT:    [[TMP8:%.*]] = load i32, ptr [[T_VAR]], align 4
 548 // CHECK9-NEXT:    [[TMP9:%.*]] = load i32, ptr [[I]], align 4
 549 // CHECK9-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP9]] to i64
 550 // CHECK9-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [2 x i32], ptr [[VEC]], i64 0, i64 [[IDXPROM]]
 551 // CHECK9-NEXT:    store i32 [[TMP8]], ptr [[ARRAYIDX]], align 4
 552 // CHECK9-NEXT:    [[TMP10:%.*]] = load ptr, ptr [[_TMP2]], align 8
 553 // CHECK9-NEXT:    [[TMP11:%.*]] = load i32, ptr [[I]], align 4
 554 // CHECK9-NEXT:    [[IDXPROM4:%.*]] = sext i32 [[TMP11]] to i64
 555 // CHECK9-NEXT:    [[ARRAYIDX5:%.*]] = getelementptr inbounds [2 x %struct.S], ptr [[S_ARR]], i64 0, i64 [[IDXPROM4]]
 556 // CHECK9-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[ARRAYIDX5]], ptr align 4 [[TMP10]], i64 4, i1 false)
 557 // CHECK9-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
 558 // CHECK9:       omp.body.continue:
 559 // CHECK9-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
 560 // CHECK9:       omp.inner.for.inc:
 561 // CHECK9-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
 562 // CHECK9-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP12]], 1
 563 // CHECK9-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4
 564 // CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND]]
 565 // CHECK9:       omp.inner.for.end:
 566 // CHECK9-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
 567 // CHECK9:       omp.loop.exit:
 568 // CHECK9-NEXT:    [[TMP13:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
 569 // CHECK9-NEXT:    [[TMP14:%.*]] = load i32, ptr [[TMP13]], align 4
 570 // CHECK9-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP14]])
 571 // CHECK9-NEXT:    call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]]) #[[ATTR4]]
 572 // CHECK9-NEXT:    [[ARRAY_BEGIN7:%.*]] = getelementptr inbounds [2 x %struct.S], ptr [[S_ARR]], i32 0, i32 0
 573 // CHECK9-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAY_BEGIN7]], i64 2
 574 // CHECK9-NEXT:    br label [[ARRAYDESTROY_BODY:%.*]]
 575 // CHECK9:       arraydestroy.body:
 576 // CHECK9-NEXT:    [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP15]], [[OMP_LOOP_EXIT]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
 577 // CHECK9-NEXT:    [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i64 -1
 578 // CHECK9-NEXT:    call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR4]]
 579 // CHECK9-NEXT:    [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN7]]
 580 // CHECK9-NEXT:    br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE8:%.*]], label [[ARRAYDESTROY_BODY]]
 581 // CHECK9:       arraydestroy.done8:
 582 // CHECK9-NEXT:    ret void
 583 //
 584 //
 585 // CHECK9-LABEL: define {{[^@]+}}@_ZN1SIfED1Ev
 586 // CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {
 587 // CHECK9-NEXT:  entry:
 588 // CHECK9-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 589 // CHECK9-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 590 // CHECK9-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 591 // CHECK9-NEXT:    call void @_ZN1SIfED2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]]) #[[ATTR4]]
 592 // CHECK9-NEXT:    ret void
 593 //
 594 //
 595 // CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l102
 596 // CHECK9-SAME: () #[[ATTR3]] {
 597 // CHECK9-NEXT:  entry:
 598 // CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l102.omp_outlined)
 599 // CHECK9-NEXT:    ret void
 600 //
 601 //
 602 // CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l102.omp_outlined
 603 // CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR3]] {
 604 // CHECK9-NEXT:  entry:
 605 // CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
 606 // CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
 607 // CHECK9-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
 608 // CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 4
 609 // CHECK9-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4
 610 // CHECK9-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4
 611 // CHECK9-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
 612 // CHECK9-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
 613 // CHECK9-NEXT:    [[I:%.*]] = alloca i32, align 4
 614 // CHECK9-NEXT:    [[I1:%.*]] = alloca i32, align 4
 615 // CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
 616 // CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
 617 // CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4
 618 // CHECK9-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 4
 619 // CHECK9-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4
 620 // CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4
 621 // CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
 622 // CHECK9-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
 623 // CHECK9-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP1]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
 624 // CHECK9-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
 625 // CHECK9-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 1
 626 // CHECK9-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
 627 // CHECK9:       cond.true:
 628 // CHECK9-NEXT:    br label [[COND_END:%.*]]
 629 // CHECK9:       cond.false:
 630 // CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
 631 // CHECK9-NEXT:    br label [[COND_END]]
 632 // CHECK9:       cond.end:
 633 // CHECK9-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]
 634 // CHECK9-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4
 635 // CHECK9-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4
 636 // CHECK9-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4
 637 // CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
 638 // CHECK9:       omp.inner.for.cond:
 639 // CHECK9-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
 640 // CHECK9-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
 641 // CHECK9-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]
 642 // CHECK9-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
 643 // CHECK9:       omp.inner.for.body:
 644 // CHECK9-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
 645 // CHECK9-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
 646 // CHECK9-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
 647 // CHECK9-NEXT:    store i32 [[ADD]], ptr [[I]], align 4
 648 // CHECK9-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
 649 // CHECK9:       omp.body.continue:
 650 // CHECK9-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
 651 // CHECK9:       omp.inner.for.inc:
 652 // CHECK9-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
 653 // CHECK9-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP8]], 1
 654 // CHECK9-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4
 655 // CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND]]
 656 // CHECK9:       omp.inner.for.end:
 657 // CHECK9-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
 658 // CHECK9:       omp.loop.exit:
 659 // CHECK9-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])
 660 // CHECK9-NEXT:    ret void
 661 //
 662 //
 663 // CHECK9-LABEL: define {{[^@]+}}@_Z5tmainIiET_v
 664 // CHECK9-SAME: () #[[ATTR5:[0-9]+]] comdat {
 665 // CHECK9-NEXT:  entry:
 666 // CHECK9-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4
 667 // CHECK9-NEXT:    [[TEST:%.*]] = alloca [[STRUCT_S_0:%.*]], align 4
 668 // CHECK9-NEXT:    [[T_VAR:%.*]] = alloca i32, align 4
 669 // CHECK9-NEXT:    [[VEC:%.*]] = alloca [2 x i32], align 4
 670 // CHECK9-NEXT:    [[S_ARR:%.*]] = alloca [2 x %struct.S.0], align 4
 671 // CHECK9-NEXT:    [[VAR:%.*]] = alloca ptr, align 8
 672 // CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 4
 673 // CHECK9-NEXT:    [[_TMP1:%.*]] = alloca ptr, align 8
 674 // CHECK9-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8
 675 // CHECK9-NEXT:    call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[TEST]])
 676 // CHECK9-NEXT:    store i32 0, ptr [[T_VAR]], align 4
 677 // CHECK9-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[VEC]], ptr align 4 @__const._Z5tmainIiET_v.vec, i64 8, i1 false)
 678 // CHECK9-NEXT:    [[ARRAYINIT_BEGIN:%.*]] = getelementptr inbounds [2 x %struct.S.0], ptr [[S_ARR]], i64 0, i64 0
 679 // CHECK9-NEXT:    call void @_ZN1SIiEC1Ei(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYINIT_BEGIN]], i32 noundef signext 1)
 680 // CHECK9-NEXT:    [[ARRAYINIT_ELEMENT:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYINIT_BEGIN]], i64 1
 681 // CHECK9-NEXT:    call void @_ZN1SIiEC1Ei(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYINIT_ELEMENT]], i32 noundef signext 2)
 682 // CHECK9-NEXT:    store ptr [[TEST]], ptr [[VAR]], align 8
 683 // CHECK9-NEXT:    store ptr undef, ptr [[_TMP1]], align 8
 684 // CHECK9-NEXT:    [[TMP0:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0
 685 // CHECK9-NEXT:    store i32 2, ptr [[TMP0]], align 4
 686 // CHECK9-NEXT:    [[TMP1:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1
 687 // CHECK9-NEXT:    store i32 0, ptr [[TMP1]], align 4
 688 // CHECK9-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2
 689 // CHECK9-NEXT:    store ptr null, ptr [[TMP2]], align 8
 690 // CHECK9-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3
 691 // CHECK9-NEXT:    store ptr null, ptr [[TMP3]], align 8
 692 // CHECK9-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4
 693 // CHECK9-NEXT:    store ptr null, ptr [[TMP4]], align 8
 694 // CHECK9-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5
 695 // CHECK9-NEXT:    store ptr null, ptr [[TMP5]], align 8
 696 // CHECK9-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6
 697 // CHECK9-NEXT:    store ptr null, ptr [[TMP6]], align 8
 698 // CHECK9-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7
 699 // CHECK9-NEXT:    store ptr null, ptr [[TMP7]], align 8
 700 // CHECK9-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8
 701 // CHECK9-NEXT:    store i64 2, ptr [[TMP8]], align 8
 702 // CHECK9-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9
 703 // CHECK9-NEXT:    store i64 0, ptr [[TMP9]], align 8
 704 // CHECK9-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10
 705 // CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP10]], align 4
 706 // CHECK9-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11
 707 // CHECK9-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP11]], align 4
 708 // CHECK9-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12
 709 // CHECK9-NEXT:    store i32 0, ptr [[TMP12]], align 4
 710 // CHECK9-NEXT:    [[TMP13:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l49.region_id, ptr [[KERNEL_ARGS]])
 711 // CHECK9-NEXT:    [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 0
 712 // CHECK9-NEXT:    br i1 [[TMP14]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
 713 // CHECK9:       omp_offload.failed:
 714 // CHECK9-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l49() #[[ATTR4]]
 715 // CHECK9-NEXT:    br label [[OMP_OFFLOAD_CONT]]
 716 // CHECK9:       omp_offload.cont:
 717 // CHECK9-NEXT:    store i32 0, ptr [[RETVAL]], align 4
 718 // CHECK9-NEXT:    [[ARRAY_BEGIN:%.*]] = getelementptr inbounds [2 x %struct.S.0], ptr [[S_ARR]], i32 0, i32 0
 719 // CHECK9-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAY_BEGIN]], i64 2
 720 // CHECK9-NEXT:    br label [[ARRAYDESTROY_BODY:%.*]]
 721 // CHECK9:       arraydestroy.body:
 722 // CHECK9-NEXT:    [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP15]], [[OMP_OFFLOAD_CONT]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
 723 // CHECK9-NEXT:    [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i64 -1
 724 // CHECK9-NEXT:    call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR4]]
 725 // CHECK9-NEXT:    [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN]]
 726 // CHECK9-NEXT:    br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE2:%.*]], label [[ARRAYDESTROY_BODY]]
 727 // CHECK9:       arraydestroy.done2:
 728 // CHECK9-NEXT:    call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[TEST]]) #[[ATTR4]]
 729 // CHECK9-NEXT:    [[TMP16:%.*]] = load i32, ptr [[RETVAL]], align 4
 730 // CHECK9-NEXT:    ret i32 [[TMP16]]
 731 //
 732 //
 733 // CHECK9-LABEL: define {{[^@]+}}@_ZN1SIfEC2Ev
 734 // CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {
 735 // CHECK9-NEXT:  entry:
 736 // CHECK9-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 737 // CHECK9-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 738 // CHECK9-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 739 // CHECK9-NEXT:    [[F:%.*]] = getelementptr inbounds [[STRUCT_S:%.*]], ptr [[THIS1]], i32 0, i32 0
 740 // CHECK9-NEXT:    store float 0.000000e+00, ptr [[F]], align 4
 741 // CHECK9-NEXT:    ret void
 742 //
 743 //
 744 // CHECK9-LABEL: define {{[^@]+}}@_ZN1SIfEC2Ef
 745 // CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]], float noundef [[A:%.*]]) unnamed_addr #[[ATTR1]] comdat {
 746 // CHECK9-NEXT:  entry:
 747 // CHECK9-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 748 // CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca float, align 4
 749 // CHECK9-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 750 // CHECK9-NEXT:    store float [[A]], ptr [[A_ADDR]], align 4
 751 // CHECK9-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 752 // CHECK9-NEXT:    [[F:%.*]] = getelementptr inbounds [[STRUCT_S:%.*]], ptr [[THIS1]], i32 0, i32 0
 753 // CHECK9-NEXT:    [[TMP0:%.*]] = load float, ptr [[A_ADDR]], align 4
 754 // CHECK9-NEXT:    store float [[TMP0]], ptr [[F]], align 4
 755 // CHECK9-NEXT:    ret void
 756 //
 757 //
 758 // CHECK9-LABEL: define {{[^@]+}}@_ZN1SIfED2Ev
 759 // CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {
 760 // CHECK9-NEXT:  entry:
 761 // CHECK9-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 762 // CHECK9-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 763 // CHECK9-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 764 // CHECK9-NEXT:    ret void
 765 //
 766 //
 767 // CHECK9-LABEL: define {{[^@]+}}@_ZN1SIiEC1Ev
 768 // CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {
 769 // CHECK9-NEXT:  entry:
 770 // CHECK9-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 771 // CHECK9-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 772 // CHECK9-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 773 // CHECK9-NEXT:    call void @_ZN1SIiEC2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]])
 774 // CHECK9-NEXT:    ret void
 775 //
 776 //
 777 // CHECK9-LABEL: define {{[^@]+}}@_ZN1SIiEC1Ei
 778 // CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]], i32 noundef signext [[A:%.*]]) unnamed_addr #[[ATTR1]] comdat {
 779 // CHECK9-NEXT:  entry:
 780 // CHECK9-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 781 // CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 4
 782 // CHECK9-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 783 // CHECK9-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 4
 784 // CHECK9-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 785 // CHECK9-NEXT:    [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 4
 786 // CHECK9-NEXT:    call void @_ZN1SIiEC2Ei(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]], i32 noundef signext [[TMP0]])
 787 // CHECK9-NEXT:    ret void
 788 //
 789 //
 790 // CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l49
 791 // CHECK9-SAME: () #[[ATTR3]] {
 792 // CHECK9-NEXT:  entry:
 793 // CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l49.omp_outlined)
 794 // CHECK9-NEXT:    ret void
 795 //
 796 //
 797 // CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l49.omp_outlined
 798 // CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR3]] {
 799 // CHECK9-NEXT:  entry:
 800 // CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
 801 // CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
 802 // CHECK9-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
 803 // CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 4
 804 // CHECK9-NEXT:    [[_TMP1:%.*]] = alloca ptr, align 8
 805 // CHECK9-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4
 806 // CHECK9-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4
 807 // CHECK9-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
 808 // CHECK9-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
 809 // CHECK9-NEXT:    [[T_VAR:%.*]] = alloca i32, align 4
 810 // CHECK9-NEXT:    [[VEC:%.*]] = alloca [2 x i32], align 4
 811 // CHECK9-NEXT:    [[S_ARR:%.*]] = alloca [2 x %struct.S.0], align 4
 812 // CHECK9-NEXT:    [[VAR:%.*]] = alloca [[STRUCT_S_0:%.*]], align 4
 813 // CHECK9-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 8
 814 // CHECK9-NEXT:    [[I:%.*]] = alloca i32, align 4
 815 // CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
 816 // CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
 817 // CHECK9-NEXT:    store ptr undef, ptr [[_TMP1]], align 8
 818 // CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4
 819 // CHECK9-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 4
 820 // CHECK9-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4
 821 // CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4
 822 // CHECK9-NEXT:    [[ARRAY_BEGIN:%.*]] = getelementptr inbounds [2 x %struct.S.0], ptr [[S_ARR]], i32 0, i32 0
 823 // CHECK9-NEXT:    [[ARRAYCTOR_END:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAY_BEGIN]], i64 2
 824 // CHECK9-NEXT:    br label [[ARRAYCTOR_LOOP:%.*]]
 825 // CHECK9:       arrayctor.loop:
 826 // CHECK9-NEXT:    [[ARRAYCTOR_CUR:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ]
 827 // CHECK9-NEXT:    call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]])
 828 // CHECK9-NEXT:    [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYCTOR_CUR]], i64 1
 829 // CHECK9-NEXT:    [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], [[ARRAYCTOR_END]]
 830 // CHECK9-NEXT:    br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]]
 831 // CHECK9:       arrayctor.cont:
 832 // CHECK9-NEXT:    call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]])
 833 // CHECK9-NEXT:    store ptr [[VAR]], ptr [[_TMP2]], align 8
 834 // CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
 835 // CHECK9-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
 836 // CHECK9-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP1]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
 837 // CHECK9-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
 838 // CHECK9-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 1
 839 // CHECK9-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
 840 // CHECK9:       cond.true:
 841 // CHECK9-NEXT:    br label [[COND_END:%.*]]
 842 // CHECK9:       cond.false:
 843 // CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
 844 // CHECK9-NEXT:    br label [[COND_END]]
 845 // CHECK9:       cond.end:
 846 // CHECK9-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]
 847 // CHECK9-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4
 848 // CHECK9-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4
 849 // CHECK9-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4
 850 // CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
 851 // CHECK9:       omp.inner.for.cond:
 852 // CHECK9-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
 853 // CHECK9-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
 854 // CHECK9-NEXT:    [[CMP3:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]
 855 // CHECK9-NEXT:    br i1 [[CMP3]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_COND_CLEANUP:%.*]]
 856 // CHECK9:       omp.inner.for.cond.cleanup:
 857 // CHECK9-NEXT:    br label [[OMP_INNER_FOR_END:%.*]]
 858 // CHECK9:       omp.inner.for.body:
 859 // CHECK9-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
 860 // CHECK9-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
 861 // CHECK9-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
 862 // CHECK9-NEXT:    store i32 [[ADD]], ptr [[I]], align 4
 863 // CHECK9-NEXT:    [[TMP8:%.*]] = load i32, ptr [[T_VAR]], align 4
 864 // CHECK9-NEXT:    [[TMP9:%.*]] = load i32, ptr [[I]], align 4
 865 // CHECK9-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP9]] to i64
 866 // CHECK9-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [2 x i32], ptr [[VEC]], i64 0, i64 [[IDXPROM]]
 867 // CHECK9-NEXT:    store i32 [[TMP8]], ptr [[ARRAYIDX]], align 4
 868 // CHECK9-NEXT:    [[TMP10:%.*]] = load ptr, ptr [[_TMP2]], align 8
 869 // CHECK9-NEXT:    [[TMP11:%.*]] = load i32, ptr [[I]], align 4
 870 // CHECK9-NEXT:    [[IDXPROM4:%.*]] = sext i32 [[TMP11]] to i64
 871 // CHECK9-NEXT:    [[ARRAYIDX5:%.*]] = getelementptr inbounds [2 x %struct.S.0], ptr [[S_ARR]], i64 0, i64 [[IDXPROM4]]
 872 // CHECK9-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[ARRAYIDX5]], ptr align 4 [[TMP10]], i64 4, i1 false)
 873 // CHECK9-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
 874 // CHECK9:       omp.body.continue:
 875 // CHECK9-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
 876 // CHECK9:       omp.inner.for.inc:
 877 // CHECK9-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
 878 // CHECK9-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP12]], 1
 879 // CHECK9-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4
 880 // CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND]]
 881 // CHECK9:       omp.inner.for.end:
 882 // CHECK9-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
 883 // CHECK9:       omp.loop.exit:
 884 // CHECK9-NEXT:    [[TMP13:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
 885 // CHECK9-NEXT:    [[TMP14:%.*]] = load i32, ptr [[TMP13]], align 4
 886 // CHECK9-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP14]])
 887 // CHECK9-NEXT:    call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]]) #[[ATTR4]]
 888 // CHECK9-NEXT:    [[ARRAY_BEGIN7:%.*]] = getelementptr inbounds [2 x %struct.S.0], ptr [[S_ARR]], i32 0, i32 0
 889 // CHECK9-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAY_BEGIN7]], i64 2
 890 // CHECK9-NEXT:    br label [[ARRAYDESTROY_BODY:%.*]]
 891 // CHECK9:       arraydestroy.body:
 892 // CHECK9-NEXT:    [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP15]], [[OMP_LOOP_EXIT]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
 893 // CHECK9-NEXT:    [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i64 -1
 894 // CHECK9-NEXT:    call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR4]]
 895 // CHECK9-NEXT:    [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN7]]
 896 // CHECK9-NEXT:    br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE8:%.*]], label [[ARRAYDESTROY_BODY]]
 897 // CHECK9:       arraydestroy.done8:
 898 // CHECK9-NEXT:    ret void
 899 //
 900 //
 901 // CHECK9-LABEL: define {{[^@]+}}@_ZN1SIiED1Ev
 902 // CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {
 903 // CHECK9-NEXT:  entry:
 904 // CHECK9-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 905 // CHECK9-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 906 // CHECK9-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 907 // CHECK9-NEXT:    call void @_ZN1SIiED2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]]) #[[ATTR4]]
 908 // CHECK9-NEXT:    ret void
 909 //
 910 //
 911 // CHECK9-LABEL: define {{[^@]+}}@_ZN1SIiEC2Ev
 912 // CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {
 913 // CHECK9-NEXT:  entry:
 914 // CHECK9-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 915 // CHECK9-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 916 // CHECK9-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 917 // CHECK9-NEXT:    [[F:%.*]] = getelementptr inbounds [[STRUCT_S_0:%.*]], ptr [[THIS1]], i32 0, i32 0
 918 // CHECK9-NEXT:    store i32 0, ptr [[F]], align 4
 919 // CHECK9-NEXT:    ret void
 920 //
 921 //
 922 // CHECK9-LABEL: define {{[^@]+}}@_ZN1SIiEC2Ei
 923 // CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]], i32 noundef signext [[A:%.*]]) unnamed_addr #[[ATTR1]] comdat {
 924 // CHECK9-NEXT:  entry:
 925 // CHECK9-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 926 // CHECK9-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 4
 927 // CHECK9-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 928 // CHECK9-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 4
 929 // CHECK9-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 930 // CHECK9-NEXT:    [[F:%.*]] = getelementptr inbounds [[STRUCT_S_0:%.*]], ptr [[THIS1]], i32 0, i32 0
 931 // CHECK9-NEXT:    [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 4
 932 // CHECK9-NEXT:    store i32 [[TMP0]], ptr [[F]], align 4
 933 // CHECK9-NEXT:    ret void
 934 //
 935 //
 936 // CHECK9-LABEL: define {{[^@]+}}@_ZN1SIiED2Ev
 937 // CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat {
 938 // CHECK9-NEXT:  entry:
 939 // CHECK9-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 940 // CHECK9-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 941 // CHECK9-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 942 // CHECK9-NEXT:    ret void
 943 //
 944 //
 945 // CHECK9-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg
 946 // CHECK9-SAME: () #[[ATTR6:[0-9]+]] {
 947 // CHECK9-NEXT:  entry:
 948 // CHECK9-NEXT:    call void @__tgt_register_requires(i64 1)
 949 // CHECK9-NEXT:    ret void
 950 //
 951 //
 952 // CHECK11-LABEL: define {{[^@]+}}@main
 953 // CHECK11-SAME: () #[[ATTR0:[0-9]+]] {
 954 // CHECK11-NEXT:  entry:
 955 // CHECK11-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4
 956 // CHECK11-NEXT:    [[G:%.*]] = alloca double, align 8
 957 // CHECK11-NEXT:    [[G1:%.*]] = alloca ptr, align 4
 958 // CHECK11-NEXT:    [[TEST:%.*]] = alloca [[STRUCT_S:%.*]], align 4
 959 // CHECK11-NEXT:    [[T_VAR:%.*]] = alloca i32, align 4
 960 // CHECK11-NEXT:    [[VEC:%.*]] = alloca [2 x i32], align 4
 961 // CHECK11-NEXT:    [[S_ARR:%.*]] = alloca [2 x %struct.S], align 4
 962 // CHECK11-NEXT:    [[VAR:%.*]] = alloca ptr, align 4
 963 // CHECK11-NEXT:    [[TMP:%.*]] = alloca i32, align 4
 964 // CHECK11-NEXT:    [[_TMP1:%.*]] = alloca ptr, align 4
 965 // CHECK11-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8
 966 // CHECK11-NEXT:    [[I:%.*]] = alloca i32, align 4
 967 // CHECK11-NEXT:    [[_TMP2:%.*]] = alloca i32, align 4
 968 // CHECK11-NEXT:    [[KERNEL_ARGS3:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8
 969 // CHECK11-NEXT:    store i32 0, ptr [[RETVAL]], align 4
 970 // CHECK11-NEXT:    store ptr [[G]], ptr [[G1]], align 4
 971 // CHECK11-NEXT:    call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[TEST]])
 972 // CHECK11-NEXT:    store i32 0, ptr [[T_VAR]], align 4
 973 // CHECK11-NEXT:    call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[VEC]], ptr align 4 @__const.main.vec, i32 8, i1 false)
 974 // CHECK11-NEXT:    [[ARRAYINIT_BEGIN:%.*]] = getelementptr inbounds [2 x %struct.S], ptr [[S_ARR]], i32 0, i32 0
 975 // CHECK11-NEXT:    call void @_ZN1SIfEC1Ef(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYINIT_BEGIN]], float noundef 1.000000e+00)
 976 // CHECK11-NEXT:    [[ARRAYINIT_ELEMENT:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYINIT_BEGIN]], i32 1
 977 // CHECK11-NEXT:    call void @_ZN1SIfEC1Ef(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYINIT_ELEMENT]], float noundef 2.000000e+00)
 978 // CHECK11-NEXT:    store ptr [[TEST]], ptr [[VAR]], align 4
 979 // CHECK11-NEXT:    store ptr undef, ptr [[_TMP1]], align 4
 980 // CHECK11-NEXT:    [[TMP0:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0
 981 // CHECK11-NEXT:    store i32 2, ptr [[TMP0]], align 4
 982 // CHECK11-NEXT:    [[TMP1:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1
 983 // CHECK11-NEXT:    store i32 0, ptr [[TMP1]], align 4
 984 // CHECK11-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2
 985 // CHECK11-NEXT:    store ptr null, ptr [[TMP2]], align 4
 986 // CHECK11-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3
 987 // CHECK11-NEXT:    store ptr null, ptr [[TMP3]], align 4
 988 // CHECK11-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4
 989 // CHECK11-NEXT:    store ptr null, ptr [[TMP4]], align 4
 990 // CHECK11-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5
 991 // CHECK11-NEXT:    store ptr null, ptr [[TMP5]], align 4
 992 // CHECK11-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6
 993 // CHECK11-NEXT:    store ptr null, ptr [[TMP6]], align 4
 994 // CHECK11-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7
 995 // CHECK11-NEXT:    store ptr null, ptr [[TMP7]], align 4
 996 // CHECK11-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8
 997 // CHECK11-NEXT:    store i64 2, ptr [[TMP8]], align 8
 998 // CHECK11-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9
 999 // CHECK11-NEXT:    store i64 0, ptr [[TMP9]], align 8
1000 // CHECK11-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10
1001 // CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP10]], align 4
1002 // CHECK11-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11
1003 // CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP11]], align 4
1004 // CHECK11-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12
1005 // CHECK11-NEXT:    store i32 0, ptr [[TMP12]], align 4
1006 // CHECK11-NEXT:    [[TMP13:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l93.region_id, ptr [[KERNEL_ARGS]])
1007 // CHECK11-NEXT:    [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 0
1008 // CHECK11-NEXT:    br i1 [[TMP14]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
1009 // CHECK11:       omp_offload.failed:
1010 // CHECK11-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l93() #[[ATTR4:[0-9]+]]
1011 // CHECK11-NEXT:    br label [[OMP_OFFLOAD_CONT]]
1012 // CHECK11:       omp_offload.cont:
1013 // CHECK11-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 0
1014 // CHECK11-NEXT:    store i32 2, ptr [[TMP15]], align 4
1015 // CHECK11-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 1
1016 // CHECK11-NEXT:    store i32 0, ptr [[TMP16]], align 4
1017 // CHECK11-NEXT:    [[TMP17:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 2
1018 // CHECK11-NEXT:    store ptr null, ptr [[TMP17]], align 4
1019 // CHECK11-NEXT:    [[TMP18:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 3
1020 // CHECK11-NEXT:    store ptr null, ptr [[TMP18]], align 4
1021 // CHECK11-NEXT:    [[TMP19:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 4
1022 // CHECK11-NEXT:    store ptr null, ptr [[TMP19]], align 4
1023 // CHECK11-NEXT:    [[TMP20:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 5
1024 // CHECK11-NEXT:    store ptr null, ptr [[TMP20]], align 4
1025 // CHECK11-NEXT:    [[TMP21:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 6
1026 // CHECK11-NEXT:    store ptr null, ptr [[TMP21]], align 4
1027 // CHECK11-NEXT:    [[TMP22:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 7
1028 // CHECK11-NEXT:    store ptr null, ptr [[TMP22]], align 4
1029 // CHECK11-NEXT:    [[TMP23:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 8
1030 // CHECK11-NEXT:    store i64 2, ptr [[TMP23]], align 8
1031 // CHECK11-NEXT:    [[TMP24:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 9
1032 // CHECK11-NEXT:    store i64 0, ptr [[TMP24]], align 8
1033 // CHECK11-NEXT:    [[TMP25:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 10
1034 // CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP25]], align 4
1035 // CHECK11-NEXT:    [[TMP26:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 11
1036 // CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP26]], align 4
1037 // CHECK11-NEXT:    [[TMP27:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS3]], i32 0, i32 12
1038 // CHECK11-NEXT:    store i32 0, ptr [[TMP27]], align 4
1039 // CHECK11-NEXT:    [[TMP28:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l102.region_id, ptr [[KERNEL_ARGS3]])
1040 // CHECK11-NEXT:    [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 0
1041 // CHECK11-NEXT:    br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED4:%.*]], label [[OMP_OFFLOAD_CONT5:%.*]]
1042 // CHECK11:       omp_offload.failed4:
1043 // CHECK11-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l102() #[[ATTR4]]
1044 // CHECK11-NEXT:    br label [[OMP_OFFLOAD_CONT5]]
1045 // CHECK11:       omp_offload.cont5:
1046 // CHECK11-NEXT:    [[CALL:%.*]] = call noundef i32 @_Z5tmainIiET_v()
1047 // CHECK11-NEXT:    store i32 [[CALL]], ptr [[RETVAL]], align 4
1048 // CHECK11-NEXT:    [[ARRAY_BEGIN:%.*]] = getelementptr inbounds [2 x %struct.S], ptr [[S_ARR]], i32 0, i32 0
1049 // CHECK11-NEXT:    [[TMP30:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAY_BEGIN]], i32 2
1050 // CHECK11-NEXT:    br label [[ARRAYDESTROY_BODY:%.*]]
1051 // CHECK11:       arraydestroy.body:
1052 // CHECK11-NEXT:    [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP30]], [[OMP_OFFLOAD_CONT5]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
1053 // CHECK11-NEXT:    [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i32 -1
1054 // CHECK11-NEXT:    call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR4]]
1055 // CHECK11-NEXT:    [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN]]
1056 // CHECK11-NEXT:    br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE6:%.*]], label [[ARRAYDESTROY_BODY]]
1057 // CHECK11:       arraydestroy.done6:
1058 // CHECK11-NEXT:    call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[TEST]]) #[[ATTR4]]
1059 // CHECK11-NEXT:    [[TMP31:%.*]] = load i32, ptr [[RETVAL]], align 4
1060 // CHECK11-NEXT:    ret i32 [[TMP31]]
1061 //
1062 //
1063 // CHECK11-LABEL: define {{[^@]+}}@_ZN1SIfEC1Ev
1064 // CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1:[0-9]+]] comdat align 2 {
1065 // CHECK11-NEXT:  entry:
1066 // CHECK11-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4
1067 // CHECK11-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
1068 // CHECK11-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
1069 // CHECK11-NEXT:    call void @_ZN1SIfEC2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]])
1070 // CHECK11-NEXT:    ret void
1071 //
1072 //
1073 // CHECK11-LABEL: define {{[^@]+}}@_ZN1SIfEC1Ef
1074 // CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]], float noundef [[A:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {
1075 // CHECK11-NEXT:  entry:
1076 // CHECK11-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4
1077 // CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca float, align 4
1078 // CHECK11-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
1079 // CHECK11-NEXT:    store float [[A]], ptr [[A_ADDR]], align 4
1080 // CHECK11-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
1081 // CHECK11-NEXT:    [[TMP0:%.*]] = load float, ptr [[A_ADDR]], align 4
1082 // CHECK11-NEXT:    call void @_ZN1SIfEC2Ef(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]], float noundef [[TMP0]])
1083 // CHECK11-NEXT:    ret void
1084 //
1085 //
1086 // CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l93
1087 // CHECK11-SAME: () #[[ATTR3:[0-9]+]] {
1088 // CHECK11-NEXT:  entry:
1089 // CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l93.omp_outlined)
1090 // CHECK11-NEXT:    ret void
1091 //
1092 //
1093 // CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l93.omp_outlined
1094 // CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR3]] {
1095 // CHECK11-NEXT:  entry:
1096 // CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
1097 // CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
1098 // CHECK11-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
1099 // CHECK11-NEXT:    [[TMP:%.*]] = alloca i32, align 4
1100 // CHECK11-NEXT:    [[_TMP1:%.*]] = alloca ptr, align 4
1101 // CHECK11-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4
1102 // CHECK11-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4
1103 // CHECK11-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
1104 // CHECK11-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
1105 // CHECK11-NEXT:    [[T_VAR:%.*]] = alloca i32, align 4
1106 // CHECK11-NEXT:    [[VEC:%.*]] = alloca [2 x i32], align 4
1107 // CHECK11-NEXT:    [[S_ARR:%.*]] = alloca [2 x %struct.S], align 4
1108 // CHECK11-NEXT:    [[VAR:%.*]] = alloca [[STRUCT_S:%.*]], align 4
1109 // CHECK11-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 4
1110 // CHECK11-NEXT:    [[SVAR:%.*]] = alloca i32, align 4
1111 // CHECK11-NEXT:    [[I:%.*]] = alloca i32, align 4
1112 // CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
1113 // CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
1114 // CHECK11-NEXT:    store ptr undef, ptr [[_TMP1]], align 4
1115 // CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4
1116 // CHECK11-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 4
1117 // CHECK11-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4
1118 // CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4
1119 // CHECK11-NEXT:    [[ARRAY_BEGIN:%.*]] = getelementptr inbounds [2 x %struct.S], ptr [[S_ARR]], i32 0, i32 0
1120 // CHECK11-NEXT:    [[ARRAYCTOR_END:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAY_BEGIN]], i32 2
1121 // CHECK11-NEXT:    br label [[ARRAYCTOR_LOOP:%.*]]
1122 // CHECK11:       arrayctor.loop:
1123 // CHECK11-NEXT:    [[ARRAYCTOR_CUR:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ]
1124 // CHECK11-NEXT:    call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]])
1125 // CHECK11-NEXT:    [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYCTOR_CUR]], i32 1
1126 // CHECK11-NEXT:    [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], [[ARRAYCTOR_END]]
1127 // CHECK11-NEXT:    br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]]
1128 // CHECK11:       arrayctor.cont:
1129 // CHECK11-NEXT:    call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]])
1130 // CHECK11-NEXT:    store ptr [[VAR]], ptr [[_TMP2]], align 4
1131 // CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
1132 // CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
1133 // CHECK11-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP1]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
1134 // CHECK11-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
1135 // CHECK11-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 1
1136 // CHECK11-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
1137 // CHECK11:       cond.true:
1138 // CHECK11-NEXT:    br label [[COND_END:%.*]]
1139 // CHECK11:       cond.false:
1140 // CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
1141 // CHECK11-NEXT:    br label [[COND_END]]
1142 // CHECK11:       cond.end:
1143 // CHECK11-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]
1144 // CHECK11-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4
1145 // CHECK11-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4
1146 // CHECK11-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4
1147 // CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
1148 // CHECK11:       omp.inner.for.cond:
1149 // CHECK11-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
1150 // CHECK11-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
1151 // CHECK11-NEXT:    [[CMP3:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]
1152 // CHECK11-NEXT:    br i1 [[CMP3]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_COND_CLEANUP:%.*]]
1153 // CHECK11:       omp.inner.for.cond.cleanup:
1154 // CHECK11-NEXT:    br label [[OMP_INNER_FOR_END:%.*]]
1155 // CHECK11:       omp.inner.for.body:
1156 // CHECK11-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
1157 // CHECK11-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
1158 // CHECK11-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
1159 // CHECK11-NEXT:    store i32 [[ADD]], ptr [[I]], align 4
1160 // CHECK11-NEXT:    [[TMP8:%.*]] = load i32, ptr [[T_VAR]], align 4
1161 // CHECK11-NEXT:    [[TMP9:%.*]] = load i32, ptr [[I]], align 4
1162 // CHECK11-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [2 x i32], ptr [[VEC]], i32 0, i32 [[TMP9]]
1163 // CHECK11-NEXT:    store i32 [[TMP8]], ptr [[ARRAYIDX]], align 4
1164 // CHECK11-NEXT:    [[TMP10:%.*]] = load ptr, ptr [[_TMP2]], align 4
1165 // CHECK11-NEXT:    [[TMP11:%.*]] = load i32, ptr [[I]], align 4
1166 // CHECK11-NEXT:    [[ARRAYIDX4:%.*]] = getelementptr inbounds [2 x %struct.S], ptr [[S_ARR]], i32 0, i32 [[TMP11]]
1167 // CHECK11-NEXT:    call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[ARRAYIDX4]], ptr align 4 [[TMP10]], i32 4, i1 false)
1168 // CHECK11-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
1169 // CHECK11:       omp.body.continue:
1170 // CHECK11-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
1171 // CHECK11:       omp.inner.for.inc:
1172 // CHECK11-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
1173 // CHECK11-NEXT:    [[ADD5:%.*]] = add nsw i32 [[TMP12]], 1
1174 // CHECK11-NEXT:    store i32 [[ADD5]], ptr [[DOTOMP_IV]], align 4
1175 // CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND]]
1176 // CHECK11:       omp.inner.for.end:
1177 // CHECK11-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
1178 // CHECK11:       omp.loop.exit:
1179 // CHECK11-NEXT:    [[TMP13:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
1180 // CHECK11-NEXT:    [[TMP14:%.*]] = load i32, ptr [[TMP13]], align 4
1181 // CHECK11-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP14]])
1182 // CHECK11-NEXT:    call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]]) #[[ATTR4]]
1183 // CHECK11-NEXT:    [[ARRAY_BEGIN6:%.*]] = getelementptr inbounds [2 x %struct.S], ptr [[S_ARR]], i32 0, i32 0
1184 // CHECK11-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAY_BEGIN6]], i32 2
1185 // CHECK11-NEXT:    br label [[ARRAYDESTROY_BODY:%.*]]
1186 // CHECK11:       arraydestroy.body:
1187 // CHECK11-NEXT:    [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP15]], [[OMP_LOOP_EXIT]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
1188 // CHECK11-NEXT:    [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i32 -1
1189 // CHECK11-NEXT:    call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR4]]
1190 // CHECK11-NEXT:    [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN6]]
1191 // CHECK11-NEXT:    br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE7:%.*]], label [[ARRAYDESTROY_BODY]]
1192 // CHECK11:       arraydestroy.done7:
1193 // CHECK11-NEXT:    ret void
1194 //
1195 //
1196 // CHECK11-LABEL: define {{[^@]+}}@_ZN1SIfED1Ev
1197 // CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {
1198 // CHECK11-NEXT:  entry:
1199 // CHECK11-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4
1200 // CHECK11-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
1201 // CHECK11-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
1202 // CHECK11-NEXT:    call void @_ZN1SIfED2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]]) #[[ATTR4]]
1203 // CHECK11-NEXT:    ret void
1204 //
1205 //
1206 // CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l102
1207 // CHECK11-SAME: () #[[ATTR3]] {
1208 // CHECK11-NEXT:  entry:
1209 // CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l102.omp_outlined)
1210 // CHECK11-NEXT:    ret void
1211 //
1212 //
1213 // CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l102.omp_outlined
1214 // CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR3]] {
1215 // CHECK11-NEXT:  entry:
1216 // CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
1217 // CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
1218 // CHECK11-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
1219 // CHECK11-NEXT:    [[TMP:%.*]] = alloca i32, align 4
1220 // CHECK11-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4
1221 // CHECK11-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4
1222 // CHECK11-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
1223 // CHECK11-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
1224 // CHECK11-NEXT:    [[I:%.*]] = alloca i32, align 4
1225 // CHECK11-NEXT:    [[I1:%.*]] = alloca i32, align 4
1226 // CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
1227 // CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
1228 // CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4
1229 // CHECK11-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 4
1230 // CHECK11-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4
1231 // CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4
1232 // CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
1233 // CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
1234 // CHECK11-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP1]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
1235 // CHECK11-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
1236 // CHECK11-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 1
1237 // CHECK11-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
1238 // CHECK11:       cond.true:
1239 // CHECK11-NEXT:    br label [[COND_END:%.*]]
1240 // CHECK11:       cond.false:
1241 // CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
1242 // CHECK11-NEXT:    br label [[COND_END]]
1243 // CHECK11:       cond.end:
1244 // CHECK11-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]
1245 // CHECK11-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4
1246 // CHECK11-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4
1247 // CHECK11-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4
1248 // CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
1249 // CHECK11:       omp.inner.for.cond:
1250 // CHECK11-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
1251 // CHECK11-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
1252 // CHECK11-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]
1253 // CHECK11-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
1254 // CHECK11:       omp.inner.for.body:
1255 // CHECK11-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
1256 // CHECK11-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
1257 // CHECK11-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
1258 // CHECK11-NEXT:    store i32 [[ADD]], ptr [[I]], align 4
1259 // CHECK11-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
1260 // CHECK11:       omp.body.continue:
1261 // CHECK11-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
1262 // CHECK11:       omp.inner.for.inc:
1263 // CHECK11-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
1264 // CHECK11-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP8]], 1
1265 // CHECK11-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4
1266 // CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND]]
1267 // CHECK11:       omp.inner.for.end:
1268 // CHECK11-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
1269 // CHECK11:       omp.loop.exit:
1270 // CHECK11-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP1]])
1271 // CHECK11-NEXT:    ret void
1272 //
1273 //
1274 // CHECK11-LABEL: define {{[^@]+}}@_Z5tmainIiET_v
1275 // CHECK11-SAME: () #[[ATTR5:[0-9]+]] comdat {
1276 // CHECK11-NEXT:  entry:
1277 // CHECK11-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4
1278 // CHECK11-NEXT:    [[TEST:%.*]] = alloca [[STRUCT_S_0:%.*]], align 4
1279 // CHECK11-NEXT:    [[T_VAR:%.*]] = alloca i32, align 4
1280 // CHECK11-NEXT:    [[VEC:%.*]] = alloca [2 x i32], align 4
1281 // CHECK11-NEXT:    [[S_ARR:%.*]] = alloca [2 x %struct.S.0], align 4
1282 // CHECK11-NEXT:    [[VAR:%.*]] = alloca ptr, align 4
1283 // CHECK11-NEXT:    [[TMP:%.*]] = alloca i32, align 4
1284 // CHECK11-NEXT:    [[_TMP1:%.*]] = alloca ptr, align 4
1285 // CHECK11-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8
1286 // CHECK11-NEXT:    call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[TEST]])
1287 // CHECK11-NEXT:    store i32 0, ptr [[T_VAR]], align 4
1288 // CHECK11-NEXT:    call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[VEC]], ptr align 4 @__const._Z5tmainIiET_v.vec, i32 8, i1 false)
1289 // CHECK11-NEXT:    [[ARRAYINIT_BEGIN:%.*]] = getelementptr inbounds [2 x %struct.S.0], ptr [[S_ARR]], i32 0, i32 0
1290 // CHECK11-NEXT:    call void @_ZN1SIiEC1Ei(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYINIT_BEGIN]], i32 noundef 1)
1291 // CHECK11-NEXT:    [[ARRAYINIT_ELEMENT:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYINIT_BEGIN]], i32 1
1292 // CHECK11-NEXT:    call void @_ZN1SIiEC1Ei(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYINIT_ELEMENT]], i32 noundef 2)
1293 // CHECK11-NEXT:    store ptr [[TEST]], ptr [[VAR]], align 4
1294 // CHECK11-NEXT:    store ptr undef, ptr [[_TMP1]], align 4
1295 // CHECK11-NEXT:    [[TMP0:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0
1296 // CHECK11-NEXT:    store i32 2, ptr [[TMP0]], align 4
1297 // CHECK11-NEXT:    [[TMP1:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1
1298 // CHECK11-NEXT:    store i32 0, ptr [[TMP1]], align 4
1299 // CHECK11-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2
1300 // CHECK11-NEXT:    store ptr null, ptr [[TMP2]], align 4
1301 // CHECK11-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3
1302 // CHECK11-NEXT:    store ptr null, ptr [[TMP3]], align 4
1303 // CHECK11-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4
1304 // CHECK11-NEXT:    store ptr null, ptr [[TMP4]], align 4
1305 // CHECK11-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5
1306 // CHECK11-NEXT:    store ptr null, ptr [[TMP5]], align 4
1307 // CHECK11-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6
1308 // CHECK11-NEXT:    store ptr null, ptr [[TMP6]], align 4
1309 // CHECK11-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7
1310 // CHECK11-NEXT:    store ptr null, ptr [[TMP7]], align 4
1311 // CHECK11-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8
1312 // CHECK11-NEXT:    store i64 2, ptr [[TMP8]], align 8
1313 // CHECK11-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9
1314 // CHECK11-NEXT:    store i64 0, ptr [[TMP9]], align 8
1315 // CHECK11-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10
1316 // CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP10]], align 4
1317 // CHECK11-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11
1318 // CHECK11-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP11]], align 4
1319 // CHECK11-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12
1320 // CHECK11-NEXT:    store i32 0, ptr [[TMP12]], align 4
1321 // CHECK11-NEXT:    [[TMP13:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l49.region_id, ptr [[KERNEL_ARGS]])
1322 // CHECK11-NEXT:    [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 0
1323 // CHECK11-NEXT:    br i1 [[TMP14]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
1324 // CHECK11:       omp_offload.failed:
1325 // CHECK11-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l49() #[[ATTR4]]
1326 // CHECK11-NEXT:    br label [[OMP_OFFLOAD_CONT]]
1327 // CHECK11:       omp_offload.cont:
1328 // CHECK11-NEXT:    store i32 0, ptr [[RETVAL]], align 4
1329 // CHECK11-NEXT:    [[ARRAY_BEGIN:%.*]] = getelementptr inbounds [2 x %struct.S.0], ptr [[S_ARR]], i32 0, i32 0
1330 // CHECK11-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAY_BEGIN]], i32 2
1331 // CHECK11-NEXT:    br label [[ARRAYDESTROY_BODY:%.*]]
1332 // CHECK11:       arraydestroy.body:
1333 // CHECK11-NEXT:    [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP15]], [[OMP_OFFLOAD_CONT]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
1334 // CHECK11-NEXT:    [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i32 -1
1335 // CHECK11-NEXT:    call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR4]]
1336 // CHECK11-NEXT:    [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN]]
1337 // CHECK11-NEXT:    br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE2:%.*]], label [[ARRAYDESTROY_BODY]]
1338 // CHECK11:       arraydestroy.done2:
1339 // CHECK11-NEXT:    call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[TEST]]) #[[ATTR4]]
1340 // CHECK11-NEXT:    [[TMP16:%.*]] = load i32, ptr [[RETVAL]], align 4
1341 // CHECK11-NEXT:    ret i32 [[TMP16]]
1342 //
1343 //
1344 // CHECK11-LABEL: define {{[^@]+}}@_ZN1SIfEC2Ev
1345 // CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {
1346 // CHECK11-NEXT:  entry:
1347 // CHECK11-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4
1348 // CHECK11-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
1349 // CHECK11-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
1350 // CHECK11-NEXT:    [[F:%.*]] = getelementptr inbounds [[STRUCT_S:%.*]], ptr [[THIS1]], i32 0, i32 0
1351 // CHECK11-NEXT:    store float 0.000000e+00, ptr [[F]], align 4
1352 // CHECK11-NEXT:    ret void
1353 //
1354 //
1355 // CHECK11-LABEL: define {{[^@]+}}@_ZN1SIfEC2Ef
1356 // CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]], float noundef [[A:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {
1357 // CHECK11-NEXT:  entry:
1358 // CHECK11-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4
1359 // CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca float, align 4
1360 // CHECK11-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
1361 // CHECK11-NEXT:    store float [[A]], ptr [[A_ADDR]], align 4
1362 // CHECK11-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
1363 // CHECK11-NEXT:    [[F:%.*]] = getelementptr inbounds [[STRUCT_S:%.*]], ptr [[THIS1]], i32 0, i32 0
1364 // CHECK11-NEXT:    [[TMP0:%.*]] = load float, ptr [[A_ADDR]], align 4
1365 // CHECK11-NEXT:    store float [[TMP0]], ptr [[F]], align 4
1366 // CHECK11-NEXT:    ret void
1367 //
1368 //
1369 // CHECK11-LABEL: define {{[^@]+}}@_ZN1SIfED2Ev
1370 // CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {
1371 // CHECK11-NEXT:  entry:
1372 // CHECK11-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4
1373 // CHECK11-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
1374 // CHECK11-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
1375 // CHECK11-NEXT:    ret void
1376 //
1377 //
1378 // CHECK11-LABEL: define {{[^@]+}}@_ZN1SIiEC1Ev
1379 // CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {
1380 // CHECK11-NEXT:  entry:
1381 // CHECK11-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4
1382 // CHECK11-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
1383 // CHECK11-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
1384 // CHECK11-NEXT:    call void @_ZN1SIiEC2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]])
1385 // CHECK11-NEXT:    ret void
1386 //
1387 //
1388 // CHECK11-LABEL: define {{[^@]+}}@_ZN1SIiEC1Ei
1389 // CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]], i32 noundef [[A:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {
1390 // CHECK11-NEXT:  entry:
1391 // CHECK11-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4
1392 // CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 4
1393 // CHECK11-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
1394 // CHECK11-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 4
1395 // CHECK11-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
1396 // CHECK11-NEXT:    [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 4
1397 // CHECK11-NEXT:    call void @_ZN1SIiEC2Ei(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]], i32 noundef [[TMP0]])
1398 // CHECK11-NEXT:    ret void
1399 //
1400 //
1401 // CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l49
1402 // CHECK11-SAME: () #[[ATTR3]] {
1403 // CHECK11-NEXT:  entry:
1404 // CHECK11-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l49.omp_outlined)
1405 // CHECK11-NEXT:    ret void
1406 //
1407 //
1408 // CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l49.omp_outlined
1409 // CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR3]] {
1410 // CHECK11-NEXT:  entry:
1411 // CHECK11-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
1412 // CHECK11-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
1413 // CHECK11-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
1414 // CHECK11-NEXT:    [[TMP:%.*]] = alloca i32, align 4
1415 // CHECK11-NEXT:    [[_TMP1:%.*]] = alloca ptr, align 4
1416 // CHECK11-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4
1417 // CHECK11-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4
1418 // CHECK11-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
1419 // CHECK11-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
1420 // CHECK11-NEXT:    [[T_VAR:%.*]] = alloca i32, align 4
1421 // CHECK11-NEXT:    [[VEC:%.*]] = alloca [2 x i32], align 4
1422 // CHECK11-NEXT:    [[S_ARR:%.*]] = alloca [2 x %struct.S.0], align 4
1423 // CHECK11-NEXT:    [[VAR:%.*]] = alloca [[STRUCT_S_0:%.*]], align 4
1424 // CHECK11-NEXT:    [[_TMP2:%.*]] = alloca ptr, align 4
1425 // CHECK11-NEXT:    [[I:%.*]] = alloca i32, align 4
1426 // CHECK11-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
1427 // CHECK11-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
1428 // CHECK11-NEXT:    store ptr undef, ptr [[_TMP1]], align 4
1429 // CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4
1430 // CHECK11-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 4
1431 // CHECK11-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4
1432 // CHECK11-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4
1433 // CHECK11-NEXT:    [[ARRAY_BEGIN:%.*]] = getelementptr inbounds [2 x %struct.S.0], ptr [[S_ARR]], i32 0, i32 0
1434 // CHECK11-NEXT:    [[ARRAYCTOR_END:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAY_BEGIN]], i32 2
1435 // CHECK11-NEXT:    br label [[ARRAYCTOR_LOOP:%.*]]
1436 // CHECK11:       arrayctor.loop:
1437 // CHECK11-NEXT:    [[ARRAYCTOR_CUR:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ]
1438 // CHECK11-NEXT:    call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]])
1439 // CHECK11-NEXT:    [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYCTOR_CUR]], i32 1
1440 // CHECK11-NEXT:    [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], [[ARRAYCTOR_END]]
1441 // CHECK11-NEXT:    br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]]
1442 // CHECK11:       arrayctor.cont:
1443 // CHECK11-NEXT:    call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]])
1444 // CHECK11-NEXT:    store ptr [[VAR]], ptr [[_TMP2]], align 4
1445 // CHECK11-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
1446 // CHECK11-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
1447 // CHECK11-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP1]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
1448 // CHECK11-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
1449 // CHECK11-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 1
1450 // CHECK11-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
1451 // CHECK11:       cond.true:
1452 // CHECK11-NEXT:    br label [[COND_END:%.*]]
1453 // CHECK11:       cond.false:
1454 // CHECK11-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
1455 // CHECK11-NEXT:    br label [[COND_END]]
1456 // CHECK11:       cond.end:
1457 // CHECK11-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]
1458 // CHECK11-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4
1459 // CHECK11-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4
1460 // CHECK11-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4
1461 // CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
1462 // CHECK11:       omp.inner.for.cond:
1463 // CHECK11-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
1464 // CHECK11-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
1465 // CHECK11-NEXT:    [[CMP3:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]
1466 // CHECK11-NEXT:    br i1 [[CMP3]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_COND_CLEANUP:%.*]]
1467 // CHECK11:       omp.inner.for.cond.cleanup:
1468 // CHECK11-NEXT:    br label [[OMP_INNER_FOR_END:%.*]]
1469 // CHECK11:       omp.inner.for.body:
1470 // CHECK11-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
1471 // CHECK11-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
1472 // CHECK11-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
1473 // CHECK11-NEXT:    store i32 [[ADD]], ptr [[I]], align 4
1474 // CHECK11-NEXT:    [[TMP8:%.*]] = load i32, ptr [[T_VAR]], align 4
1475 // CHECK11-NEXT:    [[TMP9:%.*]] = load i32, ptr [[I]], align 4
1476 // CHECK11-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [2 x i32], ptr [[VEC]], i32 0, i32 [[TMP9]]
1477 // CHECK11-NEXT:    store i32 [[TMP8]], ptr [[ARRAYIDX]], align 4
1478 // CHECK11-NEXT:    [[TMP10:%.*]] = load ptr, ptr [[_TMP2]], align 4
1479 // CHECK11-NEXT:    [[TMP11:%.*]] = load i32, ptr [[I]], align 4
1480 // CHECK11-NEXT:    [[ARRAYIDX4:%.*]] = getelementptr inbounds [2 x %struct.S.0], ptr [[S_ARR]], i32 0, i32 [[TMP11]]
1481 // CHECK11-NEXT:    call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[ARRAYIDX4]], ptr align 4 [[TMP10]], i32 4, i1 false)
1482 // CHECK11-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
1483 // CHECK11:       omp.body.continue:
1484 // CHECK11-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
1485 // CHECK11:       omp.inner.for.inc:
1486 // CHECK11-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
1487 // CHECK11-NEXT:    [[ADD5:%.*]] = add nsw i32 [[TMP12]], 1
1488 // CHECK11-NEXT:    store i32 [[ADD5]], ptr [[DOTOMP_IV]], align 4
1489 // CHECK11-NEXT:    br label [[OMP_INNER_FOR_COND]]
1490 // CHECK11:       omp.inner.for.end:
1491 // CHECK11-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
1492 // CHECK11:       omp.loop.exit:
1493 // CHECK11-NEXT:    [[TMP13:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
1494 // CHECK11-NEXT:    [[TMP14:%.*]] = load i32, ptr [[TMP13]], align 4
1495 // CHECK11-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP14]])
1496 // CHECK11-NEXT:    call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]]) #[[ATTR4]]
1497 // CHECK11-NEXT:    [[ARRAY_BEGIN6:%.*]] = getelementptr inbounds [2 x %struct.S.0], ptr [[S_ARR]], i32 0, i32 0
1498 // CHECK11-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAY_BEGIN6]], i32 2
1499 // CHECK11-NEXT:    br label [[ARRAYDESTROY_BODY:%.*]]
1500 // CHECK11:       arraydestroy.body:
1501 // CHECK11-NEXT:    [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP15]], [[OMP_LOOP_EXIT]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
1502 // CHECK11-NEXT:    [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i32 -1
1503 // CHECK11-NEXT:    call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR4]]
1504 // CHECK11-NEXT:    [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN6]]
1505 // CHECK11-NEXT:    br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE7:%.*]], label [[ARRAYDESTROY_BODY]]
1506 // CHECK11:       arraydestroy.done7:
1507 // CHECK11-NEXT:    ret void
1508 //
1509 //
1510 // CHECK11-LABEL: define {{[^@]+}}@_ZN1SIiED1Ev
1511 // CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {
1512 // CHECK11-NEXT:  entry:
1513 // CHECK11-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4
1514 // CHECK11-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
1515 // CHECK11-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
1516 // CHECK11-NEXT:    call void @_ZN1SIiED2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]]) #[[ATTR4]]
1517 // CHECK11-NEXT:    ret void
1518 //
1519 //
1520 // CHECK11-LABEL: define {{[^@]+}}@_ZN1SIiEC2Ev
1521 // CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {
1522 // CHECK11-NEXT:  entry:
1523 // CHECK11-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4
1524 // CHECK11-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
1525 // CHECK11-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
1526 // CHECK11-NEXT:    [[F:%.*]] = getelementptr inbounds [[STRUCT_S_0:%.*]], ptr [[THIS1]], i32 0, i32 0
1527 // CHECK11-NEXT:    store i32 0, ptr [[F]], align 4
1528 // CHECK11-NEXT:    ret void
1529 //
1530 //
1531 // CHECK11-LABEL: define {{[^@]+}}@_ZN1SIiEC2Ei
1532 // CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]], i32 noundef [[A:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {
1533 // CHECK11-NEXT:  entry:
1534 // CHECK11-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4
1535 // CHECK11-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 4
1536 // CHECK11-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
1537 // CHECK11-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 4
1538 // CHECK11-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
1539 // CHECK11-NEXT:    [[F:%.*]] = getelementptr inbounds [[STRUCT_S_0:%.*]], ptr [[THIS1]], i32 0, i32 0
1540 // CHECK11-NEXT:    [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 4
1541 // CHECK11-NEXT:    store i32 [[TMP0]], ptr [[F]], align 4
1542 // CHECK11-NEXT:    ret void
1543 //
1544 //
1545 // CHECK11-LABEL: define {{[^@]+}}@_ZN1SIiED2Ev
1546 // CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {
1547 // CHECK11-NEXT:  entry:
1548 // CHECK11-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4
1549 // CHECK11-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
1550 // CHECK11-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
1551 // CHECK11-NEXT:    ret void
1552 //
1553 //
1554 // CHECK11-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg
1555 // CHECK11-SAME: () #[[ATTR6:[0-9]+]] {
1556 // CHECK11-NEXT:  entry:
1557 // CHECK11-NEXT:    call void @__tgt_register_requires(i64 1)
1558 // CHECK11-NEXT:    ret void
1559 //