// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
// Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --include-generated-funcs --version 5
// And edited to fix some failures.

// This test checks a kernel argument that is an Accessor array

#include "Inputs/sycl.hpp"

using namespace sycl;

template <typename name, typename Func>
__attribute__((sycl_kernel)) void a_kernel(const Func &kernelFunc) {
  kernelFunc();
}

int main() {

  using Accessor =
      accessor<int, 1, access::mode::read_write, access::target::global_buffer>;
  Accessor acc[2];

  a_kernel<class kernel_A>(
      [=]() {
        acc[1].use();
      });
}
// CHECK-LABEL: define dso_local spir_kernel void @_ZTSZ4mainE8kernel_A(
// CHECK-SAME: ptr addrspace(1) noundef align 4 [[_ARG_ACC:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_ACC1:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_ACC2:%.*]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[_ARG_ACC3:%.*]], ptr addrspace(1) noundef align 4 [[_ARG_ACC4:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_ACC6:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_ACC7:%.*]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[_ARG_ACC8:%.*]]) #[[ATTR0:[0-9]+]]
// CHECK-NEXT:  [[ENTRY:.*:]]
// CHECK-NEXT:    [[_ARG_ACC_ADDR:%.*]] = alloca ptr addrspace(1), align 8
// CHECK-NEXT:    [[_ARG_ACC_ADDR5:%.*]] = alloca ptr addrspace(1), align 8
// CHECK-NEXT:    [[__SYCLKERNEL:%.*]] = alloca [[CLASS_ANON:%.*]], align 4
// CHECK-NEXT:    [[AGG_TMP:%.*]] = alloca %"struct.sycl::_V1::range", align 4
// CHECK-NEXT:    [[AGG_TMP10:%.*]] = alloca %"struct.sycl::_V1::range", align 4
// CHECK-NEXT:    [[AGG_TMP11:%.*]] = alloca %"struct.sycl::_V1::id", align 4
// CHECK-NEXT:    [[AGG_TMP14:%.*]] = alloca %"struct.sycl::_V1::range", align 4
// CHECK-NEXT:    [[AGG_TMP15:%.*]] = alloca %"struct.sycl::_V1::range", align 4
// CHECK-NEXT:    [[AGG_TMP16:%.*]] = alloca %"struct.sycl::_V1::id", align 4
// CHECK-NEXT:    [[_ARG_ACC_ADDR_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC_ADDR]] to ptr addrspace(4)
// CHECK-NEXT:    [[_ARG_ACC_ADDR5_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC_ADDR5]] to ptr addrspace(4)
// CHECK-NEXT:    [[__SYCLKERNEL_ASCAST:%.*]] = addrspacecast ptr [[__SYCLKERNEL]] to ptr addrspace(4)
// CHECK-NEXT:    [[AGG_TMP_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP]] to ptr addrspace(4)
// CHECK-NEXT:    [[AGG_TMP10_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP10]] to ptr addrspace(4)
// CHECK-NEXT:    [[AGG_TMP11_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP11]] to ptr addrspace(4)
// CHECK-NEXT:    [[AGG_TMP14_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP14]] to ptr addrspace(4)
// CHECK-NEXT:    [[AGG_TMP15_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP15]] to ptr addrspace(4)
// CHECK-NEXT:    [[AGG_TMP16_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP16]] to ptr addrspace(4)
// CHECK-NEXT:    store ptr addrspace(1) [[_ARG_ACC]], ptr addrspace(4) [[_ARG_ACC_ADDR_ASCAST]], align 8
// CHECK-NEXT:    [[_ARG_ACC1_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC1]] to ptr addrspace(4)
// CHECK-NEXT:    [[_ARG_ACC2_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC2]] to ptr addrspace(4)
// CHECK-NEXT:    [[_ARG_ACC3_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC3]] to ptr addrspace(4)
// CHECK-NEXT:    store ptr addrspace(1) [[_ARG_ACC4]], ptr addrspace(4) [[_ARG_ACC_ADDR5_ASCAST]], align 8
// CHECK-NEXT:    [[_ARG_ACC6_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC6]] to ptr addrspace(4)
// CHECK-NEXT:    [[_ARG_ACC7_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC7]] to ptr addrspace(4)
// CHECK-NEXT:    [[_ARG_ACC8_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC8]] to ptr addrspace(4)
// CHECK-NEXT:    [[ACC:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr addrspace(4) [[__SYCLKERNEL_ASCAST]], i32 0, i32 0
// CHECK-NEXT:    call spir_func void @_ZN4sycl3_V18accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC1Ev(ptr addrspace(4) noundef align 4 dereferenceable_or_null(12) [[ACC]]) #[[ATTR4:[0-9]+]]
// CHECK-NEXT:    [[ARRAYINIT_ELEMENT:%.*]] = getelementptr inbounds %"class.sycl::_V1::accessor", ptr addrspace(4) [[ACC]], i64 1
// CHECK-NEXT:    call spir_func void @_ZN4sycl3_V18accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC1Ev(ptr addrspace(4) noundef align 4 dereferenceable_or_null(12) [[ARRAYINIT_ELEMENT]]) #[[ATTR4]]
// CHECK-NEXT:    [[ACC9:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr addrspace(4) [[__SYCLKERNEL_ASCAST]], i32 0, i32 0
// CHECK-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds nuw [2 x %"class.sycl::_V1::accessor"], ptr addrspace(4) [[ACC9]], i64 0, i64 0
// CHECK-NEXT:    [[TMP0:%.*]] = load ptr addrspace(1), ptr addrspace(4) [[_ARG_ACC_ADDR_ASCAST]], align 8
// CHECK-NEXT:    call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP_ASCAST]], ptr addrspace(4) align 4 [[_ARG_ACC1_ASCAST]], i64 4, i1 false)
// CHECK-NEXT:    call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP10_ASCAST]], ptr addrspace(4) align 4 [[_ARG_ACC2_ASCAST]], i64 4, i1 false)
// CHECK-NEXT:    call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP11_ASCAST]], ptr addrspace(4) align 4 [[_ARG_ACC3_ASCAST]], i64 4, i1 false)
// CHECK-NEXT:    [[AGG_TMP_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP_ASCAST]] to ptr
// CHECK-NEXT:    [[AGG_TMP10_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP10_ASCAST]] to ptr
// CHECK-NEXT:    [[AGG_TMP11_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP11_ASCAST]] to ptr
// CHECK-NEXT:    call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) noundef align 4 dereferenceable_or_null(12) [[ARRAYIDX]], ptr addrspace(1) noundef [[TMP0]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[AGG_TMP_ASCAST_ASCAST]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[AGG_TMP10_ASCAST_ASCAST]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[AGG_TMP11_ASCAST_ASCAST]]) #[[ATTR4]]
// CHECK-NEXT:    [[ACC12:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr addrspace(4) [[__SYCLKERNEL_ASCAST]], i32 0, i32 0
// CHECK-NEXT:    [[ARRAYIDX13:%.*]] = getelementptr inbounds nuw [2 x %"class.sycl::_V1::accessor"], ptr addrspace(4) [[ACC12]], i64 0, i64 1
// CHECK-NEXT:    [[TMP1:%.*]] = load ptr addrspace(1), ptr addrspace(4) [[_ARG_ACC_ADDR5_ASCAST]], align 8
// CHECK-NEXT:    call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP14_ASCAST]], ptr addrspace(4) align 4 [[_ARG_ACC6_ASCAST]], i64 4, i1 false)
// CHECK-NEXT:    call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP15_ASCAST]], ptr addrspace(4) align 4 [[_ARG_ACC7_ASCAST]], i64 4, i1 false)
// CHECK-NEXT:    call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP16_ASCAST]], ptr addrspace(4) align 4 [[_ARG_ACC8_ASCAST]], i64 4, i1 false)
// CHECK-NEXT:    [[AGG_TMP14_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP14_ASCAST]] to ptr
// CHECK-NEXT:    [[AGG_TMP15_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP15_ASCAST]] to ptr
// CHECK-NEXT:    [[AGG_TMP16_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP16_ASCAST]] to ptr
// CHECK-NEXT:    call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) noundef align 4 dereferenceable_or_null(12) [[ARRAYIDX13]], ptr addrspace(1) noundef [[TMP1]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[AGG_TMP14_ASCAST_ASCAST]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[AGG_TMP15_ASCAST_ASCAST]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[AGG_TMP16_ASCAST_ASCAST]]) #[[ATTR4]]
// CHECK-NEXT:    call spir_func void @_ZZ4mainENKUlvE_clEv(ptr addrspace(4) noundef align 4 dereferenceable_or_null(24) [[__SYCLKERNEL_ASCAST]]) #[[ATTR4]]
// CHECK-NEXT:    ret void