Merge branch 'master' into add-uint64-to-bincount

Author: antonwolfy

.github/workflows/conda-package.yml

@@ -218,7 +218,7 @@ jobs:
       - name: Run tests
         if: env.rerun-tests-on-failure != 'true'
         run: |
-          if [[ ${{ matrix.python }} == ${{ env.python-ver-test-all-dtypes }} ]]; then
+          if [[ "${{ matrix.python }}" == "${{ env.python-ver-test-all-dtypes }}" ]]; then
             export DPNP_TEST_ALL_INT_TYPES=1
             python -m pytest -ra --pyargs ${{ env.package-name }}.tests
           else
@@ -238,7 +238,7 @@ jobs:
             . $CONDA/etc/profile.d/mamba.sh
             mamba activate ${{ env.test-env-name }}
 
-            if [[ ${{ matrix.python }} == ${{ env.python-ver-test-all-dtypes }} ]]; then
+            if [[ "${{ matrix.python }}" == "${{ env.python-ver-test-all-dtypes }}" ]]; then
               export DPNP_TEST_ALL_INT_TYPES=1
               python -m pytest -ra --pyargs ${{ env.package-name }}.tests
             else
@@ -383,9 +383,10 @@ jobs:
 
       - name: Run tests
         if: env.rerun-tests-on-failure != 'true'
+        shell: pwsh
         run: |
           if (${{ matrix.python }} -eq ${{ env.python-ver-test-all-dtypes }}) {
-            set DPNP_TEST_ALL_INT_TYPES=1
+            $env:DPNP_TEST_ALL_INT_TYPES=1
             python -m pytest -ra --pyargs ${{ env.package-name }}.tests
           } else {
             python -m pytest -n auto -ra --pyargs ${{ env.package-name }}.tests
@@ -399,9 +400,10 @@ jobs:
           timeout_minutes: ${{ env.rerun-tests-timeout }}
           max_attempts: ${{ env.rerun-tests-max-attempts }}
           retry_on: any
+          shell: pwsh
           command: |
             if ( ${{ matrix.python }} -eq ${{ env.python-ver-test-all-dtypes }} ) {
-              set DPNP_TEST_ALL_INT_TYPES=1
+              $env:DPNP_TEST_ALL_INT_TYPES=1
               python -m pytest -ra --pyargs ${{ env.package-name }}.tests
             } else {
               python -m pytest -n auto -ra --pyargs ${{ env.package-name }}.tests

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/conftest.py

@@ -32,6 +32,8 @@
 import numpy
 import pytest
 
+from . import config as dtype_config
+
 if numpy.lib.NumpyVersion(numpy.__version__) >= "2.0.0b1":
     from numpy.exceptions import ComplexWarning
 else:
@@ -128,20 +130,25 @@ def pytest_collection_modifyitems(config, items):
 
     dev = dpctl.select_default_device()
     is_cpu = dev.is_cpu
-    is_gpu_no_fp64 = not dev.has_aspect_fp64
+    is_gpu = dev.is_gpu
+    support_fp64 = dev.has_aspect_fp64
     is_cuda = dpnp.is_cuda_backend(dev)
 
     print("")
+    print(
+        f"DPNP Test scope includes all integer dtypes: {bool(dtype_config.all_int_types)}"
+    )
     print(f"DPNP current device is CPU: {is_cpu}")
-    print(f"DPNP current device is GPU without fp64 support: {is_gpu_no_fp64}")
+    print(f"DPNP current device is GPU: {is_gpu}")
+    print(f"DPNP current device supports fp64: {support_fp64}")
     print(f"DPNP current device is GPU with cuda backend: {is_cuda}")
     print(f"DPNP version: {dpnp.__version__}, location: {dpnp}")
     print(f"NumPy version: {numpy.__version__}, location: {numpy}")
     print(f"Python version: {sys.version}")
     print("")
-    if not is_cpu or os.getenv("DPNP_QUEUE_GPU") == "1":
+    if is_gpu or os.getenv("DPNP_QUEUE_GPU") == "1":
         excluded_tests.extend(get_excluded_tests(test_exclude_file_gpu))
-        if is_gpu_no_fp64:
+        if not support_fp64:
             excluded_tests.extend(
                 get_excluded_tests(test_exclude_file_gpu_no_fp64)
             )

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: