| 951 | |
| 952 | template <class ForwardIt> |
| 953 | sycl::event |
| 954 | min_element(sycl::queue &q, util::allocation_group &scratch_allocations, |
| 955 | ForwardIt first, ForwardIt last, |
| 956 | std::pair<ForwardIt, typename |
| 957 | std::iterator_traits<ForwardIt>::value_type> *out, |
| 958 | const std::vector<sycl::event> &deps= {}) { |
| 959 | auto problem_size = std::distance(first, last); |
| 960 | if(problem_size == 0) |
| 961 | return sycl::event{}; |
| 962 | |
| 963 | using ValueT = typename std::iterator_traits<ForwardIt>::value_type; |
| 964 | using MinPair = std::pair<ForwardIt, ValueT>; |
| 965 | |
| 966 | auto kernel = [=](sycl::id<1> idx, auto& reducer) { |
| 967 | auto input = first; |
| 968 | std::advance(input, idx[0]); |
| 969 | MinPair p = std::make_pair(input, *input); |
| 970 | reducer.combine(p); |
| 971 | }; |
| 972 | |
| 973 | auto reduce = [first] (MinPair a, MinPair b) { |
| 974 | // Preserve strict total order over two equivalent |
| 975 | // pointers, i.e. return the element that appears |
| 976 | // in the sequence nearest to first. |
| 977 | if (!(a.second < b.second) && !(b.second < a.second)) { |
| 978 | if (std::distance(first, a.first) < std::distance(first, b.first)) |
| 979 | return a; |
| 980 | else |
| 981 | return b; |
| 982 | } |
| 983 | #if __cplusplus < 202002L |
| 984 | else if (a.second < b.second) |
| 985 | #else |
| 986 | else if (std::less{}(a.second, b.second)) |
| 987 | #endif |
| 988 | return a; |
| 989 | else |
| 990 | return b; |
| 991 | }; |
| 992 | |
| 993 | MinPair init = std::make_pair(first, *first); |
| 994 | |
| 995 | return detail::transform_reduce_impl(q, scratch_allocations, out, init, |
| 996 | problem_size, kernel, reduce, deps); |
| 997 | } |
| 998 | |
| 999 | template <class ForwardIt, class Compare> |
| 1000 | sycl::event |