use sycl_adapter in krnl_linalg (#909)

Author: shssf

dpnp/backend/kernels/dpnp_krnl_linalg.cpp

@@ -29,6 +29,7 @@
 #include <dpnp_iface.hpp>
 #include "dpnp_fptr.hpp"
 #include "dpnp_utils.hpp"
+#include "dpnpc_memory_adapter.hpp"
 #include "queue_sycl.hpp"
 
 namespace mkl_blas = oneapi::mkl::blas::row_major;
@@ -39,8 +40,10 @@ void dpnp_cholesky_c(void* array1_in, void* result1, const size_t size, const si
 {
     cl::sycl::event event;
 
-    _DataType* in_array = reinterpret_cast<_DataType*>(array1_in);
-    _DataType* result = reinterpret_cast<_DataType*>(result1);
+    DPNPC_ptr_adapter<_DataType> input1_ptr(array1_in, size, true);
+    DPNPC_ptr_adapter<_DataType> result_ptr(result1, size, true, true);
+    _DataType* in_array = input1_ptr.get_ptr();
+    _DataType* result = result_ptr.get_ptr();
 
     size_t iters = size / (data_size * data_size);
 
@@ -97,8 +100,11 @@ void dpnp_cholesky_c(void* array1_in, void* result1, const size_t size, const si
 template <typename _DataType>
 void dpnp_det_c(void* array1_in, void* result1, size_t* shape, size_t ndim)
 {
-    _DataType* array_1 = reinterpret_cast<_DataType*>(array1_in);
-    _DataType* result = reinterpret_cast<_DataType*>(result1);
+    const size_t input_size = std::accumulate(shape, shape + ndim, 1, std::multiplies<size_t>());
+    if (!input_size)
+    {
+        return;
+    }
 
     size_t n = shape[ndim - 1];
     size_t size_out = 1;
@@ -110,6 +116,11 @@ void dpnp_det_c(void* array1_in, void* result1, size_t* shape, size_t ndim)
         }
     }
 
+    DPNPC_ptr_adapter<_DataType> input1_ptr(array1_in, input_size, true);
+    DPNPC_ptr_adapter<_DataType> result_ptr(result1, size_out, true, true);
+    _DataType* array_1 = input1_ptr.get_ptr();
+    _DataType* result = result_ptr.get_ptr();
+
     for (size_t i = 0; i < size_out; i++)
     {
         _DataType matrix[n][n];
@@ -194,8 +205,17 @@ template <typename _DataType, typename _ResultType>
 void dpnp_inv_c(void* array1_in, void* result1, size_t* shape, size_t ndim)
 {
     (void)ndim; // avoid warning unused variable
-    _DataType* array_1 = reinterpret_cast<_DataType*>(array1_in);
-    _ResultType* result = reinterpret_cast<_ResultType*>(result1);
+
+    const size_t input_size = std::accumulate(shape, shape + ndim, 1, std::multiplies<size_t>());
+    if (!input_size)
+    {
+        return;
+    }
+
+    DPNPC_ptr_adapter<_DataType> input1_ptr(array1_in, input_size, true);
+    DPNPC_ptr_adapter<_ResultType> result_ptr(result1, input_size, true, true);
+    _DataType* array_1 = input1_ptr.get_ptr();
+    _ResultType* result = result_ptr.get_ptr();
 
     size_t n = shape[0];
 
@@ -298,16 +318,21 @@ void dpnp_kron_c(void* array1_in,
                  size_t* res_shape,
                  size_t ndim)
 {
-    _DataType1* array1 = reinterpret_cast<_DataType1*>(array1_in);
-    _DataType2* array2 = reinterpret_cast<_DataType2*>(array2_in);
-    _ResultType* result = reinterpret_cast<_ResultType*>(result1);
-
-    size_t size = 1;
-    for (size_t i = 0; i < ndim; ++i)
+    const size_t input1_size = std::accumulate(in1_shape, in1_shape + ndim, 1, std::multiplies<size_t>());
+    const size_t input2_size = std::accumulate(in2_shape, in2_shape + ndim, 1, std::multiplies<size_t>());
+    const size_t result_size = std::accumulate(res_shape, res_shape + ndim, 1, std::multiplies<size_t>());
+    if (!(result_size && input1_size && input2_size))
     {
-        size *= res_shape[i];
+        return;
     }
 
+    DPNPC_ptr_adapter<_DataType1> input1_ptr(array1_in, input1_size);
+    DPNPC_ptr_adapter<_DataType2> input2_ptr(array2_in, input2_size);
+    DPNPC_ptr_adapter<_ResultType> result_ptr(result1, result_size);
+    _DataType1* array1 = input1_ptr.get_ptr();
+    _DataType2* array2 = input2_ptr.get_ptr();
+    _ResultType* result = result_ptr.get_ptr();
+
     size_t* _in1_shape = reinterpret_cast<size_t*>(dpnp_memory_alloc_c(ndim * sizeof(size_t)));
     size_t* _in2_shape = reinterpret_cast<size_t*>(dpnp_memory_alloc_c(ndim * sizeof(size_t)));
 
@@ -322,7 +347,7 @@ void dpnp_kron_c(void* array1_in,
     get_shape_offsets_inkernel<size_t>(in2_shape, ndim, in2_offsets);
     get_shape_offsets_inkernel<size_t>(res_shape, ndim, res_offsets);
 
-    cl::sycl::range<1> gws(size);
+    cl::sycl::range<1> gws(result_size);
     auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) {
         const size_t idx = global_id[0];
 
@@ -356,12 +381,18 @@ void dpnp_kron_c(void* array1_in,
 template <typename _DataType>
 void dpnp_matrix_rank_c(void* array1_in, void* result1, size_t* shape, size_t ndim)
 {
-    _DataType* array_1 = reinterpret_cast<_DataType*>(array1_in);
-    _DataType* result = reinterpret_cast<_DataType*>(result1);
+    const size_t input_size = std::accumulate(shape, shape + ndim, 1, std::multiplies<size_t>());
+    if (!input_size)
+    {
+        return;
+    }
+
+    DPNPC_ptr_adapter<_DataType> input1_ptr(array1_in, input_size);
+    DPNPC_ptr_adapter<_DataType> result_ptr(result1, 1);
+    _DataType* array_1 = input1_ptr.get_ptr();
+    _DataType* result = result_ptr.get_ptr();
 
     size_t elems = 1;
-    const _DataType init_val = 0;
-    dpnp_memory_memcpy_c(result, &init_val, sizeof(_DataType)); // result[0] = 0;
     if (ndim > 1)
     {
         elems = shape[0];
@@ -373,15 +404,18 @@ void dpnp_matrix_rank_c(void* array1_in, void* result1, size_t* shape, size_t nd
             }
         }
     }
+
+    _DataType acc = 0;
     for (size_t i = 0; i < elems; i++)
     {
         size_t ind = 0;
         for (size_t j = 0; j < ndim; j++)
         {
             ind += (shape[j] - 1) * i;
         }
-        result[0] += array_1[ind];
+        acc += array_1[ind];
     }
+    result[0] = acc;
 
     return;
 }
@@ -391,7 +425,8 @@ void dpnp_qr_c(void* array1_in, void* result1, void* result2, void* result3, siz
 {
     cl::sycl::event event;
 
-    _InputDT* in_array = reinterpret_cast<_InputDT*>(array1_in);
+    DPNPC_ptr_adapter<_InputDT> input1_ptr(array1_in, size_m * size_n, true);
+    _InputDT* in_array = input1_ptr.get_ptr();
 
     // math lib func overrides input
     _ComputeDT* in_a = reinterpret_cast<_ComputeDT*>(dpnp_memory_alloc_c(size_m * size_n * sizeof(_ComputeDT)));
@@ -400,13 +435,17 @@ void dpnp_qr_c(void* array1_in, void* result1, void* result2, void* result3, siz
     {
         for (size_t j = 0; j < size_n; ++j)
         {
+            // TODO transpose? use dpnp_transpose_c()
             in_a[j * size_m + i] = in_array[i * size_n + j];
         }
     }
 
-    _ComputeDT* res_q = reinterpret_cast<_ComputeDT*>(result1);
-    _ComputeDT* res_r = reinterpret_cast<_ComputeDT*>(result2);
-    _ComputeDT* tau = reinterpret_cast<_ComputeDT*>(result3);
+    DPNPC_ptr_adapter<_ComputeDT> result1_ptr(result1, size_m * size_m, true, true);
+    DPNPC_ptr_adapter<_ComputeDT> result2_ptr(result2, size_m * size_n, true, true);
+    DPNPC_ptr_adapter<_ComputeDT> result3_ptr(result3, std::min(size_m, size_n), true, true);
+    _ComputeDT* res_q = result1_ptr.get_ptr();
+    _ComputeDT* res_r = result2_ptr.get_ptr();
+    _ComputeDT* tau = result3_ptr.get_ptr();
 
     const std::int64_t lda = size_m;
 
@@ -487,18 +526,22 @@ void dpnp_svd_c(void* array1_in, void* result1, void* result2, void* result3, si
 {
     cl::sycl::event event;
 
-    _InputDT* in_array = reinterpret_cast<_InputDT*>(array1_in);
+    DPNPC_ptr_adapter<_InputDT> input1_ptr(array1_in, size_m * size_n, true); // TODO no need this if use dpnp_copy_to()
+    _InputDT* in_array = input1_ptr.get_ptr();
 
     // math lib gesvd func overrides input
     _ComputeDT* in_a = reinterpret_cast<_ComputeDT*>(dpnp_memory_alloc_c(size_m * size_n * sizeof(_ComputeDT)));
     for (size_t it = 0; it < size_m * size_n; ++it)
     {
-        in_a[it] = in_array[it];
+        in_a[it] = in_array[it]; // TODO Type conversion. memcpy can not be used directly. dpnp_copy_to() ?
     }
 
-    _ComputeDT* res_u = reinterpret_cast<_ComputeDT*>(result1);
-    _SVDT* res_s = reinterpret_cast<_SVDT*>(result2);
-    _ComputeDT* res_vt = reinterpret_cast<_ComputeDT*>(result3);
+    DPNPC_ptr_adapter<_ComputeDT> result1_ptr(result1, size_m * size_m, true, true);
+    DPNPC_ptr_adapter<_SVDT> result2_ptr(result2, std::min(size_m, size_n), true, true);
+    DPNPC_ptr_adapter<_ComputeDT> result3_ptr(result3, size_n * size_n, true, true);
+    _ComputeDT* res_u = result1_ptr.get_ptr();
+    _SVDT* res_s = result2_ptr.get_ptr();
+    _ComputeDT* res_vt = result3_ptr.get_ptr();
 
     const std::int64_t m = size_m;
     const std::int64_t n = size_n;