From 3f554d31e6feaf3f0b194de9f21962fe31ccdadf Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Wed, 16 Oct 2024 07:09:21 -0700 Subject: [PATCH 1/5] clean up legacy statistics --- dpnp/backend/include/dpnp_iface.hpp | 224 ---- dpnp/backend/include/dpnp_iface_fptr.hpp | 8 - dpnp/backend/kernels/dpnp_krnl_statistics.cpp | 1070 ----------------- dpnp/dpnp_algo/dpnp_algo.pxd | 17 - dpnp/dpnp_algo/dpnp_algo.pyx | 9 - 5 files changed, 1328 deletions(-) diff --git a/dpnp/backend/include/dpnp_iface.hpp b/dpnp/backend/include/dpnp_iface.hpp index 6ceddf889629..1c76fe7adf07 100644 --- a/dpnp/backend/include/dpnp_iface.hpp +++ b/dpnp/backend/include/dpnp_iface.hpp @@ -100,35 +100,6 @@ INP_DLLEXPORT void dpnp_memory_memcpy_c(DPCTLSyclQueueRef q_ref, INP_DLLEXPORT void dpnp_memory_memcpy_c(void *dst, const void *src, size_t size_in_bytes); -/** - * @ingroup BACKEND_API - * @brief Compute the variance along the specified axis, while ignoring NaNs. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array Input array. - * @param [in] mask_arr Input mask array when elem is nan. - * @param [out] result Output array. - * @param [in] result_size Output array size. - * @param [in] size Number of elements in input arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_nanvar_c(DPCTLSyclQueueRef q_ref, - void *array, - void *mask_arr, - void *result, - const size_t result_size, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_nanvar_c(void *array, - void *mask_arr, - void *result, - const size_t result_size, - size_t size); - /** * @ingroup BACKEND_API * @brief Custom implementation of dot function @@ -238,29 +209,6 @@ INP_DLLEXPORT void dpnp_sum_c(void *result_out, const void *initial, const long *where); -/** - * @ingroup BACKEND_API - * @brief Custom implementation of count_nonzero function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array1_in Input array. - * @param [out] result1_out Output array. - * @param [in] size Number of elements in input arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - * - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_count_nonzero_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result1_out, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void - dpnp_count_nonzero_c(void *array1_in, void *result1_out, size_t size); - /** * @ingroup BACKEND_API * @brief Return a partitioned copy of an array. @@ -541,73 +489,6 @@ INP_DLLEXPORT DPCTLSyclEventRef template INP_DLLEXPORT void dpnp_initval_c(void *result1, void *value, size_t size); -/** - * @ingroup BACKEND_API - * @brief math library implementation of max function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array1_in Input array with data. - * @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] axis Axis. - * @param [in] naxis Number of elements in axis. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_max_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result1, - const size_t result_size, - const shape_elem_type *shape, - size_t ndim, - const shape_elem_type *axis, - size_t naxis, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_max_c(void *array1_in, - void *result1, - const size_t result_size, - const shape_elem_type *shape, - size_t ndim, - const shape_elem_type *axis, - size_t naxis); - -/** - * @ingroup BACKEND_API - * @brief math library implementation of mean function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array Input array with data. - * @param [out] result Output array. - * @param [in] shape Shape of input array. - * @param [in] ndim Number of elements in shape. - * @param [in] axis Axis. - * @param [in] naxis Number of elements in axis. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_mean_c(DPCTLSyclQueueRef q_ref, - void *array, - void *result, - const shape_elem_type *shape, - size_t ndim, - const shape_elem_type *axis, - size_t naxis, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_mean_c(void *array, - void *result, - const shape_elem_type *shape, - size_t ndim, - const shape_elem_type *axis, - size_t naxis); - /** * @ingroup BACKEND_API * @brief math library implementation of median function @@ -640,41 +521,6 @@ INP_DLLEXPORT void dpnp_median_c(void *array, const shape_elem_type *axis, size_t naxis); -/** - * @ingroup BACKEND_API - * @brief math library implementation of min function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array Input array with data. - * @param [out] result 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] axis Axis. - * @param [in] naxis Number of elements in axis. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_min_c(DPCTLSyclQueueRef q_ref, - void *array, - void *result, - const size_t result_size, - const shape_elem_type *shape, - size_t ndim, - const shape_elem_type *axis, - size_t naxis, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_min_c(void *array, - void *result, - const size_t result_size, - const shape_elem_type *shape, - size_t ndim, - const shape_elem_type *axis, - size_t naxis); - /** * @ingroup BACKEND_API * @brief math library implementation of argmax function @@ -717,76 +563,6 @@ INP_DLLEXPORT DPCTLSyclEventRef template INP_DLLEXPORT void dpnp_argmin_c(void *array, void *result, size_t size); -/** - * @ingroup BACKEND_API - * @brief math library implementation of std function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array Input array with data. - * @param [out] result Output array with indices. - * @param [in] shape Shape of input array. - * @param [in] ndim Number of elements in shape. - * @param [in] axis Axis. - * @param [in] naxis Number of elements in axis. - * @param [in] ddof Delta degrees of freedom. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_std_c(DPCTLSyclQueueRef q_ref, - void *array, - void *result, - const shape_elem_type *shape, - size_t ndim, - const shape_elem_type *axis, - size_t naxis, - size_t ddof, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_std_c(void *array, - void *result, - const shape_elem_type *shape, - size_t ndim, - const shape_elem_type *axis, - size_t naxis, - size_t ddof); - -/** - * @ingroup BACKEND_API - * @brief math library implementation of var function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array Input array with data. - * @param [out] result Output array with indices. - * @param [in] shape Shape of input array. - * @param [in] ndim Number of elements in shape. - * @param [in] axis Axis. - * @param [in] naxis Number of elements in axis. - * @param [in] ddof Delta degrees of freedom. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_var_c(DPCTLSyclQueueRef q_ref, - void *array, - void *result, - const shape_elem_type *shape, - size_t ndim, - const shape_elem_type *axis, - size_t naxis, - size_t ddof, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_var_c(void *array, - void *result, - const shape_elem_type *shape, - size_t ndim, - const shape_elem_type *axis, - size_t naxis, - size_t ddof); - #define MACRO_1ARG_1TYPE_OP(__name__, __operation1__, __operation2__) \ template \ INP_DLLEXPORT DPCTLSyclEventRef __name__( \ diff --git a/dpnp/backend/include/dpnp_iface_fptr.hpp b/dpnp/backend/include/dpnp_iface_fptr.hpp index 8652a88b9b65..16be3253ec98 100644 --- a/dpnp/backend/include/dpnp_iface_fptr.hpp +++ b/dpnp/backend/include/dpnp_iface_fptr.hpp @@ -68,8 +68,6 @@ enum class DPNPFuncName : size_t DPNP_FN_CORRELATE, /**< Used in numpy.correlate() impl */ DPNP_FN_CORRELATE_EXT, /**< Used in numpy.correlate() impl, requires extra parameters */ - DPNP_FN_COUNT_NONZERO, /**< Used in numpy.count_nonzero() impl */ - DPNP_FN_COV, /**< Used in numpy.cov() impl */ DPNP_FN_DOT, /**< Used in numpy.dot() impl */ DPNP_FN_DOT_EXT, /**< Used in numpy.dot() impl, requires extra parameters */ DPNP_FN_ERF, /**< Used in scipy.special.erf impl */ @@ -79,17 +77,13 @@ enum class DPNPFuncName : size_t */ DPNP_FN_INITVAL_EXT, /**< Used in numpy ones, ones_like, zeros, zeros_like impls */ - DPNP_FN_MAX, /**< Used in numpy.max() impl */ - DPNP_FN_MEAN, /**< Used in numpy.mean() impl */ DPNP_FN_MEDIAN, /**< Used in numpy.median() impl */ DPNP_FN_MEDIAN_EXT, /**< Used in numpy.median() impl, requires extra parameters */ - DPNP_FN_MIN, /**< Used in numpy.min() impl */ DPNP_FN_MODF, /**< Used in numpy.modf() impl */ DPNP_FN_MODF_EXT, /**< Used in numpy.modf() impl, requires extra parameters */ DPNP_FN_MULTIPLY, /**< Used in numpy.multiply() impl */ - DPNP_FN_NANVAR, /**< Used in numpy.nanvar() 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 */ @@ -218,9 +212,7 @@ enum class DPNPFuncName : size_t DPNP_FN_SQRT, /**< Used in numpy.sqrt() impl */ DPNP_FN_SQRT_EXT, /**< Used in numpy.sqrt() impl, requires extra parameters */ - DPNP_FN_STD, /**< Used in numpy.std() impl */ DPNP_FN_SUM, /**< Used in numpy.sum() 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 */ DPNP_FN_LAST, /**< The latest element of the enumeration */ diff --git a/dpnp/backend/kernels/dpnp_krnl_statistics.cpp b/dpnp/backend/kernels/dpnp_krnl_statistics.cpp index a108b5200913..66a4881d7f2a 100644 --- a/dpnp/backend/kernels/dpnp_krnl_statistics.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_statistics.cpp @@ -128,503 +128,6 @@ DPCTLSyclEventRef (*dpnp_correlate_ext_c)(DPCTLSyclQueueRef, const DPCTLEventVectorRef) = dpnp_correlate_c<_DataType_output, _DataType_input1, _DataType_input2>; -template -class dpnp_cov_c_kernel1; - -template -class dpnp_cov_c_kernel2; - -template -DPCTLSyclEventRef dpnp_cov_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result1, - size_t nrows, - size_t ncols, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!nrows || !ncols) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, nrows * ncols); - _DataType *array_1 = input1_ptr.get_ptr(); - _DataType *result = reinterpret_cast<_DataType *>(result1); - - auto policy = oneapi::dpl::execution::make_device_policy< - class dpnp_cov_c_kernel1<_DataType>>(q); - - _DataType *mean = reinterpret_cast<_DataType *>( - sycl::malloc_shared(nrows * sizeof(_DataType), q)); - for (size_t i = 0; i < nrows; ++i) { - _DataType *row_start = array_1 + ncols * i; - mean[i] = std::reduce(policy, row_start, row_start + ncols, - _DataType(0), std::plus<_DataType>()) / - ncols; - } - policy.queue().wait(); - - _DataType *temp = reinterpret_cast<_DataType *>( - sycl::malloc_shared(nrows * ncols * sizeof(_DataType), q)); - for (size_t i = 0; i < nrows; ++i) { - size_t offset = ncols * i; - _DataType *row_start = array_1 + offset; - std::transform(policy, row_start, row_start + ncols, temp + offset, - [=](_DataType x) { return x - mean[i]; }); - } - policy.queue().wait(); - - sycl::event event_syrk; - - const _DataType alpha = _DataType(1) / (ncols - 1); - const _DataType beta = _DataType(0); - - event_syrk = - mkl_blas::syrk(q, // queue &exec_queue, - oneapi::mkl::uplo::upper, // uplo upper_lower, - oneapi::mkl::transpose::nontrans, // transpose trans, - nrows, // std::int64_t n, - ncols, // std::int64_t k, - alpha, // T alpha, - temp, // const T* a, - ncols, // std::int64_t lda, - beta, // T beta, - result, // T* c, - nrows); // std::int64_t ldc); - event_syrk.wait(); - - // fill lower elements - sycl::event event; - sycl::range<1> gws(nrows * nrows); - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { - const size_t idx = global_id[0]; - const size_t row_idx = idx / nrows; - const size_t col_idx = idx - row_idx * nrows; - if (col_idx < row_idx) { - result[idx] = result[col_idx * nrows + row_idx]; - } - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.parallel_for>( - gws, kernel_parallel_for_func); - }; - - event = q.submit(kernel_func); - - event.wait(); - - sycl::free(mean, q); - sycl::free(temp, q); - - return event_ref; -} - -template -void dpnp_cov_c(void *array1_in, void *result1, size_t nrows, size_t ncols) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_cov_c<_DataType>( - q_ref, array1_in, result1, nrows, ncols, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_cov_default_c)(void *, void *, size_t, size_t) = - dpnp_cov_c<_DataType>; - -template -DPCTLSyclEventRef - dpnp_count_nonzero_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result1_out, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (array1_in == nullptr) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - DPNPC_ptr_adapter<_DataType_input> input1_ptr(q_ref, array1_in, size, true); - DPNPC_ptr_adapter<_DataType_output> result_ptr(q_ref, result1_out, 1, true, - true); - _DataType_input *array1 = input1_ptr.get_ptr(); - _DataType_output *result1 = result_ptr.get_ptr(); - - result1[0] = 0; - - for (size_t i = 0; i < size; ++i) { - if (array1[i] != 0) { - result1[0] += 1; - } - } - - return event_ref; -} - -template -void dpnp_count_nonzero_c(void *array1_in, void *result1_out, size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_count_nonzero_c<_DataType_input, _DataType_output>( - q_ref, array1_in, result1_out, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_count_nonzero_default_c)(void *, void *, size_t) = - dpnp_count_nonzero_c<_DataType_input, _DataType_output>; - -template -class dpnp_max_c_kernel; - -template -DPCTLSyclEventRef dpnp_max_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result1, - const size_t result_size, - const shape_elem_type *shape, - size_t ndim, - const shape_elem_type *axis, - size_t naxis, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - const size_t size_input = std::accumulate( - shape, shape + ndim, 1, std::multiplies()); - if (!size_input) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, size_input, true); - DPNPC_ptr_adapter<_DataType> result_ptr(q_ref, result1, result_size, true, - true); - _DataType *array_1 = input1_ptr.get_ptr(); - _DataType *result = result_ptr.get_ptr(); - - if (naxis == 0) { - __attribute__((unused)) void *tmp = (void *)(axis + naxis); - - size_t size = 1; - for (size_t i = 0; i < ndim; ++i) { - size *= shape[i]; - } - - if constexpr (std::is_same<_DataType, double>::value || - std::is_same<_DataType, float>::value) - { - // Required initializing the result before call the function - result[0] = array_1[0]; - - auto dataset = - mkl_stats::make_dataset(1, size, - array_1); - - sycl::event event = mkl_stats::max(q, dataset, result); - - event.wait(); - } - else { - auto policy = oneapi::dpl::execution::make_device_policy< - class dpnp_max_c_kernel<_DataType>>(q); - - _DataType *res = std::max_element(policy, array_1, array_1 + size); - policy.queue().wait(); - - result[0] = *res; - } - } - else { - std::vector res_shape; - for (size_t i = 0; i < ndim; ++i) { - bool found = false; - for (size_t j = 0; j < naxis; ++j) { - if (static_cast(axis[j]) == i) { - found = true; - break; - } - } - if (!found) { - res_shape.push_back(shape[i]); - } - } - const size_t res_ndim = res_shape.size(); - - size_t acc = 1; - std::vector input_shape_offsets{acc}; - for (size_t i = ndim - 2; i > 0; --i) { - acc *= shape[i]; - input_shape_offsets.insert(input_shape_offsets.begin(), acc); - } - - acc = 1; - std::vector output_shape_offsets{acc}; - for (size_t i = res_ndim - 2; i > 0; --i) { - acc *= res_shape[i]; - output_shape_offsets.insert(output_shape_offsets.begin(), 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]; - size_t *source_axis = new size_t[ndim]; - size_t *result_axis = 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]; - } - - size_t result_axis_idx = 0; - for (size_t idx = 0; idx < ndim; ++idx) { - bool found = false; - for (size_t i = 0; i < naxis; ++i) { - if (static_cast(axis[i]) == idx) { - found = true; - break; - } - } - 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 < ndim; ++i) { - source_idx += input_shape_offsets[i] * source_axis[i]; - } - - result[result_idx] = array_1[source_idx]; - } - - for (size_t source_idx = 0; source_idx < size_input; ++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 / input_shape_offsets[i]; - remainder = remainder - xyz[i] * input_shape_offsets[i]; - } - - // extract result axis - size_t result_idx = 0; - for (size_t idx = 0; idx < ndim; ++idx) { - // try to find current idx in axis array - bool found = false; - for (size_t i = 0; i < naxis; ++i) { - if (static_cast(axis[i]) == idx) { - found = true; - break; - } - } - if (!found) { - result_axis[result_idx] = xyz[idx]; - result_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 (result[result_offset] < array_1[source_idx]) { - result[result_offset] = array_1[source_idx]; - } - } - - delete[] xyz; - delete[] source_axis; - delete[] result_axis; - } - - return event_ref; -} - -// Explicit instantiation of the function, since dpnp_max_c() is used by -// other template functions, but implicit instantiation is not applied anymore. -template DPCTLSyclEventRef dpnp_max_c(DPCTLSyclQueueRef q_ref, - void *, - void *, - const size_t, - const shape_elem_type *, - size_t, - const shape_elem_type *, - size_t, - const DPCTLEventVectorRef); - -template DPCTLSyclEventRef dpnp_max_c(DPCTLSyclQueueRef q_ref, - void *, - void *, - const size_t, - const shape_elem_type *, - size_t, - const shape_elem_type *, - size_t, - const DPCTLEventVectorRef); - -template DPCTLSyclEventRef dpnp_max_c(DPCTLSyclQueueRef q_ref, - void *, - void *, - const size_t, - const shape_elem_type *, - size_t, - const shape_elem_type *, - size_t, - const DPCTLEventVectorRef); - -template DPCTLSyclEventRef dpnp_max_c(DPCTLSyclQueueRef q_ref, - void *, - void *, - const size_t, - const shape_elem_type *, - size_t, - const shape_elem_type *, - size_t, - const DPCTLEventVectorRef); - -template -void dpnp_max_c(void *array1_in, - void *result1, - const size_t result_size, - const shape_elem_type *shape, - size_t ndim, - const shape_elem_type *axis, - size_t naxis) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_max_c<_DataType>(q_ref, array1_in, result1, result_size, shape, - ndim, axis, naxis, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_max_default_c)(void *, - void *, - const size_t, - const shape_elem_type *, - size_t, - const shape_elem_type *, - size_t) = dpnp_max_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_mean_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result1, - const shape_elem_type *shape, - size_t ndim, - const shape_elem_type *axis, - size_t naxis, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - __attribute__((unused)) void *tmp = (void *)(axis + naxis); - - DPCTLSyclEventRef event_ref = nullptr; - - const size_t size = std::accumulate(shape, shape + ndim, 1, - std::multiplies()); - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, size, true); - DPNPC_ptr_adapter<_ResultType> result_ptr(q_ref, result1, 1, true, true); - _DataType *array = input1_ptr.get_ptr(); - _ResultType *result = result_ptr.get_ptr(); - - if constexpr (std::is_same<_DataType, double>::value || - std::is_same<_DataType, float>::value) - { - auto dataset = mkl_stats::make_dataset< - mkl_stats::layout::row_major /*, _ResultType*/>(1, size, array); - - sycl::event event = mkl_stats::mean(q, dataset, result); - - event.wait(); - - event_ref = reinterpret_cast(&event); - - return DPCTLEvent_Copy(event_ref); - } - else { - _ResultType *sum = reinterpret_cast<_ResultType *>( - sycl::malloc_shared(1 * sizeof(_ResultType), q)); - - dpnp_sum_c<_ResultType, _DataType>(sum, array, shape, ndim, axis, naxis, - nullptr, nullptr); - - result[0] = sum[0] / static_cast<_ResultType>(size); - - sycl::free(sum, q); - - return event_ref; - } -} - -template -void dpnp_mean_c(void *array1_in, - void *result1, - const shape_elem_type *shape, - size_t ndim, - const shape_elem_type *axis, - size_t naxis) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_mean_c<_DataType, _ResultType>( - q_ref, array1_in, result1, shape, ndim, axis, naxis, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_mean_default_c)(void *, - void *, - const shape_elem_type *, - size_t, - const shape_elem_type *, - size_t) = dpnp_mean_c<_DataType, _ResultType>; - template DPCTLSyclEventRef dpnp_median_c(DPCTLSyclQueueRef q_ref, void *array1_in, @@ -707,505 +210,6 @@ DPCTLSyclEventRef (*dpnp_median_ext_c)(DPCTLSyclQueueRef, const DPCTLEventVectorRef) = dpnp_median_c<_DataType, _ResultType>; -template -class dpnp_min_c_kernel; - -template -DPCTLSyclEventRef dpnp_min_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result1, - const size_t result_size, - const shape_elem_type *shape, - size_t ndim, - const shape_elem_type *axis, - size_t naxis, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - __attribute__((unused)) void *tmp = (void *)(axis + naxis); - - DPCTLSyclEventRef event_ref = nullptr; - - const size_t size_input = std::accumulate( - shape, shape + ndim, 1, std::multiplies()); - if (!size_input) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, size_input, true); - DPNPC_ptr_adapter<_DataType> result_ptr(q_ref, result1, result_size, true, - true); - _DataType *array_1 = input1_ptr.get_ptr(); - _DataType *result = result_ptr.get_ptr(); - - if (naxis == 0) { - if constexpr (std::is_same<_DataType, double>::value || - std::is_same<_DataType, float>::value) - { - // Required initializing the result before call the function - result[0] = array_1[0]; - - auto dataset = - mkl_stats::make_dataset( - 1, size_input, array_1); - - sycl::event event = mkl_stats::min(q, dataset, result); - - event.wait(); - } - else { - auto policy = oneapi::dpl::execution::make_device_policy< - class dpnp_min_c_kernel<_DataType>>(q); - - _DataType *res = - std::min_element(policy, array_1, array_1 + size_input); - policy.queue().wait(); - - result[0] = *res; - } - } - else { - std::vector res_shape; - for (size_t i = 0; i < ndim; i++) { - bool found = false; - for (size_t j = 0; j < naxis; j++) { - if (static_cast(axis[j]) == i) { - found = true; - break; - } - } - if (!found) { - res_shape.push_back(shape[i]); - } - } - const size_t res_ndim = res_shape.size(); - - size_t acc = 1; - std::vector input_shape_offsets{acc}; - for (size_t i = ndim - 2; i > 0; --i) { - acc *= shape[i]; - input_shape_offsets.insert(input_shape_offsets.begin(), acc); - } - - acc = 1; - std::vector output_shape_offsets{acc}; - for (size_t i = res_ndim - 2; i > 0; --i) { - acc *= res_shape[i]; - output_shape_offsets.insert(output_shape_offsets.begin(), 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]; - size_t *source_axis = new size_t[ndim]; - size_t *result_axis = 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]; - } - - size_t result_axis_idx = 0; - for (size_t idx = 0; idx < ndim; ++idx) { - bool found = false; - for (size_t i = 0; i < naxis; ++i) { - if (static_cast(axis[i]) == idx) { - found = true; - break; - } - } - 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 < ndim; ++i) { - source_idx += input_shape_offsets[i] * source_axis[i]; - } - - result[result_idx] = array_1[source_idx]; - } - - for (size_t source_idx = 0; source_idx < size_input; ++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 / input_shape_offsets[i]; - remainder = remainder - xyz[i] * input_shape_offsets[i]; - } - - // extract result axis - size_t result_idx = 0; - for (size_t idx = 0; idx < ndim; ++idx) { - // try to find current idx in axis array - bool found = false; - for (size_t i = 0; i < naxis; ++i) { - if (static_cast(axis[i]) == idx) { - found = true; - break; - } - } - if (!found) { - result_axis[result_idx] = xyz[idx]; - result_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 (result[result_offset] > array_1[source_idx]) { - result[result_offset] = array_1[source_idx]; - } - } - - delete[] xyz; - delete[] source_axis; - delete[] result_axis; - } - - return event_ref; -} - -// Explicit instantiation of the function, since dpnp_min_c() is used by -// other template functions, but implicit instantiation is not applied anymore. -template DPCTLSyclEventRef dpnp_min_c(DPCTLSyclQueueRef q_ref, - void *, - void *, - const size_t, - const shape_elem_type *, - size_t, - const shape_elem_type *, - size_t, - const DPCTLEventVectorRef); - -template DPCTLSyclEventRef dpnp_min_c(DPCTLSyclQueueRef q_ref, - void *, - void *, - const size_t, - const shape_elem_type *, - size_t, - const shape_elem_type *, - size_t, - const DPCTLEventVectorRef); - -template DPCTLSyclEventRef dpnp_min_c(DPCTLSyclQueueRef q_ref, - void *, - void *, - const size_t, - const shape_elem_type *, - size_t, - const shape_elem_type *, - size_t, - const DPCTLEventVectorRef); - -template DPCTLSyclEventRef dpnp_min_c(DPCTLSyclQueueRef q_ref, - void *, - void *, - const size_t, - const shape_elem_type *, - size_t, - const shape_elem_type *, - size_t, - const DPCTLEventVectorRef); - -template -void dpnp_min_c(void *array1_in, - void *result1, - const size_t result_size, - const shape_elem_type *shape, - size_t ndim, - const shape_elem_type *axis, - size_t naxis) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_min_c<_DataType>(q_ref, array1_in, result1, result_size, shape, - ndim, axis, naxis, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_min_default_c)(void *, - void *, - const size_t, - const shape_elem_type *, - size_t, - const shape_elem_type *, - size_t) = dpnp_min_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_nanvar_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *mask_arr1, - void *result1, - const size_t result_size, - size_t arr_size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if ((array1_in == nullptr) || (mask_arr1 == nullptr) || - (result1 == nullptr)) { - return event_ref; - } - - if (arr_size == 0) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, arr_size, true); - DPNPC_ptr_adapter input2_ptr(q_ref, mask_arr1, arr_size, true); - DPNPC_ptr_adapter<_DataType> result_ptr(q_ref, result1, result_size, true, - true); - _DataType *array1 = input1_ptr.get_ptr(); - bool *mask_arr = input2_ptr.get_ptr(); - _DataType *result = result_ptr.get_ptr(); - - size_t ind = 0; - for (size_t i = 0; i < arr_size; ++i) { - if (!mask_arr[i]) { - result[ind] = array1[i]; - ind += 1; - } - } - - return event_ref; -} - -template -void dpnp_nanvar_c(void *array1_in, - void *mask_arr1, - void *result1, - const size_t result_size, - size_t arr_size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_nanvar_c<_DataType>(q_ref, array1_in, mask_arr1, result1, - result_size, arr_size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_nanvar_default_c)(void *, void *, void *, const size_t, size_t) = - dpnp_nanvar_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_std_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result1, - const shape_elem_type *shape, - size_t ndim, - const shape_elem_type *axis, - size_t naxis, - size_t ddof, - 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)); - - _ResultType *var = reinterpret_cast<_ResultType *>( - sycl::malloc_shared(1 * sizeof(_ResultType), q)); - - dpnp_var_c<_DataType, _ResultType>(array1_in, var, shape, ndim, axis, naxis, - ddof); - - const size_t result1_size = 1; - const size_t result1_ndim = 1; - const size_t result1_shape_size_in_bytes = - result1_ndim * sizeof(shape_elem_type); - const size_t result1_strides_size_in_bytes = - result1_ndim * sizeof(shape_elem_type); - shape_elem_type *result1_shape = reinterpret_cast( - sycl::malloc_shared(result1_shape_size_in_bytes, q)); - *result1_shape = 1; - shape_elem_type *result1_strides = reinterpret_cast( - sycl::malloc_shared(result1_strides_size_in_bytes, q)); - *result1_strides = 1; - - const size_t var_size = 1; - const size_t var_ndim = 1; - const size_t var_shape_size_in_bytes = var_ndim * sizeof(shape_elem_type); - const size_t var_strides_size_in_bytes = var_ndim * sizeof(shape_elem_type); - shape_elem_type *var_shape = reinterpret_cast( - sycl::malloc_shared(var_shape_size_in_bytes, q)); - *var_shape = 1; - shape_elem_type *var_strides = reinterpret_cast( - sycl::malloc_shared(var_strides_size_in_bytes, q)); - *var_strides = 1; - - DPCTLSyclEventRef e_sqrt_ref = dpnp_sqrt_c<_ResultType, _ResultType>( - q_ref, result1, result1_size, result1_ndim, result1_shape, - result1_strides, var, var_size, var_ndim, var_shape, var_strides, NULL, - NULL); - DPCTLEvent_WaitAndThrow(e_sqrt_ref); - DPCTLEvent_Delete(e_sqrt_ref); - - sycl::free(var, q); - sycl::free(result1_shape, q); - sycl::free(result1_strides, q); - sycl::free(var_shape, q); - sycl::free(var_strides, q); - - return event_ref; -} - -template -void dpnp_std_c(void *array1_in, - void *result1, - const shape_elem_type *shape, - size_t ndim, - const shape_elem_type *axis, - size_t naxis, - size_t ddof) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_std_c<_DataType, _ResultType>( - q_ref, array1_in, result1, shape, ndim, axis, naxis, ddof, - dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_std_default_c)(void *, - void *, - const shape_elem_type *, - size_t, - const shape_elem_type *, - size_t, - size_t) = dpnp_std_c<_DataType, _ResultType>; - -template -class dpnp_var_c_kernel; - -template -DPCTLSyclEventRef dpnp_var_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result1, - const shape_elem_type *shape, - size_t ndim, - const shape_elem_type *axis, - size_t naxis, - size_t ddof, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - const size_t size = std::accumulate(shape, shape + ndim, 1, - std::multiplies()); - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - sycl::event event; - - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, size); - DPNPC_ptr_adapter<_ResultType> result_ptr(q_ref, result1, 1, true, true); - _DataType *array1 = input1_ptr.get_ptr(); - _ResultType *result = result_ptr.get_ptr(); - - _ResultType *mean = reinterpret_cast<_ResultType *>( - sycl::malloc_shared(1 * sizeof(_ResultType), q)); - dpnp_mean_c<_DataType, _ResultType>(array1, mean, shape, ndim, axis, naxis); - _ResultType mean_val = mean[0]; - - _ResultType *squared_deviations = reinterpret_cast<_ResultType *>( - sycl::malloc_shared(size * sizeof(_ResultType), q)); - - sycl::range<1> gws(size); - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { - size_t i = global_id[0]; /*for (size_t i = 0; i < size; ++i)*/ - { - _ResultType deviation = - static_cast<_ResultType>(array1[i]) - mean_val; - squared_deviations[i] = deviation * deviation; - } - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.parallel_for>( - gws, kernel_parallel_for_func); - }; - - event = q.submit(kernel_func); - - event.wait(); - - dpnp_mean_c<_ResultType, _ResultType>(squared_deviations, mean, shape, ndim, - axis, naxis); - mean_val = mean[0]; - - result[0] = mean_val * size / static_cast<_ResultType>(size - ddof); - - sycl::free(mean, q); - sycl::free(squared_deviations, q); - - return event_ref; -} - -template -void dpnp_var_c(void *array1_in, - void *result1, - const shape_elem_type *shape, - size_t ndim, - const shape_elem_type *axis, - size_t naxis, - size_t ddof) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_var_c<_DataType, _ResultType>( - q_ref, array1_in, result1, shape, ndim, axis, naxis, ddof, - dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_var_default_c)(void *, - void *, - const shape_elem_type *, - size_t, - const shape_elem_type *, - size_t, - size_t) = dpnp_var_c<_DataType, _ResultType>; - void func_map_init_statistics(func_map_t &fmap) { fmap[DPNPFuncName::DPNP_FN_CORRELATE][eft_INT][eft_INT] = { @@ -1274,44 +278,6 @@ void func_map_init_statistics(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_CORRELATE_EXT][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_correlate_ext_c}; - fmap[DPNPFuncName::DPNP_FN_COUNT_NONZERO][eft_BLN][eft_BLN] = { - eft_LNG, (void *)dpnp_count_nonzero_default_c}; - fmap[DPNPFuncName::DPNP_FN_COUNT_NONZERO][eft_INT][eft_INT] = { - eft_LNG, (void *)dpnp_count_nonzero_default_c}; - fmap[DPNPFuncName::DPNP_FN_COUNT_NONZERO][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_count_nonzero_default_c}; - fmap[DPNPFuncName::DPNP_FN_COUNT_NONZERO][eft_FLT][eft_FLT] = { - eft_LNG, (void *)dpnp_count_nonzero_default_c}; - fmap[DPNPFuncName::DPNP_FN_COUNT_NONZERO][eft_DBL][eft_DBL] = { - eft_LNG, (void *)dpnp_count_nonzero_default_c}; - - fmap[DPNPFuncName::DPNP_FN_COV][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_cov_default_c}; - fmap[DPNPFuncName::DPNP_FN_COV][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_cov_default_c}; - fmap[DPNPFuncName::DPNP_FN_COV][eft_FLT][eft_FLT] = { - eft_DBL, (void *)dpnp_cov_default_c}; - fmap[DPNPFuncName::DPNP_FN_COV][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_cov_default_c}; - - fmap[DPNPFuncName::DPNP_FN_MAX][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_max_default_c}; - fmap[DPNPFuncName::DPNP_FN_MAX][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_max_default_c}; - fmap[DPNPFuncName::DPNP_FN_MAX][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_max_default_c}; - fmap[DPNPFuncName::DPNP_FN_MAX][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_max_default_c}; - - fmap[DPNPFuncName::DPNP_FN_MEAN][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_mean_default_c}; - fmap[DPNPFuncName::DPNP_FN_MEAN][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_mean_default_c}; - fmap[DPNPFuncName::DPNP_FN_MEAN][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_mean_default_c}; - fmap[DPNPFuncName::DPNP_FN_MEAN][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_mean_default_c}; - fmap[DPNPFuncName::DPNP_FN_MEDIAN][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_median_default_c}; fmap[DPNPFuncName::DPNP_FN_MEDIAN][eft_LNG][eft_LNG] = { @@ -1342,41 +308,5 @@ void func_map_init_statistics(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_MEDIAN_EXT][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_median_ext_c}; - fmap[DPNPFuncName::DPNP_FN_MIN][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_min_default_c}; - fmap[DPNPFuncName::DPNP_FN_MIN][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_min_default_c}; - fmap[DPNPFuncName::DPNP_FN_MIN][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_min_default_c}; - fmap[DPNPFuncName::DPNP_FN_MIN][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_min_default_c}; - - fmap[DPNPFuncName::DPNP_FN_NANVAR][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_nanvar_default_c}; - fmap[DPNPFuncName::DPNP_FN_NANVAR][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_nanvar_default_c}; - fmap[DPNPFuncName::DPNP_FN_NANVAR][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_nanvar_default_c}; - fmap[DPNPFuncName::DPNP_FN_NANVAR][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_nanvar_default_c}; - - fmap[DPNPFuncName::DPNP_FN_STD][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_std_default_c}; - fmap[DPNPFuncName::DPNP_FN_STD][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_std_default_c}; - fmap[DPNPFuncName::DPNP_FN_STD][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_std_default_c}; - fmap[DPNPFuncName::DPNP_FN_STD][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_std_default_c}; - - fmap[DPNPFuncName::DPNP_FN_VAR][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_var_default_c}; - fmap[DPNPFuncName::DPNP_FN_VAR][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_var_default_c}; - fmap[DPNPFuncName::DPNP_FN_VAR][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_var_default_c}; - fmap[DPNPFuncName::DPNP_FN_VAR][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_var_default_c}; - return; } diff --git a/dpnp/dpnp_algo/dpnp_algo.pxd b/dpnp/dpnp_algo/dpnp_algo.pxd index 489a20f064cb..1667e9d413a2 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pxd +++ b/dpnp/dpnp_algo/dpnp_algo.pxd @@ -129,23 +129,6 @@ ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_2in_1out_t)(c_dpctl.DPCTLSyclQueueRef, const size_t, const long * , const c_dpctl.DPCTLEventVectorRef) -ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_2in_1out_strides_t)(c_dpctl.DPCTLSyclQueueRef, - void *, - const size_t, - const size_t, - const shape_elem_type * , - const shape_elem_type * , - void *, - const size_t, - const size_t, - const shape_elem_type * , - const shape_elem_type * , - void *, - const size_t, const size_t, - const shape_elem_type * , - const shape_elem_type * , - const long * , - const c_dpctl.DPCTLEventVectorRef) except + """ diff --git a/dpnp/dpnp_algo/dpnp_algo.pyx b/dpnp/dpnp_algo/dpnp_algo.pyx index f19dcc952230..994f489b6c38 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pyx +++ b/dpnp/dpnp_algo/dpnp_algo.pyx @@ -65,15 +65,6 @@ include "dpnp_algo_special.pxi" include "dpnp_algo_statistics.pxi" -ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_flatten_t)(c_dpctl.DPCTLSyclQueueRef, - void *, const size_t, const size_t, - const shape_elem_type * , const shape_elem_type * , - void *, const size_t, const size_t, - const shape_elem_type * , const shape_elem_type * , - const long * , - const c_dpctl.DPCTLEventVectorRef) - - cpdef dpnp_queue_initialize(): """ Initialize SYCL queue which will be used for any library operations. From 88dc4e325d6beffab5d5fc4819971a9565d58b1e Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Wed, 16 Oct 2024 08:44:13 -0700 Subject: [PATCH 2/5] remove cov from dpnp_iface.hpp --- dpnp/backend/include/dpnp_iface.hpp | 24 ------------------------ 1 file changed, 24 deletions(-) diff --git a/dpnp/backend/include/dpnp_iface.hpp b/dpnp/backend/include/dpnp_iface.hpp index 1c76fe7adf07..740ef617a875 100644 --- a/dpnp/backend/include/dpnp_iface.hpp +++ b/dpnp/backend/include/dpnp_iface.hpp @@ -411,30 +411,6 @@ INP_DLLEXPORT void dpnp_correlate_c(void *result_out, const size_t input2_shape_ndim, const size_t *where); -/** - * @ingroup BACKEND_API - * @brief Custom implementation of cov function with math library and PSTL - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array Input array. - * @param [out] result Output array. - * @param [in] nrows Number of rows in input array. - * @param [in] ncols Number of columns in input array. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_cov_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result1, - size_t nrows, - size_t ncols, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void - dpnp_cov_c(void *array1_in, void *result1, size_t nrows, size_t ncols); - /** * @ingroup BACKEND_API * @brief Construct an array from an index array and a list of arrays to choose From 0e9e826812864d1dcad443899804279f72a4ebfd Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Wed, 16 Oct 2024 09:20:13 -0700 Subject: [PATCH 3/5] clean up legacy sqrt --- .../include/dpnp_gen_1arg_2type_tbl.hpp | 92 ------- dpnp/backend/include/dpnp_iface.hpp | 22 -- dpnp/backend/include/dpnp_iface_fptr.hpp | 11 +- dpnp/backend/kernels/dpnp_krnl_elemwise.cpp | 240 ------------------ 4 files changed, 4 insertions(+), 361 deletions(-) delete mode 100644 dpnp/backend/include/dpnp_gen_1arg_2type_tbl.hpp diff --git a/dpnp/backend/include/dpnp_gen_1arg_2type_tbl.hpp b/dpnp/backend/include/dpnp_gen_1arg_2type_tbl.hpp deleted file mode 100644 index 018a625a606a..000000000000 --- a/dpnp/backend/include/dpnp_gen_1arg_2type_tbl.hpp +++ /dev/null @@ -1,92 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016-2024, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -/* - * This header file contains single argument element wise functions definitions - * - * Macro `MACRO_1ARG_2TYPES_OP` must be defined before usage - * - * Parameters: - * - public name of the function and kernel name - * - operation used to calculate the result - * - mkl operation used to calculate the result - * - */ - -#ifndef MACRO_1ARG_2TYPES_OP -#error "MACRO_1ARG_2TYPES_OP is not defined" -#endif - -#ifdef _SECTION_DOCUMENTATION_GENERATION_ - -#define MACRO_1ARG_2TYPES_OP(__name__, __operation1__, __operation2__) \ - /** @ingroup BACKEND_API */ \ - /** @brief Per element operation function __name__ */ \ - /** */ \ - /** Function "__name__" executes operator "__operation1__" over each \ - * element of the array */ \ - /** */ \ - /** @param[in] q_ref Reference to SYCL queue. */ \ - /** @param[out] result_out Output array. */ \ - /** @param[in] result_size Output array size. */ \ - /** @param[in] result_ndim Number of output array dimensions. \ - */ \ - /** @param[in] result_shape Output array shape. */ \ - /** @param[in] result_strides Output array strides. */ \ - /** @param[in] input1_in Input array 1. */ \ - /** @param[in] input1_size Input array 1 size. */ \ - /** @param[in] input1_ndim Number of input array 1 dimensions. \ - */ \ - /** @param[in] input1_shape Input array 1 shape. */ \ - /** @param[in] input1_strides Input array 1 strides. */ \ - /** @param[in] where Where condition. */ \ - /** @param[in] dep_event_vec_ref Reference to vector of SYCL events. \ - */ \ - template \ - DPCTLSyclEventRef __name__( \ - DPCTLSyclQueueRef q_ref, void *result_out, const size_t result_size, \ - const size_t result_ndim, const shape_elem_type *result_shape, \ - const shape_elem_type *result_strides, const void *input1_in, \ - const size_t input1_size, const size_t input1_ndim, \ - const shape_elem_type *input1_shape, \ - const shape_elem_type *input1_strides, const size_t *where, \ - const DPCTLEventVectorRef dep_event_vec_ref); \ - \ - template \ - void __name__( \ - void *result_out, const size_t result_size, const size_t result_ndim, \ - const shape_elem_type *result_shape, \ - const shape_elem_type *result_strides, const void *input1_in, \ - const size_t input1_size, const size_t input1_ndim, \ - const shape_elem_type *input1_shape, \ - const shape_elem_type *input1_strides, const size_t *where); - -#endif - -MACRO_1ARG_2TYPES_OP(dpnp_sqrt_c, - sycl::sqrt(input_elem), - oneapi::mkl::vm::sqrt(q, input1_size, input1_data, result)) - -#undef MACRO_1ARG_2TYPES_OP diff --git a/dpnp/backend/include/dpnp_iface.hpp b/dpnp/backend/include/dpnp_iface.hpp index 740ef617a875..029a4a6d903e 100644 --- a/dpnp/backend/include/dpnp_iface.hpp +++ b/dpnp/backend/include/dpnp_iface.hpp @@ -561,28 +561,6 @@ INP_DLLEXPORT void dpnp_argmin_c(void *array, void *result, size_t size); #include -#define MACRO_1ARG_2TYPES_OP(__name__, __operation1__, __operation2__) \ - template \ - INP_DLLEXPORT DPCTLSyclEventRef __name__( \ - DPCTLSyclQueueRef q_ref, void *result_out, const size_t result_size, \ - const size_t result_ndim, const shape_elem_type *result_shape, \ - const shape_elem_type *result_strides, const void *input1_in, \ - const size_t input1_size, const size_t input1_ndim, \ - const shape_elem_type *input1_shape, \ - const shape_elem_type *input1_strides, const size_t *where, \ - const DPCTLEventVectorRef dep_event_vec_ref); \ - \ - template \ - INP_DLLEXPORT void __name__( \ - void *result_out, const size_t result_size, const size_t result_ndim, \ - const shape_elem_type *result_shape, \ - const shape_elem_type *result_strides, const void *input1_in, \ - const size_t input1_size, const size_t input1_ndim, \ - const shape_elem_type *input1_shape, \ - const shape_elem_type *input1_strides, const size_t *where); - -#include - #define MACRO_2ARG_3TYPES_OP(__name__, __operation__, __vec_operation__, \ __vec_types__, __mkl_operation__, __mkl_types__) \ template \ - class __name__##_kernel; \ - \ - template \ - class __name__##_strides_kernel; \ - \ - template \ - DPCTLSyclEventRef __name__( \ - DPCTLSyclQueueRef q_ref, void *result_out, const size_t result_size, \ - const size_t result_ndim, const shape_elem_type *result_shape, \ - const shape_elem_type *result_strides, const void *input1_in, \ - const size_t input1_size, const size_t input1_ndim, \ - const shape_elem_type *input1_shape, \ - const shape_elem_type *input1_strides, const size_t *where, \ - const DPCTLEventVectorRef dep_event_vec_ref) \ - { \ - /* avoid warning unused variable*/ \ - (void)result_shape; \ - (void)where; \ - (void)dep_event_vec_ref; \ - \ - DPCTLSyclEventRef event_ref = nullptr; \ - \ - if (!input1_size) { \ - return event_ref; \ - } \ - \ - sycl::queue q = *(reinterpret_cast(q_ref)); \ - \ - _DataType_input *input1_data = \ - static_cast<_DataType_input *>(const_cast(input1_in)); \ - _DataType_output *result = \ - static_cast<_DataType_output *>(result_out); \ - \ - shape_elem_type *input1_shape_offsets = \ - new shape_elem_type[input1_ndim]; \ - \ - get_shape_offsets_inkernel(input1_shape, input1_ndim, \ - input1_shape_offsets); \ - bool use_strides = !array_equal(input1_strides, input1_ndim, \ - input1_shape_offsets, input1_ndim); \ - delete[] input1_shape_offsets; \ - \ - sycl::event event; \ - sycl::range<1> gws(result_size); \ - \ - if (use_strides) { \ - if (result_ndim != input1_ndim) { \ - throw std::runtime_error( \ - "Result ndim=" + std::to_string(result_ndim) + \ - " mismatches with input1 ndim=" + \ - std::to_string(input1_ndim)); \ - } \ - \ - /* memory transfer optimization, use USM-host for temporary speeds \ - * up transfer to device */ \ - using usm_host_allocatorT = \ - sycl::usm_allocator; \ - \ - size_t strides_size = 2 * result_ndim; \ - shape_elem_type *dev_strides_data = \ - sycl::malloc_device(strides_size, q); \ - \ - /* create host temporary for packed strides managed by shared \ - * pointer */ \ - auto strides_host_packed = \ - std::vector( \ - strides_size, usm_host_allocatorT(q)); \ - \ - /* packed vector is concatenation of result_strides, \ - * input1_strides and input2_strides */ \ - std::copy(result_strides, result_strides + result_ndim, \ - strides_host_packed.begin()); \ - std::copy(input1_strides, input1_strides + result_ndim, \ - strides_host_packed.begin() + result_ndim); \ - \ - auto copy_strides_ev = q.copy( \ - strides_host_packed.data(), dev_strides_data, \ - strides_host_packed.size()); \ - \ - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { \ - size_t output_id = global_id[0]; /* for (size_t i = 0; i < \ - result_size; ++i) */ \ - { \ - const shape_elem_type *result_strides_data = \ - &dev_strides_data[0]; \ - const shape_elem_type *input1_strides_data = \ - &dev_strides_data[result_ndim]; \ - \ - size_t input_id = 0; \ - for (size_t i = 0; i < input1_ndim; ++i) { \ - const size_t output_xyz_id = \ - get_xyz_id_by_id_inkernel(output_id, \ - result_strides_data, \ - result_ndim, i); \ - input_id += output_xyz_id * input1_strides_data[i]; \ - } \ - \ - const _DataType_output input_elem = input1_data[input_id]; \ - result[output_id] = __operation1__; \ - } \ - }; \ - auto kernel_func = [&](sycl::handler &cgh) { \ - cgh.parallel_for>( \ - gws, kernel_parallel_for_func); \ - }; \ - \ - q.submit(kernel_func).wait(); \ - \ - sycl::free(dev_strides_data, q); \ - return event_ref; \ - } \ - else { \ - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { \ - size_t output_id = global_id[0]; /* for (size_t i = 0; i < \ - result_size; ++i) */ \ - { \ - const _DataType_output input_elem = \ - input1_data[output_id]; \ - result[output_id] = __operation1__; \ - } \ - }; \ - auto kernel_func = [&](sycl::handler &cgh) { \ - cgh.parallel_for>( \ - gws, kernel_parallel_for_func); \ - }; \ - \ - if constexpr (both_types_are_same<_DataType_input, \ - _DataType_output, float, \ - double>) \ - { \ - if (q.get_device().has(sycl::aspect::fp64)) { \ - event = __operation2__; \ - \ - event_ref = reinterpret_cast(&event); \ - return DPCTLEvent_Copy(event_ref); \ - } \ - } \ - event = q.submit(kernel_func); \ - } \ - \ - event_ref = reinterpret_cast(&event); \ - return DPCTLEvent_Copy(event_ref); \ - } \ - \ - template \ - void __name__( \ - void *result_out, const size_t result_size, const size_t result_ndim, \ - const shape_elem_type *result_shape, \ - const shape_elem_type *result_strides, const void *input1_in, \ - const size_t input1_size, const size_t input1_ndim, \ - const shape_elem_type *input1_shape, \ - const shape_elem_type *input1_strides, const size_t *where) \ - { \ - DPCTLSyclQueueRef q_ref = \ - reinterpret_cast(&DPNP_QUEUE); \ - DPCTLEventVectorRef dep_event_vec_ref = nullptr; \ - DPCTLSyclEventRef event_ref = \ - __name__<_DataType_input, _DataType_output>( \ - q_ref, result_out, result_size, result_ndim, result_shape, \ - result_strides, input1_in, input1_size, input1_ndim, \ - input1_shape, input1_strides, where, dep_event_vec_ref); \ - DPCTLEvent_WaitAndThrow(event_ref); \ - DPCTLEvent_Delete(event_ref); \ - } \ - \ - template \ - void (*__name__##_default)( \ - void *, const size_t, const size_t, const shape_elem_type *, \ - const shape_elem_type *, const void *, const size_t, const size_t, \ - const shape_elem_type *, const shape_elem_type *, const size_t *) = \ - __name__<_DataType_input, _DataType_output>; \ - \ - template \ - DPCTLSyclEventRef (*__name__##_ext)( \ - DPCTLSyclQueueRef, void *, const size_t, const size_t, \ - const shape_elem_type *, const shape_elem_type *, const void *, \ - const size_t, const size_t, const shape_elem_type *, \ - const shape_elem_type *, const size_t *, const DPCTLEventVectorRef) = \ - __name__<_DataType_input, _DataType_output>; - -#include - -static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) -{ - - fmap[DPNPFuncName::DPNP_FN_SQRT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_sqrt_c_default}; - fmap[DPNPFuncName::DPNP_FN_SQRT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_sqrt_c_default}; - fmap[DPNPFuncName::DPNP_FN_SQRT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_sqrt_c_default}; - fmap[DPNPFuncName::DPNP_FN_SQRT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_sqrt_c_default}; - - // Used in dpnp_std_c - fmap[DPNPFuncName::DPNP_FN_SQRT_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_sqrt_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SQRT_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_sqrt_c_ext}; - - return; -} - template constexpr T dispatch_erf_op(T elem) { @@ -261,36 +52,6 @@ constexpr T dispatch_erf_op(T elem) } } -template -constexpr auto dispatch_fmod_op(T elem1, T elem2) -{ - if constexpr (sycl::detail::is_integral::value) { - if constexpr (sycl::detail::is_vec::value) { - T rem; - using ElemT = typename T::element_type; -#pragma unroll - for (size_t i = 0; i < rem.size(); i++) { - if (elem2[i] == ElemT(0)) { - rem[i] = ElemT(0); - } - else { - rem[i] = elem1[i] % elem2[i]; - } - } - return rem; - } - else { - if (elem2 == T(0)) { - return T(0); - } - return elem1 % elem2; - } - } - else { - return sycl::fmod(elem1, elem2); - } -} - #define MACRO_1ARG_1TYPE_OP(__name__, __operation1__, __operation2__) \ template \ class __name__##_kernel; \ @@ -1001,7 +762,6 @@ static void func_map_init_elemwise_2arg_3type(func_map_t &fmap) void func_map_init_elemwise(func_map_t &fmap) { func_map_init_elemwise_1arg_1type(fmap); - func_map_init_elemwise_1arg_2type(fmap); func_map_init_elemwise_2arg_3type(fmap); return; From 20abff035af11bdebc473aa1d47562b0f7041427 Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Wed, 16 Oct 2024 19:04:14 -0700 Subject: [PATCH 4/5] clean up legacy backend --- dpnp/backend/CMakeLists.txt | 2 - dpnp/backend/examples/example8.cpp | 67 ---- dpnp/backend/examples/example9.cpp | 65 ---- dpnp/backend/include/dpnp_iface.hpp | 189 ---------- dpnp/backend/include/dpnp_iface_fptr.hpp | 6 - dpnp/backend/kernels/dpnp_krnl_reduction.cpp | 363 ------------------- dpnp/backend/kernels/dpnp_krnl_searching.cpp | 165 --------- dpnp/backend/kernels/dpnp_krnl_sorting.cpp | 222 ------------ dpnp/backend/src/dpnp_fptr.hpp | 53 --- dpnp/backend/src/dpnp_iface_fptr.cpp | 14 - 10 files changed, 1146 deletions(-) delete mode 100644 dpnp/backend/examples/example8.cpp delete mode 100644 dpnp/backend/examples/example9.cpp delete mode 100644 dpnp/backend/kernels/dpnp_krnl_reduction.cpp delete mode 100644 dpnp/backend/kernels/dpnp_krnl_searching.cpp diff --git a/dpnp/backend/CMakeLists.txt b/dpnp/backend/CMakeLists.txt index 64d1e81c9c28..1924ad64027f 100644 --- a/dpnp/backend/CMakeLists.txt +++ b/dpnp/backend/CMakeLists.txt @@ -30,8 +30,6 @@ set(DPNP_SRC kernels/dpnp_krnl_indexing.cpp kernels/dpnp_krnl_mathematical.cpp kernels/dpnp_krnl_random.cpp - kernels/dpnp_krnl_reduction.cpp - kernels/dpnp_krnl_searching.cpp kernels/dpnp_krnl_sorting.cpp kernels/dpnp_krnl_statistics.cpp src/constants.cpp diff --git a/dpnp/backend/examples/example8.cpp b/dpnp/backend/examples/example8.cpp deleted file mode 100644 index 19074c5c8fc3..000000000000 --- a/dpnp/backend/examples/example8.cpp +++ /dev/null @@ -1,67 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016-2024, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -/** - * Example 9. - * - * TODO explanation of the example - * - * Possible compile line: - * . /opt/intel/oneapi/setvars.sh - * g++ -g dpnp/backend/examples/example8.cpp -Idpnp -Idpnp/backend/include - * -Ldpnp -Wl,-rpath='$ORIGIN'/dpnp -ldpnp_backend_c -o example8 - * - */ -#include - -#include "dpnp_iface.hpp" - -int main(int, char **) -{ - const size_t size = 16; - - double *array = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - long *result = (long *)dpnp_memory_alloc_c(size * sizeof(long)); - - std::cout << "array" << std::endl; - for (size_t i = 0; i < size; ++i) { - array[i] = (double)(size - i) / 2; - std::cout << array[i] << ", "; - } - std::cout << std::endl; - - dpnp_argsort_c(array, result, size); - - std::cout << "array with 'sorted' indices" << std::endl; - for (size_t i = 0; i < size; ++i) { - std::cout << result[i] << ", "; - } - std::cout << std::endl; - - dpnp_memory_free_c(result); - dpnp_memory_free_c(array); - - return 0; -} diff --git a/dpnp/backend/examples/example9.cpp b/dpnp/backend/examples/example9.cpp deleted file mode 100644 index fedcb909a7af..000000000000 --- a/dpnp/backend/examples/example9.cpp +++ /dev/null @@ -1,65 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016-2024, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -/** - * Example 9. - * - * This example shows simple usage of the DPNP C++ Backend library - * to calculate sum of the given elements vector - * - * Possible compile line: - * . /opt/intel/oneapi/setvars.sh - * g++ -g dpnp/backend/examples/example9.cpp -Idpnp -Idpnp/backend/include - * -Ldpnp -Wl,-rpath='$ORIGIN'/dpnp -ldpnp_backend_c -o example9 - * - */ - -#include - -#include "dpnp_iface.hpp" - -int main(int, char **) -{ - const size_t size = 2097152; - long result = 0; - long result_verification = 0; - - long *array = - reinterpret_cast(dpnp_memory_alloc_c(size * sizeof(long))); - - for (size_t i = 0; i < size; ++i) { - array[i] = i; - result_verification += i; - } - - dpnp_sum_c(&result, array, &size, 1, NULL, 0, NULL, NULL); - - std::cout << "SUM() value: " << result - << " verification value: " << result_verification << std::endl; - - dpnp_memory_free_c(array); - - return 0; -} diff --git a/dpnp/backend/include/dpnp_iface.hpp b/dpnp/backend/include/dpnp_iface.hpp index 029a4a6d903e..5be448c446d5 100644 --- a/dpnp/backend/include/dpnp_iface.hpp +++ b/dpnp/backend/include/dpnp_iface.hpp @@ -163,52 +163,6 @@ INP_DLLEXPORT void dpnp_dot_c(void *result_out, const shape_elem_type *input2_shape, const shape_elem_type *input2_strides); -/** - * @ingroup BACKEND_API - * @brief Compute summary of input array elements. - * - * Input array is expected as @ref _DataType_input type and assume result as - * @ref _DataType_output type. The function creates no memory. - * - * Empty @ref input_shape means scalar. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result_out Output array pointer. @ref _DataType_output - * type is expected - * @param [in] input_in Input array pointer. @ref _DataType_input type - * is expected - * @param [in] input_shape Shape of @ref input_in - * @param [in] input_shape_ndim Number of elements in @ref input_shape - * @param [in] axes Array of axes to apply to @ref input_shape - * @param [in] axes_ndim Number of elements in @ref axes - * @param [in] initial Pointer to initial value for the algorithm. - * @ref _DataType_input is expected - * @param [in] where mask array - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_sum_c(DPCTLSyclQueueRef q_ref, - void *result_out, - const void *input_in, - const shape_elem_type *input_shape, - const size_t input_shape_ndim, - const shape_elem_type *axes, - const size_t axes_ndim, - const void *initial, - const long *where, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_sum_c(void *result_out, - const void *input_in, - const shape_elem_type *input_shape, - const size_t input_shape_ndim, - const shape_elem_type *axes, - const size_t axes_ndim, - const void *initial, - const long *where); - /** * @ingroup BACKEND_API * @brief Return a partitioned copy of an array. @@ -241,107 +195,6 @@ INP_DLLEXPORT void dpnp_partition_c(void *array, const shape_elem_type *shape, const size_t ndim); -/** - * @ingroup BACKEND_API - * @brief Compute Product of input array elements. - * - * Input array is expected as @ref _DataType_input type and assume result as - * @ref _DataType_output type. The function creates no memory. - * - * Empty @ref input_shape means scalar. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result_out Output array pointer. @ref _DataType_output - * type is expected - * @param [in] input_in Input array pointer. @ref _DataType_input type - * is expected - * @param [in] input_shape Shape of @ref input_in - * @param [in] input_shape_ndim Number of elements in @ref input_shape - * @param [in] axes Array of axes to apply to @ref input_shape - * @param [in] axes_ndim Number of elements in @ref axes - * @param [in] initial Pointer to initial value for the algorithm. - * @ref _DataType_input is expected - * @param [in] where mask array - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_prod_c(DPCTLSyclQueueRef q_ref, - void *result_out, - const void *input_in, - const shape_elem_type *input_shape, - const size_t input_shape_ndim, - const shape_elem_type *axes, - const size_t axes_ndim, - const void *initial, - const long *where, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_prod_c(void *result_out, - const void *input_in, - const shape_elem_type *input_shape, - const size_t input_shape_ndim, - const shape_elem_type *axes, - const size_t axes_ndim, - const void *initial, - const long *where); - -/** - * @ingroup BACKEND_API - - * @brief math library implementation of argsort function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array Input array with data. - * @param [out] result Output array with indices. - * @param [in] size Number of elements in input arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_argsort_c(DPCTLSyclQueueRef q_ref, - void *array, - void *result, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_argsort_c(void *array, void *result, size_t size); - -/** - * @ingroup BACKEND_API - * @brief math library implementation of searchsorted function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] array Input array with data. - * @param [in] v Input values to insert into array. - * @param [in] side Param for choosing a case of searching for - * elements. - * @param [in] arr_size Number of elements in input arrays. - * @param [in] v_size Number of elements in input values arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_searchsorted_c(DPCTLSyclQueueRef q_ref, - void *result, - const void *array, - const void *v, - bool side, - const size_t arr_size, - const size_t v_size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_searchsorted_c(void *result, - const void *array, - const void *v, - bool side, - const size_t arr_size, - const size_t v_size); - /** * @ingroup BACKEND_API * @brief math library implementation of sort function @@ -497,48 +350,6 @@ INP_DLLEXPORT void dpnp_median_c(void *array, const shape_elem_type *axis, size_t naxis); -/** - * @ingroup BACKEND_API - * @brief math library implementation of argmax function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array Input array with data. - * @param [out] result Output array with indices. - * @param [in] size Number of elements in input array. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_argmax_c(DPCTLSyclQueueRef q_ref, - void *array, - void *result, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_argmax_c(void *array, void *result, size_t size); - -/** - * @ingroup BACKEND_API - * @brief math library implementation of argmin function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array Input array with data. - * @param [out] result Output array with indices. - * @param [in] size Number of elements in input array. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_argmin_c(DPCTLSyclQueueRef q_ref, - void *array, - void *result, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_argmin_c(void *array, void *result, size_t size); - #define MACRO_1ARG_1TYPE_OP(__name__, __operation1__, __operation2__) \ template \ INP_DLLEXPORT DPCTLSyclEventRef __name__( \ diff --git a/dpnp/backend/include/dpnp_iface_fptr.hpp b/dpnp/backend/include/dpnp_iface_fptr.hpp index 86cfe054c6ba..f332ce29081d 100644 --- a/dpnp/backend/include/dpnp_iface_fptr.hpp +++ b/dpnp/backend/include/dpnp_iface_fptr.hpp @@ -59,9 +59,6 @@ enum class DPNPFuncName : size_t { DPNP_FN_NONE, /**< Very first element of the enumeration */ - DPNP_FN_ARGMAX, /**< Used in numpy.argmax() impl */ - DPNP_FN_ARGMIN, /**< Used in numpy.argmin() impl */ - DPNP_FN_ARGSORT, /**< Used in numpy.argsort() impl */ DPNP_FN_CHOOSE, /**< Used in numpy.choose() impl */ DPNP_FN_CHOOSE_EXT, /**< Used in numpy.choose() impl, requires extra parameters */ @@ -89,7 +86,6 @@ enum class DPNPFuncName : size_t DPNP_FN_PARTITION, /**< Used in numpy.partition() impl */ DPNP_FN_PARTITION_EXT, /**< Used in numpy.partition() impl, requires extra parameters */ - DPNP_FN_PROD, /**< Used in numpy.prod() impl */ DPNP_FN_RNG_BETA, /**< Used in numpy.random.beta() impl */ DPNP_FN_RNG_BETA_EXT, /**< Used in numpy.random.beta() impl, requires extra parameters */ @@ -207,9 +203,7 @@ enum class DPNPFuncName : size_t DPNP_FN_RNG_ZIPF, /**< Used in numpy.random.zipf() impl */ DPNP_FN_RNG_ZIPF_EXT, /**< Used in numpy.random.zipf() impl, requires extra parameters */ - DPNP_FN_SEARCHSORTED, /**< Used in numpy.searchsorted() impl */ DPNP_FN_SORT, /**< Used in numpy.sort() impl */ - DPNP_FN_SUM, /**< Used in numpy.sum() impl */ DPNP_FN_ZEROS, /**< Used in numpy.zeros() impl */ DPNP_FN_ZEROS_LIKE, /**< Used in numpy.zeros_like() impl */ DPNP_FN_LAST, /**< The latest element of the enumeration */ diff --git a/dpnp/backend/kernels/dpnp_krnl_reduction.cpp b/dpnp/backend/kernels/dpnp_krnl_reduction.cpp deleted file mode 100644 index a9383615e065..000000000000 --- a/dpnp/backend/kernels/dpnp_krnl_reduction.cpp +++ /dev/null @@ -1,363 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016-2024, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -#include -#include - -#include "dpnp_fptr.hpp" -#include "dpnp_iterator.hpp" -#include "dpnp_utils.hpp" -#include "dpnpc_memory_adapter.hpp" -#include "queue_sycl.hpp" -#include - -namespace mkl_stats = oneapi::mkl::stats; - -template -_DataType *get_array_ptr(const void *__array) -{ - void *const_ptr = const_cast(__array); - _DataType *ptr = reinterpret_cast<_DataType *>(const_ptr); - - return ptr; -} - -template -_DataType get_initial_value(const void *__initial, _DataType default_val) -{ - const _DataType *initial_ptr = - reinterpret_cast(__initial); - const _DataType init_val = - (initial_ptr == nullptr) ? default_val : *initial_ptr; - - return init_val; -} - -template -class dpnp_sum_c_kernel; - -template -DPCTLSyclEventRef - dpnp_sum_c(DPCTLSyclQueueRef q_ref, - void *result_out, - const void *input_in, - const shape_elem_type *input_shape, - const size_t input_shape_ndim, - const shape_elem_type *axes, - const size_t axes_ndim, - const void *initial, // type must be _DataType_output - const long *where, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)where; - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if ((input_in == nullptr) || (result_out == nullptr)) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - const _DataType_output init = - get_initial_value<_DataType_output>(initial, 0); - - const size_t input_size = - std::accumulate(input_shape, input_shape + input_shape_ndim, 1, - std::multiplies()); - - DPNPC_ptr_adapter<_DataType_input> input1_ptr(q_ref, input_in, input_size, - true); - _DataType_input *input = input1_ptr.get_ptr(); - _DataType_output *result = get_array_ptr<_DataType_output>(result_out); - - if (!input_shape && !input_shape_ndim) { // it is a scalar - // result[0] = input[0]; - _DataType_input input_elem = 0; - _DataType_output result_elem = 0; - q.memcpy(&input_elem, input, sizeof(_DataType_input)).wait(); - result_elem = input_elem; - q.memcpy(result, &result_elem, sizeof(_DataType_output)).wait(); - - return event_ref; - } - - if constexpr ((std::is_same<_DataType_input, double>::value || - std::is_same<_DataType_input, float>::value) && - std::is_same<_DataType_input, _DataType_output>::value) - { - // Support is limited by - // - 1D array (no axes) - // - same types for input and output - // - float64 and float32 types only - if (axes_ndim < 1) { - auto dataset = - mkl_stats::make_dataset( - 1, input_size, input); - sycl::event event = mkl_stats::raw_sum(q, dataset, result); - event.wait(); - - return event_ref; - } - } - - DPNPC_id<_DataType_input> input_it(q_ref, input, input_shape, - input_shape_ndim); - input_it.set_axes(axes, axes_ndim); - - const size_t output_size = input_it.get_output_size(); - auto policy = oneapi::dpl::execution::make_device_policy< - dpnp_sum_c_kernel<_DataType_output, _DataType_input>>(q); - for (size_t output_id = 0; output_id < output_size; ++output_id) { - // type of "init" determine internal algorithm accumulator type - _DataType_output accumulator = std::reduce( - policy, input_it.begin(output_id), input_it.end(output_id), init, - std::plus<_DataType_output>()); - policy.queue().wait(); // TODO move out of the loop - - q.memcpy(result + output_id, &accumulator, sizeof(_DataType_output)) - .wait(); // result[output_id] = accumulator; - } - - return event_ref; -} - -template -void dpnp_sum_c(void *result_out, - const void *input_in, - const shape_elem_type *input_shape, - const size_t input_shape_ndim, - const shape_elem_type *axes, - const size_t axes_ndim, - const void *initial, // type must be _DataType_output - const long *where) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_sum_c<_DataType_output, _DataType_input>( - q_ref, result_out, input_in, input_shape, input_shape_ndim, axes, - axes_ndim, initial, where, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_sum_default_c)(void *, - const void *, - const shape_elem_type *, - const size_t, - const shape_elem_type *, - const size_t, - const void *, - const long *) = - dpnp_sum_c<_DataType_output, _DataType_input>; - -template -class dpnp_prod_c_kernel; - -template -DPCTLSyclEventRef - dpnp_prod_c(DPCTLSyclQueueRef q_ref, - void *result_out, - const void *input_in, - const shape_elem_type *input_shape, - const size_t input_shape_ndim, - const shape_elem_type *axes, - const size_t axes_ndim, - const void *initial, // type must be _DataType_output - const long *where, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)where; - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if ((input_in == nullptr) || (result_out == nullptr)) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - const _DataType_output init = - get_initial_value<_DataType_output>(initial, 1); - - const size_t input_size = - std::accumulate(input_shape, input_shape + input_shape_ndim, 1, - std::multiplies()); - - DPNPC_ptr_adapter<_DataType_input> input1_ptr(q_ref, input_in, input_size, - true); - _DataType_input *input = input1_ptr.get_ptr(); - _DataType_output *result = get_array_ptr<_DataType_output>(result_out); - - if (!input_shape && !input_shape_ndim) { // it is a scalar - // result[0] = input[0]; - _DataType_input input_elem = 0; - _DataType_output result_elem = 0; - q.memcpy(&input_elem, input, sizeof(_DataType_input)).wait(); - result_elem = input_elem; - q.memcpy(result, &result_elem, sizeof(_DataType_output)).wait(); - - return event_ref; - } - - DPNPC_id<_DataType_input> input_it(q_ref, input, input_shape, - input_shape_ndim); - input_it.set_axes(axes, axes_ndim); - - const size_t output_size = input_it.get_output_size(); - auto policy = oneapi::dpl::execution::make_device_policy< - dpnp_prod_c_kernel<_DataType_output, _DataType_input>>(q); - for (size_t output_id = 0; output_id < output_size; ++output_id) { - // type of "init" determine internal algorithm accumulator type - _DataType_output accumulator = std::reduce( - policy, input_it.begin(output_id), input_it.end(output_id), init, - std::multiplies<_DataType_output>()); - policy.queue().wait(); // TODO move out of the loop - - q.memcpy(result + output_id, &accumulator, sizeof(_DataType_output)) - .wait(); // result[output_id] = accumulator; - } - - return event_ref; -} - -template -void dpnp_prod_c(void *result_out, - const void *input_in, - const shape_elem_type *input_shape, - const size_t input_shape_ndim, - const shape_elem_type *axes, - const size_t axes_ndim, - const void *initial, // type must be _DataType_output - const long *where) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_prod_c<_DataType_output, _DataType_input>( - q_ref, result_out, input_in, input_shape, input_shape_ndim, axes, - axes_ndim, initial, where, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_prod_default_c)(void *, - const void *, - const shape_elem_type *, - const size_t, - const shape_elem_type *, - const size_t, - const void *, - const long *) = - dpnp_prod_c<_DataType_output, _DataType_input>; - -void func_map_init_reduction(func_map_t &fmap) -{ - // WARNING. The meaning of the fmap is changed. Second argument represents - // RESULT_TYPE for this function handle "out" and "type" parameters require - // user selection of return type - // TODO. required refactoring of fmap to some kernelSelector - fmap[DPNPFuncName::DPNP_FN_PROD][eft_INT][eft_INT] = { - eft_LNG, (void *)dpnp_prod_default_c}; - fmap[DPNPFuncName::DPNP_FN_PROD][eft_INT][eft_LNG] = { - eft_LNG, (void *)dpnp_prod_default_c}; - fmap[DPNPFuncName::DPNP_FN_PROD][eft_INT][eft_FLT] = { - eft_FLT, (void *)dpnp_prod_default_c}; - fmap[DPNPFuncName::DPNP_FN_PROD][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_prod_default_c}; - - fmap[DPNPFuncName::DPNP_FN_PROD][eft_LNG][eft_INT] = { - eft_INT, (void *)dpnp_prod_default_c}; - fmap[DPNPFuncName::DPNP_FN_PROD][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_prod_default_c}; - fmap[DPNPFuncName::DPNP_FN_PROD][eft_LNG][eft_FLT] = { - eft_FLT, (void *)dpnp_prod_default_c}; - fmap[DPNPFuncName::DPNP_FN_PROD][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_prod_default_c}; - - fmap[DPNPFuncName::DPNP_FN_PROD][eft_FLT][eft_INT] = { - eft_INT, (void *)dpnp_prod_default_c}; - fmap[DPNPFuncName::DPNP_FN_PROD][eft_FLT][eft_LNG] = { - eft_LNG, (void *)dpnp_prod_default_c}; - fmap[DPNPFuncName::DPNP_FN_PROD][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_prod_default_c}; - fmap[DPNPFuncName::DPNP_FN_PROD][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_prod_default_c}; - - fmap[DPNPFuncName::DPNP_FN_PROD][eft_DBL][eft_INT] = { - eft_INT, (void *)dpnp_prod_default_c}; - fmap[DPNPFuncName::DPNP_FN_PROD][eft_DBL][eft_LNG] = { - eft_LNG, (void *)dpnp_prod_default_c}; - fmap[DPNPFuncName::DPNP_FN_PROD][eft_DBL][eft_FLT] = { - eft_FLT, (void *)dpnp_prod_default_c}; - fmap[DPNPFuncName::DPNP_FN_PROD][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_prod_default_c}; - - fmap[DPNPFuncName::DPNP_FN_SUM][eft_INT][eft_INT] = { - eft_LNG, (void *)dpnp_sum_default_c}; - fmap[DPNPFuncName::DPNP_FN_SUM][eft_INT][eft_LNG] = { - eft_LNG, (void *)dpnp_sum_default_c}; - fmap[DPNPFuncName::DPNP_FN_SUM][eft_INT][eft_FLT] = { - eft_FLT, (void *)dpnp_sum_default_c}; - fmap[DPNPFuncName::DPNP_FN_SUM][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_sum_default_c}; - - fmap[DPNPFuncName::DPNP_FN_SUM][eft_LNG][eft_INT] = { - eft_INT, (void *)dpnp_sum_default_c}; - fmap[DPNPFuncName::DPNP_FN_SUM][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_sum_default_c}; - fmap[DPNPFuncName::DPNP_FN_SUM][eft_LNG][eft_FLT] = { - eft_FLT, (void *)dpnp_sum_default_c}; - fmap[DPNPFuncName::DPNP_FN_SUM][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_sum_default_c}; - - fmap[DPNPFuncName::DPNP_FN_SUM][eft_FLT][eft_INT] = { - eft_INT, (void *)dpnp_sum_default_c}; - fmap[DPNPFuncName::DPNP_FN_SUM][eft_FLT][eft_LNG] = { - eft_LNG, (void *)dpnp_sum_default_c}; - fmap[DPNPFuncName::DPNP_FN_SUM][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_sum_default_c}; - fmap[DPNPFuncName::DPNP_FN_SUM][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_sum_default_c}; - - fmap[DPNPFuncName::DPNP_FN_SUM][eft_DBL][eft_INT] = { - eft_INT, (void *)dpnp_sum_default_c}; - fmap[DPNPFuncName::DPNP_FN_SUM][eft_DBL][eft_LNG] = { - eft_LNG, (void *)dpnp_sum_default_c}; - fmap[DPNPFuncName::DPNP_FN_SUM][eft_DBL][eft_FLT] = { - eft_FLT, (void *)dpnp_sum_default_c}; - fmap[DPNPFuncName::DPNP_FN_SUM][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_sum_default_c}; - - return; -} diff --git a/dpnp/backend/kernels/dpnp_krnl_searching.cpp b/dpnp/backend/kernels/dpnp_krnl_searching.cpp deleted file mode 100644 index 96ff470d7b03..000000000000 --- a/dpnp/backend/kernels/dpnp_krnl_searching.cpp +++ /dev/null @@ -1,165 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016-2024, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -#include - -#include "dpnp_fptr.hpp" -#include "dpnpc_memory_adapter.hpp" -#include "queue_sycl.hpp" -#include - -template -class dpnp_argmax_c_kernel; - -template -DPCTLSyclEventRef dpnp_argmax_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result1, - size_t 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)); - - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, size); - _DataType *array_1 = input1_ptr.get_ptr(); - _idx_DataType *result = reinterpret_cast<_idx_DataType *>(result1); - - auto policy = oneapi::dpl::execution::make_device_policy< - class dpnp_argmax_c_kernel<_DataType, _idx_DataType>>(q); - - _DataType *res = std::max_element(policy, array_1, array_1 + size); - policy.queue().wait(); - - _idx_DataType result_val = std::distance(array_1, res); - q.memcpy(result, &result_val, sizeof(_idx_DataType)) - .wait(); // result[0] = std::distance(array_1, res); - - return event_ref; -} - -template -void dpnp_argmax_c(void *array1_in, void *result1, size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_argmax_c<_DataType, _idx_DataType>( - q_ref, array1_in, result1, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_argmax_default_c)(void *, - void *, - size_t) = dpnp_argmax_c<_DataType, _idx_DataType>; - -template -class dpnp_argmin_c_kernel; - -template -DPCTLSyclEventRef dpnp_argmin_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result1, - size_t 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)); - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, size); - _DataType *array_1 = input1_ptr.get_ptr(); - _idx_DataType *result = reinterpret_cast<_idx_DataType *>(result1); - - auto policy = oneapi::dpl::execution::make_device_policy< - class dpnp_argmin_c_kernel<_DataType, _idx_DataType>>(q); - - _DataType *res = std::min_element(policy, array_1, array_1 + size); - policy.queue().wait(); - - _idx_DataType result_val = std::distance(array_1, res); - q.memcpy(result, &result_val, sizeof(_idx_DataType)) - .wait(); // result[0] = std::distance(array_1, res); - - return event_ref; -} - -template -void dpnp_argmin_c(void *array1_in, void *result1, size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_argmin_c<_DataType, _idx_DataType>( - q_ref, array1_in, result1, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_argmin_default_c)(void *, - void *, - size_t) = dpnp_argmin_c<_DataType, _idx_DataType>; - -void func_map_init_searching(func_map_t &fmap) -{ - fmap[DPNPFuncName::DPNP_FN_ARGMAX][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_argmax_default_c}; - fmap[DPNPFuncName::DPNP_FN_ARGMAX][eft_INT][eft_LNG] = { - eft_LNG, (void *)dpnp_argmax_default_c}; - fmap[DPNPFuncName::DPNP_FN_ARGMAX][eft_LNG][eft_INT] = { - eft_INT, (void *)dpnp_argmax_default_c}; - fmap[DPNPFuncName::DPNP_FN_ARGMAX][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_argmax_default_c}; - fmap[DPNPFuncName::DPNP_FN_ARGMAX][eft_FLT][eft_INT] = { - eft_INT, (void *)dpnp_argmax_default_c}; - fmap[DPNPFuncName::DPNP_FN_ARGMAX][eft_FLT][eft_LNG] = { - eft_LNG, (void *)dpnp_argmax_default_c}; - fmap[DPNPFuncName::DPNP_FN_ARGMAX][eft_DBL][eft_INT] = { - eft_INT, (void *)dpnp_argmax_default_c}; - fmap[DPNPFuncName::DPNP_FN_ARGMAX][eft_DBL][eft_LNG] = { - eft_LNG, (void *)dpnp_argmax_default_c}; - - fmap[DPNPFuncName::DPNP_FN_ARGMIN][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_argmin_default_c}; - fmap[DPNPFuncName::DPNP_FN_ARGMIN][eft_INT][eft_LNG] = { - eft_LNG, (void *)dpnp_argmin_default_c}; - fmap[DPNPFuncName::DPNP_FN_ARGMIN][eft_LNG][eft_INT] = { - eft_INT, (void *)dpnp_argmin_default_c}; - fmap[DPNPFuncName::DPNP_FN_ARGMIN][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_argmin_default_c}; - fmap[DPNPFuncName::DPNP_FN_ARGMIN][eft_FLT][eft_INT] = { - eft_INT, (void *)dpnp_argmin_default_c}; - fmap[DPNPFuncName::DPNP_FN_ARGMIN][eft_FLT][eft_LNG] = { - eft_LNG, (void *)dpnp_argmin_default_c}; - fmap[DPNPFuncName::DPNP_FN_ARGMIN][eft_DBL][eft_INT] = { - eft_INT, (void *)dpnp_argmin_default_c}; - fmap[DPNPFuncName::DPNP_FN_ARGMIN][eft_DBL][eft_LNG] = { - eft_LNG, (void *)dpnp_argmin_default_c}; - - return; -} diff --git a/dpnp/backend/kernels/dpnp_krnl_sorting.cpp b/dpnp/backend/kernels/dpnp_krnl_sorting.cpp index 2c50a490f153..73a15df0a5c8 100644 --- a/dpnp/backend/kernels/dpnp_krnl_sorting.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_sorting.cpp @@ -30,84 +30,6 @@ #include "queue_sycl.hpp" #include -template -struct _argsort_less -{ - _argsort_less(_DataType *data_ptr) - { - _data_ptr = data_ptr; - } - - inline bool operator()(const _idx_DataType &idx1, const _idx_DataType &idx2) - { - return (_data_ptr[idx1] < _data_ptr[idx2]); - } - -private: - _DataType *_data_ptr = nullptr; -}; - -template -class dpnp_argsort_c_kernel; - -template -DPCTLSyclEventRef dpnp_argsort_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result1, - size_t 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)); - - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, size, true); - DPNPC_ptr_adapter<_idx_DataType> result1_ptr(q_ref, result1, size, true, - true); - _DataType *array_1 = input1_ptr.get_ptr(); - _idx_DataType *result = result1_ptr.get_ptr(); - - std::iota(result, result + size, 0); - - auto policy = oneapi::dpl::execution::make_device_policy< - class dpnp_argsort_c_kernel<_DataType, _idx_DataType>>(q); - - std::sort(policy, result, result + size, - _argsort_less<_DataType, _idx_DataType>(array_1)); - - policy.queue().wait(); - - return event_ref; -} - -template -void dpnp_argsort_c(void *array1_in, void *result1, size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_argsort_c<_DataType, _idx_DataType>( - q_ref, array1_in, result1, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_argsort_default_c)(void *, void *, size_t) = - dpnp_argsort_c<_DataType, _idx_DataType>; - -// template void dpnp_argsort_c(void* array1_in, void* result1, -// size_t size); template void dpnp_argsort_c(void* array1_in, -// void* result1, size_t size); template void dpnp_argsort_c(void* -// array1_in, void* result1, size_t size); template void dpnp_argsort_c(void* array1_in, void* result1, size_t size); template void -// dpnp_argsort_c(void* array1_in, void* result1, size_t size); -// template void dpnp_argsort_c(void* array1_in, void* result1, -// size_t size); template void dpnp_argsort_c(void* array1_in, void* -// result1, size_t size); template void dpnp_argsort_c(void* -// array1_in, void* result1, size_t size); - template struct _sort_less { @@ -277,132 +199,6 @@ DPCTLSyclEventRef (*dpnp_partition_ext_c)(DPCTLSyclQueueRef, const DPCTLEventVectorRef) = dpnp_partition_c<_DataType>; -template -class dpnp_searchsorted_c_kernel; - -template -DPCTLSyclEventRef - dpnp_searchsorted_c(DPCTLSyclQueueRef q_ref, - void *result1, - const void *array1_in, - const void *v1_in, - bool side, - const size_t arr_size, - const size_t v_size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if ((array1_in == nullptr) || (v1_in == nullptr) || (result1 == nullptr)) { - return event_ref; - } - - if (arr_size == 0) { - return event_ref; - } - - if (v_size == 0) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, arr_size); - DPNPC_ptr_adapter<_DataType> input2_ptr(q_ref, v1_in, v_size); - const _DataType *arr = input1_ptr.get_ptr(); - const _DataType *v = input2_ptr.get_ptr(); - _IndexingType *result = reinterpret_cast<_IndexingType *>(result1); - - sycl::range<2> gws(v_size, arr_size); - auto kernel_parallel_for_func = [=](sycl::id<2> global_id) { - size_t i = global_id[0]; - size_t j = global_id[1]; - - if (j != 0) { - if (side) { - if (j == arr_size - 1) { - if (v[i] == arr[j]) { - result[i] = arr_size - 1; - } - else { - if (v[i] > arr[j]) { - result[i] = arr_size; - } - } - } - else { - if ((arr[j - 1] < v[i]) && (v[i] <= arr[j])) { - result[i] = j; - } - } - } - else { - if (j == arr_size - 1) { - if ((arr[j - 1] <= v[i]) && (v[i] < arr[j])) { - result[i] = arr_size - 1; - } - else { - if (v[i] == arr[j]) { - result[i] = arr_size; - } - else { - if (v[i] > arr[j]) { - result[i] = arr_size; - } - } - } - } - else { - if ((arr[j - 1] <= v[i]) && (v[i] < arr[j])) { - result[i] = j; - } - } - } - } - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.parallel_for< - class dpnp_searchsorted_c_kernel<_DataType, _IndexingType>>( - gws, kernel_parallel_for_func); - }; - - auto event = q.submit(kernel_func); - - event.wait(); - - return event_ref; -} - -template -void dpnp_searchsorted_c(void *result1, - const void *array1_in, - const void *v1_in, - bool side, - const size_t arr_size, - const size_t v_size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_searchsorted_c<_DataType, _IndexingType>( - q_ref, result1, array1_in, v1_in, side, arr_size, v_size, - dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_searchsorted_default_c)(void *, - const void *, - const void *, - bool, - const size_t, - const size_t) = - dpnp_searchsorted_c<_DataType, _IndexingType>; - template class dpnp_sort_c_kernel; @@ -454,15 +250,6 @@ void (*dpnp_sort_default_c)(void *, void *, size_t) = dpnp_sort_c<_DataType>; void func_map_init_sorting(func_map_t &fmap) { - fmap[DPNPFuncName::DPNP_FN_ARGSORT][eft_INT][eft_INT] = { - eft_LNG, (void *)dpnp_argsort_default_c}; - fmap[DPNPFuncName::DPNP_FN_ARGSORT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_argsort_default_c}; - fmap[DPNPFuncName::DPNP_FN_ARGSORT][eft_FLT][eft_FLT] = { - eft_LNG, (void *)dpnp_argsort_default_c}; - fmap[DPNPFuncName::DPNP_FN_ARGSORT][eft_DBL][eft_DBL] = { - eft_LNG, (void *)dpnp_argsort_default_c}; - fmap[DPNPFuncName::DPNP_FN_PARTITION][eft_INT][eft_INT] = { eft_INT, (void *)dpnp_partition_default_c}; fmap[DPNPFuncName::DPNP_FN_PARTITION][eft_LNG][eft_LNG] = { @@ -487,15 +274,6 @@ void func_map_init_sorting(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_PARTITION_EXT][eft_C128][eft_C128] = { eft_C128, (void *)dpnp_partition_ext_c>}; - fmap[DPNPFuncName::DPNP_FN_SEARCHSORTED][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_searchsorted_default_c}; - fmap[DPNPFuncName::DPNP_FN_SEARCHSORTED][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_searchsorted_default_c}; - fmap[DPNPFuncName::DPNP_FN_SEARCHSORTED][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_searchsorted_default_c}; - fmap[DPNPFuncName::DPNP_FN_SEARCHSORTED][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_searchsorted_default_c}; - fmap[DPNPFuncName::DPNP_FN_SORT][eft_INT][eft_INT] = { eft_INT, (void *)dpnp_sort_default_c}; fmap[DPNPFuncName::DPNP_FN_SORT][eft_LNG][eft_LNG] = { diff --git a/dpnp/backend/src/dpnp_fptr.hpp b/dpnp/backend/src/dpnp_fptr.hpp index f57f9cd449b5..6413224a2139 100644 --- a/dpnp/backend/src/dpnp_fptr.hpp +++ b/dpnp/backend/src/dpnp_fptr.hpp @@ -219,18 +219,6 @@ template using dpnp_remove_cvref_t = typename std::remove_cv_t>; -/** - * A helper alias template to return true value for complex types and false - * otherwise. - */ -template -struct is_complex : public std::integral_constant< - bool, - std::is_same_v<_Tp, std::complex> || - std::is_same_v<_Tp, std::complex>> -{ -}; - /** * @brief "<" comparison with complex types support. * @@ -283,45 +271,6 @@ static constexpr DPNPFuncType get_default_floating_type() : DPNPFuncType::DPNP_FT_FLOAT; } -/** - * A template function that determines the resulting floating-point type - * based on the value of the template parameter has_fp64. - */ -template -static constexpr DPNPFuncType get_floating_res_type() -{ - constexpr auto widest_type = populate_func_types(); - constexpr auto shortes_type = (widest_type == FT1) ? FT2 : FT1; - - // Return integer result type if save_int is True - if constexpr (keep_int::value) { - if constexpr (widest_type == DPNPFuncType::DPNP_FT_INT || - widest_type == DPNPFuncType::DPNP_FT_LONG) - { - return widest_type; - } - } - - // Check for double - if constexpr (widest_type == DPNPFuncType::DPNP_FT_DOUBLE) { - return widest_type; - } - - // Check for float - else if constexpr (widest_type == DPNPFuncType::DPNP_FT_FLOAT) { - // Check if the shortest type is also float - if constexpr (shortes_type == DPNPFuncType::DPNP_FT_FLOAT) { - return widest_type; - } - } - - // Default case - return get_default_floating_type(); -} - /** * FPTR interface initialization functions */ @@ -331,8 +280,6 @@ void func_map_init_indexing_func(func_map_t &fmap); void func_map_init_linalg(func_map_t &fmap); void func_map_init_mathematical(func_map_t &fmap); void func_map_init_random(func_map_t &fmap); -void func_map_init_reduction(func_map_t &fmap); -void func_map_init_searching(func_map_t &fmap); void func_map_init_sorting(func_map_t &fmap); void func_map_init_statistics(func_map_t &fmap); diff --git a/dpnp/backend/src/dpnp_iface_fptr.cpp b/dpnp/backend/src/dpnp_iface_fptr.cpp index 2e691064e23f..7bdf2ad5405d 100644 --- a/dpnp/backend/src/dpnp_iface_fptr.cpp +++ b/dpnp/backend/src/dpnp_iface_fptr.cpp @@ -111,18 +111,6 @@ size_t operator-(DPNPFuncType lhs, DPNPFuncType rhs) return result; } -void *get_dpnp_function_ptr1(DPNPFuncType &result_type, - DPNPFuncName name, - DPNPFuncType first_type, - DPNPFuncType second_type) -{ - DPNPFuncData_t result = - get_dpnp_function_ptr(name, first_type, second_type); - - result_type = result.return_type; - return result.ptr; -} - static func_map_t func_map_init() { func_map_t fmap; @@ -133,8 +121,6 @@ static func_map_t func_map_init() func_map_init_linalg(fmap); func_map_init_mathematical(fmap); func_map_init_random(fmap); - func_map_init_reduction(fmap); - func_map_init_searching(fmap); func_map_init_sorting(fmap); func_map_init_statistics(fmap); From 5609af939c4b49b89999deb7d231eb9abbb8fba5 Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Fri, 18 Oct 2024 06:30:14 -0700 Subject: [PATCH 5/5] remove functions from dpnp/backend/src/dpnp_utils.hpp --- dpnp/backend/src/dpnp_utils.hpp | 79 --------------------------------- 1 file changed, 79 deletions(-) diff --git a/dpnp/backend/src/dpnp_utils.hpp b/dpnp/backend/src/dpnp_utils.hpp index 89b8a7331537..7866ea337943 100644 --- a/dpnp/backend/src/dpnp_utils.hpp +++ b/dpnp/backend/src/dpnp_utils.hpp @@ -183,34 +183,6 @@ _DataType get_xyz_id_by_id_inkernel(size_t global_id, return xyz_id; } -/** - * @ingroup BACKEND_UTILS - * @brief Calculate linear index from ids array - * - * Calculates linear index from ids array by given offsets. This is reverse - * operation of @ref get_xyz_by_id_inkernel for example: xyz array ids should be - * [0, 1, 0] input_array_shape_offsets[20, 5, 1] global_id == 5 - * - * @param [in] xyz array with array indexes. - * @param [in] xyz_size array size for @ref xyz parameter. - * @param [in] offsets array with input offsets. - * @return linear index id of the element in multy-D array. - */ -template -size_t get_id_by_xyz_inkernel(const _DataType *xyz, - size_t xyz_size, - const _DataType *offsets) -{ - size_t global_id = 0; - - /* calculate linear index based on reconstructed [x][y][z] */ - for (size_t it = 0; it < xyz_size; ++it) { - global_id += (xyz[it] * offsets[it]); - } - - return global_id; -} - /** * @ingroup BACKEND_UTILS * @brief Check input shape is broadcastable to output one. @@ -305,57 +277,6 @@ namespace } } // namespace -/** - * @ingroup BACKEND_UTILS - * @brief Get common shape based on input shapes. - * - * Example: - * Input1 shape A[8, 1, 6, 1] - * Input2 shape B[7, 1, 5] - * Output shape will be C[8, 7, 6, 5] - * - * @param [in] input1_shape Input1 shape. - * @param [in] input1_shape_size Input1 shape size. - * @param [in] input2_shape Input2 shape. - * @param [in] input2_shape_size Input2 shape size. - * - * @exception std::domain_error Input shapes are not broadcastable. - * @return Common shape. - */ -template -static inline std::vector<_DataType> - get_result_shape(const _DataType *input1_shape, - const size_t input1_shape_size, - const _DataType *input2_shape, - const size_t input2_shape_size) -{ - const size_t result_shape_size = (input2_shape_size > input1_shape_size) - ? input2_shape_size - : input1_shape_size; - std::vector<_DataType> result_shape; - result_shape.reserve(result_shape_size); - - for (int irit1 = input1_shape_size - 1, irit2 = input2_shape_size - 1; - irit1 >= 0 || irit2 >= 0; --irit1, --irit2) - { - _DataType input1_val = (irit1 >= 0) ? input1_shape[irit1] : 1; - _DataType input2_val = (irit2 >= 0) ? input2_shape[irit2] : 1; - - if (input1_val == input2_val || input1_val == 1) { - result_shape.insert(result_shape.begin(), input2_val); - } - else if (input2_val == 1) { - result_shape.insert(result_shape.begin(), input1_val); - } - else { - throw std::domain_error("DPNP Error: get_common_shape() failed " - "with input shapes check"); - } - } - - return result_shape; -} - /** * @ingroup BACKEND_UTILS * @brief Normalizes an axes into a non-negative integer axes.