From a8a5b330a546b0d44d6617d9cf80ab1411311c4c Mon Sep 17 00:00:00 2001
From: "Fedorov, Andrey" <andrey.fedorov@intel.com>
Date: Tue, 31 Aug 2021 15:12:30 +0300
Subject: [PATCH 01/16] initial commit for group sort implementation

Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
---
 .../CL/sycl/detail/group_sort_impl.hpp        | 232 ++++++++++++++++++
 .../CL/sycl/detail/group_sort_util.hpp        |  30 +++
 sycl/include/CL/sycl/group_algorithm.hpp      |   1 +
 .../experimental/group_helpers_sorters.hpp    |  91 +++++++
 sycl/include/sycl/ext/oneapi/group_sort.hpp   | 136 ++++++++++
 5 files changed, 490 insertions(+)
 create mode 100644 sycl/include/CL/sycl/detail/group_sort_impl.hpp
 create mode 100644 sycl/include/CL/sycl/detail/group_sort_util.hpp
 create mode 100644 sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
 create mode 100644 sycl/include/sycl/ext/oneapi/group_sort.hpp

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..b9b52f8aa5008
--- /dev/null
+++ b/sycl/include/CL/sycl/detail/group_sort_impl.hpp
@@ -0,0 +1,232 @@
+//==------------ 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
+
+#include "group_sort_util.hpp"
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace detail {
+
+// ---- merge sort implementation
+
+// following to functions could be useless if std::[lower|upper]_bound worked well
+template <typename Acc, typename Value, typename Compare>
+std::size_t
+my_lower_bound(Acc acc, const std::size_t first, const std::size_t last, const Value& value, Compare comp)
+{
+    const std::size_t n = last - first;
+    const std::size_t cur = n;
+    const 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
+my_upper_bound(Acc acc, const std::size_t first, const std::size_t last, const Value& value, Compare comp)
+{
+    return my_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 my_swap(T& a, T& b)
+{
+    std::swap(a, b);
+}
+
+template<template <typename...> class TupleLike, typename T1, typename T2>
+void my_swap(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 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)
+{
+    using _Idx = std::size_t;
+    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 = my_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;
+
+        out_acc1[start_out + l_shift_1 + l_shift_2] = local_l_item_1;
+
+        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 = my_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;
+
+            out_acc1[start_out + r_shift_1 + r_shift_2] = local_r_item_1;
+        }
+
+        // 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 = my_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;
+
+            out_acc1[start_out + shift_1 + shift_2] = intermediate_item_1;
+        }
+    }
+    // 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 = my_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;
+
+        out_acc1[start_out + l_shift_1 + l_shift_2] = local_l_item_2;
+
+        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 = my_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;
+
+            out_acc1[start_out + r_shift_1 + r_shift_2] = local_r_item_2;
+        }
+
+        // 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 = my_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;
+
+            out_acc1[start_out + shift_1 + shift_2] = intermediate_item_2;
+        }
+    }
+}
+
+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]))
+                {
+                    my_swap(first[i], first[idx]);
+                }
+            }
+        }
+    }
+}
+
+template<typename Group, typename Iter, typename Compare, typename Id>
+void
+merge_sort(Group group, Iter first, const std::size_t n, Compare comp,
+        std::uint8_t* scratch, Id id)
+{
+    using T = typename std::iterator_traits<Iter>::value_type;
+    const std::size_t idx = id.get_local_id();
+    const std::size_t local = group.get_local_range(0);
+    const std::size_t chunk = (n - 1) / local + 1;
+
+    // we need to sort within work item firstly
+    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;
+    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);
+        }
+        else
+        {
+            merge(offset, temp, first, start_1, end_1, end_2, start_1, comp, chunk);
+        }
+        id.barrier();
+
+        data_in_temp = !data_in_temp;
+        sorted_size *= 2;
+    }
+
+    // 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)
diff --git a/sycl/include/CL/sycl/detail/group_sort_util.hpp b/sycl/include/CL/sycl/detail/group_sort_util.hpp
new file mode 100644
index 0000000000000..d9f3ef3b940d1
--- /dev/null
+++ b/sycl/include/CL/sycl/detail/group_sort_util.hpp
@@ -0,0 +1,30 @@
+//==------------ group_sort_util.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 utilities that are used by group sorting algorithms
+//
+
+#pragma once
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace detail {
+
+#if __cplusplus < 201703
+template <typename... _Ts>
+struct make_void_type
+{
+    using type = void;
+};
+
+template <typename... _Ts>
+using void_type = typename make_void_type<_Ts...>::type;
+#endif
+
+} // namespace detail
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
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..0638d6d1814b1
--- /dev/null
+++ b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
@@ -0,0 +1,91 @@
+//==------- 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
+
+#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::uint8_t, Extent> scratch;
+public:
+    group_with_scratchpad(Group g_, sycl::span<std::uint8_t, Extent> scratch_)
+        : g(g_), scratch(scratch_) {}
+    Group get_group() const {return g;}
+    sycl::span<std::uint8_t, Extent>
+    get_memory() const {return scratch;}
+};
+
+// ---- sorters
+template<typename Compare = std::less<>>
+class default_sorter
+{
+    Compare comp;
+    std::uint8_t* scratch;
+    std::size_t scratch_size;
+public:
+    template<std::size_t Extent>
+    default_sorter(sycl::span<std::uint8_t, Extent> scratch_, Compare comp_ = Compare())
+        : comp(comp_), scratch(scratch_.data()), scratch_size(scratch_.size())) {}
+
+    template<typename Group, typename Ptr>
+    void operator()(Group g, Ptr begin, Ptr end)
+    {
+        using T = typename std::iterator_traits<Ptr>::value_type;
+        if(scratch_size >= memory_required<T>(sycl::memory_scope::work_group/*Group::fence_scope*/, end - begin))
+            cl::sycl::detail::merge_sort(g, begin, end - begin, comp, scratch,
+                cl::sycl::detail::Builder::getNDItem<Group::dimensions>());
+        // TODO: it's better to add else branch
+    }
+
+    template<typename Group, typename T>
+    T operator()(Group g, T val)
+    {
+        auto range_size = g.get_local_range(0);
+        if(scratch_size >= memory_required<T>(sycl::memory_scope::work_group/*Group::fence_scope*/, range_size))
+        {
+            uint32_t local_id = id.get_local_id();
+            T* temp = reinterpret_cast<T*>(scratch);
+            temp[local_id] = val;
+            cl::sycl::detail::merge_sort(g, temp, range_size, comp, scratch + range_size * sizeof(T),
+                cl::sycl::detail::Builder::getNDItem<Group::dimensions>());
+            val = temp[local_id];
+        }
+        // TODO: it's better to add else branch
+        return val;
+    }
+
+    template<typename T>
+    static constexpr std::size_t
+    memory_required(sycl::memory_scope scope, std::size_t range_size)
+    {
+        return range_size * sizeof(T);
+    }
+
+    template<typename T, int dim = 1>
+    static constexpr std::size_t
+    memory_required(sycl::memory_scope scope, sycl::range<dim> r)
+    {
+        return 2 * r.size() * sizeof(T);
+    }
+};
+
+} // namespace experimental
+} // namespace oneapi
+} // namespace ext
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
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..6af32e70cc2ed
--- /dev/null
+++ b/sycl/include/sycl/ext/oneapi/group_sort.hpp
@@ -0,0 +1,136 @@
+//==--------- 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
+
+#include <type_traits>
+#include <CL/sycl/detail/defines_elementary.hpp>
+#include <CL/sycl/detail/group_sort_impl.hpp>
+#include <CL/sycl/detail/group_sort_util.hpp>
+
+#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,
+#if __cplusplus < 201703
+                        cl::sycl::detail::void_type<
+#else
+                        std::void_t<
+#endif
+                                    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 && true /*sycl::is_group<G>::value*/>{})
+    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,
+#if __cplusplus < 201703
+                        cl::sycl::detail::void_type<
+#else
+                        std::void_t<
+#endif
+                                    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::is_group<G>*/std::true_type{})
+    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)) {};
+
+// ---- 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)
+{
+    return sorter(group, value);
+}
+
+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(cl::sycl::ext::oneapi::experimental::group_with_scratchpad<Group, Extent> exec, T value, Compare comp)
+{
+    return sort_over_group(exec.get_group(), value,
+            cl::sycl::ext::oneapi::experimental::default_sorter<Compare>(exec.get_memory(),
+            cl::sycl::detail::Builder::getNDItem<Group::dimensions>(), comp));
+}
+
+template<typename Group, typename T, std::size_t Extent>
+T
+sort_over_group(cl::sycl::ext::oneapi::experimental::group_with_scratchpad<Group, Extent> exec, T value)
+{
+    return sort_over_group(exec.get_group(), value,
+            cl::sycl::ext::oneapi::experimental::default_sorter<>(exec.get_memory(), cl::sycl::detail::Builder::getNDItem<Group::dimensions>()));
+}
+
+// ---- 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)
+{
+    sorter(group, first, last);
+}
+
+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(cl::sycl::ext::oneapi::experimental::group_with_scratchpad<Group, Extent> exec, Iter first, Iter last, Compare comp)
+{
+    joint_sort(exec.get_group(), first, last, cl::sycl::ext::oneapi::experimental::default_sorter<Compare>(exec.get_memory(),
+        cl::sycl::detail::Builder::getNDItem<Group::dimensions>(), comp));
+}
+
+template<typename Group, typename Iter, std::size_t Extent>
+void
+joint_sort(cl::sycl::ext::oneapi::experimental::group_with_scratchpad<Group, Extent> exec, Iter first, Iter last)
+{
+    joint_sort(exec.get_group(), first, last, cl::sycl::ext::oneapi::experimental::default_sorter<>(exec.get_memory(),
+        cl::sycl::detail::Builder::getNDItem<Group::dimensions>()));
+}
+
+} // namespace oneapi
+} // namespace ext
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

