No all binary functions are symmetric, modified py_binary_func accordingly

Introduced contig_row_contig_matrix_broadcasting_impl_fn_ptr_t and
corresponding table. Implemented that for Add to fall back on
contig_matrix_contrig_row_broadcasting_fn.

It would be good to have a specialization for symmetric variants.

Author: oleksandr-pavlyk

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

@@ -436,6 +436,67 @@ struct AddContigMatrixContigRowBroadcastFactory
     }
 };
 
+typedef sycl::event (*add_contig_row_contig_matrix_broadcast_impl_fn_ptr_t)(
+    sycl::queue,
+    std::vector<sycl::event> &,
+    size_t,
+    size_t,
+    const char *,
+    py::ssize_t,
+    const char *,
+    py::ssize_t,
+    char *,
+    py::ssize_t,
+    const std::vector<sycl::event> &);
+
+template <typename argT1, typename argT2, typename resT>
+sycl::event add_contig_row_contig_matrix_broadcast_impl(
+    sycl::queue exec_q,
+    std::vector<sycl::event> &host_tasks,
+    size_t n0,
+    size_t n1,
+    const char *vec_p, // typeless pointer to (n1,) contiguous row
+    py::ssize_t vec_offset,
+    const char *mat_p, // typeless pointer to (n0, n1) C-contiguous matrix
+    py::ssize_t mat_offset,
+    char *res_p, // typeless pointer to (n0, n1) result C-contig. matrix,
+                 //    res[i,j] = mat[i,j] + vec[j]
+    py::ssize_t res_offset,
+    const std::vector<sycl::event> &depends = {})
+{
+    return add_contig_matrix_contig_row_broadcast_impl<argT2, argT1, resT>(
+        exec_q, host_tasks, n0, n1, mat_p, mat_offset, vec_p, vec_offset, res_p,
+        res_offset, depends);
+};
+
+template <typename fnT, typename T1, typename T2>
+struct AddContigRowContigMatrixBroadcastFactory
+{
+    fnT get()
+    {
+        if constexpr (std::is_same_v<typename AddOutputType<T1, T2>::value_type,
+                                     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)
+            {
+                fnT fn = nullptr;
+                return fn;
+            }
+            else {
+                fnT fn =
+                    add_contig_row_contig_matrix_broadcast_impl<T1, T2, resT>;
+                return fn;
+            }
+        }
+    }
+};
+
 } // namespace add
 } // namespace kernels
 } // namespace tensor

dpctl/tensor/libtensor/source/elementwise_functions.cpp

@@ -283,6 +283,7 @@ namespace fn_ns = dpctl::tensor::kernels::add;
 
 using fn_ns::add_contig_impl_fn_ptr_t;
 using fn_ns::add_contig_matrix_contig_row_broadcast_impl_fn_ptr_t;
+using fn_ns::add_contig_row_contig_matrix_broadcast_impl_fn_ptr_t;
 using fn_ns::add_strided_impl_fn_ptr_t;
 
 static add_contig_impl_fn_ptr_t add_contig_dispatch_table[td_ns::num_types]
@@ -292,35 +293,51 @@ static int add_output_id_table[td_ns::num_types][td_ns::num_types];
 static add_strided_impl_fn_ptr_t add_strided_dispatch_table[td_ns::num_types]
                                                            [td_ns::num_types];
 
+// add(matrix, row)
 static add_contig_matrix_contig_row_broadcast_impl_fn_ptr_t
     add_contig_matrix_contig_row_broadcast_dispatch_table[td_ns::num_types]
                                                          [td_ns::num_types];
 
+// add(row, matrix)
+static add_contig_row_contig_matrix_broadcast_impl_fn_ptr_t
+    add_contig_row_contig_matrix_broadcast_dispatch_table[td_ns::num_types]
+                                                         [td_ns::num_types];
+
 void populate_add_dispatch_tables(void)
 {
     using namespace td_ns;
 
-    using fn_ns::AddContigFactory;
-    DispatchTableBuilder<add_contig_impl_fn_ptr_t, AddContigFactory, num_types>
-        dtb1;
-    dtb1.populate_dispatch_table(add_contig_dispatch_table);
+    // which input types are supported, and what is the type of the result
+    using fn_ns::AddTypeMapFactory;
+    DispatchTableBuilder<int, AddTypeMapFactory, num_types> dtb1;
+    dtb1.populate_dispatch_table(add_output_id_table);
 
+    // function pointers for operation on general strided arrays
     using fn_ns::AddStridedFactory;
     DispatchTableBuilder<add_strided_impl_fn_ptr_t, AddStridedFactory,
                          num_types>
         dtb2;
     dtb2.populate_dispatch_table(add_strided_dispatch_table);
 
-    using fn_ns::AddTypeMapFactory;
-    DispatchTableBuilder<int, AddTypeMapFactory, num_types> dtb3;
-    dtb3.populate_dispatch_table(add_output_id_table);
+    // function pointers for operation on contiguous inputs and outputs
+    using fn_ns::AddContigFactory;
+    DispatchTableBuilder<add_contig_impl_fn_ptr_t, AddContigFactory, num_types>
+        dtb3;
+    dtb3.populate_dispatch_table(add_contig_dispatch_table);
 
     using fn_ns::AddContigMatrixContigRowBroadcastFactory;
     DispatchTableBuilder<add_contig_matrix_contig_row_broadcast_impl_fn_ptr_t,
                          AddContigMatrixContigRowBroadcastFactory, num_types>
         dtb4;
     dtb4.populate_dispatch_table(
         add_contig_matrix_contig_row_broadcast_dispatch_table);
+
+    using fn_ns::AddContigRowContigMatrixBroadcastFactory;
+    DispatchTableBuilder<add_contig_row_contig_matrix_broadcast_impl_fn_ptr_t,
+                         AddContigRowContigMatrixBroadcastFactory, num_types>
+        dtb5;
+    dtb5.populate_dispatch_table(
+        add_contig_row_contig_matrix_broadcast_dispatch_table);
 };
 
 } // namespace impl
