util/X86LINUX64/cuda-6.5/include/thrust/system/detail/generic/copy_if.inl

   1 /*
   2  *  Copyright 2008-2013 NVIDIA Corporation
   3  *
   4  *  Licensed under the Apache License, Version 2.0 (the "License");
   5  *  you may not use this file except in compliance with the License.
   6  *  You may obtain a copy of the License at
   7  *
   8  *      http://www.apache.org/licenses/LICENSE-2.0
   9  *
  10  *  Unless required by applicable law or agreed to in writing, software
  11  *  distributed under the License is distributed on an "AS IS" BASIS,
  12  *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
  13  *  See the License for the specific language governing permissions and
  14  *  limitations under the License.
  15  */
  16
  17 #pragma once
  18
  19 #include <thrust/detail/config.h>
  20 #include <thrust/system/detail/generic/copy_if.h>
  21 #include <thrust/detail/copy_if.h>
  22 #include <thrust/iterator/iterator_traits.h>
  23 #include <thrust/iterator/detail/minimum_system.h>
  24 #include <thrust/functional.h>
  25 #include <thrust/distance.h>
  26 #include <thrust/transform.h>
  27 #include <thrust/detail/internal_functional.h>
  28 #include <thrust/detail/temporary_array.h>
  29 #include <thrust/detail/type_traits.h>
  30 #include <thrust/scan.h>
  31 #include <thrust/scatter.h>
  32 #include <limits>
  33
  34 namespace thrust
  35 {
  36 namespace system
  37 {
  38 namespace detail
  39 {
  40 namespace generic
  41 {
  42 namespace detail
  43 {
  44
  45 template<typename IndexType,
  46          typename DerivedPolicy,
  47          typename InputIterator1,
  48          typename InputIterator2,
  49          typename OutputIterator,
  50          typename Predicate>
  51 OutputIterator copy_if(thrust::execution_policy<DerivedPolicy> &exec,
  52                        InputIterator1 first,
  53                        InputIterator1 last,
  54                        InputIterator2 stencil,
  55                        OutputIterator result,
  56                        Predicate pred)
  57 {
  58     __THRUST_DISABLE_MSVC_POSSIBLE_LOSS_OF_DATA_WARNING(IndexType n = thrust::distance(first, last));
  59
  60     // compute {0,1} predicates
  61     thrust::detail::temporary_array<IndexType, DerivedPolicy> predicates(exec, n);
  62     thrust::transform(exec,
  63                       stencil,
  64                       stencil + n,
  65                       predicates.begin(),
  66                       thrust::detail::predicate_to_integral<Predicate,IndexType>(pred));
  67
  68     // scan {0,1} predicates
  69     thrust::detail::temporary_array<IndexType, DerivedPolicy> scatter_indices(exec, n);
  70     thrust::exclusive_scan(exec,
  71                            predicates.begin(),
  72                            predicates.end(),
  73                            scatter_indices.begin(),
  74                            static_cast<IndexType>(0),
  75                            thrust::plus<IndexType>());
  76
  77     // scatter the true elements
  78     thrust::scatter_if(exec,
  79                        first,
  80                        last,
  81                        scatter_indices.begin(),
  82                        predicates.begin(),
  83                        result,
  84                        thrust::identity<IndexType>());
  85
  86     // find the end of the new sequence
  87     IndexType output_size = scatter_indices[n - 1] + predicates[n - 1];
  88
  89     return result + output_size;
  90 }
  91
  92 } // end namespace detail
  93
  94
  95 template<typename DerivedPolicy,
  96          typename InputIterator,
  97          typename OutputIterator,
  98          typename Predicate>
  99   OutputIterator copy_if(thrust::execution_policy<DerivedPolicy> &exec,
 100                          InputIterator first,
 101                          InputIterator last,
 102                          OutputIterator result,
 103                          Predicate pred)
 104 {
 105   // XXX it's potentially expensive to send [first,last) twice
 106   //     we should probably specialize this case for POD
 107   //     since we can safely keep the input in a temporary instead
 108   //     of doing two loads
 109   return thrust::copy_if(exec, first, last, first, result, pred);
 110 } // end copy_if()
 111
 112
 113 template<typename DerivedPolicy,
 114          typename InputIterator1,
 115          typename InputIterator2,
 116          typename OutputIterator,
 117          typename Predicate>
 118    OutputIterator copy_if(thrust::execution_policy<DerivedPolicy> &exec,
 119                           InputIterator1 first,
 120                           InputIterator1 last,
 121                           InputIterator2 stencil,
 122                           OutputIterator result,
 123                           Predicate pred)
 124 {
 125   typedef typename thrust::iterator_traits<InputIterator1>::difference_type difference_type;
 126
 127   // empty sequence
 128   if(first == last)
 129     return result;
 130
 131   difference_type n = thrust::distance(first, last);
 132
 133   // create an unsigned version of n (we know n is positive from the comparison above)
 134   // to avoid a warning in the compare below
 135   typename thrust::detail::make_unsigned<difference_type>::type unsigned_n(n);
 136
 137   // use 32-bit indices when possible (almost always)
 138   if(sizeof(difference_type) > sizeof(unsigned int) && unsigned_n > (std::numeric_limits<unsigned int>::max)())
 139   {
 140     result = detail::copy_if<difference_type>(exec, first, last, stencil, result, pred);
 141   } // end if
 142   else
 143   {
 144     result = detail::copy_if<unsigned int>(exec, first, last, stencil, result, pred);
 145   } // end else
 146
 147   return result;
 148 } // end copy_if()
 149
 150
 151 } // end namespace generic
 152 } // end namespace detail
 153 } // end namespace system
 154 } // end namespace thrust
 155