Add implementation of dpnp.spacing function (#2125)

* Add _spacing to ufunc extension * Add implemntation of dpnp.spacing * Unmute umath tests for dpnp.spacing * Add CFD tests * Add more tests to cover different use cases * Add dedicated sign handling for float16 dtype * Fix typo in the comment * Updated comment in test * Update dpnp/dpnp_iface_mathematical.py Co-authored-by: vtavana <[email protected]> --------- Co-authored-by: vtavana <[email protected]>
IntelPython · Nov 1, 2024 · 16d6ea1 · 16d6ea1
1 parent 2c4f3b5
commit 16d6ea1
Show file tree

Hide file tree

Showing 11 changed files with 375 additions and 4 deletions.
diff --git a/dpnp/backend/extensions/ufunc/CMakeLists.txt b/dpnp/backend/extensions/ufunc/CMakeLists.txt
@@ -38,6 +38,7 @@ set(_elementwise_sources
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/ldexp.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/logaddexp2.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/radians.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/spacing.cpp
 )
 
 set(python_module_name _ufunc_impl)

diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp
@@ -38,6 +38,7 @@
 #include "ldexp.hpp"
 #include "logaddexp2.hpp"
 #include "radians.hpp"
+#include "spacing.hpp"
 
 namespace py = pybind11;
 
@@ -61,5 +62,6 @@ void init_elementwise_functions(py::module_ m)
     init_ldexp(m);
     init_logaddexp2(m);
     init_radians(m);
+    init_spacing(m);
 }
 } // namespace dpnp::extensions::ufunc
diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/spacing.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/spacing.cpp
@@ -0,0 +1,127 @@
+//*****************************************************************************
+// 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 "kernels/elementwise_functions/spacing.hpp"
+#include "populate.hpp"
+#include "spacing.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 dpnp::extensions::ufunc
+{
+namespace py = pybind11;
+namespace py_int = dpnp::extensions::py_internal;
+
+namespace impl
+{
+namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
+namespace td_ns = dpctl::tensor::type_dispatch;
+
+/**
+ * @brief A factory to define pairs of supported types for which
+ * sycl::spacing<T> function is available.
+ *
+ * @tparam T Type of input vector `a` and of result vector `y`.
+ */
+template <typename T>
+struct OutputType
+{
+    using value_type =
+        typename std::disjunction<td_ns::TypeMapResultEntry<T, sycl::half>,
+                                  td_ns::TypeMapResultEntry<T, float>,
+                                  td_ns::TypeMapResultEntry<T, double>,
+                                  td_ns::DefaultResultEntry<void>>::result_type;
+};
+
+using dpnp::kernels::spacing::SpacingFunctor;
+
+template <typename argT,
+          typename resT = argT,
+          unsigned int vec_sz = 4,
+          unsigned int n_vecs = 2,
+          bool enable_sg_loadstore = true>
+using ContigFunctor = ew_cmn_ns::UnaryContigFunctor<argT,
+                                                    resT,
+                                                    SpacingFunctor<argT, resT>,
+                                                    vec_sz,
+                                                    n_vecs,
+                                                    enable_sg_loadstore>;
+
+template <typename argTy, typename resTy, typename IndexerT>
+using StridedFunctor = ew_cmn_ns::
+    UnaryStridedFunctor<argTy, resTy, IndexerT, SpacingFunctor<argTy, resTy>>;
+
+using ew_cmn_ns::unary_contig_impl_fn_ptr_t;
+using ew_cmn_ns::unary_strided_impl_fn_ptr_t;
+
+static unary_contig_impl_fn_ptr_t
+    spacing_contig_dispatch_vector[td_ns::num_types];
+static int spacing_output_typeid_vector[td_ns::num_types];
+static unary_strided_impl_fn_ptr_t
+    spacing_strided_dispatch_vector[td_ns::num_types];
+
+MACRO_POPULATE_DISPATCH_VECTORS(spacing);
+} // namespace impl
+
+void init_spacing(py::module_ m)
+{
+    using arrayT = dpctl::tensor::usm_ndarray;
+    using event_vecT = std::vector<sycl::event>;
+    {
+        impl::populate_spacing_dispatch_vectors();
+        using impl::spacing_contig_dispatch_vector;
+        using impl::spacing_output_typeid_vector;
+        using impl::spacing_strided_dispatch_vector;
+
+        auto spacing_pyapi = [&](const arrayT &src, const arrayT &dst,
+                                 sycl::queue &exec_q,
+                                 const event_vecT &depends = {}) {
+            return py_int::py_unary_ufunc(src, dst, exec_q, depends,
+                                          spacing_output_typeid_vector,
+                                          spacing_contig_dispatch_vector,
+                                          spacing_strided_dispatch_vector);
+        };
+        m.def("_spacing", spacing_pyapi, "", py::arg("src"), py::arg("dst"),
+              py::arg("sycl_queue"), py::arg("depends") = py::list());
+
+        auto spacing_result_type_pyapi = [&](const py::dtype &dtype) {
+            return py_int::py_unary_ufunc_result_type(
+                dtype, spacing_output_typeid_vector);
+        };
+        m.def("_spacing_result_type", spacing_result_type_pyapi);
+    }
+}
+} // namespace dpnp::extensions::ufunc
diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/spacing.hpp b/dpnp/backend/extensions/ufunc/elementwise_functions/spacing.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_spacing(py::module_ m);
+} // namespace dpnp::extensions::ufunc
diff --git a/dpnp/backend/kernels/elementwise_functions/spacing.hpp b/dpnp/backend/kernels/elementwise_functions/spacing.hpp
@@ -0,0 +1,64 @@
+//*****************************************************************************
+// 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 <sycl/sycl.hpp>
+
+namespace dpnp::kernels::spacing
+{
+template <typename argT, typename resT>
+struct SpacingFunctor
+{
+    // is function constant for given argT
+    using is_constant = typename std::false_type;
+    // constant value, if constant
+    // constexpr resT constant_value = resT{};
+    // is function defined for sycl::vec
+    using supports_vec = typename std::false_type;
+    // do both argT and resT support subgroup store/load operation
+    using supports_sg_loadstore = typename std::true_type;
+
+    resT operator()(const argT &x) const
+    {
+        if (sycl::isnan(x)) {
+            return x;
+        }
+
+        if (sycl::isinf(x)) {
+            return std::numeric_limits<resT>::quiet_NaN();
+        }
+
+        constexpr argT inf = std::numeric_limits<argT>::infinity();
+        if constexpr (std::is_same_v<argT, sycl::half>) {
+            // numpy always computes spacing towards +inf for float16 dtype
+            return sycl::nextafter(x, inf) - x;
+        }
+        else {
+            return sycl::nextafter(x, sycl::copysign(inf, x)) - x;
+        }
+    }
+};
+} // namespace dpnp::kernels::spacing
diff --git a/dpnp/dpnp_iface_mathematical.py b/dpnp/dpnp_iface_mathematical.py
@@ -130,6 +130,7 @@
     "round",
     "sign",
     "signbit",
