DPC++ Runtime
Runtime libraries for oneAPI DPC++
group_sort.hpp
Go to the documentation of this file.
1 //==--------- group_sort.hpp --- SYCL extension group sorting algorithm-----==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 #pragma once
10 
11 #if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
15 #include <type_traits>
16 
18 
20 namespace sycl {
21 namespace ext {
22 namespace oneapi {
23 namespace experimental {
24 namespace detail {
25 
26 // ---- traits
27 template <typename T, typename = void> struct has_difference_type {};
28 
29 template <typename T>
30 struct has_difference_type<T, sycl::detail::void_t<typename T::difference_type>>
31  : std::true_type {};
32 
33 template <typename T> struct has_difference_type<T *> : std::true_type {};
34 
35 template <typename Sorter, typename Group, typename Val, typename = void>
36 struct is_sorter_impl {
37  template <typename G>
38  using is_expected_return_type =
39  typename std::is_same<Val, decltype(std::declval<Sorter>()(
40  std::declval<G>(), std::declval<Val>()))>;
41 
42  template <typename G = Group>
43  static decltype(
44  std::integral_constant<bool, is_expected_return_type<G>::value &&
45  sycl::is_group_v<G>>{}) test(int);
46 
47  template <typename = Group> static std::false_type test(...);
48 };
49 
50 template <typename Sorter, typename Group,
51  typename Ptr> // multi_ptr has difference_type and don't have other
52  // iterator's fields
53 struct is_sorter_impl<
54  Sorter, Group, Ptr,
55  sycl::detail::void_t<typename has_difference_type<Ptr>::type>> {
56  template <typename G = Group>
57  static decltype(std::declval<Sorter>()(std::declval<G>(), std::declval<Ptr>(),
58  std::declval<Ptr>()),
59  sycl::detail::is_generic_group<G>{}) test(int);
60 
61  template <typename = Group> static std::false_type test(...);
62 };
63 
64 template <typename Sorter, typename Group, typename ValOrPtr>
65 struct is_sorter : decltype(is_sorter_impl<Sorter, Group, ValOrPtr>::test(0)) {
66 };
67 } // namespace detail
68 
69 // ---- sort_over_group
70 template <typename Group, typename T, typename Sorter>
71 typename std::enable_if<detail::is_sorter<Sorter, Group, T>::value, T>::type
72 sort_over_group(Group group, T value, Sorter sorter) {
73 #ifdef __SYCL_DEVICE_ONLY__
74  return sorter(group, value);
75 #else
76  (void)group;
77  (void)value;
78  (void)sorter;
79  throw sycl::exception(
80  std::error_code(PI_INVALID_DEVICE, sycl::sycl_category()),
81  "Group algorithms are not supported on host device.");
82 #endif
83 }
84 
85 template <typename Group, typename T, typename Compare, std::size_t Extent>
86 typename std::enable_if<!detail::is_sorter<Compare, Group, T>::value, T>::type
87 sort_over_group(experimental::group_with_scratchpad<Group, Extent> exec,
88  T value, Compare comp) {
89  return sort_over_group(
90  exec.get_group(), value,
91  experimental::default_sorter<Compare>(exec.get_memory(), comp));
92 }
93 
94 template <typename Group, typename T, std::size_t Extent>
95 typename std::enable_if<sycl::is_group_v<std::decay_t<Group>>, T>::type
96 sort_over_group(experimental::group_with_scratchpad<Group, Extent> exec,
97  T value) {
98  return sort_over_group(exec.get_group(), value,
99  experimental::default_sorter<>(exec.get_memory()));
100 }
101 
102 // ---- joint_sort
103 template <typename Group, typename Iter, typename Sorter>
104 typename std::enable_if<detail::is_sorter<Sorter, Group, Iter>::value,
105  void>::type
106 joint_sort(Group group, Iter first, Iter last, Sorter sorter) {
107 #ifdef __SYCL_DEVICE_ONLY__
108  sorter(group, first, last);
109 #else
110  (void)group;
111  (void)first;
112  (void)last;
113  (void)sorter;
114  throw sycl::exception(
115  std::error_code(PI_INVALID_DEVICE, sycl::sycl_category()),
116  "Group algorithms are not supported on host device.");
117 #endif
118 }
119 
120 template <typename Group, typename Iter, typename Compare, std::size_t Extent>
121 typename std::enable_if<!detail::is_sorter<Compare, Group, Iter>::value,
122  void>::type
123 joint_sort(experimental::group_with_scratchpad<Group, Extent> exec, Iter first,
124  Iter last, Compare comp) {
125  joint_sort(exec.get_group(), first, last,
126  experimental::default_sorter<Compare>(exec.get_memory(), comp));
127 }
128 
129 template <typename Group, typename Iter, std::size_t Extent>
130 typename std::enable_if<sycl::is_group_v<std::decay_t<Group>>, void>::type
131 joint_sort(experimental::group_with_scratchpad<Group, Extent> exec, Iter first,
132  Iter last) {
133  joint_sort(exec.get_group(), first, last,
134  experimental::default_sorter<>(exec.get_memory()));
135 }
136 
137 } // namespace experimental
138 } // namespace oneapi
139 } // namespace ext
140 } // namespace sycl
141 } // __SYCL_INLINE_NAMESPACE(cl)
142 #endif // __cplusplus >=201703L
T
cl::sycl::sycl_category
const std::error_category & sycl_category() noexcept
Definition: exception.cpp:116
type_traits.hpp
defines_elementary.hpp
sycl
Definition: invoke_simd.hpp:68
cl::sycl::detail::void_t
void void_t
Definition: stl_type_traits.hpp:42
group_helpers_sorters.hpp
group_sort_impl.hpp
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
PI_INVALID_DEVICE
@ PI_INVALID_DEVICE
Definition: pi.h:94
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12