// RUN: %clang_cc1 -fsycl-is-device -fsycl-int-header=%t.h -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 later to fix failures.

// This test checks a kernel with struct parameter that contains 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>;

  struct struct_acc_t {
    Accessor member_acc[2];
  } struct_acc;

  a_kernel<class kernel_C>(
      [=]() {
        struct_acc.member_acc[1].use();
      });
}

// CHECK-LABEL: define dso_local spir_kernel void @_ZTSZ4mainE8kernel_C(
// CHECK-SAME: ptr addrspace(1) noundef align 4 [[_ARG_MEMBER_ACC:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_MEMBER_ACC1:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_MEMBER_ACC2:%.*]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[_ARG_MEMBER_ACC3:%.*]], ptr addrspace(1) noundef align 4 [[_ARG_MEMBER_ACC4:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_MEMBER_ACC6:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_MEMBER_ACC7:%.*]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[_ARG_MEMBER_ACC8:%.*]]) #[[ATTR0:[0-9]+]]
// CHECK-NEXT:  [[ENTRY:.*:]]
// CHECK-NEXT:    [[_ARG_MEMBER_ACC_ADDR:%.*]] = alloca ptr addrspace(1), align 8
// CHECK-NEXT:    [[_ARG_MEMBER_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_TMP11:%.*]] = alloca %"struct.sycl::_V1::range", align 4
// CHECK-NEXT:    [[AGG_TMP12:%.*]] = alloca %"struct.sycl::_V1::id", align 4
// CHECK-NEXT:    [[AGG_TMP16:%.*]] = alloca %"struct.sycl::_V1::range", align 4
// CHECK-NEXT:    [[AGG_TMP17:%.*]] = alloca %"struct.sycl::_V1::range", align 4
// CHECK-NEXT:    [[AGG_TMP18:%.*]] = alloca %"struct.sycl::_V1::id", align 4
// CHECK-NEXT:    [[_ARG_MEMBER_ACC_ADDR_ASCAST:%.*]] = addrspacecast ptr [[_ARG_MEMBER_ACC_ADDR]] to ptr addrspace(4)
// CHECK-NEXT:    [[_ARG_MEMBER_ACC_ADDR5_ASCAST:%.*]] = addrspacecast ptr [[_ARG_MEMBER_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_TMP11_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP11]] to ptr addrspace(4)
// CHECK-NEXT:    [[AGG_TMP12_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP12]] to ptr addrspace(4)
// CHECK-NEXT:    [[AGG_TMP16_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP16]] to ptr addrspace(4)
// CHECK-NEXT:    [[AGG_TMP17_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP17]] to ptr addrspace(4)
// CHECK-NEXT:    [[AGG_TMP18_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP18]] to ptr addrspace(4)
// CHECK-NEXT:    store ptr addrspace(1) [[_ARG_MEMBER_ACC]], ptr addrspace(4) [[_ARG_MEMBER_ACC_ADDR_ASCAST]], align 8
// CHECK-NEXT:    [[_ARG_MEMBER_ACC1_ASCAST:%.*]] = addrspacecast ptr [[_ARG_MEMBER_ACC1]] to ptr addrspace(4)
// CHECK-NEXT:    [[_ARG_MEMBER_ACC2_ASCAST:%.*]] = addrspacecast ptr [[_ARG_MEMBER_ACC2]] to ptr addrspace(4)
// CHECK-NEXT:    [[_ARG_MEMBER_ACC3_ASCAST:%.*]] = addrspacecast ptr [[_ARG_MEMBER_ACC3]] to ptr addrspace(4)
// CHECK-NEXT:    store ptr addrspace(1) [[_ARG_MEMBER_ACC4]], ptr addrspace(4) [[_ARG_MEMBER_ACC_ADDR5_ASCAST]], align 8
// CHECK-NEXT:    [[_ARG_MEMBER_ACC6_ASCAST:%.*]] = addrspacecast ptr [[_ARG_MEMBER_ACC6]] to ptr addrspace(4)
// CHECK-NEXT:    [[_ARG_MEMBER_ACC7_ASCAST:%.*]] = addrspacecast ptr [[_ARG_MEMBER_ACC7]] to ptr addrspace(4)
// CHECK-NEXT:    [[_ARG_MEMBER_ACC8_ASCAST:%.*]] = addrspacecast ptr [[_ARG_MEMBER_ACC8]] to ptr addrspace(4)
// CHECK-NEXT:    [[STRUCT_ACC:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr addrspace(4) [[__SYCLKERNEL_ASCAST]], i32 0, i32 0
// CHECK-NEXT:    [[MEMBER_ACC:%.*]] = getelementptr inbounds nuw [[STRUCT_STRUCT_ACC_T:%.*]], ptr addrspace(4) [[STRUCT_ACC]], 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) [[MEMBER_ACC]]) #[[ATTR4:[0-9]+]]
// CHECK-NEXT:    [[ARRAYINIT_ELEMENT:%.*]] = getelementptr inbounds %"class.sycl::_V1::accessor", ptr addrspace(4) [[MEMBER_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:    [[STRUCT_ACC9:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr addrspace(4) [[__SYCLKERNEL_ASCAST]], i32 0, i32 0
// CHECK-NEXT:    [[MEMBER_ACC10:%.*]] = getelementptr inbounds nuw [[STRUCT_STRUCT_ACC_T]], ptr addrspace(4) [[STRUCT_ACC9]], i32 0, i32 0
// CHECK-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds nuw [2 x %"class.sycl::_V1::accessor"], ptr addrspace(4) [[MEMBER_ACC10]], i64 0, i64 0
// CHECK-NEXT:    [[TMP0:%.*]] = load ptr addrspace(1), ptr addrspace(4) [[_ARG_MEMBER_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_MEMBER_ACC1_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_MEMBER_ACC2_ASCAST]], i64 4, i1 false)
// CHECK-NEXT:    call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP12_ASCAST]], ptr addrspace(4) align 4 [[_ARG_MEMBER_ACC3_ASCAST]], i64 4, i1 false)
// CHECK-NEXT:    [[AGG_TMP_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP_ASCAST]] to ptr
// CHECK-NEXT:    [[AGG_TMP11_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP11_ASCAST]] to ptr
// CHECK-NEXT:    [[AGG_TMP12_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP12_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_TMP11_ASCAST_ASCAST]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[AGG_TMP12_ASCAST_ASCAST]]) #[[ATTR4]]
// CHECK-NEXT:    [[STRUCT_ACC13:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr addrspace(4) [[__SYCLKERNEL_ASCAST]], i32 0, i32 0
// CHECK-NEXT:    [[MEMBER_ACC14:%.*]] = getelementptr inbounds nuw [[STRUCT_STRUCT_ACC_T]], ptr addrspace(4) [[STRUCT_ACC13]], i32 0, i32 0
// CHECK-NEXT:    [[ARRAYIDX15:%.*]] = getelementptr inbounds nuw [2 x %"class.sycl::_V1::accessor"], ptr addrspace(4) [[MEMBER_ACC14]], i64 0, i64 1
// CHECK-NEXT:    [[TMP1:%.*]] = load ptr addrspace(1), ptr addrspace(4) [[_ARG_MEMBER_ACC_ADDR5_ASCAST]], align 8
// CHECK-NEXT:    call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP16_ASCAST]], ptr addrspace(4) align 4 [[_ARG_MEMBER_ACC6_ASCAST]], i64 4, i1 false)
// CHECK-NEXT:    call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP17_ASCAST]], ptr addrspace(4) align 4 [[_ARG_MEMBER_ACC7_ASCAST]], i64 4, i1 false)
// CHECK-NEXT:    call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP18_ASCAST]], ptr addrspace(4) align 4 [[_ARG_MEMBER_ACC8_ASCAST]], i64 4, i1 false)
// CHECK-NEXT:    [[AGG_TMP16_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP16_ASCAST]] to ptr
// CHECK-NEXT:    [[AGG_TMP17_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP17_ASCAST]] to ptr
// CHECK-NEXT:    [[AGG_TMP18_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP18_ASCAST]] to ptr
// CHECK-NEXT:    call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) noundef align 4 dereferenceable_or_null(12) [[ARRAYIDX15]], ptr addrspace(1) noundef [[TMP1]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[AGG_TMP16_ASCAST_ASCAST]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[AGG_TMP17_ASCAST_ASCAST]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[AGG_TMP18_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