Add implementation of dpnp.divmod

dpnp/backend/extensions/elementwise_functions/elementwise_functions.hpp

@@ -556,7 +556,7 @@ std::pair<sycl::event, sycl::event> py_binary_ufunc(
     const dpctl::tensor::usm_ndarray &src2,
     const dpctl::tensor::usm_ndarray &dst, // dst = op(src1, src2), elementwise
     sycl::queue &exec_q,
-    const std::vector<sycl::event> depends,
+    const std::vector<sycl::event> &depends,
     //
     const output_typesT &output_type_table,
     const contig_dispatchT &contig_dispatch_table,
@@ -870,7 +870,7 @@ std::pair<sycl::event, sycl::event>
                                 const dpctl::tensor::usm_ndarray &dst1,
                                 const dpctl::tensor::usm_ndarray &dst2,
                                 sycl::queue &exec_q,
-                                const std::vector<sycl::event> depends,
+                                const std::vector<sycl::event> &depends,
                                 //
                                 const output_typesT &output_types_table,
                                 const contig_dispatchT &contig_dispatch_table,
@@ -952,8 +952,10 @@ std::pair<sycl::event, sycl::event>
     auto const &same_logical_tensors =
         dpctl::tensor::overlap::SameLogicalTensors();
     if ((overlap(src1, dst1) && !same_logical_tensors(src1, dst1)) ||
+        (overlap(src1, dst2) && !same_logical_tensors(src1, dst2)) ||
         (overlap(src2, dst1) && !same_logical_tensors(src2, dst1)) ||
-        (overlap(dst1, dst2) && !same_logical_tensors(dst1, dst2)))
+        (overlap(src2, dst2) && !same_logical_tensors(src2, dst2)) ||
+        (overlap(dst1, dst2)))
     {
         throw py::value_error("Arrays index overlapping segments of memory");
     }
