Merge include-dpctl-tensor into move_tensor_elementwise_impl_unary

vlad-perevezentsev · vlad-perevezentsev · commit 5d11ff76dcd2 · 2026-03-11T04:46:07.000-07:00
diff --git a/dpctl_ext/tensor/libtensor/include/kernels/accumulators.hpp b/dpctl_ext/tensor/libtensor/include/kernels/accumulators.hpp
@@ -33,6 +33,8 @@
 //===---------------------------------------------------------------------===//
 
 #pragma once
+
+#include <algorithm>
 #include <array>
 #include <cstddef>
 #include <cstdint>
diff --git a/dpctl_ext/tensor/libtensor/include/kernels/elementwise_functions/minimum.hpp b/dpctl_ext/tensor/libtensor/include/kernels/elementwise_functions/minimum.hpp
@@ -37,6 +37,7 @@
 #include <cstddef>
 #include <cstdint>
 #include <type_traits>
+#include <vector>
 
 #include <sycl/sycl.hpp>
 
diff --git a/dpctl_ext/tensor/libtensor/include/kernels/reductions.hpp b/dpctl_ext/tensor/libtensor/include/kernels/reductions.hpp
@@ -40,7 +40,6 @@
 #include <complex>
 #include <cstddef>
 #include <cstdint>
-#include <stdexcept>
 #include <type_traits>
 #include <utility>
 #include <vector>
diff --git a/dpctl_ext/tensor/libtensor/include/kernels/sorting/radix_sort.hpp b/dpctl_ext/tensor/libtensor/include/kernels/sorting/radix_sort.hpp
@@ -37,6 +37,7 @@
 
 #include <algorithm>
 #include <array>
+#include <cassert>
 #include <cstddef>
 #include <cstdint>
 #include <limits>
diff --git a/dpctl_ext/tensor/libtensor/include/kernels/sorting/search_sorted_detail.hpp b/dpctl_ext/tensor/libtensor/include/kernels/sorting/search_sorted_detail.hpp
@@ -36,10 +36,7 @@
 
 #include <cstddef>
 
-namespace dpctl::tensor::kernels
-{
-
-namespace search_sorted_detail
+namespace dpctl::tensor::kernels::search_sorted_detail
 {
 
 template <typename T>
@@ -119,6 +116,4 @@ std::size_t upper_bound_indexed_impl(const Acc acc,
                                     acc_indexer);
 }
 
-} // namespace search_sorted_detail
-
-} // namespace dpctl::tensor::kernels
+} // namespace dpctl::tensor::kernels::search_sorted_detail
diff --git a/dpctl_ext/tensor/libtensor/include/kernels/sorting/searchsorted.hpp b/dpctl_ext/tensor/libtensor/include/kernels/sorting/searchsorted.hpp
@@ -35,7 +35,6 @@
 #pragma once
 
 #include <cstddef>
-#include <cstdint>
 #include <vector>
 
 #include <sycl/sycl.hpp>
diff --git a/dpctl_ext/tensor/libtensor/include/kernels/sorting/sort_utils.hpp b/dpctl_ext/tensor/libtensor/include/kernels/sorting/sort_utils.hpp
@@ -37,7 +37,6 @@
 #include <array>
 #include <cstddef>
 #include <cstdint>
-#include <span>
 #include <vector>
 
 #include <sycl/sycl.hpp>