From c27f0ac033db57121ed094ac10869b9d5f7bb1ae Mon Sep 17 00:00:00 2001
From: "Fedorov, Andrey" <andrey.fedorov@intel.com>
Date: Tue, 31 Aug 2021 18:07:35 +0300
Subject: [PATCH 02/16] fixed compilation issues

Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
---
 .../CL/sycl/detail/group_sort_impl.hpp        | 16 +++++----
 .../experimental/group_helpers_sorters.hpp    | 12 ++++---
 sycl/include/sycl/ext/oneapi/group_sort.hpp   | 33 ++++++++++++-------
 3 files changed, 39 insertions(+), 22 deletions(-)

diff --git a/sycl/include/CL/sycl/detail/group_sort_impl.hpp b/sycl/include/CL/sycl/detail/group_sort_impl.hpp
index b9b52f8aa5008..5312675b62bf4 100644
--- a/sycl/include/CL/sycl/detail/group_sort_impl.hpp
+++ b/sycl/include/CL/sycl/detail/group_sort_impl.hpp
@@ -10,6 +10,9 @@
 
 #pragma once
 
+#include <CL/sycl/detail/helpers.hpp>
+
+#ifdef __SYCL_DEVICE_ONLY__
 #include "group_sort_util.hpp"
 
 __SYCL_INLINE_NAMESPACE(cl) {
@@ -23,9 +26,9 @@ template <typename Acc, typename Value, typename Compare>
 std::size_t
 my_lower_bound(Acc acc, const std::size_t first, const std::size_t last, const Value& value, Compare comp)
 {
-    const std::size_t n = last - first;
-    const std::size_t cur = n;
-    const std::size_t it;
+    std::size_t n = last - first;
+    std::size_t cur = n;
+    std::size_t it;
     while (n > 0)
     {
         it = first;
@@ -68,7 +71,6 @@ 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)
 {
-    using _Idx = std::size_t;
     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);
@@ -175,12 +177,13 @@ bubble_sort(Iter first, const std::size_t begin, const std::size_t end, Compare
     }
 }
 
-template<typename Group, typename Iter, typename Compare, typename Id>
+template<typename Group, typename Iter, typename Compare>
 void
 merge_sort(Group group, Iter first, const std::size_t n, Compare comp,
-        std::uint8_t* scratch, Id id)
+        std::uint8_t* scratch)
 {
     using T = typename std::iterator_traits<Iter>::value_type;
+    auto id = cl::sycl::detail::Builder::getNDItem<Group::dimensions>()
     const std::size_t idx = id.get_local_id();
     const std::size_t local = group.get_local_range(0);
     const std::size_t chunk = (n - 1) / local + 1;
@@ -230,3 +233,4 @@ merge_sort(Group group, Iter first, const std::size_t n, Compare comp,
 } // namespace detail
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
+#endif
diff --git a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
index 0638d6d1814b1..97eb70fa9c7d2 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
@@ -40,32 +40,34 @@ class default_sorter
 public:
     template<std::size_t Extent>
     default_sorter(sycl::span<std::uint8_t, Extent> scratch_, Compare comp_ = Compare())
-        : comp(comp_), scratch(scratch_.data()), scratch_size(scratch_.size())) {}
+        : comp(comp_), scratch(scratch_.data()), scratch_size(scratch_.size()) {}
 
     template<typename Group, typename Ptr>
     void operator()(Group g, Ptr begin, Ptr end)
     {
+#ifdef __SYCL_DEVICE_ONLY__
         using T = typename std::iterator_traits<Ptr>::value_type;
         if(scratch_size >= memory_required<T>(sycl::memory_scope::work_group/*Group::fence_scope*/, end - begin))
-            cl::sycl::detail::merge_sort(g, begin, end - begin, comp, scratch,
-                cl::sycl::detail::Builder::getNDItem<Group::dimensions>());
+            cl::sycl::detail::merge_sort(g, begin, end - begin, comp, scratch);
         // TODO: it's better to add else branch
+#endif
     }
 
     template<typename Group, typename T>
     T operator()(Group g, T val)
     {
+#ifdef __SYCL_DEVICE_ONLY__
         auto range_size = g.get_local_range(0);
         if(scratch_size >= memory_required<T>(sycl::memory_scope::work_group/*Group::fence_scope*/, range_size))
         {
             uint32_t local_id = id.get_local_id();
             T* temp = reinterpret_cast<T*>(scratch);
             temp[local_id] = val;
-            cl::sycl::detail::merge_sort(g, temp, range_size, comp, scratch + range_size * sizeof(T),
-                cl::sycl::detail::Builder::getNDItem<Group::dimensions>());
+            cl::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
+#endif
         return val;
     }
 
diff --git a/sycl/include/sycl/ext/oneapi/group_sort.hpp b/sycl/include/sycl/ext/oneapi/group_sort.hpp
index 6af32e70cc2ed..d49a47e23dbde 100644
--- a/sycl/include/sycl/ext/oneapi/group_sort.hpp
+++ b/sycl/include/sycl/ext/oneapi/group_sort.hpp
@@ -10,6 +10,7 @@
 
 #include <type_traits>
 #include <CL/sycl/detail/defines_elementary.hpp>
+#include <CL/sycl/detail/type_traits.hpp>
 #include <CL/sycl/detail/group_sort_impl.hpp>
 #include <CL/sycl/detail/group_sort_util.hpp>
 
@@ -48,7 +49,7 @@ struct is_sorter_impl
 
     template<typename G = Group>
     static
-    decltype(std::integral_constant<bool, is_expected_return_type<G>::value && true /*sycl::is_group<G>::value*/>{})
+    decltype(std::integral_constant<bool, is_expected_return_type<G>::value && sycl::is_group_v<G>>{})
     test(int);
 
     template<typename = Group>
@@ -69,7 +70,7 @@ struct is_sorter_impl<Sorter, Group, Ptr,
 {
     template<typename G = Group>
     static
-    decltype(std::declval<Sorter>()(std::declval<G>(), std::declval<Ptr>(), std::declval<Ptr>()), /*sycl::is_group<G>*/std::true_type{})
+    decltype(std::declval<Sorter>()(std::declval<G>(), std::declval<Ptr>(), std::declval<Ptr>()), sycl::detail::is_generic_group<G>{})
     test(int);
 
     template<typename = Group>
@@ -80,13 +81,20 @@ struct is_sorter_impl<Sorter, Group, Ptr,
 
 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;
+    throw runtime_error("Group algorithms are not supported on host device.",
+                        PI_INVALID_DEVICE);
+#endif
 }
 
 template<typename Group, typename T, typename Compare, std::size_t Extent>
@@ -94,16 +102,15 @@ typename std::enable_if<!detail::is_sorter<Compare, Group, T>::value, T>::type
 sort_over_group(cl::sycl::ext::oneapi::experimental::group_with_scratchpad<Group, Extent> exec, T value, Compare comp)
 {
     return sort_over_group(exec.get_group(), value,
-            cl::sycl::ext::oneapi::experimental::default_sorter<Compare>(exec.get_memory(),
-            cl::sycl::detail::Builder::getNDItem<Group::dimensions>(), comp));
+            cl::sycl::ext::oneapi::experimental::default_sorter<Compare>(exec.get_memory(), comp));
 }
 
 template<typename Group, typename T, std::size_t Extent>
-T
+typename std::enable_if<sycl::is_group_v<std::decay_t<Group>>, T>::type
 sort_over_group(cl::sycl::ext::oneapi::experimental::group_with_scratchpad<Group, Extent> exec, T value)
 {
     return sort_over_group(exec.get_group(), value,
-            cl::sycl::ext::oneapi::experimental::default_sorter<>(exec.get_memory(), cl::sycl::detail::Builder::getNDItem<Group::dimensions>()));
+            cl::sycl::ext::oneapi::experimental::default_sorter<>(exec.get_memory()));
 }
 
 // ---- joint_sort
@@ -111,23 +118,27 @@ 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;
+    throw runtime_error("Group algorithms are not supported on host device.",
+                        PI_INVALID_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(cl::sycl::ext::oneapi::experimental::group_with_scratchpad<Group, Extent> exec, Iter first, Iter last, Compare comp)
 {
-    joint_sort(exec.get_group(), first, last, cl::sycl::ext::oneapi::experimental::default_sorter<Compare>(exec.get_memory(),
-        cl::sycl::detail::Builder::getNDItem<Group::dimensions>(), comp));
+    joint_sort(exec.get_group(), first, last, cl::sycl::ext::oneapi::experimental::default_sorter<Compare>(exec.get_memory(), comp));
 }
 
 template<typename Group, typename Iter, std::size_t Extent>
-void
+typename std::enable_if<sycl::is_group_v<std::decay_t<Group>>, void>::type
 joint_sort(cl::sycl::ext::oneapi::experimental::group_with_scratchpad<Group, Extent> exec, Iter first, Iter last)
 {
-    joint_sort(exec.get_group(), first, last, cl::sycl::ext::oneapi::experimental::default_sorter<>(exec.get_memory(),
-        cl::sycl::detail::Builder::getNDItem<Group::dimensions>()));
+    joint_sort(exec.get_group(), first, last, cl::sycl::ext::oneapi::experimental::default_sorter<>(exec.get_memory()));
 }
 
 } // namespace oneapi

