[SYCL] Decompose kernel parameters and add inheritance support

clang/test/CodeGenSYCL/accessor_inheritance.cpp

@@ -0,0 +1,94 @@
+// RUN:  %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+#include <sycl.hpp>
+
+struct Base {
+  int A, B;
+  cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> AccField;
+};
+
+struct Captured : Base,
+                  cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> {
+  int C;
+};
+
+int main() {
+  Captured Obj;
+  cl::sycl::kernel_single_task<class kernel>(
+      [=]() {
+        Obj.use();
+      });
+  return 0;
+}
+
+// Check kernel parameters
+// CHECK: %[[RANGE_TYPE:"struct.*cl::sycl::range"]]
+// CHECK: %[[ID_TYPE:"struct.*cl::sycl::id"]]
+// CHECK: define spir_kernel void @_ZTSZ4mainE6kernel
+// CHECK-SAME: i32 [[ARG_A:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: i32 [[ARG_B:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: i8 addrspace(1)* [[ACC1_DATA:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) align 4 [[ACC1_RANGE1:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) align 4 [[ACC1_RANGE2:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: %[[ID_TYPE]]* byval(%[[ID_TYPE]]) align 4 [[ACC1_ID:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: i8 addrspace(1)* [[ACC2_DATA:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) align 4 [[ACC2_RANGE1:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) align 4 [[ACC2_RANGE2:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: %[[ID_TYPE]]* byval(%[[ID_TYPE]]) align 4 [[ACC2_ID:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: i32 [[ARG_C:%[a-zA-Z0-9_]+]])
+
+// Allocas for kernel parameters
+// CHECK: [[ARG_A]].addr = alloca i32
+// CHECK: [[ARG_B]].addr = alloca i32
+// CHECK: [[ACC1_DATA]].addr = alloca i8 addrspace(1)*
+// CHECK: [[ACC2_DATA]].addr = alloca i8 addrspace(1)*
+// CHECK: [[ARG_C]].addr = alloca i32
+//
+// Lambda object alloca
+// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_]+]] = alloca %"class.{{.*}}.anon"
+//
+// Kernel argument stores
+// CHECK: store i32 [[ARG_A]], i32* [[ARG_A]].addr
+// CHECK: store i32 [[ARG_B]], i32* [[ARG_B]].addr
+// CHECK: store i8 addrspace(1)* [[ACC1_DATA]], i8 addrspace(1)** [[ACC1_DATA]].addr
+// CHECK: store i8 addrspace(1)* [[ACC2_DATA]], i8 addrspace(1)** [[ACC2_DATA]].addr
+// CHECK: store i32 [[ARG_C]], i32* [[ARG_C]].addr
+//
+// Check A and B scalar fields initialization
+// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class._ZTSZ4mainE3$_0.anon", %"class._ZTSZ4mainE3$_0.anon"* [[KERNEL_OBJ]], i32 0, i32 0
+// CHECK: [[BITCAST:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured* [[GEP]] to %struct{{.*}}Base*
+// CHECK: [[FIELD_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST]], i32 0, i32 0
+// CHECK: [[ARG_A_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32* [[ARG_A]].addr
+// CHECK: store i32 [[ARG_A_LOAD]], i32* [[FIELD_A]]
+// CHECK: [[FIELD_B:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST]], i32 0, i32 1
+// CHECK: [[ARG_B_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32* [[ARG_B]].addr
+// CHECK: store i32 [[ARG_B_LOAD]], i32* [[FIELD_B]]
+//
+// Check accessors initialization
+// CHECK: [[ACC_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST]], i32 0, i32 2
+// CHECK: [[ACC1_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[ACC_FIELD]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
+// Default constructor call
+// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0EEC1Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC1_AS_CAST]])
+// CHECK: [[BITCAST1:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured* [[GEP]] to i8*
+// CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, i8* [[BITCAST1]], i64 20
+// CHECK: [[BITCAST2:%[a-zA-Z0-9_]+]] = bitcast i8* [[GEP1]] to %"class{{.*}}cl::sycl::accessor"*
+// CHECK: [[ACC2_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[BITCAST2]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
+// Default constructor call
+// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0EEC2Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC2_AS_CAST]])
+
+// CHECK C field initialization
+// CHECK: [[FIELD_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Captured, %struct{{.*}}Captured* [[GEP]], i32 0, i32 2
+// CHECK: [[ARG_C_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32* [[ARG_C]].addr
+// CHECK: store i32 [[ARG_C_LOAD]], i32* [[FIELD_C]]
+//
+// Check __init method calls
+// CHECK: [[GEP2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class._ZTSZ4mainE3$_0.anon", %"class._ZTSZ4mainE3$_0.anon"* [[KERNEL_OBJ]], i32 0, i32 0
+// CHECK: [[BITCAST3:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured* [[GEP2]] to %struct{{.*}}Base*
+// CHECK: [[ACC1_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST3]], i32 0, i32 2
+// CHECK: [[ACC1_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)** [[ACC1_DATA]].addr
+// CHECK: [[ACC1_AS_CAST1:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[ACC1_FIELD]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
+// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC1_AS_CAST1]], i8 addrspace(1)* [[ACC1_DATA_LOAD]]
+//
+// CHECK: [[GEP3:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class._ZTSZ4mainE3$_0.anon", %"class._ZTSZ4mainE3$_0.anon"* [[KERNEL_OBJ]], i32 0, i32 0
+// CHECK: [[ACC2_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)** [[ACC2_DATA]].addr
+// CHECK: [[AS_CAST_CAPTURED:%[a-zA-Z0-9_]+]] = addrspacecast %struct{{.*}}Captured* [[GEP3]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
+// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[AS_CAST_CAPTURED]], i8 addrspace(1)* [[ACC2_DATA_LOAD]]

