Skip to content

Commit 00e7308

Browse files
rdeodharrdeodhar
authored
[SYCL] Add support for arrays as kernel parameters (#1841)
This patch adds support for arrays as kernel parameters. Arrays within structs are not currently supported and will be supported when struct decomposition is implemented.
1 parent 99957c5 commit 00e7308

13 files changed

+1497
-64
lines changed

‎clang/lib/Sema/SemaSYCL.cpp

+193-63
Large diffs are not rendered by default.

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

+52
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,52 @@
1+
// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only
2+
// RUN: FileCheck -input-file=%t.h %s
3+
4+
// This test checks the integration header generated when
5+
// the kernel argument is an Accessor array.
6+
7+
// CHECK: #include <CL/sycl/detail/kernel_desc.hpp>
8+
9+
// CHECK: class kernel_A;
10+
11+
// CHECK: __SYCL_INLINE_NAMESPACE(cl) {
12+
// CHECK-NEXT: namespace sycl {
13+
// CHECK-NEXT: namespace detail {
14+
15+
// CHECK: static constexpr
16+
// CHECK-NEXT: const char* const kernel_names[] = {
17+
// CHECK-NEXT: "_ZTSZ4mainE8kernel_A"
18+
// CHECK-NEXT: };
19+
20+
// CHECK: static constexpr
21+
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
22+
// CHECK-NEXT: //--- _ZTSZ4mainE8kernel_A
23+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 0 },
24+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 12 },
25+
// CHECK-EMPTY:
26+
// CHECK-NEXT: };
27+
28+
// CHECK: static constexpr
29+
// CHECK-NEXT: const unsigned kernel_signature_start[] = {
30+
// CHECK-NEXT: 0 // _ZTSZ4mainE8kernel_A
31+
// CHECK-NEXT: };
32+
33+
// CHECK: template <> struct KernelInfo<class kernel_A> {
34+
35+
#include <sycl.hpp>
36+
37+
using namespace cl::sycl;
38+
39+
template <typename name, typename Func>
40+
__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
41+
kernelFunc();
42+
}
43+
44+
int main() {
45+
46+
using Accessor =
47+
accessor<int, 1, access::mode::read_write, access::target::global_buffer>;
48+
49+
Accessor acc[2];
50+
51+
a_kernel<class kernel_A>([=]() { acc[1].use(); });
52+
}

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

