Merge pull request #867 from IntelPython/master

Merge master to gold

Author: Alexander-Makaryev

dpnp/backend/include/dpnp_iface.hpp

@@ -622,7 +622,7 @@ INP_DLLEXPORT void dpnp_initval_c(void* result1, void* value, size_t size);
  * @param [in]  shape   Shape of input array.
  * @param [in]  ndim    Number of elements in shape.
  */
-template <typename _DataType>
+template <typename _DataType, typename _ResultType>
 INP_DLLEXPORT void dpnp_inv_c(void* array1_in, void* result1, size_t* shape, size_t ndim);
 
 /**

dpnp/backend/kernels/dpnp_krnl_linalg.cpp

@@ -190,17 +190,17 @@ void dpnp_det_c(void* array1_in, void* result1, size_t* shape, size_t ndim)
     return;
 }
 
-template <typename _DataType>
+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);
-    _DataType* result = reinterpret_cast<_DataType*>(result1);
+    _ResultType* result = reinterpret_cast<_ResultType*>(result1);
 
     size_t n = shape[0];
 
-    _DataType a_arr[n][n];
-    _DataType e_arr[n][n];
+    _ResultType a_arr[n][n];
+    _ResultType e_arr[n][n];
 
     for (size_t i = 0; i < n; ++i)
     {
@@ -541,10 +541,10 @@ void func_map_init_linalg_func(func_map_t& fmap)
     fmap[DPNPFuncName::DPNP_FN_DET][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_det_c<float>};
     fmap[DPNPFuncName::DPNP_FN_DET][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_det_c<double>};
 
-    fmap[DPNPFuncName::DPNP_FN_INV][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_inv_c<int>};
-    fmap[DPNPFuncName::DPNP_FN_INV][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_inv_c<long>};
-    fmap[DPNPFuncName::DPNP_FN_INV][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_inv_c<float>};
-    fmap[DPNPFuncName::DPNP_FN_INV][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_inv_c<double>};
+    fmap[DPNPFuncName::DPNP_FN_INV][eft_INT][eft_INT] = {eft_DBL, (void*)dpnp_inv_c<int, double>};
+    fmap[DPNPFuncName::DPNP_FN_INV][eft_LNG][eft_LNG] = {eft_DBL, (void*)dpnp_inv_c<long, double>};
+    fmap[DPNPFuncName::DPNP_FN_INV][eft_FLT][eft_FLT] = {eft_DBL, (void*)dpnp_inv_c<float, double>};
+    fmap[DPNPFuncName::DPNP_FN_INV][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_inv_c<double, double>};
 
     fmap[DPNPFuncName::DPNP_FN_KRON][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_kron_c<int, int, int>};
     fmap[DPNPFuncName::DPNP_FN_KRON][eft_INT][eft_LNG] = {eft_LNG, (void*)dpnp_kron_c<int, long, long>};

dpnp/backend/kernels/dpnp_krnl_mathematical.cpp

@@ -346,50 +346,80 @@ void dpnp_remainder_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];
+        const _DataType_output input1_elem = (*input1_it)[i];
+        const _DataType_output input2_elem = (*input2_it)[i];
+        double fmod_res = cl::sycl::fmod((double)input1_elem, (double)input2_elem);
+        double add = fmod_res + input2_elem;
+        result[i] = cl::sycl::fmod(add, (double)input2_elem);
+
+    };
+    auto kernel_func = [&](cl::sycl::handler& cgh) {
+        cgh.parallel_for<class dpnp_remainder_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::fmod(DPNP_QUEUE, input1_size, input1, input2, result);
-        event.wait();
-        event = oneapi::mkl::vm::add(DPNP_QUEUE, input1_size, result, input2, result);
-        event.wait();
-        event = oneapi::mkl::vm::fmod(DPNP_QUEUE, input1_size, result, input2, 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::fmod(DPNP_QUEUE, input1_size, input1_data, input2_data, result);
+            event.wait();
+            event = oneapi::mkl::vm::add(DPNP_QUEUE, input1_size, result, input2_data, result);
+            event.wait();
+            event = oneapi::mkl::vm::fmod(DPNP_QUEUE, input1_size, result, input2_data, 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 fmod = cl::sycl::fmod((double)input_elem1, (double)input_elem2);
-                double add = fmod + input_elem2;
-                result[i] = cl::sycl::fmod(add, (double)input_elem2);
-            }
-        };
-
-        auto kernel_func = [&](cl::sycl::handler& cgh) {
-            cgh.parallel_for<class dpnp_remainder_c_kernel<_DataType_input1, _DataType_input2, _DataType_output>>(
-                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, typename _KernelNameSpecialization3>

dpnp/dparray.pyx

@@ -42,6 +42,7 @@ from dpnp.dpnp_iface import *
 # to avoid interference with Python internal functions
 from dpnp.dpnp_iface import sum as iface_sum
 from dpnp.dpnp_iface import prod as iface_prod
+from dpnp.dpnp_iface import get_dpnp_descriptor as iface_get_dpnp_descriptor
 
 from dpnp.dpnp_algo cimport *
 from dpnp.dpnp_iface_statistics import min, max  # TODO do the same as for iface_sum
@@ -386,7 +387,7 @@ cdef class dparray:
 
         if key_is_slice or key_has_slice:
             # fallback to numpy in case of slicing
-            return utils.nd2dp_array(utils.dp2nd_array(self)[key])
+            return nd2dp_array(dp2nd_array(self)[key])
 
         lin_idx = utils._get_linear_index(key, self.shape, self.ndim)
 
@@ -514,11 +515,12 @@ cdef class dparray:
                 if order == 'F':
                     return self.transpose().reshape(self.size)
 
-                return dpnp_flatten(self)
+                self_desc = iface_get_dpnp_descriptor(self)
+                return dpnp_flatten(self_desc).get_pyobj()
 
-        result = utils.dp2nd_array(self).flatten(order=order)
+        result = dp2nd_array(self).flatten(order=order)
 
-        return utils.nd2dp_array(result)
+        return nd2dp_array(result)
 
     def ravel(self, order='C'):
         """
@@ -799,11 +801,12 @@ cdef class dparray:
         elif self.dtype == numpy.complex64 or dtype == numpy.complex64:
             pass
         else:
-            return dpnp_astype(self, dtype)
+            self_desc = iface_get_dpnp_descriptor(self)
+            return dpnp_astype(self_desc, dtype).get_pyobj()
 
-        result = utils.dp2nd_array(self).astype(dtype=dtype, order=order, casting=casting, subok=subok, copy=copy)
+        result = dp2nd_array(self).astype(dtype=dtype, order=order, casting=casting, subok=subok, copy=copy)
 
-        return utils.nd2dp_array(result)
+        return nd2dp_array(result)
 
     def conj(self):
         """
@@ -1270,3 +1273,20 @@ cdef class dparray:
         """
 
         return asnumpy(self).tostring(order)
+
+
+def dp2nd_array(arr):
+    """Convert dparray to ndarray"""
+    return asnumpy(arr) if isinstance(arr, dparray) else arr
+
+
+def nd2dp_array(arr):
+    """Convert ndarray to dparray"""
+    if not isinstance(arr, numpy.ndarray):
+        return arr
+
+    result = dparray(arr.shape, dtype=arr.dtype)
+    for i in range(result.size):
+        result._setitem_scalar(i, arr.item(i))
+
+    return result

dpnp/dpnp_algo/dpnp_algo.pxd

@@ -27,7 +27,7 @@
 
 from libcpp.vector cimport vector
 from libcpp cimport bool as cpp_bool
-from dpnp.dparray cimport dparray
+
 from dpnp.dpnp_utils.dpnp_algo_utils cimport dpnp_descriptor
 
 
@@ -235,8 +235,8 @@ ctypedef void(*fptr_2in_1out_t)(void * , const void * , const size_t, const long
 ctypedef void(*fptr_blas_gemm_2in_1out_t)(void *, void * , void * , size_t, size_t, size_t)
 ctypedef void(*dpnp_reduction_c_t)(void *, const void * , const size_t*, const size_t, const long*, const size_t, const void * , const long*)
 
-cpdef dparray dpnp_astype(dparray array1, dtype_target)
-cpdef dparray dpnp_flatten(dparray array1)
+cpdef dpnp_descriptor dpnp_astype(dpnp_descriptor array1, dtype)
+cpdef dpnp_descriptor dpnp_flatten(dpnp_descriptor array1)
 
 
 """

dpnp/dpnp_algo/dpnp_algo.pyx

@@ -107,7 +107,7 @@ cpdef utils.dpnp_descriptor dpnp_array(object obj, object dtype=None):
     if not cpython.PySequence_Check(obj):
         raise TypeError(f"DPNP dpnp_array(): Unsupported non-sequence obj={type(obj)}")
 
-    obj_shape, elem_dtype = utils.get_shape_dtype(obj)
+    obj_shape, obj_dtype = utils.get_shape_dtype(obj)
     if dtype is not None:
         """ Set type from parameter. result might be empty array """
         result = utils_py.create_output_descriptor_py(obj_shape, dtype, None)
@@ -116,35 +116,37 @@ cpdef utils.dpnp_descriptor dpnp_array(object obj, object dtype=None):
             """ Empty object (ex. empty list) and no type provided """
             result = utils_py.create_output_descriptor_py(obj_shape, None, None)
         else:
-            result = utils_py.create_output_descriptor_py(obj_shape, elem_dtype, None)
+            result = utils_py.create_output_descriptor_py(obj_shape, obj_dtype, None)
 
-    utils.copy_values_to_dparray(result.get_pyobj(), obj)
+    utils.container_copy(result.get_pyobj(), obj)
 
     return result
 
 
-cpdef dparray dpnp_astype(dparray array1, dtype_target):
+cpdef utils.dpnp_descriptor dpnp_astype(utils.dpnp_descriptor array1, dtype):
     cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(array1.dtype)
-    cdef DPNPFuncType param2_type = dpnp_dtype_to_DPNPFuncType(dtype_target)
+    cdef DPNPFuncType param2_type = dpnp_dtype_to_DPNPFuncType(dtype)
 
     cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_ASTYPE, param1_type, param2_type)
 
-    result_type = dpnp_DPNPFuncType_to_dtype( < size_t > kernel_data.return_type)
-    cdef dparray result = dparray(array1.shape, dtype=result_type)
+    # ceate result array with type given by FPTR data
+    cdef shape_type_c result_shape = array1.shape
+    cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None)
 
     cdef fptr_dpnp_astype_t func = <fptr_dpnp_astype_t > kernel_data.ptr
     func(array1.get_data(), result.get_data(), array1.size)
 
     return result
 
 
-cpdef dparray dpnp_flatten(dparray array_):
+cpdef utils.dpnp_descriptor dpnp_flatten(utils.dpnp_descriptor array_):
     cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(array_.dtype)
 
     cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_FLATTEN, param1_type, param1_type)
 
-    result_type = dpnp_DPNPFuncType_to_dtype( < size_t > kernel_data.return_type)
-    cdef dparray result = dparray(array_.size, dtype=result_type)
+    # ceate result array with type given by FPTR data
+    cdef shape_type_c result_shape = (array_.size,)
+    cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None)
 
     cdef fptr_dpnp_flatten_t func = <fptr_dpnp_flatten_t > kernel_data.ptr
     func(array_.get_data(), result.get_data(), array_.size)

dpnp/dpnp_algo/dpnp_algo_indexing.pyx

@@ -65,10 +65,12 @@ ctypedef void(*custom_indexing_6in_func_ptr_t)(void *, void * , void * , const s
 ctypedef void(*fptr_dpnp_nonzero_t)(const void * , void * , const size_t * , const size_t , const size_t)
 
 
-cpdef dparray dpnp_choose(input, choices):
-    res_array = dparray(len(input), dtype=choices[0].dtype)
+cpdef utils.dpnp_descriptor dpnp_choose(object input, list choices):
+    cdef shape_type_c obj_shape = utils._object_to_tuple(len(input))
+    cdef utils.dpnp_descriptor res_array = utils_py.create_output_descriptor_py(obj_shape, choices[0].dtype, None)
+
     for i in range(len(input)):
-        res_array[i] = (choices[input[i]])[i]
+        res_array.get_pyobj()[i] = (choices[input[i]])[i]
     return res_array
 
 
@@ -169,12 +171,13 @@ cpdef object dpnp_indices(dimensions):
     return dpnp_result
 
 
-cpdef tuple dpnp_nonzero(dparray in_array1):
+cpdef tuple dpnp_nonzero(utils.dpnp_descriptor in_array1):
     cdef shape_type_c shape_arr = in_array1.shape
     res_count = in_array1.ndim
 
     # have to go through array one extra time to count size of result arrays
-    res_size = dpnp.count_nonzero(in_array1)
+    res_size_obj = dpnp_count_nonzero(in_array1)
+    cdef size_t res_size = dpnp.convert_single_elem_array_to_scalar(res_size_obj.get_pyobj())
 
     cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(in_array1.dtype)
 
@@ -183,12 +186,15 @@ cpdef tuple dpnp_nonzero(dparray in_array1):
     cdef fptr_dpnp_nonzero_t func = <fptr_dpnp_nonzero_t > kernel_data.ptr
 
     res_list = []
+    cdef utils.dpnp_descriptor res_arr
+    cdef shape_type_c result_shape
     for j in range(res_count):
-        res_arr = dparray((res_size, ), dtype=dpnp.int64)
+        result_shape = utils._object_to_tuple(res_size)
+        res_arr = utils_py.create_output_descriptor_py(result_shape, dpnp.int64, None)
 
         func(in_array1.get_data(), res_arr.get_data(), < size_t * > shape_arr.data(), in_array1.ndim, j)
 
-        res_list.append(res_arr)
+        res_list.append(res_arr.get_pyobj())
 
     result = utils._object_to_tuple(res_list)
 
@@ -263,20 +269,21 @@ cpdef dpnp_putmask(object arr, object mask, object values):
             arr[i] = values[i % values_size]
 
 
-cpdef object dpnp_select(condlist, choicelist, default):
-    size_ = condlist[0].size
-    res_array = dparray(size_, dtype=choicelist[0].dtype)
+cpdef utils.dpnp_descriptor dpnp_select(list condlist, list choicelist, default):
+    cdef size_t size_ = condlist[0].size
+    cdef utils.dpnp_descriptor res_array = utils_py.create_output_descriptor_py(condlist[0].shape, choicelist[0].dtype, None)
+
     pass_val = {a: default for a in range(size_)}
     for i in range(len(condlist)):
         for j in range(size_):
             if (condlist[i])[j]:
-                res_array[j] = (choicelist[i])[j]
+                res_array.get_pyobj()[j] = (choicelist[i])[j]
                 pass_val.pop(j)
 
     for ind, val in pass_val.items():
-        res_array[ind] = val
+        res_array.get_pyobj()[ind] = val
 
-    return res_array.reshape(condlist[0].shape)
+    return res_array
 
 
 cpdef utils.dpnp_descriptor dpnp_take(utils.dpnp_descriptor input, utils.dpnp_descriptor indices):
@@ -371,7 +378,7 @@ cpdef object dpnp_take_along_axis(object arr, object indices, int axis):
         return dpnp_result_array
 
     else:
-        result_array = dparray(shape_arr, dtype=res_type)
+        result_array = utils_py.create_output_descriptor_py(shape_arr, res_type, None).get_pyobj()
         for i in range(size_arr):
             ind = size_indices * (i // size_indices) + indices[i % size_indices]
             result_array[i] = arr[ind]