implement dpnp.hanning

Author: Vahid Tavanashad

doc/known_words.txt

@@ -57,6 +57,7 @@ Mises
 multinomial
 multivalued
 namespace
+namespaces
 namedtuple
 NaN
 NaT

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 <sycl/sycl.hpp>
+
+namespace dpnp::extensions::window::kernels
+{
+
+template <typename T>
+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 <typename fnT, typename T>
+struct HanningFactory
+{
+    fnT get()
+    {
+        if constexpr (std::is_floating_point_v<T>) {
+            return window_impl<T, HanningFunctor>;
+        }
+        else {
+            return nullptr;
+        }
+    }
+};
+
+} // namespace dpnp::extensions::window::kernels

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());
+    }
 }

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

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], {}),

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], {}),

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",
     [

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"],
         }
     )
 )