From af312f77c796155dd2f9098312669095007bc90e Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Tue, 4 Mar 2025 04:53:32 -0800 Subject: [PATCH 1/3] implement dpnp.hanning --- doc/known_words.txt | 1 + dpnp/backend/extensions/window/hanning.hpp | 66 +++++++++++ dpnp/backend/extensions/window/window_py.cpp | 17 +++ dpnp/dpnp_iface_window.py | 110 +++++++++++++++++- dpnp/tests/test_sycl_queue.py | 1 + dpnp/tests/test_usm_type.py | 1 + dpnp/tests/test_window.py | 4 +- .../cupy/math_tests/test_window.py | 4 +- 8 files changed, 199 insertions(+), 5 deletions(-) create mode 100644 dpnp/backend/extensions/window/hanning.hpp diff --git a/doc/known_words.txt b/doc/known_words.txt index f6f301907530..7de17047c721 100644 --- a/doc/known_words.txt +++ b/doc/known_words.txt @@ -57,6 +57,7 @@ Mises multinomial multivalued namespace +namespaces namedtuple NaN NaT diff --git a/dpnp/backend/extensions/window/hanning.hpp b/dpnp/backend/extensions/window/hanning.hpp new file mode 100644 index 000000000000..9d6b1ebf3e9b --- /dev/null +++ b/dpnp/backend/extensions/window/hanning.hpp @@ -0,0 +1,66 @@ +//***************************************************************************** +// 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. +// +// 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 "common.hpp" +#include + +namespace dpnp::extensions::window::kernels +{ + +template +class HanningFunctor +{ +private: + T *data = nullptr; + const std::size_t N; + +public: + HanningFunctor(T *data, const std::size_t N) : data(data), N(N) {} + + void operator()(sycl::id<1> id) const + { + const auto i = id.get(0); + + data[i] = T(0.5) - T(0.5) * sycl::cospi(T(2) * i / (N - 1)); + } +}; + +template +struct HanningFactory +{ + fnT get() + { + if constexpr (std::is_floating_point_v) { + return window_impl; + } + else { + return nullptr; + } + } +}; + +} // namespace dpnp::extensions::window::kernels diff --git a/dpnp/backend/extensions/window/window_py.cpp b/dpnp/backend/extensions/window/window_py.cpp index c1ed5a3dbc90..04020e920ab8 100644 --- a/dpnp/backend/extensions/window/window_py.cpp +++ b/dpnp/backend/extensions/window/window_py.cpp @@ -32,6 +32,7 @@ #include "common.hpp" #include "hamming.hpp" +#include "hanning.hpp" namespace window_ns = dpnp::extensions::window; namespace py = pybind11; @@ -40,6 +41,7 @@ using window_ns::window_fn_ptr_t; namespace dpctl_td_ns = dpctl::tensor::type_dispatch; static window_fn_ptr_t hamming_dispatch_vector[dpctl_td_ns::num_types]; +static window_fn_ptr_t hanning_dispatch_vector[dpctl_td_ns::num_types]; PYBIND11_MODULE(_window_impl, m) { @@ -60,4 +62,19 @@ PYBIND11_MODULE(_window_impl, m) py::arg("sycl_queue"), py::arg("result"), py::arg("depends") = py::list()); } + + { + window_ns::init_window_dispatch_vectors< + window_ns::kernels::HanningFactory>(hanning_dispatch_vector); + + auto hanning_pyapi = [&](sycl::queue &exec_q, const arrayT &result, + const event_vecT &depends = {}) { + return window_ns::py_window(exec_q, result, depends, + hanning_dispatch_vector); + }; + + m.def("_hanning", hanning_pyapi, "Call hanning kernel", + py::arg("sycl_queue"), py::arg("result"), + py::arg("depends") = py::list()); + } } diff --git a/dpnp/dpnp_iface_window.py b/dpnp/dpnp_iface_window.py index 01a3b1f0c24c..2cbd08fb5691 100644 --- a/dpnp/dpnp_iface_window.py +++ b/dpnp/dpnp_iface_window.py @@ -45,7 +45,7 @@ import dpnp import dpnp.backend.extensions.window._window_impl as wi -__all__ = ["hamming"] +__all__ = ["hamming", "hanning"] def hamming(M, device=None, usm_type=None, sycl_queue=None): @@ -154,3 +154,111 @@ def hamming(M, device=None, usm_type=None, sycl_queue=None): _manager.add_event_pair(ht_ev, win_ev) return result + + +def hanning(M, device=None, usm_type=None, sycl_queue=None): + r""" + Return the Hanning window. + + The Hanning window is a taper formed by using a weighted cosine. + + For full documentation refer to :obj:`numpy.hanning`. + + Parameters + ---------- + M : int + Number of points in the output window. If zero or less, an empty array + is returned. + device : {None, string, SyclDevice, SyclQueue, Device}, optional + An array API concept of device where the output array is created. + `device` can be ``None``, a oneAPI filter selector string, an instance + of :class:`dpctl.SyclDevice` corresponding to a non-partitioned SYCL + device, an instance of :class:`dpctl.SyclQueue`, or a + :class:`dpctl.tensor.Device` object returned by + :attr:`dpnp.ndarray.device`. + + Default: ``None``. + usm_type : {None, "device", "shared", "host"}, optional + The type of SYCL USM allocation for the output array. + + Default: ``None``. + sycl_queue : {None, SyclQueue}, optional + A SYCL queue to use for output array allocation and copying. The + `sycl_queue` can be passed as ``None`` (the default), which means + to get the SYCL queue from `device` keyword if present or to use + a default queue. + + Default: ``None``. + + Returns + ------- + out : dpnp.ndarray of shape (M,) + The window, with the maximum value normalized to one (the value one + appears only if the number of samples is odd). + + See Also + -------- + :obj:`dpnp.bartlett` : Return the Bartlett window. + :obj:`dpnp.blackman` : Return the Blackman window. + :obj:`dpnp.hamming` : Return the Hamming window. + :obj:`dpnp.kaiser` : Return the Kaiser window. + + Notes + ----- + The Hanning window is defined as + + .. math:: w(n) = 0.5 - 0.5\cos\left(\frac{2\pi{n}}{M-1}\right) + \qquad 0 \leq n \leq M-1 + + Examples + -------- + >>> import dpnp as np + >>> np.hanning(12) + array([0. , 0.07937323, 0.29229249, 0.57115742, 0.82743037, + 0.97974649, 0.97974649, 0.82743037, 0.57115742, 0.29229249, + 0.07937323, 0. ]) + + Creating the output array on a different device or with a + specified usm_type: + + >>> x = np.hanning(4) # default case + >>> x, x.device, x.usm_type + (array([0. , 0.75, 0.75, 0. ]), Device(level_zero:gpu:0), 'device') + + >>> y = np.hanning(4, device="cpu") + >>> y, y.device, y.usm_type + (array([0. , 0.75, 0.75, 0. ]), Device(opencl:cpu:0), 'device') + + >>> z = np.hanning(4, usm_type="host") + >>> z, z.device, z.usm_type + (array([0. , 0.75, 0.75, 0. ]), Device(level_zero:gpu:0), 'host') + + """ + + try: + M = int(M) + except Exception as e: + raise TypeError("M must be an integer") from e + + cfd_kwarg = { + "device": device, + "usm_type": usm_type, + "sycl_queue": sycl_queue, + } + + if M < 1: + return dpnp.empty(0, **cfd_kwarg) + if M == 1: + return dpnp.ones(1, **cfd_kwarg) + + result = dpnp.empty(int(M), **cfd_kwarg) + exec_q = result.sycl_queue + _manager = dpu.SequentialOrderManager[exec_q] + + ht_ev, win_ev = wi._hanning( + exec_q, dpnp.get_usm_ndarray(result), depends=_manager.submitted_events + ) + + _manager.add_event_pair(ht_ev, win_ev) + + return result diff --git a/dpnp/tests/test_sycl_queue.py b/dpnp/tests/test_sycl_queue.py index 7b56bb6b03f4..e5c8bbedd58c 100644 --- a/dpnp/tests/test_sycl_queue.py +++ b/dpnp/tests/test_sycl_queue.py @@ -68,6 +68,7 @@ def assert_sycl_queue_equal(result, expected): pytest.param("full", [(2, 2)], {"fill_value": 5}), pytest.param("geomspace", [1, 4, 8], {}), pytest.param("hamming", [10], {}), + pytest.param("hanning", [10], {}), pytest.param("identity", [4], {}), pytest.param("linspace", [0, 4, 8], {}), pytest.param("logspace", [0, 4, 8], {}), diff --git a/dpnp/tests/test_usm_type.py b/dpnp/tests/test_usm_type.py index 0fb563523359..0dbb5cfae1d9 100644 --- a/dpnp/tests/test_usm_type.py +++ b/dpnp/tests/test_usm_type.py @@ -192,6 +192,7 @@ def test_array_creation_from_array(func, args, usm_type_x, usm_type_y): pytest.param("full", [(2, 2)], {"fill_value": 5}), pytest.param("geomspace", [1, 4, 8], {}), pytest.param("hamming", [10], {}), + pytest.param("hanning", [10], {}), pytest.param("identity", [4], {}), pytest.param("linspace", [0, 4, 8], {}), pytest.param("logspace", [0, 4, 8], {}), diff --git a/dpnp/tests/test_window.py b/dpnp/tests/test_window.py index edb1063f6a79..3d68e577cf12 100644 --- a/dpnp/tests/test_window.py +++ b/dpnp/tests/test_window.py @@ -7,7 +7,7 @@ from .helper import assert_dtype_allclose -@pytest.mark.parametrize("func", ["hamming"]) +@pytest.mark.parametrize("func", ["hamming", "hanning"]) @pytest.mark.parametrize( "M", [ @@ -32,7 +32,7 @@ def test_window(func, M): assert_dtype_allclose(result, expected) -@pytest.mark.parametrize("func", ["hamming"]) +@pytest.mark.parametrize("func", ["hamming", "hanning"]) @pytest.mark.parametrize( "M", [ diff --git a/dpnp/tests/third_party/cupy/math_tests/test_window.py b/dpnp/tests/third_party/cupy/math_tests/test_window.py index 2cd05fec2d42..143d91f6207f 100644 --- a/dpnp/tests/third_party/cupy/math_tests/test_window.py +++ b/dpnp/tests/third_party/cupy/math_tests/test_window.py @@ -10,8 +10,8 @@ *testing.product( { "m": [0, 1, -1, 1024], - # TODO: add ["bartlett", "blackman", "hanning"] when supported - "name": ["hamming"], + # TODO: add ["bartlett", "blackman"] when supported + "name": ["hamming", "hanning"], } ) ) From 5c7f7e607e856b3fde2d63300a4f20b4ae503995 Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Tue, 4 Mar 2025 05:31:03 -0800 Subject: [PATCH 2/3] update changelog --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index e22f568bda97..e983e14a11ab 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -9,6 +9,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 ### Added * Added implementation of `dpnp.hamming` [#2341](https://github.com/IntelPython/dpnp/pull/2341), [#2357](https://github.com/IntelPython/dpnp/pull/2357) +* Added implementation of `dpnp.hanning` [#2358](https://github.com/IntelPython/dpnp/pull/2358) ### Changed From b8acf366cbd4faa086573eae1ddb28191f9a42de Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Tue, 4 Mar 2025 11:04:28 -0800 Subject: [PATCH 3/3] address comments --- dpnp/dpnp_iface_window.py | 89 ++++++++++++++++----------------------- 1 file changed, 37 insertions(+), 52 deletions(-) diff --git a/dpnp/dpnp_iface_window.py b/dpnp/dpnp_iface_window.py index 2cbd08fb5691..d3c3084951fb 100644 --- a/dpnp/dpnp_iface_window.py +++ b/dpnp/dpnp_iface_window.py @@ -48,6 +48,39 @@ __all__ = ["hamming", "hanning"] +def _call_window_kernel( + M, _window_kernel, device=None, usm_type=None, sycl_queue=None +): + + try: + M = int(M) + except Exception as e: + raise TypeError("M must be an integer") from e + + cfd_kwarg = { + "device": device, + "usm_type": usm_type, + "sycl_queue": sycl_queue, + } + + if M < 1: + return dpnp.empty(0, **cfd_kwarg) + if M == 1: + return dpnp.ones(1, **cfd_kwarg) + + result = dpnp.empty(M, **cfd_kwarg) + exec_q = result.sycl_queue + _manager = dpu.SequentialOrderManager[exec_q] + + ht_ev, win_ev = _window_kernel( + exec_q, dpnp.get_usm_ndarray(result), depends=_manager.submitted_events + ) + + _manager.add_event_pair(ht_ev, win_ev) + + return result + + def hamming(M, device=None, usm_type=None, sycl_queue=None): r""" Return the Hamming window. @@ -127,34 +160,10 @@ def hamming(M, device=None, usm_type=None, sycl_queue=None): """ - try: - M = int(M) - except Exception as e: - raise TypeError("M must be an integer") from e - - cfd_kwarg = { - "device": device, - "usm_type": usm_type, - "sycl_queue": sycl_queue, - } - - if M < 1: - return dpnp.empty(0, **cfd_kwarg) - if M == 1: - return dpnp.ones(1, **cfd_kwarg) - - result = dpnp.empty(M, **cfd_kwarg) - exec_q = result.sycl_queue - _manager = dpu.SequentialOrderManager[exec_q] - - ht_ev, win_ev = wi._hamming( - exec_q, dpnp.get_usm_ndarray(result), depends=_manager.submitted_events + return _call_window_kernel( + M, wi._hamming, device=device, usm_type=usm_type, sycl_queue=sycl_queue ) - _manager.add_event_pair(ht_ev, win_ev) - - return result - def hanning(M, device=None, usm_type=None, sycl_queue=None): r""" @@ -235,30 +244,6 @@ def hanning(M, device=None, usm_type=None, sycl_queue=None): """ - try: - M = int(M) - except Exception as e: - raise TypeError("M must be an integer") from e - - cfd_kwarg = { - "device": device, - "usm_type": usm_type, - "sycl_queue": sycl_queue, - } - - if M < 1: - return dpnp.empty(0, **cfd_kwarg) - if M == 1: - return dpnp.ones(1, **cfd_kwarg) - - result = dpnp.empty(int(M), **cfd_kwarg) - exec_q = result.sycl_queue - _manager = dpu.SequentialOrderManager[exec_q] - - ht_ev, win_ev = wi._hanning( - exec_q, dpnp.get_usm_ndarray(result), depends=_manager.submitted_events + return _call_window_kernel( + M, wi._hanning, device=device, usm_type=usm_type, sycl_queue=sycl_queue ) - - _manager.add_event_pair(ht_ev, win_ev) - - return result