diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index 9d88d04b5330..0e89bfb8bfa2 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -113,7 +113,7 @@ repos: "--disable=redefined-builtin", "--disable=unused-wildcard-import" ] - files: '^dpnp/(dpnp_iface.*|fft|linalg|dpnp_array)' + files: '^dpnp/(dpnp_iface.*|fft|linalg|special|dpnp_array)' - repo: https://github.com/macisamuele/language-formatters-pre-commit-hooks rev: v2.15.0 hooks: diff --git a/CHANGELOG.md b/CHANGELOG.md index 9ed9491d0ecc..06c7769433a7 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -34,6 +34,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 * Replaced `ci` section in `.pre-commit-config.yaml` with a new GitHub workflow with scheduled run to autoupdate the `pre-commit` configuration [#2542](https://github.com/IntelPython/dpnp/pull/2542) * FFT module is updated to perform in-place FFT in intermediate steps of ND FFT [#2543](https://github.com/IntelPython/dpnp/pull/2543) * Reused dpctl tensor include to enable experimental SYCL namespace for complex types [#2546](https://github.com/IntelPython/dpnp/pull/2546) +* Redesigned `dpnp.erf` function through pybind11 extension of OneMKL call or dedicated kernel in `ufunc` namespace [#2551](https://github.com/IntelPython/dpnp/pull/2551) ### Deprecated diff --git a/conda-recipe/meta.yaml b/conda-recipe/meta.yaml index 6f75fb4cbaef..8be4b53e60e5 100644 --- a/conda-recipe/meta.yaml +++ b/conda-recipe/meta.yaml @@ -63,6 +63,7 @@ test: requires: - pytest - setuptools + - scipy about: home: https://github.com/IntelPython/dpnp diff --git a/doc/reference/special.rst b/doc/reference/special.rst index 1261625e5e26..5e6735a6afbe 100644 --- a/doc/reference/special.rst +++ b/doc/reference/special.rst @@ -1,17 +1,20 @@ -Special Functions -================= +.. currentmodule:: dpnp.special -.. https://docs.scipy.org/doc/scipy/reference/special.html +Special functions (:mod:`dpnp.special`) +======================================= -Error Function --------------- +.. Hint:: `SciPy API Reference: Special functions (scipy.special) `_ + +Error function and Fresnel integrals +------------------------------------ .. autosummary:: :toctree: generated/ :nosignatures: - dpnp.erf - dpnp.erfc - dpnp.erfcx - dpnp.erfinv - dpnp.erfcinv + erf + erfc + erfcx + erfi + erfinv + erfcinv diff --git a/dpnp/backend/CMakeLists.txt b/dpnp/backend/CMakeLists.txt index 4a6b44d50f0d..17b986a38b92 100644 --- a/dpnp/backend/CMakeLists.txt +++ b/dpnp/backend/CMakeLists.txt @@ -26,7 +26,6 @@ set(DPNP_SRC kernels/dpnp_krnl_arraycreation.cpp kernels/dpnp_krnl_common.cpp - kernels/dpnp_krnl_elemwise.cpp kernels/dpnp_krnl_mathematical.cpp kernels/dpnp_krnl_random.cpp kernels/dpnp_krnl_sorting.cpp diff --git a/dpnp/backend/extensions/ufunc/CMakeLists.txt b/dpnp/backend/extensions/ufunc/CMakeLists.txt index 65d30d49f549..94e8077ecbee 100644 --- a/dpnp/backend/extensions/ufunc/CMakeLists.txt +++ b/dpnp/backend/extensions/ufunc/CMakeLists.txt @@ -27,6 +27,7 @@ set(_elementwise_sources ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/bitwise_count.cpp ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/common.cpp ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/degrees.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/erf.cpp ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fabs.cpp ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fix.cpp ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/float_power.cpp diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/bitwise_count.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/bitwise_count.cpp index 2fe60d2a5efe..e3d67d1ab67a 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/bitwise_count.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/bitwise_count.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include "dpctl4pybind11.hpp" diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp index c6dd3e038eb1..64feca21952b 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp @@ -27,6 +27,7 @@ #include "bitwise_count.hpp" #include "degrees.hpp" +#include "erf.hpp" #include "fabs.hpp" #include "fix.hpp" #include "float_power.hpp" @@ -56,6 +57,7 @@ void init_elementwise_functions(py::module_ m) { init_bitwise_count(m); init_degrees(m); + init_erf(m); init_fabs(m); init_fix(m); init_float_power(m); diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/degrees.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/degrees.cpp index 420a7ebf1331..9bf17aa1594a 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/degrees.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/degrees.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include "dpctl4pybind11.hpp" diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/erf.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/erf.cpp new file mode 100644 index 000000000000..567f20ac445a --- /dev/null +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/erf.cpp @@ -0,0 +1,127 @@ +//***************************************************************************** +// 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. +// +// 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 "dpctl4pybind11.hpp" + +#include "erf.hpp" +#include "kernels/elementwise_functions/erf.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" + +// 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 = dpctl::tensor::kernels::elementwise_common; +namespace td_ns = dpctl::tensor::type_dispatch; + +/** + * @brief A factory to define pairs of supported types for which + * sycl::erf function is available. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct OutputType +{ + using value_type = + typename std::disjunction, + td_ns::TypeMapResultEntry, + td_ns::DefaultResultEntry>::result_type; +}; + +using dpnp::kernels::erf::ErfFunctor; + +template +using ContigFunctor = ew_cmn_ns::UnaryContigFunctor, + vec_sz, + n_vecs, + enable_sg_loadstore>; + +template +using StridedFunctor = ew_cmn_ns:: + UnaryStridedFunctor>; + +using ew_cmn_ns::unary_contig_impl_fn_ptr_t; +using ew_cmn_ns::unary_strided_impl_fn_ptr_t; + +static unary_contig_impl_fn_ptr_t erf_contig_dispatch_vector[td_ns::num_types]; +static int erf_output_typeid_vector[td_ns::num_types]; +static unary_strided_impl_fn_ptr_t + erf_strided_dispatch_vector[td_ns::num_types]; + +MACRO_POPULATE_DISPATCH_VECTORS(erf); +} // namespace impl + +void init_erf(py::module_ m) +{ + using arrayT = dpctl::tensor::usm_ndarray; + using event_vecT = std::vector; + { + impl::populate_erf_dispatch_vectors(); + using impl::erf_contig_dispatch_vector; + using impl::erf_output_typeid_vector; + using impl::erf_strided_dispatch_vector; + + auto erf_pyapi = [&](const arrayT &src, const arrayT &dst, + sycl::queue &exec_q, + const event_vecT &depends = {}) { + return py_int::py_unary_ufunc( + src, dst, exec_q, depends, erf_output_typeid_vector, + erf_contig_dispatch_vector, erf_strided_dispatch_vector); + }; + m.def("_erf", erf_pyapi, "", py::arg("src"), py::arg("dst"), + py::arg("sycl_queue"), py::arg("depends") = py::list()); + + auto erf_result_type_pyapi = [&](const py::dtype &dtype) { + return py_int::py_unary_ufunc_result_type(dtype, + erf_output_typeid_vector); + }; + m.def("_erf_result_type", erf_result_type_pyapi); + } +} +} // namespace dpnp::extensions::ufunc diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/erf.hpp b/dpnp/backend/extensions/ufunc/elementwise_functions/erf.hpp new file mode 100644 index 000000000000..f02f7f4a6448 --- /dev/null +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/erf.hpp @@ -0,0 +1,35 @@ +//***************************************************************************** +// 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. +// +// 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_erf(py::module_ m); +} // namespace dpnp::extensions::ufunc diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/fabs.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/fabs.cpp index 4d1b061f3f58..16c5c1bddd66 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/fabs.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/fabs.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include "dpctl4pybind11.hpp" diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/fix.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/fix.cpp index 9fefbac85d41..1fe99a02ead2 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/fix.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/fix.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include "dpctl4pybind11.hpp" diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/float_power.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/float_power.cpp index bf2ca7e6fa2d..3ce168940148 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/float_power.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/float_power.cpp @@ -23,6 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include +#include + #include #include "dpctl4pybind11.hpp" diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/fmax.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/fmax.cpp index 84a0c03a1f6e..063ad57c1b34 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/fmax.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/fmax.cpp @@ -23,6 +23,8 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include + #include #include "dpctl4pybind11.hpp" diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/fmin.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/fmin.cpp index 0211be04866f..8af827f1756c 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/fmin.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/fmin.cpp @@ -23,6 +23,8 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include + #include #include "dpctl4pybind11.hpp" diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/fmod.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/fmod.cpp index cb35d50c9f98..995d2662db41 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/fmod.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/fmod.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include "dpctl4pybind11.hpp" diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/gcd.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/gcd.cpp index d9d83e273b2b..24d01c2ea3f1 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/gcd.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/gcd.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include "dpctl4pybind11.hpp" diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/heaviside.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/heaviside.cpp index 273f8a1e5d07..624d3446c007 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/heaviside.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/heaviside.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include "dpctl4pybind11.hpp" diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/i0.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/i0.cpp index c74d67c4810f..ca11b78534f2 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/i0.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/i0.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include "dpctl4pybind11.hpp" diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/interpolate.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/interpolate.cpp index 248f9be6c675..15b110e088f8 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/interpolate.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/interpolate.cpp @@ -24,8 +24,14 @@ //***************************************************************************** #include +#include +#include +#include +#include #include +#include + #include "dpctl4pybind11.hpp" #include #include diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/lcm.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/lcm.cpp index 63481a8ded74..a74d04444a4e 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/lcm.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/lcm.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include "dpctl4pybind11.hpp" diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/ldexp.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/ldexp.cpp index c59b0b1c0ef8..3da7a53c5cfa 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/ldexp.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/ldexp.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include "dpctl4pybind11.hpp" diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/logaddexp2.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/logaddexp2.cpp index d2d2df691b42..191a01b00512 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/logaddexp2.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/logaddexp2.cpp @@ -23,6 +23,8 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include + #include #include "dpctl4pybind11.hpp" diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/radians.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/radians.cpp index b2f774f271ba..cf922dd735f8 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/radians.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/radians.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include "dpctl4pybind11.hpp" diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/sinc.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/sinc.cpp index 6a59f2d82c56..f3ce28108cf9 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/sinc.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/sinc.cpp @@ -23,6 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include +#include + #include #include "dpctl4pybind11.hpp" diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/spacing.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/spacing.cpp index d28c42476b2a..c99bd043c212 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/spacing.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/spacing.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include "dpctl4pybind11.hpp" diff --git a/dpnp/backend/extensions/vm/CMakeLists.txt b/dpnp/backend/extensions/vm/CMakeLists.txt index a2b7297e72b5..057311c146c9 100644 --- a/dpnp/backend/extensions/vm/CMakeLists.txt +++ b/dpnp/backend/extensions/vm/CMakeLists.txt @@ -42,6 +42,7 @@ if(NOT _use_onemath) ${CMAKE_CURRENT_SOURCE_DIR}/cos.cpp ${CMAKE_CURRENT_SOURCE_DIR}/cosh.cpp ${CMAKE_CURRENT_SOURCE_DIR}/div.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/erf.cpp ${CMAKE_CURRENT_SOURCE_DIR}/exp.cpp ${CMAKE_CURRENT_SOURCE_DIR}/exp2.cpp ${CMAKE_CURRENT_SOURCE_DIR}/expm1.cpp diff --git a/dpnp/backend/extensions/vm/abs.cpp b/dpnp/backend/extensions/vm/abs.cpp index 0842146c04e1..5ce568543135 100644 --- a/dpnp/backend/extensions/vm/abs.cpp +++ b/dpnp/backend/extensions/vm/abs.cpp @@ -23,6 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/acos.cpp b/dpnp/backend/extensions/vm/acos.cpp index e21005208c39..9a4eecd557f5 100644 --- a/dpnp/backend/extensions/vm/acos.cpp +++ b/dpnp/backend/extensions/vm/acos.cpp @@ -23,6 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/acosh.cpp b/dpnp/backend/extensions/vm/acosh.cpp index 169c6d7f8868..efd896c905b5 100644 --- a/dpnp/backend/extensions/vm/acosh.cpp +++ b/dpnp/backend/extensions/vm/acosh.cpp @@ -23,6 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/add.cpp b/dpnp/backend/extensions/vm/add.cpp index 3cb5ac0ec816..d0f10c213753 100644 --- a/dpnp/backend/extensions/vm/add.cpp +++ b/dpnp/backend/extensions/vm/add.cpp @@ -23,7 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include #include +#include +#include #include #include diff --git a/dpnp/backend/extensions/vm/arg.cpp b/dpnp/backend/extensions/vm/arg.cpp index 174f61c72e7c..19d0f8050287 100644 --- a/dpnp/backend/extensions/vm/arg.cpp +++ b/dpnp/backend/extensions/vm/arg.cpp @@ -23,6 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/asin.cpp b/dpnp/backend/extensions/vm/asin.cpp index 51b7a0b3a219..93d220573c44 100644 --- a/dpnp/backend/extensions/vm/asin.cpp +++ b/dpnp/backend/extensions/vm/asin.cpp @@ -23,6 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/asinh.cpp b/dpnp/backend/extensions/vm/asinh.cpp index f1fc016bd3bd..9f87290f7f77 100644 --- a/dpnp/backend/extensions/vm/asinh.cpp +++ b/dpnp/backend/extensions/vm/asinh.cpp @@ -23,6 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/atan.cpp b/dpnp/backend/extensions/vm/atan.cpp index 5b68bf896342..fa13217882c8 100644 --- a/dpnp/backend/extensions/vm/atan.cpp +++ b/dpnp/backend/extensions/vm/atan.cpp @@ -23,6 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/atan2.cpp b/dpnp/backend/extensions/vm/atan2.cpp index d4a76c6fcc25..40781215b928 100644 --- a/dpnp/backend/extensions/vm/atan2.cpp +++ b/dpnp/backend/extensions/vm/atan2.cpp @@ -24,6 +24,8 @@ //***************************************************************************** #include +#include +#include #include #include diff --git a/dpnp/backend/extensions/vm/atanh.cpp b/dpnp/backend/extensions/vm/atanh.cpp index 3ef18e01cc81..42ca2e813662 100644 --- a/dpnp/backend/extensions/vm/atanh.cpp +++ b/dpnp/backend/extensions/vm/atanh.cpp @@ -23,6 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/cbrt.cpp b/dpnp/backend/extensions/vm/cbrt.cpp index 074f1ef92587..42a028a83485 100644 --- a/dpnp/backend/extensions/vm/cbrt.cpp +++ b/dpnp/backend/extensions/vm/cbrt.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/ceil.cpp b/dpnp/backend/extensions/vm/ceil.cpp index 6e85da6bd731..730d4bec9d7f 100644 --- a/dpnp/backend/extensions/vm/ceil.cpp +++ b/dpnp/backend/extensions/vm/ceil.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/common.hpp b/dpnp/backend/extensions/vm/common.hpp index b58f55ede815..3a93d7b6e8e9 100644 --- a/dpnp/backend/extensions/vm/common.hpp +++ b/dpnp/backend/extensions/vm/common.hpp @@ -25,6 +25,8 @@ #pragma once +#include + #include #include diff --git a/dpnp/backend/extensions/vm/conj.cpp b/dpnp/backend/extensions/vm/conj.cpp index 34c265b046a1..f41a1fefa495 100644 --- a/dpnp/backend/extensions/vm/conj.cpp +++ b/dpnp/backend/extensions/vm/conj.cpp @@ -23,6 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/copysign.cpp b/dpnp/backend/extensions/vm/copysign.cpp index 7a703124f641..5dcf2c5ac96d 100644 --- a/dpnp/backend/extensions/vm/copysign.cpp +++ b/dpnp/backend/extensions/vm/copysign.cpp @@ -24,6 +24,8 @@ //***************************************************************************** #include +#include +#include #include #include diff --git a/dpnp/backend/extensions/vm/cos.cpp b/dpnp/backend/extensions/vm/cos.cpp index d0eaa531df84..8c6dcab23e39 100644 --- a/dpnp/backend/extensions/vm/cos.cpp +++ b/dpnp/backend/extensions/vm/cos.cpp @@ -23,6 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/cosh.cpp b/dpnp/backend/extensions/vm/cosh.cpp index 11b671d76bea..40ca40f8c7f6 100644 --- a/dpnp/backend/extensions/vm/cosh.cpp +++ b/dpnp/backend/extensions/vm/cosh.cpp @@ -23,6 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/div.cpp b/dpnp/backend/extensions/vm/div.cpp index ebad4fadf02a..4c6201c3b775 100644 --- a/dpnp/backend/extensions/vm/div.cpp +++ b/dpnp/backend/extensions/vm/div.cpp @@ -23,7 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include #include +#include +#include #include #include diff --git a/dpnp/backend/extensions/vm/erf.cpp b/dpnp/backend/extensions/vm/erf.cpp new file mode 100644 index 000000000000..ec2f0248c604 --- /dev/null +++ b/dpnp/backend/extensions/vm/erf.cpp @@ -0,0 +1,138 @@ +//***************************************************************************** +// 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. +// +// 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 "dpctl4pybind11.hpp" + +#include "common.hpp" +#include "erf.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" + +// dpctl tensor headers +#include "kernels/elementwise_functions/common.hpp" +#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 = dpctl::tensor::kernels::elementwise_common; +namespace mkl_vm = oneapi::mkl::vm; // OneMKL namespace with VM functions +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::erf function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct OutputType +{ + using value_type = + typename std::disjunction, + td_ns::TypeMapResultEntry, + td_ns::DefaultResultEntry>::result_type; +}; + +template +static sycl::event erf_contig_impl(sycl::queue &exec_q, + std::size_t in_n, + const char *in_a, + char *out_y, + 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 resTy = typename OutputType::value_type; + resTy *y = reinterpret_cast(out_y); + + return mkl_vm::erf(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +using ew_cmn_ns::unary_contig_impl_fn_ptr_t; +using ew_cmn_ns::unary_strided_impl_fn_ptr_t; + +static int output_typeid_vector[td_ns::num_types]; +static unary_contig_impl_fn_ptr_t contig_dispatch_vector[td_ns::num_types]; + +MACRO_POPULATE_DISPATCH_VECTORS(erf); +} // namespace impl + +void init_erf(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 erf_pyapi = [&](sycl::queue &exec_q, const arrayT &src, + const arrayT &dst, const event_vecT &depends = {}) { + return py_int::py_unary_ufunc( + src, dst, exec_q, depends, output_typeid_vector, + contig_dispatch_vector, + // no support of strided implementation in OneMKL + td_ns::NullPtrVector{}); + }; + m.def("_erf", erf_pyapi, + "Call `erf` function from OneMKL VM library to compute " + "the natural (base-e) erfonential of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto erf_need_to_call_pyapi = [&](sycl::queue &exec_q, const arrayT &src, + const arrayT &dst) { + return py_internal::need_to_call_unary_ufunc( + exec_q, src, dst, output_typeid_vector, contig_dispatch_vector); + }; + m.def("_mkl_erf_to_call", erf_need_to_call_pyapi, + "Check input arguments to answer if `erf` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); +} +} // namespace dpnp::extensions::vm diff --git a/dpnp/backend/extensions/vm/erf.hpp b/dpnp/backend/extensions/vm/erf.hpp new file mode 100644 index 000000000000..25eb2b05a7cc --- /dev/null +++ b/dpnp/backend/extensions/vm/erf.hpp @@ -0,0 +1,35 @@ +//***************************************************************************** +// 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. +// +// 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_erf(py::module_ m); +} // namespace dpnp::extensions::vm diff --git a/dpnp/backend/extensions/vm/exp.cpp b/dpnp/backend/extensions/vm/exp.cpp index 83ab13300b72..bcac06905107 100644 --- a/dpnp/backend/extensions/vm/exp.cpp +++ b/dpnp/backend/extensions/vm/exp.cpp @@ -23,6 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/exp2.cpp b/dpnp/backend/extensions/vm/exp2.cpp index 9342e2c13565..d917d8469409 100644 --- a/dpnp/backend/extensions/vm/exp2.cpp +++ b/dpnp/backend/extensions/vm/exp2.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/expm1.cpp b/dpnp/backend/extensions/vm/expm1.cpp index 394456969e09..657d4ed2198b 100644 --- a/dpnp/backend/extensions/vm/expm1.cpp +++ b/dpnp/backend/extensions/vm/expm1.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/floor.cpp b/dpnp/backend/extensions/vm/floor.cpp index 87c55c2d7843..54b3754d8201 100644 --- a/dpnp/backend/extensions/vm/floor.cpp +++ b/dpnp/backend/extensions/vm/floor.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/fmax.cpp b/dpnp/backend/extensions/vm/fmax.cpp index 688a6c005c48..1700f0fda892 100644 --- a/dpnp/backend/extensions/vm/fmax.cpp +++ b/dpnp/backend/extensions/vm/fmax.cpp @@ -24,6 +24,8 @@ //***************************************************************************** #include +#include +#include #include #include diff --git a/dpnp/backend/extensions/vm/fmin.cpp b/dpnp/backend/extensions/vm/fmin.cpp index 4239667923be..ae9bab591462 100644 --- a/dpnp/backend/extensions/vm/fmin.cpp +++ b/dpnp/backend/extensions/vm/fmin.cpp @@ -24,6 +24,8 @@ //***************************************************************************** #include +#include +#include #include #include diff --git a/dpnp/backend/extensions/vm/fmod.cpp b/dpnp/backend/extensions/vm/fmod.cpp index 30361b6c8afd..e5fd67215312 100644 --- a/dpnp/backend/extensions/vm/fmod.cpp +++ b/dpnp/backend/extensions/vm/fmod.cpp @@ -24,6 +24,8 @@ //***************************************************************************** #include +#include +#include #include #include diff --git a/dpnp/backend/extensions/vm/hypot.cpp b/dpnp/backend/extensions/vm/hypot.cpp index 282257c5ac75..39339e00e31d 100644 --- a/dpnp/backend/extensions/vm/hypot.cpp +++ b/dpnp/backend/extensions/vm/hypot.cpp @@ -24,6 +24,8 @@ //***************************************************************************** #include +#include +#include #include #include diff --git a/dpnp/backend/extensions/vm/i0.cpp b/dpnp/backend/extensions/vm/i0.cpp index f9d910d85543..9d9ebd0a1a42 100644 --- a/dpnp/backend/extensions/vm/i0.cpp +++ b/dpnp/backend/extensions/vm/i0.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/inv.cpp b/dpnp/backend/extensions/vm/inv.cpp index 45cb07d3a7c2..8f1ba8483289 100644 --- a/dpnp/backend/extensions/vm/inv.cpp +++ b/dpnp/backend/extensions/vm/inv.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/ln.cpp b/dpnp/backend/extensions/vm/ln.cpp index e31250344542..63c44ca01be5 100644 --- a/dpnp/backend/extensions/vm/ln.cpp +++ b/dpnp/backend/extensions/vm/ln.cpp @@ -23,6 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/log10.cpp b/dpnp/backend/extensions/vm/log10.cpp index cb108c09bed6..b9ff909d259a 100644 --- a/dpnp/backend/extensions/vm/log10.cpp +++ b/dpnp/backend/extensions/vm/log10.cpp @@ -23,6 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/log1p.cpp b/dpnp/backend/extensions/vm/log1p.cpp index e9bac5bd6854..509fe7c99c61 100644 --- a/dpnp/backend/extensions/vm/log1p.cpp +++ b/dpnp/backend/extensions/vm/log1p.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/log2.cpp b/dpnp/backend/extensions/vm/log2.cpp index d78054047e85..94d208773996 100644 --- a/dpnp/backend/extensions/vm/log2.cpp +++ b/dpnp/backend/extensions/vm/log2.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/mul.cpp b/dpnp/backend/extensions/vm/mul.cpp index ba3e1df835e0..0913df55de4c 100644 --- a/dpnp/backend/extensions/vm/mul.cpp +++ b/dpnp/backend/extensions/vm/mul.cpp @@ -23,7 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include #include +#include +#include #include #include diff --git a/dpnp/backend/extensions/vm/nextafter.cpp b/dpnp/backend/extensions/vm/nextafter.cpp index 7a06d6aff15e..c4703b95c760 100644 --- a/dpnp/backend/extensions/vm/nextafter.cpp +++ b/dpnp/backend/extensions/vm/nextafter.cpp @@ -24,6 +24,8 @@ //***************************************************************************** #include +#include +#include #include #include diff --git a/dpnp/backend/extensions/vm/pow.cpp b/dpnp/backend/extensions/vm/pow.cpp index 8274a203c608..3a8da4f5d452 100644 --- a/dpnp/backend/extensions/vm/pow.cpp +++ b/dpnp/backend/extensions/vm/pow.cpp @@ -23,7 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include #include +#include +#include #include #include diff --git a/dpnp/backend/extensions/vm/rint.cpp b/dpnp/backend/extensions/vm/rint.cpp index 6d101807b393..d1035cc13b04 100644 --- a/dpnp/backend/extensions/vm/rint.cpp +++ b/dpnp/backend/extensions/vm/rint.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/sin.cpp b/dpnp/backend/extensions/vm/sin.cpp index 446acd67d3b7..06c89c221634 100644 --- a/dpnp/backend/extensions/vm/sin.cpp +++ b/dpnp/backend/extensions/vm/sin.cpp @@ -23,6 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/sinh.cpp b/dpnp/backend/extensions/vm/sinh.cpp index 77c5548abbf2..50506947e88a 100644 --- a/dpnp/backend/extensions/vm/sinh.cpp +++ b/dpnp/backend/extensions/vm/sinh.cpp @@ -23,6 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/sqr.cpp b/dpnp/backend/extensions/vm/sqr.cpp index efd78f4d89c0..5eb06318c6e7 100644 --- a/dpnp/backend/extensions/vm/sqr.cpp +++ b/dpnp/backend/extensions/vm/sqr.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/sqrt.cpp b/dpnp/backend/extensions/vm/sqrt.cpp index a3b42e80c255..1797d2ae936e 100644 --- a/dpnp/backend/extensions/vm/sqrt.cpp +++ b/dpnp/backend/extensions/vm/sqrt.cpp @@ -23,6 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/sub.cpp b/dpnp/backend/extensions/vm/sub.cpp index dcdba3b21a69..1a435a06a6b8 100644 --- a/dpnp/backend/extensions/vm/sub.cpp +++ b/dpnp/backend/extensions/vm/sub.cpp @@ -23,7 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include #include +#include +#include #include #include diff --git a/dpnp/backend/extensions/vm/tan.cpp b/dpnp/backend/extensions/vm/tan.cpp index 5dd9e906aa8e..dd9ad001f168 100644 --- a/dpnp/backend/extensions/vm/tan.cpp +++ b/dpnp/backend/extensions/vm/tan.cpp @@ -23,6 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/tanh.cpp b/dpnp/backend/extensions/vm/tanh.cpp index 68fcdcc375b9..5d72edb7856e 100644 --- a/dpnp/backend/extensions/vm/tanh.cpp +++ b/dpnp/backend/extensions/vm/tanh.cpp @@ -23,6 +23,10 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/trunc.cpp b/dpnp/backend/extensions/vm/trunc.cpp index 9a8984007593..d5c36fd83b9f 100644 --- a/dpnp/backend/extensions/vm/trunc.cpp +++ b/dpnp/backend/extensions/vm/trunc.cpp @@ -23,6 +23,9 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** +#include +#include + #include #include diff --git a/dpnp/backend/extensions/vm/vm_py.cpp b/dpnp/backend/extensions/vm/vm_py.cpp index ab2734794bdc..33015179b73c 100644 --- a/dpnp/backend/extensions/vm/vm_py.cpp +++ b/dpnp/backend/extensions/vm/vm_py.cpp @@ -45,6 +45,7 @@ #include "cos.hpp" #include "cosh.hpp" #include "div.hpp" +#include "erf.hpp" #include "exp.hpp" #include "exp2.hpp" #include "expm1.hpp" @@ -97,6 +98,7 @@ PYBIND11_MODULE(_vm_impl, m) vm_ns::init_cos(m); vm_ns::init_cosh(m); vm_ns::init_div(m); + vm_ns::init_erf(m); vm_ns::init_exp(m); vm_ns::init_exp2(m); vm_ns::init_expm1(m); diff --git a/dpnp/backend/include/dpnp_iface.hpp b/dpnp/backend/include/dpnp_iface.hpp index 5a0698a978af..5e892faa70c8 100644 --- a/dpnp/backend/include/dpnp_iface.hpp +++ b/dpnp/backend/include/dpnp_iface.hpp @@ -153,28 +153,6 @@ INP_DLLEXPORT DPCTLSyclEventRef template INP_DLLEXPORT void dpnp_initval_c(void *result1, void *value, size_t size); -#define MACRO_1ARG_1TYPE_OP(__name__, __operation1__, __operation2__) \ - template \ - INP_DLLEXPORT DPCTLSyclEventRef __name__( \ - DPCTLSyclQueueRef q_ref, void *result_out, const size_t result_size, \ - const size_t result_ndim, const shape_elem_type *result_shape, \ - const shape_elem_type *result_strides, const void *input1_in, \ - const size_t input1_size, const size_t input1_ndim, \ - const shape_elem_type *input1_shape, \ - const shape_elem_type *input1_strides, const size_t *where, \ - const DPCTLEventVectorRef dep_event_vec_ref); \ - \ - template \ - INP_DLLEXPORT void __name__( \ - void *result_out, const size_t result_size, const size_t result_ndim, \ - const shape_elem_type *result_shape, \ - const shape_elem_type *result_strides, const void *input1_in, \ - const size_t input1_size, const size_t input1_ndim, \ - const shape_elem_type *input1_shape, \ - const shape_elem_type *input1_strides, const size_t *where); - -#include - /** * @ingroup BACKEND_API * @brief modf function. diff --git a/dpnp/backend/include/dpnp_iface_fptr.hpp b/dpnp/backend/include/dpnp_iface_fptr.hpp index ea94d4a43b56..9f81077f6bf1 100644 --- a/dpnp/backend/include/dpnp_iface_fptr.hpp +++ b/dpnp/backend/include/dpnp_iface_fptr.hpp @@ -59,9 +59,6 @@ enum class DPNPFuncName : size_t { DPNP_FN_NONE, /**< Very first element of the enumeration */ - DPNP_FN_ERF, /**< Used in scipy.special.erf impl */ - DPNP_FN_ERF_EXT, /**< Used in scipy.special.erf impl, requires extra - parameters */ DPNP_FN_INITVAL, /**< Used in numpy ones, ones_like, zeros, zeros_like impls */ DPNP_FN_INITVAL_EXT, /**< Used in numpy ones, ones_like, zeros, zeros_like diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp deleted file mode 100644 index 2617ff29fa9c..000000000000 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ /dev/null @@ -1,255 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016-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. -// -// 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 "queue_sycl.hpp" - -template -constexpr T dispatch_erf_op(T elem) -{ - if constexpr (is_any_v) { - // TODO: need to convert to double when possible - return sycl::erf((float)elem); - } - else { - return sycl::erf(elem); - } -} - -#define MACRO_1ARG_1TYPE_OP(__name__, __operation1__, __operation2__) \ - template \ - class __name__##_kernel; \ - \ - template \ - class __name__##_strides_kernel; \ - \ - template \ - DPCTLSyclEventRef __name__( \ - DPCTLSyclQueueRef q_ref, void *result_out, const size_t result_size, \ - const size_t result_ndim, const shape_elem_type *result_shape, \ - const shape_elem_type *result_strides, const void *input1_in, \ - const size_t input1_size, const size_t input1_ndim, \ - const shape_elem_type *input1_shape, \ - const shape_elem_type *input1_strides, const size_t *where, \ - const DPCTLEventVectorRef dep_event_vec_ref) \ - { \ - /* avoid warning unused variable*/ \ - (void)result_shape; \ - (void)where; \ - (void)dep_event_vec_ref; \ - \ - DPCTLSyclEventRef event_ref = nullptr; \ - \ - if (!input1_size) { \ - return event_ref; \ - } \ - \ - sycl::queue q = *(reinterpret_cast(q_ref)); \ - \ - _DataType *input1_data = \ - static_cast<_DataType *>(const_cast(input1_in)); \ - _DataType *result = static_cast<_DataType *>(result_out); \ - \ - shape_elem_type *input1_shape_offsets = \ - new shape_elem_type[input1_ndim]; \ - \ - get_shape_offsets_inkernel(input1_shape, input1_ndim, \ - input1_shape_offsets); \ - bool use_strides = !array_equal(input1_strides, input1_ndim, \ - input1_shape_offsets, input1_ndim); \ - delete[] input1_shape_offsets; \ - \ - sycl::event event; \ - sycl::range<1> gws(result_size); \ - \ - if (use_strides) { \ - if (result_ndim != input1_ndim) { \ - throw std::runtime_error( \ - "Result ndim=" + std::to_string(result_ndim) + \ - " mismatches with input1 ndim=" + \ - std::to_string(input1_ndim)); \ - } \ - \ - /* memory transfer optimization, use USM-host for temporary speeds \ - * up transfer to device */ \ - using usm_host_allocatorT = \ - sycl::usm_allocator; \ - \ - size_t strides_size = 2 * result_ndim; \ - shape_elem_type *dev_strides_data = \ - sycl::malloc_device(strides_size, q); \ - \ - /* create host temporary for packed strides managed by shared \ - * pointer */ \ - auto strides_host_packed = \ - std::vector( \ - strides_size, usm_host_allocatorT(q)); \ - \ - /* packed vector is concatenation of result_strides, \ - * input1_strides and input2_strides */ \ - std::copy(result_strides, result_strides + result_ndim, \ - strides_host_packed.begin()); \ - std::copy(input1_strides, input1_strides + result_ndim, \ - strides_host_packed.begin() + result_ndim); \ - \ - auto copy_strides_ev = q.copy( \ - strides_host_packed.data(), dev_strides_data, \ - strides_host_packed.size()); \ - \ - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { \ - size_t output_id = global_id[0]; /* for (size_t i = 0; i < \ - result_size; ++i) */ \ - { \ - const shape_elem_type *result_strides_data = \ - &dev_strides_data[0]; \ - const shape_elem_type *input1_strides_data = \ - &dev_strides_data[result_ndim]; \ - \ - size_t input_id = 0; \ - for (size_t i = 0; i < input1_ndim; ++i) { \ - const size_t output_xyz_id = \ - get_xyz_id_by_id_inkernel(output_id, \ - result_strides_data, \ - result_ndim, i); \ - input_id += output_xyz_id * input1_strides_data[i]; \ - } \ - \ - const _DataType input_elem = input1_data[input_id]; \ - result[output_id] = __operation1__; \ - } \ - }; \ - 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 input_elem = input1_data[i]; \ - result[i] = __operation1__; \ - } \ - }; \ - auto kernel_func = [&](sycl::handler &cgh) { \ - cgh.parallel_for>( \ - gws, kernel_parallel_for_func); \ - }; \ - \ - if constexpr (is_any_v<_DataType, float, double>) { \ - if (q.get_device().has(sycl::aspect::fp64)) { \ - event = __operation2__; \ - \ - event_ref = reinterpret_cast(&event); \ - return DPCTLEvent_Copy(event_ref); \ - } \ - } \ - event = q.submit(kernel_func); \ - } \ - \ - event_ref = reinterpret_cast(&event); \ - return DPCTLEvent_Copy(event_ref); \ - } \ - \ - template \ - void __name__( \ - void *result_out, const size_t result_size, const size_t result_ndim, \ - const shape_elem_type *result_shape, \ - const shape_elem_type *result_strides, const void *input1_in, \ - const size_t input1_size, const size_t input1_ndim, \ - const shape_elem_type *input1_shape, \ - const shape_elem_type *input1_strides, const size_t *where) \ - { \ - DPCTLSyclQueueRef q_ref = \ - reinterpret_cast(&DPNP_QUEUE); \ - DPCTLEventVectorRef dep_event_vec_ref = nullptr; \ - DPCTLSyclEventRef event_ref = __name__<_DataType>( \ - q_ref, result_out, result_size, result_ndim, result_shape, \ - result_strides, input1_in, input1_size, input1_ndim, input1_shape, \ - input1_strides, where, dep_event_vec_ref); \ - DPCTLEvent_WaitAndThrow(event_ref); \ - DPCTLEvent_Delete(event_ref); \ - } \ - \ - template \ - void (*__name__##_default)( \ - void *, const size_t, const size_t, const shape_elem_type *, \ - const shape_elem_type *, const void *, const size_t, const size_t, \ - const shape_elem_type *, const shape_elem_type *, const size_t *) = \ - __name__<_DataType>; \ - \ - template \ - DPCTLSyclEventRef (*__name__##_ext)( \ - DPCTLSyclQueueRef, void *, const size_t, const size_t, \ - const shape_elem_type *, const shape_elem_type *, const void *, \ - const size_t, const size_t, const shape_elem_type *, \ - const shape_elem_type *, const size_t *, const DPCTLEventVectorRef) = \ - __name__<_DataType>; - -#include - -static void func_map_init_elemwise_1arg_1type(func_map_t &fmap) -{ - 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] = { - eft_LNG, (void *)dpnp_erf_c_default}; - fmap[DPNPFuncName::DPNP_FN_ERF][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_erf_c_default}; - fmap[DPNPFuncName::DPNP_FN_ERF][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_erf_c_default}; - - fmap[DPNPFuncName::DPNP_FN_ERF_EXT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_erf_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ERF_EXT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_erf_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ERF_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_erf_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ERF_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_erf_c_ext}; - - return; -} - -void func_map_init_elemwise(func_map_t &fmap) -{ - func_map_init_elemwise_1arg_1type(fmap); - - return; -} diff --git a/dpnp/backend/kernels/elementwise_functions/erf.hpp b/dpnp/backend/kernels/elementwise_functions/erf.hpp new file mode 100644 index 000000000000..ab69745ec485 --- /dev/null +++ b/dpnp/backend/kernels/elementwise_functions/erf.hpp @@ -0,0 +1,51 @@ +//***************************************************************************** +// 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. +// +// 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 + +#include + +namespace dpnp::kernels::erf +{ +template +struct ErfFunctor +{ + // is function constant for given argT + using is_constant = typename std::false_type; + // constant value, if constant + // constexpr Tp constant_value = Tp{}; + // is function defined for sycl::vec + using supports_vec = typename std::false_type; + // do both argT and Tp support subgroup store/load operation + using supports_sg_loadstore = typename std::true_type; + + Tp operator()(const argT &x) const + { + return sycl::erf(x); + } +}; +} // namespace dpnp::kernels::erf diff --git a/dpnp/backend/src/dpnp_fptr.hpp b/dpnp/backend/src/dpnp_fptr.hpp index 8a7d50597c12..35ac48829a30 100644 --- a/dpnp/backend/src/dpnp_fptr.hpp +++ b/dpnp/backend/src/dpnp_fptr.hpp @@ -157,7 +157,6 @@ class dpnp_less_comp * FPTR interface initialization functions */ void func_map_init_arraycreation(func_map_t &fmap); -void func_map_init_elemwise(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); diff --git a/dpnp/backend/src/dpnp_iface_fptr.cpp b/dpnp/backend/src/dpnp_iface_fptr.cpp index c921f5e9bd27..4c980a3f51c0 100644 --- a/dpnp/backend/src/dpnp_iface_fptr.cpp +++ b/dpnp/backend/src/dpnp_iface_fptr.cpp @@ -96,7 +96,6 @@ static func_map_t func_map_init() func_map_t fmap; func_map_init_arraycreation(fmap); - func_map_init_elemwise(fmap); func_map_init_linalg(fmap); func_map_init_mathematical(fmap); func_map_init_random(fmap); diff --git a/dpnp/backend/src/dpnp_utils.hpp b/dpnp/backend/src/dpnp_utils.hpp index 9c27b98d78d6..d3a0d4659c70 100644 --- a/dpnp/backend/src/dpnp_utils.hpp +++ b/dpnp/backend/src/dpnp_utils.hpp @@ -27,12 +27,12 @@ #ifndef BACKEND_UTILS_H // Cython compatibility #define BACKEND_UTILS_H -#include -#include #include #include -#include #include +#include +#include +#include #include @@ -70,99 +70,6 @@ * @} */ -/** - * @ingroup BACKEND_UTILS - * @brief Shape offset calculation used in kernels - * - * Calculates offsets of the array with given shape - * for example: - * input_array_shape[3, 4, 5] - * offsets should be [20, 5, 1] - * - * @param [in] shape array with input shape. - * @param [in] shape_size array size for @ref shape parameter. - * @param [out] offsets Result array with @ref shape_size size. - */ -template -void get_shape_offsets_inkernel(const _DataType *shape, - size_t shape_size, - _DataType *offsets) -{ - size_t dim_prod_input = 1; - for (size_t i = 0; i < shape_size; ++i) { - long i_reverse = shape_size - 1 - i; - offsets[i_reverse] = dim_prod_input; - dim_prod_input *= shape[i_reverse]; - } - - return; -} - -/** - * @ingroup BACKEND_UTILS - * @brief Calculate xyz id for given axis from linear index - * - * Calculates xyz id of the array with given shape. - * for example: - * input_array_shape_offsets[20, 5, 1] - * global_id == 5 - * axis == 1 - * xyz_id should be 1 - * - * @param [in] global_id linear index of the element in multy-D array. - * @param [in] offsets array with input offsets. - * @param [in] offsets_size array size for @ref offsets parameter. - * @param [in] axis axis. - */ -template -_DataType get_xyz_id_by_id_inkernel(size_t global_id, - const _DataType *offsets, - size_t offsets_size, - size_t axis) -{ - /* avoid warning unused variable*/ - (void)offsets_size; - - assert(axis < offsets_size); - - _DataType xyz_id = 0; - long reminder = global_id; - for (size_t i = 0; i < axis + 1; ++i) { - const _DataType axis_val = offsets[i]; - xyz_id = reminder / axis_val; - reminder = reminder % axis_val; - } - - return xyz_id; -} - -/** - * @ingroup BACKEND_UTILS - * @brief Check arrays are equal. - * - * @param [in] input1 Input1. - * @param [in] input1_size Input1 size. - * @param [in] input2 Input2. - * @param [in] input2_size Input2 size. - * - * @return Arrays are equal. - */ -template -static inline bool array_equal(const _DataType *input1, - const size_t input1_size, - const _DataType *input2, - const size_t input2_size) -{ - if (input1_size != input2_size) - return false; - - const std::vector<_DataType> input1_vec(input1, input1 + input1_size); - const std::vector<_DataType> input2_vec(input2, input2 + input2_size); - - return std::equal(std::begin(input1_vec), std::end(input1_vec), - std::begin(input2_vec)); -} - /** * @ingroup BACKEND_UTILS * @brief check support of type T by SYCL device. diff --git a/dpnp/dpnp_algo/CMakeLists.txt b/dpnp/dpnp_algo/CMakeLists.txt index df9c1e26ea7a..f734be5c30b4 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_sorting.pxi ${CMAKE_CURRENT_SOURCE_DIR}/dpnp_algo_mathematical.pxi ${CMAKE_CURRENT_SOURCE_DIR}/dpnp_algo_indexing.pxi - ${CMAKE_CURRENT_SOURCE_DIR}/dpnp_algo_special.pxi ) build_dpnp_cython_ext_with_backend( diff --git a/dpnp/dpnp_algo/dpnp_algo.pxd b/dpnp/dpnp_algo/dpnp_algo.pxd index 0dcc519d198f..84fde77fb001 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pxd +++ b/dpnp/dpnp_algo/dpnp_algo.pxd @@ -24,16 +24,10 @@ # ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF # THE POSSIBILITY OF SUCH DAMAGE. # ***************************************************************************** -cimport dpctl as c_dpctl -from libcpp cimport bool as cpp_bool - -from dpnp.dpnp_algo cimport shape_elem_type, shape_type_c -from dpnp.dpnp_utils.dpnp_algo_utils cimport dpnp_descriptor cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this namespace for Enum import cdef enum DPNPFuncName "DPNPFuncName": - DPNP_FN_ERF_EXT DPNP_FN_MODF_EXT DPNP_FN_PARTITION_EXT DPNP_FN_RNG_BETA_EXT @@ -94,15 +88,6 @@ cdef extern from "dpnp_iface_fptr.hpp": DPNPFuncData get_dpnp_function_ptr(DPNPFuncName name, DPNPFuncType first_type, DPNPFuncType second_type) except + -# C function pointer to the C library template functions -ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_1in_1out_strides_t)(c_dpctl.DPCTLSyclQueueRef, - void *, const size_t, const size_t, - const shape_elem_type * , const shape_elem_type * , - void *, const size_t, const size_t, - const shape_elem_type * , const shape_elem_type * , - const long * , - const c_dpctl.DPCTLEventVectorRef) - """ Internal functions diff --git a/dpnp/dpnp_algo/dpnp_algo.pyx b/dpnp/dpnp_algo/dpnp_algo.pyx index 718953d31449..f431cce0f1ed 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pyx +++ b/dpnp/dpnp_algo/dpnp_algo.pyx @@ -33,25 +33,13 @@ and the rest of the library """ -from libc.time cimport time, time_t -from libcpp.vector cimport vector - -import dpctl import dpnp -import dpnp.config as config -import dpnp.dpnp_container as dpnp_container -import dpnp.dpnp_utils as utils_py -from dpnp.dpnp_array import dpnp_array -cimport cpython -cimport numpy +cimport dpctl as c_dpctl cimport dpnp.dpnp_utils as utils - -import operator - -import numpy +from dpnp.dpnp_algo cimport shape_elem_type, shape_type_c __all__ = [ ] @@ -60,7 +48,6 @@ __all__ = [ include "dpnp_algo_indexing.pxi" include "dpnp_algo_mathematical.pxi" include "dpnp_algo_sorting.pxi" -include "dpnp_algo_special.pxi" """ @@ -115,78 +102,3 @@ cdef dpnp_DPNPFuncType_to_dtype(size_t type): return dpnp.bool else: utils.checker_throw_type_error("dpnp_DPNPFuncType_to_dtype", type) - - -cdef utils.dpnp_descriptor call_fptr_1in_1out_strides(DPNPFuncName fptr_name, - utils.dpnp_descriptor x1, - object dtype=None, - utils.dpnp_descriptor out=None, - object where=True, - func_name=None): - - """ Convert type (x1.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(fptr_name, param1_type, param1_type) - - x1_obj = x1.get_array() - - # get FPTR function and return type - cdef (DPNPFuncType, void *) ret_type_and_func = utils.get_ret_type_and_func(kernel_data, - x1_obj.sycl_device.has_aspect_fp64) - cdef DPNPFuncType return_type = ret_type_and_func[0] - cdef fptr_1in_1out_strides_t func = < fptr_1in_1out_strides_t > ret_type_and_func[1] - - result_type = dpnp_DPNPFuncType_to_dtype( < size_t > return_type) - - cdef shape_type_c x1_shape = x1.shape - cdef shape_type_c x1_strides = utils.strides_to_vector(x1.strides, x1_shape) - - cdef shape_type_c result_shape = x1_shape - cdef utils.dpnp_descriptor result - - if out is None: - """ Create result array with type given by FPTR data """ - result = utils.create_output_descriptor(result_shape, - return_type, - None, - device=x1_obj.sycl_device, - usm_type=x1_obj.usm_type, - sycl_queue=x1_obj.sycl_queue) - else: - if out.dtype != result_type: - utils.checker_throw_value_error(func_name, 'out.dtype', out.dtype, result_type) - if out.shape != result_shape: - utils.checker_throw_value_error(func_name, 'out.shape', out.shape, result_shape) - - result = out - - utils.get_common_usm_allocation(x1, result) # check USM allocation is common - - result_sycl_queue = result.get_array().sycl_queue - - cdef c_dpctl.SyclQueue q = result_sycl_queue - cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() - - cdef shape_type_c result_strides = utils.strides_to_vector(result.strides, result_shape) - - """ Call FPTR function """ - cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, - result.get_data(), - result.size, - result.ndim, - result_shape.data(), - result_strides.data(), - x1.get_data(), - x1.size, - x1.ndim, - x1_shape.data(), - x1_strides.data(), - NULL, - NULL) # dep_events_ref - - with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) - c_dpctl.DPCTLEvent_Delete(event_ref) - - return result diff --git a/dpnp/dpnp_iface.py b/dpnp/dpnp_iface.py index 70d1606e0d90..4b354d547025 100644 --- a/dpnp/dpnp_iface.py +++ b/dpnp/dpnp_iface.py @@ -55,6 +55,7 @@ from dpnp.linalg import * from dpnp.memory import * from dpnp.random import * +from dpnp.special import * __all__ = [ "are_same_logical_tensors", @@ -87,8 +88,6 @@ from dpnp.dpnp_iface_histograms import __all__ as __all__histograms from dpnp.dpnp_iface_indexing import * from dpnp.dpnp_iface_indexing import __all__ as __all__indexing -from dpnp.dpnp_iface_libmath import * -from dpnp.dpnp_iface_libmath import __all__ as __all__libmath from dpnp.dpnp_iface_linearalgebra import * from dpnp.dpnp_iface_linearalgebra import __all__ as __all__linearalgebra from dpnp.dpnp_iface_logic import * @@ -123,7 +122,6 @@ __all__ += __all__functional __all__ += __all__histograms __all__ += __all__indexing -__all__ += __all__libmath __all__ += __all__linearalgebra __all__ += __all__logic __all__ += __all__manipulation diff --git a/dpnp/dpnp_iface_bitwise.py b/dpnp/dpnp_iface_bitwise.py index bbc9e465f20d..f50f8dc9b872 100644 --- a/dpnp/dpnp_iface_bitwise.py +++ b/dpnp/dpnp_iface_bitwise.py @@ -37,8 +37,9 @@ """ -# pylint: disable=protected-access +# pylint: disable=duplicate-code # pylint: disable=no-name-in-module +# pylint: disable=protected-access import dpctl.tensor._tensor_elementwise_impl as ti import numpy diff --git a/dpnp/dpnp_iface_libmath.py b/dpnp/dpnp_iface_libmath.py deleted file mode 100644 index eaf6c5676a42..000000000000 --- a/dpnp/dpnp_iface_libmath.py +++ /dev/null @@ -1,97 +0,0 @@ -# -*- coding: utf-8 -*- -# ***************************************************************************** -# Copyright (c) 2016-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. -# -# 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. -# ***************************************************************************** - -""" -Interface of the function from Python Math library - -Notes ------ -This module is a face or public interface file for the library -it contains: - - Interface functions - - documentation for the functions - - The functions parameters check - -""" - -import math - -import dpnp - -# pylint: disable=no-name-in-module -from .dpnp_algo import ( - dpnp_erf, -) -from .dpnp_utils import ( - create_output_descriptor_py, -) - -__all__ = ["erf"] - - -def erf(in_array1): - """ - Returns the error function of complex argument. - - For full documentation refer to :obj:`scipy.special.erf`. - - Limitations - ----------- - Parameter `in_array1` is supported as :obj:`dpnp.ndarray`. - Otherwise the function will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. - - .. seealso:: :obj:`math.erf` - - Examples - -------- - >>> import dpnp as np - >>> x = np.linspace(2.0, 3.0, num=5) - >>> [i for i in x] - [2.0, 2.25, 2.5, 2.75, 3.0] - >>> out = np.erf(x) - >>> [i for i in out] - [0.99532227, 0.99853728, 0.99959305, 0.99989938, 0.99997791] - - """ - - x1_desc = dpnp.get_dpnp_descriptor( - in_array1, copy_when_strides=False, 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" - ) - return dpnp_erf(x1_desc).get_pyobj() - - result = create_output_descriptor_py( - in_array1.shape, in_array1.dtype, None - ).get_pyobj() - for i in range(result.size): - result[i] = math.erf(in_array1[i]) - - return result diff --git a/dpnp/dpnp_utils/dpnp_algo_utils.pxd b/dpnp/dpnp_utils/dpnp_algo_utils.pxd index f7de1c35ff23..76c5c08eb9d2 100644 --- a/dpnp/dpnp_utils/dpnp_algo_utils.pxd +++ b/dpnp/dpnp_utils/dpnp_algo_utils.pxd @@ -28,7 +28,7 @@ from libcpp cimport bool as cpp_bool from dpnp.dpnp_algo cimport shape_type_c -from dpnp.dpnp_algo.dpnp_algo cimport DPNPFuncData, DPNPFuncName, DPNPFuncType +from dpnp.dpnp_algo.dpnp_algo cimport DPNPFuncType cpdef checker_throw_value_error(function_name, param_name, param, expected) @@ -59,11 +59,6 @@ Return: cpdef tuple _object_to_tuple(object obj) -cpdef tuple get_axis_offsets(shape) -""" -Compute axis offsets in the linear array memory -""" - cdef class dpnp_descriptor: """array DPNP descriptor""" @@ -85,19 +80,7 @@ cdef dpnp_descriptor create_output_descriptor(shape_type_c output_shape, Create output dpnp_descriptor based on shape, type and 'out' parameters """ -cdef shape_type_c strides_to_vector(object strides, object shape) except * -""" -Get or calculate srtides based on shape. -""" - cdef tuple get_common_usm_allocation(dpnp_descriptor x1, dpnp_descriptor x2) """ Get common USM allocation in the form of (sycl_device, usm_type, sycl_queue) """ - -cdef (DPNPFuncType, void *) get_ret_type_and_func(DPNPFuncData kernel_data, - cpp_bool has_aspect_fp64) -""" -Get the corresponding return type and function pointer based on the -capability of the allocated result array device. -""" diff --git a/dpnp/dpnp_utils/dpnp_algo_utils.pyx b/dpnp/dpnp_utils/dpnp_algo_utils.pyx index 732b792af125..74b0f11a1b4f 100644 --- a/dpnp/dpnp_utils/dpnp_algo_utils.pyx +++ b/dpnp/dpnp_utils/dpnp_algo_utils.pyx @@ -35,7 +35,6 @@ This module contains different helpers and utilities import dpctl import dpctl.utils as dpu import numpy -from dpctl.tensor._numpy_helper import AxisError import dpnp import dpnp.config as config @@ -46,12 +45,9 @@ cimport cpython cimport cython cimport numpy from libcpp cimport bool as cpp_bool -from libcpp.complex cimport complex as cpp_complex from dpnp.dpnp_algo.dpnp_algo cimport ( dpnp_DPNPFuncType_to_dtype, - dpnp_dtype_to_DPNPFuncType, - get_dpnp_function_ptr, ) """ @@ -61,10 +57,8 @@ __all__ = [ "call_origin", "checker_throw_type_error", "checker_throw_value_error", - "create_output_descriptor_py", "convert_item", "dpnp_descriptor", - "get_axis_offsets", "get_usm_allocations", "map_dtype_to_device", "_object_to_tuple", @@ -316,39 +310,6 @@ cpdef checker_throw_value_error(function_name, param_name, param, expected): raise ValueError(err_msg) -cpdef dpnp_descriptor create_output_descriptor_py(shape_type_c output_shape, - d_type, - requested_out, - device=None, - usm_type="device", - sycl_queue=None): - py_type = dpnp.default_float_type() if d_type is None else d_type - - cdef DPNPFuncType c_type = dpnp_dtype_to_DPNPFuncType(py_type) - - return create_output_descriptor(output_shape, - c_type, - requested_out, - device=device, - usm_type=usm_type, - sycl_queue=sycl_queue) - - -cpdef tuple get_axis_offsets(shape): - """ - Compute axis offsets in the linear array memory - """ - - res_size = len(shape) - result = [0] * res_size - acc = 1 - for i in range(res_size - 1, -1, -1): - result[i] = acc - acc *= shape[i] - - return _object_to_tuple(result) - - cdef dpnp_descriptor create_output_descriptor(shape_type_c output_shape, DPNPFuncType c_type, dpnp_descriptor requested_out, @@ -416,19 +377,6 @@ cpdef cpp_bool use_origin_backend(input1=None, size_t compute_size=0): return False -cdef shape_type_c strides_to_vector(object strides, object shape) except *: - """ - Get or calculate srtides based on shape. - """ - cdef shape_type_c res - if strides is None: - res = get_axis_offsets(shape) - else: - res = strides - - return res - - cdef tuple get_common_usm_allocation(dpnp_descriptor x1, dpnp_descriptor x2): """Get common USM allocation in the form of (sycl_device, usm_type, sycl_queue).""" array1_obj = x1.get_array() @@ -449,24 +397,6 @@ cdef tuple get_common_usm_allocation(dpnp_descriptor x1, dpnp_descriptor x2): return (common_sycl_queue.sycl_device, common_usm_type, common_sycl_queue) -@cython.linetrace(False) -cdef (DPNPFuncType, void *) get_ret_type_and_func(DPNPFuncData kernel_data, - cpp_bool has_aspect_fp64): - """ - This function is responsible for determining the appropriate return type - and function pointer based on the capability of the allocated result array device. - """ - return_type = kernel_data.return_type - func = kernel_data.ptr - - if kernel_data.ptr_no_fp64 != NULL and not has_aspect_fp64: - - return_type = kernel_data.return_type_no_fp64 - func = kernel_data.ptr_no_fp64 - - return return_type, func - - cdef class dpnp_descriptor: def __init__(self, obj): """ Initialize variables """ diff --git a/dpnp/dpnp_algo/dpnp_algo_special.pxi b/dpnp/special/__init__.py similarity index 70% rename from dpnp/dpnp_algo/dpnp_algo_special.pxi rename to dpnp/special/__init__.py index 6e1900dca66a..484fe38c29ef 100644 --- a/dpnp/dpnp_algo/dpnp_algo_special.pxi +++ b/dpnp/special/__init__.py @@ -1,8 +1,6 @@ -# cython: language_level=3 -# cython: linetrace=True # -*- coding: utf-8 -*- # ***************************************************************************** -# Copyright (c) 2016-2025, Intel Corporation +# Copyright (c) 2025, Intel Corporation # All rights reserved. # # Redistribution and use in source and binary forms, with or without @@ -26,19 +24,25 @@ # THE POSSIBILITY OF SUCH DAMAGE. # ***************************************************************************** -"""Module Backend (Special part) +""" +``dpnp.special`` +================ + +The submodule provides a large collection of mathematical functions that are +widely used in science and engineering. It includes special functions of +mathematical physics (e.g., Bessel, elliptic, gamma, hypergeometric), as well +as standard functions like `erf`, `sinc`, and `logit`. -This module contains interface functions between C backend layer -and the rest of the library +The functions in the submodule invokes VM implementation from pybind11 +extension above OneMKL VM if possible or uses a dedicated SYCL kernel, or, +alternatively, is implemented through a subset of python calls. """ -# NO IMPORTs here. All imports must be placed into main "dpnp_algo.pyx" file +from ._erf import ( + erf, +) -__all__ += [ - 'dpnp_erf', +__all__ = [ + "erf", ] - - -cpdef utils.dpnp_descriptor dpnp_erf(utils.dpnp_descriptor x1): - return call_fptr_1in_1out_strides(DPNP_FN_ERF_EXT, x1) diff --git a/dpnp/special/_erf.py b/dpnp/special/_erf.py new file mode 100644 index 000000000000..7d70c6759287 --- /dev/null +++ b/dpnp/special/_erf.py @@ -0,0 +1,123 @@ +# -*- coding: utf-8 -*- +# ***************************************************************************** +# 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. +# +# 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. +# ***************************************************************************** + +""" +Interface of the Error functions + +Notes +----- +This module is a face or public interface file for the library +it contains: + - Interface functions + - documentation for the functions + +""" + +# pylint: disable=protected-access + +# pylint: disable=no-name-in-module +import dpnp.backend.extensions.ufunc._ufunc_impl as ufi +from dpnp.dpnp_algo.dpnp_elementwise_common import DPNPUnaryFunc + +__all__ = ["erf"] + + +# pylint: disable=too-few-public-methods +class DPNPErf(DPNPUnaryFunc): + """Class that implements a family of erf functions.""" + + def __init__( + self, + name, + result_type_resolver_fn, + unary_dp_impl_fn, + docs, + mkl_fn_to_call=None, + mkl_impl_fn=None, + ): + super().__init__( + name, + result_type_resolver_fn, + unary_dp_impl_fn, + docs, + mkl_fn_to_call=mkl_fn_to_call, + mkl_impl_fn=mkl_impl_fn, + ) + + def __call__(self, x, out=None): # pylint: disable=signature-differs + return super().__call__(x, out=out) + + +_ERF_DOCSTRING = r""" +Returns the error function of complex argument. + +It is defined as :math:`\frac{2}{\sqrt{\pi}} \int_{0}^{z} e^{-t^2} \, dt`. + +For full documentation refer to :obj:`scipy.special.erf`. + +Parameters +---------- +x : {dpnp.ndarray, usm_ndarray} + Input array. +out : {dpnp.ndarray, usm_ndarray}, optional + Optional output array for the function values. + +See Also +-------- +:obj:`dpnp.special.erfc` : Complementary error function. +:obj:`dpnp.special.erfinv` : Inverse of the error function. +:obj:`dpnp.special.erfcinv` : Inverse of the complementary error function. +:obj:`dpnp.special.erfcx` : Scaled complementary error function. +:obj:`dpnp.special.erfi` : Imaginary error function. + +Notes +----- +The cumulative of the unit normal distribution is given by + +.. math:: + \Phi(z) = \frac{1}{2} \left[ + 1 + \operatorname{erf} \left( + \frac{z}{\sqrt{2}} + \right) + \right] + +Examples +-------- +>>> import dpnp as np +>>> x = np.linspace(-3, 3, num=5) +>>> np.special.erf(x) +array([[-0.99997791, -0.96610515, 0. , 0.96610515, 0.99997791]) + +""" + +erf = DPNPErf( + "erf", + ufi._erf_result_type, + ufi._erf, + _ERF_DOCSTRING, + mkl_fn_to_call="_mkl_erf_to_call", + mkl_impl_fn="_erf", +) diff --git a/dpnp/tests/skipped_tests_cuda.tbl b/dpnp/tests/skipped_tests_cuda.tbl index 86c73b0e6a62..2846a3c2f5c7 100644 --- a/dpnp/tests/skipped_tests_cuda.tbl +++ b/dpnp/tests/skipped_tests_cuda.tbl @@ -807,11 +807,3 @@ tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_3_{extern tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_3_{external=True, length=20000}::test_partition_one_dim tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_3_{external=True, length=20000}::test_partition_sequence_kth tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_3_{external=True, length=20000}::test_partition_zero_dim - -# erf -tests/test_special.py::test_erf -tests/test_special.py::test_erf_fallback -tests/test_strides.py::test_erf[int32] -tests/test_strides.py::test_erf[int64] -tests/test_strides.py::test_erf[float32] -tests/test_strides.py::test_erf[float64] diff --git a/dpnp/tests/test_special.py b/dpnp/tests/test_special.py index 3c567dcd3ca6..0f52ed578e64 100644 --- a/dpnp/tests/test_special.py +++ b/dpnp/tests/test_special.py @@ -1,32 +1,51 @@ -import math - import numpy +import pytest from numpy.testing import assert_allclose import dpnp - -def test_erf(): - a = numpy.linspace(2.0, 3.0, num=10) - ia = dpnp.array(a) - - expected = numpy.empty_like(a) - for idx, val in enumerate(a): - expected[idx] = math.erf(val) - - result = dpnp.erf(ia) - - assert_allclose(result, expected) - - -def test_erf_fallback(): - a = numpy.linspace(2.0, 3.0, num=10) - dpa = dpnp.linspace(2.0, 3.0, num=10) - - expected = numpy.empty_like(a) - for idx, val in enumerate(a): - expected[idx] = math.erf(val) - - result = dpnp.erf(dpa) - - assert_allclose(result, expected) +from .helper import ( + assert_dtype_allclose, + generate_random_numpy_array, + get_all_dtypes, + get_complex_dtypes, +) +from .third_party.cupy.testing import installed, with_requires + + +@with_requires("scipy") +class TestErf: + @pytest.mark.parametrize( + "dt", get_all_dtypes(no_none=True, no_float16=False, no_complex=True) + ) + def test_basic(self, dt): + import scipy.special + + a = generate_random_numpy_array((2, 5), dtype=dt) + ia = dpnp.array(a) + + result = dpnp.special.erf(ia) + expected = scipy.special.erf(a) + + # scipy >= 0.16.0 returns float64, but dpnp returns float32 + to_float32 = dt in (dpnp.bool, dpnp.float16) + only_type_kind = installed("scipy>=0.16.0") and to_float32 + assert_dtype_allclose( + result, expected, check_only_type_kind=only_type_kind + ) + + def test_nan_inf(self): + import scipy.special + + a = numpy.array([numpy.nan, -numpy.inf, numpy.inf]) + ia = dpnp.array(a) + + result = dpnp.special.erf(ia) + expected = scipy.special.erf(a) + assert_allclose(result, expected) + + @pytest.mark.parametrize("dt", get_complex_dtypes()) + def test_complex(self, dt): + x = dpnp.empty(5, dtype=dt) + with pytest.raises(ValueError): + dpnp.special.erf(x) diff --git a/dpnp/tests/test_strides.py b/dpnp/tests/test_strides.py index 712bc7c91df3..d3657c3e0e21 100644 --- a/dpnp/tests/test_strides.py +++ b/dpnp/tests/test_strides.py @@ -12,10 +12,12 @@ get_all_dtypes, get_complex_dtypes, get_float_complex_dtypes, + get_float_dtypes, get_integer_dtypes, get_integer_float_dtypes, numpy_version, ) +from .third_party.cupy.testing import installed, with_requires @pytest.mark.usefixtures("suppress_divide_invalid_numpy_warnings") @@ -164,22 +166,21 @@ def test_reduce_hypot(dtype, stride): assert_dtype_allclose(result, expected, check_only_type_kind=flag) -@pytest.mark.parametrize( - "dtype", - get_integer_float_dtypes( - no_unsigned=True, xfail_dtypes=[dpnp.int8, dpnp.int16] - ), -) -def test_erf(dtype): - a = dpnp.linspace(-1, 1, num=10, dtype=dtype) - b = a[::2] - result = dpnp.erf(b) +@with_requires("scipy") +@pytest.mark.parametrize("dtype", get_float_dtypes(no_float16=False)) +@pytest.mark.parametrize("stride", [2, -1, -3]) +def test_erf(dtype, stride): + import scipy.special - expected = numpy.empty_like(b.asnumpy()) - for idx, val in enumerate(b): - expected[idx] = math.erf(val) + x = generate_random_numpy_array(10, dtype=dtype) + a, ia = x[::stride], dpnp.array(x)[::stride] - assert_dtype_allclose(result, expected) + result = dpnp.special.erf(ia) + expected = scipy.special.erf(a) + + # scipy >= 0.16.0 returns float64, but dpnp returns float32 + only_type_kind = installed("scipy>=0.16.0") and (dtype == dpnp.float16) + assert_dtype_allclose(result, expected, check_only_type_kind=only_type_kind) @pytest.mark.parametrize("dtype", get_float_complex_dtypes()) diff --git a/dpnp/tests/test_sycl_queue.py b/dpnp/tests/test_sycl_queue.py index 31dfb74f2cfa..cc79e1be2099 100644 --- a/dpnp/tests/test_sycl_queue.py +++ b/dpnp/tests/test_sycl_queue.py @@ -1470,6 +1470,14 @@ def test_interp(device, left, right, period): assert_sycl_queue_equal(result.sycl_queue, x.sycl_queue) +@pytest.mark.parametrize("device", valid_dev, ids=dev_ids) +def test_erf(device): + x = dpnp.linspace(-3, 3, num=5, device=device) + + result = dpnp.special.erf(x) + assert_sycl_queue_equal(result.sycl_queue, x.sycl_queue) + + @pytest.mark.parametrize("device", valid_dev, ids=dev_ids) class TestLinAlgebra: @pytest.mark.parametrize( diff --git a/dpnp/tests/test_usm_type.py b/dpnp/tests/test_usm_type.py index aed316eca533..1fed490322d7 100644 --- a/dpnp/tests/test_usm_type.py +++ b/dpnp/tests/test_usm_type.py @@ -1279,6 +1279,13 @@ def test_choose(usm_type_x, usm_type_ind): assert z.usm_type == du.get_coerced_usm_type([usm_type_x, usm_type_ind]) +@pytest.mark.parametrize("usm_type", list_of_usm_types) +def test_erf(usm_type): + x = dpnp.linspace(-3, 3, num=5, usm_type=usm_type) + y = dpnp.special.erf(x) + assert x.usm_type == y.usm_type == usm_type + + class TestInterp: @pytest.mark.parametrize("usm_type_x", list_of_usm_types) @pytest.mark.parametrize("usm_type_xp", list_of_usm_types) diff --git a/dpnp/tests/third_party/cupy/testing/_loops.py b/dpnp/tests/third_party/cupy/testing/_loops.py index f57e555ea363..946455857d7b 100644 --- a/dpnp/tests/third_party/cupy/testing/_loops.py +++ b/dpnp/tests/third_party/cupy/testing/_loops.py @@ -71,7 +71,7 @@ def _call_func_cupy(impl, args, kw, name, sp_name, scipy_name): if sp_name: kw[sp_name] = cupyx.scipy.sparse if scipy_name: - kw[scipy_name] = cupyx.scipy + kw[scipy_name] = cupy kw[name] = cupy result, error = _call_func(impl, args, kw) return result, error diff --git a/dpnp/tests/third_party/cupyx/__init__.py b/dpnp/tests/third_party/cupyx/__init__.py new file mode 100644 index 000000000000..e69de29bb2d1 diff --git a/dpnp/tests/third_party/cupyx/scipy_tests/__init__.py b/dpnp/tests/third_party/cupyx/scipy_tests/__init__.py new file mode 100644 index 000000000000..e69de29bb2d1 diff --git a/dpnp/tests/third_party/cupyx/scipy_tests/special_tests/__init__.py b/dpnp/tests/third_party/cupyx/scipy_tests/special_tests/__init__.py new file mode 100644 index 000000000000..e69de29bb2d1 diff --git a/dpnp/tests/third_party/cupyx/scipy_tests/special_tests/test_erf.py b/dpnp/tests/third_party/cupyx/scipy_tests/special_tests/test_erf.py new file mode 100644 index 000000000000..706107a45715 --- /dev/null +++ b/dpnp/tests/third_party/cupyx/scipy_tests/special_tests/test_erf.py @@ -0,0 +1,160 @@ +from __future__ import annotations + +import unittest + +import numpy +import pytest + +import dpnp as cupy +import dpnp.special +from dpnp.tests.third_party.cupy import testing + + +def _boundary_inputs(boundary, rtol, atol): + left = boundary * (1 - numpy.copysign(rtol, boundary)) - atol + right = boundary * (1 + numpy.copysign(rtol, boundary)) + atol + return [left, boundary, right] + + +@testing.with_requires("scipy") +class _TestBase: + + # @testing.with_requires('scipy>=1.16.0') + def test_erf(self): + self.check_unary("erf") + + @pytest.mark.skip("erfc() is not supported yet") + @testing.with_requires("scipy>=1.16.0") + def test_erfc(self): + self.check_unary("erfc") + + @pytest.mark.skip("erfcx() is not supported yet") + @testing.with_requires("scipy>=1.16.0") + def test_erfcx(self): + self.check_unary("erfcx") + + @pytest.mark.skip("erfinv() is not supported yet") + def test_erfinv(self): + self.check_unary("erfinv") + self.check_unary_random("erfinv", scale=2, offset=-1) + self.check_unary_boundary("erfinv", boundary=-1) + self.check_unary_boundary("erfinv", boundary=1) + + @pytest.mark.skip("erfcinv() is not supported yet") + def test_erfcinv(self): + self.check_unary("erfcinv") + self.check_unary_random("erfcinv", scale=2, offset=0) + self.check_unary_boundary("erfcinv", boundary=0) + self.check_unary_boundary("erfcinv", boundary=2) + + +@testing.with_requires("scipy") +class TestSpecial(unittest.TestCase, _TestBase): + + @testing.for_dtypes(["e", "f", "d"]) + @testing.numpy_cupy_allclose(atol=1e-5, scipy_name="scp") + def check_unary(self, name, xp, scp, dtype): + import scipy.special + + a = testing.shaped_arange((2, 3), xp, dtype) + return getattr(scp.special, name)(a) + + @testing.for_dtypes(["f", "d"]) + @testing.numpy_cupy_allclose(atol=1e-5, scipy_name="scp") + def check_unary_random(self, name, xp, scp, dtype, scale, offset): + import scipy.special + + a = testing.shaped_random((2, 3), xp, dtype, scale=scale) + offset + return getattr(scp.special, name)(a) + + @testing.for_dtypes(["f", "d"]) + @testing.numpy_cupy_allclose(atol=1e-5, scipy_name="scp") + def check_unary_boundary(self, name, xp, scp, dtype, boundary): + import scipy.special + + a = _boundary_inputs(boundary, 1.0 / 1024, 1.0 / 1024) + a = xp.array(a, dtype=dtype) + return getattr(scp.special, name)(a) + + @pytest.mark.skip("erfinv() is not supported yet") + @testing.with_requires("scipy>=1.4.0") + @testing.for_dtypes(["f", "d"]) + def test_erfinv_behavior(self, dtype): + a = cupy.empty((1,), dtype=dtype) + + a[:] = 1.0 + 1e-6 + a = cupyx.scipy.special.erfinv(a) + assert cupy.isnan(a) + a[:] = -1.0 - 1e-6 + a = cupyx.scipy.special.erfinv(a) + assert cupy.isnan(a) + a[:] = 1.0 + a = cupyx.scipy.special.erfinv(a) + assert numpy.isposinf(cupy.asnumpy(a)) + a[:] = -1.0 + a = cupyx.scipy.special.erfinv(a) + assert numpy.isneginf(cupy.asnumpy(a)) + + @pytest.mark.skip("erfcinv() is not supported yet") + @testing.with_requires("scipy>=1.4.0") + @testing.for_dtypes(["f", "d"]) + def test_erfcinv_behavior(self, dtype): + a = cupy.empty((1,), dtype=dtype) + + a[:] = 2.0 + 1e-6 + a = cupyx.scipy.special.erfcinv(a) + assert cupy.isnan(a) + a[:] = 0.0 - 1e-6 + a = cupyx.scipy.special.erfcinv(a) + assert cupy.isnan(a) + a[:] = 0.0 + a = cupyx.scipy.special.erfcinv(a) + assert numpy.isposinf(cupy.asnumpy(a)) + a[:] = 2.0 + a = cupyx.scipy.special.erfcinv(a) + assert numpy.isneginf(cupy.asnumpy(a)) + + +@pytest.mark.skip("fuse() is not supported yet") +@testing.with_requires("scipy") +class TestFusionSpecial(unittest.TestCase, _TestBase): + + @testing.for_dtypes(["e", "f", "d"]) + @testing.numpy_cupy_allclose(atol=1e-5, scipy_name="scp") + def check_unary(self, name, xp, scp, dtype): + import scipy.special + + a = testing.shaped_arange((2, 3), xp, dtype) + + @cupy.fuse() + def f(x): + return getattr(scp.special, name)(x) + + return f(a) + + @testing.for_dtypes(["f", "d"]) + @testing.numpy_cupy_allclose(atol=1e-5, scipy_name="scp") + def check_unary_random(self, name, xp, scp, dtype, scale, offset): + import scipy.special + + a = testing.shaped_random((2, 3), xp, dtype, scale=scale) + offset + + @cupy.fuse() + def f(x): + return getattr(scp.special, name)(x) + + return f(a) + + @testing.for_dtypes(["f", "d"]) + @testing.numpy_cupy_allclose(atol=1e-5, scipy_name="scp") + def check_unary_boundary(self, name, xp, scp, dtype, boundary): + import scipy.special + + a = _boundary_inputs(boundary, 1.0 / 1024, 1.0 / 1024) + a = xp.array(a, dtype=dtype) + + @cupy.fuse() + def f(x): + return getattr(scp.special, name)(x) + + return f(a) diff --git a/setup.py b/setup.py index 7285ce0cf3fd..9c798f14b37b 100644 --- a/setup.py +++ b/setup.py @@ -37,6 +37,7 @@ "dpnp.linalg", "dpnp.memory", "dpnp.random", + "dpnp.special", ], package_data={ "dpnp": [ @@ -48,6 +49,8 @@ "tests/testing/*.py", "tests/third_party/cupy/*.py", "tests/third_party/cupy/*/*.py", + "tests/third_party/cupyx/*.py", + "tests/third_party/cupyx/*/*.py", ] }, include_package_data=False,