Skip to content

Commit b0dc1ce

Browse files
dyniolsaelovikov-intel
dyniols
and
aelovikov-intel
authored
[SYCL] Add test for type aliases as kernel parameter within free function kernel (#18850)
This PR adds new test for free function kernels extension based on test plan https://github.com/intel/llvm/blob/sycl/sycl/test-e2e/FreeFunctionKernels/test-plan.md#test-type-aliases-to-allowed-kernel-paramater-types-as-kernel-parameter Extension spec: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc It removes unused alias of type `accessor<int, 1, access::mode::read_write>` from `sycl/test/extensions/free_function_errors.cpp` --------- Co-authored-by: aelovikov-intel <[email protected]>
1 parent ba68d37 commit b0dc1ce

File tree

2 files changed

+176
-2
lines changed

2 files changed

+176
-2
lines changed

‎sycl/test/extensions/free_function_errors.cpp

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -23,8 +23,6 @@ union U {
2323
float f;
2424
};
2525

26-
using accType = accessor<int, 1, access::mode::read_write>;
27-
2826
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
2927
(ext::oneapi::experimental::single_task_kernel))
3028
void ff(struct S s) {}

‎sycl/test/extensions/free_function_kernels/free_function_kernels_type_aliases.cpp

Lines changed: 176 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,176 @@
1+
// RUN: %clangxx -fsyntax-only -fsycl %s
2+
// This test verifies whether type aliases are allowed as kernel parameter to
3+
// free function kernel.
4+
#include <sycl/sycl.hpp>
5+
6+
namespace syclexp = sycl::ext::oneapi::experimental;
7+
8+
struct TestStruct {};
9+
using StructAlias = TestStruct;
10+
typedef TestStruct StructAliasTypedef;
11+
12+
class TestClass {};
13+
using ClassAlias = TestClass;
14+
typedef TestClass ClassAliasTypedef;
15+
namespace ns1 {
16+
using WriteOnlyAcc = sycl::accessor<int, 1, sycl::access::mode::write>;
17+
typedef sycl::accessor<int, 1, sycl::access::mode::write>
18+
WriteOnlyAccAliasTypedef;
19+
20+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel)
21+
void singleTaskKernelStruct(StructAlias Type) {}
22+
23+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>)
24+
void ndRangeKernelStruct(StructAlias Type) {}
25+
26+
namespace ns2 {
27+
28+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel)
29+
void singleTaskKernelStructAliasTypedef(StructAliasTypedef Type) {}
30+
31+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>)
32+
void ndRangeKernelStructAliasTypedef(StructAliasTypedef Type) {}
33+
34+
template <typename A, typename B, typename C, typename D> struct TestStructY {};
35+
36+
template <typename A, typename B>
37+
using AliasTypeY = TestStructY<A, B, int, float>;
38+
39+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel)
40+
void singleTaskKernelClass(ClassAlias Type) {}
41+
42+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>)
43+
void ndRangeKernelClass(ClassAlias Type) {}
44+
45+
namespace ns3 {
46+
47+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel)
48+
void singleTaskKernelClassAliasTypedef(ClassAliasTypedef Type) {}
49+
50+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>)
51+
void ndRangeKernelClassAliasTypedef(ClassAliasTypedef Type) {}
52+
53+
using LocalAcc = sycl::local_accessor<int, 1>;
54+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel)
55+
void singleTaskWriteOnlyAcc(WriteOnlyAcc Type) {}
56+
57+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>)
58+
void ndRangeKernelWriteOnlyAcc(WriteOnlyAcc Type) {}
59+
60+
namespace ns4 {
61+
template <typename A, typename B, typename C> struct TestStructX {};
62+
using AliasTypeX = TestStructX<int, float, int>;
63+
64+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel)
65+
void singleTaskWriteOnlyAccAliasTypedef(WriteOnlyAccAliasTypedef Type) {}
66+
67+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>)
68+
void ndRangeKernelWriteOnlyAccAliasTypedef(WriteOnlyAccAliasTypedef Type) {}
69+
70+
} // namespace ns4
71+
} // namespace ns3
72+
} // namespace ns2
73+
} // namespace ns1
74+
75+
namespace ns5 {
76+
using ReadOnlyAcc = sycl::accessor<float, 1, sycl::access::mode::read>;
77+
typedef sycl::accessor<float, 1, sycl::access::mode::read>
78+
ReadOnlyAccAliasTypedef;
79+
80+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel)
81+
void singleTaskReadOnlyAcc(ReadOnlyAcc Type) {}
82+
83+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>)
84+
void ndRangeKernelReadOnlyAcc(ReadOnlyAcc Type) {}
85+
86+
} // namespace ns5
87+
88+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel)
89+
void singleTaskReadOnlyAccAliasTypedef(ns5::ReadOnlyAccAliasTypedef Type) {}
90+
91+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>)
92+
void ndRangeKernelReadOnlyAccAliasTypedef(ns5::ReadOnlyAccAliasTypedef Type) {}
93+
94+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel)
95+
void singleTaskLocalAcc(ns1::ns2::ns3::LocalAcc Type) {}
96+
97+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>)
98+
void ndRangeKernelLocalAcc(ns1::ns2::ns3::LocalAcc Type) {}
99+
100+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel)
101+
void singleTaskTemplatedType(ns1::ns2::ns3::ns4::AliasTypeX Type) {}
102+
103+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>)
104+
void ndRangeKernelTemplatedType(ns1::ns2::ns3::ns4::AliasTypeX Type) {}
105+
106+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel)
107+
void singleTaskTemplatedTypeAliasTemplatedType(
108+
ns1::ns2::AliasTypeY<int, int> Type) {}
109+
110+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>)
111+
void ndRangeKernelTemplatedTypeAliasTemplatedType(
112+
ns1::ns2::AliasTypeY<float, float> Type) {}
113+
114+
template <auto Func, typename T> void runNdRangeCheck(T Type) {
115+
sycl::queue Queue;
116+
sycl::kernel_bundle Bundle =
117+
get_kernel_bundle<sycl::bundle_state::executable>(Queue.get_context());
118+
sycl::kernel_id Id = syclexp::get_kernel_id<Func>();
119+
sycl::kernel Kernel = Bundle.get_kernel(Id);
120+
Queue.submit([&](sycl::handler &Handler) {
121+
Handler.set_args(Type);
122+
Handler.parallel_for(sycl::nd_range{{1}, {1}}, Kernel);
123+
});
124+
}
125+
126+
template <auto Func, typename T> void runSingleTaskTest(T Type) {
127+
sycl::queue Queue;
128+
sycl::kernel_bundle Bundle =
129+
get_kernel_bundle<sycl::bundle_state::executable>(Queue.get_context());
130+
sycl::kernel_id Id = syclexp::get_kernel_id<Func>();
131+
sycl::kernel Kernel = Bundle.get_kernel(Id);
132+
Queue.submit([&](sycl::handler &Handler) {
133+
Handler.set_args(Type);
134+
Handler.single_task(Kernel);
135+
});
136+
}
137+
138+
int main() {
139+
runSingleTaskTest<ns1::singleTaskKernelStruct>(StructAlias());
140+
runSingleTaskTest<ns1::ns2::singleTaskKernelClass>(ClassAlias());
141+
runSingleTaskTest<ns5::singleTaskReadOnlyAcc>(ns5::ReadOnlyAcc());
142+
runSingleTaskTest<ns1::ns2::ns3::singleTaskWriteOnlyAcc>(ns1::WriteOnlyAcc());
143+
runSingleTaskTest<singleTaskLocalAcc>(ns1::ns2::ns3::LocalAcc());
144+
runSingleTaskTest<singleTaskTemplatedType>(ns1::ns2::ns3::ns4::AliasTypeX());
145+
runSingleTaskTest<singleTaskTemplatedTypeAliasTemplatedType>(
146+
ns1::ns2::AliasTypeY<int, int>());
147+
runSingleTaskTest<ns1::ns2::singleTaskKernelStructAliasTypedef>(
148+
StructAliasTypedef());
149+
runSingleTaskTest<ns1::ns2::ns3::singleTaskKernelClassAliasTypedef>(
150+
ClassAliasTypedef());
151+
runSingleTaskTest<ns1::ns2::ns3::ns4::singleTaskWriteOnlyAccAliasTypedef>(
152+
ns1::WriteOnlyAccAliasTypedef());
153+
runSingleTaskTest<singleTaskReadOnlyAccAliasTypedef>(
154+
ns5::ReadOnlyAccAliasTypedef());
155+
156+
runNdRangeCheck<ns1::ndRangeKernelStruct>(StructAlias());
157+
runNdRangeCheck<ns1::ns2::singleTaskKernelClass>(ClassAlias());
158+
runNdRangeCheck<ns5::ndRangeKernelReadOnlyAcc>(ns5::ReadOnlyAcc());
159+
runNdRangeCheck<ns1::ns2::ns3::ndRangeKernelWriteOnlyAcc>(
160+
ns1::WriteOnlyAcc());
161+
runNdRangeCheck<ndRangeKernelLocalAcc>(ns1::ns2::ns3::LocalAcc());
162+
runNdRangeCheck<ndRangeKernelTemplatedType>(ns1::ns2::ns3::ns4::AliasTypeX());
163+
runNdRangeCheck<ndRangeKernelTemplatedTypeAliasTemplatedType>(
164+
ns1::ns2::AliasTypeY<float, float>());
165+
166+
runNdRangeCheck<ns1::ns2::ndRangeKernelStructAliasTypedef>(
167+
StructAliasTypedef());
168+
runNdRangeCheck<ns1::ns2::ns3::ndRangeKernelClassAliasTypedef>(
169+
ClassAliasTypedef());
170+
runNdRangeCheck<ns1::ns2::ns3::ns4::ndRangeKernelWriteOnlyAccAliasTypedef>(
171+
ns1::WriteOnlyAccAliasTypedef());
172+
runNdRangeCheck<ndRangeKernelReadOnlyAccAliasTypedef>(
173+
ns5::ReadOnlyAccAliasTypedef());
174+
175+
return 0;
176+
}

0 commit comments

Comments
 (0)