Add dedicated windows functors

Author: antonwolfy

dpnp/backend/extensions/indexing/choose.cpp

@@ -56,9 +56,10 @@
 
 namespace dpnp::extensions::indexing
 {
+namespace py = pybind11;
+
 namespace impl
 {
-namespace py = pybind11;
 namespace td_ns = dpctl::tensor::type_dispatch;
 
 using dpctl::tensor::ssize_t;

dpnp/backend/extensions/ufunc/elementwise_functions/interpolate.cpp

@@ -51,28 +51,19 @@
 #include "ext/common.hpp"
 #include "ext/validation_utils.hpp"
 
-namespace py = pybind11;
-namespace td_ns = dpctl::tensor::type_dispatch;
-namespace type_utils = dpctl::tensor::type_utils;
-
-using ext::common::value_type_of;
-using ext::validation::array_names;
-
-using ext::common::dtype_from_typenum;
-using ext::validation::check_has_dtype;
-using ext::validation::check_num_dims;
-using ext::validation::check_same_dtype;
-using ext::validation::check_same_size;
-using ext::validation::common_checks;
-
 namespace dpnp::extensions::ufunc
 {
+namespace py = pybind11;
+
 namespace impl
 {
-using ext::common::init_dispatch_vector;
+namespace td_ns = dpctl::tensor::type_dispatch;
+namespace type_utils = dpctl::tensor::type_utils;
 
 template <typename T>
-using value_type_of_t = typename value_type_of<T>::type;
+using value_type_of_t = typename ext::common::value_type_of<T>::type;
+
+using ext::common::dtype_from_typenum;
 
 typedef sycl::event (*interpolate_fn_ptr_t)(sycl::queue &,
                                             const void *,      // x
@@ -163,6 +154,13 @@ struct InterpolateFactory
 
 namespace detail
 {
+using ext::validation::array_names;
+using ext::validation::check_has_dtype;
+using ext::validation::check_num_dims;
+using ext::validation::check_same_dtype;
+using ext::validation::check_same_size;
+using ext::validation::common_checks;
+
 void validate(const dpctl::tensor::usm_ndarray &x,
               const dpctl::tensor::usm_ndarray &idx,
               const dpctl::tensor::usm_ndarray &xp,
@@ -258,6 +256,7 @@ std::pair<sycl::event, sycl::event>
 
 static void init_interpolate_dispatch_vectors()
 {
+    using ext::common::init_dispatch_vector;
     init_dispatch_vector<interpolate_fn_ptr_t, impl::InterpolateFactory>(
         interpolate_dispatch_vector);
 }

dpnp/backend/extensions/window/common.hpp

@@ -28,11 +28,11 @@
 
 #pragma once
 
-#include <pybind11/pybind11.h>
-#include <pybind11/stl.h>
 #include <sycl/sycl.hpp>
 
 #include "dpctl4pybind11.hpp"
+#include <pybind11/pybind11.h>
+#include <pybind11/stl.h>
 
 // dpctl tensor headers
 #include "utils/output_validation.hpp"
@@ -41,10 +41,8 @@
 
 namespace dpnp::extensions::window
 {
-
-namespace dpctl_td_ns = dpctl::tensor::type_dispatch;
-
 namespace py = pybind11;
+namespace td_ns = dpctl::tensor::type_dispatch;
 
 typedef sycl::event (*window_fn_ptr_t)(sycl::queue &,
                                        char *,
@@ -72,6 +70,20 @@ sycl::event window_impl(sycl::queue &exec_q,
     return window_ev;
 }
 
+template <typename fnT, typename T, template <typename> typename FunctorT>
+struct Factory
+{
+    fnT get()
+    {
+        if constexpr (std::is_floating_point_v<T>) {
+            return window_impl<T, FunctorT>;
+        }
+        else {
+            return nullptr;
+        }
+    }
+};
+
 template <typename funcPtrT>
 std::tuple<size_t, char *, funcPtrT>
     window_fn(sycl::queue &exec_q,
@@ -101,7 +113,7 @@ std::tuple<size_t, char *, funcPtrT>
     }
 
     const int result_typenum = result.get_typenum();
-    auto array_types = dpctl_td_ns::usm_ndarray_types();
+    auto array_types = td_ns::usm_ndarray_types();
     const int result_type_id = array_types.typenum_to_lookup_id(result_typenum);
     funcPtrT fn = window_dispatch_vector[result_type_id];
 

dpnp/backend/extensions/window/kaiser.cpp

@@ -29,6 +29,8 @@
 #include "kaiser.hpp"
 #include "common.hpp"
 
+#include "kernels/window/kaiser.hpp"
+
 // utils extension header
 #include "ext/common.hpp"
 
@@ -39,48 +41,21 @@
 
 #include <sycl/sycl.hpp>
 
-#include "kernels/elementwise_functions/i0.hpp"
-
 namespace dpnp::extensions::window
 {
-namespace dpctl_td_ns = dpctl::tensor::type_dispatch;
-
-using ext::common::init_dispatch_vector;
+namespace py = pybind11;
+namespace td_ns = dpctl::tensor::type_dispatch;
 
 typedef sycl::event (*kaiser_fn_ptr_t)(sycl::queue &,
                                        char *,
                                        const std::size_t,
                                        const py::object &,
                                        const std::vector<sycl::event> &);
 
-static kaiser_fn_ptr_t kaiser_dispatch_vector[dpctl_td_ns::num_types];
+static kaiser_fn_ptr_t kaiser_dispatch_vector[td_ns::num_types];
 
-template <typename T>
-class KaiserFunctor
+namespace impl
 {
-private:
-    T *res = nullptr;
-    const std::size_t N;
-    const T beta;
-
-public:
-    KaiserFunctor(T *res, const std::size_t N, const T beta)
-        : res(res), N(N), beta(beta)
-    {
-    }
-
-    void operator()(sycl::id<1> id) const
-    {
-        using dpnp::kernels::i0::cyl_bessel_i0;
-
-        const auto i = id.get(0);
-        const T alpha = (N - 1) / T(2);
-        const T tmp = (i - alpha) / alpha;
-        res[i] = cyl_bessel_i0(beta * sycl::sqrt(1 - tmp * tmp)) /
-                 cyl_bessel_i0(beta);
-    }
-};
-
 template <typename T>
 sycl::event kaiser_impl(sycl::queue &exec_q,
                         char *result,
@@ -96,7 +71,7 @@ sycl::event kaiser_impl(sycl::queue &exec_q,
     sycl::event kaiser_ev = exec_q.submit([&](sycl::handler &cgh) {
         cgh.depends_on(depends);
 
-        using KaiserKernel = KaiserFunctor<T>;
+        using KaiserKernel = dpnp::kernels::kaiser::KaiserFunctor<T>;
         cgh.parallel_for<KaiserKernel>(sycl::range<1>(nelems),
                                        KaiserKernel(res, nelems, beta));
     });
@@ -138,11 +113,12 @@ std::pair<sycl::event, sycl::event>
 
     return std::make_pair(args_ev, kaiser_ev);
 }
+} // namespace impl
 
 void init_kaiser_dispatch_vectors()
 {
-    init_dispatch_vector<kaiser_fn_ptr_t, KaiserFactory>(
+    using ext::common::init_dispatch_vector;
+    init_dispatch_vector<kaiser_fn_ptr_t, impl::KaiserFactory>(
         kaiser_dispatch_vector);
 }
-
 } // namespace dpnp::extensions::window

dpnp/backend/extensions/window/kaiser.hpp

@@ -28,17 +28,20 @@
 
 #pragma once
 
-#include <dpctl4pybind11.hpp>
 #include <sycl/sycl.hpp>
 
+#include <dpctl4pybind11.hpp>
+#include <pybind11/pybind11.h>
+
 namespace dpnp::extensions::window
 {
+namespace py = pybind11;
+
 extern std::pair<sycl::event, sycl::event>
     py_kaiser(sycl::queue &exec_q,
               const py::object &beta,
               const dpctl::tensor::usm_ndarray &result,
               const std::vector<sycl::event> &depends);
 
 extern void init_kaiser_dispatch_vectors(void);
-
 } // namespace dpnp::extensions::window

dpnp/backend/extensions/window/window_py.cpp

@@ -33,11 +33,12 @@
 #include <pybind11/pybind11.h>
 #include <pybind11/stl.h>
 
-#include "bartlett.hpp"
-#include "blackman.hpp"
+#include "kernels/window/bartlett.hpp"
+#include "kernels/window/blackman.hpp"
+#include "kernels/window/hamming.hpp"
+#include "kernels/window/hanning.hpp"
+
 #include "common.hpp"
-#include "hamming.hpp"
-#include "hanning.hpp"
 #include "kaiser.hpp"
 
 // utils extension header
@@ -51,6 +52,22 @@ using window_ns::window_fn_ptr_t;
 
 namespace dpctl_td_ns = dpctl::tensor::type_dispatch;
 
+template <typename fnT, typename T>
+using BartlettFactory =
+    window_ns::Factory<fnT, T, dpnp::kernels::bartlett::BartlettFunctor>;
+
+template <typename fnT, typename T>
+using BlackmanFactory =
+    window_ns::Factory<fnT, T, dpnp::kernels::blackman::BlackmanFunctor>;
+
+template <typename fnT, typename T>
+using HammingFactory =
+    window_ns::Factory<fnT, T, dpnp::kernels::hamming::HammingFunctor>;
+
+template <typename fnT, typename T>
+using HanningFactory =
+    window_ns::Factory<fnT, T, dpnp::kernels::hanning::HanningFunctor>;
+
 static window_fn_ptr_t bartlett_dispatch_vector[dpctl_td_ns::num_types];
 static window_fn_ptr_t blackman_dispatch_vector[dpctl_td_ns::num_types];
 static window_fn_ptr_t hamming_dispatch_vector[dpctl_td_ns::num_types];
@@ -62,8 +79,7 @@ PYBIND11_MODULE(_window_impl, m)
     using event_vecT = std::vector<sycl::event>;
 
     {
-        init_dispatch_vector<window_ns::window_fn_ptr_t,
-                             window_ns::kernels::BartlettFactory>(
+        init_dispatch_vector<window_ns::window_fn_ptr_t, BartlettFactory>(
             bartlett_dispatch_vector);
 
         auto bartlett_pyapi = [&](sycl::queue &exec_q, const arrayT &result,
@@ -78,8 +94,7 @@ PYBIND11_MODULE(_window_impl, m)
     }
 
     {
-        init_dispatch_vector<window_ns::window_fn_ptr_t,
-                             window_ns::kernels::BlackmanFactory>(
+        init_dispatch_vector<window_ns::window_fn_ptr_t, BlackmanFactory>(
             blackman_dispatch_vector);
 
         auto blackman_pyapi = [&](sycl::queue &exec_q, const arrayT &result,
@@ -94,8 +109,7 @@ PYBIND11_MODULE(_window_impl, m)
     }
 
     {
-        init_dispatch_vector<window_ns::window_fn_ptr_t,
-                             window_ns::kernels::HammingFactory>(
+        init_dispatch_vector<window_ns::window_fn_ptr_t, HammingFactory>(
             hamming_dispatch_vector);
 
         auto hamming_pyapi = [&](sycl::queue &exec_q, const arrayT &result,
@@ -110,8 +124,7 @@ PYBIND11_MODULE(_window_impl, m)
     }
 
     {
-        init_dispatch_vector<window_ns::window_fn_ptr_t,
-                             window_ns::kernels::HanningFactory>(
+        init_dispatch_vector<window_ns::window_fn_ptr_t, HanningFactory>(
             hanning_dispatch_vector);
 
         auto hanning_pyapi = [&](sycl::queue &exec_q, const arrayT &result,

dpnp/backend/kernels/indexing/choose.hpp

@@ -1,5 +1,5 @@
 //*****************************************************************************
-// Copyright (c) 2024, Intel Corporation
+// Copyright (c) 2026, Intel Corporation
 // All rights reserved.
 //
 // Redistribution and use in source and binary forms, with or without

…p/backend/extensions/window/bartlett.hpp dpnp/backend/kernels/window/bartlett.hppdpnp/backend/extensions/window/bartlett.hpp renamed to dpnp/backend/kernels/window/bartlett.hpp dpnp/backend/extensions/window/bartlett.hpp renamed to dpnp/backend/kernels/window/bartlett.hpp

@@ -1,5 +1,5 @@
 //*****************************************************************************
-// Copyright (c) 2025, Intel Corporation
+// Copyright (c) 2026, Intel Corporation
 // All rights reserved.
 //
 // Redistribution and use in source and binary forms, with or without
@@ -19,7 +19,7 @@
 // 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, RES, OR PROFITS; OR BUSINESS
+// 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
@@ -28,12 +28,10 @@
 
 #pragma once
 
-#include "common.hpp"
 #include <sycl/sycl.hpp>
 
-namespace dpnp::extensions::window::kernels
+namespace dpnp::kernels::bartlett
 {
-
 template <typename T>
 class BartlettFunctor
 {
@@ -52,19 +50,4 @@ class BartlettFunctor
         res[i] = T(1) - sycl::fabs(i - alpha) / alpha;
     }
 };
-
-template <typename fnT, typename T>
-struct BartlettFactory
-{
-    fnT get()
-    {
-        if constexpr (std::is_floating_point_v<T>) {
-            return window_impl<T, BartlettFunctor>;
-        }
-        else {
-            return nullptr;
-        }
-    }
-};
-
-} // namespace dpnp::extensions::window::kernels
+} // namespace dpnp::kernels::bartlett