diff --git a/dpnp/backend/extensions/ufunc/CMakeLists.txt b/dpnp/backend/extensions/ufunc/CMakeLists.txt index 4fe714dbeb5b..1d7bc2f4ac1b 100644 --- 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 index c5b931d3d21d..e29def5df4ac 100644 --- 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 new file mode 100644 index 000000000000..1320e8bb6374 --- /dev/null +++ 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 + +#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 function is available. + * + * @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; +}; + +using dpnp::kernels::spacing::SpacingFunctor; + +template +using ContigFunctor = ew_cmn_ns::UnaryContigFunctor, + vec_sz, + n_vecs, + enable_sg_loadstore>; + +template +using StridedFunctor = ew_cmn_ns:: + UnaryStridedFunctor>; + +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; + { + 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 new file mode 100644 index 000000000000..db44697d7b55 --- /dev/null +++ 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 + +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 new file mode 100644 index 000000000000..d7b20b629dee --- /dev/null +++ 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 + +namespace dpnp::kernels::spacing +{ +template +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::quiet_NaN(); + } + + constexpr argT inf = std::numeric_limits::infinity(); + if constexpr (std::is_same_v) { + // 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 index 4a6bd9edc8ef..ecbe35c741cf 100644 --- 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 index 1248c32c9acf..d0ee8bb0b617 100644 --- 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 index 2ab14190fa78..82e7c2abee13 100644 --- 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 diff --git a/tests/test_mathematical.py b/tests/test_mathematical.py index 3d85022459ba..01555df2718e 100644 --- a/tests/test_mathematical.py +++ b/tests/test_mathematical.py @@ -1897,6 +1897,93 @@ def test_wrong_tol_type(self, xp, tol_val): assert_raises(TypeError, xp.real_if_close, a, tol=tol_val) +class TestSpacing: + @pytest.mark.parametrize("sign", [1, -1]) + @pytest.mark.parametrize("dt", get_float_dtypes()) + def test_basic(self, sign, dt): + a = numpy.array( + [1, numpy.nan, numpy.inf, 1e10, 1e-5, 1000, 10500], dtype=dt + ) + a *= sign + ia = dpnp.array(a) + + result = dpnp.spacing(ia) + expected = numpy.spacing(a) + assert_equal(result, expected) + + # switch to negatives + result = dpnp.spacing(-ia) + expected = numpy.spacing(-a) + assert_equal(result, expected) + + @pytest.mark.parametrize("dt", get_float_dtypes()) + def test_zeros(self, dt): + a = numpy.array([0.0, -0.0], dtype=dt) + ia = dpnp.array(a) + + result = dpnp.spacing(ia) + expected = numpy.spacing(a) + if numpy.lib.NumpyVersion(numpy.__version__) < "2.0.0": + assert_equal(result, expected) + else: + # numpy.spacing(-0.0) == numpy.spacing(0.0), i.e. NumPy returns + # positive value (looks as a bug in NumPy), because for any other + # negative input the NumPy result will be also a negative value. + expected[1] *= -1 + assert_equal(result, expected) + + @pytest.mark.parametrize("dt", get_float_dtypes(no_float16=False)) + @pytest.mark.parametrize("val", [1, 1e-5, 1000]) + @pytest.mark.parametrize("xp", [numpy, dpnp]) + def test_vs_nextafter(self, val, dt, xp): + a = xp.array(val, dtype=dt) + a1 = xp.array(val + 1, dtype=dt) + assert (xp.nextafter(a, a1) - a) == xp.spacing(a) + + @pytest.mark.filterwarnings("ignore::RuntimeWarning") + @pytest.mark.skipif(not has_support_aspect16(), reason="no fp16 support") + def test_fp16(self): + a = numpy.arange(0x7C00, dtype=numpy.uint16) + + # all values are non-negative finites + b = a.view(dtype=numpy.float16) + ib = dpnp.array(b) + + result = dpnp.spacing(ib) + expected = numpy.spacing(b) + assert_equal(result, expected) + + # switch to negatives + a |= 0x8000 + ib = dpnp.array(b) + + result = dpnp.spacing(ib) + expected = numpy.spacing(b) + assert_equal(result, expected) + + @pytest.mark.parametrize("dt", get_integer_dtypes()) + def test_integer(self, dt): + a = numpy.array([1, 0, -3], dtype=dt) + ia = dpnp.array(a) + + result = dpnp.spacing(ia) + expected = numpy.spacing(a) + assert_dtype_allclose(result, expected) + + def test_bool(self): + a = numpy.array([True, False]) + ia = dpnp.array(a) + + result = dpnp.spacing(ia) + expected = numpy.spacing(a) + assert_dtype_allclose(result, expected) + + @pytest.mark.parametrize("xp", [numpy, dpnp]) + def test_complex(self, xp): + a = xp.array([2.1 + 4e-14j, 5.2 + 3e-15j]) + assert_raises((TypeError, ValueError), xp.spacing, a) + + class TestTrapezoid: def get_numpy_func(self): if numpy.lib.NumpyVersion(numpy.__version__) < "2.0.0": diff --git a/tests/test_sycl_queue.py b/tests/test_sycl_queue.py index e7bf0403ea55..e8c0969f8d0c 100644 --- a/tests/test_sycl_queue.py +++ b/tests/test_sycl_queue.py @@ -505,6 +505,7 @@ def test_meshgrid(device): pytest.param("sinh", [-5.0, -3.5, 0.0, 3.5, 5.0]), pytest.param("sort", [2.0, 1.0, 7.0, 4.0]), pytest.param("sort_complex", [1 + 2j, 2 - 1j, 3 - 2j, 3 - 3j, 3 + 5j]), + pytest.param("spacing", [1, 2, -3, 0]), pytest.param("sqrt", [1.0, 3.0, 9.0]), pytest.param("square", [1.0, 3.0, 9.0]), pytest.param("std", [1.0, 2.0, 4.0, 7.0]), diff --git a/tests/test_usm_type.py b/tests/test_usm_type.py index cebb530050b7..13752a377dbe 100644 --- a/tests/test_usm_type.py +++ b/tests/test_usm_type.py @@ -644,6 +644,7 @@ def test_norm(usm_type, ord, axis): pytest.param("sinh", [-5.0, -3.5, 0.0, 3.5, 5.0]), pytest.param("sort", [2.0, 1.0, 7.0, 4.0]), pytest.param("sort_complex", [1 + 2j, 2 - 1j, 3 - 2j, 3 - 3j, 3 + 5j]), + pytest.param("spacing", [1, 2, -3, 0]), pytest.param("sqrt", [1.0, 3.0, 9.0]), pytest.param("square", [1.0, 3.0, 9.0]), pytest.param("std", [1.0, 2.0, 4.0, 7.0]),