| 1043 | |
| 1044 | template <class ForwardIt> |
| 1045 | sycl::event |
| 1046 | max_element(sycl::queue &q, util::allocation_group &scratch_allocations, |
| 1047 | ForwardIt first, ForwardIt last, |
| 1048 | std::pair<ForwardIt, typename |
| 1049 | std::iterator_traits<ForwardIt>::value_type> *out, |
| 1050 | const std::vector<sycl::event> &deps= {}) { |
| 1051 | auto problem_size = std::distance(first, last); |
| 1052 | if(problem_size == 0) |
| 1053 | return sycl::event{}; |
| 1054 | |
| 1055 | using ValueT = typename std::iterator_traits<ForwardIt>::value_type; |
| 1056 | using MaxPair = std::pair<ForwardIt, ValueT>; |
| 1057 | |
| 1058 | auto kernel = [=](sycl::id<1> idx, auto& reducer) { |
| 1059 | auto input = first; |
| 1060 | std::advance(input, idx[0]); |
| 1061 | MaxPair p = std::make_pair(input, *input); |
| 1062 | reducer.combine(p); |
| 1063 | }; |
| 1064 | |
| 1065 | auto reduce = [first] (MaxPair a, MaxPair b) { |
| 1066 | // Preserve strict total order over two equivalent |
| 1067 | // pointers, i.e. return the element that appears |
| 1068 | // in the sequence nearest to first. |
| 1069 | if (!(a.second < b.second) && !(b.second < a.second)) { |
| 1070 | if (std::distance(first, a.first) < std::distance(first, b.first)) |
| 1071 | return a; |
| 1072 | else |
| 1073 | return b; |
| 1074 | } |
| 1075 | #if __cplusplus < 202002L |
| 1076 | else if (a.second < b.second) |
| 1077 | #else |
| 1078 | else if (std::less{}(a.second, b.second)) |
| 1079 | #endif |
| 1080 | return b; |
| 1081 | else |
| 1082 | return a; |
| 1083 | }; |
| 1084 | |
| 1085 | MaxPair init = std::make_pair(first, *first); |
| 1086 | |
| 1087 | return detail::transform_reduce_impl(q, scratch_allocations, out, init, |
| 1088 | problem_size, kernel, reduce, deps); |
| 1089 | } |
| 1090 | |
| 1091 | template <class ForwardIt, class Compare> |
| 1092 | sycl::event |