dpcpp_kernels.cpp exports functions that compile functions to sycl::kernel

test_sycl_queue_submit.cpp uses dpcpp_kernel to get kernels and submit them
using DPCTLQueue_SubmitRange, and  _SubmitNDRange in the future.

This does not work with coverage turned off due to LLVM version incompatibility
between IGC and DPCPP.

Author: oleksandr-pavlyk

dpctl-capi/tests/CMakeLists.txt

@@ -21,6 +21,11 @@ foreach(tf ${spirv-test-files})
     file(COPY ${tf} DESTINATION ${CMAKE_CURRENT_BINARY_DIR})
 endforeach()
 
+add_library(dpcpp_kernels
+    STATIC
+    ${CMAKE_CURRENT_SOURCE_DIR}/dpcpp_kernels.cpp
+)
+
 if(DPCTL_GENERATE_COVERAGE)
     file(GLOB_RECURSE
         sources ${CMAKE_CURRENT_SOURCE_DIR}/*.cpp
@@ -35,11 +40,12 @@ if(DPCTL_GENERATE_COVERAGE)
     list(REMOVE_ITEM
         dpctl_sources
         "${CMAKE_CURRENT_SOURCE_DIR}/../source/dpctl_vector_templ.cpp"
+        "${CMAKE_CURRENT_SOURCE_DIR}/../source/dpcpp_kernels.cpp"
     )
 
     # Add profiling flags
     set(CMAKE_CXX_FLAGS
-        "${CMAKE_CXX_FLAGS} -fprofile-instr-generate -fcoverage-mapping"
+        "${CMAKE_CXX_FLAGS} -fprofile-instr-generate -fcoverage-mapping -DDPCTL_COVERAGE"
     )
 
     # Add all dpctl sources into a single executable so that we can run coverage
@@ -55,6 +61,7 @@ if(DPCTL_GENERATE_COVERAGE)
         GTest::GTest
         ${IntelSycl_OPENCL_LIBRARY}
         ${CMAKE_DL_LIBS}
+	    dpcpp_kernels
     )
     add_custom_target(llvm-cov
         COMMAND ${CMAKE_MAKE_PROGRAM} dpctl_c_api_tests
@@ -96,6 +103,7 @@ else()
         GTest::GTest
         DPCTLSyclInterface
         ${LEVEL_ZERO_LIBRARY}
+	    dpcpp_kernels
     )
 endif()
 

dpctl-capi/tests/dpcpp_kernels.cpp

@@ -0,0 +1,16 @@
+#include "dpcpp_kernels.hpp"
+#include <CL/sycl.hpp>
+#include <cstddef>
+
+template sycl::kernel
+dpcpp_kernels::get_fill_kernel<int>(sycl::queue &, size_t, int *, int);
+
+template sycl::kernel
+dpcpp_kernels::get_range_kernel<int>(sycl::queue &, size_t, int *);
+
+template sycl::kernel dpcpp_kernels::get_mad_kernel<int, int>(sycl::queue &,
+                                                              size_t,
+                                                              int *,
+                                                              int *,
+                                                              int *,
+                                                              int);

dpctl-capi/tests/dpcpp_kernels.hpp

@@ -0,0 +1,108 @@
+#pragma once
+#include <CL/sycl.hpp>
+
+namespace dpcpp_kernels
+{
+
+namespace
+{
+template <typename T> class populate_a;
+
+template <typename T> class populate_b;
+
+template <typename T, typename scT> class mad_kern;
+
+template <typename name, class kernelFunc>
+auto make_cgh_function(int n, kernelFunc func)
+{
+    auto Kernel = [&](sycl::handler &cgh) {
+        cgh.parallel_for<name>(sycl::range<1>(n), func);
+    };
+    return Kernel;
+};
+
+template <typename Ty, typename scT> struct MadFunc
+{
+    const Ty *in1, *in2;
+    Ty *out;
+    scT val;
+    MadFunc(const Ty *a, const Ty *b, Ty *c, scT d)
+        : in1(a), in2(b), out(c), val(d)
+    {
+    }
+    void operator()(sycl::id<1> myId) const
+    {
+        auto gid = myId[0];
+        out[gid] = in1[gid] + val * in2[gid];
+        return;
+    }
+};
+
+template <typename T> struct FillFunc
+{
+    T *out;
+    T val;
+    FillFunc(T *a, T val) : out(a), val(val) {}
+    void operator()(sycl::id<1> myId) const
+    {
+        auto gid = myId[0];
+        out[gid] = val;
+        return;
+    };
+};
+
+template <typename T> struct RangeFunc
+{
+    T *out;
+    RangeFunc(T *b) : out(b) {}
+    void operator()(sycl::id<1> myId) const
+    {
+        auto gid = myId[0];
+        out[gid] = T(gid);
+        return;
+    };
+};
+
+} // namespace
+
+template <typename T>
+sycl::kernel get_fill_kernel(sycl::queue &q, size_t n, T *out, T fill_val)
+{
+    // out[i] = fill_val
+    sycl::program program(q.get_context());
+
+    [[maybe_unused]] auto cgh_fn =
+        make_cgh_function<class populate_a<T>>(n, FillFunc<T>(out, fill_val));
+
+    program.build_with_kernel_type<populate_a<T>>();
+    return program.get_kernel<populate_a<T>>();
+};
+
+template <typename T>
+sycl::kernel get_range_kernel(sycl::queue &q, size_t n, T *b)
+{
+    // b[i] = i
+    sycl::program program(q.get_context());
+
+    [[maybe_unused]] auto cgh_fn =
+        make_cgh_function<class populate_b<T>>(n, RangeFunc<T>(b));
+
+    program.build_with_kernel_type<populate_b<T>>();
+    return program.get_kernel<populate_b<T>>();
+};
+
+template <typename T, typename scT>
+sycl::kernel
+get_mad_kernel(sycl::queue &q, size_t n, T *in1, T *in2, T *out, scT val)
+{
+    // c[i] = a[i] + b[i] * val
+    sycl::program program(q.get_context());
+
+    [[maybe_unused]] auto cgh_fn = make_cgh_function<class mad_kern<T, scT>>(
+        n, MadFunc<T, scT>(in1, in2, out, val));
+
+    program.build_with_kernel_type<mad_kern<T, scT>>();
+    return program.get_kernel<mad_kern<T, scT>>();
+};
+
+} // namespace dpcpp_kernels

dpctl-capi/tests/test_sycl_queue_submit.cpp

@@ -24,6 +24,7 @@
 //===----------------------------------------------------------------------===//
 
 #include "Support/CBindingWrapping.h"
+#include "dpcpp_kernels.hpp"
 #include "dpctl_sycl_context_interface.h"
 #include "dpctl_sycl_device_interface.h"
 #include "dpctl_sycl_device_selector_interface.h"
@@ -123,3 +124,96 @@ TEST_F(TestQueueSubmit, CheckSubmitRange_saxpy)
     DPCTLDevice_Delete(DRef);
     DPCTLDeviceSelector_Delete(DSRef);
 }
+
+#ifndef DPCTL_COVERAGE
+namespace
+{
+
+template <typename T,
+          DPCTLKernelArgType katT,
+          typename scT,
+          DPCTLKernelArgType katscT>
+bool common_submit_range_fn(sycl::queue &q, size_t n, scT val)
+{
+    T *a = sycl::malloc_device<T>(n, q);
+    T *b = sycl::malloc_device<T>(n, q);
+    T *c = sycl::malloc_device<T>(n, q);
+    T fill_val = 1;
+    size_t Range[] = {n};
+
+    auto popA_kernel = dpcpp_kernels::get_fill_kernel<T>(q, n, a, fill_val);
+    auto popB_kernel = dpcpp_kernels::get_range_kernel<T>(q, n, b);
+    auto mad_kernel = dpcpp_kernels::get_mad_kernel<T, scT>(q, n, a, b, c, val);
+
+    DPCTLSyclKernelRef popAKernRef =
+        reinterpret_cast<DPCTLSyclKernelRef>(&popA_kernel);
+    DPCTLSyclKernelRef popBKernRef =
+        reinterpret_cast<DPCTLSyclKernelRef>(&popB_kernel);
+    DPCTLSyclKernelRef madKernRef =
+        reinterpret_cast<DPCTLSyclKernelRef>(&mad_kernel);
+
+    DPCTLSyclQueueRef QRef = reinterpret_cast<DPCTLSyclQueueRef>(&q);
+    void *popAArgs[] = {reinterpret_cast<void *>(a),
+                        reinterpret_cast<void *>(&fill_val)};
+    DPCTLKernelArgType popAKernelArgTypes[] = {DPCTL_VOID_PTR, katT};
+
+    DPCTLSyclEventRef popAERef =
+        DPCTLQueue_SubmitRange(popAKernRef, QRef, popAArgs, popAKernelArgTypes,
+                               2, Range, 1, nullptr, 0);
+
+    void *popBArgs[] = {reinterpret_cast<void *>(b)};
+    DPCTLKernelArgType popBKernelArgTypes[] = {DPCTL_VOID_PTR};
+
+    DPCTLSyclEventRef popBERef =
+        DPCTLQueue_SubmitRange(popBKernRef, QRef, popBArgs, popBKernelArgTypes,
+                               1, Range, 1, nullptr, 0);
+
+    void *madArgs[] = {reinterpret_cast<void *>(a), reinterpret_cast<void *>(b),
+                       reinterpret_cast<void *>(c),
+                       reinterpret_cast<void *>(&val)};
+    DPCTLKernelArgType madKernelArgTypes[] = {DPCTL_VOID_PTR, DPCTL_VOID_PTR,
+                                              DPCTL_VOID_PTR, katscT};
+
+    DPCTLSyclEventRef deps[2] = {popAERef, popBERef};
+    DPCTLSyclEventRef madRef = DPCTLQueue_SubmitRange(
+        madKernRef, QRef, madArgs, madKernelArgTypes, 4, Range, 1, deps, 2);
+
+    DPCTLQueue_Wait(QRef);
+    DPCTLEvent_Delete(madRef);
+    DPCTLEvent_Delete(popBERef);
+    DPCTLEvent_Delete(popAERef);
+
+    bool worked = true;
+    T *host_data = new T[n];
+    q.memcpy(host_data, c, n * sizeof(T)).wait();
+    for (size_t i = 0; i < n; ++i) {
+        worked = worked && (host_data[i] == T(fill_val) + T(i) * T(val));
+    }
+    delete[] host_data;
+
+    sycl::free(c, q);
+    sycl::free(b, q);
+    sycl::free(a, q);
+
+    return worked;
+}
+
+} // namespace
+
+struct TestQueueSubmitRange : public ::testing::Test
+{
+    sycl::queue q;
+    size_t n_elems = 512;
+
+    TestQueueSubmitRange() : q(sycl::default_selector{}) {}
+    ~TestQueueSubmitRange() {}
+};
+
+TEST_F(TestQueueSubmitRange, ChkSubmitRangeInt)
+{
+    bool worked = false;
+    worked = common_submit_range_fn<int, DPCTL_INT, int, DPCTL_INT>(q, n_elems,
+                                                                    int(-1));
+    EXPECT_TRUE(worked);
+}
+#endif