Add license headers

Add common typedef for sort_contig_fn_ptr_t from merge_sort header
and radix_sort header files into a new file.
Used it in cpp files.

Author: oleksandr-pavlyk

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

@@ -703,17 +703,6 @@ merge_sorted_block_contig_impl(sycl::queue &q,
 
 } // end of namespace merge_sort_detail
 
-typedef sycl::event (*sort_contig_fn_ptr_t)(sycl::queue &,
-                                            size_t,
-                                            size_t,
-                                            const char *,
-                                            char *,
-                                            ssize_t,
-                                            ssize_t,
-                                            ssize_t,
-                                            ssize_t,
-                                            const std::vector<sycl::event> &);
-
 template <typename argTy, typename Comp = std::less<argTy>>
 sycl::event stable_sort_axis1_contig_impl(
     sycl::queue &exec_q,

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

@@ -1,3 +1,30 @@
+//
+//                      Data Parallel Control (dpctl)
+//
+// Copyright 2020-2024 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+//===--------------------------------------------------------------------===//
+///
+/// \file
+/// This file defines functions of dpctl.tensor._tensor_sorting_impl
+/// extension.
+//===--------------------------------------------------------------------===//
+
+// Implementation in this file were adapted from oneDPL's radix sort
+// implementation, license Apache-2.0 WITH LLVM-exception
+
 #pragma once
 
 #include <cstdint>
@@ -217,7 +244,7 @@ radix_sort_count_submit(sycl::queue &exec_q,
 
     // iteration space info
     const std::size_t n = n_values;
-    // Each segment is processed by a work-group
+    // each segment is processed by a work-group
     const std::size_t elems_per_segment = (n + n_segments - 1) / n_segments;
     const std::size_t no_op_flag_id = n_counts - 1;
 
@@ -372,8 +399,7 @@ sycl::event radix_sort_scan_submit(sycl::queue &exec_q,
             const auto lid = ndit.get_local_linear_id();
 
             // NB: No race condition here, because the condition may ever be
-            // true
-            //     for only on one WG, one WI.
+            // true for only on one WG, one WI.
             if ((lid == wg_size - 1) && (begin_ptr[scan_size - 1] == n_values))
             {
                 // set flag, since all the values got into one
@@ -613,7 +639,6 @@ radix_sort_reorder_submit(sycl::queue &exec_q,
                           const ProjT &proj_op,
                           const std::vector<sycl::event> dependency_events)
 {
-    // typedefs
     using ValueT = InputT;
     using PeerHelper = peer_prefix_helper<OffsetT, PeerAlgo>;
 
@@ -649,7 +674,6 @@ radix_sort_reorder_submit(sycl::queue &exec_q,
 
         sycl::nd_range<1> ndRange{gRange, lRange};
 
-        // Each work-group processes one segment ?
         cgh.parallel_for<KernelName>(ndRange, [=](sycl::nd_item<1> ndit) {
             const std::size_t group_id = ndit.get_group(0);
             const std::size_t iter_id = group_id / n_segments;
@@ -670,9 +694,9 @@ radix_sort_reorder_submit(sycl::queue &exec_q,
                 return;
             }
 
-            // 1. create a private array for storing offset values
-            //    and add total offset and offset for compute unit for a certain
-            //    radix state
+            // create a private array for storing offset values
+            // and add total offset and offset for compute unit
+            // for a certain radix state
             std::array<OffsetT, radix_states> offset_arr{};
             const std::size_t scan_size = n_segments + 1;
 
@@ -688,7 +712,7 @@ radix_sort_reorder_submit(sycl::queue &exec_q,
                 const std::uint32_t local_offset_id =
                     segment_id + scan_size * radix_state_id;
 
-                // scan bins (serial)
+                // scan bins serially
                 const std::size_t last_segment_bucket_id =
                     radix_state_id * scan_size - 1;
                 scanned_bin += b_offset_ptr[last_segment_bucket_id];
@@ -739,7 +763,7 @@ radix_sort_reorder_submit(sycl::queue &exec_q,
             if (tail_size > 0) {
                 ValueT in_val;
 
-                // greater than any actual radix state
+                // default: is greater than any actual radix state
                 std::uint32_t bucket_id = radix_states;
                 if (lid < tail_size) {
                     in_val = std::move(b_input_ptr[seg_end + lid]);
@@ -749,6 +773,7 @@ radix_sort_reorder_submit(sycl::queue &exec_q,
                     bucket_id =
                         get_bucket_id<radix_mask>(mapped_val, radix_offset);
                 }
+
                 OffsetT new_offset_id = 0;
                 for (std::uint32_t radix_state_id = 0;
                      radix_state_id < radix_states; ++radix_state_id)
@@ -761,6 +786,7 @@ radix_sort_reorder_submit(sycl::queue &exec_q,
 
                     offset_arr[radix_state_id] += sg_total_offset;
                 }
+
                 if (lid < tail_size) {
                     b_output_ptr[new_offset_id] = std::move(in_val);
                 }
@@ -1589,19 +1615,6 @@ template <typename IndexT, typename ValueT, typename ProjT> struct IndexedProj
 
 } // end of namespace radix_sort_details
 
-// same signature as sort_contig_fn_ptr_t
-typedef sycl::event (*radix_sort_contig_fn_ptr_t)(
-    sycl::queue &,
-    size_t,
-    size_t,
-    const char *,
-    char *,
-    ssize_t,
-    ssize_t,
-    ssize_t,
-    ssize_t,
-    const std::vector<sycl::event> &);
-
 template <typename argTy, bool sort_ascending>
 sycl::event
 radix_sort_axis1_contig_impl(sycl::queue &exec_q,

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

@@ -0,0 +1,50 @@
+//
+//                      Data Parallel Control (dpctl)
+//
+// Copyright 2020-2024 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+//===--------------------------------------------------------------------===//
+///
+/// \file
+/// This file defines functions of dpctl.tensor._tensor_sorting_impl
+/// extension.
+//===--------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/sycl.hpp>
+#include <vector>
+
+namespace dpctl
+{
+namespace tensor
+{
+namespace kernels
+{
+
+typedef sycl::event (*sort_contig_fn_ptr_t)(sycl::queue &,
+                                            size_t,
+                                            size_t,
+                                            const char *,
+                                            char *,
+                                            ssize_t,
+                                            ssize_t,
+                                            ssize_t,
+                                            ssize_t,
+                                            const std::vector<sycl::event> &);
+
+}
+} // namespace tensor
+} // namespace dpctl

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

@@ -33,6 +33,7 @@
 #include "utils/type_dispatch.hpp"
 
 #include "kernels/sorting/merge_sort.hpp"
+#include "kernels/sorting/sort_impl_fn_ptr_t.hpp"
 #include "rich_comparisons.hpp"
 
 #include "argsort.hpp"

dpctl/tensor/libtensor/source/sorting/py_argsort_common.hpp

@@ -32,10 +32,6 @@
 #include "utils/output_validation.hpp"
 #include "utils/type_dispatch.hpp"
 
-#include "argsort.hpp"
-#include "kernels/sorting/merge_sort.hpp"
-#include "rich_comparisons.hpp"
-
 namespace td_ns = dpctl::tensor::type_dispatch;
 
 namespace dpctl

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

@@ -1,3 +1,27 @@
+//
+//                      Data Parallel Control (dpctl)
+//
+// Copyright 2020-2024 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+//===--------------------------------------------------------------------===//
+///
+/// \file
+/// This file defines functions of dpctl.tensor._tensor_sorting_impl
+/// extension.
+//===--------------------------------------------------------------------===//
+
 #include <cstdint>
 #include <utility>
 #include <vector>
@@ -15,10 +39,11 @@
 #include "utils/type_dispatch.hpp"
 
 #include "kernels/sorting/radix_sort.hpp"
-#include "radix_sort_support.hpp"
+#include "kernels/sorting/sort_impl_fn_ptr_t.hpp"
 
 #include "py_argsort_common.hpp"
 #include "radix_argsort.hpp"
+#include "radix_sort_support.hpp"
 
 namespace dpctl
 {
@@ -30,12 +55,12 @@ namespace py_internal
 namespace td_ns = dpctl::tensor::type_dispatch;
 namespace impl_ns = dpctl::tensor::kernels::radix_sort_details;
 
-using dpctl::tensor::kernels::radix_sort_contig_fn_ptr_t;
+using dpctl::tensor::kernels::sort_contig_fn_ptr_t;
 
-static radix_sort_contig_fn_ptr_t
+static sort_contig_fn_ptr_t
     ascending_radix_argsort_contig_dispatch_table[td_ns::num_types]
                                                  [td_ns::num_types];
-static radix_sort_contig_fn_ptr_t
+static sort_contig_fn_ptr_t
     descending_radix_argsort_contig_dispatch_table[td_ns::num_types]
                                                   [td_ns::num_types];
 
@@ -79,15 +104,15 @@ struct DescendingRadixArgSortContigFactory
 
 void init_radix_argsort_dispatch_tables(void)
 {
-    using dpctl::tensor::kernels::radix_sort_contig_fn_ptr_t;
+    using dpctl::tensor::kernels::sort_contig_fn_ptr_t;
 
-    td_ns::DispatchTableBuilder<radix_sort_contig_fn_ptr_t,
+    td_ns::DispatchTableBuilder<sort_contig_fn_ptr_t,
                                 AscendingRadixArgSortContigFactory,
                                 td_ns::num_types>
         dtb1;
     dtb1.populate_dispatch_table(ascending_radix_argsort_contig_dispatch_table);
 
-    td_ns::DispatchTableBuilder<radix_sort_contig_fn_ptr_t,
+    td_ns::DispatchTableBuilder<sort_contig_fn_ptr_t,
                                 DescendingRadixArgSortContigFactory,
                                 td_ns::num_types>
         dtb2;

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

@@ -1,3 +1,27 @@
+//
+//                      Data Parallel Control (dpctl)
+//
+// Copyright 2020-2024 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+//===--------------------------------------------------------------------===//
+///
+/// \file
+/// This file defines functions of dpctl.tensor._tensor_sorting_impl
+/// extension.
+//===--------------------------------------------------------------------===//
+
 #include <cstdint>
 #include <utility>
 #include <vector>
@@ -15,10 +39,11 @@
 #include "utils/type_dispatch.hpp"
 
 #include "kernels/sorting/radix_sort.hpp"
-#include "radix_sort_support.hpp"
+#include "kernels/sorting/sort_impl_fn_ptr_t.hpp"
 
 #include "py_sort_common.hpp"
 #include "radix_sort.hpp"
+#include "radix_sort_support.hpp"
 
 namespace dpctl
 {
@@ -30,10 +55,10 @@ namespace py_internal
 namespace td_ns = dpctl::tensor::type_dispatch;
 namespace impl_ns = dpctl::tensor::kernels::radix_sort_details;
 
-using dpctl::tensor::kernels::radix_sort_contig_fn_ptr_t;
-static radix_sort_contig_fn_ptr_t
+using dpctl::tensor::kernels::sort_contig_fn_ptr_t;
+static sort_contig_fn_ptr_t
     ascending_radix_sort_contig_dispatch_vector[td_ns::num_types];
-static radix_sort_contig_fn_ptr_t
+static sort_contig_fn_ptr_t
     descending_radix_sort_contig_dispatch_vector[td_ns::num_types];
 
 template <typename fnT, typename argTy> struct AscendingRadixSortContigFactory
@@ -66,15 +91,14 @@ template <typename fnT, typename argTy> struct DescendingRadixSortContigFactory
 
 void init_radix_sort_dispatch_vectors(void)
 {
-    using dpctl::tensor::kernels::radix_sort_contig_fn_ptr_t;
+    using dpctl::tensor::kernels::sort_contig_fn_ptr_t;
 
-    td_ns::DispatchVectorBuilder<radix_sort_contig_fn_ptr_t,
-                                 AscendingRadixSortContigFactory,
-                                 td_ns::num_types>
+    td_ns::DispatchVectorBuilder<
+        sort_contig_fn_ptr_t, AscendingRadixSortContigFactory, td_ns::num_types>
         dtv1;
     dtv1.populate_dispatch_vector(ascending_radix_sort_contig_dispatch_vector);
 
-    td_ns::DispatchVectorBuilder<radix_sort_contig_fn_ptr_t,
+    td_ns::DispatchVectorBuilder<sort_contig_fn_ptr_t,
                                  DescendingRadixSortContigFactory,
                                  td_ns::num_types>
         dtv2;

dpctl/tensor/libtensor/source/sorting/radix_sort_support.hpp

@@ -1,3 +1,27 @@
+//
+//                      Data Parallel Control (dpctl)
+//
+// Copyright 2020-2024 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+//===--------------------------------------------------------------------===//
+///
+/// \file
+/// This file defines functions of dpctl.tensor._tensor_sorting_impl
+/// extension.
+//===--------------------------------------------------------------------===//
+
 #pragma once
 
 #include <type_traits>