From 0159ff8c51be4a5e442b3123171ac5e84f03680f Mon Sep 17 00:00:00 2001
From: "Fedorov, Andrey" <andrey.fedorov@intel.com>
Date: Wed, 1 Sep 2021 11:20:12 +0300
Subject: [PATCH 03/16] applied clang-format

Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
---
 .../CL/sycl/detail/group_sort_impl.hpp        | 351 +++++++++---------
 .../CL/sycl/detail/group_sort_util.hpp        |   6 +-
 .../experimental/group_helpers_sorters.hpp    | 106 +++---
 sycl/include/sycl/ext/oneapi/group_sort.hpp   | 161 ++++----
 4 files changed, 309 insertions(+), 315 deletions(-)

diff --git a/sycl/include/CL/sycl/detail/group_sort_impl.hpp b/sycl/include/CL/sycl/detail/group_sort_impl.hpp
index 5312675b62bf4..55ec9f5645237 100644
--- a/sycl/include/CL/sycl/detail/group_sort_impl.hpp
+++ b/sycl/include/CL/sycl/detail/group_sort_impl.hpp
@@ -21,213 +21,202 @@ namespace detail {
 
 // ---- merge sort implementation
 
-// following to functions could be useless if std::[lower|upper]_bound worked well
+// following to functions could be useless if std::[lower|upper]_bound worked
+// well
 template <typename Acc, typename Value, typename Compare>
-std::size_t
-my_lower_bound(Acc acc, const std::size_t first, const 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;
+std::size_t my_lower_bound(Acc acc, const std::size_t first,
+                           const 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
-my_upper_bound(Acc acc, const std::size_t first, const std::size_t last, const Value& value, Compare comp)
-{
-    return my_lower_bound(acc, first, last, value, [comp](auto x, auto y){return !comp(y, x);});
+std::size_t my_upper_bound(Acc acc, const std::size_t first,
+                           const std::size_t last, const Value &value,
+                           Compare comp) {
+  return my_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 my_swap(T& a, T& b)
-{
-    std::swap(a, b);
-}
+template <typename T> void my_swap(T &a, T &b) { std::swap(a, b); }
 
-template<template <typename...> class TupleLike, typename T1, typename T2>
-void my_swap(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 <template <typename...> class TupleLike, typename T1, typename T2>
+void my_swap(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 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)
-{
-    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 = my_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;
-
-        out_acc1[start_out + l_shift_1 + l_shift_2] = local_l_item_1;
-
-        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 = my_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;
-
-            out_acc1[start_out + r_shift_1 + r_shift_2] = local_r_item_1;
-        }
-
-        // 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 = my_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;
-
-            out_acc1[start_out + shift_1 + shift_2] = intermediate_item_1;
-        }
+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) {
+  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 =
+        my_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;
+
+    out_acc1[start_out + l_shift_1 + l_shift_2] = local_l_item_1;
+
+    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 = my_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;
+
+      out_acc1[start_out + r_shift_1 + r_shift_2] = local_r_item_1;
     }
-    // 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 = my_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;
-
-        out_acc1[start_out + l_shift_1 + l_shift_2] = local_l_item_2;
-
-        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 = my_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;
-
-            out_acc1[start_out + r_shift_1 + r_shift_2] = local_r_item_2;
-        }
 
-        // 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 = my_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;
+    // 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 =
+          my_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;
+
+      out_acc1[start_out + shift_1 + shift_2] = intermediate_item_1;
+    }
+  }
+  // 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 =
+        my_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;
+
+    out_acc1[start_out + l_shift_1 + l_shift_2] = local_l_item_2;
+
+    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 = my_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;
+
+      out_acc1[start_out + r_shift_1 + r_shift_2] = local_r_item_2;
+    }
 
-            out_acc1[start_out + shift_1 + shift_2] = intermediate_item_2;
-        }
+    // 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 =
+          my_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;
+
+      out_acc1[start_out + shift_1 + shift_2] = intermediate_item_2;
     }
+  }
 }
 
 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]))
-                {
-                    my_swap(first[i], first[idx]);
-                }
-            }
+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])) {
+          my_swap(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::uint8_t* scratch)
-{
-    using T = typename std::iterator_traits<Iter>::value_type;
-    auto id = cl::sycl::detail::Builder::getNDItem<Group::dimensions>()
-    const std::size_t idx = id.get_local_id();
-    const std::size_t local = group.get_local_range(0);
-    const std::size_t chunk = (n - 1) / local + 1;
-
-    // we need to sort within work item firstly
-    bubble_sort(first, idx * chunk, sycl::min((idx + 1) * chunk, n), comp);
+template <typename Group, typename Iter, typename Compare>
+void merge_sort(Group group, Iter first, const std::size_t n, Compare comp,
+                std::uint8_t *scratch) {
+  using T = typename std::iterator_traits<Iter>::value_type;
+  auto id = cl::sycl::detail::Builder::getNDItem<Group::dimensions>()
+      const std::size_t idx = id.get_local_id();
+  const std::size_t local = group.get_local_range(0);
+  const std::size_t chunk = (n - 1) / local + 1;
+
+  // we need to sort within work item firstly
+  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;
+  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);
+    } else {
+      merge(offset, temp, first, start_1, end_1, end_2, start_1, comp, chunk);
+    }
     id.barrier();
 
-    T* temp = reinterpret_cast<T*>(scratch);
-    bool data_in_temp = false;
-    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);
-        }
-        else
-        {
-            merge(offset, temp, first, start_1, end_1, end_2, start_1, comp, chunk);
-        }
-        id.barrier();
+    data_in_temp = !data_in_temp;
+    sorted_size *= 2;
+  }
 