+78
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,78 @@
1+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
2+
3+
// This test checks a kernel argument that is an Accessor array
4+
5+
#include <sycl.hpp>
6+
7+
using namespace cl::sycl;
8+
9+
template <typename name, typename Func>
10+
__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
11+
kernelFunc();
12+
}
13+
14+
int main() {
15+
16+
using Accessor =
17+
accessor<int, 1, access::mode::read_write, access::target::global_buffer>;
18+
Accessor acc[2];
19+
20+
a_kernel<class kernel_A>(
21+
[=]() {
22+
acc[1].use();
23+
});
24+
}
25+
26+
// Check kernel_A parameters
27+
// CHECK: define spir_kernel void @{{.*}}kernel_A
28+
// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG1:%[a-zA-Z0-9_]+]],
29+
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+_1]],
30+
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+_2]],
31+
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+_3]],
32+
// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG2:%[a-zA-Z0-9_]+_4]],
33+
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+_6]],
34+
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+_7]],
35+
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET2:%[a-zA-Z0-9_]+_8]])
36+
37+
// CHECK alloca for pointer arguments
38+
// CHECK: [[MEM_ARG1:%[a-zA-Z0-9_.]+]] = alloca i32 addrspace(1)*, align 8
39+
// CHECK: [[MEM_ARG2:%[a-zA-Z0-9_.]+]] = alloca i32 addrspace(1)*, align 8
40+
41+
// CHECK lambda object alloca
42+
// CHECK: [[LOCAL_OBJECT:%0]] = alloca %"class.{{.*}}.anon", align 4
43+
44+
// CHECK allocas for ranges
45+
// CHECK: [[ACC_RANGE1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
46+
// CHECK: [[MEM_RANGE1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
47+
// CHECK: [[OFFSET1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id"
48+
// CHECK: [[ACC_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
49+
// CHECK: [[MEM_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
50+
// CHECK: [[OFFSET2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id"
51+
52+
// CHECK accessor array default inits
53+
// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
54+
// CHECK: [[BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], [2 x [[ACCESSOR]]]* [[ACCESSOR_ARRAY1]], i64 0, i64 0
55+
// CHECK: [[END:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR:.*]], [[ACCESSOR]]* [[BEGIN]], i64 2
56+
// CHECK: [[NEXT0:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* {{.*}}, i64 1
57+
// CHECK: [[ELEMENT:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* {{.*}}, i64 1
58+
// CHECK: [[ELEMENT:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* {{.*}}, i64 2
59+
// CHECK: [[NEXT1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* {{.*}}, i64 1
60+
61+
// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
62+
// CHECK: [[INDEX:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[ACCESSOR_ARRAY2]], i64 0, i64 0
63+
64+
// CHECK load from kernel pointer argument alloca
65+
// CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]]
66+
67+
// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast [[ACCESSOR]]* {{.*}} to [[ACCESSOR]] addrspace(4)*
68+
69+
// CHECK acc[0] __init method call
70+
// 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]])
71+
72+
// CHECK load from kernel pointer argument alloca
73+
// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG2]]
74+
75+
// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast [[ACCESSOR]]* {{.*}} to [[ACCESSOR]] addrspace(4)*
76+
77+
// CHECK acc[1] __init method call
78+
// 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]])

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

+60
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,60 @@
1+
// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only
2+
// RUN: FileCheck -input-file=%t.h %s
3+
// XFAIL for now due to : https://github.com/intel/llvm/issues/2018
4+
// XFAIL: *
5+
6+
// This test checks the integration header when kernel argument
7+
// is a struct containing an Accessor array.
8+
9+
// CHECK: #include <CL/sycl/detail/kernel_desc.hpp>
10+
11+
// CHECK: class kernel_C;
12+
13+
// CHECK: __SYCL_INLINE_NAMESPACE(cl) {
14+
// CHECK-NEXT: namespace sycl {
15+
// CHECK-NEXT: namespace detail {
16+
17+
// CHECK: static constexpr
18+
// CHECK-NEXT: const char* const kernel_names[] = {
19+
// CHECK-NEXT: "_ZTSZ4mainE8kernel_C"
20+
// CHECK-NEXT: };
21+
22+
// CHECK: static constexpr
23+
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
24+
// CHECK-NEXT: //--- _ZTSZ4mainE8kernel_C
25+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 24, 0 },
26+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 0 },
27+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 12 },
28+
// CHECK-EMPTY:
29+
// CHECK-NEXT: };
30+
31+
// CHECK: static constexpr
32+
// CHECK-NEXT: const unsigned kernel_signature_start[] = {
33+
// CHECK-NEXT: 0 // _ZTSZ4mainE8kernel_C
34+
// CHECK-NEXT: };
35+
36+
// CHECK: template <> struct KernelInfo<class kernel_C> {
37+
38+
#include <sycl.hpp>
39+
40+
using namespace cl::sycl;
41+
42+
template <typename name, typename Func>
43+
__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
44+
kernelFunc();
45+
}
46+
47+
int main() {
48+
49+
using Accessor =
50+
accessor<int, 1, access::mode::read_write, access::target::global_buffer>;
51+
52+
struct struct_acc_t {
53+
Accessor member_acc[2];
54+
} struct_acc;
55+
56+
a_kernel<class kernel_C>(
57+
[=]() {
58+
struct_acc.member_acc[1].use();
59+
});
60+
}

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

