/* * 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 #include #include #include #include #include #include #include THRUST_NAMESPACE_BEGIN namespace system { namespace detail { namespace generic { template __host__ __device__ OutputIterator set_difference(thrust::execution_policy &exec, InputIterator1 first1, InputIterator1 last1, InputIterator2 first2, InputIterator2 last2, OutputIterator result) { typedef typename thrust::iterator_value::type value_type; return thrust::set_difference(exec, first1, last1, first2, last2, result, thrust::less()); } // end set_difference() template __host__ __device__ thrust::pair set_difference_by_key(thrust::execution_policy &exec, InputIterator1 keys_first1, InputIterator1 keys_last1, InputIterator2 keys_first2, InputIterator2 keys_last2, InputIterator3 values_first1, InputIterator4 values_first2, OutputIterator1 keys_result, OutputIterator2 values_result) { typedef typename thrust::iterator_value::type value_type; return thrust::set_difference_by_key(exec, keys_first1, keys_last1, keys_first2, keys_last2, values_first1, values_first2, keys_result, values_result, thrust::less()); } // end set_difference_by_key() template __host__ __device__ thrust::pair set_difference_by_key(thrust::execution_policy &exec, InputIterator1 keys_first1, InputIterator1 keys_last1, InputIterator2 keys_first2, InputIterator2 keys_last2, InputIterator3 values_first1, InputIterator4 values_first2, OutputIterator1 keys_result, OutputIterator2 values_result, StrictWeakOrdering comp) { typedef thrust::tuple iterator_tuple1; typedef thrust::tuple iterator_tuple2; typedef thrust::tuple iterator_tuple3; typedef thrust::zip_iterator zip_iterator1; typedef thrust::zip_iterator zip_iterator2; typedef thrust::zip_iterator zip_iterator3; zip_iterator1 zipped_first1 = thrust::make_zip_iterator(thrust::make_tuple(keys_first1, values_first1)); zip_iterator1 zipped_last1 = thrust::make_zip_iterator(thrust::make_tuple(keys_last1, values_first1)); zip_iterator2 zipped_first2 = thrust::make_zip_iterator(thrust::make_tuple(keys_first2, values_first2)); zip_iterator2 zipped_last2 = thrust::make_zip_iterator(thrust::make_tuple(keys_last2, values_first2)); zip_iterator3 zipped_result = thrust::make_zip_iterator(thrust::make_tuple(keys_result, values_result)); thrust::detail::compare_first comp_first(comp); iterator_tuple3 result = thrust::set_difference(exec, zipped_first1, zipped_last1, zipped_first2, zipped_last2, zipped_result, comp_first).get_iterator_tuple(); return thrust::make_pair(thrust::get<0>(result), thrust::get<1>(result)); } // end set_difference_by_key() template __host__ __device__ OutputIterator set_intersection(thrust::execution_policy &exec, InputIterator1 first1, InputIterator1 last1, InputIterator2 first2, InputIterator2 last2, OutputIterator result) { typedef typename thrust::iterator_value::type value_type; return thrust::set_intersection(exec, first1, last1, first2, last2, result, thrust::less()); } // end set_intersection() template __host__ __device__ thrust::pair set_intersection_by_key(thrust::execution_policy &exec, InputIterator1 keys_first1, InputIterator1 keys_last1, InputIterator2 keys_first2, InputIterator2 keys_last2, InputIterator3 values_first1, OutputIterator1 keys_result, OutputIterator2 values_result) { typedef typename thrust::iterator_value::type value_type; return thrust::set_intersection_by_key(exec, keys_first1, keys_last1, keys_first2, keys_last2, values_first1, keys_result, values_result, thrust::less()); } // end set_intersection_by_key() template __host__ __device__ thrust::pair set_intersection_by_key(thrust::execution_policy &exec, InputIterator1 keys_first1, InputIterator1 keys_last1, InputIterator2 keys_first2, InputIterator2 keys_last2, InputIterator3 values_first1, OutputIterator1 keys_result, OutputIterator2 values_result, StrictWeakOrdering comp) { typedef typename thrust::iterator_value::type value_type1; typedef thrust::constant_iterator constant_iterator; typedef thrust::tuple iterator_tuple1; typedef thrust::tuple iterator_tuple2; typedef thrust::tuple iterator_tuple3; typedef thrust::zip_iterator zip_iterator1; typedef thrust::zip_iterator zip_iterator2; typedef thrust::zip_iterator zip_iterator3; // fabricate a values_first2 by repeating a default-constructed value_type1 // XXX assumes value_type1 is default-constructible constant_iterator values_first2 = thrust::make_constant_iterator(value_type1()); zip_iterator1 zipped_first1 = thrust::make_zip_iterator(thrust::make_tuple(keys_first1, values_first1)); zip_iterator1 zipped_last1 = thrust::make_zip_iterator(thrust::make_tuple(keys_last1, values_first1)); zip_iterator2 zipped_first2 = thrust::make_zip_iterator(thrust::make_tuple(keys_first2, values_first2)); zip_iterator2 zipped_last2 = thrust::make_zip_iterator(thrust::make_tuple(keys_last2, values_first2)); zip_iterator3 zipped_result = thrust::make_zip_iterator(thrust::make_tuple(keys_result, values_result)); thrust::detail::compare_first comp_first(comp); iterator_tuple3 result = thrust::set_intersection(exec, zipped_first1, zipped_last1, zipped_first2, zipped_last2, zipped_result, comp_first).get_iterator_tuple(); return thrust::make_pair(thrust::get<0>(result), thrust::get<1>(result)); } // end set_intersection_by_key() template __host__ __device__ OutputIterator set_symmetric_difference(thrust::execution_policy &exec, InputIterator1 first1, InputIterator1 last1, InputIterator2 first2, InputIterator2 last2, OutputIterator result) { typedef typename thrust::iterator_value::type value_type; return thrust::set_symmetric_difference(exec, first1, last1, first2, last2, result, thrust::less()); } // end set_symmetric_difference() template __host__ __device__ thrust::pair set_symmetric_difference_by_key(thrust::execution_policy &exec, InputIterator1 keys_first1, InputIterator1 keys_last1, InputIterator2 keys_first2, InputIterator2 keys_last2, InputIterator3 values_first1, InputIterator4 values_first2, OutputIterator1 keys_result, OutputIterator2 values_result) { typedef typename thrust::iterator_value::type value_type; return thrust::set_symmetric_difference_by_key(exec, keys_first1, keys_last1, keys_first2, keys_last2, values_first1, values_first2, keys_result, values_result, thrust::less()); } // end set_symmetric_difference_by_key() template __host__ __device__ thrust::pair set_symmetric_difference_by_key(thrust::execution_policy &exec, InputIterator1 keys_first1, InputIterator1 keys_last1, InputIterator2 keys_first2, InputIterator2 keys_last2, InputIterator3 values_first1, InputIterator4 values_first2, OutputIterator1 keys_result, OutputIterator2 values_result, StrictWeakOrdering comp) { typedef thrust::tuple iterator_tuple1; typedef thrust::tuple iterator_tuple2; typedef thrust::tuple iterator_tuple3; typedef thrust::zip_iterator zip_iterator1; typedef thrust::zip_iterator zip_iterator2; typedef thrust::zip_iterator zip_iterator3; zip_iterator1 zipped_first1 = thrust::make_zip_iterator(thrust::make_tuple(keys_first1, values_first1)); zip_iterator1 zipped_last1 = thrust::make_zip_iterator(thrust::make_tuple(keys_last1, values_first1)); zip_iterator2 zipped_first2 = thrust::make_zip_iterator(thrust::make_tuple(keys_first2, values_first2)); zip_iterator2 zipped_last2 = thrust::make_zip_iterator(thrust::make_tuple(keys_last2, values_first2)); zip_iterator3 zipped_result = thrust::make_zip_iterator(thrust::make_tuple(keys_result, values_result)); thrust::detail::compare_first comp_first(comp); iterator_tuple3 result = thrust::set_symmetric_difference(exec, zipped_first1, zipped_last1, zipped_first2, zipped_last2, zipped_result, comp_first).get_iterator_tuple(); return thrust::make_pair(thrust::get<0>(result), thrust::get<1>(result)); } // end set_symmetric_difference_by_key() template __host__ __device__ OutputIterator set_union(thrust::execution_policy &exec, InputIterator1 first1, InputIterator1 last1, InputIterator2 first2, InputIterator2 last2, OutputIterator result) { typedef typename thrust::iterator_value::type value_type; return thrust::set_union(exec, first1, last1, first2, last2, result, thrust::less()); } // end set_union() template __host__ __device__ thrust::pair set_union_by_key(thrust::execution_policy &exec, InputIterator1 keys_first1, InputIterator1 keys_last1, InputIterator2 keys_first2, InputIterator2 keys_last2, InputIterator3 values_first1, InputIterator4 values_first2, OutputIterator1 keys_result, OutputIterator2 values_result) { typedef typename thrust::iterator_value::type value_type; return thrust::set_union_by_key(exec, keys_first1, keys_last1, keys_first2, keys_last2, values_first1, values_first2, keys_result, values_result, thrust::less()); } // end set_union_by_key() template __host__ __device__ thrust::pair set_union_by_key(thrust::execution_policy &exec, InputIterator1 keys_first1, InputIterator1 keys_last1, InputIterator2 keys_first2, InputIterator2 keys_last2, InputIterator3 values_first1, InputIterator4 values_first2, OutputIterator1 keys_result, OutputIterator2 values_result, StrictWeakOrdering comp) { typedef thrust::tuple iterator_tuple1; typedef thrust::tuple iterator_tuple2; typedef thrust::tuple iterator_tuple3; typedef thrust::zip_iterator zip_iterator1; typedef thrust::zip_iterator zip_iterator2; typedef thrust::zip_iterator zip_iterator3; zip_iterator1 zipped_first1 = thrust::make_zip_iterator(thrust::make_tuple(keys_first1, values_first1)); zip_iterator1 zipped_last1 = thrust::make_zip_iterator(thrust::make_tuple(keys_last1, values_first1)); zip_iterator2 zipped_first2 = thrust::make_zip_iterator(thrust::make_tuple(keys_first2, values_first2)); zip_iterator2 zipped_last2 = thrust::make_zip_iterator(thrust::make_tuple(keys_last2, values_first2)); zip_iterator3 zipped_result = thrust::make_zip_iterator(thrust::make_tuple(keys_result, values_result)); thrust::detail::compare_first comp_first(comp); iterator_tuple3 result = thrust::set_union(exec, zipped_first1, zipped_last1, zipped_first2, zipped_last2, zipped_result, comp_first).get_iterator_tuple(); return thrust::make_pair(thrust::get<0>(result), thrust::get<1>(result)); } // end set_union_by_key() template __host__ __device__ OutputIterator set_difference(thrust::execution_policy &, InputIterator1, InputIterator1, InputIterator2, InputIterator2, OutputIterator result, StrictWeakOrdering) { THRUST_STATIC_ASSERT_MSG( (thrust::detail::depend_on_instantiation::value) , "unimplemented for this system" ); return result; } // end set_difference() template __host__ __device__ OutputIterator set_intersection(thrust::execution_policy &, InputIterator1, InputIterator1, InputIterator2, InputIterator2, OutputIterator result, StrictWeakOrdering) { THRUST_STATIC_ASSERT_MSG( (thrust::detail::depend_on_instantiation::value) , "unimplemented for this system" ); return result; } // end set_intersection() template __host__ __device__ OutputIterator set_symmetric_difference(thrust::execution_policy &, InputIterator1, InputIterator1, InputIterator2, InputIterator2, OutputIterator result, StrictWeakOrdering) { THRUST_STATIC_ASSERT_MSG( (thrust::detail::depend_on_instantiation::value) , "unimplemented for this system" ); return result; } // end set_symmetric_difference() template __host__ __device__ OutputIterator set_union(thrust::execution_policy &, InputIterator1, InputIterator1, InputIterator2, InputIterator2, OutputIterator result, StrictWeakOrdering) { THRUST_STATIC_ASSERT_MSG( (thrust::detail::depend_on_instantiation::value) , "unimplemented for this system" ); return result; } // end set_union() } // end namespace generic } // end namespace detail } // end namespace system THRUST_NAMESPACE_END