[Clang][OpenMP] Update tests using update_cc_test_checks.py

[tianshilei1992]

Summary

Make preparation for other patches

[tianshilei1992]

unify check prefix(es)

[tianshilei1992]

remove hash

[jdoerfert]

LG, assuming the bots are happy

[tianshilei1992]

update

clang/test/OpenMP/amdgcn_target_codegen.cpp

+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs
 // REQUIRES: amdgpu-registered-target
 
 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc
 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s
 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
 
 #define N 1000
 
 int test_amdgcn_target_tid_threads() {
-// CHECK-LABEL: define weak_odr protected amdgpu_kernel void @{{.*}}test_amdgcn_target_tid_threads
-
   int arr[N];
-
-// CHECK: call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @1 to ptr), i8 1, i1 true)
 #pragma omp target
   for (int i = 0; i < N; i++) {
     arr[i] = 1;
   }
-
   return arr[0];
 }
 
 int test_amdgcn_target_tid_threads_simd() {
-// CHECK-LABEL: define weak_odr protected amdgpu_kernel void @{{.*}}test_amdgcn_target_tid_threads_simd
-
   int arr[N];
-
-// CHECK: call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @1 to ptr), i8 2, i1 false)
 #pragma omp target simd
   for (int i = 0; i < N; i++) {
     arr[i] = 1;
   }
   return arr[0];
 }
 
 #endif
+// CHECK-LABEL: define {{[^@]+}}@__omp_offloading_30_7d6f1507__Z30test_amdgcn_target_tid_threadsv_l14
+// CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(4000) [[ARR:%.*]]) #[[ATTR0:[0-9]+]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[ARR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT:    [[I:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT:    [[ARR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ARR_ADDR]] to ptr
+// CHECK-NEXT:    [[I_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[I]] to ptr
+// CHECK-NEXT:    store ptr [[ARR]], ptr [[ARR_ADDR_ASCAST]], align 8
+// CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[ARR_ADDR_ASCAST]], align 8
+// CHECK-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @[[GLOB1:[0-9]+]] to ptr), i8 1, i1 true)
+// CHECK-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK:       user_code.entry:
+// CHECK-NEXT:    store i32 0, ptr [[I_ASCAST]], align 4
+// CHECK-NEXT:    br label [[FOR_COND:%.*]]
+// CHECK:       for.cond:
+// CHECK-NEXT:    [[TMP2:%.*]] = load i32, ptr [[I_ASCAST]], align 4
+// CHECK-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP2]], 1000
+// CHECK-NEXT:    br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]]
+// CHECK:       for.body:
+// CHECK-NEXT:    [[TMP3:%.*]] = load i32, ptr [[I_ASCAST]], align 4
+// CHECK-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP3]] to i64
+// CHECK-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
+// CHECK-NEXT:    store i32 1, ptr [[ARRAYIDX]], align 4
+// CHECK-NEXT:    br label [[FOR_INC:%.*]]
+// CHECK:       for.inc:
+// CHECK-NEXT:    [[TMP4:%.*]] = load i32, ptr [[I_ASCAST]], align 4
+// CHECK-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP4]], 1
+// CHECK-NEXT:    store i32 [[INC]], ptr [[I_ASCAST]], align 4
+// CHECK-NEXT:    br label [[FOR_COND]], !llvm.loop [[LOOP9:![0-9]+]]
+// CHECK:       worker.exit:
+// CHECK-NEXT:    ret void
+// CHECK:       for.end:
+// CHECK-NEXT:    call void @__kmpc_target_deinit(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i8 1)
+// CHECK-NEXT:    ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@__omp_offloading_30_7d6f1507__Z35test_amdgcn_target_tid_threads_simdv_l23
+// CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(4000) [[ARR:%.*]]) #[[ATTR1:[0-9]+]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[ARR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT:    [[TMP:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT:    [[I:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT:    [[ARR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ARR_ADDR]] to ptr
+// CHECK-NEXT:    [[TMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[TMP]] to ptr
+// CHECK-NEXT:    [[DOTOMP_IV_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IV]] to ptr
+// CHECK-NEXT:    [[I_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[I]] to ptr
+// CHECK-NEXT:    store ptr [[ARR]], ptr [[ARR_ADDR_ASCAST]], align 8
+// CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[ARR_ADDR_ASCAST]], align 8
+// CHECK-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i8 2, i1 false)
+// CHECK-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK:       user_code.entry:
+// CHECK-NEXT:    store i32 0, ptr [[DOTOMP_IV_ASCAST]], align 4
+// CHECK-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK:       omp.inner.for.cond:
+// CHECK-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4, !llvm.access.group [[ACC_GRP11:![0-9]+]]
+// CHECK-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP2]], 1000
+// CHECK-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK:       omp.inner.for.body:
+// CHECK-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4, !llvm.access.group [[ACC_GRP11]]
+// CHECK-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1
+// CHECK-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK-NEXT:    store i32 [[ADD]], ptr [[I_ASCAST]], align 4, !llvm.access.group [[ACC_GRP11]]
+// CHECK-NEXT:    [[TMP4:%.*]] = load i32, ptr [[I_ASCAST]], align 4, !llvm.access.group [[ACC_GRP11]]
+// CHECK-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP4]] to i64
+// CHECK-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
+// CHECK-NEXT:    store i32 1, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP11]]
+// CHECK-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK:       omp.body.continue:
+// CHECK-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK:       omp.inner.for.inc:
+// CHECK-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4, !llvm.access.group [[ACC_GRP11]]
+// CHECK-NEXT:    [[ADD1:%.*]] = add nsw i32 [[TMP5]], 1
+// CHECK-NEXT:    store i32 [[ADD1]], ptr [[DOTOMP_IV_ASCAST]], align 4, !llvm.access.group [[ACC_GRP11]]
+// CHECK-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP12:![0-9]+]]
+// CHECK:       worker.exit:
+// CHECK-NEXT:    ret void
+// CHECK:       omp.inner.for.end:
+// CHECK-NEXT:    store i32 1000, ptr [[I_ASCAST]], align 4
+// CHECK-NEXT:    call void @__kmpc_target_deinit(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i8 2)
+// CHECK-NEXT:    ret void
+//

clang/test/OpenMP/nvptx_target_parallel_proc_bind_codegen.cpp

+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs
 // Test target codegen - host bc file has to be created first.
 // 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.bc
-// 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-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
+// 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-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK45-64
 // 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.bc
-// 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-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
-// 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-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
+// 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-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix=CHECK45-32
+// 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-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefixes=CHECK45-32,CHECK45-32-EX
 
 // 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.bc
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK-64
 // 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.bc
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
-// 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-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix=CHECK-32
+// 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-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefixes=CHECK-32,CHECK-32-EX
 
 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
 
 // Check that the execution mode of all 3 target regions on the gpu is set to SPMD Mode.
