Redesigned dpnp.modf function to be a part of ufunc and vm pybind11 extensions

CHANGELOG.md

@@ -28,6 +28,7 @@ Also, that release drops support for Python 3.9, making Python 3.10 the minimum
 * Changed the license from `BSD-2-Clause` to `BSD-3-Clause` [#2593](https://github.com/IntelPython/dpnp/pull/2593)
 * Defined explicit versions range of the Python interpreter which is needed during the build [#2634](https://github.com/IntelPython/dpnp/pull/2634)
 * Aligned documentation with NumPy and CuPy style by using short function names [#2633](https://github.com/IntelPython/dpnp/pull/2633)
+* Redesigned `dpnp.modf` function to be a part of `ufunc` and `vm` pybind11 extensions [#2654](https://github.com/IntelPython/dpnp/pull/2654)
 
 ### Deprecated
 

dpnp/backend/CMakeLists.txt

@@ -30,7 +30,6 @@
 set(DPNP_SRC
     kernels/dpnp_krnl_arraycreation.cpp
     kernels/dpnp_krnl_common.cpp
-    kernels/dpnp_krnl_mathematical.cpp
     kernels/dpnp_krnl_random.cpp
     kernels/dpnp_krnl_sorting.cpp
     src/dpnp_iface_fptr.cpp

dpnp/backend/extensions/ufunc/CMakeLists.txt

@@ -47,6 +47,7 @@ set(_elementwise_sources
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/lcm.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/ldexp.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/logaddexp2.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/modf.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/nan_to_num.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/radians.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/sinc.cpp

dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp

@@ -46,6 +46,7 @@
 #include "lcm.hpp"
 #include "ldexp.hpp"
 #include "logaddexp2.hpp"
+#include "modf.hpp"
 #include "nan_to_num.hpp"
 #include "radians.hpp"
 #include "sinc.hpp"
@@ -78,6 +79,7 @@ void init_elementwise_functions(py::module_ m)
     init_lcm(m);
     init_ldexp(m);
     init_logaddexp2(m);
+    init_modf(m);
     init_nan_to_num(m);
     init_radians(m);
     init_sinc(m);

dpnp/backend/extensions/ufunc/elementwise_functions/frexp.cpp

@@ -37,6 +37,7 @@
 
 #include "frexp.hpp"
 #include "kernels/elementwise_functions/frexp.hpp"
+#include "populate.hpp"
 
 // include a local copy of elementwise common header from dpctl tensor:
 // dpctl/tensor/libtensor/source/elementwise_functions/elementwise_functions.hpp
@@ -65,10 +66,9 @@ namespace td_int_ns = py_int::type_dispatch;
 namespace td_ns = dpctl::tensor::type_dispatch;
 
 using dpnp::kernels::frexp::FrexpFunctor;
-using ext::common::init_dispatch_vector;
 
 template <typename T>
-struct FrexpOutputType
+struct OutputType
 {
     using table_type = std::disjunction< // disjunction is C++17
                                          // feature, supported by DPC++
@@ -81,15 +81,13 @@ struct FrexpOutputType
     using value_type2 = typename table_type::result_type2;
 };
 
-// contiguous implementation
-
 template <typename argTy,
           typename resTy1 = argTy,
           typename resTy2 = argTy,
           std::uint8_t vec_sz = 4u,
           std::uint8_t n_vecs = 2u,
           bool enable_sg_loadstore = true>
-using FrexpContigFunctor =
+using ContigFunctor =
     ew_cmn_ns::UnaryTwoOutputsContigFunctor<argTy,
                                             resTy1,
                                             resTy2,
@@ -98,115 +96,14 @@ using FrexpContigFunctor =
                                             n_vecs,
                                             enable_sg_loadstore>;
 
-// strided implementation
-
 template <typename argTy, typename resTy1, typename resTy2, typename IndexerT>
-using FrexpStridedFunctor = ew_cmn_ns::UnaryTwoOutputsStridedFunctor<
+using StridedFunctor = ew_cmn_ns::UnaryTwoOutputsStridedFunctor<
     argTy,
     resTy1,
     resTy2,
     IndexerT,
     FrexpFunctor<argTy, resTy1, resTy2>>;
 
-template <typename T1,
-          typename T2,
-          typename T3,
-          unsigned int vec_sz,
-          unsigned int n_vecs>
-class frexp_contig_kernel;
-
-template <typename argTy>
-sycl::event frexp_contig_impl(sycl::queue &exec_q,
-                              size_t nelems,
-                              const char *arg_p,
-                              char *res1_p,
-                              char *res2_p,
-                              const std::vector<sycl::event> &depends = {})
-{
-    return ew_cmn_ns::unary_two_outputs_contig_impl<
-        argTy, FrexpOutputType, FrexpContigFunctor, frexp_contig_kernel>(
-        exec_q, nelems, arg_p, res1_p, res2_p, depends);
-}
-
-template <typename fnT, typename T>
-struct FrexpContigFactory
-{
-    fnT get()
-    {
-        if constexpr (std::is_same_v<typename FrexpOutputType<T>::value_type1,
-                                     void> ||
-                      std::is_same_v<typename FrexpOutputType<T>::value_type2,
-                                     void>)
-        {
-            fnT fn = nullptr;
-            return fn;
-        }
-        else {
-            fnT fn = frexp_contig_impl<T>;
-            return fn;
-        }
-    }
-};
-
-template <typename T1, typename T2, typename T3, typename T4>
-class frexp_strided_kernel;
-
-template <typename argTy>
-sycl::event
-    frexp_strided_impl(sycl::queue &exec_q,
-                       size_t nelems,
-                       int nd,
-                       const ssize_t *shape_and_strides,
-                       const char *arg_p,
-                       ssize_t arg_offset,
-                       char *res1_p,
-                       ssize_t res1_offset,
-                       char *res2_p,
-                       ssize_t res2_offset,
-                       const std::vector<sycl::event> &depends,
-                       const std::vector<sycl::event> &additional_depends)
-{
-    return ew_cmn_ns::unary_two_outputs_strided_impl<
-        argTy, FrexpOutputType, FrexpStridedFunctor, frexp_strided_kernel>(
-        exec_q, nelems, nd, shape_and_strides, arg_p, arg_offset, res1_p,
-        res1_offset, res2_p, res2_offset, depends, additional_depends);
-}
-
-template <typename fnT, typename T>
-struct FrexpStridedFactory
-{
-    fnT get()
-    {
-        if constexpr (std::is_same_v<typename FrexpOutputType<T>::value_type1,
-                                     void> ||
-                      std::is_same_v<typename FrexpOutputType<T>::value_type2,
-                                     void>)
-        {
-            fnT fn = nullptr;
-            return fn;
-        }
-        else {
-            fnT fn = frexp_strided_impl<T>;
-            return fn;
-        }
-    }
-};
-
-template <typename fnT, typename T>
-struct FrexpTypeMapFactory
-{
-    /*! @brief get typeid for output type of sycl::frexp(T x) */
-    std::enable_if_t<std::is_same<fnT, std::pair<int, int>>::value,
-                     std::pair<int, int>>
-        get()
-    {
-        using rT1 = typename FrexpOutputType<T>::value_type1;
-        using rT2 = typename FrexpOutputType<T>::value_type2;
-        return std::make_pair(td_ns::GetTypeid<rT1>{}.get(),
-                              td_ns::GetTypeid<rT2>{}.get());
-    }
-};
-
 using ew_cmn_ns::unary_two_outputs_contig_impl_fn_ptr_t;
 using ew_cmn_ns::unary_two_outputs_strided_impl_fn_ptr_t;
 
@@ -216,15 +113,7 @@ static std::pair<int, int> frexp_output_typeid_vector[td_ns::num_types];
 static unary_two_outputs_strided_impl_fn_ptr_t
     frexp_strided_dispatch_vector[td_ns::num_types];
 
-void populate_frexp_dispatch_vectors(void)
-{
-    init_dispatch_vector<unary_two_outputs_contig_impl_fn_ptr_t,
-                         FrexpContigFactory>(frexp_contig_dispatch_vector);
-    init_dispatch_vector<unary_two_outputs_strided_impl_fn_ptr_t,
-                         FrexpStridedFactory>(frexp_strided_dispatch_vector);
-    init_dispatch_vector<std::pair<int, int>, FrexpTypeMapFactory>(
-        frexp_output_typeid_vector);
-};
+MACRO_POPULATE_DISPATCH_2OUTS_VECTORS(frexp);
 } // namespace impl
 
 void init_frexp(py::module_ m)

dpnp/backend/extensions/ufunc/elementwise_functions/modf.cpp

@@ -0,0 +1,146 @@
+//*****************************************************************************
+// Copyright (c) 2025, Intel Corporation
+// All rights reserved.
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+// - Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// - Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+// - Neither the name of the copyright holder nor the names of its contributors
+//   may be used to endorse or promote products derived from this software
+//   without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
+// THE POSSIBILITY OF SUCH DAMAGE.
+//*****************************************************************************
+
+#include <cstdint>
+#include <type_traits>
+#include <utility>
+#include <vector>
+
+#include <sycl/sycl.hpp>
+
+#include "dpctl4pybind11.hpp"
+
+#include "kernels/elementwise_functions/modf.hpp"
+#include "modf.hpp"
+#include "populate.hpp"
+
+// include a local copy of elementwise common header from dpctl tensor:
+// dpctl/tensor/libtensor/source/elementwise_functions/elementwise_functions.hpp
+// TODO: replace by including dpctl header once available
+#include "../../elementwise_functions/elementwise_functions.hpp"
+
+#include "../../elementwise_functions/common.hpp"
+#include "../../elementwise_functions/type_dispatch_building.hpp"
+
+// utils extension header
+#include "ext/common.hpp"
+
+// dpctl tensor headers
+#include "kernels/elementwise_functions/common.hpp"
+#include "utils/type_dispatch.hpp"
+
+namespace dpnp::extensions::ufunc
+{
+namespace py = pybind11;
+namespace py_int = dpnp::extensions::py_internal;
+
+namespace impl
+{
+namespace ew_cmn_ns = dpnp::extensions::py_internal::elementwise_common;
+namespace td_int_ns = py_int::type_dispatch;
+namespace td_ns = dpctl::tensor::type_dispatch;
+
+using dpnp::kernels::modf::ModfFunctor;
+
+template <typename T>
+struct OutputType
+{
+    using table_type = std::disjunction< // disjunction is C++17
+                                         // feature, supported by DPC++
+        td_int_ns::TypeMapTwoResultsEntry<T, sycl::half>,
+        td_int_ns::TypeMapTwoResultsEntry<T, float>,
+        td_int_ns::TypeMapTwoResultsEntry<T, double>,
+        td_int_ns::DefaultTwoResultsEntry<void>>;
+    using value_type1 = typename table_type::result_type1;
+    using value_type2 = typename table_type::result_type2;
+};
+
+template <typename argTy,
+          typename resTy1 = argTy,
+          typename resTy2 = argTy,
+          std::uint8_t vec_sz = 4u,
+          std::uint8_t n_vecs = 2u,
+          bool enable_sg_loadstore = true>
+using ContigFunctor =
+    ew_cmn_ns::UnaryTwoOutputsContigFunctor<argTy,
+                                            resTy1,
+                                            resTy2,
+                                            ModfFunctor<argTy, resTy1, resTy2>,
+                                            vec_sz,
+                                            n_vecs,
+                                            enable_sg_loadstore>;
+
+template <typename argTy, typename resTy1, typename resTy2, typename IndexerT>
+using StridedFunctor = ew_cmn_ns::UnaryTwoOutputsStridedFunctor<
+    argTy,
+    resTy1,
+    resTy2,
+    IndexerT,
+    ModfFunctor<argTy, resTy1, resTy2>>;
+
+using ew_cmn_ns::unary_two_outputs_contig_impl_fn_ptr_t;
+using ew_cmn_ns::unary_two_outputs_strided_impl_fn_ptr_t;
+
+static unary_two_outputs_contig_impl_fn_ptr_t
+    modf_contig_dispatch_vector[td_ns::num_types];
+static std::pair<int, int> modf_output_typeid_vector[td_ns::num_types];
+static unary_two_outputs_strided_impl_fn_ptr_t
+    modf_strided_dispatch_vector[td_ns::num_types];
+
+MACRO_POPULATE_DISPATCH_2OUTS_VECTORS(modf);
+} // namespace impl
+
+void init_modf(py::module_ m)
+{
+    using arrayT = dpctl::tensor::usm_ndarray;
+    using event_vecT = std::vector<sycl::event>;
+    {
+        impl::populate_modf_dispatch_vectors();
+        using impl::modf_contig_dispatch_vector;
+        using impl::modf_output_typeid_vector;
+        using impl::modf_strided_dispatch_vector;
+
+        auto modf_pyapi = [&](const arrayT &src, const arrayT &dst1,
+                              const arrayT &dst2, sycl::queue &exec_q,
+                              const event_vecT &depends = {}) {
+            return py_int::py_unary_two_outputs_ufunc(
+                src, dst1, dst2, exec_q, depends, modf_output_typeid_vector,
+                modf_contig_dispatch_vector, modf_strided_dispatch_vector);
+        };
+        m.def("_modf", modf_pyapi, "", py::arg("src"), py::arg("dst1"),
+              py::arg("dst2"), py::arg("sycl_queue"),
+              py::arg("depends") = py::list());
+
+        auto modf_result_type_pyapi = [&](const py::dtype &dtype) {
+            return py_int::py_unary_two_outputs_ufunc_result_type(
+                dtype, modf_output_typeid_vector);
+        };
+        m.def("_modf_result_type", modf_result_type_pyapi);
+    }
+}
+} // namespace dpnp::extensions::ufunc

dpnp/backend/extensions/ufunc/elementwise_functions/modf.hpp

@@ -0,0 +1,38 @@
+//*****************************************************************************
+// Copyright (c) 2025, Intel Corporation
+// All rights reserved.
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+// - Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// - Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+// - Neither the name of the copyright holder nor the names of its contributors
+//   may be used to endorse or promote products derived from this software
+//   without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
+// THE POSSIBILITY OF SUCH DAMAGE.
+//*****************************************************************************
+
+#pragma once
+
+#include <pybind11/pybind11.h>
+
+namespace py = pybind11;
+
+namespace dpnp::extensions::ufunc
+{
+void init_modf(py::module_ m);
+} // namespace dpnp::extensions::ufunc