diff --git a/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp b/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp new file mode 100644 index 0000000000000..260c7d9b0203c --- /dev/null +++ b/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp @@ -0,0 +1,189 @@ +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// This test verifies whether sycl::accessor can be used with free function +// kernels extension. + +#include + +#include "helpers.hpp" + +template +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel)) +void globalScopeSingleFreeFunc( + sycl::accessor + Accessor, + int Value) { + for (auto &Elem : Accessor) + Elem = Value; +} +namespace ns { +template +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel)) +void nsNdRangeFreeFunc(sycl::accessor + Accessor, + int Value) { + auto Item = syclext::this_work_item::get_nd_item().get_global_id(); + Accessor[Item] = Value; +} +} // namespace ns + +template +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel)) +void ndRangeFreeFuncMultipleParameters( + sycl::accessor + InputAAcc, + sycl::accessor + InputBAcc, + sycl::accessor + ResultAcc) { + auto Item = syclext::this_work_item::get_nd_item().get_global_id(); + ResultAcc[Item] = InputAAcc[Item] + InputBAcc[Item]; +} + +// TODO: Need to add checks for a static member functions of a class as free +// function kernel. + +template +int runSingleTaskTest(sycl::queue &Queue, sycl::context &Context, + sycl::range NumOfElementsPerDim, + std::string_view ErrorMessage, + const int ExpectedResultValue) { + sycl::kernel UsedKernel = getKernel(Context); + std::vector ResultData(NumOfElementsPerDim.size(), 0); + { + sycl::buffer Buffer(ResultData.data(), NumOfElementsPerDim); + Queue.submit([&](sycl::handler &Handler) { + sycl::accessor Accessor{Buffer, Handler}; + Handler.set_args(Accessor, ExpectedResultValue); + Handler.single_task(UsedKernel); + }); + } + return performResultCheck(NumOfElementsPerDim.size(), ResultData.data(), + ErrorMessage, ExpectedResultValue); +} + +template +int runNdRangeTest(sycl::queue &Queue, sycl::context &Context, + sycl::nd_range NdRange, std::string_view ErrorMessage, + const int ExpectedResultValue) { + sycl::kernel UsedKernel = getKernel(Context); + std::vector ResultData(NdRange.get_global_range().size(), 0); + { + sycl::buffer Buffer(ResultData.data(), + NdRange.get_global_range()); + Queue.submit([&](sycl::handler &Handler) { + sycl::accessor Accessor{Buffer, Handler}; + Handler.set_args(Accessor, ExpectedResultValue); + Handler.parallel_for(NdRange, UsedKernel); + }); + } + return performResultCheck(NdRange.get_global_range().size(), + ResultData.data(), ErrorMessage, + ExpectedResultValue); +} + +template +int runNdRangeTestMultipleParameters(sycl::queue &Queue, sycl::context &Context, + sycl::nd_range NdRange, + std::string_view ErrorMessage, + sycl::range<3> Values) { + sycl::kernel UsedKernel = getKernel(Context); + std::vector InputAData(NdRange.get_global_range().size(), Values[0]); + std::vector InputBData(NdRange.get_global_range().size(), Values[1]); + std::vector ResultData(NdRange.get_global_range().size(), 0); + + { + sycl::buffer InputABuffer(InputAData.data(), + NdRange.get_global_range()); + sycl::buffer InputBBuffer(InputBData.data(), + NdRange.get_global_range()); + sycl::buffer ResultBuffer(ResultData.data(), + NdRange.get_global_range()); + Queue.submit([&](sycl::handler &Handler) { + sycl::accessor + InputAAcc{InputABuffer, Handler}; + sycl::accessor + InputBAcc{InputBBuffer, Handler}; + sycl::accessor ResultAcc{ + ResultBuffer, Handler}; + Handler.set_args(InputAAcc, InputBAcc, ResultAcc); + Handler.parallel_for(NdRange, UsedKernel); + }); + } + return performResultCheck(NdRange.get_global_range().size(), + ResultData.data(), ErrorMessage, Values[2]); +} + +int main() { + + int Failed = 0; + sycl::queue Queue; + sycl::context Context = Queue.get_context(); + { + // Check that sycl::accessor is supported inside single_task free function + // kernel + Failed += runSingleTaskTest, 1>( + Queue, Context, sycl::range<1>{10}, + "globalScopeSingleFreeFunc with sycl::accessor<1>", 1); + Failed += runSingleTaskTest, 2>( + Queue, Context, sycl::range<2>{10, 10}, + "globalScopeSingleFreeFunc with sycl::accessor<2>", 2); + Failed += runSingleTaskTest, 3>( + Queue, Context, sycl::range<3>{5, 5, 5}, + "globalScopeSingleFreeFunc with sycl::accessor<3>", 3); + } + + { + // Check that sycl::accessor is supported inside nd_range free function + // kernel + Failed += runNdRangeTest, 1>( + Queue, Context, sycl::nd_range{sycl::range{10}, sycl::range{2}}, + "ns::nsNdRangeFreeFunc with sycl::accessor<1>", 4); + Failed += runNdRangeTest, 2>( + Queue, Context, sycl::nd_range{sycl::range{16, 16}, sycl::range{4, 4}}, + "ns::nsNdRangeFreeFunc with sycl::accessor<2>", 5); + Failed += runNdRangeTest, 3>( + Queue, Context, + sycl::nd_range{sycl::range{10, 10, 10}, sycl::range{2, 2, 2}}, + "ns::nsNdRangeFreeFunc with sycl::accessor<3>", 6); + } + + { + // Check that multiple sycl::accessor are supported inside nd_range free + // function kernel + Failed += + runNdRangeTestMultipleParameters, + 1>( + Queue, Context, sycl::nd_range{sycl::range{10}, sycl::range{2}}, + "ndRangeFreeFuncMultipleParameters with multiple sycl::accessor<1>", + sycl::range{111, 111, 222}); + Failed += + runNdRangeTestMultipleParameters, + 2>( + Queue, Context, + sycl::nd_range{sycl::range{16, 16}, sycl::range{4, 4}}, + "ndRangeFreeFuncMultipleParameters with multiple sycl::accessor<2>", + sycl::range{222, 222, 444}); + Failed += + runNdRangeTestMultipleParameters, + 3>( + Queue, Context, + sycl::nd_range{sycl::range{10, 10, 10}, sycl::range{2, 2, 2}}, + "ndRangeFreeFuncMultipleParameters with multiple sycl::accessor<3>", + sycl::range{444, 444, 888}); + } + return Failed; +} diff --git a/sycl/test-e2e/FreeFunctionKernels/helpers.hpp b/sycl/test-e2e/FreeFunctionKernels/helpers.hpp index e2bd18e6ffab4..bb561589f9711 100644 --- a/sycl/test-e2e/FreeFunctionKernels/helpers.hpp +++ b/sycl/test-e2e/FreeFunctionKernels/helpers.hpp @@ -6,19 +6,19 @@ namespace syclext = sycl::ext::oneapi; namespace syclexp = sycl::ext::oneapi::experimental; -template +template static int performResultCheck(size_t NumberOfElements, const T *ResultPtr, std::string_view TestName, - T ExpectedResultValue) { - int IsSuccessful{0}; + S ExpectedResultValue) { + int Failed{0}; for (size_t i = 0; i < NumberOfElements; i++) { if (ResultPtr[i] != ExpectedResultValue) { std::cerr << "Failed " << TestName << " : " << ResultPtr[i] << " != " << ExpectedResultValue << std::endl; - ++IsSuccessful; + ++Failed; } } - return IsSuccessful; + return Failed; } template static sycl::kernel getKernel(sycl::context &Context) { diff --git a/sycl/test-e2e/FreeFunctionKernels/local_accessor_as_kernel_parameter.cpp b/sycl/test-e2e/FreeFunctionKernels/local_accessor_as_kernel_parameter.cpp new file mode 100644 index 0000000000000..74167ea08917a --- /dev/null +++ b/sycl/test-e2e/FreeFunctionKernels/local_accessor_as_kernel_parameter.cpp @@ -0,0 +1,94 @@ +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// This test verifies whether sycl::local_accessor can be used with free +// function kernels extension. + +#include +#include +#include + +#include "helpers.hpp" + +constexpr size_t BIN_SIZE = 4; +constexpr size_t NUM_BINS = 4; +constexpr size_t INPUT_SIZE = 1024; + +namespace ns { +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void nsNdRangeFreeFunc(sycl::accessor InputAccessor, + sycl::accessor ResultAccessor, + sycl::local_accessor LocalAccessor) { + + size_t LocalWorkItemId = + syclext::this_work_item::get_nd_item<1>().get_local_id(); + size_t GlobalWorkItemId = + syclext::this_work_item::get_nd_item<1>().get_global_id(); + sycl::group<1> WorkGroup = syclext::this_work_item::get_work_group<1>(); + + if (LocalWorkItemId < BIN_SIZE) + LocalAccessor[LocalWorkItemId] = 0; + + sycl::group_barrier(WorkGroup); + + int Value = InputAccessor[GlobalWorkItemId]; + sycl::atomic_ref + AtomicRefLocal(LocalAccessor[Value]); + AtomicRefLocal++; + sycl::group_barrier(WorkGroup); + + if (LocalWorkItemId < BIN_SIZE) { + sycl::atomic_ref + AtomicRefGlobal(ResultAccessor[LocalWorkItemId]); + AtomicRefGlobal.fetch_add(LocalAccessor[LocalWorkItemId]); + } +} +} // namespace ns + +// TODO: Need to add checks for a static member functions of a class as free +// function kerenl + +void FillWithData(std::vector &Data, std::vector &Values) { + constexpr size_t Offset = INPUT_SIZE / NUM_BINS; + for (size_t i = 0; i < NUM_BINS; ++i) { + std::fill(Data.begin() + i * Offset, Data.begin() + (i + 1) * Offset, + Values[i]); + } +} + +int main() { + + int Failed = 0; + sycl::queue Queue; + sycl::context Context = Queue.get_context(); + { + // Check that sycl::local_accesor is supported inside nd_range free function + // kernel. + std::vector ExpectedHistogramNumbers = {0, 1, 2, 3}; + std::vector ResultData(BIN_SIZE, 0); + + std::vector InputData(INPUT_SIZE); + FillWithData(InputData, ExpectedHistogramNumbers); + { + sycl::buffer InputBuffer(InputData); + sycl::buffer ResultBuffer(ResultData); + sycl::kernel UsedKernel = getKernel(Context); + Queue.submit([&](sycl::handler &Handler) { + sycl::accessor InputAccessor{InputBuffer, Handler}; + sycl::accessor ResultsAccessor{ResultBuffer, Handler}; + sycl::local_accessor LocalMemPerWG(sycl::range<1>(BIN_SIZE), + Handler); + Handler.set_args(InputAccessor, ResultsAccessor, LocalMemPerWG); + sycl::nd_range<1> Ndr{INPUT_SIZE, INPUT_SIZE / NUM_BINS}; + Handler.parallel_for(Ndr, UsedKernel); + }); + } + Failed += + performResultCheck(NUM_BINS, ResultData.data(), + "sycl::nd_range_kernel with sycl::local_accessor", + INPUT_SIZE / NUM_BINS); + } + return Failed; +}