made changes for multi-dimensional groups

Signed-off-by: Fedorov, Andrey <[email protected]>

Author: andreyfe1

sycl/include/CL/sycl/detail/group_sort_impl.hpp

@@ -67,8 +67,8 @@ struct GetValueType<sycl::multi_ptr<ElementType, Space>> {
   using type = ElementType;
 };
 
-// since we couldn't assign data to raw memory, it's better to use placement for
-// first assignment
+// since we couldn't assign data to raw memory, it's better to use placement
+// for first assignment
 template <typename Acc, typename T>
 void set_value(Acc ptr, const std::size_t idx, const T &val, bool is_first) {
   if (is_first) {
@@ -97,8 +97,8 @@ void merge(const std::size_t offset, InAcc &in_acc1, OutAcc &out_acc1,
   const std::size_t local_size_1 = local_end_1 - local_start_1;
   const std::size_t local_size_2 = local_end_2 - local_start_2;
 
-  // TODO: process cases where all elements of 1st sequence > 2nd, 2nd > 1st to
-  // improve performance
+  // TODO: process cases where all elements of 1st sequence > 2nd, 2nd > 1st
+  // to improve performance
 
   // Process 1st sequence
   if (local_start_1 < local_end_1) {
@@ -204,7 +204,7 @@ void merge_sort(Group group, Iter first, const std::size_t n, Compare comp,
                 std::byte *scratch) {
   using T = typename GetValueType<Iter>::type;
   auto id = sycl::detail::Builder::getNDItem<Group::dimensions>();
-  const std::size_t idx = id.get_local_id();
+  const std::size_t idx = id.get_local_linear_id();
   const std::size_t local = group.get_local_range().size();
   const std::size_t chunk = (n - 1) / local + 1;
 

sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp

@@ -63,7 +63,7 @@ template <typename Compare = std::less<>> class default_sorter {
     auto range_size = g.get_local_range().size();
     if (scratch_size >= memory_required<T>(Group::fence_scope, range_size)) {
       auto id = sycl::detail::Builder::getNDItem<Group::dimensions>();
-      uint32_t local_id = id.get_local_id();
+      std::size_t local_id = id.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,