Skip to content

Commit 605a700

Browse files
dyniolsdyniols
committed
added local_accessor and accessor test for free_func_kernels extension
1 parent d877bb9 commit 605a700

File tree

3 files changed

+175
-2
lines changed

3 files changed

+175
-2
lines changed

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

Lines changed: 80 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,80 @@
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+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel))
12+
void globalScopeSingleFreeFunc(sycl::accessor<int, 1> Accessor,
13+
size_t NumOfElements, int Value) {
14+
for (size_t i = 0; i < NumOfElements; ++i) {
15+
Accessor[i] = Value;
16+
}
17+
}
18+
19+
namespace ns {
20+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<3>))
21+
void nsNdRangeFreeFunc(sycl::accessor<int, 1> Accessor, int Value) {
22+
size_t Item =
23+
syclext::this_work_item::get_nd_item<3>().get_global_linear_id();
24+
Accessor[Item] = Value;
25+
}
26+
} // namespace ns
27+
28+
// TODO: Need to add checks for a static member functions of a class as free
29+
// function kerenl
30+
31+
int main() {
32+
33+
int Failed = 0;
34+
sycl::queue Queue;
35+
sycl::context Context = Queue.get_context();
36+
constexpr size_t NumOfElements = 1024;
37+
{
38+
// Check that sycl::accessor is supported inside nd_range free function
39+
// kernel.
40+
std::vector<int> ResultHostData(NumOfElements, 0);
41+
constexpr int ExpectedResultValue = 111;
42+
{
43+
sycl::buffer<int, 1> Buffer(ResultHostData);
44+
sycl::kernel UsedKernel = getKernel<ns::nsNdRangeFreeFunc>(Context);
45+
Queue.submit([&](sycl::handler &Handler) {
46+
sycl::accessor<int, 1> Accessor{Buffer, Handler};
47+
Handler.set_args(Accessor, ExpectedResultValue);
48+
sycl::nd_range<3> Ndr{{4, 4, NumOfElements / 16}, {4, 4, 4}};
49+
Handler.parallel_for(Ndr, UsedKernel);
50+
});
51+
}
52+
53+
Failed += performResultCheck(NumOfElements, ResultHostData.data(),
54+
"ns::nsNdRangeFreeFunc with sycl::accessor",
55+
ExpectedResultValue);
56+
}
57+
58+
{
59+
// Check that sycl::accessor is supported inside single_task free function
60+
// kernel.
61+
std::vector<int> ResultHostData(NumOfElements, 0);
62+
constexpr int ExpectedResultValue = 222;
63+
{
64+
sycl::buffer<int, 1> Buffer(ResultHostData);
65+
sycl::kernel UsedKernel = getKernel<globalScopeSingleFreeFunc>(Context);
66+
Queue.submit([&](sycl::handler &Handler) {
67+
sycl::accessor<int, 1> Accessor{Buffer, Handler};
68+
Handler.set_arg(0, Accessor);
69+
Handler.set_arg(1, NumOfElements);
70+
Handler.set_arg(2, ExpectedResultValue);
71+
Handler.single_task(UsedKernel);
72+
});
73+
}
74+
Failed += performResultCheck(
75+
NumOfElements, ResultHostData.data(),
76+
"globalScopeSingleFreeFunc with sycl::accessor", ExpectedResultValue);
77+
}
78+
79+
return Failed;
80+
}

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -6,10 +6,10 @@
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) {
12+
S ExpectedResultValue) {
1313
int IsSuccessful{0};
1414
for (size_t i = 0; i < NumberOfElements; i++) {
1515
if (ResultPtr[i] != ExpectedResultValue) {

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

Lines changed: 93 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,93 @@
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+
int main() {
61+
62+
int Failed = 0;
63+
sycl::queue Queue;
64+
sycl::context Context = Queue.get_context();
65+
{
66+
// Check that sycl::local_accesor is supported inside nd_range free function
67+
// kernel.
68+
std::vector<int> ExpectedHistogramNumbers = {0, 1, 2, 3};
69+
std::vector<int> ResultData(BIN_SIZE, 0);
70+
71+
std::vector<int> InputData(INPUT_SIZE);
72+
FillWithData(InputData, ExpectedHistogramNumbers);
73+
{
74+
sycl::buffer<int, 1> InputBuffer(InputData);
75+
sycl::buffer<int, 1> ResultBuffer(ResultData);
76+
sycl::kernel UsedKernel = getKernel<ns::nsNdRangeFreeFunc>(Context);
77+
Queue.submit([&](sycl::handler &Handler) {
78+
sycl::accessor<int, 1> InputAccessor{InputBuffer, Handler};
79+
sycl::accessor<int, 1> ResultsAccessor{ResultBuffer, Handler};
80+
sycl::local_accessor<int> LocalMemPerWG(sycl::range<1>(BIN_SIZE),
81+
Handler);
82+
Handler.set_args(InputAccessor, ResultsAccessor, LocalMemPerWG);
83+
sycl::nd_range<1> Ndr{INPUT_SIZE, INPUT_SIZE / NUM_BINS};
84+
Handler.parallel_for(Ndr, UsedKernel);
85+
});
86+
}
87+
Failed +=
88+
performResultCheck(NUM_BINS, ResultData.data(),
89+
"sycl::nd_range_kernel with sycl::local_accessor",
90+
INPUT_SIZE / NUM_BINS);
91+
}
92+
return Failed;
93+
}

0 commit comments

Comments
 (0)