@@ -78,10 +77,8 @@ sycl::event iota_impl(sycl::queue &exec_q,
             }
 
             if (offset + n_wi * max_sgSize < nelems) {
-                static constexpr auto group_ls_props = syclexp::properties{
-                    syclexp::data_placement_striped
-                    // , syclexp::full_group
-                };
+                static constexpr auto group_ls_props =
+                    syclexp::properties{syclexp::data_placement_striped};
 
                 auto out_multi_ptr = sycl::address_space_cast<
                     sycl::access::address_space::global_space,
diff --git a/dpctl_ext/tensor/libtensor/include/kernels/sorting/topk.hpp b/dpctl_ext/tensor/libtensor/include/kernels/sorting/topk.hpp
@@ -34,12 +34,14 @@
 
 #pragma once
 
+#include <algorithm>
+#include <cassert>
 #include <cstddef>
 #include <cstdint>
+#include <functional>
 #include <stdexcept>
 #include <vector>
 
-#include <sycl/ext/oneapi/sub_group_mask.hpp>
 #include <sycl/sycl.hpp>
 
 #include "kernels/sorting/merge_sort.hpp"
diff --git a/dpctl_ext/tensor/libtensor/include/utils/sycl_utils.hpp b/dpctl_ext/tensor/libtensor/include/utils/sycl_utils.hpp
@@ -33,9 +33,12 @@
 #pragma once
 
 #include <algorithm>
+#include <cmath>
 #include <complex>
 #include <cstddef>
 #include <cstdint>
+#include <iterator>
+#include <limits>
 #include <type_traits>
 #include <vector>
 
diff --git a/dpctl_ext/tensor/libtensor/include/utils/type_dispatch.hpp b/dpctl_ext/tensor/libtensor/include/utils/type_dispatch.hpp
@@ -32,6 +32,7 @@
 
 #pragma once
 
+#include <cassert>
 #include <cstdint>
 #include <stdexcept>
 #include <string>
diff --git a/dpctl_ext/tensor/libtensor/include/utils/type_dispatch_building.hpp b/dpctl_ext/tensor/libtensor/include/utils/type_dispatch_building.hpp
@@ -32,6 +32,7 @@
 
 #pragma once
 
+#include <cassert>
 #include <complex>
 #include <cstdint>
 #include <type_traits>
diff --git a/dpctl_ext/tensor/libtensor/source/accumulators/cumulative_logsumexp.cpp b/dpctl_ext/tensor/libtensor/source/accumulators/cumulative_logsumexp.cpp
@@ -337,7 +337,6 @@ void init_cumulative_logsumexp(py::module_ m)
 
     auto cumlogsumexp_dtype_supported = [&](const py::dtype &input_dtype,
                                             const py::dtype &output_dtype) {
-        using dpctl::tensor::py_internal::py_accumulate_dtype_supported;
         return py_accumulate_dtype_supported(
             input_dtype, output_dtype, cumlogsumexp_strided_dispatch_table);
     };
diff --git a/dpctl_ext/tensor/libtensor/source/accumulators/cumulative_prod.cpp b/dpctl_ext/tensor/libtensor/source/accumulators/cumulative_prod.cpp
@@ -346,7 +346,6 @@ void init_cumulative_prod(py::module_ m)
 
     auto cumprod_dtype_supported = [&](const py::dtype &input_dtype,
                                        const py::dtype &output_dtype) {
-        using dpctl::tensor::py_internal::py_accumulate_dtype_supported;
         return py_accumulate_dtype_supported(input_dtype, output_dtype,
                                              cumprod_strided_dispatch_table);
     };
diff --git a/dpctl_ext/tensor/libtensor/source/accumulators/cumulative_sum.cpp b/dpctl_ext/tensor/libtensor/source/accumulators/cumulative_sum.cpp
@@ -344,7 +344,6 @@ void init_cumulative_sum(py::module_ m)
 
     auto cumsum_dtype_supported = [&](const py::dtype &input_dtype,
                                       const py::dtype &output_dtype) {
-        using dpctl::tensor::py_internal::py_accumulate_dtype_supported;
         return py_accumulate_dtype_supported(input_dtype, output_dtype,
                                              cumsum_strided_dispatch_table);
     };
diff --git a/dpctl_ext/tensor/libtensor/source/boolean_advanced_indexing.cpp b/dpctl_ext/tensor/libtensor/source/boolean_advanced_indexing.cpp
@@ -35,6 +35,7 @@
 
 #include <algorithm>
 #include <array>
+#include <cassert>
 #include <cstddef>
 #include <cstdint>
 #include <functional>
diff --git a/dpctl_ext/tensor/libtensor/source/copy_and_cast_usm_to_usm.cpp b/dpctl_ext/tensor/libtensor/source/copy_and_cast_usm_to_usm.cpp
@@ -33,6 +33,7 @@
 //===----------------------------------------------------------------------===//
 
 #include <array>
+#include <cassert>
 #include <cstddef>
 #include <sycl/sycl.hpp>
 #include <tuple>
diff --git a/dpctl_ext/tensor/libtensor/source/reductions/argmax.cpp b/dpctl_ext/tensor/libtensor/source/reductions/argmax.cpp
@@ -264,7 +264,6 @@ void init_argmax(py::module_ m)
         auto argmax_pyapi = [&](const arrayT &src, int trailing_dims_to_reduce,
                                 const arrayT &dst, sycl::queue &exec_q,
                                 const event_vecT &depends = {}) {
-            using dpctl::tensor::py_internal::py_search_over_axis;
             return py_search_over_axis(
                 src, trailing_dims_to_reduce, dst, exec_q, depends,
                 argmax_over_axis_strided_temps_dispatch_table,
diff --git a/dpctl_ext/tensor/libtensor/source/reductions/argmin.cpp b/dpctl_ext/tensor/libtensor/source/reductions/argmin.cpp
@@ -264,7 +264,6 @@ void init_argmin(py::module_ m)
         auto argmin_pyapi = [&](const arrayT &src, int trailing_dims_to_reduce,
                                 const arrayT &dst, sycl::queue &exec_q,
                                 const event_vecT &depends = {}) {
-            using dpctl::tensor::py_internal::py_search_over_axis;
             return py_search_over_axis(
                 src, trailing_dims_to_reduce, dst, exec_q, depends,
                 argmin_over_axis_strided_temps_dispatch_table,
diff --git a/dpctl_ext/tensor/libtensor/source/reductions/logsumexp.cpp b/dpctl_ext/tensor/libtensor/source/reductions/logsumexp.cpp
@@ -40,6 +40,7 @@
 #include <sycl/sycl.hpp>
 
 #include "dpnp4pybind11.hpp"
+#include <pybind11/numpy.h>
 #include <pybind11/pybind11.h>
 #include <pybind11/stl.h>
 
@@ -233,7 +234,6 @@ void init_logsumexp(py::module_ m)
                                    int trailing_dims_to_reduce,
                                    const arrayT &dst, sycl::queue &exec_q,
                                    const event_vecT &depends = {}) {
-            using dpctl::tensor::py_internal::py_tree_reduction_over_axis;
             return py_tree_reduction_over_axis(
                 src, trailing_dims_to_reduce, dst, exec_q, depends,
                 logsumexp_over_axis_strided_temps_dispatch_table,
@@ -246,7 +246,6 @@ void init_logsumexp(py::module_ m)
 
         auto logsumexp_dtype_supported = [&](const py::dtype &input_dtype,
                                              const py::dtype &output_dtype) {
-            using dpctl::tensor::py_internal::py_tree_reduction_dtype_supported;
             return py_tree_reduction_dtype_supported(
                 input_dtype, output_dtype,
                 logsumexp_over_axis_strided_temps_dispatch_table);
diff --git a/dpctl_ext/tensor/libtensor/source/reductions/max.cpp b/dpctl_ext/tensor/libtensor/source/reductions/max.cpp
@@ -391,7 +391,6 @@ void init_max(py::module_ m)
         auto max_pyapi = [&](const arrayT &src, int trailing_dims_to_reduce,
                              const arrayT &dst, sycl::queue &exec_q,
                              const event_vecT &depends = {}) {
-            using dpctl::tensor::py_internal::py_reduction_over_axis;
             return py_reduction_over_axis(
                 src, trailing_dims_to_reduce, dst, exec_q, depends,
                 max_over_axis_strided_atomic_dispatch_table,
diff --git a/dpctl_ext/tensor/libtensor/source/reductions/min.cpp b/dpctl_ext/tensor/libtensor/source/reductions/min.cpp
@@ -393,7 +393,6 @@ void init_min(py::module_ m)
         auto min_pyapi = [&](const arrayT &src, int trailing_dims_to_reduce,
                              const arrayT &dst, sycl::queue &exec_q,
                              const event_vecT &depends = {}) {
-            using dpctl::tensor::py_internal::py_reduction_over_axis;
             return py_reduction_over_axis(
                 src, trailing_dims_to_reduce, dst, exec_q, depends,
                 min_over_axis_strided_atomic_dispatch_table,
diff --git a/dpctl_ext/tensor/libtensor/source/reductions/prod.cpp b/dpctl_ext/tensor/libtensor/source/reductions/prod.cpp
@@ -35,12 +35,14 @@
 
 #include <complex>
 #include <cstdint>
+#include <string>
 #include <type_traits>
 #include <vector>
 
 #include <sycl/sycl.hpp>
 
 #include "dpnp4pybind11.hpp"
+#include <pybind11/numpy.h>
 #include <pybind11/pybind11.h>
 #include <pybind11/stl.h>
 
@@ -432,7 +434,6 @@ void init_prod(py::module_ m)
         auto prod_pyapi = [&](const arrayT &src, int trailing_dims_to_reduce,
                               const arrayT &dst, sycl::queue &exec_q,
                               const event_vecT &depends = {}) {
-            using dpctl::tensor::py_internal::py_reduction_over_axis;
             return py_reduction_over_axis(
                 src, trailing_dims_to_reduce, dst, exec_q, depends,
                 prod_over_axis_strided_atomic_dispatch_table,
@@ -450,7 +451,6 @@ void init_prod(py::module_ m)
         auto prod_dtype_supported =
             [&](const py::dtype &input_dtype, const py::dtype &output_dtype,
                 const std::string &dst_usm_type, sycl::queue &q) {
-                using dpctl::tensor::py_internal::py_reduction_dtype_supported;
                 return py_reduction_dtype_supported(
                     input_dtype, output_dtype, dst_usm_type, q,
                     prod_over_axis_strided_atomic_dispatch_table,
diff --git a/dpctl_ext/tensor/libtensor/source/reductions/reduce_hypot.cpp b/dpctl_ext/tensor/libtensor/source/reductions/reduce_hypot.cpp
@@ -40,13 +40,14 @@
 #include <sycl/sycl.hpp>
 
 #include "dpnp4pybind11.hpp"
+#include <pybind11/numpy.h>
 #include <pybind11/pybind11.h>
 #include <pybind11/stl.h>
 
 #include "kernels/reductions.hpp"
+#include "utils/sycl_utils.hpp"
 #include "utils/type_dispatch_building.hpp"
 
-#include "reduction_atomic_support.hpp"
 #include "reduction_over_axis.hpp"
 
 namespace dpctl::tensor::py_internal
@@ -229,7 +230,6 @@ void init_reduce_hypot(py::module_ m)
         auto hypot_pyapi = [&](const arrayT &src, int trailing_dims_to_reduce,
                                const arrayT &dst, sycl::queue &exec_q,
                                const event_vecT &depends = {}) {
-            using dpctl::tensor::py_internal::py_tree_reduction_over_axis;
             return py_tree_reduction_over_axis(
                 src, trailing_dims_to_reduce, dst, exec_q, depends,
                 hypot_over_axis_strided_temps_dispatch_table,
@@ -242,7 +242,6 @@ void init_reduce_hypot(py::module_ m)
 
         auto hypot_dtype_supported = [&](const py::dtype &input_dtype,
                                          const py::dtype &output_dtype) {
-            using dpctl::tensor::py_internal::py_tree_reduction_dtype_supported;
             return py_tree_reduction_dtype_supported(
                 input_dtype, output_dtype,
                 hypot_over_axis_strided_temps_dispatch_table);
diff --git a/dpctl_ext/tensor/libtensor/source/reductions/reduction_atomic_support.hpp b/dpctl_ext/tensor/libtensor/source/reductions/reduction_atomic_support.hpp
@@ -34,7 +34,6 @@
 //===---------------------------------------------------------------------===//
 
 #pragma once
-#include <complex>
 #include <type_traits>
 
 #include <sycl/sycl.hpp>
diff --git a/dpctl_ext/tensor/libtensor/source/reductions/reduction_over_axis.hpp b/dpctl_ext/tensor/libtensor/source/reductions/reduction_over_axis.hpp
@@ -37,9 +37,10 @@
 
 #include <algorithm>
 #include <cstddef>
-#include <cstdint>
 #include <exception>
+#include <iterator>
 #include <stdexcept>
+#include <string>
 #include <type_traits>
 #include <utility>
 #include <vector>
@@ -49,7 +50,6 @@
 #include "dpnp4pybind11.hpp"
 #include <pybind11/numpy.h>
 #include <pybind11/pybind11.h>
-#include <pybind11/stl.h>
 
 #include "kernels/reductions.hpp"
 #include "simplify_iteration_space.hpp"
@@ -325,9 +325,6 @@ std::pair<sycl::event, sycl::event> py_reduction_over_axis(
         }
     }
 
-    using dpctl::tensor::py_internal::simplify_iteration_space;
-    using dpctl::tensor::py_internal::simplify_iteration_space_1;
-
     // TODO: not used anywhere
     auto const &src_shape_vecs = src.get_shape_vector();
     auto const &src_strides_vecs = src.get_strides_vector();
@@ -637,9 +634,6 @@ std::pair<sycl::event, sycl::event> py_tree_reduction_over_axis(
         }
     }
 
-    using dpctl::tensor::py_internal::simplify_iteration_space;
-    using dpctl::tensor::py_internal::simplify_iteration_space_1;
-
     auto const &src_shape_vecs = src.get_shape_vector();
     auto const &src_strides_vecs = src.get_strides_vector();
     auto const &dst_strides_vecs = dst.get_strides_vector();
@@ -912,8 +906,6 @@ std::pair<sycl::event, sycl::event> py_search_over_axis(
         }
     }
 
-    using dpctl::tensor::py_internal::simplify_iteration_space;
-
     auto const &src_shape_vecs = src.get_shape_vector();
     auto const &src_strides_vecs = src.get_strides_vector();
     auto const &dst_strides_vecs = dst.get_strides_vector();
@@ -1205,7 +1197,6 @@ std::pair<sycl::event, sycl::event>
     shT simplified_red_src_strides;
     py::ssize_t red_src_offset(0);
 
-    using dpctl::tensor::py_internal::simplify_iteration_space_1;
     simplify_iteration_space_1(
         simplified_red_nd, src_shape_ptr + dst_nd, red_src_strides,
         // output
@@ -1231,7 +1222,6 @@ std::pair<sycl::event, sycl::event>
         simplified_iter_dst_strides.push_back(0);
     }
     else {
-        using dpctl::tensor::py_internal::simplify_iteration_space;
         simplify_iteration_space(
             iter_nd, src_shape_ptr, iter_src_strides, iter_dst_strides,
             // output
diff --git a/dpctl_ext/tensor/libtensor/source/reductions/sum.cpp b/dpctl_ext/tensor/libtensor/source/reductions/sum.cpp
@@ -35,12 +35,14 @@
 
 #include <complex>
 #include <cstdint>
+#include <string>
 #include <type_traits>
 #include <vector>
 
 #include <sycl/sycl.hpp>
 
 #include "dpnp4pybind11.hpp"
+#include <pybind11/numpy.h>
 #include <pybind11/pybind11.h>
 #include <pybind11/stl.h>
 
@@ -429,7 +431,6 @@ void init_sum(py::module_ m)
         auto sum_pyapi = [&](const arrayT &src, int trailing_dims_to_reduce,
                              const arrayT &dst, sycl::queue &exec_q,
                              const event_vecT &depends = {}) {
-            using dpctl::tensor::py_internal::py_reduction_over_axis;
             return py_reduction_over_axis(
                 src, trailing_dims_to_reduce, dst, exec_q, depends,
                 sum_over_axis_strided_atomic_dispatch_table,
@@ -447,7 +448,6 @@ void init_sum(py::module_ m)
         auto sum_dtype_supported =
             [&](const py::dtype &input_dtype, const py::dtype &output_dtype,
                 const std::string &dst_usm_type, sycl::queue &q) {
-                using dpctl::tensor::py_internal::py_reduction_dtype_supported;
                 return py_reduction_dtype_supported(
                     input_dtype, output_dtype, dst_usm_type, q,
                     sum_over_axis_strided_atomic_dispatch_table,
diff --git a/dpctl_ext/tensor/libtensor/source/simplify_iteration_space.cpp b/dpctl_ext/tensor/libtensor/source/simplify_iteration_space.cpp
diff --git a/dpctl_ext/tensor/libtensor/source/sorting/isin.cpp b/dpctl_ext/tensor/libtensor/source/sorting/isin.cpp
diff --git a/dpctl_ext/tensor/libtensor/source/sorting/merge_argsort.cpp b/dpctl_ext/tensor/libtensor/source/sorting/merge_argsort.cpp
diff --git a/dpctl_ext/tensor/libtensor/source/sorting/merge_sort.cpp b/dpctl_ext/tensor/libtensor/source/sorting/merge_sort.cpp
diff --git a/dpctl_ext/tensor/libtensor/source/sorting/py_argsort_common.hpp b/dpctl_ext/tensor/libtensor/source/sorting/py_argsort_common.hpp
diff --git a/dpctl_ext/tensor/libtensor/source/sorting/py_sort_common.hpp b/dpctl_ext/tensor/libtensor/source/sorting/py_sort_common.hpp
diff --git a/dpctl_ext/tensor/libtensor/source/sorting/radix_argsort.cpp b/dpctl_ext/tensor/libtensor/source/sorting/radix_argsort.cpp
diff --git a/dpctl_ext/tensor/libtensor/source/sorting/radix_sort.cpp b/dpctl_ext/tensor/libtensor/source/sorting/radix_sort.cpp
diff --git a/dpctl_ext/tensor/libtensor/source/sorting/searchsorted.cpp b/dpctl_ext/tensor/libtensor/source/sorting/searchsorted.cpp
diff --git a/dpctl_ext/tensor/libtensor/source/sorting/topk.cpp b/dpctl_ext/tensor/libtensor/source/sorting/topk.cpp
diff --git a/dpctl_ext/tensor/libtensor/source/tensor_reductions.cpp b/dpctl_ext/tensor/libtensor/source/tensor_reductions.cpp
diff --git a/dpctl_ext/tensor/libtensor/source/tensor_sorting.cpp b/dpctl_ext/tensor/libtensor/source/tensor_sorting.cpp
diff --git a/dpnp/backend/include/dpnp4pybind11.hpp b/dpnp/backend/include/dpnp4pybind11.hpp