-        data_in_temp = !data_in_temp;
-        sorted_size *= 2;
-    }
-
-    // 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();
+  // 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
diff --git a/sycl/include/CL/sycl/detail/group_sort_util.hpp b/sycl/include/CL/sycl/detail/group_sort_util.hpp
index d9f3ef3b940d1..9765088f6597f 100644
--- a/sycl/include/CL/sycl/detail/group_sort_util.hpp
+++ b/sycl/include/CL/sycl/detail/group_sort_util.hpp
@@ -15,11 +15,7 @@ namespace sycl {
 namespace detail {
 
 #if __cplusplus < 201703
-template <typename... _Ts>
-struct make_void_type
-{
-    using type = void;
-};
+template <typename... _Ts> struct make_void_type { using type = void; };
 
 template <typename... _Ts>
 using void_type = typename make_void_type<_Ts...>::type;
diff --git a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
index 97eb70fa9c7d2..e4fa348762838 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
@@ -17,73 +17,71 @@ namespace oneapi {
 namespace experimental {
 
 // ---- group helpers
-template<typename Group, std::size_t Extent>
-class group_with_scratchpad
-{
-    Group g;
-    sycl::span<std::uint8_t, Extent> scratch;
+template <typename Group, std::size_t Extent> class group_with_scratchpad {
+  Group g;
+  sycl::span<std::uint8_t, Extent> scratch;
+
 public:
-    group_with_scratchpad(Group g_, sycl::span<std::uint8_t, Extent> scratch_)
-        : g(g_), scratch(scratch_) {}
-    Group get_group() const {return g;}
-    sycl::span<std::uint8_t, Extent>
-    get_memory() const {return scratch;}
+  group_with_scratchpad(Group g_, sycl::span<std::uint8_t, Extent> scratch_)
+      : g(g_), scratch(scratch_) {}
+  Group get_group() const { return g; }
+  sycl::span<std::uint8_t, Extent> get_memory() const { return scratch; }
 };
 
 // ---- sorters
-template<typename Compare = std::less<>>
-class default_sorter
-{
-    Compare comp;
-    std::uint8_t* scratch;
-    std::size_t scratch_size;
+template <typename Compare = std::less<>> class default_sorter {
+  Compare comp;
+  std::uint8_t *scratch;
+  std::size_t scratch_size;
+
 public:
-    template<std::size_t Extent>
-    default_sorter(sycl::span<std::uint8_t, Extent> scratch_, Compare comp_ = Compare())
-        : comp(comp_), scratch(scratch_.data()), scratch_size(scratch_.size()) {}
+  template <std::size_t Extent>
+  default_sorter(sycl::span<std::uint8_t, Extent> scratch_,
+                 Compare comp_ = Compare())
+      : comp(comp_), scratch(scratch_.data()), scratch_size(scratch_.size()) {}
 
-    template<typename Group, typename Ptr>
-    void operator()(Group g, Ptr begin, Ptr end)
-    {
+  template <typename Group, typename Ptr>
+  void operator()(Group g, Ptr begin, Ptr end) {
 #ifdef __SYCL_DEVICE_ONLY__
-        using T = typename std::iterator_traits<Ptr>::value_type;
-        if(scratch_size >= memory_required<T>(sycl::memory_scope::work_group/*Group::fence_scope*/, end - begin))
-            cl::sycl::detail::merge_sort(g, begin, end - begin, comp, scratch);
-        // TODO: it's better to add else branch
+    using T = typename std::iterator_traits<Ptr>::value_type;
+    if (scratch_size >=
+        memory_required<T>(
+            sycl::memory_scope::work_group /*Group::fence_scope*/, end - begin))
+      cl::sycl::detail::merge_sort(g, begin, end - begin, comp, scratch);
+      // TODO: it's better to add else branch
 #endif
-    }
+  }
 
-    template<typename Group, typename T>
-    T operator()(Group g, T val)
-    {
+  template <typename Group, typename T> T operator()(Group g, T val) {
 #ifdef __SYCL_DEVICE_ONLY__
-        auto range_size = g.get_local_range(0);
-        if(scratch_size >= memory_required<T>(sycl::memory_scope::work_group/*Group::fence_scope*/, range_size))
-        {
-            uint32_t local_id = id.get_local_id();
-            T* temp = reinterpret_cast<T*>(scratch);
-            temp[local_id] = val;
-            cl::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
-#endif
-        return val;
+    auto range_size = g.get_local_range(0);
+    if (scratch_size >=
+        memory_required<T>(
+            sycl::memory_scope::work_group /*Group::fence_scope*/,
+            range_size)) {
+      uint32_t local_id = id.get_local_id();
+      T *temp = reinterpret_cast<T *>(scratch);
+      temp[local_id] = val;
+      cl::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
+#endif
+    return val;
+  }
 
-    template<typename T>
-    static constexpr std::size_t
-    memory_required(sycl::memory_scope scope, std::size_t range_size)
-    {
-        return range_size * sizeof(T);
-    }
+  template <typename T>
+  static constexpr std::size_t memory_required(sycl::memory_scope scope,
+                                               std::size_t range_size) {
+    return range_size * sizeof(T);
+  }
 
-    template<typename T, int dim = 1>
-    static constexpr std::size_t
-    memory_required(sycl::memory_scope scope, sycl::range<dim> r)
-    {
-        return 2 * r.size() * sizeof(T);
-    }
+  template <typename T, int dim = 1>
+  static constexpr std::size_t memory_required(sycl::memory_scope scope,
+                                               sycl::range<dim> r) {
+    return 2 * r.size() * sizeof(T);
+  }
 };
 
 } // namespace experimental
diff --git a/sycl/include/sycl/ext/oneapi/group_sort.hpp b/sycl/include/sycl/ext/oneapi/group_sort.hpp
index d49a47e23dbde..b919882ad1b9b 100644
--- a/sycl/include/sycl/ext/oneapi/group_sort.hpp
+++ b/sycl/include/sycl/ext/oneapi/group_sort.hpp
@@ -8,11 +8,11 @@
 
 #pragma once
 
-#include <type_traits>
 #include <CL/sycl/detail/defines_elementary.hpp>
-#include <CL/sycl/detail/type_traits.hpp>
 #include <CL/sycl/detail/group_sort_impl.hpp>
 #include <CL/sycl/detail/group_sort_util.hpp>
+#include <CL/sycl/detail/type_traits.hpp>
+#include <type_traits>
 
 #include "experimental/group_helpers_sorters.hpp"
 
@@ -26,119 +26,130 @@ namespace oneapi {
 namespace detail {
 
 // ---- traits
-template< typename T, typename = void >
-struct has_difference_type{ };
+template <typename T, typename = void> struct has_difference_type {};
 
-template< typename T >
+template <typename T>
 struct has_difference_type<T,
 #if __cplusplus < 201703
-                        cl::sycl::detail::void_type<
+                           cl::sycl::detail::void_type<
 #else
-                        std::void_t<
+                           std::void_t<
 #endif
-                                    typename T::difference_type>> : std::true_type {};
-
-template< typename T >
-struct has_difference_type<T*> : std::true_type {};
+                               typename T::difference_type>> : 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 T> struct has_difference_type<T *> : std::true_type {};
 
-    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 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 = Group>
-    static
-    std::false_type
-    test(...);
+  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,
+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,
 #if __cplusplus < 201703
-                        cl::sycl::detail::void_type<
+                            cl::sycl::detail::void_type<
 #else
-                        std::void_t<
+                            std::void_t<
 #endif
-                                    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(...);
+                                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)) {};
+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>
+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)
-{
+sort_over_group(Group group, T value, Sorter sorter) {
 #ifdef __SYCL_DEVICE_ONLY__
-    return sorter(group, value);
+  return sorter(group, value);
 #else
-    (void)group;
-    throw runtime_error("Group algorithms are not supported on host device.",
-                        PI_INVALID_DEVICE);
+  (void)group;
+  throw runtime_error("Group algorithms are not supported on host device.",
+                      PI_INVALID_DEVICE);
 #endif
 }
 
-template<typename Group, typename T, typename Compare, std::size_t Extent>
+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(cl::sycl::ext::oneapi::experimental::group_with_scratchpad<Group, Extent> exec, T value, Compare comp)
-{
-    return sort_over_group(exec.get_group(), value,
-            cl::sycl::ext::oneapi::experimental::default_sorter<Compare>(exec.get_memory(), comp));
+sort_over_group(
+    cl::sycl::ext::oneapi::experimental::group_with_scratchpad<Group, Extent>
+        exec,
+    T value, Compare comp) {
+  return sort_over_group(
+      exec.get_group(), value,
+      cl::sycl::ext::oneapi::experimental::default_sorter<Compare>(
+          exec.get_memory(), comp));
 }
 
-template<typename Group, typename T, std::size_t Extent>
+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(cl::sycl::ext::oneapi::experimental::group_with_scratchpad<Group, Extent> exec, T value)
-{
-    return sort_over_group(exec.get_group(), value,
-            cl::sycl::ext::oneapi::experimental::default_sorter<>(exec.get_memory()));
+sort_over_group(
+    cl::sycl::ext::oneapi::experimental::group_with_scratchpad<Group, Extent>
+        exec,
+    T value) {
+  return sort_over_group(
+      exec.get_group(), value,
+      cl::sycl::ext::oneapi::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)
-{
+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);
+  sorter(group, first, last);
 #else
-    (void)group;
-    throw runtime_error("Group algorithms are not supported on host device.",
-                        PI_INVALID_DEVICE);
+  (void)group;
+  throw runtime_error("Group algorithms are not supported on host device.",
+                      PI_INVALID_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(cl::sycl::ext::oneapi::experimental::group_with_scratchpad<Group, Extent> exec, Iter first, Iter last, Compare comp)
-{
-    joint_sort(exec.get_group(), first, last, cl::sycl::ext::oneapi::experimental::default_sorter<Compare>(exec.get_memory(), comp));
+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(
+    cl::sycl::ext::oneapi::experimental::group_with_scratchpad<Group, Extent>
+        exec,
+    Iter first, Iter last, Compare comp) {
+  joint_sort(exec.get_group(), first, last,
+             cl::sycl::ext::oneapi::experimental::default_sorter<Compare>(
+                 exec.get_memory(), comp));
 }
 
-template<typename Group, typename Iter, std::size_t Extent>
+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(cl::sycl::ext::oneapi::experimental::group_with_scratchpad<Group, Extent> exec, Iter first, Iter last)
-{
-    joint_sort(exec.get_group(), first, last, cl::sycl::ext::oneapi::experimental::default_sorter<>(exec.get_memory()));
+joint_sort(
+    cl::sycl::ext::oneapi::experimental::group_with_scratchpad<Group, Extent>
+        exec,
+    Iter first, Iter last) {
+  joint_sort(
+      exec.get_group(), first, last,
+      cl::sycl::ext::oneapi::experimental::default_sorter<>(exec.get_memory()));
 }
 
 } // namespace oneapi

From d8aab753854e2e0c487dcc67cf70389cb497fb06 Mon Sep 17 00:00:00 2001
From: "Fedorov, Andrey" <andrey.fedorov@intel.com>
Date: Wed, 1 Sep 2021 11:42:30 +0300
Subject: [PATCH 04/16] applied clang-format

Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
---
 sycl/include/sycl/ext/oneapi/group_sort.hpp | 9 ++++-----
 1 file changed, 4 insertions(+), 5 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/group_sort.hpp b/sycl/include/sycl/ext/oneapi/group_sort.hpp
index b919882ad1b9b..f3d794696be84 100644
--- a/sycl/include/sycl/ext/oneapi/group_sort.hpp
+++ b/sycl/include/sycl/ext/oneapi/group_sort.hpp
@@ -58,14 +58,13 @@ struct is_sorter_impl {
 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,
+struct is_sorter_impl<Sorter, Group, Ptr,
 #if __cplusplus < 201703
-                            cl::sycl::detail::void_type<
+                      cl::sycl::detail::void_type<
 #else
-                            std::void_t<
+                      std::void_t<
 #endif
-                                typename has_difference_type<Ptr>::type>> {
+                          typename has_difference_type<Ptr>::type>> {
   template <typename G = Group>
   static decltype(std::declval<Sorter>()(std::declval<G>(), std::declval<Ptr>(),
                                          std::declval<Ptr>()),

From f40e9cc065d530b742355e4f1698f8d1b7b5e920 Mon Sep 17 00:00:00 2001
From: "Fedorov, Andrey" <andrey.fedorov@intel.com>
Date: Wed, 1 Sep 2021 12:08:51 +0300
Subject: [PATCH 05/16] applied clang-format

Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
---
 sycl/include/sycl/ext/oneapi/group_sort.hpp | 6 ++++--
 1 file changed, 4 insertions(+), 2 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/group_sort.hpp b/sycl/include/sycl/ext/oneapi/group_sort.hpp
index f3d794696be84..d2cf7a8ac4550 100644
--- a/sycl/include/sycl/ext/oneapi/group_sort.hpp
+++ b/sycl/include/sycl/ext/oneapi/group_sort.hpp
@@ -50,7 +50,8 @@ struct is_sorter_impl {
   template <typename G = Group>
   static decltype(
       std::integral_constant<bool, is_expected_return_type<G>::value &&
-                                       sycl::is_group_v<G>>{}) test(int);
+                                       sycl::is_group_v<G>>{})
+  test(int);
 
   template <typename = Group> static std::false_type test(...);
 };
@@ -68,7 +69,8 @@ struct is_sorter_impl<Sorter, Group, Ptr,
   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);
+                  sycl::detail::is_generic_group<G>{})
+  test(int);
 
   template <typename = Group> static std::false_type test(...);
 };

From 6b557617ed15943fb024dd216e53416f223eb970 Mon Sep 17 00:00:00 2001
From: "Fedorov, Andrey" <andrey.fedorov@intel.com>
Date: Thu, 2 Sep 2021 15:21:01 +0300
Subject: [PATCH 06/16] fixed when tests are passed

Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
---
 sycl/include/CL/sycl/detail/group_sort_impl.hpp       | 11 +++++------
 .../ext/oneapi/experimental/group_helpers_sorters.hpp |  1 +
 2 files changed, 6 insertions(+), 6 deletions(-)

diff --git a/sycl/include/CL/sycl/detail/group_sort_impl.hpp b/sycl/include/CL/sycl/detail/group_sort_impl.hpp
index 55ec9f5645237..4a9e0d1bfa1cb 100644
--- a/sycl/include/CL/sycl/detail/group_sort_impl.hpp
+++ b/sycl/include/CL/sycl/detail/group_sort_impl.hpp
@@ -24,9 +24,8 @@ namespace detail {
 // following to functions could be useless if std::[lower|upper]_bound worked
 // well
 template <typename Acc, typename Value, typename Compare>
-std::size_t my_lower_bound(Acc acc, const std::size_t first,
-                           const std::size_t last, const Value &value,
-                           Compare comp) {
+std::size_t my_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;
@@ -178,9 +177,9 @@ template <typename Group, typename Iter, typename Compare>
 void merge_sort(Group group, Iter first, const std::size_t n, Compare comp,
                 std::uint8_t *scratch) {
   using T = typename std::iterator_traits<Iter>::value_type;
-  auto id = cl::sycl::detail::Builder::getNDItem<Group::dimensions>()
-      const std::size_t idx = id.get_local_id();
-  const std::size_t local = group.get_local_range(0);
+  auto id = cl::sycl::detail::Builder::getNDItem<Group::dimensions>();
+  const std::size_t idx = id.get_local_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 firstly
diff --git a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
index e4fa348762838..0a8a0c5c8bd12 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
@@ -59,6 +59,7 @@ template <typename Compare = std::less<>> class default_sorter {
         memory_required<T>(
             sycl::memory_scope::work_group /*Group::fence_scope*/,
             range_size)) {
+      auto id = cl::sycl::detail::Builder::getNDItem<Group::dimensions>();
       uint32_t local_id = id.get_local_id();
       T *temp = reinterpret_cast<T *>(scratch);
       temp[local_id] = val;

From cb54021649940006bf55f826b6bf840f6dd31921 Mon Sep 17 00:00:00 2001
From: "Fedorov, Andrey" <andrey.fedorov@intel.com>
Date: Thu, 2 Sep 2021 17:33:42 +0300
Subject: [PATCH 07/16] found that fence_scope is implemented

Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
---
 .../ext/oneapi/experimental/group_helpers_sorters.hpp | 11 +++--------
 1 file changed, 3 insertions(+), 8 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
index 0a8a0c5c8bd12..8546976e6b77c 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
@@ -44,9 +44,7 @@ template <typename Compare = std::less<>> class default_sorter {
   void operator()(Group g, Ptr begin, Ptr end) {
 #ifdef __SYCL_DEVICE_ONLY__
     using T = typename std::iterator_traits<Ptr>::value_type;
-    if (scratch_size >=
-        memory_required<T>(
-            sycl::memory_scope::work_group /*Group::fence_scope*/, end - begin))
+    if (scratch_size >= memory_required<T>(Group::fence_scope, end - begin))
       cl::sycl::detail::merge_sort(g, begin, end - begin, comp, scratch);
       // TODO: it's better to add else branch
 #endif
@@ -54,11 +52,8 @@ template <typename Compare = std::less<>> class default_sorter {
 
   template <typename Group, typename T> T operator()(Group g, T val) {
 #ifdef __SYCL_DEVICE_ONLY__
-    auto range_size = g.get_local_range(0);
-    if (scratch_size >=
-        memory_required<T>(
-            sycl::memory_scope::work_group /*Group::fence_scope*/,
-            range_size)) {
+    auto range_size = g.get_local_range().size();
+    if (scratch_size >= memory_required<T>(Group::fence_scope, range_size)) {
       auto id = cl::sycl::detail::Builder::getNDItem<Group::dimensions>();
       uint32_t local_id = id.get_local_id();
       T *temp = reinterpret_cast<T *>(scratch);

From 68f2d8e35f1d05c32176d1a56a45625a006114dd Mon Sep 17 00:00:00 2001
From: "Fedorov, Andrey" <andrey.fedorov@intel.com>
Date: Fri, 3 Sep 2021 15:09:07 +0300
Subject: [PATCH 08/16] applied comments from reviewers

Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
---
 .../CL/sycl/detail/group_sort_impl.hpp        | 47 +++++++-------
 .../CL/sycl/detail/group_sort_util.hpp        | 26 --------
 .../experimental/group_helpers_sorters.hpp    | 25 ++++++--
 sycl/include/sycl/ext/oneapi/group_sort.hpp   | 61 ++++++-------------
 4 files changed, 61 insertions(+), 98 deletions(-)
 delete mode 100644 sycl/include/CL/sycl/detail/group_sort_util.hpp

diff --git a/sycl/include/CL/sycl/detail/group_sort_impl.hpp b/sycl/include/CL/sycl/detail/group_sort_impl.hpp
index 4a9e0d1bfa1cb..b79a1d1ae93ad 100644
--- a/sycl/include/CL/sycl/detail/group_sort_impl.hpp
+++ b/sycl/include/CL/sycl/detail/group_sort_impl.hpp
@@ -13,7 +13,6 @@
 #include <CL/sycl/detail/helpers.hpp>
 
 #ifdef __SYCL_DEVICE_ONLY__
-#include "group_sort_util.hpp"
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
@@ -21,11 +20,11 @@ namespace detail {
 
 // ---- merge sort implementation
 
-// following to functions could be useless if std::[lower|upper]_bound worked
+// following two functions could be useless if std::[lower|upper]_bound worked
 // well
 template <typename Acc, typename Value, typename Compare>
-std::size_t my_lower_bound(Acc acc, std::size_t first, std::size_t last,
-                           const Value &value, Compare comp) {
+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;
@@ -42,18 +41,18 @@ std::size_t my_lower_bound(Acc acc, std::size_t first, std::size_t last,
 }
 
 template <typename Acc, typename Value, typename Compare>
-std::size_t my_upper_bound(Acc acc, const std::size_t first,
-                           const std::size_t last, const Value &value,
-                           Compare comp) {
-  return my_lower_bound(acc, first, last, value,
-                        [comp](auto x, auto y) { return !comp(y, x); });
+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 my_swap(T &a, T &b) { std::swap(a, b); }
+template <typename T> void swap_tuples(T &a, T &b) { std::swap(a, b); }
 
 template <template <typename...> class TupleLike, typename T1, typename T2>
-void my_swap(TupleLike<T1, T2> &&a, TupleLike<T1, T2> &&b) {
+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));
 }
@@ -86,7 +85,7 @@ void merge(const std::size_t offset, InAcc &in_acc1, OutAcc &out_acc1,
     // 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 =
-        my_lower_bound(in_acc1, start_2, end_2, local_l_item_1, comp);
+        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;
 
@@ -96,8 +95,8 @@ void merge(const std::size_t offset, InAcc &in_acc1, OutAcc &out_acc1,
     // 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 = my_lower_bound(in_acc1, l_search_bound_2, end_2,
-                                        local_r_item_1, comp);
+      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;
 
@@ -110,8 +109,8 @@ void merge(const std::size_t offset, InAcc &in_acc1, OutAcc &out_acc1,
       // we shouldn't seek in whole 2nd sequence. Just for the part where the
       // 1st sequence should be
       l_search_bound_2 =
-          my_lower_bound(in_acc1, l_search_bound_2, r_search_bound_2,
-                         intermediate_item_1, comp);
+          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;
 
@@ -124,7 +123,7 @@ void merge(const std::size_t offset, InAcc &in_acc1, OutAcc &out_acc1,
     // 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 =
-        my_upper_bound(in_acc1, start_1, end_1, local_l_item_2, comp);
+        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;
 
@@ -134,8 +133,8 @@ void merge(const std::size_t offset, InAcc &in_acc1, OutAcc &out_acc1,
     // 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 = my_upper_bound(in_acc1, l_search_bound_1, end_1,
-                                        local_r_item_2, comp);
+      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;
 
@@ -148,8 +147,8 @@ void merge(const std::size_t offset, InAcc &in_acc1, OutAcc &out_acc1,
       // we shouldn't seek in whole 1st sequence. Just for the part where the
       // 2nd sequence should be
       l_search_bound_1 =
-          my_upper_bound(in_acc1, l_search_bound_1, r_search_bound_1,
-                         intermediate_item_2, comp);
+          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;
 
@@ -166,7 +165,7 @@ void bubble_sort(Iter first, const std::size_t begin, const std::size_t end,
       // Handle intermediate items
       for (std::size_t idx = i + 1; idx < end; ++idx) {
         if (comp(first[idx], first[i])) {
-          my_swap(first[i], first[idx]);
+          detail::swap_tuples(first[i], first[idx]);
         }
       }
     }
@@ -177,12 +176,12 @@ template <typename Group, typename Iter, typename Compare>
 void merge_sort(Group group, Iter first, const std::size_t n, Compare comp,
                 std::uint8_t *scratch) {
   using T = typename std::iterator_traits<Iter>::value_type;
-  auto id = cl::sycl::detail::Builder::getNDItem<Group::dimensions>();
+  auto id = sycl::detail::Builder::getNDItem<Group::dimensions>();
   const std::size_t idx = id.get_local_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 firstly
+  // we need to sort within work item first
   bubble_sort(first, idx * chunk, sycl::min((idx + 1) * chunk, n), comp);
   id.barrier();
 
diff --git a/sycl/include/CL/sycl/detail/group_sort_util.hpp b/sycl/include/CL/sycl/detail/group_sort_util.hpp
deleted file mode 100644
index 9765088f6597f..0000000000000
--- a/sycl/include/CL/sycl/detail/group_sort_util.hpp
+++ /dev/null
@@ -1,26 +0,0 @@
-//==------------ group_sort_util.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 utilities that are used by group sorting algorithms
-//
-
-#pragma once
-
-__SYCL_INLINE_NAMESPACE(cl) {
-namespace sycl {
-namespace detail {
-
-#if __cplusplus < 201703
-template <typename... _Ts> struct make_void_type { using type = void; };
-
-template <typename... _Ts>
-using void_type = typename make_void_type<_Ts...>::type;
-#endif
-
-} // namespace detail
-} // namespace sycl
-} // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
index 8546976e6b77c..5550ce4166d6c 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
@@ -41,12 +41,19 @@ template <typename Compare = std::less<>> class default_sorter {
       : comp(comp_), scratch(scratch_.data()), scratch_size(scratch_.size()) {}
 
   template <typename Group, typename Ptr>
-  void operator()(Group g, Ptr begin, Ptr end) {
+  void operator()(Group g, Ptr first, Ptr last) {
 #ifdef __SYCL_DEVICE_ONLY__
     using T = typename std::iterator_traits<Ptr>::value_type;
-    if (scratch_size >= memory_required<T>(Group::fence_scope, end - begin))
-      cl::sycl::detail::merge_sort(g, begin, end - begin, comp, scratch);
+    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 runtime_error(
+        "default_sorter constructor is not supported on host device.",
+        PI_INVALID_DEVICE);
 #endif
   }
 
@@ -54,15 +61,21 @@ template <typename Compare = std::less<>> class default_sorter {
 #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 = cl::sycl::detail::Builder::getNDItem<Group::dimensions>();
+      auto id = sycl::detail::Builder::getNDItem<Group::dimensions>();
       uint32_t local_id = id.get_local_id();
       T *temp = reinterpret_cast<T *>(scratch);
       temp[local_id] = val;
-      cl::sycl::detail::merge_sort(g, temp, range_size, comp,
-                                   scratch + range_size * sizeof(T));
+      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;
+    (void)val;
+    throw runtime_error(
+        "default_sorter operator() is not supported on host device.",
+        PI_INVALID_DEVICE);
 #endif
     return val;
   }
diff --git a/sycl/include/sycl/ext/oneapi/group_sort.hpp b/sycl/include/sycl/ext/oneapi/group_sort.hpp
index d2cf7a8ac4550..4a1d6a8068726 100644
--- a/sycl/include/sycl/ext/oneapi/group_sort.hpp
+++ b/sycl/include/sycl/ext/oneapi/group_sort.hpp
@@ -10,7 +10,6 @@
 
 #include <CL/sycl/detail/defines_elementary.hpp>
 #include <CL/sycl/detail/group_sort_impl.hpp>
-#include <CL/sycl/detail/group_sort_util.hpp>
 #include <CL/sycl/detail/type_traits.hpp>
 #include <type_traits>
 
@@ -29,14 +28,8 @@ namespace detail {
 template <typename T, typename = void> struct has_difference_type {};
 
 template <typename T>
-struct has_difference_type<T,
-#if __cplusplus < 201703
-                           cl::sycl::detail::void_type<
-#else
-                           std::void_t<
-#endif
-                               typename T::difference_type>> : std::true_type {
-};
+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 {};
 
@@ -59,13 +52,9 @@ struct is_sorter_impl {
 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,
-#if __cplusplus < 201703
-                      cl::sycl::detail::void_type<
-#else
-                      std::void_t<
-#endif
-                          typename has_difference_type<Ptr>::type>> {
+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>()),
@@ -95,25 +84,19 @@ sort_over_group(Group group, T value, Sorter sorter) {
 
 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(
-    cl::sycl::ext::oneapi::experimental::group_with_scratchpad<Group, Extent>
-        exec,
-    T value, Compare comp) {
+sort_over_group(experimental::group_with_scratchpad<Group, Extent> exec,
+                T value, Compare comp) {
   return sort_over_group(
       exec.get_group(), value,
-      cl::sycl::ext::oneapi::experimental::default_sorter<Compare>(
-          exec.get_memory(), comp));
+      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(
-    cl::sycl::ext::oneapi::experimental::group_with_scratchpad<Group, Extent>
-        exec,
-    T value) {
-  return sort_over_group(
-      exec.get_group(), value,
-      cl::sycl::ext::oneapi::experimental::default_sorter<>(exec.get_memory()));
+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
@@ -133,24 +116,18 @@ joint_sort(Group group, Iter first, Iter last, Sorter sorter) {
 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(
-    cl::sycl::ext::oneapi::experimental::group_with_scratchpad<Group, Extent>
-        exec,
-    Iter first, Iter last, Compare comp) {
+joint_sort(experimental::group_with_scratchpad<Group, Extent> exec, Iter first,
+           Iter last, Compare comp) {
   joint_sort(exec.get_group(), first, last,
-             cl::sycl::ext::oneapi::experimental::default_sorter<Compare>(
-                 exec.get_memory(), comp));
+             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(
-    cl::sycl::ext::oneapi::experimental::group_with_scratchpad<Group, Extent>
-        exec,
-    Iter first, Iter last) {
-  joint_sort(
-      exec.get_group(), first, last,
-      cl::sycl::ext::oneapi::experimental::default_sorter<>(exec.get_memory()));
+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

From 339d043651ba8a65416dd70252bd3a6eab9a798d Mon Sep 17 00:00:00 2001
From: "Fedorov, Andrey" <andrey.fedorov@intel.com>
Date: Fri, 3 Sep 2021 15:39:05 +0300
Subject: [PATCH 09/16] fix for get_pointer usage

Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
---
 sycl/include/CL/sycl/detail/group_sort_impl.hpp       | 11 ++++++++++-
 .../ext/oneapi/experimental/group_helpers_sorters.hpp |  2 +-
 sycl/include/sycl/ext/oneapi/group_sort.hpp           |  6 ++----
 3 files changed, 13 insertions(+), 6 deletions(-)

diff --git a/sycl/include/CL/sycl/detail/group_sort_impl.hpp b/sycl/include/CL/sycl/detail/group_sort_impl.hpp
index b79a1d1ae93ad..0924efc8ff9a5 100644
--- a/sycl/include/CL/sycl/detail/group_sort_impl.hpp
+++ b/sycl/include/CL/sycl/detail/group_sort_impl.hpp
@@ -57,6 +57,15 @@ void swap_tuples(TupleLike<T1, T2> &&a, TupleLike<T1, T2> &&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;
+};
+
 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,
@@ -175,7 +184,7 @@ void bubble_sort(Iter first, const std::size_t begin, const std::size_t end,
 template <typename Group, typename Iter, typename Compare>
 void merge_sort(Group group, Iter first, const std::size_t n, Compare comp,
                 std::uint8_t *scratch) {
-  using T = typename std::iterator_traits<Iter>::value_type;
+  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 local = group.get_local_range().size();
diff --git a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
index 5550ce4166d6c..4df655861fb1d 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
@@ -43,7 +43,7 @@ template <typename Compare = std::less<>> class default_sorter {
   template <typename Group, typename Ptr>
   void operator()(Group g, Ptr first, Ptr last) {
 #ifdef __SYCL_DEVICE_ONLY__
-    using T = typename std::iterator_traits<Ptr>::value_type;
+    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
diff --git a/sycl/include/sycl/ext/oneapi/group_sort.hpp b/sycl/include/sycl/ext/oneapi/group_sort.hpp
index 4a1d6a8068726..8d927d479a32a 100644
--- a/sycl/include/sycl/ext/oneapi/group_sort.hpp
+++ b/sycl/include/sycl/ext/oneapi/group_sort.hpp
@@ -43,8 +43,7 @@ struct is_sorter_impl {
   template <typename G = Group>
   static decltype(
       std::integral_constant<bool, is_expected_return_type<G>::value &&
-                                       sycl::is_group_v<G>>{})
-  test(int);
+                                       sycl::is_group_v<G>>{}) test(int);
 
   template <typename = Group> static std::false_type test(...);
 };
@@ -58,8 +57,7 @@ struct is_sorter_impl<
   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);
+                  sycl::detail::is_generic_group<G>{}) test(int);
 
   template <typename = Group> static std::false_type test(...);
 };

From eddd7a31215c37532adf2098b28864d8f4f69857 Mon Sep 17 00:00:00 2001
From: "Fedorov, Andrey" <andrey.fedorov@intel.com>
Date: Fri, 3 Sep 2021 15:42:10 +0300
Subject: [PATCH 10/16] align with clang-format from infrastructure

Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
---
 sycl/include/sycl/ext/oneapi/group_sort.hpp | 6 ++++--
 1 file changed, 4 insertions(+), 2 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/group_sort.hpp b/sycl/include/sycl/ext/oneapi/group_sort.hpp
index 8d927d479a32a..4a1d6a8068726 100644
--- a/sycl/include/sycl/ext/oneapi/group_sort.hpp
+++ b/sycl/include/sycl/ext/oneapi/group_sort.hpp
@@ -43,7 +43,8 @@ struct is_sorter_impl {
   template <typename G = Group>
   static decltype(
       std::integral_constant<bool, is_expected_return_type<G>::value &&
-                                       sycl::is_group_v<G>>{}) test(int);
+                                       sycl::is_group_v<G>>{})
+  test(int);
 
   template <typename = Group> static std::false_type test(...);
 };
@@ -57,7 +58,8 @@ struct is_sorter_impl<
   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);
+                  sycl::detail::is_generic_group<G>{})
+  test(int);
 
   template <typename = Group> static std::false_type test(...);
 };

From b4aa9c3e1149dc439b0cf1e2919a062b7ce8ac00 Mon Sep 17 00:00:00 2001
From: "Fedorov, Andrey" <andrey.fedorov@intel.com>
Date: Mon, 6 Sep 2021 12:32:04 +0300
Subject: [PATCH 11/16] apply some comments from review

Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
---
 .../SYCL_INTEL_group_sort.asciidoc            | 24 +++++++++----------
 .../CL/sycl/detail/group_sort_impl.hpp        |  2 +-
 .../experimental/group_helpers_sorters.hpp    | 22 ++++++++---------
 3 files changed, 23 insertions(+), 25 deletions(-)

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
index 0924efc8ff9a5..5a236d0a7739a 100644
--- a/sycl/include/CL/sycl/detail/group_sort_impl.hpp
+++ b/sycl/include/CL/sycl/detail/group_sort_impl.hpp
@@ -183,7 +183,7 @@ void bubble_sort(Iter first, const std::size_t begin, const std::size_t end,
 
 template <typename Group, typename Iter, typename Compare>
 void merge_sort(Group group, Iter first, const std::size_t n, Compare comp,
-                std::uint8_t *scratch) {
+                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();
diff --git a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
index 4df655861fb1d..d697d9233149f 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
@@ -19,24 +19,24 @@ namespace experimental {
 // ---- group helpers
 template <typename Group, std::size_t Extent> class group_with_scratchpad {
   Group g;
-  sycl::span<std::uint8_t, Extent> scratch;
+  sycl::span<std::byte, Extent> scratch;
 
 public:
-  group_with_scratchpad(Group g_, sycl::span<std::uint8_t, Extent> scratch_)
+  group_with_scratchpad(Group g_, sycl::span<std::byte, Extent> scratch_)
       : g(g_), scratch(scratch_) {}
   Group get_group() const { return g; }
-  sycl::span<std::uint8_t, Extent> get_memory() const { return scratch; }
+  sycl::span<std::byte, Extent> get_memory() const { return scratch; }
 };
 
 // ---- sorters
 template <typename Compare = std::less<>> class default_sorter {
   Compare comp;
-  std::uint8_t *scratch;
+  std::byte *scratch;
   std::size_t scratch_size;
 
 public:
   template <std::size_t Extent>
-  default_sorter(sycl::span<std::uint8_t, Extent> scratch_,
+  default_sorter(sycl::span<std::byte, Extent> scratch_,
                  Compare comp_ = Compare())
       : comp(comp_), scratch(scratch_.data()), scratch_size(scratch_.size()) {}
 
@@ -51,9 +51,9 @@ template <typename Compare = std::less<>> class default_sorter {
     (void)g;
     (void)first;
     (void)last;
-    throw runtime_error(
-        "default_sorter constructor is not supported on host device.",
-        PI_INVALID_DEVICE);
+    throw sycl::exception(std::error_code(PI_INVALID_DEVICE, sycl::sycl_category()),
+        "default_sorter constructor is not supported on host device.");
+
 #endif
   }
 
@@ -72,10 +72,8 @@ template <typename Compare = std::less<>> class default_sorter {
     // TODO: it's better to add else branch
 #else
     (void)g;
-    (void)val;
-    throw runtime_error(
-        "default_sorter operator() is not supported on host device.",
-        PI_INVALID_DEVICE);
+    throw sycl::exception(std::error_code(PI_INVALID_DEVICE, sycl::sycl_category()),
+        "default_sorter operator() is not supported on host device.");
 #endif
     return val;
   }

From 768d3a33b0f7e901be216d4688a640441112b3a4 Mon Sep 17 00:00:00 2001
From: "Fedorov, Andrey" <andrey.fedorov@intel.com>
Date: Mon, 6 Sep 2021 12:40:03 +0300
Subject: [PATCH 12/16] replace more runtime_errors with sycl::exception

Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
---
 sycl/include/sycl/ext/oneapi/group_sort.hpp | 13 +++++++++----
 1 file changed, 9 insertions(+), 4 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/group_sort.hpp b/sycl/include/sycl/ext/oneapi/group_sort.hpp
index 4a1d6a8068726..b165151de3335 100644
--- a/sycl/include/sycl/ext/oneapi/group_sort.hpp
+++ b/sycl/include/sycl/ext/oneapi/group_sort.hpp
@@ -77,8 +77,10 @@ sort_over_group(Group group, T value, Sorter sorter) {
   return sorter(group, value);
 #else
   (void)group;
-  throw runtime_error("Group algorithms are not supported on host device.",
-                      PI_INVALID_DEVICE);
+  (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
 }
 
@@ -108,8 +110,11 @@ joint_sort(Group group, Iter first, Iter last, Sorter sorter) {
   sorter(group, first, last);
 #else
   (void)group;
-  throw runtime_error("Group algorithms are not supported on host device.",
-                      PI_INVALID_DEVICE);
+  (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
 }
 

From 972f6efcb2179a462b076f337343caaeaa812ccf Mon Sep 17 00:00:00 2001
From: "Fedorov, Andrey" <andrey.fedorov@intel.com>
Date: Mon, 6 Sep 2021 13:26:53 +0300
Subject: [PATCH 13/16] add placement new instead of assignment

Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
---
 .../CL/sycl/detail/group_sort_impl.hpp        | 42 +++++++++++++++----
 .../experimental/group_helpers_sorters.hpp    | 15 +++----
 sycl/include/sycl/ext/oneapi/group_sort.hpp   | 10 +++--
 3 files changed, 47 insertions(+), 20 deletions(-)

diff --git a/sycl/include/CL/sycl/detail/group_sort_impl.hpp b/sycl/include/CL/sycl/detail/group_sort_impl.hpp
index 5a236d0a7739a..af1cb03a1569a 100644
--- a/sycl/include/CL/sycl/detail/group_sort_impl.hpp
+++ b/sycl/include/CL/sycl/detail/group_sort_impl.hpp
@@ -66,11 +66,22 @@ 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) {
+           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 =
@@ -98,7 +109,9 @@ void merge(const std::size_t offset, InAcc &in_acc1, OutAcc &out_acc1,
     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;
 
-    out_acc1[start_out + l_shift_1 + l_shift_2] = local_l_item_1;
+    // out_acc1[start_out + l_shift_1 + l_shift_2] = local_l_item_1;
+    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
@@ -109,7 +122,9 @@ void merge(const std::size_t offset, InAcc &in_acc1, OutAcc &out_acc1,
       const auto r_shift_1 = local_end_1 - 1 - start_1;
       const auto r_shift_2 = r_search_bound_2 - start_2;
 
-      out_acc1[start_out + r_shift_1 + r_shift_2] = local_r_item_1;
+      // out_acc1[start_out + r_shift_1 + r_shift_2] = local_r_item_1;
+      set_value(out_acc1, start_out + r_shift_1 + r_shift_2, local_r_item_1,
+                is_first);
     }
 
     // Handle intermediate items
@@ -123,7 +138,8 @@ void merge(const std::size_t offset, InAcc &in_acc1, OutAcc &out_acc1,
       const std::size_t shift_1 = idx - start_1;
       const std::size_t shift_2 = l_search_bound_2 - start_2;
 
-      out_acc1[start_out + shift_1 + shift_2] = intermediate_item_1;
+      set_value(out_acc1, start_out + shift_1 + shift_2, intermediate_item_1,
+                is_first);
     }
   }
   // Process 2nd sequence
@@ -136,7 +152,8 @@ void merge(const std::size_t offset, InAcc &in_acc1, OutAcc &out_acc1,
     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;
 
-    out_acc1[start_out + l_shift_1 + l_shift_2] = local_l_item_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
@@ -147,7 +164,8 @@ void merge(const std::size_t offset, InAcc &in_acc1, OutAcc &out_acc1,
       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;
 
-      out_acc1[start_out + r_shift_1 + r_shift_2] = local_r_item_2;
+      set_value(out_acc1, start_out + r_shift_1 + r_shift_2, local_r_item_2,
+                is_first);
     }
 
     // Handle intermediate items
@@ -161,7 +179,8 @@ void merge(const std::size_t offset, InAcc &in_acc1, OutAcc &out_acc1,
       const std::size_t shift_1 = l_search_bound_1 - start_1;
       const std::size_t shift_2 = idx - start_2;
 
-      out_acc1[start_out + shift_1 + shift_2] = intermediate_item_2;
+      set_value(out_acc1, start_out + shift_1 + shift_2, intermediate_item_2,
+                is_first);
     }
   }
 }
@@ -196,6 +215,7 @@ void merge_sort(Group group, Iter first, const std::size_t n, Compare comp,
 
   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 =
@@ -205,14 +225,18 @@ void merge_sort(Group group, Iter first, const std::size_t n, Compare comp,
     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);
+      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);
+      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
diff --git a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
index d697d9233149f..a498f5495a88e 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
@@ -51,9 +51,9 @@ template <typename Compare = std::less<>> class default_sorter {
     (void)g;
     (void)first;
     (void)last;
-    throw sycl::exception(std::error_code(PI_INVALID_DEVICE, sycl::sycl_category()),
+    throw sycl::exception(
+        std::error_code(PI_INVALID_DEVICE, sycl::sycl_category()),
         "default_sorter constructor is not supported on host device.");
-
 #endif
   }
 
@@ -64,7 +64,7 @@ template <typename Compare = std::less<>> class default_sorter {
       auto id = sycl::detail::Builder::getNDItem<Group::dimensions>();
       uint32_t local_id = id.get_local_id();
       T *temp = reinterpret_cast<T *>(scratch);
-      temp[local_id] = val;
+      ::new (temp + local_id) T(val);
       sycl::detail::merge_sort(g, temp, range_size, comp,
                                scratch + range_size * sizeof(T));
       val = temp[local_id];
@@ -72,22 +72,23 @@ template <typename Compare = std::less<>> class default_sorter {
     // TODO: it's better to add else branch
 #else
     (void)g;
-    throw sycl::exception(std::error_code(PI_INVALID_DEVICE, sycl::sycl_category()),
+    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 scope,
+  static constexpr std::size_t memory_required(sycl::memory_scope,
                                                std::size_t range_size) {
-    return range_size * sizeof(T);
+    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 * r.size() * sizeof(T);
+    return 2 * memory_required<T>(scope, r.size());
   }
 };
 
diff --git a/sycl/include/sycl/ext/oneapi/group_sort.hpp b/sycl/include/sycl/ext/oneapi/group_sort.hpp
index b165151de3335..1b14615f1a811 100644
--- a/sycl/include/sycl/ext/oneapi/group_sort.hpp
+++ b/sycl/include/sycl/ext/oneapi/group_sort.hpp
@@ -79,8 +79,9 @@ sort_over_group(Group group, T value, Sorter sorter) {
   (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.");
+  throw sycl::exception(
+      std::error_code(PI_INVALID_DEVICE, sycl::sycl_category()),
+      "Group algorithms are not supported on host device.");
 #endif
 }
 
@@ -113,8 +114,9 @@ joint_sort(Group group, Iter first, Iter last, Sorter sorter) {
   (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.");
+  throw sycl::exception(
+      std::error_code(PI_INVALID_DEVICE, sycl::sycl_category()),
+      "Group algorithms are not supported on host device.");
 #endif
 }
 

From 8a07d5920761699d8776c2302bd1b30d7a3b5ac3 Mon Sep 17 00:00:00 2001
From: "Fedorov, Andrey" <andrey.fedorov@intel.com>
Date: Mon, 6 Sep 2021 17:18:57 +0300
Subject: [PATCH 14/16] remove comments

Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
---
 sycl/include/CL/sycl/detail/group_sort_impl.hpp | 2 --
 1 file changed, 2 deletions(-)

diff --git a/sycl/include/CL/sycl/detail/group_sort_impl.hpp b/sycl/include/CL/sycl/detail/group_sort_impl.hpp
index af1cb03a1569a..cfc5a26a7c4eb 100644
--- a/sycl/include/CL/sycl/detail/group_sort_impl.hpp
+++ b/sycl/include/CL/sycl/detail/group_sort_impl.hpp
@@ -109,7 +109,6 @@ void merge(const std::size_t offset, InAcc &in_acc1, OutAcc &out_acc1,
     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;
 
-    // out_acc1[start_out + l_shift_1 + l_shift_2] = local_l_item_1;
     set_value(out_acc1, start_out + l_shift_1 + l_shift_2, local_l_item_1,
               is_first);
 
@@ -122,7 +121,6 @@ void merge(const std::size_t offset, InAcc &in_acc1, OutAcc &out_acc1,
       const auto r_shift_1 = local_end_1 - 1 - start_1;
       const auto r_shift_2 = r_search_bound_2 - start_2;
 
-      // out_acc1[start_out + r_shift_1 + r_shift_2] = local_r_item_1;
       set_value(out_acc1, start_out + r_shift_1 + r_shift_2, local_r_item_1,
                 is_first);
     }

From b6ae2e5bd5e4e86066ae6cfe20e0df3c113f092d Mon Sep 17 00:00:00 2001
From: "Fedorov, Andrey" <andrey.fedorov@intel.com>
Date: Thu, 9 Sep 2021 16:12:36 +0300
Subject: [PATCH 15/16] add more macros

Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
---
 sycl/include/CL/sycl/detail/group_sort_impl.hpp                 | 2 ++
 .../sycl/ext/oneapi/experimental/group_helpers_sorters.hpp      | 2 ++
 sycl/include/sycl/ext/oneapi/group_sort.hpp                     | 2 ++
 3 files changed, 6 insertions(+)

diff --git a/sycl/include/CL/sycl/detail/group_sort_impl.hpp b/sycl/include/CL/sycl/detail/group_sort_impl.hpp
index cfc5a26a7c4eb..03cc65f3942ed 100644
--- a/sycl/include/CL/sycl/detail/group_sort_impl.hpp
+++ b/sycl/include/CL/sycl/detail/group_sort_impl.hpp
@@ -10,6 +10,7 @@
 
 #pragma once
 
+#if __cplusplus >= 201703L
 #include <CL/sycl/detail/helpers.hpp>
 
 #ifdef __SYCL_DEVICE_ONLY__
@@ -252,3 +253,4 @@ void merge_sort(Group group, Iter first, const std::size_t n, Compare comp,
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
 #endif
+#endif // __cplusplus >=201703L
\ No newline at end of file
diff --git a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
index a498f5495a88e..392e9043051d3 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
@@ -8,6 +8,7 @@
 
 #pragma once
 
+#if __cplusplus >= 201703L
 #include <CL/sycl/detail/group_sort_impl.hpp>
 
 __SYCL_INLINE_NAMESPACE(cl) {
@@ -97,3 +98,4 @@ template <typename Compare = std::less<>> class default_sorter {
 } // 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
index 1b14615f1a811..03f1e6500966c 100644
--- a/sycl/include/sycl/ext/oneapi/group_sort.hpp
+++ b/sycl/include/sycl/ext/oneapi/group_sort.hpp
@@ -8,6 +8,7 @@
 
 #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>
@@ -141,3 +142,4 @@ joint_sort(experimental::group_with_scratchpad<Group, Extent> exec, Iter first,
 } // namespace ext
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
+#endif // __cplusplus >=201703L

From e162168c3ff77c4816bd1f47ae677db894548f8e Mon Sep 17 00:00:00 2001
From: "Fedorov, Andrey" <andrey.fedorov@intel.com>
Date: Wed, 15 Sep 2021 11:04:03 +0300
Subject: [PATCH 16/16] made changes for multi-dimensional groups

Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
---
 sycl/include/CL/sycl/detail/group_sort_impl.hpp        | 10 +++++-----
 .../ext/oneapi/experimental/group_helpers_sorters.hpp  |  2 +-
 2 files changed, 6 insertions(+), 6 deletions(-)

diff --git a/sycl/include/CL/sycl/detail/group_sort_impl.hpp b/sycl/include/CL/sycl/detail/group_sort_impl.hpp
index 03cc65f3942ed..e80f2fd131e77 100644
--- a/sycl/include/CL/sycl/detail/group_sort_impl.hpp
+++ b/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;
 
diff --git a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
index 392e9043051d3..932d0a29fb888 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp
+++ b/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,