Skip to content

Commit 37124f8

Browse files
normallytangentilluhad
normallytangent
and
illuhad
authored
[stdpar][algorithms] Add support for is_sorted and is_sorted_until (AdaptiveCpp#1781)
* [stdpar][algorithms] Add support for `is_sorted` and `is_sorted_until` * Do not test parallel policy --------- Co-authored-by: Aksel Alpay <[email protected]>
1 parent 56ca3c2 commit 37124f8

File tree

9 files changed

+699
-0
lines changed

9 files changed

+699
-0
lines changed

‎doc/algorithms.md‎

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -396,6 +396,38 @@ sycl::event sort(sycl::queue &q, RandomIt first, RandomIt last,
396396
Compare comp = std::less<>{},
397397
const std::vector<sycl::event>& deps = {});
398398

399+
/// The result of the operation will be stored in out.
400+
///
401+
/// out must point to device-accessible memory, and will be set to 0
402+
/// for a negative result, and 1 for a positive result.
403+
template <class ForwardIt>
404+
sycl::event is_sorted(sycl::queue &q, ForwardIt first, ForwardIt last,
405+
detail::early_exit_flag_t* out,
406+
const std::vector<sycl::event>& deps = {});
407+
408+
/// The result of the operation will be stored in out.
409+
///
410+
/// out must point to device-accessible memory, and will be set to 0
411+
/// for a negative result, and 1 for a positive result.
412+
template <class ForwardIt, class Compare>
413+
sycl::event is_sorted(sycl::queue &q, ForwardIt first, ForwardIt last,
414+
detail::early_exit_flag_t* out,
415+
Compare comp,
416+
const std::vector<sycl::event>& deps = {});
417+
418+
template<class ForwardIt>
419+
sycl::event is_sorted_until(sycl::queue &q, util::allocation_group &scratch_allocations,
420+
ForwardIt first, ForwardIt last,
421+
typename std::iterator_traits<ForwardIt>::difference_type *out,
422+
const std::vector<sycl::event>& deps = {});
423+
424+
template<class ForwardIt, class Compare>
425+
sycl::event is_sorted_until(sycl::queue &q, util::allocation_group &scratch_allocations,
426+
ForwardIt first, ForwardIt last,
427+
typename std::iterator_traits<ForwardIt>::difference_type *out,
428+
Compare comp,
429+
const std::vector<sycl::event>& deps = {});
430+
399431
template< class ForwardIt1, class ForwardIt2,
400432
class ForwardIt3, class Compare >
401433
sycl::event merge(sycl::queue& q,

‎doc/stdpar.md‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -59,6 +59,8 @@ Offloading is implemented for the following STL algorithms:
5959
|`equal` | |
6060
|`merge` | |
6161
|`sort` | may not scale optimally for large problems |
62+
|`is_sorted_until` | both overloads |
63+
|`is_sorted` | both overloads |
6264
|`inclusive_scan` | |
6365
|`exclusive_scan` | |
6466
|`transform_inclusive_scan` | |

‎include/hipSYCL/algorithms/algorithm.hpp‎

Lines changed: 117 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -970,6 +970,123 @@ sycl::event sort(sycl::queue &q, RandomIt first, RandomIt last,
970970
return sorting::bitonic_sort(q, first, last, comp, deps);
971971
}
972972

973+
template <class ForwardIt>
974+
sycl::event is_sorted(sycl::queue &q, ForwardIt first, ForwardIt last,
975+
detail::early_exit_flag_t* out,
976+
const std::vector<sycl::event>& deps = {}) {
977+
std::size_t problem_size = std::distance(first, last);
978+
if(problem_size == 0)
979+
return sycl::event{};
980+
981+
auto evt = detail::early_exit_for_each(q, problem_size, out,
982+
[=](sycl::id<1> idx) -> bool {
983+
auto it = first;
984+
std::advance(it, idx[0]);
985+
if(it != last-1) {
986+
auto next = std::next(it, 1);
987+
#if __cplusplus < 202002L
988+
return (*next < *it);
989+
#else
990+
return std::less{}(*next, *it);
991+
#endif
992+
}
993+
return false;
994+
}, deps);
995+
return q.single_task(evt, [=](){
996+
*out = static_cast<detail::early_exit_flag_t>(!(*out));
997+
});
998+
}
999+
1000+
template <class ForwardIt, class Compare>
1001+
sycl::event is_sorted(sycl::queue &q, ForwardIt first, ForwardIt last,
1002+
detail::early_exit_flag_t* out,
1003+
Compare comp,
1004+
const std::vector<sycl::event>& deps = {}) {
1005+
std::size_t problem_size = std::distance(first, last);
1006+
if(problem_size == 0)
1007+
return sycl::event{};
1008+
1009+
auto evt = detail::early_exit_for_each(q, problem_size, out,
1010+
[=](sycl::id<1> idx) -> bool {
1011+
auto it = std::next(first, idx[0]);
1012+
if(it != last-1) {
1013+
auto next = std::next(it, 1);
1014+
return comp(*next, *it);
1015+
}
1016+
return false;
1017+
}, deps);
1018+
return q.single_task(evt, [=](){
1019+
*out = static_cast<detail::early_exit_flag_t>(!(*out));
1020+
});
1021+
}
1022+
1023+
template<class ForwardIt>
1024+
sycl::event is_sorted_until(sycl::queue &q, util::allocation_group &scratch_allocations,
1025+
ForwardIt first, ForwardIt last,
1026+
typename std::iterator_traits<ForwardIt>::difference_type *out,
1027+
const std::vector<sycl::event>& deps = {}) {
1028+
if(first == last)
1029+
return sycl::event{};
1030+
1031+
using DiffT = typename std::iterator_traits<ForwardIt>::difference_type;
1032+
DiffT problem_size = std::distance(first, last);
1033+
1034+
auto transform = [=] (ForwardIt input) {
1035+
auto next = std::next(input, 1);
1036+
#if __cplusplus < 202002L
1037+
if((*next < *input))
1038+
#else
1039+
if(std::less{}(*next, *input))
1040+
#endif
1041+
return std::distance(first, input);
1042+
1043+
return problem_size;
1044+
};
1045+
1046+
auto kernel = [=](sycl::id<1> idx, auto& reducer) {
1047+
auto input = first;
1048+
std::advance(input, idx[0]);
1049+
reducer.combine(transform(input));
1050+
};
1051+
1052+
auto reduce = sycl::minimum<DiffT>{};
1053+
1054+
return detail::transform_reduce_impl(q, scratch_allocations, out, std::numeric_limits<DiffT>::max(),
1055+
problem_size, kernel, reduce, deps);
1056+
}
1057+
1058+
template<class ForwardIt, class Compare>
1059+
sycl::event is_sorted_until(sycl::queue &q, util::allocation_group &scratch_allocations,
1060+
ForwardIt first, ForwardIt last,
1061+
typename std::iterator_traits<ForwardIt>::difference_type *out,
1062+
Compare comp,
1063+
const std::vector<sycl::event>& deps = {}) {
1064+
if(first == last)
1065+
return sycl::event{};
1066+
1067+
using DiffT = typename std::iterator_traits<ForwardIt>::difference_type;
1068+
DiffT problem_size = std::distance(first, last);
1069+
1070+
auto transform = [=] (ForwardIt input) {
1071+
auto next = std::next(input, 1);
1072+
if(comp(*next, *input))
1073+
return std::distance(first, input);
1074+
1075+
return problem_size;
1076+
};
1077+
1078+
auto kernel = [=](sycl::id<1> idx, auto& reducer) {
1079+
auto input = first;
1080+
std::advance(input, idx[0]);
1081+
reducer.combine(transform(input));
1082+
};
1083+
1084+
auto reduce = sycl::minimum<DiffT>{};
1085+
1086+
return detail::transform_reduce_impl(q, scratch_allocations, out, std::numeric_limits<DiffT>::max(),
1087+
problem_size, kernel, reduce, deps);
1088+
}
1089+
9731090
template< class ForwardIt1, class ForwardIt2,
9741091
class ForwardIt3, class Compare >
9751092
sycl::event merge(sycl::queue& q,

‎include/hipSYCL/std/stdpar/detail/algorithm_fwd.hpp‎

Lines changed: 38 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -191,6 +191,25 @@ typename std::iterator_traits<ForwardIt>::difference_type
191191
count_if( hipsycl::stdpar::par_unseq, ForwardIt first, ForwardIt last,
192192
UnaryPredicate p );
193193

194+
template<class ForwardIt>
195+
HIPSYCL_STDPAR_ENTRYPOINT
196+
bool is_sorted(hipsycl::stdpar::par_unseq, ForwardIt first, ForwardIt last);
197+
198+
template<class ForwardIt, class Compare>
199+
HIPSYCL_STDPAR_ENTRYPOINT
200+
bool is_sorted(hipsycl::stdpar::par_unseq, ForwardIt first, ForwardIt last,
201+
Compare comp);
202+
203+
template<class ForwardIt>
204+
HIPSYCL_STDPAR_ENTRYPOINT
205+
ForwardIt is_sorted_until(hipsycl::stdpar::par_unseq, ForwardIt first,
206+
ForwardIt last);
207+
208+
template<class ForwardIt, class Compare>
209+
HIPSYCL_STDPAR_ENTRYPOINT
210+
ForwardIt is_sorted_until(hipsycl::stdpar::par_unseq, ForwardIt first,
211+
ForwardIt last, Compare comp);
212+
194213
template <class ForwardIt1, class ForwardIt2>
195214
HIPSYCL_STDPAR_ENTRYPOINT
196215
bool equal(hipsycl::stdpar::par_unseq, ForwardIt1 first1, ForwardIt1 last1,
@@ -238,6 +257,25 @@ typename std::iterator_traits<ForwardIt>::difference_type
238257
count_if( hipsycl::stdpar::par, ForwardIt first, ForwardIt last,
239258
UnaryPredicate p );
240259

260+
template<class ForwardIt>
261+
HIPSYCL_STDPAR_ENTRYPOINT
262+
bool is_sorted(hipsycl::stdpar::par, ForwardIt first, ForwardIt last);
263+
264+
template<class ForwardIt, class Compare>
265+
HIPSYCL_STDPAR_ENTRYPOINT
266+
bool is_sorted(hipsycl::stdpar::par, ForwardIt first, ForwardIt last,
267+
Compare comp);
268+
269+
template<class ForwardIt>
270+
HIPSYCL_STDPAR_ENTRYPOINT
271+
ForwardIt is_sorted_until(hipsycl::stdpar::par, ForwardIt first,
272+
ForwardIt last);
273+
274+
template<class ForwardIt, class Compare>
275+
HIPSYCL_STDPAR_ENTRYPOINT
276+
ForwardIt is_sorted_until(hipsycl::stdpar::par, ForwardIt first,
277+
ForwardIt last, Compare comp);
278+
241279
template <class ForwardIt1, class ForwardIt2, class T>
242280
HIPSYCL_STDPAR_ENTRYPOINT ForwardIt2 remove_copy(hipsycl::stdpar::par,
243281
ForwardIt1 first, ForwardIt1 last,

‎include/hipSYCL/std/stdpar/detail/offload_heuristic_db.hpp‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -63,6 +63,8 @@ struct count{};
6363
struct count_if{};
6464
struct equal {};
6565
struct sort {};
66+
struct is_sorted {};
67+
struct is_sorted_until {};
6668
struct merge {};
6769
struct inclusive_scan {};
6870
struct exclusive_scan {};

0 commit comments

Comments
 (0)