From c8a2822dc775e30fb0fbccd5c8a31533d42acb78 Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Fri, 7 Mar 2025 07:04:14 -0800 Subject: [PATCH 1/7] implement dpnp.kaiser --- dpnp/backend/extensions/window/CMakeLists.txt | 1 + dpnp/backend/extensions/window/kaiser.cpp | 180 ++++++++++++++++++ dpnp/backend/extensions/window/kaiser.hpp | 41 ++++ dpnp/backend/extensions/window/window_py.cpp | 9 + dpnp/dpnp_iface_window.py | 153 ++++++++++++++- dpnp/tests/test_sycl_queue.py | 1 + dpnp/tests/test_usm_type.py | 1 + dpnp/tests/test_window.py | 61 +++++- .../cupy/math_tests/test_window.py | 6 +- 9 files changed, 439 insertions(+), 14 deletions(-) create mode 100644 dpnp/backend/extensions/window/kaiser.cpp create mode 100644 dpnp/backend/extensions/window/kaiser.hpp diff --git a/dpnp/backend/extensions/window/CMakeLists.txt b/dpnp/backend/extensions/window/CMakeLists.txt index 2744ceb027c..a1c329f8a20 100644 --- a/dpnp/backend/extensions/window/CMakeLists.txt +++ b/dpnp/backend/extensions/window/CMakeLists.txt @@ -26,6 +26,7 @@ set(python_module_name _window_impl) set(_module_src + ${CMAKE_CURRENT_SOURCE_DIR}/kaiser.cpp ${CMAKE_CURRENT_SOURCE_DIR}/window_py.cpp ) diff --git a/dpnp/backend/extensions/window/kaiser.cpp b/dpnp/backend/extensions/window/kaiser.cpp new file mode 100644 index 00000000000..85332fcbc16 --- /dev/null +++ b/dpnp/backend/extensions/window/kaiser.cpp @@ -0,0 +1,180 @@ +//***************************************************************************** +// 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. +//***************************************************************************** + +#include "kaiser.hpp" +#include "utils/output_validation.hpp" +#include "utils/type_dispatch.hpp" +#include "utils/type_utils.hpp" +#include + +/** + * Version of SYCL DPC++ 2025.1 compiler where an issue with + * sycl::ext::intel::math::cyl_bessel_i0(x) is fully resolved. + */ +#ifndef __SYCL_COMPILER_BESSEL_I0_SUPPORT +#define __SYCL_COMPILER_BESSEL_I0_SUPPORT 20241208L +#endif + +#if __SYCL_COMPILER_VERSION >= __SYCL_COMPILER_BESSEL_I0_SUPPORT +#include +#endif + +#include "../kernels/elementwise_functions/i0.hpp" + +namespace dpnp::extensions::window +{ +namespace dpctl_td_ns = dpctl::tensor::type_dispatch; + +typedef sycl::event (*kaiser_fn_ptr_t)(sycl::queue &, + char *, + const std::size_t, + const float, + const std::vector &); + +static kaiser_fn_ptr_t kaiser_dispatch_vector[dpctl_td_ns::num_types]; + +template +class KaiserFunctor +{ +private: + T *data = nullptr; + const std::size_t N; + const float beta; + +public: + KaiserFunctor(T *data, const std::size_t N, const float beta) + : data(data), N(N), beta(beta) + { + } + + void operator()(sycl::id<1> id) const + { +#if __SYCL_COMPILER_VERSION >= __SYCL_COMPILER_BESSEL_I0_SUPPORT + using sycl::ext::intel::math::cyl_bessel_i0; +#else + using dpnp::kernels::i0::impl::cyl_bessel_i0; +#endif + + const auto i = id.get(0); + const T alpha = (N - 1) / T(2); + const T tmp = (i - alpha) / alpha; + data[i] = cyl_bessel_i0(beta * sycl::sqrt(1 - tmp * tmp)) / + cyl_bessel_i0(beta); + } +}; + +template class Functor> +sycl::event kaiser_impl(sycl::queue &q, + char *result, + const std::size_t nelems, + const float beta, + const std::vector &depends) +{ + dpctl::tensor::type_utils::validate_type_for_device(q); + + T *res = reinterpret_cast(result); + + sycl::event kaiser_ev = q.submit([&](sycl::handler &cgh) { + cgh.depends_on(depends); + + using KaiserKernel = Functor; + cgh.parallel_for(sycl::range<1>(nelems), + KaiserKernel(res, nelems, beta)); + }); + + return kaiser_ev; +} + +template +struct KaiserFactory +{ + fnT get() + { + if constexpr (std::is_floating_point_v) { + return kaiser_impl; + } + else { + return nullptr; + } + } +}; + +std::pair + py_kaiser(sycl::queue &exec_q, + const float beta, + const dpctl::tensor::usm_ndarray &result, + const std::vector &depends) +{ + dpctl::tensor::validation::CheckWritable::throw_if_not_writable(result); + + int nd = result.get_ndim(); + if (nd != 1) { + throw py::value_error("Array should be 1d"); + } + + if (!dpctl::utils::queues_are_compatible(exec_q, {result.get_queue()})) { + throw py::value_error( + "Execution queue is not compatible with allocation queue."); + } + + const bool is_result_c_contig = result.is_c_contiguous(); + if (!is_result_c_contig) { + throw py::value_error("The result input array is not c-contiguous."); + } + + size_t nelems = result.get_size(); + if (nelems == 0) { + return std::make_pair(sycl::event{}, sycl::event{}); + } + + int result_typenum = result.get_typenum(); + auto array_types = dpctl_td_ns::usm_ndarray_types(); + int result_type_id = array_types.typenum_to_lookup_id(result_typenum); + auto fn = kaiser_dispatch_vector[result_type_id]; + + if (fn == nullptr) { + throw std::runtime_error("Type of given array is not supported"); + } + + char *result_typeless_ptr = result.get_data(); + sycl::event kaiser_ev = + fn(exec_q, result_typeless_ptr, nelems, beta, depends); + sycl::event args_ev = + dpctl::utils::keep_args_alive(exec_q, {result}, {kaiser_ev}); + + return std::make_pair(args_ev, kaiser_ev); +} + +void init_kaiser_dispatch_vectors() +{ + dpctl_td_ns::DispatchVectorBuilder + contig; + contig.populate_dispatch_vector(kaiser_dispatch_vector); + + return; +} + +} // namespace dpnp::extensions::window diff --git a/dpnp/backend/extensions/window/kaiser.hpp b/dpnp/backend/extensions/window/kaiser.hpp new file mode 100644 index 00000000000..830f88e42c0 --- /dev/null +++ b/dpnp/backend/extensions/window/kaiser.hpp @@ -0,0 +1,41 @@ +//***************************************************************************** +// 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 +#include + +namespace dpnp::extensions::window +{ +extern std::pair + py_kaiser(sycl::queue &exec_q, + const float beta, + const dpctl::tensor::usm_ndarray &result, + const std::vector &depends); + +extern void init_kaiser_dispatch_vectors(void); + +} // namespace dpnp::extensions::window diff --git a/dpnp/backend/extensions/window/window_py.cpp b/dpnp/backend/extensions/window/window_py.cpp index 3e4ff884907..7c471d4d648 100644 --- a/dpnp/backend/extensions/window/window_py.cpp +++ b/dpnp/backend/extensions/window/window_py.cpp @@ -35,6 +35,7 @@ #include "common.hpp" #include "hamming.hpp" #include "hanning.hpp" +#include "kaiser.hpp" namespace window_ns = dpnp::extensions::window; namespace py = pybind11; @@ -111,4 +112,12 @@ PYBIND11_MODULE(_window_impl, m) py::arg("sycl_queue"), py::arg("result"), py::arg("depends") = py::list()); } + + { + window_ns::init_kaiser_dispatch_vectors(); + + m.def("_kaiser", window_ns::py_kaiser, "Call Kaiser kernel", + py::arg("sycl_queue"), py::arg("beta"), py::arg("result"), + py::arg("depends") = py::list()); + } } diff --git a/dpnp/dpnp_iface_window.py b/dpnp/dpnp_iface_window.py index 3a53162c187..87f329c173e 100644 --- a/dpnp/dpnp_iface_window.py +++ b/dpnp/dpnp_iface_window.py @@ -45,11 +45,11 @@ import dpnp import dpnp.backend.extensions.window._window_impl as wi -__all__ = ["bartlett", "blackman", "hamming", "hanning"] +__all__ = ["bartlett", "blackman", "hamming", "hanning", "kaiser"] def _call_window_kernel( - M, _window_kernel, device=None, usm_type=None, sycl_queue=None + M, _window_kernel, device=None, usm_type=None, sycl_queue=None, beta=None ): try: @@ -72,16 +72,26 @@ def _call_window_kernel( exec_q = result.sycl_queue _manager = dpu.SequentialOrderManager[exec_q] - ht_ev, win_ev = _window_kernel( - exec_q, dpnp.get_usm_ndarray(result), depends=_manager.submitted_events - ) + if beta is None: + ht_ev, win_ev = _window_kernel( + exec_q, + dpnp.get_usm_ndarray(result), + depends=_manager.submitted_events, + ) + else: + ht_ev, win_ev = _window_kernel( + exec_q, + beta, + dpnp.get_usm_ndarray(result), + depends=_manager.submitted_events, + ) _manager.add_event_pair(ht_ev, win_ev) return result -def bartlett(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. @@ -175,7 +185,7 @@ def bartlett(M, device=None, usm_type=None, sycl_queue=None): ) -def blackman(M, device=None, usm_type=None, sycl_queue=None): +def blackman(M, *, device=None, usm_type=None, sycl_queue=None): r""" Return the Blackman window. @@ -268,7 +278,7 @@ def blackman(M, device=None, usm_type=None, sycl_queue=None): ) -def hamming(M, device=None, usm_type=None, sycl_queue=None): +def hamming(M, *, device=None, usm_type=None, sycl_queue=None): r""" Return the Hamming window. @@ -352,7 +362,7 @@ def hamming(M, device=None, usm_type=None, sycl_queue=None): ) -def hanning(M, device=None, usm_type=None, sycl_queue=None): +def hanning(M, *, device=None, usm_type=None, sycl_queue=None): r""" Return the Hanning window. @@ -434,3 +444,128 @@ def hanning(M, device=None, usm_type=None, sycl_queue=None): return _call_window_kernel( M, wi._hanning, device=device, usm_type=usm_type, sycl_queue=sycl_queue ) + + +def kaiser(M, beta, *, device=None, usm_type=None, sycl_queue=None): + r""" + Return the Kaiser window. + + The Kaiser window is a taper formed by using a Bessel function. + + For full documentation refer to :obj:`numpy.kaiser`. + + Parameters + ---------- + M : int + Number of points in the output window. If zero or less, an empty array + is returned. + beta : float + Shape parameter for window. + 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.hanning` : Return the Hanning window. + + Notes + ----- + The Kaiser window is defined as + + .. math:: w(n) = I_0\left( \beta \sqrt{1-\frac{4n^2}{(M-1)^2}} + \right)/I_0(\beta) + + with + + .. math:: \quad -\frac{M-1}{2} \leq n \leq \frac{M-1}{2}, + + where :math:`I_0` is the modified zeroth-order Bessel function. + + The Kaiser can approximate many other windows by varying the beta + parameter. + + ==== ======================= + beta Window shape + ==== ======================= + 0 Rectangular + 5 Similar to a Hamming + 6 Similar to a Hanning + 8.6 Similar to a Blackman + ==== ======================= + + A beta value of 14 is probably a good starting point. Note that as beta + gets large, the window narrows, and so the number of samples needs to be + large enough to sample the increasingly narrow spike, otherwise NaNs will + get returned. + + Examples + -------- + >>> import dpnp as np + >>> np.kaiser(12, 14) + array([7.72686638e-06, 3.46009173e-03, 4.65200161e-02, 2.29737107e-01, + 5.99885281e-01, 9.45674843e-01, 9.45674843e-01, 5.99885281e-01, + 2.29737107e-01, 4.65200161e-02, 3.46009173e-03, 7.72686638e-06]) + + Creating the output array on a different device or with a + specified usm_type: + + >>> x = np.kaiser(3, 14) # default case + >>> x, x.device, x.usm_type + (array([7.72686638e-06, 9.99999941e-01, 7.72686638e-06]), + Device(level_zero:gpu:0), + 'device') + + >>> y = np.kaiser(3, 14, device="cpu") + >>> y, y.device, y.usm_type + (array([7.72686638e-06, 9.99999941e-01, 7.72686638e-06]), + Device(opencl:cpu:0), + 'device') + + >>> z = np.kaiser(3, 14, usm_type="host") + >>> z, z.device, z.usm_type + (array([7.72686638e-06, 9.99999941e-01, 7.72686638e-06]), + Device(level_zero:gpu:0), + 'host') + + """ + + try: + beta = float(beta) + except Exception as e: + raise TypeError("beta must be a float") from e + + return _call_window_kernel( + M, + wi._kaiser, + device=device, + usm_type=usm_type, + sycl_queue=sycl_queue, + beta=beta, + ) diff --git a/dpnp/tests/test_sycl_queue.py b/dpnp/tests/test_sycl_queue.py index 133832fa207..b6ab0ace40e 100644 --- a/dpnp/tests/test_sycl_queue.py +++ b/dpnp/tests/test_sycl_queue.py @@ -72,6 +72,7 @@ def assert_sycl_queue_equal(result, expected): pytest.param("hamming", [10], {}), pytest.param("hanning", [10], {}), pytest.param("identity", [4], {}), + pytest.param("kaiser", [10], {"beta": 14}), pytest.param("linspace", [0, 4, 8], {}), pytest.param("logspace", [0, 4, 8], {}), pytest.param("logspace", [0, 4, 8], {"base": [10]}), diff --git a/dpnp/tests/test_usm_type.py b/dpnp/tests/test_usm_type.py index eabd07daf84..3f922dbbdac 100644 --- a/dpnp/tests/test_usm_type.py +++ b/dpnp/tests/test_usm_type.py @@ -196,6 +196,7 @@ def test_array_creation_from_array(func, args, usm_type_x, usm_type_y): pytest.param("hamming", [10], {}), pytest.param("hanning", [10], {}), pytest.param("identity", [4], {}), + pytest.param("kaiser", [10], {"beta": 14}), pytest.param("linspace", [0, 4, 8], {}), pytest.param("logspace", [0, 4, 8], {}), pytest.param("logspace", [0, 4, 8], {"base": [10]}), diff --git a/dpnp/tests/test_window.py b/dpnp/tests/test_window.py index 3fa22625ade..4f8838f3923 100644 --- a/dpnp/tests/test_window.py +++ b/dpnp/tests/test_window.py @@ -1,6 +1,6 @@ import numpy import pytest -from numpy.testing import assert_raises +from numpy.testing import assert_allclose, assert_raises import dpnp @@ -46,3 +46,62 @@ def test_window(func, M): ) def test_window_error(func, M): assert_raises(TypeError, getattr(dpnp, func), M) + + +class TestKaiser: + @pytest.mark.parametrize( + "M", + [ + True, + False, + 0, + dpnp.int32(1), + 4, + 5.0, + dpnp.float32(6), + dpnp.array(7), + numpy.array(8.0), + ], + ) + def test_kaiser_M(self, M): + result = dpnp.kaiser(M, 14) + + if isinstance(M, dpnp.ndarray): + M = M.asnumpy() + expected = numpy.kaiser(M, 14) + + assert_allclose(result, expected, rtol=1e-6, atol=1e-6) + + @pytest.mark.filterwarnings("ignore::RuntimeWarning") + @pytest.mark.parametrize( + "beta", + [ + True, + False, + 14, + dpnp.int32(14), + 14, + -14.0, + dpnp.float32(14), + dpnp.array(-14), + numpy.array(14.0), + numpy.inf, + numpy.array(-numpy.inf), + dpnp.array(dpnp.nan), + ], + ) + def test_kaiser_beta(self, beta): + result = dpnp.kaiser(4, beta) + + if isinstance(beta, dpnp.ndarray): + beta = beta.asnumpy() + expected = numpy.kaiser(4, beta) + + assert_allclose(result, expected, rtol=1e-6) + + @pytest.mark.parametrize( + "beta", + [5 + 4j, numpy.array(5 + 4j), dpnp.array([5, 3])], + ) + def test_kaiser_error(self, beta): + assert_raises(TypeError, dpnp.kaiser, 4, beta) 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 fa59e19b932..f6fdd05855e 100644 --- a/dpnp/tests/third_party/cupy/math_tests/test_window.py +++ b/dpnp/tests/third_party/cupy/math_tests/test_window.py @@ -32,8 +32,7 @@ def test_window(self, xp): ) class TestKaiser(unittest.TestCase): - @pytest.mark.skip("kaiser function is not supported yet") - @testing.numpy_cupy_allclose(atol=1e-5) + @testing.numpy_cupy_allclose(atol=1e-5, type_check=has_support_aspect64()) def test_kaiser_parametric(self, xp): return getattr(xp, self.name)(self.m, self.beta) @@ -41,7 +40,6 @@ def test_kaiser_parametric(self, xp): @testing.parameterize(*testing.product({"m": [-1, 0, 1]})) class TestKaiserBoundary(unittest.TestCase): - @pytest.mark.skip("kaiser function is not supported yet") - @testing.numpy_cupy_allclose(atol=1e-5) + @testing.numpy_cupy_allclose(atol=1e-5, type_check=has_support_aspect64()) def test_kaiser(self, xp): return xp.kaiser(self.m, 1.5) From 15ce51d67ba7651fef7178ce1abf54abaa210014 Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Tue, 25 Mar 2025 13:21:03 -0700 Subject: [PATCH 2/7] update changelog --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 96c51515161..a29d4c1205c 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -12,6 +12,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 * 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) +* Added implementation of `dpnp.kaiser` [#2387](https://github.com/IntelPython/dpnp/pull/2387) ### Changed From a3d191c7c2bc292a68e3c8df9a584f4de267eacf Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad <120411540+vtavana@users.noreply.github.com> Date: Mon, 31 Mar 2025 11:57:28 -0500 Subject: [PATCH 3/7] Update dpnp/dpnp_iface_window.py Co-authored-by: Anton <100830759+antonwolfy@users.noreply.github.com> --- dpnp/dpnp_iface_window.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dpnp/dpnp_iface_window.py b/dpnp/dpnp_iface_window.py index 87f329c173e..6850197443e 100644 --- a/dpnp/dpnp_iface_window.py +++ b/dpnp/dpnp_iface_window.py @@ -520,7 +520,7 @@ def kaiser(M, beta, *, device=None, usm_type=None, sycl_queue=None): 8.6 Similar to a Blackman ==== ======================= - A beta value of 14 is probably a good starting point. Note that as beta + A beta value of ``14`` is probably a good starting point. Note that as beta gets large, the window narrows, and so the number of samples needs to be large enough to sample the increasingly narrow spike, otherwise NaNs will get returned. From 6fe1e2bca432a4a7be0f4b231ef1717901982b0d Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Mon, 31 Mar 2025 10:17:36 -0700 Subject: [PATCH 4/7] pass beta as python object --- dpnp/backend/extensions/window/kaiser.cpp | 13 +++++++------ dpnp/backend/extensions/window/kaiser.hpp | 2 +- 2 files changed, 8 insertions(+), 7 deletions(-) diff --git a/dpnp/backend/extensions/window/kaiser.cpp b/dpnp/backend/extensions/window/kaiser.cpp index 85332fcbc16..d0d747d0421 100644 --- a/dpnp/backend/extensions/window/kaiser.cpp +++ b/dpnp/backend/extensions/window/kaiser.cpp @@ -50,7 +50,7 @@ namespace dpctl_td_ns = dpctl::tensor::type_dispatch; typedef sycl::event (*kaiser_fn_ptr_t)(sycl::queue &, char *, const std::size_t, - const float, + const py::object &, const std::vector &); static kaiser_fn_ptr_t kaiser_dispatch_vector[dpctl_td_ns::num_types]; @@ -61,10 +61,10 @@ class KaiserFunctor private: T *data = nullptr; const std::size_t N; - const float beta; + const T beta; public: - KaiserFunctor(T *data, const std::size_t N, const float beta) + KaiserFunctor(T *data, const std::size_t N, const T beta) : data(data), N(N), beta(beta) { } @@ -89,12 +89,13 @@ template class Functor> sycl::event kaiser_impl(sycl::queue &q, char *result, const std::size_t nelems, - const float beta, + const py::object &py_beta, const std::vector &depends) { dpctl::tensor::type_utils::validate_type_for_device(q); T *res = reinterpret_cast(result); + const T beta = py::cast(py_beta); sycl::event kaiser_ev = q.submit([&](sycl::handler &cgh) { cgh.depends_on(depends); @@ -123,7 +124,7 @@ struct KaiserFactory std::pair py_kaiser(sycl::queue &exec_q, - const float beta, + const py::object &py_beta, const dpctl::tensor::usm_ndarray &result, const std::vector &depends) { @@ -160,7 +161,7 @@ std::pair char *result_typeless_ptr = result.get_data(); sycl::event kaiser_ev = - fn(exec_q, result_typeless_ptr, nelems, beta, depends); + fn(exec_q, result_typeless_ptr, nelems, py_beta, depends); sycl::event args_ev = dpctl::utils::keep_args_alive(exec_q, {result}, {kaiser_ev}); diff --git a/dpnp/backend/extensions/window/kaiser.hpp b/dpnp/backend/extensions/window/kaiser.hpp index 830f88e42c0..421616a8768 100644 --- a/dpnp/backend/extensions/window/kaiser.hpp +++ b/dpnp/backend/extensions/window/kaiser.hpp @@ -32,7 +32,7 @@ namespace dpnp::extensions::window { extern std::pair py_kaiser(sycl::queue &exec_q, - const float beta, + const py::object &beta, const dpctl::tensor::usm_ndarray &result, const std::vector &depends); From 672d0f8d4fd497206caacc327dd105c4a95d71ed Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Mon, 31 Mar 2025 10:44:28 -0700 Subject: [PATCH 5/7] reuse init_window_dispatch_vectors for kaiser func --- dpnp/backend/extensions/window/common.hpp | 10 ++++++---- dpnp/backend/extensions/window/kaiser.cpp | 11 +++++------ dpnp/backend/extensions/window/window_py.cpp | 12 ++++++++---- 3 files changed, 19 insertions(+), 14 deletions(-) diff --git a/dpnp/backend/extensions/window/common.hpp b/dpnp/backend/extensions/window/common.hpp index e86c52c6081..d604497b2ee 100644 --- a/dpnp/backend/extensions/window/common.hpp +++ b/dpnp/backend/extensions/window/common.hpp @@ -67,7 +67,7 @@ sycl::event window_impl(sycl::queue &q, return window_ev; } -std::pair +inline std::pair py_window(sycl::queue &exec_q, const dpctl::tensor::usm_ndarray &result, const std::vector &depends, @@ -112,10 +112,12 @@ std::pair return std::make_pair(args_ev, window_ev); } -template