@@ -1142,7 +1144,7 @@ std::pair<sycl::event, sycl::event>
     py_binary_inplace_ufunc(const dpctl::tensor::usm_ndarray &lhs,
                             const dpctl::tensor::usm_ndarray &rhs,
                             sycl::queue &exec_q,
-                            const std::vector<sycl::event> depends,
+                            const std::vector<sycl::event> &depends,
                             //
                             const output_typesT &output_type_table,
                             const contig_dispatchT &contig_dispatch_table,

dpnp/backend/kernels/elementwise_functions/divmod.hpp

@@ -28,6 +28,8 @@
 
 #pragma once
 
+#include <type_traits>
+
 #include <sycl/sycl.hpp>
 
 namespace dpnp::kernels::divmod

dpnp/dpnp_algo/dpnp_elementwise_common.py

@@ -510,7 +510,7 @@ class DPNPBinaryFunc(BinaryElementwiseFunc):
 
     Args:
     name : {str}
-        Name of the unary function
+        Name of the binary function
     result_type_resovle_fn : {callable}
         Function that takes dtype of the input and returns the dtype of
         the result if the implementation functions supports it, or
@@ -530,7 +530,7 @@ class DPNPBinaryFunc(BinaryElementwiseFunc):
         corresponds to computational tasks associated with function
         evaluation.
     docs : {str}
-        Documentation string for the unary function.
+        Documentation string for the binary function.
     mkl_fn_to_call : {None, str}
         Check input arguments to answer if function from OneMKL VM library
         can be used.
@@ -828,17 +828,17 @@ def __call__(self, *args, **kwargs):
 
 class DPNPBinaryTwoOutputsFunc(BinaryElementwiseFunc):
     """
-    Class that implements unary element-wise functions with two output arrays.
+    Class that implements binary element-wise functions with two output arrays.
 
     Parameters
     ----------
     name : {str}
-        Name of the unary function
+        Name of the binary function
     result_type_resolver_fn : {callable}
         Function that takes dtype of the input and returns the dtype of
         the result if the implementation functions supports it, or
         returns `None` otherwise.
-    unary_dp_impl_fn : {callable}
+    binary_dp_impl_fn : {callable}
         Data-parallel implementation function with signature
         `impl_fn(src: usm_ndarray, dst: usm_ndarray,
             sycl_queue: SyclQueue, depends: Optional[List[SyclEvent]])`
@@ -852,7 +852,7 @@ class DPNPBinaryTwoOutputsFunc(BinaryElementwiseFunc):
         computational tasks complete execution, while the second event
         corresponds to computational tasks associated with function evaluation.
     docs : {str}
-        Documentation string for the unary function.
+        Documentation string for the binary function.
 
     """
 
@@ -1020,6 +1020,14 @@ def __call__(
             if not res.flags.writable:
                 raise ValueError("output array is read-only")
 
+            for other_out in out[:i]:
+                if other_out is None:
+                    continue
+
+                other_out = dpnp.get_usm_ndarray(other_out)
+                if dti._array_overlap(res, other_out):
+                    raise ValueError("Output arrays cannot overlap")
+
             if res.shape != res_shape:
                 raise ValueError(
                     "The shape of input and output arrays are inconsistent. "
@@ -1042,19 +1050,20 @@ def __call__(
                 # Allocate a temporary buffer with the required dtype
                 out[i] = dpt.empty_like(res, dtype=res_dt)
             else:
-                for x, dt in zip([x1, x2], buf_dts):
-                    if dpnp.isscalar(x):
-                        pass
-                    elif dt is not None:
-                        pass
-                    elif not dti._array_overlap(x, res):
-                        pass
-                    elif dti._same_logical_tensors(x, res):
-                        pass
-
-                    # Allocate a temporary buffer to avoid memory overlapping.
-                    # Note if `dt` is not None, a temporary copy of `x` will be
-                    # created, so the array overlap check isn't needed.
+                # If `dt` is not None, a temporary copy of `x` will be created,
+                # so the array overlap check isn't needed.
+                x_to_check = [
+                    x
+                    for x, dt in zip([x1, x2], buf_dts)
+                    if not dpnp.isscalar(x) and dt is None
+                ]
+
+                if any(
+                    dti._array_overlap(x, res)
+                    and not dti._same_logical_tensors(x, res)
+                    for x in x_to_check
+                ):
+                    # allocate a temporary buffer to avoid memory overlapping
                     out[i] = dpt.empty_like(res)
 
         x1 = dpnp.as_usm_ndarray(x1, dtype=x1_dt, sycl_queue=exec_q)

dpnp/dpnp_iface_mathematical.py

@@ -1581,19 +1581,19 @@ def diff(a, n=1, axis=-1, prepend=None, append=None):
     Divisor input array, expected to have a real-valued floating-point data
     type.
 out1 : {None, dpnp.ndarray, usm_ndarray}, optional
-    Output array for the quotient to populate. Array must have the same shape
-    as `x` and the expected data type.
+    Output array for the quotient to populate. Array must have a shape that
+    the inputs broadcast to and the expected data type.
 
     Default: ``None``.
 out2 : {None, dpnp.ndarray, usm_ndarray}, optional
-    Output array for the remainder to populate. Array must have the same shape
-    as `x` and the expected data type.
+    Output array for the remainder to populate. Array must have a shape that
+    the inputs broadcast to and the expected data type.
 
     Default: ``None``.
 out : tuple of None, dpnp.ndarray, or usm_ndarray, optional
     A location into which the result is stored. If provided, it must be a tuple
     and have length equal to the number of outputs. Each provided array must
-    have the same shape as `x` and the expected data type.
+    have a shape that the inputs broadcast to and the expected data type.
     It is prohibited to pass output arrays through `out` keyword when either
     `out1` or `out2` is passed.
 

dpnp/tests/test_binary_two_outputs_ufuncs.py

@@ -217,6 +217,16 @@ def test_out1_overlap(self, func, dt):
         _ = getattr(numpy, func)(a[size::], 1, a[::2])
         assert_array_equal(ia, a)
 
+    def test_out_arrays_overlap(self, func):
+        a = dpnp.arange(7)
+        out = dpnp.zeros_like(a)
+
+        with pytest.raises(
+            ValueError,
+            match="Output arrays cannot overlap",
+        ):
+            _ = getattr(dpnp, func)(a, 2, out=(out, out))
+
     def test_out_scalar_input(self, func):
         a = generate_random_numpy_array((3, 7), low=1, dtype=int)
         out = numpy.zeros_like(a)