diff --git a/dpnp/backend/include/dpnp_iface.hpp b/dpnp/backend/include/dpnp_iface.hpp index 0fc5595041c6..4efea15a38b4 100644 --- a/dpnp/backend/include/dpnp_iface.hpp +++ b/dpnp/backend/include/dpnp_iface.hpp @@ -205,38 +205,6 @@ INP_DLLEXPORT void dpnp_nanvar_c(void *array, const size_t result_size, size_t size); -/** - * @ingroup BACKEND_API - * @brief Return the indices of the elements that are non-zero. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array1 Input array. - * @param [out] result1 Output array. - * @param [in] result_size Output array size. - * @param [in] shape Shape of input array. - * @param [in] ndim Number of elements in shape. - * @param [in] j Number input array. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_nonzero_c(DPCTLSyclQueueRef q_ref, - const void *array1, - void *result1, - const size_t result_size, - const shape_elem_type *shape, - const size_t ndim, - const size_t j, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_nonzero_c(const void *array1, - void *result1, - const size_t result_size, - const shape_elem_type *shape, - const size_t ndim, - const size_t j); - /** * @ingroup BACKEND_API * @brief Custom implementation of dot function @@ -448,35 +416,6 @@ INP_DLLEXPORT void dpnp_partition_c(void *array, const shape_elem_type *shape, const size_t ndim); -/** - * @ingroup BACKEND_API - * @brief Place of array elements - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] arr Input array. - * @param [in] mask Mask array. - * @param [in] vals Vals array. - * @param [in] arr_size Number of input elements in `arr`. - * @param [in] vals_size Number of input elements in `vals`. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_place_c(DPCTLSyclQueueRef q_ref, - void *arr, - long *mask, - void *vals, - const size_t arr_size, - const size_t vals_size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_place_c(void *arr, - long *mask, - void *vals, - const size_t arr_size, - const size_t vals_size); - /** * @ingroup BACKEND_API * @brief Compute Product of input array elements. @@ -523,78 +462,6 @@ INP_DLLEXPORT void dpnp_prod_c(void *result_out, const void *initial, const long *where); -/** - * @ingroup BACKEND_API - * @brief Replaces specified elements of an array with given values. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array Input array. - * @param [in] ind Target indices, interpreted as integers. - * @param [in] v Values to place in array at target indices. - * @param [in] size Number of input elements in `array`. - * @param [in] size_ind Number of input elements in `ind`. - * @param [in] size_v Number of input elements in `v`. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_put_c(DPCTLSyclQueueRef q_ref, - void *array, - void *ind, - void *v, - const size_t size, - const size_t size_ind, - const size_t size_v, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_put_c(void *array, - void *ind, - void *v, - const size_t size, - const size_t size_ind, - const size_t size_v); - -/** - * @ingroup BACKEND_API - * @brief Put values into the destination array by matching 1d index and data - * slices. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] arr_in Input array. - * @param [in] indices_in Indices to change along each 1d slice of - * arr. - * @param [in] values_in Values to insert at those indices. - * @param [in] axis The axis to take 1d slices along. - * @param [in] shape Shape of input array. - * @param [in] ndim Number of input array dimensions. - * @param [in] size_indices Size of indices. - * @param [in] values_size Size of values. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_put_along_axis_c(DPCTLSyclQueueRef q_ref, - void *arr_in, - long *indices_in, - void *values_in, - size_t axis, - const shape_elem_type *shape, - size_t ndim, - size_t size_indices, - size_t values_size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_put_along_axis_c(void *arr_in, - long *indices_in, - void *values_in, - size_t axis, - const shape_elem_type *shape, - size_t ndim, - size_t size_indices, - size_t values_size); - /** * @ingroup BACKEND_API @@ -776,42 +643,6 @@ INP_DLLEXPORT void dpnp_choose_c(void *result1, size_t choices_size, size_t choice_size); -/** - * @ingroup BACKEND_API - * @brief math library implementation of diagonal function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array1_in Input array with data. - * @param [in] input1_size Input1 data size. - * @param [out] result1 Output array. - * @param [in] offset Offset of the diagonal from the main - * diagonal. - * @param [in] shape Shape of input array. - * @param [in] res_shape Shape of output array. - * @param [in] res_ndim Number of elements in shape. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_diagonal_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - const size_t input1_size, - void *result1, - const size_t offset, - shape_elem_type *shape, - shape_elem_type *res_shape, - const size_t res_ndim, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_diagonal_c(void *array1_in, - const size_t input1_size, - void *result1, - const size_t offset, - shape_elem_type *shape, - shape_elem_type *res_shape, - const size_t res_ndim); - /** * @ingroup BACKEND_API * @brief implementation of creating filled with value array function @@ -1044,35 +875,6 @@ INP_DLLEXPORT void dpnp_std_c(void *array, size_t naxis, size_t ddof); -/** - * @ingroup BACKEND_API - * @brief math library implementation of take function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array Input array with data. - * @param [in] array1_size Input array size. - * @param [in] indices Input array with indices. - * @param [out] result Output array. - * @param [in] size Number of elements in the input array. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_take_c(DPCTLSyclQueueRef q_ref, - void *array, - const size_t array1_size, - void *indices, - void *result, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_take_c(void *array, - const size_t array1_size, - void *indices, - void *result, - size_t size); - /** * @ingroup BACKEND_API * @brief math library implementation of var function @@ -1183,32 +985,6 @@ INP_DLLEXPORT void dpnp_var_c(void *array, #include -/** - * @ingroup BACKEND_API - * @brief fill_diagonal function. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array1_in Input array. - * @param [in] val Value to write on the diagonal. - * @param [in] shape Input shape. - * @param [in] ndim Number of elements in shape. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_fill_diagonal_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *val, - shape_elem_type *shape, - const size_t ndim, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_fill_diagonal_c(void *array1_in, - void *val, - shape_elem_type *shape, - const size_t ndim); - /** * @ingroup BACKEND_API * @brief modf function. diff --git a/dpnp/backend/include/dpnp_iface_fptr.hpp b/dpnp/backend/include/dpnp_iface_fptr.hpp index d62e5998583e..aaaf90c27bb0 100644 --- a/dpnp/backend/include/dpnp_iface_fptr.hpp +++ b/dpnp/backend/include/dpnp_iface_fptr.hpp @@ -81,22 +81,20 @@ enum class DPNPFuncName : size_t DPNP_FN_DEGREES, /**< Used in numpy.degrees() impl */ DPNP_FN_DEGREES_EXT, /**< Used in numpy.degrees() impl, requires extra parameters */ - DPNP_FN_DIAGONAL, /**< Used in numpy.diagonal() impl */ DPNP_FN_DOT, /**< Used in numpy.dot() impl */ DPNP_FN_DOT_EXT, /**< Used in numpy.dot() impl, requires extra parameters */ DPNP_FN_EDIFF1D, /**< Used in numpy.ediff1d() impl */ - DPNP_FN_EDIFF1D_EXT, /**< Used in numpy.ediff1d() impl, requires extra - parameters */ - DPNP_FN_ERF, /**< Used in scipy.special.erf impl */ - DPNP_FN_ERF_EXT, /**< Used in scipy.special.erf impl, requires extra - parameters */ - DPNP_FN_FFT_FFT, /**< Used in numpy.fft.fft() impl */ - DPNP_FN_FFT_FFT_EXT, /**< Used in numpy.fft.fft() impl, requires extra - parameters */ - DPNP_FN_FFT_RFFT, /**< Used in numpy.fft.rfft() impl */ - DPNP_FN_FFT_RFFT_EXT, /**< Used in numpy.fft.rfft() impl, requires extra - parameters */ - DPNP_FN_FILL_DIAGONAL, /**< Used in numpy.fill_diagonal() impl */ + DPNP_FN_EDIFF1D_EXT, /**< Used in numpy.ediff1d() impl, requires extra + parameters */ + DPNP_FN_ERF, /**< Used in scipy.special.erf impl */ + DPNP_FN_ERF_EXT, /**< Used in scipy.special.erf impl, requires extra + parameters */ + DPNP_FN_FFT_FFT, /**< Used in numpy.fft.fft() impl */ + DPNP_FN_FFT_FFT_EXT, /**< Used in numpy.fft.fft() impl, requires extra + parameters */ + DPNP_FN_FFT_RFFT, /**< Used in numpy.fft.rfft() impl */ + DPNP_FN_FFT_RFFT_EXT, /**< Used in numpy.fft.rfft() impl, requires extra + parameters */ DPNP_FN_INITVAL, /**< Used in numpy ones, ones_like, zeros, zeros_like impls */ DPNP_FN_INITVAL_EXT, /**< Used in numpy ones, ones_like, zeros, zeros_like @@ -116,23 +114,19 @@ enum class DPNPFuncName : size_t */ DPNP_FN_MULTIPLY, /**< Used in numpy.multiply() impl */ DPNP_FN_NANVAR, /**< Used in numpy.nanvar() impl */ - DPNP_FN_NONZERO, /**< Used in numpy.nonzero() impl */ DPNP_FN_ONES, /**< Used in numpy.ones() impl */ DPNP_FN_ONES_LIKE, /**< Used in numpy.ones_like() impl */ DPNP_FN_PARTITION, /**< Used in numpy.partition() impl */ - DPNP_FN_PARTITION_EXT, /**< Used in numpy.partition() impl, requires extra - parameters */ - DPNP_FN_PLACE, /**< Used in numpy.place() impl */ - DPNP_FN_PROD, /**< Used in numpy.prod() impl */ - DPNP_FN_PUT, /**< Used in numpy.put() impl */ - DPNP_FN_PUT_ALONG_AXIS, /**< Used in numpy.put_along_axis() impl */ - DPNP_FN_RADIANS, /**< Used in numpy.radians() impl */ - DPNP_FN_RADIANS_EXT, /**< Used in numpy.radians() impl, requires extra - parameters */ - DPNP_FN_RNG_BETA, /**< Used in numpy.random.beta() impl */ - DPNP_FN_RNG_BETA_EXT, /**< Used in numpy.random.beta() impl, requires extra - parameters */ - DPNP_FN_RNG_BINOMIAL, /**< Used in numpy.random.binomial() impl */ + DPNP_FN_PARTITION_EXT, /**< Used in numpy.partition() impl, requires extra + parameters */ + DPNP_FN_PROD, /**< Used in numpy.prod() impl */ + DPNP_FN_RADIANS, /**< Used in numpy.radians() impl */ + DPNP_FN_RADIANS_EXT, /**< Used in numpy.radians() impl, requires extra + parameters */ + DPNP_FN_RNG_BETA, /**< Used in numpy.random.beta() impl */ + DPNP_FN_RNG_BETA_EXT, /**< Used in numpy.random.beta() impl, requires extra + parameters */ + DPNP_FN_RNG_BINOMIAL, /**< Used in numpy.random.binomial() impl */ DPNP_FN_RNG_BINOMIAL_EXT, /**< Used in numpy.random.binomial() impl, requires extra parameters */ DPNP_FN_RNG_CHISQUARE, /**< Used in numpy.random.chisquare() impl */ @@ -253,7 +247,6 @@ enum class DPNPFuncName : size_t */ DPNP_FN_STD, /**< Used in numpy.std() impl */ DPNP_FN_SUM, /**< Used in numpy.sum() impl */ - DPNP_FN_TAKE, /**< Used in numpy.take() impl */ DPNP_FN_VAR, /**< Used in numpy.var() impl */ DPNP_FN_ZEROS, /**< Used in numpy.zeros() impl */ DPNP_FN_ZEROS_LIKE, /**< Used in numpy.zeros_like() impl */ diff --git a/dpnp/backend/kernels/dpnp_krnl_indexing.cpp b/dpnp/backend/kernels/dpnp_krnl_indexing.cpp index 523acd447c64..5400da817581 100644 --- a/dpnp/backend/kernels/dpnp_krnl_indexing.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_indexing.cpp @@ -125,693 +125,6 @@ DPCTLSyclEventRef (*dpnp_choose_ext_c)(DPCTLSyclQueueRef, const DPCTLEventVectorRef) = dpnp_choose_c<_DataType1, _DataType2>; -template -DPCTLSyclEventRef dpnp_diagonal_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - const size_t input1_size, - void *result1, - const size_t offset, - shape_elem_type *shape, - shape_elem_type *res_shape, - const size_t res_ndim, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - const size_t res_size = std::accumulate(res_shape, res_shape + res_ndim, 1, - std::multiplies()); - if (!(res_size && input1_size)) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, input1_size, - true); - DPNPC_ptr_adapter<_DataType> result_ptr(q_ref, result1, res_size, true, - true); - _DataType *array_1 = input1_ptr.get_ptr(); - _DataType *result = result_ptr.get_ptr(); - - const size_t res_shape_ndim_sub_1 = - static_cast(res_shape[res_ndim - 1]); - - if (res_ndim <= 1) { - for (size_t i = 0; i < res_shape_ndim_sub_1; ++i) { - result[i] = array_1[i * shape[res_ndim] + i + offset]; - } - } - else { - std::map> xyz; - for (size_t i = 0; i < static_cast(res_shape[0]); i++) { - xyz[i] = {i}; - } - - size_t index = 1; - while (index < res_ndim - 1) { - size_t shape_element = res_shape[index]; - std::map> new_shape_array; - size_t ind = 0; - for (size_t i = 0; i < shape_element; i++) { - for (size_t j = 0; j < xyz.size(); j++) { - std::vector new_shape; - std::vector list_ind = xyz[j]; - for (size_t k = 0; k < list_ind.size(); k++) { - new_shape.push_back(list_ind.at(k)); - } - new_shape.push_back(i); - new_shape_array[ind] = new_shape; - ind += 1; - } - } - size_t len_new_shape_array = new_shape_array.size() * (index + 1); - - for (size_t k = 0; k < len_new_shape_array; k++) { - xyz[k] = new_shape_array[k]; - } - index += 1; - } - - for (size_t i = 0; i < res_shape_ndim_sub_1; i++) { - for (size_t j = 0; j < xyz.size(); j++) { - std::vector ind_list = xyz[j]; - if (ind_list.size() == 0) { - continue; - } - else { - std::vector ind_input_{i, i + offset}; - ind_input_.insert(ind_input_.end(), ind_list.begin(), - ind_list.end()); - - std::vector ind_output_ = ind_list; - ind_output_.push_back(i); - - const size_t ind_output_size = ind_output_.size(); - size_t ind_output = 0; - size_t n = 1; - for (size_t k = 0; k < ind_output_size; k++) { - size_t ind = ind_output_size - 1 - k; - ind_output += n * ind_output_[ind]; - n *= res_shape[ind]; - } - - const size_t ind_input_size = ind_input_.size(); - size_t ind_input = 0; - size_t m = 1; - for (size_t k = 0; k < ind_input_size; k++) { - size_t ind = ind_input_size - 1 - k; - ind_input += m * ind_input_[ind]; - m *= shape[ind]; - } - - result[ind_output] = array_1[ind_input]; - } - } - } - } - - return event_ref; -} - -template -void dpnp_diagonal_c(void *array1_in, - const size_t input1_size, - void *result1, - const size_t offset, - shape_elem_type *shape, - shape_elem_type *res_shape, - const size_t res_ndim) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_diagonal_c<_DataType>( - q_ref, array1_in, input1_size, result1, offset, shape, res_shape, - res_ndim, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_diagonal_default_c)(void *, - const size_t, - void *, - const size_t, - shape_elem_type *, - shape_elem_type *, - const size_t) = dpnp_diagonal_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_fill_diagonal_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *val_in, - shape_elem_type *shape, - const size_t ndim, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - const size_t result_size = std::accumulate( - shape, shape + ndim, 1, std::multiplies()); - if (!(result_size && array1_in)) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - DPNPC_ptr_adapter<_DataType> result_ptr(q_ref, array1_in, result_size, true, - true); - DPNPC_ptr_adapter<_DataType> val_ptr(q_ref, val_in, 1, true); - _DataType *array_1 = result_ptr.get_ptr(); - _DataType *val_arr = val_ptr.get_ptr(); - - shape_elem_type min_shape = shape[0]; - for (size_t i = 0; i < ndim; ++i) { - if (shape[i] < min_shape) { - min_shape = shape[i]; - } - } - - _DataType val = val_arr[0]; - - for (size_t i = 0; i < static_cast(min_shape); ++i) { - size_t ind = 0; - size_t n = 1; - for (size_t k = 0; k < ndim; k++) { - size_t ind_ = ndim - 1 - k; - ind += n * i; - n *= shape[ind_]; - } - array_1[ind] = val; - } - - return event_ref; -} - -template -void dpnp_fill_diagonal_c(void *array1_in, - void *val_in, - shape_elem_type *shape, - const size_t ndim) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_fill_diagonal_c<_DataType>( - q_ref, array1_in, val_in, shape, ndim, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_fill_diagonal_default_c)(void *, - void *, - shape_elem_type *, - const size_t) = - dpnp_fill_diagonal_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_nonzero_c(DPCTLSyclQueueRef q_ref, - const void *in_array1, - void *result1, - const size_t result_size, - const shape_elem_type *shape, - const size_t ndim, - const size_t j, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if ((in_array1 == nullptr) || (result1 == nullptr)) { - return event_ref; - } - - if (ndim == 0) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - const size_t input1_size = std::accumulate( - shape, shape + ndim, 1, std::multiplies()); - - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, in_array1, input1_size, - true); - DPNPC_ptr_adapter result_ptr(q_ref, result1, result_size, true, true); - const _DataType *arr = input1_ptr.get_ptr(); - long *result = result_ptr.get_ptr(); - - size_t idx = 0; - size_t *ids = new size_t[ndim]; - - for (size_t i = 0; i < input1_size; ++i) { - if (arr[i] != 0) { - size_t ind1 = input1_size; - size_t ind2 = i; - - for (size_t k = 0; k < ndim; ++k) { - ind1 = ind1 / shape[k]; - ids[k] = ind2 / ind1; - ind2 = ind2 % ind1; - } - - result[idx] = ids[j]; - idx += 1; - } - } - delete[] ids; - - return event_ref; -} - -template -void dpnp_nonzero_c(const void *in_array1, - void *result1, - const size_t result_size, - const shape_elem_type *shape, - const size_t ndim, - const size_t j) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_nonzero_c<_DataType>(q_ref, in_array1, result1, result_size, shape, - ndim, j, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_nonzero_default_c)(const void *, - void *, - const size_t, - const shape_elem_type *, - const size_t, - const size_t) = dpnp_nonzero_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_place_c(DPCTLSyclQueueRef q_ref, - void *arr_in, - long *mask_in, - void *vals_in, - const size_t arr_size, - const size_t vals_size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!arr_size) { - return event_ref; - } - - if (!vals_size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, vals_in, vals_size, true); - DPNPC_ptr_adapter<_DataType> result_ptr(q_ref, arr_in, arr_size, true, - true); - _DataType *vals = input1_ptr.get_ptr(); - _DataType *arr = result_ptr.get_ptr(); - - DPNPC_ptr_adapter mask_ptr(q_ref, mask_in, arr_size, true); - long *mask = mask_ptr.get_ptr(); - - size_t counter = 0; - for (size_t i = 0; i < arr_size; ++i) { - if (mask[i]) { - arr[i] = vals[counter % vals_size]; - counter += 1; - } - } - - return event_ref; -} - -template -void dpnp_place_c(void *arr_in, - long *mask_in, - void *vals_in, - const size_t arr_size, - const size_t vals_size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_place_c<_DataType>(q_ref, arr_in, mask_in, vals_in, arr_size, - vals_size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_place_default_c)(void *, - long *, - void *, - const size_t, - const size_t) = dpnp_place_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_put_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *ind_in, - void *v_in, - const size_t size, - const size_t size_ind, - const size_t size_v, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if ((array1_in == nullptr) || (ind_in == nullptr) || (v_in == nullptr)) { - return event_ref; - } - - if (size_v == 0) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - DPNPC_ptr_adapter input1_ptr(q_ref, ind_in, size_ind, true); - DPNPC_ptr_adapter<_DataType> input2_ptr(q_ref, v_in, size_v, true); - DPNPC_ptr_adapter<_DataType> result_ptr(q_ref, array1_in, size, true, true); - size_t *ind = input1_ptr.get_ptr(); - _DataType *v = input2_ptr.get_ptr(); - _DataType *array_1 = result_ptr.get_ptr(); - - for (size_t i = 0; i < size; ++i) { - for (size_t j = 0; j < size_ind; ++j) { - if (i == ind[j] || (i == (size + ind[j]))) { - array_1[i] = v[j % size_v]; - } - } - } - - return event_ref; -} - -template -void dpnp_put_c(void *array1_in, - void *ind_in, - void *v_in, - const size_t size, - const size_t size_ind, - const size_t size_v) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_put_c<_DataType, _IndecesType, _ValueType>( - q_ref, array1_in, ind_in, v_in, size, size_ind, size_v, - dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_put_default_c)(void *, - void *, - void *, - const size_t, - const size_t, - const size_t) = - dpnp_put_c<_DataType, _IndecesType, _ValueType>; - -template -DPCTLSyclEventRef - dpnp_put_along_axis_c(DPCTLSyclQueueRef q_ref, - void *arr_in, - long *indices_in, - void *values_in, - size_t axis, - const shape_elem_type *shape, - size_t ndim, - size_t size_indices, - size_t values_size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - sycl::queue q = *(reinterpret_cast(q_ref)); - - const size_t size_arr = std::accumulate(shape, shape + ndim, 1, - std::multiplies()); - - DPNPC_ptr_adapter input1_ptr(q_ref, indices_in, size_indices, true); - DPNPC_ptr_adapter<_DataType> input2_ptr(q_ref, values_in, values_size, - true); - DPNPC_ptr_adapter<_DataType> result_ptr(q_ref, arr_in, size_arr, true, - true); - size_t *indices = input1_ptr.get_ptr(); - _DataType *values = input2_ptr.get_ptr(); - _DataType *arr = result_ptr.get_ptr(); - - if (axis != (ndim - 1)) { - std::vector res_shape; - for (size_t i = 0; i < ndim; i++) { - if (axis != i) { - res_shape.push_back(shape[i]); - } - } - size_t res_ndim = res_shape.size(); - - size_t prod = 1; - for (size_t i = 0; i < res_ndim; ++i) { - if (res_shape[i] != 0) { - prod *= res_shape[i]; - } - } - - size_t *ind_array = new size_t[prod]; - bool *bool_ind_array = new bool[prod]; - for (size_t i = 0; i < prod; ++i) { - bool_ind_array[i] = true; - } - - size_t *arr_shape_offsets = new size_t[ndim]; - size_t acc = 1; - for (size_t i = ndim - 1; i > 0; --i) { - arr_shape_offsets[i] = acc; - acc *= shape[i]; - } - arr_shape_offsets[0] = acc; - - size_t *output_shape_offsets = new size_t[res_ndim]; - acc = 1; - if (res_ndim > 0) { - for (size_t i = res_ndim - 1; i > 0; --i) { - output_shape_offsets[i] = acc; - acc *= res_shape[i]; - } - } - output_shape_offsets[0] = acc; - - size_t size_result = 1; - for (size_t i = 0; i < res_ndim; ++i) { - size_result *= res_shape[i]; - } - - // init result array - size_t *xyz = new size_t[res_ndim]; - for (size_t result_idx = 0; result_idx < size_result; ++result_idx) { - size_t remainder = result_idx; - for (size_t i = 0; i < res_ndim; ++i) { - xyz[i] = remainder / output_shape_offsets[i]; - remainder = remainder - xyz[i] * output_shape_offsets[i]; - } - - // FIXME: computed and unused. Commented out per compiler warning - // size_t source_axis[ndim]; - // size_t result_axis_idx = 0; - // for (size_t idx = 0; idx < ndim; ++idx) { - // bool found = false; - // if (axis == idx) { - // found = true; - // } - // if (found) { - // source_axis[idx] = 0; - // } - // else { - // source_axis[idx] = xyz[result_axis_idx]; - // result_axis_idx++; - // } - // } - - // size_t source_idx = 0; - // for (size_t i = 0; i < static_cast(ndim); ++i) - // { - // source_idx += arr_shape_offsets[i] * source_axis[i]; - // } - } - - for (size_t source_idx = 0; source_idx < size_arr; ++source_idx) { - // reconstruct x,y,z from linear source_idx - size_t remainder = source_idx; - for (size_t i = 0; i < ndim; ++i) { - xyz[i] = remainder / arr_shape_offsets[i]; - remainder = remainder - xyz[i] * arr_shape_offsets[i]; - } - - // extract result axis - std::vector result_axis; - for (size_t idx = 0; idx < ndim; ++idx) { - // try to find current idx in axis array - if (axis != idx) { - result_axis.push_back(xyz[idx]); - } - } - - // Construct result offset - size_t result_offset = 0; - for (size_t i = 0; i < res_ndim; ++i) { - result_offset += output_shape_offsets[i] * result_axis[i]; - } - - if (bool_ind_array[result_offset]) { - ind_array[result_offset] = 0; - bool_ind_array[result_offset] = false; - } - else { - ind_array[result_offset] += 1; - } - - if ((ind_array[result_offset] % size_indices) == - indices[result_offset % size_indices]) - { - arr[source_idx] = values[source_idx % values_size]; - } - } - - delete[] ind_array; - delete[] bool_ind_array; - delete[] arr_shape_offsets; - delete[] output_shape_offsets; - delete[] xyz; - } - else { - for (size_t i = 0; i < size_arr; ++i) { - size_t ind = - size_indices * (i / size_indices) + indices[i % size_indices]; - arr[ind] = values[i % values_size]; - } - } - return event_ref; -} - -template -void dpnp_put_along_axis_c(void *arr_in, - long *indices_in, - void *values_in, - size_t axis, - const shape_elem_type *shape, - size_t ndim, - size_t size_indices, - size_t values_size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_put_along_axis_c<_DataType>( - q_ref, arr_in, indices_in, values_in, axis, shape, ndim, size_indices, - values_size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_put_along_axis_default_c)(void *, - long *, - void *, - size_t, - const shape_elem_type *, - size_t, - size_t, - size_t) = - dpnp_put_along_axis_c<_DataType>; - -template -class dpnp_take_c_kernel; - -template -DPCTLSyclEventRef dpnp_take_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - const size_t array1_size, - void *indices1, - void *result1, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)array1_size; - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - sycl::queue q = *(reinterpret_cast(q_ref)); - - _DataType *array_1 = reinterpret_cast<_DataType *>(array1_in); - _IndecesType *indices = reinterpret_cast<_IndecesType *>(indices1); - _DataType *result = reinterpret_cast<_DataType *>(result1); - - sycl::range<1> gws(size); - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { - const size_t idx = global_id[0]; - result[idx] = array_1[indices[idx]]; - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.parallel_for>( - gws, kernel_parallel_for_func); - }; - - sycl::event event = q.submit(kernel_func); - - event_ref = reinterpret_cast(&event); - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_take_c(void *array1_in, - const size_t array1_size, - void *indices1, - void *result1, - size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_take_c<_DataType, _IndecesType>( - q_ref, array1_in, array1_size, indices1, result1, size, - dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_take_default_c)(void *, const size_t, void *, void *, size_t) = - dpnp_take_c<_DataType, _IndecesType>; - -template -DPCTLSyclEventRef (*dpnp_take_ext_c)(DPCTLSyclQueueRef, - void *, - const size_t, - void *, - void *, - size_t, - const DPCTLEventVectorRef) = - dpnp_take_c<_DataType, _IndecesType>; - void func_map_init_indexing_func(func_map_t &fmap) { fmap[DPNPFuncName::DPNP_FN_CHOOSE][eft_INT][eft_INT] = { @@ -847,85 +160,5 @@ void func_map_init_indexing_func(func_map_t &fmap) eft_FLT, (void *)dpnp_choose_ext_c}; fmap[DPNPFuncName::DPNP_FN_CHOOSE_EXT][eft_LNG][eft_DBL] = { eft_DBL, (void *)dpnp_choose_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_DIAGONAL][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_diagonal_default_c}; - fmap[DPNPFuncName::DPNP_FN_DIAGONAL][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_diagonal_default_c}; - fmap[DPNPFuncName::DPNP_FN_DIAGONAL][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_diagonal_default_c}; - fmap[DPNPFuncName::DPNP_FN_DIAGONAL][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_diagonal_default_c}; - - fmap[DPNPFuncName::DPNP_FN_FILL_DIAGONAL][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_fill_diagonal_default_c}; - fmap[DPNPFuncName::DPNP_FN_FILL_DIAGONAL][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_fill_diagonal_default_c}; - fmap[DPNPFuncName::DPNP_FN_FILL_DIAGONAL][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_fill_diagonal_default_c}; - fmap[DPNPFuncName::DPNP_FN_FILL_DIAGONAL][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_fill_diagonal_default_c}; - - fmap[DPNPFuncName::DPNP_FN_NONZERO][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_nonzero_default_c}; - fmap[DPNPFuncName::DPNP_FN_NONZERO][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_nonzero_default_c}; - fmap[DPNPFuncName::DPNP_FN_NONZERO][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_nonzero_default_c}; - fmap[DPNPFuncName::DPNP_FN_NONZERO][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_nonzero_default_c}; - - fmap[DPNPFuncName::DPNP_FN_PLACE][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_place_default_c}; - fmap[DPNPFuncName::DPNP_FN_PLACE][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_place_default_c}; - fmap[DPNPFuncName::DPNP_FN_PLACE][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_place_default_c}; - fmap[DPNPFuncName::DPNP_FN_PLACE][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_place_default_c}; - - fmap[DPNPFuncName::DPNP_FN_PUT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_put_default_c}; - fmap[DPNPFuncName::DPNP_FN_PUT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_put_default_c}; - fmap[DPNPFuncName::DPNP_FN_PUT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_put_default_c}; - fmap[DPNPFuncName::DPNP_FN_PUT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_put_default_c}; - - fmap[DPNPFuncName::DPNP_FN_PUT_ALONG_AXIS][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_put_along_axis_default_c}; - fmap[DPNPFuncName::DPNP_FN_PUT_ALONG_AXIS][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_put_along_axis_default_c}; - fmap[DPNPFuncName::DPNP_FN_PUT_ALONG_AXIS][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_put_along_axis_default_c}; - fmap[DPNPFuncName::DPNP_FN_PUT_ALONG_AXIS][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_put_along_axis_default_c}; - - fmap[DPNPFuncName::DPNP_FN_TAKE][eft_BLN][eft_INT] = { - eft_BLN, (void *)dpnp_take_default_c}; - fmap[DPNPFuncName::DPNP_FN_TAKE][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_take_default_c}; - fmap[DPNPFuncName::DPNP_FN_TAKE][eft_LNG][eft_INT] = { - eft_LNG, (void *)dpnp_take_default_c}; - fmap[DPNPFuncName::DPNP_FN_TAKE][eft_FLT][eft_INT] = { - eft_FLT, (void *)dpnp_take_default_c}; - fmap[DPNPFuncName::DPNP_FN_TAKE][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_take_default_c}; - fmap[DPNPFuncName::DPNP_FN_TAKE][eft_C128][eft_INT] = { - eft_C128, (void *)dpnp_take_default_c, int32_t>}; - fmap[DPNPFuncName::DPNP_FN_TAKE][eft_BLN][eft_LNG] = { - eft_BLN, (void *)dpnp_take_default_c}; - fmap[DPNPFuncName::DPNP_FN_TAKE][eft_INT][eft_LNG] = { - eft_INT, (void *)dpnp_take_default_c}; - fmap[DPNPFuncName::DPNP_FN_TAKE][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_take_default_c}; - fmap[DPNPFuncName::DPNP_FN_TAKE][eft_FLT][eft_LNG] = { - eft_FLT, (void *)dpnp_take_default_c}; - fmap[DPNPFuncName::DPNP_FN_TAKE][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_take_default_c}; - fmap[DPNPFuncName::DPNP_FN_TAKE][eft_C128][eft_LNG] = { - eft_C128, (void *)dpnp_take_default_c, int64_t>}; - return; }