intel · bader · Jul 2, 2020 · Jun 9, 2020 · Jun 9, 2020 · Jun 9, 2020
@@ -0,0 +1,52 @@
+// 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 when
+// the kernel argument is an Accessor array.
+
+// CHECK: #include <CL/sycl/detail/kernel_desc.hpp>
+
+// CHECK: class kernel_A;
+
+// CHECK: __SYCL_INLINE_NAMESPACE(cl) {
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: namespace detail {
+
+// CHECK: static constexpr
+// CHECK-NEXT: const char* const kernel_names[] = {
+// CHECK-NEXT:   "_ZTSZ4mainE8kernel_A"
+// CHECK-NEXT: };
+
+// CHECK: static constexpr
+// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
+// CHECK-NEXT:   //--- _ZTSZ4mainE8kernel_A
+// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 0 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 12 },
+// CHECK-EMPTY:
+// CHECK-NEXT: };
+
+// CHECK: static constexpr
+// CHECK-NEXT: const unsigned kernel_signature_start[] = {
+// CHECK-NEXT:  0 // _ZTSZ4mainE8kernel_A
+// CHECK-NEXT: };
+
+// CHECK: template <> struct KernelInfo<class kernel_A> {
+
+#include <sycl.hpp>
+
+using namespace cl::sycl;
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void a_kernel(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(); });
+}
@@ -0,0 +1,78 @@
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+
+// This test checks a kernel argument that is an Accessor array
+
+#include <sycl.hpp>
+
+using namespace cl::sycl;
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void a_kernel(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 kernel_A parameters
+// CHECK: define spir_kernel void @{{.*}}kernel_A
+// 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]],
+// CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+_3]],
+// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG2:%[a-zA-Z0-9_]+_4]],
+// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+_6]],
+// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+_7]],
+// CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET2:%[a-zA-Z0-9_]+_8]])
+
+// CHECK alloca for pointer arguments
+// CHECK: [[MEM_ARG1:%[a-zA-Z0-9_.]+]] = alloca i32 addrspace(1)*, align 8
+// CHECK: [[MEM_ARG2:%[a-zA-Z0-9_.]+]] = alloca i32 addrspace(1)*, align 8
+
+// CHECK lambda object alloca
+// CHECK: [[LOCAL_OBJECT:%0]] = alloca %"class.{{.*}}.anon", align 4
+
+// CHECK allocas for ranges
+// CHECK: [[ACC_RANGE1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
+// CHECK: [[MEM_RANGE1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
+// CHECK: [[OFFSET1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id"
+// CHECK: [[ACC_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
+// CHECK: [[MEM_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
+// CHECK: [[OFFSET2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id"
+
+// CHECK accessor array default inits
+// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: [[BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], [2 x [[ACCESSOR]]]* [[ACCESSOR_ARRAY1]], i64 0, i64 0
+// CHECK: [[END:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR:.*]], [[ACCESSOR]]* [[BEGIN]], i64 2
+// CHECK: [[NEXT0:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* {{.*}}, i64 1
+// CHECK: [[ELEMENT:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* {{.*}}, i64 1
+// CHECK: [[ELEMENT:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* {{.*}}, i64 2
+// CHECK: [[NEXT1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* {{.*}}, i64 1
+
+// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: [[INDEX:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[ACCESSOR_ARRAY2]], 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]]
+
+// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast [[ACCESSOR]]* {{.*}} to [[ACCESSOR]] addrspace(4)*
+
+// CHECK acc[0] __init method call
+// 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 load from kernel pointer argument alloca
+// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG2]]
+
+// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast [[ACCESSOR]]* {{.*}} to [[ACCESSOR]] addrspace(4)*
+
+// CHECK acc[1] __init method call
+// 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]])
@@ -0,0 +1,60 @@
+// 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.
+
+// CHECK: #include <CL/sycl/detail/kernel_desc.hpp>
+
+// CHECK: class kernel_C;
+
+// CHECK: __SYCL_INLINE_NAMESPACE(cl) {
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: namespace detail {
+
+// CHECK: static constexpr
+// CHECK-NEXT: const char* const kernel_names[] = {
+// CHECK-NEXT:   "_ZTSZ4mainE8kernel_C"
+// CHECK-NEXT: };
+
+// 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:
+// CHECK-NEXT: };
+
+// CHECK: static constexpr
+// CHECK-NEXT: const unsigned kernel_signature_start[] = {
+// CHECK-NEXT:  0 // _ZTSZ4mainE8kernel_C
+// CHECK-NEXT: };
+
+// CHECK: template <> struct KernelInfo<class kernel_C> {
+
+#include <sycl.hpp>
+
+using namespace cl::sycl;
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void a_kernel(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();
+      });
+}
@@ -0,0 +1,85 @@
+// 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.
+
+#include <sycl.hpp>
+
+using namespace cl::sycl;
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void a_kernel(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 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]],
+// CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]],
+// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG2:%[a-zA-Z0-9_]+4]],
+// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+6]],
+// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+7]],
+// CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2:%[a-zA-Z0-9_]+8]])
+
+// Check alloca for pointer arguments
+// CHECK: [[MEM_ARG1]].addr{{[0-9]*}} = alloca i32 addrspace(1)*, align 8
+// CHECK: [[MEM_ARG1]].addr{{[0-9]*}} = alloca i32 addrspace(1)*, align 8
+
+// Check lambda object alloca
+// CHECK: [[LOCAL_OBJECT:%0]] = alloca %"class.{{.*}}.anon", align 4
+
+// Check allocas for ranges
+// CHECK: [[ACC_RANGE1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
+// CHECK: [[MEM_RANGE1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
+// CHECK: [[OFFSET1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id"
+// CHECK: [[ACC_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
+// 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 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 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]])
@@ -0,0 +1,55 @@
+// 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.
+
+// CHECK: #include <CL/sycl/detail/kernel_desc.hpp>
+
+// CHECK: class kernel_B;
+
+// CHECK: __SYCL_INLINE_NAMESPACE(cl) {
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: namespace detail {
+
+// CHECK: static constexpr
+// CHECK-NEXT: const char* const kernel_names[] = {
+// CHECK-NEXT:   "_ZTSZ4mainE8kernel_B"
+// CHECK-NEXT: };
+
+// CHECK: static constexpr
+// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
+// CHECK-NEXT: //--- _ZTSZ4mainE8kernel_B
+// 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_std_layout, 4, 8 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 12 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 16 },
+// CHECK-EMPTY:
+// CHECK-NEXT: };
+
+// CHECK: static constexpr
+// CHECK-NEXT: const unsigned kernel_signature_start[] = {
+// CHECK-NEXT:  0 // _ZTSZ4mainE8kernel_B
+// CHECK-NEXT: };
+
+// CHECK: template <> struct KernelInfo<class kernel_B> {
+
+#include <sycl.hpp>
+
+using namespace cl::sycl;
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
+  kernelFunc();
+}
+
+int main() {
+
+  int a[5];
+
+  a_kernel<class kernel_B>(
+      [=]() {
+        int local = a[3];
+      });
+}
@@ -0,0 +1,43 @@
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+
+// This test checks a kernel with an argument that is a POD array.
+
+#include <sycl.hpp>
+
+using namespace cl::sycl;
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
+  kernelFunc();
+}
+
+int main() {
+
+  int a[2];
+
+  a_kernel<class kernel_B>(
+      [=]() {
+        int local = a[1];
+      });
+}
+
+// Check kernel_B parameters
+// CHECK: define spir_kernel void @{{.*}}kernel_B
+// CHECK-SAME: i32 [[ELEM_ARG0:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: i32 [[ELEM_ARG1:%[a-zA-Z_]+_[0-9]+]])
+
+// Check local lambda object alloca
+// CHECK: [[LOCAL_OBJECT:%[0-9]+]] = alloca %"class.{{.*}}.anon", align 4
+
+// Check local variables created for parameters
+// CHECK: store i32 [[ELEM_ARG0]], i32* [[ELEM_L0:%[a-zA-Z_]+.addr]], align 4
+// CHECK: store i32 [[ELEM_ARG1]], i32* [[ELEM_L1:%[a-zA-Z_]+.addr[0-9]*]], align 4
+
+// Check init of local array
+// CHECK: [[ARRAY:%[0-9]*]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: [[ARRAY_BEGIN:%[a-zA-Z_.]+]] = getelementptr inbounds [2 x i32], [2 x i32]* [[ARRAY]], i64 0, i64 0
+// CHECK: [[ARRAY0:%[0-9]*]] = load i32, i32* [[ELEM_L0]], align 4
+// CHECK: store i32 [[ARRAY0]], i32* [[ARRAY_BEGIN]], align 4
+// CHECK: [[ARRAY_ELEMENT:%[a-zA-Z_.]+]] = getelementptr inbounds i32, i32* %arrayinit.begin, i64 1
+// CHECK: [[ARRAY1:%[0-9]*]] = load i32, i32* [[ELEM_L1]], align 4
+// CHECK: store i32 [[ARRAY1]], i32* [[ARRAY_ELEMENT]], align 4