Add dedicated HistogramFunctor kernel

Author: antonwolfy

dpnp/backend/extensions/statistics/CMakeLists.txt

@@ -67,7 +67,7 @@ set_target_properties(
 
 target_include_directories(
     ${python_module_name}
-    PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../common
+    PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../ ${CMAKE_CURRENT_SOURCE_DIR}/../common
 )
 
 # treat below headers as system to suppress the warnings there during the build

dpnp/backend/extensions/statistics/histogram_common.hpp

@@ -28,9 +28,15 @@
 
 #pragma once
 
+#include <cstddef>
+#include <cstdint>
+#include <optional>
+#include <type_traits>
+
 #include <sycl/sycl.hpp>
 
 #include "ext/common.hpp"
+#include "kernels/statistics/histogram.hpp"
 
 namespace dpctl::tensor
 {
@@ -277,9 +283,6 @@ bool check_in_bounds(const dT &val, const dT &min, const dT &max)
     return !_less(val, min) && !_less(max, val) && !IsNan<dT>::isnan(val);
 }
 
-template <typename T, typename HistImpl, typename Edges, typename Weights>
-class histogram_kernel;
-
 template <typename T, typename HistImpl, typename Edges, typename Weights>
 void submit_histogram(const T *in,
                       const size_t size,
@@ -291,41 +294,12 @@ void submit_histogram(const T *in,
                       sycl::nd_range<1> nd_range,
                       sycl::handler &cgh)
 {
-    cgh.parallel_for<histogram_kernel<T, HistImpl, Edges, Weights>>(
-        nd_range, [=](sycl::nd_item<1> item) {
-            auto id = item.get_group_linear_id();
-            auto lid = item.get_local_linear_id();
-            auto group = item.get_group();
-            auto local_size = item.get_local_range(0);
-
-            hist.init(item);
-            edges.init(item);
-
-            if constexpr (HistImpl::sync_after_init || Edges::sync_after_init) {
-                sycl::group_barrier(group, sycl::memory_scope::work_group);
-            }
-
-            auto bounds = edges.get_bounds();
-
-            for (uint32_t i = 0; i < WorkPI; ++i) {
-                auto data_idx = id * WorkPI * local_size + i * local_size + lid;
-                if (data_idx < size) {
-                    auto *d = &in[data_idx * dims];
-
-                    if (edges.in_bounds(d, bounds)) {
-                        auto bin = edges.get_bin(item, d, bounds);
-                        auto weight = weights.get(data_idx);
-                        hist.add(item, bin, weight);
-                    }
-                }
-            }
-
-            if constexpr (HistImpl::sync_before_finalize) {
-                sycl::group_barrier(group, sycl::memory_scope::work_group);
-            }
+    using HistogramKernel =
+        dpnp::kernels::histogram::HistogramFunctor<T, HistImpl, Edges, Weights>;
 
-            hist.finalize(item);
-        });
+    cgh.parallel_for<HistogramKernel>(
+        nd_range,
+        HistogramKernel(in, size, dims, WorkPI, hist, edges, weights));
 }
 
 void validate(const usm_ndarray &sample,

dpnp/backend/kernels/statistics/histogram.hpp

@@ -0,0 +1,99 @@
+//*****************************************************************************
+// Copyright (c) 2026, 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.
+// - Neither the name of the copyright holder nor the names of its contributors
+//   may be used to endorse or promote products derived from this software
+//   without specific prior written permission.
+//
+// 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 <cstddef>
+#include <cstdint>
+
+#include <sycl/sycl.hpp>
+
+namespace dpnp::kernels::histogram
+{
+template <typename T, typename HistImpl, typename Edges, typename Weights>
+class HistogramFunctor
+{
+private:
+    const T *in = nullptr;
+    const std::size_t size;
+    const std::size_t dims;
+    const std::uint32_t WorkPI;
+    const HistImpl hist;
+    const Edges edges;
+    const Weights weights;
+
+public:
+    HistogramFunctor(const T *in_,
+                     const std::size_t size_,
+                     const std::size_t dims_,
+                     const std::uint32_t WorkPI_,
+                     const HistImpl &hist_,
+                     const Edges &edges_,
+                     const Weights &weights_)
+        : in(in_), size(size_), dims(dims_), WorkPI(WorkPI_), hist(hist_),
+          edges(edges_), weights(weights_)
+    {
+    }
+
+    void operator()(sycl::nd_item<1> item) const
+    {
+        auto id = item.get_group_linear_id();
+        auto lid = item.get_local_linear_id();
+        auto group = item.get_group();
+        auto local_size = item.get_local_range(0);
+
+        hist.init(item);
+        edges.init(item);
+
+        if constexpr (HistImpl::sync_after_init || Edges::sync_after_init) {
+            sycl::group_barrier(group, sycl::memory_scope::work_group);
+        }
+
+        auto bounds = edges.get_bounds();
+
+        for (uint32_t i = 0; i < WorkPI; ++i) {
+            auto data_idx = id * WorkPI * local_size + i * local_size + lid;
+            if (data_idx < size) {
+                auto *d = &in[data_idx * dims];
+
+                if (edges.in_bounds(d, bounds)) {
+                    auto bin = edges.get_bin(item, d, bounds);
+                    auto weight = weights.get(data_idx);
+                    hist.add(item, bin, weight);
+                }
+            }
+        }
+
+        if constexpr (HistImpl::sync_before_finalize) {
+            sycl::group_barrier(group, sycl::memory_scope::work_group);
+        }
+
+        hist.finalize(item);
+    }
+};
+} // namespace dpnp::kernels::histogram