+85
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,85 @@
1+
// 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
2+
// XFAIL: *
3+
4+
// This test checks a kernel with struct parameter that contains an Accessor array.
5+
6+
#include <sycl.hpp>
7+
8+
using namespace cl::sycl;
9+
10+
template <typename name, typename Func>
11+
__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
12+
kernelFunc();
13+
}
14+
15+
int main() {
16+
17+
using Accessor =
18+
accessor<int, 1, access::mode::read_write, access::target::global_buffer>;
19+
20+
struct struct_acc_t {
21+
Accessor member_acc[2];
22+
} struct_acc;
23+
24+
a_kernel<class kernel_C>(
25+
[=]() {
26+
struct_acc.member_acc[1].use();
27+
});
28+
}
29+
30+
// CHECK kernel_C parameters
31+
// CHECK: define spir_kernel void @{{.*}}kernel_C
32+
// CHECK-SAME: %struct.{{.*}}.struct_acc_t* byval(%struct.{{.*}}.struct_acc_t) align 4 [[STRUCT:%[a-zA-Z0-9_]+]],
33+
// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG1:%[a-zA-Z0-9_]+]],
34+
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
35+
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
36+
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]],
37+
// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG2:%[a-zA-Z0-9_]+4]],
38+
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+6]],
39+
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+7]],
40+
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2:%[a-zA-Z0-9_]+8]])
41+
42+
// Check alloca for pointer arguments
43+
// CHECK: [[MEM_ARG1]].addr{{[0-9]*}} = alloca i32 addrspace(1)*, align 8
44+
// CHECK: [[MEM_ARG1]].addr{{[0-9]*}} = alloca i32 addrspace(1)*, align 8
45+
46+
// Check lambda object alloca
47+
// CHECK: [[LOCAL_OBJECT:%0]] = alloca %"class.{{.*}}.anon", align 4
48+
49+
// Check allocas for ranges
50+
// CHECK: [[ACC_RANGE1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
51+
// CHECK: [[MEM_RANGE1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
52+
// CHECK: [[OFFSET1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id"
53+
// CHECK: [[ACC_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
54+
// CHECK: [[MEM_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
55+
// CHECK: [[OFFSET2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id"
56+
57+
// Check init of local struct
58+
// CHECK: [[L_STRUCT_ADDR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
59+
// CHECK: [[MEMCPY_DST:%[0-9a-zA-Z_]+]] = bitcast %struct.{{.*}}struct_acc_t* [[L_STRUCT_ADDR]] to i8*
60+
// CHECK: [[MEMCPY_SRC:%[0-9a-zA-Z_]+]] = bitcast %struct.{{.*}}struct_acc_t* %{{[0-9a-zA-Z_]+}} to i8*
61+
// CHECK: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 [[MEMCPY_DST]], i8* align 4 [[MEMCPY_SRC]], i64 24, i1 false)
62+
63+
// Check accessor array GEP for member_acc[0]
64+
// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
65+
// CHECK: [[MEMBER1:%[a-zA-Z_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[ACCESSOR_ARRAY1]], i32 0, i32 0
66+
// 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
67+
68+
// Check load from kernel pointer argument alloca
69+
// CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr{{[0-9]*}}
70+
71+
// Check acc[0] __init method call
72+
// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[Z0]] to %"class{{.*}}accessor" addrspace(4)*
73+
// 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]])
74+
75+
// Check accessor array GEP for member_acc[1]
76+
// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
77+
// CHECK: [[MEMBER2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[ACCESSOR_ARRAY2]], i32 0, i32 0
78+
// 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
79+
80+
// Check load from kernel pointer argument alloca
81+
// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr{{[0-9]*}}
82+
83+
// Check acc[1] __init method call
84+
// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[Z1]] to %"class{{.*}}accessor" addrspace(4)*
85+
// 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]])

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

