Merge branch 'master' into fix-all-int-windows

Author: vtavana

CHANGELOG.md

@@ -10,6 +10,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0
 
 * 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)
+* Added implementation of `dpnp.blackman` [#2363](https://github.com/IntelPython/dpnp/pull/2363)
 
 ### Changed
 

dpnp/backend/extensions/window/blackman.hpp

@@ -0,0 +1,67 @@
+//*****************************************************************************
+// 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 BlackmanFunctor
+{
+private:
+    T *data = nullptr;
+    const std::size_t N;
+
+public:
+    BlackmanFunctor(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.42) - T(0.5) * sycl::cospi(T(2) * i / (N - 1)) +
+                  T(0.08) * sycl::cospi(T(4) * i / (N - 1));
+    }
+};
+
+template <typename fnT, typename T>
+struct BlackmanFactory
+{
+    fnT get()
+    {
+        if constexpr (std::is_floating_point_v<T>) {
+            return window_impl<T, BlackmanFunctor>;
+        }
+        else {
+            return nullptr;
+        }
+    }
+};
+
+} // namespace dpnp::extensions::window::kernels

dpnp/backend/extensions/window/window_py.cpp

@@ -30,6 +30,7 @@
 #include <pybind11/pybind11.h>
 #include <pybind11/stl.h>
 
+#include "blackman.hpp"
 #include "common.hpp"
 #include "hamming.hpp"
 #include "hanning.hpp"
@@ -40,6 +41,7 @@ using window_ns::window_fn_ptr_t;
 
 namespace dpctl_td_ns = dpctl::tensor::type_dispatch;
 
+static window_fn_ptr_t blackman_dispatch_vector[dpctl_td_ns::num_types];
 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];
 
@@ -48,6 +50,21 @@ PYBIND11_MODULE(_window_impl, m)
     using arrayT = dpctl::tensor::usm_ndarray;
     using event_vecT = std::vector<sycl::event>;
 
+    {
+        window_ns::init_window_dispatch_vectors<
+            window_ns::kernels::BlackmanFactory>(blackman_dispatch_vector);
+
+        auto blackman_pyapi = [&](sycl::queue &exec_q, const arrayT &result,
+                                  const event_vecT &depends = {}) {
+            return window_ns::py_window(exec_q, result, depends,
+                                        blackman_dispatch_vector);
+        };
+
+        m.def("_blackman", blackman_pyapi, "Call Blackman kernel",
+              py::arg("sycl_queue"), py::arg("result"),
+              py::arg("depends") = py::list());
+    }
+
     {
         window_ns::init_window_dispatch_vectors<
             window_ns::kernels::HammingFactory>(hamming_dispatch_vector);
@@ -58,7 +75,7 @@ PYBIND11_MODULE(_window_impl, m)
                                         hamming_dispatch_vector);
         };
 
-        m.def("_hamming", hamming_pyapi, "Call hamming kernel",
+        m.def("_hamming", hamming_pyapi, "Call Hamming kernel",
               py::arg("sycl_queue"), py::arg("result"),
               py::arg("depends") = py::list());
     }
@@ -73,7 +90,7 @@ PYBIND11_MODULE(_window_impl, m)
                                         hanning_dispatch_vector);
         };
 
-        m.def("_hanning", hanning_pyapi, "Call hanning kernel",
+        m.def("_hanning", hanning_pyapi, "Call Hanning kernel",
               py::arg("sycl_queue"), py::arg("result"),
               py::arg("depends") = py::list());
     }

dpnp/dpnp_iface_counting.py

@@ -58,16 +58,19 @@ def count_nonzero(a, axis=None, *, keepdims=False, out=None):
         Axis or tuple of axes along which to count non-zeros.
         Default value means that non-zeros will be counted along a flattened
         version of `a`.
+
         Default: ``None``.
-    keepdims : bool, optional
+    keepdims : {None, bool}, optional
         If this is set to ``True``, the axes that are counted are left in the
         result as dimensions with size one. With this option, the result will
         broadcast correctly against the input array.
+
         Default: ``False``.
     out : {None, dpnp.ndarray, usm_ndarray}, optional
         The array into which the result is written. The data type of `out` must
         match the expected shape and the expected data type of the result.
         If ``None`` then a new array is returned.
+
         Default: ``None``.
 
     Returns

dpnp/dpnp_iface_window.py

@@ -45,7 +45,7 @@
 import dpnp
 import dpnp.backend.extensions.window._window_impl as wi
 
