[SYCL] Cleanup SYCL sorters implementation (#13962)

againull · web-flow · commit 673ce2ce742c · 2024-05-30T06:26:13.000-07:00
diff --git a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
@@ -51,46 +51,43 @@ template <typename Group, size_t Extent> class group_with_scratchpad {
 // ---- sorters
 template <typename Compare = std::less<>> class default_sorter {
   Compare comp;
-  std::byte *scratch;
-  size_t scratch_size;
+  sycl::span<std::byte> scratch;
 
 public:
   template <size_t Extent>
   default_sorter(sycl::span<std::byte, Extent> scratch_,
                  Compare comp_ = Compare())
-      : comp(comp_), scratch(scratch_.data()), scratch_size(scratch_.size()) {}
+      : comp(comp_), scratch(scratch_) {}
 
   template <typename Group, typename Ptr>
-  void operator()(Group g, Ptr first, Ptr last) {
+  void operator()([[maybe_unused]] Group g, [[maybe_unused]] Ptr first,
+                  [[maybe_unused]] Ptr last) {
 #ifdef __SYCL_DEVICE_ONLY__
-    using T = typename sycl::detail::GetValueType<Ptr>::type;
-    if (scratch_size >= memory_required<T>(Group::fence_scope, last - first))
-      sycl::detail::merge_sort(g, first, last - first, comp, scratch);
-      // TODO: it's better to add else branch
+    // Per extension specification if scratch size is less than the value
+    // returned by memory_required then behavior is undefined, so we don't check
+    // that the scratch size statisfies the requirement.
+    sycl::detail::merge_sort(g, first, last - first, comp, scratch.data());
 #else
-    (void)g;
-    (void)first;
-    (void)last;
     throw sycl::exception(
         std::error_code(PI_ERROR_INVALID_DEVICE, sycl::sycl_category()),
         "default_sorter constructor is not supported on host device.");
 #endif
   }
 
-  template <typename Group, typename T> T operator()(Group g, T val) {
+  template <typename Group, typename T>
+  T operator()([[maybe_unused]] Group g, T val) {
 #ifdef __SYCL_DEVICE_ONLY__
+    // Per extension specification if scratch size is less than the value
+    // returned by memory_required then behavior is undefined, so we don't check
+    // that the scratch size statisfies the requirement.
     auto range_size = g.get_local_range().size();
-    if (scratch_size >= memory_required<T>(Group::fence_scope, range_size)) {
-      size_t local_id = g.get_local_linear_id();
-      T *temp = reinterpret_cast<T *>(scratch);
-      ::new (temp + local_id) T(val);
-      sycl::detail::merge_sort(g, temp, range_size, comp,
-                               scratch + range_size * sizeof(T));
-      val = temp[local_id];
-    }
-    // TODO: it's better to add else branch
+    size_t local_id = g.get_local_linear_id();
+    T *temp = reinterpret_cast<T *>(scratch.data());
+    ::new (temp + local_id) T(val);
+    sycl::detail::merge_sort(g, temp, range_size, comp,
+                             scratch.data() + range_size * sizeof(T));
+    val = temp[local_id];
 #else
-    (void)g;
     throw sycl::exception(
         std::error_code(PI_ERROR_INVALID_DEVICE, sycl::sycl_category()),
         "default_sorter operator() is not supported on host device.");
@@ -129,62 +126,56 @@ template <typename ValT, sorting_order OrderT = sorting_order::ascending,
           unsigned int BitsPerPass = 4>
 class radix_sorter {
 
-  std::byte *scratch = nullptr;
+  sycl::span<std::byte> scratch;
   uint32_t first_bit = 0;
   uint32_t last_bit = 0;
-  size_t scratch_size = 0;
 
   static constexpr uint32_t bits = BitsPerPass;
+  using bitset_t = std::bitset<sizeof(ValT) * CHAR_BIT>;
 
 public:
   template <size_t Extent>
   radix_sorter(sycl::span<std::byte, Extent> scratch_,
-               const std::bitset<sizeof(ValT) *CHAR_BIT> mask =
-                   std::bitset<sizeof(ValT) * CHAR_BIT>(
-                       (std::numeric_limits<unsigned long long>::max)()))
-      : scratch(scratch_.data()), scratch_size(scratch_.size()) {
+               const bitset_t mask = bitset_t{}.set())
+      : scratch(scratch_) {
     static_assert((std::is_arithmetic<ValT>::value ||
                    std::is_same<ValT, sycl::half>::value ||
                    std::is_same<ValT, sycl::ext::oneapi::bfloat16>::value),
                   "radix sort is not usable");
 
-    first_bit = 0;
-    while (first_bit < mask.size() && !mask[first_bit])
-      ++first_bit;
-
-    last_bit = first_bit;
-    while (last_bit < mask.size() && mask[last_bit])
-      ++last_bit;
+    for (first_bit = 0; first_bit < mask.size() && !mask[first_bit];
+         ++first_bit)
+      ;
+    for (last_bit = first_bit; last_bit < mask.size() && mask[last_bit];
+         ++last_bit)
+      ;
   }
 
   template <typename GroupT, typename PtrT>
-  void operator()(GroupT g, PtrT first, PtrT last) {
-    (void)g;
-    (void)first;
-    (void)last;
+  void operator()([[maybe_unused]] GroupT g, [[maybe_unused]] PtrT first,
+                  [[maybe_unused]] PtrT last) {
 #ifdef __SYCL_DEVICE_ONLY__
     sycl::detail::privateDynamicSort</*is_key_value=*/false,
                                      OrderT == sorting_order::ascending,
                                      /*empty*/ 1, BitsPerPass>(
-        g, first, /*empty*/ first, (last - first) > 0 ? (last - first) : 0,
-        scratch, first_bit, last_bit);
+        g, first, /*empty*/ first, last - first, scratch.data(), first_bit,
+        last_bit);
 #else
     throw sycl::exception(
         std::error_code(PI_ERROR_INVALID_DEVICE, sycl::sycl_category()),
         "radix_sorter is not supported on host device.");
 #endif
   }
 
-  template <typename GroupT> ValT operator()(GroupT g, ValT val) {
-    (void)g;
-    (void)val;
+  template <typename GroupT>
+  ValT operator()([[maybe_unused]] GroupT g, [[maybe_unused]] ValT val) {
 #ifdef __SYCL_DEVICE_ONLY__
     ValT result[]{val};
     sycl::detail::privateStaticSort</*is_key_value=*/false,
                                     /*is_blocked=*/true,
                                     OrderT == sorting_order::ascending,
                                     /*items_per_work_item=*/1, bits>(
-        g, result, /*empty*/ result, scratch, first_bit, last_bit);
+        g, result, /*empty*/ result, scratch.data(), first_bit, last_bit);
     return result[0];
 #else
     throw sycl::exception(
@@ -193,20 +184,16 @@ class radix_sorter {
 #endif
   }
 
-  static constexpr size_t memory_required(sycl::memory_scope scope,
+  static constexpr size_t memory_required(sycl::memory_scope,
                                           size_t range_size) {
-    // Scope is not important so far
-    (void)scope;
     return range_size * sizeof(ValT) +
            (1 << bits) * range_size * sizeof(uint32_t) + alignof(uint32_t);
   }
 
   // memory_helpers
   template <int dimensions = 1>
-  static constexpr size_t memory_required(sycl::memory_scope scope,
+  static constexpr size_t memory_required(sycl::memory_scope,
                                           sycl::range<dimensions> local_range) {
-    // Scope is not important so far
-    (void)scope;
     return (std::max)(local_range.size() * sizeof(ValT),
                       local_range.size() * (1 << bits) * sizeof(uint32_t));
   }
diff --git a/sycl/include/sycl/ext/oneapi/experimental/group_sort.hpp b/sycl/include/sycl/ext/oneapi/experimental/group_sort.hpp
@@ -73,13 +73,11 @@ struct is_sorter : decltype(is_sorter_impl<Sorter, Group, ValOrPtr>::test(0)) {
 // ---- sort_over_group
 template <typename Group, typename T, typename Sorter>
 std::enable_if_t<detail::is_sorter<Sorter, Group, T>::value, T>
-sort_over_group(Group group, T value, Sorter sorter) {
+sort_over_group([[maybe_unused]] Group group, [[maybe_unused]] T value,
+                [[maybe_unused]] Sorter sorter) {
 #ifdef __SYCL_DEVICE_ONLY__
   return sorter(group, value);
 #else
-  (void)group;
-  (void)value;
-  (void)sorter;
   throw sycl::exception(
       std::error_code(PI_ERROR_INVALID_DEVICE, sycl::sycl_category()),
       "Group algorithms are not supported on host device.");
@@ -106,14 +104,11 @@ sort_over_group(experimental::group_with_scratchpad<Group, Extent> exec,
 // ---- joint_sort
 template <typename Group, typename Iter, typename Sorter>
 std::enable_if_t<detail::is_sorter<Sorter, Group, Iter>::value, void>
-joint_sort(Group group, Iter first, Iter last, Sorter sorter) {
+joint_sort([[maybe_unused]] Group group, [[maybe_unused]] Iter first,
+           [[maybe_unused]] Iter last, [[maybe_unused]] Sorter sorter) {
 #ifdef __SYCL_DEVICE_ONLY__
   sorter(group, first, last);
 #else
-  (void)group;
-  (void)first;
-  (void)last;
-  (void)sorter;
   throw sycl::exception(
       std::error_code(PI_ERROR_INVALID_DEVICE, sycl::sycl_category()),
       "Group algorithms are not supported on host device.");
diff --git a/sycl/test-e2e/GroupAlgorithm/SYCL2020/sort.cpp b/sycl/test-e2e/GroupAlgorithm/SYCL2020/sort.cpp
@@ -187,6 +187,9 @@ void RunJointSort(sycl::queue &Q, const std::vector<T> &DataToSort,
              const size_t EndIdx =
                  std::min(ChunkSize * (PartID + 1), NumOfElements);
 
+             if (EndIdx <= StartIdx)
+               return;
+
              // This version of API always sorts in ascending order
              if constexpr (std::is_same_v<Compare, std::less<T>>)
                oneapi_exp::joint_sort(
