[SYCL][E2E] Add free function kernel tests as device function and host function (#18824)

This PR adds new e2e tests for free function kernels extension based on
test plan
https://github.com/intel/llvm/blob/sycl/sycl/test-e2e/FreeFunctionKernels/test-plan.md#perform-test-that-free-function-kernel-can-be-used-as-device-function-within-another-kernel

Extension spec:
https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc

The overall idea behind this test to verify if free function kernel even
if marked with one of the properties (`nd_range_kernel` and
`single_task_kernel`) can still be used as device or host function.

Author: dyniols

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 <numeric>
+
+#include <sycl/usm.hpp>
+
+#include <sycl/ext/oneapi/free_function_queries.hpp>
+
+#include "helpers.hpp"
+
+template <typename T, int Dims>
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<Dims>))
+void setValues(T *DataPtr, size_t N, T ExpectedResult) {
+#if __SYCL_DEVICE_ONLY__
+  auto GlobalLinId =
+      syclext::this_work_item::get_nd_item<Dims>().get_global_linear_id();
+  DataPtr[GlobalLinId] = ExpectedResult;
+#else
+  for (size_t I = 0; I < N; ++I)
+    DataPtr[I] = ExpectedResult;
+#endif
+}
+
+template <typename T>
+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 <typename T>
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel))
+void singleTaskKernel(T *DataPtr, size_t N) {
+  performReverse(DataPtr, N);
+}
+} // namespace ns
+
+template <typename T, int Dims>
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<Dims>))
+void ndRangekernel(T *DataPtr, size_t N, T ExpectedResult) {
+  setValues<T, Dims>(DataPtr, N, ExpectedResult);
+}
+
+int main() {
+  int Failed = 0;
+  constexpr size_t N = 256;
+  {
+    constexpr int ExpectedResultValue = 111;
+    std::array<int, N> Numbers;
+    std::fill(Numbers.begin(), Numbers.end(), 0);
+    setValues<int, 1>(Numbers.data(), Numbers.size(), ExpectedResultValue);
+    Failed += performResultCheck(
+        N, Numbers.data(),
+        "setValues() free function kernel used as normal host function",
+        ExpectedResultValue);
+  }
+
+  {
+    std::array<int, N> Numbers;
+    std::iota(Numbers.begin(), Numbers.end(), 0);
+    std::array<int, N> ExpectedResultValues;
+    std::iota(ExpectedResultValues.begin(), ExpectedResultValues.end(), 0);
+    std::reverse(ExpectedResultValues.begin(), ExpectedResultValues.end());
+    performReverse(Numbers.data(), Numbers.size());
+    Failed += performResultCheck<N>(
+        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<ns::singleTaskKernel<int>>(Context);
+    std::array<int, N> ExpectedResultValues;
+    std::iota(ExpectedResultValues.begin(), ExpectedResultValues.end(), 0);
+    std::reverse(ExpectedResultValues.begin(), ExpectedResultValues.end());
+
+    int *DataPtr = sycl::malloc_shared<int>(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<N>(
+        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<ndRangekernel<int, 3>>(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<int>(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;
+}

sycl/test-e2e/FreeFunctionKernels/helpers.hpp

@@ -21,6 +21,21 @@ static int performResultCheck(size_t NumberOfElements, const T *ResultPtr,
   return Failed;
 }
 
+template <size_t NumOfElements, typename T, typename S>
+static int
+performResultCheck(const T *ResultPtr, std::string_view TestName,
+                   std::array<S, NumOfElements> 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 <auto *Func> static sycl::kernel getKernel(sycl::context &Context) {
   sycl::kernel_bundle<sycl::bundle_state::executable> KernelBundle =
       syclexp::get_kernel_bundle<Func, sycl::bundle_state::executable>(Context);