Skip to content

Commit 09600fa

Browse files
dyniolsdyniols
authored
[SYCL][E2E] Add new tests for free function kernels extension and move e2e tests for that feature to common directory (#18256)
This PR adds new tests for free function kernels extension based on test plan https://github.com/intel/llvm/blob/sycl/sycl/test-e2e/FreeFunctionKernels/test-plan.md and moves tests from `sycl/test-e2e/Experimental/free_functions` into `sycl/test-e2e/FreeFunctionKernels`
1 parent a153a85 commit 09600fa

11 files changed

+690
-0
lines changed

‎sycl/test-e2e/Experimental/free_functions/address_test.cpp renamed to ‎sycl/test-e2e/FreeFunctionKernels/address_test.cpp

File renamed without changes.

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

Lines changed: 29 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,29 @@
1+
#include <string_view>
2+
3+
#include <sycl/detail/core.hpp>
4+
#include <sycl/kernel_bundle.hpp>
5+
6+
namespace syclext = sycl::ext::oneapi;
7+
namespace syclexp = sycl::ext::oneapi::experimental;
8+
9+
template <typename T>
10+
static int performResultCheck(size_t NumberOfElements, const T *ResultPtr,
11+
std::string_view TestName,
12+
T ExpectedResultValue) {
13+
int IsSuccessful{0};
14+
for (size_t i = 0; i < NumberOfElements; i++) {
15+
if (ResultPtr[i] != ExpectedResultValue) {
16+
std::cerr << "Failed " << TestName << " : " << ResultPtr[i]
17+
<< " != " << ExpectedResultValue << std::endl;
18+
++IsSuccessful;
19+
}
20+
}
21+
return IsSuccessful;
22+
}
23+
24+
template <auto *Func> static sycl::kernel getKernel(sycl::context &Context) {
25+
sycl::kernel_bundle<sycl::bundle_state::executable> KernelBundle =
26+
syclexp::get_kernel_bundle<Func, sycl::bundle_state::executable>(Context);
27+
sycl::kernel KernelId = KernelBundle.ext_oneapi_get_kernel<Func>();
28+
return KernelId;
29+
}

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

Lines changed: 132 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,132 @@
1+
// REQUIRES: aspect-usm_shared_allocations
2+
// RUN: %{build} -o %t.out
3+
// RUN: %{run} %t.out
4+
5+
// This test verifies whether an id<Dimensions> can be passed as a kernel
6+
// parameter to a free function kernel.
7+
8+
#include <sycl/usm.hpp>
9+
10+
#include <sycl/ext/oneapi/free_function_queries.hpp>
11+
12+
#include "helpers.hpp"
13+
14+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel))
15+
void globalScopeSingleFreeFuncKernel(int *Ptr, size_t NumOfElements,
16+
sycl::id<1> Id) {
17+
for (size_t i = 0; i < NumOfElements; ++i) {
18+
Ptr[i] = static_cast<int>(Id[0]);
19+
}
20+
}
21+
22+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<2>))
23+
void globalScopeNdRangeFreeFuncKernel(int *Ptr, sycl::id<2> Id) {
24+
size_t Item =
25+
syclext::this_work_item::get_nd_item<2>().get_global_linear_id();
26+
Ptr[Item] = static_cast<int>(Id[0] + Id[1]);
27+
}
28+
29+
namespace ns {
30+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel))
31+
void nsSingleFreeFuncKernel(int *Ptr, size_t NumOfElements, sycl::id<1> Id) {
32+
for (size_t i = 0; i < NumOfElements; ++i) {
33+
Ptr[i] = static_cast<int>(Id[0]);
34+
}
35+
}
36+
37+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<3>))
38+
void nsNdRangeFreeFuncKernel(int *Ptr, sycl::id<3> Id) {
39+
size_t Item =
40+
syclext::this_work_item::get_nd_item<3>().get_global_linear_id();
41+
Ptr[Item] = static_cast<int>(Id[0] + Id[1] + Id[2]);
42+
}
43+
} // namespace ns
44+
45+
// TODO: Need to add checks for a static member functions of a class as free
46+
// function kerenl
47+
48+
int main() {
49+
int Failed = 0;
50+
sycl::queue Queue;
51+
sycl::context Context = Queue.get_context();
52+
constexpr size_t NumOfElements = 1024;
53+
int *Data = sycl::malloc_shared<int>(NumOfElements, Queue);
54+
55+
{
56+
std::fill(Data, Data + NumOfElements, 0);
57+
sycl::kernel UsedKernel = getKernel<ns::nsSingleFreeFuncKernel>(Context);
58+
59+
sycl::id<1> Id{11};
60+
int ExpectedResultValue = static_cast<int>(Id[0]);
61+
Queue
62+
.submit([&](sycl::handler &Handler) {
63+
Handler.set_args(Data, NumOfElements, Id);
64+
Handler.single_task(UsedKernel);
65+
})
66+
.wait();
67+
68+
Failed += performResultCheck(
69+
NumOfElements, Data, "ns::nsSingleFreeFuncKernel", ExpectedResultValue);
70+
}
71+
72+
{
73+
std::fill(Data, Data + NumOfElements, 0);
74+
sycl::kernel UsedKernel = getKernel<ns::nsNdRangeFreeFuncKernel>(Context);
75+
sycl::id<3> Id{22, 22, 22};
76+
int ExpectedResultValue = static_cast<int>(Id[0] + Id[1] + Id[2]);
77+
Queue
78+
.submit([&](sycl::handler &Handler) {
79+
Handler.set_arg(0, Data);
80+
Handler.set_arg(1, Id);
81+
sycl::nd_range<3> Ndr{{4, 4, NumOfElements / 16}, {4, 4, 4}};
82+
Handler.parallel_for(Ndr, UsedKernel);
83+
})
84+
.wait();
85+
86+
Failed +=
87+
performResultCheck(NumOfElements, Data, "ns::nsNdRangeFreeFuncKernel",
88+
ExpectedResultValue);
89+
}
90+
91+
{
92+
std::fill(Data, Data + NumOfElements, 0);
93+
sycl::kernel UsedKernel =
94+
getKernel<globalScopeSingleFreeFuncKernel>(Context);
95+
sycl::id<1> Id{33};
96+
int ExpectedResultValue = static_cast<int>(Id[0]);
97+
Queue
98+
.submit([&](sycl::handler &Handler) {
99+
Handler.set_arg(0, Data);
100+
Handler.set_arg(1, NumOfElements);
101+
Handler.set_arg(2, Id);
102+
Handler.single_task(UsedKernel);
103+
})
104+
.wait();
105+
106+
Failed += performResultCheck(NumOfElements, Data,
107+
"globalScopeSingleFreeFuncKernel",
108+
ExpectedResultValue);
109+
}
110+
111+
{
112+
std::fill(Data, Data + NumOfElements, 0);
113+
sycl::kernel UsedKernel =
114+
getKernel<globalScopeNdRangeFreeFuncKernel>(Context);
115+
sycl::id<2> Id{44, 44};
116+
int ExpectedResultValue = static_cast<int>(Id[0] + Id[1]);
117+
Queue
118+
.submit([&](sycl::handler &Handler) {
119+
Handler.set_args(Data, Id);
120+
sycl::nd_range<2> Ndr{{8, NumOfElements / 8}, {8, 8}};
121+
Handler.parallel_for(Ndr, UsedKernel);
122+
})
123+
.wait();
124+
125+
Failed += performResultCheck(NumOfElements, Data,
126+
"globalScopeNdRangeFreeFuncKernel",
127+
ExpectedResultValue);
128+
}
129+
130+
sycl::free(Data, Queue);
131+
return Failed;
132+
}

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

