/* * Copyright 2008-2013 NVIDIA Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. * You may obtain a copy of the License at * * http://www.apache.org/licenses/LICENSE-2.0 * * Unless required by applicable law or agreed to in writing, software * distributed under the License is distributed on an "AS IS" BASIS, * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. * See the License for the specific language governing permissions and * limitations under the License. */ #pragma once #include #if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) # pragma GCC system_header #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) # pragma clang system_header #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) # pragma system_header #endif // no system header #include #include #include #include #include #include #include #include THRUST_NAMESPACE_BEGIN namespace system { namespace detail { namespace generic { template _CCCL_HOST_DEVICE void sort(thrust::execution_policy& exec, RandomAccessIterator first, RandomAccessIterator last) { typedef typename thrust::iterator_value::type value_type; thrust::sort(exec, first, last, thrust::less()); } // end sort() template _CCCL_HOST_DEVICE void sort(thrust::execution_policy& exec, RandomAccessIterator first, RandomAccessIterator last, StrictWeakOrdering comp) { // implement with stable_sort thrust::stable_sort(exec, first, last, comp); } // end sort() template _CCCL_HOST_DEVICE void sort_by_key( thrust::execution_policy& exec, RandomAccessIterator1 keys_first, RandomAccessIterator1 keys_last, RandomAccessIterator2 values_first) { typedef typename thrust::iterator_value::type value_type; thrust::sort_by_key(exec, keys_first, keys_last, values_first, thrust::less()); } // end sort_by_key() template _CCCL_HOST_DEVICE void sort_by_key( thrust::execution_policy& exec, RandomAccessIterator1 keys_first, RandomAccessIterator1 keys_last, RandomAccessIterator2 values_first, StrictWeakOrdering comp) { // implement with stable_sort_by_key thrust::stable_sort_by_key(exec, keys_first, keys_last, values_first, comp); } // end sort_by_key() template _CCCL_HOST_DEVICE void stable_sort(thrust::execution_policy& exec, RandomAccessIterator first, RandomAccessIterator last) { typedef typename thrust::iterator_value::type value_type; thrust::stable_sort(exec, first, last, thrust::less()); } // end stable_sort() template _CCCL_HOST_DEVICE void stable_sort_by_key( thrust::execution_policy& exec, RandomAccessIterator1 keys_first, RandomAccessIterator1 keys_last, RandomAccessIterator2 values_first) { typedef typename iterator_value::type value_type; thrust::stable_sort_by_key(exec, keys_first, keys_last, values_first, thrust::less()); } // end stable_sort_by_key() template _CCCL_HOST_DEVICE bool is_sorted(thrust::execution_policy& exec, ForwardIterator first, ForwardIterator last) { return thrust::is_sorted_until(exec, first, last) == last; } // end is_sorted() template _CCCL_HOST_DEVICE bool is_sorted(thrust::execution_policy& exec, ForwardIterator first, ForwardIterator last, Compare comp) { return thrust::is_sorted_until(exec, first, last, comp) == last; } // end is_sorted() template _CCCL_HOST_DEVICE ForwardIterator is_sorted_until(thrust::execution_policy& exec, ForwardIterator first, ForwardIterator last) { typedef typename thrust::iterator_value::type InputType; return thrust::is_sorted_until(exec, first, last, thrust::less()); } // end is_sorted_until() template _CCCL_HOST_DEVICE ForwardIterator is_sorted_until( thrust::execution_policy& exec, ForwardIterator first, ForwardIterator last, Compare comp) { if (thrust::distance(first, last) < 2) { return last; } typedef thrust::tuple IteratorTuple; typedef thrust::zip_iterator ZipIterator; ForwardIterator first_plus_one = first; thrust::advance(first_plus_one, 1); ZipIterator zipped_first = thrust::make_zip_iterator(thrust::make_tuple(first_plus_one, first)); ZipIterator zipped_last = thrust::make_zip_iterator(thrust::make_tuple(last, first)); return thrust::get<0>( thrust::find_if(exec, zipped_first, zipped_last, thrust::detail::tuple_binary_predicate(comp)) .get_iterator_tuple()); } // end is_sorted_until() template _CCCL_HOST_DEVICE void stable_sort(thrust::execution_policy&, RandomAccessIterator, RandomAccessIterator, StrictWeakOrdering) { THRUST_STATIC_ASSERT_MSG((thrust::detail::depend_on_instantiation::value), "unimplemented for this system"); } // end stable_sort() template _CCCL_HOST_DEVICE void stable_sort_by_key( thrust::execution_policy&, RandomAccessIterator1, RandomAccessIterator1, RandomAccessIterator2, StrictWeakOrdering) { THRUST_STATIC_ASSERT_MSG((thrust::detail::depend_on_instantiation::value), "unimplemented for this system"); } // end stable_sort_by_key() } // namespace generic } // namespace detail } // namespace system THRUST_NAMESPACE_END