clang/test/OpenMP/reduction_implicit_map.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 -verify -fopenmp -fopenmp-cuda-mode -x c++ \
   3 // RUN:  -triple powerpc64le-unknown-unknown -DCUDA \
   4 // RUN:  -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o \
   5 // RUN:  %t-ppc-host.bc
   6
   7 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ \
   8 // RUN:  -triple nvptx64-unknown-unknown -DCUA \
   9 // RUN:  -fopenmp-targets=nvptx64-nvidia-cuda -DCUDA -emit-llvm %s \
  10 // RUN:  -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc \
  11 // RUN:  -o - | FileCheck %s --check-prefix CHECK
  12
  13 // RUN: %clang_cc1 -verify -fopenmp -x c++ \
  14 // RUN:   -triple powerpc64le-unknown-unknown -DDIAG\
  15 // RUN:   -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm \
  16 // RUN:   %s -o - | FileCheck  %s \
  17 // RUN:   --check-prefix=CHECK1
  18
  19 // RUN: %clang_cc1 -verify -fopenmp -x c++ \
  20 // RUN:   -triple i386-unknown-unknown \
  21 // RUN:   -fopenmp-targets=i386-pc-linux-gnu -emit-llvm \
  22 // RUN:   %s -o - | FileCheck  %s \
  23 // RUN:   --check-prefix=CHECK2
  24
  25
  26 #if defined(CUDA)
  27 // expected-no-diagnostics
  28
  29 int foo(int n) {
  30   double *e;
  31   //no error and no implicit map generated for e[:1]
  32   #pragma omp target parallel reduction(+: e[:1])
  33     *e=10;
  34   ;
  35   return 0;
  36 }
  37 // CHECK-NOT @.offload_maptypes
  38 #elif defined(DIAG)
  39 class S2 {
  40   mutable int a;
  41 public:
  42   S2():a(0) { }
  43   S2(S2 &s2):a(s2.a) { }
  44   S2 &operator +(S2 &s);
  45 };
  46 int bar() {
  47  S2 o[5];
  48   //warnig "copyable and not guaranteed to be mapped correctly" and
  49   //implicit map generated.
  50 #pragma omp target parallel reduction(+:o[0]) //expected-warning {{type 'S2' is not trivially copyable and not guaranteed to be mapped correctly}}
  51   for (int i = 0; i < 10; i++);
  52   double b[10][10][10];
  53   //no error no implicit map generated, the map for b is generated but not
  54   //for b[0:2][2:4][1].
  55 #pragma omp target parallel for reduction(task, +: b[0:2][2:4][1])
  56   for (long long i = 0; i < 10; ++i);
  57   return 0;
  58 }
  59 // map for variable o
  60 // map for b:
  61 #else
  62 // expected-no-diagnostics
  63
  64 // generate implicit map for array elements or array sections in reduction
  65 // clause. In following case: the implicit map is generate for output[0]
  66 // with map size 4 and output[:3] with map size 12.
  67 void sum(int* input, int size, int* output)
  68 {
  69 #pragma omp target teams distribute parallel for reduction(+: output[0]) \
  70                                                  map(to: input [0:size])
  71   for (int i = 0; i < size; i++)
  72     output[0] += input[i];
  73 #pragma omp target teams distribute parallel for reduction(+: output[:3])  \
  74                                                  map(to: input [0:size])
  75   for (int i = 0; i < size; i++)
  76     output[0] += input[i];
  77   int a[10];
  78 #pragma omp target parallel reduction(+: a[:2])
  79   for (int i = 0; i < size; i++)
  80     ;
  81 #pragma omp target parallel reduction(+: a[3])
  82   for (int i = 0; i < size; i++)
  83     ;
  84 }
  85 #endif
  86 int main()
  87 {
  88 #if defined(CUDA)
  89   int a = foo(10);
  90 #elif defined(DIAG)
  91   int a = bar();
  92 #else
  93   const int size = 100;
  94   int *array = new int[size];
  95   int result = 0;
  96   sum(array, size, &result);
  97 #endif
  98   return 0;
  99 }
 100 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooi_l32
 101 // CHECK-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[E:%.*]]) #[[ATTR0:[0-9]+]] {
 102 // CHECK-NEXT:  entry:
 103 // CHECK-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8
 104 // CHECK-NEXT:    [[E_ADDR:%.*]] = alloca ptr, align 8
 105 // CHECK-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 8
 106 // CHECK-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8
 107 // CHECK-NEXT:    store ptr [[E]], ptr [[E_ADDR]], align 8
 108 // CHECK-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooi_l32_kernel_environment, ptr [[DYN_PTR]])
 109 // CHECK-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
 110 // CHECK-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
 111 // CHECK:       user_code.entry:
 112 // CHECK-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])
 113 // CHECK-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[E_ADDR]], align 8
 114 // CHECK-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 0
 115 // CHECK-NEXT:    store ptr [[TMP2]], ptr [[TMP3]], align 8
 116 // CHECK-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooi_l32_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i64 1)
 117 // CHECK-NEXT:    call void @__kmpc_target_deinit()
 118 // CHECK-NEXT:    ret void
 119 // CHECK:       worker.exit:
 120 // CHECK-NEXT:    ret void
 121 //
 122 //
 123 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooi_l32_omp_outlined
 124 // CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[E:%.*]]) #[[ATTR1:[0-9]+]] {
 125 // CHECK-NEXT:  entry:
 126 // CHECK-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
 127 // CHECK-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
 128 // CHECK-NEXT:    [[E_ADDR:%.*]] = alloca ptr, align 8
 129 // CHECK-NEXT:    [[E2:%.*]] = alloca double, align 8
 130 // CHECK-NEXT:    [[TMP:%.*]] = alloca ptr, align 8
 131 // CHECK-NEXT:    [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 8
 132 // CHECK-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
 133 // CHECK-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
 134 // CHECK-NEXT:    store ptr [[E]], ptr [[E_ADDR]], align 8
 135 // CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[E_ADDR]], align 8
 136 // CHECK-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds nuw double, ptr [[TMP0]], i64 0
 137 // CHECK-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[E_ADDR]], align 8
 138 // CHECK-NEXT:    [[ARRAYIDX1:%.*]] = getelementptr inbounds nuw double, ptr [[TMP1]], i64 0
 139 // CHECK-NEXT:    store double 0.000000e+00, ptr [[E2]], align 8
 140 // CHECK-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[E_ADDR]], align 8
 141 // CHECK-NEXT:    [[TMP3:%.*]] = ptrtoint ptr [[TMP2]] to i64
 142 // CHECK-NEXT:    [[TMP4:%.*]] = ptrtoint ptr [[ARRAYIDX]] to i64
 143 // CHECK-NEXT:    [[TMP5:%.*]] = sub i64 [[TMP3]], [[TMP4]]
 144 // CHECK-NEXT:    [[TMP6:%.*]] = sdiv exact i64 [[TMP5]], ptrtoint (ptr getelementptr (double, ptr null, i32 1) to i64)
 145 // CHECK-NEXT:    [[TMP7:%.*]] = getelementptr double, ptr [[E2]], i64 [[TMP6]]
 146 // CHECK-NEXT:    store ptr [[TMP7]], ptr [[TMP]], align 8
 147 // CHECK-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[TMP]], align 8
 148 // CHECK-NEXT:    store double 1.000000e+01, ptr [[TMP8]], align 8
 149 // CHECK-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0
 150 // CHECK-NEXT:    store ptr [[E2]], ptr [[TMP9]], align 8
 151 // CHECK-NEXT:    [[TMP10:%.*]] = call i32 @__kmpc_nvptx_parallel_reduce_nowait_v2(ptr @[[GLOB1]], i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @_omp_reduction_shuffle_and_reduce_func, ptr @_omp_reduction_inter_warp_copy_func)
 152 // CHECK-NEXT:    [[TMP11:%.*]] = icmp eq i32 [[TMP10]], 1
 153 // CHECK-NEXT:    br i1 [[TMP11]], label [[DOTOMP_REDUCTION_THEN:%.*]], label [[DOTOMP_REDUCTION_DONE:%.*]]
 154 // CHECK:       .omp.reduction.then:
 155 // CHECK-NEXT:    [[TMP12:%.*]] = load double, ptr [[ARRAYIDX]], align 8
 156 // CHECK-NEXT:    [[TMP13:%.*]] = load double, ptr [[E2]], align 8
 157 // CHECK-NEXT:    [[ADD:%.*]] = fadd double [[TMP12]], [[TMP13]]
 158 // CHECK-NEXT:    store double [[ADD]], ptr [[ARRAYIDX]], align 8
 159 // CHECK-NEXT:    br label [[DOTOMP_REDUCTION_DONE]]
 160 // CHECK:       .omp.reduction.done:
 161 // CHECK-NEXT:    ret void
 162 //
 163 //
 164 // CHECK-LABEL: define {{[^@]+}}@_omp_reduction_shuffle_and_reduce_func
 165 // CHECK-SAME: (ptr noundef [[TMP0:%.*]], i16 noundef signext [[TMP1:%.*]], i16 noundef signext [[TMP2:%.*]], i16 noundef signext [[TMP3:%.*]]) #[[ATTR2:[0-9]+]] {
 166 // CHECK-NEXT:  entry:
 167 // CHECK-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 8
 168 // CHECK-NEXT:    [[DOTADDR1:%.*]] = alloca i16, align 2
 169 // CHECK-NEXT:    [[DOTADDR2:%.*]] = alloca i16, align 2
 170 // CHECK-NEXT:    [[DOTADDR3:%.*]] = alloca i16, align 2
 171 // CHECK-NEXT:    [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST:%.*]] = alloca [1 x ptr], align 8
 172 // CHECK-NEXT:    [[DOTOMP_REDUCTION_ELEMENT:%.*]] = alloca double, align 8
 173 // CHECK-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 8
 174 // CHECK-NEXT:    store i16 [[TMP1]], ptr [[DOTADDR1]], align 2
 175 // CHECK-NEXT:    store i16 [[TMP2]], ptr [[DOTADDR2]], align 2
 176 // CHECK-NEXT:    store i16 [[TMP3]], ptr [[DOTADDR3]], align 2
 177 // CHECK-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTADDR]], align 8
 178 // CHECK-NEXT:    [[TMP5:%.*]] = load i16, ptr [[DOTADDR1]], align 2
 179 // CHECK-NEXT:    [[TMP6:%.*]] = load i16, ptr [[DOTADDR2]], align 2
 180 // CHECK-NEXT:    [[TMP7:%.*]] = load i16, ptr [[DOTADDR3]], align 2
 181 // CHECK-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP4]], i64 0, i64 0
 182 // CHECK-NEXT:    [[TMP9:%.*]] = load ptr, ptr [[TMP8]], align 8
 183 // CHECK-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i64 0, i64 0
 184 // CHECK-NEXT:    [[TMP11:%.*]] = getelementptr double, ptr [[TMP9]], i64 1
 185 // CHECK-NEXT:    [[TMP12:%.*]] = load i64, ptr [[TMP9]], align 8
 186 // CHECK-NEXT:    [[TMP13:%.*]] = call i32 @__kmpc_get_warp_size()
 187 // CHECK-NEXT:    [[TMP14:%.*]] = trunc i32 [[TMP13]] to i16
 188 // CHECK-NEXT:    [[TMP15:%.*]] = call i64 @__kmpc_shuffle_int64(i64 [[TMP12]], i16 [[TMP6]], i16 [[TMP14]])
 189 // CHECK-NEXT:    store i64 [[TMP15]], ptr [[DOTOMP_REDUCTION_ELEMENT]], align 8
 190 // CHECK-NEXT:    [[TMP16:%.*]] = getelementptr i64, ptr [[TMP9]], i64 1
 191 // CHECK-NEXT:    [[TMP17:%.*]] = getelementptr i64, ptr [[DOTOMP_REDUCTION_ELEMENT]], i64 1
 192 // CHECK-NEXT:    store ptr [[DOTOMP_REDUCTION_ELEMENT]], ptr [[TMP10]], align 8
 193 // CHECK-NEXT:    [[TMP18:%.*]] = icmp eq i16 [[TMP7]], 0
 194 // CHECK-NEXT:    [[TMP19:%.*]] = icmp eq i16 [[TMP7]], 1
 195 // CHECK-NEXT:    [[TMP20:%.*]] = icmp ult i16 [[TMP5]], [[TMP6]]
 196 // CHECK-NEXT:    [[TMP21:%.*]] = and i1 [[TMP19]], [[TMP20]]
 197 // CHECK-NEXT:    [[TMP22:%.*]] = icmp eq i16 [[TMP7]], 2
 198 // CHECK-NEXT:    [[TMP23:%.*]] = and i16 [[TMP5]], 1
 199 // CHECK-NEXT:    [[TMP24:%.*]] = icmp eq i16 [[TMP23]], 0
 200 // CHECK-NEXT:    [[TMP25:%.*]] = and i1 [[TMP22]], [[TMP24]]
 201 // CHECK-NEXT:    [[TMP26:%.*]] = icmp sgt i16 [[TMP6]], 0
 202 // CHECK-NEXT:    [[TMP27:%.*]] = and i1 [[TMP25]], [[TMP26]]
 203 // CHECK-NEXT:    [[TMP28:%.*]] = or i1 [[TMP18]], [[TMP21]]
 204 // CHECK-NEXT:    [[TMP29:%.*]] = or i1 [[TMP28]], [[TMP27]]
 205 // CHECK-NEXT:    br i1 [[TMP29]], label [[THEN:%.*]], label [[ELSE:%.*]]
 206 // CHECK:       then:
 207 // CHECK-NEXT:    call void @"{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooi_l32_omp_outlined_omp$reduction$reduction_func"(ptr [[TMP4]], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]]) #[[ATTR3:[0-9]+]]
 208 // CHECK-NEXT:    br label [[IFCONT:%.*]]
 209 // CHECK:       else:
 210 // CHECK-NEXT:    br label [[IFCONT]]
 211 // CHECK:       ifcont:
 212 // CHECK-NEXT:    [[TMP30:%.*]] = icmp eq i16 [[TMP7]], 1
 213 // CHECK-NEXT:    [[TMP31:%.*]] = icmp uge i16 [[TMP5]], [[TMP6]]
 214 // CHECK-NEXT:    [[TMP32:%.*]] = and i1 [[TMP30]], [[TMP31]]
 215 // CHECK-NEXT:    br i1 [[TMP32]], label [[THEN4:%.*]], label [[ELSE5:%.*]]
 216 // CHECK:       then4:
 217 // CHECK-NEXT:    [[TMP33:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i64 0, i64 0
 218 // CHECK-NEXT:    [[TMP34:%.*]] = load ptr, ptr [[TMP33]], align 8
 219 // CHECK-NEXT:    [[TMP35:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP4]], i64 0, i64 0
 220 // CHECK-NEXT:    [[TMP36:%.*]] = load ptr, ptr [[TMP35]], align 8
 221 // CHECK-NEXT:    [[TMP37:%.*]] = load double, ptr [[TMP34]], align 8
 222 // CHECK-NEXT:    store double [[TMP37]], ptr [[TMP36]], align 8
 223 // CHECK-NEXT:    br label [[IFCONT6:%.*]]
 224 // CHECK:       else5:
 225 // CHECK-NEXT:    br label [[IFCONT6]]
 226 // CHECK:       ifcont6:
 227 // CHECK-NEXT:    ret void
 228 //
 229 //
 230 // CHECK-LABEL: define {{[^@]+}}@_omp_reduction_inter_warp_copy_func
 231 // CHECK-SAME: (ptr noundef [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR2]] {
 232 // CHECK-NEXT:  entry:
 233 // CHECK-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 8
 234 // CHECK-NEXT:    [[DOTADDR1:%.*]] = alloca i32, align 4
 235 // CHECK-NEXT:    [[DOTCNT_ADDR:%.*]] = alloca i32, align 4
 236 // CHECK-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 8
 237 // CHECK-NEXT:    store i32 [[TMP1]], ptr [[DOTADDR1]], align 4
 238 // CHECK-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()
 239 // CHECK-NEXT:    [[TMP3:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()
 240 // CHECK-NEXT:    [[NVPTX_LANE_ID:%.*]] = and i32 [[TMP3]], 31
 241 // CHECK-NEXT:    [[TMP4:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()
 242 // CHECK-NEXT:    [[NVPTX_WARP_ID:%.*]] = ashr i32 [[TMP4]], 5
 243 // CHECK-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[DOTADDR]], align 8
 244 // CHECK-NEXT:    store i32 0, ptr [[DOTCNT_ADDR]], align 4
 245 // CHECK-NEXT:    br label [[PRECOND:%.*]]
 246 // CHECK:       precond:
 247 // CHECK-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCNT_ADDR]], align 4
 248 // CHECK-NEXT:    [[TMP7:%.*]] = icmp ult i32 [[TMP6]], 2
 249 // CHECK-NEXT:    br i1 [[TMP7]], label [[BODY:%.*]], label [[EXIT:%.*]]
 250 // CHECK:       body:
 251 // CHECK-NEXT:    [[OMP_GLOBAL_THREAD_NUM:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
 252 // CHECK-NEXT:    call void @__kmpc_barrier(ptr @[[GLOB2:[0-9]+]], i32 [[OMP_GLOBAL_THREAD_NUM]])
 253 // CHECK-NEXT:    [[WARP_MASTER:%.*]] = icmp eq i32 [[NVPTX_LANE_ID]], 0
 254 // CHECK-NEXT:    br i1 [[WARP_MASTER]], label [[THEN:%.*]], label [[ELSE:%.*]]
 255 // CHECK:       then:
 256 // CHECK-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP5]], i64 0, i64 0
 257 // CHECK-NEXT:    [[TMP9:%.*]] = load ptr, ptr [[TMP8]], align 8
 258 // CHECK-NEXT:    [[TMP10:%.*]] = getelementptr i32, ptr [[TMP9]], i32 [[TMP6]]
 259 // CHECK-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[NVPTX_WARP_ID]]
 260 // CHECK-NEXT:    [[TMP12:%.*]] = load i32, ptr [[TMP10]], align 4
 261 // CHECK-NEXT:    store volatile i32 [[TMP12]], ptr addrspace(3) [[TMP11]], align 4
 262 // CHECK-NEXT:    br label [[IFCONT:%.*]]
 263 // CHECK:       else:
 264 // CHECK-NEXT:    br label [[IFCONT]]
 265 // CHECK:       ifcont:
 266 // CHECK-NEXT:    [[OMP_GLOBAL_THREAD_NUM2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
 267 // CHECK-NEXT:    call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[OMP_GLOBAL_THREAD_NUM2]])
 268 // CHECK-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTADDR1]], align 4
 269 // CHECK-NEXT:    [[IS_ACTIVE_THREAD:%.*]] = icmp ult i32 [[TMP2]], [[TMP13]]
 270 // CHECK-NEXT:    br i1 [[IS_ACTIVE_THREAD]], label [[THEN3:%.*]], label [[ELSE4:%.*]]
 271 // CHECK:       then3:
 272 // CHECK-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[TMP2]]
 273 // CHECK-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP5]], i64 0, i64 0
 274 // CHECK-NEXT:    [[TMP16:%.*]] = load ptr, ptr [[TMP15]], align 8
 275 // CHECK-NEXT:    [[TMP17:%.*]] = getelementptr i32, ptr [[TMP16]], i32 [[TMP6]]
 276 // CHECK-NEXT:    [[TMP18:%.*]] = load volatile i32, ptr addrspace(3) [[TMP14]], align 4
 277 // CHECK-NEXT:    store i32 [[TMP18]], ptr [[TMP17]], align 4
 278 // CHECK-NEXT:    br label [[IFCONT5:%.*]]
 279 // CHECK:       else4:
 280 // CHECK-NEXT:    br label [[IFCONT5]]
 281 // CHECK:       ifcont5:
 282 // CHECK-NEXT:    [[TMP19:%.*]] = add nsw i32 [[TMP6]], 1
 283 // CHECK-NEXT:    store i32 [[TMP19]], ptr [[DOTCNT_ADDR]], align 4
 284 // CHECK-NEXT:    br label [[PRECOND]]
 285 // CHECK:       exit:
 286 // CHECK-NEXT:    ret void
 287 //
 288 //
 289 // CHECK1-LABEL: define {{[^@]+}}@_Z3barv
 290 // CHECK1-SAME: () #[[ATTR0:[0-9]+]] {
 291 // CHECK1-NEXT:  entry:
 292 // CHECK1-NEXT:    [[O:%.*]] = alloca [5 x %class.S2], align 4
 293 // CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8
 294 // CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8
 295 // CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8
 296 // CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8
 297 // CHECK1-NEXT:    [[B:%.*]] = alloca [10 x [10 x [10 x double]]], align 8
 298 // CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS1:%.*]] = alloca [1 x ptr], align 8
 299 // CHECK1-NEXT:    [[DOTOFFLOAD_PTRS2:%.*]] = alloca [1 x ptr], align 8
 300 // CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS3:%.*]] = alloca [1 x ptr], align 8
 301 // CHECK1-NEXT:    [[KERNEL_ARGS4:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8
 302 // CHECK1-NEXT:    [[ARRAY_BEGIN:%.*]] = getelementptr inbounds [5 x %class.S2], ptr [[O]], i32 0, i32 0
 303 // CHECK1-NEXT:    [[ARRAYCTOR_END:%.*]] = getelementptr inbounds [[CLASS_S2:%.*]], ptr [[ARRAY_BEGIN]], i64 5
 304 // CHECK1-NEXT:    br label [[ARRAYCTOR_LOOP:%.*]]
 305 // CHECK1:       arrayctor.loop:
 306 // CHECK1-NEXT:    [[ARRAYCTOR_CUR:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ]
 307 // CHECK1-NEXT:    call void @_ZN2S2C1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]])
 308 // CHECK1-NEXT:    [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[CLASS_S2]], ptr [[ARRAYCTOR_CUR]], i64 1
 309 // CHECK1-NEXT:    [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], [[ARRAYCTOR_END]]
 310 // CHECK1-NEXT:    br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]]
 311 // CHECK1:       arrayctor.cont:
 312 // CHECK1-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [5 x %class.S2], ptr [[O]], i64 0, i64 0
 313 // CHECK1-NEXT:    [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
 314 // CHECK1-NEXT:    store ptr [[O]], ptr [[TMP0]], align 8
 315 // CHECK1-NEXT:    [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
 316 // CHECK1-NEXT:    store ptr [[ARRAYIDX]], ptr [[TMP1]], align 8
 317 // CHECK1-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
 318 // CHECK1-NEXT:    store ptr null, ptr [[TMP2]], align 8
 319 // CHECK1-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
 320 // CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
 321 // CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0
 322 // CHECK1-NEXT:    store i32 3, ptr [[TMP5]], align 4
 323 // CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1
 324 // CHECK1-NEXT:    store i32 1, ptr [[TMP6]], align 4
 325 // CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2
 326 // CHECK1-NEXT:    store ptr [[TMP3]], ptr [[TMP7]], align 8
 327 // CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3
 328 // CHECK1-NEXT:    store ptr [[TMP4]], ptr [[TMP8]], align 8
 329 // CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4
 330 // CHECK1-NEXT:    store ptr @.offload_sizes, ptr [[TMP9]], align 8
 331 // CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5
 332 // CHECK1-NEXT:    store ptr @.offload_maptypes, ptr [[TMP10]], align 8
 333 // CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6
 334 // CHECK1-NEXT:    store ptr null, ptr [[TMP11]], align 8
 335 // CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7
 336 // CHECK1-NEXT:    store ptr null, ptr [[TMP12]], align 8
 337 // CHECK1-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8
 338 // CHECK1-NEXT:    store i64 0, ptr [[TMP13]], align 8
 339 // CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9
 340 // CHECK1-NEXT:    store i64 0, ptr [[TMP14]], align 8
 341 // CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10
 342 // CHECK1-NEXT:    store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP15]], align 4
 343 // CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11
 344 // CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4
 345 // CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12
 346 // CHECK1-NEXT:    store i32 0, ptr [[TMP17]], align 4
 347 // CHECK1-NEXT:    [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2:[0-9]+]], i64 -1, i32 1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barv_l50.region_id, ptr [[KERNEL_ARGS]])
 348 // CHECK1-NEXT:    [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0
 349 // CHECK1-NEXT:    br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
 350 // CHECK1:       omp_offload.failed:
 351 // CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barv_l50(ptr [[O]]) #[[ATTR6:[0-9]+]]
 352 // CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]
 353 // CHECK1:       omp_offload.cont:
 354 // CHECK1-NEXT:    [[TMP20:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0
 355 // CHECK1-NEXT:    store ptr [[B]], ptr [[TMP20]], align 8
 356 // CHECK1-NEXT:    [[TMP21:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 0
 357 // CHECK1-NEXT:    store ptr [[B]], ptr [[TMP21]], align 8
 358 // CHECK1-NEXT:    [[TMP22:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS3]], i64 0, i64 0
 359 // CHECK1-NEXT:    store ptr null, ptr [[TMP22]], align 8
 360 // CHECK1-NEXT:    [[TMP23:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0
 361 // CHECK1-NEXT:    [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 0
 362 // CHECK1-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 0
 363 // CHECK1-NEXT:    store i32 3, ptr [[TMP25]], align 4
 364 // CHECK1-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 1
 365 // CHECK1-NEXT:    store i32 1, ptr [[TMP26]], align 4
 366 // CHECK1-NEXT:    [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 2
 367 // CHECK1-NEXT:    store ptr [[TMP23]], ptr [[TMP27]], align 8
 368 // CHECK1-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 3
 369 // CHECK1-NEXT:    store ptr [[TMP24]], ptr [[TMP28]], align 8
 370 // CHECK1-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 4
 371 // CHECK1-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP29]], align 8
 372 // CHECK1-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 5
 373 // CHECK1-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP30]], align 8
 374 // CHECK1-NEXT:    [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 6
 375 // CHECK1-NEXT:    store ptr null, ptr [[TMP31]], align 8
 376 // CHECK1-NEXT:    [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 7
 377 // CHECK1-NEXT:    store ptr null, ptr [[TMP32]], align 8
 378 // CHECK1-NEXT:    [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 8
 379 // CHECK1-NEXT:    store i64 0, ptr [[TMP33]], align 8
 380 // CHECK1-NEXT:    [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 9
 381 // CHECK1-NEXT:    store i64 0, ptr [[TMP34]], align 8
 382 // CHECK1-NEXT:    [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 10
 383 // CHECK1-NEXT:    store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP35]], align 4
 384 // CHECK1-NEXT:    [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 11
 385 // CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP36]], align 4
 386 // CHECK1-NEXT:    [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 12
 387 // CHECK1-NEXT:    store i32 0, ptr [[TMP37]], align 4
 388 // CHECK1-NEXT:    [[TMP38:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barv_l55.region_id, ptr [[KERNEL_ARGS4]])
 389 // CHECK1-NEXT:    [[TMP39:%.*]] = icmp ne i32 [[TMP38]], 0
 390 // CHECK1-NEXT:    br i1 [[TMP39]], label [[OMP_OFFLOAD_FAILED5:%.*]], label [[OMP_OFFLOAD_CONT6:%.*]]
 391 // CHECK1:       omp_offload.failed5:
 392 // CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barv_l55(ptr [[B]]) #[[ATTR6]]
 393 // CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT6]]
 394 // CHECK1:       omp_offload.cont6:
 395 // CHECK1-NEXT:    ret i32 0
 396 //
 397 //
 398 // CHECK1-LABEL: define {{[^@]+}}@_ZN2S2C1Ev
 399 // CHECK1-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR0]] comdat {
 400 // CHECK1-NEXT:  entry:
 401 // CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 402 // CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 403 // CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 404 // CHECK1-NEXT:    call void @_ZN2S2C2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]])
 405 // CHECK1-NEXT:    ret void
 406 //
 407 //
 408 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barv_l50
 409 // CHECK1-SAME: (ptr noundef nonnull align 4 dereferenceable(20) [[O:%.*]]) #[[ATTR1:[0-9]+]] {
 410 // CHECK1-NEXT:  entry:
 411 // CHECK1-NEXT:    [[O_ADDR:%.*]] = alloca ptr, align 8
 412 // CHECK1-NEXT:    store ptr [[O]], ptr [[O_ADDR]], align 8
 413 // CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[O_ADDR]], align 8
 414 // CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barv_l50.omp_outlined, ptr [[TMP0]])
 415 // CHECK1-NEXT:    ret void
 416 //
 417 //
 418 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barv_l50.omp_outlined
 419 // CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(20) [[O:%.*]]) #[[ATTR1]] {
 420 // CHECK1-NEXT:  entry:
 421 // CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
 422 // CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
 423 // CHECK1-NEXT:    [[O_ADDR:%.*]] = alloca ptr, align 8
 424 // CHECK1-NEXT:    [[O1:%.*]] = alloca [[CLASS_S2:%.*]], align 4
 425 // CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4
 426 // CHECK1-NEXT:    [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 8
 427 // CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
 428 // CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
 429 // CHECK1-NEXT:    store ptr [[O]], ptr [[O_ADDR]], align 8
 430 // CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[O_ADDR]], align 8
 431 // CHECK1-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [5 x %class.S2], ptr [[TMP0]], i64 0, i64 0
 432 // CHECK1-NEXT:    call void @_ZN2S2C1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[O1]])
 433 // CHECK1-NEXT:    [[TMP1:%.*]] = ptrtoint ptr [[TMP0]] to i64
 434 // CHECK1-NEXT:    [[TMP2:%.*]] = ptrtoint ptr [[ARRAYIDX]] to i64
 435 // CHECK1-NEXT:    [[TMP3:%.*]] = sub i64 [[TMP1]], [[TMP2]]
 436 // CHECK1-NEXT:    [[TMP4:%.*]] = sdiv exact i64 [[TMP3]], ptrtoint (ptr getelementptr ([[CLASS_S2]], ptr null, i32 1) to i64)
 437 // CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr [[CLASS_S2]], ptr [[O1]], i64 [[TMP4]]
 438 // CHECK1-NEXT:    store i32 0, ptr [[I]], align 4
 439 // CHECK1-NEXT:    br label [[FOR_COND:%.*]]
 440 // CHECK1:       for.cond:
 441 // CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[I]], align 4
 442 // CHECK1-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP6]], 10
 443 // CHECK1-NEXT:    br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]]
 444 // CHECK1:       for.body:
 445 // CHECK1-NEXT:    br label [[FOR_INC:%.*]]
 446 // CHECK1:       for.inc:
 447 // CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[I]], align 4
 448 // CHECK1-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP7]], 1
 449 // CHECK1-NEXT:    store i32 [[INC]], ptr [[I]], align 4
 450 // CHECK1-NEXT:    br label [[FOR_COND]], !llvm.loop [[LOOP5:![0-9]+]]
 451 // CHECK1:       for.end:
 452 // CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0
 453 // CHECK1-NEXT:    store ptr [[O1]], ptr [[TMP8]], align 8
 454 // CHECK1-NEXT:    [[TMP9:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
 455 // CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[TMP9]], align 4
 456 // CHECK1-NEXT:    [[TMP11:%.*]] = call i32 @__kmpc_reduce_nowait(ptr @[[GLOB1:[0-9]+]], i32 [[TMP10]], i32 1, i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barv_l50.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)
 457 // CHECK1-NEXT:    switch i32 [[TMP11]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [
 458 // CHECK1-NEXT:      i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]
 459 // CHECK1-NEXT:      i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]
 460 // CHECK1-NEXT:    ]
 461 // CHECK1:       .omp.reduction.case1:
 462 // CHECK1-NEXT:    [[CALL:%.*]] = call noundef nonnull align 4 dereferenceable(4) ptr @_ZN2S2plERS_(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYIDX]], ptr noundef nonnull align 4 dereferenceable(4) [[O1]])
 463 // CHECK1-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[ARRAYIDX]], ptr align 4 [[CALL]], i64 4, i1 false)
 464 // CHECK1-NEXT:    call void @__kmpc_end_reduce_nowait(ptr @[[GLOB1]], i32 [[TMP10]], ptr @.gomp_critical_user_.reduction.var)
 465 // CHECK1-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]
 466 // CHECK1:       .omp.reduction.case2:
 467 // CHECK1-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
 468 // CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[TMP12]], align 4
 469 // CHECK1-NEXT:    call void @__kmpc_critical(ptr @[[GLOB2]], i32 [[TMP13]], ptr @.gomp_critical_user_.atomic_reduction.var)
 470 // CHECK1-NEXT:    [[CALL2:%.*]] = call noundef nonnull align 4 dereferenceable(4) ptr @_ZN2S2plERS_(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYIDX]], ptr noundef nonnull align 4 dereferenceable(4) [[O1]])
 471 // CHECK1-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[ARRAYIDX]], ptr align 4 [[CALL2]], i64 4, i1 false)
 472 // CHECK1-NEXT:    call void @__kmpc_end_critical(ptr @[[GLOB2]], i32 [[TMP13]], ptr @.gomp_critical_user_.atomic_reduction.var)
 473 // CHECK1-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]
 474 // CHECK1:       .omp.reduction.default:
 475 // CHECK1-NEXT:    ret void
 476 //
 477 //
 478 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barv_l50.omp_outlined.omp.reduction.reduction_func
 479 // CHECK1-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR2:[0-9]+]] {
 480 // CHECK1-NEXT:  entry:
 481 // CHECK1-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 8
 482 // CHECK1-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 8
 483 // CHECK1-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 8
 484 // CHECK1-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 8
 485 // CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 8
 486 // CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 8
 487 // CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i64 0, i64 0
 488 // CHECK1-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 8
 489 // CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i64 0, i64 0
 490 // CHECK1-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 8
 491 // CHECK1-NEXT:    [[CALL:%.*]] = call noundef nonnull align 4 dereferenceable(4) ptr @_ZN2S2plERS_(ptr noundef nonnull align 4 dereferenceable(4) [[TMP7]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP5]])
 492 // CHECK1-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[TMP7]], ptr align 4 [[CALL]], i64 4, i1 false)
 493 // CHECK1-NEXT:    ret void
 494 //
 495 //
 496 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barv_l55
 497 // CHECK1-SAME: (ptr noundef nonnull align 8 dereferenceable(8000) [[B:%.*]]) #[[ATTR1]] {
 498 // CHECK1-NEXT:  entry:
 499 // CHECK1-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8
 500 // CHECK1-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8
 501 // CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 8
 502 // CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barv_l55.omp_outlined, ptr [[TMP0]])
 503 // CHECK1-NEXT:    ret void
 504 //
 505 //
 506 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barv_l55.omp_outlined
 507 // CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 8 dereferenceable(8000) [[B:%.*]]) #[[ATTR1]] {
 508 // CHECK1-NEXT:  entry:
 509 // CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
 510 // CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
 511 // CHECK1-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8
 512 // CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i64, align 8
 513 // CHECK1-NEXT:    [[TMP:%.*]] = alloca i64, align 8
 514 // CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i64, align 8
 515 // CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i64, align 8
 516 // CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i64, align 8
 517 // CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
 518 // CHECK1-NEXT:    [[SAVED_STACK:%.*]] = alloca ptr, align 8
 519 // CHECK1-NEXT:    [[__VLA_EXPR0:%.*]] = alloca i64, align 8
 520 // CHECK1-NEXT:    [[DOTRD_INPUT_:%.*]] = alloca [1 x %struct.kmp_taskred_input_t], align 8
 521 // CHECK1-NEXT:    [[DOTTASK_RED_:%.*]] = alloca ptr, align 8
 522 // CHECK1-NEXT:    [[I:%.*]] = alloca i64, align 8
 523 // CHECK1-NEXT:    [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [2 x ptr], align 8
 524 // CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
 525 // CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
 526 // CHECK1-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8
 527 // CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 8
 528 // CHECK1-NEXT:    store i64 0, ptr [[DOTOMP_LB]], align 8
 529 // CHECK1-NEXT:    store i64 9, ptr [[DOTOMP_UB]], align 8
 530 // CHECK1-NEXT:    store i64 1, ptr [[DOTOMP_STRIDE]], align 8
 531 // CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4
 532 // CHECK1-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds nuw [10 x [10 x [10 x double]]], ptr [[TMP0]], i64 0, i64 0
 533 // CHECK1-NEXT:    [[ARRAYDECAY:%.*]] = getelementptr inbounds [10 x [10 x double]], ptr [[ARRAYIDX]], i64 0, i64 0
 534 // CHECK1-NEXT:    [[ARRAYIDX1:%.*]] = getelementptr inbounds nuw [10 x double], ptr [[ARRAYDECAY]], i64 2
 535 // CHECK1-NEXT:    [[ARRAYDECAY2:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX1]], i64 0, i64 0
 536 // CHECK1-NEXT:    [[ARRAYIDX3:%.*]] = getelementptr inbounds nuw double, ptr [[ARRAYDECAY2]], i64 1
 537 // CHECK1-NEXT:    [[ARRAYIDX4:%.*]] = getelementptr inbounds nuw [10 x [10 x [10 x double]]], ptr [[TMP0]], i64 0, i64 1
 538 // CHECK1-NEXT:    [[ARRAYDECAY5:%.*]] = getelementptr inbounds [10 x [10 x double]], ptr [[ARRAYIDX4]], i64 0, i64 0
 539 // CHECK1-NEXT:    [[ARRAYIDX6:%.*]] = getelementptr inbounds nuw [10 x double], ptr [[ARRAYDECAY5]], i64 5
 540 // CHECK1-NEXT:    [[ARRAYDECAY7:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX6]], i64 0, i64 0
 541 // CHECK1-NEXT:    [[ARRAYIDX8:%.*]] = getelementptr inbounds nuw double, ptr [[ARRAYDECAY7]], i64 1
 542 // CHECK1-NEXT:    [[TMP1:%.*]] = ptrtoint ptr [[ARRAYIDX8]] to i64
 543 // CHECK1-NEXT:    [[TMP2:%.*]] = ptrtoint ptr [[ARRAYIDX3]] to i64
 544 // CHECK1-NEXT:    [[TMP3:%.*]] = sub i64 [[TMP1]], [[TMP2]]
 545 // CHECK1-NEXT:    [[TMP4:%.*]] = sdiv exact i64 [[TMP3]], ptrtoint (ptr getelementptr (double, ptr null, i32 1) to i64)
 546 // CHECK1-NEXT:    [[TMP5:%.*]] = add nuw i64 [[TMP4]], 1
 547 // CHECK1-NEXT:    [[TMP6:%.*]] = mul nuw i64 [[TMP5]], ptrtoint (ptr getelementptr (double, ptr null, i32 1) to i64)
 548 // CHECK1-NEXT:    [[TMP7:%.*]] = call ptr @llvm.stacksave.p0()
 549 // CHECK1-NEXT:    store ptr [[TMP7]], ptr [[SAVED_STACK]], align 8
 550 // CHECK1-NEXT:    [[VLA:%.*]] = alloca double, i64 [[TMP5]], align 8
 551 // CHECK1-NEXT:    store i64 [[TMP5]], ptr [[__VLA_EXPR0]], align 8
 552 // CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr double, ptr [[VLA]], i64 [[TMP5]]
 553 // CHECK1-NEXT:    [[OMP_ARRAYINIT_ISEMPTY:%.*]] = icmp eq ptr [[VLA]], [[TMP8]]
 554 // CHECK1-NEXT:    br i1 [[OMP_ARRAYINIT_ISEMPTY]], label [[OMP_ARRAYINIT_DONE:%.*]], label [[OMP_ARRAYINIT_BODY:%.*]]
 555 // CHECK1:       omp.arrayinit.body:
 556 // CHECK1-NEXT:    [[OMP_ARRAYCPY_DESTELEMENTPAST:%.*]] = phi ptr [ [[VLA]], [[ENTRY:%.*]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT:%.*]], [[OMP_ARRAYINIT_BODY]] ]
 557 // CHECK1-NEXT:    store double 0.000000e+00, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], align 8
 558 // CHECK1-NEXT:    [[OMP_ARRAYCPY_DEST_ELEMENT]] = getelementptr double, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], i32 1
 559 // CHECK1-NEXT:    [[OMP_ARRAYCPY_DONE:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT]], [[TMP8]]
 560 // CHECK1-NEXT:    br i1 [[OMP_ARRAYCPY_DONE]], label [[OMP_ARRAYINIT_DONE]], label [[OMP_ARRAYINIT_BODY]]
 561 // CHECK1:       omp.arrayinit.done:
 562 // CHECK1-NEXT:    [[TMP9:%.*]] = ptrtoint ptr [[TMP0]] to i64
 563 // CHECK1-NEXT:    [[TMP10:%.*]] = ptrtoint ptr [[ARRAYIDX3]] to i64
 564 // CHECK1-NEXT:    [[TMP11:%.*]] = sub i64 [[TMP9]], [[TMP10]]
 565 // CHECK1-NEXT:    [[TMP12:%.*]] = sdiv exact i64 [[TMP11]], ptrtoint (ptr getelementptr (double, ptr null, i32 1) to i64)
 566 // CHECK1-NEXT:    [[TMP13:%.*]] = getelementptr double, ptr [[VLA]], i64 [[TMP12]]
 567 // CHECK1-NEXT:    [[DOTRD_INPUT_GEP_:%.*]] = getelementptr inbounds nuw [1 x %struct.kmp_taskred_input_t], ptr [[DOTRD_INPUT_]], i64 0, i64 0
 568 // CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASKRED_INPUT_T:%.*]], ptr [[DOTRD_INPUT_GEP_]], i32 0, i32 0
 569 // CHECK1-NEXT:    [[ARRAYIDX9:%.*]] = getelementptr inbounds nuw [10 x [10 x [10 x double]]], ptr [[TMP0]], i64 0, i64 0
 570 // CHECK1-NEXT:    [[ARRAYDECAY10:%.*]] = getelementptr inbounds [10 x [10 x double]], ptr [[ARRAYIDX9]], i64 0, i64 0
 571 // CHECK1-NEXT:    [[ARRAYIDX11:%.*]] = getelementptr inbounds nuw [10 x double], ptr [[ARRAYDECAY10]], i64 2
 572 // CHECK1-NEXT:    [[ARRAYDECAY12:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX11]], i64 0, i64 0
 573 // CHECK1-NEXT:    [[ARRAYIDX13:%.*]] = getelementptr inbounds nuw double, ptr [[ARRAYDECAY12]], i64 1
 574 // CHECK1-NEXT:    [[ARRAYIDX14:%.*]] = getelementptr inbounds nuw [10 x [10 x [10 x double]]], ptr [[TMP0]], i64 0, i64 1
 575 // CHECK1-NEXT:    [[ARRAYDECAY15:%.*]] = getelementptr inbounds [10 x [10 x double]], ptr [[ARRAYIDX14]], i64 0, i64 0
 576 // CHECK1-NEXT:    [[ARRAYIDX16:%.*]] = getelementptr inbounds nuw [10 x double], ptr [[ARRAYDECAY15]], i64 5
 577 // CHECK1-NEXT:    [[ARRAYDECAY17:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX16]], i64 0, i64 0
 578 // CHECK1-NEXT:    [[ARRAYIDX18:%.*]] = getelementptr inbounds nuw double, ptr [[ARRAYDECAY17]], i64 1
 579 // CHECK1-NEXT:    store ptr [[VLA]], ptr [[TMP14]], align 8
 580 // CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASKRED_INPUT_T]], ptr [[DOTRD_INPUT_GEP_]], i32 0, i32 1
 581 // CHECK1-NEXT:    store ptr [[ARRAYIDX13]], ptr [[TMP15]], align 8
 582 // CHECK1-NEXT:    [[TMP16:%.*]] = ptrtoint ptr [[ARRAYIDX18]] to i64
 583 // CHECK1-NEXT:    [[TMP17:%.*]] = ptrtoint ptr [[ARRAYIDX13]] to i64
 584 // CHECK1-NEXT:    [[TMP18:%.*]] = sub i64 [[TMP16]], [[TMP17]]
 585 // CHECK1-NEXT:    [[TMP19:%.*]] = sdiv exact i64 [[TMP18]], ptrtoint (ptr getelementptr (double, ptr null, i32 1) to i64)
 586 // CHECK1-NEXT:    [[TMP20:%.*]] = add nuw i64 [[TMP19]], 1
 587 // CHECK1-NEXT:    [[TMP21:%.*]] = mul nuw i64 [[TMP20]], ptrtoint (ptr getelementptr (double, ptr null, i32 1) to i64)
 588 // CHECK1-NEXT:    [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASKRED_INPUT_T]], ptr [[DOTRD_INPUT_GEP_]], i32 0, i32 2
 589 // CHECK1-NEXT:    store i64 [[TMP21]], ptr [[TMP22]], align 8
 590 // CHECK1-NEXT:    [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASKRED_INPUT_T]], ptr [[DOTRD_INPUT_GEP_]], i32 0, i32 3
 591 // CHECK1-NEXT:    store ptr @.red_init., ptr [[TMP23]], align 8
 592 // CHECK1-NEXT:    [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASKRED_INPUT_T]], ptr [[DOTRD_INPUT_GEP_]], i32 0, i32 4
 593 // CHECK1-NEXT:    store ptr null, ptr [[TMP24]], align 8
 594 // CHECK1-NEXT:    [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASKRED_INPUT_T]], ptr [[DOTRD_INPUT_GEP_]], i32 0, i32 5
 595 // CHECK1-NEXT:    store ptr @.red_comb., ptr [[TMP25]], align 8
 596 // CHECK1-NEXT:    [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASKRED_INPUT_T]], ptr [[DOTRD_INPUT_GEP_]], i32 0, i32 6
 597 // CHECK1-NEXT:    store i32 1, ptr [[TMP26]], align 8
 598 // CHECK1-NEXT:    [[TMP27:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
 599 // CHECK1-NEXT:    [[TMP28:%.*]] = load i32, ptr [[TMP27]], align 4
 600 // CHECK1-NEXT:    [[TMP29:%.*]] = call ptr @__kmpc_taskred_modifier_init(ptr @[[GLOB2]], i32 [[TMP28]], i32 1, i32 1, ptr [[DOTRD_INPUT_]])
 601 // CHECK1-NEXT:    store ptr [[TMP29]], ptr [[DOTTASK_RED_]], align 8
 602 // CHECK1-NEXT:    [[TMP30:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
 603 // CHECK1-NEXT:    [[TMP31:%.*]] = load i32, ptr [[TMP30]], align 4
 604 // CHECK1-NEXT:    call void @__kmpc_for_static_init_8(ptr @[[GLOB3:[0-9]+]], i32 [[TMP31]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i64 1, i64 1)
 605 // CHECK1-NEXT:    [[TMP32:%.*]] = load i64, ptr [[DOTOMP_UB]], align 8
 606 // CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i64 [[TMP32]], 9
 607 // CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
 608 // CHECK1:       cond.true:
 609 // CHECK1-NEXT:    br label [[COND_END:%.*]]
 610 // CHECK1:       cond.false:
 611 // CHECK1-NEXT:    [[TMP33:%.*]] = load i64, ptr [[DOTOMP_UB]], align 8
 612 // CHECK1-NEXT:    br label [[COND_END]]
 613 // CHECK1:       cond.end:
 614 // CHECK1-NEXT:    [[COND:%.*]] = phi i64 [ 9, [[COND_TRUE]] ], [ [[TMP33]], [[COND_FALSE]] ]
 615 // CHECK1-NEXT:    store i64 [[COND]], ptr [[DOTOMP_UB]], align 8
 616 // CHECK1-NEXT:    [[TMP34:%.*]] = load i64, ptr [[DOTOMP_LB]], align 8
 617 // CHECK1-NEXT:    store i64 [[TMP34]], ptr [[DOTOMP_IV]], align 8
 618 // CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
 619 // CHECK1:       omp.inner.for.cond:
 620 // CHECK1-NEXT:    [[TMP35:%.*]] = load i64, ptr [[DOTOMP_IV]], align 8
 621 // CHECK1-NEXT:    [[TMP36:%.*]] = load i64, ptr [[DOTOMP_UB]], align 8
 622 // CHECK1-NEXT:    [[CMP19:%.*]] = icmp sle i64 [[TMP35]], [[TMP36]]
 623 // CHECK1-NEXT:    br i1 [[CMP19]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_COND_CLEANUP:%.*]]
 624 // CHECK1:       omp.inner.for.cond.cleanup:
 625 // CHECK1-NEXT:    br label [[OMP_INNER_FOR_END:%.*]]
 626 // CHECK1:       omp.inner.for.body:
 627 // CHECK1-NEXT:    [[TMP37:%.*]] = load i64, ptr [[DOTOMP_IV]], align 8
 628 // CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i64 [[TMP37]], 1
 629 // CHECK1-NEXT:    [[ADD:%.*]] = add nsw i64 0, [[MUL]]
 630 // CHECK1-NEXT:    store i64 [[ADD]], ptr [[I]], align 8
 631 // CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
 632 // CHECK1:       omp.body.continue:
 633 // CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
 634 // CHECK1:       omp.inner.for.inc:
 635 // CHECK1-NEXT:    [[TMP38:%.*]] = load i64, ptr [[DOTOMP_IV]], align 8
 636 // CHECK1-NEXT:    [[ADD20:%.*]] = add nsw i64 [[TMP38]], 1
 637 // CHECK1-NEXT:    store i64 [[ADD20]], ptr [[DOTOMP_IV]], align 8
 638 // CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]]
 639 // CHECK1:       omp.inner.for.end:
 640 // CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
 641 // CHECK1:       omp.loop.exit:
 642 // CHECK1-NEXT:    [[TMP39:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
 643 // CHECK1-NEXT:    [[TMP40:%.*]] = load i32, ptr [[TMP39]], align 4
 644 // CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB3]], i32 [[TMP40]])
 645 // CHECK1-NEXT:    [[TMP41:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
 646 // CHECK1-NEXT:    [[TMP42:%.*]] = load i32, ptr [[TMP41]], align 4
 647 // CHECK1-NEXT:    call void @__kmpc_task_reduction_modifier_fini(ptr @[[GLOB2]], i32 [[TMP42]], i32 1)
 648 // CHECK1-NEXT:    [[TMP43:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0
 649 // CHECK1-NEXT:    store ptr [[VLA]], ptr [[TMP43]], align 8
 650 // CHECK1-NEXT:    [[TMP44:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 1
 651 // CHECK1-NEXT:    [[TMP45:%.*]] = inttoptr i64 [[TMP5]] to ptr
 652 // CHECK1-NEXT:    store ptr [[TMP45]], ptr [[TMP44]], align 8
 653 // CHECK1-NEXT:    [[TMP46:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
 654 // CHECK1-NEXT:    [[TMP47:%.*]] = load i32, ptr [[TMP46]], align 4
 655 // CHECK1-NEXT:    [[TMP48:%.*]] = call i32 @__kmpc_reduce_nowait(ptr @[[GLOB1]], i32 [[TMP47]], i32 1, i64 16, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barv_l55.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)
 656 // CHECK1-NEXT:    switch i32 [[TMP48]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [
 657 // CHECK1-NEXT:      i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]
 658 // CHECK1-NEXT:      i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]
 659 // CHECK1-NEXT:    ]
 660 // CHECK1:       .omp.reduction.case1:
 661 // CHECK1-NEXT:    [[TMP49:%.*]] = getelementptr double, ptr [[ARRAYIDX3]], i64 [[TMP5]]
 662 // CHECK1-NEXT:    [[OMP_ARRAYCPY_ISEMPTY:%.*]] = icmp eq ptr [[ARRAYIDX3]], [[TMP49]]
 663 // CHECK1-NEXT:    br i1 [[OMP_ARRAYCPY_ISEMPTY]], label [[OMP_ARRAYCPY_DONE25:%.*]], label [[OMP_ARRAYCPY_BODY:%.*]]
 664 // CHECK1:       omp.arraycpy.body:
 665 // CHECK1-NEXT:    [[OMP_ARRAYCPY_SRCELEMENTPAST:%.*]] = phi ptr [ [[VLA]], [[DOTOMP_REDUCTION_CASE1]] ], [ [[OMP_ARRAYCPY_SRC_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
 666 // CHECK1-NEXT:    [[OMP_ARRAYCPY_DESTELEMENTPAST21:%.*]] = phi ptr [ [[ARRAYIDX3]], [[DOTOMP_REDUCTION_CASE1]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT23:%.*]], [[OMP_ARRAYCPY_BODY]] ]
 667 // CHECK1-NEXT:    [[TMP50:%.*]] = load double, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST21]], align 8
 668 // CHECK1-NEXT:    [[TMP51:%.*]] = load double, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], align 8
 669 // CHECK1-NEXT:    [[ADD22:%.*]] = fadd double [[TMP50]], [[TMP51]]
 670 // CHECK1-NEXT:    store double [[ADD22]], ptr [[OMP_ARRAYCPY_DESTELEMENTPAST21]], align 8
 671 // CHECK1-NEXT:    [[OMP_ARRAYCPY_DEST_ELEMENT23]] = getelementptr double, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST21]], i32 1
 672 // CHECK1-NEXT:    [[OMP_ARRAYCPY_SRC_ELEMENT]] = getelementptr double, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], i32 1
 673 // CHECK1-NEXT:    [[OMP_ARRAYCPY_DONE24:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT23]], [[TMP49]]
 674 // CHECK1-NEXT:    br i1 [[OMP_ARRAYCPY_DONE24]], label [[OMP_ARRAYCPY_DONE25]], label [[OMP_ARRAYCPY_BODY]]
 675 // CHECK1:       omp.arraycpy.done25:
 676 // CHECK1-NEXT:    call void @__kmpc_end_reduce_nowait(ptr @[[GLOB1]], i32 [[TMP47]], ptr @.gomp_critical_user_.reduction.var)
 677 // CHECK1-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]
 678 // CHECK1:       .omp.reduction.case2:
 679 // CHECK1-NEXT:    [[TMP52:%.*]] = getelementptr double, ptr [[ARRAYIDX3]], i64 [[TMP5]]
 680 // CHECK1-NEXT:    [[OMP_ARRAYCPY_ISEMPTY26:%.*]] = icmp eq ptr [[ARRAYIDX3]], [[TMP52]]
 681 // CHECK1-NEXT:    br i1 [[OMP_ARRAYCPY_ISEMPTY26]], label [[OMP_ARRAYCPY_DONE33:%.*]], label [[OMP_ARRAYCPY_BODY27:%.*]]
 682 // CHECK1:       omp.arraycpy.body27:
 683 // CHECK1-NEXT:    [[OMP_ARRAYCPY_SRCELEMENTPAST28:%.*]] = phi ptr [ [[VLA]], [[DOTOMP_REDUCTION_CASE2]] ], [ [[OMP_ARRAYCPY_SRC_ELEMENT31:%.*]], [[OMP_ARRAYCPY_BODY27]] ]
 684 // CHECK1-NEXT:    [[OMP_ARRAYCPY_DESTELEMENTPAST29:%.*]] = phi ptr [ [[ARRAYIDX3]], [[DOTOMP_REDUCTION_CASE2]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT30:%.*]], [[OMP_ARRAYCPY_BODY27]] ]
 685 // CHECK1-NEXT:    [[TMP53:%.*]] = load double, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST28]], align 8
 686 // CHECK1-NEXT:    [[TMP54:%.*]] = atomicrmw fadd ptr [[OMP_ARRAYCPY_DESTELEMENTPAST29]], double [[TMP53]] monotonic, align 8
 687 // CHECK1-NEXT:    [[OMP_ARRAYCPY_DEST_ELEMENT30]] = getelementptr double, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST29]], i32 1
 688 // CHECK1-NEXT:    [[OMP_ARRAYCPY_SRC_ELEMENT31]] = getelementptr double, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST28]], i32 1
 689 // CHECK1-NEXT:    [[OMP_ARRAYCPY_DONE32:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT30]], [[TMP52]]
 690 // CHECK1-NEXT:    br i1 [[OMP_ARRAYCPY_DONE32]], label [[OMP_ARRAYCPY_DONE33]], label [[OMP_ARRAYCPY_BODY27]]
 691 // CHECK1:       omp.arraycpy.done33:
 692 // CHECK1-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]
 693 // CHECK1:       .omp.reduction.default:
 694 // CHECK1-NEXT:    [[TMP55:%.*]] = load ptr, ptr [[SAVED_STACK]], align 8
 695 // CHECK1-NEXT:    call void @llvm.stackrestore.p0(ptr [[TMP55]])
 696 // CHECK1-NEXT:    ret void
 697 //
 698 //
 699 // CHECK1-LABEL: define {{[^@]+}}@.red_init.
 700 // CHECK1-SAME: (ptr noalias noundef [[TMP0:%.*]], ptr noalias noundef [[TMP1:%.*]]) #[[ATTR2]] {
 701 // CHECK1-NEXT:  entry:
 702 // CHECK1-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 8
 703 // CHECK1-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 8
 704 // CHECK1-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 8
 705 // CHECK1-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 8
 706 // CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 8
 707 // CHECK1-NEXT:    [[TMP3:%.*]] = call align 8 ptr @llvm.threadlocal.address.p0(ptr align 8 @{{reduction_size[.].+[.]}})
 708 // CHECK1-NEXT:    [[TMP4:%.*]] = load i64, ptr [[TMP3]], align 8
 709 // CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr double, ptr [[TMP2]], i64 [[TMP4]]
 710 // CHECK1-NEXT:    [[OMP_ARRAYINIT_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP5]]
 711 // CHECK1-NEXT:    br i1 [[OMP_ARRAYINIT_ISEMPTY]], label [[OMP_ARRAYINIT_DONE:%.*]], label [[OMP_ARRAYINIT_BODY:%.*]]
 712 // CHECK1:       omp.arrayinit.body:
 713 // CHECK1-NEXT:    [[OMP_ARRAYCPY_DESTELEMENTPAST:%.*]] = phi ptr [ [[TMP2]], [[ENTRY:%.*]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT:%.*]], [[OMP_ARRAYINIT_BODY]] ]
 714 // CHECK1-NEXT:    store double 0.000000e+00, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], align 8
 715 // CHECK1-NEXT:    [[OMP_ARRAYCPY_DEST_ELEMENT]] = getelementptr double, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], i32 1
 716 // CHECK1-NEXT:    [[OMP_ARRAYCPY_DONE:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT]], [[TMP5]]
 717 // CHECK1-NEXT:    br i1 [[OMP_ARRAYCPY_DONE]], label [[OMP_ARRAYINIT_DONE]], label [[OMP_ARRAYINIT_BODY]]
 718 // CHECK1:       omp.arrayinit.done:
 719 // CHECK1-NEXT:    ret void
 720 //
 721 //
 722 // CHECK1-LABEL: define {{[^@]+}}@.red_comb.
 723 // CHECK1-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR2]] {
 724 // CHECK1-NEXT:  entry:
 725 // CHECK1-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 8
 726 // CHECK1-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 8
 727 // CHECK1-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 8
 728 // CHECK1-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 8
 729 // CHECK1-NEXT:    [[TMP2:%.*]] = call align 8 ptr @llvm.threadlocal.address.p0(ptr align 8 @{{reduction_size[.].+[.]}})
 730 // CHECK1-NEXT:    [[TMP3:%.*]] = load i64, ptr [[TMP2]], align 8
 731 // CHECK1-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTADDR]], align 8
 732 // CHECK1-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[DOTADDR1]], align 8
 733 // CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr double, ptr [[TMP4]], i64 [[TMP3]]
 734 // CHECK1-NEXT:    [[OMP_ARRAYCPY_ISEMPTY:%.*]] = icmp eq ptr [[TMP4]], [[TMP6]]
 735 // CHECK1-NEXT:    br i1 [[OMP_ARRAYCPY_ISEMPTY]], label [[OMP_ARRAYCPY_DONE2:%.*]], label [[OMP_ARRAYCPY_BODY:%.*]]
 736 // CHECK1:       omp.arraycpy.body:
 737 // CHECK1-NEXT:    [[OMP_ARRAYCPY_SRCELEMENTPAST:%.*]] = phi ptr [ [[TMP5]], [[ENTRY:%.*]] ], [ [[OMP_ARRAYCPY_SRC_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
 738 // CHECK1-NEXT:    [[OMP_ARRAYCPY_DESTELEMENTPAST:%.*]] = phi ptr [ [[TMP4]], [[ENTRY]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
 739 // CHECK1-NEXT:    [[TMP7:%.*]] = load double, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], align 8
 740 // CHECK1-NEXT:    [[TMP8:%.*]] = load double, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], align 8
 741 // CHECK1-NEXT:    [[ADD:%.*]] = fadd double [[TMP7]], [[TMP8]]
 742 // CHECK1-NEXT:    store double [[ADD]], ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], align 8
 743 // CHECK1-NEXT:    [[OMP_ARRAYCPY_DEST_ELEMENT]] = getelementptr double, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], i32 1
 744 // CHECK1-NEXT:    [[OMP_ARRAYCPY_SRC_ELEMENT]] = getelementptr double, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], i32 1
 745 // CHECK1-NEXT:    [[OMP_ARRAYCPY_DONE:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT]], [[TMP6]]
 746 // CHECK1-NEXT:    br i1 [[OMP_ARRAYCPY_DONE]], label [[OMP_ARRAYCPY_DONE2]], label [[OMP_ARRAYCPY_BODY]]
 747 // CHECK1:       omp.arraycpy.done2:
 748 // CHECK1-NEXT:    ret void
 749 //
 750 //
 751 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barv_l55.omp_outlined.omp.reduction.reduction_func
 752 // CHECK1-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR2]] {
 753 // CHECK1-NEXT:  entry:
 754 // CHECK1-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 8
 755 // CHECK1-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 8
 756 // CHECK1-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 8
 757 // CHECK1-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 8
 758 // CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 8
 759 // CHECK1-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 8
 760 // CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP3]], i64 0, i64 0
 761 // CHECK1-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 8
 762 // CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP2]], i64 0, i64 0
 763 // CHECK1-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 8
 764 // CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP2]], i64 0, i64 1
 765 // CHECK1-NEXT:    [[TMP9:%.*]] = load ptr, ptr [[TMP8]], align 8
 766 // CHECK1-NEXT:    [[TMP10:%.*]] = ptrtoint ptr [[TMP9]] to i64
 767 // CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr double, ptr [[TMP7]], i64 [[TMP10]]
 768 // CHECK1-NEXT:    [[OMP_ARRAYCPY_ISEMPTY:%.*]] = icmp eq ptr [[TMP7]], [[TMP11]]
 769 // CHECK1-NEXT:    br i1 [[OMP_ARRAYCPY_ISEMPTY]], label [[OMP_ARRAYCPY_DONE2:%.*]], label [[OMP_ARRAYCPY_BODY:%.*]]
 770 // CHECK1:       omp.arraycpy.body:
 771 // CHECK1-NEXT:    [[OMP_ARRAYCPY_SRCELEMENTPAST:%.*]] = phi ptr [ [[TMP5]], [[ENTRY:%.*]] ], [ [[OMP_ARRAYCPY_SRC_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
 772 // CHECK1-NEXT:    [[OMP_ARRAYCPY_DESTELEMENTPAST:%.*]] = phi ptr [ [[TMP7]], [[ENTRY]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
 773 // CHECK1-NEXT:    [[TMP12:%.*]] = load double, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], align 8
 774 // CHECK1-NEXT:    [[TMP13:%.*]] = load double, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], align 8
 775 // CHECK1-NEXT:    [[ADD:%.*]] = fadd double [[TMP12]], [[TMP13]]
 776 // CHECK1-NEXT:    store double [[ADD]], ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], align 8
 777 // CHECK1-NEXT:    [[OMP_ARRAYCPY_DEST_ELEMENT]] = getelementptr double, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], i32 1
 778 // CHECK1-NEXT:    [[OMP_ARRAYCPY_SRC_ELEMENT]] = getelementptr double, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], i32 1
 779 // CHECK1-NEXT:    [[OMP_ARRAYCPY_DONE:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT]], [[TMP11]]
 780 // CHECK1-NEXT:    br i1 [[OMP_ARRAYCPY_DONE]], label [[OMP_ARRAYCPY_DONE2]], label [[OMP_ARRAYCPY_BODY]]
 781 // CHECK1:       omp.arraycpy.done2:
 782 // CHECK1-NEXT:    ret void
 783 //
 784 //
 785 // CHECK1-LABEL: define {{[^@]+}}@_ZN2S2C2Ev
 786 // CHECK1-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR0]] comdat {
 787 // CHECK1-NEXT:  entry:
 788 // CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 789 // CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 790 // CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 791 // CHECK1-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[CLASS_S2:%.*]], ptr [[THIS1]], i32 0, i32 0
 792 // CHECK1-NEXT:    store i32 0, ptr [[A]], align 4
 793 // CHECK1-NEXT:    ret void
 794 //
 795 //
 796 // CHECK1-LABEL: define {{[^@]+}}@main
 797 // CHECK1-SAME: () #[[ATTR9:[0-9]+]] {
 798 // CHECK1-NEXT:  entry:
 799 // CHECK1-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4
 800 // CHECK1-NEXT:    [[A:%.*]] = alloca i32, align 4
 801 // CHECK1-NEXT:    store i32 0, ptr [[RETVAL]], align 4
 802 // CHECK1-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_Z3barv()
 803 // CHECK1-NEXT:    store i32 [[CALL]], ptr [[A]], align 4
 804 // CHECK1-NEXT:    ret i32 0
 805 //
 806 //
 807 // CHECK2-LABEL: define {{[^@]+}}@_Z3sumPiiS_
 808 // CHECK2-SAME: (ptr noundef [[INPUT:%.*]], i32 noundef [[SIZE:%.*]], ptr noundef [[OUTPUT:%.*]]) #[[ATTR0:[0-9]+]] {
 809 // CHECK2-NEXT:  entry:
 810 // CHECK2-NEXT:    [[INPUT_ADDR:%.*]] = alloca ptr, align 4
 811 // CHECK2-NEXT:    [[SIZE_ADDR:%.*]] = alloca i32, align 4
 812 // CHECK2-NEXT:    [[OUTPUT_ADDR:%.*]] = alloca ptr, align 4
 813 // CHECK2-NEXT:    [[SIZE_CASTED:%.*]] = alloca i32, align 4
 814 // CHECK2-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 4
 815 // CHECK2-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 4
 816 // CHECK2-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 4
 817 // CHECK2-NEXT:    [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 4
 818 // CHECK2-NEXT:    [[TMP:%.*]] = alloca i32, align 4
 819 // CHECK2-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4
 820 // CHECK2-NEXT:    [[DOTCAPTURE_EXPR_2:%.*]] = alloca i32, align 4
 821 // CHECK2-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8
 822 // CHECK2-NEXT:    [[SIZE_CASTED4:%.*]] = alloca i32, align 4
 823 // CHECK2-NEXT:    [[DOTOFFLOAD_BASEPTRS7:%.*]] = alloca [3 x ptr], align 4
 824 // CHECK2-NEXT:    [[DOTOFFLOAD_PTRS8:%.*]] = alloca [3 x ptr], align 4
 825 // CHECK2-NEXT:    [[DOTOFFLOAD_MAPPERS9:%.*]] = alloca [3 x ptr], align 4
 826 // CHECK2-NEXT:    [[DOTOFFLOAD_SIZES10:%.*]] = alloca [3 x i64], align 4
 827 // CHECK2-NEXT:    [[_TMP11:%.*]] = alloca i32, align 4
 828 // CHECK2-NEXT:    [[DOTCAPTURE_EXPR_12:%.*]] = alloca i32, align 4
 829 // CHECK2-NEXT:    [[DOTCAPTURE_EXPR_13:%.*]] = alloca i32, align 4
 830 // CHECK2-NEXT:    [[KERNEL_ARGS18:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8
 831 // CHECK2-NEXT:    [[A:%.*]] = alloca [10 x i32], align 4
 832 // CHECK2-NEXT:    [[SIZE_CASTED21:%.*]] = alloca i32, align 4
 833 // CHECK2-NEXT:    [[DOTOFFLOAD_BASEPTRS23:%.*]] = alloca [2 x ptr], align 4
 834 // CHECK2-NEXT:    [[DOTOFFLOAD_PTRS24:%.*]] = alloca [2 x ptr], align 4
 835 // CHECK2-NEXT:    [[DOTOFFLOAD_MAPPERS25:%.*]] = alloca [2 x ptr], align 4
 836 // CHECK2-NEXT:    [[KERNEL_ARGS26:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8
 837 // CHECK2-NEXT:    [[SIZE_CASTED29:%.*]] = alloca i32, align 4
 838 // CHECK2-NEXT:    [[DOTOFFLOAD_BASEPTRS31:%.*]] = alloca [2 x ptr], align 4
 839 // CHECK2-NEXT:    [[DOTOFFLOAD_PTRS32:%.*]] = alloca [2 x ptr], align 4
 840 // CHECK2-NEXT:    [[DOTOFFLOAD_MAPPERS33:%.*]] = alloca [2 x ptr], align 4
 841 // CHECK2-NEXT:    [[KERNEL_ARGS34:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8
 842 // CHECK2-NEXT:    store ptr [[INPUT]], ptr [[INPUT_ADDR]], align 4
 843 // CHECK2-NEXT:    store i32 [[SIZE]], ptr [[SIZE_ADDR]], align 4
 844 // CHECK2-NEXT:    store ptr [[OUTPUT]], ptr [[OUTPUT_ADDR]], align 4
 845 // CHECK2-NEXT:    [[TMP0:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
 846 // CHECK2-NEXT:    store i32 [[TMP0]], ptr [[SIZE_CASTED]], align 4
 847 // CHECK2-NEXT:    [[TMP1:%.*]] = load i32, ptr [[SIZE_CASTED]], align 4
 848 // CHECK2-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[OUTPUT_ADDR]], align 4
 849 // CHECK2-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[INPUT_ADDR]], align 4
 850 // CHECK2-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[OUTPUT_ADDR]], align 4
 851 // CHECK2-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[OUTPUT_ADDR]], align 4
 852 // CHECK2-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP5]], i32 0
 853 // CHECK2-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[INPUT_ADDR]], align 4
 854 // CHECK2-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[INPUT_ADDR]], align 4
 855 // CHECK2-NEXT:    [[ARRAYIDX1:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP7]], i32 0
 856 // CHECK2-NEXT:    [[TMP8:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
 857 // CHECK2-NEXT:    [[TMP9:%.*]] = mul nuw i32 [[TMP8]], 4
 858 // CHECK2-NEXT:    [[TMP10:%.*]] = sext i32 [[TMP9]] to i64
 859 // CHECK2-NEXT:    call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[DOTOFFLOAD_SIZES]], ptr align 4 @.offload_sizes, i32 24, i1 false)
 860 // CHECK2-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
 861 // CHECK2-NEXT:    store i32 [[TMP1]], ptr [[TMP11]], align 4
 862 // CHECK2-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
 863 // CHECK2-NEXT:    store i32 [[TMP1]], ptr [[TMP12]], align 4
 864 // CHECK2-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0
 865 // CHECK2-NEXT:    store ptr null, ptr [[TMP13]], align 4
 866 // CHECK2-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
 867 // CHECK2-NEXT:    store ptr [[TMP4]], ptr [[TMP14]], align 4
 868 // CHECK2-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1
 869 // CHECK2-NEXT:    store ptr [[ARRAYIDX]], ptr [[TMP15]], align 4
 870 // CHECK2-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 1
 871 // CHECK2-NEXT:    store ptr null, ptr [[TMP16]], align 4
 872 // CHECK2-NEXT:    [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2
 873 // CHECK2-NEXT:    store ptr [[TMP6]], ptr [[TMP17]], align 4
 874 // CHECK2-NEXT:    [[TMP18:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2
 875 // CHECK2-NEXT:    store ptr [[ARRAYIDX1]], ptr [[TMP18]], align 4
 876 // CHECK2-NEXT:    [[TMP19:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 2
 877 // CHECK2-NEXT:    store i64 [[TMP10]], ptr [[TMP19]], align 4
 878 // CHECK2-NEXT:    [[TMP20:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 2
 879 // CHECK2-NEXT:    store ptr null, ptr [[TMP20]], align 4
 880 // CHECK2-NEXT:    [[TMP21:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
 881 // CHECK2-NEXT:    [[TMP22:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
 882 // CHECK2-NEXT:    [[TMP23:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0
 883 // CHECK2-NEXT:    [[TMP24:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
 884 // CHECK2-NEXT:    store i32 [[TMP24]], ptr [[DOTCAPTURE_EXPR_]], align 4
 885 // CHECK2-NEXT:    [[TMP25:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
 886 // CHECK2-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP25]], 0
 887 // CHECK2-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1
 888 // CHECK2-NEXT:    [[SUB3:%.*]] = sub nsw i32 [[DIV]], 1
 889 // CHECK2-NEXT:    store i32 [[SUB3]], ptr [[DOTCAPTURE_EXPR_2]], align 4
 890 // CHECK2-NEXT:    [[TMP26:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 4
 891 // CHECK2-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP26]], 1
 892 // CHECK2-NEXT:    [[TMP27:%.*]] = zext i32 [[ADD]] to i64
 893 // CHECK2-NEXT:    [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0
 894 // CHECK2-NEXT:    store i32 3, ptr [[TMP28]], align 4
 895 // CHECK2-NEXT:    [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1
 896 // CHECK2-NEXT:    store i32 3, ptr [[TMP29]], align 4
 897 // CHECK2-NEXT:    [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2
 898 // CHECK2-NEXT:    store ptr [[TMP21]], ptr [[TMP30]], align 4
 899 // CHECK2-NEXT:    [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3
 900 // CHECK2-NEXT:    store ptr [[TMP22]], ptr [[TMP31]], align 4
 901 // CHECK2-NEXT:    [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4
 902 // CHECK2-NEXT:    store ptr [[TMP23]], ptr [[TMP32]], align 4
 903 // CHECK2-NEXT:    [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5
 904 // CHECK2-NEXT:    store ptr @.offload_maptypes, ptr [[TMP33]], align 4
 905 // CHECK2-NEXT:    [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6
 906 // CHECK2-NEXT:    store ptr null, ptr [[TMP34]], align 4
 907 // CHECK2-NEXT:    [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7
 908 // CHECK2-NEXT:    store ptr null, ptr [[TMP35]], align 4
 909 // CHECK2-NEXT:    [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8
 910 // CHECK2-NEXT:    store i64 [[TMP27]], ptr [[TMP36]], align 8
 911 // CHECK2-NEXT:    [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9
 912 // CHECK2-NEXT:    store i64 0, ptr [[TMP37]], align 8
 913 // CHECK2-NEXT:    [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10
 914 // CHECK2-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP38]], align 4
 915 // CHECK2-NEXT:    [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11
 916 // CHECK2-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP39]], align 4
 917 // CHECK2-NEXT:    [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12
 918 // CHECK2-NEXT:    store i32 0, ptr [[TMP40]], align 4
 919 // CHECK2-NEXT:    [[TMP41:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB4:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l69.region_id, ptr [[KERNEL_ARGS]])
 920 // CHECK2-NEXT:    [[TMP42:%.*]] = icmp ne i32 [[TMP41]], 0
 921 // CHECK2-NEXT:    br i1 [[TMP42]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
 922 // CHECK2:       omp_offload.failed:
 923 // CHECK2-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l69(i32 [[TMP1]], ptr [[TMP2]], ptr [[TMP3]]) #[[ATTR2:[0-9]+]]
 924 // CHECK2-NEXT:    br label [[OMP_OFFLOAD_CONT]]
 925 // CHECK2:       omp_offload.cont:
 926 // CHECK2-NEXT:    [[TMP43:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
 927 // CHECK2-NEXT:    store i32 [[TMP43]], ptr [[SIZE_CASTED4]], align 4
 928 // CHECK2-NEXT:    [[TMP44:%.*]] = load i32, ptr [[SIZE_CASTED4]], align 4
 929 // CHECK2-NEXT:    [[TMP45:%.*]] = load ptr, ptr [[OUTPUT_ADDR]], align 4
 930 // CHECK2-NEXT:    [[TMP46:%.*]] = load ptr, ptr [[INPUT_ADDR]], align 4
 931 // CHECK2-NEXT:    [[TMP47:%.*]] = load ptr, ptr [[OUTPUT_ADDR]], align 4
 932 // CHECK2-NEXT:    [[TMP48:%.*]] = load ptr, ptr [[OUTPUT_ADDR]], align 4
 933 // CHECK2-NEXT:    [[ARRAYIDX5:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP48]], i32 0
 934 // CHECK2-NEXT:    [[TMP49:%.*]] = load ptr, ptr [[INPUT_ADDR]], align 4
 935 // CHECK2-NEXT:    [[TMP50:%.*]] = load ptr, ptr [[INPUT_ADDR]], align 4
 936 // CHECK2-NEXT:    [[ARRAYIDX6:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP50]], i32 0
 937 // CHECK2-NEXT:    [[TMP51:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
 938 // CHECK2-NEXT:    [[TMP52:%.*]] = mul nuw i32 [[TMP51]], 4
 939 // CHECK2-NEXT:    [[TMP53:%.*]] = sext i32 [[TMP52]] to i64
 940 // CHECK2-NEXT:    call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[DOTOFFLOAD_SIZES10]], ptr align 4 @.offload_sizes.1, i32 24, i1 false)
 941 // CHECK2-NEXT:    [[TMP54:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 0
 942 // CHECK2-NEXT:    store i32 [[TMP44]], ptr [[TMP54]], align 4
 943 // CHECK2-NEXT:    [[TMP55:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 0
 944 // CHECK2-NEXT:    store i32 [[TMP44]], ptr [[TMP55]], align 4
 945 // CHECK2-NEXT:    [[TMP56:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS9]], i32 0, i32 0
 946 // CHECK2-NEXT:    store ptr null, ptr [[TMP56]], align 4
 947 // CHECK2-NEXT:    [[TMP57:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 1
 948 // CHECK2-NEXT:    store ptr [[TMP47]], ptr [[TMP57]], align 4
 949 // CHECK2-NEXT:    [[TMP58:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 1
 950 // CHECK2-NEXT:    store ptr [[ARRAYIDX5]], ptr [[TMP58]], align 4
 951 // CHECK2-NEXT:    [[TMP59:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS9]], i32 0, i32 1
 952 // CHECK2-NEXT:    store ptr null, ptr [[TMP59]], align 4
 953 // CHECK2-NEXT:    [[TMP60:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 2
 954 // CHECK2-NEXT:    store ptr [[TMP49]], ptr [[TMP60]], align 4
 955 // CHECK2-NEXT:    [[TMP61:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 2
 956 // CHECK2-NEXT:    store ptr [[ARRAYIDX6]], ptr [[TMP61]], align 4
 957 // CHECK2-NEXT:    [[TMP62:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES10]], i32 0, i32 2
 958 // CHECK2-NEXT:    store i64 [[TMP53]], ptr [[TMP62]], align 4
 959 // CHECK2-NEXT:    [[TMP63:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS9]], i32 0, i32 2
 960 // CHECK2-NEXT:    store ptr null, ptr [[TMP63]], align 4
 961 // CHECK2-NEXT:    [[TMP64:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 0
 962 // CHECK2-NEXT:    [[TMP65:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 0
 963 // CHECK2-NEXT:    [[TMP66:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES10]], i32 0, i32 0
 964 // CHECK2-NEXT:    [[TMP67:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
 965 // CHECK2-NEXT:    store i32 [[TMP67]], ptr [[DOTCAPTURE_EXPR_12]], align 4
 966 // CHECK2-NEXT:    [[TMP68:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_12]], align 4
 967 // CHECK2-NEXT:    [[SUB14:%.*]] = sub nsw i32 [[TMP68]], 0
 968 // CHECK2-NEXT:    [[DIV15:%.*]] = sdiv i32 [[SUB14]], 1
 969 // CHECK2-NEXT:    [[SUB16:%.*]] = sub nsw i32 [[DIV15]], 1
 970 // CHECK2-NEXT:    store i32 [[SUB16]], ptr [[DOTCAPTURE_EXPR_13]], align 4
 971 // CHECK2-NEXT:    [[TMP69:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_13]], align 4
 972 // CHECK2-NEXT:    [[ADD17:%.*]] = add nsw i32 [[TMP69]], 1
 973 // CHECK2-NEXT:    [[TMP70:%.*]] = zext i32 [[ADD17]] to i64
 974 // CHECK2-NEXT:    [[TMP71:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 0
 975 // CHECK2-NEXT:    store i32 3, ptr [[TMP71]], align 4
 976 // CHECK2-NEXT:    [[TMP72:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 1
 977 // CHECK2-NEXT:    store i32 3, ptr [[TMP72]], align 4
 978 // CHECK2-NEXT:    [[TMP73:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 2
 979 // CHECK2-NEXT:    store ptr [[TMP64]], ptr [[TMP73]], align 4
 980 // CHECK2-NEXT:    [[TMP74:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 3
 981 // CHECK2-NEXT:    store ptr [[TMP65]], ptr [[TMP74]], align 4
 982 // CHECK2-NEXT:    [[TMP75:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 4
 983 // CHECK2-NEXT:    store ptr [[TMP66]], ptr [[TMP75]], align 4
 984 // CHECK2-NEXT:    [[TMP76:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 5
 985 // CHECK2-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP76]], align 4
 986 // CHECK2-NEXT:    [[TMP77:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 6
 987 // CHECK2-NEXT:    store ptr null, ptr [[TMP77]], align 4
 988 // CHECK2-NEXT:    [[TMP78:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 7
 989 // CHECK2-NEXT:    store ptr null, ptr [[TMP78]], align 4
 990 // CHECK2-NEXT:    [[TMP79:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 8
 991 // CHECK2-NEXT:    store i64 [[TMP70]], ptr [[TMP79]], align 8
 992 // CHECK2-NEXT:    [[TMP80:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 9
 993 // CHECK2-NEXT:    store i64 0, ptr [[TMP80]], align 8
 994 // CHECK2-NEXT:    [[TMP81:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 10
 995 // CHECK2-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP81]], align 4
 996 // CHECK2-NEXT:    [[TMP82:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 11
 997 // CHECK2-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP82]], align 4
 998 // CHECK2-NEXT:    [[TMP83:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS18]], i32 0, i32 12
 999 // CHECK2-NEXT:    store i32 0, ptr [[TMP83]], align 4