Lines changed: 147 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,147 @@
1+
// REQUIRES: aspect-usm_shared_allocations
2+
// RUN: %{build} -o %t.out
3+
// RUN: %{run} %t.out
4+
5+
// This test verifies whether a marray<T, NumElements> can be passed as a kernel
6+
// parameter to a free function kernel.
7+
8+
#include <sycl/usm.hpp>
9+
10+
#include <sycl/ext/oneapi/free_function_queries.hpp>
11+
12+
#include "helpers.hpp"
13+
14+
static constexpr size_t M_ARRAY_SIZE = 5;
15+
16+
static float sumMArray(sycl::marray<float, M_ARRAY_SIZE> MArray) {
17+
float SumOfMArray = 0.0;
18+
for (const auto Value : MArray) {
19+
SumOfMArray += Value;
20+
}
21+
return SumOfMArray;
22+
}
23+
24+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel))
25+
void globalScopeSingleFreeFuncKernel(float *Ptr, size_t NumOfElements,
26+
sycl::marray<float, M_ARRAY_SIZE> MArray) {
27+
for (size_t i = 0; i < NumOfElements; ++i) {
28+
Ptr[i] = sumMArray(MArray);
29+
}
30+
}
31+
32+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<2>))
33+
void globalScopeNdRangeFreeFuncKernel(
34+
float *Ptr, sycl::marray<float, M_ARRAY_SIZE> MArray) {
35+
size_t Item =
36+
syclext::this_work_item::get_nd_item<2>().get_global_linear_id();
37+
Ptr[Item] = sumMArray(MArray);
38+
}
39+
40+
namespace ns {
41+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel))
42+
void nsSingleFreeFuncKernel(float *Ptr, size_t NumOfElements,
43+
sycl::marray<float, M_ARRAY_SIZE> MArray) {
44+
for (size_t i = 0; i < NumOfElements; ++i) {
45+
Ptr[i] = sumMArray(MArray);
46+
}
47+
}
48+
49+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<3>))
50+
void nsNdRangeFreeFuncKernel(float *Ptr,
51+
sycl::marray<float, M_ARRAY_SIZE> MArray) {
52+
size_t Item =
53+
syclext::this_work_item::get_nd_item<3>().get_global_linear_id();
54+
Ptr[Item] = sumMArray(MArray);
55+
}
56+
} // namespace ns
57+
58+
// TODO: Need to add checks for a static member functions of a class as free
59+
// function kerenl
60+
61+
int main() {
62+
int Failed = 0;
63+
sycl::queue Queue;
64+
sycl::context Context = Queue.get_context();
65+
constexpr size_t NumOfElements = 1024;
66+
float *Data = sycl::malloc_shared<float>(NumOfElements, Queue);
67+
68+
{
69+
std::fill(Data, Data + NumOfElements, 0);
70+
sycl::kernel UsedKernel = getKernel<ns::nsSingleFreeFuncKernel>(Context);
71+
72+
sycl::marray<float, M_ARRAY_SIZE> MArray{1.0, 2.0, 3.0, 4.0, 5.0};
73+
float ExpectedResultValue = sumMArray(MArray);
74+
Queue
75+
.submit([&](sycl::handler &Handler) {
76+
Handler.set_args(Data, NumOfElements, MArray);
77+
Handler.single_task(UsedKernel);
78+
})
79+
.wait();
80+
81+
Failed += performResultCheck(
82+
NumOfElements, Data, "ns::nsSingleFreeFuncKernel", ExpectedResultValue);
83+
}
84+
85+
{
86+
std::fill(Data, Data + NumOfElements, 0);
87+
sycl::kernel UsedKernel = getKernel<ns::nsNdRangeFreeFuncKernel>(Context);
88+
89+
sycl::marray<float, M_ARRAY_SIZE> MArray{100.0, 100.0, 100.0, 100.0, 100.0};
90+
float ExpectedResultValue = sumMArray(MArray);
91+
Queue
92+
.submit([&](sycl::handler &Handler) {
93+
Handler.set_arg(0, Data);
94+
Handler.set_arg(1, MArray);
95+
sycl::nd_range<3> Ndr{{4, 4, NumOfElements / 16}, {4, 4, 4}};
96+
Handler.parallel_for(Ndr, UsedKernel);
97+
})
98+
.wait();
99+
100+
Failed +=
101+
performResultCheck(NumOfElements, Data, "ns::nsNdRangeFreeFuncKernel",
102+
ExpectedResultValue);
103+
}
104+
105+
{
106+
std::fill(Data, Data + NumOfElements, 0);
107+
sycl::kernel UsedKernel =
108+
getKernel<globalScopeSingleFreeFuncKernel>(Context);
109+
sycl::marray<float, M_ARRAY_SIZE> MArray{500.0, 500.0, 500.0, 500.0, 500.0};
110+
float ExpectedResultValue = sumMArray(MArray);
111+
Queue
112+
.submit([&](sycl::handler &Handler) {
113+
Handler.set_arg(0, Data);
114+
Handler.set_arg(1, NumOfElements);
115+
Handler.set_arg(2, MArray);
116+
Handler.single_task(UsedKernel);
117+
})
118+
.wait();
119+
120+
Failed += performResultCheck(NumOfElements, Data,
121+
"globalScopeSingleFreeFuncKernel",
122+
ExpectedResultValue);
123+
}
124+
125+
{
126+
std::fill(Data, Data + NumOfElements, 0);
127+
sycl::kernel UsedKernel =
128+
getKernel<globalScopeNdRangeFreeFuncKernel>(Context);
129+
sycl::marray<float, M_ARRAY_SIZE> MArray{1000.0, 1000.0, 1000.0, 1000.0,
130+
1000.0};
131+
float ExpectedResultValue = sumMArray(MArray);
132+
Queue
133+
.submit([&](sycl::handler &Handler) {
134+
Handler.set_args(Data, MArray);
135+
sycl::nd_range<2> Ndr{{8, NumOfElements / 8}, {8, 8}};
136+
Handler.parallel_for(Ndr, UsedKernel);
137+
})
138+
.wait();
139+
140+
Failed += performResultCheck(NumOfElements, Data,
141+
"globalScopeNdRangeFreeFuncKernel",
142+
ExpectedResultValue);
143+
}
144+
145+
sycl::free(Data, Queue);
146+
return Failed;
147+
}

‎sycl/test-e2e/Experimental/free_functions/namespace.cpp renamed to ‎sycl/test-e2e/FreeFunctionKernels/namespace.cpp

File renamed without changes.

0 commit comments

Comments
 (0)