clang/test/OpenMP/target_has_device_addr_codegen_01.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]+" --prefix-filecheck-ir-name _
   2 // RUN: %clang_cc1 -DCK1 -verify -Wno-vla -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s
   3
   4 // RUN: %clang_cc1 -DCK1 -verify -Wno-vla -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
   5 // expected-no-diagnostics
   6
   7 struct S {
   8   int a = 0;
   9   int *ptr = &a;
  10   int &ref = a;
  11   int arr[4];
  12   S() {}
  13   void foo() {
  14 #pragma omp target has_device_addr(a, ref, ptr[0:4], arr[:a])
  15     ++a, ++*ptr, ++ref, ++arr[0];
  16   }
  17 };
  18
  19 int main() {
  20   float a = 0;
  21   float *ptr = &a;
  22   float &ref = a;
  23   float arr[4];
  24   float vla[(int)a];
  25   S s;
  26   s.foo();
  27 #pragma omp target has_device_addr(a, ref, ptr[0:4], arr[:(int)a], vla[0])
  28   ++a, ++*ptr, ++ref, ++arr[0], ++vla[0];
  29   return a;
  30 }
  31
  32
  33 // CHECK-LABEL: define {{[^@]+}}@main
  34 // CHECK-SAME: () #[[ATTR0:[0-9]+]] {
  35 // CHECK-NEXT:  entry:
  36 // CHECK-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4
  37 // CHECK-NEXT:    [[A:%.*]] = alloca float, align 4
  38 // CHECK-NEXT:    [[PTR:%.*]] = alloca ptr, align 8
  39 // CHECK-NEXT:    [[REF:%.*]] = alloca ptr, align 8
  40 // CHECK-NEXT:    [[ARR:%.*]] = alloca [4 x float], align 4
  41 // CHECK-NEXT:    [[SAVED_STACK:%.*]] = alloca ptr, align 8
  42 // CHECK-NEXT:    [[__VLA_EXPR0:%.*]] = alloca i64, align 8
  43 // CHECK-NEXT:    [[S:%.*]] = alloca [[STRUCT_S:%.*]], align 8
  44 // CHECK-NEXT:    [[TMP:%.*]] = alloca ptr, align 8
  45 // CHECK-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [6 x ptr], align 8
  46 // CHECK-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [6 x ptr], align 8
  47 // CHECK-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [6 x ptr], align 8
  48 // CHECK-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8
  49 // CHECK-NEXT:    store i32 0, ptr [[RETVAL]], align 4
  50 // CHECK-NEXT:    store float 0.000000e+00, ptr [[A]], align 4
  51 // CHECK-NEXT:    store ptr [[A]], ptr [[PTR]], align 8
  52 // CHECK-NEXT:    store ptr [[A]], ptr [[REF]], align 8
  53 // CHECK-NEXT:    [[TMP0:%.*]] = load float, ptr [[A]], align 4
  54 // CHECK-NEXT:    [[CONV:%.*]] = fptosi float [[TMP0]] to i32
  55 // CHECK-NEXT:    [[TMP1:%.*]] = zext i32 [[CONV]] to i64
  56 // CHECK-NEXT:    [[TMP2:%.*]] = call ptr @llvm.stacksave.p0()
  57 // CHECK-NEXT:    store ptr [[TMP2]], ptr [[SAVED_STACK]], align 8
  58 // CHECK-NEXT:    [[VLA:%.*]] = alloca float, i64 [[TMP1]], align 4
  59 // CHECK-NEXT:    store i64 [[TMP1]], ptr [[__VLA_EXPR0]], align 8
  60 // CHECK-NEXT:    call void @_ZN1SC1Ev(ptr noundef nonnull align 8 dereferenceable(40) [[S]])
  61 // CHECK-NEXT:    call void @_ZN1S3fooEv(ptr noundef nonnull align 8 dereferenceable(40) [[S]])
  62 // CHECK-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[REF]], align 8
  63 // CHECK-NEXT:    store ptr [[TMP3]], ptr [[TMP]], align 8
  64 // CHECK-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[PTR]], align 8
  65 // CHECK-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[TMP]], align 8
  66 // CHECK-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
  67 // CHECK-NEXT:    store ptr [[A]], ptr [[TMP6]], align 8
  68 // CHECK-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
  69 // CHECK-NEXT:    store ptr [[A]], ptr [[TMP7]], align 8
  70 // CHECK-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
  71 // CHECK-NEXT:    store ptr null, ptr [[TMP8]], align 8
  72 // CHECK-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
  73 // CHECK-NEXT:    store ptr [[TMP4]], ptr [[TMP9]], align 8
  74 // CHECK-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1
  75 // CHECK-NEXT:    store ptr [[TMP4]], ptr [[TMP10]], align 8
  76 // CHECK-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
  77 // CHECK-NEXT:    store ptr null, ptr [[TMP11]], align 8
  78 // CHECK-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2
  79 // CHECK-NEXT:    store ptr [[TMP5]], ptr [[TMP12]], align 8
  80 // CHECK-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2
  81 // CHECK-NEXT:    store ptr [[TMP5]], ptr [[TMP13]], align 8
  82 // CHECK-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2
  83 // CHECK-NEXT:    store ptr null, ptr [[TMP14]], align 8
  84 // CHECK-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3
  85 // CHECK-NEXT:    store ptr [[ARR]], ptr [[TMP15]], align 8
  86 // CHECK-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3
  87 // CHECK-NEXT:    store ptr [[ARR]], ptr [[TMP16]], align 8
  88 // CHECK-NEXT:    [[TMP17:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3
  89 // CHECK-NEXT:    store ptr null, ptr [[TMP17]], align 8
  90 // CHECK-NEXT:    [[TMP18:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 4
  91 // CHECK-NEXT:    store i64 [[TMP1]], ptr [[TMP18]], align 8
  92 // CHECK-NEXT:    [[TMP19:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 4
  93 // CHECK-NEXT:    store i64 [[TMP1]], ptr [[TMP19]], align 8
  94 // CHECK-NEXT:    [[TMP20:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 4
  95 // CHECK-NEXT:    store ptr null, ptr [[TMP20]], align 8
  96 // CHECK-NEXT:    [[TMP21:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 5
  97 // CHECK-NEXT:    store ptr [[VLA]], ptr [[TMP21]], align 8
  98 // CHECK-NEXT:    [[TMP22:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 5
  99 // CHECK-NEXT:    store ptr [[VLA]], ptr [[TMP22]], align 8
 100 // CHECK-NEXT:    [[TMP23:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 5
 101 // CHECK-NEXT:    store ptr null, ptr [[TMP23]], align 8
 102 // CHECK-NEXT:    [[TMP24:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
 103 // CHECK-NEXT:    [[TMP25:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
 104 // CHECK-NEXT:    [[TMP26:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0
 105 // CHECK-NEXT:    store i32 2, ptr [[TMP26]], align 4
 106 // CHECK-NEXT:    [[TMP27:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1
 107 // CHECK-NEXT:    store i32 6, ptr [[TMP27]], align 4
 108 // CHECK-NEXT:    [[TMP28:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2
 109 // CHECK-NEXT:    store ptr [[TMP24]], ptr [[TMP28]], align 8
 110 // CHECK-NEXT:    [[TMP29:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3
 111 // CHECK-NEXT:    store ptr [[TMP25]], ptr [[TMP29]], align 8
 112 // CHECK-NEXT:    [[TMP30:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4
 113 // CHECK-NEXT:    store ptr @.offload_sizes, ptr [[TMP30]], align 8
 114 // CHECK-NEXT:    [[TMP31:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5
 115 // CHECK-NEXT:    store ptr @.offload_maptypes, ptr [[TMP31]], align 8
 116 // CHECK-NEXT:    [[TMP32:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6
 117 // CHECK-NEXT:    store ptr null, ptr [[TMP32]], align 8
 118 // CHECK-NEXT:    [[TMP33:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7
 119 // CHECK-NEXT:    store ptr null, ptr [[TMP33]], align 8
 120 // CHECK-NEXT:    [[TMP34:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8
 121 // CHECK-NEXT:    store i64 0, ptr [[TMP34]], align 8
 122 // CHECK-NEXT:    [[TMP35:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9
 123 // CHECK-NEXT:    store i64 0, ptr [[TMP35]], align 8
 124 // CHECK-NEXT:    [[TMP36:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10
 125 // CHECK-NEXT:    store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP36]], align 4
 126 // CHECK-NEXT:    [[TMP37:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11
 127 // CHECK-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP37]], align 4
 128 // CHECK-NEXT:    [[TMP38:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12
 129 // CHECK-NEXT:    store i32 0, ptr [[TMP38]], align 4
 130 // CHECK-NEXT:    [[TMP39:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l27.region_id, ptr [[KERNEL_ARGS]])
 131 // CHECK-NEXT:    [[TMP40:%.*]] = icmp ne i32 [[TMP39]], 0
 132 // CHECK-NEXT:    br i1 [[TMP40]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
 133 // CHECK:       omp_offload.failed:
 134 // CHECK-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l27(ptr [[A]], ptr [[TMP4]], ptr [[TMP5]], ptr [[ARR]], i64 [[TMP1]], ptr [[VLA]]) #[[ATTR5:[0-9]+]]
 135 // CHECK-NEXT:    br label [[OMP_OFFLOAD_CONT]]
 136 // CHECK:       omp_offload.cont:
 137 // CHECK-NEXT:    [[TMP41:%.*]] = load float, ptr [[A]], align 4
 138 // CHECK-NEXT:    [[CONV1:%.*]] = fptosi float [[TMP41]] to i32
 139 // CHECK-NEXT:    store i32 [[CONV1]], ptr [[RETVAL]], align 4
 140 // CHECK-NEXT:    [[TMP42:%.*]] = load ptr, ptr [[SAVED_STACK]], align 8
 141 // CHECK-NEXT:    call void @llvm.stackrestore.p0(ptr [[TMP42]])
 142 // CHECK-NEXT:    [[TMP43:%.*]] = load i32, ptr [[RETVAL]], align 4
 143 // CHECK-NEXT:    ret i32 [[TMP43]]
 144 //
 145 //
 146 // CHECK-LABEL: define {{[^@]+}}@_ZN1SC1Ev
 147 // CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(40) [[THIS:%.*]]) unnamed_addr #[[ATTR2:[0-9]+]] comdat {
 148 // CHECK-NEXT:  entry:
 149 // CHECK-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 150 // CHECK-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 151 // CHECK-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 152 // CHECK-NEXT:    call void @_ZN1SC2Ev(ptr noundef nonnull align 8 dereferenceable(40) [[THIS1]])
 153 // CHECK-NEXT:    ret void
 154 //
 155 //
 156 // CHECK-LABEL: define {{[^@]+}}@_ZN1S3fooEv
 157 // CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(40) [[THIS:%.*]]) #[[ATTR3:[0-9]+]] comdat {
 158 // CHECK-NEXT:  entry:
 159 // CHECK-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 160 // CHECK-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [5 x ptr], align 8
 161 // CHECK-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [5 x ptr], align 8
 162 // CHECK-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [5 x ptr], align 8
 163 // CHECK-NEXT:    [[DOTOFFLOAD_SIZES:%.*]] = alloca [5 x i64], align 8
 164 // CHECK-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8
 165 // CHECK-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 166 // CHECK-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 167 // CHECK-NEXT:    [[A:%.*]] = getelementptr inbounds [[STRUCT_S:%.*]], ptr [[THIS1]], i32 0, i32 0
 168 // CHECK-NEXT:    [[REF:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 2
 169 // CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[REF]], align 8
 170 // CHECK-NEXT:    [[PTR:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
 171 // CHECK-NEXT:    [[ARR:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 3
 172 // CHECK-NEXT:    [[TMP1:%.*]] = getelementptr [4 x i32], ptr [[ARR]], i32 1
 173 // CHECK-NEXT:    [[TMP2:%.*]] = ptrtoint ptr [[TMP1]] to i64
 174 // CHECK-NEXT:    [[TMP3:%.*]] = ptrtoint ptr [[A]] to i64
 175 // CHECK-NEXT:    [[TMP4:%.*]] = sub i64 [[TMP2]], [[TMP3]]
 176 // CHECK-NEXT:    [[TMP5:%.*]] = sdiv exact i64 [[TMP4]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
 177 // CHECK-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES]], ptr align 8 @.offload_sizes.1, i64 40, i1 false)
 178 // CHECK-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
 179 // CHECK-NEXT:    store ptr [[THIS1]], ptr [[TMP6]], align 8
 180 // CHECK-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
 181 // CHECK-NEXT:    store ptr [[A]], ptr [[TMP7]], align 8
 182 // CHECK-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [5 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0
 183 // CHECK-NEXT:    store i64 [[TMP5]], ptr [[TMP8]], align 8
 184 // CHECK-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
 185 // CHECK-NEXT:    store ptr null, ptr [[TMP9]], align 8
 186 // CHECK-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
 187 // CHECK-NEXT:    store ptr [[THIS1]], ptr [[TMP10]], align 8
 188 // CHECK-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1
 189 // CHECK-NEXT:    store ptr [[A]], ptr [[TMP11]], align 8
 190 // CHECK-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
 191 // CHECK-NEXT:    store ptr null, ptr [[TMP12]], align 8
 192 // CHECK-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2
 193 // CHECK-NEXT:    store ptr [[THIS1]], ptr [[TMP13]], align 8
 194 // CHECK-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2
 195 // CHECK-NEXT:    store ptr [[TMP0]], ptr [[TMP14]], align 8
 196 // CHECK-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2
 197 // CHECK-NEXT:    store ptr null, ptr [[TMP15]], align 8
 198 // CHECK-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3
 199 // CHECK-NEXT:    store ptr [[THIS1]], ptr [[TMP16]], align 8
 200 // CHECK-NEXT:    [[TMP17:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3
 201 // CHECK-NEXT:    store ptr [[PTR]], ptr [[TMP17]], align 8
 202 // CHECK-NEXT:    [[TMP18:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3
 203 // CHECK-NEXT:    store ptr null, ptr [[TMP18]], align 8
 204 // CHECK-NEXT:    [[TMP19:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 4
 205 // CHECK-NEXT:    store ptr [[THIS1]], ptr [[TMP19]], align 8
 206 // CHECK-NEXT:    [[TMP20:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 4
 207 // CHECK-NEXT:    store ptr [[ARR]], ptr [[TMP20]], align 8
 208 // CHECK-NEXT:    [[TMP21:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 4
 209 // CHECK-NEXT:    store ptr null, ptr [[TMP21]], align 8
 210 // CHECK-NEXT:    [[TMP22:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
 211 // CHECK-NEXT:    [[TMP23:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
 212 // CHECK-NEXT:    [[TMP24:%.*]] = getelementptr inbounds [5 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0
 213 // CHECK-NEXT:    [[TMP25:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0
 214 // CHECK-NEXT:    store i32 2, ptr [[TMP25]], align 4
 215 // CHECK-NEXT:    [[TMP26:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1
 216 // CHECK-NEXT:    store i32 5, ptr [[TMP26]], align 4
 217 // CHECK-NEXT:    [[TMP27:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2
 218 // CHECK-NEXT:    store ptr [[TMP22]], ptr [[TMP27]], align 8
 219 // CHECK-NEXT:    [[TMP28:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3
 220 // CHECK-NEXT:    store ptr [[TMP23]], ptr [[TMP28]], align 8
 221 // CHECK-NEXT:    [[TMP29:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4
 222 // CHECK-NEXT:    store ptr [[TMP24]], ptr [[TMP29]], align 8
 223 // CHECK-NEXT:    [[TMP30:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5
 224 // CHECK-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP30]], align 8
 225 // CHECK-NEXT:    [[TMP31:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6
 226 // CHECK-NEXT:    store ptr null, ptr [[TMP31]], align 8
 227 // CHECK-NEXT:    [[TMP32:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7
 228 // CHECK-NEXT:    store ptr null, ptr [[TMP32]], align 8
 229 // CHECK-NEXT:    [[TMP33:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8
 230 // CHECK-NEXT:    store i64 0, ptr [[TMP33]], align 8
 231 // CHECK-NEXT:    [[TMP34:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9
 232 // CHECK-NEXT:    store i64 0, ptr [[TMP34]], align 8
 233 // CHECK-NEXT:    [[TMP35:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10
 234 // CHECK-NEXT:    store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP35]], align 4
 235 // CHECK-NEXT:    [[TMP36:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11
 236 // CHECK-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP36]], align 4
 237 // CHECK-NEXT:    [[TMP37:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12
 238 // CHECK-NEXT:    store i32 0, ptr [[TMP37]], align 4
 239 // CHECK-NEXT:    [[TMP38:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l14.region_id, ptr [[KERNEL_ARGS]])
 240 // CHECK-NEXT:    [[TMP39:%.*]] = icmp ne i32 [[TMP38]], 0
 241 // CHECK-NEXT:    br i1 [[TMP39]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
 242 // CHECK:       omp_offload.failed:
 243 // CHECK-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l14(ptr [[THIS1]]) #[[ATTR5]]
 244 // CHECK-NEXT:    br label [[OMP_OFFLOAD_CONT]]
 245 // CHECK:       omp_offload.cont:
 246 // CHECK-NEXT:    ret void
 247 //
 248 //
 249 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l27
 250 // CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef [[PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[REF:%.*]], ptr noundef nonnull align 4 dereferenceable(16) [[ARR:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[VLA1:%.*]]) #[[ATTR4:[0-9]+]] {
 251 // CHECK-NEXT:  entry:
 252 // CHECK-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8
 253 // CHECK-NEXT:    [[PTR_ADDR:%.*]] = alloca ptr, align 8
 254 // CHECK-NEXT:    [[REF_ADDR:%.*]] = alloca ptr, align 8
 255 // CHECK-NEXT:    [[ARR_ADDR:%.*]] = alloca ptr, align 8
 256 // CHECK-NEXT:    [[VLA_ADDR:%.*]] = alloca i64, align 8
 257 // CHECK-NEXT:    [[VLA_ADDR2:%.*]] = alloca ptr, align 8
 258 // CHECK-NEXT:    [[TMP:%.*]] = alloca ptr, align 8
 259 // CHECK-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8
 260 // CHECK-NEXT:    store ptr [[PTR]], ptr [[PTR_ADDR]], align 8
 261 // CHECK-NEXT:    store ptr [[REF]], ptr [[REF_ADDR]], align 8
 262 // CHECK-NEXT:    store ptr [[ARR]], ptr [[ARR_ADDR]], align 8
 263 // CHECK-NEXT:    store i64 [[VLA]], ptr [[VLA_ADDR]], align 8
 264 // CHECK-NEXT:    store ptr [[VLA1]], ptr [[VLA_ADDR2]], align 8
 265 // CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
 266 // CHECK-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[REF_ADDR]], align 8
 267 // CHECK-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[ARR_ADDR]], align 8
 268 // CHECK-NEXT:    [[TMP3:%.*]] = load i64, ptr [[VLA_ADDR]], align 8
 269 // CHECK-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[VLA_ADDR2]], align 8
 270 // CHECK-NEXT:    store ptr [[TMP1]], ptr [[TMP]], align 8
 271 // CHECK-NEXT:    [[TMP5:%.*]] = load float, ptr [[TMP0]], align 4
 272 // CHECK-NEXT:    [[INC:%.*]] = fadd float [[TMP5]], 1.000000e+00
 273 // CHECK-NEXT:    store float [[INC]], ptr [[TMP0]], align 4
 274 // CHECK-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8
 275 // CHECK-NEXT:    [[TMP7:%.*]] = load float, ptr [[TMP6]], align 4
 276 // CHECK-NEXT:    [[INC3:%.*]] = fadd float [[TMP7]], 1.000000e+00
 277 // CHECK-NEXT:    store float [[INC3]], ptr [[TMP6]], align 4
 278 // CHECK-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[TMP]], align 8
 279 // CHECK-NEXT:    [[TMP9:%.*]] = load float, ptr [[TMP8]], align 4
 280 // CHECK-NEXT:    [[INC4:%.*]] = fadd float [[TMP9]], 1.000000e+00
 281 // CHECK-NEXT:    store float [[INC4]], ptr [[TMP8]], align 4
 282 // CHECK-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [4 x float], ptr [[TMP2]], i64 0, i64 0
 283 // CHECK-NEXT:    [[TMP10:%.*]] = load float, ptr [[ARRAYIDX]], align 4
 284 // CHECK-NEXT:    [[INC5:%.*]] = fadd float [[TMP10]], 1.000000e+00
 285 // CHECK-NEXT:    store float [[INC5]], ptr [[ARRAYIDX]], align 4
 286 // CHECK-NEXT:    [[ARRAYIDX6:%.*]] = getelementptr inbounds float, ptr [[TMP4]], i64 0
 287 // CHECK-NEXT:    [[TMP11:%.*]] = load float, ptr [[ARRAYIDX6]], align 4
 288 // CHECK-NEXT:    [[INC7:%.*]] = fadd float [[TMP11]], 1.000000e+00
 289 // CHECK-NEXT:    store float [[INC7]], ptr [[ARRAYIDX6]], align 4
 290 // CHECK-NEXT:    ret void
 291 //
 292 //
 293 // CHECK-LABEL: define {{[^@]+}}@_ZN1SC2Ev
 294 // CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(40) [[THIS:%.*]]) unnamed_addr #[[ATTR2]] comdat {
 295 // CHECK-NEXT:  entry:
 296 // CHECK-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 297 // CHECK-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 298 // CHECK-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 299 // CHECK-NEXT:    [[A:%.*]] = getelementptr inbounds [[STRUCT_S:%.*]], ptr [[THIS1]], i32 0, i32 0
 300 // CHECK-NEXT:    store i32 0, ptr [[A]], align 8
 301 // CHECK-NEXT:    [[PTR:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
 302 // CHECK-NEXT:    [[A2:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 0
 303 // CHECK-NEXT:    store ptr [[A2]], ptr [[PTR]], align 8
 304 // CHECK-NEXT:    [[REF:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 2
 305 // CHECK-NEXT:    [[A3:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 0
 306 // CHECK-NEXT:    store ptr [[A3]], ptr [[REF]], align 8
 307 // CHECK-NEXT:    ret void
 308 //
 309 //
 310 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l14
 311 // CHECK-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR4]] {
 312 // CHECK-NEXT:  entry:
 313 // CHECK-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 314 // CHECK-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 315 // CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 316 // CHECK-NEXT:    [[A:%.*]] = getelementptr inbounds [[STRUCT_S:%.*]], ptr [[TMP0]], i32 0, i32 0
 317 // CHECK-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A]], align 8
 318 // CHECK-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP1]], 1
 319 // CHECK-NEXT:    store i32 [[INC]], ptr [[A]], align 8
 320 // CHECK-NEXT:    [[PTR:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[TMP0]], i32 0, i32 1
 321 // CHECK-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[PTR]], align 8
 322 // CHECK-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4
 323 // CHECK-NEXT:    [[INC1:%.*]] = add nsw i32 [[TMP3]], 1
 324 // CHECK-NEXT:    store i32 [[INC1]], ptr [[TMP2]], align 4
 325 // CHECK-NEXT:    [[REF:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[TMP0]], i32 0, i32 2
 326 // CHECK-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[REF]], align 8
 327 // CHECK-NEXT:    [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 4
 328 // CHECK-NEXT:    [[INC2:%.*]] = add nsw i32 [[TMP5]], 1
 329 // CHECK-NEXT:    store i32 [[INC2]], ptr [[TMP4]], align 4
 330 // CHECK-NEXT:    [[ARR:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[TMP0]], i32 0, i32 3
 331 // CHECK-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [4 x i32], ptr [[ARR]], i64 0, i64 0
 332 // CHECK-NEXT:    [[TMP6:%.*]] = load i32, ptr [[ARRAYIDX]], align 8
 333 // CHECK-NEXT:    [[INC3:%.*]] = add nsw i32 [[TMP6]], 1
 334 // CHECK-NEXT:    store i32 [[INC3]], ptr [[ARRAYIDX]], align 8
 335 // CHECK-NEXT:    ret void
 336 //
 337 //
 338 // CHECK-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg
 339 // CHECK-SAME: () #[[ATTR7:[0-9]+]] {
 340 // CHECK-NEXT:  entry:
 341 // CHECK-NEXT:    call void @__tgt_register_requires(i64 1)
 342 // CHECK-NEXT:    ret void
 343 //
 344 //
 345 // SIMD-ONLY0-LABEL: define {{[^@]+}}@main
 346 // SIMD-ONLY0-SAME: () #[[ATTR0:[0-9]+]] {
 347 // SIMD-ONLY0-NEXT:  entry:
 348 // SIMD-ONLY0-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4
 349 // SIMD-ONLY0-NEXT:    [[A:%.*]] = alloca float, align 4
 350 // SIMD-ONLY0-NEXT:    [[PTR:%.*]] = alloca ptr, align 8
 351 // SIMD-ONLY0-NEXT:    [[REF:%.*]] = alloca ptr, align 8
 352 // SIMD-ONLY0-NEXT:    [[ARR:%.*]] = alloca [4 x float], align 4
 353 // SIMD-ONLY0-NEXT:    [[SAVED_STACK:%.*]] = alloca ptr, align 8
 354 // SIMD-ONLY0-NEXT:    [[__VLA_EXPR0:%.*]] = alloca i64, align 8
 355 // SIMD-ONLY0-NEXT:    [[S:%.*]] = alloca [[STRUCT_S:%.*]], align 8
 356 // SIMD-ONLY0-NEXT:    [[TMP:%.*]] = alloca ptr, align 8
 357 // SIMD-ONLY0-NEXT:    store i32 0, ptr [[RETVAL]], align 4
 358 // SIMD-ONLY0-NEXT:    store float 0.000000e+00, ptr [[A]], align 4
 359 // SIMD-ONLY0-NEXT:    store ptr [[A]], ptr [[PTR]], align 8
 360 // SIMD-ONLY0-NEXT:    store ptr [[A]], ptr [[REF]], align 8
 361 // SIMD-ONLY0-NEXT:    [[TMP0:%.*]] = load float, ptr [[A]], align 4
 362 // SIMD-ONLY0-NEXT:    [[CONV:%.*]] = fptosi float [[TMP0]] to i32
 363 // SIMD-ONLY0-NEXT:    [[TMP1:%.*]] = zext i32 [[CONV]] to i64
 364 // SIMD-ONLY0-NEXT:    [[TMP2:%.*]] = call ptr @llvm.stacksave.p0()
 365 // SIMD-ONLY0-NEXT:    store ptr [[TMP2]], ptr [[SAVED_STACK]], align 8
 366 // SIMD-ONLY0-NEXT:    [[VLA:%.*]] = alloca float, i64 [[TMP1]], align 4
 367 // SIMD-ONLY0-NEXT:    store i64 [[TMP1]], ptr [[__VLA_EXPR0]], align 8
 368 // SIMD-ONLY0-NEXT:    call void @_ZN1SC1Ev(ptr noundef nonnull align 8 dereferenceable(40) [[S]])
 369 // SIMD-ONLY0-NEXT:    call void @_ZN1S3fooEv(ptr noundef nonnull align 8 dereferenceable(40) [[S]])
 370 // SIMD-ONLY0-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[REF]], align 8
 371 // SIMD-ONLY0-NEXT:    store ptr [[TMP3]], ptr [[TMP]], align 8
 372 // SIMD-ONLY0-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[REF]], align 8
 373 // SIMD-ONLY0-NEXT:    [[TMP5:%.*]] = load float, ptr [[A]], align 4
 374 // SIMD-ONLY0-NEXT:    [[INC:%.*]] = fadd float [[TMP5]], 1.000000e+00
 375 // SIMD-ONLY0-NEXT:    store float [[INC]], ptr [[A]], align 4
 376 // SIMD-ONLY0-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[PTR]], align 8
 377 // SIMD-ONLY0-NEXT:    [[TMP7:%.*]] = load float, ptr [[TMP6]], align 4
 378 // SIMD-ONLY0-NEXT:    [[INC1:%.*]] = fadd float [[TMP7]], 1.000000e+00
 379 // SIMD-ONLY0-NEXT:    store float [[INC1]], ptr [[TMP6]], align 4
 380 // SIMD-ONLY0-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[TMP]], align 8
 381 // SIMD-ONLY0-NEXT:    [[TMP9:%.*]] = load float, ptr [[TMP8]], align 4
 382 // SIMD-ONLY0-NEXT:    [[INC2:%.*]] = fadd float [[TMP9]], 1.000000e+00
 383 // SIMD-ONLY0-NEXT:    store float [[INC2]], ptr [[TMP8]], align 4
 384 // SIMD-ONLY0-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [4 x float], ptr [[ARR]], i64 0, i64 0
 385 // SIMD-ONLY0-NEXT:    [[TMP10:%.*]] = load float, ptr [[ARRAYIDX]], align 4
 386 // SIMD-ONLY0-NEXT:    [[INC3:%.*]] = fadd float [[TMP10]], 1.000000e+00
 387 // SIMD-ONLY0-NEXT:    store float [[INC3]], ptr [[ARRAYIDX]], align 4
 388 // SIMD-ONLY0-NEXT:    [[ARRAYIDX4:%.*]] = getelementptr inbounds float, ptr [[VLA]], i64 0
 389 // SIMD-ONLY0-NEXT:    [[TMP11:%.*]] = load float, ptr [[ARRAYIDX4]], align 4
 390 // SIMD-ONLY0-NEXT:    [[INC5:%.*]] = fadd float [[TMP11]], 1.000000e+00
 391 // SIMD-ONLY0-NEXT:    store float [[INC5]], ptr [[ARRAYIDX4]], align 4
 392 // SIMD-ONLY0-NEXT:    [[TMP12:%.*]] = load float, ptr [[A]], align 4
 393 // SIMD-ONLY0-NEXT:    [[CONV6:%.*]] = fptosi float [[TMP12]] to i32
 394 // SIMD-ONLY0-NEXT:    store i32 [[CONV6]], ptr [[RETVAL]], align 4
 395 // SIMD-ONLY0-NEXT:    [[TMP13:%.*]] = load ptr, ptr [[SAVED_STACK]], align 8
 396 // SIMD-ONLY0-NEXT:    call void @llvm.stackrestore.p0(ptr [[TMP13]])
 397 // SIMD-ONLY0-NEXT:    [[TMP14:%.*]] = load i32, ptr [[RETVAL]], align 4
 398 // SIMD-ONLY0-NEXT:    ret i32 [[TMP14]]
 399 //
 400 //
 401 // SIMD-ONLY0-LABEL: define {{[^@]+}}@_ZN1SC1Ev
 402 // SIMD-ONLY0-SAME: (ptr noundef nonnull align 8 dereferenceable(40) [[THIS:%.*]]) unnamed_addr #[[ATTR2:[0-9]+]] comdat {
 403 // SIMD-ONLY0-NEXT:  entry:
 404 // SIMD-ONLY0-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 405 // SIMD-ONLY0-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 406 // SIMD-ONLY0-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 407 // SIMD-ONLY0-NEXT:    call void @_ZN1SC2Ev(ptr noundef nonnull align 8 dereferenceable(40) [[THIS1]])
 408 // SIMD-ONLY0-NEXT:    ret void
 409 //
 410 //
 411 // SIMD-ONLY0-LABEL: define {{[^@]+}}@_ZN1S3fooEv
 412 // SIMD-ONLY0-SAME: (ptr noundef nonnull align 8 dereferenceable(40) [[THIS:%.*]]) #[[ATTR3:[0-9]+]] comdat {
 413 // SIMD-ONLY0-NEXT:  entry:
 414 // SIMD-ONLY0-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 415 // SIMD-ONLY0-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 416 // SIMD-ONLY0-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 417 // SIMD-ONLY0-NEXT:    [[A:%.*]] = getelementptr inbounds [[STRUCT_S:%.*]], ptr [[THIS1]], i32 0, i32 0
 418 // SIMD-ONLY0-NEXT:    [[TMP0:%.*]] = load i32, ptr [[A]], align 8
 419 // SIMD-ONLY0-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP0]], 1
 420 // SIMD-ONLY0-NEXT:    store i32 [[INC]], ptr [[A]], align 8
 421 // SIMD-ONLY0-NEXT:    [[PTR:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
 422 // SIMD-ONLY0-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[PTR]], align 8
 423 // SIMD-ONLY0-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4
 424 // SIMD-ONLY0-NEXT:    [[INC2:%.*]] = add nsw i32 [[TMP2]], 1
 425 // SIMD-ONLY0-NEXT:    store i32 [[INC2]], ptr [[TMP1]], align 4
 426 // SIMD-ONLY0-NEXT:    [[REF:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 2
 427 // SIMD-ONLY0-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[REF]], align 8
 428 // SIMD-ONLY0-NEXT:    [[TMP4:%.*]] = load i32, ptr [[TMP3]], align 4
 429 // SIMD-ONLY0-NEXT:    [[INC3:%.*]] = add nsw i32 [[TMP4]], 1
 430 // SIMD-ONLY0-NEXT:    store i32 [[INC3]], ptr [[TMP3]], align 4
 431 // SIMD-ONLY0-NEXT:    [[ARR:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 3
 432 // SIMD-ONLY0-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [4 x i32], ptr [[ARR]], i64 0, i64 0
 433 // SIMD-ONLY0-NEXT:    [[TMP5:%.*]] = load i32, ptr [[ARRAYIDX]], align 8
 434 // SIMD-ONLY0-NEXT:    [[INC4:%.*]] = add nsw i32 [[TMP5]], 1
 435 // SIMD-ONLY0-NEXT:    store i32 [[INC4]], ptr [[ARRAYIDX]], align 8
 436 // SIMD-ONLY0-NEXT:    ret void
 437 //
 438 //
 439 // SIMD-ONLY0-LABEL: define {{[^@]+}}@_ZN1SC2Ev
 440 // SIMD-ONLY0-SAME: (ptr noundef nonnull align 8 dereferenceable(40) [[THIS:%.*]]) unnamed_addr #[[ATTR2]] comdat {
 441 // SIMD-ONLY0-NEXT:  entry:
 442 // SIMD-ONLY0-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 443 // SIMD-ONLY0-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 444 // SIMD-ONLY0-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 445 // SIMD-ONLY0-NEXT:    [[A:%.*]] = getelementptr inbounds [[STRUCT_S:%.*]], ptr [[THIS1]], i32 0, i32 0
 446 // SIMD-ONLY0-NEXT:    store i32 0, ptr [[A]], align 8
 447 // SIMD-ONLY0-NEXT:    [[PTR:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
 448 // SIMD-ONLY0-NEXT:    [[A2:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 0
 449 // SIMD-ONLY0-NEXT:    store ptr [[A2]], ptr [[PTR]], align 8
 450 // SIMD-ONLY0-NEXT:    [[REF:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 2
 451 // SIMD-ONLY0-NEXT:    [[A3:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 0
 452 // SIMD-ONLY0-NEXT:    store ptr [[A3]], ptr [[REF]], align 8
 453 // SIMD-ONLY0-NEXT:    ret void
 454 //