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); 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.