From 61ea04ac0d907ad2927388a9ec21204037bec410 Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Fri, 21 Feb 2025 10:04:00 -0800 Subject: [PATCH 1/7] implement dpnp.hamming --- doc/reference/routines.rst | 1 + doc/reference/window.rst | 17 ++ dpnp/CMakeLists.txt | 8 +- .../extensions/indexing/CMakeLists.txt | 2 +- dpnp/backend/extensions/indexing/choose.cpp | 2 +- dpnp/backend/extensions/indexing/choose.hpp | 2 +- .../extensions/indexing/choose_kernel.hpp | 2 +- .../extensions/indexing/indexing_py.cpp | 3 +- .../extensions/statistics/CMakeLists.txt | 2 +- .../extensions/statistics/statistics_py.cpp | 1 - dpnp/backend/extensions/ufunc/ufunc_py.cpp | 1 - dpnp/backend/extensions/window/CMakeLists.txt | 90 +++++++++ dpnp/backend/extensions/window/hamming.cpp | 116 ++++++++++++ dpnp/backend/extensions/window/hamming.hpp | 35 ++++ .../extensions/window/hamming_kernel.hpp | 82 ++++++++ dpnp/backend/extensions/window/window_py.cpp | 37 ++++ dpnp/dpnp_iface.py | 3 + dpnp/dpnp_iface_window.py | 176 ++++++++++++++++++ dpnp/tests/test_sycl_queue.py | 15 ++ dpnp/tests/test_usm_type.py | 11 ++ dpnp/tests/test_window.py | 50 +++++ .../cupy/math_tests/test_window.py | 10 +- 22 files changed, 650 insertions(+), 16 deletions(-) create mode 100644 doc/reference/window.rst create mode 100644 dpnp/backend/extensions/window/CMakeLists.txt create mode 100644 dpnp/backend/extensions/window/hamming.cpp create mode 100644 dpnp/backend/extensions/window/hamming.hpp create mode 100644 dpnp/backend/extensions/window/hamming_kernel.hpp create mode 100644 dpnp/backend/extensions/window/window_py.cpp create mode 100644 dpnp/dpnp_iface_window.py create mode 100644 dpnp/tests/test_window.py diff --git a/doc/reference/routines.rst b/doc/reference/routines.rst index 3b91283c8e07..1dd4a205b0cf 100644 --- a/doc/reference/routines.rst +++ b/doc/reference/routines.rst @@ -27,3 +27,4 @@ These functions cover a subset of set sort statistics + window diff --git a/doc/reference/window.rst b/doc/reference/window.rst new file mode 100644 index 000000000000..9b943b1ff5a8 --- /dev/null +++ b/doc/reference/window.rst @@ -0,0 +1,17 @@ +Window functions +================ + +.. https://numpy.org/doc/stable/reference/routines.window.html + +Various windows +--------------- + +.. autosummary:: + :toctree: generated/ + :nosignatures: + + dpnp.bartlett + dpnp.blackman + dpnp.hamming + dpnp.hanning + dpnp.kaiser diff --git a/dpnp/CMakeLists.txt b/dpnp/CMakeLists.txt index 298648378256..6be90d849dc4 100644 --- a/dpnp/CMakeLists.txt +++ b/dpnp/CMakeLists.txt @@ -58,11 +58,13 @@ endfunction() add_subdirectory(backend) add_subdirectory(backend/extensions/blas) add_subdirectory(backend/extensions/fft) +add_subdirectory(backend/extensions/indexing) add_subdirectory(backend/extensions/lapack) -add_subdirectory(backend/extensions/vm) -add_subdirectory(backend/extensions/ufunc) add_subdirectory(backend/extensions/statistics) -add_subdirectory(backend/extensions/indexing) +add_subdirectory(backend/extensions/ufunc) +add_subdirectory(backend/extensions/vm) +add_subdirectory(backend/extensions/window) + add_subdirectory(dpnp_algo) add_subdirectory(dpnp_utils) diff --git a/dpnp/backend/extensions/indexing/CMakeLists.txt b/dpnp/backend/extensions/indexing/CMakeLists.txt index 34f4fe380748..df90ea36f256 100644 --- a/dpnp/backend/extensions/indexing/CMakeLists.txt +++ b/dpnp/backend/extensions/indexing/CMakeLists.txt @@ -1,5 +1,5 @@ # ***************************************************************************** -# Copyright (c) 2016-2024, Intel Corporation +# Copyright (c) 2025, Intel Corporation # All rights reserved. # # Redistribution and use in source and binary forms, with or without diff --git a/dpnp/backend/extensions/indexing/choose.cpp b/dpnp/backend/extensions/indexing/choose.cpp index 83008a6795ad..e0fd0767aca2 100644 --- a/dpnp/backend/extensions/indexing/choose.cpp +++ b/dpnp/backend/extensions/indexing/choose.cpp @@ -1,5 +1,5 @@ //***************************************************************************** -// Copyright (c) 2024, Intel Corporation +// Copyright (c) 2025, Intel Corporation // All rights reserved. // // Redistribution and use in source and binary forms, with or without diff --git a/dpnp/backend/extensions/indexing/choose.hpp b/dpnp/backend/extensions/indexing/choose.hpp index 63b155ec6463..92298890f5ab 100644 --- a/dpnp/backend/extensions/indexing/choose.hpp +++ b/dpnp/backend/extensions/indexing/choose.hpp @@ -1,5 +1,5 @@ //***************************************************************************** -// Copyright (c) 2024, Intel Corporation +// Copyright (c) 2025, Intel Corporation // All rights reserved. // // Redistribution and use in source and binary forms, with or without diff --git a/dpnp/backend/extensions/indexing/choose_kernel.hpp b/dpnp/backend/extensions/indexing/choose_kernel.hpp index 0798239eac14..2f20ea5913f3 100644 --- a/dpnp/backend/extensions/indexing/choose_kernel.hpp +++ b/dpnp/backend/extensions/indexing/choose_kernel.hpp @@ -1,5 +1,5 @@ //***************************************************************************** -// Copyright (c) 2024, Intel Corporation +// Copyright (c) 2025, Intel Corporation // All rights reserved. // // Redistribution and use in source and binary forms, with or without diff --git a/dpnp/backend/extensions/indexing/indexing_py.cpp b/dpnp/backend/extensions/indexing/indexing_py.cpp index 143202477751..e4c0bcee10ee 100644 --- a/dpnp/backend/extensions/indexing/indexing_py.cpp +++ b/dpnp/backend/extensions/indexing/indexing_py.cpp @@ -1,5 +1,5 @@ //***************************************************************************** -// Copyright (c) 2024, Intel Corporation +// Copyright (c) 2025, Intel Corporation // All rights reserved. // // Redistribution and use in source and binary forms, with or without @@ -28,7 +28,6 @@ //***************************************************************************** #include -#include #include "choose.hpp" diff --git a/dpnp/backend/extensions/statistics/CMakeLists.txt b/dpnp/backend/extensions/statistics/CMakeLists.txt index e2529cd04859..b11714849ffd 100644 --- a/dpnp/backend/extensions/statistics/CMakeLists.txt +++ b/dpnp/backend/extensions/statistics/CMakeLists.txt @@ -1,5 +1,5 @@ # ***************************************************************************** -# Copyright (c) 2016-2025, Intel Corporation +# Copyright (c) 2024-2025, Intel Corporation # All rights reserved. # # Redistribution and use in source and binary forms, with or without diff --git a/dpnp/backend/extensions/statistics/statistics_py.cpp b/dpnp/backend/extensions/statistics/statistics_py.cpp index 3694cff4e0f3..1ec760a7c9f0 100644 --- a/dpnp/backend/extensions/statistics/statistics_py.cpp +++ b/dpnp/backend/extensions/statistics/statistics_py.cpp @@ -28,7 +28,6 @@ //***************************************************************************** #include -#include #include "bincount.hpp" #include "histogram.hpp" diff --git a/dpnp/backend/extensions/ufunc/ufunc_py.cpp b/dpnp/backend/extensions/ufunc/ufunc_py.cpp index 771605468c1e..d7824c30b286 100644 --- a/dpnp/backend/extensions/ufunc/ufunc_py.cpp +++ b/dpnp/backend/extensions/ufunc/ufunc_py.cpp @@ -27,7 +27,6 @@ #include "elementwise_functions/common.hpp" -namespace py = pybind11; namespace ufunc_ns = dpnp::extensions::ufunc; PYBIND11_MODULE(_ufunc_impl, m) diff --git a/dpnp/backend/extensions/window/CMakeLists.txt b/dpnp/backend/extensions/window/CMakeLists.txt new file mode 100644 index 000000000000..d4c392ef7d19 --- /dev/null +++ b/dpnp/backend/extensions/window/CMakeLists.txt @@ -0,0 +1,90 @@ +# ***************************************************************************** +# 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. +# ***************************************************************************** + + +set(python_module_name _window_impl) +set(_module_src + ${CMAKE_CURRENT_SOURCE_DIR}/hamming.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/window_py.cpp +) + +pybind11_add_module(${python_module_name} MODULE ${_module_src}) +add_sycl_to_target(TARGET ${python_module_name} SOURCES ${_module_src}) + +if(_dpnp_sycl_targets) + # make fat binary + target_compile_options( + ${python_module_name} + PRIVATE + -fsycl-targets=${_dpnp_sycl_targets} + ) + target_link_options( + ${python_module_name} + PRIVATE + -fsycl-targets=${_dpnp_sycl_targets} + ) +endif() + +if (WIN32) + if (${CMAKE_VERSION} VERSION_LESS "3.27") + # this is a work-around for target_link_options inserting option after -link option, cause + # linker to ignore it. + set(CMAKE_CXX_LINK_FLAGS "${CMAKE_CXX_LINK_FLAGS} -fsycl-device-code-split=per_kernel") + endif() +endif() + +set_target_properties(${python_module_name} PROPERTIES CMAKE_POSITION_INDEPENDENT_CODE ON) + +target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../include) +target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../src) + +target_include_directories(${python_module_name} PUBLIC ${Dpctl_INCLUDE_DIR}) +target_include_directories(${python_module_name} PUBLIC ${Dpctl_TENSOR_INCLUDE_DIR}) + +if (WIN32) + target_compile_options(${python_module_name} PRIVATE + /clang:-fno-approx-func + /clang:-fno-finite-math-only + ) +else() + target_compile_options(${python_module_name} PRIVATE + -fno-approx-func + -fno-finite-math-only + ) +endif() + +target_link_options(${python_module_name} PUBLIC -fsycl-device-code-split=per_kernel) + +if (DPNP_GENERATE_COVERAGE) + target_link_options(${python_module_name} PRIVATE -fprofile-instr-generate -fcoverage-mapping) +endif() + +if (DPNP_WITH_REDIST) + set_target_properties(${python_module_name} PROPERTIES INSTALL_RPATH "$ORIGIN/../../../../../../") +endif() + +install(TARGETS ${python_module_name} + DESTINATION "dpnp/backend/extensions/window" +) diff --git a/dpnp/backend/extensions/window/hamming.cpp b/dpnp/backend/extensions/window/hamming.cpp new file mode 100644 index 000000000000..b7c5cc178304 --- /dev/null +++ b/dpnp/backend/extensions/window/hamming.cpp @@ -0,0 +1,116 @@ +//***************************************************************************** +// 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 +#include +#include + +#include "dpctl4pybind11.hpp" +#include "hamming_kernel.hpp" +#include "utils/output_validation.hpp" +#include "utils/type_dispatch.hpp" + +namespace dpnp::extensions::window +{ + +namespace dpctl_td_ns = dpctl::tensor::type_dispatch; + +static kernels::hamming_fn_ptr_t hamming_dispatch_table[dpctl_td_ns::num_types]; + +namespace py = pybind11; + +std::pair + py_hamming(sycl::queue &exec_q, + 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"); + } + + size_t nelems = result.get_size(); + dpctl::tensor::validation::AmpleMemory::throw_if_not_ample(result, nelems); + 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 = hamming_dispatch_table[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 hamming_ev = fn(exec_q, result_typeless_ptr, nelems, depends); + sycl::event args_ev = + dpctl::utils::keep_args_alive(exec_q, {result}, {hamming_ev}); + + return std::make_pair(args_ev, hamming_ev); +} + +template +struct HammingFactory +{ + fnT get() + { + if constexpr (std::is_same::value || + std::is_same::value) { + return kernels::hamming_impl; + } + else { + return nullptr; + } + } +}; + +void init_hamming_dispatch_tables(void) +{ + using kernels::hamming_fn_ptr_t; + + dpctl_td_ns::DispatchVectorBuilder + contig; + contig.populate_dispatch_vector(hamming_dispatch_table); + + return; +} + +void init_hamming(py::module_ m) +{ + dpnp::extensions::window::init_hamming_dispatch_tables(); + + m.def("_hamming", &py_hamming, "Call hamming kernel", py::arg("sycl_queue"), + py::arg("result"), py::arg("depends") = py::list()); + + return; +} + +} // namespace dpnp::extensions::window diff --git a/dpnp/backend/extensions/window/hamming.hpp b/dpnp/backend/extensions/window/hamming.hpp new file mode 100644 index 000000000000..70fbc58506e4 --- /dev/null +++ b/dpnp/backend/extensions/window/hamming.hpp @@ -0,0 +1,35 @@ +//***************************************************************************** +// 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 + +namespace py = pybind11; + +namespace dpnp::extensions::window +{ +void init_hamming(py::module_ m); +} // namespace dpnp::extensions::window diff --git a/dpnp/backend/extensions/window/hamming_kernel.hpp b/dpnp/backend/extensions/window/hamming_kernel.hpp new file mode 100644 index 000000000000..6c1f4151c3e7 --- /dev/null +++ b/dpnp/backend/extensions/window/hamming_kernel.hpp @@ -0,0 +1,82 @@ +//***************************************************************************** +// 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 "kernels/dpctl_tensor_types.hpp" + +#include "utils/type_utils.hpp" + +namespace dpnp::extensions::window::kernels +{ + +template +class HammingFunctor +{ +private: + T *data; + size_t N; + +public: + HammingFunctor(T *data, size_t N) : data(data), N(N) {} + + void operator()(sycl::id<1> id) const + { + dpctl::tensor::ssize_t i = id[0]; + + data[i] = static_cast(0.54) - + static_cast(0.46) * + sycl::cospi((static_cast(2.0) * i) / (N - 1)); + } +}; + +typedef sycl::event (*hamming_fn_ptr_t)(sycl::queue &, + char *, + size_t, + const std::vector &); + +template +sycl::event hamming_impl(sycl::queue &q, + char *result, + size_t nelems, + const std::vector &depends) +{ + dpctl::tensor::type_utils::validate_type_for_device(q); + + T *res = reinterpret_cast(result); + + sycl::event hamming_ev = q.submit([&](sycl::handler &cgh) { + cgh.depends_on(depends); + + cgh.parallel_for(sycl::range<1>(nelems), + HammingFunctor(res, nelems)); + }); + + return hamming_ev; +} + +} // namespace dpnp::extensions::window::kernels diff --git a/dpnp/backend/extensions/window/window_py.cpp b/dpnp/backend/extensions/window/window_py.cpp new file mode 100644 index 000000000000..e27f840caac5 --- /dev/null +++ b/dpnp/backend/extensions/window/window_py.cpp @@ -0,0 +1,37 @@ +//***************************************************************************** +// 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. +//***************************************************************************** +// +// This file defines functions of dpnp.backend._window_impl extensions +// +//***************************************************************************** + +#include + +#include "hamming.hpp" + +PYBIND11_MODULE(_window_impl, m) +{ + dpnp::extensions::window::init_hamming(m); +} diff --git a/dpnp/dpnp_iface.py b/dpnp/dpnp_iface.py index c81ba27fde4e..f6382aea4133 100644 --- a/dpnp/dpnp_iface.py +++ b/dpnp/dpnp_iface.py @@ -106,6 +106,8 @@ from dpnp.dpnp_iface_statistics import __all__ as __all__statistics from dpnp.dpnp_iface_trigonometric import * from dpnp.dpnp_iface_trigonometric import __all__ as __all__trigonometric +from dpnp.dpnp_iface_window import * +from dpnp.dpnp_iface_window import __all__ as __all__window # pylint: disable=no-name-in-module from .dpnp_utils import ( @@ -129,6 +131,7 @@ __all__ += __all__searching __all__ += __all__sorting __all__ += __all__statistics +__all__ += __all__window __all__ += __all__trigonometric diff --git a/dpnp/dpnp_iface_window.py b/dpnp/dpnp_iface_window.py new file mode 100644 index 000000000000..edb3f4f611b5 --- /dev/null +++ b/dpnp/dpnp_iface_window.py @@ -0,0 +1,176 @@ +# ***************************************************************************** +# 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. +# ***************************************************************************** + +""" +Interface of the window functions of dpnp + +Notes +----- +This module is a face or public interface file for the library +it contains: + - Interface functions + - documentation for the functions + - The functions parameters check + +""" + +# pylint: disable=no-name-in-module +# pylint: disable=invalid-name +# pylint: disable=protected-access + +import dpctl.utils as dpu +import numpy + +import dpnp +import dpnp.backend.extensions.window._window_impl as wi + +__all__ = ["hamming"] + + +def _validate_input(val): + + is_numpy_array = isinstance(val, numpy.ndarray) + is_array = dpnp.is_supported_array_type(val) or is_numpy_array + if is_array: + is_0d_arr = val.ndim == 0 + is_int_float = dpnp.issubdtype(val.dtype, (dpnp.integer, dpnp.floating)) + raise_error = not (is_0d_arr and is_int_float) + if not raise_error: + is_nan = numpy.isnan(val) if is_numpy_array else dpnp.isnan(val) + is_inf = numpy.isinf(val) if is_numpy_array else dpnp.isinf(val) + raise_error = is_nan or is_inf + else: + is_int = isinstance(val, (int, numpy.integer, dpnp.integer)) + is_float = isinstance(val, (float, numpy.floating, dpnp.floating)) + raise_error = not (is_int or is_float) + if not raise_error: + raise_error = val in [numpy.inf, -numpy.inf, numpy.nan] + + if raise_error: + raise TypeError("M must be an integer") + + +def hamming(M, device=None, usm_type=None, sycl_queue=None): + r""" + Return the Hamming window. + + The Hamming window is a taper formed by using a weighted cosine. + + For full documentation refer to :obj:`numpy.hamming`. + + 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 + 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.hanning` : Return the Hanning window. + :obj:`dpnp.kaiser` : Return the Kaiser window. + + Notes + ----- + The Hamming window is defined as + + .. math:: w(n) = 0.54 - 0.46\cos\left(\frac{2\pi{n}}{M-1}\right) + \qquad 0 \leq n \leq M-1 + + Examples + -------- + >>> import dpnp as np + >>> np.hamming(12) + array([0.08 , 0.15302337, 0.34890909, 0.60546483, 0.84123594, + 0.98136677, 0.98136677, 0.84123594, 0.60546483, 0.34890909, + 0.15302337, 0.08 ]) # may vary + + Creating the output array on a different device or with a + specified usm_type: + + >>> x = np.hamming(4) # default case + >>> x, x.device, x.usm_type + (array([0.08, 0.77, 0.77, 0.08]), Device(level_zero:gpu:0), 'device') + + >>> y = np.hamming(4, device="cpu") + >>> y, y.device, y.usm_type + (array([0.08, 0.77, 0.77, 0.08]), Device(opencl:cpu:0), 'device') + + >>> z = np.hamming(4, usm_type="host") + >>> z, z.device, z.usm_type + (array([0.08, 0.77, 0.77, 0.08]), Device(level_zero:gpu:0), 'host') + + """ + + _validate_input(M) + + cfd_kwarg = { + "device": device, + "usm_type": usm_type, + "sycl_queue": sycl_queue, + } + if M < 1: + return dpnp.empty(0, **cfd_kwarg) + if M == 1: + return dpnp.ones(1, **cfd_kwarg) + + result = dpnp.empty(int(M), **cfd_kwarg) + exec_q = result.sycl_queue + _manager = dpu.SequentialOrderManager[exec_q] + + ht_ev, win_ev = wi._hamming( + exec_q, dpnp.get_usm_ndarray(result), depends=_manager.submitted_events + ) + + _manager.add_event_pair(ht_ev, win_ev) + + return result diff --git a/dpnp/tests/test_sycl_queue.py b/dpnp/tests/test_sycl_queue.py index c2eb62a52ff7..51eb4e1e574a 100644 --- a/dpnp/tests/test_sycl_queue.py +++ b/dpnp/tests/test_sycl_queue.py @@ -2120,3 +2120,18 @@ def test_choose(device): inds = dpnp.array([0, 1, 3], dtype="i4", device=device) result = dpnp.choose(inds, chc) assert_sycl_queue_equal(result.sycl_queue, chc.sycl_queue) + + +@pytest.mark.parametrize("func", ["hamming"]) +@pytest.mark.parametrize( + "device", + valid_dev + [None], + ids=[device.filter_string for device in valid_dev] + [None], +) +def test_window(func, device): + result = getattr(dpnp, func)(10, device=device) + if device is None: + # assert against default device + device = dpctl.select_default_device() + + assert result.sycl_device == device diff --git a/dpnp/tests/test_usm_type.py b/dpnp/tests/test_usm_type.py index 213f618f187d..ef87912b0941 100644 --- a/dpnp/tests/test_usm_type.py +++ b/dpnp/tests/test_usm_type.py @@ -1845,3 +1845,14 @@ def test_choose(usm_type_x, usm_type_ind): assert chc.usm_type == usm_type_x assert ind.usm_type == usm_type_ind assert z.usm_type == du.get_coerced_usm_type([usm_type_x, usm_type_ind]) + + +@pytest.mark.parametrize("func", ["hamming"]) +@pytest.mark.parametrize("usm_type", list_of_usm_types + [None]) +def test_window(func, usm_type): + result = getattr(dp, func)(10, usm_type=usm_type) + if usm_type is None: + # assert against default USM type + assert result.usm_type == "device" + else: + assert result.usm_type == usm_type diff --git a/dpnp/tests/test_window.py b/dpnp/tests/test_window.py new file mode 100644 index 000000000000..da7deca038e5 --- /dev/null +++ b/dpnp/tests/test_window.py @@ -0,0 +1,50 @@ +import numpy +import pytest +from numpy.testing import assert_raises + +import dpnp + +from .helper import assert_dtype_allclose + + +@pytest.mark.filterwarnings("ignore::DeprecationWarning") +@pytest.mark.parametrize("func", ["hamming"]) +@pytest.mark.parametrize( + "M", + [ + True, + False, + 0, + dpnp.int32(1), + 4, + 5.0, + dpnp.float32(6), + dpnp.array(7), + numpy.array(8), + ], +) +def test_window(func, M): + result = getattr(dpnp, func)(M) + + if isinstance(M, dpnp.ndarray): + M = M.asnumpy() + expected = getattr(numpy, func)(M) + + assert_dtype_allclose(result, expected) + + +@pytest.mark.filterwarnings("ignore::DeprecationWarning") +@pytest.mark.parametrize("func", ["hamming"]) +@pytest.mark.parametrize( + "M", + [ + 5 + 4j, + numpy.array(5 + 4j), + dpnp.array([5]), + numpy.inf, + numpy.array(-numpy.inf), + dpnp.array(dpnp.nan), + ], +) +def test_window_error(func, M): + assert_raises(TypeError, getattr(dpnp, func), 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 1d7e3642cf93..2cd05fec2d42 100644 --- a/dpnp/tests/third_party/cupy/math_tests/test_window.py +++ b/dpnp/tests/third_party/cupy/math_tests/test_window.py @@ -2,22 +2,22 @@ import pytest +from dpnp.tests.helper import has_support_aspect64 from dpnp.tests.third_party.cupy import testing -pytest.skip("window functions are not supported yet", allow_module_level=True) - @testing.parameterize( *testing.product( { "m": [0, 1, -1, 1024], - "name": ["bartlett", "blackman", "hamming", "hanning"], + # TODO: add ["bartlett", "blackman", "hanning"] when supported + "name": ["hamming"], } ) ) class TestWindow(unittest.TestCase): - @testing.numpy_cupy_allclose(atol=1e-5) + @testing.numpy_cupy_allclose(atol=1e-5, type_check=has_support_aspect64()) def test_window(self, xp): return getattr(xp, self.name)(self.m) @@ -33,6 +33,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) def test_kaiser_parametric(self, xp): return getattr(xp, self.name)(self.m, self.beta) @@ -41,6 +42,7 @@ 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) def test_kaiser(self, xp): return xp.kaiser(self.m, 1.5) From 6d4c81247e20ba8b314304e4b429e9026aace9e8 Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Tue, 25 Feb 2025 18:33:04 -0800 Subject: [PATCH 2/7] update-test --- dpnp/dpnp_iface_window.py | 5 +++-- dpnp/tests/test_window.py | 4 +--- 2 files changed, 4 insertions(+), 5 deletions(-) diff --git a/dpnp/dpnp_iface_window.py b/dpnp/dpnp_iface_window.py index edb3f4f611b5..a591a921a56c 100644 --- a/dpnp/dpnp_iface_window.py +++ b/dpnp/dpnp_iface_window.py @@ -55,8 +55,9 @@ def _validate_input(val): is_array = dpnp.is_supported_array_type(val) or is_numpy_array if is_array: is_0d_arr = val.ndim == 0 - is_int_float = dpnp.issubdtype(val.dtype, (dpnp.integer, dpnp.floating)) - raise_error = not (is_0d_arr and is_int_float) + is_int = dpnp.issubdtype(val.dtype, dpnp.integer) + is_float = dpnp.issubdtype(val.dtype, dpnp.floating) + raise_error = not (is_0d_arr and is_int or is_float) if not raise_error: is_nan = numpy.isnan(val) if is_numpy_array else dpnp.isnan(val) is_inf = numpy.isinf(val) if is_numpy_array else dpnp.isinf(val) diff --git a/dpnp/tests/test_window.py b/dpnp/tests/test_window.py index da7deca038e5..9e6c9e509f27 100644 --- a/dpnp/tests/test_window.py +++ b/dpnp/tests/test_window.py @@ -7,7 +7,6 @@ from .helper import assert_dtype_allclose -@pytest.mark.filterwarnings("ignore::DeprecationWarning") @pytest.mark.parametrize("func", ["hamming"]) @pytest.mark.parametrize( "M", @@ -20,7 +19,7 @@ 5.0, dpnp.float32(6), dpnp.array(7), - numpy.array(8), + numpy.array(8.0), ], ) def test_window(func, M): @@ -33,7 +32,6 @@ def test_window(func, M): assert_dtype_allclose(result, expected) -@pytest.mark.filterwarnings("ignore::DeprecationWarning") @pytest.mark.parametrize("func", ["hamming"]) @pytest.mark.parametrize( "M", From d19569f7acd00b854545afc6d273d54f6f22c887 Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Wed, 26 Feb 2025 06:49:29 -0800 Subject: [PATCH 3/7] update changelog --- CHANGELOG.md | 2 ++ .../extensions/indexing/indexing_py.cpp | 2 +- .../extensions/statistics/statistics_py.cpp | 2 +- dpnp/backend/extensions/window/hamming.hpp | 2 +- dpnp/dpnp_iface_indexing.py | 5 --- dpnp/dpnp_iface_manipulation.py | 1 - dpnp/dpnp_iface_window.py | 4 +-- dpnp/fft/dpnp_iface_fft.py | 36 +++++++------------ 8 files changed, 20 insertions(+), 34 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index ab22e25e74f1..ca23d4d7ddf7 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -8,6 +8,8 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 ### Added +* Added implementation of `dpnp.hamming` [#2341](https://github.com/IntelPython/dpnp/pull/2341) + ### Changed ### Fixed diff --git a/dpnp/backend/extensions/indexing/indexing_py.cpp b/dpnp/backend/extensions/indexing/indexing_py.cpp index e4c0bcee10ee..ce19dd684fb3 100644 --- a/dpnp/backend/extensions/indexing/indexing_py.cpp +++ b/dpnp/backend/extensions/indexing/indexing_py.cpp @@ -23,7 +23,7 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** // -// This file defines functions of dpnp.backend._lapack_impl extensions +// This file defines functions of dpnp.backend._indexing_impl extensions // //***************************************************************************** diff --git a/dpnp/backend/extensions/statistics/statistics_py.cpp b/dpnp/backend/extensions/statistics/statistics_py.cpp index 1ec760a7c9f0..6636d3f7d531 100644 --- a/dpnp/backend/extensions/statistics/statistics_py.cpp +++ b/dpnp/backend/extensions/statistics/statistics_py.cpp @@ -23,7 +23,7 @@ // THE POSSIBILITY OF SUCH DAMAGE. //***************************************************************************** // -// This file defines functions of dpnp.backend._lapack_impl extensions +// This file defines functions of dpnp.backend._statistics_impl extensions // //***************************************************************************** diff --git a/dpnp/backend/extensions/window/hamming.hpp b/dpnp/backend/extensions/window/hamming.hpp index 70fbc58506e4..8a900d1c3381 100644 --- a/dpnp/backend/extensions/window/hamming.hpp +++ b/dpnp/backend/extensions/window/hamming.hpp @@ -32,4 +32,4 @@ namespace py = pybind11; namespace dpnp::extensions::window { void init_hamming(py::module_ m); -} // namespace dpnp::extensions::window +} diff --git a/dpnp/dpnp_iface_indexing.py b/dpnp/dpnp_iface_indexing.py index 630f2d7a52e5..998579430626 100644 --- a/dpnp/dpnp_iface_indexing.py +++ b/dpnp/dpnp_iface_indexing.py @@ -476,7 +476,6 @@ def diag_indices(n, ndim=2, device=None, usm_type="device", sycl_queue=None): :attr:`dpnp.ndarray.device`. Default: ``None``. - usm_type : {"device", "shared", "host"}, optional The type of SYCL USM allocation for the output array. @@ -1084,7 +1083,6 @@ def indices( :attr:`dpnp.ndarray.device`. Default: ``None``. - usm_type : {"device", "shared", "host"}, optional The type of SYCL USM allocation for the output array. @@ -1349,7 +1347,6 @@ def mask_indices( :attr:`dpnp.ndarray.device`. Default: ``None``. - usm_type : {"device", "shared", "host"}, optional The type of SYCL USM allocation for the output array. @@ -2360,7 +2357,6 @@ def tril_indices( :attr:`dpnp.ndarray.device`. Default: ``None``. - usm_type : {"device", "shared", "host"}, optional The type of SYCL USM allocation for the output array. @@ -2570,7 +2566,6 @@ def triu_indices( :attr:`dpnp.ndarray.device`. Default: ``None``. - usm_type : {"device", "shared", "host"}, optional The type of SYCL USM allocation for the output array. diff --git a/dpnp/dpnp_iface_manipulation.py b/dpnp/dpnp_iface_manipulation.py index 1755adac6978..4a164c0e45d6 100644 --- a/dpnp/dpnp_iface_manipulation.py +++ b/dpnp/dpnp_iface_manipulation.py @@ -752,7 +752,6 @@ def asarray_chkfinite( :attr:`dpnp.ndarray.device`. Default: ``None``. - usm_type : {None, "device", "shared", "host"}, optional The type of SYCL USM allocation for the output array. diff --git a/dpnp/dpnp_iface_window.py b/dpnp/dpnp_iface_window.py index a591a921a56c..d34828840d09 100644 --- a/dpnp/dpnp_iface_window.py +++ b/dpnp/dpnp_iface_window.py @@ -109,7 +109,7 @@ def hamming(M, device=None, usm_type=None, sycl_queue=None): Returns ------- - out : dpnp.ndarray + 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). @@ -153,12 +153,12 @@ def hamming(M, device=None, usm_type=None, sycl_queue=None): """ _validate_input(M) - cfd_kwarg = { "device": device, "usm_type": usm_type, "sycl_queue": sycl_queue, } + if M < 1: return dpnp.empty(0, **cfd_kwarg) if M == 1: diff --git a/dpnp/fft/dpnp_iface_fft.py b/dpnp/fft/dpnp_iface_fft.py index bc9aa3ac9df3..c82daad70d0a 100644 --- a/dpnp/fft/dpnp_iface_fft.py +++ b/dpnp/fft/dpnp_iface_fft.py @@ -341,32 +341,22 @@ def fftfreq(n, d=1.0, device=None, usm_type=None, sycl_queue=None): raise ValueError("`n` should be an integer") if not dpnp.isscalar(d): raise ValueError("`d` should be an scalar") + + cfd_kwarg = { + "device": device, + "usm_type": usm_type, + "sycl_queue": sycl_queue, + } + val = 1.0 / (n * d) - results = dpnp.empty( - n, - dtype=dpnp.intp, - device=device, - usm_type=usm_type, - sycl_queue=sycl_queue, - ) + results = dpnp.empty(n, dtype=dpnp.intp, **cfd_kwarg) + m = (n - 1) // 2 + 1 - p1 = dpnp.arange( - 0, - m, - dtype=dpnp.intp, - device=device, - usm_type=usm_type, - sycl_queue=sycl_queue, - ) + p1 = dpnp.arange(0, m, dtype=dpnp.intp, **cfd_kwarg) + results[:m] = p1 - p2 = dpnp.arange( - m - n, - 0, - dtype=dpnp.intp, - device=device, - usm_type=usm_type, - sycl_queue=sycl_queue, - ) + p2 = dpnp.arange(m - n, 0, dtype=dpnp.intp, **cfd_kwarg) + results[m:] = p2 return results * val From 0fca5e49310e341a741621faabe34e6daee7bc89 Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Thu, 27 Feb 2025 09:46:54 -0800 Subject: [PATCH 4/7] address comments --- dpnp/backend/extensions/window/hamming.cpp | 14 +++++-- .../extensions/window/hamming_kernel.hpp | 23 +++++------ dpnp/dpnp_iface.py | 2 +- dpnp/dpnp_iface_window.py | 33 +++------------ dpnp/tests/test_sycl_queue.py | 41 +++++++++++++++---- dpnp/tests/test_window.py | 3 +- 6 files changed, 62 insertions(+), 54 deletions(-) diff --git a/dpnp/backend/extensions/window/hamming.cpp b/dpnp/backend/extensions/window/hamming.cpp index b7c5cc178304..1e6fc5edf9cd 100644 --- a/dpnp/backend/extensions/window/hamming.cpp +++ b/dpnp/backend/extensions/window/hamming.cpp @@ -53,8 +53,17 @@ std::pair 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(); - dpctl::tensor::validation::AmpleMemory::throw_if_not_ample(result, nelems); if (nelems == 0) { return std::make_pair(sycl::event{}, sycl::event{}); } @@ -81,8 +90,7 @@ struct HammingFactory { fnT get() { - if constexpr (std::is_same::value || - std::is_same::value) { + if constexpr (std::is_floating_point_v) { return kernels::hamming_impl; } else { diff --git a/dpnp/backend/extensions/window/hamming_kernel.hpp b/dpnp/backend/extensions/window/hamming_kernel.hpp index 6c1f4151c3e7..ebe92e693600 100644 --- a/dpnp/backend/extensions/window/hamming_kernel.hpp +++ b/dpnp/backend/extensions/window/hamming_kernel.hpp @@ -27,8 +27,6 @@ #include -#include "kernels/dpctl_tensor_types.hpp" - #include "utils/type_utils.hpp" namespace dpnp::extensions::window::kernels @@ -38,31 +36,29 @@ template class HammingFunctor { private: - T *data; - size_t N; + T *data = nullptr; + const std::size_t N; public: - HammingFunctor(T *data, size_t N) : data(data), N(N) {} + HammingFunctor(T *data, const std::size_t N) : data(data), N(N) {} void operator()(sycl::id<1> id) const { - dpctl::tensor::ssize_t i = id[0]; + const auto i = id.get(0); - data[i] = static_cast(0.54) - - static_cast(0.46) * - sycl::cospi((static_cast(2.0) * i) / (N - 1)); + data[i] = T(0.54) - T(0.46) * sycl::cospi(T(2) * i / (N - 1)); } }; typedef sycl::event (*hamming_fn_ptr_t)(sycl::queue &, char *, - size_t, + const std::size_t, const std::vector &); template sycl::event hamming_impl(sycl::queue &q, char *result, - size_t nelems, + const std::size_t nelems, const std::vector &depends) { dpctl::tensor::type_utils::validate_type_for_device(q); @@ -72,8 +68,9 @@ sycl::event hamming_impl(sycl::queue &q, sycl::event hamming_ev = q.submit([&](sycl::handler &cgh) { cgh.depends_on(depends); - cgh.parallel_for(sycl::range<1>(nelems), - HammingFunctor(res, nelems)); + using HammingKernel = HammingFunctor; + cgh.parallel_for(sycl::range<1>(nelems), + HammingKernel(res, nelems)); }); return hamming_ev; diff --git a/dpnp/dpnp_iface.py b/dpnp/dpnp_iface.py index f6382aea4133..95cce5e242ca 100644 --- a/dpnp/dpnp_iface.py +++ b/dpnp/dpnp_iface.py @@ -131,8 +131,8 @@ __all__ += __all__searching __all__ += __all__sorting __all__ += __all__statistics -__all__ += __all__window __all__ += __all__trigonometric +__all__ += __all__window def are_same_logical_tensors(ar1, ar2): diff --git a/dpnp/dpnp_iface_window.py b/dpnp/dpnp_iface_window.py index d34828840d09..01a3b1f0c24c 100644 --- a/dpnp/dpnp_iface_window.py +++ b/dpnp/dpnp_iface_window.py @@ -41,7 +41,6 @@ # pylint: disable=protected-access import dpctl.utils as dpu -import numpy import dpnp import dpnp.backend.extensions.window._window_impl as wi @@ -49,30 +48,6 @@ __all__ = ["hamming"] -def _validate_input(val): - - is_numpy_array = isinstance(val, numpy.ndarray) - is_array = dpnp.is_supported_array_type(val) or is_numpy_array - if is_array: - is_0d_arr = val.ndim == 0 - is_int = dpnp.issubdtype(val.dtype, dpnp.integer) - is_float = dpnp.issubdtype(val.dtype, dpnp.floating) - raise_error = not (is_0d_arr and is_int or is_float) - if not raise_error: - is_nan = numpy.isnan(val) if is_numpy_array else dpnp.isnan(val) - is_inf = numpy.isinf(val) if is_numpy_array else dpnp.isinf(val) - raise_error = is_nan or is_inf - else: - is_int = isinstance(val, (int, numpy.integer, dpnp.integer)) - is_float = isinstance(val, (float, numpy.floating, dpnp.floating)) - raise_error = not (is_int or is_float) - if not raise_error: - raise_error = val in [numpy.inf, -numpy.inf, numpy.nan] - - if raise_error: - raise TypeError("M must be an integer") - - def hamming(M, device=None, usm_type=None, sycl_queue=None): r""" Return the Hamming window. @@ -152,7 +127,11 @@ def hamming(M, device=None, usm_type=None, sycl_queue=None): """ - _validate_input(M) + try: + M = int(M) + except Exception as e: + raise TypeError("M must be an integer") from e + cfd_kwarg = { "device": device, "usm_type": usm_type, @@ -164,7 +143,7 @@ def hamming(M, device=None, usm_type=None, sycl_queue=None): if M == 1: return dpnp.ones(1, **cfd_kwarg) - result = dpnp.empty(int(M), **cfd_kwarg) + result = dpnp.empty(M, **cfd_kwarg) exec_q = result.sycl_queue _manager = dpu.SequentialOrderManager[exec_q] diff --git a/dpnp/tests/test_sycl_queue.py b/dpnp/tests/test_sycl_queue.py index 51eb4e1e574a..7d4b9215aa66 100644 --- a/dpnp/tests/test_sycl_queue.py +++ b/dpnp/tests/test_sycl_queue.py @@ -51,6 +51,8 @@ def assert_sycl_queue_equal(result, expected): "func, arg, kwargs", [ pytest.param("arange", [-25.7], {"stop": 10**8, "step": 15}), + pytest.param("eye", [4, 2], {}), + pytest.param("empty", [(2, 2)], {}), pytest.param( "frombuffer", [b"\x01\x02\x03\x04"], {"dtype": dpnp.int32} ), @@ -62,9 +64,8 @@ def assert_sycl_queue_equal(result, expected): pytest.param("fromiter", [[1, 2, 3, 4]], {"dtype": dpnp.int64}), pytest.param("fromstring", ["1 2"], {"dtype": int, "sep": " "}), pytest.param("full", [(2, 2)], {"fill_value": 5}), - pytest.param("eye", [4, 2], {}), - pytest.param("empty", [(2, 2)], {}), pytest.param("geomspace", [1, 4, 8], {}), + pytest.param("hamming", [10], {}), pytest.param("identity", [4], {}), pytest.param("linspace", [0, 4, 8], {}), pytest.param("logspace", [0, 4, 8], {}), @@ -75,12 +76,18 @@ def assert_sycl_queue_equal(result, expected): ], ) @pytest.mark.parametrize( - "device", valid_dev, ids=[dev.filter_string for dev in valid_dev] + "device", + valid_dev + [None], + ids=[dev.filter_string for dev in valid_dev] + [None], ) def test_array_creation(func, arg, kwargs, device): - dpnp_kwargs = dict(kwargs) - dpnp_kwargs["device"] = device - x = getattr(dpnp, func)(*arg, **dpnp_kwargs) + kwargs = dict(kwargs) + kwargs["device"] = device + x = getattr(dpnp, func)(*arg, **kwargs) + + if device is None: + # assert against default device + device = dpctl.select_default_device() assert x.sycl_device == device @@ -211,7 +218,9 @@ def test_array_creation_cross_device_2d_array( @pytest.mark.parametrize( - "device", valid_dev, ids=[dev.filter_string for dev in valid_dev] + "device", + valid_dev + [None], + ids=[dev.filter_string for dev in valid_dev] + [None], ) def test_array_creation_from_file(device): with tempfile.TemporaryFile() as fh: @@ -221,11 +230,16 @@ def test_array_creation_from_file(device): fh.seek(0) x = dpnp.fromfile(fh, device=device) + if device is None: + # assert against default device + device = dpctl.select_default_device() assert x.sycl_device == device @pytest.mark.parametrize( - "device", valid_dev, ids=[dev.filter_string for dev in valid_dev] + "device", + valid_dev + [None], + ids=[dev.filter_string for dev in valid_dev] + [None], ) def test_array_creation_load_txt(device): with tempfile.TemporaryFile() as fh: @@ -235,6 +249,9 @@ def test_array_creation_load_txt(device): fh.seek(0) x = dpnp.loadtxt(fh, device=device) + if device is None: + # assert against default device + device = dpctl.select_default_device() assert x.sycl_device == device @@ -1056,10 +1073,16 @@ def test_rfftn(self, device): @pytest.mark.parametrize("func", ["fftfreq", "rfftfreq"]) @pytest.mark.parametrize( - "device", valid_dev, ids=[dev.filter_string for dev in valid_dev] + "device", + valid_dev + [None], + ids=[dev.filter_string for dev in valid_dev] + [None], ) def test_fftfreq(self, func, device): result = getattr(dpnp.fft, func)(10, 0.5, device=device) + + if device is None: + # assert against default device + device = dpctl.select_default_device() assert result.sycl_device == device @pytest.mark.parametrize("func", ["fftshift", "ifftshift"]) diff --git a/dpnp/tests/test_window.py b/dpnp/tests/test_window.py index 9e6c9e509f27..dde27c45e4a2 100644 --- a/dpnp/tests/test_window.py +++ b/dpnp/tests/test_window.py @@ -38,7 +38,7 @@ def test_window(func, M): [ 5 + 4j, numpy.array(5 + 4j), - dpnp.array([5]), + dpnp.array([5, 3]), numpy.inf, numpy.array(-numpy.inf), dpnp.array(dpnp.nan), @@ -46,3 +46,4 @@ def test_window(func, M): ) def test_window_error(func, M): assert_raises(TypeError, getattr(dpnp, func), M) + # assert_raises(TypeError, getattr(numpy, func), M) From c4da47b5a29b02ab41cc7571ddc592cc00e79bd3 Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Thu, 27 Feb 2025 09:50:42 -0800 Subject: [PATCH 5/7] remove leftover --- dpnp/tests/test_window.py | 1 - 1 file changed, 1 deletion(-) diff --git a/dpnp/tests/test_window.py b/dpnp/tests/test_window.py index dde27c45e4a2..edb1063f6a79 100644 --- a/dpnp/tests/test_window.py +++ b/dpnp/tests/test_window.py @@ -46,4 +46,3 @@ def test_window(func, M): ) def test_window_error(func, M): assert_raises(TypeError, getattr(dpnp, func), M) - # assert_raises(TypeError, getattr(numpy, func), M) From 3c2fb7221d611f276da3783835500c560618e1f0 Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Thu, 27 Feb 2025 10:24:04 -0800 Subject: [PATCH 6/7] remove redundant test in test_sycl_queue.py --- dpnp/tests/test_sycl_queue.py | 15 --------------- 1 file changed, 15 deletions(-) diff --git a/dpnp/tests/test_sycl_queue.py b/dpnp/tests/test_sycl_queue.py index 7d4b9215aa66..3541a4fe584a 100644 --- a/dpnp/tests/test_sycl_queue.py +++ b/dpnp/tests/test_sycl_queue.py @@ -2143,18 +2143,3 @@ def test_choose(device): inds = dpnp.array([0, 1, 3], dtype="i4", device=device) result = dpnp.choose(inds, chc) assert_sycl_queue_equal(result.sycl_queue, chc.sycl_queue) - - -@pytest.mark.parametrize("func", ["hamming"]) -@pytest.mark.parametrize( - "device", - valid_dev + [None], - ids=[device.filter_string for device in valid_dev] + [None], -) -def test_window(func, device): - result = getattr(dpnp, func)(10, device=device) - if device is None: - # assert against default device - device = dpctl.select_default_device() - - assert result.sycl_device == device From b0aeb6fb3bd914e77d86e3a5b6e47eab93e0b0db Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Thu, 27 Feb 2025 10:39:28 -0800 Subject: [PATCH 7/7] add TODO --- dpnp/tests/test_usm_type.py | 1 + 1 file changed, 1 insertion(+) diff --git a/dpnp/tests/test_usm_type.py b/dpnp/tests/test_usm_type.py index ef87912b0941..d51041ee51fa 100644 --- a/dpnp/tests/test_usm_type.py +++ b/dpnp/tests/test_usm_type.py @@ -1847,6 +1847,7 @@ def test_choose(usm_type_x, usm_type_ind): assert z.usm_type == du.get_coerced_usm_type([usm_type_x, usm_type_ind]) +# TODO: add it as part of `test_array_creation_from_scratch` function @pytest.mark.parametrize("func", ["hamming"]) @pytest.mark.parametrize("usm_type", list_of_usm_types + [None]) def test_window(func, usm_type):