diff --git a/dpnp/backend/CMakeLists.txt b/dpnp/backend/CMakeLists.txt index f9eb1a35d288..d96320bf0acd 100644 --- a/dpnp/backend/CMakeLists.txt +++ b/dpnp/backend/CMakeLists.txt @@ -25,7 +25,6 @@ set(DPNP_SRC kernels/dpnp_krnl_arraycreation.cpp - kernels/dpnp_krnl_bitwise.cpp kernels/dpnp_krnl_common.cpp kernels/dpnp_krnl_elemwise.cpp kernels/dpnp_krnl_fft.cpp diff --git a/dpnp/backend/include/dpnp_gen_1arg_1type_tbl.hpp b/dpnp/backend/include/dpnp_gen_1arg_1type_tbl.hpp index ea1c477173f6..32df8aeda723 100644 --- a/dpnp/backend/include/dpnp_gen_1arg_1type_tbl.hpp +++ b/dpnp/backend/include/dpnp_gen_1arg_1type_tbl.hpp @@ -83,24 +83,9 @@ #endif // _SECTION_DOCUMENTATION_GENERATION_ -MACRO_1ARG_1TYPE_OP(dpnp_conjugate_c, - std::conj(input_elem), - q.submit(kernel_func)) -MACRO_1ARG_1TYPE_OP(dpnp_copy_c, input_elem, q.submit(kernel_func)) MACRO_1ARG_1TYPE_OP(dpnp_erf_c, dispatch_erf_op(input_elem), oneapi::mkl::vm::erf(q, input1_size, input1_data, result)) -MACRO_1ARG_1TYPE_OP(dpnp_negative_c, -input_elem, q.submit(kernel_func)) -MACRO_1ARG_1TYPE_OP( - dpnp_recip_c, - _DataType(1) / input_elem, - q.submit(kernel_func)) // error: no member named 'recip' in namespace 'sycl' -MACRO_1ARG_1TYPE_OP(dpnp_sign_c, - dispatch_sign_op(input_elem), - q.submit(kernel_func)) // no sycl::sign for int and long -MACRO_1ARG_1TYPE_OP(dpnp_square_c, - input_elem *input_elem, - oneapi::mkl::vm::sqr(q, input1_size, input1_data, result)) #undef MACRO_1ARG_1TYPE_OP diff --git a/dpnp/backend/include/dpnp_gen_1arg_2type_tbl.hpp b/dpnp/backend/include/dpnp_gen_1arg_2type_tbl.hpp index 3abc54c7212b..a27353866d27 100644 --- a/dpnp/backend/include/dpnp_gen_1arg_2type_tbl.hpp +++ b/dpnp/backend/include/dpnp_gen_1arg_2type_tbl.hpp @@ -85,95 +85,15 @@ #endif -MACRO_1ARG_2TYPES_OP(dpnp_acos_c, - sycl::acos(input_elem), - oneapi::mkl::vm::acos(q, input1_size, input1_data, result)) -MACRO_1ARG_2TYPES_OP( - dpnp_acosh_c, - sycl::acosh(input_elem), - oneapi::mkl::vm::acosh(q, input1_size, input1_data, result)) -MACRO_1ARG_2TYPES_OP(dpnp_asin_c, - sycl::asin(input_elem), - oneapi::mkl::vm::asin(q, input1_size, input1_data, result)) -MACRO_1ARG_2TYPES_OP( - dpnp_asinh_c, - sycl::asinh(input_elem), - oneapi::mkl::vm::asinh(q, input1_size, input1_data, result)) -MACRO_1ARG_2TYPES_OP(dpnp_atan_c, - sycl::atan(input_elem), - oneapi::mkl::vm::atan(q, input1_size, input1_data, result)) -MACRO_1ARG_2TYPES_OP( - dpnp_atanh_c, - sycl::atanh(input_elem), - oneapi::mkl::vm::atanh(q, input1_size, input1_data, result)) -MACRO_1ARG_2TYPES_OP(dpnp_cbrt_c, - sycl::cbrt(input_elem), - oneapi::mkl::vm::cbrt(q, input1_size, input1_data, result)) -MACRO_1ARG_2TYPES_OP(dpnp_ceil_c, - sycl::ceil(input_elem), - oneapi::mkl::vm::ceil(q, input1_size, input1_data, result)) MACRO_1ARG_2TYPES_OP(dpnp_copyto_c, input_elem, q.submit(kernel_func)) -MACRO_1ARG_2TYPES_OP(dpnp_cos_c, - sycl::cos(input_elem), - oneapi::mkl::vm::cos(q, input1_size, input1_data, result)) -MACRO_1ARG_2TYPES_OP(dpnp_cosh_c, - sycl::cosh(input_elem), - oneapi::mkl::vm::cosh(q, input1_size, input1_data, result)) MACRO_1ARG_2TYPES_OP(dpnp_degrees_c, sycl::degrees(input_elem), q.submit(kernel_func)) -MACRO_1ARG_2TYPES_OP(dpnp_exp2_c, - sycl::exp2(input_elem), - oneapi::mkl::vm::exp2(q, input1_size, input1_data, result)) -MACRO_1ARG_2TYPES_OP(dpnp_exp_c, - sycl::exp(input_elem), - oneapi::mkl::vm::exp(q, input1_size, input1_data, result)) -MACRO_1ARG_2TYPES_OP( - dpnp_expm1_c, - sycl::expm1(input_elem), - oneapi::mkl::vm::expm1(q, input1_size, input1_data, result)) -MACRO_1ARG_2TYPES_OP(dpnp_fabs_c, - sycl::fabs(input_elem), - oneapi::mkl::vm::abs(q, input1_size, input1_data, result)) -MACRO_1ARG_2TYPES_OP( - dpnp_floor_c, - sycl::floor(input_elem), - oneapi::mkl::vm::floor(q, input1_size, input1_data, result)) -MACRO_1ARG_2TYPES_OP( - dpnp_log10_c, - sycl::log10(input_elem), - oneapi::mkl::vm::log10(q, input1_size, input1_data, result)) -MACRO_1ARG_2TYPES_OP( - dpnp_log1p_c, - sycl::log1p(input_elem), - oneapi::mkl::vm::log1p(q, input1_size, input1_data, result)) -MACRO_1ARG_2TYPES_OP(dpnp_log2_c, - sycl::log2(input_elem), - oneapi::mkl::vm::log2(q, input1_size, input1_data, result)) -MACRO_1ARG_2TYPES_OP(dpnp_log_c, - sycl::log(input_elem), - oneapi::mkl::vm::ln(q, input1_size, input1_data, result)) MACRO_1ARG_2TYPES_OP(dpnp_radians_c, sycl::radians(input_elem), q.submit(kernel_func)) -MACRO_1ARG_2TYPES_OP(dpnp_sin_c, - sycl::sin(input_elem), - oneapi::mkl::vm::sin(q, input1_size, input1_data, result)) -MACRO_1ARG_2TYPES_OP(dpnp_sinh_c, - sycl::sinh(input_elem), - oneapi::mkl::vm::sinh(q, input1_size, input1_data, result)) MACRO_1ARG_2TYPES_OP(dpnp_sqrt_c, sycl::sqrt(input_elem), oneapi::mkl::vm::sqrt(q, input1_size, input1_data, result)) -MACRO_1ARG_2TYPES_OP(dpnp_tan_c, - sycl::tan(input_elem), - oneapi::mkl::vm::tan(q, input1_size, input1_data, result)) -MACRO_1ARG_2TYPES_OP(dpnp_tanh_c, - sycl::tanh(input_elem), - oneapi::mkl::vm::tanh(q, input1_size, input1_data, result)) -MACRO_1ARG_2TYPES_OP( - dpnp_trunc_c, - sycl::trunc(input_elem), - oneapi::mkl::vm::trunc(q, input1_size, input1_data, result)) #undef MACRO_1ARG_2TYPES_OP diff --git a/dpnp/backend/include/dpnp_gen_2arg_1type_tbl.hpp b/dpnp/backend/include/dpnp_gen_2arg_1type_tbl.hpp deleted file mode 100644 index 130283e5834d..000000000000 --- a/dpnp/backend/include/dpnp_gen_2arg_1type_tbl.hpp +++ /dev/null @@ -1,105 +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 bitwise functions definitions - * - * Macro `MACRO_2ARG_1TYPE_OP` must be defined before usage - * - * Parameters: - * - public name of the function and kernel name - * - operation used to calculate the result - * - */ - -#ifndef MACRO_2ARG_1TYPE_OP -#error "MACRO_2ARG_1TYPE_OP is not defined" -#endif - -#ifdef _SECTION_DOCUMENTATION_GENERATION_ - -#define MACRO_2ARG_1TYPE_OP(__name__, __operation__) \ - /** @ingroup BACKEND_API */ \ - /** @brief Element wise operation function __name__ */ \ - /** */ \ - /** Function "__name__" executes operator "__operation__" over \ - * corresponding elements of input arrays */ \ - /** */ \ - /** @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] input2_in Input array 2. */ \ - /** @param[in] input2_size Input array 2 size. */ \ - /** @param[in] input2_ndim Number of input array 2 dimensions. \ - */ \ - /** @param[in] input2_shape Input array 2 shape. */ \ - /** @param[in] input2_strides Input array 2 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 void *input2_in, \ - const size_t input2_size, const size_t input2_ndim, \ - const shape_elem_type *input2_shape, \ - const shape_elem_type *input2_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 void *input2_in, \ - const size_t input2_size, const size_t input2_ndim, \ - const shape_elem_type *input2_shape, \ - const shape_elem_type *input2_strides, const size_t *where); - -#endif - -MACRO_2ARG_1TYPE_OP(dpnp_bitwise_and_c, input1_elem &input2_elem) -MACRO_2ARG_1TYPE_OP(dpnp_bitwise_or_c, input1_elem | input2_elem) -MACRO_2ARG_1TYPE_OP(dpnp_bitwise_xor_c, input1_elem ^ input2_elem) -MACRO_2ARG_1TYPE_OP(dpnp_left_shift_c, input1_elem << input2_elem) -MACRO_2ARG_1TYPE_OP(dpnp_right_shift_c, input1_elem >> input2_elem) - -#undef MACRO_2ARG_1TYPE_OP diff --git a/dpnp/backend/include/dpnp_gen_2arg_2type_tbl.hpp b/dpnp/backend/include/dpnp_gen_2arg_2type_tbl.hpp deleted file mode 100644 index d84accb0757d..000000000000 --- a/dpnp/backend/include/dpnp_gen_2arg_2type_tbl.hpp +++ /dev/null @@ -1,94 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2023-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_2ARG_2TYPES_LOGIC_OP` must be defined before usage - * - * Parameters: - * - public name of the function and kernel name - * - operation used to calculate the result - * - */ - -#ifndef MACRO_2ARG_2TYPES_LOGIC_OP -#error "MACRO_2ARG_2TYPES_LOGIC_OP is not defined" -#endif - -#ifdef _SECTION_DOCUMENTATION_GENERATION_ - -#define MACRO_2ARG_2TYPES_LOGIC_OP(__name__, __operation__) \ - /** @ingroup BACKEND_API */ \ - /** @brief Per element operation function __name__ */ \ - /** */ \ - /** Function "__name__" executes operator "__operation__" over \ - * corresponding elements of input arrays */ \ - /** */ \ - /** @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] input2_in Input array 2. */ \ - /** @param[in] input2_size Input array 2 size. */ \ - /** @param[in] input2_ndim Number of input array 2 dimensions. \ - */ \ - /** @param[in] input2_shape Input array 2 shape. */ \ - /** @param[in] input2_strides Input array 2 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 void *input2_in, \ - const size_t input2_size, const size_t input2_ndim, \ - const shape_elem_type *input2_shape, \ - const shape_elem_type *input2_strides, const size_t *where, \ - const DPCTLEventVectorRef dep_event_vec_ref); - -#endif - -MACRO_2ARG_2TYPES_LOGIC_OP(dpnp_equal_c, input1_elem == input2_elem) -MACRO_2ARG_2TYPES_LOGIC_OP(dpnp_greater_c, input1_elem > input2_elem) -MACRO_2ARG_2TYPES_LOGIC_OP(dpnp_greater_equal_c, input1_elem >= input2_elem) -MACRO_2ARG_2TYPES_LOGIC_OP(dpnp_less_c, input1_elem < input2_elem) -MACRO_2ARG_2TYPES_LOGIC_OP(dpnp_less_equal_c, input1_elem <= input2_elem) -MACRO_2ARG_2TYPES_LOGIC_OP(dpnp_not_equal_c, input1_elem != input2_elem) - -#undef MACRO_2ARG_2TYPES_LOGIC_OP diff --git a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp index 7423085d6593..dcec3f8192bb 100644 --- a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp +++ b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp @@ -103,40 +103,6 @@ #endif -MACRO_2ARG_3TYPES_OP(dpnp_add_c, - input1_elem + input2_elem, - x1 + x2, - MACRO_UNPACK_TYPES(bool, std::int32_t, std::int64_t), - oneapi::mkl::vm::add, - MACRO_UNPACK_TYPES(float, - double, - std::complex, - std::complex)) - -MACRO_2ARG_3TYPES_OP(dpnp_arctan2_c, - sycl::atan2(input1_elem, input2_elem), - sycl::atan2(x1, x2), - MACRO_UNPACK_TYPES(float, double), - oneapi::mkl::vm::atan2, - MACRO_UNPACK_TYPES(float, double)) - -MACRO_2ARG_3TYPES_OP(dpnp_copysign_c, - sycl::copysign(input1_elem, input2_elem), - sycl::copysign(x1, x2), - MACRO_UNPACK_TYPES(float, double), - oneapi::mkl::vm::copysign, - MACRO_UNPACK_TYPES(float, double)) - -MACRO_2ARG_3TYPES_OP(dpnp_divide_c, - input1_elem / input2_elem, - x1 / x2, - MACRO_UNPACK_TYPES(bool, std::int32_t, std::int64_t), - oneapi::mkl::vm::div, - MACRO_UNPACK_TYPES(float, - double, - std::complex, - std::complex)) - MACRO_2ARG_3TYPES_OP( dpnp_fmod_c, dispatch_fmod_op(input1_elem, input2_elem), @@ -145,13 +111,6 @@ MACRO_2ARG_3TYPES_OP( oneapi::mkl::vm::fmod, MACRO_UNPACK_TYPES(float, double)) -MACRO_2ARG_3TYPES_OP(dpnp_hypot_c, - sycl::hypot(input1_elem, input2_elem), - sycl::hypot(x1, x2), - MACRO_UNPACK_TYPES(float, double), - oneapi::mkl::vm::hypot, - MACRO_UNPACK_TYPES(float, double)) - MACRO_2ARG_3TYPES_OP(dpnp_maximum_c, sycl::max(input1_elem, input2_elem), nullptr, @@ -181,17 +140,6 @@ MACRO_2ARG_3TYPES_OP(dpnp_multiply_c, std::complex, std::complex)) -MACRO_2ARG_3TYPES_OP(dpnp_power_c, - static_cast<_DataType_output>(std::pow(input1_elem, - input2_elem)), - sycl::pow(x1, x2), - MACRO_UNPACK_TYPES(float, double), - oneapi::mkl::vm::pow, - MACRO_UNPACK_TYPES(float, - double, - std::complex, - std::complex)) - MACRO_2ARG_3TYPES_OP(dpnp_subtract_c, input1_elem - input2_elem, x1 - x2, diff --git a/dpnp/backend/include/dpnp_iface.hpp b/dpnp/backend/include/dpnp_iface.hpp index ccbf6fa85361..324e7a612b1a 100644 --- a/dpnp/backend/include/dpnp_iface.hpp +++ b/dpnp/backend/include/dpnp_iface.hpp @@ -202,28 +202,6 @@ template INP_DLLEXPORT void dpnp_arange_c(size_t start, size_t step, void *result1, size_t size); -/** - * @ingroup BACKEND_API - * @brief Copy of the array, cast to a specified type. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array Input array. - * @param [out] result Output array. - * @param [in] size Number of input elements in `array`. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_astype_c(DPCTLSyclQueueRef q_ref, - const void *array, - void *result, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void - dpnp_astype_c(const void *array, void *result, const size_t size); - /** * @ingroup BACKEND_API * @brief Implementation of full function @@ -266,67 +244,6 @@ INP_DLLEXPORT DPCTLSyclEventRef template INP_DLLEXPORT void dpnp_full_like_c(void *array_in, void *result, size_t size); -/** - * @ingroup BACKEND_API - * @brief Matrix multiplication. - * - * Matrix multiplication procedure. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result_out Output array. - * @param [in] result_size Size of output array. - * @param [in] result_ndim Number of output array dimensions. - * @param [in] result_shape Shape of output array. - * @param [in] result_strides Strides of output array. - * @param [in] input1_in First input array. - * @param [in] input1_size Size of first input array. - * @param [in] input1_ndim Number of first input array dimensions. - * @param [in] input1_shape Shape of first input array. - * @param [in] input1_strides Strides of first input array. - * @param [in] input2_in Second input array. - * @param [in] input2_size Size of second input array. - * @param [in] input2_ndim Number of second input array dimensions. - * @param [in] input2_shape Shape of second input array. - * @param [in] input2_strides Strides of second input array. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_matmul_c(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 void *input2_in, - const size_t input2_size, - const size_t input2_ndim, - const shape_elem_type *input2_shape, - const shape_elem_type *input2_strides, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_matmul_c(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 void *input2_in, - const size_t input2_size, - const size_t input2_ndim, - const shape_elem_type *input2_shape, - const shape_elem_type *input2_strides); - /** * @ingroup BACKEND_API * @brief Compute the variance along the specified axis, while ignoring NaNs. @@ -388,28 +305,6 @@ INP_DLLEXPORT void dpnp_nonzero_c(const void *array1, const size_t ndim, const size_t j); -/** - * @ingroup BACKEND_API - * @brief absolute function. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] input1_in Input array. - * @param [out] result1 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_elemwise_absolute_c(DPCTLSyclQueueRef q_ref, - const void *input1_in, - void *result1, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void - dpnp_elemwise_absolute_c(const void *input1_in, void *result1, size_t size); - /** * @ingroup BACKEND_API * @brief Custom implementation of dot function @@ -473,98 +368,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 Custom implementation of cross function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result_out Output array. - * @param [in] input1_in First input array. - * @param [in] input1_size Size of first input array. - * @param [in] input1_shape Shape of first input array. - * @param [in] input1_shape_ndim Number of first array dimensions. - * @param [in] input2_in Second input array. - * @param [in] input2_size Shape of second input array. - * @param [in] input2_shape Shape of first input array. - * @param [in] input2_shape_ndim Number of second array dimensions. - * @param [in] where Mask array. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_cross_c(DPCTLSyclQueueRef q_ref, - void *result_out, - const void *input1_in, - const size_t input1_size, - const shape_elem_type *input1_shape, - const size_t input1_shape_ndim, - const void *input2_in, - const size_t input2_size, - const shape_elem_type *input2_shape, - const size_t input2_shape_ndim, - const size_t *where, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_cross_c(void *result_out, - const void *input1_in, - const size_t input1_size, - const shape_elem_type *input1_shape, - const size_t input1_shape_ndim, - const void *input2_in, - const size_t input2_size, - const shape_elem_type *input2_shape, - const size_t input2_shape_ndim, - const size_t *where); - -/** - * @ingroup BACKEND_API - * @brief Custom implementation of cumprod function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array1_in Input array. - * @param [out] result1 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_cumprod_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result1, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_cumprod_c(void *array1_in, void *result1, size_t size); - -/** - * @ingroup BACKEND_API - * @brief Custom implementation of cumsum function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array1_in Input array. - * @param [out] result1 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_cumsum_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result1, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_cumsum_c(void *array1_in, void *result1, size_t size); - /** * @ingroup BACKEND_API * @brief The differences between consecutive elements of an array. @@ -910,54 +713,6 @@ INP_DLLEXPORT void dpnp_put_along_axis_c(void *arr_in, size_t size_indices, size_t values_size); -/** - * @ingroup BACKEND_API - * @brief Compute the eigenvalues and right eigenvectors of a square array. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array_in Input array[size][size] - * @param [out] result1 The eigenvalues, each repeated according to - * its multiplicity - * @param [out] result2 The normalized (unit "length") eigenvectors - * @param [in] size One dimension of square [size][size] array - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_eig_c(DPCTLSyclQueueRef q_ref, - const void *array_in, - void *result1, - void *result2, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void - dpnp_eig_c(const void *array_in, void *result1, void *result2, size_t size); - -/** - * @ingroup BACKEND_API - * @brief Compute the eigenvalues of a square array. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array_in Input array[size][size] - * @param [out] result1 The eigenvalues, each repeated according to - * its multiplicity - * @param [in] size One dimension of square [size][size] array - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_eigvals_c(DPCTLSyclQueueRef q_ref, - const void *array_in, - void *result1, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void - dpnp_eigvals_c(const void *array_in, void *result1, size_t size); - /** * @ingroup BACKEND_API * @brief Return a 2-D array with ones on the diagonal and zeros elsewhere. @@ -1056,32 +811,6 @@ INP_DLLEXPORT DPCTLSyclEventRef template INP_DLLEXPORT void dpnp_sort_c(void *array, void *result, size_t size); -/** - * @ingroup BACKEND_API - * @brief math library implementation of cholesky function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array Input array with data. - * @param [out] result Output array. - * @param [in] size Number of elements in input arrays. - * @param [in] data_size Last element of shape arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_cholesky_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result1, - const size_t size, - const size_t data_size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_cholesky_c(void *array1_in, - void *result1, - const size_t size, - const size_t data_size); - /** * @ingroup BACKEND_API * @brief correlate function @@ -1154,32 +883,6 @@ template INP_DLLEXPORT void dpnp_cov_c(void *array1_in, void *result1, size_t nrows, size_t ncols); -/** - * @ingroup BACKEND_API - * @brief math library implementation of det 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] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_det_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result1, - shape_elem_type *shape, - size_t ndim, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_det_c(void *array1_in, - void *result1, - shape_elem_type *shape, - size_t ndim); - /** * @ingroup BACKEND_API * @brief Construct an array from an index array and a list of arrays to choose @@ -1344,58 +1047,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 inv function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array1_in Input array with data. - * @param [out] result1 Output array. - * @param [in] shape Shape of input array. - * @param [in] ndim Number of elements in shape. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_inv_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result1, - shape_elem_type *shape, - size_t ndim, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_inv_c(void *array1_in, - void *result1, - shape_elem_type *shape, - size_t ndim); - -/** - * @ingroup BACKEND_API - * @brief math library implementation of matrix_rank function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array1_in Input array with data. - * @param [out] result1 Output array. - * @param [in] shape Shape of input array. - * @param [in] ndim Number of elements in shape. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_matrix_rank_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result1, - shape_elem_type *shape, - size_t ndim, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_matrix_rank_c(void *array1_in, - void *result1, - shape_elem_type *shape, - size_t ndim); - /** * @ingroup BACKEND_API * @brief math library implementation of max function @@ -1572,33 +1223,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 around function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] input_in Input array with data. - * @param [out] result_out Output array with indices. - * @param [in] input_size Number of elements in input arrays. - * @param [in] decimals Number of decimal places to round. Support - * only with default value 0. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_around_c(DPCTLSyclQueueRef q_ref, - const void *input_in, - void *result_out, - const size_t input_size, - const int decimals, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_around_c(const void *input_in, - void *result_out, - const size_t input_size, - const int decimals); - /** * @ingroup BACKEND_API * @brief math library implementation of std function @@ -1820,55 +1444,6 @@ INP_DLLEXPORT void dpnp_var_c(void *array, size_t naxis, size_t ddof); -/** - * @ingroup BACKEND_API - * @brief Implementation of invert function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array1_in Input array. - * @param [out] result1 Output array. - * @param [in] size Number of elements in the input array. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_invert_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_invert_c(void *array1_in, void *result, size_t size); - -#define MACRO_2ARG_1TYPE_OP(__name__, __operation__) \ - 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 void *input2_in, \ - const size_t input2_size, const size_t input2_ndim, \ - const shape_elem_type *input2_shape, \ - const shape_elem_type *input2_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 void *input2_in, \ - const size_t input2_size, const size_t input2_ndim, \ - const shape_elem_type *input2_shape, \ - const shape_elem_type *input2_strides, const size_t *where); - -#include - #define MACRO_1ARG_1TYPE_OP(__name__, __operation1__, __operation2__) \ template \ INP_DLLEXPORT DPCTLSyclEventRef __name__( \ @@ -1913,23 +1488,6 @@ INP_DLLEXPORT void dpnp_invert_c(void *array1_in, void *result, size_t size); #include -#define MACRO_2ARG_2TYPES_LOGIC_OP(__name__, __operation__) \ - 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 void *input2_in, \ - const size_t input2_size, const size_t input2_ndim, \ - const shape_elem_type *input2_shape, \ - const shape_elem_type *input2_strides, const size_t *where, \ - const DPCTLEventVectorRef dep_event_vec_ref); - -#include - #define MACRO_2ARG_3TYPES_OP(__name__, __operation__, __vec_operation__, \ __vec_types__, __mkl_operation__, __mkl_types__) \ template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_floor_divide_c(DPCTLSyclQueueRef q_ref, - void *result_out, - const void *input1_in, - const size_t input1_size, - const shape_elem_type *input1_shape, - const size_t input1_shape_ndim, - const void *input2_in, - const size_t input2_size, - const shape_elem_type *input2_shape, - const size_t input2_shape_ndim, - const size_t *where, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_floor_divide_c(void *result_out, - const void *input1_in, - const size_t input1_size, - const shape_elem_type *input1_shape, - const size_t input1_shape_ndim, - const void *input2_in, - const size_t input2_size, - const shape_elem_type *input2_shape, - const size_t input2_shape_ndim, - const size_t *where); - /** * @ingroup BACKEND_API * @brief modf function. @@ -2099,54 +1609,6 @@ INP_DLLEXPORT DPCTLSyclEventRef template INP_DLLEXPORT void dpnp_ones_like_c(void *result, size_t size); -/** - * @ingroup BACKEND_API - * @brief remainder function. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result_out Output array. - * @param [in] input1_in First input array. - * @param [in] input1_size Size of first input array. - * @param [in] input1_shape Shape of first input array. - * @param [in] input1_shape_ndim Number of first array dimensions. - * @param [in] input2_in Second input array. - * @param [in] input2_size Shape of second input array. - * @param [in] input2_shape Shape of first input array. - * @param [in] input2_shape_ndim Number of second array dimensions. - * @param [in] where Mask array. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_remainder_c(DPCTLSyclQueueRef q_ref, - void *result_out, - const void *input1_in, - const size_t input1_size, - const shape_elem_type *input1_shape, - const size_t input1_shape_ndim, - const void *input2_in, - const size_t input2_size, - const shape_elem_type *input2_shape, - const size_t input2_shape_ndim, - const size_t *where, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_remainder_c(void *result_out, - const void *input1_in, - const size_t input1_size, - const shape_elem_type *input1_shape, - const size_t input1_shape_ndim, - const void *input2_in, - const size_t input2_size, - const shape_elem_type *input2_shape, - const size_t input2_shape_ndim, - const size_t *where); - /** * @ingroup BACKEND_API * @brief repeat elements of an array. @@ -2173,81 +1635,6 @@ INP_DLLEXPORT void dpnp_repeat_c(const void *array_in, const size_t repeats, const size_t size); -/** - * @ingroup BACKEND_API - * @brief transpose function. Permute axes of the input to the output with - * elements permutation. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array1_in Input array. - * @param [in] input_shape Input shape. - * @param [in] result_shape Output shape. - * @param [in] permute_axes Order of axis by it's id as it should be - * presented in output. - * @param [in] ndim Number of elements in shapes and axes. - * @param [out] result1 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_elemwise_transpose_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - const shape_elem_type *input_shape, - const shape_elem_type *result_shape, - const shape_elem_type *permute_axes, - size_t ndim, - void *result1, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void - dpnp_elemwise_transpose_c(void *array1_in, - const shape_elem_type *input_shape, - const shape_elem_type *result_shape, - const shape_elem_type *permute_axes, - size_t ndim, - void *result1, - size_t size); - -/** - * @ingroup BACKEND_API - * @brief Custom implementation of trapz function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array1_in First input array. - * @param [in] array2_in Second input array. - * @param [out] result1 Output array. - * @param [in] dx The spacing between sample points. - * @param [in] array1_size Number of elements in first input array. - * @param [in] array2_size Number of elements in second input arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - * - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_trapz_c(DPCTLSyclQueueRef q_ref, - const void *array1_in, - const void *array2_in, - void *result1, - double dx, - size_t array1_size, - size_t array2_size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_trapz_c(const void *array1_in, - const void *array2_in, - void *result1, - double dx, - size_t array1_size, - size_t array2_size); - /** * @ingroup BACKEND_API * @brief Implementation of vander function diff --git a/dpnp/backend/include/dpnp_iface_fptr.hpp b/dpnp/backend/include/dpnp_iface_fptr.hpp index 1172bcbe4f5f..a39174931fec 100644 --- a/dpnp/backend/include/dpnp_iface_fptr.hpp +++ b/dpnp/backend/include/dpnp_iface_fptr.hpp @@ -58,77 +58,42 @@ */ enum class DPNPFuncName : size_t { - DPNP_FN_NONE, /**< Very first element of the enumeration */ - DPNP_FN_ABSOLUTE, /**< Used in numpy.absolute() impl */ - DPNP_FN_ADD, /**< Used in numpy.add() impl */ - DPNP_FN_ALL, /**< Used in numpy.all() impl */ - DPNP_FN_ALLCLOSE, /**< Used in numpy.allclose() impl */ - DPNP_FN_ALLCLOSE_EXT, /**< Used in numpy.allclose() impl, requires extra - parameters */ - DPNP_FN_ANY, /**< Used in numpy.any() impl */ - DPNP_FN_ARANGE, /**< Used in numpy.arange() impl */ - DPNP_FN_ARCCOS, /**< Used in numpy.arccos() impl */ - DPNP_FN_ARCCOSH, /**< Used in numpy.arccosh() impl */ - DPNP_FN_ARCSIN, /**< Used in numpy.arcsin() impl */ - DPNP_FN_ARCSINH, /**< Used in numpy.arcsinh() impl */ - DPNP_FN_ARCTAN, /**< Used in numpy.arctan() impl */ - DPNP_FN_ARCTAN2, /**< Used in numpy.arctan2() impl */ - DPNP_FN_ARCTANH, /**< Used in numpy.arctanh() impl */ - 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_AROUND, /**< Used in numpy.around() impl */ - DPNP_FN_ASTYPE, /**< Used in numpy.astype() impl */ - DPNP_FN_BITWISE_AND, /**< Used in numpy.bitwise_and() impl */ - DPNP_FN_BITWISE_OR, /**< Used in numpy.bitwise_or() impl */ - DPNP_FN_BITWISE_XOR, /**< Used in numpy.bitwise_xor() impl */ - DPNP_FN_CBRT, /**< Used in numpy.cbrt() impl */ - DPNP_FN_CEIL, /**< Used in numpy.ceil() impl */ - DPNP_FN_CHOLESKY, /**< Used in numpy.linalg.cholesky() impl */ - DPNP_FN_CONJUGATE, /**< Used in numpy.conjugate() impl */ - DPNP_FN_CHOOSE, /**< Used in numpy.choose() impl */ - DPNP_FN_CHOOSE_EXT, /**< Used in numpy.choose() impl, requires extra - parameters */ - DPNP_FN_COPY, /**< Used in numpy.copy() impl */ - DPNP_FN_COPY_EXT, /**< Used in numpy.copy() impl, requires extra parameters - */ - DPNP_FN_COPYSIGN, /**< Used in numpy.copysign() impl */ - DPNP_FN_COPYTO, /**< Used in numpy.copyto() impl */ + DPNP_FN_NONE, /**< Very first element of the enumeration */ + DPNP_FN_ALL, /**< Used in numpy.all() impl */ + DPNP_FN_ALLCLOSE, /**< Used in numpy.allclose() impl */ + DPNP_FN_ALLCLOSE_EXT, /**< Used in numpy.allclose() impl, requires extra + parameters */ + DPNP_FN_ANY, /**< Used in numpy.any() impl */ + DPNP_FN_ARANGE, /**< Used in numpy.arange() impl */ + 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 */ + DPNP_FN_COPYTO, /**< Used in numpy.copyto() impl */ DPNP_FN_COPYTO_EXT, /**< Used in numpy.copyto() impl, requires extra parameters */ DPNP_FN_CORRELATE, /**< Used in numpy.correlate() impl */ DPNP_FN_CORRELATE_EXT, /**< Used in numpy.correlate() impl, requires extra parameters */ - DPNP_FN_COS, /**< Used in numpy.cos() impl */ - DPNP_FN_COSH, /**< Used in numpy.cosh() impl */ DPNP_FN_COUNT_NONZERO, /**< Used in numpy.count_nonzero() impl */ DPNP_FN_COV, /**< Used in numpy.cov() impl */ - DPNP_FN_CROSS, /**< Used in numpy.cross() impl */ - DPNP_FN_CUMPROD, /**< Used in numpy.cumprod() impl */ - DPNP_FN_CUMSUM, /**< Used in numpy.cumsum() impl */ DPNP_FN_DEGREES, /**< Used in numpy.degrees() impl */ DPNP_FN_DEGREES_EXT, /**< Used in numpy.degrees() impl, requires extra parameters */ - DPNP_FN_DET, /**< Used in numpy.linalg.det() impl */ DPNP_FN_DIAG, /**< Used in numpy.diag() impl */ DPNP_FN_DIAG_INDICES, /**< Used in numpy.diag_indices() impl */ DPNP_FN_DIAGONAL, /**< Used in numpy.diagonal() impl */ - DPNP_FN_DIVIDE, /**< Used in numpy.divide() impl */ DPNP_FN_DOT, /**< Used in numpy.dot() impl */ DPNP_FN_DOT_EXT, /**< Used in numpy.dot() impl, requires extra parameters */ DPNP_FN_EDIFF1D, /**< Used in numpy.ediff1d() impl */ DPNP_FN_EDIFF1D_EXT, /**< Used in numpy.ediff1d() impl, requires extra parameters */ - DPNP_FN_EIG, /**< Used in numpy.linalg.eig() impl */ - DPNP_FN_EIGVALS, /**< Used in numpy.linalg.eigvals() impl */ DPNP_FN_ERF, /**< Used in scipy.special.erf impl */ DPNP_FN_ERF_EXT, /**< Used in scipy.special.erf impl, requires extra parameters */ DPNP_FN_EYE, /**< Used in numpy.eye() impl */ - DPNP_FN_EXP, /**< Used in numpy.exp() impl */ - DPNP_FN_EXP2, /**< Used in numpy.exp2() impl */ - DPNP_FN_EXPM1, /**< Used in numpy.expm1() impl */ - DPNP_FN_FABS, /**< Used in numpy.fabs() impl */ DPNP_FN_FFT_FFT, /**< Used in numpy.fft.fft() impl */ DPNP_FN_FFT_FFT_EXT, /**< Used in numpy.fft.fft() impl, requires extra parameters */ @@ -136,30 +101,15 @@ enum class DPNPFuncName : size_t DPNP_FN_FFT_RFFT_EXT, /**< Used in numpy.fft.rfft() impl, requires extra parameters */ DPNP_FN_FILL_DIAGONAL, /**< Used in numpy.fill_diagonal() impl */ - DPNP_FN_FLATTEN, /**< Used in numpy.flatten() impl */ - DPNP_FN_FLOOR, /**< Used in numpy.floor() impl */ - DPNP_FN_FLOOR_DIVIDE, /**< Used in numpy.floor_divide() impl */ - DPNP_FN_FMOD, /**< Used in numpy.fmod() impl */ DPNP_FN_FULL, /**< Used in numpy.full() impl */ DPNP_FN_FULL_LIKE, /**< Used in numpy.full_like() impl */ - DPNP_FN_HYPOT, /**< Used in numpy.hypot() impl */ DPNP_FN_IDENTITY, /**< Used in numpy.identity() impl */ DPNP_FN_INITVAL, /**< Used in numpy ones, ones_like, zeros, zeros_like impls */ DPNP_FN_INITVAL_EXT, /**< Used in numpy ones, ones_like, zeros, zeros_like impls */ - DPNP_FN_INV, /**< Used in numpy.linalg.inv() impl */ DPNP_FN_INVERT, /**< Used in numpy.invert() impl */ - DPNP_FN_KRON, /**< Used in numpy.kron() impl */ - DPNP_FN_LEFT_SHIFT, /**< Used in numpy.left_shift() impl */ - DPNP_FN_LOG, /**< Used in numpy.log() impl */ - DPNP_FN_LOG10, /**< Used in numpy.log10() impl */ - DPNP_FN_LOG2, /**< Used in numpy.log2() impl */ - DPNP_FN_LOG1P, /**< Used in numpy.log1p() impl */ - DPNP_FN_MATMUL, /**< Used in numpy.matmul() impl */ - DPNP_FN_MATRIX_RANK, /**< Used in numpy.linalg.matrix_rank() impl */ DPNP_FN_MAX, /**< Used in numpy.max() impl */ - DPNP_FN_MAXIMUM, /**< Used in numpy.fmax() impl */ DPNP_FN_MAXIMUM_EXT, /**< Used in numpy.fmax() impl , requires extra parameters */ DPNP_FN_MEAN, /**< Used in numpy.mean() impl */ @@ -167,7 +117,6 @@ enum class DPNPFuncName : size_t DPNP_FN_MEDIAN_EXT, /**< Used in numpy.median() impl, requires extra parameters */ DPNP_FN_MIN, /**< Used in numpy.min() impl */ - DPNP_FN_MINIMUM, /**< Used in numpy.fmin() impl */ DPNP_FN_MINIMUM_EXT, /**< Used in numpy.fmax() impl, requires extra parameters */ DPNP_FN_MODF, /**< Used in numpy.modf() impl */ @@ -175,7 +124,6 @@ enum class DPNPFuncName : size_t */ DPNP_FN_MULTIPLY, /**< Used in numpy.multiply() impl */ DPNP_FN_NANVAR, /**< Used in numpy.nanvar() impl */ - DPNP_FN_NEGATIVE, /**< Used in numpy.negative() impl */ DPNP_FN_NONZERO, /**< Used in numpy.nonzero() impl */ DPNP_FN_ONES, /**< Used in numpy.ones() impl */ DPNP_FN_ONES_LIKE, /**< Used in numpy.ones_like() impl */ @@ -183,19 +131,14 @@ enum class DPNPFuncName : size_t DPNP_FN_PARTITION_EXT, /**< Used in numpy.partition() impl, requires extra parameters */ DPNP_FN_PLACE, /**< Used in numpy.place() impl */ - DPNP_FN_POWER, /**< Used in numpy.power() impl */ DPNP_FN_PROD, /**< Used in numpy.prod() impl */ DPNP_FN_PTP, /**< Used in numpy.ptp() impl */ DPNP_FN_PUT, /**< Used in numpy.put() impl */ DPNP_FN_PUT_ALONG_AXIS, /**< Used in numpy.put_along_axis() impl */ - DPNP_FN_QR, /**< Used in numpy.linalg.qr() impl */ DPNP_FN_RADIANS, /**< Used in numpy.radians() impl */ DPNP_FN_RADIANS_EXT, /**< Used in numpy.radians() impl, requires extra parameters */ - DPNP_FN_REMAINDER, /**< Used in numpy.remainder() impl */ - DPNP_FN_RECIP, /**< Used in numpy.recip() impl */ DPNP_FN_REPEAT, /**< Used in numpy.repeat() impl */ - DPNP_FN_RIGHT_SHIFT, /**< Used in numpy.right_shift() 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 */ @@ -314,32 +257,22 @@ enum class DPNPFuncName : size_t DPNP_FN_RNG_ZIPF_EXT, /**< Used in numpy.random.zipf() impl, requires extra parameters */ DPNP_FN_SEARCHSORTED, /**< Used in numpy.searchsorted() impl */ - DPNP_FN_SIGN, /**< Used in numpy.sign() impl */ - DPNP_FN_SIN, /**< Used in numpy.sin() impl */ - DPNP_FN_SINH, /**< Used in numpy.sinh() impl */ DPNP_FN_SORT, /**< Used in numpy.sort() impl */ DPNP_FN_SQRT, /**< Used in numpy.sqrt() impl */ DPNP_FN_SQRT_EXT, /**< Used in numpy.sqrt() impl, requires extra parameters */ - DPNP_FN_SQUARE, /**< Used in numpy.square() impl */ DPNP_FN_STD, /**< Used in numpy.std() impl */ - DPNP_FN_SUBTRACT, /**< Used in numpy.subtract() impl */ DPNP_FN_SUBTRACT_EXT, /**< Used in numpy.subtract() impl, requires extra parameters */ DPNP_FN_SUM, /**< Used in numpy.sum() impl */ - DPNP_FN_SVD, /**< Used in numpy.linalg.svd() impl */ DPNP_FN_TAKE, /**< Used in numpy.take() impl */ - DPNP_FN_TAN, /**< Used in numpy.tan() impl */ - DPNP_FN_TANH, /**< Used in numpy.tanh() impl */ DPNP_FN_TRANSPOSE, /**< Used in numpy.transpose() impl */ DPNP_FN_TRACE, /**< Used in numpy.trace() impl */ - DPNP_FN_TRAPZ, /**< Used in numpy.trapz() impl */ DPNP_FN_TRAPZ_EXT, /**< Used in numpy.trapz() impl, requires extra parameters */ DPNP_FN_TRI, /**< Used in numpy.tri() impl */ DPNP_FN_TRIL, /**< Used in numpy.tril() impl */ DPNP_FN_TRIU, /**< Used in numpy.triu() impl */ - DPNP_FN_TRUNC, /**< Used in numpy.trunc() impl */ DPNP_FN_VANDER, /**< Used in numpy.vander() impl */ DPNP_FN_VAR, /**< Used in numpy.var() impl */ DPNP_FN_ZEROS, /**< Used in numpy.zeros() impl */ diff --git a/dpnp/backend/kernels/dpnp_krnl_bitwise.cpp b/dpnp/backend/kernels/dpnp_krnl_bitwise.cpp deleted file mode 100644 index 9db8425f6de4..000000000000 --- a/dpnp/backend/kernels/dpnp_krnl_bitwise.cpp +++ /dev/null @@ -1,438 +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 "dpnp_iface.hpp" -#include "dpnp_iterator.hpp" -#include "dpnp_utils.hpp" -#include "dpnpc_memory_adapter.hpp" -#include "queue_sycl.hpp" - -// dpctl tensor headers -#include "kernels/alignment.hpp" - -using dpctl::tensor::kernels::alignment_utils::is_aligned; -using dpctl::tensor::kernels::alignment_utils::required_alignment; - -template -class dpnp_invert_c_kernel; - -template -DPCTLSyclEventRef dpnp_invert_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)); - sycl::event event; - - _DataType *input_data = static_cast<_DataType *>(array1_in); - _DataType *result = static_cast<_DataType *>(result1); - - constexpr size_t lws = 64; - constexpr unsigned int vec_sz = 8; - - auto gws_range = - sycl::range<1>(((size + lws * vec_sz - 1) / (lws * vec_sz)) * lws); - auto lws_range = sycl::range<1>(lws); - - auto kernel_parallel_for_func = [=](sycl::nd_item<1> nd_it) { - auto sg = nd_it.get_sub_group(); - const auto max_sg_size = sg.get_max_local_range()[0]; - const size_t start = - vec_sz * (nd_it.get_group(0) * nd_it.get_local_range(0) + - sg.get_group_id()[0] * max_sg_size); - - if (is_aligned(input_data) && - is_aligned(result) && - (start + static_cast(vec_sz) * max_sg_size < size)) - { - auto input_multi_ptr = sycl::address_space_cast< - sycl::access::address_space::global_space, - sycl::access::decorated::yes>(&input_data[start]); - auto result_multi_ptr = sycl::address_space_cast< - sycl::access::address_space::global_space, - sycl::access::decorated::yes>(&result[start]); - - sycl::vec<_DataType, vec_sz> x = sg.load(input_multi_ptr); - sycl::vec<_DataType, vec_sz> res_vec; - - if constexpr (std::is_same_v<_DataType, bool>) { -#pragma unroll - for (size_t k = 0; k < vec_sz; ++k) { - res_vec[k] = !(x[k]); - } - } - else { - res_vec = ~x; - } - - sg.store(result_multi_ptr, res_vec); - } - else { - for (size_t k = start + sg.get_local_id()[0]; k < size; - k += max_sg_size) { - if constexpr (std::is_same_v<_DataType, bool>) { - result[k] = !(input_data[k]); - } - else { - result[k] = ~(input_data[k]); - } - } - } - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.parallel_for>( - sycl::nd_range<1>(gws_range, lws_range), kernel_parallel_for_func); - }; - event = q.submit(kernel_func); - - event_ref = reinterpret_cast(&event); - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_invert_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_invert_c<_DataType>( - q_ref, array1_in, result1, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_invert_default_c)(void *, - void *, - size_t) = dpnp_invert_c<_DataType>; - -static void func_map_init_bitwise_1arg_1type(func_map_t &fmap) -{ - fmap[DPNPFuncName::DPNP_FN_INVERT][eft_BLN][eft_BLN] = { - eft_BLN, (void *)dpnp_invert_default_c}; - fmap[DPNPFuncName::DPNP_FN_INVERT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_invert_default_c}; - fmap[DPNPFuncName::DPNP_FN_INVERT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_invert_default_c}; - - return; -} - -#define MACRO_2ARG_1TYPE_OP(__name__, __operation__) \ - template \ - class __name__##_kernel; \ - \ - template \ - class __name__##_strides_kernel; \ - \ - template \ - class __name__##_broadcast_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 void *input2_in, \ - const size_t input2_size, const size_t input2_ndim, \ - const shape_elem_type *input2_shape, \ - const shape_elem_type *input2_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 || !input2_size) { \ - return event_ref; \ - } \ - \ - sycl::queue q = *(reinterpret_cast(q_ref)); \ - \ - _DataType *input1_data = \ - static_cast<_DataType *>(const_cast(input1_in)); \ - _DataType *input2_data = \ - static_cast<_DataType *>(const_cast(input2_in)); \ - _DataType *result = static_cast<_DataType *>(result_out); \ - \ - bool use_broadcasting = !array_equal(input1_shape, input1_ndim, \ - input2_shape, input2_ndim); \ - \ - 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; \ - \ - shape_elem_type *input2_shape_offsets = \ - new shape_elem_type[input2_ndim]; \ - \ - get_shape_offsets_inkernel(input2_shape, input2_ndim, \ - input2_shape_offsets); \ - use_strides = \ - use_strides || !array_equal(input2_strides, input2_ndim, \ - input2_shape_offsets, input2_ndim); \ - delete[] input2_shape_offsets; \ - \ - sycl::event event; \ - sycl::range<1> gws(result_size); \ - \ - if (use_broadcasting) { \ - DPNPC_id<_DataType> *input1_it; \ - const size_t input1_it_size_in_bytes = \ - sizeof(DPNPC_id<_DataType>); \ - input1_it = reinterpret_cast *>( \ - dpnp_memory_alloc_c(q_ref, input1_it_size_in_bytes)); \ - new (input1_it) \ - DPNPC_id<_DataType>(q_ref, input1_data, input1_shape, \ - input1_strides, input1_ndim); \ - \ - input1_it->broadcast_to_shape(result_shape, result_ndim); \ - \ - DPNPC_id<_DataType> *input2_it; \ - const size_t input2_it_size_in_bytes = \ - sizeof(DPNPC_id<_DataType>); \ - input2_it = reinterpret_cast *>( \ - dpnp_memory_alloc_c(q_ref, input2_it_size_in_bytes)); \ - new (input2_it) \ - DPNPC_id<_DataType>(q_ref, input2_data, input2_shape, \ - input2_strides, input2_ndim); \ - \ - input2_it->broadcast_to_shape(result_shape, result_ndim); \ - \ - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { \ - const size_t i = global_id[0]; /* for (size_t i = 0; i < \ - result_size; ++i) */ \ - { \ - const _DataType input1_elem = (*input1_it)[i]; \ - const _DataType input2_elem = (*input2_it)[i]; \ - result[i] = __operation__; \ - } \ - }; \ - auto kernel_func = [&](sycl::handler &cgh) { \ - cgh.parallel_for< \ - class __name__##_broadcast_kernel<_DataType>>( \ - gws, kernel_parallel_for_func); \ - }; \ - \ - q.submit(kernel_func).wait(); \ - \ - input1_it->~DPNPC_id(); \ - input2_it->~DPNPC_id(); \ - \ - return event_ref; \ - } \ - else if (use_strides) { \ - if ((result_ndim != input1_ndim) || (result_ndim != input2_ndim)) \ - { \ - throw std::runtime_error( \ - "Result ndim=" + std::to_string(result_ndim) + \ - " mismatches with either input1 ndim=" + \ - std::to_string(input1_ndim) + \ - " or input2 ndim=" + std::to_string(input2_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 = 3 * 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); \ - std::copy(input2_strides, input2_strides + result_ndim, \ - strides_host_packed.begin() + 2 * 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) { \ - const 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]; \ - const shape_elem_type *input2_strides_data = \ - &dev_strides_data[2 * result_ndim]; \ - \ - size_t input1_id = 0; \ - size_t input2_id = 0; \ - \ - for (size_t i = 0; i < result_ndim; ++i) { \ - const size_t output_xyz_id = \ - get_xyz_id_by_id_inkernel(output_id, \ - result_strides_data, \ - result_ndim, i); \ - input1_id += output_xyz_id * input1_strides_data[i]; \ - input2_id += output_xyz_id * input2_strides_data[i]; \ - } \ - \ - const _DataType input1_elem = \ - (input1_size == 1) ? input1_data[0] \ - : input1_data[input1_id]; \ - const _DataType input2_elem = \ - (input2_size == 1) ? input2_data[0] \ - : input2_data[input2_id]; \ - result[output_id] = __operation__; \ - } \ - }; \ - auto kernel_func = [&](sycl::handler &cgh) { \ - cgh.depends_on(copy_strides_ev); \ - 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 i = global_id[0]; /* for (size_t i = 0; i < \ - result_size; ++i) */ \ - const _DataType input1_elem = \ - (input1_size == 1) ? input1_data[0] : input1_data[i]; \ - const _DataType input2_elem = \ - (input2_size == 1) ? input2_data[0] : input2_data[i]; \ - result[i] = __operation__; \ - }; \ - auto kernel_func = [&](sycl::handler &cgh) { \ - cgh.parallel_for>( \ - gws, kernel_parallel_for_func); \ - }; \ - 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 void *input2_in, \ - const size_t input2_size, const size_t input2_ndim, \ - const shape_elem_type *input2_shape, \ - const shape_elem_type *input2_strides, const size_t *where) \ - { \ - DPCTLSyclQueueRef q_ref = \ - reinterpret_cast(&DPNP_QUEUE); \ - DPCTLEventVectorRef dep_event_vec_ref = nullptr; \ - DPCTLSyclEventRef event_ref = __name__<_DataType>( \ - q_ref, result_out, result_size, result_ndim, result_shape, \ - result_strides, input1_in, input1_size, input1_ndim, input1_shape, \ - input1_strides, input2_in, input2_size, input2_ndim, input2_shape, \ - input2_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 void *, \ - const size_t, const size_t, const shape_elem_type *, \ - const shape_elem_type *, const size_t *) = __name__<_DataType>; - -#include - -static void func_map_init_bitwise_2arg_1type(func_map_t &fmap) -{ - fmap[DPNPFuncName::DPNP_FN_BITWISE_AND][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_bitwise_and_c_default}; - fmap[DPNPFuncName::DPNP_FN_BITWISE_AND][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_bitwise_and_c_default}; - - fmap[DPNPFuncName::DPNP_FN_BITWISE_OR][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_bitwise_or_c_default}; - fmap[DPNPFuncName::DPNP_FN_BITWISE_OR][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_bitwise_or_c_default}; - - fmap[DPNPFuncName::DPNP_FN_BITWISE_XOR][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_bitwise_xor_c_default}; - fmap[DPNPFuncName::DPNP_FN_BITWISE_XOR][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_bitwise_xor_c_default}; - - fmap[DPNPFuncName::DPNP_FN_LEFT_SHIFT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_left_shift_c_default}; - fmap[DPNPFuncName::DPNP_FN_LEFT_SHIFT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_left_shift_c_default}; - - fmap[DPNPFuncName::DPNP_FN_RIGHT_SHIFT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_right_shift_c_default}; - fmap[DPNPFuncName::DPNP_FN_RIGHT_SHIFT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_right_shift_c_default}; - - return; -} - -void func_map_init_bitwise(func_map_t &fmap) -{ - func_map_init_bitwise_1arg_1type(fmap); - func_map_init_bitwise_2arg_1type(fmap); - - return; -} diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp index 486851516dcf..20be65f53cab 100644 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp @@ -230,77 +230,6 @@ using dpctl::tensor::kernels::alignment_utils::required_alignment; static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) { - fmap[DPNPFuncName::DPNP_FN_ARCCOS][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_acos_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCCOS][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_acos_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCCOS][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_acos_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCCOS][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_acos_c_default}; - - fmap[DPNPFuncName::DPNP_FN_ARCCOSH][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_acosh_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCCOSH][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_acosh_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCCOSH][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_acosh_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCCOSH][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_acosh_c_default}; - - fmap[DPNPFuncName::DPNP_FN_ARCSIN][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_asin_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCSIN][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_asin_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCSIN][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_asin_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCSIN][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_asin_c_default}; - - fmap[DPNPFuncName::DPNP_FN_ARCSINH][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_asinh_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCSINH][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_asinh_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCSINH][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_asinh_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCSINH][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_asinh_c_default}; - - fmap[DPNPFuncName::DPNP_FN_ARCTAN][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_atan_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_atan_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_atan_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_atan_c_default}; - - fmap[DPNPFuncName::DPNP_FN_ARCTANH][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_atanh_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTANH][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_atanh_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTANH][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_atanh_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTANH][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_atanh_c_default}; - - fmap[DPNPFuncName::DPNP_FN_CBRT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_cbrt_c_default}; - fmap[DPNPFuncName::DPNP_FN_CBRT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_cbrt_c_default}; - fmap[DPNPFuncName::DPNP_FN_CBRT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_cbrt_c_default}; - fmap[DPNPFuncName::DPNP_FN_CBRT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_cbrt_c_default}; - - fmap[DPNPFuncName::DPNP_FN_CEIL][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_ceil_c_default}; - fmap[DPNPFuncName::DPNP_FN_CEIL][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_ceil_c_default}; - fmap[DPNPFuncName::DPNP_FN_CEIL][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_ceil_c_default}; - fmap[DPNPFuncName::DPNP_FN_CEIL][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_ceil_c_default}; fmap[DPNPFuncName::DPNP_FN_COPYTO][eft_BLN][eft_BLN] = { eft_BLN, (void *)dpnp_copyto_c_default}; @@ -378,24 +307,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_COPYTO_EXT][eft_FLT][eft_FLT] = { eft_FLT, (void *)dpnp_copyto_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COS][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_cos_c_default}; - fmap[DPNPFuncName::DPNP_FN_COS][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_cos_c_default}; - fmap[DPNPFuncName::DPNP_FN_COS][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_cos_c_default}; - fmap[DPNPFuncName::DPNP_FN_COS][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_cos_c_default}; - - fmap[DPNPFuncName::DPNP_FN_COSH][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_cosh_c_default}; - fmap[DPNPFuncName::DPNP_FN_COSH][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_cosh_c_default}; - fmap[DPNPFuncName::DPNP_FN_COSH][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_cosh_c_default}; - fmap[DPNPFuncName::DPNP_FN_COSH][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_cosh_c_default}; - fmap[DPNPFuncName::DPNP_FN_DEGREES][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_degrees_c_default}; fmap[DPNPFuncName::DPNP_FN_DEGREES][eft_LNG][eft_LNG] = { @@ -426,87 +337,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_DEGREES_EXT][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_degrees_c_ext}; - fmap[DPNPFuncName::DPNP_FN_EXP2][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_exp2_c_default}; - fmap[DPNPFuncName::DPNP_FN_EXP2][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_exp2_c_default}; - fmap[DPNPFuncName::DPNP_FN_EXP2][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_exp2_c_default}; - fmap[DPNPFuncName::DPNP_FN_EXP2][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_exp2_c_default}; - - fmap[DPNPFuncName::DPNP_FN_EXP][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_exp_c_default}; - fmap[DPNPFuncName::DPNP_FN_EXP][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_exp_c_default}; - fmap[DPNPFuncName::DPNP_FN_EXP][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_exp_c_default}; - fmap[DPNPFuncName::DPNP_FN_EXP][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_exp_c_default}; - - fmap[DPNPFuncName::DPNP_FN_EXPM1][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_expm1_c_default}; - fmap[DPNPFuncName::DPNP_FN_EXPM1][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_expm1_c_default}; - fmap[DPNPFuncName::DPNP_FN_EXPM1][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_expm1_c_default}; - fmap[DPNPFuncName::DPNP_FN_EXPM1][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_expm1_c_default}; - - fmap[DPNPFuncName::DPNP_FN_FABS][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_fabs_c_default}; - fmap[DPNPFuncName::DPNP_FN_FABS][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_fabs_c_default}; - fmap[DPNPFuncName::DPNP_FN_FABS][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_fabs_c_default}; - fmap[DPNPFuncName::DPNP_FN_FABS][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_fabs_c_default}; - - fmap[DPNPFuncName::DPNP_FN_FLOOR][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_floor_c_default}; - fmap[DPNPFuncName::DPNP_FN_FLOOR][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_floor_c_default}; - fmap[DPNPFuncName::DPNP_FN_FLOOR][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_floor_c_default}; - fmap[DPNPFuncName::DPNP_FN_FLOOR][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_floor_c_default}; - - fmap[DPNPFuncName::DPNP_FN_LOG10][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_log10_c_default}; - fmap[DPNPFuncName::DPNP_FN_LOG10][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_log10_c_default}; - fmap[DPNPFuncName::DPNP_FN_LOG10][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_log10_c_default}; - fmap[DPNPFuncName::DPNP_FN_LOG10][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_log10_c_default}; - - fmap[DPNPFuncName::DPNP_FN_LOG1P][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_log1p_c_default}; - fmap[DPNPFuncName::DPNP_FN_LOG1P][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_log1p_c_default}; - fmap[DPNPFuncName::DPNP_FN_LOG1P][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_log1p_c_default}; - fmap[DPNPFuncName::DPNP_FN_LOG1P][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_log1p_c_default}; - - fmap[DPNPFuncName::DPNP_FN_LOG2][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_log2_c_default}; - fmap[DPNPFuncName::DPNP_FN_LOG2][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_log2_c_default}; - fmap[DPNPFuncName::DPNP_FN_LOG2][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_log2_c_default}; - fmap[DPNPFuncName::DPNP_FN_LOG2][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_log2_c_default}; - - fmap[DPNPFuncName::DPNP_FN_LOG][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_log_c_default}; - fmap[DPNPFuncName::DPNP_FN_LOG][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_log_c_default}; - fmap[DPNPFuncName::DPNP_FN_LOG][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_log_c_default}; - fmap[DPNPFuncName::DPNP_FN_LOG][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_log_c_default}; - fmap[DPNPFuncName::DPNP_FN_RADIANS][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_radians_c_default}; fmap[DPNPFuncName::DPNP_FN_RADIANS][eft_LNG][eft_LNG] = { @@ -537,24 +367,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_RADIANS_EXT][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_radians_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SIN][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_sin_c_default}; - fmap[DPNPFuncName::DPNP_FN_SIN][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_sin_c_default}; - fmap[DPNPFuncName::DPNP_FN_SIN][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_sin_c_default}; - fmap[DPNPFuncName::DPNP_FN_SIN][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_sin_c_default}; - - fmap[DPNPFuncName::DPNP_FN_SINH][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_sinh_c_default}; - fmap[DPNPFuncName::DPNP_FN_SINH][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_sinh_c_default}; - fmap[DPNPFuncName::DPNP_FN_SINH][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_sinh_c_default}; - fmap[DPNPFuncName::DPNP_FN_SINH][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_sinh_c_default}; - 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] = { @@ -570,33 +382,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_SQRT_EXT][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_sqrt_c_ext}; - fmap[DPNPFuncName::DPNP_FN_TAN][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_tan_c_default}; - fmap[DPNPFuncName::DPNP_FN_TAN][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_tan_c_default}; - fmap[DPNPFuncName::DPNP_FN_TAN][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_tan_c_default}; - fmap[DPNPFuncName::DPNP_FN_TAN][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_tan_c_default}; - - fmap[DPNPFuncName::DPNP_FN_TANH][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_tanh_c_default}; - fmap[DPNPFuncName::DPNP_FN_TANH][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_tanh_c_default}; - fmap[DPNPFuncName::DPNP_FN_TANH][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_tanh_c_default}; - fmap[DPNPFuncName::DPNP_FN_TANH][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_tanh_c_default}; - - fmap[DPNPFuncName::DPNP_FN_TRUNC][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_trunc_c_default}; - fmap[DPNPFuncName::DPNP_FN_TRUNC][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_trunc_c_default}; - fmap[DPNPFuncName::DPNP_FN_TRUNC][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_trunc_c_default}; - fmap[DPNPFuncName::DPNP_FN_TRUNC][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_trunc_c_default}; - return; } @@ -612,21 +397,6 @@ constexpr T dispatch_erf_op(T elem) } } -template -constexpr T dispatch_sign_op(T elem) -{ - if constexpr (is_any_v) { - if (elem > 0) - return T(1); - if (elem < 0) - return T(-1); - return elem; // elem is 0 - } - else { - return sycl::sign(elem); - } -} - template constexpr auto dispatch_fmod_op(T elem1, T elem2) { @@ -837,45 +607,6 @@ constexpr auto dispatch_fmod_op(T elem1, T elem2) static void func_map_init_elemwise_1arg_1type(func_map_t &fmap) { - fmap[DPNPFuncName::DPNP_FN_CONJUGATE][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_copy_c_default}; - fmap[DPNPFuncName::DPNP_FN_CONJUGATE][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_copy_c_default}; - fmap[DPNPFuncName::DPNP_FN_CONJUGATE][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_copy_c_default}; - fmap[DPNPFuncName::DPNP_FN_CONJUGATE][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_copy_c_default}; - fmap[DPNPFuncName::DPNP_FN_CONJUGATE][eft_C128][eft_C128] = { - eft_C128, (void *)dpnp_conjugate_c_default>}; - - fmap[DPNPFuncName::DPNP_FN_COPY][eft_BLN][eft_BLN] = { - eft_BLN, (void *)dpnp_copy_c_default}; - fmap[DPNPFuncName::DPNP_FN_COPY][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_copy_c_default}; - fmap[DPNPFuncName::DPNP_FN_COPY][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_copy_c_default}; - fmap[DPNPFuncName::DPNP_FN_COPY][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_copy_c_default}; - fmap[DPNPFuncName::DPNP_FN_COPY][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_copy_c_default}; - fmap[DPNPFuncName::DPNP_FN_COPY][eft_C128][eft_C128] = { - eft_C128, (void *)dpnp_copy_c_default>}; - - fmap[DPNPFuncName::DPNP_FN_COPY_EXT][eft_BLN][eft_BLN] = { - eft_BLN, (void *)dpnp_copy_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COPY_EXT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_copy_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COPY_EXT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_copy_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COPY_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_copy_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COPY_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_copy_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COPY_EXT][eft_C64][eft_C64] = { - eft_C64, (void *)dpnp_copy_c_ext>}; - fmap[DPNPFuncName::DPNP_FN_COPY_EXT][eft_C128][eft_C128] = { - eft_C128, (void *)dpnp_copy_c_ext>}; - fmap[DPNPFuncName::DPNP_FN_ERF][eft_INT][eft_INT] = { eft_INT, (void *)dpnp_erf_c_default}; fmap[DPNPFuncName::DPNP_FN_ERF][eft_LNG][eft_LNG] = { @@ -894,55 +625,6 @@ static void func_map_init_elemwise_1arg_1type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_ERF_EXT][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_erf_c_ext}; - fmap[DPNPFuncName::DPNP_FN_FLATTEN][eft_BLN][eft_BLN] = { - eft_BLN, (void *)dpnp_copy_c_default}; - fmap[DPNPFuncName::DPNP_FN_FLATTEN][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_copy_c_default}; - fmap[DPNPFuncName::DPNP_FN_FLATTEN][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_copy_c_default}; - fmap[DPNPFuncName::DPNP_FN_FLATTEN][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_copy_c_default}; - fmap[DPNPFuncName::DPNP_FN_FLATTEN][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_copy_c_default}; - fmap[DPNPFuncName::DPNP_FN_FLATTEN][eft_C128][eft_C128] = { - eft_C128, (void *)dpnp_copy_c_default>}; - - fmap[DPNPFuncName::DPNP_FN_NEGATIVE][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_negative_c_default}; - fmap[DPNPFuncName::DPNP_FN_NEGATIVE][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_negative_c_default}; - fmap[DPNPFuncName::DPNP_FN_NEGATIVE][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_negative_c_default}; - fmap[DPNPFuncName::DPNP_FN_NEGATIVE][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_negative_c_default}; - - fmap[DPNPFuncName::DPNP_FN_RECIP][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_recip_c_default}; - fmap[DPNPFuncName::DPNP_FN_RECIP][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_recip_c_default}; - fmap[DPNPFuncName::DPNP_FN_RECIP][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_recip_c_default}; - fmap[DPNPFuncName::DPNP_FN_RECIP][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_recip_c_default}; - - fmap[DPNPFuncName::DPNP_FN_SIGN][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_sign_c_default}; - fmap[DPNPFuncName::DPNP_FN_SIGN][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_sign_c_default}; - fmap[DPNPFuncName::DPNP_FN_SIGN][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_sign_c_default}; - fmap[DPNPFuncName::DPNP_FN_SIGN][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_sign_c_default}; - - fmap[DPNPFuncName::DPNP_FN_SQUARE][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_square_c_default}; - fmap[DPNPFuncName::DPNP_FN_SQUARE][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_square_c_default}; - fmap[DPNPFuncName::DPNP_FN_SQUARE][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_square_c_default}; - fmap[DPNPFuncName::DPNP_FN_SQUARE][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_square_c_default}; - return; } @@ -1344,47 +1026,6 @@ static void func_map_init_elemwise_1arg_1type(func_map_t &fmap) #include -template -static constexpr DPNPFuncType get_divide_res_type() -{ - constexpr auto widest_type = populate_func_types(); - constexpr auto shortes_type = (widest_type == FT1) ? FT2 : FT1; - - if constexpr (widest_type == DPNPFuncType::DPNP_FT_CMPLX128 || - widest_type == DPNPFuncType::DPNP_FT_DOUBLE) - { - return widest_type; - } - else if constexpr (widest_type == DPNPFuncType::DPNP_FT_CMPLX64) { - if constexpr (shortes_type == DPNPFuncType::DPNP_FT_DOUBLE) { - return DPNPFuncType::DPNP_FT_CMPLX128; - } - else if constexpr (has_fp64::value && - (shortes_type == DPNPFuncType::DPNP_FT_INT || - shortes_type == DPNPFuncType::DPNP_FT_LONG)) - { - return DPNPFuncType::DPNP_FT_CMPLX128; - } - } - else if constexpr (widest_type == DPNPFuncType::DPNP_FT_FLOAT) { - if constexpr (has_fp64::value && - (shortes_type == DPNPFuncType::DPNP_FT_INT || - shortes_type == DPNPFuncType::DPNP_FT_LONG)) - { - return DPNPFuncType::DPNP_FT_DOUBLE; - } - } - else if constexpr (has_fp64::value) { - return DPNPFuncType::DPNP_FT_DOUBLE; - } - else { - return DPNPFuncType::DPNP_FT_FLOAT; - } - return widest_type; -} - template static void func_map_elemwise_2arg_3type_core(func_map_t &fmap) { @@ -1445,270 +1086,7 @@ static void func_map_elemwise_2arg_3type_short_helper(func_map_t &fmap) static void func_map_init_elemwise_2arg_3type(func_map_t &fmap) { - fmap[DPNPFuncName::DPNP_FN_ADD][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_add_c_default}; - fmap[DPNPFuncName::DPNP_FN_ADD][eft_INT][eft_LNG] = { - eft_LNG, (void *)dpnp_add_c_default}; - fmap[DPNPFuncName::DPNP_FN_ADD][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_add_c_default}; - fmap[DPNPFuncName::DPNP_FN_ADD][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_add_c_default}; - fmap[DPNPFuncName::DPNP_FN_ADD][eft_LNG][eft_INT] = { - eft_LNG, (void *)dpnp_add_c_default}; - fmap[DPNPFuncName::DPNP_FN_ADD][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_add_c_default}; - fmap[DPNPFuncName::DPNP_FN_ADD][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_add_c_default}; - fmap[DPNPFuncName::DPNP_FN_ADD][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_add_c_default}; - fmap[DPNPFuncName::DPNP_FN_ADD][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_add_c_default}; - fmap[DPNPFuncName::DPNP_FN_ADD][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_add_c_default}; - fmap[DPNPFuncName::DPNP_FN_ADD][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_add_c_default}; - fmap[DPNPFuncName::DPNP_FN_ADD][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_add_c_default}; - fmap[DPNPFuncName::DPNP_FN_ADD][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_add_c_default}; - fmap[DPNPFuncName::DPNP_FN_ADD][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_add_c_default}; - fmap[DPNPFuncName::DPNP_FN_ADD][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_add_c_default}; - fmap[DPNPFuncName::DPNP_FN_ADD][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_add_c_default}; - - fmap[DPNPFuncName::DPNP_FN_ARCTAN2][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_arctan2_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2][eft_INT][eft_LNG] = { - eft_DBL, (void *)dpnp_arctan2_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_arctan2_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_arctan2_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2][eft_LNG][eft_INT] = { - eft_DBL, (void *)dpnp_arctan2_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_arctan2_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_arctan2_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_arctan2_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_arctan2_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_arctan2_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_arctan2_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_arctan2_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_arctan2_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_arctan2_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_arctan2_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_arctan2_c_default}; - - fmap[DPNPFuncName::DPNP_FN_COPYSIGN][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_copysign_c_default}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN][eft_INT][eft_LNG] = { - eft_DBL, (void *)dpnp_copysign_c_default}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_copysign_c_default}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_copysign_c_default}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN][eft_LNG][eft_INT] = { - eft_DBL, (void *)dpnp_copysign_c_default}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_copysign_c_default}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_copysign_c_default}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_copysign_c_default}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_copysign_c_default}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_copysign_c_default}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_copysign_c_default}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_copysign_c_default}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_copysign_c_default}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_copysign_c_default}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_copysign_c_default}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_copysign_c_default}; - - fmap[DPNPFuncName::DPNP_FN_DIVIDE][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_divide_c_default}; - fmap[DPNPFuncName::DPNP_FN_DIVIDE][eft_INT][eft_LNG] = { - eft_DBL, (void *)dpnp_divide_c_default}; - fmap[DPNPFuncName::DPNP_FN_DIVIDE][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_divide_c_default}; - fmap[DPNPFuncName::DPNP_FN_DIVIDE][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_divide_c_default}; - fmap[DPNPFuncName::DPNP_FN_DIVIDE][eft_LNG][eft_INT] = { - eft_DBL, (void *)dpnp_divide_c_default}; - fmap[DPNPFuncName::DPNP_FN_DIVIDE][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_divide_c_default}; - fmap[DPNPFuncName::DPNP_FN_DIVIDE][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_divide_c_default}; - fmap[DPNPFuncName::DPNP_FN_DIVIDE][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_divide_c_default}; - fmap[DPNPFuncName::DPNP_FN_DIVIDE][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_divide_c_default}; - fmap[DPNPFuncName::DPNP_FN_DIVIDE][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_divide_c_default}; - fmap[DPNPFuncName::DPNP_FN_DIVIDE][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_divide_c_default}; - fmap[DPNPFuncName::DPNP_FN_DIVIDE][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_divide_c_default}; - fmap[DPNPFuncName::DPNP_FN_DIVIDE][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_divide_c_default}; - fmap[DPNPFuncName::DPNP_FN_DIVIDE][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_divide_c_default}; - fmap[DPNPFuncName::DPNP_FN_DIVIDE][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_divide_c_default}; - fmap[DPNPFuncName::DPNP_FN_DIVIDE][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_divide_c_default}; - - fmap[DPNPFuncName::DPNP_FN_FMOD][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_fmod_c_default}; - fmap[DPNPFuncName::DPNP_FN_FMOD][eft_INT][eft_LNG] = { - eft_LNG, (void *)dpnp_fmod_c_default}; - fmap[DPNPFuncName::DPNP_FN_FMOD][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_fmod_c_default}; - fmap[DPNPFuncName::DPNP_FN_FMOD][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_fmod_c_default}; - fmap[DPNPFuncName::DPNP_FN_FMOD][eft_LNG][eft_INT] = { - eft_LNG, (void *)dpnp_fmod_c_default}; - fmap[DPNPFuncName::DPNP_FN_FMOD][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_fmod_c_default}; - fmap[DPNPFuncName::DPNP_FN_FMOD][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_fmod_c_default}; - fmap[DPNPFuncName::DPNP_FN_FMOD][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_fmod_c_default}; - fmap[DPNPFuncName::DPNP_FN_FMOD][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_fmod_c_default}; - fmap[DPNPFuncName::DPNP_FN_FMOD][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_fmod_c_default}; - fmap[DPNPFuncName::DPNP_FN_FMOD][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_fmod_c_default}; - fmap[DPNPFuncName::DPNP_FN_FMOD][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_fmod_c_default}; - fmap[DPNPFuncName::DPNP_FN_FMOD][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_fmod_c_default}; - fmap[DPNPFuncName::DPNP_FN_FMOD][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_fmod_c_default}; - fmap[DPNPFuncName::DPNP_FN_FMOD][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_fmod_c_default}; - fmap[DPNPFuncName::DPNP_FN_FMOD][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_fmod_c_default}; - - fmap[DPNPFuncName::DPNP_FN_HYPOT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_hypot_c_default}; - fmap[DPNPFuncName::DPNP_FN_HYPOT][eft_INT][eft_LNG] = { - eft_DBL, (void *)dpnp_hypot_c_default}; - fmap[DPNPFuncName::DPNP_FN_HYPOT][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_hypot_c_default}; - fmap[DPNPFuncName::DPNP_FN_HYPOT][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_hypot_c_default}; - fmap[DPNPFuncName::DPNP_FN_HYPOT][eft_LNG][eft_INT] = { - eft_DBL, (void *)dpnp_hypot_c_default}; - fmap[DPNPFuncName::DPNP_FN_HYPOT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_hypot_c_default}; - fmap[DPNPFuncName::DPNP_FN_HYPOT][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_hypot_c_default}; - fmap[DPNPFuncName::DPNP_FN_HYPOT][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_hypot_c_default}; - fmap[DPNPFuncName::DPNP_FN_HYPOT][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_hypot_c_default}; - fmap[DPNPFuncName::DPNP_FN_HYPOT][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_hypot_c_default}; - fmap[DPNPFuncName::DPNP_FN_HYPOT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_hypot_c_default}; - fmap[DPNPFuncName::DPNP_FN_HYPOT][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_hypot_c_default}; - fmap[DPNPFuncName::DPNP_FN_HYPOT][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_hypot_c_default}; - fmap[DPNPFuncName::DPNP_FN_HYPOT][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_hypot_c_default}; - fmap[DPNPFuncName::DPNP_FN_HYPOT][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_hypot_c_default}; - fmap[DPNPFuncName::DPNP_FN_HYPOT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_hypot_c_default}; - - fmap[DPNPFuncName::DPNP_FN_MAXIMUM][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_maximum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM][eft_INT][eft_LNG] = { - eft_LNG, (void *)dpnp_maximum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_maximum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_maximum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM][eft_LNG][eft_INT] = { - eft_LNG, (void *)dpnp_maximum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_maximum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_maximum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_maximum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_maximum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_maximum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_maximum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_maximum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_maximum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_maximum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_maximum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_maximum_c_default}; - - fmap[DPNPFuncName::DPNP_FN_MINIMUM][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_minimum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM][eft_INT][eft_LNG] = { - eft_LNG, (void *)dpnp_minimum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_minimum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_minimum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM][eft_LNG][eft_INT] = { - eft_LNG, (void *)dpnp_minimum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_minimum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_minimum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_minimum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_minimum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_minimum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_minimum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_minimum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_minimum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_minimum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_minimum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_minimum_c_default}; - + // Used in dpnp_dot_c fmap[DPNPFuncName::DPNP_FN_MULTIPLY][eft_BLN][eft_BLN] = { eft_BLN, (void *)dpnp_multiply_c_default}; fmap[DPNPFuncName::DPNP_FN_MULTIPLY][eft_BLN][eft_INT] = { @@ -1811,72 +1189,6 @@ static void func_map_init_elemwise_2arg_3type(func_map_t &fmap) (void *)dpnp_multiply_c_default< std::complex, std::complex, std::complex>}; - fmap[DPNPFuncName::DPNP_FN_POWER][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_power_c_default}; - fmap[DPNPFuncName::DPNP_FN_POWER][eft_INT][eft_LNG] = { - eft_LNG, (void *)dpnp_power_c_default}; - fmap[DPNPFuncName::DPNP_FN_POWER][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_power_c_default}; - fmap[DPNPFuncName::DPNP_FN_POWER][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_power_c_default}; - fmap[DPNPFuncName::DPNP_FN_POWER][eft_LNG][eft_INT] = { - eft_LNG, (void *)dpnp_power_c_default}; - fmap[DPNPFuncName::DPNP_FN_POWER][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_power_c_default}; - fmap[DPNPFuncName::DPNP_FN_POWER][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_power_c_default}; - fmap[DPNPFuncName::DPNP_FN_POWER][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_power_c_default}; - fmap[DPNPFuncName::DPNP_FN_POWER][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_power_c_default}; - fmap[DPNPFuncName::DPNP_FN_POWER][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_power_c_default}; - fmap[DPNPFuncName::DPNP_FN_POWER][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_power_c_default}; - fmap[DPNPFuncName::DPNP_FN_POWER][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_power_c_default}; - fmap[DPNPFuncName::DPNP_FN_POWER][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_power_c_default}; - fmap[DPNPFuncName::DPNP_FN_POWER][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_power_c_default}; - fmap[DPNPFuncName::DPNP_FN_POWER][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_power_c_default}; - fmap[DPNPFuncName::DPNP_FN_POWER][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_power_c_default}; - - fmap[DPNPFuncName::DPNP_FN_SUBTRACT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_subtract_c_default}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT][eft_INT][eft_LNG] = { - eft_LNG, (void *)dpnp_subtract_c_default}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_subtract_c_default}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_subtract_c_default}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT][eft_LNG][eft_INT] = { - eft_LNG, (void *)dpnp_subtract_c_default}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_subtract_c_default}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_subtract_c_default}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_subtract_c_default}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_subtract_c_default}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_subtract_c_default}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_subtract_c_default}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_subtract_c_default}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_subtract_c_default}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_subtract_c_default}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_subtract_c_default}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_subtract_c_default}; - func_map_elemwise_2arg_3type_helper(fmap); diff --git a/dpnp/backend/kernels/dpnp_krnl_logic.cpp b/dpnp/backend/kernels/dpnp_krnl_logic.cpp index 0174b47339a8..c8197d728898 100644 --- a/dpnp/backend/kernels/dpnp_krnl_logic.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_logic.cpp @@ -405,278 +405,6 @@ DPCTLSyclEventRef (*dpnp_any_ext_c)(DPCTLSyclQueueRef, const DPCTLEventVectorRef) = dpnp_any_c<_DataType, _ResultType>; -#define MACRO_2ARG_2TYPES_LOGIC_OP(__name__, __operation__) \ - template \ - class __name__##_kernel; \ - \ - template \ - class __name__##_broadcast_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 void *input2_in, \ - const size_t input2_size, const size_t input2_ndim, \ - const shape_elem_type *input2_shape, \ - const shape_elem_type *input2_strides, const size_t *where, \ - const DPCTLEventVectorRef dep_event_vec_ref) \ - { \ - /* avoid warning unused variable*/ \ - (void)where; \ - (void)dep_event_vec_ref; \ - \ - DPCTLSyclEventRef event_ref = nullptr; \ - \ - if (!input1_size || !input2_size) { \ - return event_ref; \ - } \ - \ - sycl::queue q = *(reinterpret_cast(q_ref)); \ - \ - _DataType_input1 *input1_data = \ - static_cast<_DataType_input1 *>(const_cast(input1_in)); \ - _DataType_input2 *input2_data = \ - static_cast<_DataType_input2 *>(const_cast(input2_in)); \ - bool *result = static_cast(result_out); \ - \ - bool use_broadcasting = !array_equal(input1_shape, input1_ndim, \ - input2_shape, input2_ndim); \ - \ - 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; \ - \ - shape_elem_type *input2_shape_offsets = \ - new shape_elem_type[input2_ndim]; \ - \ - get_shape_offsets_inkernel(input2_shape, input2_ndim, \ - input2_shape_offsets); \ - use_strides = \ - use_strides || !array_equal(input2_strides, input2_ndim, \ - input2_shape_offsets, input2_ndim); \ - delete[] input2_shape_offsets; \ - \ - sycl::event event; \ - sycl::range<1> gws(result_size); /* used only when use_broadcasting or \ - use_strides is true */ \ - \ - if (use_broadcasting) { \ - DPNPC_id<_DataType_input1> *input1_it; \ - const size_t input1_it_size_in_bytes = \ - sizeof(DPNPC_id<_DataType_input1>); \ - input1_it = reinterpret_cast *>( \ - dpnp_memory_alloc_c(q_ref, input1_it_size_in_bytes)); \ - new (input1_it) \ - DPNPC_id<_DataType_input1>(q_ref, input1_data, input1_shape, \ - input1_strides, input1_ndim); \ - \ - input1_it->broadcast_to_shape(result_shape, result_ndim); \ - \ - DPNPC_id<_DataType_input2> *input2_it; \ - const size_t input2_it_size_in_bytes = \ - sizeof(DPNPC_id<_DataType_input2>); \ - input2_it = reinterpret_cast *>( \ - dpnp_memory_alloc_c(q_ref, input2_it_size_in_bytes)); \ - new (input2_it) \ - DPNPC_id<_DataType_input2>(q_ref, input2_data, input2_shape, \ - input2_strides, input2_ndim); \ - \ - input2_it->broadcast_to_shape(result_shape, result_ndim); \ - \ - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { \ - const size_t i = global_id[0]; /* for (size_t i = 0; i < \ - result_size; ++i) */ \ - { \ - const _DataType_input1 input1_elem = (*input1_it)[i]; \ - const _DataType_input2 input2_elem = (*input2_it)[i]; \ - result[i] = __operation__; \ - } \ - }; \ - auto kernel_func = [&](sycl::handler &cgh) { \ - cgh.parallel_for>( \ - gws, kernel_parallel_for_func); \ - }; \ - \ - q.submit(kernel_func).wait(); \ - \ - input1_it->~DPNPC_id(); \ - input2_it->~DPNPC_id(); \ - \ - return event_ref; \ - } \ - else if (use_strides) { \ - if ((result_ndim != input1_ndim) || (result_ndim != input2_ndim)) \ - { \ - throw std::runtime_error( \ - "Result ndim=" + std::to_string(result_ndim) + \ - " mismatches with either input1 ndim=" + \ - std::to_string(input1_ndim) + \ - " or input2 ndim=" + std::to_string(input2_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 = 3 * 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); \ - std::copy(input2_strides, input2_strides + result_ndim, \ - strides_host_packed.begin() + 2 * 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) { \ - const 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]; \ - const shape_elem_type *input2_strides_data = \ - &dev_strides_data[2 * result_ndim]; \ - \ - size_t input1_id = 0; \ - size_t input2_id = 0; \ - \ - for (size_t i = 0; i < result_ndim; ++i) { \ - const size_t output_xyz_id = \ - get_xyz_id_by_id_inkernel(output_id, \ - result_strides_data, \ - result_ndim, i); \ - input1_id += output_xyz_id * input1_strides_data[i]; \ - input2_id += output_xyz_id * input2_strides_data[i]; \ - } \ - \ - const _DataType_input1 input1_elem = \ - input1_data[input1_id]; \ - const _DataType_input2 input2_elem = \ - input2_data[input2_id]; \ - result[output_id] = __operation__; \ - } \ - }; \ - auto kernel_func = [&](sycl::handler &cgh) { \ - cgh.depends_on(copy_strides_ev); \ - cgh.parallel_for>( \ - gws, kernel_parallel_for_func); \ - }; \ - \ - q.submit(kernel_func).wait(); \ - \ - sycl::free(dev_strides_data, q); \ - return event_ref; \ - } \ - else { \ - constexpr size_t lws = 64; \ - constexpr unsigned int vec_sz = 8; \ - \ - auto gws_range = sycl::range<1>( \ - ((result_size + lws * vec_sz - 1) / (lws * vec_sz)) * lws); \ - auto lws_range = sycl::range<1>(lws); \ - \ - auto kernel_parallel_for_func = [=](sycl::nd_item<1> nd_it) { \ - auto sg = nd_it.get_sub_group(); \ - const auto max_sg_size = sg.get_max_local_range()[0]; \ - const size_t start = \ - vec_sz * (nd_it.get_group(0) * nd_it.get_local_range(0) + \ - sg.get_group_id()[0] * max_sg_size); \ - \ - if (is_aligned(input1_data) && \ - is_aligned(input2_data) && \ - is_aligned(result) && \ - (start + static_cast(vec_sz) * max_sg_size < \ - result_size)) \ - { \ - auto input1_multi_ptr = sycl::address_space_cast< \ - sycl::access::address_space::global_space, \ - sycl::access::decorated::yes>(&input1_data[start]); \ - auto input2_multi_ptr = sycl::address_space_cast< \ - sycl::access::address_space::global_space, \ - sycl::access::decorated::yes>(&input2_data[start]); \ - auto result_multi_ptr = sycl::address_space_cast< \ - sycl::access::address_space::global_space, \ - sycl::access::decorated::yes>(&result[start]); \ - \ - sycl::vec<_DataType_input1, vec_sz> x1 = \ - sg.load(input1_multi_ptr); \ - sycl::vec<_DataType_input2, vec_sz> x2 = \ - sg.load(input2_multi_ptr); \ - sycl::vec res_vec; \ - \ - for (size_t k = 0; k < vec_sz; ++k) { \ - const _DataType_input1 input1_elem = x1[k]; \ - const _DataType_input2 input2_elem = x2[k]; \ - res_vec[k] = __operation__; \ - } \ - sg.store(result_multi_ptr, res_vec); \ - } \ - else { \ - for (size_t k = start; k < result_size; ++k) { \ - const _DataType_input1 input1_elem = input1_data[k]; \ - const _DataType_input2 input2_elem = input2_data[k]; \ - result[k] = __operation__; \ - } \ - } \ - }; \ - \ - auto kernel_func = [&](sycl::handler &cgh) { \ - cgh.parallel_for>( \ - sycl::nd_range<1>(gws_range, lws_range), \ - kernel_parallel_for_func); \ - }; \ - event = q.submit(kernel_func); \ - } \ - \ - event_ref = reinterpret_cast(&event); \ - return DPCTLEvent_Copy(event_ref); \ - } \ - \ - 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 void *, const size_t, const size_t, \ - const shape_elem_type *, const shape_elem_type *, const size_t *, \ - const DPCTLEventVectorRef) = \ - __name__<_DataType_input1, _DataType_input2>; - void func_map_init_logic(func_map_t &fmap) { fmap[DPNPFuncName::DPNP_FN_ALL][eft_BLN][eft_BLN] = { diff --git a/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp b/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp index b485701154cc..44cd91854df4 100644 --- a/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp @@ -44,379 +44,6 @@ using dpctl::tensor::kernels::alignment_utils::required_alignment; static_assert(__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_VECTOR_ABS_CHANGED, "SYCL DPC++ compiler does not meet minimum version requirement"); -template -class dpnp_around_c_kernel; - -template -DPCTLSyclEventRef dpnp_around_c(DPCTLSyclQueueRef q_ref, - const void *input_in, - void *result_out, - const size_t input_size, - const int decimals, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - (void)decimals; - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!input_size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - sycl::event event; - - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, input_in, input_size); - _DataType *input = input1_ptr.get_ptr(); - _DataType *result = reinterpret_cast<_DataType *>(result_out); - - if constexpr (std::is_same<_DataType, double>::value || - std::is_same<_DataType, float>::value) - { - event = oneapi::mkl::vm::rint(q, input_size, input, result); - } - else { - sycl::range<1> gws(input_size); - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { - size_t i = global_id[0]; - { - result[i] = std::rint(input[i]); - } - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.parallel_for>( - gws, kernel_parallel_for_func); - }; - - event = q.submit(kernel_func); - } - - event_ref = reinterpret_cast(&event); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_around_c(const void *input_in, - void *result_out, - const size_t input_size, - const int decimals) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_around_c<_DataType>( - q_ref, input_in, result_out, input_size, decimals, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_around_default_c)(const void *, void *, const size_t, const int) = - dpnp_around_c<_DataType>; - -template -class dpnp_elemwise_absolute_c_kernel; - -template -DPCTLSyclEventRef - dpnp_elemwise_absolute_c(DPCTLSyclQueueRef q_ref, - const void *input1_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; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - sycl::event event; - - _DataType_input *array1 = - static_cast<_DataType_input *>(const_cast(input1_in)); - _DataType_output *result = static_cast<_DataType_output *>(result1); - - if constexpr (is_any_v<_DataType_input, float, double, std::complex, - std::complex>) - { - event = oneapi::mkl::vm::abs(q, size, array1, result); - } - else { - static_assert( - is_any_v<_DataType_input, int32_t, int64_t>, - "Integer types are only expected to pass in 'abs' kernel"); - static_assert(std::is_same_v<_DataType_input, _DataType_output>, - "Result type must match a type of input data"); - - constexpr size_t lws = 64; - constexpr unsigned int vec_sz = 8; - - auto gws_range = - sycl::range<1>(((size + lws * vec_sz - 1) / (lws * vec_sz)) * lws); - auto lws_range = sycl::range<1>(lws); - - auto kernel_parallel_for_func = [=](sycl::nd_item<1> nd_it) { - auto sg = nd_it.get_sub_group(); - const auto max_sg_size = sg.get_max_local_range()[0]; - const size_t start = - vec_sz * (nd_it.get_group(0) * nd_it.get_local_range(0) + - sg.get_group_id()[0] * max_sg_size); - - if (is_aligned(array1) && - is_aligned(result) && - (start + static_cast(vec_sz) * max_sg_size < size)) - { - auto array_multi_ptr = sycl::address_space_cast< - sycl::access::address_space::global_space, - sycl::access::decorated::yes>(&array1[start]); - auto result_multi_ptr = sycl::address_space_cast< - sycl::access::address_space::global_space, - sycl::access::decorated::yes>(&result[start]); - - sycl::vec<_DataType_input, vec_sz> data_vec = - sg.load(array_multi_ptr); - - sycl::vec<_DataType_output, vec_sz> res_vec = - sycl::abs(data_vec); - - sg.store(result_multi_ptr, res_vec); - } - else { - for (size_t k = start + sg.get_local_id()[0]; k < size; - k += max_sg_size) { - result[k] = std::abs(array1[k]); - } - } - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.parallel_for>( - sycl::nd_range<1>(gws_range, lws_range), - kernel_parallel_for_func); - }; - event = q.submit(kernel_func); - } - - event_ref = reinterpret_cast(&event); - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_elemwise_absolute_c(const void *input1_in, void *result1, size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_elemwise_absolute_c<_DataType, _DataType>( - q_ref, input1_in, result1, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_elemwise_absolute_default_c)(const void *, void *, size_t) = - dpnp_elemwise_absolute_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_cross_c(DPCTLSyclQueueRef q_ref, - void *result_out, - const void *input1_in, - const size_t input1_size, - const shape_elem_type *input1_shape, - const size_t input1_shape_ndim, - const void *input2_in, - const size_t input2_size, - const shape_elem_type *input2_shape, - const size_t input2_shape_ndim, - const size_t *where, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - (void)input1_size; // avoid warning unused variable - (void)input1_shape; - (void)input1_shape_ndim; - (void)input2_size; - (void)input2_shape; - (void)input2_shape_ndim; - (void)where; - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - sycl::queue q = *(reinterpret_cast(q_ref)); - - DPNPC_ptr_adapter<_DataType_input1> input1_ptr(q_ref, input1_in, - input1_size, true); - DPNPC_ptr_adapter<_DataType_input2> input2_ptr(q_ref, input2_in, - input2_size, true); - DPNPC_ptr_adapter<_DataType_output> result_ptr(q_ref, result_out, - input1_size, true, true); - const _DataType_input1 *input1 = input1_ptr.get_ptr(); - const _DataType_input2 *input2 = input2_ptr.get_ptr(); - _DataType_output *result = result_ptr.get_ptr(); - - result[0] = input1[1] * input2[2] - input1[2] * input2[1]; - - result[1] = input1[2] * input2[0] - input1[0] * input2[2]; - - result[2] = input1[0] * input2[1] - input1[1] * input2[0]; - - return event_ref; -} - -template -void dpnp_cross_c(void *result_out, - const void *input1_in, - const size_t input1_size, - const shape_elem_type *input1_shape, - const size_t input1_shape_ndim, - const void *input2_in, - const size_t input2_size, - const shape_elem_type *input2_shape, - const size_t input2_shape_ndim, - const size_t *where) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_cross_c<_DataType_output, _DataType_input1, _DataType_input2>( - q_ref, result_out, input1_in, input1_size, input1_shape, - input1_shape_ndim, input2_in, input2_size, input2_shape, - input2_shape_ndim, where, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_cross_default_c)(void *, - const void *, - const size_t, - const shape_elem_type *, - const size_t, - const void *, - const size_t, - const shape_elem_type *, - const size_t, - const size_t *) = - dpnp_cross_c<_DataType_output, _DataType_input1, _DataType_input2>; - -template -class dpnp_cumprod_c_kernel; - -template -DPCTLSyclEventRef dpnp_cumprod_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; - - if (!size) { - 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, size, true, - true); - _DataType_input *array1 = input1_ptr.get_ptr(); - _DataType_output *result = result_ptr.get_ptr(); - - _DataType_output cur_res = 1; - - for (size_t i = 0; i < size; ++i) { - cur_res *= array1[i]; - result[i] = cur_res; - } - - return event_ref; -} - -template -void dpnp_cumprod_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_cumprod_c<_DataType_input, _DataType_output>( - q_ref, array1_in, result1, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_cumprod_default_c)(void *, void *, size_t) = - dpnp_cumprod_c<_DataType_input, _DataType_output>; - -template -class dpnp_cumsum_c_kernel; - -template -DPCTLSyclEventRef dpnp_cumsum_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; - - if (!size) { - 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, size, true, - true); - _DataType_input *array1 = input1_ptr.get_ptr(); - _DataType_output *result = result_ptr.get_ptr(); - - _DataType_output cur_res = 0; - - for (size_t i = 0; i < size; ++i) { - cur_res += array1[i]; - result[i] = cur_res; - } - - return event_ref; -} - -template -void dpnp_cumsum_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_cumsum_c<_DataType_input, _DataType_output>( - q_ref, array1_in, result1, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_cumsum_default_c)(void *, void *, size_t) = - dpnp_cumsum_c<_DataType_input, _DataType_output>; - template class dpnp_ediff1d_c_kernel; @@ -541,176 +168,6 @@ DPCTLSyclEventRef (*dpnp_ediff1d_ext_c)(DPCTLSyclQueueRef, const DPCTLEventVectorRef) = dpnp_ediff1d_c<_DataType_input, _DataType_output>; -template -class dpnp_floor_divide_c_kernel; - -template -DPCTLSyclEventRef - dpnp_floor_divide_c(DPCTLSyclQueueRef q_ref, - void *result_out, - const void *input1_in, - const size_t input1_size, - const shape_elem_type *input1_shape, - const size_t input1_shape_ndim, - const void *input2_in, - const size_t input2_size, - const shape_elem_type *input2_shape, - const size_t input2_shape_ndim, - const size_t *where, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)where; - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!input1_size || !input2_size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - DPNPC_ptr_adapter<_DataType_input1> input1_ptr(q_ref, input1_in, - input1_size); - DPNPC_ptr_adapter<_DataType_input2> input2_ptr(q_ref, input2_in, - input2_size); - _DataType_input1 *input1_data = input1_ptr.get_ptr(); - _DataType_input2 *input2_data = input2_ptr.get_ptr(); - _DataType_output *result = reinterpret_cast<_DataType_output *>(result_out); - - std::vector result_shape = get_result_shape( - input1_shape, input1_shape_ndim, input2_shape, input2_shape_ndim); - - DPNPC_id<_DataType_input1> *input1_it; - const size_t input1_it_size_in_bytes = sizeof(DPNPC_id<_DataType_input1>); - input1_it = reinterpret_cast *>( - dpnp_memory_alloc_c(q_ref, input1_it_size_in_bytes)); - new (input1_it) DPNPC_id<_DataType_input1>(q_ref, input1_data, input1_shape, - input1_shape_ndim); - - input1_it->broadcast_to_shape(result_shape); - - DPNPC_id<_DataType_input2> *input2_it; - const size_t input2_it_size_in_bytes = sizeof(DPNPC_id<_DataType_input2>); - input2_it = reinterpret_cast *>( - dpnp_memory_alloc_c(q_ref, input2_it_size_in_bytes)); - new (input2_it) DPNPC_id<_DataType_input2>(q_ref, input2_data, input2_shape, - input2_shape_ndim); - - input2_it->broadcast_to_shape(result_shape); - - const size_t result_size = input1_it->get_output_size(); - - sycl::range<1> gws(result_size); - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { - const size_t i = - global_id[0]; /* for (size_t i = 0; i < result_size; ++i) */ - const _DataType_output input1_elem = (*input1_it)[i]; - const _DataType_output input2_elem = (*input2_it)[i]; - - double div = (double)input1_elem / (double)input2_elem; - result[i] = static_cast<_DataType_output>(sycl::floor(div)); - }; - auto kernel_func = [&](sycl::handler &cgh) { - cgh.parallel_for>( - gws, kernel_parallel_for_func); - }; - - sycl::event event; - - if (input1_size == input2_size) { - if constexpr ((std::is_same<_DataType_input1, double>::value || - std::is_same<_DataType_input1, float>::value) && - std::is_same<_DataType_input2, _DataType_input1>::value) - { - event = oneapi::mkl::vm::div(q, input1_size, input1_data, - input2_data, result); - event.wait(); - event = oneapi::mkl::vm::floor(q, input1_size, result, result); - } - else { - event = q.submit(kernel_func); - } - } - else { - event = q.submit(kernel_func); - } - - event.wait(); - - input1_it->~DPNPC_id(); - input2_it->~DPNPC_id(); - - sycl::free(input1_it, q); - sycl::free(input2_it, q); - - return event_ref; -} - -template -void dpnp_floor_divide_c(void *result_out, - const void *input1_in, - const size_t input1_size, - const shape_elem_type *input1_shape, - const size_t input1_shape_ndim, - const void *input2_in, - const size_t input2_size, - const shape_elem_type *input2_shape, - const size_t input2_shape_ndim, - const size_t *where) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_floor_divide_c<_DataType_output, _DataType_input1, - _DataType_input2>( - q_ref, result_out, input1_in, input1_size, input1_shape, - input1_shape_ndim, input2_in, input2_size, input2_shape, - input2_shape_ndim, where, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_floor_divide_default_c)(void *, - const void *, - const size_t, - const shape_elem_type *, - const size_t, - const void *, - const size_t, - const shape_elem_type *, - const size_t, - const size_t *) = - dpnp_floor_divide_c<_DataType_output, _DataType_input1, _DataType_input2>; - -template -DPCTLSyclEventRef (*dpnp_floor_divide_ext_c)(DPCTLSyclQueueRef, - void *, - const void *, - const size_t, - const shape_elem_type *, - const size_t, - const void *, - const size_t, - const shape_elem_type *, - const size_t, - const size_t *, - const DPCTLEventVectorRef) = - dpnp_floor_divide_c<_DataType_output, _DataType_input1, _DataType_input2>; - template class dpnp_modf_c_kernel; @@ -796,363 +253,8 @@ DPCTLSyclEventRef (*dpnp_modf_ext_c)(DPCTLSyclQueueRef, const DPCTLEventVectorRef) = dpnp_modf_c<_DataType_input, _DataType_output>; -template -class dpnp_remainder_c_kernel; - -template -DPCTLSyclEventRef dpnp_remainder_c(DPCTLSyclQueueRef q_ref, - void *result_out, - const void *input1_in, - const size_t input1_size, - const shape_elem_type *input1_shape, - const size_t input1_shape_ndim, - const void *input2_in, - const size_t input2_size, - const shape_elem_type *input2_shape, - const size_t input2_shape_ndim, - const size_t *where, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)where; - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!input1_size || !input2_size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - DPNPC_ptr_adapter<_DataType_input1> input1_ptr(q_ref, input1_in, - input1_size); - DPNPC_ptr_adapter<_DataType_input2> input2_ptr(q_ref, input2_in, - input2_size); - _DataType_input1 *input1_data = input1_ptr.get_ptr(); - _DataType_input2 *input2_data = input2_ptr.get_ptr(); - _DataType_output *result = reinterpret_cast<_DataType_output *>(result_out); - - std::vector result_shape = get_result_shape( - input1_shape, input1_shape_ndim, input2_shape, input2_shape_ndim); - - DPNPC_id<_DataType_input1> *input1_it; - const size_t input1_it_size_in_bytes = sizeof(DPNPC_id<_DataType_input1>); - input1_it = reinterpret_cast *>( - dpnp_memory_alloc_c(q_ref, input1_it_size_in_bytes)); - new (input1_it) DPNPC_id<_DataType_input1>(q_ref, input1_data, input1_shape, - input1_shape_ndim); - - input1_it->broadcast_to_shape(result_shape); - - DPNPC_id<_DataType_input2> *input2_it; - const size_t input2_it_size_in_bytes = sizeof(DPNPC_id<_DataType_input2>); - input2_it = reinterpret_cast *>( - dpnp_memory_alloc_c(q_ref, input2_it_size_in_bytes)); - new (input2_it) DPNPC_id<_DataType_input2>(q_ref, input2_data, input2_shape, - input2_shape_ndim); - - input2_it->broadcast_to_shape(result_shape); - - const size_t result_size = input1_it->get_output_size(); - - sycl::range<1> gws(result_size); - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { - const size_t i = global_id[0]; - const _DataType_output input1_elem = (*input1_it)[i]; - const _DataType_output input2_elem = (*input2_it)[i]; - double fmod_res = sycl::fmod((double)input1_elem, (double)input2_elem); - double add = fmod_res + input2_elem; - result[i] = sycl::fmod(add, (double)input2_elem); - }; - auto kernel_func = [&](sycl::handler &cgh) { - cgh.parallel_for>( - gws, kernel_parallel_for_func); - }; - - sycl::event event; - - if (input1_size == input2_size) { - if constexpr ((std::is_same<_DataType_input1, double>::value || - std::is_same<_DataType_input1, float>::value) && - std::is_same<_DataType_input2, _DataType_input1>::value) - { - event = oneapi::mkl::vm::fmod(q, input1_size, input1_data, - input2_data, result); - event.wait(); - event = oneapi::mkl::vm::add(q, input1_size, result, input2_data, - result); - event.wait(); - event = oneapi::mkl::vm::fmod(q, input1_size, result, input2_data, - result); - } - else { - event = q.submit(kernel_func); - } - } - else { - event = q.submit(kernel_func); - } - - event.wait(); - - input1_it->~DPNPC_id(); - input2_it->~DPNPC_id(); - - return event_ref; -} - -template -void dpnp_remainder_c(void *result_out, - const void *input1_in, - const size_t input1_size, - const shape_elem_type *input1_shape, - const size_t input1_shape_ndim, - const void *input2_in, - const size_t input2_size, - const shape_elem_type *input2_shape, - const size_t input2_shape_ndim, - const size_t *where) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_remainder_c<_DataType_output, _DataType_input1, _DataType_input2>( - q_ref, result_out, input1_in, input1_size, input1_shape, - input1_shape_ndim, input2_in, input2_size, input2_shape, - input2_shape_ndim, where, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_remainder_default_c)(void *, - const void *, - const size_t, - const shape_elem_type *, - const size_t, - const void *, - const size_t, - const shape_elem_type *, - const size_t, - const size_t *) = - dpnp_remainder_c<_DataType_output, _DataType_input1, _DataType_input2>; - -template -class dpnp_trapz_c_kernel; - -template -DPCTLSyclEventRef dpnp_trapz_c(DPCTLSyclQueueRef q_ref, - const void *array1_in, - const void *array2_in, - void *result1, - double dx, - size_t array1_size, - size_t array2_size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if ((array1_in == nullptr) || (array2_in == nullptr && array2_size > 1)) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - sycl::event event; - - DPNPC_ptr_adapter<_DataType_input1> input1_ptr(q_ref, array1_in, - array1_size); - DPNPC_ptr_adapter<_DataType_input2> input2_ptr(q_ref, array2_in, - array2_size); - _DataType_input1 *array1 = input1_ptr.get_ptr(); - _DataType_input2 *array2 = input2_ptr.get_ptr(); - _DataType_output *result = reinterpret_cast<_DataType_output *>(result1); - - if (array1_size < 2) { - const _DataType_output init_val = 0; - q.memcpy(result, &init_val, sizeof(_DataType_output)) - .wait(); // result[0] = 0; - - return event_ref; - } - - if (array1_size == array2_size) { - size_t cur_res_size = array1_size - 2; - - _DataType_output *cur_res = reinterpret_cast<_DataType_output *>( - sycl::malloc_shared((cur_res_size) * sizeof(_DataType_output), q)); - - sycl::range<1> gws(cur_res_size); - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { - size_t i = global_id[0]; - { - cur_res[i] = array1[i + 1] * (array2[i + 2] - array2[i]); - } - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.parallel_for>( - gws, kernel_parallel_for_func); - }; - - event = q.submit(kernel_func); - - event.wait(); - - shape_elem_type _shape = cur_res_size; - dpnp_sum_c<_DataType_output, _DataType_output>(result, cur_res, &_shape, - 1, NULL, 0, NULL, NULL); - - sycl::free(cur_res, q); - - result[0] += array1[0] * (array2[1] - array2[0]) + - array1[array1_size - 1] * - (array2[array2_size - 1] - array2[array2_size - 2]); - - result[0] *= 0.5; - } - else { - shape_elem_type _shape = array1_size; - dpnp_sum_c<_DataType_output, _DataType_input1>(result, array1, &_shape, - 1, NULL, 0, NULL, NULL); - - result[0] -= (array1[0] + array1[array1_size - 1]) * 0.5; - result[0] *= dx; - } - return event_ref; -} - -template -void dpnp_trapz_c(const void *array1_in, - const void *array2_in, - void *result1, - double dx, - size_t array1_size, - size_t array2_size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_trapz_c<_DataType_input1, _DataType_input2, _DataType_output>( - q_ref, array1_in, array2_in, result1, dx, array1_size, array2_size, - dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_trapz_default_c)(const void *, - const void *, - void *, - double, - size_t, - size_t) = - dpnp_trapz_c<_DataType_input1, _DataType_input2, _DataType_output>; - -template -DPCTLSyclEventRef (*dpnp_trapz_ext_c)(DPCTLSyclQueueRef, - const void *, - const void *, - void *, - double, - size_t, - size_t, - const DPCTLEventVectorRef) = - dpnp_trapz_c<_DataType_input1, _DataType_input2, _DataType_output>; - void func_map_init_mathematical(func_map_t &fmap) { - fmap[DPNPFuncName::DPNP_FN_ABSOLUTE][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_elemwise_absolute_default_c}; - fmap[DPNPFuncName::DPNP_FN_ABSOLUTE][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_elemwise_absolute_default_c}; - fmap[DPNPFuncName::DPNP_FN_ABSOLUTE][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_elemwise_absolute_default_c}; - fmap[DPNPFuncName::DPNP_FN_ABSOLUTE][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_elemwise_absolute_default_c}; - - fmap[DPNPFuncName::DPNP_FN_AROUND][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_around_default_c}; - fmap[DPNPFuncName::DPNP_FN_AROUND][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_around_default_c}; - fmap[DPNPFuncName::DPNP_FN_AROUND][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_around_default_c}; - fmap[DPNPFuncName::DPNP_FN_AROUND][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_around_default_c}; - - fmap[DPNPFuncName::DPNP_FN_CROSS][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_cross_default_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS][eft_INT][eft_LNG] = { - eft_LNG, (void *)dpnp_cross_default_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_cross_default_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_cross_default_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS][eft_LNG][eft_INT] = { - eft_LNG, (void *)dpnp_cross_default_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_cross_default_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_cross_default_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_cross_default_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_cross_default_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_cross_default_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_cross_default_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_cross_default_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_cross_default_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_cross_default_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_cross_default_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_cross_default_c}; - - fmap[DPNPFuncName::DPNP_FN_CUMPROD][eft_INT][eft_INT] = { - eft_LNG, (void *)dpnp_cumprod_default_c}; - fmap[DPNPFuncName::DPNP_FN_CUMPROD][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_cumprod_default_c}; - fmap[DPNPFuncName::DPNP_FN_CUMPROD][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_cumprod_default_c}; - fmap[DPNPFuncName::DPNP_FN_CUMPROD][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_cumprod_default_c}; - - fmap[DPNPFuncName::DPNP_FN_CUMSUM][eft_INT][eft_INT] = { - eft_LNG, (void *)dpnp_cumsum_default_c}; - fmap[DPNPFuncName::DPNP_FN_CUMSUM][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_cumsum_default_c}; - fmap[DPNPFuncName::DPNP_FN_CUMSUM][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_cumsum_default_c}; - fmap[DPNPFuncName::DPNP_FN_CUMSUM][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_cumsum_default_c}; fmap[DPNPFuncName::DPNP_FN_EDIFF1D][eft_INT][eft_INT] = { eft_LNG, (void *)dpnp_ediff1d_default_c}; @@ -1172,43 +274,6 @@ void func_map_init_mathematical(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_EDIFF1D_EXT][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_ediff1d_ext_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE][eft_INT][eft_INT] = { - eft_INT, - (void *)dpnp_floor_divide_default_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE][eft_INT][eft_LNG] = { - eft_LNG, - (void *)dpnp_floor_divide_default_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_floor_divide_default_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_floor_divide_default_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE][eft_LNG][eft_INT] = { - eft_LNG, - (void *)dpnp_floor_divide_default_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE][eft_LNG][eft_LNG] = { - eft_LNG, - (void *)dpnp_floor_divide_default_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_floor_divide_default_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_floor_divide_default_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_floor_divide_default_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_floor_divide_default_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_floor_divide_default_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_floor_divide_default_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_floor_divide_default_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_floor_divide_default_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_floor_divide_default_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_floor_divide_default_c}; - fmap[DPNPFuncName::DPNP_FN_MODF][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_modf_default_c}; fmap[DPNPFuncName::DPNP_FN_MODF][eft_LNG][eft_LNG] = { @@ -1227,104 +292,5 @@ void func_map_init_mathematical(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_MODF_EXT][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_modf_ext_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_remainder_default_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER][eft_INT][eft_LNG] = { - eft_LNG, (void *)dpnp_remainder_default_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_remainder_default_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_remainder_default_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER][eft_LNG][eft_INT] = { - eft_LNG, (void *)dpnp_remainder_default_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_remainder_default_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_remainder_default_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_remainder_default_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_remainder_default_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_remainder_default_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_remainder_default_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_remainder_default_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_remainder_default_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_remainder_default_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_remainder_default_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_remainder_default_c}; - - fmap[DPNPFuncName::DPNP_FN_TRAPZ][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_trapz_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ][eft_INT][eft_LNG] = { - eft_DBL, (void *)dpnp_trapz_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_trapz_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_trapz_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ][eft_LNG][eft_INT] = { - eft_DBL, (void *)dpnp_trapz_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_trapz_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_trapz_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_trapz_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_trapz_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_trapz_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_trapz_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_trapz_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_trapz_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_trapz_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_trapz_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_trapz_default_c}; - - fmap[DPNPFuncName::DPNP_FN_TRAPZ_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_trapz_ext_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ_EXT][eft_INT][eft_LNG] = { - eft_DBL, (void *)dpnp_trapz_ext_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ_EXT][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_trapz_ext_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ_EXT][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_trapz_ext_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ_EXT][eft_LNG][eft_INT] = { - eft_DBL, (void *)dpnp_trapz_ext_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_trapz_ext_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ_EXT][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_trapz_ext_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ_EXT][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_trapz_ext_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ_EXT][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_trapz_ext_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ_EXT][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_trapz_ext_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_trapz_ext_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ_EXT][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_trapz_ext_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ_EXT][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_trapz_ext_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ_EXT][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_trapz_ext_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ_EXT][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_trapz_ext_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_trapz_ext_c}; - return; } diff --git a/dpnp/backend/src/dpnp_fptr.hpp b/dpnp/backend/src/dpnp_fptr.hpp index 20fc5305e9a2..2a9c42eb1720 100644 --- a/dpnp/backend/src/dpnp_fptr.hpp +++ b/dpnp/backend/src/dpnp_fptr.hpp @@ -326,7 +326,6 @@ static constexpr DPNPFuncType get_floating_res_type() * FPTR interface initialization functions */ void func_map_init_arraycreation(func_map_t &fmap); -void func_map_init_bitwise(func_map_t &fmap); void func_map_init_elemwise(func_map_t &fmap); void func_map_init_fft_func(func_map_t &fmap); void func_map_init_indexing_func(func_map_t &fmap); diff --git a/dpnp/backend/src/dpnp_iface_fptr.cpp b/dpnp/backend/src/dpnp_iface_fptr.cpp index 460896bfa2dd..f8214212728d 100644 --- a/dpnp/backend/src/dpnp_iface_fptr.cpp +++ b/dpnp/backend/src/dpnp_iface_fptr.cpp @@ -167,7 +167,6 @@ static func_map_t func_map_init() func_map_t fmap; func_map_init_arraycreation(fmap); - func_map_init_bitwise(fmap); func_map_init_elemwise(fmap); func_map_init_fft_func(fmap); func_map_init_indexing_func(fmap); diff --git a/dpnp/dpnp_algo/CMakeLists.txt b/dpnp/dpnp_algo/CMakeLists.txt index 2c3a49c6be4b..1aea452c5d93 100644 --- a/dpnp/dpnp_algo/CMakeLists.txt +++ b/dpnp/dpnp_algo/CMakeLists.txt @@ -3,7 +3,6 @@ set(dpnp_algo_pyx_deps ${CMAKE_CURRENT_SOURCE_DIR}/dpnp_algo_statistics.pxi ${CMAKE_CURRENT_SOURCE_DIR}/dpnp_algo_trigonometric.pxi ${CMAKE_CURRENT_SOURCE_DIR}/dpnp_algo_sorting.pxi - ${CMAKE_CURRENT_SOURCE_DIR}/dpnp_algo_arraycreation.pxi ${CMAKE_CURRENT_SOURCE_DIR}/dpnp_algo_mathematical.pxi ${CMAKE_CURRENT_SOURCE_DIR}/dpnp_algo_indexing.pxi ${CMAKE_CURRENT_SOURCE_DIR}/dpnp_algo_logic.pxi diff --git a/dpnp/dpnp_algo/dpnp_algo.pxd b/dpnp/dpnp_algo/dpnp_algo.pxd index 4e91151697c0..37663bee8343 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pxd +++ b/dpnp/dpnp_algo/dpnp_algo.pxd @@ -35,7 +35,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na cdef enum DPNPFuncName "DPNPFuncName": DPNP_FN_ALLCLOSE_EXT DPNP_FN_CHOOSE_EXT - DPNP_FN_COPY_EXT DPNP_FN_CORRELATE_EXT DPNP_FN_DEGREES_EXT DPNP_FN_EDIFF1D_EXT @@ -172,11 +171,6 @@ cpdef dpnp_descriptor dpnp_isclose(dpnp_descriptor input1, dpnp_descriptor input double rtol=*, double atol=*, cpp_bool equal_nan=*) -""" -Array creation routines -""" -cpdef dpnp_descriptor dpnp_copy(dpnp_descriptor x1) - """ Mathematical functions """ diff --git a/dpnp/dpnp_algo/dpnp_algo.pyx b/dpnp/dpnp_algo/dpnp_algo.pyx index c8d99c56912c..4c560d50e0b3 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pyx +++ b/dpnp/dpnp_algo/dpnp_algo.pyx @@ -58,7 +58,6 @@ __all__ = [ ] -include "dpnp_algo_arraycreation.pxi" include "dpnp_algo_indexing.pxi" include "dpnp_algo_logic.pxi" include "dpnp_algo_mathematical.pxi" diff --git a/dpnp/dpnp_algo/dpnp_algo_arraycreation.pxi b/dpnp/dpnp_algo/dpnp_algo_arraycreation.pxi deleted file mode 100644 index bd86a4618484..000000000000 --- a/dpnp/dpnp_algo/dpnp_algo_arraycreation.pxi +++ /dev/null @@ -1,78 +0,0 @@ -# cython: language_level=3 -# cython: linetrace=True -# -*- coding: utf-8 -*- -# ***************************************************************************** -# 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. -# ***************************************************************************** - -"""Module Backend (array creation part) - -This module contains interface functions between C backend layer -and the rest of the library - -""" - -# NO IMPORTs here. All imports must be placed into main "dpnp_algo.pyx" file - -__all__ += [ - "dpnp_copy", -] - - -ctypedef c_dpctl.DPCTLSyclEventRef(*custom_1in_1out_func_ptr_t)(c_dpctl.DPCTLSyclQueueRef, - void *, - void * , - const int , - shape_elem_type * , - shape_elem_type * , - const size_t, - const size_t, - const c_dpctl.DPCTLEventVectorRef) -ctypedef c_dpctl.DPCTLSyclEventRef(*ftpr_custom_vander_1in_1out_t)(c_dpctl.DPCTLSyclQueueRef, - void * , void * , size_t, size_t, int, - const c_dpctl.DPCTLEventVectorRef) except + -ctypedef c_dpctl.DPCTLSyclEventRef(*custom_arraycreation_1in_1out_func_ptr_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 shape_elem_type *, - const size_t, - const c_dpctl.DPCTLEventVectorRef) -ctypedef c_dpctl.DPCTLSyclEventRef(*custom_indexing_1out_func_ptr_t)(c_dpctl.DPCTLSyclQueueRef, - void * , - const size_t , - const size_t , - const int, - const c_dpctl.DPCTLEventVectorRef) except + - - -cpdef utils.dpnp_descriptor dpnp_copy(utils.dpnp_descriptor x1): - return call_fptr_1in_1out_strides(DPNP_FN_COPY_EXT, x1) diff --git a/dpnp/dpnp_algo/dpnp_algo_sorting.pxi b/dpnp/dpnp_algo/dpnp_algo_sorting.pxi index 4947fa9e41d1..5da472a246bd 100644 --- a/dpnp/dpnp_algo/dpnp_algo_sorting.pxi +++ b/dpnp/dpnp_algo/dpnp_algo_sorting.pxi @@ -58,7 +58,7 @@ cpdef utils.dpnp_descriptor dpnp_partition(utils.dpnp_descriptor arr, int kth, a cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_PARTITION_EXT, param1_type, param1_type) - cdef utils.dpnp_descriptor arr2 = dpnp_copy(arr) + cdef utils.dpnp_descriptor arr2 = dpnp.get_dpnp_descriptor(arr.get_pyobj().copy(), copy_when_nondefault_queue=False) arr_obj = arr.get_array()