Merge master into impl_nan_to_num

Author: vlad-perevezentsev

.github/workflows/CODEOWNERS

@@ -0,0 +1 @@
+* @antonwolfy @npolina4 @vlad-perevezentsev @vtavana

.pre-commit-config.yaml

@@ -96,6 +96,7 @@ repos:
             [
             "-rn", # Only display messages
             "-sn", # Don't display the score
+            "--disable=c-extension-no-member",
             "--disable=import-error",
             "--disable=redefined-builtin",
             "--disable=unused-wildcard-import"

doc/reference/logic.rst

@@ -68,6 +68,7 @@ Comparison
    dpnp.allclose
    dpnp.isclose
    dpnp.array_equal
+   dpnp.array_equiv
    dpnp.greater
    dpnp.greater_equal
    dpnp.less

doc/reference/ufunc.rst

@@ -20,21 +20,27 @@ Math operations
    dpnp.add
    dpnp.subtract
    dpnp.multiply
+   dpnp.matmul
    dpnp.divide
    dpnp.logaddexp
    dpnp.logaddexp2
    dpnp.true_divide
    dpnp.floor_divide
    dpnp.negative
+   dpnp.positive
    dpnp.power
+   dpnp.float_power
    dpnp.remainder
    dpnp.mod
    dpnp.fmod
-   dpnp.abs
+   dpnp.divmod
    dpnp.absolute
    dpnp.fabs
    dpnp.rint
    dpnp.sign
+   dpnp.heaviside
+   dpnp.conj
+   dpnp.conjugate
    dpnp.exp
    dpnp.exp2
    dpnp.log
@@ -44,13 +50,24 @@ Math operations
    dpnp.log1p
    dpnp.proj
    dpnp.sqrt
-   dpnp.cbrt
    dpnp.square
+   dpnp.cbrt
    dpnp.reciprocal
    dpnp.rsqrt
    dpnp.gcd
    dpnp.lcm
 
+.. tip::
+
+   The optional output arguments can be used to help you save memory
+   for large calculations. If your arrays are large, complicated
+   expressions can take longer than absolutely necessary due to the
+   creation and (later) destruction of temporary calculation
+   spaces. For example, the expression ``G = A * B + C`` is equivalent to
+   ``T1 = A * B; G = T1 + C; del T1``. It will be more quickly executed
+   as ``G = A * B; add(G, C, G)`` which is the same as
+   ``G = A * B; G += C``.
+
 
 Trigonometric functions
 ~~~~~~~~~~~~~~~~~~~~~~~

dpnp/backend/extensions/lapack/CMakeLists.txt

@@ -30,6 +30,7 @@ set(_module_src
     ${CMAKE_CURRENT_SOURCE_DIR}/geqrf.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/geqrf_batch.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/gesv.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/gesv_batch.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/gesvd.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/getrf.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/getrf_batch.cpp

dpnp/backend/extensions/lapack/common_helpers.hpp

@@ -24,9 +24,11 @@
 //*****************************************************************************
 
 #pragma once
+#include <pybind11/pybind11.h>
+#include <sycl/sycl.hpp>
+
 #include <complex>
 #include <cstring>
-#include <pybind11/pybind11.h>
 #include <stdexcept>
 
 namespace dpnp::extensions::lapack::helper
@@ -63,4 +65,89 @@ inline bool check_zeros_shape(int ndim, const py::ssize_t *shape)
     }
     return src_nelems == 0;
 }
