2
* Copyright 2008-2011 NVIDIA Corporation
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
8
* http://www.apache.org/licenses/LICENSE-2.0
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.
19
#include <thrust/iterator/iterator_traits.h>
20
#include <thrust/iterator/detail/minimum_space.h>
21
#include <thrust/functional.h>
22
#include <thrust/distance.h>
23
#include <thrust/transform.h>
25
#include <thrust/detail/internal_functional.h>
26
#include <thrust/detail/raw_buffer.h>
27
#include <thrust/detail/type_traits.h>
29
#include <thrust/detail/device/scan.h>
30
#include <thrust/scatter.h>
41
// XXX WAR circular #inclusion with forward declaration
42
template<typename InputIterator,
43
typename OutputIterator,
45
typename AssociativeOperator>
46
OutputIterator exclusive_scan(InputIterator first,
48
OutputIterator result,
50
AssociativeOperator binary_op);
53
template<typename InputIterator1,
54
typename InputIterator2,
55
typename InputIterator3,
56
typename RandomAccessIterator,
58
void scatter_if(InputIterator1 first,
61
InputIterator3 stencil,
62
RandomAccessIterator output,
71
template<typename IndexType,
72
typename InputIterator1,
73
typename InputIterator2,
74
typename OutputIterator,
76
OutputIterator copy_if(InputIterator1 first,
78
InputIterator2 stencil,
79
OutputIterator result,
82
typedef typename thrust::detail::minimum_space<
83
typename thrust::iterator_space<InputIterator1>::type,
84
typename thrust::iterator_space<InputIterator2>::type,
85
typename thrust::iterator_space<OutputIterator>::type
88
__THRUST_DISABLE_MSVC_POSSIBLE_LOSS_OF_DATA_WARNING(IndexType n = thrust::distance(first, last));
90
// compute {0,1} predicates
91
thrust::detail::raw_buffer<IndexType, Space> predicates(n);
92
thrust::transform(stencil,
95
thrust::detail::predicate_to_integral<Predicate,IndexType>(pred));
97
// scan {0,1} predicates
98
thrust::detail::raw_buffer<IndexType, Space> scatter_indices(n);
99
thrust::detail::device::exclusive_scan(predicates.begin(),
101
scatter_indices.begin(),
102
static_cast<IndexType>(0),
103
thrust::plus<IndexType>());
105
// scatter the true elements
106
thrust::scatter_if(first,
108
scatter_indices.begin(),
111
thrust::identity<IndexType>());
113
// find the end of the new sequence
114
IndexType output_size = scatter_indices[n - 1] + predicates[n - 1];
116
return result + output_size;
119
} // end namespace detail
121
template<typename InputIterator1,
122
typename InputIterator2,
123
typename OutputIterator,
125
OutputIterator copy_if(InputIterator1 first,
127
InputIterator2 stencil,
128
OutputIterator result,
131
typedef typename thrust::iterator_traits<InputIterator1>::difference_type difference_type;
137
difference_type n = thrust::distance(first, last);
139
// create an unsigned version of n (we know n is positive from the comparison above)
140
// to avoid a warning in the compare below
141
typename thrust::detail::make_unsigned<difference_type>::type unsigned_n(n);
143
// use 32-bit indices when possible (almost always)
144
if (sizeof(difference_type) > sizeof(unsigned int) && unsigned_n > (std::numeric_limits<unsigned int>::max)())
145
return detail::copy_if<difference_type>(first, last, stencil, result, pred);
147
return detail::copy_if<unsigned int>(first, last, stencil, result, pred);
150
} // end namespace generic
151
} // end namespace device
152
} // end namespace detail
153
} // end namespace thrust
added
removed
detail/device/generic/copy_if.inl