clang/test/CodeGenSYCL/inheritance.cpp

@@ -0,0 +1,84 @@
+// RUN:  %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+
+#include <sycl.hpp>
+
+class second_base {
+public:
+  int e;
+};
+
+class InnerFieldBase {
+public:
+  int d;
+};
+class InnerField : public InnerFieldBase {
+  int c;
+};
+
+struct base {
+public:
+  int b;
+  InnerField obj;
+};
+
+struct derived : base, second_base {
+  int a;
+
+  void operator()() {
+  }
+};
+
+int main() {
+  cl::sycl::queue q;
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    derived f{};
+    cgh.single_task(f);
+  });
+
+  return 0;
+}
+
+// Check kernel paramters
+// CHECK: define spir_kernel void @{{.*}}derived(i32 %_arg_b, i32 %_arg_d, i32 %_arg_c, i32 %_arg_e, i32 %_arg_a)
+
+// Check alloca for kernel paramters
+// CHECK: %[[ARG_B:[a-zA-Z0-9_.]+]] = alloca i32, align 4
+// CHECK: %[[ARG_D:[a-zA-Z0-9_.]+]] = alloca i32, align 4
+// CHECK: %[[ARG_C:[a-zA-Z0-9_.]+]] = alloca i32, align 4
+// CHECK: %[[ARG_E:[a-zA-Z0-9_.]+]] = alloca i32, align 4
+// CHECK: %[[ARG_A:[a-zA-Z0-9_.]+]] = alloca i32, align 4
+
+// Check alloca for local functor object
+// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = alloca %struct.{{.*}}.derived, align 4
+
+// Initialize field 'b'
+// CHECK: %[[BITCAST1:[0-9]+]] = bitcast %struct.{{.*}}.derived* %[[LOCAL_OBJECT]] to %struct.{{.*}}.base*
+// CHECK: %[[GEP_B:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.base, %struct.{{.*}}.base* %[[BITCAST1]], i32 0, i32 0
+// CHECK: %[[LOAD_B:[0-9]+]] = load i32, i32* %[[ARG_B]], align 4
+// CHECK: store i32 %[[LOAD_B]], i32* %[[GEP_B]], align 4
+
+// Initialize field 'd'
+// CHECK: %[[GEP_OBJ:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.base, %struct.{{.*}}.base* %[[BITCAST1]], i32 0, i32 1
+// CHECK: %[[BITCAST2:[0-9]+]] = bitcast %class.{{.*}}.InnerField* %[[GEP_OBJ]] to %class.{{.*}}.InnerFieldBase*
+// CHECK: %[[GEP_D:[a-zA-Z0-9]+]] = getelementptr inbounds %class.{{.*}}.InnerFieldBase, %class.{{.*}}.InnerFieldBase* %[[BITCAST2]], i32 0, i32 0
+// CHECK: %[[LOAD_D:[0-9]+]] = load i32, i32* %[[ARG_D]], align 4
+// CHECK: store i32 %[[LOAD_D]], i32* %[[GEP_D]], align 4
+
+// Initialize field 'c'
+// CHECK: %[[GEP_C:[a-zA-Z0-9]+]] = getelementptr inbounds %class.{{.*}}.InnerField, %class.{{.*}}.InnerField* %[[GEP_OBJ]], i32 0, i32 1
+// CHECK: %[[LOAD_C:[0-9]+]] = load i32, i32* %[[ARG_C]], align 4
+// CHECK: store i32 %[[LOAD_C]], i32* %[[GEP_C]], align 4
+
+// Initialize field 'e'
+// CHECK: %[[BITCAST3:[0-9]+]] = bitcast %struct.{{.*}}.derived* %[[LOCAL_OBJECT]] to i8*
+// CHECK: %[[GEP_DERIVED:[a-zA-Z0-9]+]] = getelementptr inbounds i8, i8* %[[BITCAST3]], i64 12
+// CHECK: %[[BITCAST4:[0-9]+]] = bitcast i8* %[[GEP_DERIVED]] to %class.{{.*}}.second_base*
+// CHECK: %[[GEP_E:[a-zA-Z0-9]+]] = getelementptr inbounds %class.{{.*}}.second_base, %class.{{.*}}.second_base* %[[BITCAST4]], i32 0, i32 0
+// CHECK: %[[LOAD_E:[0-9]+]] = load i32, i32* %[[ARG_E]], align 4
+// CHECK: store i32 %[[LOAD_E]], i32* %[[GEP_E]], align 4
+
+// Initialize field 'a'
+// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.derived, %struct.{{.*}}.derived* %[[LOCAL_OBJECT]], i32 0, i32 2
+// CHECK: %[[LOAD_A:[0-9]+]] = load i32, i32* %[[ARG_A]], align 4
+// CHECK: store i32 %[[LOAD_A]], i32* %[[GEP_A]], align 4

