add broadcasting for floor_divide func (#702)

* add broadcasting for floor_divide

Author: KsanaKozlova

dpnp/backend/kernels/dpnp_krnl_mathematical.cpp

@@ -29,6 +29,7 @@
 
 #include <dpnp_iface.hpp>
 #include "dpnp_fptr.hpp"
+#include "dpnp_iterator.hpp"
 #include "dpnp_utils.hpp"
 #include "queue_sycl.hpp"
 
@@ -226,47 +227,78 @@ void dpnp_floor_divide_c(void* result_out,
                          const size_t input2_shape_ndim,
                          const size_t* where)
 {
-    (void)input1_shape;
-    (void)input1_shape_ndim;
-    (void)input2_size;
-    (void)input2_shape;
-    (void)input2_shape_ndim;
     (void)where;
 
-    cl::sycl::event event;
-    _DataType_input1* input1 = reinterpret_cast<_DataType_input1*>(const_cast<void*>(input1_in));
-    _DataType_input2* input2 = reinterpret_cast<_DataType_input2*>(const_cast<void*>(input2_in));
+    if (!input1_size || !input2_size)
+    {
+        return;
+    }
+
+    _DataType_input1* input1_data = reinterpret_cast<_DataType_input1*>(const_cast<void*>(input1_in));
+    _DataType_input2* input2_data = reinterpret_cast<_DataType_input2*>(const_cast<void*>(input2_in));
     _DataType_output* result = reinterpret_cast<_DataType_output*>(result_out);
 
-    if constexpr ((std::is_same<_DataType_input1, double>::value || std::is_same<_DataType_input1, float>::value) &&
-                  std::is_same<_DataType_input2, _DataType_input1>::value)
+    std::vector<size_t> result_shape = get_result_shape(input1_shape, input1_shape_ndim,
+                                                            input2_shape, input2_shape_ndim); 
+
+    DPNPC_id<_DataType_input1>* input1_it;
+    const size_t input1_it_size_in_bytes = sizeof(DPNPC_id<_DataType_input1>);
+    input1_it = reinterpret_cast<DPNPC_id<_DataType_input1>*>(dpnp_memory_alloc_c(input1_it_size_in_bytes));
+    new (input1_it) DPNPC_id<_DataType_input1>(input1_data, input1_shape, input1_shape_ndim);
+
+    input1_it->broadcast_to_shape(result_shape);
+
+    DPNPC_id<_DataType_input2>* input2_it;
+    const size_t input2_it_size_in_bytes = sizeof(DPNPC_id<_DataType_input2>);
+    input2_it = reinterpret_cast<DPNPC_id<_DataType_input2>*>(dpnp_memory_alloc_c(input2_it_size_in_bytes));
+    new (input2_it) DPNPC_id<_DataType_input2>(input2_data, input2_shape, input2_shape_ndim);
+
+    input2_it->broadcast_to_shape(result_shape);
+    
+    const size_t result_size = input1_it->get_output_size();
+
+
+    cl::sycl::range<1> gws(result_size); 
+    auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) {
+        const size_t i = global_id[0]; /* for (size_t i = 0; i < result_size; ++i) */
+        const _DataType_output input1_elem = (*input1_it)[i];
+        const _DataType_output input2_elem = (*input2_it)[i];
+
+        double div = (double)input1_elem / (double)input2_elem;
+        result[i] = static_cast<_DataType_output>(cl::sycl::floor(div));
+    };
+    auto kernel_func = [&](cl::sycl::handler& cgh) {
+        cgh.parallel_for<class dpnp_floor_divide_c_kernel<_DataType_output, _DataType_input1, _DataType_input2>>(
+            gws, kernel_parallel_for_func);
+    };
+
+    cl::sycl::event event;
+
+    if (input1_size == input2_size)
     {
-        event = oneapi::mkl::vm::div(DPNP_QUEUE, input1_size, input1, input2, result);
-        event.wait();
-        event = oneapi::mkl::vm::floor(DPNP_QUEUE, input1_size, result, result);
+        if constexpr ((std::is_same<_DataType_input1, double>::value ||
+                        std::is_same<_DataType_input1, float>::value) &&
+                        std::is_same<_DataType_input2, _DataType_input1>::value)
+        {
+            event = oneapi::mkl::vm::div(DPNP_QUEUE, input1_size, input1_data, input2_data, result);
+            event.wait();
+            event = oneapi::mkl::vm::floor(DPNP_QUEUE, input1_size, result, result);
+        }
+        else
+        {
+            event = DPNP_QUEUE.submit(kernel_func);
+        }
     }
     else
     {
-        cl::sycl::range<1> gws(input1_size);
-        auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) {
-            size_t i = global_id[0]; /*for (size_t i = 0; i < size; ++i)*/
-            {
-                _DataType_input1 input_elem1 = input1[i];
-                _DataType_input2 input_elem2 = input2[i];
-                double div = (double)input_elem1 / (double)input_elem2;
-                result[i] = static_cast<_DataType_output>(cl::sycl::floor(div));
-            }
-        };
-
-        auto kernel_func = [&](cl::sycl::handler& cgh) {
-            cgh.parallel_for<class dpnp_floor_divide_c_kernel<_DataType_output, _DataType_input1, _DataType_input2>>(
-                gws, kernel_parallel_for_func);
-        };
-
         event = DPNP_QUEUE.submit(kernel_func);
     }
 
     event.wait();
