Attempt to make use of DPC++ generated kernels in testing of DPCTLQueue_SubmitRange

Author: oleksandr-pavlyk

dpctl-capi/tests/test_sycl_queue_submit.cpp

@@ -123,3 +123,306 @@ TEST_F(TestQueueSubmit, CheckSubmitRange_saxpy)
     DPCTLDevice_Delete(DRef);
     DPCTLDeviceSelector_Delete(DSRef);
 }
+
+namespace
+{
+
+template <typename T> class populate_a;
+
+template <typename T> class populate_b;
+
+template <typename T, typename scT> class add_a_and_b;
+
+template <typename T> struct kernel_arg_t
+{
+    static constexpr DPCTLKernelArgType value = DPCTL_VOID_PTR;
+};
+
+/*
+template <>
+struct kernel_arg_t<char> {
+    static constexpr DPCTLKernelArgType value = DPCTL_CHAR;
+};
+
+template <>
+struct kernel_arg_t<signed char> {
+    static constexpr DPCTLKernelArgType value = DPCTL_SIGNED_CHAR;
+};
+
+template <>
+struct kernel_arg_t<unsigned char> {
+    static constexpr DPCTLKernelArgType value = DPCTL_UNSIGNED_CHAR;
+};
+*/
+template <> struct kernel_arg_t<short>
+{
+    static constexpr DPCTLKernelArgType value = DPCTL_SHORT;
+};
+
+template <> struct kernel_arg_t<int>
+{
+    static constexpr DPCTLKernelArgType value = DPCTL_INT;
+};
+
+template <> struct kernel_arg_t<unsigned int>
+{
+    static constexpr DPCTLKernelArgType value = DPCTL_UNSIGNED_INT;
+};
+
+template <> struct kernel_arg_t<long>
+{
+    static constexpr DPCTLKernelArgType value = DPCTL_LONG;
+};
+
+template <> struct kernel_arg_t<unsigned long>
+{
+    static constexpr DPCTLKernelArgType value = DPCTL_UNSIGNED_LONG;
+};
+
+template <> struct kernel_arg_t<long long>
+{
+    static constexpr DPCTLKernelArgType value = DPCTL_LONG_LONG;
+};
+
+template <> struct kernel_arg_t<unsigned long long>
+{
+    static constexpr DPCTLKernelArgType value = DPCTL_UNSIGNED_LONG_LONG;
+};
+
+/*
+template <>
+struct kernel_arg_t<size_t> {
+    static constexpr DPCTLKernelArgType value = DPCTL_SIZE_T;
+};
+*/
+
+template <> struct kernel_arg_t<float>
+{
+    static constexpr DPCTLKernelArgType value = DPCTL_FLOAT;
+};
+
+template <> struct kernel_arg_t<double>
+{
+    static constexpr DPCTLKernelArgType value = DPCTL_DOUBLE;
+};
+
+/*
+template <>
+struct kernel_arg_t<long double> {
+static constexpr DPCTLKernelArgType value = DPCTL_LONG_DOUBLE;
+};
+*/
+
+#ifdef USE_FUNCTOR
+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 Add
+{
+    const Ty *in1, *in2;
+    Ty *out;
+    scT val;
+    Add(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 PopulateA
+{
+    T *out;
+    PopulateA(T *a) : out(a) {}
+    void operator()(sycl::id<1> myId) const
+    {
+        auto gid = myId[0];
+        out[gid] = T(1);
+        return;
+    };
+};
+
+template <typename T> struct PopulateB
+{
+    T *out;
+    PopulateB(T *b) : out(b) {}
+    void operator()(sycl::id<1> myId) const
+    {
+        auto gid = myId[0];
+        out[gid] = T(gid);
+        return;
+    };
+};
+#endif
+
+template <typename T, typename scT>
+bool common_submit_range_fn(sycl::queue &q, size_t n, scT val)
+{
+    sycl::program program(q.get_context());
+
+    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 *d = sycl::malloc_device<T>(n, q);
+
+#ifdef USE_FUNCTOR
+    auto popa_fn = make_cgh_function<class populate_a<T>>(n, PopulateA<T>(a));
+#else
+    auto popa_fn = [&](sycl::handler &cgh) {
+        cgh.parallel_for<populate_a<T>>(
+            n, [=](sycl::id<1> idx) { a[idx[0]] = T(1); });
+    };
+#endif
+    std::cout << q.get_context().get_platform().get_backend() << std::endl;
+    sycl::event popa_ev = q.submit(popa_fn);
+
+#ifdef USE_FUNCTOR
+    auto popb_fn = make_cgh_function<class populate_b<T>>(n, PopulateB<T>(b));
+#else
+    auto popb_fn = [&](sycl::handler &cgh) {
+        cgh.parallel_for<populate_b<T>>(
+            n, [=](sycl::id<1> idx) { b[idx[0]] = T(idx[0]); });
+    };
+#endif
+    sycl::event popb_ev = q.submit(popb_fn);
+
+#ifdef USE_FUNCTOR
+    auto add_fn = make_cgh_function<class add_a_and_b<T, scT>>(
+        n, Add<T, scT>(a, b, c, val));
+#else
+    auto add_fn = [&](sycl::handler &cgh) {
+        cgh.depends_on({popa_ev, popb_ev});
+        cgh.parallel_for<add_a_and_b<T, scT>>(n, [=](sycl::id<1> idx) {
+            const auto gid = idx[0];
+            const T va = a[gid];
+            const T vb = b[gid];
+            c[gid] = va + val * vb;
+        });
+    };
+#endif
+
+    sycl::event add_ev = q.submit(add_fn);
+    add_ev.wait_and_throw();
+
+    program.build_with_kernel_type<add_a_and_b<T, scT>>();
+    auto kern = program.get_kernel<add_a_and_b<T, scT>>();
+
+    DPCTLSyclQueueRef QRef = reinterpret_cast<DPCTLSyclQueueRef>(&q);
+    DPCTLSyclKernelRef KRef = reinterpret_cast<DPCTLSyclKernelRef>(&kern);
+    DPCTLSyclEventRef PopAERef = reinterpret_cast<DPCTLSyclEventRef>(&popa_ev);
+    DPCTLSyclEventRef PopBERef = reinterpret_cast<DPCTLSyclEventRef>(&popb_ev);
+
+    void *args2[4] = {reinterpret_cast<void *>(a), reinterpret_cast<void *>(b),
+                      reinterpret_cast<void *>(d),
+                      reinterpret_cast<void *>(&val)};
+    DPCTLKernelArgType addKernelArgTypes[] = {DPCTL_VOID_PTR, DPCTL_VOID_PTR,
+                                              DPCTL_VOID_PTR,
+                                              kernel_arg_t<scT>::value};
+    size_t Range[] = {n};
+    DPCTLSyclEventRef events[2] = {PopAERef, PopBERef};
+    auto ERef = DPCTLQueue_SubmitRange(KRef, QRef, args2, addKernelArgTypes, 4,
+                                       Range, 1, events, 2);
+    DPCTLQueue_Wait(QRef);
+    DPCTLEvent_Delete(ERef);
+
+    T *host_data = new T[n];
+    q.memcpy(host_data, d, n * sizeof(T));
+    q.wait_and_throw();
+
+    bool worked = true;
+    for (size_t i = 0; i < n; ++i) {
+        worked = worked && (host_data[i] == T(1) + val * T(i));
+    }
+
+    sycl::free(a, q);
+    sycl::free(b, q);
+    sycl::free(c, q);
+    sycl::free(d, q);
+
+    return worked;
+};
+
+} // end of anonymous 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, int>(q, n_elems, int(-1));
+    EXPECT_TRUE(worked);
+}
+
+TEST_F(TestQueueSubmitRange, ChkSubmitRangeUnsignedInt)
+{
+    bool worked = false;
+    worked = common_submit_range_fn<unsigned int, unsigned int>(q, n_elems, 2);
+    EXPECT_TRUE(worked);
+}
+
+TEST_F(TestQueueSubmitRange, ChkSubmitRangeShort)
+{
+    bool worked = false;
+    worked = common_submit_range_fn<short, short>(q, n_elems, short(-1));
+    EXPECT_TRUE(worked);
+}
+
+TEST_F(TestQueueSubmitRange, ChkSubmitRangeLong)
+{
+    bool worked = false;
+    worked = common_submit_range_fn<long, long>(q, n_elems, -1);
+    EXPECT_TRUE(worked);
+}
+
+TEST_F(TestQueueSubmitRange, ChkSubmitRangeUnsignedLong)
+{
+    bool worked = false;
+    worked =
+        common_submit_range_fn<unsigned long, unsigned long>(q, n_elems, 2);
+    EXPECT_TRUE(worked);
+}
+
+TEST_F(TestQueueSubmitRange, ChkSubmitRangeLongLong)
+{
+    bool worked = false;
+    worked = common_submit_range_fn<long long, long long>(q, n_elems, -1);
+    EXPECT_TRUE(worked);
+}
+
+TEST_F(TestQueueSubmitRange, ChkSubmitRangeUnsignedLongLong)
+{
+    bool worked = false;
+    worked = common_submit_range_fn<unsigned long long, unsigned long long>(
+        q, n_elems, 2);
+    EXPECT_TRUE(worked);
+}
+
+TEST_F(TestQueueSubmitRange, ChkSubmitRangeFloat)
+{
+    bool worked = false;
+    worked = common_submit_range_fn<float, float>(q, n_elems, 0.5);
+    EXPECT_TRUE(worked);
+}
+
+TEST_F(TestQueueSubmitRange, ChkSubmitRangeDouble)
+{
+    bool worked = false;
+    worked = common_submit_range_fn<double, double>(q, n_elems, 0.5);
+    EXPECT_TRUE(worked);
+}