diff --git a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc
index b050bb3a372e0..fe4a76abe62d1 100755
--- a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc
+++ b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc
@@ -146,7 +146,7 @@ namespace sycl::ext::oneapi::experimental {
   class default_sorter {
   public:
     template<std::size_t Extent>
-    default_sorter(sycl::span<uint8_t, Extent> scratch, Compare comp = Compare());
+    default_sorter(sycl::span<std::byte, Extent> scratch, Compare comp = Compare());
 
     template<typename Group, typename Ptr>
     void operator()(Group g, Ptr first, Ptr last);
@@ -167,7 +167,7 @@ namespace sycl::ext::oneapi::experimental {
   class radix_sorter {
   public:
     template<std::size_t Extent>
-    radix_sorter(sycl::span<uint8_t, Extent> scratch,
+    radix_sorter(sycl::span<std::byte, Extent> scratch,
                  const std::bitset<sizeof(T) * CHAR_BIT> mask =
                      std::bitset<sizeof(T) * CHAR_BIT> (std::numeric_limits<unsigned long long>::max()));
 
@@ -215,7 +215,7 @@ Table 4. Constructors of the `default_sorter` class.
 |Constructor|Description
 
 |`template<std::size_t Extent>
-default_sorter(sycl::span<uint8_t, Extent> scratch, Compare comp = Compare())`
+default_sorter(sycl::span<std::byte, Extent> scratch, Compare comp = Compare())`
 |Creates the `default_sorter` object using `comp`.
 Additional memory for the algorithm is provided using `scratch`.
 If `scratch.size()` is less than the value returned by
@@ -264,7 +264,7 @@ Table 6. Constructors of the `radix_sorter` class.
 |Constructor|Description
 
 |`template<std::size_t Extent>
-radix_sorter(sycl::span<uint8_t, Extent> scratch, const std::bitset<sizeof(T) * CHAR_BIT> mask = std::bitset<sizeof(T) * CHAR_BIT>
+radix_sorter(sycl::span<std::byte, Extent> scratch, const std::bitset<sizeof(T) * CHAR_BIT> mask = std::bitset<sizeof(T) * CHAR_BIT>
 (std::numeric_limits<unsigned long long>::max()))`
 |Creates the `radix_sorter` object to sort values considering only bits
 that corresponds to 1 in `mask`.
@@ -350,16 +350,16 @@ namespace sycl::ext::oneapi::experimental {
   class group_with_scratchpad
   {
   public:
-    group_with_scratchpad(Group group, sycl::span<uint8_t, Extent> scratch);
+    group_with_scratchpad(Group group, sycl::span<std::byte, Extent> scratch);
     Group get_group() const;
 
-    sycl::span<uint8_t, Extent>
+    sycl::span<std::byte, Extent>
     get_memory() const;
   };
 
   // Deduction guides
   template<typename Group, std::size_t Extent>
-  group_with_scratchpad(Group, sycl::span<uint8_t, Extent>)
+  group_with_scratchpad(Group, sycl::span<std::byte, Extent>)
     -> group_with_scratchpad<Group, Extent>;
 
 }
@@ -372,7 +372,7 @@ Table 9. Constructors of the `group_with_scratchpad` class.
 |===
 |Constructor|Description
 
-|`group_with_scratchpad(Group group, sycl::span<uint8_t, Extent> scratch)`
+|`group_with_scratchpad(Group group, sycl::span<std::byte, Extent> scratch)`
 |Creates the `group_with_scratchpad` object using `group` and `scratch`.
 `sycl::is_group_v<std::decay_t<Group>>` must be true.
 `scratch.size()` must not be less than value returned by the `memory_required` method
@@ -388,7 +388,7 @@ Table 10. Member functions of the `group_with_scratchpad` class.
 |`Group get_group() const`
 |Returns the `Group` class object that is handled by the `group_with_scratchpad` object.
 
-|`sycl::span<uint8_t, Extent>
+|`sycl::span<std::byte, Extent>
 get_memory() const`
 |Returns `sycl::span` that represents an additional memory
 that is handled by the `group_with_scratchpad` object.
@@ -508,7 +508,7 @@ size_t temp_memory_size =
 
 q.submit([&](sycl::handler& h) {
   auto acc = sycl::accessor(buf, h);
-  auto scratch = sycl::local_accessor<uint8_t, 1>( {temp_memory_size}, h );
+  auto scratch = sycl::local_accessor<std::byte, 1>( {temp_memory_size}, h );
 
   h.parallel_for(
     sycl::nd_range<1>{ /*global_size = */ {256}, /*local_size = */ {256} },
@@ -546,7 +546,7 @@ size_t temp_memory_size =
 
 q.submit([&](sycl::handler& h) {
   auto acc = sycl::accessor(buf, h);
-  auto scratch = sycl::local_accessor<uint8_t, 1>( {temp_memory_size}, h);
+  auto scratch = sycl::local_accessor<std::byte, 1>( {temp_memory_size}, h);
 
   h.parallel_for(
     sycl::nd_range<1>{ local_range, local_range },
@@ -583,7 +583,7 @@ size_t temp_memory_size =
 q.submit([&](sycl::handler& h) {
   auto keys_acc = sycl::accessor(keys_buf, h);
   auto vals_acc = sycl::accessor(vals_buf, h);
-  auto scratch = sycl::local_accessor<uint8_t, 1>( {temp_memory_size}, h);
+  auto scratch = sycl::local_accessor<std::byte, 1>( {temp_memory_size}, h);
 
   h.parallel_for(
     sycl::nd_range<1>{ /*global_size = */ {1024}, /*local_size = */ {256} },
diff --git a/sycl/include/CL/sycl/detail/group_sort_impl.hpp b/sycl/include/CL/sycl/detail/group_sort_impl.hpp
new file mode 100644
index 0000000000000..e80f2fd131e77
--- /dev/null
+++ b/sycl/include/CL/sycl/detail/group_sort_impl.hpp
@@ -0,0 +1,256 @@
+//==------------ group_sort_impl.hpp ---------------------------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// This file includes some functions for group sorting algorithm implementations
+//
+
+#pragma once
+
+#if __cplusplus >= 201703L
+#include <CL/sycl/detail/helpers.hpp>
+
+#ifdef __SYCL_DEVICE_ONLY__
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace detail {
+
+// ---- merge sort implementation
+
+// following two functions could be useless if std::[lower|upper]_bound worked
+// well
+template <typename Acc, typename Value, typename Compare>
+std::size_t lower_bound(Acc acc, std::size_t first, std::size_t last,
+                        const Value &value, Compare comp) {
+  std::size_t n = last - first;
+  std::size_t cur = n;
+  std::size_t it;
+  while (n > 0) {
+    it = first;
+    cur = n / 2;
+    it += cur;
+    if (comp(acc[it], value)) {
+      n -= cur + 1, first = ++it;
+    } else
+      n = cur;
+  }
+  return first;
+}
+
+template <typename Acc, typename Value, typename Compare>
+std::size_t upper_bound(Acc acc, const std::size_t first,
+                        const std::size_t last, const Value &value,
+                        Compare comp) {
+  return detail::lower_bound(acc, first, last, value,
+                             [comp](auto x, auto y) { return !comp(y, x); });
+}
+
+// swap for all data types including tuple-like types
+template <typename T> void swap_tuples(T &a, T &b) { std::swap(a, b); }
+
+template <template <typename...> class TupleLike, typename T1, typename T2>
+void swap_tuples(TupleLike<T1, T2> &&a, TupleLike<T1, T2> &&b) {
+  std::swap(std::get<0>(a), std::get<0>(b));
+  std::swap(std::get<1>(a), std::get<1>(b));
+}
+
+template <typename Iter> struct GetValueType {
+  using type = typename std::iterator_traits<Iter>::value_type;
+};
+
+template <typename ElementType, access::address_space Space>
+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
+template <typename Acc, typename T>
+void set_value(Acc ptr, const std::size_t idx, const T &val, bool is_first) {
+  if (is_first) {
+    ::new (ptr + idx) T(val);
+  } else {
+    ptr[idx] = val;
+  }
+}
+
+template <typename InAcc, typename OutAcc, typename Compare>
+void merge(const std::size_t offset, InAcc &in_acc1, OutAcc &out_acc1,
+           const std::size_t start_1, const std::size_t end_1,
+           const std::size_t end_2, const std::size_t start_out, Compare comp,
+           const std::size_t chunk, bool is_first) {
+  const std::size_t start_2 = end_1;
+  // Borders of the sequences to merge within this call
+  const std::size_t local_start_1 =
+      sycl::min(static_cast<std::size_t>(offset + start_1), end_1);
+  const std::size_t local_end_1 =
+      sycl::min(static_cast<std::size_t>(local_start_1 + chunk), end_1);
+  const std::size_t local_start_2 =
+      sycl::min(static_cast<std::size_t>(offset + start_2), end_2);
+  const std::size_t local_end_2 =
+      sycl::min(static_cast<std::size_t>(local_start_2 + chunk), end_2);
+
+  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
+
+  // Process 1st sequence
+  if (local_start_1 < local_end_1) {
+    // Reduce the range for searching within the 2nd sequence and handle bound
+    // items find left border in 2nd sequence
+    const auto local_l_item_1 = in_acc1[local_start_1];
+    std::size_t l_search_bound_2 =
+        detail::lower_bound(in_acc1, start_2, end_2, local_l_item_1, comp);
+    const std::size_t l_shift_1 = local_start_1 - start_1;
+    const std::size_t l_shift_2 = l_search_bound_2 - start_2;
+
+    set_value(out_acc1, start_out + l_shift_1 + l_shift_2, local_l_item_1,
+              is_first);
+
+    std::size_t r_search_bound_2{};
+    // find right border in 2nd sequence
+    if (local_size_1 > 1) {
+      const auto local_r_item_1 = in_acc1[local_end_1 - 1];
+      r_search_bound_2 = detail::lower_bound(in_acc1, l_search_bound_2, end_2,
+                                             local_r_item_1, comp);
+      const auto r_shift_1 = local_end_1 - 1 - start_1;
+      const auto r_shift_2 = r_search_bound_2 - start_2;
+
+      set_value(out_acc1, start_out + r_shift_1 + r_shift_2, local_r_item_1,
+                is_first);
+    }
+
+    // Handle intermediate items
+    for (std::size_t idx = local_start_1 + 1; idx < local_end_1 - 1; ++idx) {
+      const auto intermediate_item_1 = in_acc1[idx];
+      // we shouldn't seek in whole 2nd sequence. Just for the part where the
+      // 1st sequence should be
+      l_search_bound_2 =
+          detail::lower_bound(in_acc1, l_search_bound_2, r_search_bound_2,
+                              intermediate_item_1, comp);
+      const std::size_t shift_1 = idx - start_1;
+      const std::size_t shift_2 = l_search_bound_2 - start_2;
+
+      set_value(out_acc1, start_out + shift_1 + shift_2, intermediate_item_1,
+                is_first);
+    }
+  }
+  // Process 2nd sequence
+  if (local_start_2 < local_end_2) {
+    // Reduce the range for searching within the 1st sequence and handle bound
+    // items find left border in 1st sequence
+    const auto local_l_item_2 = in_acc1[local_start_2];
+    std::size_t l_search_bound_1 =
+        detail::upper_bound(in_acc1, start_1, end_1, local_l_item_2, comp);
+    const std::size_t l_shift_1 = l_search_bound_1 - start_1;
+    const std::size_t l_shift_2 = local_start_2 - start_2;
+
+    set_value(out_acc1, start_out + l_shift_1 + l_shift_2, local_l_item_2,
+              is_first);
+
+    std::size_t r_search_bound_1{};
+    // find right border in 1st sequence
+    if (local_size_2 > 1) {
+      const auto local_r_item_2 = in_acc1[local_end_2 - 1];
+      r_search_bound_1 = detail::upper_bound(in_acc1, l_search_bound_1, end_1,
+                                             local_r_item_2, comp);
+      const std::size_t r_shift_1 = r_search_bound_1 - start_1;
+      const std::size_t r_shift_2 = local_end_2 - 1 - start_2;
+
+      set_value(out_acc1, start_out + r_shift_1 + r_shift_2, local_r_item_2,
+                is_first);
+    }
+
+    // Handle intermediate items
+    for (auto idx = local_start_2 + 1; idx < local_end_2 - 1; ++idx) {
+      const auto intermediate_item_2 = in_acc1[idx];
+      // we shouldn't seek in whole 1st sequence. Just for the part where the
+      // 2nd sequence should be
+      l_search_bound_1 =
+          detail::upper_bound(in_acc1, l_search_bound_1, r_search_bound_1,
+                              intermediate_item_2, comp);
+      const std::size_t shift_1 = l_search_bound_1 - start_1;
+      const std::size_t shift_2 = idx - start_2;
+
+      set_value(out_acc1, start_out + shift_1 + shift_2, intermediate_item_2,
+                is_first);
+    }
+  }
+}
+
+template <typename Iter, typename Compare>
+void bubble_sort(Iter first, const std::size_t begin, const std::size_t end,
+                 Compare comp) {
+  if (begin < end) {
+    for (std::size_t i = begin; i < end; ++i) {
+      // Handle intermediate items
+      for (std::size_t idx = i + 1; idx < end; ++idx) {
+        if (comp(first[idx], first[i])) {
+          detail::swap_tuples(first[i], first[idx]);
+        }
+      }
+    }
+  }
+}
+
+template <typename Group, typename Iter, typename Compare>
+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_linear_id();
+  const std::size_t local = group.get_local_range().size();
+  const std::size_t chunk = (n - 1) / local + 1;
+
+  // we need to sort within work item first
+  bubble_sort(first, idx * chunk, sycl::min((idx + 1) * chunk, n), comp);
+  id.barrier();
+
+  T *temp = reinterpret_cast<T *>(scratch);
+  bool data_in_temp = false;
+  bool is_first = true;
+  std::size_t sorted_size = 1;
+  while (sorted_size * chunk < n) {
+    const std::size_t start_1 =
+        sycl::min(2 * sorted_size * chunk * (idx / sorted_size), n);
+    const std::size_t end_1 = sycl::min(start_1 + sorted_size * chunk, n);
+    const std::size_t end_2 = sycl::min(end_1 + sorted_size * chunk, n);
+    const std::size_t offset = chunk * (idx % sorted_size);
+
+    if (!data_in_temp) {
+      merge(offset, first, temp, start_1, end_1, end_2, start_1, comp, chunk,
+            is_first);
+    } else {
+      merge(offset, temp, first, start_1, end_1, end_2, start_1, comp, chunk,
+            /*is_first*/ false);
+    }
+    id.barrier();
+
+    data_in_temp = !data_in_temp;
+    sorted_size *= 2;
+    if (is_first)
+      is_first = false;
+  }
+
+  // copy back if data is in a temporary storage
+  if (data_in_temp) {
+    for (std::size_t i = 0; i < chunk; ++i) {
+      if (idx * chunk + i < n) {
+        first[idx * chunk + i] = temp[idx * chunk + i];
+      }
+    }
+    id.barrier();
+  }
+}
+
+} // namespace detail
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
+#endif
+#endif // __cplusplus >=201703L
\ No newline at end of file
diff --git a/sycl/include/CL/sycl/group_algorithm.hpp b/sycl/include/CL/sycl/group_algorithm.hpp
index 8277f2d0037be..170b255a0f18e 100644
--- a/sycl/include/CL/sycl/group_algorithm.hpp
+++ b/sycl/include/CL/sycl/group_algorithm.hpp
@@ -18,6 +18,7 @@
 #include <CL/sycl/nd_item.hpp>
 #include <CL/sycl/sub_group.hpp>
 #include <sycl/ext/oneapi/functional.hpp>
+#include <sycl/ext/oneapi/group_sort.hpp>
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
diff --git a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
new file mode 100644
index 0000000000000..932d0a29fb888
--- /dev/null
+++ b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
@@ -0,0 +1,101 @@
+//==------- group_helpers_sorters.hpp - SYCL sorters and group helpers -----==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#if __cplusplus >= 201703L
+#include <CL/sycl/detail/group_sort_impl.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace ext {
+namespace oneapi {
+namespace experimental {
+
+// ---- group helpers
+template <typename Group, std::size_t Extent> class group_with_scratchpad {
+  Group g;
+  sycl::span<std::byte, Extent> scratch;
+
+public:
+  group_with_scratchpad(Group g_, sycl::span<std::byte, Extent> scratch_)
+      : g(g_), scratch(scratch_) {}
+  Group get_group() const { return g; }
+  sycl::span<std::byte, Extent> get_memory() const { return scratch; }
+};
+
+// ---- sorters
+template <typename Compare = std::less<>> class default_sorter {
+  Compare comp;
+  std::byte *scratch;
+  std::size_t scratch_size;
+
+public:
+  template <std::size_t Extent>
+  default_sorter(sycl::span<std::byte, Extent> scratch_,
+                 Compare comp_ = Compare())
+      : comp(comp_), scratch(scratch_.data()), scratch_size(scratch_.size()) {}
+
+  template <typename Group, typename Ptr>
+  void operator()(Group g, Ptr first, 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
+#else
+    (void)g;
+    (void)first;
+    (void)last;
+    throw sycl::exception(
+        std::error_code(PI_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) {
+#ifdef __SYCL_DEVICE_ONLY__
+    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>();
+      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,
+                               scratch + range_size * sizeof(T));
+      val = temp[local_id];
+    }
+    // TODO: it's better to add else branch
+#else
+    (void)g;
+    throw sycl::exception(
+        std::error_code(PI_INVALID_DEVICE, sycl::sycl_category()),
+        "default_sorter operator() is not supported on host device.");
+#endif
+    return val;
+  }
+
+  template <typename T>
+  static constexpr std::size_t memory_required(sycl::memory_scope,
+                                               std::size_t range_size) {
+    return range_size * sizeof(T) + alignof(T);
+  }
+
+  template <typename T, int dim = 1>
+  static constexpr std::size_t memory_required(sycl::memory_scope scope,
+                                               sycl::range<dim> r) {
+    return 2 * memory_required<T>(scope, r.size());
+  }
+};
+
+} // namespace experimental
+} // namespace oneapi
+} // namespace ext
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
+#endif // __cplusplus >=201703L
diff --git a/sycl/include/sycl/ext/oneapi/group_sort.hpp b/sycl/include/sycl/ext/oneapi/group_sort.hpp
new file mode 100644
index 0000000000000..03f1e6500966c
--- /dev/null
+++ b/sycl/include/sycl/ext/oneapi/group_sort.hpp
@@ -0,0 +1,145 @@
+//==--------- group_sort.hpp --- SYCL extension group sorting algorithm-----==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#if __cplusplus >= 201703L
+#include <CL/sycl/detail/defines_elementary.hpp>
+#include <CL/sycl/detail/group_sort_impl.hpp>
+#include <CL/sycl/detail/type_traits.hpp>
+#include <type_traits>
+
+#include "experimental/group_helpers_sorters.hpp"
+
+// define feature macro for group sorting algorithms
+#define SYCL_EXT_ONEAPI_GROUP_SORT 1
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace ext {
+namespace oneapi {
+namespace detail {
+
+// ---- traits
+template <typename T, typename = void> struct has_difference_type {};
+
+template <typename T>
+struct has_difference_type<T, sycl::detail::void_t<typename T::difference_type>>
+    : std::true_type {};
+
+template <typename T> struct has_difference_type<T *> : std::true_type {};
+
+template <typename Sorter, typename Group, typename Val, typename = void>
+struct is_sorter_impl {
+  template <typename G>
+  using is_expected_return_type =
+      typename std::is_same<Val, decltype(std::declval<Sorter>()(
+                                     std::declval<G>(), std::declval<Val>()))>;
+
+  template <typename G = Group>
+  static decltype(
+      std::integral_constant<bool, is_expected_return_type<G>::value &&
+                                       sycl::is_group_v<G>>{})
+  test(int);
+
+  template <typename = Group> static std::false_type test(...);
+};
+
+template <typename Sorter, typename Group,
+          typename Ptr> // multi_ptr has difference_type and don't have other
+                        // iterator's fields
+struct is_sorter_impl<
+    Sorter, Group, Ptr,
+    sycl::detail::void_t<typename has_difference_type<Ptr>::type>> {
+  template <typename G = Group>
+  static decltype(std::declval<Sorter>()(std::declval<G>(), std::declval<Ptr>(),
+                                         std::declval<Ptr>()),
+                  sycl::detail::is_generic_group<G>{})
+  test(int);
+
+  template <typename = Group> static std::false_type test(...);
+};
+
+template <typename Sorter, typename Group, typename ValOrPtr>
+struct is_sorter : decltype(is_sorter_impl<Sorter, Group, ValOrPtr>::test(0)) {
+};
+} // namespace detail
+
+// ---- sort_over_group
+template <typename Group, typename T, typename Sorter>
+typename std::enable_if<detail::is_sorter<Sorter, Group, T>::value, T>::type
+sort_over_group(Group group, T value, Sorter sorter) {
+#ifdef __SYCL_DEVICE_ONLY__
+  return sorter(group, value);
+#else
+  (void)group;
+  (void)value;
+  (void)sorter;
+  throw sycl::exception(
+      std::error_code(PI_INVALID_DEVICE, sycl::sycl_category()),
+      "Group algorithms are not supported on host device.");
+#endif
+}
+
+template <typename Group, typename T, typename Compare, std::size_t Extent>
+typename std::enable_if<!detail::is_sorter<Compare, Group, T>::value, T>::type
+sort_over_group(experimental::group_with_scratchpad<Group, Extent> exec,
+                T value, Compare comp) {
+  return sort_over_group(
+      exec.get_group(), value,
+      experimental::default_sorter<Compare>(exec.get_memory(), comp));
+}
+
+template <typename Group, typename T, std::size_t Extent>
+typename std::enable_if<sycl::is_group_v<std::decay_t<Group>>, T>::type
+sort_over_group(experimental::group_with_scratchpad<Group, Extent> exec,
+                T value) {
+  return sort_over_group(exec.get_group(), value,
+                         experimental::default_sorter<>(exec.get_memory()));
+}
+
+// ---- joint_sort
+template <typename Group, typename Iter, typename Sorter>
+typename std::enable_if<detail::is_sorter<Sorter, Group, Iter>::value,
+                        void>::type
+joint_sort(Group group, Iter first, Iter last, 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_INVALID_DEVICE, sycl::sycl_category()),
+      "Group algorithms are not supported on host device.");
+#endif
+}
+
+template <typename Group, typename Iter, typename Compare, std::size_t Extent>
+typename std::enable_if<!detail::is_sorter<Compare, Group, Iter>::value,
+                        void>::type
+joint_sort(experimental::group_with_scratchpad<Group, Extent> exec, Iter first,
+           Iter last, Compare comp) {
+  joint_sort(exec.get_group(), first, last,
+             experimental::default_sorter<Compare>(exec.get_memory(), comp));
+}
+
+template <typename Group, typename Iter, std::size_t Extent>
+typename std::enable_if<sycl::is_group_v<std::decay_t<Group>>, void>::type
+joint_sort(experimental::group_with_scratchpad<Group, Extent> exec, Iter first,
+           Iter last) {
+  joint_sort(exec.get_group(), first, last,
+             experimental::default_sorter<>(exec.get_memory()));
+}
+
+} // namespace oneapi
+} // namespace ext
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
+#endif // __cplusplus >=201703L