@@ -365,6 +382,7 @@ void init_elementwise_functions(py::module_ m)
         impl::populate_add_dispatch_tables();
         using impl::add_contig_dispatch_table;
         using impl::add_contig_matrix_contig_row_broadcast_dispatch_table;
+        using impl::add_contig_row_contig_matrix_broadcast_dispatch_table;
         using impl::add_output_id_table;
         using impl::add_strided_dispatch_table;
 
@@ -382,7 +400,10 @@ void init_elementwise_functions(py::module_ m)
                 add_strided_dispatch_table,
                 // function pointers to handle operation of c-contig matrix and
                 // c-contig row with broadcasting (may be nullptr)
-                add_contig_matrix_contig_row_broadcast_dispatch_table);
+                add_contig_matrix_contig_row_broadcast_dispatch_table,
+                // function pointers to handle operation of c-contig matrix and
+                // c-contig row with broadcasting (may be nullptr)
+                add_contig_row_contig_matrix_broadcast_dispatch_table);
         };
         auto add_result_type_pyapi = [&](py::dtype dtype1, py::dtype dtype2) {
             return py_binary_ufunc_result_type(dtype1, dtype2,

dpctl/tensor/libtensor/source/elementwise_functions.hpp

@@ -297,7 +297,8 @@ bool isEqual(Container const &c, std::initializer_list<T> const &l)
 template <typename output_typesT,
           typename contig_dispatchT,
           typename strided_dispatchT,
-          typename matrix_row_dispatchT>
+          typename contig_matrix_row_dispatchT,
+          typename contig_row_matrix_dispatchT>
 std::pair<sycl::event, sycl::event> py_binary_ufunc(
     dpctl::tensor::usm_ndarray src1,
     dpctl::tensor::usm_ndarray src2,
@@ -308,7 +309,10 @@ std::pair<sycl::event, sycl::event> py_binary_ufunc(
     const output_typesT &output_type_table,
     const contig_dispatchT &contig_dispatch_table,
     const strided_dispatchT &strided_dispatch_table,
-    const matrix_row_dispatchT &contig_matrix_row_broadcast_dispatch_table)
+    const contig_matrix_row_dispatchT
+        &contig_matrix_row_broadcast_dispatch_table,
+    const contig_row_matrix_dispatchT
+        &contig_row_matrix_broadcast_dispatch_table)
 {
     // check type_nums
     int src1_typenum = src1.get_typenum();
@@ -507,15 +511,15 @@ std::pair<sycl::event, sycl::event> py_binary_ufunc(
                 isEqual(simplified_src2_strides, {one, simplified_shape[0]}) &&
                 isEqual(simplified_dst_strides, {one, simplified_shape[0]}))
             {
-                auto matrix_row_broadcast_fn =
-                    contig_matrix_row_broadcast_dispatch_table[src2_typeid]
-                                                              [src1_typeid];
-                if (matrix_row_broadcast_fn != nullptr) {
+                auto row_matrix_broadcast_fn =
+                    contig_row_matrix_broadcast_dispatch_table[src1_typeid]
+                                                              [src2_typeid];
+                if (row_matrix_broadcast_fn != nullptr) {
                     size_t n0 = simplified_shape[1];
                     size_t n1 = simplified_shape[0];
-                    sycl::event comp_ev = matrix_row_broadcast_fn(
-                        exec_q, host_tasks, n0, n1, src2_data, src2_offset,
-                        src1_data, src1_offset, dst_data, dst_offset, depends);
+                    sycl::event comp_ev = row_matrix_broadcast_fn(
+                        exec_q, host_tasks, n0, n1, src1_data, src1_offset,
+                        src2_data, src2_offset, dst_data, dst_offset, depends);
 
                     return std::make_pair(
                         dpctl::utils::keep_args_alive(exec_q, {src1, src2, dst},