Added BinaryContigMatrixContigRowBroadcastFunctor and RowMatrix variant

Added to common.hpp templated callable to generate kernels for arbitrary
binary operator. Applied that to addition code.

Also implemented true_divide operator, exported as _tensor_impl._divide

Author: oleksandr-pavlyk

dpctl/tensor/libtensor/include/kernels/elementwise_functions/add.hpp

@@ -163,7 +163,7 @@ sycl::event add_contig_impl(sycl::queue exec_q,
                             py::ssize_t res_offset,
                             const std::vector<sycl::event> &depends = {})
 {
-    sycl::event add_ev = exec_q.submit([&](sycl::handler &cgh) {
+    sycl::event comp_ev = exec_q.submit([&](sycl::handler &cgh) {
         cgh.depends_on(depends);
 
         size_t lws = 64;
@@ -188,7 +188,7 @@ sycl::event add_contig_impl(sycl::queue exec_q,
             AddContigFunctor<argTy1, argTy2, resTy, vec_sz, n_vecs>(
                 arg1_tp, arg2_tp, res_tp, nelems));
     });
-    return add_ev;
+    return comp_ev;
 }
 
 template <typename fnT, typename T1, typename T2> struct AddContigFactory
@@ -249,7 +249,7 @@ sycl::event add_strided_impl(sycl::queue exec_q,
                              const std::vector<sycl::event> &depends,
                              const std::vector<sycl::event> &additional_depends)
 {
-    sycl::event abs_ev = exec_q.submit([&](sycl::handler &cgh) {
+    sycl::event comp_ev = exec_q.submit([&](sycl::handler &cgh) {
         cgh.depends_on(depends);
         cgh.depends_on(additional_depends);
 
@@ -270,7 +270,7 @@ sycl::event add_strided_impl(sycl::queue exec_q,
             {nelems}, AddStridedFunctor<argTy1, argTy2, resTy, IndexerT>(
                           arg1_tp, arg2_tp, res_tp, indexer));
     });
-    return abs_ev;
+    return comp_ev;
 }
 
 template <typename fnT, typename T1, typename T2> struct AddStridedFactory
@@ -290,7 +290,7 @@ template <typename fnT, typename T1, typename T2> struct AddStridedFactory
 };
 
 template <typename argT1, typename argT2, typename resT>
-class add_matrix_vector_broadcast_sg_krn;
+class add_matrix_row_broadcast_sg_krn;
 
 typedef sycl::event (*add_contig_matrix_contig_row_broadcast_impl_fn_ptr_t)(
     sycl::queue,
@@ -305,6 +305,14 @@ typedef sycl::event (*add_contig_matrix_contig_row_broadcast_impl_fn_ptr_t)(
     py::ssize_t,
     const std::vector<sycl::event> &);
 
+template <typename argT1, typename argT2, typename resT>
+using AddContigMatrixContigRowBroadcastingFunctor =
+    elementwise_common::BinaryContigMatrixContigRowBroadcastingFunctor<
+        argT1,
+        argT2,
+        resT,
+        AddFunctor<argT1, argT2, resT>>;
+
 template <typename argT1, typename argT2, typename resT>
 sycl::event add_contig_matrix_contig_row_broadcast_impl(
     sycl::queue exec_q,
@@ -361,41 +369,11 @@ sycl::event add_contig_matrix_contig_row_broadcast_impl(
         size_t n_groups = (n_elems + lws - 1) / lws;
         auto gwsRange = sycl::range<1>(n_groups * lws);
 
-        cgh.parallel_for<class add_matrix_vector_broadcast_sg_krn<argT1, argT2, resT>>(
+        cgh.parallel_for<
+            class add_matrix_row_broadcast_sg_krn<argT1, argT2, resT>>(
             sycl::nd_range<1>(gwsRange, lwsRange),
-            [=](sycl::nd_item<1> ndit)
-        {
-            auto sg = ndit.get_sub_group();
-            size_t gid = ndit.get_global_linear_id();
-
-            std::uint8_t sgSize = sg.get_local_range()[0];
-            size_t base = gid - sg.get_local_id()[0];
-
-            if (base + sgSize < n_elems) {
-                using in_ptrT1 =
-                    sycl::multi_ptr<const argT1,
-                                    sycl::access::address_space::global_space>;
-                using in_ptrT2 =
-                    sycl::multi_ptr<const argT2,
-                                    sycl::access::address_space::global_space>;
-                using res_ptrT =
-                    sycl::multi_ptr<resT,
-                                    sycl::access::address_space::global_space>;
-
-                const argT1 mat_el = sg.load(in_ptrT1(&mat[base]));
-                const argT2 vec_el = sg.load(in_ptrT2(&padded_vec[base % n1]));
-
-                resT res_el = mat_el + vec_el;
-
-                sg.store(res_ptrT(&res[base]), res_el);
-            }
-            else {
-                for (size_t k = base + sg.get_local_id()[0]; k < n_elems;
-                     k += sgSize) {
-                    res[k] = mat[k] + padded_vec[k % n1];
-                }
-            }
-        });
+            AddContigMatrixContigRowBroadcastingFunctor<argT1, argT2, resT>(
+                mat, padded_vec, res, n_elems, n1));
     });
 
     sycl::event tmp_cleanup_ev = exec_q.submit([&](sycl::handler &cgh) {
@@ -413,13 +391,12 @@ struct AddContigMatrixContigRowBroadcastFactory
 {
     fnT get()
     {
-        if constexpr (std::is_same_v<typename AddOutputType<T1, T2>::value_type,
-                                     void>) {
+        using resT = typename AddOutputType<T1, T2>::value_type;
+        if constexpr (std::is_same_v<resT, void>) {
             fnT fn = nullptr;
             return fn;
         }
         else {
-            using resT = typename AddOutputType<T1, T2>::value_type;
             if constexpr (dpctl::tensor::type_utils::is_complex<T1>::value ||
                           dpctl::tensor::type_utils::is_complex<T2>::value ||
                           dpctl::tensor::type_utils::is_complex<resT>::value)
@@ -474,13 +451,12 @@ struct AddContigRowContigMatrixBroadcastFactory
 {
     fnT get()
     {
-        if constexpr (std::is_same_v<typename AddOutputType<T1, T2>::value_type,
-                                     void>) {
+        using resT = typename AddOutputType<T1, T2>::value_type;
+        if constexpr (std::is_same_v<resT, void>) {
             fnT fn = nullptr;
             return fn;
         }
         else {
-            using resT = typename AddOutputType<T1, T2>::value_type;
             if constexpr (dpctl::tensor::type_utils::is_complex<T1>::value ||
                           dpctl::tensor::type_utils::is_complex<T2>::value ||
                           dpctl::tensor::type_utils::is_complex<resT>::value)

dpctl/tensor/libtensor/include/kernels/elementwise_functions/common.hpp

@@ -420,6 +420,130 @@ struct BinaryStridedFunctor
     }
 };
 
+template <typename argT1,
+          typename argT2,
+          typename resT,
+          typename BinaryOperatorT>
+struct BinaryContigMatrixContigRowBroadcastingFunctor
+{
+private:
+    const argT1 *mat;
+    const argT2 *padded_vec;
+    resT *res;
+    size_t n_elems;
+    size_t n1;
+
+public:
+    BinaryContigMatrixContigRowBroadcastingFunctor(const argT1 *mat_tp,
+                                                   const argT2 *row_tp,
+                                                   resT *res_tp,
+                                                   size_t n_elems_in_mat,
+                                                   size_t n_elems_in_row)
+        : mat(mat_tp), padded_vec(row_tp), res(res_tp), n_elems(n_elems_in_mat),
+          n1(n_elems_in_row)
+    {
+    }
+
+    void operator()(sycl::nd_item<1> ndit) const
+    {
+        BinaryOperatorT op{};
+        static_assert(BinaryOperatorT::supports_sg_loadstore::value);
+
+        auto sg = ndit.get_sub_group();
+        size_t gid = ndit.get_global_linear_id();
+
+        std::uint8_t sgSize = sg.get_local_range()[0];
+        size_t base = gid - sg.get_local_id()[0];
+
+        if (base + sgSize < n_elems) {
+            using in_ptrT1 =
+                sycl::multi_ptr<const argT1,
+                                sycl::access::address_space::global_space>;
+            using in_ptrT2 =
+                sycl::multi_ptr<const argT2,
+                                sycl::access::address_space::global_space>;
+            using res_ptrT =
+                sycl::multi_ptr<resT,
+                                sycl::access::address_space::global_space>;
+
+            const argT1 mat_el = sg.load(in_ptrT1(&mat[base]));
+            const argT2 vec_el = sg.load(in_ptrT2(&padded_vec[base % n1]));
+
+            resT res_el = op(mat_el, vec_el);
+
+            sg.store(res_ptrT(&res[base]), res_el);
+        }
+        else {
+            for (size_t k = base + sg.get_local_id()[0]; k < n_elems;
+                 k += sgSize) {
+                res[k] = op(mat[k], padded_vec[k % n1]);
+            }
+        }
+    }
+};
+
+template <typename argT1,
+          typename argT2,
+          typename resT,
+          typename BinaryOperatorT>
+struct BinaryContigRowContigMatrixBroadcastingFunctor
+{
+private:
+    const argT1 *padded_vec;
+    const argT2 *mat;
+    resT *res;
+    size_t n_elems;
+    size_t n1;
+
+public:
+    BinaryContigRowContigMatrixBroadcastingFunctor(const argT1 *row_tp,
+                                                   const argT2 *mat_tp,
+                                                   resT *res_tp,
+                                                   size_t n_elems_in_mat,
+                                                   size_t n_elems_in_row)
+        : padded_vec(row_tp), mat(mat_tp), res(res_tp), n_elems(n_elems_in_mat),
+          n1(n_elems_in_row)
+    {
+    }
+
+    void operator()(sycl::nd_item<1> ndit) const
+    {
+        BinaryOperatorT op{};
+        static_assert(BinaryOperatorT::supports_sg_loadstore::value);
+
+        auto sg = ndit.get_sub_group();
+        size_t gid = ndit.get_global_linear_id();
+
+        std::uint8_t sgSize = sg.get_local_range()[0];
+        size_t base = gid - sg.get_local_id()[0];
+
+        if (base + sgSize < n_elems) {
+            using in_ptrT1 =
+                sycl::multi_ptr<const argT1,
+                                sycl::access::address_space::global_space>;
+            using in_ptrT2 =
+                sycl::multi_ptr<const argT2,
+                                sycl::access::address_space::global_space>;
+            using res_ptrT =
+                sycl::multi_ptr<resT,
+                                sycl::access::address_space::global_space>;
+
+            const argT2 mat_el = sg.load(in_ptrT2(&mat[base]));
+            const argT1 vec_el = sg.load(in_ptrT1(&padded_vec[base % n1]));
+
+            resT res_el = op(vec_el, mat_el);
+
+            sg.store(res_ptrT(&res[base]), res_el);
+        }
+        else {
+            for (size_t k = base + sg.get_local_id()[0]; k < n_elems;
+                 k += sgSize) {
+                res[k] = op(padded_vec[k % n1], mat[k]);
+            }
+        }
+    }
+};
+
 } // namespace elementwise_common
 } // namespace kernels
 } // namespace tensor