diff --git a/CHANGELOG.md b/CHANGELOG.md index 41e83d9f4714..32bb55d37ce9 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/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..ce19dd684fb3 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 @@ -23,12 +23,11 @@ // 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 // //***************************************************************************** #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..6636d3f7d531 100644 --- a/dpnp/backend/extensions/statistics/statistics_py.cpp +++ b/dpnp/backend/extensions/statistics/statistics_py.cpp @@ -23,12 +23,11 @@ // 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 // //***************************************************************************** #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..1e6fc5edf9cd --- /dev/null +++ b/dpnp/backend/extensions/window/hamming.cpp @@ -0,0 +1,124 @@ +//***************************************************************************** +// 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"); + } + + 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 = 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_floating_point_v) { + 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..8a900d1c3381 --- /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); +} diff --git a/dpnp/backend/extensions/window/hamming_kernel.hpp b/dpnp/backend/extensions/window/hamming_kernel.hpp new file mode 100644 index 000000000000..ebe92e693600 --- /dev/null +++ b/dpnp/backend/extensions/window/hamming_kernel.hpp @@ -0,0 +1,79 @@ +//***************************************************************************** +// 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 "utils/type_utils.hpp" + +namespace dpnp::extensions::window::kernels +{ + +template +class HammingFunctor +{ +private: + T *data = nullptr; + const std::size_t N; + +public: + HammingFunctor(T *data, const std::size_t N) : data(data), N(N) {} + + void operator()(sycl::id<1> id) const + { + const auto i = id.get(0); + + data[i] = T(0.54) - T(0.46) * sycl::cospi(T(2) * i / (N - 1)); + } +}; + +typedef sycl::event (*hamming_fn_ptr_t)(sycl::queue &, + char *, + const std::size_t, + const std::vector &); + +template +sycl::event hamming_impl(sycl::queue &q, + char *result, + const std::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); + + using HammingKernel = HammingFunctor; + cgh.parallel_for(sycl::range<1>(nelems), + HammingKernel(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..95cce5e242ca 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 ( @@ -130,6 +132,7 @@ __all__ += __all__sorting __all__ += __all__statistics __all__ += __all__trigonometric +__all__ += __all__window def are_same_logical_tensors(ar1, ar2): 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 new file mode 100644 index 000000000000..01a3b1f0c24c --- /dev/null +++ b/dpnp/dpnp_iface_window.py @@ -0,0 +1,156 @@ +# ***************************************************************************** +# 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 dpnp +import dpnp.backend.extensions.window._window_impl as wi + +__all__ = ["hamming"] + + +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 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.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') + + """ + + 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, + "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(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/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 diff --git a/dpnp/tests/test_sycl_queue.py b/dpnp/tests/test_sycl_queue.py index c2eb62a52ff7..3541a4fe584a 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_usm_type.py b/dpnp/tests/test_usm_type.py index 213f618f187d..d51041ee51fa 100644 --- a/dpnp/tests/test_usm_type.py +++ b/dpnp/tests/test_usm_type.py @@ -1845,3 +1845,15 @@ 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]) + + +# 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): + 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..edb1063f6a79 --- /dev/null +++ b/dpnp/tests/test_window.py @@ -0,0 +1,48 @@ +import numpy +import pytest +from numpy.testing import assert_raises + +import dpnp + +from .helper import assert_dtype_allclose + + +@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.0), + ], +) +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.parametrize("func", ["hamming"]) +@pytest.mark.parametrize( + "M", + [ + 5 + 4j, + numpy.array(5 + 4j), + dpnp.array([5, 3]), + 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)