Merge branch 'master' into build_target_cuda

Author: npolina4

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