Rework implementation of dpnp.fmod function

dpnp/backend/extensions/ufunc/CMakeLists.txt

@@ -26,6 +26,7 @@
 set(_elementwise_sources
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/common.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fabs.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmod.cpp
 )
 
 set(python_module_name _ufunc_impl)

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

@@ -26,6 +26,7 @@
 #include <pybind11/pybind11.h>
 
 #include "fabs.hpp"
+#include "fmod.hpp"
 
 namespace py = pybind11;
 
@@ -37,5 +38,6 @@ namespace dpnp::extensions::ufunc
 void init_elementwise_functions(py::module_ m)
 {
     init_fabs(m);
+    init_fmod(m);
 }
 } // namespace dpnp::extensions::ufunc

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

@@ -0,0 +1,193 @@
+//*****************************************************************************
+// Copyright (c) 2024, 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 <sycl/sycl.hpp>
+
+#include "dpctl4pybind11.hpp"
+
+#include "fmod.hpp"
+#include "kernels/elementwise_functions/fmod.hpp"
+#include "populate.hpp"
+
+// include a local copy of elementwise common header from dpctl tensor:
+// dpctl/tensor/libtensor/source/elementwise_functions/elementwise_functions.hpp
+// TODO: replace by including dpctl header once available
+#include "../../elementwise_functions/elementwise_functions.hpp"
+
+// dpctl tensor headers
+#include "kernels/elementwise_functions/common.hpp"
+#include "utils/type_dispatch.hpp"
+
+namespace py = pybind11;
+
+namespace dpnp::extensions::ufunc
+{
+namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
+namespace py_int = dpnp::extensions::py_internal;
+namespace td_ns = dpctl::tensor::type_dispatch;
+
+using ew_cmn_ns::unary_contig_impl_fn_ptr_t;
+using ew_cmn_ns::unary_strided_impl_fn_ptr_t;
+
+namespace impl
+{
+/**
+ * @brief A factory to define pairs of supported types for which
+ * sycl::fmod<T> function is available.
+ *
+ * @tparam T1 Type of input vectors `a`
+ * @tparam T2 Type of input vectors `b`
+ */
+template <typename T1, typename T2>
+struct OutputType
+{
+    using value_type = typename std::disjunction<
+        td_ns::BinaryTypeMapResultEntry<T1, bool, T2, bool, std::int8_t>,
+        td_ns::BinaryTypeMapResultEntry<T1,
+                                        std::uint8_t,
+                                        T2,
+                                        std::uint8_t,
+                                        std::uint8_t>,
+        td_ns::BinaryTypeMapResultEntry<T1,
+                                        std::int8_t,
+                                        T2,
+                                        std::int8_t,
+                                        std::int8_t>,
+        td_ns::BinaryTypeMapResultEntry<T1,
+                                        std::uint16_t,
+                                        T2,
+                                        std::uint16_t,
+                                        std::uint16_t>,
+        td_ns::BinaryTypeMapResultEntry<T1,
+                                        std::int16_t,
+                                        T2,
+                                        std::int16_t,
+                                        std::int16_t>,
+        td_ns::BinaryTypeMapResultEntry<T1,
+                                        std::uint32_t,
+                                        T2,
+                                        std::uint32_t,
+                                        std::uint32_t>,
+        td_ns::BinaryTypeMapResultEntry<T1,
+                                        std::int32_t,
+                                        T2,
+                                        std::int32_t,
+                                        std::int32_t>,
+        td_ns::BinaryTypeMapResultEntry<T1,
+                                        std::uint64_t,
+                                        T2,
+                                        std::uint64_t,
+                                        std::uint64_t>,
+        td_ns::BinaryTypeMapResultEntry<T1,
+                                        std::int64_t,
+                                        T2,
+                                        std::int64_t,
+                                        std::int64_t>,
+        td_ns::BinaryTypeMapResultEntry<T1,
+                                        sycl::half,
+                                        T2,
+                                        sycl::half,
+                                        sycl::half>,
+        td_ns::BinaryTypeMapResultEntry<T1, float, T2, float, float>,
+        td_ns::BinaryTypeMapResultEntry<T1, double, T2, double, double>,
+        td_ns::DefaultResultEntry<void>>::result_type;
+};
+
+using dpnp::kernels::fmod::FmodFunctor;
+
+template <typename argT1,
+          typename argT2,
+          typename resT,
+          unsigned int vec_sz = 4,
+          unsigned int n_vecs = 2,
+          bool enable_sg_loadstore = true>
+using ContigFunctor =
+    ew_cmn_ns::BinaryContigFunctor<argT1,
+                                   argT2,
+                                   resT,
+                                   FmodFunctor<argT1, argT2, resT>,
+                                   vec_sz,
+                                   n_vecs,
+                                   enable_sg_loadstore>;
+
+template <typename argT1, typename argT2, typename resT, typename IndexerT>
+using StridedFunctor =
+    ew_cmn_ns::BinaryStridedFunctor<argT1,
+                                    argT2,
+                                    resT,
+                                    IndexerT,
+                                    FmodFunctor<argT1, argT2, resT>>;
+
+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 binary_contig_impl_fn_ptr_t fmod_contig_dispatch_table[td_ns::num_types]
+                                                             [td_ns::num_types];
+static int fmod_output_typeid_table[td_ns::num_types][td_ns::num_types];
+static binary_strided_impl_fn_ptr_t
+    fmod_strided_dispatch_table[td_ns::num_types][td_ns::num_types];
+
+MACRO_POPULATE_DISPATCH_TABLES(fmod);
+} // namespace impl
+
+void init_fmod(py::module_ m)
+{
+    using arrayT = dpctl::tensor::usm_ndarray;
+    using event_vecT = std::vector<sycl::event>;
+    {
+        impl::populate_fmod_dispatch_tables();
+        using impl::fmod_contig_dispatch_table;
+        using impl::fmod_output_typeid_table;
+        using impl::fmod_strided_dispatch_table;
+
+        auto fmod_pyapi = [&](const arrayT &src1, const arrayT &src2,
+                              const arrayT &dst, sycl::queue &exec_q,
+                              const event_vecT &depends = {}) {
+            return py_int::py_binary_ufunc(
+                src1, src2, dst, exec_q, depends, fmod_output_typeid_table,
+                fmod_contig_dispatch_table, fmod_strided_dispatch_table,
+                // 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("_fmod", fmod_pyapi, "", py::arg("src1"), py::arg("src2"),
+              py::arg("dst"), py::arg("sycl_queue"),
+              py::arg("depends") = py::list());
+
+        auto fmod_result_type_pyapi = [&](const py::dtype &dtype1,
+                                          const py::dtype &dtype2) {
+            return py_int::py_binary_ufunc_result_type(
+                dtype1, dtype2, fmod_output_typeid_table);
+        };
+        m.def("_fmod_result_type", fmod_result_type_pyapi);
+    }
+}
+} // namespace dpnp::extensions::ufunc

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

@@ -0,0 +1,35 @@
+//*****************************************************************************
+// Copyright (c) 2024, 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 <pybind11/pybind11.h>
+
+namespace py = pybind11;
+
+namespace dpnp::extensions::ufunc
+{
+void init_fmod(py::module_ m);
+} // namespace dpnp::extensions::ufunc

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

@@ -26,7 +26,8 @@
 #pragma once
 
 /**
- * @brief A macro used to define factories and a populating universal functions.
+ * @brief A macro used to define factories and a populating unary universal
+ * functions.
  */
 #define MACRO_POPULATE_DISPATCH_VECTORS(__name__)                              \
     template <typename T1, typename T2, unsigned int vec_sz,                   \
@@ -120,3 +121,110 @@
             dvb3;                                                              \
         dvb3.populate_dispatch_vector(__name__##_output_typeid_vector);        \
     };
+
+/**
+ * @brief A macro used to define factories and a populating binary universal
+ * functions.
+ */
+#define MACRO_POPULATE_DISPATCH_TABLES(__name__)                               \
+    template <typename argT1, typename argT2, typename resT,                   \
+              unsigned int vec_sz, unsigned int n_vecs>                        \
+    class __name__##_contig_kernel;                                            \
+                                                                               \
+    template <typename argTy1, typename argTy2>                                \
+    sycl::event __name__##_contig_impl(                                        \
+        sycl::queue &exec_q, size_t nelems, const char *arg1_p,                \
+        py::ssize_t arg1_offset, const char *arg2_p, py::ssize_t arg2_offset,  \
+        char *res_p, py::ssize_t res_offset,                                   \
+        const std::vector<sycl::event> &depends = {})                          \
+    {                                                                          \
+        return ew_cmn_ns::binary_contig_impl<argTy1, argTy2, OutputType,       \
+                                             ContigFunctor,                    \
+                                             __name__##_contig_kernel>(        \
+            exec_q, nelems, arg1_p, arg1_offset, arg2_p, arg2_offset, res_p,   \
+            res_offset, depends);                                              \
+    }                                                                          \
+                                                                               \
+    template <typename fnT, typename T1, typename T2>                          \
+    struct ContigFactory                                                       \
+    {                                                                          \
+        fnT get()                                                              \
+        {                                                                      \
+            if constexpr (std::is_same_v<                                      \
+                              typename OutputType<T1, T2>::value_type, void>)  \
+            {                                                                  \
+                                                                               \
+                fnT fn = nullptr;                                              \
+                return fn;                                                     \
+            }                                                                  \
+            else {                                                             \
+                fnT fn = __name__##_contig_impl<T1, T2>;                       \
+                return fn;                                                     \
+            }                                                                  \
+        }                                                                      \
+    };                                                                         \
+                                                                               \
+    template <typename fnT, typename T1, typename T2>                          \
+    struct TypeMapFactory                                                      \
+    {                                                                          \
+        std::enable_if_t<std::is_same<fnT, int>::value, int> get()             \
+        {                                                                      \
+            using rT = typename OutputType<T1, T2>::value_type;                \
+            return td_ns::GetTypeid<rT>{}.get();                               \
+        }                                                                      \
+    };                                                                         \
+                                                                               \
+    template <typename T1, typename T2, typename resT, typename IndexerT>      \
+    class __name__##_strided_kernel;                                           \
+                                                                               \
+    template <typename argTy1, typename argTy2>                                \
+    sycl::event __name__##_strided_impl(                                       \
+        sycl::queue &exec_q, size_t nelems, int nd,                            \
+        const py::ssize_t *shape_and_strides, const char *arg1_p,              \
+        py::ssize_t arg1_offset, const char *arg2_p, py::ssize_t arg2_offset,  \
+        char *res_p, py::ssize_t res_offset,                                   \
+        const std::vector<sycl::event> &depends,                               \
+        const std::vector<sycl::event> &additional_depends)                    \
+    {                                                                          \
+        return ew_cmn_ns::binary_strided_impl<argTy1, argTy2, OutputType,      \
+                                              StridedFunctor,                  \
+                                              __name__##_strided_kernel>(      \
+            exec_q, nelems, nd, shape_and_strides, arg1_p, arg1_offset,        \
+            arg2_p, arg2_offset, res_p, res_offset, depends,                   \
+            additional_depends);                                               \
+    }                                                                          \
+                                                                               \
+    template <typename fnT, typename T1, typename T2>                          \
+    struct StridedFactory                                                      \
+    {                                                                          \
+        fnT get()                                                              \
+        {                                                                      \
+            if constexpr (std::is_same_v<                                      \
+                              typename OutputType<T1, T2>::value_type, void>)  \
+            {                                                                  \
+                fnT fn = nullptr;                                              \
+                return fn;                                                     \
+            }                                                                  \
+            else {                                                             \
+                fnT fn = __name__##_strided_impl<T1, T2>;                      \
+                return fn;                                                     \
+            }                                                                  \
+        }                                                                      \
+    };                                                                         \
+                                                                               \
+    void populate_##__name__##_dispatch_tables(void)                           \
+    {                                                                          \
+        td_ns::DispatchTableBuilder<binary_contig_impl_fn_ptr_t,               \
+                                    ContigFactory, td_ns::num_types>           \
+            dvb1;                                                              \
+        dvb1.populate_dispatch_table(__name__##_contig_dispatch_table);        \
+                                                                               \
+        td_ns::DispatchTableBuilder<binary_strided_impl_fn_ptr_t,              \
+                                    StridedFactory, td_ns::num_types>          \
+            dvb2;                                                              \
+        dvb2.populate_dispatch_table(__name__##_strided_dispatch_table);       \
+                                                                               \
+        td_ns::DispatchTableBuilder<int, TypeMapFactory, td_ns::num_types>     \
+            dvb3;                                                              \
+        dvb3.populate_dispatch_table(__name__##_output_typeid_table);          \
+    };

dpnp/backend/extensions/vm/CMakeLists.txt

@@ -43,6 +43,7 @@ set(_elementwise_sources
     ${CMAKE_CURRENT_SOURCE_DIR}/exp2.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/expm1.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/floor.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/fmod.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/hypot.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/ln.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/log10.cpp