clang/test/CodeGenSYCL/integration_header.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -emit-llvm
 // RUN: FileCheck -input-file=%t.h %s
 //
 // CHECK: #include <CL/sycl/detail/kernel_desc.hpp>
@@ -28,9 +28,11 @@
 // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
 // CHECK-NEXT:   //--- _ZTSZ4mainE12first_kernel
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 0 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 4 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 6112, 16 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_sampler, 8, 32 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 1, 4 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 8 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 12 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 6112, 24 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_sampler, 8, 40 },
 // CHECK-EMPTY:
 // CHECK-NEXT:   //--- _ZTSN16second_namespace13second_kernelIcEE
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 0 },
@@ -46,12 +48,15 @@
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 0 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 6112, 4 },
 // CHECK-EMPTY:
-// CHECK-NEXT:   //--- _ZTSZ4mainE16accessor_in_base
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 64, 0 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 8 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 24 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 40 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 52 },
+// CHECK-NEXT:  //--- _ZTSZ4mainE16accessor_in_base
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 0 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 4 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_accessor, 4062, 8 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 20 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_accessor, 4062, 24 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 36 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_accessor, 4062, 40 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_accessor, 4062, 52 },
 // CHECK-EMPTY:
 // CHECK-NEXT: };
 //
@@ -116,15 +121,13 @@ int main() {
       acc2;
   int i = 13;
   cl::sycl::sampler smplr;
-  // TODO: Uncomemnt when structures in kernel arguments are correctly processed
-  //       by SYCL compiler
-  /*  struct {
+  struct {
     char c;
     int i;
   } test_s;
-  test_s.c = 14;*/
+  test_s.c = 14;
   kernel_single_task<class first_kernel>([=]() {
-    if (i == 13 /*&& test_s.c == 14*/) {
+    if (i == 13 && test_s.c == 14) {
 
       acc1.use();
       acc2.use();
@@ -151,10 +154,9 @@ int main() {
       }
   });
 
-  // FIXME: We cannot use the member-capture because all the handlers except the
-  // integration header handler in SemaSYCL don't handle base types right.
   accessor_in_base::captured c;
-  kernel_single_task<class accessor_in_base>([c]() {
+  kernel_single_task<class accessor_in_base>([=]() {
+    c.use();
   });
 
   return 0;

clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp

@@ -1,7 +1,5 @@
 // RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only
 // RUN: FileCheck -input-file=%t.h %s
-// XFAIL for now due to : https://github.com/intel/llvm/issues/2018
-// XFAIL: *
 
 // This test checks the integration header when kernel argument
 // is a struct containing an Accessor array.
@@ -22,7 +20,6 @@
 // CHECK: static constexpr
 // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
 // CHECK-NEXT: //--- _ZTSZ4mainE8kernel_C
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 24, 0 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 0 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 12 },
 // CHECK-EMPTY:

clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp

@@ -1,5 +1,4 @@
 // RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -fsycl-int-header=%t.h -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
-// XFAIL: *
 
 // This test checks a kernel with struct parameter that contains an Accessor array.
 
@@ -29,7 +28,6 @@ int main() {
 
 // CHECK kernel_C parameters
 // CHECK: define spir_kernel void @{{.*}}kernel_C
-// CHECK-SAME: %struct.{{.*}}.struct_acc_t* byval(%struct.{{.*}}.struct_acc_t) align 4 [[STRUCT:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: i32 addrspace(1)* [[MEM_ARG1:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
 // CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
@@ -54,32 +52,26 @@ int main() {
 // CHECK: [[MEM_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
 // CHECK: [[OFFSET2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id"
 
-// Check init of local struct
-// CHECK: [[L_STRUCT_ADDR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
-// CHECK: [[MEMCPY_DST:%[0-9a-zA-Z_]+]] = bitcast %struct.{{.*}}struct_acc_t* [[L_STRUCT_ADDR]] to i8*
-// CHECK: [[MEMCPY_SRC:%[0-9a-zA-Z_]+]] = bitcast %struct.{{.*}}struct_acc_t* %{{[0-9a-zA-Z_]+}} to i8*
-// CHECK: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 [[MEMCPY_DST]], i8* align 4 [[MEMCPY_SRC]], i64 24, i1 false)
-
-// Check accessor array GEP for member_acc[0]
-// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
-// CHECK: [[MEMBER1:%[a-zA-Z_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[ACCESSOR_ARRAY1]], i32 0, i32 0
-// CHECK: [[Z0:%[a-zA-Z0-9_]*]] = getelementptr inbounds [2 x %"class.{{.*}}.cl::sycl::accessor"], [2 x %"class.{{.*}}.cl::sycl::accessor"]* [[MEMBER1]], i64 0, i64 0
-
-// Check load from kernel pointer argument alloca
-// CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr{{[0-9]*}}
+// Check loop which calls the default constructor for each element of accessor array is emitted.
+// CHECK: [[GEP_LAMBDA:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: [[GEP_MEMBER_ACC:%[a-zA-Z_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[GEP_LAMBDA]], i32 0, i32 0
+// CHECK: [[ARRAY_BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], [2 x [[ACCESSOR]]]* [[GEP_MEMBER_ACC]], i64 0, i64 0
+// CHECK: [[ARRAY_END:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* [[ARRAY_BEGIN]], i64 2
+// CHECK: br label %arrayctor.loop
+// CHECK: arrayctor.loop:
 
 // Check acc[0] __init method call
-// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[Z0]] to %"class{{.*}}accessor" addrspace(4)*
-// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACC_CAST1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]])
-
-// Check accessor array GEP for member_acc[1]
-// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
-// CHECK: [[MEMBER2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[ACCESSOR_ARRAY2]], i32 0, i32 0
-// CHECK: [[Z1:%[a-zA-Z0-9_]*]] = getelementptr inbounds [2 x %"class.{{.*}}.cl::sycl::accessor"], [2 x %"class.{{.*}}.cl::sycl::accessor"]* [[MEMBER2]], i64 0, i64 1
-
-// Check load from kernel pointer argument alloca
-// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr{{[0-9]*}}
+// CHECK: [[GEP_LAMBDA1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: [[GEP_MEMBER_ACC1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[GEP_LAMBDA1]], i32 0, i32 0
+// CHECK: [[ARRAY_IDX1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[GEP_MEMBER_ACC1]], i64 0, i64 0
+// CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr
+// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast [[ACCESSOR]]* [[ARRAY_IDX1]] to [[ACCESSOR]] addrspace(4)*
+// CHECK: call spir_func void @{{.*}}__init{{.*}}([[ACCESSOR]] addrspace(4)* [[ACC_CAST1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]])
 
 // Check acc[1] __init method call
-// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[Z1]] to %"class{{.*}}accessor" addrspace(4)*
-// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACC_CAST2]], i32 addrspace(1)* [[MEM_LOAD2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2]])
+// CHECK: [[GEP_LAMBDA2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: [[GEP_MEMBER_ACC2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[GEP_LAMBDA2]], i32 0, i32 0
+// CHECK: [[ARRAY_IDX2:%[a-zA-Z0-9_]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[GEP_MEMBER_ACC2]], i64 0, i64 1
+// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr
+// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast [[ACCESSOR]]* [[ARRAY_IDX2]] to [[ACCESSOR]] addrspace(4)*
+// CHECK: call spir_func void @{{.*}}__init{{.*}}([[ACCESSOR]] addrspace(4)* [[ACC_CAST2]], i32 addrspace(1)* [[MEM_LOAD2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2]])

clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp

@@ -1,6 +1,5 @@
 // RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only
 // RUN: FileCheck -input-file=%t.h %s
-
 // This test checks the integration header generated for a kernel
 // with an argument that is a POD array.