+    "spacing",
     "subtract",
     "sum",
     "trapezoid",
@@ -3749,6 +3750,62 @@ def real_if_close(a, tol=100):
 )
 
 
+_SPACING_DOCSTRING = """
+Return the distance between `x` and the nearest adjacent number.
+
+For full documentation refer to :obj:`numpy.spacing`.
+
+Parameters
+----------
+x : {dpnp.ndarray, usm_ndarray}
+    The array of values to find the spacing of, expected to have a real-valued
+    data type.
+out : {None, dpnp.ndarray, usm_ndarray}, optional
+    Output array to populate.
+    Array must have the correct shape and the expected data type.
+    Default: ``None``.
+order : {"C", "F", "A", "K"}, optional
+    Memory layout of the newly output array, if parameter `out` is ``None``.
+    Default: ``"K"``.
+
+Returns
+-------
+out : dpnp.ndarray
+    The spacing of values of `x`. The data type of the returned array is
+    determined by the Type Promotion Rules.
+
+Limitations
+-----------
+Parameters `where` and `subok` are supported with their default values.
+Keyword argument `kwargs` is currently unsupported.
+Otherwise ``NotImplementedError`` exception will be raised.
+
+Notes
+-----
+It can be considered as a generalization of EPS:
+``dpnp.spacing(dpnp.float64(1)) == dpnp.finfo(dpnp.float64).eps``, and there
+should not be any representable number between ``x + spacing(x)`` and `x` for
+any finite `x`.
+
+Spacing of +- inf and NaN is ``NaN``.
+
+Examples
+--------
+>>> import dpnp as np
+>>> a = np.array(1)
+>>> b = np.spacing(a)
+>>> b == np.finfo(b.dtype).eps
+array(True)
+"""
+
+spacing = DPNPUnaryFunc(
+    "spacing",
+    ufi._spacing_result_type,
+    ufi._spacing,
+    _SPACING_DOCSTRING,
+)
+
+
 _SUBTRACT_DOCSTRING = """
 Calculates the difference between each element `x1_i` of the input
 array `x1` and the respective element `x2_i` of the input array `x2`.

diff --git a/tests/skipped_tests.tbl b/tests/skipped_tests.tbl
@@ -17,8 +17,6 @@ tests/test_umath.py::test_umaths[('divmod', 'ff')]
 tests/test_umath.py::test_umaths[('divmod', 'dd')]
 tests/test_umath.py::test_umaths[('frexp', 'f')]
 tests/test_umath.py::test_umaths[('frexp', 'd')]
-tests/test_umath.py::test_umaths[('spacing', 'f')]
-tests/test_umath.py::test_umaths[('spacing', 'd')]
 
 tests/third_party/cupy/core_tests/test_ndarray_conversion.py::TestNdarrayToBytes_param_0_{shape=()}::test_item
 tests/third_party/cupy/core_tests/test_ndarray_conversion.py::TestNdarrayToBytes_param_1_{shape=(1,)}::test_item

diff --git a/tests/skipped_tests_gpu.tbl b/tests/skipped_tests_gpu.tbl
@@ -24,8 +24,6 @@ tests/test_umath.py::test_umaths[('divmod', 'dd')]
 tests/test_umath.py::test_umaths[('floor_divide', 'ff')]
 tests/test_umath.py::test_umaths[('frexp', 'f')]
 tests/test_umath.py::test_umaths[('frexp', 'd')]
-tests/test_umath.py::test_umaths[('spacing', 'f')]
-tests/test_umath.py::test_umaths[('spacing', 'd')]
 
 tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsGeometric_param_2_{p_shape=(3, 2), shape=(4, 3, 2)}::test_geometric
 tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsGeometric_param_3_{p_shape=(3, 2), shape=(3, 2)}::test_geometric