1000 // CHECK2-NEXT:    [[TMP84:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB4]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l73.region_id, ptr [[KERNEL_ARGS18]])
1001 // CHECK2-NEXT:    [[TMP85:%.*]] = icmp ne i32 [[TMP84]], 0
1002 // CHECK2-NEXT:    br i1 [[TMP85]], label [[OMP_OFFLOAD_FAILED19:%.*]], label [[OMP_OFFLOAD_CONT20:%.*]]
1003 // CHECK2:       omp_offload.failed19:
1004 // CHECK2-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l73(i32 [[TMP44]], ptr [[TMP45]], ptr [[TMP46]]) #[[ATTR2]]
1005 // CHECK2-NEXT:    br label [[OMP_OFFLOAD_CONT20]]
1006 // CHECK2:       omp_offload.cont20:
1007 // CHECK2-NEXT:    [[TMP86:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
1008 // CHECK2-NEXT:    store i32 [[TMP86]], ptr [[SIZE_CASTED21]], align 4
1009 // CHECK2-NEXT:    [[TMP87:%.*]] = load i32, ptr [[SIZE_CASTED21]], align 4
1010 // CHECK2-NEXT:    [[ARRAYIDX22:%.*]] = getelementptr inbounds nuw [10 x i32], ptr [[A]], i32 0, i32 0
1011 // CHECK2-NEXT:    [[TMP88:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS23]], i32 0, i32 0
1012 // CHECK2-NEXT:    store i32 [[TMP87]], ptr [[TMP88]], align 4
1013 // CHECK2-NEXT:    [[TMP89:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS24]], i32 0, i32 0
1014 // CHECK2-NEXT:    store i32 [[TMP87]], ptr [[TMP89]], align 4
1015 // CHECK2-NEXT:    [[TMP90:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS25]], i32 0, i32 0
1016 // CHECK2-NEXT:    store ptr null, ptr [[TMP90]], align 4
1017 // CHECK2-NEXT:    [[TMP91:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS23]], i32 0, i32 1
1018 // CHECK2-NEXT:    store ptr [[A]], ptr [[TMP91]], align 4
1019 // CHECK2-NEXT:    [[TMP92:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS24]], i32 0, i32 1
1020 // CHECK2-NEXT:    store ptr [[ARRAYIDX22]], ptr [[TMP92]], align 4
1021 // CHECK2-NEXT:    [[TMP93:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS25]], i32 0, i32 1
1022 // CHECK2-NEXT:    store ptr null, ptr [[TMP93]], align 4
1023 // CHECK2-NEXT:    [[TMP94:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS23]], i32 0, i32 0
1024 // CHECK2-NEXT:    [[TMP95:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS24]], i32 0, i32 0
1025 // CHECK2-NEXT:    [[TMP96:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS26]], i32 0, i32 0
1026 // CHECK2-NEXT:    store i32 3, ptr [[TMP96]], align 4
1027 // CHECK2-NEXT:    [[TMP97:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS26]], i32 0, i32 1
1028 // CHECK2-NEXT:    store i32 2, ptr [[TMP97]], align 4
1029 // CHECK2-NEXT:    [[TMP98:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS26]], i32 0, i32 2
1030 // CHECK2-NEXT:    store ptr [[TMP94]], ptr [[TMP98]], align 4
1031 // CHECK2-NEXT:    [[TMP99:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS26]], i32 0, i32 3
1032 // CHECK2-NEXT:    store ptr [[TMP95]], ptr [[TMP99]], align 4
1033 // CHECK2-NEXT:    [[TMP100:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS26]], i32 0, i32 4
1034 // CHECK2-NEXT:    store ptr @.offload_sizes.3, ptr [[TMP100]], align 4
1035 // CHECK2-NEXT:    [[TMP101:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS26]], i32 0, i32 5
1036 // CHECK2-NEXT:    store ptr @.offload_maptypes.4, ptr [[TMP101]], align 4
1037 // CHECK2-NEXT:    [[TMP102:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS26]], i32 0, i32 6
1038 // CHECK2-NEXT:    store ptr null, ptr [[TMP102]], align 4
1039 // CHECK2-NEXT:    [[TMP103:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS26]], i32 0, i32 7
1040 // CHECK2-NEXT:    store ptr null, ptr [[TMP103]], align 4
1041 // CHECK2-NEXT:    [[TMP104:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS26]], i32 0, i32 8
1042 // CHECK2-NEXT:    store i64 0, ptr [[TMP104]], align 8
1043 // CHECK2-NEXT:    [[TMP105:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS26]], i32 0, i32 9
1044 // CHECK2-NEXT:    store i64 0, ptr [[TMP105]], align 8
1045 // CHECK2-NEXT:    [[TMP106:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS26]], i32 0, i32 10
1046 // CHECK2-NEXT:    store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP106]], align 4
1047 // CHECK2-NEXT:    [[TMP107:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS26]], i32 0, i32 11
1048 // CHECK2-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP107]], align 4
1049 // CHECK2-NEXT:    [[TMP108:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS26]], i32 0, i32 12
1050 // CHECK2-NEXT:    store i32 0, ptr [[TMP108]], align 4
1051 // CHECK2-NEXT:    [[TMP109:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB4]], i64 -1, i32 1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l78.region_id, ptr [[KERNEL_ARGS26]])
1052 // CHECK2-NEXT:    [[TMP110:%.*]] = icmp ne i32 [[TMP109]], 0
1053 // CHECK2-NEXT:    br i1 [[TMP110]], label [[OMP_OFFLOAD_FAILED27:%.*]], label [[OMP_OFFLOAD_CONT28:%.*]]
1054 // CHECK2:       omp_offload.failed27:
1055 // CHECK2-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l78(i32 [[TMP87]], ptr [[A]]) #[[ATTR2]]
1056 // CHECK2-NEXT:    br label [[OMP_OFFLOAD_CONT28]]
1057 // CHECK2:       omp_offload.cont28:
1058 // CHECK2-NEXT:    [[TMP111:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
1059 // CHECK2-NEXT:    store i32 [[TMP111]], ptr [[SIZE_CASTED29]], align 4
1060 // CHECK2-NEXT:    [[TMP112:%.*]] = load i32, ptr [[SIZE_CASTED29]], align 4
1061 // CHECK2-NEXT:    [[ARRAYIDX30:%.*]] = getelementptr inbounds [10 x i32], ptr [[A]], i32 0, i32 3
1062 // CHECK2-NEXT:    [[TMP113:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS31]], i32 0, i32 0
1063 // CHECK2-NEXT:    store i32 [[TMP112]], ptr [[TMP113]], align 4
1064 // CHECK2-NEXT:    [[TMP114:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS32]], i32 0, i32 0
1065 // CHECK2-NEXT:    store i32 [[TMP112]], ptr [[TMP114]], align 4
1066 // CHECK2-NEXT:    [[TMP115:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS33]], i32 0, i32 0
1067 // CHECK2-NEXT:    store ptr null, ptr [[TMP115]], align 4
1068 // CHECK2-NEXT:    [[TMP116:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS31]], i32 0, i32 1
1069 // CHECK2-NEXT:    store ptr [[A]], ptr [[TMP116]], align 4
1070 // CHECK2-NEXT:    [[TMP117:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS32]], i32 0, i32 1
1071 // CHECK2-NEXT:    store ptr [[ARRAYIDX30]], ptr [[TMP117]], align 4
1072 // CHECK2-NEXT:    [[TMP118:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS33]], i32 0, i32 1
1073 // CHECK2-NEXT:    store ptr null, ptr [[TMP118]], align 4
1074 // CHECK2-NEXT:    [[TMP119:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS31]], i32 0, i32 0
1075 // CHECK2-NEXT:    [[TMP120:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS32]], i32 0, i32 0
1076 // CHECK2-NEXT:    [[TMP121:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS34]], i32 0, i32 0
1077 // CHECK2-NEXT:    store i32 3, ptr [[TMP121]], align 4
1078 // CHECK2-NEXT:    [[TMP122:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS34]], i32 0, i32 1
1079 // CHECK2-NEXT:    store i32 2, ptr [[TMP122]], align 4
1080 // CHECK2-NEXT:    [[TMP123:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS34]], i32 0, i32 2
1081 // CHECK2-NEXT:    store ptr [[TMP119]], ptr [[TMP123]], align 4
1082 // CHECK2-NEXT:    [[TMP124:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS34]], i32 0, i32 3
1083 // CHECK2-NEXT:    store ptr [[TMP120]], ptr [[TMP124]], align 4
1084 // CHECK2-NEXT:    [[TMP125:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS34]], i32 0, i32 4
1085 // CHECK2-NEXT:    store ptr @.offload_sizes.5, ptr [[TMP125]], align 4
1086 // CHECK2-NEXT:    [[TMP126:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS34]], i32 0, i32 5
1087 // CHECK2-NEXT:    store ptr @.offload_maptypes.6, ptr [[TMP126]], align 4
1088 // CHECK2-NEXT:    [[TMP127:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS34]], i32 0, i32 6
1089 // CHECK2-NEXT:    store ptr null, ptr [[TMP127]], align 4
1090 // CHECK2-NEXT:    [[TMP128:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS34]], i32 0, i32 7
1091 // CHECK2-NEXT:    store ptr null, ptr [[TMP128]], align 4
1092 // CHECK2-NEXT:    [[TMP129:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS34]], i32 0, i32 8
1093 // CHECK2-NEXT:    store i64 0, ptr [[TMP129]], align 8
1094 // CHECK2-NEXT:    [[TMP130:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS34]], i32 0, i32 9
1095 // CHECK2-NEXT:    store i64 0, ptr [[TMP130]], align 8
1096 // CHECK2-NEXT:    [[TMP131:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS34]], i32 0, i32 10
1097 // CHECK2-NEXT:    store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP131]], align 4
1098 // CHECK2-NEXT:    [[TMP132:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS34]], i32 0, i32 11
1099 // CHECK2-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP132]], align 4
1100 // CHECK2-NEXT:    [[TMP133:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS34]], i32 0, i32 12
1101 // CHECK2-NEXT:    store i32 0, ptr [[TMP133]], align 4
1102 // CHECK2-NEXT:    [[TMP134:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB4]], i64 -1, i32 1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l81.region_id, ptr [[KERNEL_ARGS34]])
1103 // CHECK2-NEXT:    [[TMP135:%.*]] = icmp ne i32 [[TMP134]], 0
1104 // CHECK2-NEXT:    br i1 [[TMP135]], label [[OMP_OFFLOAD_FAILED35:%.*]], label [[OMP_OFFLOAD_CONT36:%.*]]
1105 // CHECK2:       omp_offload.failed35:
1106 // CHECK2-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l81(i32 [[TMP112]], ptr [[A]]) #[[ATTR2]]
1107 // CHECK2-NEXT:    br label [[OMP_OFFLOAD_CONT36]]
1108 // CHECK2:       omp_offload.cont36:
1109 // CHECK2-NEXT:    ret void
1110 //
1111 //
1112 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l69
1113 // CHECK2-SAME: (i32 noundef [[SIZE:%.*]], ptr noundef [[OUTPUT:%.*]], ptr noundef [[INPUT:%.*]]) #[[ATTR1:[0-9]+]] {
1114 // CHECK2-NEXT:  entry:
1115 // CHECK2-NEXT:    [[SIZE_ADDR:%.*]] = alloca i32, align 4
1116 // CHECK2-NEXT:    [[OUTPUT_ADDR:%.*]] = alloca ptr, align 4
1117 // CHECK2-NEXT:    [[INPUT_ADDR:%.*]] = alloca ptr, align 4
1118 // CHECK2-NEXT:    [[SIZE_CASTED:%.*]] = alloca i32, align 4
1119 // CHECK2-NEXT:    store i32 [[SIZE]], ptr [[SIZE_ADDR]], align 4
1120 // CHECK2-NEXT:    store ptr [[OUTPUT]], ptr [[OUTPUT_ADDR]], align 4
1121 // CHECK2-NEXT:    store ptr [[INPUT]], ptr [[INPUT_ADDR]], align 4
1122 // CHECK2-NEXT:    [[TMP0:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
1123 // CHECK2-NEXT:    store i32 [[TMP0]], ptr [[SIZE_CASTED]], align 4
1124 // CHECK2-NEXT:    [[TMP1:%.*]] = load i32, ptr [[SIZE_CASTED]], align 4
1125 // CHECK2-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[OUTPUT_ADDR]], align 4
1126 // CHECK2-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[INPUT_ADDR]], align 4
1127 // CHECK2-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB4]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l69.omp_outlined, i32 [[TMP1]], ptr [[TMP2]], ptr [[TMP3]])
1128 // CHECK2-NEXT:    ret void
1129 //
1130 //
1131 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l69.omp_outlined
1132 // CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[SIZE:%.*]], ptr noundef [[OUTPUT:%.*]], ptr noundef [[INPUT:%.*]]) #[[ATTR1]] {
1133 // CHECK2-NEXT:  entry:
1134 // CHECK2-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
1135 // CHECK2-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
1136 // CHECK2-NEXT:    [[SIZE_ADDR:%.*]] = alloca i32, align 4
1137 // CHECK2-NEXT:    [[OUTPUT_ADDR:%.*]] = alloca ptr, align 4
1138 // CHECK2-NEXT:    [[INPUT_ADDR:%.*]] = alloca ptr, align 4
1139 // CHECK2-NEXT:    [[OUTPUT1:%.*]] = alloca i32, align 4
1140 // CHECK2-NEXT:    [[TMP:%.*]] = alloca ptr, align 4
1141 // CHECK2-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
1142 // CHECK2-NEXT:    [[_TMP2:%.*]] = alloca i32, align 4
1143 // CHECK2-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4
1144 // CHECK2-NEXT:    [[DOTCAPTURE_EXPR_3:%.*]] = alloca i32, align 4
1145 // CHECK2-NEXT:    [[I:%.*]] = alloca i32, align 4
1146 // CHECK2-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4
1147 // CHECK2-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4
1148 // CHECK2-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
1149 // CHECK2-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
1150 // CHECK2-NEXT:    [[I5:%.*]] = alloca i32, align 4
1151 // CHECK2-NEXT:    [[SIZE_CASTED:%.*]] = alloca i32, align 4
1152 // CHECK2-NEXT:    [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 4
1153 // CHECK2-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
1154 // CHECK2-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
1155 // CHECK2-NEXT:    store i32 [[SIZE]], ptr [[SIZE_ADDR]], align 4
1156 // CHECK2-NEXT:    store ptr [[OUTPUT]], ptr [[OUTPUT_ADDR]], align 4
1157 // CHECK2-NEXT:    store ptr [[INPUT]], ptr [[INPUT_ADDR]], align 4
1158 // CHECK2-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[OUTPUT_ADDR]], align 4
1159 // CHECK2-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i32 0
1160 // CHECK2-NEXT:    store i32 0, ptr [[OUTPUT1]], align 4
1161 // CHECK2-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[OUTPUT_ADDR]], align 4
1162 // CHECK2-NEXT:    [[TMP2:%.*]] = ptrtoint ptr [[TMP1]] to i64
1163 // CHECK2-NEXT:    [[TMP3:%.*]] = ptrtoint ptr [[ARRAYIDX]] to i64
1164 // CHECK2-NEXT:    [[TMP4:%.*]] = sub i64 [[TMP2]], [[TMP3]]
1165 // CHECK2-NEXT:    [[TMP5:%.*]] = sdiv exact i64 [[TMP4]], ptrtoint (ptr getelementptr (i32, ptr null, i32 1) to i64)
1166 // CHECK2-NEXT:    [[TMP6:%.*]] = getelementptr i32, ptr [[OUTPUT1]], i64 [[TMP5]]
1167 // CHECK2-NEXT:    store ptr [[TMP6]], ptr [[TMP]], align 4
1168 // CHECK2-NEXT:    [[TMP7:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
1169 // CHECK2-NEXT:    store i32 [[TMP7]], ptr [[DOTCAPTURE_EXPR_]], align 4
1170 // CHECK2-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
1171 // CHECK2-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP8]], 0
1172 // CHECK2-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1
1173 // CHECK2-NEXT:    [[SUB4:%.*]] = sub nsw i32 [[DIV]], 1
1174 // CHECK2-NEXT:    store i32 [[SUB4]], ptr [[DOTCAPTURE_EXPR_3]], align 4
1175 // CHECK2-NEXT:    store i32 0, ptr [[I]], align 4
1176 // CHECK2-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
1177 // CHECK2-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP9]]
1178 // CHECK2-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]
1179 // CHECK2:       omp.precond.then:
1180 // CHECK2-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 4
1181 // CHECK2-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_3]], align 4
1182 // CHECK2-NEXT:    store i32 [[TMP10]], ptr [[DOTOMP_COMB_UB]], align 4
1183 // CHECK2-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4
1184 // CHECK2-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4
1185 // CHECK2-NEXT:    [[TMP11:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
1186 // CHECK2-NEXT:    [[TMP12:%.*]] = load i32, ptr [[TMP11]], align 4
1187 // CHECK2-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP12]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
1188 // CHECK2-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4
1189 // CHECK2-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_3]], align 4
1190 // CHECK2-NEXT:    [[CMP6:%.*]] = icmp sgt i32 [[TMP13]], [[TMP14]]
1191 // CHECK2-NEXT:    br i1 [[CMP6]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
1192 // CHECK2:       cond.true:
1193 // CHECK2-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_3]], align 4
1194 // CHECK2-NEXT:    br label [[COND_END:%.*]]
1195 // CHECK2:       cond.false:
1196 // CHECK2-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4
1197 // CHECK2-NEXT:    br label [[COND_END]]
1198 // CHECK2:       cond.end:
1199 // CHECK2-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP15]], [[COND_TRUE]] ], [ [[TMP16]], [[COND_FALSE]] ]
1200 // CHECK2-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4
1201 // CHECK2-NEXT:    [[TMP17:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4
1202 // CHECK2-NEXT:    store i32 [[TMP17]], ptr [[DOTOMP_IV]], align 4
1203 // CHECK2-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
1204 // CHECK2:       omp.inner.for.cond:
1205 // CHECK2-NEXT:    [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
1206 // CHECK2-NEXT:    [[TMP19:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4
1207 // CHECK2-NEXT:    [[CMP7:%.*]] = icmp sle i32 [[TMP18]], [[TMP19]]
1208 // CHECK2-NEXT:    br i1 [[CMP7]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
1209 // CHECK2:       omp.inner.for.body:
1210 // CHECK2-NEXT:    [[TMP20:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4
1211 // CHECK2-NEXT:    [[TMP21:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4
1212 // CHECK2-NEXT:    [[TMP22:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
1213 // CHECK2-NEXT:    store i32 [[TMP22]], ptr [[SIZE_CASTED]], align 4
1214 // CHECK2-NEXT:    [[TMP23:%.*]] = load i32, ptr [[SIZE_CASTED]], align 4
1215 // CHECK2-NEXT:    [[TMP24:%.*]] = load ptr, ptr [[TMP]], align 4
1216 // CHECK2-NEXT:    [[TMP25:%.*]] = load ptr, ptr [[INPUT_ADDR]], align 4
1217 // CHECK2-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB4]], i32 5, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l69.omp_outlined.omp_outlined, i32 [[TMP20]], i32 [[TMP21]], i32 [[TMP23]], ptr [[TMP24]], ptr [[TMP25]])
1218 // CHECK2-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
1219 // CHECK2:       omp.inner.for.inc:
1220 // CHECK2-NEXT:    [[TMP26:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
1221 // CHECK2-NEXT:    [[TMP27:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4
1222 // CHECK2-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP26]], [[TMP27]]
1223 // CHECK2-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4
1224 // CHECK2-NEXT:    br label [[OMP_INNER_FOR_COND]]
1225 // CHECK2:       omp.inner.for.end:
1226 // CHECK2-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
1227 // CHECK2:       omp.loop.exit:
1228 // CHECK2-NEXT:    [[TMP28:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
1229 // CHECK2-NEXT:    [[TMP29:%.*]] = load i32, ptr [[TMP28]], align 4
1230 // CHECK2-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP29]])
1231 // CHECK2-NEXT:    br label [[OMP_PRECOND_END]]
1232 // CHECK2:       omp.precond.end:
1233 // CHECK2-NEXT:    [[TMP30:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 0
1234 // CHECK2-NEXT:    store ptr [[OUTPUT1]], ptr [[TMP30]], align 4
1235 // CHECK2-NEXT:    [[TMP31:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
1236 // CHECK2-NEXT:    [[TMP32:%.*]] = load i32, ptr [[TMP31]], align 4
1237 // CHECK2-NEXT:    [[TMP33:%.*]] = call i32 @__kmpc_reduce_nowait(ptr @[[GLOB3:[0-9]+]], i32 [[TMP32]], i32 1, i32 4, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l69.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)
1238 // CHECK2-NEXT:    switch i32 [[TMP33]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [
1239 // CHECK2-NEXT:      i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]
1240 // CHECK2-NEXT:      i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]
1241 // CHECK2-NEXT:    ]
1242 // CHECK2:       .omp.reduction.case1:
1243 // CHECK2-NEXT:    [[TMP34:%.*]] = load i32, ptr [[ARRAYIDX]], align 4
1244 // CHECK2-NEXT:    [[TMP35:%.*]] = load i32, ptr [[OUTPUT1]], align 4
1245 // CHECK2-NEXT:    [[ADD8:%.*]] = add nsw i32 [[TMP34]], [[TMP35]]
1246 // CHECK2-NEXT:    store i32 [[ADD8]], ptr [[ARRAYIDX]], align 4
1247 // CHECK2-NEXT:    call void @__kmpc_end_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP32]], ptr @.gomp_critical_user_.reduction.var)
1248 // CHECK2-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]
1249 // CHECK2:       .omp.reduction.case2:
1250 // CHECK2-NEXT:    [[TMP36:%.*]] = load i32, ptr [[OUTPUT1]], align 4
1251 // CHECK2-NEXT:    [[TMP37:%.*]] = atomicrmw add ptr [[ARRAYIDX]], i32 [[TMP36]] monotonic, align 4
1252 // CHECK2-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]
1253 // CHECK2:       .omp.reduction.default:
1254 // CHECK2-NEXT:    ret void
1255 //
1256 //
1257 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l69.omp_outlined.omp_outlined
1258 // CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[DOTPREVIOUS_LB_:%.*]], i32 noundef [[DOTPREVIOUS_UB_:%.*]], i32 noundef [[SIZE:%.*]], ptr noundef [[OUTPUT:%.*]], ptr noundef [[INPUT:%.*]]) #[[ATTR1]] {
1259 // CHECK2-NEXT:  entry:
1260 // CHECK2-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
1261 // CHECK2-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
1262 // CHECK2-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i32, align 4
1263 // CHECK2-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i32, align 4
1264 // CHECK2-NEXT:    [[SIZE_ADDR:%.*]] = alloca i32, align 4
1265 // CHECK2-NEXT:    [[OUTPUT_ADDR:%.*]] = alloca ptr, align 4
1266 // CHECK2-NEXT:    [[INPUT_ADDR:%.*]] = alloca ptr, align 4
1267 // CHECK2-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
1268 // CHECK2-NEXT:    [[TMP:%.*]] = alloca i32, align 4
1269 // CHECK2-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4
1270 // CHECK2-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4
1271 // CHECK2-NEXT:    [[I:%.*]] = alloca i32, align 4
1272 // CHECK2-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4
1273 // CHECK2-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4
1274 // CHECK2-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
1275 // CHECK2-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
1276 // CHECK2-NEXT:    [[OUTPUT3:%.*]] = alloca i32, align 4
1277 // CHECK2-NEXT:    [[_TMP4:%.*]] = alloca ptr, align 4
1278 // CHECK2-NEXT:    [[I5:%.*]] = alloca i32, align 4
1279 // CHECK2-NEXT:    [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 4
1280 // CHECK2-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
1281 // CHECK2-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
1282 // CHECK2-NEXT:    store i32 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 4
1283 // CHECK2-NEXT:    store i32 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 4
1284 // CHECK2-NEXT:    store i32 [[SIZE]], ptr [[SIZE_ADDR]], align 4
1285 // CHECK2-NEXT:    store ptr [[OUTPUT]], ptr [[OUTPUT_ADDR]], align 4
1286 // CHECK2-NEXT:    store ptr [[INPUT]], ptr [[INPUT_ADDR]], align 4
1287 // CHECK2-NEXT:    [[TMP0:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
1288 // CHECK2-NEXT:    store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 4
1289 // CHECK2-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
1290 // CHECK2-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP1]], 0
1291 // CHECK2-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1
1292 // CHECK2-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1
1293 // CHECK2-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4
1294 // CHECK2-NEXT:    store i32 0, ptr [[I]], align 4
1295 // CHECK2-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
1296 // CHECK2-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP2]]
1297 // CHECK2-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]
1298 // CHECK2:       omp.precond.then:
1299 // CHECK2-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4
1300 // CHECK2-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4
1301 // CHECK2-NEXT:    store i32 [[TMP3]], ptr [[DOTOMP_UB]], align 4
1302 // CHECK2-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTPREVIOUS_LB__ADDR]], align 4
1303 // CHECK2-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTPREVIOUS_UB__ADDR]], align 4
1304 // CHECK2-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_LB]], align 4
1305 // CHECK2-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_UB]], align 4
1306 // CHECK2-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4
1307 // CHECK2-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4
1308 // CHECK2-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[OUTPUT_ADDR]], align 4
1309 // CHECK2-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP6]], i32 0
1310 // CHECK2-NEXT:    store i32 0, ptr [[OUTPUT3]], align 4
1311 // CHECK2-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[OUTPUT_ADDR]], align 4
1312 // CHECK2-NEXT:    [[TMP8:%.*]] = ptrtoint ptr [[TMP7]] to i64
1313 // CHECK2-NEXT:    [[TMP9:%.*]] = ptrtoint ptr [[ARRAYIDX]] to i64
1314 // CHECK2-NEXT:    [[TMP10:%.*]] = sub i64 [[TMP8]], [[TMP9]]
1315 // CHECK2-NEXT:    [[TMP11:%.*]] = sdiv exact i64 [[TMP10]], ptrtoint (ptr getelementptr (i32, ptr null, i32 1) to i64)
1316 // CHECK2-NEXT:    [[TMP12:%.*]] = getelementptr i32, ptr [[OUTPUT3]], i64 [[TMP11]]
1317 // CHECK2-NEXT:    store ptr [[TMP12]], ptr [[_TMP4]], align 4
1318 // CHECK2-NEXT:    [[TMP13:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
1319 // CHECK2-NEXT:    [[TMP14:%.*]] = load i32, ptr [[TMP13]], align 4
1320 // CHECK2-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP14]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
1321 // CHECK2-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
1322 // CHECK2-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4
1323 // CHECK2-NEXT:    [[CMP6:%.*]] = icmp sgt i32 [[TMP15]], [[TMP16]]
1324 // CHECK2-NEXT:    br i1 [[CMP6]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
1325 // CHECK2:       cond.true:
1326 // CHECK2-NEXT:    [[TMP17:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4
1327 // CHECK2-NEXT:    br label [[COND_END:%.*]]
1328 // CHECK2:       cond.false:
1329 // CHECK2-NEXT:    [[TMP18:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
1330 // CHECK2-NEXT:    br label [[COND_END]]
1331 // CHECK2:       cond.end:
1332 // CHECK2-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP17]], [[COND_TRUE]] ], [ [[TMP18]], [[COND_FALSE]] ]
1333 // CHECK2-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4
1334 // CHECK2-NEXT:    [[TMP19:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4
1335 // CHECK2-NEXT:    store i32 [[TMP19]], ptr [[DOTOMP_IV]], align 4
1336 // CHECK2-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
1337 // CHECK2:       omp.inner.for.cond:
1338 // CHECK2-NEXT:    [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
1339 // CHECK2-NEXT:    [[TMP21:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
1340 // CHECK2-NEXT:    [[CMP7:%.*]] = icmp sle i32 [[TMP20]], [[TMP21]]
1341 // CHECK2-NEXT:    br i1 [[CMP7]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
1342 // CHECK2:       omp.inner.for.body:
1343 // CHECK2-NEXT:    [[TMP22:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
1344 // CHECK2-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP22]], 1
1345 // CHECK2-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
1346 // CHECK2-NEXT:    store i32 [[ADD]], ptr [[I5]], align 4
1347 // CHECK2-NEXT:    [[TMP23:%.*]] = load ptr, ptr [[INPUT_ADDR]], align 4
1348 // CHECK2-NEXT:    [[TMP24:%.*]] = load i32, ptr [[I5]], align 4
1349 // CHECK2-NEXT:    [[ARRAYIDX8:%.*]] = getelementptr inbounds i32, ptr [[TMP23]], i32 [[TMP24]]
1350 // CHECK2-NEXT:    [[TMP25:%.*]] = load i32, ptr [[ARRAYIDX8]], align 4
1351 // CHECK2-NEXT:    [[TMP26:%.*]] = load ptr, ptr [[_TMP4]], align 4
1352 // CHECK2-NEXT:    [[ARRAYIDX9:%.*]] = getelementptr inbounds i32, ptr [[TMP26]], i32 0
1353 // CHECK2-NEXT:    [[TMP27:%.*]] = load i32, ptr [[ARRAYIDX9]], align 4
1354 // CHECK2-NEXT:    [[ADD10:%.*]] = add nsw i32 [[TMP27]], [[TMP25]]
1355 // CHECK2-NEXT:    store i32 [[ADD10]], ptr [[ARRAYIDX9]], align 4
1356 // CHECK2-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
1357 // CHECK2:       omp.body.continue:
1358 // CHECK2-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
1359 // CHECK2:       omp.inner.for.inc:
1360 // CHECK2-NEXT:    [[TMP28:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
1361 // CHECK2-NEXT:    [[ADD11:%.*]] = add nsw i32 [[TMP28]], 1
1362 // CHECK2-NEXT:    store i32 [[ADD11]], ptr [[DOTOMP_IV]], align 4
1363 // CHECK2-NEXT:    br label [[OMP_INNER_FOR_COND]]
1364 // CHECK2:       omp.inner.for.end:
1365 // CHECK2-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
1366 // CHECK2:       omp.loop.exit:
1367 // CHECK2-NEXT:    [[TMP29:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
1368 // CHECK2-NEXT:    [[TMP30:%.*]] = load i32, ptr [[TMP29]], align 4
1369 // CHECK2-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP30]])
1370 // CHECK2-NEXT:    [[TMP31:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 0
1371 // CHECK2-NEXT:    store ptr [[OUTPUT3]], ptr [[TMP31]], align 4
1372 // CHECK2-NEXT:    [[TMP32:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
1373 // CHECK2-NEXT:    [[TMP33:%.*]] = load i32, ptr [[TMP32]], align 4
1374 // CHECK2-NEXT:    [[TMP34:%.*]] = call i32 @__kmpc_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP33]], i32 1, i32 4, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l69.omp_outlined.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)
1375 // CHECK2-NEXT:    switch i32 [[TMP34]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [
1376 // CHECK2-NEXT:      i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]
1377 // CHECK2-NEXT:      i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]
1378 // CHECK2-NEXT:    ]
1379 // CHECK2:       .omp.reduction.case1:
1380 // CHECK2-NEXT:    [[TMP35:%.*]] = load i32, ptr [[ARRAYIDX]], align 4
1381 // CHECK2-NEXT:    [[TMP36:%.*]] = load i32, ptr [[OUTPUT3]], align 4
1382 // CHECK2-NEXT:    [[ADD12:%.*]] = add nsw i32 [[TMP35]], [[TMP36]]
1383 // CHECK2-NEXT:    store i32 [[ADD12]], ptr [[ARRAYIDX]], align 4
1384 // CHECK2-NEXT:    call void @__kmpc_end_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP33]], ptr @.gomp_critical_user_.reduction.var)
1385 // CHECK2-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]
1386 // CHECK2:       .omp.reduction.case2:
1387 // CHECK2-NEXT:    [[TMP37:%.*]] = load i32, ptr [[OUTPUT3]], align 4
1388 // CHECK2-NEXT:    [[TMP38:%.*]] = atomicrmw add ptr [[ARRAYIDX]], i32 [[TMP37]] monotonic, align 4
1389 // CHECK2-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]
1390 // CHECK2:       .omp.reduction.default:
1391 // CHECK2-NEXT:    br label [[OMP_PRECOND_END]]
1392 // CHECK2:       omp.precond.end:
1393 // CHECK2-NEXT:    ret void
1394 //
1395 //
1396 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l69.omp_outlined.omp_outlined.omp.reduction.reduction_func
1397 // CHECK2-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3:[0-9]+]] {
1398 // CHECK2-NEXT:  entry:
1399 // CHECK2-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 4
1400 // CHECK2-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 4
1401 // CHECK2-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 4
1402 // CHECK2-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 4
1403 // CHECK2-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 4
1404 // CHECK2-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 4
1405 // CHECK2-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i32 0, i32 0
1406 // CHECK2-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 4
1407 // CHECK2-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i32 0, i32 0
1408 // CHECK2-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 4
1409 // CHECK2-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4
1410 // CHECK2-NEXT:    [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4
1411 // CHECK2-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]
1412 // CHECK2-NEXT:    store i32 [[ADD]], ptr [[TMP7]], align 4
1413 // CHECK2-NEXT:    ret void
1414 //
1415 //
1416 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l69.omp_outlined.omp.reduction.reduction_func
1417 // CHECK2-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3]] {
1418 // CHECK2-NEXT:  entry:
1419 // CHECK2-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 4
1420 // CHECK2-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 4
1421 // CHECK2-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 4
1422 // CHECK2-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 4
1423 // CHECK2-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 4
1424 // CHECK2-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 4
1425 // CHECK2-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i32 0, i32 0
1426 // CHECK2-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 4
1427 // CHECK2-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i32 0, i32 0
1428 // CHECK2-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 4
1429 // CHECK2-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4
1430 // CHECK2-NEXT:    [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4
1431 // CHECK2-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]
1432 // CHECK2-NEXT:    store i32 [[ADD]], ptr [[TMP7]], align 4
1433 // CHECK2-NEXT:    ret void
1434 //
1435 //
1436 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l73
1437 // CHECK2-SAME: (i32 noundef [[SIZE:%.*]], ptr noundef [[OUTPUT:%.*]], ptr noundef [[INPUT:%.*]]) #[[ATTR1]] {
1438 // CHECK2-NEXT:  entry:
1439 // CHECK2-NEXT:    [[SIZE_ADDR:%.*]] = alloca i32, align 4
1440 // CHECK2-NEXT:    [[OUTPUT_ADDR:%.*]] = alloca ptr, align 4
1441 // CHECK2-NEXT:    [[INPUT_ADDR:%.*]] = alloca ptr, align 4
1442 // CHECK2-NEXT:    [[SIZE_CASTED:%.*]] = alloca i32, align 4
1443 // CHECK2-NEXT:    store i32 [[SIZE]], ptr [[SIZE_ADDR]], align 4
1444 // CHECK2-NEXT:    store ptr [[OUTPUT]], ptr [[OUTPUT_ADDR]], align 4
1445 // CHECK2-NEXT:    store ptr [[INPUT]], ptr [[INPUT_ADDR]], align 4
1446 // CHECK2-NEXT:    [[TMP0:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
1447 // CHECK2-NEXT:    store i32 [[TMP0]], ptr [[SIZE_CASTED]], align 4
1448 // CHECK2-NEXT:    [[TMP1:%.*]] = load i32, ptr [[SIZE_CASTED]], align 4
1449 // CHECK2-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[OUTPUT_ADDR]], align 4
1450 // CHECK2-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[INPUT_ADDR]], align 4
1451 // CHECK2-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB4]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l73.omp_outlined, i32 [[TMP1]], ptr [[TMP2]], ptr [[TMP3]])
1452 // CHECK2-NEXT:    ret void
1453 //
1454 //
1455 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l73.omp_outlined
1456 // CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[SIZE:%.*]], ptr noundef [[OUTPUT:%.*]], ptr noundef [[INPUT:%.*]]) #[[ATTR1]] {
1457 // CHECK2-NEXT:  entry:
1458 // CHECK2-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
1459 // CHECK2-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
1460 // CHECK2-NEXT:    [[SIZE_ADDR:%.*]] = alloca i32, align 4
1461 // CHECK2-NEXT:    [[OUTPUT_ADDR:%.*]] = alloca ptr, align 4
1462 // CHECK2-NEXT:    [[INPUT_ADDR:%.*]] = alloca ptr, align 4
1463 // CHECK2-NEXT:    [[OUTPUT2:%.*]] = alloca [3 x i32], align 4
1464 // CHECK2-NEXT:    [[TMP:%.*]] = alloca ptr, align 4
1465 // CHECK2-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
1466 // CHECK2-NEXT:    [[_TMP3:%.*]] = alloca i32, align 4
1467 // CHECK2-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4
1468 // CHECK2-NEXT:    [[DOTCAPTURE_EXPR_4:%.*]] = alloca i32, align 4
1469 // CHECK2-NEXT:    [[I:%.*]] = alloca i32, align 4
1470 // CHECK2-NEXT:    [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4
1471 // CHECK2-NEXT:    [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4
1472 // CHECK2-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
1473 // CHECK2-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
1474 // CHECK2-NEXT:    [[I6:%.*]] = alloca i32, align 4
1475 // CHECK2-NEXT:    [[SIZE_CASTED:%.*]] = alloca i32, align 4
1476 // CHECK2-NEXT:    [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 4
1477 // CHECK2-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
1478 // CHECK2-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
1479 // CHECK2-NEXT:    store i32 [[SIZE]], ptr [[SIZE_ADDR]], align 4
1480 // CHECK2-NEXT:    store ptr [[OUTPUT]], ptr [[OUTPUT_ADDR]], align 4
1481 // CHECK2-NEXT:    store ptr [[INPUT]], ptr [[INPUT_ADDR]], align 4
1482 // CHECK2-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[OUTPUT_ADDR]], align 4
1483 // CHECK2-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP0]], i32 0
1484 // CHECK2-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[OUTPUT_ADDR]], align 4
1485 // CHECK2-NEXT:    [[ARRAYIDX1:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP1]], i32 2
1486 // CHECK2-NEXT:    [[ARRAY_BEGIN:%.*]] = getelementptr inbounds [3 x i32], ptr [[OUTPUT2]], i32 0, i32 0
1487 // CHECK2-NEXT:    [[TMP2:%.*]] = getelementptr i32, ptr [[ARRAY_BEGIN]], i32 3
1488 // CHECK2-NEXT:    [[OMP_ARRAYINIT_ISEMPTY:%.*]] = icmp eq ptr [[ARRAY_BEGIN]], [[TMP2]]
1489 // CHECK2-NEXT:    br i1 [[OMP_ARRAYINIT_ISEMPTY]], label [[OMP_ARRAYINIT_DONE:%.*]], label [[OMP_ARRAYINIT_BODY:%.*]]
1490 // CHECK2:       omp.arrayinit.body:
1491 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DESTELEMENTPAST:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY:%.*]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT:%.*]], [[OMP_ARRAYINIT_BODY]] ]
1492 // CHECK2-NEXT:    store i32 0, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], align 4
1493 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DEST_ELEMENT]] = getelementptr i32, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], i32 1
1494 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DONE:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT]], [[TMP2]]
1495 // CHECK2-NEXT:    br i1 [[OMP_ARRAYCPY_DONE]], label [[OMP_ARRAYINIT_DONE]], label [[OMP_ARRAYINIT_BODY]]
1496 // CHECK2:       omp.arrayinit.done:
1497 // CHECK2-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[OUTPUT_ADDR]], align 4
1498 // CHECK2-NEXT:    [[TMP4:%.*]] = ptrtoint ptr [[TMP3]] to i64
1499 // CHECK2-NEXT:    [[TMP5:%.*]] = ptrtoint ptr [[ARRAYIDX]] to i64
1500 // CHECK2-NEXT:    [[TMP6:%.*]] = sub i64 [[TMP4]], [[TMP5]]
1501 // CHECK2-NEXT:    [[TMP7:%.*]] = sdiv exact i64 [[TMP6]], ptrtoint (ptr getelementptr (i32, ptr null, i32 1) to i64)
1502 // CHECK2-NEXT:    [[TMP8:%.*]] = getelementptr i32, ptr [[OUTPUT2]], i64 [[TMP7]]
1503 // CHECK2-NEXT:    store ptr [[TMP8]], ptr [[TMP]], align 4
1504 // CHECK2-NEXT:    [[TMP9:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
1505 // CHECK2-NEXT:    store i32 [[TMP9]], ptr [[DOTCAPTURE_EXPR_]], align 4
1506 // CHECK2-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
1507 // CHECK2-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP10]], 0
1508 // CHECK2-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1
1509 // CHECK2-NEXT:    [[SUB5:%.*]] = sub nsw i32 [[DIV]], 1
1510 // CHECK2-NEXT:    store i32 [[SUB5]], ptr [[DOTCAPTURE_EXPR_4]], align 4
1511 // CHECK2-NEXT:    store i32 0, ptr [[I]], align 4
1512 // CHECK2-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
1513 // CHECK2-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP11]]
1514 // CHECK2-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]
1515 // CHECK2:       omp.precond.then:
1516 // CHECK2-NEXT:    store i32 0, ptr [[DOTOMP_COMB_LB]], align 4
1517 // CHECK2-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 4
1518 // CHECK2-NEXT:    store i32 [[TMP12]], ptr [[DOTOMP_COMB_UB]], align 4
1519 // CHECK2-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4
1520 // CHECK2-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4
1521 // CHECK2-NEXT:    [[TMP13:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
1522 // CHECK2-NEXT:    [[TMP14:%.*]] = load i32, ptr [[TMP13]], align 4
1523 // CHECK2-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP14]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
1524 // CHECK2-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4
1525 // CHECK2-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 4
1526 // CHECK2-NEXT:    [[CMP7:%.*]] = icmp sgt i32 [[TMP15]], [[TMP16]]
1527 // CHECK2-NEXT:    br i1 [[CMP7]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
1528 // CHECK2:       cond.true:
1529 // CHECK2-NEXT:    [[TMP17:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 4
1530 // CHECK2-NEXT:    br label [[COND_END:%.*]]
1531 // CHECK2:       cond.false:
1532 // CHECK2-NEXT:    [[TMP18:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4
1533 // CHECK2-NEXT:    br label [[COND_END]]
1534 // CHECK2:       cond.end:
1535 // CHECK2-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP17]], [[COND_TRUE]] ], [ [[TMP18]], [[COND_FALSE]] ]
1536 // CHECK2-NEXT:    store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4
1537 // CHECK2-NEXT:    [[TMP19:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4
1538 // CHECK2-NEXT:    store i32 [[TMP19]], ptr [[DOTOMP_IV]], align 4
1539 // CHECK2-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
1540 // CHECK2:       omp.inner.for.cond:
1541 // CHECK2-NEXT:    [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
1542 // CHECK2-NEXT:    [[TMP21:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4
1543 // CHECK2-NEXT:    [[CMP8:%.*]] = icmp sle i32 [[TMP20]], [[TMP21]]
1544 // CHECK2-NEXT:    br i1 [[CMP8]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
1545 // CHECK2:       omp.inner.for.body:
1546 // CHECK2-NEXT:    [[TMP22:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4
1547 // CHECK2-NEXT:    [[TMP23:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4
1548 // CHECK2-NEXT:    [[TMP24:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
1549 // CHECK2-NEXT:    store i32 [[TMP24]], ptr [[SIZE_CASTED]], align 4
1550 // CHECK2-NEXT:    [[TMP25:%.*]] = load i32, ptr [[SIZE_CASTED]], align 4
1551 // CHECK2-NEXT:    [[TMP26:%.*]] = load ptr, ptr [[TMP]], align 4
1552 // CHECK2-NEXT:    [[TMP27:%.*]] = load ptr, ptr [[INPUT_ADDR]], align 4
1553 // CHECK2-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB4]], i32 5, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l73.omp_outlined.omp_outlined, i32 [[TMP22]], i32 [[TMP23]], i32 [[TMP25]], ptr [[TMP26]], ptr [[TMP27]])
1554 // CHECK2-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
1555 // CHECK2:       omp.inner.for.inc:
1556 // CHECK2-NEXT:    [[TMP28:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
1557 // CHECK2-NEXT:    [[TMP29:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4
1558 // CHECK2-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP28]], [[TMP29]]
1559 // CHECK2-NEXT:    store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4
1560 // CHECK2-NEXT:    br label [[OMP_INNER_FOR_COND]]
1561 // CHECK2:       omp.inner.for.end:
1562 // CHECK2-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
1563 // CHECK2:       omp.loop.exit:
1564 // CHECK2-NEXT:    [[TMP30:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
1565 // CHECK2-NEXT:    [[TMP31:%.*]] = load i32, ptr [[TMP30]], align 4
1566 // CHECK2-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP31]])
1567 // CHECK2-NEXT:    br label [[OMP_PRECOND_END]]
1568 // CHECK2:       omp.precond.end:
1569 // CHECK2-NEXT:    [[TMP32:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 0
1570 // CHECK2-NEXT:    store ptr [[OUTPUT2]], ptr [[TMP32]], align 4
1571 // CHECK2-NEXT:    [[TMP33:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
1572 // CHECK2-NEXT:    [[TMP34:%.*]] = load i32, ptr [[TMP33]], align 4
1573 // CHECK2-NEXT:    [[TMP35:%.*]] = call i32 @__kmpc_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP34]], i32 1, i32 4, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l73.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)
1574 // CHECK2-NEXT:    switch i32 [[TMP35]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [
1575 // CHECK2-NEXT:      i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]
1576 // CHECK2-NEXT:      i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]
1577 // CHECK2-NEXT:    ]
1578 // CHECK2:       .omp.reduction.case1:
1579 // CHECK2-NEXT:    [[TMP36:%.*]] = getelementptr i32, ptr [[ARRAYIDX]], i32 3
1580 // CHECK2-NEXT:    [[OMP_ARRAYCPY_ISEMPTY:%.*]] = icmp eq ptr [[ARRAYIDX]], [[TMP36]]
1581 // CHECK2-NEXT:    br i1 [[OMP_ARRAYCPY_ISEMPTY]], label [[OMP_ARRAYCPY_DONE13:%.*]], label [[OMP_ARRAYCPY_BODY:%.*]]
1582 // CHECK2:       omp.arraycpy.body:
1583 // CHECK2-NEXT:    [[OMP_ARRAYCPY_SRCELEMENTPAST:%.*]] = phi ptr [ [[OUTPUT2]], [[DOTOMP_REDUCTION_CASE1]] ], [ [[OMP_ARRAYCPY_SRC_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
1584 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DESTELEMENTPAST9:%.*]] = phi ptr [ [[ARRAYIDX]], [[DOTOMP_REDUCTION_CASE1]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT11:%.*]], [[OMP_ARRAYCPY_BODY]] ]
1585 // CHECK2-NEXT:    [[TMP37:%.*]] = load i32, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST9]], align 4
1586 // CHECK2-NEXT:    [[TMP38:%.*]] = load i32, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], align 4
1587 // CHECK2-NEXT:    [[ADD10:%.*]] = add nsw i32 [[TMP37]], [[TMP38]]
1588 // CHECK2-NEXT:    store i32 [[ADD10]], ptr [[OMP_ARRAYCPY_DESTELEMENTPAST9]], align 4
1589 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DEST_ELEMENT11]] = getelementptr i32, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST9]], i32 1
1590 // CHECK2-NEXT:    [[OMP_ARRAYCPY_SRC_ELEMENT]] = getelementptr i32, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], i32 1
1591 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DONE12:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT11]], [[TMP36]]
1592 // CHECK2-NEXT:    br i1 [[OMP_ARRAYCPY_DONE12]], label [[OMP_ARRAYCPY_DONE13]], label [[OMP_ARRAYCPY_BODY]]
1593 // CHECK2:       omp.arraycpy.done13:
1594 // CHECK2-NEXT:    call void @__kmpc_end_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP34]], ptr @.gomp_critical_user_.reduction.var)
1595 // CHECK2-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]
1596 // CHECK2:       .omp.reduction.case2:
1597 // CHECK2-NEXT:    [[TMP39:%.*]] = getelementptr i32, ptr [[ARRAYIDX]], i32 3
1598 // CHECK2-NEXT:    [[OMP_ARRAYCPY_ISEMPTY14:%.*]] = icmp eq ptr [[ARRAYIDX]], [[TMP39]]
1599 // CHECK2-NEXT:    br i1 [[OMP_ARRAYCPY_ISEMPTY14]], label [[OMP_ARRAYCPY_DONE21:%.*]], label [[OMP_ARRAYCPY_BODY15:%.*]]
1600 // CHECK2:       omp.arraycpy.body15:
1601 // CHECK2-NEXT:    [[OMP_ARRAYCPY_SRCELEMENTPAST16:%.*]] = phi ptr [ [[OUTPUT2]], [[DOTOMP_REDUCTION_CASE2]] ], [ [[OMP_ARRAYCPY_SRC_ELEMENT19:%.*]], [[OMP_ARRAYCPY_BODY15]] ]
1602 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DESTELEMENTPAST17:%.*]] = phi ptr [ [[ARRAYIDX]], [[DOTOMP_REDUCTION_CASE2]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT18:%.*]], [[OMP_ARRAYCPY_BODY15]] ]
1603 // CHECK2-NEXT:    [[TMP40:%.*]] = load i32, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST16]], align 4
1604 // CHECK2-NEXT:    [[TMP41:%.*]] = atomicrmw add ptr [[OMP_ARRAYCPY_DESTELEMENTPAST17]], i32 [[TMP40]] monotonic, align 4
1605 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DEST_ELEMENT18]] = getelementptr i32, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST17]], i32 1
1606 // CHECK2-NEXT:    [[OMP_ARRAYCPY_SRC_ELEMENT19]] = getelementptr i32, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST16]], i32 1
1607 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DONE20:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT18]], [[TMP39]]
1608 // CHECK2-NEXT:    br i1 [[OMP_ARRAYCPY_DONE20]], label [[OMP_ARRAYCPY_DONE21]], label [[OMP_ARRAYCPY_BODY15]]
1609 // CHECK2:       omp.arraycpy.done21:
1610 // CHECK2-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]
1611 // CHECK2:       .omp.reduction.default:
1612 // CHECK2-NEXT:    ret void
1613 //
1614 //
1615 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l73.omp_outlined.omp_outlined
1616 // CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[DOTPREVIOUS_LB_:%.*]], i32 noundef [[DOTPREVIOUS_UB_:%.*]], i32 noundef [[SIZE:%.*]], ptr noundef [[OUTPUT:%.*]], ptr noundef [[INPUT:%.*]]) #[[ATTR1]] {
1617 // CHECK2-NEXT:  entry:
1618 // CHECK2-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
1619 // CHECK2-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
1620 // CHECK2-NEXT:    [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i32, align 4
1621 // CHECK2-NEXT:    [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i32, align 4
1622 // CHECK2-NEXT:    [[SIZE_ADDR:%.*]] = alloca i32, align 4
1623 // CHECK2-NEXT:    [[OUTPUT_ADDR:%.*]] = alloca ptr, align 4
1624 // CHECK2-NEXT:    [[INPUT_ADDR:%.*]] = alloca ptr, align 4
1625 // CHECK2-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
1626 // CHECK2-NEXT:    [[TMP:%.*]] = alloca i32, align 4
1627 // CHECK2-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4
1628 // CHECK2-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4
1629 // CHECK2-NEXT:    [[I:%.*]] = alloca i32, align 4
1630 // CHECK2-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4
1631 // CHECK2-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4
1632 // CHECK2-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
1633 // CHECK2-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
1634 // CHECK2-NEXT:    [[OUTPUT4:%.*]] = alloca [3 x i32], align 4
1635 // CHECK2-NEXT:    [[_TMP5:%.*]] = alloca ptr, align 4
1636 // CHECK2-NEXT:    [[I6:%.*]] = alloca i32, align 4
1637 // CHECK2-NEXT:    [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 4
1638 // CHECK2-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
1639 // CHECK2-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
1640 // CHECK2-NEXT:    store i32 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 4
1641 // CHECK2-NEXT:    store i32 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 4
1642 // CHECK2-NEXT:    store i32 [[SIZE]], ptr [[SIZE_ADDR]], align 4
1643 // CHECK2-NEXT:    store ptr [[OUTPUT]], ptr [[OUTPUT_ADDR]], align 4
1644 // CHECK2-NEXT:    store ptr [[INPUT]], ptr [[INPUT_ADDR]], align 4
1645 // CHECK2-NEXT:    [[TMP0:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
1646 // CHECK2-NEXT:    store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 4
1647 // CHECK2-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
1648 // CHECK2-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP1]], 0
1649 // CHECK2-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1
1650 // CHECK2-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1
1651 // CHECK2-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4
1652 // CHECK2-NEXT:    store i32 0, ptr [[I]], align 4
1653 // CHECK2-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
1654 // CHECK2-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP2]]
1655 // CHECK2-NEXT:    br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]
1656 // CHECK2:       omp.precond.then:
1657 // CHECK2-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4
1658 // CHECK2-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4
1659 // CHECK2-NEXT:    store i32 [[TMP3]], ptr [[DOTOMP_UB]], align 4
1660 // CHECK2-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTPREVIOUS_LB__ADDR]], align 4
1661 // CHECK2-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTPREVIOUS_UB__ADDR]], align 4
1662 // CHECK2-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_LB]], align 4
1663 // CHECK2-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_UB]], align 4
1664 // CHECK2-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4
1665 // CHECK2-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4
1666 // CHECK2-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[OUTPUT_ADDR]], align 4
1667 // CHECK2-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP6]], i32 0
1668 // CHECK2-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[OUTPUT_ADDR]], align 4
1669 // CHECK2-NEXT:    [[ARRAYIDX3:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP7]], i32 2
1670 // CHECK2-NEXT:    [[ARRAY_BEGIN:%.*]] = getelementptr inbounds [3 x i32], ptr [[OUTPUT4]], i32 0, i32 0
1671 // CHECK2-NEXT:    [[TMP8:%.*]] = getelementptr i32, ptr [[ARRAY_BEGIN]], i32 3
1672 // CHECK2-NEXT:    [[OMP_ARRAYINIT_ISEMPTY:%.*]] = icmp eq ptr [[ARRAY_BEGIN]], [[TMP8]]
1673 // CHECK2-NEXT:    br i1 [[OMP_ARRAYINIT_ISEMPTY]], label [[OMP_ARRAYINIT_DONE:%.*]], label [[OMP_ARRAYINIT_BODY:%.*]]
1674 // CHECK2:       omp.arrayinit.body:
1675 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DESTELEMENTPAST:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[OMP_PRECOND_THEN]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT:%.*]], [[OMP_ARRAYINIT_BODY]] ]
1676 // CHECK2-NEXT:    store i32 0, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], align 4
1677 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DEST_ELEMENT]] = getelementptr i32, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], i32 1
1678 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DONE:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT]], [[TMP8]]
1679 // CHECK2-NEXT:    br i1 [[OMP_ARRAYCPY_DONE]], label [[OMP_ARRAYINIT_DONE]], label [[OMP_ARRAYINIT_BODY]]
1680 // CHECK2:       omp.arrayinit.done:
1681 // CHECK2-NEXT:    [[TMP9:%.*]] = load ptr, ptr [[OUTPUT_ADDR]], align 4
1682 // CHECK2-NEXT:    [[TMP10:%.*]] = ptrtoint ptr [[TMP9]] to i64
1683 // CHECK2-NEXT:    [[TMP11:%.*]] = ptrtoint ptr [[ARRAYIDX]] to i64
1684 // CHECK2-NEXT:    [[TMP12:%.*]] = sub i64 [[TMP10]], [[TMP11]]
1685 // CHECK2-NEXT:    [[TMP13:%.*]] = sdiv exact i64 [[TMP12]], ptrtoint (ptr getelementptr (i32, ptr null, i32 1) to i64)
1686 // CHECK2-NEXT:    [[TMP14:%.*]] = getelementptr i32, ptr [[OUTPUT4]], i64 [[TMP13]]
1687 // CHECK2-NEXT:    store ptr [[TMP14]], ptr [[_TMP5]], align 4
1688 // CHECK2-NEXT:    [[TMP15:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
1689 // CHECK2-NEXT:    [[TMP16:%.*]] = load i32, ptr [[TMP15]], align 4
1690 // CHECK2-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB2]], i32 [[TMP16]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
1691 // CHECK2-NEXT:    [[TMP17:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
1692 // CHECK2-NEXT:    [[TMP18:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4
1693 // CHECK2-NEXT:    [[CMP7:%.*]] = icmp sgt i32 [[TMP17]], [[TMP18]]
1694 // CHECK2-NEXT:    br i1 [[CMP7]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
1695 // CHECK2:       cond.true:
1696 // CHECK2-NEXT:    [[TMP19:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4
1697 // CHECK2-NEXT:    br label [[COND_END:%.*]]
1698 // CHECK2:       cond.false:
1699 // CHECK2-NEXT:    [[TMP20:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
1700 // CHECK2-NEXT:    br label [[COND_END]]
1701 // CHECK2:       cond.end:
1702 // CHECK2-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP19]], [[COND_TRUE]] ], [ [[TMP20]], [[COND_FALSE]] ]
1703 // CHECK2-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4
1704 // CHECK2-NEXT:    [[TMP21:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4
1705 // CHECK2-NEXT:    store i32 [[TMP21]], ptr [[DOTOMP_IV]], align 4
1706 // CHECK2-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
1707 // CHECK2:       omp.inner.for.cond:
1708 // CHECK2-NEXT:    [[TMP22:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
1709 // CHECK2-NEXT:    [[TMP23:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
1710 // CHECK2-NEXT:    [[CMP8:%.*]] = icmp sle i32 [[TMP22]], [[TMP23]]
1711 // CHECK2-NEXT:    br i1 [[CMP8]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
1712 // CHECK2:       omp.inner.for.body:
1713 // CHECK2-NEXT:    [[TMP24:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
1714 // CHECK2-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP24]], 1
1715 // CHECK2-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
1716 // CHECK2-NEXT:    store i32 [[ADD]], ptr [[I6]], align 4
1717 // CHECK2-NEXT:    [[TMP25:%.*]] = load ptr, ptr [[INPUT_ADDR]], align 4
1718 // CHECK2-NEXT:    [[TMP26:%.*]] = load i32, ptr [[I6]], align 4
1719 // CHECK2-NEXT:    [[ARRAYIDX9:%.*]] = getelementptr inbounds i32, ptr [[TMP25]], i32 [[TMP26]]
1720 // CHECK2-NEXT:    [[TMP27:%.*]] = load i32, ptr [[ARRAYIDX9]], align 4
1721 // CHECK2-NEXT:    [[TMP28:%.*]] = load ptr, ptr [[_TMP5]], align 4
1722 // CHECK2-NEXT:    [[ARRAYIDX10:%.*]] = getelementptr inbounds i32, ptr [[TMP28]], i32 0
1723 // CHECK2-NEXT:    [[TMP29:%.*]] = load i32, ptr [[ARRAYIDX10]], align 4
1724 // CHECK2-NEXT:    [[ADD11:%.*]] = add nsw i32 [[TMP29]], [[TMP27]]
1725 // CHECK2-NEXT:    store i32 [[ADD11]], ptr [[ARRAYIDX10]], align 4
1726 // CHECK2-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
1727 // CHECK2:       omp.body.continue:
1728 // CHECK2-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
1729 // CHECK2:       omp.inner.for.inc:
1730 // CHECK2-NEXT:    [[TMP30:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
1731 // CHECK2-NEXT:    [[ADD12:%.*]] = add nsw i32 [[TMP30]], 1
1732 // CHECK2-NEXT:    store i32 [[ADD12]], ptr [[DOTOMP_IV]], align 4
1733 // CHECK2-NEXT:    br label [[OMP_INNER_FOR_COND]]
1734 // CHECK2:       omp.inner.for.end:
1735 // CHECK2-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
1736 // CHECK2:       omp.loop.exit:
1737 // CHECK2-NEXT:    [[TMP31:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
1738 // CHECK2-NEXT:    [[TMP32:%.*]] = load i32, ptr [[TMP31]], align 4
1739 // CHECK2-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP32]])
1740 // CHECK2-NEXT:    [[TMP33:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 0
1741 // CHECK2-NEXT:    store ptr [[OUTPUT4]], ptr [[TMP33]], align 4
1742 // CHECK2-NEXT:    [[TMP34:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
1743 // CHECK2-NEXT:    [[TMP35:%.*]] = load i32, ptr [[TMP34]], align 4
1744 // CHECK2-NEXT:    [[TMP36:%.*]] = call i32 @__kmpc_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP35]], i32 1, i32 4, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l73.omp_outlined.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)
1745 // CHECK2-NEXT:    switch i32 [[TMP36]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [
1746 // CHECK2-NEXT:      i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]
1747 // CHECK2-NEXT:      i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]
1748 // CHECK2-NEXT:    ]
1749 // CHECK2:       .omp.reduction.case1:
1750 // CHECK2-NEXT:    [[TMP37:%.*]] = getelementptr i32, ptr [[ARRAYIDX]], i32 3
1751 // CHECK2-NEXT:    [[OMP_ARRAYCPY_ISEMPTY:%.*]] = icmp eq ptr [[ARRAYIDX]], [[TMP37]]
1752 // CHECK2-NEXT:    br i1 [[OMP_ARRAYCPY_ISEMPTY]], label [[OMP_ARRAYCPY_DONE17:%.*]], label [[OMP_ARRAYCPY_BODY:%.*]]
1753 // CHECK2:       omp.arraycpy.body:
1754 // CHECK2-NEXT:    [[OMP_ARRAYCPY_SRCELEMENTPAST:%.*]] = phi ptr [ [[OUTPUT4]], [[DOTOMP_REDUCTION_CASE1]] ], [ [[OMP_ARRAYCPY_SRC_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
1755 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DESTELEMENTPAST13:%.*]] = phi ptr [ [[ARRAYIDX]], [[DOTOMP_REDUCTION_CASE1]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT15:%.*]], [[OMP_ARRAYCPY_BODY]] ]
1756 // CHECK2-NEXT:    [[TMP38:%.*]] = load i32, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST13]], align 4
1757 // CHECK2-NEXT:    [[TMP39:%.*]] = load i32, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], align 4
1758 // CHECK2-NEXT:    [[ADD14:%.*]] = add nsw i32 [[TMP38]], [[TMP39]]
1759 // CHECK2-NEXT:    store i32 [[ADD14]], ptr [[OMP_ARRAYCPY_DESTELEMENTPAST13]], align 4
1760 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DEST_ELEMENT15]] = getelementptr i32, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST13]], i32 1
1761 // CHECK2-NEXT:    [[OMP_ARRAYCPY_SRC_ELEMENT]] = getelementptr i32, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], i32 1
1762 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DONE16:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT15]], [[TMP37]]
1763 // CHECK2-NEXT:    br i1 [[OMP_ARRAYCPY_DONE16]], label [[OMP_ARRAYCPY_DONE17]], label [[OMP_ARRAYCPY_BODY]]
1764 // CHECK2:       omp.arraycpy.done17:
1765 // CHECK2-NEXT:    call void @__kmpc_end_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP35]], ptr @.gomp_critical_user_.reduction.var)
1766 // CHECK2-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]
1767 // CHECK2:       .omp.reduction.case2:
1768 // CHECK2-NEXT:    [[TMP40:%.*]] = getelementptr i32, ptr [[ARRAYIDX]], i32 3
1769 // CHECK2-NEXT:    [[OMP_ARRAYCPY_ISEMPTY18:%.*]] = icmp eq ptr [[ARRAYIDX]], [[TMP40]]
1770 // CHECK2-NEXT:    br i1 [[OMP_ARRAYCPY_ISEMPTY18]], label [[OMP_ARRAYCPY_DONE25:%.*]], label [[OMP_ARRAYCPY_BODY19:%.*]]
1771 // CHECK2:       omp.arraycpy.body19:
1772 // CHECK2-NEXT:    [[OMP_ARRAYCPY_SRCELEMENTPAST20:%.*]] = phi ptr [ [[OUTPUT4]], [[DOTOMP_REDUCTION_CASE2]] ], [ [[OMP_ARRAYCPY_SRC_ELEMENT23:%.*]], [[OMP_ARRAYCPY_BODY19]] ]
1773 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DESTELEMENTPAST21:%.*]] = phi ptr [ [[ARRAYIDX]], [[DOTOMP_REDUCTION_CASE2]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT22:%.*]], [[OMP_ARRAYCPY_BODY19]] ]
1774 // CHECK2-NEXT:    [[TMP41:%.*]] = load i32, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST20]], align 4
1775 // CHECK2-NEXT:    [[TMP42:%.*]] = atomicrmw add ptr [[OMP_ARRAYCPY_DESTELEMENTPAST21]], i32 [[TMP41]] monotonic, align 4
1776 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DEST_ELEMENT22]] = getelementptr i32, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST21]], i32 1
1777 // CHECK2-NEXT:    [[OMP_ARRAYCPY_SRC_ELEMENT23]] = getelementptr i32, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST20]], i32 1
1778 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DONE24:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT22]], [[TMP40]]
1779 // CHECK2-NEXT:    br i1 [[OMP_ARRAYCPY_DONE24]], label [[OMP_ARRAYCPY_DONE25]], label [[OMP_ARRAYCPY_BODY19]]
1780 // CHECK2:       omp.arraycpy.done25:
1781 // CHECK2-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]
1782 // CHECK2:       .omp.reduction.default:
1783 // CHECK2-NEXT:    br label [[OMP_PRECOND_END]]
1784 // CHECK2:       omp.precond.end:
1785 // CHECK2-NEXT:    ret void
1786 //
1787 //
1788 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l73.omp_outlined.omp_outlined.omp.reduction.reduction_func
1789 // CHECK2-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3]] {
1790 // CHECK2-NEXT:  entry:
1791 // CHECK2-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 4
1792 // CHECK2-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 4
1793 // CHECK2-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 4
1794 // CHECK2-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 4
1795 // CHECK2-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 4
1796 // CHECK2-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 4
1797 // CHECK2-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i32 0, i32 0
1798 // CHECK2-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 4
1799 // CHECK2-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i32 0, i32 0
1800 // CHECK2-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 4
1801 // CHECK2-NEXT:    [[TMP8:%.*]] = getelementptr i32, ptr [[TMP7]], i32 3
1802 // CHECK2-NEXT:    [[OMP_ARRAYCPY_ISEMPTY:%.*]] = icmp eq ptr [[TMP7]], [[TMP8]]
1803 // CHECK2-NEXT:    br i1 [[OMP_ARRAYCPY_ISEMPTY]], label [[OMP_ARRAYCPY_DONE2:%.*]], label [[OMP_ARRAYCPY_BODY:%.*]]
1804 // CHECK2:       omp.arraycpy.body:
1805 // CHECK2-NEXT:    [[OMP_ARRAYCPY_SRCELEMENTPAST:%.*]] = phi ptr [ [[TMP5]], [[ENTRY:%.*]] ], [ [[OMP_ARRAYCPY_SRC_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
1806 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DESTELEMENTPAST:%.*]] = phi ptr [ [[TMP7]], [[ENTRY]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
1807 // CHECK2-NEXT:    [[TMP9:%.*]] = load i32, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], align 4
1808 // CHECK2-NEXT:    [[TMP10:%.*]] = load i32, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], align 4
1809 // CHECK2-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP9]], [[TMP10]]
1810 // CHECK2-NEXT:    store i32 [[ADD]], ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], align 4
1811 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DEST_ELEMENT]] = getelementptr i32, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], i32 1
1812 // CHECK2-NEXT:    [[OMP_ARRAYCPY_SRC_ELEMENT]] = getelementptr i32, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], i32 1
1813 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DONE:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT]], [[TMP8]]
1814 // CHECK2-NEXT:    br i1 [[OMP_ARRAYCPY_DONE]], label [[OMP_ARRAYCPY_DONE2]], label [[OMP_ARRAYCPY_BODY]]
1815 // CHECK2:       omp.arraycpy.done2:
1816 // CHECK2-NEXT:    ret void
1817 //
1818 //
1819 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l73.omp_outlined.omp.reduction.reduction_func
1820 // CHECK2-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3]] {
1821 // CHECK2-NEXT:  entry:
1822 // CHECK2-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 4
1823 // CHECK2-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 4
1824 // CHECK2-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 4
1825 // CHECK2-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 4
1826 // CHECK2-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 4
1827 // CHECK2-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 4
1828 // CHECK2-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i32 0, i32 0
1829 // CHECK2-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 4
1830 // CHECK2-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i32 0, i32 0
1831 // CHECK2-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 4
1832 // CHECK2-NEXT:    [[TMP8:%.*]] = getelementptr i32, ptr [[TMP7]], i32 3
1833 // CHECK2-NEXT:    [[OMP_ARRAYCPY_ISEMPTY:%.*]] = icmp eq ptr [[TMP7]], [[TMP8]]
1834 // CHECK2-NEXT:    br i1 [[OMP_ARRAYCPY_ISEMPTY]], label [[OMP_ARRAYCPY_DONE2:%.*]], label [[OMP_ARRAYCPY_BODY:%.*]]
1835 // CHECK2:       omp.arraycpy.body:
1836 // CHECK2-NEXT:    [[OMP_ARRAYCPY_SRCELEMENTPAST:%.*]] = phi ptr [ [[TMP5]], [[ENTRY:%.*]] ], [ [[OMP_ARRAYCPY_SRC_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
1837 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DESTELEMENTPAST:%.*]] = phi ptr [ [[TMP7]], [[ENTRY]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
1838 // CHECK2-NEXT:    [[TMP9:%.*]] = load i32, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], align 4
1839 // CHECK2-NEXT:    [[TMP10:%.*]] = load i32, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], align 4
1840 // CHECK2-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP9]], [[TMP10]]
1841 // CHECK2-NEXT:    store i32 [[ADD]], ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], align 4
1842 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DEST_ELEMENT]] = getelementptr i32, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], i32 1
1843 // CHECK2-NEXT:    [[OMP_ARRAYCPY_SRC_ELEMENT]] = getelementptr i32, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], i32 1
1844 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DONE:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT]], [[TMP8]]
1845 // CHECK2-NEXT:    br i1 [[OMP_ARRAYCPY_DONE]], label [[OMP_ARRAYCPY_DONE2]], label [[OMP_ARRAYCPY_BODY]]
1846 // CHECK2:       omp.arraycpy.done2:
1847 // CHECK2-NEXT:    ret void
1848 //
1849 //
1850 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l78
1851 // CHECK2-SAME: (i32 noundef [[SIZE:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR1]] {
1852 // CHECK2-NEXT:  entry:
1853 // CHECK2-NEXT:    [[SIZE_ADDR:%.*]] = alloca i32, align 4
1854 // CHECK2-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 4
1855 // CHECK2-NEXT:    [[SIZE_CASTED:%.*]] = alloca i32, align 4
1856 // CHECK2-NEXT:    store i32 [[SIZE]], ptr [[SIZE_ADDR]], align 4
1857 // CHECK2-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 4
1858 // CHECK2-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 4
1859 // CHECK2-NEXT:    [[TMP1:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
1860 // CHECK2-NEXT:    store i32 [[TMP1]], ptr [[SIZE_CASTED]], align 4
1861 // CHECK2-NEXT:    [[TMP2:%.*]] = load i32, ptr [[SIZE_CASTED]], align 4
1862 // CHECK2-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB4]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l78.omp_outlined, i32 [[TMP2]], ptr [[TMP0]])
1863 // CHECK2-NEXT:    ret void
1864 //
1865 //
1866 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l78.omp_outlined
1867 // CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[SIZE:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR1]] {
1868 // CHECK2-NEXT:  entry:
1869 // CHECK2-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
1870 // CHECK2-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
1871 // CHECK2-NEXT:    [[SIZE_ADDR:%.*]] = alloca i32, align 4
1872 // CHECK2-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 4
1873 // CHECK2-NEXT:    [[A2:%.*]] = alloca [2 x i32], align 4
1874 // CHECK2-NEXT:    [[I:%.*]] = alloca i32, align 4
1875 // CHECK2-NEXT:    [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 4
1876 // CHECK2-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
1877 // CHECK2-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
1878 // CHECK2-NEXT:    store i32 [[SIZE]], ptr [[SIZE_ADDR]], align 4
1879 // CHECK2-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 4
1880 // CHECK2-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 4
1881 // CHECK2-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds nuw [10 x i32], ptr [[TMP0]], i32 0, i32 0
1882 // CHECK2-NEXT:    [[ARRAYIDX1:%.*]] = getelementptr inbounds nuw [10 x i32], ptr [[TMP0]], i32 0, i32 1
1883 // CHECK2-NEXT:    [[ARRAY_BEGIN:%.*]] = getelementptr inbounds [2 x i32], ptr [[A2]], i32 0, i32 0
1884 // CHECK2-NEXT:    [[TMP1:%.*]] = getelementptr i32, ptr [[ARRAY_BEGIN]], i32 2
1885 // CHECK2-NEXT:    [[OMP_ARRAYINIT_ISEMPTY:%.*]] = icmp eq ptr [[ARRAY_BEGIN]], [[TMP1]]
1886 // CHECK2-NEXT:    br i1 [[OMP_ARRAYINIT_ISEMPTY]], label [[OMP_ARRAYINIT_DONE:%.*]], label [[OMP_ARRAYINIT_BODY:%.*]]
1887 // CHECK2:       omp.arrayinit.body:
1888 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DESTELEMENTPAST:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY:%.*]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT:%.*]], [[OMP_ARRAYINIT_BODY]] ]
1889 // CHECK2-NEXT:    store i32 0, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], align 4
1890 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DEST_ELEMENT]] = getelementptr i32, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], i32 1
1891 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DONE:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT]], [[TMP1]]
1892 // CHECK2-NEXT:    br i1 [[OMP_ARRAYCPY_DONE]], label [[OMP_ARRAYINIT_DONE]], label [[OMP_ARRAYINIT_BODY]]
1893 // CHECK2:       omp.arrayinit.done:
1894 // CHECK2-NEXT:    [[TMP2:%.*]] = ptrtoint ptr [[TMP0]] to i64
1895 // CHECK2-NEXT:    [[TMP3:%.*]] = ptrtoint ptr [[ARRAYIDX]] to i64
1896 // CHECK2-NEXT:    [[TMP4:%.*]] = sub i64 [[TMP2]], [[TMP3]]
1897 // CHECK2-NEXT:    [[TMP5:%.*]] = sdiv exact i64 [[TMP4]], ptrtoint (ptr getelementptr (i32, ptr null, i32 1) to i64)
1898 // CHECK2-NEXT:    [[TMP6:%.*]] = getelementptr i32, ptr [[A2]], i64 [[TMP5]]
1899 // CHECK2-NEXT:    store i32 0, ptr [[I]], align 4
1900 // CHECK2-NEXT:    br label [[FOR_COND:%.*]]
1901 // CHECK2:       for.cond:
1902 // CHECK2-NEXT:    [[TMP7:%.*]] = load i32, ptr [[I]], align 4
1903 // CHECK2-NEXT:    [[TMP8:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
1904 // CHECK2-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP7]], [[TMP8]]
1905 // CHECK2-NEXT:    br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]]
1906 // CHECK2:       for.body:
1907 // CHECK2-NEXT:    br label [[FOR_INC:%.*]]
1908 // CHECK2:       for.inc:
1909 // CHECK2-NEXT:    [[TMP9:%.*]] = load i32, ptr [[I]], align 4
1910 // CHECK2-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP9]], 1
1911 // CHECK2-NEXT:    store i32 [[INC]], ptr [[I]], align 4
1912 // CHECK2-NEXT:    br label [[FOR_COND]], !llvm.loop [[LOOP10:![0-9]+]]
1913 // CHECK2:       for.end:
1914 // CHECK2-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 0
1915 // CHECK2-NEXT:    store ptr [[A2]], ptr [[TMP10]], align 4
1916 // CHECK2-NEXT:    [[TMP11:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
1917 // CHECK2-NEXT:    [[TMP12:%.*]] = load i32, ptr [[TMP11]], align 4
1918 // CHECK2-NEXT:    [[TMP13:%.*]] = call i32 @__kmpc_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP12]], i32 1, i32 4, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l78.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)
1919 // CHECK2-NEXT:    switch i32 [[TMP13]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [
1920 // CHECK2-NEXT:      i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]
1921 // CHECK2-NEXT:      i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]
1922 // CHECK2-NEXT:    ]
1923 // CHECK2:       .omp.reduction.case1:
1924 // CHECK2-NEXT:    [[TMP14:%.*]] = getelementptr i32, ptr [[ARRAYIDX]], i32 2
1925 // CHECK2-NEXT:    [[OMP_ARRAYCPY_ISEMPTY:%.*]] = icmp eq ptr [[ARRAYIDX]], [[TMP14]]
1926 // CHECK2-NEXT:    br i1 [[OMP_ARRAYCPY_ISEMPTY]], label [[OMP_ARRAYCPY_DONE6:%.*]], label [[OMP_ARRAYCPY_BODY:%.*]]
1927 // CHECK2:       omp.arraycpy.body:
1928 // CHECK2-NEXT:    [[OMP_ARRAYCPY_SRCELEMENTPAST:%.*]] = phi ptr [ [[A2]], [[DOTOMP_REDUCTION_CASE1]] ], [ [[OMP_ARRAYCPY_SRC_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
1929 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DESTELEMENTPAST3:%.*]] = phi ptr [ [[ARRAYIDX]], [[DOTOMP_REDUCTION_CASE1]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT4:%.*]], [[OMP_ARRAYCPY_BODY]] ]
1930 // CHECK2-NEXT:    [[TMP15:%.*]] = load i32, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST3]], align 4
1931 // CHECK2-NEXT:    [[TMP16:%.*]] = load i32, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], align 4
1932 // CHECK2-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP15]], [[TMP16]]
1933 // CHECK2-NEXT:    store i32 [[ADD]], ptr [[OMP_ARRAYCPY_DESTELEMENTPAST3]], align 4
1934 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DEST_ELEMENT4]] = getelementptr i32, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST3]], i32 1
1935 // CHECK2-NEXT:    [[OMP_ARRAYCPY_SRC_ELEMENT]] = getelementptr i32, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], i32 1
1936 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DONE5:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT4]], [[TMP14]]
1937 // CHECK2-NEXT:    br i1 [[OMP_ARRAYCPY_DONE5]], label [[OMP_ARRAYCPY_DONE6]], label [[OMP_ARRAYCPY_BODY]]
1938 // CHECK2:       omp.arraycpy.done6:
1939 // CHECK2-NEXT:    call void @__kmpc_end_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP12]], ptr @.gomp_critical_user_.reduction.var)
1940 // CHECK2-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]
1941 // CHECK2:       .omp.reduction.case2:
1942 // CHECK2-NEXT:    [[TMP17:%.*]] = getelementptr i32, ptr [[ARRAYIDX]], i32 2
1943 // CHECK2-NEXT:    [[OMP_ARRAYCPY_ISEMPTY7:%.*]] = icmp eq ptr [[ARRAYIDX]], [[TMP17]]
1944 // CHECK2-NEXT:    br i1 [[OMP_ARRAYCPY_ISEMPTY7]], label [[OMP_ARRAYCPY_DONE14:%.*]], label [[OMP_ARRAYCPY_BODY8:%.*]]
1945 // CHECK2:       omp.arraycpy.body8:
1946 // CHECK2-NEXT:    [[OMP_ARRAYCPY_SRCELEMENTPAST9:%.*]] = phi ptr [ [[A2]], [[DOTOMP_REDUCTION_CASE2]] ], [ [[OMP_ARRAYCPY_SRC_ELEMENT12:%.*]], [[OMP_ARRAYCPY_BODY8]] ]
1947 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DESTELEMENTPAST10:%.*]] = phi ptr [ [[ARRAYIDX]], [[DOTOMP_REDUCTION_CASE2]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT11:%.*]], [[OMP_ARRAYCPY_BODY8]] ]
1948 // CHECK2-NEXT:    [[TMP18:%.*]] = load i32, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST9]], align 4
1949 // CHECK2-NEXT:    [[TMP19:%.*]] = atomicrmw add ptr [[OMP_ARRAYCPY_DESTELEMENTPAST10]], i32 [[TMP18]] monotonic, align 4
1950 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DEST_ELEMENT11]] = getelementptr i32, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST10]], i32 1
1951 // CHECK2-NEXT:    [[OMP_ARRAYCPY_SRC_ELEMENT12]] = getelementptr i32, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST9]], i32 1
1952 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DONE13:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT11]], [[TMP17]]
1953 // CHECK2-NEXT:    br i1 [[OMP_ARRAYCPY_DONE13]], label [[OMP_ARRAYCPY_DONE14]], label [[OMP_ARRAYCPY_BODY8]]
1954 // CHECK2:       omp.arraycpy.done14:
1955 // CHECK2-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]
1956 // CHECK2:       .omp.reduction.default:
1957 // CHECK2-NEXT:    ret void
1958 //
1959 //
1960 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l78.omp_outlined.omp.reduction.reduction_func
1961 // CHECK2-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3]] {
1962 // CHECK2-NEXT:  entry:
1963 // CHECK2-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 4
1964 // CHECK2-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 4
1965 // CHECK2-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 4
1966 // CHECK2-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 4
1967 // CHECK2-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 4
1968 // CHECK2-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 4
1969 // CHECK2-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i32 0, i32 0
1970 // CHECK2-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 4
1971 // CHECK2-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i32 0, i32 0
1972 // CHECK2-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 4
1973 // CHECK2-NEXT:    [[TMP8:%.*]] = getelementptr i32, ptr [[TMP7]], i32 2
1974 // CHECK2-NEXT:    [[OMP_ARRAYCPY_ISEMPTY:%.*]] = icmp eq ptr [[TMP7]], [[TMP8]]
1975 // CHECK2-NEXT:    br i1 [[OMP_ARRAYCPY_ISEMPTY]], label [[OMP_ARRAYCPY_DONE2:%.*]], label [[OMP_ARRAYCPY_BODY:%.*]]
1976 // CHECK2:       omp.arraycpy.body:
1977 // CHECK2-NEXT:    [[OMP_ARRAYCPY_SRCELEMENTPAST:%.*]] = phi ptr [ [[TMP5]], [[ENTRY:%.*]] ], [ [[OMP_ARRAYCPY_SRC_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
1978 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DESTELEMENTPAST:%.*]] = phi ptr [ [[TMP7]], [[ENTRY]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
1979 // CHECK2-NEXT:    [[TMP9:%.*]] = load i32, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], align 4
1980 // CHECK2-NEXT:    [[TMP10:%.*]] = load i32, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], align 4
1981 // CHECK2-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP9]], [[TMP10]]
1982 // CHECK2-NEXT:    store i32 [[ADD]], ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], align 4
1983 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DEST_ELEMENT]] = getelementptr i32, ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], i32 1
1984 // CHECK2-NEXT:    [[OMP_ARRAYCPY_SRC_ELEMENT]] = getelementptr i32, ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], i32 1
1985 // CHECK2-NEXT:    [[OMP_ARRAYCPY_DONE:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT]], [[TMP8]]
1986 // CHECK2-NEXT:    br i1 [[OMP_ARRAYCPY_DONE]], label [[OMP_ARRAYCPY_DONE2]], label [[OMP_ARRAYCPY_BODY]]
1987 // CHECK2:       omp.arraycpy.done2:
1988 // CHECK2-NEXT:    ret void
1989 //
1990 //
1991 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l81
1992 // CHECK2-SAME: (i32 noundef [[SIZE:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR1]] {
1993 // CHECK2-NEXT:  entry:
1994 // CHECK2-NEXT:    [[SIZE_ADDR:%.*]] = alloca i32, align 4
1995 // CHECK2-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 4
1996 // CHECK2-NEXT:    [[SIZE_CASTED:%.*]] = alloca i32, align 4
1997 // CHECK2-NEXT:    store i32 [[SIZE]], ptr [[SIZE_ADDR]], align 4
1998 // CHECK2-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 4
1999 // CHECK2-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 4
2000 // CHECK2-NEXT:    [[TMP1:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
2001 // CHECK2-NEXT:    store i32 [[TMP1]], ptr [[SIZE_CASTED]], align 4
2002 // CHECK2-NEXT:    [[TMP2:%.*]] = load i32, ptr [[SIZE_CASTED]], align 4
2003 // CHECK2-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB4]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l81.omp_outlined, i32 [[TMP2]], ptr [[TMP0]])
2004 // CHECK2-NEXT:    ret void
2005 //
2006 //
2007 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l81.omp_outlined
2008 // CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[SIZE:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[A:%.*]]) #[[ATTR1]] {
2009 // CHECK2-NEXT:  entry:
2010 // CHECK2-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
2011 // CHECK2-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
2012 // CHECK2-NEXT:    [[SIZE_ADDR:%.*]] = alloca i32, align 4
2013 // CHECK2-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 4
2014 // CHECK2-NEXT:    [[A1:%.*]] = alloca i32, align 4
2015 // CHECK2-NEXT:    [[I:%.*]] = alloca i32, align 4
2016 // CHECK2-NEXT:    [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 4
2017 // CHECK2-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
2018 // CHECK2-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
2019 // CHECK2-NEXT:    store i32 [[SIZE]], ptr [[SIZE_ADDR]], align 4
2020 // CHECK2-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 4
2021 // CHECK2-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 4
2022 // CHECK2-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 3
2023 // CHECK2-NEXT:    store i32 0, ptr [[A1]], align 4
2024 // CHECK2-NEXT:    [[TMP1:%.*]] = ptrtoint ptr [[TMP0]] to i64
2025 // CHECK2-NEXT:    [[TMP2:%.*]] = ptrtoint ptr [[ARRAYIDX]] to i64
2026 // CHECK2-NEXT:    [[TMP3:%.*]] = sub i64 [[TMP1]], [[TMP2]]
2027 // CHECK2-NEXT:    [[TMP4:%.*]] = sdiv exact i64 [[TMP3]], ptrtoint (ptr getelementptr (i32, ptr null, i32 1) to i64)
2028 // CHECK2-NEXT:    [[TMP5:%.*]] = getelementptr i32, ptr [[A1]], i64 [[TMP4]]
2029 // CHECK2-NEXT:    store i32 0, ptr [[I]], align 4
2030 // CHECK2-NEXT:    br label [[FOR_COND:%.*]]
2031 // CHECK2:       for.cond:
2032 // CHECK2-NEXT:    [[TMP6:%.*]] = load i32, ptr [[I]], align 4
2033 // CHECK2-NEXT:    [[TMP7:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
2034 // CHECK2-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP6]], [[TMP7]]
2035 // CHECK2-NEXT:    br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]]
2036 // CHECK2:       for.body:
2037 // CHECK2-NEXT:    br label [[FOR_INC:%.*]]
2038 // CHECK2:       for.inc:
2039 // CHECK2-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I]], align 4
2040 // CHECK2-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP8]], 1
2041 // CHECK2-NEXT:    store i32 [[INC]], ptr [[I]], align 4
2042 // CHECK2-NEXT:    br label [[FOR_COND]], !llvm.loop [[LOOP12:![0-9]+]]
2043 // CHECK2:       for.end:
2044 // CHECK2-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 0
2045 // CHECK2-NEXT:    store ptr [[A1]], ptr [[TMP9]], align 4
2046 // CHECK2-NEXT:    [[TMP10:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
2047 // CHECK2-NEXT:    [[TMP11:%.*]] = load i32, ptr [[TMP10]], align 4
2048 // CHECK2-NEXT:    [[TMP12:%.*]] = call i32 @__kmpc_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP11]], i32 1, i32 4, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l81.omp_outlined.omp.reduction.reduction_func, ptr @.gomp_critical_user_.reduction.var)
2049 // CHECK2-NEXT:    switch i32 [[TMP12]], label [[DOTOMP_REDUCTION_DEFAULT:%.*]] [
2050 // CHECK2-NEXT:      i32 1, label [[DOTOMP_REDUCTION_CASE1:%.*]]
2051 // CHECK2-NEXT:      i32 2, label [[DOTOMP_REDUCTION_CASE2:%.*]]
2052 // CHECK2-NEXT:    ]
2053 // CHECK2:       .omp.reduction.case1:
2054 // CHECK2-NEXT:    [[TMP13:%.*]] = load i32, ptr [[ARRAYIDX]], align 4
2055 // CHECK2-NEXT:    [[TMP14:%.*]] = load i32, ptr [[A1]], align 4
2056 // CHECK2-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP13]], [[TMP14]]
2057 // CHECK2-NEXT:    store i32 [[ADD]], ptr [[ARRAYIDX]], align 4
2058 // CHECK2-NEXT:    call void @__kmpc_end_reduce_nowait(ptr @[[GLOB3]], i32 [[TMP11]], ptr @.gomp_critical_user_.reduction.var)
2059 // CHECK2-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]
2060 // CHECK2:       .omp.reduction.case2:
2061 // CHECK2-NEXT:    [[TMP15:%.*]] = load i32, ptr [[A1]], align 4
2062 // CHECK2-NEXT:    [[TMP16:%.*]] = atomicrmw add ptr [[ARRAYIDX]], i32 [[TMP15]] monotonic, align 4
2063 // CHECK2-NEXT:    br label [[DOTOMP_REDUCTION_DEFAULT]]
2064 // CHECK2:       .omp.reduction.default:
2065 // CHECK2-NEXT:    ret void
2066 //
2067 //
2068 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3sumPiiS__l81.omp_outlined.omp.reduction.reduction_func
2069 // CHECK2-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]]) #[[ATTR3]] {
2070 // CHECK2-NEXT:  entry:
2071 // CHECK2-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 4
2072 // CHECK2-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 4
2073 // CHECK2-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 4
2074 // CHECK2-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 4
2075 // CHECK2-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DOTADDR]], align 4
2076 // CHECK2-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 4
2077 // CHECK2-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP3]], i32 0, i32 0
2078 // CHECK2-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 4
2079 // CHECK2-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP2]], i32 0, i32 0
2080 // CHECK2-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 4
2081 // CHECK2-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4
2082 // CHECK2-NEXT:    [[TMP9:%.*]] = load i32, ptr [[TMP5]], align 4
2083 // CHECK2-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]
2084 // CHECK2-NEXT:    store i32 [[ADD]], ptr [[TMP7]], align 4
2085 // CHECK2-NEXT:    ret void
2086 //
2087 //
2088 // CHECK2-LABEL: define {{[^@]+}}@main
2089 // CHECK2-SAME: () #[[ATTR6:[0-9]+]] {
2090 // CHECK2-NEXT:  entry:
2091 // CHECK2-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4
2092 // CHECK2-NEXT:    [[SIZE:%.*]] = alloca i32, align 4
2093 // CHECK2-NEXT:    [[ARRAY:%.*]] = alloca ptr, align 4
2094 // CHECK2-NEXT:    [[RESULT:%.*]] = alloca i32, align 4
2095 // CHECK2-NEXT:    store i32 0, ptr [[RETVAL]], align 4
2096 // CHECK2-NEXT:    store i32 100, ptr [[SIZE]], align 4
2097 // CHECK2-NEXT:    [[CALL:%.*]] = call noalias noundef nonnull ptr @_Znaj(i32 noundef 400) #[[ATTR8:[0-9]+]]
2098 // CHECK2-NEXT:    store ptr [[CALL]], ptr [[ARRAY]], align 4
2099 // CHECK2-NEXT:    store i32 0, ptr [[RESULT]], align 4
2100 // CHECK2-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[ARRAY]], align 4
2101 // CHECK2-NEXT:    call void @_Z3sumPiiS_(ptr noundef [[TMP0]], i32 noundef 100, ptr noundef [[RESULT]])
2102 // CHECK2-NEXT:    ret i32 0
2103 //