Rework implementation of dpnp.fmax and dpnp.fmin functions

doc/reference/ufunc.rst

@@ -105,10 +105,12 @@ Comparison functions
    dpnp.less_equal
    dpnp.not_equal
    dpnp.equal
+
    dpnp.logical_and
    dpnp.logical_or
    dpnp.logical_xor
    dpnp.logical_not
+
    dpnp.maximum
    dpnp.minimum
    dpnp.fmax

dpnp/backend/extensions/ufunc/CMakeLists.txt

@@ -26,6 +26,8 @@
 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/fmax.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmin.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmod.cpp
 )
 

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

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

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

@@ -0,0 +1,137 @@
+//*****************************************************************************
+// Copyright (c) 2024, Intel Corporation
+// All rights reserved.
+//
+// Redistribution and use in source and binary forms, with or without
+// maxification, 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 "fmax.hpp"
+#include "kernels/elementwise_functions/fmax.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 "kernels/elementwise_functions/maximum.hpp"
+#include "utils/type_dispatch.hpp"
+
+namespace py = pybind11;
+
+namespace dpnp::extensions::ufunc
+{
+namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
+namespace max_ns = dpctl::tensor::kernels::maximum;
+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
+{
+// Supports the same types table as for maximum function in dpctl
+template <typename T1, typename T2>
+using OutputType = max_ns::MaximumOutputType<T1, T2>;
+
+using dpnp::kernels::fmax::FmaxFunctor;
+
+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,
+                                   FmaxFunctor<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,
+                                    FmaxFunctor<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 fmax_contig_dispatch_table[td_ns::num_types]
+                                                             [td_ns::num_types];
+static int fmax_output_typeid_table[td_ns::num_types][td_ns::num_types];
+static binary_strided_impl_fn_ptr_t
+    fmax_strided_dispatch_table[td_ns::num_types][td_ns::num_types];
+
+MACRO_POPULATE_DISPATCH_TABLES(fmax);
+} // namespace impl
+
+void init_fmax(py::module_ m)
+{
+    using arrayT = dpctl::tensor::usm_ndarray;
+    using event_vecT = std::vector<sycl::event>;
+    {
+        impl::populate_fmax_dispatch_tables();
+        using impl::fmax_contig_dispatch_table;
+        using impl::fmax_output_typeid_table;
+        using impl::fmax_strided_dispatch_table;
+
+        auto fmax_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, fmax_output_typeid_table,
+                fmax_contig_dispatch_table, fmax_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("_fmax", fmax_pyapi, "", py::arg("src1"), py::arg("src2"),
+              py::arg("dst"), py::arg("sycl_queue"),
+              py::arg("depends") = py::list());
+
+        auto fmax_result_type_pyapi = [&](const py::dtype &dtype1,
+                                          const py::dtype &dtype2) {
+            return py_int::py_binary_ufunc_result_type(
+                dtype1, dtype2, fmax_output_typeid_table);
+        };
+        m.def("_fmax_result_type", fmax_result_type_pyapi);
+    }
+}
+} // namespace dpnp::extensions::ufunc

dpnp/backend/extensions/ufunc/elementwise_functions/fmax.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_fmax(py::module_ m);
+} // namespace dpnp::extensions::ufunc

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

@@ -0,0 +1,137 @@
+//*****************************************************************************
+// Copyright (c) 2024, Intel Corporation
+// All rights reserved.
+//
+// Redistribution and use in source and binary forms, with or without
+// maxification, 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 "fmin.hpp"
+#include "kernels/elementwise_functions/fmin.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 "kernels/elementwise_functions/minimum.hpp"
+#include "utils/type_dispatch.hpp"
+
+namespace py = pybind11;
+
+namespace dpnp::extensions::ufunc
+{
+namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
+namespace min_ns = dpctl::tensor::kernels::minimum;
+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
+{
+// Supports the same types table as for minimum function in dpctl
+template <typename T1, typename T2>
+using OutputType = min_ns::MinimumOutputType<T1, T2>;
+
+using dpnp::kernels::fmin::FminFunctor;
+
+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,
+                                   FminFunctor<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,
+                                    FminFunctor<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 fmin_contig_dispatch_table[td_ns::num_types]
+                                                             [td_ns::num_types];
+static int fmin_output_typeid_table[td_ns::num_types][td_ns::num_types];
+static binary_strided_impl_fn_ptr_t
+    fmin_strided_dispatch_table[td_ns::num_types][td_ns::num_types];
+
+MACRO_POPULATE_DISPATCH_TABLES(fmin);
+} // namespace impl
+
+void init_fmin(py::module_ m)
+{
+    using arrayT = dpctl::tensor::usm_ndarray;
+    using event_vecT = std::vector<sycl::event>;
+    {
+        impl::populate_fmin_dispatch_tables();
+        using impl::fmin_contig_dispatch_table;
+        using impl::fmin_output_typeid_table;
+        using impl::fmin_strided_dispatch_table;
+
+        auto fmin_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, fmin_output_typeid_table,
+                fmin_contig_dispatch_table, fmin_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("_fmin", fmin_pyapi, "", py::arg("src1"), py::arg("src2"),
+              py::arg("dst"), py::arg("sycl_queue"),
+              py::arg("depends") = py::list());
+
+        auto fmin_result_type_pyapi = [&](const py::dtype &dtype1,
+                                          const py::dtype &dtype2) {
+            return py_int::py_binary_ufunc_result_type(
+                dtype1, dtype2, fmin_output_typeid_table);
+        };
+        m.def("_fmin_result_type", fmin_result_type_pyapi);
+    }
+}
+} // namespace dpnp::extensions::ufunc

dpnp/backend/extensions/ufunc/elementwise_functions/fmin.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_fmin(py::module_ m);
+} // namespace dpnp::extensions::ufunc

dpnp/backend/extensions/vm/CMakeLists.txt

@@ -43,6 +43,8 @@ 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}/fmax.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/fmin.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/fmod.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/hypot.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/ln.cpp