move multiply() to descriptor usage (#740)

Author: shssf

dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp

@@ -66,7 +66,12 @@ MACRO_2ARG_3TYPES_OP(dpnp_fmod_c, cl::sycl::fmod((double)input1_elem, (double)in
 MACRO_2ARG_3TYPES_OP(dpnp_hypot_c, cl::sycl::hypot((double)input1_elem, (double)input2_elem), oneapi::mkl::vm::hypot)
 MACRO_2ARG_3TYPES_OP(dpnp_maximum_c, cl::sycl::max(input1_elem, input2_elem), oneapi::mkl::vm::fmax)
 MACRO_2ARG_3TYPES_OP(dpnp_minimum_c, cl::sycl::min(input1_elem, input2_elem), oneapi::mkl::vm::fmin)
+
+// "multiply" needs to be standalone kernel (not autogenerated) due to complex algorithm. This is not an element wise.
+// pytest "tests/third_party/cupy/creation_tests/test_ranges.py::TestMgrid::test_mgrid3"
+// requires multiplication shape1[10] with shape2[10,1] and result expected as shape[10,10]
 MACRO_2ARG_3TYPES_OP(dpnp_multiply_c, input1_elem* input2_elem, oneapi::mkl::vm::mul)
+
 MACRO_2ARG_3TYPES_OP(dpnp_power_c, cl::sycl::pow((double)input1_elem, (double)input2_elem), oneapi::mkl::vm::pow)
 MACRO_2ARG_3TYPES_OP(dpnp_subtract_c, input1_elem - input2_elem, oneapi::mkl::vm::sub)
 

dpnp/dpnp_algo/dpnp_algo.pyx

@@ -409,3 +409,26 @@ cdef dparray call_fptr_2in_1out(DPNPFuncName fptr_name, object x1_obj, object x2
          x2_dparray.get_data(), x2_dparray.size, x2_shape.data(), x2_shape.size(), NULL)
 
     return result
+
+# this is replacement for "call_fptr_2in_1out". original function must be deleted after transotion.
+cdef dparray call_fptr_2in_1out_new(DPNPFuncName fptr_name, dpnp_descriptor x1_obj, dpnp_descriptor x2_obj,
+                                    object dtype=None, dparray out=None, object where=True):
+    # Convert string type names (dparray.dtype) to C enum DPNPFuncType
+    cdef DPNPFuncType x1_c_type = dpnp_dtype_to_DPNPFuncType(x1_obj.dtype)
+    cdef DPNPFuncType x2_c_type = dpnp_dtype_to_DPNPFuncType(x2_obj.dtype)
+
+    # get the FPTR data structure
+    cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(fptr_name, x1_c_type, x2_c_type)
+
+    # Create result array
+    cdef dparray_shape_type x1_shape = x1_obj.shape
+    cdef dparray_shape_type x2_shape = x2_obj.shape
+    cdef dparray_shape_type result_shape = get_common_shape(x1_shape, x2_shape)
+    cdef dparray result = create_output_array(result_shape, kernel_data.return_type, out)
+
+    """ Call FPTR function """
+    cdef fptr_2in_1out_t func = <fptr_2in_1out_t > kernel_data.ptr
+    func(result.get_data(), x1_obj.get_data(), x1_obj.size, x1_shape.data(), x1_shape.size(),
+         x2_obj.get_data(), x2_obj.size, x2_shape.data(), x2_shape.size(), NULL)
+
+    return result

dpnp/dpnp_algo/dpnp_algo_linearalgebra.pyx

@@ -64,7 +64,9 @@ cpdef dparray dpnp_dot(dparray in_array1, dparray in_array2):
 
     # scalar
     if dim1 == 0 or dim2 == 0:
-        return dpnp_multiply(in_array1, in_array2)
+        x1_desc = dpnp.get_dpnp_descriptor(in_array1)
+        x2_desc = dpnp.get_dpnp_descriptor(in_array2)
+        return dpnp_multiply(x1_desc, x2_desc)
 
     cdef size_t size1 = 0
     cdef size_t size2 = 0

dpnp/dpnp_algo/dpnp_algo_mathematical.pyx

@@ -277,7 +277,7 @@ cpdef tuple dpnp_modf(dpnp_descriptor x1):
 
 
 cpdef dparray dpnp_multiply(object x1_obj, object x2_obj, object dtype=None, dparray out=None, object where=True):
-    return call_fptr_2in_1out(DPNP_FN_MULTIPLY, x1_obj, x2_obj, dtype=dtype, out=out, where=where)
+    return call_fptr_2in_1out_new(DPNP_FN_MULTIPLY, x1_obj, x2_obj, dtype=dtype, out=out, where=where)
 
 
 cpdef dparray dpnp_nancumprod(dparray x1):

dpnp/dpnp_iface_mathematical.py

@@ -1071,19 +1071,20 @@ def multiply(x1, x2, dtype=None, out=None, where=True, **kwargs):
     [1, 4, 9, 16, 25]
 
     """
-    x1_is_scalar, x2_is_scalar = dpnp.isscalar(x1), dpnp.isscalar(x2)
-    x1_is_dparray, x2_is_dparray = isinstance(x1, dparray), isinstance(x2, dparray)
 
-    if not use_origin_backend(x1) and not kwargs:
-        if not x1_is_dparray and not x1_is_scalar:
-            pass
-        elif not x2_is_dparray and not x2_is_scalar:
+    x1_is_scalar = dpnp.isscalar(x1)
+    x2_is_scalar = dpnp.isscalar(x2)
+    x1_desc = dpnp.get_dpnp_descriptor(x1)
+    x2_desc = dpnp.get_dpnp_descriptor(x2)
+
+    if x1_desc and x2_desc and not kwargs:
+        if not x2_desc and not x2_is_scalar:
             pass
         elif x1_is_scalar and x2_is_scalar:
             pass
-        elif x1_is_dparray and x1.ndim == 0:
+        elif x1_desc and x1_desc.ndim == 0:
             pass
-        elif x2_is_dparray and x2.ndim == 0:
+        elif x2_desc and x2_desc.ndim == 0:
             pass
         elif dtype is not None:
             pass
@@ -1092,7 +1093,7 @@ def multiply(x1, x2, dtype=None, out=None, where=True, **kwargs):
         elif not where:
             pass
         else:
-            return dpnp_multiply(x1, x2, dtype=dtype, out=out, where=where)
+            return dpnp_multiply(x1_desc, x2_desc, dtype=dtype, out=out, where=where)
 
     return call_origin(numpy.multiply, x1, x2, dtype=dtype, out=out, where=where, **kwargs)
 

tests/skipped_tests.tbl

@@ -241,6 +241,8 @@ tests/third_party/cupy/creation_tests/test_ranges.py::TestMeshgrid_param_7_{copy
 tests/third_party/cupy/creation_tests/test_ranges.py::TestMeshgrid_param_7_{copy=True, indexing='ij', sparse=True}::test_meshgrid1
 tests/third_party/cupy/creation_tests/test_ranges.py::TestMeshgrid_param_7_{copy=True, indexing='ij', sparse=True}::test_meshgrid2
 tests/third_party/cupy/creation_tests/test_ranges.py::TestMeshgrid_param_7_{copy=True, indexing='ij', sparse=True}::test_meshgrid3
+tests/third_party/cupy/creation_tests/test_ranges.py::TestMgrid::test_mgrid3
+tests/third_party/cupy/creation_tests/test_ranges.py::TestOgrid::test_ogrid3
 tests/third_party/cupy/creation_tests/test_ranges.py::TestOgrid::test_ogrid4
 tests/third_party/cupy/creation_tests/test_ranges.py::TestOgrid::test_ogrid5
 tests/third_party/cupy/creation_tests/test_ranges.py::TestRanges::test_linspace_array_start_stop

tests/skipped_tests_gpu.tbl

@@ -252,6 +252,8 @@ tests/third_party/cupy/creation_tests/test_ranges.py::TestMeshgrid_param_7_{copy
 tests/third_party/cupy/creation_tests/test_ranges.py::TestMeshgrid_param_7_{copy=True, indexing='ij', sparse=True}::test_meshgrid1
 tests/third_party/cupy/creation_tests/test_ranges.py::TestMeshgrid_param_7_{copy=True, indexing='ij', sparse=True}::test_meshgrid2
 tests/third_party/cupy/creation_tests/test_ranges.py::TestMeshgrid_param_7_{copy=True, indexing='ij', sparse=True}::test_meshgrid3
+tests/third_party/cupy/creation_tests/test_ranges.py::TestMgrid::test_mgrid3
+tests/third_party/cupy/creation_tests/test_ranges.py::TestOgrid::test_ogrid3
 tests/third_party/cupy/creation_tests/test_ranges.py::TestOgrid::test_ogrid4
 tests/third_party/cupy/creation_tests/test_ranges.py::TestOgrid::test_ogrid5
 tests/third_party/cupy/creation_tests/test_ranges.py::TestRanges::test_linspace_array_start_stop