Renamed some files in sorting folders, in preparation for introduction of radix sort

Author: oleksandr-pavlyk

dpctl/tensor/libtensor/include/kernels/sorting/sort.hpp renamed to dpctl/tensor/libtensor/include/kernels/sorting/merge_sort.hpp

@@ -32,7 +32,7 @@
 #include <vector>
 
 #include "kernels/dpctl_tensor_types.hpp"
-#include "kernels/sorting/sort_detail.hpp"
+#include "kernels/sorting/search_sorted_detail.hpp"
 
 namespace dpctl
 {
@@ -41,9 +41,11 @@ namespace tensor
 namespace kernels
 {
 
-namespace sort_detail
+namespace merge_sort_detail
 {
 
+using namespace dpctl::tensor::kernels::search_sorted_detail;
+
 /*! @brief Merge two contiguous sorted segments */
 template <typename InAcc, typename OutAcc, typename Compare>
 void merge_impl(const std::size_t offset,
@@ -699,7 +701,7 @@ merge_sorted_block_contig_impl(sycl::queue &q,
     return dep_ev;
 }
 
-} // end of namespace sort_detail
+} // end of namespace merge_sort_detail
 
 typedef sycl::event (*sort_contig_fn_ptr_t)(sycl::queue &,
                                             size_t,
@@ -741,8 +743,8 @@ sycl::event stable_sort_axis1_contig_impl(
     if (sort_nelems < sequential_sorting_threshold) {
         // equal work-item sorts entire row
         sycl::event sequential_sorting_ev =
-            sort_detail::sort_base_step_contig_impl<const argTy *, argTy *,
-                                                    Comp>(
+            merge_sort_detail::sort_base_step_contig_impl<const argTy *,
+                                                          argTy *, Comp>(
                 exec_q, iter_nelems, sort_nelems, arg_tp, res_tp, comp,
                 sort_nelems, depends);
 
@@ -753,16 +755,16 @@ sycl::event stable_sort_axis1_contig_impl(
 
         // Sort segments of the array
         sycl::event base_sort_ev =
-            sort_detail::sort_over_work_group_contig_impl<const argTy *,
-                                                          argTy *, Comp>(
+            merge_sort_detail::sort_over_work_group_contig_impl<const argTy *,
+                                                                argTy *, Comp>(
                 exec_q, iter_nelems, sort_nelems, arg_tp, res_tp, comp,
                 sorted_block_size, // modified in place with size of sorted
                                    // block size
                 depends);
 
         // Merge segments in parallel until all elements are sorted
         sycl::event merges_ev =
-            sort_detail::merge_sorted_block_contig_impl<argTy *, Comp>(
+            merge_sort_detail::merge_sorted_block_contig_impl<argTy *, Comp>(
                 exec_q, iter_nelems, sort_nelems, res_tp, comp,
                 sorted_block_size, {base_sort_ev});
 
@@ -837,21 +839,24 @@ sycl::event stable_argsort_axis1_contig_impl(
         });
 
     // Sort segments of the array
-    sycl::event base_sort_ev = sort_detail::sort_over_work_group_contig_impl(
-        exec_q, iter_nelems, sort_nelems, res_tp, res_tp, index_comp,
-        sorted_block_size, // modified in place with size of sorted block size
-        {populate_indexed_data_ev});
+    sycl::event base_sort_ev =
+        merge_sort_detail::sort_over_work_group_contig_impl(
+            exec_q, iter_nelems, sort_nelems, res_tp, res_tp, index_comp,
+            sorted_block_size, // modified in place with size of sorted block
+                               // size
+            {populate_indexed_data_ev});
 
     // Merge segments in parallel until all elements are sorted
-    sycl::event merges_ev = sort_detail::merge_sorted_block_contig_impl(
+    sycl::event merges_ev = merge_sort_detail::merge_sorted_block_contig_impl(
         exec_q, iter_nelems, sort_nelems, res_tp, index_comp, sorted_block_size,
         {base_sort_ev});
 
     sycl::event write_out_ev = exec_q.submit([&](sycl::handler &cgh) {
         cgh.depends_on(merges_ev);
 
         auto temp_acc =
-            sort_detail::GetReadOnlyAccess<decltype(res_tp)>{}(res_tp, cgh);
+            merge_sort_detail::GetReadOnlyAccess<decltype(res_tp)>{}(res_tp,
+                                                                     cgh);
 
         using KernelName = index_map_to_rows_krn<argTy, IndexTy, ValueComp>;
 

dpctl/tensor/libtensor/include/kernels/sorting/sort_detail.hpp renamed to dpctl/tensor/libtensor/include/kernels/sorting/search_sorted_detail.hpp

@@ -35,7 +35,7 @@ namespace tensor
 namespace kernels
 {
 
-namespace sort_detail
+namespace search_sorted_detail
 {
 
 template <typename T> T quotient_ceil(T n, T m) { return (n + m - 1) / m; }
@@ -111,7 +111,7 @@ std::size_t upper_bound_indexed_impl(const Acc acc,
                                     acc_indexer);
 }
 
-} // namespace sort_detail
+} // namespace search_sorted_detail
 
 } // namespace kernels
 } // namespace tensor

dpctl/tensor/libtensor/include/kernels/sorting/searchsorted.hpp

@@ -31,7 +31,7 @@
 #include <vector>
 
 #include "kernels/dpctl_tensor_types.hpp"
-#include "kernels/sorting/sort_detail.hpp"
+#include "kernels/sorting/search_sorted_detail.hpp"
 #include "utils/offset_utils.hpp"
 
 namespace dpctl
@@ -91,7 +91,7 @@ struct SearchSortedFunctor
 
             // lower_bound returns the first pos such that bool(hay[pos] <
             // needle_v) is false, i.e. needle_v <= hay[pos]
-            pos = sort_detail::lower_bound_indexed_impl(
+            pos = search_sorted_detail::lower_bound_indexed_impl(
                 hay_tp, zero, hay_nelems, needle_v, comp, hay_indexer);
         }
         else {
@@ -100,7 +100,7 @@ struct SearchSortedFunctor
 
             // upper_bound returns the first pos such that bool(needle_v <
             // hay[pos]) is true, i.e. needle_v < hay[pos]
-            pos = sort_detail::upper_bound_indexed_impl(
+            pos = search_sorted_detail::upper_bound_indexed_impl(
                 hay_tp, zero, hay_nelems, needle_v, comp, hay_indexer);
         }
 

dpctl/tensor/libtensor/source/sorting/argsort.cpp

@@ -33,7 +33,7 @@
 #include "utils/type_dispatch.hpp"
 
 #include "argsort.hpp"
-#include "kernels/sorting/sort.hpp"
+#include "kernels/sorting/merge_sort.hpp"
 #include "rich_comparisons.hpp"
 
 namespace td_ns = dpctl::tensor::type_dispatch;

dpctl/tensor/libtensor/source/sorting/sort.cpp

@@ -33,7 +33,7 @@
 #include "utils/output_validation.hpp"
 #include "utils/type_dispatch.hpp"
 
-#include "kernels/sorting/sort.hpp"
+#include "kernels/sorting/merge_sort.hpp"
 #include "rich_comparisons.hpp"
 #include "sort.hpp"