Merge move_tensor_sorting_impl into move_tensor_reductions_impl_ext

Author: vlad-perevezentsev

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

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>

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,

dpctl_ext/tensor/libtensor/include/kernels/sorting/topk.hpp

@@ -34,12 +34,13 @@
 
 #pragma once
 
+#include <algorithm>
 #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"

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);
     };

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);
     };

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);
     };

dpctl_ext/tensor/libtensor/source/sorting/isin.cpp

@@ -35,6 +35,7 @@
 
 #include <cstddef>
 #include <stdexcept>
+#include <tuple>
 #include <utility>
 #include <vector>
 
@@ -46,6 +47,7 @@
 
 #include "kernels/sorting/isin.hpp"
 #include "utils/memory_overlap.hpp"
+#include "utils/offset_utils.hpp"
 #include "utils/output_validation.hpp"
 #include "utils/sycl_alloc_utils.hpp"
 #include "utils/type_dispatch.hpp"
@@ -254,7 +256,7 @@ std::pair<sycl::event, sycl::event>
         simplified_dst_strides.push_back(0);
     }
     else {
-        dpctl::tensor::py_internal::simplify_iteration_space(
+        simplify_iteration_space(
             // modified by reference
             simplified_nd,
             // read-only inputs
@@ -313,9 +315,8 @@ std::pair<sycl::event, sycl::event>
 
 void init_isin_functions(py::module_ m)
 {
-    dpctl::tensor::py_internal::detail::init_isin_dispatch_vector();
+    detail::init_isin_dispatch_vector();
 
-    using dpctl::tensor::py_internal::py_isin;
     m.def("_isin", &py_isin, py::arg("needles"), py::arg("hay"), py::arg("dst"),
           py::arg("sycl_queue"), py::arg("invert"),
           py::arg("depends") = py::list());

dpctl_ext/tensor/libtensor/source/sorting/merge_argsort.cpp

@@ -35,6 +35,8 @@
 
 #include <cstdint>
 #include <type_traits>
+#include <utility>
+#include <vector>
 
 #include <sycl/sycl.hpp>
 
@@ -121,18 +123,16 @@ void init_merge_argsort_dispatch_tables(void)
 
 void init_merge_argsort_functions(py::module_ m)
 {
-    dpctl::tensor::py_internal::init_merge_argsort_dispatch_tables();
+    init_merge_argsort_dispatch_tables();
 
     auto py_argsort_ascending = [](const dpctl::tensor::usm_ndarray &src,
                                    const int trailing_dims_to_sort,
                                    const dpctl::tensor::usm_ndarray &dst,
                                    sycl::queue &exec_q,
                                    const std::vector<sycl::event> &depends)
         -> std::pair<sycl::event, sycl::event> {
-        return dpctl::tensor::py_internal::py_argsort(
-            src, trailing_dims_to_sort, dst, exec_q, depends,
-            dpctl::tensor::py_internal::
-                ascending_argsort_contig_dispatch_table);
+        return py_argsort(src, trailing_dims_to_sort, dst, exec_q, depends,
+                          ascending_argsort_contig_dispatch_table);
     };
     m.def("_argsort_ascending", py_argsort_ascending, py::arg("src"),
           py::arg("trailing_dims_to_sort"), py::arg("dst"),
@@ -144,10 +144,8 @@ void init_merge_argsort_functions(py::module_ m)
                                     sycl::queue &exec_q,
                                     const std::vector<sycl::event> &depends)
         -> std::pair<sycl::event, sycl::event> {
-        return dpctl::tensor::py_internal::py_argsort(
-            src, trailing_dims_to_sort, dst, exec_q, depends,
-            dpctl::tensor::py_internal::
-                descending_argsort_contig_dispatch_table);
+        return py_argsort(src, trailing_dims_to_sort, dst, exec_q, depends,
+                          descending_argsort_contig_dispatch_table);
     };
     m.def("_argsort_descending", py_argsort_descending, py::arg("src"),
           py::arg("trailing_dims_to_sort"), py::arg("dst"),

dpctl_ext/tensor/libtensor/source/sorting/merge_sort.cpp

@@ -33,6 +33,9 @@
 /// extension.
 //===----------------------------------------------------------------------===//
 
+#include <utility>
+#include <vector>
+
 #include <sycl/sycl.hpp>
 
 #include "dpnp4pybind11.hpp"
@@ -102,17 +105,16 @@ void init_merge_sort_dispatch_vectors(void)
 
 void init_merge_sort_functions(py::module_ m)
 {
-    dpctl::tensor::py_internal::init_merge_sort_dispatch_vectors();
+    init_merge_sort_dispatch_vectors();
 
     auto py_sort_ascending = [](const dpctl::tensor::usm_ndarray &src,
                                 const int trailing_dims_to_sort,
                                 const dpctl::tensor::usm_ndarray &dst,
                                 sycl::queue &exec_q,
                                 const std::vector<sycl::event> &depends)
         -> std::pair<sycl::event, sycl::event> {
-        return dpctl::tensor::py_internal::py_sort(
-            src, trailing_dims_to_sort, dst, exec_q, depends,
-            dpctl::tensor::py_internal::ascending_sort_contig_dispatch_vector);
+        return py_sort(src, trailing_dims_to_sort, dst, exec_q, depends,
+                       ascending_sort_contig_dispatch_vector);
     };
     m.def("_sort_ascending", py_sort_ascending, py::arg("src"),
           py::arg("trailing_dims_to_sort"), py::arg("dst"),
@@ -124,9 +126,8 @@ void init_merge_sort_functions(py::module_ m)
                                  sycl::queue &exec_q,
                                  const std::vector<sycl::event> &depends)
         -> std::pair<sycl::event, sycl::event> {
-        return dpctl::tensor::py_internal::py_sort(
-            src, trailing_dims_to_sort, dst, exec_q, depends,
-            dpctl::tensor::py_internal::descending_sort_contig_dispatch_vector);
+        return py_sort(src, trailing_dims_to_sort, dst, exec_q, depends,
+                       descending_sort_contig_dispatch_vector);
     };
     m.def("_sort_descending", py_sort_descending, py::arg("src"),
           py::arg("trailing_dims_to_sort"), py::arg("dst"),