brintos / llvm-project-archived public Read only

1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _2// Test target codegen - host bc file has to be created first.3// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc4// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK45-645// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc6// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix=CHECK45-327// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix=CHECK45-32-EX8 9// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc10// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK-6411// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc12// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix=CHECK-3213// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix=CHECK-32-EX14 15// expected-no-diagnostics16#ifndef HEADER17#define HEADER18 19// Check that the execution mode of all 2 target regions on the gpu is set to NonSPMD Mode.20 21#define N 100022 23template<typename tx>24tx ftemplate(int n) {25  tx a[N];26  short aa[N];27  tx b[10];28 29  #pragma omp target simd30  for(int i = 0; i < n; i++) {31    a[i] = 1;32  }33 34  #pragma omp target simd35  for (int i = 0; i < n; i++) {36    aa[i] += 1;37  }38 39  #pragma omp target simd40  for(int i = 0; i < 10; i++) {41    b[i] += 1;42  }43 44  #pragma omp target simd reduction(+:n)45  for(int i = 0; i < 10; i++) {46    b[i] += 1;47  }48 49  return a[0];50}51 52int bar(int n){53  int a = 0;54 55  a += ftemplate<int>(n);56 57  return a;58}59 60#endif61// CHECK45-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l2962// CHECK45-64-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR0:[0-9]+]] {63// CHECK45-64-NEXT:  entry:64// CHECK45-64-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 865// CHECK45-64-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 866// CHECK45-64-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 867// CHECK45-64-NEXT:    [[TMP:%.*]] = alloca i32, align 468// CHECK45-64-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 469// CHECK45-64-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 470// CHECK45-64-NEXT:    [[I:%.*]] = alloca i32, align 471// CHECK45-64-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 472// CHECK45-64-NEXT:    [[I3:%.*]] = alloca i32, align 473// CHECK45-64-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 874// CHECK45-64-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 875// CHECK45-64-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 876// CHECK45-64-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 877// CHECK45-64-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l29_kernel_environment, ptr [[DYN_PTR]])78// CHECK45-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -179// CHECK45-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]80// CHECK45-64:       user_code.entry:81// CHECK45-64-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 482// CHECK45-64-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 483// CHECK45-64-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 484// CHECK45-64-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 085// CHECK45-64-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 186// CHECK45-64-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 187// CHECK45-64-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 488// CHECK45-64-NEXT:    store i32 0, ptr [[I]], align 489// CHECK45-64-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 490// CHECK45-64-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]91// CHECK45-64-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]92// CHECK45-64:       simd.if.then:93// CHECK45-64-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 494// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]95// CHECK45-64:       omp.inner.for.cond:96// CHECK45-64-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24:![0-9]+]]97// CHECK45-64-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP24]]98// CHECK45-64-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 199// CHECK45-64-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]100// CHECK45-64-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]101// CHECK45-64:       omp.inner.for.body:102// CHECK45-64-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24]]103// CHECK45-64-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1104// CHECK45-64-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]105// CHECK45-64-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP24]]106// CHECK45-64-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP24]]107// CHECK45-64-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP8]] to i64108// CHECK45-64-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]109// CHECK45-64-NEXT:    store i32 1, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP24]]110// CHECK45-64-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]111// CHECK45-64:       omp.body.continue:112// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]113// CHECK45-64:       omp.inner.for.inc:114// CHECK45-64-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24]]115// CHECK45-64-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP9]], 1116// CHECK45-64-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24]]117// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP25:![0-9]+]]118// CHECK45-64:       worker.exit:119// CHECK45-64-NEXT:    ret void120// CHECK45-64:       omp.inner.for.end:121// CHECK45-64-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4122// CHECK45-64-NEXT:    [[SUB7:%.*]] = sub nsw i32 [[TMP10]], 0123// CHECK45-64-NEXT:    [[DIV8:%.*]] = sdiv i32 [[SUB7]], 1124// CHECK45-64-NEXT:    [[MUL9:%.*]] = mul nsw i32 [[DIV8]], 1125// CHECK45-64-NEXT:    [[ADD10:%.*]] = add nsw i32 0, [[MUL9]]126// CHECK45-64-NEXT:    store i32 [[ADD10]], ptr [[I3]], align 4127// CHECK45-64-NEXT:    br label [[SIMD_IF_END]]128// CHECK45-64:       simd.if.end:129// CHECK45-64-NEXT:    call void @__kmpc_target_deinit()130// CHECK45-64-NEXT:    ret void131//132//133// CHECK45-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l34134// CHECK45-64-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 2 dereferenceable(2000) [[AA:%.*]]) #[[ATTR0]] {135// CHECK45-64-NEXT:  entry:136// CHECK45-64-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8137// CHECK45-64-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 8138// CHECK45-64-NEXT:    [[AA_ADDR:%.*]] = alloca ptr, align 8139// CHECK45-64-NEXT:    [[TMP:%.*]] = alloca i32, align 4140// CHECK45-64-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4141// CHECK45-64-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4142// CHECK45-64-NEXT:    [[I:%.*]] = alloca i32, align 4143// CHECK45-64-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4144// CHECK45-64-NEXT:    [[I3:%.*]] = alloca i32, align 4145// CHECK45-64-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8146// CHECK45-64-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 8147// CHECK45-64-NEXT:    store ptr [[AA]], ptr [[AA_ADDR]], align 8148// CHECK45-64-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 8149// CHECK45-64-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l34_kernel_environment, ptr [[DYN_PTR]])150// CHECK45-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1151// CHECK45-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]152// CHECK45-64:       user_code.entry:153// CHECK45-64-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4154// CHECK45-64-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4155// CHECK45-64-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4156// CHECK45-64-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0157// CHECK45-64-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1158// CHECK45-64-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1159// CHECK45-64-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4160// CHECK45-64-NEXT:    store i32 0, ptr [[I]], align 4161// CHECK45-64-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4162// CHECK45-64-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]163// CHECK45-64-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]164// CHECK45-64:       simd.if.then:165// CHECK45-64-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4166// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]167// CHECK45-64:       omp.inner.for.cond:168// CHECK45-64-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28:![0-9]+]]169// CHECK45-64-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP28]]170// CHECK45-64-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 1171// CHECK45-64-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]172// CHECK45-64-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]173// CHECK45-64:       omp.inner.for.body:174// CHECK45-64-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28]]175// CHECK45-64-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1176// CHECK45-64-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]177// CHECK45-64-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP28]]178// CHECK45-64-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP28]]179// CHECK45-64-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP8]] to i64180// CHECK45-64-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i16], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]181// CHECK45-64-NEXT:    [[TMP9:%.*]] = load i16, ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP28]]182// CHECK45-64-NEXT:    [[CONV:%.*]] = sext i16 [[TMP9]] to i32183// CHECK45-64-NEXT:    [[ADD6:%.*]] = add nsw i32 [[CONV]], 1184// CHECK45-64-NEXT:    [[CONV7:%.*]] = trunc i32 [[ADD6]] to i16185// CHECK45-64-NEXT:    store i16 [[CONV7]], ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP28]]186// CHECK45-64-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]187// CHECK45-64:       omp.body.continue:188// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]189// CHECK45-64:       omp.inner.for.inc:190// CHECK45-64-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28]]191// CHECK45-64-NEXT:    [[ADD8:%.*]] = add nsw i32 [[TMP10]], 1192// CHECK45-64-NEXT:    store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28]]193// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP29:![0-9]+]]194// CHECK45-64:       worker.exit:195// CHECK45-64-NEXT:    ret void196// CHECK45-64:       omp.inner.for.end:197// CHECK45-64-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4198// CHECK45-64-NEXT:    [[SUB9:%.*]] = sub nsw i32 [[TMP11]], 0199// CHECK45-64-NEXT:    [[DIV10:%.*]] = sdiv i32 [[SUB9]], 1200// CHECK45-64-NEXT:    [[MUL11:%.*]] = mul nsw i32 [[DIV10]], 1201// CHECK45-64-NEXT:    [[ADD12:%.*]] = add nsw i32 0, [[MUL11]]202// CHECK45-64-NEXT:    store i32 [[ADD12]], ptr [[I3]], align 4203// CHECK45-64-NEXT:    br label [[SIMD_IF_END]]204// CHECK45-64:       simd.if.end:205// CHECK45-64-NEXT:    call void @__kmpc_target_deinit()206// CHECK45-64-NEXT:    ret void207//208//209// CHECK45-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l39210// CHECK45-64-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR0]] {211// CHECK45-64-NEXT:  entry:212// CHECK45-64-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8213// CHECK45-64-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8214// CHECK45-64-NEXT:    [[TMP:%.*]] = alloca i32, align 4215// CHECK45-64-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4216// CHECK45-64-NEXT:    [[I:%.*]] = alloca i32, align 4217// CHECK45-64-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8218// CHECK45-64-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8219// CHECK45-64-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 8220// CHECK45-64-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l39_kernel_environment, ptr [[DYN_PTR]])221// CHECK45-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1222// CHECK45-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]223// CHECK45-64:       user_code.entry:224// CHECK45-64-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4225// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]226// CHECK45-64:       omp.inner.for.cond:227// CHECK45-64-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31:![0-9]+]]228// CHECK45-64-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP2]], 10229// CHECK45-64-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]230// CHECK45-64:       omp.inner.for.body:231// CHECK45-64-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31]]232// CHECK45-64-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1233// CHECK45-64-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]234// CHECK45-64-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP31]]235// CHECK45-64-NEXT:    [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP31]]236// CHECK45-64-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP4]] to i64237// CHECK45-64-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]238// CHECK45-64-NEXT:    [[TMP5:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP31]]239// CHECK45-64-NEXT:    [[ADD1:%.*]] = add nsw i32 [[TMP5]], 1240// CHECK45-64-NEXT:    store i32 [[ADD1]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP31]]241// CHECK45-64-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]242// CHECK45-64:       omp.body.continue:243// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]244// CHECK45-64:       omp.inner.for.inc:245// CHECK45-64-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31]]246// CHECK45-64-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1247// CHECK45-64-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31]]248// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP32:![0-9]+]]249// CHECK45-64:       worker.exit:250// CHECK45-64-NEXT:    ret void251// CHECK45-64:       omp.inner.for.end:252// CHECK45-64-NEXT:    store i32 10, ptr [[I]], align 4253// CHECK45-64-NEXT:    call void @__kmpc_target_deinit()254// CHECK45-64-NEXT:    ret void255//256//257// CHECK45-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l44258// CHECK45-64-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]]) #[[ATTR0]] {259// CHECK45-64-NEXT:  entry:260// CHECK45-64-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8261// CHECK45-64-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8262// CHECK45-64-NEXT:    [[N_ADDR:%.*]] = alloca ptr, align 8263// CHECK45-64-NEXT:    [[TMP:%.*]] = alloca i32, align 4264// CHECK45-64-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4265// CHECK45-64-NEXT:    [[I:%.*]] = alloca i32, align 4266// CHECK45-64-NEXT:    [[N1:%.*]] = alloca i32, align 4267// CHECK45-64-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8268// CHECK45-64-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8269// CHECK45-64-NEXT:    store ptr [[N]], ptr [[N_ADDR]], align 8270// CHECK45-64-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 8271// CHECK45-64-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[N_ADDR]], align 8272// CHECK45-64-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l44_kernel_environment, ptr [[DYN_PTR]])273// CHECK45-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -1274// CHECK45-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]275// CHECK45-64:       user_code.entry:276// CHECK45-64-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4277// CHECK45-64-NEXT:    store i32 0, ptr [[N1]], align 4278// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]279// CHECK45-64:       omp.inner.for.cond:280// CHECK45-64-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34:![0-9]+]]281// CHECK45-64-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP3]], 10282// CHECK45-64-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]283// CHECK45-64:       omp.inner.for.body:284// CHECK45-64-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34]]285// CHECK45-64-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP4]], 1286// CHECK45-64-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]287// CHECK45-64-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP34]]288// CHECK45-64-NEXT:    [[TMP5:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP34]]289// CHECK45-64-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP5]] to i64290// CHECK45-64-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]291// CHECK45-64-NEXT:    [[TMP6:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP34]]292// CHECK45-64-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1293// CHECK45-64-NEXT:    store i32 [[ADD2]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP34]]294// CHECK45-64-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]295// CHECK45-64:       omp.body.continue:296// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]297// CHECK45-64:       omp.inner.for.inc:298// CHECK45-64-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34]]299// CHECK45-64-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP7]], 1300// CHECK45-64-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34]]301// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP35:![0-9]+]]302// CHECK45-64:       worker.exit:303// CHECK45-64-NEXT:    ret void304// CHECK45-64:       omp.inner.for.end:305// CHECK45-64-NEXT:    store i32 10, ptr [[I]], align 4306// CHECK45-64-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP1]], align 4307// CHECK45-64-NEXT:    [[TMP9:%.*]] = load i32, ptr [[N1]], align 4308// CHECK45-64-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]309// CHECK45-64-NEXT:    store i32 [[ADD4]], ptr [[TMP1]], align 4310// CHECK45-64-NEXT:    call void @__kmpc_target_deinit()311// CHECK45-64-NEXT:    ret void312//313//314// CHECK45-32-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l29315// CHECK45-32-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR0:[0-9]+]] {316// CHECK45-32-NEXT:  entry:317// CHECK45-32-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4318// CHECK45-32-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4319// CHECK45-32-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 4320// CHECK45-32-NEXT:    [[TMP:%.*]] = alloca i32, align 4321// CHECK45-32-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4322// CHECK45-32-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4323// CHECK45-32-NEXT:    [[I:%.*]] = alloca i32, align 4324// CHECK45-32-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4325// CHECK45-32-NEXT:    [[I3:%.*]] = alloca i32, align 4326// CHECK45-32-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4327// CHECK45-32-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4328// CHECK45-32-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 4329// CHECK45-32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 4330// CHECK45-32-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l29_kernel_environment, ptr [[DYN_PTR]])331// CHECK45-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1332// CHECK45-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]333// CHECK45-32:       user_code.entry:334// CHECK45-32-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4335// CHECK45-32-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4336// CHECK45-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4337// CHECK45-32-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0338// CHECK45-32-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1339// CHECK45-32-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1340// CHECK45-32-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4341// CHECK45-32-NEXT:    store i32 0, ptr [[I]], align 4342// CHECK45-32-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4343// CHECK45-32-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]344// CHECK45-32-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]345// CHECK45-32:       simd.if.then:346// CHECK45-32-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4347// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]348// CHECK45-32:       omp.inner.for.cond:349// CHECK45-32-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24:![0-9]+]]350// CHECK45-32-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP24]]351// CHECK45-32-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 1352// CHECK45-32-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]353// CHECK45-32-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]354// CHECK45-32:       omp.inner.for.body:355// CHECK45-32-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24]]356// CHECK45-32-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1357// CHECK45-32-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]358// CHECK45-32-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP24]]359// CHECK45-32-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP24]]360// CHECK45-32-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i32 0, i32 [[TMP8]]361// CHECK45-32-NEXT:    store i32 1, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP24]]362// CHECK45-32-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]363// CHECK45-32:       omp.body.continue:364// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]365// CHECK45-32:       omp.inner.for.inc:366// CHECK45-32-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24]]367// CHECK45-32-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP9]], 1368// CHECK45-32-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24]]369// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP25:![0-9]+]]370// CHECK45-32:       worker.exit:371// CHECK45-32-NEXT:    ret void372// CHECK45-32:       omp.inner.for.end:373// CHECK45-32-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4374// CHECK45-32-NEXT:    [[SUB7:%.*]] = sub nsw i32 [[TMP10]], 0375// CHECK45-32-NEXT:    [[DIV8:%.*]] = sdiv i32 [[SUB7]], 1376// CHECK45-32-NEXT:    [[MUL9:%.*]] = mul nsw i32 [[DIV8]], 1377// CHECK45-32-NEXT:    [[ADD10:%.*]] = add nsw i32 0, [[MUL9]]378// CHECK45-32-NEXT:    store i32 [[ADD10]], ptr [[I3]], align 4379// CHECK45-32-NEXT:    br label [[SIMD_IF_END]]380// CHECK45-32:       simd.if.end:381// CHECK45-32-NEXT:    call void @__kmpc_target_deinit()382// CHECK45-32-NEXT:    ret void383//384//385// CHECK45-32-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l34386// CHECK45-32-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 2 dereferenceable(2000) [[AA:%.*]]) #[[ATTR0]] {387// CHECK45-32-NEXT:  entry:388// CHECK45-32-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4389// CHECK45-32-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4390// CHECK45-32-NEXT:    [[AA_ADDR:%.*]] = alloca ptr, align 4391// CHECK45-32-NEXT:    [[TMP:%.*]] = alloca i32, align 4392// CHECK45-32-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4393// CHECK45-32-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4394// CHECK45-32-NEXT:    [[I:%.*]] = alloca i32, align 4395// CHECK45-32-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4396// CHECK45-32-NEXT:    [[I3:%.*]] = alloca i32, align 4397// CHECK45-32-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4398// CHECK45-32-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4399// CHECK45-32-NEXT:    store ptr [[AA]], ptr [[AA_ADDR]], align 4400// CHECK45-32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 4401// CHECK45-32-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l34_kernel_environment, ptr [[DYN_PTR]])402// CHECK45-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1403// CHECK45-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]404// CHECK45-32:       user_code.entry:405// CHECK45-32-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4406// CHECK45-32-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4407// CHECK45-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4408// CHECK45-32-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0409// CHECK45-32-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1410// CHECK45-32-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1411// CHECK45-32-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4412// CHECK45-32-NEXT:    store i32 0, ptr [[I]], align 4413// CHECK45-32-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4414// CHECK45-32-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]415// CHECK45-32-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]416// CHECK45-32:       simd.if.then:417// CHECK45-32-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4418// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]419// CHECK45-32:       omp.inner.for.cond:420// CHECK45-32-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28:![0-9]+]]421// CHECK45-32-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP28]]422// CHECK45-32-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 1423// CHECK45-32-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]424// CHECK45-32-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]425// CHECK45-32:       omp.inner.for.body:426// CHECK45-32-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28]]427// CHECK45-32-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1428// CHECK45-32-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]429// CHECK45-32-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP28]]430// CHECK45-32-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP28]]431// CHECK45-32-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i16], ptr [[TMP0]], i32 0, i32 [[TMP8]]432// CHECK45-32-NEXT:    [[TMP9:%.*]] = load i16, ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP28]]433// CHECK45-32-NEXT:    [[CONV:%.*]] = sext i16 [[TMP9]] to i32434// CHECK45-32-NEXT:    [[ADD6:%.*]] = add nsw i32 [[CONV]], 1435// CHECK45-32-NEXT:    [[CONV7:%.*]] = trunc i32 [[ADD6]] to i16436// CHECK45-32-NEXT:    store i16 [[CONV7]], ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP28]]437// CHECK45-32-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]438// CHECK45-32:       omp.body.continue:439// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]440// CHECK45-32:       omp.inner.for.inc:441// CHECK45-32-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28]]442// CHECK45-32-NEXT:    [[ADD8:%.*]] = add nsw i32 [[TMP10]], 1443// CHECK45-32-NEXT:    store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28]]444// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP29:![0-9]+]]445// CHECK45-32:       worker.exit:446// CHECK45-32-NEXT:    ret void447// CHECK45-32:       omp.inner.for.end:448// CHECK45-32-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4449// CHECK45-32-NEXT:    [[SUB9:%.*]] = sub nsw i32 [[TMP11]], 0450// CHECK45-32-NEXT:    [[DIV10:%.*]] = sdiv i32 [[SUB9]], 1451// CHECK45-32-NEXT:    [[MUL11:%.*]] = mul nsw i32 [[DIV10]], 1452// CHECK45-32-NEXT:    [[ADD12:%.*]] = add nsw i32 0, [[MUL11]]453// CHECK45-32-NEXT:    store i32 [[ADD12]], ptr [[I3]], align 4454// CHECK45-32-NEXT:    br label [[SIMD_IF_END]]455// CHECK45-32:       simd.if.end:456// CHECK45-32-NEXT:    call void @__kmpc_target_deinit()457// CHECK45-32-NEXT:    ret void458//459//460// CHECK45-32-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l39461// CHECK45-32-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR0]] {462// CHECK45-32-NEXT:  entry:463// CHECK45-32-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4464// CHECK45-32-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4465// CHECK45-32-NEXT:    [[TMP:%.*]] = alloca i32, align 4466// CHECK45-32-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4467// CHECK45-32-NEXT:    [[I:%.*]] = alloca i32, align 4468// CHECK45-32-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4469// CHECK45-32-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4470// CHECK45-32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4471// CHECK45-32-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l39_kernel_environment, ptr [[DYN_PTR]])472// CHECK45-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1473// CHECK45-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]474// CHECK45-32:       user_code.entry:475// CHECK45-32-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4476// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]477// CHECK45-32:       omp.inner.for.cond:478// CHECK45-32-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31:![0-9]+]]479// CHECK45-32-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP2]], 10480// CHECK45-32-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]481// CHECK45-32:       omp.inner.for.body:482// CHECK45-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31]]483// CHECK45-32-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1484// CHECK45-32-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]485// CHECK45-32-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP31]]486// CHECK45-32-NEXT:    [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP31]]487// CHECK45-32-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP4]]488// CHECK45-32-NEXT:    [[TMP5:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP31]]489// CHECK45-32-NEXT:    [[ADD1:%.*]] = add nsw i32 [[TMP5]], 1490// CHECK45-32-NEXT:    store i32 [[ADD1]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP31]]491// CHECK45-32-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]492// CHECK45-32:       omp.body.continue:493// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]494// CHECK45-32:       omp.inner.for.inc:495// CHECK45-32-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31]]496// CHECK45-32-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1497// CHECK45-32-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31]]498// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP32:![0-9]+]]499// CHECK45-32:       worker.exit:500// CHECK45-32-NEXT:    ret void501// CHECK45-32:       omp.inner.for.end:502// CHECK45-32-NEXT:    store i32 10, ptr [[I]], align 4503// CHECK45-32-NEXT:    call void @__kmpc_target_deinit()504// CHECK45-32-NEXT:    ret void505//506//507// CHECK45-32-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l44508// CHECK45-32-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]]) #[[ATTR0]] {509// CHECK45-32-NEXT:  entry:510// CHECK45-32-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4511// CHECK45-32-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4512// CHECK45-32-NEXT:    [[N_ADDR:%.*]] = alloca ptr, align 4513// CHECK45-32-NEXT:    [[TMP:%.*]] = alloca i32, align 4514// CHECK45-32-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4515// CHECK45-32-NEXT:    [[I:%.*]] = alloca i32, align 4516// CHECK45-32-NEXT:    [[N1:%.*]] = alloca i32, align 4517// CHECK45-32-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4518// CHECK45-32-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4519// CHECK45-32-NEXT:    store ptr [[N]], ptr [[N_ADDR]], align 4520// CHECK45-32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4521// CHECK45-32-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[N_ADDR]], align 4522// CHECK45-32-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l44_kernel_environment, ptr [[DYN_PTR]])523// CHECK45-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -1524// CHECK45-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]525// CHECK45-32:       user_code.entry:526// CHECK45-32-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4527// CHECK45-32-NEXT:    store i32 0, ptr [[N1]], align 4528// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]529// CHECK45-32:       omp.inner.for.cond:530// CHECK45-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34:![0-9]+]]531// CHECK45-32-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP3]], 10532// CHECK45-32-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]533// CHECK45-32:       omp.inner.for.body:534// CHECK45-32-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34]]535// CHECK45-32-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP4]], 1536// CHECK45-32-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]537// CHECK45-32-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP34]]538// CHECK45-32-NEXT:    [[TMP5:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP34]]539// CHECK45-32-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP5]]540// CHECK45-32-NEXT:    [[TMP6:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP34]]541// CHECK45-32-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1542// CHECK45-32-NEXT:    store i32 [[ADD2]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP34]]543// CHECK45-32-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]544// CHECK45-32:       omp.body.continue:545// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]546// CHECK45-32:       omp.inner.for.inc:547// CHECK45-32-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34]]548// CHECK45-32-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP7]], 1549// CHECK45-32-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34]]550// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP35:![0-9]+]]551// CHECK45-32:       worker.exit:552// CHECK45-32-NEXT:    ret void553// CHECK45-32:       omp.inner.for.end:554// CHECK45-32-NEXT:    store i32 10, ptr [[I]], align 4555// CHECK45-32-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP1]], align 4556// CHECK45-32-NEXT:    [[TMP9:%.*]] = load i32, ptr [[N1]], align 4557// CHECK45-32-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]558// CHECK45-32-NEXT:    store i32 [[ADD4]], ptr [[TMP1]], align 4559// CHECK45-32-NEXT:    call void @__kmpc_target_deinit()560// CHECK45-32-NEXT:    ret void561//562//563// CHECK45-32-EX-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l29564// CHECK45-32-EX-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR0:[0-9]+]] {565// CHECK45-32-EX-NEXT:  entry:566// CHECK45-32-EX-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4567// CHECK45-32-EX-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4568// CHECK45-32-EX-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 4569// CHECK45-32-EX-NEXT:    [[TMP:%.*]] = alloca i32, align 4570// CHECK45-32-EX-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4571// CHECK45-32-EX-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4572// CHECK45-32-EX-NEXT:    [[I:%.*]] = alloca i32, align 4573// CHECK45-32-EX-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4574// CHECK45-32-EX-NEXT:    [[I3:%.*]] = alloca i32, align 4575// CHECK45-32-EX-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4576// CHECK45-32-EX-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4577// CHECK45-32-EX-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 4578// CHECK45-32-EX-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 4579// CHECK45-32-EX-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l29_kernel_environment, ptr [[DYN_PTR]])580// CHECK45-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1581// CHECK45-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]582// CHECK45-32-EX:       user_code.entry:583// CHECK45-32-EX-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4584// CHECK45-32-EX-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4585// CHECK45-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4586// CHECK45-32-EX-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0587// CHECK45-32-EX-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1588// CHECK45-32-EX-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1589// CHECK45-32-EX-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4590// CHECK45-32-EX-NEXT:    store i32 0, ptr [[I]], align 4591// CHECK45-32-EX-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4592// CHECK45-32-EX-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]593// CHECK45-32-EX-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]594// CHECK45-32-EX:       simd.if.then:595// CHECK45-32-EX-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4596// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]597// CHECK45-32-EX:       omp.inner.for.cond:598// CHECK45-32-EX-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24:![0-9]+]]599// CHECK45-32-EX-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP24]]600// CHECK45-32-EX-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 1601// CHECK45-32-EX-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]602// CHECK45-32-EX-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]603// CHECK45-32-EX:       omp.inner.for.body:604// CHECK45-32-EX-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24]]605// CHECK45-32-EX-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1606// CHECK45-32-EX-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]607// CHECK45-32-EX-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP24]]608// CHECK45-32-EX-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP24]]609// CHECK45-32-EX-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i32 0, i32 [[TMP8]]610// CHECK45-32-EX-NEXT:    store i32 1, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP24]]611// CHECK45-32-EX-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]612// CHECK45-32-EX:       omp.body.continue:613// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]614// CHECK45-32-EX:       omp.inner.for.inc:615// CHECK45-32-EX-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24]]616// CHECK45-32-EX-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP9]], 1617// CHECK45-32-EX-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24]]618// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP25:![0-9]+]]619// CHECK45-32-EX:       worker.exit:620// CHECK45-32-EX-NEXT:    ret void621// CHECK45-32-EX:       omp.inner.for.end:622// CHECK45-32-EX-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4623// CHECK45-32-EX-NEXT:    [[SUB7:%.*]] = sub nsw i32 [[TMP10]], 0624// CHECK45-32-EX-NEXT:    [[DIV8:%.*]] = sdiv i32 [[SUB7]], 1625// CHECK45-32-EX-NEXT:    [[MUL9:%.*]] = mul nsw i32 [[DIV8]], 1626// CHECK45-32-EX-NEXT:    [[ADD10:%.*]] = add nsw i32 0, [[MUL9]]627// CHECK45-32-EX-NEXT:    store i32 [[ADD10]], ptr [[I3]], align 4628// CHECK45-32-EX-NEXT:    br label [[SIMD_IF_END]]629// CHECK45-32-EX:       simd.if.end:630// CHECK45-32-EX-NEXT:    call void @__kmpc_target_deinit()631// CHECK45-32-EX-NEXT:    ret void632//633//634// CHECK45-32-EX-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l34635// CHECK45-32-EX-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 2 dereferenceable(2000) [[AA:%.*]]) #[[ATTR0]] {636// CHECK45-32-EX-NEXT:  entry:637// CHECK45-32-EX-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4638// CHECK45-32-EX-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4639// CHECK45-32-EX-NEXT:    [[AA_ADDR:%.*]] = alloca ptr, align 4640// CHECK45-32-EX-NEXT:    [[TMP:%.*]] = alloca i32, align 4641// CHECK45-32-EX-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4642// CHECK45-32-EX-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4643// CHECK45-32-EX-NEXT:    [[I:%.*]] = alloca i32, align 4644// CHECK45-32-EX-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4645// CHECK45-32-EX-NEXT:    [[I3:%.*]] = alloca i32, align 4646// CHECK45-32-EX-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4647// CHECK45-32-EX-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4648// CHECK45-32-EX-NEXT:    store ptr [[AA]], ptr [[AA_ADDR]], align 4649// CHECK45-32-EX-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 4650// CHECK45-32-EX-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l34_kernel_environment, ptr [[DYN_PTR]])651// CHECK45-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1652// CHECK45-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]653// CHECK45-32-EX:       user_code.entry:654// CHECK45-32-EX-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4655// CHECK45-32-EX-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4656// CHECK45-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4657// CHECK45-32-EX-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0658// CHECK45-32-EX-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1659// CHECK45-32-EX-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1660// CHECK45-32-EX-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4661// CHECK45-32-EX-NEXT:    store i32 0, ptr [[I]], align 4662// CHECK45-32-EX-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4663// CHECK45-32-EX-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]664// CHECK45-32-EX-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]665// CHECK45-32-EX:       simd.if.then:666// CHECK45-32-EX-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4667// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]668// CHECK45-32-EX:       omp.inner.for.cond:669// CHECK45-32-EX-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28:![0-9]+]]670// CHECK45-32-EX-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP28]]671// CHECK45-32-EX-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 1672// CHECK45-32-EX-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]673// CHECK45-32-EX-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]674// CHECK45-32-EX:       omp.inner.for.body:675// CHECK45-32-EX-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28]]676// CHECK45-32-EX-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1677// CHECK45-32-EX-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]678// CHECK45-32-EX-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP28]]679// CHECK45-32-EX-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP28]]680// CHECK45-32-EX-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i16], ptr [[TMP0]], i32 0, i32 [[TMP8]]681// CHECK45-32-EX-NEXT:    [[TMP9:%.*]] = load i16, ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP28]]682// CHECK45-32-EX-NEXT:    [[CONV:%.*]] = sext i16 [[TMP9]] to i32683// CHECK45-32-EX-NEXT:    [[ADD6:%.*]] = add nsw i32 [[CONV]], 1684// CHECK45-32-EX-NEXT:    [[CONV7:%.*]] = trunc i32 [[ADD6]] to i16685// CHECK45-32-EX-NEXT:    store i16 [[CONV7]], ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP28]]686// CHECK45-32-EX-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]687// CHECK45-32-EX:       omp.body.continue:688// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]689// CHECK45-32-EX:       omp.inner.for.inc:690// CHECK45-32-EX-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28]]691// CHECK45-32-EX-NEXT:    [[ADD8:%.*]] = add nsw i32 [[TMP10]], 1692// CHECK45-32-EX-NEXT:    store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28]]693// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP29:![0-9]+]]694// CHECK45-32-EX:       worker.exit:695// CHECK45-32-EX-NEXT:    ret void696// CHECK45-32-EX:       omp.inner.for.end:697// CHECK45-32-EX-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4698// CHECK45-32-EX-NEXT:    [[SUB9:%.*]] = sub nsw i32 [[TMP11]], 0699// CHECK45-32-EX-NEXT:    [[DIV10:%.*]] = sdiv i32 [[SUB9]], 1700// CHECK45-32-EX-NEXT:    [[MUL11:%.*]] = mul nsw i32 [[DIV10]], 1701// CHECK45-32-EX-NEXT:    [[ADD12:%.*]] = add nsw i32 0, [[MUL11]]702// CHECK45-32-EX-NEXT:    store i32 [[ADD12]], ptr [[I3]], align 4703// CHECK45-32-EX-NEXT:    br label [[SIMD_IF_END]]704// CHECK45-32-EX:       simd.if.end:705// CHECK45-32-EX-NEXT:    call void @__kmpc_target_deinit()706// CHECK45-32-EX-NEXT:    ret void707//708//709// CHECK45-32-EX-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l39710// CHECK45-32-EX-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR0]] {711// CHECK45-32-EX-NEXT:  entry:712// CHECK45-32-EX-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4713// CHECK45-32-EX-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4714// CHECK45-32-EX-NEXT:    [[TMP:%.*]] = alloca i32, align 4715// CHECK45-32-EX-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4716// CHECK45-32-EX-NEXT:    [[I:%.*]] = alloca i32, align 4717// CHECK45-32-EX-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4718// CHECK45-32-EX-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4719// CHECK45-32-EX-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4720// CHECK45-32-EX-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l39_kernel_environment, ptr [[DYN_PTR]])721// CHECK45-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1722// CHECK45-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]723// CHECK45-32-EX:       user_code.entry:724// CHECK45-32-EX-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4725// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]726// CHECK45-32-EX:       omp.inner.for.cond:727// CHECK45-32-EX-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31:![0-9]+]]728// CHECK45-32-EX-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP2]], 10729// CHECK45-32-EX-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]730// CHECK45-32-EX:       omp.inner.for.body:731// CHECK45-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31]]732// CHECK45-32-EX-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1733// CHECK45-32-EX-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]734// CHECK45-32-EX-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP31]]735// CHECK45-32-EX-NEXT:    [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP31]]736// CHECK45-32-EX-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP4]]737// CHECK45-32-EX-NEXT:    [[TMP5:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP31]]738// CHECK45-32-EX-NEXT:    [[ADD1:%.*]] = add nsw i32 [[TMP5]], 1739// CHECK45-32-EX-NEXT:    store i32 [[ADD1]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP31]]740// CHECK45-32-EX-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]741// CHECK45-32-EX:       omp.body.continue:742// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]743// CHECK45-32-EX:       omp.inner.for.inc:744// CHECK45-32-EX-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31]]745// CHECK45-32-EX-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1746// CHECK45-32-EX-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31]]747// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP32:![0-9]+]]748// CHECK45-32-EX:       worker.exit:749// CHECK45-32-EX-NEXT:    ret void750// CHECK45-32-EX:       omp.inner.for.end:751// CHECK45-32-EX-NEXT:    store i32 10, ptr [[I]], align 4752// CHECK45-32-EX-NEXT:    call void @__kmpc_target_deinit()753// CHECK45-32-EX-NEXT:    ret void754//755//756// CHECK45-32-EX-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l44757// CHECK45-32-EX-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]]) #[[ATTR0]] {758// CHECK45-32-EX-NEXT:  entry:759// CHECK45-32-EX-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4760// CHECK45-32-EX-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4761// CHECK45-32-EX-NEXT:    [[N_ADDR:%.*]] = alloca ptr, align 4762// CHECK45-32-EX-NEXT:    [[TMP:%.*]] = alloca i32, align 4763// CHECK45-32-EX-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4764// CHECK45-32-EX-NEXT:    [[I:%.*]] = alloca i32, align 4765// CHECK45-32-EX-NEXT:    [[N1:%.*]] = alloca i32, align 4766// CHECK45-32-EX-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4767// CHECK45-32-EX-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4768// CHECK45-32-EX-NEXT:    store ptr [[N]], ptr [[N_ADDR]], align 4769// CHECK45-32-EX-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4770// CHECK45-32-EX-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[N_ADDR]], align 4771// CHECK45-32-EX-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l44_kernel_environment, ptr [[DYN_PTR]])772// CHECK45-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -1773// CHECK45-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]774// CHECK45-32-EX:       user_code.entry:775// CHECK45-32-EX-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4776// CHECK45-32-EX-NEXT:    store i32 0, ptr [[N1]], align 4777// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]778// CHECK45-32-EX:       omp.inner.for.cond:779// CHECK45-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34:![0-9]+]]780// CHECK45-32-EX-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP3]], 10781// CHECK45-32-EX-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]782// CHECK45-32-EX:       omp.inner.for.body:783// CHECK45-32-EX-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34]]784// CHECK45-32-EX-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP4]], 1785// CHECK45-32-EX-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]786// CHECK45-32-EX-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP34]]787// CHECK45-32-EX-NEXT:    [[TMP5:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP34]]788// CHECK45-32-EX-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP5]]789// CHECK45-32-EX-NEXT:    [[TMP6:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP34]]790// CHECK45-32-EX-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1791// CHECK45-32-EX-NEXT:    store i32 [[ADD2]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP34]]792// CHECK45-32-EX-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]793// CHECK45-32-EX:       omp.body.continue:794// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]795// CHECK45-32-EX:       omp.inner.for.inc:796// CHECK45-32-EX-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34]]797// CHECK45-32-EX-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP7]], 1798// CHECK45-32-EX-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34]]799// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP35:![0-9]+]]800// CHECK45-32-EX:       worker.exit:801// CHECK45-32-EX-NEXT:    ret void802// CHECK45-32-EX:       omp.inner.for.end:803// CHECK45-32-EX-NEXT:    store i32 10, ptr [[I]], align 4804// CHECK45-32-EX-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP1]], align 4805// CHECK45-32-EX-NEXT:    [[TMP9:%.*]] = load i32, ptr [[N1]], align 4806// CHECK45-32-EX-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]807// CHECK45-32-EX-NEXT:    store i32 [[ADD4]], ptr [[TMP1]], align 4808// CHECK45-32-EX-NEXT:    call void @__kmpc_target_deinit()809// CHECK45-32-EX-NEXT:    ret void810//811//812// CHECK-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l29813// CHECK-64-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR0:[0-9]+]] {814// CHECK-64-NEXT:  entry:815// CHECK-64-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8816// CHECK-64-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 8817// CHECK-64-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8818// CHECK-64-NEXT:    [[TMP:%.*]] = alloca i32, align 4819// CHECK-64-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4820// CHECK-64-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4821// CHECK-64-NEXT:    [[I:%.*]] = alloca i32, align 4822// CHECK-64-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4823// CHECK-64-NEXT:    [[I3:%.*]] = alloca i32, align 4824// CHECK-64-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8825// CHECK-64-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 8826// CHECK-64-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8827// CHECK-64-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8828// CHECK-64-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l29_kernel_environment, ptr [[DYN_PTR]])829// CHECK-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1830// CHECK-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]831// CHECK-64:       user_code.entry:832// CHECK-64-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4833// CHECK-64-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4834// CHECK-64-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4835// CHECK-64-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0836// CHECK-64-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1837// CHECK-64-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1838// CHECK-64-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4839// CHECK-64-NEXT:    store i32 0, ptr [[I]], align 4840// CHECK-64-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4841// CHECK-64-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]842// CHECK-64-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]843// CHECK-64:       simd.if.then:844// CHECK-64-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4845// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]846// CHECK-64:       omp.inner.for.cond:847// CHECK-64-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24:![0-9]+]]848// CHECK-64-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP24]]849// CHECK-64-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 1850// CHECK-64-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]851// CHECK-64-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]852// CHECK-64:       omp.inner.for.body:853// CHECK-64-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24]]854// CHECK-64-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1855// CHECK-64-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]856// CHECK-64-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP24]]857// CHECK-64-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP24]]858// CHECK-64-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP8]] to i64859// CHECK-64-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]860// CHECK-64-NEXT:    store i32 1, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP24]]861// CHECK-64-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]862// CHECK-64:       omp.body.continue:863// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]864// CHECK-64:       omp.inner.for.inc:865// CHECK-64-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24]]866// CHECK-64-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP9]], 1867// CHECK-64-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24]]868// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP25:![0-9]+]]869// CHECK-64:       worker.exit:870// CHECK-64-NEXT:    ret void871// CHECK-64:       omp.inner.for.end:872// CHECK-64-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4873// CHECK-64-NEXT:    [[SUB7:%.*]] = sub nsw i32 [[TMP10]], 0874// CHECK-64-NEXT:    [[DIV8:%.*]] = sdiv i32 [[SUB7]], 1875// CHECK-64-NEXT:    [[MUL9:%.*]] = mul nsw i32 [[DIV8]], 1876// CHECK-64-NEXT:    [[ADD10:%.*]] = add nsw i32 0, [[MUL9]]877// CHECK-64-NEXT:    store i32 [[ADD10]], ptr [[I3]], align 4878// CHECK-64-NEXT:    br label [[SIMD_IF_END]]879// CHECK-64:       simd.if.end:880// CHECK-64-NEXT:    call void @__kmpc_target_deinit()881// CHECK-64-NEXT:    ret void882//883//884// CHECK-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l34885// CHECK-64-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[N:%.*]], ptr noundef nonnull align 2 dereferenceable(2000) [[AA:%.*]]) #[[ATTR0]] {886// CHECK-64-NEXT:  entry:887// CHECK-64-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8888// CHECK-64-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 8889// CHECK-64-NEXT:    [[AA_ADDR:%.*]] = alloca ptr, align 8890// CHECK-64-NEXT:    [[TMP:%.*]] = alloca i32, align 4891// CHECK-64-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4892// CHECK-64-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4893// CHECK-64-NEXT:    [[I:%.*]] = alloca i32, align 4894// CHECK-64-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4895// CHECK-64-NEXT:    [[I3:%.*]] = alloca i32, align 4896// CHECK-64-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8897// CHECK-64-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 8898// CHECK-64-NEXT:    store ptr [[AA]], ptr [[AA_ADDR]], align 8899// CHECK-64-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 8900// CHECK-64-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l34_kernel_environment, ptr [[DYN_PTR]])901// CHECK-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1902// CHECK-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]903// CHECK-64:       user_code.entry:904// CHECK-64-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4905// CHECK-64-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4906// CHECK-64-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4907// CHECK-64-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0908// CHECK-64-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1909// CHECK-64-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1910// CHECK-64-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4911// CHECK-64-NEXT:    store i32 0, ptr [[I]], align 4912// CHECK-64-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4913// CHECK-64-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]914// CHECK-64-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]915// CHECK-64:       simd.if.then:916// CHECK-64-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4917// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]918// CHECK-64:       omp.inner.for.cond:919// CHECK-64-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28:![0-9]+]]920// CHECK-64-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP28]]921// CHECK-64-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 1922// CHECK-64-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]923// CHECK-64-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]924// CHECK-64:       omp.inner.for.body:925// CHECK-64-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28]]926// CHECK-64-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1927// CHECK-64-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]928// CHECK-64-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP28]]929// CHECK-64-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP28]]930// CHECK-64-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP8]] to i64931// CHECK-64-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i16], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]932// CHECK-64-NEXT:    [[TMP9:%.*]] = load i16, ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP28]]933// CHECK-64-NEXT:    [[CONV:%.*]] = sext i16 [[TMP9]] to i32934// CHECK-64-NEXT:    [[ADD6:%.*]] = add nsw i32 [[CONV]], 1935// CHECK-64-NEXT:    [[CONV7:%.*]] = trunc i32 [[ADD6]] to i16936// CHECK-64-NEXT:    store i16 [[CONV7]], ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP28]]937// CHECK-64-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]938// CHECK-64:       omp.body.continue:939// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]940// CHECK-64:       omp.inner.for.inc:941// CHECK-64-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28]]942// CHECK-64-NEXT:    [[ADD8:%.*]] = add nsw i32 [[TMP10]], 1943// CHECK-64-NEXT:    store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28]]944// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP29:![0-9]+]]945// CHECK-64:       worker.exit:946// CHECK-64-NEXT:    ret void947// CHECK-64:       omp.inner.for.end:948// CHECK-64-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4949// CHECK-64-NEXT:    [[SUB9:%.*]] = sub nsw i32 [[TMP11]], 0950// CHECK-64-NEXT:    [[DIV10:%.*]] = sdiv i32 [[SUB9]], 1951// CHECK-64-NEXT:    [[MUL11:%.*]] = mul nsw i32 [[DIV10]], 1952// CHECK-64-NEXT:    [[ADD12:%.*]] = add nsw i32 0, [[MUL11]]953// CHECK-64-NEXT:    store i32 [[ADD12]], ptr [[I3]], align 4954// CHECK-64-NEXT:    br label [[SIMD_IF_END]]955// CHECK-64:       simd.if.end:956// CHECK-64-NEXT:    call void @__kmpc_target_deinit()957// CHECK-64-NEXT:    ret void958//959//960// CHECK-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l39961// CHECK-64-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR0]] {962// CHECK-64-NEXT:  entry:963// CHECK-64-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8964// CHECK-64-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8965// CHECK-64-NEXT:    [[TMP:%.*]] = alloca i32, align 4966// CHECK-64-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4967// CHECK-64-NEXT:    [[I:%.*]] = alloca i32, align 4968// CHECK-64-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8969// CHECK-64-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8970// CHECK-64-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 8971// CHECK-64-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l39_kernel_environment, ptr [[DYN_PTR]])972// CHECK-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1973// CHECK-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]974// CHECK-64:       user_code.entry:975// CHECK-64-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4976// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]977// CHECK-64:       omp.inner.for.cond:978// CHECK-64-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31:![0-9]+]]979// CHECK-64-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP2]], 10980// CHECK-64-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]981// CHECK-64:       omp.inner.for.body:982// CHECK-64-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31]]983// CHECK-64-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1984// CHECK-64-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]985// CHECK-64-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP31]]986// CHECK-64-NEXT:    [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP31]]987// CHECK-64-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP4]] to i64988// CHECK-64-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]989// CHECK-64-NEXT:    [[TMP5:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP31]]990// CHECK-64-NEXT:    [[ADD1:%.*]] = add nsw i32 [[TMP5]], 1991// CHECK-64-NEXT:    store i32 [[ADD1]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP31]]992// CHECK-64-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]993// CHECK-64:       omp.body.continue:994// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]995// CHECK-64:       omp.inner.for.inc:996// CHECK-64-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31]]997// CHECK-64-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1998// CHECK-64-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31]]999// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP32:![0-9]+]]1000// CHECK-64:       worker.exit:1001// CHECK-64-NEXT:    ret void1002// CHECK-64:       omp.inner.for.end:1003// CHECK-64-NEXT:    store i32 10, ptr [[I]], align 41004// CHECK-64-NEXT:    call void @__kmpc_target_deinit()1005// CHECK-64-NEXT:    ret void1006//1007//1008// CHECK-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l441009// CHECK-64-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]]) #[[ATTR0]] {1010// CHECK-64-NEXT:  entry:1011// CHECK-64-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81012// CHECK-64-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 81013// CHECK-64-NEXT:    [[N_ADDR:%.*]] = alloca ptr, align 81014// CHECK-64-NEXT:    [[TMP:%.*]] = alloca i32, align 41015// CHECK-64-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41016// CHECK-64-NEXT:    [[I:%.*]] = alloca i32, align 41017// CHECK-64-NEXT:    [[N1:%.*]] = alloca i32, align 41018// CHECK-64-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81019// CHECK-64-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 81020// CHECK-64-NEXT:    store ptr [[N]], ptr [[N_ADDR]], align 81021// CHECK-64-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 81022// CHECK-64-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[N_ADDR]], align 81023// CHECK-64-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l44_kernel_environment, ptr [[DYN_PTR]])1024// CHECK-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -11025// CHECK-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]1026// CHECK-64:       user_code.entry:1027// CHECK-64-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 41028// CHECK-64-NEXT:    store i32 0, ptr [[N1]], align 41029// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1030// CHECK-64:       omp.inner.for.cond:1031// CHECK-64-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34:![0-9]+]]1032// CHECK-64-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP3]], 101033// CHECK-64-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1034// CHECK-64:       omp.inner.for.body:1035// CHECK-64-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34]]1036// CHECK-64-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP4]], 11037// CHECK-64-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1038// CHECK-64-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP34]]1039// CHECK-64-NEXT:    [[TMP5:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP34]]1040// CHECK-64-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP5]] to i641041// CHECK-64-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]1042// CHECK-64-NEXT:    [[TMP6:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP34]]1043// CHECK-64-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 11044// CHECK-64-NEXT:    store i32 [[ADD2]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP34]]1045// CHECK-64-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1046// CHECK-64:       omp.body.continue:1047// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1048// CHECK-64:       omp.inner.for.inc:1049// CHECK-64-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34]]1050// CHECK-64-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP7]], 11051// CHECK-64-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34]]1052// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP35:![0-9]+]]1053// CHECK-64:       worker.exit:1054// CHECK-64-NEXT:    ret void1055// CHECK-64:       omp.inner.for.end:1056// CHECK-64-NEXT:    store i32 10, ptr [[I]], align 41057// CHECK-64-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP1]], align 41058// CHECK-64-NEXT:    [[TMP9:%.*]] = load i32, ptr [[N1]], align 41059// CHECK-64-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]1060// CHECK-64-NEXT:    store i32 [[ADD4]], ptr [[TMP1]], align 41061// CHECK-64-NEXT:    call void @__kmpc_target_deinit()1062// CHECK-64-NEXT:    ret void1063//1064//1065// CHECK-32-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l291066// CHECK-32-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR0:[0-9]+]] {1067// CHECK-32-NEXT:  entry:1068// CHECK-32-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41069// CHECK-32-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41070// CHECK-32-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 41071// CHECK-32-NEXT:    [[TMP:%.*]] = alloca i32, align 41072// CHECK-32-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41073// CHECK-32-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41074// CHECK-32-NEXT:    [[I:%.*]] = alloca i32, align 41075// CHECK-32-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41076// CHECK-32-NEXT:    [[I3:%.*]] = alloca i32, align 41077// CHECK-32-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41078// CHECK-32-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41079// CHECK-32-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 41080// CHECK-32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41081// CHECK-32-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l29_kernel_environment, ptr [[DYN_PTR]])1082// CHECK-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -11083// CHECK-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]1084// CHECK-32:       user_code.entry:1085// CHECK-32-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 41086// CHECK-32-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 41087// CHECK-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41088// CHECK-32-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 01089// CHECK-32-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 11090// CHECK-32-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11091// CHECK-32-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41092// CHECK-32-NEXT:    store i32 0, ptr [[I]], align 41093// CHECK-32-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41094// CHECK-32-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]1095// CHECK-32-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]1096// CHECK-32:       simd.if.then:1097// CHECK-32-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 41098// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1099// CHECK-32:       omp.inner.for.cond:1100// CHECK-32-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24:![0-9]+]]1101// CHECK-32-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP24]]1102// CHECK-32-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 11103// CHECK-32-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]1104// CHECK-32-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1105// CHECK-32:       omp.inner.for.body:1106// CHECK-32-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24]]1107// CHECK-32-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 11108// CHECK-32-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]1109// CHECK-32-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP24]]1110// CHECK-32-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP24]]1111// CHECK-32-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i32 0, i32 [[TMP8]]1112// CHECK-32-NEXT:    store i32 1, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP24]]1113// CHECK-32-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1114// CHECK-32:       omp.body.continue:1115// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1116// CHECK-32:       omp.inner.for.inc:1117// CHECK-32-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24]]1118// CHECK-32-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP9]], 11119// CHECK-32-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24]]1120// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP25:![0-9]+]]1121// CHECK-32:       worker.exit:1122// CHECK-32-NEXT:    ret void1123// CHECK-32:       omp.inner.for.end:1124// CHECK-32-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41125// CHECK-32-NEXT:    [[SUB7:%.*]] = sub nsw i32 [[TMP10]], 01126// CHECK-32-NEXT:    [[DIV8:%.*]] = sdiv i32 [[SUB7]], 11127// CHECK-32-NEXT:    [[MUL9:%.*]] = mul nsw i32 [[DIV8]], 11128// CHECK-32-NEXT:    [[ADD10:%.*]] = add nsw i32 0, [[MUL9]]1129// CHECK-32-NEXT:    store i32 [[ADD10]], ptr [[I3]], align 41130// CHECK-32-NEXT:    br label [[SIMD_IF_END]]1131// CHECK-32:       simd.if.end:1132// CHECK-32-NEXT:    call void @__kmpc_target_deinit()1133// CHECK-32-NEXT:    ret void1134//1135//1136// CHECK-32-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l341137// CHECK-32-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 2 dereferenceable(2000) [[AA:%.*]]) #[[ATTR0]] {1138// CHECK-32-NEXT:  entry:1139// CHECK-32-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41140// CHECK-32-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41141// CHECK-32-NEXT:    [[AA_ADDR:%.*]] = alloca ptr, align 41142// CHECK-32-NEXT:    [[TMP:%.*]] = alloca i32, align 41143// CHECK-32-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41144// CHECK-32-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41145// CHECK-32-NEXT:    [[I:%.*]] = alloca i32, align 41146// CHECK-32-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41147// CHECK-32-NEXT:    [[I3:%.*]] = alloca i32, align 41148// CHECK-32-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41149// CHECK-32-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41150// CHECK-32-NEXT:    store ptr [[AA]], ptr [[AA_ADDR]], align 41151// CHECK-32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 41152// CHECK-32-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l34_kernel_environment, ptr [[DYN_PTR]])1153// CHECK-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -11154// CHECK-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]1155// CHECK-32:       user_code.entry:1156// CHECK-32-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 41157// CHECK-32-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 41158// CHECK-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41159// CHECK-32-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 01160// CHECK-32-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 11161// CHECK-32-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11162// CHECK-32-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41163// CHECK-32-NEXT:    store i32 0, ptr [[I]], align 41164// CHECK-32-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41165// CHECK-32-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]1166// CHECK-32-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]1167// CHECK-32:       simd.if.then:1168// CHECK-32-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 41169// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1170// CHECK-32:       omp.inner.for.cond:1171// CHECK-32-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28:![0-9]+]]1172// CHECK-32-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP28]]1173// CHECK-32-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 11174// CHECK-32-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]1175// CHECK-32-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1176// CHECK-32:       omp.inner.for.body:1177// CHECK-32-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28]]1178// CHECK-32-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 11179// CHECK-32-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]1180// CHECK-32-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP28]]1181// CHECK-32-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP28]]1182// CHECK-32-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i16], ptr [[TMP0]], i32 0, i32 [[TMP8]]1183// CHECK-32-NEXT:    [[TMP9:%.*]] = load i16, ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP28]]1184// CHECK-32-NEXT:    [[CONV:%.*]] = sext i16 [[TMP9]] to i321185// CHECK-32-NEXT:    [[ADD6:%.*]] = add nsw i32 [[CONV]], 11186// CHECK-32-NEXT:    [[CONV7:%.*]] = trunc i32 [[ADD6]] to i161187// CHECK-32-NEXT:    store i16 [[CONV7]], ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP28]]1188// CHECK-32-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1189// CHECK-32:       omp.body.continue:1190// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1191// CHECK-32:       omp.inner.for.inc:1192// CHECK-32-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28]]1193// CHECK-32-NEXT:    [[ADD8:%.*]] = add nsw i32 [[TMP10]], 11194// CHECK-32-NEXT:    store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28]]1195// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP29:![0-9]+]]1196// CHECK-32:       worker.exit:1197// CHECK-32-NEXT:    ret void1198// CHECK-32:       omp.inner.for.end:1199// CHECK-32-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41200// CHECK-32-NEXT:    [[SUB9:%.*]] = sub nsw i32 [[TMP11]], 01201// CHECK-32-NEXT:    [[DIV10:%.*]] = sdiv i32 [[SUB9]], 11202// CHECK-32-NEXT:    [[MUL11:%.*]] = mul nsw i32 [[DIV10]], 11203// CHECK-32-NEXT:    [[ADD12:%.*]] = add nsw i32 0, [[MUL11]]1204// CHECK-32-NEXT:    store i32 [[ADD12]], ptr [[I3]], align 41205// CHECK-32-NEXT:    br label [[SIMD_IF_END]]1206// CHECK-32:       simd.if.end:1207// CHECK-32-NEXT:    call void @__kmpc_target_deinit()1208// CHECK-32-NEXT:    ret void1209//1210//1211// CHECK-32-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l391212// CHECK-32-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR0]] {1213// CHECK-32-NEXT:  entry:1214// CHECK-32-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41215// CHECK-32-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 41216// CHECK-32-NEXT:    [[TMP:%.*]] = alloca i32, align 41217// CHECK-32-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41218// CHECK-32-NEXT:    [[I:%.*]] = alloca i32, align 41219// CHECK-32-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41220// CHECK-32-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 41221// CHECK-32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 41222// CHECK-32-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l39_kernel_environment, ptr [[DYN_PTR]])1223// CHECK-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -11224// CHECK-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]1225// CHECK-32:       user_code.entry:1226// CHECK-32-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 41227// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1228// CHECK-32:       omp.inner.for.cond:1229// CHECK-32-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31:![0-9]+]]1230// CHECK-32-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP2]], 101231// CHECK-32-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1232// CHECK-32:       omp.inner.for.body:1233// CHECK-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31]]1234// CHECK-32-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP3]], 11235// CHECK-32-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1236// CHECK-32-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP31]]1237// CHECK-32-NEXT:    [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP31]]1238// CHECK-32-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP4]]1239// CHECK-32-NEXT:    [[TMP5:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP31]]1240// CHECK-32-NEXT:    [[ADD1:%.*]] = add nsw i32 [[TMP5]], 11241// CHECK-32-NEXT:    store i32 [[ADD1]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP31]]1242// CHECK-32-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1243// CHECK-32:       omp.body.continue:1244// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1245// CHECK-32:       omp.inner.for.inc:1246// CHECK-32-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31]]1247// CHECK-32-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 11248// CHECK-32-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31]]1249// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP32:![0-9]+]]1250// CHECK-32:       worker.exit:1251// CHECK-32-NEXT:    ret void1252// CHECK-32:       omp.inner.for.end:1253// CHECK-32-NEXT:    store i32 10, ptr [[I]], align 41254// CHECK-32-NEXT:    call void @__kmpc_target_deinit()1255// CHECK-32-NEXT:    ret void1256//1257//1258// CHECK-32-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l441259// CHECK-32-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]]) #[[ATTR0]] {1260// CHECK-32-NEXT:  entry:1261// CHECK-32-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41262// CHECK-32-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 41263// CHECK-32-NEXT:    [[N_ADDR:%.*]] = alloca ptr, align 41264// CHECK-32-NEXT:    [[TMP:%.*]] = alloca i32, align 41265// CHECK-32-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41266// CHECK-32-NEXT:    [[I:%.*]] = alloca i32, align 41267// CHECK-32-NEXT:    [[N1:%.*]] = alloca i32, align 41268// CHECK-32-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41269// CHECK-32-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 41270// CHECK-32-NEXT:    store ptr [[N]], ptr [[N_ADDR]], align 41271// CHECK-32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 41272// CHECK-32-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[N_ADDR]], align 41273// CHECK-32-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l44_kernel_environment, ptr [[DYN_PTR]])1274// CHECK-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -11275// CHECK-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]1276// CHECK-32:       user_code.entry:1277// CHECK-32-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 41278// CHECK-32-NEXT:    store i32 0, ptr [[N1]], align 41279// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1280// CHECK-32:       omp.inner.for.cond:1281// CHECK-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34:![0-9]+]]1282// CHECK-32-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP3]], 101283// CHECK-32-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1284// CHECK-32:       omp.inner.for.body:1285// CHECK-32-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34]]1286// CHECK-32-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP4]], 11287// CHECK-32-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1288// CHECK-32-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP34]]1289// CHECK-32-NEXT:    [[TMP5:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP34]]1290// CHECK-32-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP5]]1291// CHECK-32-NEXT:    [[TMP6:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP34]]1292// CHECK-32-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 11293// CHECK-32-NEXT:    store i32 [[ADD2]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP34]]1294// CHECK-32-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1295// CHECK-32:       omp.body.continue:1296// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1297// CHECK-32:       omp.inner.for.inc:1298// CHECK-32-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34]]1299// CHECK-32-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP7]], 11300// CHECK-32-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34]]1301// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP35:![0-9]+]]1302// CHECK-32:       worker.exit:1303// CHECK-32-NEXT:    ret void1304// CHECK-32:       omp.inner.for.end:1305// CHECK-32-NEXT:    store i32 10, ptr [[I]], align 41306// CHECK-32-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP1]], align 41307// CHECK-32-NEXT:    [[TMP9:%.*]] = load i32, ptr [[N1]], align 41308// CHECK-32-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]1309// CHECK-32-NEXT:    store i32 [[ADD4]], ptr [[TMP1]], align 41310// CHECK-32-NEXT:    call void @__kmpc_target_deinit()1311// CHECK-32-NEXT:    ret void1312//1313//1314// CHECK-32-EX-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l291315// CHECK-32-EX-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR0:[0-9]+]] {1316// CHECK-32-EX-NEXT:  entry:1317// CHECK-32-EX-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41318// CHECK-32-EX-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41319// CHECK-32-EX-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 41320// CHECK-32-EX-NEXT:    [[TMP:%.*]] = alloca i32, align 41321// CHECK-32-EX-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41322// CHECK-32-EX-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41323// CHECK-32-EX-NEXT:    [[I:%.*]] = alloca i32, align 41324// CHECK-32-EX-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41325// CHECK-32-EX-NEXT:    [[I3:%.*]] = alloca i32, align 41326// CHECK-32-EX-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41327// CHECK-32-EX-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41328// CHECK-32-EX-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 41329// CHECK-32-EX-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41330// CHECK-32-EX-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l29_kernel_environment, ptr [[DYN_PTR]])1331// CHECK-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -11332// CHECK-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]1333// CHECK-32-EX:       user_code.entry:1334// CHECK-32-EX-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 41335// CHECK-32-EX-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 41336// CHECK-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41337// CHECK-32-EX-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 01338// CHECK-32-EX-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 11339// CHECK-32-EX-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11340// CHECK-32-EX-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41341// CHECK-32-EX-NEXT:    store i32 0, ptr [[I]], align 41342// CHECK-32-EX-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41343// CHECK-32-EX-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]1344// CHECK-32-EX-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]1345// CHECK-32-EX:       simd.if.then:1346// CHECK-32-EX-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 41347// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1348// CHECK-32-EX:       omp.inner.for.cond:1349// CHECK-32-EX-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24:![0-9]+]]1350// CHECK-32-EX-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP24]]1351// CHECK-32-EX-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 11352// CHECK-32-EX-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]1353// CHECK-32-EX-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1354// CHECK-32-EX:       omp.inner.for.body:1355// CHECK-32-EX-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24]]1356// CHECK-32-EX-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 11357// CHECK-32-EX-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]1358// CHECK-32-EX-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP24]]1359// CHECK-32-EX-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP24]]1360// CHECK-32-EX-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i32 0, i32 [[TMP8]]1361// CHECK-32-EX-NEXT:    store i32 1, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP24]]1362// CHECK-32-EX-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1363// CHECK-32-EX:       omp.body.continue:1364// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1365// CHECK-32-EX:       omp.inner.for.inc:1366// CHECK-32-EX-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24]]1367// CHECK-32-EX-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP9]], 11368// CHECK-32-EX-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP24]]1369// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP25:![0-9]+]]1370// CHECK-32-EX:       worker.exit:1371// CHECK-32-EX-NEXT:    ret void1372// CHECK-32-EX:       omp.inner.for.end:1373// CHECK-32-EX-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41374// CHECK-32-EX-NEXT:    [[SUB7:%.*]] = sub nsw i32 [[TMP10]], 01375// CHECK-32-EX-NEXT:    [[DIV8:%.*]] = sdiv i32 [[SUB7]], 11376// CHECK-32-EX-NEXT:    [[MUL9:%.*]] = mul nsw i32 [[DIV8]], 11377// CHECK-32-EX-NEXT:    [[ADD10:%.*]] = add nsw i32 0, [[MUL9]]1378// CHECK-32-EX-NEXT:    store i32 [[ADD10]], ptr [[I3]], align 41379// CHECK-32-EX-NEXT:    br label [[SIMD_IF_END]]1380// CHECK-32-EX:       simd.if.end:1381// CHECK-32-EX-NEXT:    call void @__kmpc_target_deinit()1382// CHECK-32-EX-NEXT:    ret void1383//1384//1385// CHECK-32-EX-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l341386// CHECK-32-EX-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[N:%.*]], ptr noundef nonnull align 2 dereferenceable(2000) [[AA:%.*]]) #[[ATTR0]] {1387// CHECK-32-EX-NEXT:  entry:1388// CHECK-32-EX-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41389// CHECK-32-EX-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 41390// CHECK-32-EX-NEXT:    [[AA_ADDR:%.*]] = alloca ptr, align 41391// CHECK-32-EX-NEXT:    [[TMP:%.*]] = alloca i32, align 41392// CHECK-32-EX-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41393// CHECK-32-EX-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41394// CHECK-32-EX-NEXT:    [[I:%.*]] = alloca i32, align 41395// CHECK-32-EX-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41396// CHECK-32-EX-NEXT:    [[I3:%.*]] = alloca i32, align 41397// CHECK-32-EX-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41398// CHECK-32-EX-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 41399// CHECK-32-EX-NEXT:    store ptr [[AA]], ptr [[AA_ADDR]], align 41400// CHECK-32-EX-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 41401// CHECK-32-EX-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l34_kernel_environment, ptr [[DYN_PTR]])1402// CHECK-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -11403// CHECK-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]1404// CHECK-32-EX:       user_code.entry:1405// CHECK-32-EX-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 41406// CHECK-32-EX-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 41407// CHECK-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41408// CHECK-32-EX-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 01409// CHECK-32-EX-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 11410// CHECK-32-EX-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11411// CHECK-32-EX-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41412// CHECK-32-EX-NEXT:    store i32 0, ptr [[I]], align 41413// CHECK-32-EX-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41414// CHECK-32-EX-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]1415// CHECK-32-EX-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]1416// CHECK-32-EX:       simd.if.then:1417// CHECK-32-EX-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 41418// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1419// CHECK-32-EX:       omp.inner.for.cond:1420// CHECK-32-EX-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28:![0-9]+]]1421// CHECK-32-EX-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP28]]1422// CHECK-32-EX-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 11423// CHECK-32-EX-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]1424// CHECK-32-EX-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1425// CHECK-32-EX:       omp.inner.for.body:1426// CHECK-32-EX-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28]]1427// CHECK-32-EX-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 11428// CHECK-32-EX-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]1429// CHECK-32-EX-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP28]]1430// CHECK-32-EX-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP28]]1431// CHECK-32-EX-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i16], ptr [[TMP0]], i32 0, i32 [[TMP8]]1432// CHECK-32-EX-NEXT:    [[TMP9:%.*]] = load i16, ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP28]]1433// CHECK-32-EX-NEXT:    [[CONV:%.*]] = sext i16 [[TMP9]] to i321434// CHECK-32-EX-NEXT:    [[ADD6:%.*]] = add nsw i32 [[CONV]], 11435// CHECK-32-EX-NEXT:    [[CONV7:%.*]] = trunc i32 [[ADD6]] to i161436// CHECK-32-EX-NEXT:    store i16 [[CONV7]], ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP28]]1437// CHECK-32-EX-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1438// CHECK-32-EX:       omp.body.continue:1439// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1440// CHECK-32-EX:       omp.inner.for.inc:1441// CHECK-32-EX-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28]]1442// CHECK-32-EX-NEXT:    [[ADD8:%.*]] = add nsw i32 [[TMP10]], 11443// CHECK-32-EX-NEXT:    store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP28]]1444// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP29:![0-9]+]]1445// CHECK-32-EX:       worker.exit:1446// CHECK-32-EX-NEXT:    ret void1447// CHECK-32-EX:       omp.inner.for.end:1448// CHECK-32-EX-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41449// CHECK-32-EX-NEXT:    [[SUB9:%.*]] = sub nsw i32 [[TMP11]], 01450// CHECK-32-EX-NEXT:    [[DIV10:%.*]] = sdiv i32 [[SUB9]], 11451// CHECK-32-EX-NEXT:    [[MUL11:%.*]] = mul nsw i32 [[DIV10]], 11452// CHECK-32-EX-NEXT:    [[ADD12:%.*]] = add nsw i32 0, [[MUL11]]1453// CHECK-32-EX-NEXT:    store i32 [[ADD12]], ptr [[I3]], align 41454// CHECK-32-EX-NEXT:    br label [[SIMD_IF_END]]1455// CHECK-32-EX:       simd.if.end:1456// CHECK-32-EX-NEXT:    call void @__kmpc_target_deinit()1457// CHECK-32-EX-NEXT:    ret void1458//1459//1460// CHECK-32-EX-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l391461// CHECK-32-EX-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR0]] {1462// CHECK-32-EX-NEXT:  entry:1463// CHECK-32-EX-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41464// CHECK-32-EX-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 41465// CHECK-32-EX-NEXT:    [[TMP:%.*]] = alloca i32, align 41466// CHECK-32-EX-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41467// CHECK-32-EX-NEXT:    [[I:%.*]] = alloca i32, align 41468// CHECK-32-EX-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41469// CHECK-32-EX-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 41470// CHECK-32-EX-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 41471// CHECK-32-EX-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l39_kernel_environment, ptr [[DYN_PTR]])1472// CHECK-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -11473// CHECK-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]1474// CHECK-32-EX:       user_code.entry:1475// CHECK-32-EX-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 41476// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1477// CHECK-32-EX:       omp.inner.for.cond:1478// CHECK-32-EX-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31:![0-9]+]]1479// CHECK-32-EX-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP2]], 101480// CHECK-32-EX-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1481// CHECK-32-EX:       omp.inner.for.body:1482// CHECK-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31]]1483// CHECK-32-EX-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP3]], 11484// CHECK-32-EX-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1485// CHECK-32-EX-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP31]]1486// CHECK-32-EX-NEXT:    [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP31]]1487// CHECK-32-EX-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP4]]1488// CHECK-32-EX-NEXT:    [[TMP5:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP31]]1489// CHECK-32-EX-NEXT:    [[ADD1:%.*]] = add nsw i32 [[TMP5]], 11490// CHECK-32-EX-NEXT:    store i32 [[ADD1]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP31]]1491// CHECK-32-EX-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1492// CHECK-32-EX:       omp.body.continue:1493// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1494// CHECK-32-EX:       omp.inner.for.inc:1495// CHECK-32-EX-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31]]1496// CHECK-32-EX-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 11497// CHECK-32-EX-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP31]]1498// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP32:![0-9]+]]1499// CHECK-32-EX:       worker.exit:1500// CHECK-32-EX-NEXT:    ret void1501// CHECK-32-EX:       omp.inner.for.end:1502// CHECK-32-EX-NEXT:    store i32 10, ptr [[I]], align 41503// CHECK-32-EX-NEXT:    call void @__kmpc_target_deinit()1504// CHECK-32-EX-NEXT:    ret void1505//1506//1507// CHECK-32-EX-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l441508// CHECK-32-EX-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]]) #[[ATTR0]] {1509// CHECK-32-EX-NEXT:  entry:1510// CHECK-32-EX-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41511// CHECK-32-EX-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 41512// CHECK-32-EX-NEXT:    [[N_ADDR:%.*]] = alloca ptr, align 41513// CHECK-32-EX-NEXT:    [[TMP:%.*]] = alloca i32, align 41514// CHECK-32-EX-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 41515// CHECK-32-EX-NEXT:    [[I:%.*]] = alloca i32, align 41516// CHECK-32-EX-NEXT:    [[N1:%.*]] = alloca i32, align 41517// CHECK-32-EX-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41518// CHECK-32-EX-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 41519// CHECK-32-EX-NEXT:    store ptr [[N]], ptr [[N_ADDR]], align 41520// CHECK-32-EX-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 41521// CHECK-32-EX-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[N_ADDR]], align 41522// CHECK-32-EX-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l44_kernel_environment, ptr [[DYN_PTR]])1523// CHECK-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -11524// CHECK-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]1525// CHECK-32-EX:       user_code.entry:1526// CHECK-32-EX-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 41527// CHECK-32-EX-NEXT:    store i32 0, ptr [[N1]], align 41528// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]1529// CHECK-32-EX:       omp.inner.for.cond:1530// CHECK-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34:![0-9]+]]1531// CHECK-32-EX-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP3]], 101532// CHECK-32-EX-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1533// CHECK-32-EX:       omp.inner.for.body:1534// CHECK-32-EX-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34]]1535// CHECK-32-EX-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP4]], 11536// CHECK-32-EX-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]1537// CHECK-32-EX-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP34]]1538// CHECK-32-EX-NEXT:    [[TMP5:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP34]]1539// CHECK-32-EX-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP5]]1540// CHECK-32-EX-NEXT:    [[TMP6:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP34]]1541// CHECK-32-EX-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 11542// CHECK-32-EX-NEXT:    store i32 [[ADD2]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP34]]1543// CHECK-32-EX-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]1544// CHECK-32-EX:       omp.body.continue:1545// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]1546// CHECK-32-EX:       omp.inner.for.inc:1547// CHECK-32-EX-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34]]1548// CHECK-32-EX-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP7]], 11549// CHECK-32-EX-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP34]]1550// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP35:![0-9]+]]1551// CHECK-32-EX:       worker.exit:1552// CHECK-32-EX-NEXT:    ret void1553// CHECK-32-EX:       omp.inner.for.end:1554// CHECK-32-EX-NEXT:    store i32 10, ptr [[I]], align 41555// CHECK-32-EX-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP1]], align 41556// CHECK-32-EX-NEXT:    [[TMP9:%.*]] = load i32, ptr [[N1]], align 41557// CHECK-32-EX-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]1558// CHECK-32-EX-NEXT:    store i32 [[ADD4]], ptr [[TMP1]], align 41559// CHECK-32-EX-NEXT:    call void @__kmpc_target_deinit()1560// CHECK-32-EX-NEXT:    ret void1561//1562