+
+// Allocate the memory for the pivot indices
+inline std::int64_t *alloc_ipiv(const std::int64_t n, sycl::queue &exec_q)
+{
+    std::int64_t *ipiv = nullptr;
+
+    try {
+        ipiv = sycl::malloc_device<std::int64_t>(n, exec_q);
+        if (!ipiv) {
+            throw std::runtime_error("Device allocation for ipiv failed");
+        }
+    } catch (sycl::exception const &e) {
+        if (ipiv != nullptr)
+            sycl::free(ipiv, exec_q);
+        throw std::runtime_error(
+            std::string(
+                "Unexpected SYCL exception caught during ipiv allocation: ") +
+            e.what());
+    }
+
+    return ipiv;
+}
+
+// Allocate the total memory for the total pivot indices with proper alignment
+// for batch implementations
+template <typename T>
+inline std::int64_t *alloc_ipiv_batch(const std::int64_t n,
+                                      std::int64_t n_linear_streams,
+                                      sycl::queue &exec_q)
+{
+    // Get padding size to ensure memory allocations are aligned to 256 bytes
+    // for better performance
+    const std::int64_t padding = 256 / sizeof(T);
+
+    // Calculate the total size needed for the pivot indices array for all
+    // linear streams with proper alignment
+    size_t alloc_ipiv_size = round_up_mult(n_linear_streams * n, padding);
+
+    return alloc_ipiv(alloc_ipiv_size, exec_q);
+}
+
+// Allocate the memory for the scratchpad
+template <typename T>
+inline T *alloc_scratchpad(std::int64_t scratchpad_size, sycl::queue &exec_q)
+{
+    T *scratchpad = nullptr;
+
+    try {
+        if (scratchpad_size > 0) {
+            scratchpad = sycl::malloc_device<T>(scratchpad_size, exec_q);
+            if (!scratchpad) {
+                throw std::runtime_error(
+                    "Device allocation for scratchpad failed");
+            }
+        }
+    } catch (sycl::exception const &e) {
+        if (scratchpad != nullptr) {
+            sycl::free(scratchpad, exec_q);
+        }
+        throw std::runtime_error(std::string("Unexpected SYCL exception caught "
+                                             "during scratchpad allocation: ") +
+                                 e.what());
+    }
+
+    return scratchpad;
+}
+
+// Allocate the total scratchpad memory with proper alignment for batch
+// implementations
+template <typename T>
+inline T *alloc_scratchpad_batch(std::int64_t scratchpad_size,
+                                 std::int64_t n_linear_streams,
+                                 sycl::queue &exec_q)
+{
+    // Get padding size to ensure memory allocations are aligned to 256 bytes
+    // for better performance
+    const std::int64_t padding = 256 / sizeof(T);
+
+    // Calculate the total scratchpad memory size needed for all linear
+    // streams with proper alignment
+    const size_t alloc_scratch_size =
+        round_up_mult(n_linear_streams * scratchpad_size, padding);
+
+    return alloc_scratchpad<T>(alloc_scratch_size, exec_q);
+}
 } // namespace dpnp::extensions::lapack::helper

dpnp/backend/extensions/lapack/evd_batch_common.hpp

@@ -119,34 +119,4 @@ std::pair<sycl::event, sycl::event>
 
     return std::make_pair(ht_ev, evd_batch_ev);
 }
-
-template <typename T>
-inline T *alloc_scratchpad(std::int64_t scratchpad_size,
-                           std::int64_t n_linear_streams,
-                           sycl::queue &exec_q)
-{
-    // Get padding size to ensure memory allocations are aligned to 256 bytes
-    // for better performance
-    const std::int64_t padding = 256 / sizeof(T);
-
-    if (scratchpad_size <= 0) {
-        throw std::runtime_error(
-            "Invalid scratchpad size: must be greater than zero."
-            " Calculated scratchpad size: " +
-            std::to_string(scratchpad_size));
-    }
-
-    // Calculate the total scratchpad memory size needed for all linear
-    // streams with proper alignment
-    const size_t alloc_scratch_size =
-        helper::round_up_mult(n_linear_streams * scratchpad_size, padding);
-
-    // Allocate memory for the total scratchpad
-    T *scratchpad = sycl::malloc_device<T>(alloc_scratch_size, exec_q);
-    if (!scratchpad) {
-        throw std::runtime_error("Device allocation for scratchpad failed");
-    }
-
-    return scratchpad;
-}
 } // namespace dpnp::extensions::lapack::evd