Add dedicated SlidingWindow1dFunctor and SlidingWindow1dSmallFunctor kernels

antonwolfy · antonwolfy · commit 6295ca3aa04b · 2026-03-17T15:06:32.000+01:00
diff --git a/dpnp/backend/extensions/statistics/histogram_common.hpp b/dpnp/backend/extensions/statistics/histogram_common.hpp
@@ -35,23 +35,19 @@
 
 #include <sycl/sycl.hpp>
 
+#include "dpctl4pybind11.hpp"
+
 #include "ext/common.hpp"
 #include "kernels/statistics/histogram.hpp"
 
-namespace dpctl::tensor
+namespace statistics::histogram
 {
-class usm_ndarray;
-}
-
 using dpctl::tensor::usm_ndarray;
 
 using ext::common::AtomicOp;
 using ext::common::IsNan;
 using ext::common::Less;
 
-namespace statistics::histogram
-{
-
 template <typename T, int Dims>
 struct CachedData
 {
diff --git a/dpnp/backend/extensions/statistics/sliding_window1d.hpp b/dpnp/backend/extensions/statistics/sliding_window1d.hpp
@@ -28,23 +28,19 @@
 
 #pragma once
 
-#include <algorithm>
-
-#include "utils/math_utils.hpp"
-#include <sycl/sycl.hpp>
+#include <cstddef>
+#include <cstdint>
 #include <type_traits>
 
-#include <stdio.h>
-
-#include "ext/common.hpp"
+#include <sycl/sycl.hpp>
 
-using dpctl::tensor::usm_ndarray;
+#include "dpctl4pybind11.hpp"
 
-using ext::common::Align;
-using ext::common::CeilDiv;
+#include "kernels/statistics/sliding_window1d.hpp"
 
 namespace statistics::sliding_window1d
 {
+using dpctl::tensor::usm_ndarray;
 
 template <typename T, uint32_t Size>
 class _RegistryDataStorage
@@ -464,60 +460,6 @@ PaddedSpan<T, SizeT>
     return PaddedSpan<T, SizeT>(data, size, offset);
 }
 
-template <typename Results,
-          typename AData,
-          typename VData,
-          typename Op,
-          typename Red>
-void process_block(Results &results,
-                   uint32_t r_size,
-                   AData &a_data,
-                   VData &v_data,
-                   uint32_t block_size,
-                   Op op,
-                   Red red)
-{
-    for (uint32_t i = 0; i < block_size; ++i) {
-        auto v_val = v_data.broadcast(i);
-        for (uint32_t r = 0; r < r_size; ++r) {
-            results[r] = red(results[r], op(a_data[r], v_val));
-        }
-        a_data.advance_left();
-    }
-}
-
-template <typename SizeT>
-SizeT get_global_linear_id(const uint32_t wpi, const sycl::nd_item<1> &item)
-{
-    auto sbgroup = item.get_sub_group();
-    const auto sg_loc_id = sbgroup.get_local_linear_id();
-
-    const SizeT sg_base_id = wpi * (item.get_global_linear_id() - sg_loc_id);
-    const SizeT id = sg_base_id + sg_loc_id;
-
-    return id;
-}
-
-template <typename SizeT>
-uint32_t get_results_num(const uint32_t wpi,
-                         const SizeT size,
-                         const SizeT global_id,
-                         const sycl::nd_item<1> &item)
-{
-    auto sbgroup = item.get_sub_group();
-
-    const auto sbg_size = sbgroup.get_max_local_range()[0];
-    const auto size_ = sycl::sub_sat(size, global_id);
-    return std::min(SizeT(wpi), CeilDiv(size_, sbg_size));
-}
-
-template <uint32_t WorkPI,
-          typename T,
-          typename SizeT,
-          typename Op,
-          typename Red>
-class sliding_window1d_kernel;
-
 template <uint32_t WorkPI,
           typename T,
           typename SizeT,
@@ -531,76 +473,15 @@ void submit_sliding_window1d(const PaddedSpan<const T, SizeT> &a,
                              sycl::nd_range<1> nd_range,
                              sycl::handler &cgh)
 {
-    cgh.parallel_for<sliding_window1d_kernel<WorkPI, T, SizeT, Op, Red>>(
-        nd_range, [=](sycl::nd_item<1> item) {
-            auto glid = get_global_linear_id<SizeT>(WorkPI, item);
-
-            auto results = RegistryData<T, WorkPI>(item);
-            results.fill(0);
-
-            auto results_num = get_results_num(WorkPI, out.size(), glid, item);
-
-            const auto *a_begin = a.begin();
-            const auto *a_end = a.end();
+    using SlidingWindow1dKernel =
+        dpnp::kernels::sliding_window1d::SlidingWindow1dFunctor<
+            WorkPI, PaddedSpan<const T, SizeT>, Span<const T, SizeT>, Op, Red,
+            Span<T, SizeT>, RegistryData, RegistryWindow>;
 
-            auto sbgroup = item.get_sub_group();
-
-            const auto chunks_count =
-                CeilDiv(v.size(), sbgroup.get_max_local_range()[0]);
-
-            const auto *a_ptr = &a.padded_begin()[glid];
-
-            auto _a_load_cond = [a_begin, a_end](auto &&ptr) {
-                return ptr >= a_begin && ptr < a_end;
-            };
-
-            auto a_data = RegistryWindow<const T, WorkPI + 1>(item);
-            a_ptr = a_data.load(a_ptr, _a_load_cond, 0);
-
-            const auto *v_ptr = &v.begin()[sbgroup.get_local_linear_id()];
-            auto v_size = v.size();
-
-            for (uint32_t b = 0; b < chunks_count; ++b) {
-                auto v_data = RegistryData<const T>(item);
-                v_ptr = v_data.load(v_ptr, v_data.x() < v_size, 0);
-
-                uint32_t chunk_size_ =
-                    std::min(v_size, SizeT(v_data.total_size()));
-                process_block(results, results_num, a_data, v_data, chunk_size_,
-                              op, red);
-
-                if (b != chunks_count - 1) {
-                    a_ptr = a_data.load_lane(a_data.size_y() - 1, a_ptr,
-                                             _a_load_cond, 0);
-                    v_size -= v_data.total_size();
-                }
-            }
-
-            auto *const out_ptr = out.begin();
-            // auto *const out_end = out.end();
-
-            auto y_start = glid;
-            auto y_stop =
-                std::min(y_start + WorkPI * results.size_x(), out.size());
-            uint32_t i = 0;
-            for (uint32_t y = y_start; y < y_stop; y += results.size_x()) {
-                out_ptr[y] = results[i++];
-            }
-            // while the code itself seems to be valid, inside correlate
-            // kernel it results in memory corruption. Further investigation
-            // is needed. SAT-7693
-            // corruption results.store(&out_ptr[glid],
-            //               [out_end](auto &&ptr) { return ptr < out_end; });
-        });
+    cgh.parallel_for<SlidingWindow1dKernel>(
+        nd_range, SlidingWindow1dKernel(a, v, op, red, out));
 }
 
-template <uint32_t WorkPI,
-          typename T,
-          typename SizeT,
-          typename Op,
-          typename Red>
-class sliding_window1d_small_kernel;
-
 template <uint32_t WorkPI,
           typename T,
           typename SizeT,
@@ -614,56 +495,13 @@ void submit_sliding_window1d_small_kernel(const PaddedSpan<const T, SizeT> &a,
                                           sycl::nd_range<1> nd_range,
                                           sycl::handler &cgh)
 {
-    cgh.parallel_for<sliding_window1d_small_kernel<WorkPI, T, SizeT, Op, Red>>(
-        nd_range, [=](sycl::nd_item<1> item) {
-            auto glid = get_global_linear_id<SizeT>(WorkPI, item);
-
-            auto results = RegistryData<T, WorkPI>(item);
-            results.fill(0);
-
-            auto sbgroup = item.get_sub_group();
-            auto sg_size = sbgroup.get_max_local_range()[0];
-
-            const uint32_t to_read = WorkPI * sg_size + v.size();
-            const auto *a_begin = a.begin();
-
-            const auto *a_ptr = &a.padded_begin()[glid];
-            const auto *a_end = std::min(a_ptr + to_read, a.end());
-
-            auto _a_load_cond = [a_begin, a_end](auto &&ptr) {
-                return ptr >= a_begin && ptr < a_end;
-            };
+    using SlidingWindow1dSmallKernel =
+        dpnp::kernels::sliding_window1d::SlidingWindow1dSmallFunctor<
+            WorkPI, PaddedSpan<const T, SizeT>, Span<const T, SizeT>, Op, Red,
+            Span<T, SizeT>, RegistryData, RegistryWindow>;
 
-            auto a_data = RegistryWindow<const T, WorkPI + 1>(item);
-            a_data.load(a_ptr, _a_load_cond, 0);
-
-            const auto *v_ptr = &v.begin()[sbgroup.get_local_linear_id()];
-            auto v_size = v.size();
-
-            auto v_data = RegistryData<const T>(item);
-            v_ptr = v_data.load(v_ptr, v_data.x() < v_size, 0);
-
-            auto results_num = get_results_num(WorkPI, out.size(), glid, item);
-
-            process_block(results, results_num, a_data, v_data, v_size, op,
-                          red);
-
-            auto *const out_ptr = out.begin();
-            // auto *const out_end = out.end();
-
-            auto y_start = glid;
-            auto y_stop =
-                std::min(y_start + WorkPI * results.size_x(), out.size());
-            uint32_t i = 0;
-            for (uint32_t y = y_start; y < y_stop; y += results.size_x()) {
-                out_ptr[y] = results[i++];
-            }
-            // while the code itself seems to be valid, inside correlate
-            // kernel it results in memory corruption. Further investigation
-            // is needed. SAT-7693
-            // corruption results.store(&out_ptr[glid],
-            //               [out_end](auto &&ptr) { return ptr < out_end; });
-        });
+    cgh.parallel_for<SlidingWindow1dSmallKernel>(
+        nd_range, SlidingWindow1dSmallKernel(a, v, op, red, out));
 }
 
 void validate(const usm_ndarray &a,
diff --git a/dpnp/backend/kernels/statistics/sliding_window1d.hpp b/dpnp/backend/kernels/statistics/sliding_window1d.hpp
