diff --git a/CHANGELOG.md b/CHANGELOG.md index e983e14a11ab..5972b12767fa 100644 --- a/CHANGELOG.md +++ b/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 diff --git a/dpnp/backend/extensions/window/blackman.hpp b/dpnp/backend/extensions/window/blackman.hpp new file mode 100644 index 000000000000..4185decb03b9 --- /dev/null +++ b/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 + +namespace dpnp::extensions::window::kernels +{ + +template +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 +struct BlackmanFactory +{ + 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 04020e920ab8..5f5a1a50db78 100644 --- a/dpnp/backend/extensions/window/window_py.cpp +++ b/dpnp/backend/extensions/window/window_py.cpp @@ -30,6 +30,7 @@ #include #include +#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; + { + 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()); } diff --git a/dpnp/dpnp_iface_window.py b/dpnp/dpnp_iface_window.py index d3c3084951fb..fd754a7b4946 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", "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. diff --git a/dpnp/tests/test_sycl_queue.py b/dpnp/tests/test_sycl_queue.py index 21ff91950e1b..4accd1123083 100644 --- a/dpnp/tests/test_sycl_queue.py +++ b/dpnp/tests/test_sycl_queue.py @@ -53,6 +53,7 @@ def assert_sycl_queue_equal(result, expected): "func, arg, kwargs", [ pytest.param("arange", [-25.7], {"stop": 10**8, "step": 15}), + pytest.param("blackman", [10], {}), pytest.param("eye", [4, 2], {}), pytest.param("empty", [(2, 2)], {}), pytest.param( diff --git a/dpnp/tests/test_usm_type.py b/dpnp/tests/test_usm_type.py index 5e29b843ac4e..ff5e0036ef81 100644 --- a/dpnp/tests/test_usm_type.py +++ b/dpnp/tests/test_usm_type.py @@ -177,6 +177,7 @@ def test_array_creation_from_array(func, args, usm_type_x, usm_type_y): "func, arg, kwargs", [ pytest.param("arange", [-25.7], {"stop": 10**8, "step": 15}), + pytest.param("blackman", [10], {}), pytest.param("eye", [4, 2], {}), pytest.param("empty", [(3, 4)], {}), pytest.param( diff --git a/dpnp/tests/test_window.py b/dpnp/tests/test_window.py index 3d68e577cf12..9529a882beb6 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", "hanning"]) +@pytest.mark.parametrize("func", ["blackman", "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", "hanning"]) +@pytest.mark.parametrize("func", ["blackman", "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 143d91f6207f..27ce810bdb4e 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"] when supported - "name": ["hamming", "hanning"], + # TODO: add ["bartlett"] when supported + "name": ["blackman", "hamming", "hanning"], } ) )