Skip to content

Commit f4c367b

Browse files
dyniolssteffenlarsen
dyniols
and
steffenlarsen
authored
[SYCL][E2E] Add local_accessor and accessor tests for sycl_ext_oneapi_free_function_kernels extension (#18672)
This PR adds new e2e tests for free function kernels extension based on test plan https://github.com/intel/llvm/blob/sycl/sycl/test-e2e/FreeFunctionKernels/test-plan.md --------- Co-authored-by: Steffen Larsen <[email protected]>
1 parent 98a87f3 commit f4c367b

File tree

3 files changed

+288
-5
lines changed

3 files changed

+288
-5
lines changed

‎sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp

Lines changed: 189 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,189 @@
1+
// RUN: %{build} -o %t.out
2+
// RUN: %{run} %t.out
3+
4+
// This test verifies whether sycl::accessor can be used with free function
5+
// kernels extension.
6+
7+
#include <sycl/ext/oneapi/free_function_queries.hpp>
8+
9+
#include "helpers.hpp"
10+
11+
template <int Dims>
12+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel))
13+
void globalScopeSingleFreeFunc(
14+
sycl::accessor<int, Dims, sycl::access::mode::read_write,
15+
sycl::access::target::device,
16+
sycl::access::placeholder::false_t>
17+
Accessor,
18+
int Value) {
19+
for (auto &Elem : Accessor)
20+
Elem = Value;
21+
}
22+
namespace ns {
23+
template <int Dims>
24+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<Dims>))
25+
void nsNdRangeFreeFunc(sycl::accessor<int, Dims, sycl::access::mode::read_write,
26+
sycl::access::target::device,
27+
sycl::access::placeholder::false_t>
28+
Accessor,
29+
int Value) {
30+
auto Item = syclext::this_work_item::get_nd_item<Dims>().get_global_id();
31+
Accessor[Item] = Value;
32+
}
33+
} // namespace ns
34+
35+
template <int Dims>
36+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<Dims>))
37+
void ndRangeFreeFuncMultipleParameters(
38+
sycl::accessor<int, Dims, sycl::access::mode::read,
39+
sycl::access::target::device,
40+
sycl::access::placeholder::false_t>
41+
InputAAcc,
42+
sycl::accessor<int, Dims, sycl::access::mode::read,
43+
sycl::access::target::device,
44+
sycl::access::placeholder::false_t>
45+
InputBAcc,
46+
sycl::accessor<int, Dims, sycl::access::mode::write,
47+
sycl::access::target::device,
48+
sycl::access::placeholder::false_t>
49+
ResultAcc) {
50+
auto Item = syclext::this_work_item::get_nd_item<Dims>().get_global_id();
51+
ResultAcc[Item] = InputAAcc[Item] + InputBAcc[Item];
52+
}
53+
54+
// TODO: Need to add checks for a static member functions of a class as free
55+
// function kernel.
56+
57+
template <auto Func, size_t Dims>
58+
int runSingleTaskTest(sycl::queue &Queue, sycl::context &Context,
59+
sycl::range<Dims> NumOfElementsPerDim,
60+
std::string_view ErrorMessage,
61+
const int ExpectedResultValue) {
62+
sycl::kernel UsedKernel = getKernel<Func>(Context);
63+
std::vector<int> ResultData(NumOfElementsPerDim.size(), 0);
64+
{
65+
sycl::buffer<int, Dims> Buffer(ResultData.data(), NumOfElementsPerDim);
66+
Queue.submit([&](sycl::handler &Handler) {
67+
sycl::accessor<int, Dims> Accessor{Buffer, Handler};
68+
Handler.set_args(Accessor, ExpectedResultValue);
69+
Handler.single_task(UsedKernel);
70+
});
71+
}
72+
return performResultCheck(NumOfElementsPerDim.size(), ResultData.data(),
73+
ErrorMessage, ExpectedResultValue);
74+
}
75+
76+
template <auto Func, size_t Dims>
77+
int runNdRangeTest(sycl::queue &Queue, sycl::context &Context,
78+
sycl::nd_range<Dims> NdRange, std::string_view ErrorMessage,
79+
const int ExpectedResultValue) {
80+
sycl::kernel UsedKernel = getKernel<Func>(Context);
81+
std::vector<int> ResultData(NdRange.get_global_range().size(), 0);
82+
{
83+
sycl::buffer<int, Dims> Buffer(ResultData.data(),
84+
NdRange.get_global_range());
85+
Queue.submit([&](sycl::handler &Handler) {
86+
sycl::accessor<int, Dims> Accessor{Buffer, Handler};
87+
Handler.set_args(Accessor, ExpectedResultValue);
88+
Handler.parallel_for(NdRange, UsedKernel);
89+
});
90+
}
91+
return performResultCheck(NdRange.get_global_range().size(),
92+
ResultData.data(), ErrorMessage,
93+
ExpectedResultValue);
94+
}
95+
96+
template <auto Func, size_t Dims>
97+
int runNdRangeTestMultipleParameters(sycl::queue &Queue, sycl::context &Context,
98+
sycl::nd_range<Dims> NdRange,
99+
std::string_view ErrorMessage,
100+
sycl::range<3> Values) {
101+
sycl::kernel UsedKernel = getKernel<Func>(Context);
102+
std::vector<int> InputAData(NdRange.get_global_range().size(), Values[0]);
103+
std::vector<int> InputBData(NdRange.get_global_range().size(), Values[1]);
104+
std::vector<int> ResultData(NdRange.get_global_range().size(), 0);
105+
106+
{
107+
sycl::buffer<int, Dims> InputABuffer(InputAData.data(),
108+
NdRange.get_global_range());
109+
sycl::buffer<int, Dims> InputBBuffer(InputBData.data(),
110+
NdRange.get_global_range());
111+
sycl::buffer<int, Dims> ResultBuffer(ResultData.data(),
112+
NdRange.get_global_range());
113+
Queue.submit([&](sycl::handler &Handler) {
114+
sycl::accessor<int, Dims, sycl::access::mode::read,
115+
sycl::access::target::device>
116+
InputAAcc{InputABuffer, Handler};
117+
sycl::accessor<int, Dims, sycl::access::mode::read,
118+
sycl::access::target::device>
119+
InputBAcc{InputBBuffer, Handler};
120+
sycl::accessor<int, Dims, sycl::access::mode::write> ResultAcc{
121+
ResultBuffer, Handler};
122+
Handler.set_args(InputAAcc, InputBAcc, ResultAcc);
123+
Handler.parallel_for(NdRange, UsedKernel);
124+
});
125+
}
126+
return performResultCheck(NdRange.get_global_range().size(),
127+
ResultData.data(), ErrorMessage, Values[2]);
128+
}
129+
130+
int main() {
131+
132+
int Failed = 0;
133+
sycl::queue Queue;
134+
sycl::context Context = Queue.get_context();
135+
{
136+
// Check that sycl::accessor is supported inside single_task free function
137+
// kernel
138+
Failed += runSingleTaskTest<globalScopeSingleFreeFunc<1>, 1>(
139+
Queue, Context, sycl::range<1>{10},
140+
"globalScopeSingleFreeFunc with sycl::accessor<1>", 1);
141+
Failed += runSingleTaskTest<globalScopeSingleFreeFunc<2>, 2>(
142+
Queue, Context, sycl::range<2>{10, 10},
143+
"globalScopeSingleFreeFunc with sycl::accessor<2>", 2);
144+
Failed += runSingleTaskTest<globalScopeSingleFreeFunc<3>, 3>(
145+
Queue, Context, sycl::range<3>{5, 5, 5},
146+
"globalScopeSingleFreeFunc with sycl::accessor<3>", 3);
147+
}
148+
149+
{
150+
// Check that sycl::accessor is supported inside nd_range free function
151+
// kernel
152+
Failed += runNdRangeTest<ns::nsNdRangeFreeFunc<1>, 1>(
153+
Queue, Context, sycl::nd_range{sycl::range{10}, sycl::range{2}},
154+
"ns::nsNdRangeFreeFunc with sycl::accessor<1>", 4);
155+
Failed += runNdRangeTest<ns::nsNdRangeFreeFunc<2>, 2>(
156+
Queue, Context, sycl::nd_range{sycl::range{16, 16}, sycl::range{4, 4}},
157+
"ns::nsNdRangeFreeFunc with sycl::accessor<2>", 5);
158+
Failed += runNdRangeTest<ns::nsNdRangeFreeFunc<3>, 3>(
159+
Queue, Context,
160+
sycl::nd_range{sycl::range{10, 10, 10}, sycl::range{2, 2, 2}},
161+
"ns::nsNdRangeFreeFunc with sycl::accessor<3>", 6);
162+
}
163+
164+
{
165+
// Check that multiple sycl::accessor are supported inside nd_range free
166+
// function kernel
167+
Failed +=
168+
runNdRangeTestMultipleParameters<ndRangeFreeFuncMultipleParameters<1>,
169+
1>(
170+
Queue, Context, sycl::nd_range{sycl::range{10}, sycl::range{2}},
171+
"ndRangeFreeFuncMultipleParameters with multiple sycl::accessor<1>",
172+
sycl::range{111, 111, 222});
173+
Failed +=
174+
runNdRangeTestMultipleParameters<ndRangeFreeFuncMultipleParameters<2>,
175+
2>(
176+
Queue, Context,
177+
sycl::nd_range{sycl::range{16, 16}, sycl::range{4, 4}},
178+
"ndRangeFreeFuncMultipleParameters with multiple sycl::accessor<2>",
179+
sycl::range{222, 222, 444});
180+
Failed +=
181+
runNdRangeTestMultipleParameters<ndRangeFreeFuncMultipleParameters<3>,
182+
3>(
183+
Queue, Context,
184+
sycl::nd_range{sycl::range{10, 10, 10}, sycl::range{2, 2, 2}},
185+
"ndRangeFreeFuncMultipleParameters with multiple sycl::accessor<3>",
186+
sycl::range{444, 444, 888});
187+
}
188+
return Failed;
189+
}

‎sycl/test-e2e/FreeFunctionKernels/helpers.hpp

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -6,19 +6,19 @@
66
namespace syclext = sycl::ext::oneapi;
77
namespace syclexp = sycl::ext::oneapi::experimental;
88

9-
template <typename T>
9+
template <typename T, typename S>
1010
static int performResultCheck(size_t NumberOfElements, const T *ResultPtr,
1111
std::string_view TestName,
12-
T ExpectedResultValue) {
13-
int IsSuccessful{0};
12+
S ExpectedResultValue) {
13+
int Failed{0};
1414
for (size_t i = 0; i < NumberOfElements; i++) {
1515
if (ResultPtr[i] != ExpectedResultValue) {
1616
std::cerr << "Failed " << TestName << " : " << ResultPtr[i]
1717
<< " != " << ExpectedResultValue << std::endl;
18-
++IsSuccessful;
18+
++Failed;
1919
}
2020
}
21-
return IsSuccessful;
21+
return Failed;
2222
}
2323

2424
template <auto *Func> static sycl::kernel getKernel(sycl::context &Context) {

‎sycl/test-e2e/FreeFunctionKernels/local_accessor_as_kernel_parameter.cpp

Lines changed: 94 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,94 @@
1+
// RUN: %{build} -o %t.out
2+
// RUN: %{run} %t.out
3+
4+
// This test verifies whether sycl::local_accessor can be used with free
5+
// function kernels extension.
6+
7+
#include <sycl/atomic_ref.hpp>
8+
#include <sycl/ext/oneapi/free_function_queries.hpp>
9+
#include <sycl/group_barrier.hpp>
10+
11+
#include "helpers.hpp"
12+
13+
constexpr size_t BIN_SIZE = 4;
14+
constexpr size_t NUM_BINS = 4;
15+
constexpr size_t INPUT_SIZE = 1024;
16+
17+
namespace ns {
18+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
19+
void nsNdRangeFreeFunc(sycl::accessor<int, 1> InputAccessor,
20+
sycl::accessor<int, 1> ResultAccessor,
21+
sycl::local_accessor<int, 1> LocalAccessor) {
22+
23+
size_t LocalWorkItemId =
24+
syclext::this_work_item::get_nd_item<1>().get_local_id();
25+
size_t GlobalWorkItemId =
26+
syclext::this_work_item::get_nd_item<1>().get_global_id();
27+
sycl::group<1> WorkGroup = syclext::this_work_item::get_work_group<1>();
28+
29+
if (LocalWorkItemId < BIN_SIZE)
30+
LocalAccessor[LocalWorkItemId] = 0;
31+
32+
sycl::group_barrier(WorkGroup);
33+
34+
int Value = InputAccessor[GlobalWorkItemId];
35+
sycl::atomic_ref<int, sycl::memory_order::relaxed,
36+
sycl::memory_scope::work_group>
37+
AtomicRefLocal(LocalAccessor[Value]);
38+
AtomicRefLocal++;
39+
sycl::group_barrier(WorkGroup);
40+
41+
if (LocalWorkItemId < BIN_SIZE) {
42+
sycl::atomic_ref<int, sycl::memory_order::relaxed,
43+
sycl::memory_scope::device>
44+
AtomicRefGlobal(ResultAccessor[LocalWorkItemId]);
45+
AtomicRefGlobal.fetch_add(LocalAccessor[LocalWorkItemId]);
46+
}
47+
}
48+
} // namespace ns
49+
50+
// TODO: Need to add checks for a static member functions of a class as free
51+
// function kerenl
52+
53+
void FillWithData(std::vector<int> &Data, std::vector<int> &Values) {
54+
constexpr size_t Offset = INPUT_SIZE / NUM_BINS;
55+
for (size_t i = 0; i < NUM_BINS; ++i) {
56+
std::fill(Data.begin() + i * Offset, Data.begin() + (i + 1) * Offset,
57+
Values[i]);
58+
}
59+
}
60+
61+
int main() {
62+
63+
int Failed = 0;
64+
sycl::queue Queue;
65+
sycl::context Context = Queue.get_context();
66+
{
67+
// Check that sycl::local_accesor is supported inside nd_range free function
68+
// kernel.
69+
std::vector<int> ExpectedHistogramNumbers = {0, 1, 2, 3};
70+
std::vector<int> ResultData(BIN_SIZE, 0);
71+
72+
std::vector<int> InputData(INPUT_SIZE);
73+
FillWithData(InputData, ExpectedHistogramNumbers);
74+
{
75+
sycl::buffer<int, 1> InputBuffer(InputData);
76+
sycl::buffer<int, 1> ResultBuffer(ResultData);
77+
sycl::kernel UsedKernel = getKernel<ns::nsNdRangeFreeFunc>(Context);
78+
Queue.submit([&](sycl::handler &Handler) {
79+
sycl::accessor<int, 1> InputAccessor{InputBuffer, Handler};
80+
sycl::accessor<int, 1> ResultsAccessor{ResultBuffer, Handler};
81+
sycl::local_accessor<int> LocalMemPerWG(sycl::range<1>(BIN_SIZE),
82+
Handler);
83+
Handler.set_args(InputAccessor, ResultsAccessor, LocalMemPerWG);
84+
sycl::nd_range<1> Ndr{INPUT_SIZE, INPUT_SIZE / NUM_BINS};
85+
Handler.parallel_for(Ndr, UsedKernel);
86+
});
87+
}
88+
Failed +=
89+
performResultCheck(NUM_BINS, ResultData.data(),
90+
"sycl::nd_range_kernel with sycl::local_accessor",
91+
INPUT_SIZE / NUM_BINS);
92+
}
93+
return Failed;
94+
}

0 commit comments

Comments
 (0)