implement dpnp.hanning

IntelPython · vtavana · Mar 4, 2025 · Feb 21, 2025 · Mar 4, 2025 · Mar 4, 2025
commit 9fe7829d97c13a2a2fd6bd7d6e1c45be115fd8de
@@ -2,7 +2,9 @@ al
 ary
 backend
 bandlimited
+bincount
 bitwise
+Blackman
 boolean
 broadcastable
 broadcasted
@@ -36,6 +38,7 @@ fs
 getter
 Golub
 Hadamard
+Hanning
 histogrammed
 Hypergeometric
 kwargs
@@ -49,9 +52,12 @@ Lanczos
 Lomax
 Mersenne
 meshgrid
+minlength
 Mises
 multinomial
 multivalued
+namespace
+namedtuple
 NaN
 NaT
 nd
@@ -69,6 +75,7 @@ Nyquist
 oneAPI
 ord
 orthonormal
+radix
 Penrose
 Polyutils
 pre
@@ -79,6 +86,7 @@ representable
 resampling
 runtimes
 scikit
+se
 signbit
 signum
 sinc

@@ -97,7 +97,7 @@ Other matrix operations

   dpnp.diagonal
   dpnp.linalg.diagonal (Array API compatible)
-   dpnp.linalg.matrix_tranpose (Array API compatible)
+   dpnp.linalg.matrix_transpose (Array API compatible)

 Exceptions
 ----------

@@ -44,7 +44,7 @@ Array type testing


 Logical operations
----------------
+------------------

 .. autosummary::
   :toctree: generated/

@@ -18,7 +18,7 @@ Constructing arrays
 -------------------

 New arrays can be constructed using the routines detailed in
-:ref:`Array Creation Routines <routines.creation>`, and also by using the low-level
+:ref:`Array Creation Routines <routines.array-creation>`, and also by using the low-level
 :class:`dpnp.ndarray` constructor:

 .. autosummary::

@@ -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._blas_impl extensions
 //
 //*****************************************************************************


@@ -26,7 +26,6 @@

 set(python_module_name _window_impl)
 set(_module_src
-    ${CMAKE_CURRENT_SOURCE_DIR}/hamming.cpp
    ${CMAKE_CURRENT_SOURCE_DIR}/window_py.cpp
 )


@@ -23,28 +23,56 @@
 // THE POSSIBILITY OF SUCH DAMAGE.
 //*****************************************************************************

+#pragma once
+
 #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"
+#include "utils/type_utils.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;

+typedef sycl::event (*window_fn_ptr_t)(sycl::queue &,
+                                       char *,
+                                       const std::size_t,
+                                       const std::vector<sycl::event> &);
+
+template <typename T, template <typename> class Functor>
+sycl::event window_impl(sycl::queue &q,
+                        char *result,
+                        const std::size_t nelems,
+                        const std::vector<sycl::event> &depends)
+{
+    dpctl::tensor::type_utils::validate_type_for_device<T>(q);
+
+    T *res = reinterpret_cast<T *>(result);
+
+    sycl::event window_ev = q.submit([&](sycl::handler &cgh) {
+        cgh.depends_on(depends);
+
+        using WindowKernel = Functor<T>;
+        cgh.parallel_for<WindowKernel>(sycl::range<1>(nelems),
+                                       WindowKernel(res, nelems));
+    });
+
+    return window_ev;
+}
+
+template <typename dispatchT>
 std::pair<sycl::event, sycl::event>
-    py_hamming(sycl::queue &exec_q,
-               const dpctl::tensor::usm_ndarray &result,
-               const std::vector<sycl::event> &depends)
+    py_window(sycl::queue &exec_q,
+              const dpctl::tensor::usm_ndarray &result,
+              const std::vector<sycl::event> &depends,
+              const dispatchT &window_dispatch_vector)
 {
    dpctl::tensor::validation::CheckWritable::throw_if_not_writable(result);

@@ -71,52 +99,27 @@ std::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];
+    auto fn = window_dispatch_vector[result_type_id];

    if (fn == nullptr) {
        throw std::runtime_error("Type of given array is not supported");
    }

    char *result_typeless_ptr = result.get_data();
-    sycl::event hamming_ev = fn(exec_q, result_typeless_ptr, nelems, depends);
+    sycl::event window_ev = fn(exec_q, result_typeless_ptr, nelems, depends);
    sycl::event args_ev =
-        dpctl::utils::keep_args_alive(exec_q, {result}, {hamming_ev});
+        dpctl::utils::keep_args_alive(exec_q, {result}, {window_ev});

-    return std::make_pair(args_ev, hamming_ev);
+    return std::make_pair(args_ev, window_ev);
 }

-template <typename fnT, typename T>
-struct HammingFactory
+template <template <typename fnT, typename T> typename factoryT>
+void init_window_dispatch_vectors(window_fn_ptr_t window_dispatch_vector[])
 {
-    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::DispatchVectorBuilder<window_fn_ptr_t, factoryT,
                                       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());
+    contig.populate_dispatch_vector(window_dispatch_vector);

    return;
 }

@@ -25,11 +25,42 @@

 #pragma once

-#include <pybind11/pybind11.h>
+#include "common.hpp"
+#include <sycl/sycl.hpp>

-namespace py = pybind11;
+namespace dpnp::extensions::window::kernels
+{

-namespace dpnp::extensions::window
+template <typename T>
+class HammingFunctor
 {
-void init_hamming(py::module_ m);
-}
+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));
+    }
+};
+
+template <typename fnT, typename T>
+struct HammingFactory
+{
+    fnT get()
+    {
+        if constexpr (std::is_floating_point_v<T>) {
+            return window_impl<T, HammingFunctor>;
+        }
+        else {
+            return nullptr;
+        }
+    }
+};
+
+} // namespace dpnp::extensions::window::kernels
@@ -28,10 +28,36 @@
 //*****************************************************************************

 #include <pybind11/pybind11.h>
+#include <pybind11/stl.h>

+#include "common.hpp"
 #include "hamming.hpp"

+namespace window_ns = dpnp::extensions::window;
+namespace py = pybind11;
+using window_ns::window_fn_ptr_t;
+
+namespace dpctl_td_ns = dpctl::tensor::type_dispatch;
+
+static window_fn_ptr_t hamming_dispatch_vector[dpctl_td_ns::num_types];
+
 PYBIND11_MODULE(_window_impl, m)
 {
-    dpnp::extensions::window::init_hamming(m);
+    using arrayT = dpctl::tensor::usm_ndarray;
+    using event_vecT = std::vector<sycl::event>;
+
+    {
+        window_ns::init_window_dispatch_vectors<
+            window_ns::kernels::HammingFactory>(hamming_dispatch_vector);
+
+        auto hamming_pyapi = [&](sycl::queue &exec_q, const arrayT &result,
+                                 const event_vecT &depends = {}) {
+            return window_ns::py_window(exec_q, result, depends,
+                                        hamming_dispatch_vector);
+        };
+
+        m.def("_hamming", hamming_pyapi, "Call hamming kernel",
+              py::arg("sycl_queue"), py::arg("result"),
+              py::arg("depends") = py::list());
+    }
 }
@@ -108,7 +108,7 @@ def _replace_nan(a, val):

    Returns
    -------
-    out : {dpnp.ndarray}
+    out : dpnp.ndarray
        If `a` is of inexact type, return a copy of `a` with the NaNs
        replaced by the fill value, otherwise return `a`.
    mask: {bool, None}