From 31573f9e6122a02924fb5ad5834109f7fec381eb Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Tue, 13 May 2025 12:00:41 +0200 Subject: [PATCH 1/9] Add mkl_vm::arg() extnesion to be used in dpnp.angle --- dpnp/backend/extensions/vm/CMakeLists.txt | 1 + dpnp/backend/extensions/vm/arg.cpp | 135 ++++++++++++++++++++++ dpnp/backend/extensions/vm/arg.hpp | 35 ++++++ dpnp/backend/extensions/vm/vm_py.cpp | 2 + dpnp/dpnp_algo/dpnp_elementwise_common.py | 4 + dpnp/dpnp_iface_mathematical.py | 2 + 6 files changed, 179 insertions(+) create mode 100644 dpnp/backend/extensions/vm/arg.cpp create mode 100644 dpnp/backend/extensions/vm/arg.hpp diff --git a/dpnp/backend/extensions/vm/CMakeLists.txt b/dpnp/backend/extensions/vm/CMakeLists.txt index 60886b3a2ce8..4d15b3444a91 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 diff --git a/dpnp/backend/extensions/vm/arg.cpp b/dpnp/backend/extensions/vm/arg.cpp new file mode 100644 index 000000000000..441d308ec7f4 --- /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 inverse tangent 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/vm_py.cpp b/dpnp/backend/extensions/vm/vm_py.cpp index c5a8d39c9b67..d53aa578573a 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" @@ -80,6 +81,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); diff --git a/dpnp/dpnp_algo/dpnp_elementwise_common.py b/dpnp/dpnp_algo/dpnp_elementwise_common.py index 0e232e0a2bcc..7f6d49820a14 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"): diff --git a/dpnp/dpnp_iface_mathematical.py b/dpnp/dpnp_iface_mathematical.py index 3613c9bffff6..d42623213221 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", ) From 6c9037df57ca460c37628c1b0e04d1930476f025 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Tue, 13 May 2025 12:39:36 +0200 Subject: [PATCH 2/9] Add mkl_vm::copysign() extnesion to be used in dpnp.copysign --- dpnp/backend/extensions/vm/CMakeLists.txt | 1 + dpnp/backend/extensions/vm/copysign.cpp | 167 ++++++++++++++++++++++ dpnp/backend/extensions/vm/copysign.hpp | 35 +++++ dpnp/backend/extensions/vm/vm_py.cpp | 2 + dpnp/dpnp_iface_mathematical.py | 2 + 5 files changed, 207 insertions(+) create mode 100644 dpnp/backend/extensions/vm/copysign.cpp create mode 100644 dpnp/backend/extensions/vm/copysign.hpp diff --git a/dpnp/backend/extensions/vm/CMakeLists.txt b/dpnp/backend/extensions/vm/CMakeLists.txt index 4d15b3444a91..2996a61bd33f 100644 --- a/dpnp/backend/extensions/vm/CMakeLists.txt +++ b/dpnp/backend/extensions/vm/CMakeLists.txt @@ -38,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 diff --git a/dpnp/backend/extensions/vm/copysign.cpp b/dpnp/backend/extensions/vm/copysign.cpp new file mode 100644 index 000000000000..69e6fecb4629 --- /dev/null +++ b/dpnp/backend/extensions/vm/copysign.cpp @@ -0,0 +1,167 @@ +//***************************************************************************** +// 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` of " + "elements containing the next representable floating-point values " + "following the values from the elements of `src1` in the direction 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/vm_py.cpp b/dpnp/backend/extensions/vm/vm_py.cpp index d53aa578573a..d107748ec04d 100644 --- a/dpnp/backend/extensions/vm/vm_py.cpp +++ b/dpnp/backend/extensions/vm/vm_py.cpp @@ -41,6 +41,7 @@ #include "cbrt.hpp" #include "ceil.hpp" #include "conj.hpp" +#include "copysign.hpp" #include "cos.hpp" #include "cosh.hpp" #include "div.hpp" @@ -90,6 +91,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); diff --git a/dpnp/dpnp_iface_mathematical.py b/dpnp/dpnp_iface_mathematical.py index d42623213221..07288cf7ec85 100644 --- a/dpnp/dpnp_iface_mathematical.py +++ b/dpnp/dpnp_iface_mathematical.py @@ -874,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", ) From 5682a942d947b475571ffa0995314d2ee6a2ed02 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Tue, 13 May 2025 13:29:41 +0200 Subject: [PATCH 3/9] Add mkl_vm::i0() extnesion to be used in dpnp.i0 --- dpnp/backend/extensions/vm/CMakeLists.txt | 1 + dpnp/backend/extensions/vm/i0.cpp | 136 ++++++++++++++++++++++ dpnp/backend/extensions/vm/i0.hpp | 35 ++++++ dpnp/backend/extensions/vm/vm_py.cpp | 2 + dpnp/dpnp_algo/dpnp_elementwise_common.py | 4 + dpnp/dpnp_iface_mathematical.py | 2 + 6 files changed, 180 insertions(+) create mode 100644 dpnp/backend/extensions/vm/i0.cpp create mode 100644 dpnp/backend/extensions/vm/i0.hpp diff --git a/dpnp/backend/extensions/vm/CMakeLists.txt b/dpnp/backend/extensions/vm/CMakeLists.txt index 2996a61bd33f..e25d00f1ce88 100644 --- a/dpnp/backend/extensions/vm/CMakeLists.txt +++ b/dpnp/backend/extensions/vm/CMakeLists.txt @@ -50,6 +50,7 @@ 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}/ln.cpp ${CMAKE_CURRENT_SOURCE_DIR}/log10.cpp ${CMAKE_CURRENT_SOURCE_DIR}/log1p.cpp diff --git a/dpnp/backend/extensions/vm/i0.cpp b/dpnp/backend/extensions/vm/i0.cpp new file mode 100644 index 000000000000..a5f4ee965de1 --- /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::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 inverse tangent of vector elements", + 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/vm_py.cpp b/dpnp/backend/extensions/vm/vm_py.cpp index d107748ec04d..69ad02c3f20d 100644 --- a/dpnp/backend/extensions/vm/vm_py.cpp +++ b/dpnp/backend/extensions/vm/vm_py.cpp @@ -53,6 +53,7 @@ #include "fmin.hpp" #include "fmod.hpp" #include "hypot.hpp" +#include "i0.hpp" #include "ln.hpp" #include "log10.hpp" #include "log1p.hpp" @@ -103,6 +104,7 @@ 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_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 7f6d49820a14..519e6cec55ee 100644 --- a/dpnp/dpnp_algo/dpnp_elementwise_common.py +++ b/dpnp/dpnp_algo/dpnp_elementwise_common.py @@ -567,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 07288cf7ec85..58ecb1bf9d13 100644 --- a/dpnp/dpnp_iface_mathematical.py +++ b/dpnp/dpnp_iface_mathematical.py @@ -2681,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", ) From eb9fc5c7239c8ec5674588d5c15211b9869356ec Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Tue, 13 May 2025 13:44:21 +0200 Subject: [PATCH 4/9] Add mkl_vm::inv() extnesion to be used in dpnp.reciprocal --- dpnp/backend/extensions/vm/CMakeLists.txt | 1 + dpnp/backend/extensions/vm/inv.cpp | 136 ++++++++++++++++++++++ dpnp/backend/extensions/vm/inv.hpp | 35 ++++++ dpnp/backend/extensions/vm/vm_py.cpp | 2 + dpnp/dpnp_iface_trigonometric.py | 2 + 5 files changed, 176 insertions(+) create mode 100644 dpnp/backend/extensions/vm/inv.cpp create mode 100644 dpnp/backend/extensions/vm/inv.hpp diff --git a/dpnp/backend/extensions/vm/CMakeLists.txt b/dpnp/backend/extensions/vm/CMakeLists.txt index e25d00f1ce88..9ad513ccbff6 100644 --- a/dpnp/backend/extensions/vm/CMakeLists.txt +++ b/dpnp/backend/extensions/vm/CMakeLists.txt @@ -51,6 +51,7 @@ if(NOT _use_onemkl_interfaces) ${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/inv.cpp b/dpnp/backend/extensions/vm/inv.cpp new file mode 100644 index 000000000000..599edfcb86a1 --- /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::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 inverse tangent 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 69ad02c3f20d..ad52e499419b 100644 --- a/dpnp/backend/extensions/vm/vm_py.cpp +++ b/dpnp/backend/extensions/vm/vm_py.cpp @@ -54,6 +54,7 @@ #include "fmod.hpp" #include "hypot.hpp" #include "i0.hpp" +#include "inv.hpp" #include "ln.hpp" #include "log10.hpp" #include "log1p.hpp" @@ -105,6 +106,7 @@ PYBIND11_MODULE(_vm_impl, 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_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", ) From a9c2ca72b6e3e260e86a6f34d9dce043592b2073 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Tue, 13 May 2025 16:31:09 +0200 Subject: [PATCH 5/9] Add a note why mkl_vm::remainder() is not intended to be used --- dpnp/dpnp_iface_mathematical.py | 2 ++ 1 file changed, 2 insertions(+) diff --git a/dpnp/dpnp_iface_mathematical.py b/dpnp/dpnp_iface_mathematical.py index 58ecb1bf9d13..c5f890837216 100644 --- a/dpnp/dpnp_iface_mathematical.py +++ b/dpnp/dpnp_iface_mathematical.py @@ -3969,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 From 869adfffe8def8839c32f256ece2040b9ed49a71 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Tue, 13 May 2025 21:46:11 +0200 Subject: [PATCH 6/9] Do not use sycl::half due to precision reason --- dpnp/backend/extensions/vm/i0.cpp | 1 - dpnp/backend/extensions/vm/inv.cpp | 1 - 2 files changed, 2 deletions(-) diff --git a/dpnp/backend/extensions/vm/i0.cpp b/dpnp/backend/extensions/vm/i0.cpp index a5f4ee965de1..38d65fda1d66 100644 --- a/dpnp/backend/extensions/vm/i0.cpp +++ b/dpnp/backend/extensions/vm/i0.cpp @@ -65,7 +65,6 @@ struct OutputType using value_type = typename std::disjunction, td_ns::TypeMapResultEntry, - td_ns::TypeMapResultEntry, td_ns::DefaultResultEntry>::result_type; }; diff --git a/dpnp/backend/extensions/vm/inv.cpp b/dpnp/backend/extensions/vm/inv.cpp index 599edfcb86a1..3c7b5d4071a3 100644 --- a/dpnp/backend/extensions/vm/inv.cpp +++ b/dpnp/backend/extensions/vm/inv.cpp @@ -65,7 +65,6 @@ struct OutputType using value_type = typename std::disjunction, td_ns::TypeMapResultEntry, - td_ns::TypeMapResultEntry, td_ns::DefaultResultEntry>::result_type; }; From 16053147e293d6a25a1857139c49bf5eb2d1855c Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Wed, 14 May 2025 00:04:34 +0200 Subject: [PATCH 7/9] Populate changelog --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) 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 From fa4baa18c2a04df284b2c8485972b6ea2d8588bf Mon Sep 17 00:00:00 2001 From: Anton <100830759+antonwolfy@users.noreply.github.com> Date: Wed, 14 May 2025 11:22:30 +0200 Subject: [PATCH 8/9] Apply suggestions from code review Co-authored-by: Vahid Tavanashad <120411540+vtavana@users.noreply.github.com> --- dpnp/backend/extensions/vm/arg.cpp | 2 +- dpnp/backend/extensions/vm/i0.cpp | 2 +- dpnp/backend/extensions/vm/inv.cpp | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/dpnp/backend/extensions/vm/arg.cpp b/dpnp/backend/extensions/vm/arg.cpp index 441d308ec7f4..174f61c72e7c 100644 --- a/dpnp/backend/extensions/vm/arg.cpp +++ b/dpnp/backend/extensions/vm/arg.cpp @@ -118,7 +118,7 @@ void init_arg(py::module_ m) }; m.def("_arg", arg_pyapi, "Call `arg` function from OneMKL VM library to compute " - "the inverse tangent of vector elements", + "the argument of vector elements", py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), py::arg("depends") = py::list()); diff --git a/dpnp/backend/extensions/vm/i0.cpp b/dpnp/backend/extensions/vm/i0.cpp index 38d65fda1d66..d4ea6f392684 100644 --- a/dpnp/backend/extensions/vm/i0.cpp +++ b/dpnp/backend/extensions/vm/i0.cpp @@ -118,7 +118,7 @@ void init_i0(py::module_ m) }; m.def("_i0", i0_pyapi, "Call `i0` function from OneMKL VM library to compute " - "the inverse tangent of vector elements", + "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()); diff --git a/dpnp/backend/extensions/vm/inv.cpp b/dpnp/backend/extensions/vm/inv.cpp index 3c7b5d4071a3..68952d396573 100644 --- a/dpnp/backend/extensions/vm/inv.cpp +++ b/dpnp/backend/extensions/vm/inv.cpp @@ -118,7 +118,7 @@ void init_inv(py::module_ m) }; m.def("_inv", inv_pyapi, "Call `inv` function from OneMKL VM library to compute " - "the inverse tangent of vector elements", + "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()); From 957233275212f293d5f204d10160e6177ea08e1a Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Wed, 14 May 2025 11:27:42 +0200 Subject: [PATCH 9/9] Update description of _copysign in vm extension --- dpnp/backend/extensions/vm/copysign.cpp | 7 +++---- dpnp/backend/extensions/vm/i0.cpp | 5 +++-- dpnp/backend/extensions/vm/inv.cpp | 3 ++- 3 files changed, 8 insertions(+), 7 deletions(-) diff --git a/dpnp/backend/extensions/vm/copysign.cpp b/dpnp/backend/extensions/vm/copysign.cpp index 69e6fecb4629..7a703124f641 100644 --- a/dpnp/backend/extensions/vm/copysign.cpp +++ b/dpnp/backend/extensions/vm/copysign.cpp @@ -143,10 +143,9 @@ void init_copysign(py::module_ m) }; m.def( "_copysign", copysign_pyapi, - "Call `copysign` function from OneMKL VM library to return `dst` of " - "elements containing the next representable floating-point values " - "following the values from the elements of `src1` in the direction of " - "the corresponding elements of `src2`", + "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()); diff --git a/dpnp/backend/extensions/vm/i0.cpp b/dpnp/backend/extensions/vm/i0.cpp index d4ea6f392684..f9d910d85543 100644 --- a/dpnp/backend/extensions/vm/i0.cpp +++ b/dpnp/backend/extensions/vm/i0.cpp @@ -117,8 +117,9 @@ void init_i0(py::module_ m) 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", + "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()); diff --git a/dpnp/backend/extensions/vm/inv.cpp b/dpnp/backend/extensions/vm/inv.cpp index 68952d396573..45cb07d3a7c2 100644 --- a/dpnp/backend/extensions/vm/inv.cpp +++ b/dpnp/backend/extensions/vm/inv.cpp @@ -118,7 +118,8 @@ void init_inv(py::module_ m) }; m.def("_inv", inv_pyapi, "Call `inv` function from OneMKL VM library to compute " - "the element-wise multiplicative inverse (or reciprocal) of vector elements", + "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());