diff --git a/CHANGELOG.md b/CHANGELOG.md index 5519ebb6774d..f2533b207f97 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -28,6 +28,7 @@ Also, that release drops support for Python 3.9, making Python 3.10 the minimum * Changed the license from `BSD-2-Clause` to `BSD-3-Clause` [#2593](https://github.com/IntelPython/dpnp/pull/2593) * Defined explicit versions range of the Python interpreter which is needed during the build [#2634](https://github.com/IntelPython/dpnp/pull/2634) * Aligned documentation with NumPy and CuPy style by using short function names [#2633](https://github.com/IntelPython/dpnp/pull/2633) +* Redesigned `dpnp.modf` function to be a part of `ufunc` and `vm` pybind11 extensions [#2654](https://github.com/IntelPython/dpnp/pull/2654) ### Deprecated diff --git a/dpnp/backend/CMakeLists.txt b/dpnp/backend/CMakeLists.txt index f88e52e352f5..18085c20543c 100644 --- a/dpnp/backend/CMakeLists.txt +++ b/dpnp/backend/CMakeLists.txt @@ -30,7 +30,6 @@ set(DPNP_SRC kernels/dpnp_krnl_arraycreation.cpp kernels/dpnp_krnl_common.cpp - kernels/dpnp_krnl_mathematical.cpp kernels/dpnp_krnl_random.cpp kernels/dpnp_krnl_sorting.cpp src/dpnp_iface_fptr.cpp diff --git a/dpnp/backend/extensions/ufunc/CMakeLists.txt b/dpnp/backend/extensions/ufunc/CMakeLists.txt index 62e8249b508f..5be39dfa8477 100644 --- a/dpnp/backend/extensions/ufunc/CMakeLists.txt +++ b/dpnp/backend/extensions/ufunc/CMakeLists.txt @@ -47,6 +47,7 @@ set(_elementwise_sources ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/lcm.cpp ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/ldexp.cpp ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/logaddexp2.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/modf.cpp ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/nan_to_num.cpp ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/radians.cpp ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/sinc.cpp diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp index f59f6c709a74..2cf81aead4e1 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp @@ -46,6 +46,7 @@ #include "lcm.hpp" #include "ldexp.hpp" #include "logaddexp2.hpp" +#include "modf.hpp" #include "nan_to_num.hpp" #include "radians.hpp" #include "sinc.hpp" @@ -78,6 +79,7 @@ void init_elementwise_functions(py::module_ m) init_lcm(m); init_ldexp(m); init_logaddexp2(m); + init_modf(m); init_nan_to_num(m); init_radians(m); init_sinc(m); diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/frexp.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/frexp.cpp index acb576aeba86..4439f1e76993 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/frexp.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/frexp.cpp @@ -37,6 +37,7 @@ #include "frexp.hpp" #include "kernels/elementwise_functions/frexp.hpp" +#include "populate.hpp" // include a local copy of elementwise common header from dpctl tensor: // dpctl/tensor/libtensor/source/elementwise_functions/elementwise_functions.hpp @@ -65,10 +66,9 @@ namespace td_int_ns = py_int::type_dispatch; namespace td_ns = dpctl::tensor::type_dispatch; using dpnp::kernels::frexp::FrexpFunctor; -using ext::common::init_dispatch_vector; template -struct FrexpOutputType +struct OutputType { using table_type = std::disjunction< // disjunction is C++17 // feature, supported by DPC++ @@ -81,15 +81,13 @@ struct FrexpOutputType using value_type2 = typename table_type::result_type2; }; -// contiguous implementation - template -using FrexpContigFunctor = +using ContigFunctor = ew_cmn_ns::UnaryTwoOutputsContigFunctor; -// strided implementation - template -using FrexpStridedFunctor = ew_cmn_ns::UnaryTwoOutputsStridedFunctor< +using StridedFunctor = ew_cmn_ns::UnaryTwoOutputsStridedFunctor< argTy, resTy1, resTy2, IndexerT, FrexpFunctor>; -template -class frexp_contig_kernel; - -template -sycl::event frexp_contig_impl(sycl::queue &exec_q, - size_t nelems, - const char *arg_p, - char *res1_p, - char *res2_p, - const std::vector &depends = {}) -{ - return ew_cmn_ns::unary_two_outputs_contig_impl< - argTy, FrexpOutputType, FrexpContigFunctor, frexp_contig_kernel>( - exec_q, nelems, arg_p, res1_p, res2_p, depends); -} - -template -struct FrexpContigFactory -{ - fnT get() - { - if constexpr (std::is_same_v::value_type1, - void> || - std::is_same_v::value_type2, - void>) - { - fnT fn = nullptr; - return fn; - } - else { - fnT fn = frexp_contig_impl; - return fn; - } - } -}; - -template -class frexp_strided_kernel; - -template -sycl::event - frexp_strided_impl(sycl::queue &exec_q, - size_t nelems, - int nd, - const ssize_t *shape_and_strides, - const char *arg_p, - ssize_t arg_offset, - char *res1_p, - ssize_t res1_offset, - char *res2_p, - ssize_t res2_offset, - const std::vector &depends, - const std::vector &additional_depends) -{ - return ew_cmn_ns::unary_two_outputs_strided_impl< - argTy, FrexpOutputType, FrexpStridedFunctor, frexp_strided_kernel>( - exec_q, nelems, nd, shape_and_strides, arg_p, arg_offset, res1_p, - res1_offset, res2_p, res2_offset, depends, additional_depends); -} - -template -struct FrexpStridedFactory -{ - fnT get() - { - if constexpr (std::is_same_v::value_type1, - void> || - std::is_same_v::value_type2, - void>) - { - fnT fn = nullptr; - return fn; - } - else { - fnT fn = frexp_strided_impl; - return fn; - } - } -}; - -template -struct FrexpTypeMapFactory -{ - /*! @brief get typeid for output type of sycl::frexp(T x) */ - std::enable_if_t>::value, - std::pair> - get() - { - using rT1 = typename FrexpOutputType::value_type1; - using rT2 = typename FrexpOutputType::value_type2; - return std::make_pair(td_ns::GetTypeid{}.get(), - td_ns::GetTypeid{}.get()); - } -}; - using ew_cmn_ns::unary_two_outputs_contig_impl_fn_ptr_t; using ew_cmn_ns::unary_two_outputs_strided_impl_fn_ptr_t; @@ -216,15 +113,7 @@ static std::pair frexp_output_typeid_vector[td_ns::num_types]; static unary_two_outputs_strided_impl_fn_ptr_t frexp_strided_dispatch_vector[td_ns::num_types]; -void populate_frexp_dispatch_vectors(void) -{ - init_dispatch_vector(frexp_contig_dispatch_vector); - init_dispatch_vector(frexp_strided_dispatch_vector); - init_dispatch_vector, FrexpTypeMapFactory>( - frexp_output_typeid_vector); -}; +MACRO_POPULATE_DISPATCH_2OUTS_VECTORS(frexp); } // namespace impl void init_frexp(py::module_ m) diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/modf.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/modf.cpp new file mode 100644 index 000000000000..f8aab23d5630 --- /dev/null +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/modf.cpp @@ -0,0 +1,146 @@ +//***************************************************************************** +// Copyright (c) 2025, 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. +// - Neither the name of the copyright holder nor the names of its contributors +// may be used to endorse or promote products derived from this software +// without specific prior written permission. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE +// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE +// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR +// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF +// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF +// THE POSSIBILITY OF SUCH DAMAGE. +//***************************************************************************** + +#include +#include +#include +#include + +#include + +#include "dpctl4pybind11.hpp" + +#include "kernels/elementwise_functions/modf.hpp" +#include "modf.hpp" +#include "populate.hpp" + +// include a local copy of elementwise common header from dpctl tensor: +// dpctl/tensor/libtensor/source/elementwise_functions/elementwise_functions.hpp +// TODO: replace by including dpctl header once available +#include "../../elementwise_functions/elementwise_functions.hpp" + +#include "../../elementwise_functions/common.hpp" +#include "../../elementwise_functions/type_dispatch_building.hpp" + +// utils extension header +#include "ext/common.hpp" + +// dpctl tensor headers +#include "kernels/elementwise_functions/common.hpp" +#include "utils/type_dispatch.hpp" + +namespace dpnp::extensions::ufunc +{ +namespace py = pybind11; +namespace py_int = dpnp::extensions::py_internal; + +namespace impl +{ +namespace ew_cmn_ns = dpnp::extensions::py_internal::elementwise_common; +namespace td_int_ns = py_int::type_dispatch; +namespace td_ns = dpctl::tensor::type_dispatch; + +using dpnp::kernels::modf::ModfFunctor; + +template +struct OutputType +{ + using table_type = std::disjunction< // disjunction is C++17 + // feature, supported by DPC++ + td_int_ns::TypeMapTwoResultsEntry, + td_int_ns::TypeMapTwoResultsEntry, + td_int_ns::TypeMapTwoResultsEntry, + td_int_ns::DefaultTwoResultsEntry>; + using value_type1 = typename table_type::result_type1; + using value_type2 = typename table_type::result_type2; +}; + +template +using ContigFunctor = + ew_cmn_ns::UnaryTwoOutputsContigFunctor, + vec_sz, + n_vecs, + enable_sg_loadstore>; + +template +using StridedFunctor = ew_cmn_ns::UnaryTwoOutputsStridedFunctor< + argTy, + resTy1, + resTy2, + IndexerT, + ModfFunctor>; + +using ew_cmn_ns::unary_two_outputs_contig_impl_fn_ptr_t; +using ew_cmn_ns::unary_two_outputs_strided_impl_fn_ptr_t; + +static unary_two_outputs_contig_impl_fn_ptr_t + modf_contig_dispatch_vector[td_ns::num_types]; +static std::pair modf_output_typeid_vector[td_ns::num_types]; +static unary_two_outputs_strided_impl_fn_ptr_t + modf_strided_dispatch_vector[td_ns::num_types]; + +MACRO_POPULATE_DISPATCH_2OUTS_VECTORS(modf); +} // namespace impl + +void init_modf(py::module_ m) +{ + using arrayT = dpctl::tensor::usm_ndarray; + using event_vecT = std::vector; + { + impl::populate_modf_dispatch_vectors(); + using impl::modf_contig_dispatch_vector; + using impl::modf_output_typeid_vector; + using impl::modf_strided_dispatch_vector; + + auto modf_pyapi = [&](const arrayT &src, const arrayT &dst1, + const arrayT &dst2, sycl::queue &exec_q, + const event_vecT &depends = {}) { + return py_int::py_unary_two_outputs_ufunc( + src, dst1, dst2, exec_q, depends, modf_output_typeid_vector, + modf_contig_dispatch_vector, modf_strided_dispatch_vector); + }; + m.def("_modf", modf_pyapi, "", py::arg("src"), py::arg("dst1"), + py::arg("dst2"), py::arg("sycl_queue"), + py::arg("depends") = py::list()); + + auto modf_result_type_pyapi = [&](const py::dtype &dtype) { + return py_int::py_unary_two_outputs_ufunc_result_type( + dtype, modf_output_typeid_vector); + }; + m.def("_modf_result_type", modf_result_type_pyapi); + } +} +} // namespace dpnp::extensions::ufunc diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/modf.hpp b/dpnp/backend/extensions/ufunc/elementwise_functions/modf.hpp new file mode 100644 index 000000000000..3c57ba6834f4 --- /dev/null +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/modf.hpp @@ -0,0 +1,38 @@ +//***************************************************************************** +// Copyright (c) 2025, 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. +// - Neither the name of the copyright holder nor the names of its contributors +// may be used to endorse or promote products derived from this software +// without specific prior written permission. +// +// 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. +//***************************************************************************** + +#pragma once + +#include + +namespace py = pybind11; + +namespace dpnp::extensions::ufunc +{ +void init_modf(py::module_ m); +} // namespace dpnp::extensions::ufunc diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/populate.hpp b/dpnp/backend/extensions/ufunc/elementwise_functions/populate.hpp index 83d265bd6186..9715b1cc96d4 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/populate.hpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/populate.hpp @@ -125,6 +125,109 @@ namespace ext_ns = ext::common; __name__##_output_typeid_vector); \ }; +/** + * @brief A macro used to define factories and a populating unary universal + * functions with two output arrays. + */ +#define MACRO_POPULATE_DISPATCH_2OUTS_VECTORS(__name__) \ + template \ + class __name__##_contig_kernel; \ + \ + template \ + sycl::event __name__##_contig_impl( \ + sycl::queue &exec_q, size_t nelems, const char *arg_p, char *res1_p, \ + char *res2_p, const std::vector &depends = {}) \ + { \ + return ew_cmn_ns::unary_two_outputs_contig_impl< \ + argTy, OutputType, ContigFunctor, __name__##_contig_kernel>( \ + exec_q, nelems, arg_p, res1_p, res2_p, depends); \ + } \ + \ + template \ + struct ContigFactory \ + { \ + fnT get() \ + { \ + if constexpr (std::is_same_v::value_type1, \ + void> || \ + std::is_same_v::value_type2, \ + void>) \ + { \ + fnT fn = nullptr; \ + return fn; \ + } \ + else { \ + fnT fn = __name__##_contig_impl; \ + return fn; \ + } \ + } \ + }; \ + \ + template \ + struct TypeMapFactory \ + { \ + std::enable_if_t>::value, \ + std::pair> \ + get() \ + { \ + using rT1 = typename OutputType::value_type1; \ + using rT2 = typename OutputType::value_type2; \ + return std::make_pair(td_ns::GetTypeid{}.get(), \ + td_ns::GetTypeid{}.get()); \ + } \ + }; \ + \ + template \ + class __name__##_strided_kernel; \ + \ + template \ + sycl::event __name__##_strided_impl( \ + sycl::queue &exec_q, size_t nelems, int nd, \ + const py::ssize_t *shape_and_strides, const char *arg_p, \ + py::ssize_t arg_offset, char *res1_p, py::ssize_t res1_offset, \ + char *res2_p, py::ssize_t res2_offset, \ + const std::vector &depends, \ + const std::vector &additional_depends) \ + { \ + return ew_cmn_ns::unary_two_outputs_strided_impl< \ + argTy, OutputType, StridedFunctor, __name__##_strided_kernel>( \ + exec_q, nelems, nd, shape_and_strides, arg_p, arg_offset, res1_p, \ + res1_offset, res2_p, res2_offset, depends, additional_depends); \ + } \ + \ + template \ + struct StridedFactory \ + { \ + fnT get() \ + { \ + if constexpr (std::is_same_v::value_type1, \ + void> || \ + std::is_same_v::value_type2, \ + void>) \ + { \ + fnT fn = nullptr; \ + return fn; \ + } \ + else { \ + fnT fn = __name__##_strided_impl; \ + return fn; \ + } \ + } \ + }; \ + \ + void populate_##__name__##_dispatch_vectors(void) \ + { \ + ext_ns::init_dispatch_vector( \ + __name__##_contig_dispatch_vector); \ + ext_ns::init_dispatch_vector( \ + __name__##_strided_dispatch_vector); \ + ext_ns::init_dispatch_vector, TypeMapFactory>( \ + __name__##_output_typeid_vector); \ + }; + /** * @brief A macro used to define factories and a populating binary universal * functions. diff --git a/dpnp/backend/extensions/vm/CMakeLists.txt b/dpnp/backend/extensions/vm/CMakeLists.txt index 2ea7e29ddd3f..8a463f857460 100644 --- a/dpnp/backend/extensions/vm/CMakeLists.txt +++ b/dpnp/backend/extensions/vm/CMakeLists.txt @@ -61,6 +61,7 @@ if(NOT _use_onemath) ${CMAKE_CURRENT_SOURCE_DIR}/log10.cpp ${CMAKE_CURRENT_SOURCE_DIR}/log1p.cpp ${CMAKE_CURRENT_SOURCE_DIR}/log2.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/modf.cpp ${CMAKE_CURRENT_SOURCE_DIR}/mul.cpp ${CMAKE_CURRENT_SOURCE_DIR}/nextafter.cpp ${CMAKE_CURRENT_SOURCE_DIR}/pow.cpp diff --git a/dpnp/backend/extensions/vm/common.hpp b/dpnp/backend/extensions/vm/common.hpp index 0dd9ff9209cf..c5f88491a72a 100644 --- a/dpnp/backend/extensions/vm/common.hpp +++ b/dpnp/backend/extensions/vm/common.hpp @@ -155,6 +155,115 @@ bool need_to_call_unary_ufunc(sycl::queue &exec_q, return true; } +template +bool need_to_call_unary_two_outputs_ufunc( + sycl::queue &exec_q, + const dpctl::tensor::usm_ndarray &src, + const dpctl::tensor::usm_ndarray &dst1, + const dpctl::tensor::usm_ndarray &dst2, + const output_typesT &output_type_vec, + const contig_dispatchT &contig_dispatch_vector) +{ + // check type_nums + int src_typenum = src.get_typenum(); + int dst1_typenum = dst1.get_typenum(); + int dst2_typenum = dst2.get_typenum(); + + const auto &array_types = td_ns::usm_ndarray_types(); + int src_typeid = array_types.typenum_to_lookup_id(src_typenum); + int dst1_typeid = array_types.typenum_to_lookup_id(dst1_typenum); + int dst2_typeid = array_types.typenum_to_lookup_id(dst2_typenum); + + std::pair func_output_typeids = output_type_vec[src_typeid]; + + // check that types are supported + if (dst1_typeid != func_output_typeids.first || + dst2_typeid != func_output_typeids.second) + { + return false; + } + + // OneMKL VM functions perform a copy on host if no double type support + if (!exec_q.get_device().has(sycl::aspect::fp64)) { + return false; + } + + // check that queues are compatible + if (!dpctl::utils::queues_are_compatible(exec_q, {src, dst1, dst2})) { + return false; + } + + // dimensions must be the same + int src_nd = src.get_ndim(); + int dst1_nd = dst1.get_ndim(); + int dst2_nd = dst2.get_ndim(); + if (src_nd != dst1_nd || src_nd != dst2_nd) { + return false; + } + else if (dst1_nd == 0 || dst2_nd == 0) { + // don't call OneMKL for 0d arrays + return false; + } + + // shapes must be the same + const py::ssize_t *src_shape = src.get_shape_raw(); + const py::ssize_t *dst1_shape = dst1.get_shape_raw(); + const py::ssize_t *dst2_shape = dst2.get_shape_raw(); + bool shapes_equal(true); + size_t src_nelems(1); + + for (int i = 0; i < src_nd; ++i) { + src_nelems *= static_cast(src_shape[i]); + shapes_equal = shapes_equal && (src_shape[i] == dst1_shape[i]) && + (src_shape[i] == dst2_shape[i]); + } + if (!shapes_equal) { + return false; + } + + // if nelems is zero, return false + if (src_nelems == 0) { + return false; + } + + // ensure that outputs are ample enough to accommodate all elements + auto dst1_offsets = dst1.get_minmax_offsets(); + auto dst2_offsets = dst2.get_minmax_offsets(); + // destinations must be ample enough to accommodate all elements + { + size_t range1 = + static_cast(dst1_offsets.second - dst1_offsets.first); + size_t range2 = + static_cast(dst2_offsets.second - dst2_offsets.first); + if ((range1 + 1 < src_nelems) || (range2 + 1 < src_nelems)) { + return false; + } + } + + // check memory overlap + auto const &overlap = dpctl::tensor::overlap::MemoryOverlap(); + if (overlap(src, dst1) || overlap(src, dst2) || overlap(dst1, dst2)) { + return false; + } + + // support only contiguous inputs + bool is_src_c_contig = src.is_c_contiguous(); + bool is_dst1_c_contig = dst1.is_c_contiguous(); + bool is_dst2_c_contig = dst2.is_c_contiguous(); + + bool all_c_contig = + (is_src_c_contig && is_dst1_c_contig && is_dst2_c_contig); + if (!all_c_contig) { + return false; + } + + // MKL function is not defined for the type + if (contig_dispatch_vector[src_typeid] == nullptr) { + return false; + } + return true; +} + template bool need_to_call_binary_ufunc(sycl::queue &exec_q, const dpctl::tensor::usm_ndarray &src1, @@ -299,6 +408,54 @@ bool need_to_call_binary_ufunc(sycl::queue &exec_q, ContigFactory>(contig_dispatch_vector); \ }; +/** + * @brief A macro used to define factories and a populating unary functions + * with two output arrays to dispatch to a callback with proper OneMKL function + * within VM extension scope. + */ +#define MACRO_POPULATE_DISPATCH_2OUTS_VECTORS(__name__) \ + template \ + struct ContigFactory \ + { \ + fnT get() \ + { \ + if constexpr (std::is_same_v::value_type1, \ + void> || \ + std::is_same_v::value_type2, \ + void>) \ + { \ + fnT fn = nullptr; \ + return fn; \ + } \ + else { \ + fnT fn = __name__##_contig_impl; \ + return fn; \ + } \ + } \ + }; \ + \ + template \ + struct TypeMapFactory \ + { \ + std::enable_if_t>::value, \ + std::pair> \ + get() \ + { \ + using rT1 = typename OutputType::value_type1; \ + using rT2 = typename OutputType::value_type2; \ + return std::make_pair(td_ns::GetTypeid{}.get(), \ + td_ns::GetTypeid{}.get()); \ + } \ + }; \ + \ + static void populate_dispatch_vectors(void) \ + { \ + ext_ns::init_dispatch_vector, TypeMapFactory>( \ + output_typeid_vector); \ + ext_ns::init_dispatch_vector(contig_dispatch_vector); \ + }; + /** * @brief A macro used to define factories and a populating binary functions * to dispatch to a callback with proper OneMKL function within VM extension diff --git a/dpnp/backend/extensions/vm/modf.cpp b/dpnp/backend/extensions/vm/modf.cpp new file mode 100644 index 000000000000..ef68c79d8b42 --- /dev/null +++ b/dpnp/backend/extensions/vm/modf.cpp @@ -0,0 +1,160 @@ +//***************************************************************************** +// Copyright (c) 2025, 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. +// - Neither the name of the copyright holder nor the names of its contributors +// may be used to endorse or promote products derived from this software +// without specific prior written permission. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE +// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE +// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR +// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF +// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF +// THE POSSIBILITY OF SUCH DAMAGE. +//***************************************************************************** + +#include +#include +#include +#include +#include + +#include +#include + +#include "dpctl4pybind11.hpp" + +#include "common.hpp" +#include "modf.hpp" + +// include a local copy of elementwise common header from dpctl tensor: +// dpctl/tensor/libtensor/source/elementwise_functions/elementwise_functions.hpp +// TODO: replace by including dpctl header once available +#include "../elementwise_functions/elementwise_functions.hpp" + +#include "../elementwise_functions/common.hpp" +#include "../elementwise_functions/type_dispatch_building.hpp" + +// dpctl tensor headers +#include "utils/type_dispatch.hpp" +#include "utils/type_utils.hpp" + +namespace dpnp::extensions::vm +{ +namespace py = pybind11; +namespace py_int = dpnp::extensions::py_internal; +namespace td_ns = dpctl::tensor::type_dispatch; + +namespace impl +{ +namespace ew_cmn_ns = dpnp::extensions::py_internal::elementwise_common; +namespace mkl_vm = oneapi::mkl::vm; // OneMKL namespace with VM functions +namespace td_int_ns = py_int::type_dispatch; +namespace tu_ns = dpctl::tensor::type_utils; + +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::modf function. + * + * @tparam T Type of input vector `a` and of result vectors `y` and `z`. + */ +template +struct OutputType +{ + using table_type = + std::disjunction, + td_int_ns::TypeMapTwoResultsEntry, + td_int_ns::TypeMapTwoResultsEntry, + td_int_ns::DefaultTwoResultsEntry>; + using value_type1 = typename table_type::result_type1; + using value_type2 = typename table_type::result_type2; +}; + +template +static sycl::event modf_contig_impl(sycl::queue &exec_q, + std::size_t in_n, + const char *in_a, + char *out_y, + char *out_z, + const std::vector &depends) +{ + tu_ns::validate_type_for_device(exec_q); + + std::int64_t n = static_cast(in_n); + const T *a = reinterpret_cast(in_a); + + using fractT = typename OutputType::value_type1; + using intT = typename OutputType::value_type2; + fractT *y = reinterpret_cast(out_y); + intT *z = reinterpret_cast(out_z); + + return mkl_vm::modf(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + z, // pointer `z` to the output truncated integer values + y, // pointer `y` to the output remaining fraction parts + depends); +} + +using ew_cmn_ns::unary_two_outputs_contig_impl_fn_ptr_t; +using ew_cmn_ns::unary_two_outputs_strided_impl_fn_ptr_t; + +static std::pair output_typeid_vector[td_ns::num_types]; +static unary_two_outputs_contig_impl_fn_ptr_t + contig_dispatch_vector[td_ns::num_types]; + +MACRO_POPULATE_DISPATCH_2OUTS_VECTORS(modf); +} // namespace impl + +void init_modf(py::module_ m) +{ + using arrayT = dpctl::tensor::usm_ndarray; + using event_vecT = std::vector; + + impl::populate_dispatch_vectors(); + using impl::contig_dispatch_vector; + using impl::output_typeid_vector; + + auto modf_pyapi = [&](sycl::queue &exec_q, const arrayT &src, + const arrayT &dst1, const arrayT &dst2, + const event_vecT &depends = {}) { + return py_int::py_unary_two_outputs_ufunc( + src, dst1, dst2, exec_q, depends, output_typeid_vector, + contig_dispatch_vector, + // no support of strided implementation in OneMKL + td_ns::NullPtrVector< + impl::unary_two_outputs_strided_impl_fn_ptr_t>{}); + }; + m.def("_modf", modf_pyapi, + "Call `modf` function from OneMKL VM library to compute " + "a truncated integer value and the remaining fraction part " + "for each vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst1"), + py::arg("dst2"), py::arg("depends") = py::list()); + + auto modf_need_to_call_pyapi = [&](sycl::queue &exec_q, const arrayT &src, + const arrayT &dst1, const arrayT &dst2) { + return py_internal::need_to_call_unary_two_outputs_ufunc( + exec_q, src, dst1, dst2, output_typeid_vector, + contig_dispatch_vector); + }; + m.def("_mkl_modf_to_call", modf_need_to_call_pyapi, + "Check input arguments to answer if `modf` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst1"), + py::arg("dst2")); +} +} // namespace dpnp::extensions::vm diff --git a/dpnp/backend/extensions/vm/modf.hpp b/dpnp/backend/extensions/vm/modf.hpp new file mode 100644 index 000000000000..885d9b277738 --- /dev/null +++ b/dpnp/backend/extensions/vm/modf.hpp @@ -0,0 +1,38 @@ +//***************************************************************************** +// Copyright (c) 2025, 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. +// - Neither the name of the copyright holder nor the names of its contributors +// may be used to endorse or promote products derived from this software +// without specific prior written permission. +// +// 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. +//***************************************************************************** + +#pragma once + +#include + +namespace py = pybind11; + +namespace dpnp::extensions::vm +{ +void init_modf(py::module_ m); +} // namespace dpnp::extensions::vm diff --git a/dpnp/backend/extensions/vm/vm_py.cpp b/dpnp/backend/extensions/vm/vm_py.cpp index bcd0434c84a1..985f0bbe25c0 100644 --- a/dpnp/backend/extensions/vm/vm_py.cpp +++ b/dpnp/backend/extensions/vm/vm_py.cpp @@ -63,6 +63,7 @@ #include "log10.hpp" #include "log1p.hpp" #include "log2.hpp" +#include "modf.hpp" #include "mul.hpp" #include "nextafter.hpp" #include "pow.hpp" @@ -116,6 +117,7 @@ PYBIND11_MODULE(_vm_impl, m) vm_ns::init_log10(m); vm_ns::init_log1p(m); vm_ns::init_log2(m); + vm_ns::init_modf(m); vm_ns::init_mul(m); vm_ns::init_nextafter(m); vm_ns::init_pow(m); diff --git a/dpnp/backend/include/dpnp_iface.hpp b/dpnp/backend/include/dpnp_iface.hpp index 7e080659d49f..b7627101aab4 100644 --- a/dpnp/backend/include/dpnp_iface.hpp +++ b/dpnp/backend/include/dpnp_iface.hpp @@ -156,32 +156,6 @@ INP_DLLEXPORT DPCTLSyclEventRef template INP_DLLEXPORT void dpnp_initval_c(void *result1, void *value, size_t size); -/** - * @ingroup BACKEND_API - * @brief modf function. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array1_in Input array. - * @param [out] result1_out Output array 1. - * @param [out] result2_out Output array 2. - * @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_modf_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result1_out, - void *result2_out, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_modf_c(void *array1_in, - void *result1_out, - void *result2_out, - size_t size); - /** * @ingroup BACKEND_API * @brief Implementation of ones function diff --git a/dpnp/backend/include/dpnp_iface_fptr.hpp b/dpnp/backend/include/dpnp_iface_fptr.hpp index 617ef163e55e..868a31bdaac5 100644 --- a/dpnp/backend/include/dpnp_iface_fptr.hpp +++ b/dpnp/backend/include/dpnp_iface_fptr.hpp @@ -64,14 +64,11 @@ enum class DPNPFuncName : size_t DPNP_FN_NONE, /**< Very first element of the enumeration */ 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_MODF, /**< Used in numpy.modf() impl */ - DPNP_FN_MODF_EXT, /**< Used in numpy.modf() impl, requires extra parameters - */ - DPNP_FN_ONES, /**< Used in numpy.ones() impl */ - DPNP_FN_ONES_LIKE, /**< Used in numpy.ones_like() impl */ - DPNP_FN_PARTITION, /**< Used in numpy.partition() impl */ + DPNP_FN_INITVAL_EXT, /**< Used in numpy ones, ones_like, zeros, zeros_like + impls */ + DPNP_FN_ONES, /**< Used in numpy.ones() impl */ + DPNP_FN_ONES_LIKE, /**< Used in numpy.ones_like() impl */ + DPNP_FN_PARTITION, /**< Used in numpy.partition() impl */ DPNP_FN_PARTITION_EXT, /**< Used in numpy.partition() impl, requires extra parameters */ DPNP_FN_RNG_BETA, /**< Used in numpy.random.beta() impl */ diff --git a/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp b/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp deleted file mode 100644 index c7101ac9ba67..000000000000 --- a/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp +++ /dev/null @@ -1,160 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016, 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. -// - Neither the name of the copyright holder nor the names of its contributors -// may be used to endorse or promote products derived from this software -// without specific prior written permission. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -#include -#include -#include - -#include - -#include "dpnp_fptr.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; - -static_assert(__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_VECTOR_ABS_CHANGED, - "SYCL DPC++ compiler does not meet minimum version requirement"); - -template -class dpnp_modf_c_kernel; - -template -DPCTLSyclEventRef dpnp_modf_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - void *result1_out, - void *result2_out, - 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; - - DPNPC_ptr_adapter<_DataType_input> input1_ptr(q_ref, array1_in, size); - _DataType_input *array1 = input1_ptr.get_ptr(); - _DataType_output *result1 = - reinterpret_cast<_DataType_output *>(result1_out); - _DataType_output *result2 = - reinterpret_cast<_DataType_output *>(result2_out); - - if constexpr (std::is_same<_DataType_input, double>::value || - std::is_same<_DataType_input, float>::value) - { - event = oneapi::mkl::vm::modf(q, size, array1, result2, result1); - } - else { - sycl::range<1> gws(size); - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { - size_t i = global_id[0]; /*for (size_t i = 0; i < size; ++i)*/ - { - double input_elem1 = static_cast(array1[i]); - auto res_multi_ptr = sycl::address_space_cast< - sycl::access::address_space::global_space, - sycl::access::decorated::yes>(&result1[i]); - - result2[i] = sycl::modf(input_elem1, res_multi_ptr); - } - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.parallel_for< - class dpnp_modf_c_kernel<_DataType_input, _DataType_output>>( - gws, kernel_parallel_for_func); - }; - - event = q.submit(kernel_func); - } - - event_ref = reinterpret_cast(&event); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_modf_c(void *array1_in, - void *result1_out, - void *result2_out, - size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_modf_c<_DataType_input, _DataType_output>(q_ref, array1_in, - result1_out, result2_out, - size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_modf_default_c)(void *, void *, void *, size_t) = - dpnp_modf_c<_DataType_input, _DataType_output>; - -template -DPCTLSyclEventRef (*dpnp_modf_ext_c)(DPCTLSyclQueueRef, - void *, - void *, - void *, - size_t, - const DPCTLEventVectorRef) = - dpnp_modf_c<_DataType_input, _DataType_output>; - -void func_map_init_mathematical(func_map_t &fmap) -{ - - 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] = { - eft_DBL, (void *)dpnp_modf_default_c}; - fmap[DPNPFuncName::DPNP_FN_MODF][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_modf_default_c}; - fmap[DPNPFuncName::DPNP_FN_MODF][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_modf_default_c}; - - fmap[DPNPFuncName::DPNP_FN_MODF_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_modf_ext_c}; - fmap[DPNPFuncName::DPNP_FN_MODF_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_modf_ext_c}; - fmap[DPNPFuncName::DPNP_FN_MODF_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_modf_ext_c}; - fmap[DPNPFuncName::DPNP_FN_MODF_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_modf_ext_c}; - - return; -} diff --git a/dpnp/backend/kernels/elementwise_functions/modf.hpp b/dpnp/backend/kernels/elementwise_functions/modf.hpp new file mode 100644 index 000000000000..07a8ef65e55e --- /dev/null +++ b/dpnp/backend/kernels/elementwise_functions/modf.hpp @@ -0,0 +1,53 @@ +//***************************************************************************** +// Copyright (c) 2025, 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. +// - Neither the name of the copyright holder nor the names of its contributors +// may be used to endorse or promote products derived from this software +// without specific prior written permission. +// +// 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. +//***************************************************************************** + +#pragma once + +#include + +namespace dpnp::kernels::modf +{ +template +struct ModfFunctor +{ + // is function constant for given argT + using is_constant = typename std::false_type; + // constant value, if constant + // constexpr resT1 constant_value1 = resT1{}; + // constexpr resT2 constant_value2 = resT2{}; + // is function defined for sycl::vec + using supports_vec = typename std::false_type; + // do both argT and fractT, intT support subgroup store/load operation + using supports_sg_loadstore = typename std::true_type; + + fractT operator()(const argT &in, intT &integral) const + { + return sycl::modf(in, &integral); + } +}; +} // namespace dpnp::kernels::modf diff --git a/dpnp/backend/src/dpnp_fptr.hpp b/dpnp/backend/src/dpnp_fptr.hpp index 397dd84d1068..9fd2dedb4a2d 100644 --- a/dpnp/backend/src/dpnp_fptr.hpp +++ b/dpnp/backend/src/dpnp_fptr.hpp @@ -161,7 +161,6 @@ class dpnp_less_comp */ void func_map_init_arraycreation(func_map_t &fmap); void func_map_init_linalg(func_map_t &fmap); -void func_map_init_mathematical(func_map_t &fmap); void func_map_init_random(func_map_t &fmap); void func_map_init_sorting(func_map_t &fmap); diff --git a/dpnp/backend/src/dpnp_iface_fptr.cpp b/dpnp/backend/src/dpnp_iface_fptr.cpp index 190e62dcf8ea..ab1769c7e56c 100644 --- a/dpnp/backend/src/dpnp_iface_fptr.cpp +++ b/dpnp/backend/src/dpnp_iface_fptr.cpp @@ -100,7 +100,6 @@ static func_map_t func_map_init() func_map_init_arraycreation(fmap); func_map_init_linalg(fmap); - func_map_init_mathematical(fmap); func_map_init_random(fmap); func_map_init_sorting(fmap); diff --git a/dpnp/backend/src/dpnp_utils.hpp b/dpnp/backend/src/dpnp_utils.hpp index 8c7a71b05b45..03e56e0716aa 100644 --- a/dpnp/backend/src/dpnp_utils.hpp +++ b/dpnp/backend/src/dpnp_utils.hpp @@ -41,14 +41,6 @@ #include -/** - * Version of SYCL DPC++ 2023 compiler where a return type of sycl::abs() is - * changed from unsigned integer to signed one of input vector. - */ -#ifndef __SYCL_COMPILER_VECTOR_ABS_CHANGED -#define __SYCL_COMPILER_VECTOR_ABS_CHANGED 20230503L -#endif - /** * Version of Intel MKL at which transition to OneMKL release 2023.0.0 occurs. */ diff --git a/dpnp/dpnp_algo/CMakeLists.txt b/dpnp/dpnp_algo/CMakeLists.txt index 359d72b7a7ba..cb2a89cc4a19 100644 --- a/dpnp/dpnp_algo/CMakeLists.txt +++ b/dpnp/dpnp_algo/CMakeLists.txt @@ -29,7 +29,6 @@ set(dpnp_algo_pyx_deps ${CMAKE_CURRENT_SOURCE_DIR}/dpnp_algo_sorting.pxi - ${CMAKE_CURRENT_SOURCE_DIR}/dpnp_algo_mathematical.pxi ${CMAKE_CURRENT_SOURCE_DIR}/dpnp_algo_indexing.pxi ) diff --git a/dpnp/dpnp_algo/dpnp_algo.pxd b/dpnp/dpnp_algo/dpnp_algo.pxd index 22b0686b3c6a..1e755770893c 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pxd +++ b/dpnp/dpnp_algo/dpnp_algo.pxd @@ -31,7 +31,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this namespace for Enum import cdef enum DPNPFuncName "DPNPFuncName": - DPNP_FN_MODF_EXT DPNP_FN_PARTITION_EXT DPNP_FN_RNG_BETA_EXT DPNP_FN_RNG_BINOMIAL_EXT diff --git a/dpnp/dpnp_algo/dpnp_algo.pyx b/dpnp/dpnp_algo/dpnp_algo.pyx index 1ecd3cf0646f..53abcad11986 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pyx +++ b/dpnp/dpnp_algo/dpnp_algo.pyx @@ -49,7 +49,6 @@ __all__ = [ include "dpnp_algo_indexing.pxi" -include "dpnp_algo_mathematical.pxi" include "dpnp_algo_sorting.pxi" diff --git a/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi b/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi deleted file mode 100644 index fbe388b6df10..000000000000 --- a/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi +++ /dev/null @@ -1,92 +0,0 @@ -# cython: language_level=3 -# cython: linetrace=True -# -*- coding: utf-8 -*- -# ***************************************************************************** -# Copyright (c) 2016, 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. -# - Neither the name of the copyright holder nor the names of its contributors -# may be used to endorse or promote products derived from this software -# without specific prior written permission. -# -# 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 (Mathematical 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_modf", -] - - -ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_1in_2out_t)(c_dpctl.DPCTLSyclQueueRef, - void * , void * , void * , size_t, - const c_dpctl.DPCTLEventVectorRef) - - -cpdef tuple dpnp_modf(utils.dpnp_descriptor x1): - """ Convert string type names (array.dtype) to C enum DPNPFuncType """ - cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(x1.dtype) - - """ get the FPTR data structure """ - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_MODF_EXT, param1_type, DPNP_FT_NONE) - - x1_obj = x1.get_array() - - # create result array with type given by FPTR data - cdef shape_type_c result_shape = x1.shape - cdef utils.dpnp_descriptor result1 = utils.create_output_descriptor(result_shape, - kernel_data.return_type, - None, - device=x1_obj.sycl_device, - usm_type=x1_obj.usm_type, - sycl_queue=x1_obj.sycl_queue) - cdef utils.dpnp_descriptor result2 = utils.create_output_descriptor(result_shape, - kernel_data.return_type, - None, - device=x1_obj.sycl_device, - usm_type=x1_obj.usm_type, - sycl_queue=x1_obj.sycl_queue) - - _, _, result_sycl_queue = utils.get_common_usm_allocation(result1, result2) - - cdef c_dpctl.SyclQueue q = result_sycl_queue - cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() - - cdef fptr_1in_2out_t func = kernel_data.ptr - """ Call FPTR function """ - cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, - x1.get_data(), - result1.get_data(), - result2.get_data(), - x1.size, - NULL) # dep_events_ref - - with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) - c_dpctl.DPCTLEvent_Delete(event_ref) - - return (result1.get_pyobj(), result2.get_pyobj()) diff --git a/dpnp/dpnp_algo/dpnp_elementwise_common.py b/dpnp/dpnp_algo/dpnp_elementwise_common.py index 5bb5cc715d8e..7490d4e03e43 100644 --- a/dpnp/dpnp_algo/dpnp_elementwise_common.py +++ b/dpnp/dpnp_algo/dpnp_elementwise_common.py @@ -92,12 +92,12 @@ class DPNPUnaryFunc(UnaryElementwiseFunc): corresponds to computational tasks associated with function evaluation. docs : {str} Documentation string for the unary function. - mkl_fn_to_call : {callable} + mkl_fn_to_call : {None, str} Check input arguments to answer if function from OneMKL VM library can be used. - mkl_impl_fn : {callable} + mkl_impl_fn : {None, str} Function from OneMKL VM library to call. - acceptance_fn : {callable}, optional + acceptance_fn : {None, callable}, optional Function to influence type promotion behavior of this unary function. The function takes 4 arguments: arg_dtype - Data type of the first argument @@ -130,7 +130,9 @@ def _call_func(src, dst, sycl_queue, depends=None): if depends is None: depends = [] - if vmi._is_available() and mkl_fn_to_call is not None: + if vmi._is_available() and not ( + mkl_impl_fn is None or mkl_fn_to_call is None + ): if getattr(vmi, mkl_fn_to_call)(sycl_queue, src, dst): # call pybind11 extension for unary function from OneMKL VM return getattr(vmi, mkl_impl_fn)( @@ -231,6 +233,11 @@ class DPNPUnaryTwoOutputsFunc(UnaryElementwiseFunc): corresponds to computational tasks associated with function evaluation. docs : {str} Documentation string for the unary function. + mkl_fn_to_call : {None, str} + Check input arguments to answer if function from OneMKL VM library + can be used. + mkl_impl_fn : {None, str} + Function from OneMKL VM library to call. """ @@ -240,11 +247,29 @@ def __init__( result_type_resolver_fn, unary_dp_impl_fn, docs, + mkl_fn_to_call=None, + mkl_impl_fn=None, ): + def _call_func(src, dst1, dst2, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class.""" + + if depends is None: + depends = [] + + if vmi._is_available() and not ( + mkl_impl_fn is None or mkl_fn_to_call is None + ): + if getattr(vmi, mkl_fn_to_call)(sycl_queue, src, dst1, dst2): + # call pybind11 extension for unary function from OneMKL VM + return getattr(vmi, mkl_impl_fn)( + sycl_queue, src, dst1, dst2, depends + ) + return unary_dp_impl_fn(src, dst1, dst2, sycl_queue, depends) + super().__init__( name, result_type_resolver_fn, - unary_dp_impl_fn, + _call_func, docs, ) self.__name__ = "DPNPUnaryTwoOutputsFunc" @@ -428,7 +453,7 @@ def __call__( if not isinstance(res, dpnp_array): # Always return dpnp.ndarray out[i] = dpnp_array._create_from_usm_ndarray(res) - return out + return tuple(out) class DPNPBinaryFunc(BinaryElementwiseFunc): @@ -458,12 +483,12 @@ class DPNPBinaryFunc(BinaryElementwiseFunc): evaluation. docs : {str} Documentation string for the unary function. - mkl_fn_to_call : {callable} + mkl_fn_to_call : {None, str} Check input arguments to answer if function from OneMKL VM library can be used. - mkl_impl_fn : {callable} + mkl_impl_fn : {None, str} Function from OneMKL VM library to call. - binary_inplace_fn : {callable}, optional + binary_inplace_fn : {None, callable}, optional Data-parallel implementation function with signature `impl_fn(src: usm_ndarray, dst: usm_ndarray, sycl_queue: SyclQueue, depends: Optional[List[SyclEvent]])` @@ -475,7 +500,7 @@ class DPNPBinaryFunc(BinaryElementwiseFunc): including async lifetime management of Python arguments, while the second event corresponds to computational tasks associated with function evaluation. - acceptance_fn : {callable}, optional + acceptance_fn : {None, callable}, optional Function to influence type promotion behavior of this binary function. The function takes 6 arguments: arg1_dtype - Data type of the first argument @@ -488,7 +513,7 @@ class DPNPBinaryFunc(BinaryElementwiseFunc): The function is only called when both arguments of the binary function require casting, e.g. both arguments of `dpctl.tensor.logaddexp` are arrays with integral data type. - weak_type_resolver : {callable}, optional + weak_type_resolver : {None, callable}, optional Function to influence type promotion behavior for Python scalar types of this binary function. The function takes 3 arguments: o1_dtype - Data type or Python scalar type of the first argument @@ -520,7 +545,9 @@ def _call_func(src1, src2, dst, sycl_queue, depends=None): if depends is None: depends = [] - if vmi._is_available() and mkl_fn_to_call is not None: + if vmi._is_available() and not ( + mkl_impl_fn is None or mkl_fn_to_call is None + ): if getattr(vmi, mkl_fn_to_call)(sycl_queue, src1, src2, dst): # call pybind11 extension for binary function from OneMKL VM return getattr(vmi, mkl_impl_fn)( diff --git a/dpnp/dpnp_iface_mathematical.py b/dpnp/dpnp_iface_mathematical.py index bd7e83ae3a8a..023e3ce96e07 100644 --- a/dpnp/dpnp_iface_mathematical.py +++ b/dpnp/dpnp_iface_mathematical.py @@ -62,7 +62,6 @@ import dpnp import dpnp.backend.extensions.ufunc._ufunc_impl as ufi -from .dpnp_algo import dpnp_modf from .dpnp_algo.dpnp_elementwise_common import ( DPNPI0, DPNPAngle, @@ -82,7 +81,7 @@ resolve_weak_types_2nd_arg_int, ) from .dpnp_array import dpnp_array -from .dpnp_utils import call_origin, get_usm_allocations +from .dpnp_utils import get_usm_allocations from .dpnp_utils.dpnp_utils_linearalgebra import dpnp_cross from .dpnp_utils.dpnp_utils_reduction import dpnp_wrap_reduction_call @@ -2366,7 +2365,6 @@ def ediff1d(ary, to_end=None, to_begin=None): as `x` and the expected data type. Default: ``None``. - out : tuple of None, dpnp.ndarray, or usm_ndarray, optional A location into which the result is stored. If provided, it must be a tuple and have length equal to the number of outputs. Each provided array must @@ -3352,42 +3350,84 @@ def interp(x, xp, fp, left=None, right=None, period=None): ) -def modf(x1, **kwargs): - """ - Return the fractional and integral parts of an array, element-wise. +_MODF_DOCSTRING = """ +Decompose each element :math:`x_i` of the input array `x` into the fractional +and the integral parts. - For full documentation refer to :obj:`numpy.modf`. +The fractional and integral parts are negative if the given :math:`x_i` is +negative. - Limitations - ----------- - Parameter `x` is supported as :obj:`dpnp.ndarray`. - Keyword argument `kwargs` is currently unsupported. - Otherwise the function will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. +For full documentation refer to :obj:`numpy.modf`. - Examples - -------- - >>> import dpnp as np - >>> a = np.array([1, 2]) - >>> result = np.modf(a) - >>> [[x for x in y] for y in result ] - [[1.0, 2.0], [0.0, 0.0]] +Parameters +---------- +x : {dpnp.ndarray, usm_ndarray} + Array of numbers to be decomposed, expected to have a real-valued + floating-point data type. +out1 : {None, dpnp.ndarray, usm_ndarray}, optional + Output array for the fractional parts to populate. Array must have the same + shape as `x` and the expected data type. - """ + Default: ``None``. +out2 : {None, dpnp.ndarray, usm_ndarray}, optional + Output array for the integral parts to populate. Array must have the same + shape as `x` and the expected data type. - x1_desc = dpnp.get_dpnp_descriptor(x1, copy_when_nondefault_queue=False) - if x1_desc: - if dpnp.is_cuda_backend(x1_desc.get_array()): # pragma: no cover - raise NotImplementedError( - "Running on CUDA is currently not supported" - ) + Default: ``None``. +out : tuple of None, dpnp.ndarray, or usm_ndarray, optional + A location into which the result is stored. If provided, it must be a tuple + and have length equal to the number of outputs. Each provided array must + have the same shape as `x` and the expected data type. + It is prohibited to pass output arrays through `out` keyword when either + `out1` or `out2` is passed. - if kwargs: - pass - else: - return dpnp_modf(x1_desc) + Default: ``(None, None)``. +order : {None, "C", "F", "A", "K"}, optional + Memory layout of the newly output array, if parameter `out` is ``None``. + + Default: ``"K"``. + +Returns +------- +y1 : dpnp.ndarray + Fractional part of `x`. +y2 : dpnp.ndarray + Integral part of `x`. + +Limitations +----------- +Parameters `where`, `dtype` and `subok` are supported with their default values. +Keyword argument `kwargs` is currently unsupported. +Otherwise ``NotImplementedError`` exception will be raised. - return call_origin(numpy.modf, x1, **kwargs) +See Also +-------- +:obj:`dpnp.divmod` : ``divmod(x, 1)`` is an equivalent to ``modf(x)`` with the + return values switched, except it always has a positive remainder. + +Notes +----- +For integer input the return values are floats. + +Examples +-------- +>>> import dpnp as np +>>> x = np.array([0, 3.5]) +>>> np.modf(x) +(array([0. , 0.5]), array([0., 3.])) +>>> np.modf(np.array(-0.5)) +(array(-0.5), array(-0.)) + +""" + +modf = DPNPUnaryTwoOutputsFunc( + "_modf", + ufi._modf_result_type, + ufi._modf, + _MODF_DOCSTRING, + mkl_fn_to_call="_mkl_modf_to_call", + mkl_impl_fn="_modf", +) _MULTIPLY_DOCSTRING = """ diff --git a/dpnp/tests/skipped_tests_cuda.tbl b/dpnp/tests/skipped_tests_cuda.tbl index 2846a3c2f5c7..f035fba3302f 100644 --- a/dpnp/tests/skipped_tests_cuda.tbl +++ b/dpnp/tests/skipped_tests_cuda.tbl @@ -1,11 +1,5 @@ # NotImplementedError -# modf -tests/test_arithmetic.py::TestArithmetic::test_modf_part1 -tests/test_arithmetic.py::TestArithmetic::test_modf_part2 -tests/test_sycl_queue.py::test_modf[cuda:gpu:0] -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticModf::test_modf - # random tests/test_random_state.py::TestNormal::test_distr[host-float32] tests/test_random_state.py::TestNormal::test_distr[host-float64] diff --git a/dpnp/tests/test_arithmetic.py b/dpnp/tests/test_arithmetic.py deleted file mode 100644 index 2d1741b9d229..000000000000 --- a/dpnp/tests/test_arithmetic.py +++ /dev/null @@ -1,38 +0,0 @@ -import unittest - -from .third_party.cupy import testing - - -class TestArithmetic(unittest.TestCase): - @testing.for_float_dtypes() - @testing.numpy_cupy_allclose() - def test_modf_part1(self, xp, dtype): - a = xp.array([-2.5, -1.5, -0.5, 0, 0.5, 1.5, 2.5], dtype=dtype) - b, _ = xp.modf(a) - return b - - @testing.for_float_dtypes() - @testing.numpy_cupy_allclose() - def test_modf_part2(self, xp, dtype): - a = xp.array([-2.5, -1.5, -0.5, 0, 0.5, 1.5, 2.5], dtype=dtype) - _, c = xp.modf(a) - return c - - @testing.for_float_dtypes() - @testing.numpy_cupy_allclose(type_check=False) - def test_nanprod(self, xp, dtype): - a = xp.array([-2.5, -1.5, xp.nan, 10.5, 1.5, xp.nan], dtype=dtype) - return xp.nanprod(a) - - @testing.for_float_dtypes() - @testing.numpy_cupy_allclose() - def test_nansum(self, xp, dtype): - a = xp.array([-2.5, -1.5, xp.nan, 10.5, 1.5, xp.nan], dtype=dtype) - return xp.nansum(a) - - @testing.for_float_dtypes() - @testing.numpy_cupy_allclose() - def test_remainder(self, xp, dtype): - a = xp.array([5, -3, -2, -1, -5], dtype=dtype) - b = xp.full(a.size, 3, dtype=dtype) - return xp.remainder(a, b) diff --git a/dpnp/tests/test_binary_ufuncs.py b/dpnp/tests/test_binary_ufuncs.py index 782541ab41f9..ab1146201b4d 100644 --- a/dpnp/tests/test_binary_ufuncs.py +++ b/dpnp/tests/test_binary_ufuncs.py @@ -27,7 +27,7 @@ """ The scope includes tests with only functions which are instances of -`DPNPUnaryFunc` class. +`DPNPBinaryFunc` class. """ diff --git a/dpnp/tests/test_mathematical.py b/dpnp/tests/test_mathematical.py index 8d4cdb426ae8..5a1626360a57 100644 --- a/dpnp/tests/test_mathematical.py +++ b/dpnp/tests/test_mathematical.py @@ -33,7 +33,6 @@ has_support_aspect16, has_support_aspect64, is_intel_numpy, - is_win_platform, numpy_version, ) from .third_party.cupy import testing @@ -711,133 +710,6 @@ def test_errors(self): assert_raises(ExecutionPlacementError, dpnp.ediff1d, ia, to_end=to_end) -class TestFrexp: - ALL_DTYPES = get_all_dtypes(no_none=True) - ALL_DTYPES_NO_COMPLEX = get_all_dtypes( - no_none=True, no_float16=False, no_complex=True - ) - ALL_FLOAT_DTYPES = get_float_dtypes(no_float16=False) - - @pytest.mark.parametrize("dt", ALL_DTYPES_NO_COMPLEX) - def test_basic(self, dt): - a = get_abs_array([-2, 5, 1, 4, 3], dtype=dt) - ia = dpnp.array(a) - - res1, res2 = dpnp.frexp(ia) - exp1, exp2 = numpy.frexp(a) - assert_array_equal(res1, exp1) - assert_array_equal(res2, exp2) - - @pytest.mark.parametrize("dt", ALL_FLOAT_DTYPES) - def test_out(self, dt): - a = numpy.array(5.7, dtype=dt) - ia = dpnp.array(a) - - out1 = numpy.empty((), dtype=dt) - out2 = numpy.empty((), dtype=numpy.int32) - iout1, iout2 = dpnp.array(out1), dpnp.array(out2) - - res1, res2 = dpnp.frexp(ia, iout1) - exp1, exp2 = numpy.frexp(a, out1) - assert_array_equal(res1, exp1) - assert_array_equal(res2, exp2) - assert res1 is iout1 - - res1, res2 = dpnp.frexp(ia, None, iout2) - exp1, exp2 = numpy.frexp(a, None, out2) - assert_array_equal(res1, exp1) - assert_array_equal(res2, exp2) - assert res2 is iout2 - - res1, res2 = dpnp.frexp(ia, iout1, iout2) - exp1, exp2 = numpy.frexp(a, out1, out2) - assert_array_equal(res1, exp1) - assert_array_equal(res2, exp2) - assert res1 is iout1 - assert res2 is iout2 - - @pytest.mark.parametrize("dt", ALL_DTYPES_NO_COMPLEX) - @pytest.mark.parametrize("out1_dt", ALL_DTYPES) - @pytest.mark.parametrize("out2_dt", ALL_DTYPES) - def test_out_all_dtypes(self, dt, out1_dt, out2_dt): - a = numpy.ones(9, dtype=dt) - ia = dpnp.array(a) - - out1 = numpy.zeros(9, dtype=out1_dt) - out2 = numpy.zeros(9, dtype=out2_dt) - iout1, iout2 = dpnp.array(out1), dpnp.array(out2) - - try: - res1, res2 = dpnp.frexp(ia, out=(iout1, iout2)) - except TypeError: - # expect numpy to fail with the same reason - with pytest.raises(TypeError): - _ = numpy.frexp(a, out=(out1, out2)) - else: - exp1, exp2 = numpy.frexp(a, out=(out1, out2)) - assert_array_equal(res1, exp1) - assert_array_equal(res2, exp2) - assert res1 is iout1 - assert res2 is iout2 - - @pytest.mark.skipif( - is_win_platform(), - reason="numpy.frexp gives different answers for NAN/INF on Windows and Linux", - ) - @pytest.mark.parametrize("stride", [-4, -2, -1, 1, 2, 4]) - @pytest.mark.parametrize("dt", ALL_FLOAT_DTYPES) - def test_strides_out(self, stride, dt): - a = numpy.array( - [numpy.nan, numpy.nan, numpy.inf, -numpy.inf, 0.0, -0.0, 1.0, -1.0], - dtype=dt, - ) - ia = dpnp.array(a) - - out_mant = numpy.ones_like(a) - out_exp = 2 * numpy.ones_like(a, dtype="i") - iout_mant, iout_exp = dpnp.array(out_mant), dpnp.array(out_exp) - - res1, res2 = dpnp.frexp( - ia[::stride], out=(iout_mant[::stride], iout_exp[::stride]) - ) - exp1, exp2 = numpy.frexp( - a[::stride], out=(out_mant[::stride], out_exp[::stride]) - ) - assert_array_equal(res1, exp1) - assert_array_equal(res2, exp2) - - assert_array_equal(iout_mant, out_mant) - assert_array_equal(iout_exp, out_exp) - - @pytest.mark.parametrize("dt", ALL_FLOAT_DTYPES) - def test_out1_overlap(self, dt): - size = 15 - a = numpy.ones(2 * size, dtype=dt) - ia = dpnp.array(a) - - # out1 overlaps memory of input array - _ = dpnp.frexp(ia[size::], ia[::2]) - _ = numpy.frexp(a[size::], a[::2]) - assert_array_equal(ia, a) - - @pytest.mark.parametrize("dt", ALL_FLOAT_DTYPES) - def test_empty(self, dt): - a = numpy.empty(0, dtype=dt) - ia = dpnp.array(a) - - res1, res2 = dpnp.frexp(ia) - exp1, exp2 = numpy.frexp(a) - assert_array_equal(res1, exp1, strict=True) - assert_array_equal(res2, exp2, strict=True) - - @pytest.mark.parametrize("xp", [numpy, dpnp]) - @pytest.mark.parametrize("dt", get_complex_dtypes()) - def test_complex_dtype(self, xp, dt): - a = xp.array([-2, 5, 1, 4, 3], dtype=dt) - with pytest.raises((TypeError, ValueError)): - _ = xp.frexp(a) - - class TestGradient: @pytest.mark.parametrize("dt", get_all_dtypes(no_none=True, no_bool=True)) def test_basic(self, dt): diff --git a/dpnp/tests/test_sycl_queue.py b/dpnp/tests/test_sycl_queue.py index 434ec48c3436..dfd090dbc5d1 100644 --- a/dpnp/tests/test_sycl_queue.py +++ b/dpnp/tests/test_sycl_queue.py @@ -360,6 +360,7 @@ def test_1in_1out(func, data, device): "func, data", [ pytest.param("frexp", numpy.arange(9)), + pytest.param("modf", [0, 3.5]), ], ) @pytest.mark.parametrize("device", valid_dev, ids=dev_ids) @@ -850,16 +851,6 @@ def test_random_state(func, args, kwargs, device, usm_type): assert_sycl_queue_equal(res_array.sycl_queue, sycl_queue) -@pytest.mark.parametrize("device", valid_dev, ids=dev_ids) -def test_modf(device): - x = dpnp.array([0, 3.5], device=device) - result1, result2 = dpnp.modf(x) - - expected_queue = x.sycl_queue - assert_sycl_queue_equal(result1.sycl_queue, expected_queue) - assert_sycl_queue_equal(result2.sycl_queue, expected_queue) - - @pytest.mark.parametrize("device", valid_dev, ids=dev_ids) def test_einsum(device): array_list = [] diff --git a/dpnp/tests/test_umath.py b/dpnp/tests/test_umath.py index 31b55a204a10..2e272327b8a6 100644 --- a/dpnp/tests/test_umath.py +++ b/dpnp/tests/test_umath.py @@ -117,11 +117,6 @@ def test_umaths(test_cases): pytest.skip("dpctl-2031") elif umath in ["divmod"]: pytest.skip("Not implemented umath") - elif umath == "modf": - if args[0].dtype == dpnp.float16: - pytest.skip("dpnp.modf is not supported with dpnp.float16") - elif is_cuda_device(): - pytest.skip("dpnp.modf is not supported on CUDA device") elif umath in ["vecmat", "matvec"]: if is_win_platform() and not is_gpu_device(): pytest.skip("SAT-8073") diff --git a/dpnp/tests/test_unary_two_outputs_ufuncs.py b/dpnp/tests/test_unary_two_outputs_ufuncs.py new file mode 100644 index 000000000000..006b846eb236 --- /dev/null +++ b/dpnp/tests/test_unary_two_outputs_ufuncs.py @@ -0,0 +1,157 @@ +import numpy +import pytest +from numpy.testing import ( + assert_array_equal, +) + +import dpnp + +from .helper import ( + generate_random_numpy_array, + get_all_dtypes, + get_complex_dtypes, + get_float_dtypes, + is_win_platform, +) + +""" +The scope includes tests with only functions which are instances of +`DPNPUnaryTwoOutputsFunc` class. + +""" + + +@pytest.mark.parametrize("func", ["frexp", "modf"]) +class TestUnaryTwoOutputs: + ALL_DTYPES = get_all_dtypes(no_none=True) + ALL_DTYPES_NO_COMPLEX = get_all_dtypes( + no_none=True, no_float16=False, no_complex=True + ) + ALL_FLOAT_DTYPES = get_float_dtypes(no_float16=False) + + def _get_out_dtypes(self, func, dt): + if func == "frexp": + return (dt, numpy.int32) + return (dt, dt) + + @pytest.mark.parametrize("dt", ALL_DTYPES_NO_COMPLEX) + def test_basic(self, func, dt): + a = generate_random_numpy_array((2, 5), dtype=dt) + ia = dpnp.array(a) + + res1, res2 = getattr(dpnp, func)(ia) + exp1, exp2 = getattr(numpy, func)(a) + assert_array_equal(res1, exp1) + assert_array_equal(res2, exp2) + + @pytest.mark.parametrize("dt", ALL_FLOAT_DTYPES) + def test_out(self, func, dt): + a = numpy.array(5.7, dtype=dt) + ia = dpnp.array(a) + + dt1, dt2 = self._get_out_dtypes(func, dt) + out1 = numpy.empty((), dtype=dt1) + out2 = numpy.empty((), dtype=dt2) + iout1, iout2 = dpnp.array(out1), dpnp.array(out2) + + res1, res2 = getattr(dpnp, func)(ia, iout1) + exp1, exp2 = getattr(numpy, func)(a, out1) + assert_array_equal(res1, exp1) + assert_array_equal(res2, exp2) + assert res1 is iout1 + + res1, res2 = getattr(dpnp, func)(ia, None, iout2) + exp1, exp2 = getattr(numpy, func)(a, None, out2) + assert_array_equal(res1, exp1) + assert_array_equal(res2, exp2) + assert res2 is iout2 + + res1, res2 = getattr(dpnp, func)(ia, iout1, iout2) + exp1, exp2 = getattr(numpy, func)(a, out1, out2) + assert_array_equal(res1, exp1) + assert_array_equal(res2, exp2) + assert res1 is iout1 + assert res2 is iout2 + + @pytest.mark.parametrize("dt", ALL_DTYPES_NO_COMPLEX) + @pytest.mark.parametrize("out1_dt", ALL_DTYPES) + @pytest.mark.parametrize("out2_dt", ALL_DTYPES) + def test_out_all_dtypes(self, func, dt, out1_dt, out2_dt): + a = numpy.ones(9, dtype=dt) + ia = dpnp.array(a) + + out1 = numpy.zeros(9, dtype=out1_dt) + out2 = numpy.zeros(9, dtype=out2_dt) + iout1, iout2 = dpnp.array(out1), dpnp.array(out2) + + try: + res1, res2 = getattr(dpnp, func)(ia, out=(iout1, iout2)) + except TypeError: + # expect numpy to fail with the same reason + with pytest.raises(TypeError): + _ = getattr(numpy, func)(a, out=(out1, out2)) + else: + exp1, exp2 = getattr(numpy, func)(a, out=(out1, out2)) + assert_array_equal(res1, exp1) + assert_array_equal(res2, exp2) + assert res1 is iout1 + assert res2 is iout2 + + @pytest.mark.parametrize("stride", [-4, -2, -1, 1, 2, 4]) + @pytest.mark.parametrize("dt", ALL_FLOAT_DTYPES) + def test_strides_out(self, func, stride, dt): + if func == "frexp" and is_win_platform(): + pytest.skip( + "numpy.frexp gives different answers for NAN/INF on Windows and Linux" + ) + + a = numpy.array( + [numpy.nan, numpy.nan, numpy.inf, -numpy.inf, 0.0, -0.0, 1.0, -1.0], + dtype=dt, + ) + ia = dpnp.array(a) + + dt1, dt2 = self._get_out_dtypes(func, dt) + out_mant = numpy.ones_like(a, dtype=dt1) + out_exp = 2 * numpy.ones_like(a, dtype=dt2) + iout_mant, iout_exp = dpnp.array(out_mant), dpnp.array(out_exp) + + res1, res2 = getattr(dpnp, func)( + ia[::stride], out=(iout_mant[::stride], iout_exp[::stride]) + ) + exp1, exp2 = getattr(numpy, func)( + a[::stride], out=(out_mant[::stride], out_exp[::stride]) + ) + assert_array_equal(res1, exp1) + assert_array_equal(res2, exp2) + + assert_array_equal(iout_mant, out_mant) + assert_array_equal(iout_exp, out_exp) + + @pytest.mark.parametrize("dt", ALL_FLOAT_DTYPES) + def test_out1_overlap(self, func, dt): + size = 15 + a = numpy.ones(2 * size, dtype=dt) + ia = dpnp.array(a) + + # out1 overlaps memory of input array + _ = getattr(dpnp, func)(ia[size::], ia[::2]) + _ = getattr(numpy, func)(a[size::], a[::2]) + assert_array_equal(ia, a) + + @pytest.mark.parametrize("dt", ALL_FLOAT_DTYPES) + def test_empty(self, func, dt): + a = numpy.empty(0, dtype=dt) + ia = dpnp.array(a) + + res1, res2 = getattr(dpnp, func)(ia) + exp1, exp2 = getattr(numpy, func)(a) + assert_array_equal(res1, exp1, strict=True) + assert_array_equal(res2, exp2, strict=True) + + @pytest.mark.parametrize("xp", [numpy, dpnp]) + @pytest.mark.parametrize("dt", get_complex_dtypes()) + def test_complex_dtype(self, func, xp, dt): + a = xp.array([-2, 5, 1, 4, 3], dtype=dt) + with pytest.raises((TypeError, ValueError)): + _ = getattr(xp, func)(a) diff --git a/dpnp/tests/test_usm_type.py b/dpnp/tests/test_usm_type.py index 582b371e2f6c..90292bd1ce1f 100644 --- a/dpnp/tests/test_usm_type.py +++ b/dpnp/tests/test_usm_type.py @@ -653,6 +653,7 @@ def test_1in_1out(func, data, usm_type): "func, data", [ pytest.param("frexp", numpy.arange(9)), + pytest.param("modf", [0, 3.5]), ], ) @pytest.mark.parametrize("usm_type", list_of_usm_types) diff --git a/dpnp/tests/third_party/cupy/math_tests/test_arithmetic.py b/dpnp/tests/third_party/cupy/math_tests/test_arithmetic.py index 1e233bddada1..02ebae462398 100644 --- a/dpnp/tests/third_party/cupy/math_tests/test_arithmetic.py +++ b/dpnp/tests/third_party/cupy/math_tests/test_arithmetic.py @@ -1,3 +1,5 @@ +from __future__ import annotations + import itertools import warnings