implement dpnp.hamming

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

doc/reference/routines.rst

@@ -27,3 +27,4 @@ These functions cover a subset of
    set
    sort
    statistics
+   window

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

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)

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

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

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

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

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 <pybind11/pybind11.h>
-#include <pybind11/stl.h>
 
 #include "choose.hpp"
 

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

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 <pybind11/pybind11.h>
-#include <pybind11/stl.h>
 
 #include "bincount.hpp"
 #include "histogram.hpp"

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)

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"
+)

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 <pybind11/pybind11.h>
+#include <pybind11/stl.h>
+#include <sycl/sycl.hpp>
+
+#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<sycl::event, sycl::event>
+    py_hamming(sycl::queue &exec_q,
+               const dpctl::tensor::usm_ndarray &result,
+               const std::vector<sycl::event> &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 <typename fnT, typename T>
+struct HammingFactory
+{
+    fnT get()
+    {
+        if constexpr (std::is_floating_point_v<T>) {
+            return kernels::hamming_impl<T>;
+        }
+        else {
+            return nullptr;
+        }
+    }
+};
+
+void init_hamming_dispatch_tables(void)
+{
+    using kernels::hamming_fn_ptr_t;
+
+    dpctl_td_ns::DispatchVectorBuilder<hamming_fn_ptr_t, HammingFactory,
+                                       dpctl_td_ns::num_types>
+        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

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 <pybind11/pybind11.h>
+
+namespace py = pybind11;
+
+namespace dpnp::extensions::window
+{
+void init_hamming(py::module_ m);
+}