+
+    input1_it->~DPNPC_id();
+    input2_it->~DPNPC_id();
+
 }
 
 template <typename _KernelNameSpecialization1, typename _KernelNameSpecialization2>

dpnp/dpnp_iface_mathematical.py

@@ -681,15 +681,16 @@ def floor(x1, **kwargs):
     return call_origin(numpy.floor, x1, **kwargs)
 
 
-def floor_divide(x1, x2, **kwargs):
+def floor_divide(x1, x2, dtype=None, out=None, where=True, **kwargs):
     """
     Compute the largest integer smaller or equal to the division of the inputs.
 
     For full documentation refer to :obj:`numpy.floor_divide`.
 
     Limitations
     -----------
-        Parameters ``x1`` and ``x2`` are supported as :obj:`dpnp.ndarray`.
+        Parameters ``x1`` and ``x2`` are supported as either :obj:`dpnp.ndarray` or scalar.
+        Parameters ``dtype``, ``out`` and ``where`` are supported with their default values.
         Keyword arguments ``kwargs`` are currently unsupported.
         Otherwise the functions will be executed sequentially on CPU.
         Input array data types are limited by supported DPNP :ref:`Data types`.
@@ -710,13 +711,40 @@ def floor_divide(x1, x2, **kwargs):
 
     """
 
-    is_x1_dparray = isinstance(x1, dparray)
-    is_x2_dparray = isinstance(x2, dparray)
-
-    if not use_origin_backend(x1) and is_x1_dparray and is_x2_dparray and not kwargs:
-        return dpnp_floor_divide(x1, x2)
+    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:
+            pass
+        elif x1_is_scalar and x2_is_scalar:
+            pass
+        elif x1_is_dparray and x1.ndim == 0:
+            pass
+        elif x2_is_dparray and x2.ndim == 0:
+            pass
+        elif x2_is_scalar and x2 == 0:
+            pass
+        elif x1_is_dparray and x2_is_dparray and x1.size != x2.size:
+            pass
+        elif x1_is_dparray and x2_is_dparray and x1.shape != x2.shape:
+            pass
+        elif out is not None and not isinstance(out, dparray):
+            pass
+        elif dtype is not None:
+            pass
+        elif out is not None:
+            pass
+        elif not where:
+            pass
+        elif x1_is_scalar and x2.ndim > 1:
+            pass
+        else:
+            return dpnp_floor_divide(x1, x2, out=out, where=where, dtype=dtype)
 
-    return call_origin(numpy.floor_divide, x1, x2, **kwargs)
+    return call_origin(numpy.floor_divide, x1, x2, out=out, where=where, dtype=dtype, **kwargs)
 
 
 def fmax(*args, **kwargs):

tests/test_mathematical.py

@@ -105,6 +105,9 @@ def test_divide(self, dtype, lhs, rhs):
     def test_fmod(self, dtype, lhs, rhs):
         self._test_mathematical('fmod', dtype, lhs, rhs)
 
+    def test_floor_divide(self, dtype, lhs, rhs):
+        self._test_mathematical('floor_divide', dtype, lhs, rhs)
+
     def test_hypot(self, dtype, lhs, rhs):
         self._test_mathematical('hypot', dtype, lhs, rhs)