diff --git a/CHANGELOG.md b/CHANGELOG.md index 033da8fd0822..2cd53f693ee0 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -29,6 +29,7 @@ This release achieves 100% compliance with Python Array API specification (revis * Removed `einsum_call` keyword from `dpnp.einsum_path` signature [#2421](https://github.com/IntelPython/dpnp/pull/2421) * Changed `"max dimensions"` to `None` in array API capabilities [#2432](https://github.com/IntelPython/dpnp/pull/2432) * Updated kernel header `i0.hpp` to expose `cyl_bessel_i0` function depending on build target [#2440](https://github.com/IntelPython/dpnp/pull/2440) +* Added MKL functions `arg`, `copysign`, `i0`, and `inv` from VM namespace to be used by implementation of the appropriate element-wise functions [#2445](https://github.com/IntelPython/dpnp/pull/2445) ### Fixed diff --git a/dpnp/backend/extensions/vm/CMakeLists.txt b/dpnp/backend/extensions/vm/CMakeLists.txt index 60886b3a2ce8..9ad513ccbff6 100644 --- a/dpnp/backend/extensions/vm/CMakeLists.txt +++ b/dpnp/backend/extensions/vm/CMakeLists.txt @@ -29,6 +29,7 @@ if(NOT _use_onemkl_interfaces) ${CMAKE_CURRENT_SOURCE_DIR}/acos.cpp ${CMAKE_CURRENT_SOURCE_DIR}/acosh.cpp ${CMAKE_CURRENT_SOURCE_DIR}/add.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/arg.cpp ${CMAKE_CURRENT_SOURCE_DIR}/asin.cpp ${CMAKE_CURRENT_SOURCE_DIR}/asinh.cpp ${CMAKE_CURRENT_SOURCE_DIR}/atan.cpp @@ -37,6 +38,7 @@ if(NOT _use_onemkl_interfaces) ${CMAKE_CURRENT_SOURCE_DIR}/cbrt.cpp ${CMAKE_CURRENT_SOURCE_DIR}/ceil.cpp ${CMAKE_CURRENT_SOURCE_DIR}/conj.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/copysign.cpp ${CMAKE_CURRENT_SOURCE_DIR}/cos.cpp ${CMAKE_CURRENT_SOURCE_DIR}/cosh.cpp ${CMAKE_CURRENT_SOURCE_DIR}/div.cpp @@ -48,6 +50,8 @@ if(NOT _use_onemkl_interfaces) ${CMAKE_CURRENT_SOURCE_DIR}/fmin.cpp ${CMAKE_CURRENT_SOURCE_DIR}/fmod.cpp ${CMAKE_CURRENT_SOURCE_DIR}/hypot.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/i0.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/inv.cpp ${CMAKE_CURRENT_SOURCE_DIR}/ln.cpp ${CMAKE_CURRENT_SOURCE_DIR}/log10.cpp ${CMAKE_CURRENT_SOURCE_DIR}/log1p.cpp diff --git a/dpnp/backend/extensions/vm/arg.cpp b/dpnp/backend/extensions/vm/arg.cpp new file mode 100644 index 000000000000..174f61c72e7c --- /dev/null +++ b/dpnp/backend/extensions/vm/arg.cpp @@ -0,0 +1,135 @@ +//***************************************************************************** +// 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 "dpctl4pybind11.hpp" + +#include "arg.hpp" +#include "common.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::arg 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, double>, + td_ns::TypeMapResultEntry, float>, + td_ns::DefaultResultEntry>::result_type; +}; + +template +static sycl::event arg_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::arg(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(arg); +} // namespace impl + +void init_arg(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 arg_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("_arg", arg_pyapi, + "Call `arg` function from OneMKL VM library to compute " + "the argument of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto arg_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_arg_to_call", arg_need_to_call_pyapi, + "Check input arguments to answer if `arg` 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/arg.hpp b/dpnp/backend/extensions/vm/arg.hpp new file mode 100644 index 000000000000..fe9bbc9bff5b --- /dev/null +++ b/dpnp/backend/extensions/vm/arg.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_arg(py::module_ m); +} // namespace dpnp::extensions::vm diff --git a/dpnp/backend/extensions/vm/copysign.cpp b/dpnp/backend/extensions/vm/copysign.cpp new file mode 100644 index 000000000000..7a703124f641 --- /dev/null +++ b/dpnp/backend/extensions/vm/copysign.cpp @@ -0,0 +1,166 @@ +//***************************************************************************** +// 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 "common.hpp" +#include "copysign.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::copysign function. + * + * @tparam T Type of input vectors `a` and `b` and of result vector `y`. + */ +template +struct OutputType +{ + using value_type = typename std::disjunction< + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::DefaultResultEntry>::result_type; +}; + +template +static sycl::event copysign_contig_impl(sycl::queue &exec_q, + std::size_t in_n, + const char *in_a, + py::ssize_t a_offset, + const char *in_b, + py::ssize_t b_offset, + char *out_y, + py::ssize_t out_offset, + const std::vector &depends) +{ + tu_ns::validate_type_for_device(exec_q); + tu_ns::validate_type_for_device(exec_q); + + if ((a_offset != 0) || (b_offset != 0) || (out_offset != 0)) { + throw std::runtime_error("Arrays offsets have to be equals to 0"); + } + + std::int64_t n = static_cast(in_n); + const T1 *a = reinterpret_cast(in_a); + const T2 *b = reinterpret_cast(in_b); + + using resTy = typename OutputType::value_type; + resTy *y = reinterpret_cast(out_y); + + return mkl_vm::copysign( + exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing 1st input vector of size n + b, // pointer `b` containing 2nd input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +using ew_cmn_ns::binary_contig_impl_fn_ptr_t; +using ew_cmn_ns::binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t; +using ew_cmn_ns::binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t; +using ew_cmn_ns::binary_strided_impl_fn_ptr_t; + +static int output_typeid_vector[td_ns::num_types][td_ns::num_types]; +static binary_contig_impl_fn_ptr_t contig_dispatch_vector[td_ns::num_types] + [td_ns::num_types]; + +MACRO_POPULATE_DISPATCH_TABLES(copysign); +} // namespace impl + +void init_copysign(py::module_ m) +{ + using arrayT = dpctl::tensor::usm_ndarray; + using event_vecT = std::vector; + + impl::populate_dispatch_tables(); + using impl::contig_dispatch_vector; + using impl::output_typeid_vector; + + auto copysign_pyapi = [&](sycl::queue &exec_q, const arrayT &src1, + const arrayT &src2, const arrayT &dst, + const event_vecT &depends = {}) { + return py_int::py_binary_ufunc( + src1, src2, dst, exec_q, depends, output_typeid_vector, + contig_dispatch_vector, + // no support of strided implementation in OneMKL + td_ns::NullPtrTable{}, + // no support of C-contig row with broadcasting in OneMKL + td_ns::NullPtrTable< + impl:: + binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t>{}, + td_ns::NullPtrTable< + impl:: + binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t>{}); + }; + m.def( + "_copysign", copysign_pyapi, + "Call `copysign` function from OneMKL VM library to return `dst` with " + "the elements of `src1` with the sign changed to match the sign " + "of the corresponding elements of `src2`", + py::arg("sycl_queue"), py::arg("src1"), py::arg("src2"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto copysign_need_to_call_pyapi = [&](sycl::queue &exec_q, + const arrayT &src1, + const arrayT &src2, + const arrayT &dst) { + return py_internal::need_to_call_binary_ufunc(exec_q, src1, src2, dst, + output_typeid_vector, + contig_dispatch_vector); + }; + m.def("_mkl_copysign_to_call", copysign_need_to_call_pyapi, + "Check input arguments to answer if `copysign` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src1"), py::arg("src2"), + py::arg("dst")); +} +} // namespace dpnp::extensions::vm diff --git a/dpnp/backend/extensions/vm/copysign.hpp b/dpnp/backend/extensions/vm/copysign.hpp new file mode 100644 index 000000000000..c52d1bee75ae --- /dev/null +++ b/dpnp/backend/extensions/vm/copysign.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_copysign(py::module_ m); +} // namespace dpnp::extensions::vm diff --git a/dpnp/backend/extensions/vm/i0.cpp b/dpnp/backend/extensions/vm/i0.cpp new file mode 100644 index 000000000000..f9d910d85543 --- /dev/null +++ b/dpnp/backend/extensions/vm/i0.cpp @@ -0,0 +1,136 @@ +//***************************************************************************** +// 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 "dpctl4pybind11.hpp" + +#include "common.hpp" +#include "i0.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::i0 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 i0_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::i0(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(i0); +} // namespace impl + +void init_i0(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 i0_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("_i0", i0_pyapi, + "Call `i0` function from OneMKL VM library to compute the " + "element-wise regular modified cylindrical Bessel function " + "of order 0", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto i0_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_i0_to_call", i0_need_to_call_pyapi, + "Check input arguments to answer if `i0` 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/i0.hpp b/dpnp/backend/extensions/vm/i0.hpp new file mode 100644 index 000000000000..296a54bda60b --- /dev/null +++ b/dpnp/backend/extensions/vm/i0.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_i0(py::module_ m); +} // namespace dpnp::extensions::vm diff --git a/dpnp/backend/extensions/vm/inv.cpp b/dpnp/backend/extensions/vm/inv.cpp new file mode 100644 index 000000000000..45cb07d3a7c2 --- /dev/null +++ b/dpnp/backend/extensions/vm/inv.cpp @@ -0,0 +1,136 @@ +//***************************************************************************** +// 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 "dpctl4pybind11.hpp" + +#include "common.hpp" +#include "inv.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::inv 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 inv_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::inv(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(inv); +} // namespace impl + +void init_inv(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 inv_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("_inv", inv_pyapi, + "Call `inv` function from OneMKL VM library to compute " + "the element-wise multiplicative inverse (or reciprocal) " + "of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto inv_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_inv_to_call", inv_need_to_call_pyapi, + "Check input arguments to answer if `inv` 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/inv.hpp b/dpnp/backend/extensions/vm/inv.hpp new file mode 100644 index 000000000000..5f967974ca70 --- /dev/null +++ b/dpnp/backend/extensions/vm/inv.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_inv(py::module_ m); +} // namespace dpnp::extensions::vm diff --git a/dpnp/backend/extensions/vm/vm_py.cpp b/dpnp/backend/extensions/vm/vm_py.cpp index c5a8d39c9b67..ad52e499419b 100644 --- a/dpnp/backend/extensions/vm/vm_py.cpp +++ b/dpnp/backend/extensions/vm/vm_py.cpp @@ -32,6 +32,7 @@ #include "acos.hpp" #include "acosh.hpp" #include "add.hpp" +#include "arg.hpp" #include "asin.hpp" #include "asinh.hpp" #include "atan.hpp" @@ -40,6 +41,7 @@ #include "cbrt.hpp" #include "ceil.hpp" #include "conj.hpp" +#include "copysign.hpp" #include "cos.hpp" #include "cosh.hpp" #include "div.hpp" @@ -51,6 +53,8 @@ #include "fmin.hpp" #include "fmod.hpp" #include "hypot.hpp" +#include "i0.hpp" +#include "inv.hpp" #include "ln.hpp" #include "log10.hpp" #include "log1p.hpp" @@ -80,6 +84,7 @@ PYBIND11_MODULE(_vm_impl, m) vm_ns::init_acos(m); vm_ns::init_acosh(m); vm_ns::init_add(m); + vm_ns::init_arg(m); vm_ns::init_asin(m); vm_ns::init_asinh(m); vm_ns::init_atan(m); @@ -88,6 +93,7 @@ PYBIND11_MODULE(_vm_impl, m) vm_ns::init_cbrt(m); vm_ns::init_ceil(m); vm_ns::init_conj(m); + vm_ns::init_copysign(m); vm_ns::init_cos(m); vm_ns::init_cosh(m); vm_ns::init_div(m); @@ -99,6 +105,8 @@ PYBIND11_MODULE(_vm_impl, m) vm_ns::init_fmin(m); vm_ns::init_fmod(m); vm_ns::init_hypot(m); + vm_ns::init_i0(m); + vm_ns::init_inv(m); vm_ns::init_ln(m); vm_ns::init_log10(m); vm_ns::init_log1p(m); diff --git a/dpnp/dpnp_algo/dpnp_elementwise_common.py b/dpnp/dpnp_algo/dpnp_elementwise_common.py index 0e232e0a2bcc..519e6cec55ee 100644 --- a/dpnp/dpnp_algo/dpnp_elementwise_common.py +++ b/dpnp/dpnp_algo/dpnp_elementwise_common.py @@ -497,12 +497,16 @@ def __init__( 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, deg=False, out=None, order="K"): @@ -563,12 +567,16 @@ def __init__( 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, order="K"): diff --git a/dpnp/dpnp_iface_mathematical.py b/dpnp/dpnp_iface_mathematical.py index 3613c9bffff6..c5f890837216 100644 --- a/dpnp/dpnp_iface_mathematical.py +++ b/dpnp/dpnp_iface_mathematical.py @@ -561,6 +561,8 @@ def _process_ediff1d_args(arg, arg_name, ary_dtype, ary_sycl_queue, usm_type): ti._angle_result_type, ti._angle, _ANGLE_DOCSTRING, + mkl_fn_to_call="_mkl_arg_to_call", + mkl_impl_fn="_arg", ) @@ -872,6 +874,8 @@ def clip(a, /, min=None, max=None, *, out=None, order="K", **kwargs): ti._copysign_result_type, ti._copysign, _COPYSIGN_DOCSTRING, + mkl_fn_to_call="_mkl_copysign_to_call", + mkl_impl_fn="_copysign", ) @@ -2677,6 +2681,8 @@ def gradient(f, *varargs, axis=None, edge_order=1): ufi._i0_result_type, ufi._i0, _I0_DOCSTRING, + mkl_fn_to_call="_mkl_i0_to_call", + mkl_impl_fn="_i0", ) @@ -3963,6 +3969,8 @@ def real_if_close(a, tol=100): ti._remainder, _REMAINDER_DOCSTRING, binary_inplace_fn=ti._remainder_inplace, + # mkl_vm::remainder() isn't implemented, because it follows C's modulo + # operator, but Python's one is expected acc to Python Array API spec ) mod = remainder diff --git a/dpnp/dpnp_iface_trigonometric.py b/dpnp/dpnp_iface_trigonometric.py index fff45f92f96b..3207ab1c3986 100644 --- a/dpnp/dpnp_iface_trigonometric.py +++ b/dpnp/dpnp_iface_trigonometric.py @@ -2102,6 +2102,8 @@ def logsumexp(x, /, *, axis=None, dtype=None, keepdims=False, out=None): ti._reciprocal_result_type, ti._reciprocal, _RECIPROCAL_DOCSTRING, + mkl_fn_to_call="_mkl_inv_to_call", + mkl_impl_fn="_inv", )