diff --git a/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_as_device_host_functions.cpp b/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_as_device_host_functions.cpp new file mode 100644 index 0000000000000..81d714c0d6ddd --- /dev/null +++ b/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_as_device_host_functions.cpp @@ -0,0 +1,127 @@ +// REQUIRES: aspect-usm_shared_allocations +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// This test verifies whether free function kernel can be used as device +// function within another kernel or can be used as normal host function. + +#include + +#include + +#include + +#include "helpers.hpp" + +template +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel)) +void setValues(T *DataPtr, size_t N, T ExpectedResult) { +#if __SYCL_DEVICE_ONLY__ + auto GlobalLinId = + syclext::this_work_item::get_nd_item().get_global_linear_id(); + DataPtr[GlobalLinId] = ExpectedResult; +#else + for (size_t I = 0; I < N; ++I) + DataPtr[I] = ExpectedResult; +#endif +} + +template +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel)) +void performReverse(T *DataPtr, size_t N) { + for (size_t I = 0, J = N - 1; I < J; ++I, --J) { + std::swap(DataPtr[I], DataPtr[J]); + } +} + +namespace ns { +template +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel)) +void singleTaskKernel(T *DataPtr, size_t N) { + performReverse(DataPtr, N); +} +} // namespace ns + +template +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel)) +void ndRangekernel(T *DataPtr, size_t N, T ExpectedResult) { + setValues(DataPtr, N, ExpectedResult); +} + +int main() { + int Failed = 0; + constexpr size_t N = 256; + { + constexpr int ExpectedResultValue = 111; + std::array Numbers; + std::fill(Numbers.begin(), Numbers.end(), 0); + setValues(Numbers.data(), Numbers.size(), ExpectedResultValue); + Failed += performResultCheck( + N, Numbers.data(), + "setValues() free function kernel used as normal host function", + ExpectedResultValue); + } + + { + std::array Numbers; + std::iota(Numbers.begin(), Numbers.end(), 0); + std::array ExpectedResultValues; + std::iota(ExpectedResultValues.begin(), ExpectedResultValues.end(), 0); + std::reverse(ExpectedResultValues.begin(), ExpectedResultValues.end()); + performReverse(Numbers.data(), Numbers.size()); + Failed += performResultCheck( + Numbers.data(), + "performReverse() free function kernel used as normal host function", + ExpectedResultValues); + } + + sycl::queue Queue; + sycl::context Context = Queue.get_context(); + + { + sycl::kernel UsedKernel = getKernel>(Context); + std::array ExpectedResultValues; + std::iota(ExpectedResultValues.begin(), ExpectedResultValues.end(), 0); + std::reverse(ExpectedResultValues.begin(), ExpectedResultValues.end()); + + int *DataPtr = sycl::malloc_shared(N, Queue); + std::iota(DataPtr, DataPtr + N, 0); + + Queue + .submit([&](sycl::handler &Handler) { + Handler.set_args(DataPtr, N); + Handler.single_task(UsedKernel); + }) + .wait(); + Failed += performResultCheck( + DataPtr, + "performReverse() free function kernel used as device function within " + "another kernel", + ExpectedResultValues); + sycl::free(DataPtr, Queue); + } + + { + constexpr int ExpectedResultValue = 222; + + sycl::kernel UsedKernel = getKernel>(Context); + + sycl::nd_range NdRange{sycl::range{16, 4, 4}, sycl::range{2, 2, 2}}; + size_t NumberOfElements = NdRange.get_global_range().size(); + int *DataPtr = sycl::malloc_shared(NumberOfElements, Queue); + std::fill(DataPtr, DataPtr + NumberOfElements, 0); + Queue + .submit([&](sycl::handler &Handler) { + Handler.set_args(DataPtr, NumberOfElements, ExpectedResultValue); + Handler.parallel_for(NdRange, UsedKernel); + }) + .wait(); + + Failed += performResultCheck(NumberOfElements, DataPtr, + "setValues() free function kernel used as " + "device function within another kernel", + ExpectedResultValue); + sycl::free(DataPtr, Queue); + } + return Failed; +} diff --git a/sycl/test-e2e/FreeFunctionKernels/helpers.hpp b/sycl/test-e2e/FreeFunctionKernels/helpers.hpp index bb561589f9711..63d0ca7802825 100644 --- a/sycl/test-e2e/FreeFunctionKernels/helpers.hpp +++ b/sycl/test-e2e/FreeFunctionKernels/helpers.hpp @@ -21,6 +21,21 @@ static int performResultCheck(size_t NumberOfElements, const T *ResultPtr, return Failed; } +template +static int +performResultCheck(const T *ResultPtr, std::string_view TestName, + std::array ExpectedResultValue) { + int Failed{0}; + for (size_t i = 0; i < NumOfElements; i++) { + if (ResultPtr[i] != ExpectedResultValue[i]) { + std::cerr << "Failed " << TestName << " : " << ResultPtr[i] + << " != " << ExpectedResultValue[i] << std::endl; + ++Failed; + } + } + return Failed; +} + template static sycl::kernel getKernel(sycl::context &Context) { sycl::kernel_bundle KernelBundle = syclexp::get_kernel_bundle(Context);