-// CHECK-DAG: {{@__omp_offloading_.+l29}}_exec_mode = weak protected constant i8 2
-// CHECK-DAG: {{@__omp_offloading_.+l33}}_exec_mode = weak protected constant i8 2
-// CHECK-DAG: {{@__omp_offloading_.+l38}}_exec_mode = weak protected constant i8 2
 
 template<typename tx>
 tx ftemplate(int n) {
   tx a = 0;
   short aa = 0;
   tx b[10];
 
   #pragma omp target parallel proc_bind(master)

 int bar(int n){
   int a = 0;
 
   a += ftemplate<int>(n);
 
   return a;
 }
 
-// CHECK-NOT: call void @__kmpc_push_proc_bind
 #endif
+// CHECK45-64-LABEL: define {{[^@]+}}@__omp_offloading_30_700c39a7__Z9ftemplateIiET_i_l27
+// CHECK45-64-SAME: () #[[ATTR0:[0-9]+]] {
+// CHECK45-64-NEXT:  entry:
+// CHECK45-64-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x ptr], align 8
+// CHECK45-64-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1:[0-9]+]], i8 2, i1 false)
+// CHECK45-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
+// CHECK45-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK45-64:       user_code.entry:
+// CHECK45-64-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK45-64-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i64 0)
+// CHECK45-64-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK45-64-NEXT:    ret void
+// CHECK45-64:       worker.exit:
+// CHECK45-64-NEXT:    ret void
+//
+//
+// CHECK45-64-LABEL: define {{[^@]+}}@__omp_outlined__
+// CHECK45-64-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] {
+// CHECK45-64-NEXT:  entry:
+// CHECK45-64-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK45-64-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK45-64-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
+// CHECK45-64-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
+// CHECK45-64-NEXT:    ret void
+//
+//
+// CHECK45-64-LABEL: define {{[^@]+}}@__omp_offloading_30_700c39a7__Z9ftemplateIiET_i_l31
+// CHECK45-64-SAME: (i64 noundef [[AA:%.*]]) #[[ATTR0]] {
+// CHECK45-64-NEXT:  entry:
+// CHECK45-64-NEXT:    [[AA_ADDR:%.*]] = alloca i64, align 8
+// CHECK45-64-NEXT:    [[AA_CASTED:%.*]] = alloca i64, align 8
+// CHECK45-64-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 8
+// CHECK45-64-NEXT:    store i64 [[AA]], ptr [[AA_ADDR]], align 8
+// CHECK45-64-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK45-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
+// CHECK45-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK45-64:       user_code.entry:
+// CHECK45-64-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK45-64-NEXT:    [[TMP2:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK45-64-NEXT:    store i16 [[TMP2]], ptr [[AA_CASTED]], align 2
+// CHECK45-64-NEXT:    [[TMP3:%.*]] = load i64, ptr [[AA_CASTED]], align 8
+// CHECK45-64-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 0
+// CHECK45-64-NEXT:    [[TMP5:%.*]] = inttoptr i64 [[TMP3]] to ptr
+// CHECK45-64-NEXT:    store ptr [[TMP5]], ptr [[TMP4]], align 8
+// CHECK45-64-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__1, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i64 1)
+// CHECK45-64-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK45-64-NEXT:    ret void
+// CHECK45-64:       worker.exit:
+// CHECK45-64-NEXT:    ret void
+//
+//
+// CHECK45-64-LABEL: define {{[^@]+}}@__omp_outlined__1
+// CHECK45-64-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[AA:%.*]]) #[[ATTR1]] {
+// CHECK45-64-NEXT:  entry:
+// CHECK45-64-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK45-64-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK45-64-NEXT:    [[AA_ADDR:%.*]] = alloca i64, align 8
+// CHECK45-64-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
+// CHECK45-64-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
+// CHECK45-64-NEXT:    store i64 [[AA]], ptr [[AA_ADDR]], align 8
+// CHECK45-64-NEXT:    [[TMP0:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK45-64-NEXT:    [[CONV:%.*]] = sext i16 [[TMP0]] to i32
+// CHECK45-64-NEXT:    [[ADD:%.*]] = add nsw i32 [[CONV]], 1
+// CHECK45-64-NEXT:    [[CONV1:%.*]] = trunc i32 [[ADD]] to i16
+// CHECK45-64-NEXT:    store i16 [[CONV1]], ptr [[AA_ADDR]], align 2
+// CHECK45-64-NEXT:    ret void
+//
+//
+// CHECK45-64-LABEL: define {{[^@]+}}@__omp_offloading_30_700c39a7__Z9ftemplateIiET_i_l36
+// CHECK45-64-SAME: (i64 noundef [[A:%.*]], i64 noundef [[AA:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR0]] {
+// CHECK45-64-NEXT:  entry:
+// CHECK45-64-NEXT:    [[A_ADDR:%.*]] = alloca i64, align 8
+// CHECK45-64-NEXT:    [[AA_ADDR:%.*]] = alloca i64, align 8
+// CHECK45-64-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8
+// CHECK45-64-NEXT:    [[A_CASTED:%.*]] = alloca i64, align 8
+// CHECK45-64-NEXT:    [[AA_CASTED:%.*]] = alloca i64, align 8
+// CHECK45-64-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [3 x ptr], align 8
+// CHECK45-64-NEXT:    store i64 [[A]], ptr [[A_ADDR]], align 8
+// CHECK45-64-NEXT:    store i64 [[AA]], ptr [[AA_ADDR]], align 8
+// CHECK45-64-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8
+// CHECK45-64-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 8
+// CHECK45-64-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK45-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK45-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK45-64:       user_code.entry:
+// CHECK45-64-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK45-64-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A_ADDR]], align 4
+// CHECK45-64-NEXT:    store i32 [[TMP3]], ptr [[A_CASTED]], align 4
+// CHECK45-64-NEXT:    [[TMP4:%.*]] = load i64, ptr [[A_CASTED]], align 8
+// CHECK45-64-NEXT:    [[TMP5:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK45-64-NEXT:    store i16 [[TMP5]], ptr [[AA_CASTED]], align 2
+// CHECK45-64-NEXT:    [[TMP6:%.*]] = load i64, ptr [[AA_CASTED]], align 8
+// CHECK45-64-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 0
+// CHECK45-64-NEXT:    [[TMP8:%.*]] = inttoptr i64 [[TMP4]] to ptr
+// CHECK45-64-NEXT:    store ptr [[TMP8]], ptr [[TMP7]], align 8
+// CHECK45-64-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 1
+// CHECK45-64-NEXT:    [[TMP10:%.*]] = inttoptr i64 [[TMP6]] to ptr
+// CHECK45-64-NEXT:    store ptr [[TMP10]], ptr [[TMP9]], align 8
+// CHECK45-64-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 2
+// CHECK45-64-NEXT:    store ptr [[TMP0]], ptr [[TMP11]], align 8
+// CHECK45-64-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP2]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__2, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i64 3)
+// CHECK45-64-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK45-64-NEXT:    ret void
+// CHECK45-64:       worker.exit:
+// CHECK45-64-NEXT:    ret void
+//
+//
+// CHECK45-64-LABEL: define {{[^@]+}}@__omp_outlined__2
+// CHECK45-64-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[A:%.*]], i64 noundef [[AA:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR1]] {
+// CHECK45-64-NEXT:  entry:
+// CHECK45-64-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK45-64-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK45-64-NEXT:    [[A_ADDR:%.*]] = alloca i64, align 8
+// CHECK45-64-NEXT:    [[AA_ADDR:%.*]] = alloca i64, align 8
+// CHECK45-64-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8
+// CHECK45-64-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
+// CHECK45-64-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
+// CHECK45-64-NEXT:    store i64 [[A]], ptr [[A_ADDR]], align 8
+// CHECK45-64-NEXT:    store i64 [[AA]], ptr [[AA_ADDR]], align 8
+// CHECK45-64-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8
+// CHECK45-64-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 8
+// CHECK45-64-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 4
+// CHECK45-64-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], 1
+// CHECK45-64-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 4
+// CHECK45-64-NEXT:    [[TMP2:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK45-64-NEXT:    [[CONV:%.*]] = sext i16 [[TMP2]] to i32
+// CHECK45-64-NEXT:    [[ADD1:%.*]] = add nsw i32 [[CONV]], 1
+// CHECK45-64-NEXT:    [[CONV2:%.*]] = trunc i32 [[ADD1]] to i16
+// CHECK45-64-NEXT:    store i16 [[CONV2]], ptr [[AA_ADDR]], align 2
+// CHECK45-64-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 2
+// CHECK45-64-NEXT:    [[TMP3:%.*]] = load i32, ptr [[ARRAYIDX]], align 4
+// CHECK45-64-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP3]], 1
+// CHECK45-64-NEXT:    store i32 [[ADD3]], ptr [[ARRAYIDX]], align 4
+// CHECK45-64-NEXT:    ret void
+//
+//
+// CHECK45-32-LABEL: define {{[^@]+}}@__omp_offloading_30_700c39a7__Z9ftemplateIiET_i_l27
+// CHECK45-32-SAME: () #[[ATTR0:[0-9]+]] {
+// CHECK45-32-NEXT:  entry:
+// CHECK45-32-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x ptr], align 4
+// CHECK45-32-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1:[0-9]+]], i8 2, i1 false)
+// CHECK45-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
+// CHECK45-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK45-32:       user_code.entry:
+// CHECK45-32-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK45-32-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 0)
+// CHECK45-32-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK45-32-NEXT:    ret void
+// CHECK45-32:       worker.exit:
+// CHECK45-32-NEXT:    ret void
+//
+//
+// CHECK45-32-LABEL: define {{[^@]+}}@__omp_outlined__
+// CHECK45-32-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] {
+// CHECK45-32-NEXT:  entry:
+// CHECK45-32-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
+// CHECK45-32-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
+// CHECK45-32-NEXT:    ret void
+//
+//
+// CHECK45-32-LABEL: define {{[^@]+}}@__omp_offloading_30_700c39a7__Z9ftemplateIiET_i_l31
+// CHECK45-32-SAME: (i32 noundef [[AA:%.*]]) #[[ATTR0]] {
+// CHECK45-32-NEXT:  entry:
+// CHECK45-32-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[AA_CASTED:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 4
+// CHECK45-32-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK45-32-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK45-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
+// CHECK45-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK45-32:       user_code.entry:
+// CHECK45-32-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK45-32-NEXT:    [[TMP2:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK45-32-NEXT:    store i16 [[TMP2]], ptr [[AA_CASTED]], align 2
+// CHECK45-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[AA_CASTED]], align 4
+// CHECK45-32-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 0
+// CHECK45-32-NEXT:    [[TMP5:%.*]] = inttoptr i32 [[TMP3]] to ptr
+// CHECK45-32-NEXT:    store ptr [[TMP5]], ptr [[TMP4]], align 4
+// CHECK45-32-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__1, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 1)
+// CHECK45-32-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK45-32-NEXT:    ret void
+// CHECK45-32:       worker.exit:
+// CHECK45-32-NEXT:    ret void
+//
+//
+// CHECK45-32-LABEL: define {{[^@]+}}@__omp_outlined__1
+// CHECK45-32-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[AA:%.*]]) #[[ATTR1]] {
+// CHECK45-32-NEXT:  entry:
+// CHECK45-32-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
+// CHECK45-32-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
+// CHECK45-32-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK45-32-NEXT:    [[TMP0:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK45-32-NEXT:    [[CONV:%.*]] = sext i16 [[TMP0]] to i32
+// CHECK45-32-NEXT:    [[ADD:%.*]] = add nsw i32 [[CONV]], 1
+// CHECK45-32-NEXT:    [[CONV1:%.*]] = trunc i32 [[ADD]] to i16
+// CHECK45-32-NEXT:    store i16 [[CONV1]], ptr [[AA_ADDR]], align 2
+// CHECK45-32-NEXT:    ret void
+//
+//
+// CHECK45-32-LABEL: define {{[^@]+}}@__omp_offloading_30_700c39a7__Z9ftemplateIiET_i_l36
+// CHECK45-32-SAME: (i32 noundef [[A:%.*]], i32 noundef [[AA:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR0]] {
+// CHECK45-32-NEXT:  entry:
+// CHECK45-32-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-NEXT:    [[A_CASTED:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[AA_CASTED:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [3 x ptr], align 4
+// CHECK45-32-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 4
+// CHECK45-32-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK45-32-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4
+// CHECK45-32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4
+// CHECK45-32-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK45-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK45-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK45-32:       user_code.entry:
+// CHECK45-32-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK45-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A_ADDR]], align 4
+// CHECK45-32-NEXT:    store i32 [[TMP3]], ptr [[A_CASTED]], align 4
+// CHECK45-32-NEXT:    [[TMP4:%.*]] = load i32, ptr [[A_CASTED]], align 4
+// CHECK45-32-NEXT:    [[TMP5:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK45-32-NEXT:    store i16 [[TMP5]], ptr [[AA_CASTED]], align 2
+// CHECK45-32-NEXT:    [[TMP6:%.*]] = load i32, ptr [[AA_CASTED]], align 4
+// CHECK45-32-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 0
+// CHECK45-32-NEXT:    [[TMP8:%.*]] = inttoptr i32 [[TMP4]] to ptr
+// CHECK45-32-NEXT:    store ptr [[TMP8]], ptr [[TMP7]], align 4
+// CHECK45-32-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 1
+// CHECK45-32-NEXT:    [[TMP10:%.*]] = inttoptr i32 [[TMP6]] to ptr
+// CHECK45-32-NEXT:    store ptr [[TMP10]], ptr [[TMP9]], align 4
+// CHECK45-32-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 2
+// CHECK45-32-NEXT:    store ptr [[TMP0]], ptr [[TMP11]], align 4
+// CHECK45-32-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP2]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__2, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 3)
+// CHECK45-32-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK45-32-NEXT:    ret void
+// CHECK45-32:       worker.exit:
+// CHECK45-32-NEXT:    ret void
+//
+//
+// CHECK45-32-LABEL: define {{[^@]+}}@__omp_outlined__2
+// CHECK45-32-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[A:%.*]], i32 noundef [[AA:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR1]] {
+// CHECK45-32-NEXT:  entry:
+// CHECK45-32-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
+// CHECK45-32-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
+// CHECK45-32-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 4
+// CHECK45-32-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK45-32-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4
+// CHECK45-32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4
+// CHECK45-32-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 4
+// CHECK45-32-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], 1
+// CHECK45-32-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 4
+// CHECK45-32-NEXT:    [[TMP2:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK45-32-NEXT:    [[CONV:%.*]] = sext i16 [[TMP2]] to i32
+// CHECK45-32-NEXT:    [[ADD1:%.*]] = add nsw i32 [[CONV]], 1
+// CHECK45-32-NEXT:    [[CONV2:%.*]] = trunc i32 [[ADD1]] to i16
+// CHECK45-32-NEXT:    store i16 [[CONV2]], ptr [[AA_ADDR]], align 2
+// CHECK45-32-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 2
+// CHECK45-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[ARRAYIDX]], align 4
+// CHECK45-32-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP3]], 1
+// CHECK45-32-NEXT:    store i32 [[ADD3]], ptr [[ARRAYIDX]], align 4
+// CHECK45-32-NEXT:    ret void
+//
+//
+// CHECK45-32-EX-LABEL: define {{[^@]+}}@__omp_offloading_30_700c39a7__Z9ftemplateIiET_i_l27
+// CHECK45-32-EX-SAME: () #[[ATTR0:[0-9]+]] {
+// CHECK45-32-EX-NEXT:  entry:
+// CHECK45-32-EX-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x ptr], align 4
+// CHECK45-32-EX-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1:[0-9]+]], i8 2, i1 false)
+// CHECK45-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
+// CHECK45-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK45-32-EX:       user_code.entry:
+// CHECK45-32-EX-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK45-32-EX-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 0)
+// CHECK45-32-EX-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK45-32-EX-NEXT:    ret void
+// CHECK45-32-EX:       worker.exit:
+// CHECK45-32-EX-NEXT:    ret void
+//
+//
+// CHECK45-32-EX-LABEL: define {{[^@]+}}@__omp_outlined__
+// CHECK45-32-EX-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] {
+// CHECK45-32-EX-NEXT:  entry:
+// CHECK45-32-EX-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-EX-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-EX-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
+// CHECK45-32-EX-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
+// CHECK45-32-EX-NEXT:    ret void
+//
+//
+// CHECK45-32-EX-LABEL: define {{[^@]+}}@__omp_offloading_30_700c39a7__Z9ftemplateIiET_i_l31
+// CHECK45-32-EX-SAME: (i32 noundef [[AA:%.*]]) #[[ATTR0]] {
+// CHECK45-32-EX-NEXT:  entry:
+// CHECK45-32-EX-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[AA_CASTED:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 4
+// CHECK45-32-EX-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK45-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
+// CHECK45-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK45-32-EX:       user_code.entry:
+// CHECK45-32-EX-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK45-32-EX-NEXT:    [[TMP2:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK45-32-EX-NEXT:    store i16 [[TMP2]], ptr [[AA_CASTED]], align 2
+// CHECK45-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[AA_CASTED]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 0
+// CHECK45-32-EX-NEXT:    [[TMP5:%.*]] = inttoptr i32 [[TMP3]] to ptr
+// CHECK45-32-EX-NEXT:    store ptr [[TMP5]], ptr [[TMP4]], align 4
+// CHECK45-32-EX-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__1, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 1)
+// CHECK45-32-EX-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK45-32-EX-NEXT:    ret void
+// CHECK45-32-EX:       worker.exit:
+// CHECK45-32-EX-NEXT:    ret void
+//
+//
+// CHECK45-32-EX-LABEL: define {{[^@]+}}@__omp_outlined__1
+// CHECK45-32-EX-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[AA:%.*]]) #[[ATTR1]] {
+// CHECK45-32-EX-NEXT:  entry:
+// CHECK45-32-EX-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-EX-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-EX-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
+// CHECK45-32-EX-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
+// CHECK45-32-EX-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP0:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK45-32-EX-NEXT:    [[CONV:%.*]] = sext i16 [[TMP0]] to i32
+// CHECK45-32-EX-NEXT:    [[ADD:%.*]] = add nsw i32 [[CONV]], 1
+// CHECK45-32-EX-NEXT:    [[CONV1:%.*]] = trunc i32 [[ADD]] to i16
+// CHECK45-32-EX-NEXT:    store i16 [[CONV1]], ptr [[AA_ADDR]], align 2
+// CHECK45-32-EX-NEXT:    ret void
+//
+//
+// CHECK45-32-EX-LABEL: define {{[^@]+}}@__omp_offloading_30_700c39a7__Z9ftemplateIiET_i_l36
+// CHECK45-32-EX-SAME: (i32 noundef [[A:%.*]], i32 noundef [[AA:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR0]] {
+// CHECK45-32-EX-NEXT:  entry:
+// CHECK45-32-EX-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-EX-NEXT:    [[A_CASTED:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[AA_CASTED:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [3 x ptr], align 4
+// CHECK45-32-EX-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK45-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK45-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK45-32-EX:       user_code.entry:
+// CHECK45-32-EX-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK45-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    store i32 [[TMP3]], ptr [[A_CASTED]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP4:%.*]] = load i32, ptr [[A_CASTED]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP5:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK45-32-EX-NEXT:    store i16 [[TMP5]], ptr [[AA_CASTED]], align 2
+// CHECK45-32-EX-NEXT:    [[TMP6:%.*]] = load i32, ptr [[AA_CASTED]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 0
+// CHECK45-32-EX-NEXT:    [[TMP8:%.*]] = inttoptr i32 [[TMP4]] to ptr
+// CHECK45-32-EX-NEXT:    store ptr [[TMP8]], ptr [[TMP7]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 1
+// CHECK45-32-EX-NEXT:    [[TMP10:%.*]] = inttoptr i32 [[TMP6]] to ptr
+// CHECK45-32-EX-NEXT:    store ptr [[TMP10]], ptr [[TMP9]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 2
+// CHECK45-32-EX-NEXT:    store ptr [[TMP0]], ptr [[TMP11]], align 4
+// CHECK45-32-EX-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP2]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__2, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 3)
+// CHECK45-32-EX-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK45-32-EX-NEXT:    ret void
+// CHECK45-32-EX:       worker.exit:
+// CHECK45-32-EX-NEXT:    ret void
+//
+//
+// CHECK45-32-EX-LABEL: define {{[^@]+}}@__omp_outlined__2
+// CHECK45-32-EX-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[A:%.*]], i32 noundef [[AA:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR1]] {
+// CHECK45-32-EX-NEXT:  entry:
+// CHECK45-32-EX-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-EX-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-EX-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-EX-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
+// CHECK45-32-EX-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
+// CHECK45-32-EX-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], 1
+// CHECK45-32-EX-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP2:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK45-32-EX-NEXT:    [[CONV:%.*]] = sext i16 [[TMP2]] to i32
+// CHECK45-32-EX-NEXT:    [[ADD1:%.*]] = add nsw i32 [[CONV]], 1
+// CHECK45-32-EX-NEXT:    [[CONV2:%.*]] = trunc i32 [[ADD1]] to i16
+// CHECK45-32-EX-NEXT:    store i16 [[CONV2]], ptr [[AA_ADDR]], align 2
+// CHECK45-32-EX-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 2
+// CHECK45-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[ARRAYIDX]], align 4
+// CHECK45-32-EX-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP3]], 1
+// CHECK45-32-EX-NEXT:    store i32 [[ADD3]], ptr [[ARRAYIDX]], align 4
+// CHECK45-32-EX-NEXT:    ret void
+//
+//
+// CHECK-64-LABEL: define {{[^@]+}}@__omp_offloading_30_700c39a7__Z9ftemplateIiET_i_l27
+// CHECK-64-SAME: () #[[ATTR0:[0-9]+]] {
+// CHECK-64-NEXT:  entry:
+// CHECK-64-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x ptr], align 8
+// CHECK-64-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1:[0-9]+]], i8 2, i1 false)
+// CHECK-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
+// CHECK-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK-64:       user_code.entry:
+// CHECK-64-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK-64-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i64 0)
+// CHECK-64-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-64-NEXT:    ret void
+// CHECK-64:       worker.exit:
+// CHECK-64-NEXT:    ret void
+//
+//
+// CHECK-64-LABEL: define {{[^@]+}}@__omp_outlined__
+// CHECK-64-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] {
+// CHECK-64-NEXT:  entry:
+// CHECK-64-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-64-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-64-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
+// CHECK-64-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
+// CHECK-64-NEXT:    ret void
+//
+//
+// CHECK-64-LABEL: define {{[^@]+}}@__omp_offloading_30_700c39a7__Z9ftemplateIiET_i_l31
+// CHECK-64-SAME: (i64 noundef [[AA:%.*]]) #[[ATTR0]] {
+// CHECK-64-NEXT:  entry:
+// CHECK-64-NEXT:    [[AA_ADDR:%.*]] = alloca i64, align 8
+// CHECK-64-NEXT:    [[AA_CASTED:%.*]] = alloca i64, align 8
+// CHECK-64-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 8
+// CHECK-64-NEXT:    store i64 [[AA]], ptr [[AA_ADDR]], align 8
+// CHECK-64-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
+// CHECK-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK-64:       user_code.entry:
+// CHECK-64-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK-64-NEXT:    [[TMP2:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK-64-NEXT:    store i16 [[TMP2]], ptr [[AA_CASTED]], align 2
+// CHECK-64-NEXT:    [[TMP3:%.*]] = load i64, ptr [[AA_CASTED]], align 8
+// CHECK-64-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 0
+// CHECK-64-NEXT:    [[TMP5:%.*]] = inttoptr i64 [[TMP3]] to ptr
+// CHECK-64-NEXT:    store ptr [[TMP5]], ptr [[TMP4]], align 8
+// CHECK-64-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__1, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i64 1)
+// CHECK-64-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-64-NEXT:    ret void
+// CHECK-64:       worker.exit:
+// CHECK-64-NEXT:    ret void
+//
+//
+// CHECK-64-LABEL: define {{[^@]+}}@__omp_outlined__1
+// CHECK-64-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[AA:%.*]]) #[[ATTR1]] {
+// CHECK-64-NEXT:  entry:
+// CHECK-64-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-64-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-64-NEXT:    [[AA_ADDR:%.*]] = alloca i64, align 8
+// CHECK-64-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
+// CHECK-64-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
+// CHECK-64-NEXT:    store i64 [[AA]], ptr [[AA_ADDR]], align 8
+// CHECK-64-NEXT:    [[TMP0:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK-64-NEXT:    [[CONV:%.*]] = sext i16 [[TMP0]] to i32
+// CHECK-64-NEXT:    [[ADD:%.*]] = add nsw i32 [[CONV]], 1
+// CHECK-64-NEXT:    [[CONV1:%.*]] = trunc i32 [[ADD]] to i16
+// CHECK-64-NEXT:    store i16 [[CONV1]], ptr [[AA_ADDR]], align 2
+// CHECK-64-NEXT:    ret void
+//
+//
+// CHECK-64-LABEL: define {{[^@]+}}@__omp_offloading_30_700c39a7__Z9ftemplateIiET_i_l36
+// CHECK-64-SAME: (i64 noundef [[A:%.*]], i64 noundef [[AA:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR0]] {
+// CHECK-64-NEXT:  entry:
+// CHECK-64-NEXT:    [[A_ADDR:%.*]] = alloca i64, align 8
+// CHECK-64-NEXT:    [[AA_ADDR:%.*]] = alloca i64, align 8
+// CHECK-64-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8
+// CHECK-64-NEXT:    [[A_CASTED:%.*]] = alloca i64, align 8
+// CHECK-64-NEXT:    [[AA_CASTED:%.*]] = alloca i64, align 8
+// CHECK-64-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [3 x ptr], align 8
+// CHECK-64-NEXT:    store i64 [[A]], ptr [[A_ADDR]], align 8
+// CHECK-64-NEXT:    store i64 [[AA]], ptr [[AA_ADDR]], align 8
+// CHECK-64-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8
+// CHECK-64-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 8
+// CHECK-64-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK-64:       user_code.entry:
+// CHECK-64-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK-64-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A_ADDR]], align 4
+// CHECK-64-NEXT:    store i32 [[TMP3]], ptr [[A_CASTED]], align 4
+// CHECK-64-NEXT:    [[TMP4:%.*]] = load i64, ptr [[A_CASTED]], align 8
+// CHECK-64-NEXT:    [[TMP5:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK-64-NEXT:    store i16 [[TMP5]], ptr [[AA_CASTED]], align 2
+// CHECK-64-NEXT:    [[TMP6:%.*]] = load i64, ptr [[AA_CASTED]], align 8
+// CHECK-64-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 0
+// CHECK-64-NEXT:    [[TMP8:%.*]] = inttoptr i64 [[TMP4]] to ptr
+// CHECK-64-NEXT:    store ptr [[TMP8]], ptr [[TMP7]], align 8
+// CHECK-64-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 1
+// CHECK-64-NEXT:    [[TMP10:%.*]] = inttoptr i64 [[TMP6]] to ptr
+// CHECK-64-NEXT:    store ptr [[TMP10]], ptr [[TMP9]], align 8
+// CHECK-64-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 2
+// CHECK-64-NEXT:    store ptr [[TMP0]], ptr [[TMP11]], align 8
+// CHECK-64-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP2]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__2, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i64 3)
+// CHECK-64-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-64-NEXT:    ret void
+// CHECK-64:       worker.exit:
+// CHECK-64-NEXT:    ret void
+//
+//
+// CHECK-64-LABEL: define {{[^@]+}}@__omp_outlined__2
+// CHECK-64-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[A:%.*]], i64 noundef [[AA:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR1]] {
+// CHECK-64-NEXT:  entry:
+// CHECK-64-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-64-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-64-NEXT:    [[A_ADDR:%.*]] = alloca i64, align 8
+// CHECK-64-NEXT:    [[AA_ADDR:%.*]] = alloca i64, align 8
+// CHECK-64-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8
+// CHECK-64-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
+// CHECK-64-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
+// CHECK-64-NEXT:    store i64 [[A]], ptr [[A_ADDR]], align 8
+// CHECK-64-NEXT:    store i64 [[AA]], ptr [[AA_ADDR]], align 8
+// CHECK-64-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8
+// CHECK-64-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 8
+// CHECK-64-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 4
+// CHECK-64-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], 1
+// CHECK-64-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 4
+// CHECK-64-NEXT:    [[TMP2:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK-64-NEXT:    [[CONV:%.*]] = sext i16 [[TMP2]] to i32
+// CHECK-64-NEXT:    [[ADD1:%.*]] = add nsw i32 [[CONV]], 1
+// CHECK-64-NEXT:    [[CONV2:%.*]] = trunc i32 [[ADD1]] to i16
+// CHECK-64-NEXT:    store i16 [[CONV2]], ptr [[AA_ADDR]], align 2
+// CHECK-64-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 2
+// CHECK-64-NEXT:    [[TMP3:%.*]] = load i32, ptr [[ARRAYIDX]], align 4
+// CHECK-64-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP3]], 1
+// CHECK-64-NEXT:    store i32 [[ADD3]], ptr [[ARRAYIDX]], align 4
+// CHECK-64-NEXT:    ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@__omp_offloading_30_700c39a7__Z9ftemplateIiET_i_l27
+// CHECK-SAME: () #[[ATTR0:[0-9]+]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x ptr], align 4
+// CHECK-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1:[0-9]+]], i8 2, i1 false)
+// CHECK-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
+// CHECK-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK:       user_code.entry:
+// CHECK-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 0)
+// CHECK-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-NEXT:    ret void
+// CHECK:       worker.exit:
+// CHECK-NEXT:    ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@__omp_outlined__
+// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
+// CHECK-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
+// CHECK-NEXT:    ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@__omp_offloading_30_700c39a7__Z9ftemplateIiET_i_l31
+// CHECK-SAME: (i32 noundef [[AA:%.*]]) #[[ATTR0]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    [[AA_CASTED:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 4
+// CHECK-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
+// CHECK-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK:       user_code.entry:
+// CHECK-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK-NEXT:    [[TMP2:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK-NEXT:    store i16 [[TMP2]], ptr [[AA_CASTED]], align 2
+// CHECK-NEXT:    [[TMP3:%.*]] = load i32, ptr [[AA_CASTED]], align 4
+// CHECK-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 0
+// CHECK-NEXT:    [[TMP5:%.*]] = inttoptr i32 [[TMP3]] to ptr
+// CHECK-NEXT:    store ptr [[TMP5]], ptr [[TMP4]], align 4
+// CHECK-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__1, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 1)
+// CHECK-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-NEXT:    ret void
+// CHECK:       worker.exit:
+// CHECK-NEXT:    ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@__omp_outlined__1
+// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[AA:%.*]]) #[[ATTR1]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
+// CHECK-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
+// CHECK-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK-NEXT:    [[TMP0:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK-NEXT:    [[CONV:%.*]] = sext i16 [[TMP0]] to i32
+// CHECK-NEXT:    [[ADD:%.*]] = add nsw i32 [[CONV]], 1
+// CHECK-NEXT:    [[CONV1:%.*]] = trunc i32 [[ADD]] to i16
+// CHECK-NEXT:    store i16 [[CONV1]], ptr [[AA_ADDR]], align 2
+// CHECK-NEXT:    ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@__omp_offloading_30_700c39a7__Z9ftemplateIiET_i_l36
+// CHECK-SAME: (i32 noundef [[A:%.*]], i32 noundef [[AA:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR0]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4
+// CHECK-NEXT:    [[A_CASTED:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    [[AA_CASTED:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [3 x ptr], align 4
+// CHECK-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 4
+// CHECK-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4
+// CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4
+// CHECK-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK:       user_code.entry:
+// CHECK-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A_ADDR]], align 4
+// CHECK-NEXT:    store i32 [[TMP3]], ptr [[A_CASTED]], align 4
+// CHECK-NEXT:    [[TMP4:%.*]] = load i32, ptr [[A_CASTED]], align 4
+// CHECK-NEXT:    [[TMP5:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK-NEXT:    store i16 [[TMP5]], ptr [[AA_CASTED]], align 2
+// CHECK-NEXT:    [[TMP6:%.*]] = load i32, ptr [[AA_CASTED]], align 4
+// CHECK-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 0
+// CHECK-NEXT:    [[TMP8:%.*]] = inttoptr i32 [[TMP4]] to ptr
+// CHECK-NEXT:    store ptr [[TMP8]], ptr [[TMP7]], align 4
+// CHECK-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 1
+// CHECK-NEXT:    [[TMP10:%.*]] = inttoptr i32 [[TMP6]] to ptr
+// CHECK-NEXT:    store ptr [[TMP10]], ptr [[TMP9]], align 4
+// CHECK-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 2
+// CHECK-NEXT:    store ptr [[TMP0]], ptr [[TMP11]], align 4
+// CHECK-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP2]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__2, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 3)
+// CHECK-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-NEXT:    ret void
+// CHECK:       worker.exit:
+// CHECK-NEXT:    ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@__omp_outlined__2
+// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[A:%.*]], i32 noundef [[AA:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR1]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4
+// CHECK-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
+// CHECK-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
+// CHECK-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 4
+// CHECK-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4
+// CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4
+// CHECK-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 4
+// CHECK-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], 1
+// CHECK-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 4
+// CHECK-NEXT:    [[TMP2:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK-NEXT:    [[CONV:%.*]] = sext i16 [[TMP2]] to i32
+// CHECK-NEXT:    [[ADD1:%.*]] = add nsw i32 [[CONV]], 1
+// CHECK-NEXT:    [[CONV2:%.*]] = trunc i32 [[ADD1]] to i16
+// CHECK-NEXT:    store i16 [[CONV2]], ptr [[AA_ADDR]], align 2
+// CHECK-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 2
+// CHECK-NEXT:    [[TMP3:%.*]] = load i32, ptr [[ARRAYIDX]], align 4
+// CHECK-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP3]], 1
+// CHECK-NEXT:    store i32 [[ADD3]], ptr [[ARRAYIDX]], align 4
+// CHECK-NEXT:    ret void
+//
+//
+// CHECK-32-LABEL: define {{[^@]+}}@__omp_offloading_30_700c39a7__Z9ftemplateIiET_i_l27
+// CHECK-32-SAME: () #[[ATTR0:[0-9]+]] {
+// CHECK-32-NEXT:  entry:
+// CHECK-32-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x ptr], align 4
+// CHECK-32-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1:[0-9]+]], i8 2, i1 false)
+// CHECK-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
+// CHECK-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK-32:       user_code.entry:
+// CHECK-32-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK-32-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 0)
+// CHECK-32-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-32-NEXT:    ret void
+// CHECK-32:       worker.exit:
+// CHECK-32-NEXT:    ret void
+//
+//
+// CHECK-32-LABEL: define {{[^@]+}}@__omp_outlined__
+// CHECK-32-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] {
+// CHECK-32-NEXT:  entry:
+// CHECK-32-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
+// CHECK-32-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
+// CHECK-32-NEXT:    ret void
+//
+//
+// CHECK-32-LABEL: define {{[^@]+}}@__omp_offloading_30_700c39a7__Z9ftemplateIiET_i_l31
+// CHECK-32-SAME: (i32 noundef [[AA:%.*]]) #[[ATTR0]] {
+// CHECK-32-NEXT:  entry:
+// CHECK-32-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[AA_CASTED:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 4
+// CHECK-32-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK-32-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
+// CHECK-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK-32:       user_code.entry:
+// CHECK-32-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK-32-NEXT:    [[TMP2:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK-32-NEXT:    store i16 [[TMP2]], ptr [[AA_CASTED]], align 2
+// CHECK-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[AA_CASTED]], align 4
+// CHECK-32-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 0
+// CHECK-32-NEXT:    [[TMP5:%.*]] = inttoptr i32 [[TMP3]] to ptr
+// CHECK-32-NEXT:    store ptr [[TMP5]], ptr [[TMP4]], align 4
+// CHECK-32-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__1, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 1)
+// CHECK-32-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-32-NEXT:    ret void
+// CHECK-32:       worker.exit:
+// CHECK-32-NEXT:    ret void
+//
+//
+// CHECK-32-LABEL: define {{[^@]+}}@__omp_outlined__1
+// CHECK-32-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[AA:%.*]]) #[[ATTR1]] {
+// CHECK-32-NEXT:  entry:
+// CHECK-32-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
+// CHECK-32-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
+// CHECK-32-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK-32-NEXT:    [[TMP0:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK-32-NEXT:    [[CONV:%.*]] = sext i16 [[TMP0]] to i32
+// CHECK-32-NEXT:    [[ADD:%.*]] = add nsw i32 [[CONV]], 1
+// CHECK-32-NEXT:    [[CONV1:%.*]] = trunc i32 [[ADD]] to i16
+// CHECK-32-NEXT:    store i16 [[CONV1]], ptr [[AA_ADDR]], align 2
+// CHECK-32-NEXT:    ret void
+//
+//
+// CHECK-32-LABEL: define {{[^@]+}}@__omp_offloading_30_700c39a7__Z9ftemplateIiET_i_l36
+// CHECK-32-SAME: (i32 noundef [[A:%.*]], i32 noundef [[AA:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR0]] {
+// CHECK-32-NEXT:  entry:
+// CHECK-32-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-NEXT:    [[A_CASTED:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[AA_CASTED:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [3 x ptr], align 4
+// CHECK-32-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 4
+// CHECK-32-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK-32-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4
+// CHECK-32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4
+// CHECK-32-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK-32:       user_code.entry:
+// CHECK-32-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A_ADDR]], align 4
+// CHECK-32-NEXT:    store i32 [[TMP3]], ptr [[A_CASTED]], align 4
+// CHECK-32-NEXT:    [[TMP4:%.*]] = load i32, ptr [[A_CASTED]], align 4
+// CHECK-32-NEXT:    [[TMP5:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK-32-NEXT:    store i16 [[TMP5]], ptr [[AA_CASTED]], align 2
+// CHECK-32-NEXT:    [[TMP6:%.*]] = load i32, ptr [[AA_CASTED]], align 4
+// CHECK-32-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 0
+// CHECK-32-NEXT:    [[TMP8:%.*]] = inttoptr i32 [[TMP4]] to ptr
+// CHECK-32-NEXT:    store ptr [[TMP8]], ptr [[TMP7]], align 4
+// CHECK-32-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 1
+// CHECK-32-NEXT:    [[TMP10:%.*]] = inttoptr i32 [[TMP6]] to ptr
+// CHECK-32-NEXT:    store ptr [[TMP10]], ptr [[TMP9]], align 4
+// CHECK-32-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 2
+// CHECK-32-NEXT:    store ptr [[TMP0]], ptr [[TMP11]], align 4
+// CHECK-32-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP2]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__2, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 3)
+// CHECK-32-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-32-NEXT:    ret void
+// CHECK-32:       worker.exit:
+// CHECK-32-NEXT:    ret void
+//
+//
+// CHECK-32-LABEL: define {{[^@]+}}@__omp_outlined__2
+// CHECK-32-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[A:%.*]], i32 noundef [[AA:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR1]] {
+// CHECK-32-NEXT:  entry:
+// CHECK-32-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
+// CHECK-32-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
+// CHECK-32-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 4
+// CHECK-32-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK-32-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4
+// CHECK-32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4
+// CHECK-32-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 4
+// CHECK-32-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], 1
+// CHECK-32-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 4
+// CHECK-32-NEXT:    [[TMP2:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK-32-NEXT:    [[CONV:%.*]] = sext i16 [[TMP2]] to i32
+// CHECK-32-NEXT:    [[ADD1:%.*]] = add nsw i32 [[CONV]], 1
+// CHECK-32-NEXT:    [[CONV2:%.*]] = trunc i32 [[ADD1]] to i16
+// CHECK-32-NEXT:    store i16 [[CONV2]], ptr [[AA_ADDR]], align 2
+// CHECK-32-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 2
+// CHECK-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[ARRAYIDX]], align 4
+// CHECK-32-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP3]], 1
+// CHECK-32-NEXT:    store i32 [[ADD3]], ptr [[ARRAYIDX]], align 4
+// CHECK-32-NEXT:    ret void
+//
+//
+// CHECK-32-EX-LABEL: define {{[^@]+}}@__omp_offloading_30_700c39a7__Z9ftemplateIiET_i_l27
+// CHECK-32-EX-SAME: () #[[ATTR0:[0-9]+]] {
+// CHECK-32-EX-NEXT:  entry:
+// CHECK-32-EX-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x ptr], align 4
+// CHECK-32-EX-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1:[0-9]+]], i8 2, i1 false)
+// CHECK-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
+// CHECK-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK-32-EX:       user_code.entry:
+// CHECK-32-EX-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK-32-EX-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 0)
+// CHECK-32-EX-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-32-EX-NEXT:    ret void
+// CHECK-32-EX:       worker.exit:
+// CHECK-32-EX-NEXT:    ret void
+//
+//
+// CHECK-32-EX-LABEL: define {{[^@]+}}@__omp_outlined__
+// CHECK-32-EX-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] {
+// CHECK-32-EX-NEXT:  entry:
+// CHECK-32-EX-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-EX-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-EX-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
+// CHECK-32-EX-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
+// CHECK-32-EX-NEXT:    ret void
+//
+//
+// CHECK-32-EX-LABEL: define {{[^@]+}}@__omp_offloading_30_700c39a7__Z9ftemplateIiET_i_l31
+// CHECK-32-EX-SAME: (i32 noundef [[AA:%.*]]) #[[ATTR0]] {
+// CHECK-32-EX-NEXT:  entry:
+// CHECK-32-EX-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[AA_CASTED:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 4
+// CHECK-32-EX-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK-32-EX-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
+// CHECK-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK-32-EX:       user_code.entry:
+// CHECK-32-EX-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK-32-EX-NEXT:    [[TMP2:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK-32-EX-NEXT:    store i16 [[TMP2]], ptr [[AA_CASTED]], align 2
+// CHECK-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[AA_CASTED]], align 4
+// CHECK-32-EX-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 0
+// CHECK-32-EX-NEXT:    [[TMP5:%.*]] = inttoptr i32 [[TMP3]] to ptr
+// CHECK-32-EX-NEXT:    store ptr [[TMP5]], ptr [[TMP4]], align 4
+// CHECK-32-EX-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__1, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 1)
+// CHECK-32-EX-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-32-EX-NEXT:    ret void
+// CHECK-32-EX:       worker.exit:
+// CHECK-32-EX-NEXT:    ret void
+//
+//
+// CHECK-32-EX-LABEL: define {{[^@]+}}@__omp_outlined__1
+// CHECK-32-EX-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[AA:%.*]]) #[[ATTR1]] {
+// CHECK-32-EX-NEXT:  entry:
+// CHECK-32-EX-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-EX-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-EX-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
+// CHECK-32-EX-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
+// CHECK-32-EX-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK-32-EX-NEXT:    [[TMP0:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK-32-EX-NEXT:    [[CONV:%.*]] = sext i16 [[TMP0]] to i32
+// CHECK-32-EX-NEXT:    [[ADD:%.*]] = add nsw i32 [[CONV]], 1
+// CHECK-32-EX-NEXT:    [[CONV1:%.*]] = trunc i32 [[ADD]] to i16
+// CHECK-32-EX-NEXT:    store i16 [[CONV1]], ptr [[AA_ADDR]], align 2
+// CHECK-32-EX-NEXT:    ret void
+//
+//
+// CHECK-32-EX-LABEL: define {{[^@]+}}@__omp_offloading_30_700c39a7__Z9ftemplateIiET_i_l36
+// CHECK-32-EX-SAME: (i32 noundef [[A:%.*]], i32 noundef [[AA:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR0]] {
+// CHECK-32-EX-NEXT:  entry:
+// CHECK-32-EX-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-EX-NEXT:    [[A_CASTED:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[AA_CASTED:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [3 x ptr], align 4
+// CHECK-32-EX-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 4
+// CHECK-32-EX-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK-32-EX-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4
+// CHECK-32-EX-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4
+// CHECK-32-EX-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK-32-EX:       user_code.entry:
+// CHECK-32-EX-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A_ADDR]], align 4
+// CHECK-32-EX-NEXT:    store i32 [[TMP3]], ptr [[A_CASTED]], align 4
+// CHECK-32-EX-NEXT:    [[TMP4:%.*]] = load i32, ptr [[A_CASTED]], align 4
+// CHECK-32-EX-NEXT:    [[TMP5:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK-32-EX-NEXT:    store i16 [[TMP5]], ptr [[AA_CASTED]], align 2
+// CHECK-32-EX-NEXT:    [[TMP6:%.*]] = load i32, ptr [[AA_CASTED]], align 4
+// CHECK-32-EX-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 0
+// CHECK-32-EX-NEXT:    [[TMP8:%.*]] = inttoptr i32 [[TMP4]] to ptr
+// CHECK-32-EX-NEXT:    store ptr [[TMP8]], ptr [[TMP7]], align 4
+// CHECK-32-EX-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 1
+// CHECK-32-EX-NEXT:    [[TMP10:%.*]] = inttoptr i32 [[TMP6]] to ptr
+// CHECK-32-EX-NEXT:    store ptr [[TMP10]], ptr [[TMP9]], align 4
+// CHECK-32-EX-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 2
+// CHECK-32-EX-NEXT:    store ptr [[TMP0]], ptr [[TMP11]], align 4
+// CHECK-32-EX-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP2]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__2, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 3)
+// CHECK-32-EX-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-32-EX-NEXT:    ret void
+// CHECK-32-EX:       worker.exit:
+// CHECK-32-EX-NEXT:    ret void
+//
+//
+// CHECK-32-EX-LABEL: define {{[^@]+}}@__omp_outlined__2
+// CHECK-32-EX-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[A:%.*]], i32 noundef [[AA:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR1]] {
+// CHECK-32-EX-NEXT:  entry:
+// CHECK-32-EX-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-EX-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-EX-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[AA_ADDR:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-EX-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
+// CHECK-32-EX-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
+// CHECK-32-EX-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 4
+// CHECK-32-EX-NEXT:    store i32 [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK-32-EX-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4
+// CHECK-32-EX-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4
+// CHECK-32-EX-NEXT:    [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 4
+// CHECK-32-EX-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP1]], 1
+// CHECK-32-EX-NEXT:    store i32 [[ADD]], ptr [[A_ADDR]], align 4
+// CHECK-32-EX-NEXT:    [[TMP2:%.*]] = load i16, ptr [[AA_ADDR]], align 2
+// CHECK-32-EX-NEXT:    [[CONV:%.*]] = sext i16 [[TMP2]] to i32
+// CHECK-32-EX-NEXT:    [[ADD1:%.*]] = add nsw i32 [[CONV]], 1
+// CHECK-32-EX-NEXT:    [[CONV2:%.*]] = trunc i32 [[ADD1]] to i16
+// CHECK-32-EX-NEXT:    store i16 [[CONV2]], ptr [[AA_ADDR]], align 2
+// CHECK-32-EX-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 2
+// CHECK-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[ARRAYIDX]], align 4
+// CHECK-32-EX-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP3]], 1
+// CHECK-32-EX-NEXT:    store i32 [[ADD3]], ptr [[ARRAYIDX]], align 4
+// CHECK-32-EX-NEXT:    ret void
+//

clang/test/OpenMP/nvptx_target_simd_codegen.cpp

+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs
 // Test target codegen - host bc file has to be created first.
 // 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.bc
-// 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-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
+// 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-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK45-64
 // 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.bc
-// 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-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
-// 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-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
+// 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-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix=CHECK45-32
+// 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-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefixes=CHECK45-32,CHECK45-32-EX
 
 // 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.bc
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK-64
 // 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.bc
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
-// 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-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix=CHECK-32
+// 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-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefixes=CHECK-32,CHECK-32-EX
 
 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
 
 // Check that the execution mode of all 2 target regions on the gpu is set to NonSPMD Mode.
-// CHECK-DAG: {{@__omp_offloading_.+l32}}_exec_mode = weak protected constant i8 2
-// CHECK-DAG: {{@__omp_offloading_.+l37}}_exec_mode = weak protected constant i8 2
-// CHECK-DAG: {{@__omp_offloading_.+l42}}_exec_mode = weak protected constant i8 2
-// CHECK-DAG: {{@__omp_offloading_.+l47}}_exec_mode = weak protected constant i8 2
 
 #define N 1000
 
 template<typename tx>
 tx ftemplate(int n) {
   tx a[N];
   short aa[N];
   tx b[10];

 int bar(int n){
   int a = 0;
 
   a += ftemplate<int>(n);
 
   return a;
 }
 
-// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+l32}}(
-// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false)
-// CHECK-NOT: call void @__kmpc_for_static_init
-// CHECK-NOT: call void @__kmpc_for_static_fini
-// CHECK: call void @__kmpc_target_deinit({{.*}}, i8 2)
-// CHECK: ret void
-
-// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+l37}}(
-// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false)
-// CHECK-NOT: call void @__kmpc_for_static_init
-// CHECK-NOT: call void @__kmpc_for_static_fini
-// CHECK: call void @__kmpc_target_deinit({{.*}}, i8 2)
-// CHECK: ret void
-
-// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+l42}}(
-// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false)
-// CHECK-NOT: call void @__kmpc_for_static_init
-// CHECK-NOT: call void @__kmpc_for_static_fini
-// CHECK: call void @__kmpc_target_deinit({{.*}}, i8 2)
-// CHECK: ret void
-
-// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+l47}}(
-// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false)
-// CHECK-NOT: call void @__kmpc_for_static_init
-// CHECK-NOT: call void @__kmpc_for_static_fini
-// CHECK-NOT: call void @__kmpc_nvptx_end_reduce_nowait(
-// CHECK: call void @__kmpc_target_deinit({{.*}}, i8 2)
-// CHECK: ret void
-
 #endif
+// CHECK45-64-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l29
+// CHECK45-64-SAME: (i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR0:[0-9]+]] {
+// CHECK45-64-NEXT:  entry:
+// CHECK45-64-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 8
+// CHECK45-64-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8
+// CHECK45-64-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK45-64-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4
+// CHECK45-64-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4
+// CHECK45-64-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK45-64-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK45-64-NEXT:    [[I3:%.*]] = alloca i32, align 4
+// CHECK45-64-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 8
+// CHECK45-64-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8
+// CHECK45-64-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
+// CHECK45-64-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1:[0-9]+]], i8 2, i1 false)
+// CHECK45-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK45-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK45-64:       user_code.entry:
+// CHECK45-64-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4
+// CHECK45-64-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-64-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-64-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0
+// CHECK45-64-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1
+// CHECK45-64-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1
+// CHECK45-64-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4
+// CHECK45-64-NEXT:    store i32 0, ptr [[I]], align 4
+// CHECK45-64-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-64-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]
+// CHECK45-64-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]
+// CHECK45-64:       simd.if.then:
+// CHECK45-64-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK45-64:       omp.inner.for.cond:
+// CHECK45-64-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12:![0-9]+]]
+// CHECK45-64-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK45-64-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK45-64-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]
+// CHECK45-64-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK45-64:       omp.inner.for.body:
+// CHECK45-64-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK45-64-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
+// CHECK45-64-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK45-64-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK45-64-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK45-64-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP8]] to i64
+// CHECK45-64-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
+// CHECK45-64-NEXT:    store i32 1, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK45-64-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK45-64:       omp.body.continue:
+// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK45-64:       omp.inner.for.inc:
+// CHECK45-64-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK45-64-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP9]], 1
+// CHECK45-64-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP13:![0-9]+]]
+// CHECK45-64:       worker.exit:
+// CHECK45-64-NEXT:    ret void
+// CHECK45-64:       omp.inner.for.end:
+// CHECK45-64-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-64-NEXT:    [[SUB7:%.*]] = sub nsw i32 [[TMP10]], 0
+// CHECK45-64-NEXT:    [[DIV8:%.*]] = sdiv i32 [[SUB7]], 1
+// CHECK45-64-NEXT:    [[MUL9:%.*]] = mul nsw i32 [[DIV8]], 1
+// CHECK45-64-NEXT:    [[ADD10:%.*]] = add nsw i32 0, [[MUL9]]
+// CHECK45-64-NEXT:    store i32 [[ADD10]], ptr [[I3]], align 4
+// CHECK45-64-NEXT:    br label [[SIMD_IF_END]]
+// CHECK45-64:       simd.if.end:
+// CHECK45-64-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK45-64-NEXT:    ret void
+//
+//
+// CHECK45-64-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l34
+// CHECK45-64-SAME: (i64 noundef [[N:%.*]], ptr noundef nonnull align 2 dereferenceable(2000) [[AA:%.*]]) #[[ATTR0]] {
+// CHECK45-64-NEXT:  entry:
+// CHECK45-64-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 8
+// CHECK45-64-NEXT:    [[AA_ADDR:%.*]] = alloca ptr, align 8
+// CHECK45-64-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK45-64-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4
+// CHECK45-64-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4
+// CHECK45-64-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK45-64-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK45-64-NEXT:    [[I3:%.*]] = alloca i32, align 4
+// CHECK45-64-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 8
+// CHECK45-64-NEXT:    store ptr [[AA]], ptr [[AA_ADDR]], align 8
+// CHECK45-64-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 8
+// CHECK45-64-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK45-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK45-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK45-64:       user_code.entry:
+// CHECK45-64-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4
+// CHECK45-64-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-64-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-64-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0
+// CHECK45-64-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1
+// CHECK45-64-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1
+// CHECK45-64-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4
+// CHECK45-64-NEXT:    store i32 0, ptr [[I]], align 4
+// CHECK45-64-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-64-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]
+// CHECK45-64-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]
+// CHECK45-64:       simd.if.then:
+// CHECK45-64-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK45-64:       omp.inner.for.cond:
+// CHECK45-64-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16:![0-9]+]]
+// CHECK45-64-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-64-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK45-64-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]
+// CHECK45-64-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK45-64:       omp.inner.for.body:
+// CHECK45-64-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-64-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
+// CHECK45-64-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK45-64-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-64-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-64-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP8]] to i64
+// CHECK45-64-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i16], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
+// CHECK45-64-NEXT:    [[TMP9:%.*]] = load i16, ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-64-NEXT:    [[CONV:%.*]] = sext i16 [[TMP9]] to i32
+// CHECK45-64-NEXT:    [[ADD6:%.*]] = add nsw i32 [[CONV]], 1
+// CHECK45-64-NEXT:    [[CONV7:%.*]] = trunc i32 [[ADD6]] to i16
+// CHECK45-64-NEXT:    store i16 [[CONV7]], ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-64-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK45-64:       omp.body.continue:
+// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK45-64:       omp.inner.for.inc:
+// CHECK45-64-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-64-NEXT:    [[ADD8:%.*]] = add nsw i32 [[TMP10]], 1
+// CHECK45-64-NEXT:    store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP17:![0-9]+]]
+// CHECK45-64:       worker.exit:
+// CHECK45-64-NEXT:    ret void
+// CHECK45-64:       omp.inner.for.end:
+// CHECK45-64-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-64-NEXT:    [[SUB9:%.*]] = sub nsw i32 [[TMP11]], 0
+// CHECK45-64-NEXT:    [[DIV10:%.*]] = sdiv i32 [[SUB9]], 1
+// CHECK45-64-NEXT:    [[MUL11:%.*]] = mul nsw i32 [[DIV10]], 1
+// CHECK45-64-NEXT:    [[ADD12:%.*]] = add nsw i32 0, [[MUL11]]
+// CHECK45-64-NEXT:    store i32 [[ADD12]], ptr [[I3]], align 4
+// CHECK45-64-NEXT:    br label [[SIMD_IF_END]]
+// CHECK45-64:       simd.if.end:
+// CHECK45-64-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK45-64-NEXT:    ret void
+//
+//
+// CHECK45-64-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l39
+// CHECK45-64-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR0]] {
+// CHECK45-64-NEXT:  entry:
+// CHECK45-64-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8
+// CHECK45-64-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK45-64-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK45-64-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK45-64-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8
+// CHECK45-64-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 8
+// CHECK45-64-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK45-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK45-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK45-64:       user_code.entry:
+// CHECK45-64-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK45-64:       omp.inner.for.cond:
+// CHECK45-64-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19:![0-9]+]]
+// CHECK45-64-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP2]], 10
+// CHECK45-64-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK45-64:       omp.inner.for.body:
+// CHECK45-64-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK45-64-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1
+// CHECK45-64-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK45-64-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK45-64-NEXT:    [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK45-64-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP4]] to i64
+// CHECK45-64-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
+// CHECK45-64-NEXT:    [[TMP5:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK45-64-NEXT:    [[ADD1:%.*]] = add nsw i32 [[TMP5]], 1
+// CHECK45-64-NEXT:    store i32 [[ADD1]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK45-64-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK45-64:       omp.body.continue:
+// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK45-64:       omp.inner.for.inc:
+// CHECK45-64-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK45-64-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK45-64-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP20:![0-9]+]]
+// CHECK45-64:       worker.exit:
+// CHECK45-64-NEXT:    ret void
+// CHECK45-64:       omp.inner.for.end:
+// CHECK45-64-NEXT:    store i32 10, ptr [[I]], align 4
+// CHECK45-64-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK45-64-NEXT:    ret void
+//
+//
+// CHECK45-64-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l44
+// CHECK45-64-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]]) #[[ATTR0]] {
+// CHECK45-64-NEXT:  entry:
+// CHECK45-64-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8
+// CHECK45-64-NEXT:    [[N_ADDR:%.*]] = alloca ptr, align 8
+// CHECK45-64-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK45-64-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK45-64-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK45-64-NEXT:    [[N1:%.*]] = alloca i32, align 4
+// CHECK45-64-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8
+// CHECK45-64-NEXT:    store ptr [[N]], ptr [[N_ADDR]], align 8
+// CHECK45-64-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 8
+// CHECK45-64-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[N_ADDR]], align 8
+// CHECK45-64-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK45-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -1
+// CHECK45-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK45-64:       user_code.entry:
+// CHECK45-64-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK45-64-NEXT:    store i32 0, ptr [[N1]], align 4
+// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK45-64:       omp.inner.for.cond:
+// CHECK45-64-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22:![0-9]+]]
+// CHECK45-64-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP3]], 10
+// CHECK45-64-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK45-64:       omp.inner.for.body:
+// CHECK45-64-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK45-64-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP4]], 1
+// CHECK45-64-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK45-64-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK45-64-NEXT:    [[TMP5:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK45-64-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP5]] to i64
+// CHECK45-64-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
+// CHECK45-64-NEXT:    [[TMP6:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK45-64-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK45-64-NEXT:    store i32 [[ADD2]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK45-64-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK45-64:       omp.body.continue:
+// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK45-64:       omp.inner.for.inc:
+// CHECK45-64-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK45-64-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP7]], 1
+// CHECK45-64-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK45-64-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP23:![0-9]+]]
+// CHECK45-64:       worker.exit:
+// CHECK45-64-NEXT:    ret void
+// CHECK45-64:       omp.inner.for.end:
+// CHECK45-64-NEXT:    store i32 10, ptr [[I]], align 4
+// CHECK45-64-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP1]], align 4
+// CHECK45-64-NEXT:    [[TMP9:%.*]] = load i32, ptr [[N1]], align 4
+// CHECK45-64-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]
+// CHECK45-64-NEXT:    store i32 [[ADD4]], ptr [[TMP1]], align 4
+// CHECK45-64-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK45-64-NEXT:    ret void
+//
+//
+// CHECK45-32-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l29
+// CHECK45-32-SAME: (i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR0:[0-9]+]] {
+// CHECK45-32-NEXT:  entry:
+// CHECK45-32-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[I3:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4
+// CHECK45-32-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 4
+// CHECK45-32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 4
+// CHECK45-32-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1:[0-9]+]], i8 2, i1 false)
+// CHECK45-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK45-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK45-32:       user_code.entry:
+// CHECK45-32-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4
+// CHECK45-32-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-32-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0
+// CHECK45-32-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1
+// CHECK45-32-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1
+// CHECK45-32-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4
+// CHECK45-32-NEXT:    store i32 0, ptr [[I]], align 4
+// CHECK45-32-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-32-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]
+// CHECK45-32-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]
+// CHECK45-32:       simd.if.then:
+// CHECK45-32-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK45-32:       omp.inner.for.cond:
+// CHECK45-32-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12:![0-9]+]]
+// CHECK45-32-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK45-32-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK45-32-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]
+// CHECK45-32-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK45-32:       omp.inner.for.body:
+// CHECK45-32-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK45-32-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
+// CHECK45-32-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK45-32-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK45-32-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK45-32-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i32 0, i32 [[TMP8]]
+// CHECK45-32-NEXT:    store i32 1, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK45-32-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK45-32:       omp.body.continue:
+// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK45-32:       omp.inner.for.inc:
+// CHECK45-32-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK45-32-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP9]], 1
+// CHECK45-32-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP13:![0-9]+]]
+// CHECK45-32:       worker.exit:
+// CHECK45-32-NEXT:    ret void
+// CHECK45-32:       omp.inner.for.end:
+// CHECK45-32-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-32-NEXT:    [[SUB7:%.*]] = sub nsw i32 [[TMP10]], 0
+// CHECK45-32-NEXT:    [[DIV8:%.*]] = sdiv i32 [[SUB7]], 1
+// CHECK45-32-NEXT:    [[MUL9:%.*]] = mul nsw i32 [[DIV8]], 1
+// CHECK45-32-NEXT:    [[ADD10:%.*]] = add nsw i32 0, [[MUL9]]
+// CHECK45-32-NEXT:    store i32 [[ADD10]], ptr [[I3]], align 4
+// CHECK45-32-NEXT:    br label [[SIMD_IF_END]]
+// CHECK45-32:       simd.if.end:
+// CHECK45-32-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK45-32-NEXT:    ret void
+//
+//
+// CHECK45-32-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l34
+// CHECK45-32-SAME: (i32 noundef [[N:%.*]], ptr noundef nonnull align 2 dereferenceable(2000) [[AA:%.*]]) #[[ATTR0]] {
+// CHECK45-32-NEXT:  entry:
+// CHECK45-32-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[AA_ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[I3:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4
+// CHECK45-32-NEXT:    store ptr [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK45-32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 4
+// CHECK45-32-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK45-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK45-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK45-32:       user_code.entry:
+// CHECK45-32-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4
+// CHECK45-32-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-32-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0
+// CHECK45-32-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1
+// CHECK45-32-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1
+// CHECK45-32-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4
+// CHECK45-32-NEXT:    store i32 0, ptr [[I]], align 4
+// CHECK45-32-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-32-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]
+// CHECK45-32-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]
+// CHECK45-32:       simd.if.then:
+// CHECK45-32-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK45-32:       omp.inner.for.cond:
+// CHECK45-32-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16:![0-9]+]]
+// CHECK45-32-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-32-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK45-32-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]
+// CHECK45-32-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK45-32:       omp.inner.for.body:
+// CHECK45-32-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-32-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
+// CHECK45-32-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK45-32-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-32-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-32-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i16], ptr [[TMP0]], i32 0, i32 [[TMP8]]
+// CHECK45-32-NEXT:    [[TMP9:%.*]] = load i16, ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-32-NEXT:    [[CONV:%.*]] = sext i16 [[TMP9]] to i32
+// CHECK45-32-NEXT:    [[ADD6:%.*]] = add nsw i32 [[CONV]], 1
+// CHECK45-32-NEXT:    [[CONV7:%.*]] = trunc i32 [[ADD6]] to i16
+// CHECK45-32-NEXT:    store i16 [[CONV7]], ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-32-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK45-32:       omp.body.continue:
+// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK45-32:       omp.inner.for.inc:
+// CHECK45-32-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-32-NEXT:    [[ADD8:%.*]] = add nsw i32 [[TMP10]], 1
+// CHECK45-32-NEXT:    store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP17:![0-9]+]]
+// CHECK45-32:       worker.exit:
+// CHECK45-32-NEXT:    ret void
+// CHECK45-32:       omp.inner.for.end:
+// CHECK45-32-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-32-NEXT:    [[SUB9:%.*]] = sub nsw i32 [[TMP11]], 0
+// CHECK45-32-NEXT:    [[DIV10:%.*]] = sdiv i32 [[SUB9]], 1
+// CHECK45-32-NEXT:    [[MUL11:%.*]] = mul nsw i32 [[DIV10]], 1
+// CHECK45-32-NEXT:    [[ADD12:%.*]] = add nsw i32 0, [[MUL11]]
+// CHECK45-32-NEXT:    store i32 [[ADD12]], ptr [[I3]], align 4
+// CHECK45-32-NEXT:    br label [[SIMD_IF_END]]
+// CHECK45-32:       simd.if.end:
+// CHECK45-32-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK45-32-NEXT:    ret void
+//
+//
+// CHECK45-32-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l39
+// CHECK45-32-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR0]] {
+// CHECK45-32-NEXT:  entry:
+// CHECK45-32-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4
+// CHECK45-32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4
+// CHECK45-32-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK45-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK45-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK45-32:       user_code.entry:
+// CHECK45-32-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK45-32:       omp.inner.for.cond:
+// CHECK45-32-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19:![0-9]+]]
+// CHECK45-32-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP2]], 10
+// CHECK45-32-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK45-32:       omp.inner.for.body:
+// CHECK45-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK45-32-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1
+// CHECK45-32-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK45-32-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK45-32-NEXT:    [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK45-32-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP4]]
+// CHECK45-32-NEXT:    [[TMP5:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK45-32-NEXT:    [[ADD1:%.*]] = add nsw i32 [[TMP5]], 1
+// CHECK45-32-NEXT:    store i32 [[ADD1]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK45-32-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK45-32:       omp.body.continue:
+// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK45-32:       omp.inner.for.inc:
+// CHECK45-32-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK45-32-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK45-32-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP20:![0-9]+]]
+// CHECK45-32:       worker.exit:
+// CHECK45-32-NEXT:    ret void
+// CHECK45-32:       omp.inner.for.end:
+// CHECK45-32-NEXT:    store i32 10, ptr [[I]], align 4
+// CHECK45-32-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK45-32-NEXT:    ret void
+//
+//
+// CHECK45-32-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l44
+// CHECK45-32-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]]) #[[ATTR0]] {
+// CHECK45-32-NEXT:  entry:
+// CHECK45-32-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-NEXT:    [[N_ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    [[N1:%.*]] = alloca i32, align 4
+// CHECK45-32-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4
+// CHECK45-32-NEXT:    store ptr [[N]], ptr [[N_ADDR]], align 4
+// CHECK45-32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4
+// CHECK45-32-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[N_ADDR]], align 4
+// CHECK45-32-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK45-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -1
+// CHECK45-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK45-32:       user_code.entry:
+// CHECK45-32-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK45-32-NEXT:    store i32 0, ptr [[N1]], align 4
+// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK45-32:       omp.inner.for.cond:
+// CHECK45-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22:![0-9]+]]
+// CHECK45-32-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP3]], 10
+// CHECK45-32-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK45-32:       omp.inner.for.body:
+// CHECK45-32-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK45-32-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP4]], 1
+// CHECK45-32-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK45-32-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK45-32-NEXT:    [[TMP5:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK45-32-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP5]]
+// CHECK45-32-NEXT:    [[TMP6:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK45-32-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK45-32-NEXT:    store i32 [[ADD2]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK45-32-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK45-32:       omp.body.continue:
+// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK45-32:       omp.inner.for.inc:
+// CHECK45-32-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK45-32-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP7]], 1
+// CHECK45-32-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK45-32-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP23:![0-9]+]]
+// CHECK45-32:       worker.exit:
+// CHECK45-32-NEXT:    ret void
+// CHECK45-32:       omp.inner.for.end:
+// CHECK45-32-NEXT:    store i32 10, ptr [[I]], align 4
+// CHECK45-32-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP1]], align 4
+// CHECK45-32-NEXT:    [[TMP9:%.*]] = load i32, ptr [[N1]], align 4
+// CHECK45-32-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]
+// CHECK45-32-NEXT:    store i32 [[ADD4]], ptr [[TMP1]], align 4
+// CHECK45-32-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK45-32-NEXT:    ret void
+//
+//
+// CHECK45-32-EX-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l29
+// CHECK45-32-EX-SAME: (i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR0:[0-9]+]] {
+// CHECK45-32-EX-NEXT:  entry:
+// CHECK45-32-EX-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-EX-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[I3:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1:[0-9]+]], i8 2, i1 false)
+// CHECK45-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK45-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK45-32-EX:       user_code.entry:
+// CHECK45-32-EX-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-32-EX-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0
+// CHECK45-32-EX-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1
+// CHECK45-32-EX-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1
+// CHECK45-32-EX-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4
+// CHECK45-32-EX-NEXT:    store i32 0, ptr [[I]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-32-EX-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]
+// CHECK45-32-EX-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]
+// CHECK45-32-EX:       simd.if.then:
+// CHECK45-32-EX-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK45-32-EX:       omp.inner.for.cond:
+// CHECK45-32-EX-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12:![0-9]+]]
+// CHECK45-32-EX-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK45-32-EX-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK45-32-EX-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]
+// CHECK45-32-EX-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK45-32-EX:       omp.inner.for.body:
+// CHECK45-32-EX-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK45-32-EX-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
+// CHECK45-32-EX-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK45-32-EX-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK45-32-EX-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK45-32-EX-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i32 0, i32 [[TMP8]]
+// CHECK45-32-EX-NEXT:    store i32 1, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK45-32-EX-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK45-32-EX:       omp.body.continue:
+// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK45-32-EX:       omp.inner.for.inc:
+// CHECK45-32-EX-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK45-32-EX-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP9]], 1
+// CHECK45-32-EX-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP13:![0-9]+]]
+// CHECK45-32-EX:       worker.exit:
+// CHECK45-32-EX-NEXT:    ret void
+// CHECK45-32-EX:       omp.inner.for.end:
+// CHECK45-32-EX-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-32-EX-NEXT:    [[SUB7:%.*]] = sub nsw i32 [[TMP10]], 0
+// CHECK45-32-EX-NEXT:    [[DIV8:%.*]] = sdiv i32 [[SUB7]], 1
+// CHECK45-32-EX-NEXT:    [[MUL9:%.*]] = mul nsw i32 [[DIV8]], 1
+// CHECK45-32-EX-NEXT:    [[ADD10:%.*]] = add nsw i32 0, [[MUL9]]
+// CHECK45-32-EX-NEXT:    store i32 [[ADD10]], ptr [[I3]], align 4
+// CHECK45-32-EX-NEXT:    br label [[SIMD_IF_END]]
+// CHECK45-32-EX:       simd.if.end:
+// CHECK45-32-EX-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK45-32-EX-NEXT:    ret void
+//
+//
+// CHECK45-32-EX-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l34
+// CHECK45-32-EX-SAME: (i32 noundef [[N:%.*]], ptr noundef nonnull align 2 dereferenceable(2000) [[AA:%.*]]) #[[ATTR0]] {
+// CHECK45-32-EX-NEXT:  entry:
+// CHECK45-32-EX-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[AA_ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-EX-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[I3:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    store ptr [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK45-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK45-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK45-32-EX:       user_code.entry:
+// CHECK45-32-EX-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-32-EX-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0
+// CHECK45-32-EX-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1
+// CHECK45-32-EX-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1
+// CHECK45-32-EX-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4
+// CHECK45-32-EX-NEXT:    store i32 0, ptr [[I]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-32-EX-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]
+// CHECK45-32-EX-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]
+// CHECK45-32-EX:       simd.if.then:
+// CHECK45-32-EX-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK45-32-EX:       omp.inner.for.cond:
+// CHECK45-32-EX-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16:![0-9]+]]
+// CHECK45-32-EX-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-32-EX-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK45-32-EX-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]
+// CHECK45-32-EX-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK45-32-EX:       omp.inner.for.body:
+// CHECK45-32-EX-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-32-EX-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
+// CHECK45-32-EX-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK45-32-EX-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-32-EX-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-32-EX-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i16], ptr [[TMP0]], i32 0, i32 [[TMP8]]
+// CHECK45-32-EX-NEXT:    [[TMP9:%.*]] = load i16, ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-32-EX-NEXT:    [[CONV:%.*]] = sext i16 [[TMP9]] to i32
+// CHECK45-32-EX-NEXT:    [[ADD6:%.*]] = add nsw i32 [[CONV]], 1
+// CHECK45-32-EX-NEXT:    [[CONV7:%.*]] = trunc i32 [[ADD6]] to i16
+// CHECK45-32-EX-NEXT:    store i16 [[CONV7]], ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-32-EX-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK45-32-EX:       omp.body.continue:
+// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK45-32-EX:       omp.inner.for.inc:
+// CHECK45-32-EX-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-32-EX-NEXT:    [[ADD8:%.*]] = add nsw i32 [[TMP10]], 1
+// CHECK45-32-EX-NEXT:    store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP17:![0-9]+]]
+// CHECK45-32-EX:       worker.exit:
+// CHECK45-32-EX-NEXT:    ret void
+// CHECK45-32-EX:       omp.inner.for.end:
+// CHECK45-32-EX-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK45-32-EX-NEXT:    [[SUB9:%.*]] = sub nsw i32 [[TMP11]], 0
+// CHECK45-32-EX-NEXT:    [[DIV10:%.*]] = sdiv i32 [[SUB9]], 1
+// CHECK45-32-EX-NEXT:    [[MUL11:%.*]] = mul nsw i32 [[DIV10]], 1
+// CHECK45-32-EX-NEXT:    [[ADD12:%.*]] = add nsw i32 0, [[MUL11]]
+// CHECK45-32-EX-NEXT:    store i32 [[ADD12]], ptr [[I3]], align 4
+// CHECK45-32-EX-NEXT:    br label [[SIMD_IF_END]]
+// CHECK45-32-EX:       simd.if.end:
+// CHECK45-32-EX-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK45-32-EX-NEXT:    ret void
+//
+//
+// CHECK45-32-EX-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l39
+// CHECK45-32-EX-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR0]] {
+// CHECK45-32-EX-NEXT:  entry:
+// CHECK45-32-EX-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-EX-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK45-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK45-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK45-32-EX:       user_code.entry:
+// CHECK45-32-EX-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK45-32-EX:       omp.inner.for.cond:
+// CHECK45-32-EX-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19:![0-9]+]]
+// CHECK45-32-EX-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP2]], 10
+// CHECK45-32-EX-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK45-32-EX:       omp.inner.for.body:
+// CHECK45-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK45-32-EX-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1
+// CHECK45-32-EX-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK45-32-EX-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK45-32-EX-NEXT:    [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK45-32-EX-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP4]]
+// CHECK45-32-EX-NEXT:    [[TMP5:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK45-32-EX-NEXT:    [[ADD1:%.*]] = add nsw i32 [[TMP5]], 1
+// CHECK45-32-EX-NEXT:    store i32 [[ADD1]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK45-32-EX-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK45-32-EX:       omp.body.continue:
+// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK45-32-EX:       omp.inner.for.inc:
+// CHECK45-32-EX-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK45-32-EX-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK45-32-EX-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP20:![0-9]+]]
+// CHECK45-32-EX:       worker.exit:
+// CHECK45-32-EX-NEXT:    ret void
+// CHECK45-32-EX:       omp.inner.for.end:
+// CHECK45-32-EX-NEXT:    store i32 10, ptr [[I]], align 4
+// CHECK45-32-EX-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK45-32-EX-NEXT:    ret void
+//
+//
+// CHECK45-32-EX-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l44
+// CHECK45-32-EX-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]]) #[[ATTR0]] {
+// CHECK45-32-EX-NEXT:  entry:
+// CHECK45-32-EX-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-EX-NEXT:    [[N_ADDR:%.*]] = alloca ptr, align 4
+// CHECK45-32-EX-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    [[N1:%.*]] = alloca i32, align 4
+// CHECK45-32-EX-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    store ptr [[N]], ptr [[N_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[N_ADDR]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK45-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -1
+// CHECK45-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK45-32-EX:       user_code.entry:
+// CHECK45-32-EX-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK45-32-EX-NEXT:    store i32 0, ptr [[N1]], align 4
+// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK45-32-EX:       omp.inner.for.cond:
+// CHECK45-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22:![0-9]+]]
+// CHECK45-32-EX-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP3]], 10
+// CHECK45-32-EX-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK45-32-EX:       omp.inner.for.body:
+// CHECK45-32-EX-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK45-32-EX-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP4]], 1
+// CHECK45-32-EX-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK45-32-EX-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK45-32-EX-NEXT:    [[TMP5:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK45-32-EX-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP5]]
+// CHECK45-32-EX-NEXT:    [[TMP6:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK45-32-EX-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK45-32-EX-NEXT:    store i32 [[ADD2]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK45-32-EX-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK45-32-EX:       omp.body.continue:
+// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK45-32-EX:       omp.inner.for.inc:
+// CHECK45-32-EX-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK45-32-EX-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP7]], 1
+// CHECK45-32-EX-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK45-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP23:![0-9]+]]
+// CHECK45-32-EX:       worker.exit:
+// CHECK45-32-EX-NEXT:    ret void
+// CHECK45-32-EX:       omp.inner.for.end:
+// CHECK45-32-EX-NEXT:    store i32 10, ptr [[I]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP1]], align 4
+// CHECK45-32-EX-NEXT:    [[TMP9:%.*]] = load i32, ptr [[N1]], align 4
+// CHECK45-32-EX-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]
+// CHECK45-32-EX-NEXT:    store i32 [[ADD4]], ptr [[TMP1]], align 4
+// CHECK45-32-EX-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK45-32-EX-NEXT:    ret void
+//
+//
+// CHECK-64-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l29
+// CHECK-64-SAME: (i64 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR0:[0-9]+]] {
+// CHECK-64-NEXT:  entry:
+// CHECK-64-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 8
+// CHECK-64-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8
+// CHECK-64-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK-64-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4
+// CHECK-64-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4
+// CHECK-64-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK-64-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK-64-NEXT:    [[I3:%.*]] = alloca i32, align 4
+// CHECK-64-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 8
+// CHECK-64-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8
+// CHECK-64-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
+// CHECK-64-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1:[0-9]+]], i8 2, i1 false)
+// CHECK-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK-64:       user_code.entry:
+// CHECK-64-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4
+// CHECK-64-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-64-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-64-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0
+// CHECK-64-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1
+// CHECK-64-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1
+// CHECK-64-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4
+// CHECK-64-NEXT:    store i32 0, ptr [[I]], align 4
+// CHECK-64-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-64-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]
+// CHECK-64-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]
+// CHECK-64:       simd.if.then:
+// CHECK-64-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK-64:       omp.inner.for.cond:
+// CHECK-64-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12:![0-9]+]]
+// CHECK-64-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK-64-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK-64-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]
+// CHECK-64-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK-64:       omp.inner.for.body:
+// CHECK-64-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK-64-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
+// CHECK-64-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK-64-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK-64-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK-64-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP8]] to i64
+// CHECK-64-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
+// CHECK-64-NEXT:    store i32 1, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK-64-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK-64:       omp.body.continue:
+// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK-64:       omp.inner.for.inc:
+// CHECK-64-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK-64-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP9]], 1
+// CHECK-64-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP13:![0-9]+]]
+// CHECK-64:       worker.exit:
+// CHECK-64-NEXT:    ret void
+// CHECK-64:       omp.inner.for.end:
+// CHECK-64-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-64-NEXT:    [[SUB7:%.*]] = sub nsw i32 [[TMP10]], 0
+// CHECK-64-NEXT:    [[DIV8:%.*]] = sdiv i32 [[SUB7]], 1
+// CHECK-64-NEXT:    [[MUL9:%.*]] = mul nsw i32 [[DIV8]], 1
+// CHECK-64-NEXT:    [[ADD10:%.*]] = add nsw i32 0, [[MUL9]]
+// CHECK-64-NEXT:    store i32 [[ADD10]], ptr [[I3]], align 4
+// CHECK-64-NEXT:    br label [[SIMD_IF_END]]
+// CHECK-64:       simd.if.end:
+// CHECK-64-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-64-NEXT:    ret void
+//
+//
+// CHECK-64-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l34
+// CHECK-64-SAME: (i64 noundef [[N:%.*]], ptr noundef nonnull align 2 dereferenceable(2000) [[AA:%.*]]) #[[ATTR0]] {
+// CHECK-64-NEXT:  entry:
+// CHECK-64-NEXT:    [[N_ADDR:%.*]] = alloca i64, align 8
+// CHECK-64-NEXT:    [[AA_ADDR:%.*]] = alloca ptr, align 8
+// CHECK-64-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK-64-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4
+// CHECK-64-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4
+// CHECK-64-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK-64-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK-64-NEXT:    [[I3:%.*]] = alloca i32, align 4
+// CHECK-64-NEXT:    store i64 [[N]], ptr [[N_ADDR]], align 8
+// CHECK-64-NEXT:    store ptr [[AA]], ptr [[AA_ADDR]], align 8
+// CHECK-64-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 8
+// CHECK-64-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK-64:       user_code.entry:
+// CHECK-64-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4
+// CHECK-64-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-64-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-64-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0
+// CHECK-64-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1
+// CHECK-64-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1
+// CHECK-64-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4
+// CHECK-64-NEXT:    store i32 0, ptr [[I]], align 4
+// CHECK-64-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-64-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]
+// CHECK-64-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]
+// CHECK-64:       simd.if.then:
+// CHECK-64-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK-64:       omp.inner.for.cond:
+// CHECK-64-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16:![0-9]+]]
+// CHECK-64-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK-64-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK-64-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]
+// CHECK-64-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK-64:       omp.inner.for.body:
+// CHECK-64-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK-64-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
+// CHECK-64-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK-64-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK-64-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK-64-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP8]] to i64
+// CHECK-64-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i16], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
+// CHECK-64-NEXT:    [[TMP9:%.*]] = load i16, ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP16]]
+// CHECK-64-NEXT:    [[CONV:%.*]] = sext i16 [[TMP9]] to i32
+// CHECK-64-NEXT:    [[ADD6:%.*]] = add nsw i32 [[CONV]], 1
+// CHECK-64-NEXT:    [[CONV7:%.*]] = trunc i32 [[ADD6]] to i16
+// CHECK-64-NEXT:    store i16 [[CONV7]], ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP16]]
+// CHECK-64-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK-64:       omp.body.continue:
+// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK-64:       omp.inner.for.inc:
+// CHECK-64-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK-64-NEXT:    [[ADD8:%.*]] = add nsw i32 [[TMP10]], 1
+// CHECK-64-NEXT:    store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP17:![0-9]+]]
+// CHECK-64:       worker.exit:
+// CHECK-64-NEXT:    ret void
+// CHECK-64:       omp.inner.for.end:
+// CHECK-64-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-64-NEXT:    [[SUB9:%.*]] = sub nsw i32 [[TMP11]], 0
+// CHECK-64-NEXT:    [[DIV10:%.*]] = sdiv i32 [[SUB9]], 1
+// CHECK-64-NEXT:    [[MUL11:%.*]] = mul nsw i32 [[DIV10]], 1
+// CHECK-64-NEXT:    [[ADD12:%.*]] = add nsw i32 0, [[MUL11]]
+// CHECK-64-NEXT:    store i32 [[ADD12]], ptr [[I3]], align 4
+// CHECK-64-NEXT:    br label [[SIMD_IF_END]]
+// CHECK-64:       simd.if.end:
+// CHECK-64-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-64-NEXT:    ret void
+//
+//
+// CHECK-64-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l39
+// CHECK-64-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR0]] {
+// CHECK-64-NEXT:  entry:
+// CHECK-64-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8
+// CHECK-64-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK-64-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK-64-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK-64-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8
+// CHECK-64-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 8
+// CHECK-64-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK-64:       user_code.entry:
+// CHECK-64-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK-64:       omp.inner.for.cond:
+// CHECK-64-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19:![0-9]+]]
+// CHECK-64-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP2]], 10
+// CHECK-64-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK-64:       omp.inner.for.body:
+// CHECK-64-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK-64-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1
+// CHECK-64-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK-64-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK-64-NEXT:    [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK-64-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP4]] to i64
+// CHECK-64-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
+// CHECK-64-NEXT:    [[TMP5:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK-64-NEXT:    [[ADD1:%.*]] = add nsw i32 [[TMP5]], 1
+// CHECK-64-NEXT:    store i32 [[ADD1]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK-64-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK-64:       omp.body.continue:
+// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK-64:       omp.inner.for.inc:
+// CHECK-64-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK-64-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK-64-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP20:![0-9]+]]
+// CHECK-64:       worker.exit:
+// CHECK-64-NEXT:    ret void
+// CHECK-64:       omp.inner.for.end:
+// CHECK-64-NEXT:    store i32 10, ptr [[I]], align 4
+// CHECK-64-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-64-NEXT:    ret void
+//
+//
+// CHECK-64-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l44
+// CHECK-64-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]]) #[[ATTR0]] {
+// CHECK-64-NEXT:  entry:
+// CHECK-64-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 8
+// CHECK-64-NEXT:    [[N_ADDR:%.*]] = alloca ptr, align 8
+// CHECK-64-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK-64-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK-64-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK-64-NEXT:    [[N1:%.*]] = alloca i32, align 4
+// CHECK-64-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 8
+// CHECK-64-NEXT:    store ptr [[N]], ptr [[N_ADDR]], align 8
+// CHECK-64-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 8
+// CHECK-64-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[N_ADDR]], align 8
+// CHECK-64-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -1
+// CHECK-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK-64:       user_code.entry:
+// CHECK-64-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK-64-NEXT:    store i32 0, ptr [[N1]], align 4
+// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK-64:       omp.inner.for.cond:
+// CHECK-64-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22:![0-9]+]]
+// CHECK-64-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP3]], 10
+// CHECK-64-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK-64:       omp.inner.for.body:
+// CHECK-64-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK-64-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP4]], 1
+// CHECK-64-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK-64-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK-64-NEXT:    [[TMP5:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK-64-NEXT:    [[IDXPROM:%.*]] = sext i32 [[TMP5]] to i64
+// CHECK-64-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
+// CHECK-64-NEXT:    [[TMP6:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK-64-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK-64-NEXT:    store i32 [[ADD2]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK-64-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK-64:       omp.body.continue:
+// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK-64:       omp.inner.for.inc:
+// CHECK-64-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK-64-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP7]], 1
+// CHECK-64-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK-64-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP23:![0-9]+]]
+// CHECK-64:       worker.exit:
+// CHECK-64-NEXT:    ret void
+// CHECK-64:       omp.inner.for.end:
+// CHECK-64-NEXT:    store i32 10, ptr [[I]], align 4
+// CHECK-64-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP1]], align 4
+// CHECK-64-NEXT:    [[TMP9:%.*]] = load i32, ptr [[N1]], align 4
+// CHECK-64-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]
+// CHECK-64-NEXT:    store i32 [[ADD4]], ptr [[TMP1]], align 4
+// CHECK-64-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-64-NEXT:    ret void
+//
+//
+// CHECK-32-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l29
+// CHECK-32-SAME: (i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR0:[0-9]+]] {
+// CHECK-32-NEXT:  entry:
+// CHECK-32-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[I3:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4
+// CHECK-32-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 4
+// CHECK-32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 4
+// CHECK-32-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1:[0-9]+]], i8 2, i1 false)
+// CHECK-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK-32:       user_code.entry:
+// CHECK-32-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4
+// CHECK-32-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-32-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0
+// CHECK-32-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1
+// CHECK-32-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1
+// CHECK-32-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4
+// CHECK-32-NEXT:    store i32 0, ptr [[I]], align 4
+// CHECK-32-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-32-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]
+// CHECK-32-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]
+// CHECK-32:       simd.if.then:
+// CHECK-32-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK-32:       omp.inner.for.cond:
+// CHECK-32-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12:![0-9]+]]
+// CHECK-32-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK-32-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK-32-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]
+// CHECK-32-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK-32:       omp.inner.for.body:
+// CHECK-32-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK-32-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
+// CHECK-32-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK-32-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK-32-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK-32-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i32 0, i32 [[TMP8]]
+// CHECK-32-NEXT:    store i32 1, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK-32-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK-32:       omp.body.continue:
+// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK-32:       omp.inner.for.inc:
+// CHECK-32-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK-32-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP9]], 1
+// CHECK-32-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP13:![0-9]+]]
+// CHECK-32:       worker.exit:
+// CHECK-32-NEXT:    ret void
+// CHECK-32:       omp.inner.for.end:
+// CHECK-32-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-32-NEXT:    [[SUB7:%.*]] = sub nsw i32 [[TMP10]], 0
+// CHECK-32-NEXT:    [[DIV8:%.*]] = sdiv i32 [[SUB7]], 1
+// CHECK-32-NEXT:    [[MUL9:%.*]] = mul nsw i32 [[DIV8]], 1
+// CHECK-32-NEXT:    [[ADD10:%.*]] = add nsw i32 0, [[MUL9]]
+// CHECK-32-NEXT:    store i32 [[ADD10]], ptr [[I3]], align 4
+// CHECK-32-NEXT:    br label [[SIMD_IF_END]]
+// CHECK-32:       simd.if.end:
+// CHECK-32-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-32-NEXT:    ret void
+//
+//
+// CHECK-32-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l34
+// CHECK-32-SAME: (i32 noundef [[N:%.*]], ptr noundef nonnull align 2 dereferenceable(2000) [[AA:%.*]]) #[[ATTR0]] {
+// CHECK-32-NEXT:  entry:
+// CHECK-32-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[AA_ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[I3:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4
+// CHECK-32-NEXT:    store ptr [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK-32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 4
+// CHECK-32-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK-32:       user_code.entry:
+// CHECK-32-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4
+// CHECK-32-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-32-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0
+// CHECK-32-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1
+// CHECK-32-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1
+// CHECK-32-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4
+// CHECK-32-NEXT:    store i32 0, ptr [[I]], align 4
+// CHECK-32-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-32-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]
+// CHECK-32-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]
+// CHECK-32:       simd.if.then:
+// CHECK-32-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK-32:       omp.inner.for.cond:
+// CHECK-32-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16:![0-9]+]]
+// CHECK-32-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK-32-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK-32-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]
+// CHECK-32-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK-32:       omp.inner.for.body:
+// CHECK-32-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK-32-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
+// CHECK-32-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK-32-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK-32-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK-32-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i16], ptr [[TMP0]], i32 0, i32 [[TMP8]]
+// CHECK-32-NEXT:    [[TMP9:%.*]] = load i16, ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP16]]
+// CHECK-32-NEXT:    [[CONV:%.*]] = sext i16 [[TMP9]] to i32
+// CHECK-32-NEXT:    [[ADD6:%.*]] = add nsw i32 [[CONV]], 1
+// CHECK-32-NEXT:    [[CONV7:%.*]] = trunc i32 [[ADD6]] to i16
+// CHECK-32-NEXT:    store i16 [[CONV7]], ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP16]]
+// CHECK-32-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK-32:       omp.body.continue:
+// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK-32:       omp.inner.for.inc:
+// CHECK-32-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK-32-NEXT:    [[ADD8:%.*]] = add nsw i32 [[TMP10]], 1
+// CHECK-32-NEXT:    store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP17:![0-9]+]]
+// CHECK-32:       worker.exit:
+// CHECK-32-NEXT:    ret void
+// CHECK-32:       omp.inner.for.end:
+// CHECK-32-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-32-NEXT:    [[SUB9:%.*]] = sub nsw i32 [[TMP11]], 0
+// CHECK-32-NEXT:    [[DIV10:%.*]] = sdiv i32 [[SUB9]], 1
+// CHECK-32-NEXT:    [[MUL11:%.*]] = mul nsw i32 [[DIV10]], 1
+// CHECK-32-NEXT:    [[ADD12:%.*]] = add nsw i32 0, [[MUL11]]
+// CHECK-32-NEXT:    store i32 [[ADD12]], ptr [[I3]], align 4
+// CHECK-32-NEXT:    br label [[SIMD_IF_END]]
+// CHECK-32:       simd.if.end:
+// CHECK-32-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-32-NEXT:    ret void
+//
+//
+// CHECK-32-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l39
+// CHECK-32-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR0]] {
+// CHECK-32-NEXT:  entry:
+// CHECK-32-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4
+// CHECK-32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4
+// CHECK-32-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK-32:       user_code.entry:
+// CHECK-32-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK-32:       omp.inner.for.cond:
+// CHECK-32-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19:![0-9]+]]
+// CHECK-32-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP2]], 10
+// CHECK-32-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK-32:       omp.inner.for.body:
+// CHECK-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK-32-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1
+// CHECK-32-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK-32-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK-32-NEXT:    [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK-32-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP4]]
+// CHECK-32-NEXT:    [[TMP5:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK-32-NEXT:    [[ADD1:%.*]] = add nsw i32 [[TMP5]], 1
+// CHECK-32-NEXT:    store i32 [[ADD1]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK-32-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK-32:       omp.body.continue:
+// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK-32:       omp.inner.for.inc:
+// CHECK-32-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK-32-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK-32-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP20:![0-9]+]]
+// CHECK-32:       worker.exit:
+// CHECK-32-NEXT:    ret void
+// CHECK-32:       omp.inner.for.end:
+// CHECK-32-NEXT:    store i32 10, ptr [[I]], align 4
+// CHECK-32-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-32-NEXT:    ret void
+//
+//
+// CHECK-32-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l44
+// CHECK-32-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]]) #[[ATTR0]] {
+// CHECK-32-NEXT:  entry:
+// CHECK-32-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-NEXT:    [[N_ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    [[N1:%.*]] = alloca i32, align 4
+// CHECK-32-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4
+// CHECK-32-NEXT:    store ptr [[N]], ptr [[N_ADDR]], align 4
+// CHECK-32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4
+// CHECK-32-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[N_ADDR]], align 4
+// CHECK-32-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -1
+// CHECK-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK-32:       user_code.entry:
+// CHECK-32-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK-32-NEXT:    store i32 0, ptr [[N1]], align 4
+// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK-32:       omp.inner.for.cond:
+// CHECK-32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22:![0-9]+]]
+// CHECK-32-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP3]], 10
+// CHECK-32-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK-32:       omp.inner.for.body:
+// CHECK-32-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK-32-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP4]], 1
+// CHECK-32-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK-32-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK-32-NEXT:    [[TMP5:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK-32-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP5]]
+// CHECK-32-NEXT:    [[TMP6:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK-32-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK-32-NEXT:    store i32 [[ADD2]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK-32-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK-32:       omp.body.continue:
+// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK-32:       omp.inner.for.inc:
+// CHECK-32-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK-32-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP7]], 1
+// CHECK-32-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK-32-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP23:![0-9]+]]
+// CHECK-32:       worker.exit:
+// CHECK-32-NEXT:    ret void
+// CHECK-32:       omp.inner.for.end:
+// CHECK-32-NEXT:    store i32 10, ptr [[I]], align 4
+// CHECK-32-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP1]], align 4
+// CHECK-32-NEXT:    [[TMP9:%.*]] = load i32, ptr [[N1]], align 4
+// CHECK-32-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]
+// CHECK-32-NEXT:    store i32 [[ADD4]], ptr [[TMP1]], align 4
+// CHECK-32-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-32-NEXT:    ret void
+//
+//
+// CHECK-32-EX-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l29
+// CHECK-32-EX-SAME: (i32 noundef [[N:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[A:%.*]]) #[[ATTR0:[0-9]+]] {
+// CHECK-32-EX-NEXT:  entry:
+// CHECK-32-EX-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-EX-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[I3:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4
+// CHECK-32-EX-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 4
+// CHECK-32-EX-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 4
+// CHECK-32-EX-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1:[0-9]+]], i8 2, i1 false)
+// CHECK-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK-32-EX:       user_code.entry:
+// CHECK-32-EX-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4
+// CHECK-32-EX-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-32-EX-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0
+// CHECK-32-EX-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1
+// CHECK-32-EX-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1
+// CHECK-32-EX-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4
+// CHECK-32-EX-NEXT:    store i32 0, ptr [[I]], align 4
+// CHECK-32-EX-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-32-EX-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]
+// CHECK-32-EX-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]
+// CHECK-32-EX:       simd.if.then:
+// CHECK-32-EX-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK-32-EX:       omp.inner.for.cond:
+// CHECK-32-EX-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12:![0-9]+]]
+// CHECK-32-EX-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK-32-EX-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK-32-EX-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]
+// CHECK-32-EX-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK-32-EX:       omp.inner.for.body:
+// CHECK-32-EX-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK-32-EX-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
+// CHECK-32-EX-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK-32-EX-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK-32-EX-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK-32-EX-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i32 0, i32 [[TMP8]]
+// CHECK-32-EX-NEXT:    store i32 1, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK-32-EX-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK-32-EX:       omp.body.continue:
+// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK-32-EX:       omp.inner.for.inc:
+// CHECK-32-EX-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK-32-EX-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP9]], 1
+// CHECK-32-EX-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]
+// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP13:![0-9]+]]
+// CHECK-32-EX:       worker.exit:
+// CHECK-32-EX-NEXT:    ret void
+// CHECK-32-EX:       omp.inner.for.end:
+// CHECK-32-EX-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-32-EX-NEXT:    [[SUB7:%.*]] = sub nsw i32 [[TMP10]], 0
+// CHECK-32-EX-NEXT:    [[DIV8:%.*]] = sdiv i32 [[SUB7]], 1
+// CHECK-32-EX-NEXT:    [[MUL9:%.*]] = mul nsw i32 [[DIV8]], 1
+// CHECK-32-EX-NEXT:    [[ADD10:%.*]] = add nsw i32 0, [[MUL9]]
+// CHECK-32-EX-NEXT:    store i32 [[ADD10]], ptr [[I3]], align 4
+// CHECK-32-EX-NEXT:    br label [[SIMD_IF_END]]
+// CHECK-32-EX:       simd.if.end:
+// CHECK-32-EX-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-32-EX-NEXT:    ret void
+//
+//
+// CHECK-32-EX-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l34
+// CHECK-32-EX-SAME: (i32 noundef [[N:%.*]], ptr noundef nonnull align 2 dereferenceable(2000) [[AA:%.*]]) #[[ATTR0]] {
+// CHECK-32-EX-NEXT:  entry:
+// CHECK-32-EX-NEXT:    [[N_ADDR:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[AA_ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-EX-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[I3:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    store i32 [[N]], ptr [[N_ADDR]], align 4
+// CHECK-32-EX-NEXT:    store ptr [[AA]], ptr [[AA_ADDR]], align 4
+// CHECK-32-EX-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 4
+// CHECK-32-EX-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK-32-EX:       user_code.entry:
+// CHECK-32-EX-NEXT:    [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4
+// CHECK-32-EX-NEXT:    store i32 [[TMP2]], ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-32-EX-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP3]], 0
+// CHECK-32-EX-NEXT:    [[DIV:%.*]] = sdiv i32 [[SUB]], 1
+// CHECK-32-EX-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1
+// CHECK-32-EX-NEXT:    store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4
+// CHECK-32-EX-NEXT:    store i32 0, ptr [[I]], align 4
+// CHECK-32-EX-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-32-EX-NEXT:    [[CMP:%.*]] = icmp slt i32 0, [[TMP4]]
+// CHECK-32-EX-NEXT:    br i1 [[CMP]], label [[SIMD_IF_THEN:%.*]], label [[SIMD_IF_END:%.*]]
+// CHECK-32-EX:       simd.if.then:
+// CHECK-32-EX-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK-32-EX:       omp.inner.for.cond:
+// CHECK-32-EX-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16:![0-9]+]]
+// CHECK-32-EX-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK-32-EX-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK-32-EX-NEXT:    [[CMP4:%.*]] = icmp slt i32 [[TMP5]], [[ADD]]
+// CHECK-32-EX-NEXT:    br i1 [[CMP4]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK-32-EX:       omp.inner.for.body:
+// CHECK-32-EX-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK-32-EX-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
+// CHECK-32-EX-NEXT:    [[ADD5:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK-32-EX-NEXT:    store i32 [[ADD5]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK-32-EX-NEXT:    [[TMP8:%.*]] = load i32, ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK-32-EX-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i16], ptr [[TMP0]], i32 0, i32 [[TMP8]]
+// CHECK-32-EX-NEXT:    [[TMP9:%.*]] = load i16, ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP16]]
+// CHECK-32-EX-NEXT:    [[CONV:%.*]] = sext i16 [[TMP9]] to i32
+// CHECK-32-EX-NEXT:    [[ADD6:%.*]] = add nsw i32 [[CONV]], 1
+// CHECK-32-EX-NEXT:    [[CONV7:%.*]] = trunc i32 [[ADD6]] to i16
+// CHECK-32-EX-NEXT:    store i16 [[CONV7]], ptr [[ARRAYIDX]], align 2, !llvm.access.group [[ACC_GRP16]]
+// CHECK-32-EX-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK-32-EX:       omp.body.continue:
+// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK-32-EX:       omp.inner.for.inc:
+// CHECK-32-EX-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK-32-EX-NEXT:    [[ADD8:%.*]] = add nsw i32 [[TMP10]], 1
+// CHECK-32-EX-NEXT:    store i32 [[ADD8]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]
+// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP17:![0-9]+]]
+// CHECK-32-EX:       worker.exit:
+// CHECK-32-EX-NEXT:    ret void
+// CHECK-32-EX:       omp.inner.for.end:
+// CHECK-32-EX-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4
+// CHECK-32-EX-NEXT:    [[SUB9:%.*]] = sub nsw i32 [[TMP11]], 0
+// CHECK-32-EX-NEXT:    [[DIV10:%.*]] = sdiv i32 [[SUB9]], 1
+// CHECK-32-EX-NEXT:    [[MUL11:%.*]] = mul nsw i32 [[DIV10]], 1
+// CHECK-32-EX-NEXT:    [[ADD12:%.*]] = add nsw i32 0, [[MUL11]]
+// CHECK-32-EX-NEXT:    store i32 [[ADD12]], ptr [[I3]], align 4
+// CHECK-32-EX-NEXT:    br label [[SIMD_IF_END]]
+// CHECK-32-EX:       simd.if.end:
+// CHECK-32-EX-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-32-EX-NEXT:    ret void
+//
+//
+// CHECK-32-EX-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l39
+// CHECK-32-EX-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR0]] {
+// CHECK-32-EX-NEXT:  entry:
+// CHECK-32-EX-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-EX-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4
+// CHECK-32-EX-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4
+// CHECK-32-EX-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK-32-EX:       user_code.entry:
+// CHECK-32-EX-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK-32-EX:       omp.inner.for.cond:
+// CHECK-32-EX-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19:![0-9]+]]
+// CHECK-32-EX-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP2]], 10
+// CHECK-32-EX-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK-32-EX:       omp.inner.for.body:
+// CHECK-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK-32-EX-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1
+// CHECK-32-EX-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK-32-EX-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK-32-EX-NEXT:    [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK-32-EX-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP4]]
+// CHECK-32-EX-NEXT:    [[TMP5:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK-32-EX-NEXT:    [[ADD1:%.*]] = add nsw i32 [[TMP5]], 1
+// CHECK-32-EX-NEXT:    store i32 [[ADD1]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK-32-EX-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK-32-EX:       omp.body.continue:
+// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK-32-EX:       omp.inner.for.inc:
+// CHECK-32-EX-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK-32-EX-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK-32-EX-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP19]]
+// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP20:![0-9]+]]
+// CHECK-32-EX:       worker.exit:
+// CHECK-32-EX-NEXT:    ret void
+// CHECK-32-EX:       omp.inner.for.end:
+// CHECK-32-EX-NEXT:    store i32 10, ptr [[I]], align 4
+// CHECK-32-EX-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-32-EX-NEXT:    ret void
+//
+//
+// CHECK-32-EX-LABEL: define {{[^@]+}}@__omp_offloading_30_700c3a28__Z9ftemplateIiET_i_l44
+// CHECK-32-EX-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]]) #[[ATTR0]] {
+// CHECK-32-EX-NEXT:  entry:
+// CHECK-32-EX-NEXT:    [[B_ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-EX-NEXT:    [[N_ADDR:%.*]] = alloca ptr, align 4
+// CHECK-32-EX-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[I:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    [[N1:%.*]] = alloca i32, align 4
+// CHECK-32-EX-NEXT:    store ptr [[B]], ptr [[B_ADDR]], align 4
+// CHECK-32-EX-NEXT:    store ptr [[N]], ptr [[N_ADDR]], align 4
+// CHECK-32-EX-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4
+// CHECK-32-EX-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[N_ADDR]], align 4
+// CHECK-32-EX-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1]], i8 2, i1 false)
+// CHECK-32-EX-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -1
+// CHECK-32-EX-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK-32-EX:       user_code.entry:
+// CHECK-32-EX-NEXT:    store i32 0, ptr [[DOTOMP_IV]], align 4
+// CHECK-32-EX-NEXT:    store i32 0, ptr [[N1]], align 4
+// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// CHECK-32-EX:       omp.inner.for.cond:
+// CHECK-32-EX-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22:![0-9]+]]
+// CHECK-32-EX-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP3]], 10
+// CHECK-32-EX-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// CHECK-32-EX:       omp.inner.for.body:
+// CHECK-32-EX-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK-32-EX-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP4]], 1
+// CHECK-32-EX-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
+// CHECK-32-EX-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK-32-EX-NEXT:    [[TMP5:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK-32-EX-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 [[TMP5]]
+// CHECK-32-EX-NEXT:    [[TMP6:%.*]] = load i32, ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK-32-EX-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK-32-EX-NEXT:    store i32 [[ADD2]], ptr [[ARRAYIDX]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK-32-EX-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// CHECK-32-EX:       omp.body.continue:
+// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// CHECK-32-EX:       omp.inner.for.inc:
+// CHECK-32-EX-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK-32-EX-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP7]], 1
+// CHECK-32-EX-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22]]
+// CHECK-32-EX-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP23:![0-9]+]]
+// CHECK-32-EX:       worker.exit:
+// CHECK-32-EX-NEXT:    ret void
+// CHECK-32-EX:       omp.inner.for.end:
+// CHECK-32-EX-NEXT:    store i32 10, ptr [[I]], align 4
+// CHECK-32-EX-NEXT:    [[TMP8:%.*]] = load i32, ptr [[TMP1]], align 4
+// CHECK-32-EX-NEXT:    [[TMP9:%.*]] = load i32, ptr [[N1]], align 4
+// CHECK-32-EX-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP8]], [[TMP9]]
+// CHECK-32-EX-NEXT:    store i32 [[ADD4]], ptr [[TMP1]], align 4
+// CHECK-32-EX-NEXT:    call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2)
+// CHECK-32-EX-NEXT:    ret void
+//