-__all__ = ["hamming", "hanning"]
+__all__ = ["blackman", "hamming", "hanning"]
 
 
 def _call_window_kernel(
@@ -81,6 +81,99 @@ def _call_window_kernel(
     return result
 
 
+def blackman(M, device=None, usm_type=None, sycl_queue=None):
+    r"""
+    Return the Blackman window.
+
+    The Blackman window is a taper formed by using the first three terms of a
+    summation of cosines. It was designed to have close to the minimal leakage
+    possible. It is close to optimal, only slightly worse than a Kaiser window.
+
+    For full documentation refer to :obj:`numpy.blackman`.
+
+    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.hamming` : Return the Hamming window.
+    :obj:`dpnp.hanning` : Return the Hanning window.
+    :obj:`dpnp.kaiser` : Return the Kaiser window.
+
+    Notes
+    -----
+    The Blackman window is defined as
+
+    .. math::  w(n) = 0.42 - 0.5\cos\left(\frac{2\pi{n}}{M-1}\right)
+               + 0.08\cos\left(\frac{4\pi{n}}{M-1}\right)
+               \qquad 0 \leq n \leq M-1
+
+    Examples
+    --------
+    >>> import dpnp as np
+    >>> np.blackman(12)
+    array([-1.38777878e-17,  3.26064346e-02,  1.59903635e-01,  4.14397981e-01,
+            7.36045180e-01,  9.67046769e-01,  9.67046769e-01,  7.36045180e-01,
+            4.14397981e-01,  1.59903635e-01,  3.26064346e-02, -1.38777878e-17])
+
+    Creating the output array on a different device or with a
+    specified usm_type:
+
+    >>> x = np.blackman(3) # default case
+    >>> x, x.device, x.usm_type
+    (array([-1.38777878e-17,  1.00000000e+00, -1.38777878e-17]),
+     Device(level_zero:gpu:0),
+     'device')
+
+    >>> y = np.blackman(3, device="cpu")
+    >>> y, y.device, y.usm_type
+    (array([-1.38777878e-17,  1.00000000e+00, -1.38777878e-17]),
+     Device(opencl:cpu:0),
+     'device')
+
+    >>> z = np.blackman(3, usm_type="host")
+    >>> z, z.device, z.usm_type
+    (array([-1.38777878e-17,  1.00000000e+00, -1.38777878e-17]),
+     Device(level_zero:gpu:0),
+     'host')
+
+    """
+
+    return _call_window_kernel(
+        M, wi._blackman, device=device, usm_type=usm_type, sycl_queue=sycl_queue
+    )
+
+
 def hamming(M, device=None, usm_type=None, sycl_queue=None):
     r"""
     Return the Hamming window.

dpnp/tests/test_indexing.py

@@ -1200,15 +1200,17 @@ def test_empty_indices_error(self):
         )
 
     def test_empty_indices(self):
-        assert_equal(
-            dpnp.ravel_multi_index(
-                (dpnp.array([], dtype=int), dpnp.array([], dtype=int)), (5, 3)
-            ),
-            [],
-        )
-        assert_equal(
-            dpnp.ravel_multi_index(dpnp.array([[], []], dtype=int), (5, 3)), []
-        )
+        a = numpy.array([], dtype=int)
+        ia = dpnp.array(a)
+        result = dpnp.ravel_multi_index((ia, ia), (5, 3))
+        expected = numpy.ravel_multi_index((a, a), (5, 3))
+        assert_equal(result, expected)
+
+        a = numpy.array([[], []], dtype=int)
+        ia = dpnp.array(a)
+        result = dpnp.ravel_multi_index(ia, (5, 3))
+        expected = numpy.ravel_multi_index(a, (5, 3))
+        assert_equal(result, expected)
 
 
 class TestUnravelIndex:

dpnp/tests/test_linalg.py

@@ -293,7 +293,9 @@ def test_empty(self, shape, p):
         expected = numpy.linalg.cond(a, p=p)
         assert_dtype_allclose(result, expected)
 
-    @pytest.mark.parametrize("dtype", get_all_dtypes(no_bool=True))
+    @pytest.mark.parametrize(
+        "dtype", get_all_dtypes(no_none=True, no_bool=True)
+    )
     @pytest.mark.parametrize(
         "shape", [(4, 4), (2, 4, 3, 3)], ids=["(4, 4)", "(2, 4, 3, 3)"]
     )
@@ -2408,10 +2410,8 @@ def test_qr(self, dtype, shape, mode):
             if mode in ("complete", "reduced"):
                 result = dpnp.linalg.qr(ia, mode)
                 dpnp_q, dpnp_r = result.Q, result.R
-                assert_almost_equal(
-                    dpnp.matmul(dpnp_q, dpnp_r),
-                    a,
-                    decimal=5,
+                assert dpnp.allclose(
+                    dpnp.matmul(dpnp_q, dpnp_r), ia, atol=1e-05
                 )
             else:  # mode=="raw"
                 dpnp_q, dpnp_r = dpnp.linalg.qr(ia, mode)
@@ -2424,15 +2424,13 @@ def test_qr(self, dtype, shape, mode):
     @pytest.mark.parametrize(
         "shape",
         [(32, 32), (8, 16, 16)],
-        ids=[
-            "(32, 32)",
-            "(8, 16, 16)",
-        ],
+        ids=["(32, 32)", "(8, 16, 16)"],
     )
     @pytest.mark.parametrize("mode", ["r", "raw", "complete", "reduced"])
     def test_qr_large(self, dtype, shape, mode):
         a = generate_random_numpy_array(shape, dtype, seed_value=81)
         ia = dpnp.array(a)
+
         if mode == "r":
             np_r = numpy.linalg.qr(a, mode)
             dpnp_r = dpnp.linalg.qr(ia, mode)
@@ -2443,11 +2441,7 @@ def test_qr_large(self, dtype, shape, mode):
             if mode in ("complete", "reduced"):
                 result = dpnp.linalg.qr(ia, mode)
                 dpnp_q, dpnp_r = result.Q, result.R
-                assert_almost_equal(
-                    dpnp.matmul(dpnp_q, dpnp_r),
-                    a,
-                    decimal=5,
-                )
+                assert dpnp.allclose(dpnp.matmul(dpnp_q, dpnp_r), ia, atol=1e-5)
             else:  # mode=="raw"
                 dpnp_q, dpnp_r = dpnp.linalg.qr(ia, mode)
                 assert_allclose(dpnp_q, np_q, atol=1e-4)
@@ -3169,9 +3163,7 @@ def test_test_tensorinv_errors(self):
 class TestTensorsolve:
     @pytest.mark.parametrize("dtype", get_all_dtypes())
     @pytest.mark.parametrize(
-        "axes",
-        [None, (1,), (2,)],
-        ids=["None", "(1,)", "(2,)"],
+        "axes", [None, (1,), (2,)], ids=["None", "(1,)", "(2,)"]
     )
     def test_tensorsolve_axes(self, dtype, axes):
         a = numpy.eye(12).reshape(12, 3, 4).astype(dtype)