From 605a7005bc2a879c164c83282016e0036b07f8eb Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Mon, 26 May 2025 17:07:13 +0200 Subject: [PATCH 01/10] added local_accessor and accessor test for free_func_kernels extension --- .../accessor_as_kernel_parameter.cpp | 80 ++++++++++++++++ sycl/test-e2e/FreeFunctionKernels/helpers.hpp | 4 +- .../local_accessor_as_kernel_parameter.cpp | 93 +++++++++++++++++++ 3 files changed, 175 insertions(+), 2 deletions(-) create mode 100644 sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp create mode 100644 sycl/test-e2e/FreeFunctionKernels/local_accessor_as_kernel_parameter.cpp 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..941a3c0edeefe --- /dev/null +++ b/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp @@ -0,0 +1,80 @@ +// 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" + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel)) +void globalScopeSingleFreeFunc(sycl::accessor Accessor, + size_t NumOfElements, int Value) { + for (size_t i = 0; i < NumOfElements; ++i) { + Accessor[i] = Value; + } +} + +namespace ns { +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<3>)) +void nsNdRangeFreeFunc(sycl::accessor Accessor, int Value) { + size_t Item = + syclext::this_work_item::get_nd_item<3>().get_global_linear_id(); + Accessor[Item] = Value; +} +} // namespace ns + +// TODO: Need to add checks for a static member functions of a class as free +// function kerenl + +int main() { + + int Failed = 0; + sycl::queue Queue; + sycl::context Context = Queue.get_context(); + constexpr size_t NumOfElements = 1024; + { + // Check that sycl::accessor is supported inside nd_range free function + // kernel. + std::vector ResultHostData(NumOfElements, 0); + constexpr int ExpectedResultValue = 111; + { + sycl::buffer Buffer(ResultHostData); + sycl::kernel UsedKernel = getKernel(Context); + Queue.submit([&](sycl::handler &Handler) { + sycl::accessor Accessor{Buffer, Handler}; + Handler.set_args(Accessor, ExpectedResultValue); + sycl::nd_range<3> Ndr{{4, 4, NumOfElements / 16}, {4, 4, 4}}; + Handler.parallel_for(Ndr, UsedKernel); + }); + } + + Failed += performResultCheck(NumOfElements, ResultHostData.data(), + "ns::nsNdRangeFreeFunc with sycl::accessor", + ExpectedResultValue); + } + + { + // Check that sycl::accessor is supported inside single_task free function + // kernel. + std::vector ResultHostData(NumOfElements, 0); + constexpr int ExpectedResultValue = 222; + { + sycl::buffer Buffer(ResultHostData); + sycl::kernel UsedKernel = getKernel(Context); + Queue.submit([&](sycl::handler &Handler) { + sycl::accessor Accessor{Buffer, Handler}; + Handler.set_arg(0, Accessor); + Handler.set_arg(1, NumOfElements); + Handler.set_arg(2, ExpectedResultValue); + Handler.single_task(UsedKernel); + }); + } + Failed += performResultCheck( + NumOfElements, ResultHostData.data(), + "globalScopeSingleFreeFunc with sycl::accessor", ExpectedResultValue); + } + + return Failed; +} diff --git a/sycl/test-e2e/FreeFunctionKernels/helpers.hpp b/sycl/test-e2e/FreeFunctionKernels/helpers.hpp index e2bd18e6ffab4..9c97c9e2a798e 100644 --- a/sycl/test-e2e/FreeFunctionKernels/helpers.hpp +++ b/sycl/test-e2e/FreeFunctionKernels/helpers.hpp @@ -6,10 +6,10 @@ 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) { + S ExpectedResultValue) { int IsSuccessful{0}; for (size_t i = 0; i < NumberOfElements; i++) { if (ResultPtr[i] != ExpectedResultValue) { 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..e9b410486e35e --- /dev/null +++ b/sycl/test-e2e/FreeFunctionKernels/local_accessor_as_kernel_parameter.cpp @@ -0,0 +1,93 @@ +// 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; +} From f3e6a0e693633daa043e3d011a191c794bb15012 Mon Sep 17 00:00:00 2001 From: Daniel Skrobot Date: Mon, 2 Jun 2025 14:01:48 +0200 Subject: [PATCH 02/10] Update sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp Co-authored-by: Steffen Larsen --- .../FreeFunctionKernels/accessor_as_kernel_parameter.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp b/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp index 941a3c0edeefe..a694b717e4117 100644 --- a/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp +++ b/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp @@ -11,8 +11,8 @@ SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel)) void globalScopeSingleFreeFunc(sycl::accessor Accessor, size_t NumOfElements, int Value) { - for (size_t i = 0; i < NumOfElements; ++i) { - Accessor[i] = Value; + for (size_t I = 0; I < NumOfElements; ++I) { + Accessor[I] = Value; } } From 6028b4513205ef9ece009604ee8513abdfef87df Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Mon, 2 Jun 2025 14:07:23 +0200 Subject: [PATCH 03/10] added 2 and 3 dimensional checks for accessor --- .../accessor_as_kernel_parameter.cpp | 64 +++++++++++++++++++ sycl/test-e2e/FreeFunctionKernels/helpers.hpp | 6 +- 2 files changed, 67 insertions(+), 3 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp b/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp index a694b717e4117..468807f925edd 100644 --- a/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp +++ b/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp @@ -23,8 +23,24 @@ void nsNdRangeFreeFunc(sycl::accessor Accessor, int Value) { syclext::this_work_item::get_nd_item<3>().get_global_linear_id(); Accessor[Item] = Value; } + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<3>)) +void nsNdRangeFreeFuncWith3DimAccessor(sycl::accessor Accessor, + int Value) { + sycl::nd_item<3> NdItem = syclext::this_work_item::get_nd_item<3>(); + Accessor[NdItem.get_global_id(0)][NdItem.get_global_id(1)] + [NdItem.get_global_id(2)] = Value; +} + } // namespace ns +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<2>)) +void globalndRangeFreeFuncWith2DimAccessor(sycl::accessor Accessor, + int Value) { + sycl::nd_item<2> NdItem = syclext::this_work_item::get_nd_item<2>(); + Accessor[NdItem.get_group_linear_id()][NdItem.get_local_linear_id()] = Value; +} + // TODO: Need to add checks for a static member functions of a class as free // function kerenl @@ -76,5 +92,53 @@ int main() { "globalScopeSingleFreeFunc with sycl::accessor", ExpectedResultValue); } + { + // Check that sycl::accessor<2> is supported inside single_task free + // function kernel. + std::vector ResultHostData(NumOfElements, 0); + constexpr int ExpectedResultValue = 333; + { + sycl::range<2> BufRange{8, NumOfElements / 8}; + sycl::buffer Buffer(ResultHostData.data(), BufRange); + sycl::kernel UsedKernel = + getKernel(Context); + Queue.submit([&](sycl::handler &Handler) { + sycl::accessor Accessor{Buffer, Handler}; + Handler.set_arg(0, Accessor); + Handler.set_arg(1, ExpectedResultValue); + sycl::nd_range<2> Ndr{{128, 8}, {16, 8}}; + Handler.parallel_for(Ndr, UsedKernel); + }); + } + Failed += performResultCheck( + NumOfElements, ResultHostData.data(), + "globalndRangeFreeFuncWith2DimAccessor with sycl::accessor<2>", + ExpectedResultValue); + } + + { + // Check that sycl::accessor<3> is supported inside single_task free + // function kernel. + std::vector ResultHostData(NumOfElements, 0); + constexpr int ExpectedResultValue = 444; + { + sycl::range<3> BufRange{64, 4, 4}; + sycl::buffer Buffer(ResultHostData.data(), BufRange); + sycl::kernel UsedKernel = + getKernel(Context); + Queue.submit([&](sycl::handler &Handler) { + sycl::accessor Accessor{Buffer, Handler}; + Handler.set_arg(0, Accessor); + Handler.set_arg(1, ExpectedResultValue); + sycl::nd_range<3> Ndr{{64, 4, 4}, {16, 4, 2}}; + Handler.parallel_for(Ndr, UsedKernel); + }); + } + Failed += performResultCheck( + NumOfElements, ResultHostData.data(), + "ns::nsNdRangeFreeFuncWith3DimAccessor with sycl::accessor<3>", + ExpectedResultValue); + } + return Failed; } diff --git a/sycl/test-e2e/FreeFunctionKernels/helpers.hpp b/sycl/test-e2e/FreeFunctionKernels/helpers.hpp index 9c97c9e2a798e..bb561589f9711 100644 --- a/sycl/test-e2e/FreeFunctionKernels/helpers.hpp +++ b/sycl/test-e2e/FreeFunctionKernels/helpers.hpp @@ -10,15 +10,15 @@ template static int performResultCheck(size_t NumberOfElements, const T *ResultPtr, std::string_view TestName, S ExpectedResultValue) { - int IsSuccessful{0}; + 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) { From 86c8a659f68a0931f4e526ed184dd9e81491ff15 Mon Sep 17 00:00:00 2001 From: Daniel Skrobot Date: Mon, 2 Jun 2025 14:58:24 +0200 Subject: [PATCH 04/10] Update sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp Co-authored-by: Steffen Larsen --- .../FreeFunctionKernels/accessor_as_kernel_parameter.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp b/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp index 468807f925edd..c0c2812a55429 100644 --- a/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp +++ b/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp @@ -42,7 +42,7 @@ void globalndRangeFreeFuncWith2DimAccessor(sycl::accessor Accessor, } // TODO: Need to add checks for a static member functions of a class as free -// function kerenl +// function kernel. int main() { From 28ad0e829cbea187a7a41915a87c6c022d8d2b59 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Wed, 4 Jun 2025 15:55:24 +0200 Subject: [PATCH 05/10] reworked to support multiple dims for free function kernels --- .../accessor_as_kernel_parameter.cpp | 270 +++++++++++------- 1 file changed, 165 insertions(+), 105 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp b/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp index c0c2812a55429..1e321c4286c0d 100644 --- a/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp +++ b/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp @@ -8,137 +8,197 @@ #include "helpers.hpp" +template SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel)) -void globalScopeSingleFreeFunc(sycl::accessor Accessor, - size_t NumOfElements, int Value) { - for (size_t I = 0; I < NumOfElements; ++I) { - Accessor[I] = Value; +void globalScopeSingleFreeFunc( + sycl::accessor + Accessor, + sycl::range NumOfElements, int Value) { + if constexpr (Dims == 1) { + for (size_t I = 0; I < NumOfElements[0]; ++I) + Accessor[I] = Value; + } else if constexpr (Dims == 2) { + for (size_t I = 0; I < NumOfElements[0]; ++I) { + for (size_t J = 0; J < NumOfElements[1]; ++J) { + Accessor[I][J] = Value; + } + } + } else if constexpr (Dims == 3) { + for (size_t I = 0; I < NumOfElements[0]; ++I) { + for (size_t J = 0; J < NumOfElements[1]; ++J) { + for (size_t K = 0; K < NumOfElements[2]; ++K) { + Accessor[I][J][K] = Value; + } + } + } } } - namespace ns { -SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<3>)) -void nsNdRangeFreeFunc(sycl::accessor Accessor, int Value) { - size_t Item = - syclext::this_work_item::get_nd_item<3>().get_global_linear_id(); +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; } - -SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<3>)) -void nsNdRangeFreeFuncWith3DimAccessor(sycl::accessor Accessor, - int Value) { - sycl::nd_item<3> NdItem = syclext::this_work_item::get_nd_item<3>(); - Accessor[NdItem.get_global_id(0)][NdItem.get_global_id(1)] - [NdItem.get_global_id(2)] = Value; -} - } // namespace ns -SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<2>)) -void globalndRangeFreeFuncWith2DimAccessor(sycl::accessor Accessor, - int Value) { - sycl::nd_item<2> NdItem = syclext::this_work_item::get_nd_item<2>(); - Accessor[NdItem.get_group_linear_id()][NdItem.get_local_linear_id()] = Value; +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, NumOfElementsPerDim, 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(); - constexpr size_t NumOfElements = 1024; - { - // Check that sycl::accessor is supported inside nd_range free function - // kernel. - std::vector ResultHostData(NumOfElements, 0); - constexpr int ExpectedResultValue = 111; - { - sycl::buffer Buffer(ResultHostData); - sycl::kernel UsedKernel = getKernel(Context); - Queue.submit([&](sycl::handler &Handler) { - sycl::accessor Accessor{Buffer, Handler}; - Handler.set_args(Accessor, ExpectedResultValue); - sycl::nd_range<3> Ndr{{4, 4, NumOfElements / 16}, {4, 4, 4}}; - Handler.parallel_for(Ndr, UsedKernel); - }); - } - - Failed += performResultCheck(NumOfElements, ResultHostData.data(), - "ns::nsNdRangeFreeFunc with sycl::accessor", - ExpectedResultValue); - } - { // Check that sycl::accessor is supported inside single_task free function - // kernel. - std::vector ResultHostData(NumOfElements, 0); - constexpr int ExpectedResultValue = 222; - { - sycl::buffer Buffer(ResultHostData); - sycl::kernel UsedKernel = getKernel(Context); - Queue.submit([&](sycl::handler &Handler) { - sycl::accessor Accessor{Buffer, Handler}; - Handler.set_arg(0, Accessor); - Handler.set_arg(1, NumOfElements); - Handler.set_arg(2, ExpectedResultValue); - Handler.single_task(UsedKernel); - }); - } - Failed += performResultCheck( - NumOfElements, ResultHostData.data(), - "globalScopeSingleFreeFunc with sycl::accessor", ExpectedResultValue); + // 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>{10, 10, 10}, + "globalScopeSingleFreeFunc with sycl::accessor<3>", 3); } { - // Check that sycl::accessor<2> is supported inside single_task free - // function kernel. - std::vector ResultHostData(NumOfElements, 0); - constexpr int ExpectedResultValue = 333; - { - sycl::range<2> BufRange{8, NumOfElements / 8}; - sycl::buffer Buffer(ResultHostData.data(), BufRange); - sycl::kernel UsedKernel = - getKernel(Context); - Queue.submit([&](sycl::handler &Handler) { - sycl::accessor Accessor{Buffer, Handler}; - Handler.set_arg(0, Accessor); - Handler.set_arg(1, ExpectedResultValue); - sycl::nd_range<2> Ndr{{128, 8}, {16, 8}}; - Handler.parallel_for(Ndr, UsedKernel); - }); - } - Failed += performResultCheck( - NumOfElements, ResultHostData.data(), - "globalndRangeFreeFuncWith2DimAccessor with sycl::accessor<2>", - ExpectedResultValue); + // 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{10}}, + "ns::nsNdRangeFreeFunc with sycl::accessor<1>", 4); + Failed += runNdRangeTest, 2>( + Queue, Context, + sycl::nd_range{sycl::range{10, 10}, sycl::range{10, 10}}, + "ns::nsNdRangeFreeFunc with sycl::accessor<2>", 5); + Failed += runNdRangeTest, 3>( + Queue, Context, + sycl::nd_range{sycl::range{10, 10, 10}, sycl::range{10, 10, 10}}, + "ns::nsNdRangeFreeFunc with sycl::accessor<3>", 5); } { - // Check that sycl::accessor<3> is supported inside single_task free - // function kernel. - std::vector ResultHostData(NumOfElements, 0); - constexpr int ExpectedResultValue = 444; - { - sycl::range<3> BufRange{64, 4, 4}; - sycl::buffer Buffer(ResultHostData.data(), BufRange); - sycl::kernel UsedKernel = - getKernel(Context); - Queue.submit([&](sycl::handler &Handler) { - sycl::accessor Accessor{Buffer, Handler}; - Handler.set_arg(0, Accessor); - Handler.set_arg(1, ExpectedResultValue); - sycl::nd_range<3> Ndr{{64, 4, 4}, {16, 4, 2}}; - Handler.parallel_for(Ndr, UsedKernel); - }); - } - Failed += performResultCheck( - NumOfElements, ResultHostData.data(), - "ns::nsNdRangeFreeFuncWith3DimAccessor with sycl::accessor<3>", - ExpectedResultValue); + // 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{10}}, + "ndRangeFreeFuncMultipleParameters with multiple sycl::accessor<1>", + sycl::range{111, 111, 222}); + Failed += + runNdRangeTestMultipleParameters, + 2>( + Queue, Context, + sycl::nd_range{sycl::range{10, 10}, sycl::range{10, 10}}, + "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{10, 10, 10}}, + "ndRangeFreeFuncMultipleParameters with multiple sycl::accessor<3>", + sycl::range{444, 444, 888}); } - return Failed; } From eb0388761c08c43f08592100ae51f2ae3bca4211 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Wed, 4 Jun 2025 17:05:16 +0200 Subject: [PATCH 06/10] fix --- .../accessor_as_kernel_parameter.cpp | 24 ++++--------------- 1 file changed, 4 insertions(+), 20 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp b/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp index 1e321c4286c0d..cf83c744fb457 100644 --- a/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp +++ b/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp @@ -15,25 +15,9 @@ void globalScopeSingleFreeFunc( sycl::access::target::device, sycl::access::placeholder::false_t> Accessor, - sycl::range NumOfElements, int Value) { - if constexpr (Dims == 1) { - for (size_t I = 0; I < NumOfElements[0]; ++I) - Accessor[I] = Value; - } else if constexpr (Dims == 2) { - for (size_t I = 0; I < NumOfElements[0]; ++I) { - for (size_t J = 0; J < NumOfElements[1]; ++J) { - Accessor[I][J] = Value; - } - } - } else if constexpr (Dims == 3) { - for (size_t I = 0; I < NumOfElements[0]; ++I) { - for (size_t J = 0; J < NumOfElements[1]; ++J) { - for (size_t K = 0; K < NumOfElements[2]; ++K) { - Accessor[I][J][K] = Value; - } - } - } - } + int Value) { + for (auto &Elem : Accessor) + Elem = Value; } namespace ns { template @@ -81,7 +65,7 @@ int runSingleTaskTest(sycl::queue &Queue, sycl::context &Context, sycl::buffer Buffer(ResultData.data(), NumOfElementsPerDim); Queue.submit([&](sycl::handler &Handler) { sycl::accessor Accessor{Buffer, Handler}; - Handler.set_args(Accessor, NumOfElementsPerDim, ExpectedResultValue); + Handler.set_args(Accessor, ExpectedResultValue); Handler.single_task(UsedKernel); }); } From 70cf4b9cc436bbeeb408c335e65973d0dfb605b7 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Wed, 4 Jun 2025 17:58:33 +0200 Subject: [PATCH 07/10] fix --- .../accessor_as_kernel_parameter.cpp | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp b/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp index cf83c744fb457..e82816d1025b0 100644 --- a/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp +++ b/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp @@ -140,7 +140,7 @@ int main() { Queue, Context, sycl::range<2>{10, 10}, "globalScopeSingleFreeFunc with sycl::accessor<2>", 2); Failed += runSingleTaskTest, 3>( - Queue, Context, sycl::range<3>{10, 10, 10}, + Queue, Context, sycl::range<3>{5, 5, 5}, "globalScopeSingleFreeFunc with sycl::accessor<3>", 3); } @@ -148,16 +148,16 @@ int main() { // 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{10}}, + 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{10, 10}, sycl::range{10, 10}}, + 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{10, 10, 10}}, - "ns::nsNdRangeFreeFunc with sycl::accessor<3>", 5); + sycl::nd_range{sycl::range{10, 10, 10}, sycl::range{2, 2, 2}}, + "ns::nsNdRangeFreeFunc with sycl::accessor<3>", 6); } { @@ -166,21 +166,21 @@ int main() { Failed += runNdRangeTestMultipleParameters, 1>( - Queue, Context, sycl::nd_range{sycl::range{10}, sycl::range{10}}, + 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{10, 10}, sycl::range{10, 10}}, + sycl::nd_range{sycl::range{10, 10}, 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{10, 10, 10}}, + sycl::nd_range{sycl::range{10, 10, 10}, sycl::range{2, 2, 2}}, "ndRangeFreeFuncMultipleParameters with multiple sycl::accessor<3>", sycl::range{444, 444, 888}); } From 7d852fbdb49fcec890b0950fc91224a0c5e7377f Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Wed, 4 Jun 2025 19:46:17 +0200 Subject: [PATCH 08/10] fixed formatting --- .../FreeFunctionKernels/accessor_as_kernel_parameter.cpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp b/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp index e82816d1025b0..3041ce0d11199 100644 --- a/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp +++ b/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp @@ -151,8 +151,7 @@ int main() { 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}}, + 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, @@ -173,7 +172,7 @@ int main() { runNdRangeTestMultipleParameters, 2>( Queue, Context, - sycl::nd_range{sycl::range{10, 10}, sycl::range{4, 4}}, + sycl::nd_range{sycl::range{16, 16}, sycl::range{4, 4}}, "ndRangeFreeFuncMultipleParameters with multiple sycl::accessor<2>", sycl::range{222, 222, 444}); Failed += From 6821b904d71ad8d40273ac1406a041af38774324 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Thu, 5 Jun 2025 10:38:15 +0200 Subject: [PATCH 09/10] updated spacing between functions --- .../FreeFunctionKernels/accessor_as_kernel_parameter.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp b/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp index 3041ce0d11199..260c7d9b0203c 100644 --- a/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp +++ b/sycl/test-e2e/FreeFunctionKernels/accessor_as_kernel_parameter.cpp @@ -72,6 +72,7 @@ int runSingleTaskTest(sycl::queue &Queue, sycl::context &Context, 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, @@ -125,6 +126,7 @@ int runNdRangeTestMultipleParameters(sycl::queue &Queue, sycl::context &Context, return performResultCheck(NdRange.get_global_range().size(), ResultData.data(), ErrorMessage, Values[2]); } + int main() { int Failed = 0; From 5d2ddc30ef493442eced4594ce60f7d9e020f26f Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Thu, 5 Jun 2025 10:40:44 +0200 Subject: [PATCH 10/10] fix --- .../FreeFunctionKernels/local_accessor_as_kernel_parameter.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/test-e2e/FreeFunctionKernels/local_accessor_as_kernel_parameter.cpp b/sycl/test-e2e/FreeFunctionKernels/local_accessor_as_kernel_parameter.cpp index e9b410486e35e..74167ea08917a 100644 --- a/sycl/test-e2e/FreeFunctionKernels/local_accessor_as_kernel_parameter.cpp +++ b/sycl/test-e2e/FreeFunctionKernels/local_accessor_as_kernel_parameter.cpp @@ -57,6 +57,7 @@ void FillWithData(std::vector &Data, std::vector &Values) { Values[i]); } } + int main() { int Failed = 0;