From e22975bec905511aecc9aae57e3457f2e3b2adbf Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Wed, 5 Mar 2025 09:51:06 -0800 Subject: [PATCH 1/4] implement dpnp.bartlett --- dpnp/backend/extensions/window/bartlett.hpp | 67 +++++++++++++ dpnp/backend/extensions/window/window_py.cpp | 17 ++++ dpnp/dpnp_iface_window.py | 95 ++++++++++++++++++- 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 | 3 +- 7 files changed, 183 insertions(+), 5 deletions(-) create mode 100644 dpnp/backend/extensions/window/bartlett.hpp diff --git a/dpnp/backend/extensions/window/bartlett.hpp b/dpnp/backend/extensions/window/bartlett.hpp new file mode 100644 index 000000000000..0ade4722b8d9 --- /dev/null +++ b/dpnp/backend/extensions/window/bartlett.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 BartlettFunctor +{ +private: + T *data = nullptr; + const std::size_t N; + +public: + BartlettFunctor(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(2) / (N - 1) * ((N - 1) / T(2) - sycl::fabs(i - (N - 1) / T(2))); + } +}; + +template +struct BartlettFactory +{ + 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 5f5a1a50db78..3e4ff8849070 100644 --- a/dpnp/backend/extensions/window/window_py.cpp +++ b/dpnp/backend/extensions/window/window_py.cpp @@ -30,6 +30,7 @@ #include #include +#include "bartlett.hpp" #include "blackman.hpp" #include "common.hpp" #include "hamming.hpp" @@ -41,6 +42,7 @@ using window_ns::window_fn_ptr_t; namespace dpctl_td_ns = dpctl::tensor::type_dispatch; +static window_fn_ptr_t bartlett_dispatch_vector[dpctl_td_ns::num_types]; 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]; @@ -50,6 +52,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::BartlettFactory>(bartlett_dispatch_vector); + + auto bartlett_pyapi = [&](sycl::queue &exec_q, const arrayT &result, + const event_vecT &depends = {}) { + return window_ns::py_window(exec_q, result, depends, + bartlett_dispatch_vector); + }; + + m.def("_bartlett", bartlett_pyapi, "Call Bartlett kernel", + py::arg("sycl_queue"), py::arg("result"), + py::arg("depends") = py::list()); + } + { window_ns::init_window_dispatch_vectors< window_ns::kernels::BlackmanFactory>(blackman_dispatch_vector); diff --git a/dpnp/dpnp_iface_window.py b/dpnp/dpnp_iface_window.py index fd754a7b4946..16349c76a7cc 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__ = ["blackman", "hamming", "hanning"] +__all__ = ["bartlett", "blackman", "hamming", "hanning"] def _call_window_kernel( @@ -174,6 +174,99 @@ def blackman(M, device=None, usm_type=None, sycl_queue=None): ) +def bartlett(M, device=None, usm_type=None, sycl_queue=None): + r""" + Return the Bartlett window. + + The Bartlett window is very similar to a triangular window, except that the + end points are at zero. It is often used in signal processing for tapering + a signal, without generating too much ripple in the frequency domain. + + For full documentation refer to :obj:`numpy.bartlett`. + + 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 triangular window, with the maximum value normalized to one + (the value one appears only if the number of samples is odd), with the + first and last samples equal to zero. + + See Also + -------- + :obj:`dpnp.blackman` : Return the Blackman window. + :obj:`dpnp.hamming` : Return the Hamming window. + :obj:`dpnp.hanning` : Return the Hanning window. + :obj:`dpnp.kaiser` : Return the Kaiser window. + + Notes + ----- + The Bartlett window is defined as + + .. math:: w(n) = \frac{2}{M-1} \left(\frac{M-1}{2} - + \left|n - \frac{M-1}{2}\right|\right) + + Examples + -------- + >>> import dpnp as np + >>> np.bartlett(12) + array([0. , 0.18181818, 0.36363636, 0.54545455, 0.72727273, + 0.90909091, 0.90909091, 0.72727273, 0.54545455, 0.36363636, + 0.18181818, 0. ]) + + Creating the output array on a different device or with a + specified usm_type: + + >>> x = np.bartlett(4) # default case + >>> x, x.device, x.usm_type + (array([0. , 0.66666667, 0.66666667, 0. ]), + Device(level_zero:gpu:0), + 'device') + + >>> y = np.bartlett(4, device="cpu") + >>> y, y.device, y.usm_type + (array([0. , 0.66666667, 0.66666667, 0. ]), + Device(opencl:cpu:0), + 'device') + + >>> z = np.bartlett(4, usm_type="host") + >>> z, z.device, z.usm_type + (array([0. , 0.66666667, 0.66666667, 0. ]), + Device(level_zero:gpu:0), + 'host') + + """ + + return _call_window_kernel( + M, wi._bartlett, 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 4accd1123083..133832fa2078 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("bartlett", [10], {}), pytest.param("blackman", [10], {}), pytest.param("eye", [4, 2], {}), pytest.param("empty", [(2, 2)], {}), diff --git a/dpnp/tests/test_usm_type.py b/dpnp/tests/test_usm_type.py index ff5e0036ef81..eabd07daf840 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("bartlett", [10], {}), pytest.param("blackman", [10], {}), pytest.param("eye", [4, 2], {}), pytest.param("empty", [(3, 4)], {}), diff --git a/dpnp/tests/test_window.py b/dpnp/tests/test_window.py index 9529a882beb6..3fa22625aded 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", ["blackman", "hamming", "hanning"]) +@pytest.mark.parametrize("func", ["bartlett", "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", ["blackman", "hamming", "hanning"]) +@pytest.mark.parametrize("func", ["bartlett", "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 27ce810bdb4e..fa59e19b9328 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,7 @@ *testing.product( { "m": [0, 1, -1, 1024], - # TODO: add ["bartlett"] when supported - "name": ["blackman", "hamming", "hanning"], + "name": ["bartlett", "blackman", "hamming", "hanning"], } ) ) From 3f2576e9ae8a7f7b39b84b5b2ecb0a0d00982154 Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Thu, 6 Mar 2025 13:53:34 -0800 Subject: [PATCH 2/4] update changelog --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 5972b12767fa..d4f0756730d5 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -11,6 +11,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) +* Added implementation of `dpnp.bartlett` [#2366](https://github.com/IntelPython/dpnp/pull/2366) ### Changed From 5c6570df360c47a78b4bc020188e15a5c788f234 Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Thu, 6 Mar 2025 14:21:45 -0800 Subject: [PATCH 3/4] update docstring and change order --- dpnp/dpnp_iface_window.py | 95 ++++++++++++++++++++------------------- 1 file changed, 48 insertions(+), 47 deletions(-) diff --git a/dpnp/dpnp_iface_window.py b/dpnp/dpnp_iface_window.py index 16349c76a7cc..3a53162c1875 100644 --- a/dpnp/dpnp_iface_window.py +++ b/dpnp/dpnp_iface_window.py @@ -81,15 +81,15 @@ def _call_window_kernel( return result -def blackman(M, device=None, usm_type=None, sycl_queue=None): +def bartlett(M, device=None, usm_type=None, sycl_queue=None): r""" - Return the Blackman window. + Return the Bartlett 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. + The Bartlett window is very similar to a triangular window, except that the + end points are at zero. It is often used in signal processing for tapering + a signal, without generating too much ripple in the frequency domain. - For full documentation refer to :obj:`numpy.blackman`. + For full documentation refer to :obj:`numpy.bartlett`. Parameters ---------- @@ -120,69 +120,70 @@ def blackman(M, device=None, usm_type=None, sycl_queue=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). + The triangular window, with the maximum value normalized to one + (the value one appears only if the number of samples is odd), with the + first and last samples equal to zero. 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.hanning` : Return the Hanning window. :obj:`dpnp.kaiser` : Return the Kaiser window. Notes ----- - The Blackman window is defined as + The Bartlett 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) + .. math:: w(n) = \frac{2}{M-1} \left(\frac{M-1}{2} - + \left|n - \frac{M-1}{2}\right|\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]) + >>> np.bartlett(12) + array([0. , 0.18181818, 0.36363636, 0.54545455, 0.72727273, + 0.90909091, 0.90909091, 0.72727273, 0.54545455, 0.36363636, + 0.18181818, 0. ]) Creating the output array on a different device or with a specified usm_type: - >>> x = np.blackman(3) # default case + >>> x = np.bartlett(4) # default case >>> x, x.device, x.usm_type - (array([-1.38777878e-17, 1.00000000e+00, -1.38777878e-17]), + (array([0. , 0.66666667, 0.66666667, 0. ]), Device(level_zero:gpu:0), 'device') - >>> y = np.blackman(3, device="cpu") + >>> y = np.bartlett(4, device="cpu") >>> y, y.device, y.usm_type - (array([-1.38777878e-17, 1.00000000e+00, -1.38777878e-17]), + (array([0. , 0.66666667, 0.66666667, 0. ]), Device(opencl:cpu:0), 'device') - >>> z = np.blackman(3, usm_type="host") + >>> z = np.bartlett(4, usm_type="host") >>> z, z.device, z.usm_type - (array([-1.38777878e-17, 1.00000000e+00, -1.38777878e-17]), + (array([0. , 0.66666667, 0.66666667, 0. ]), Device(level_zero:gpu:0), 'host') """ return _call_window_kernel( - M, wi._blackman, device=device, usm_type=usm_type, sycl_queue=sycl_queue + M, wi._bartlett, device=device, usm_type=usm_type, sycl_queue=sycl_queue ) -def bartlett(M, device=None, usm_type=None, sycl_queue=None): +def blackman(M, device=None, usm_type=None, sycl_queue=None): r""" - Return the Bartlett window. + Return the Blackman window. - The Bartlett window is very similar to a triangular window, except that the - end points are at zero. It is often used in signal processing for tapering - a signal, without generating too much ripple in the frequency domain. + 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.bartlett`. + For full documentation refer to :obj:`numpy.blackman`. Parameters ---------- @@ -213,57 +214,57 @@ def bartlett(M, device=None, usm_type=None, sycl_queue=None): Returns ------- out : dpnp.ndarray of shape (M,) - The triangular window, with the maximum value normalized to one - (the value one appears only if the number of samples is odd), with the - first and last samples equal to zero. + 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.blackman` : Return the Blackman window. + :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 Bartlett window is defined as + The Blackman window is defined as - .. math:: w(n) = \frac{2}{M-1} \left(\frac{M-1}{2} - - \left|n - \frac{M-1}{2}\right|\right) + .. 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.bartlett(12) - array([0. , 0.18181818, 0.36363636, 0.54545455, 0.72727273, - 0.90909091, 0.90909091, 0.72727273, 0.54545455, 0.36363636, - 0.18181818, 0. ]) + >>> 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.bartlett(4) # default case + >>> x = np.blackman(3) # default case >>> x, x.device, x.usm_type - (array([0. , 0.66666667, 0.66666667, 0. ]), + (array([-1.38777878e-17, 1.00000000e+00, -1.38777878e-17]), Device(level_zero:gpu:0), 'device') - >>> y = np.bartlett(4, device="cpu") + >>> y = np.blackman(3, device="cpu") >>> y, y.device, y.usm_type - (array([0. , 0.66666667, 0.66666667, 0. ]), + (array([-1.38777878e-17, 1.00000000e+00, -1.38777878e-17]), Device(opencl:cpu:0), 'device') - >>> z = np.bartlett(4, usm_type="host") + >>> z = np.blackman(3, usm_type="host") >>> z, z.device, z.usm_type - (array([0. , 0.66666667, 0.66666667, 0. ]), + (array([-1.38777878e-17, 1.00000000e+00, -1.38777878e-17]), Device(level_zero:gpu:0), 'host') """ return _call_window_kernel( - M, wi._bartlett, device=device, usm_type=usm_type, sycl_queue=sycl_queue + M, wi._blackman, device=device, usm_type=usm_type, sycl_queue=sycl_queue ) From 9434ebd8445a59eea7f7d4f7a65588731142b3bb Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Fri, 7 Mar 2025 05:40:30 -0800 Subject: [PATCH 4/4] address comment --- dpnp/backend/extensions/window/bartlett.hpp | 4 ++-- dpnp/backend/extensions/window/blackman.hpp | 5 +++-- 2 files changed, 5 insertions(+), 4 deletions(-) diff --git a/dpnp/backend/extensions/window/bartlett.hpp b/dpnp/backend/extensions/window/bartlett.hpp index 0ade4722b8d9..00c24bb48521 100644 --- a/dpnp/backend/extensions/window/bartlett.hpp +++ b/dpnp/backend/extensions/window/bartlett.hpp @@ -45,8 +45,8 @@ class BartlettFunctor { const auto i = id.get(0); - data[i] = - T(2) / (N - 1) * ((N - 1) / T(2) - sycl::fabs(i - (N - 1) / T(2))); + const T alpha = (N - 1) / T(2); + data[i] = T(1) - sycl::fabs(i - alpha) / alpha; } }; diff --git a/dpnp/backend/extensions/window/blackman.hpp b/dpnp/backend/extensions/window/blackman.hpp index 4185decb03b9..82b6243a9434 100644 --- a/dpnp/backend/extensions/window/blackman.hpp +++ b/dpnp/backend/extensions/window/blackman.hpp @@ -45,8 +45,9 @@ class BlackmanFunctor { 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)); + const T alpha = T(2) * i / (N - 1); + data[i] = T(0.42) - T(0.5) * sycl::cospi(alpha) + + T(0.08) * sycl::cospi(T(2) * alpha); } };