+55
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,55 @@
1+
// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only
2+
// RUN: FileCheck -input-file=%t.h %s
3+
4+
// This test checks the integration header generated for a kernel
5+
// with an argument that is a POD array.
6+
7+
// CHECK: #include <CL/sycl/detail/kernel_desc.hpp>
8+
9+
// CHECK: class kernel_B;
10+
11+
// CHECK: __SYCL_INLINE_NAMESPACE(cl) {
12+
// CHECK-NEXT: namespace sycl {
13+
// CHECK-NEXT: namespace detail {
14+
15+
// CHECK: static constexpr
16+
// CHECK-NEXT: const char* const kernel_names[] = {
17+
// CHECK-NEXT: "_ZTSZ4mainE8kernel_B"
18+
// CHECK-NEXT: };
19+
20+
// CHECK: static constexpr
21+
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
22+
// CHECK-NEXT: //--- _ZTSZ4mainE8kernel_B
23+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
24+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 4 },
25+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 8 },
26+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 12 },
27+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 16 },
28+
// CHECK-EMPTY:
29+
// CHECK-NEXT: };
30+
31+
// CHECK: static constexpr
32+
// CHECK-NEXT: const unsigned kernel_signature_start[] = {
33+
// CHECK-NEXT: 0 // _ZTSZ4mainE8kernel_B
34+
// CHECK-NEXT: };
35+
36+
// CHECK: template <> struct KernelInfo<class kernel_B> {
37+
38+
#include <sycl.hpp>
39+
40+
using namespace cl::sycl;
41+
42+
template <typename name, typename Func>
43+
__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
44+
kernelFunc();
45+
}
46+
47+
int main() {
48+
49+
int a[5];
50+
51+
a_kernel<class kernel_B>(
52+
[=]() {
53+
int local = a[3];
54+
});
55+
}

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

+43
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,43 @@
1+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
2+
3+
// This test checks a kernel with an argument that is a POD array.
4+
5+
#include <sycl.hpp>
6+
7+
using namespace cl::sycl;
8+
9+
template <typename name, typename Func>
10+
__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
11+
kernelFunc();
12+
}
13+
14+
int main() {
15+
16+
int a[2];
17+
18+
a_kernel<class kernel_B>(
19+
[=]() {
20+
int local = a[1];
21+
});
22+
}
23+
24+
// Check kernel_B parameters
25+
// CHECK: define spir_kernel void @{{.*}}kernel_B
26+
// CHECK-SAME: i32 [[ELEM_ARG0:%[a-zA-Z0-9_]+]],
27+
// CHECK-SAME: i32 [[ELEM_ARG1:%[a-zA-Z_]+_[0-9]+]])
28+
29+
// Check local lambda object alloca
30+
// CHECK: [[LOCAL_OBJECT:%[0-9]+]] = alloca %"class.{{.*}}.anon", align 4
31+
32+
// Check local variables created for parameters
33+
// CHECK: store i32 [[ELEM_ARG0]], i32* [[ELEM_L0:%[a-zA-Z_]+.addr]], align 4
34+
// CHECK: store i32 [[ELEM_ARG1]], i32* [[ELEM_L1:%[a-zA-Z_]+.addr[0-9]*]], align 4
35+
36+
// Check init of local array
37+
// CHECK: [[ARRAY:%[0-9]*]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
38+
// CHECK: [[ARRAY_BEGIN:%[a-zA-Z_.]+]] = getelementptr inbounds [2 x i32], [2 x i32]* [[ARRAY]], i64 0, i64 0
39+
// CHECK: [[ARRAY0:%[0-9]*]] = load i32, i32* [[ELEM_L0]], align 4
40+
// CHECK: store i32 [[ARRAY0]], i32* [[ARRAY_BEGIN]], align 4
41+
// CHECK: [[ARRAY_ELEMENT:%[a-zA-Z_.]+]] = getelementptr inbounds i32, i32* %arrayinit.begin, i64 1
42+
// CHECK: [[ARRAY1:%[0-9]*]] = load i32, i32* [[ELEM_L1]], align 4
43+
// CHECK: store i32 [[ARRAY1]], i32* [[ARRAY_ELEMENT]], align 4

0 commit comments

Comments
 (0)