1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
|
/*
* 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.
*/
/*! \file distance.h
* \brief Device implementations for distance.
*/
#pragma once
#include <thrust/detail/config.h>
#include <thrust/detail/get_iterator_value.h>
#include <thrust/extrema.h>
#include <thrust/functional.h>
#include <thrust/pair.h>
#include <thrust/reduce.h>
#include <thrust/transform_reduce.h>
#include <thrust/iterator/iterator_traits.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/zip_iterator.h>
THRUST_NAMESPACE_BEGIN
namespace system
{
namespace detail
{
namespace generic
{
namespace detail
{
//////////////
// Functors //
//////////////
//
// return the smaller/larger element making sure to prefer the
// first occurance of the minimum/maximum element
template <typename InputType, typename IndexType, typename BinaryPredicate>
struct min_element_reduction
{
BinaryPredicate comp;
__host__ __device__
min_element_reduction(BinaryPredicate comp) : comp(comp){}
__host__ __device__
thrust::tuple<InputType, IndexType>
operator()(const thrust::tuple<InputType, IndexType>& lhs,
const thrust::tuple<InputType, IndexType>& rhs )
{
if(comp(thrust::get<0>(lhs), thrust::get<0>(rhs)))
return lhs;
if(comp(thrust::get<0>(rhs), thrust::get<0>(lhs)))
return rhs;
// values are equivalent, prefer value with smaller index
if(thrust::get<1>(lhs) < thrust::get<1>(rhs))
return lhs;
else
return rhs;
} // end operator()()
}; // end min_element_reduction
template <typename InputType, typename IndexType, typename BinaryPredicate>
struct max_element_reduction
{
BinaryPredicate comp;
__host__ __device__
max_element_reduction(BinaryPredicate comp) : comp(comp){}
__host__ __device__
thrust::tuple<InputType, IndexType>
operator()(const thrust::tuple<InputType, IndexType>& lhs,
const thrust::tuple<InputType, IndexType>& rhs )
{
if(comp(thrust::get<0>(lhs), thrust::get<0>(rhs)))
return rhs;
if(comp(thrust::get<0>(rhs), thrust::get<0>(lhs)))
return lhs;
// values are equivalent, prefer value with smaller index
if(thrust::get<1>(lhs) < thrust::get<1>(rhs))
return lhs;
else
return rhs;
} // end operator()()
}; // end max_element_reduction
// return the smaller & larger element making sure to prefer the
// first occurance of the minimum/maximum element
template <typename InputType, typename IndexType, typename BinaryPredicate>
struct minmax_element_reduction
{
BinaryPredicate comp;
__host__ __device__
minmax_element_reduction(BinaryPredicate comp) : comp(comp){}
__host__ __device__
thrust::tuple< thrust::tuple<InputType,IndexType>, thrust::tuple<InputType,IndexType> >
operator()(const thrust::tuple< thrust::tuple<InputType,IndexType>, thrust::tuple<InputType,IndexType> >& lhs,
const thrust::tuple< thrust::tuple<InputType,IndexType>, thrust::tuple<InputType,IndexType> >& rhs )
{
return thrust::make_tuple(min_element_reduction<InputType, IndexType, BinaryPredicate>(comp)(thrust::get<0>(lhs), thrust::get<0>(rhs)),
max_element_reduction<InputType, IndexType, BinaryPredicate>(comp)(thrust::get<1>(lhs), thrust::get<1>(rhs)));
} // end operator()()
}; // end minmax_element_reduction
template <typename InputType, typename IndexType>
struct duplicate_tuple
{
__host__ __device__
thrust::tuple< thrust::tuple<InputType,IndexType>, thrust::tuple<InputType,IndexType> >
operator()(const thrust::tuple<InputType,IndexType>& t)
{
return thrust::make_tuple(t, t);
}
}; // end duplicate_tuple
} // end namespace detail
template <typename DerivedPolicy, typename ForwardIterator>
__host__ __device__
ForwardIterator min_element(thrust::execution_policy<DerivedPolicy> &exec,
ForwardIterator first,
ForwardIterator last)
{
typedef typename thrust::iterator_value<ForwardIterator>::type value_type;
return thrust::min_element(exec, first, last, thrust::less<value_type>());
} // end min_element()
template <typename DerivedPolicy, typename ForwardIterator, typename BinaryPredicate>
__host__ __device__
ForwardIterator min_element(thrust::execution_policy<DerivedPolicy> &exec,
ForwardIterator first,
ForwardIterator last,
BinaryPredicate comp)
{
if (first == last)
return last;
typedef typename thrust::iterator_traits<ForwardIterator>::value_type InputType;
typedef typename thrust::iterator_traits<ForwardIterator>::difference_type IndexType;
thrust::tuple<InputType, IndexType> result =
thrust::reduce
(exec,
thrust::make_zip_iterator(thrust::make_tuple(first, thrust::counting_iterator<IndexType>(0))),
thrust::make_zip_iterator(thrust::make_tuple(first, thrust::counting_iterator<IndexType>(0))) + (last - first),
thrust::tuple<InputType, IndexType>(thrust::detail::get_iterator_value(derived_cast(exec), first), 0),
detail::min_element_reduction<InputType, IndexType, BinaryPredicate>(comp));
return first + thrust::get<1>(result);
} // end min_element()
template <typename DerivedPolicy, typename ForwardIterator>
__host__ __device__
ForwardIterator max_element(thrust::execution_policy<DerivedPolicy> &exec,
ForwardIterator first,
ForwardIterator last)
{
typedef typename thrust::iterator_value<ForwardIterator>::type value_type;
return thrust::max_element(exec, first, last, thrust::less<value_type>());
} // end max_element()
template <typename DerivedPolicy, typename ForwardIterator, typename BinaryPredicate>
__host__ __device__
ForwardIterator max_element(thrust::execution_policy<DerivedPolicy> &exec,
ForwardIterator first,
ForwardIterator last,
BinaryPredicate comp)
{
if (first == last)
return last;
typedef typename thrust::iterator_traits<ForwardIterator>::value_type InputType;
typedef typename thrust::iterator_traits<ForwardIterator>::difference_type IndexType;
thrust::tuple<InputType, IndexType> result =
thrust::reduce
(exec,
thrust::make_zip_iterator(thrust::make_tuple(first, thrust::counting_iterator<IndexType>(0))),
thrust::make_zip_iterator(thrust::make_tuple(first, thrust::counting_iterator<IndexType>(0))) + (last - first),
thrust::tuple<InputType, IndexType>(thrust::detail::get_iterator_value(derived_cast(exec),first), 0),
detail::max_element_reduction<InputType, IndexType, BinaryPredicate>(comp));
return first + thrust::get<1>(result);
} // end max_element()
template <typename DerivedPolicy, typename ForwardIterator>
__host__ __device__
thrust::pair<ForwardIterator,ForwardIterator> minmax_element(thrust::execution_policy<DerivedPolicy> &exec,
ForwardIterator first,
ForwardIterator last)
{
typedef typename thrust::iterator_value<ForwardIterator>::type value_type;
return thrust::minmax_element(exec, first, last, thrust::less<value_type>());
} // end minmax_element()
template <typename DerivedPolicy, typename ForwardIterator, typename BinaryPredicate>
__host__ __device__
thrust::pair<ForwardIterator,ForwardIterator> minmax_element(thrust::execution_policy<DerivedPolicy> &exec,
ForwardIterator first,
ForwardIterator last,
BinaryPredicate comp)
{
if (first == last)
return thrust::make_pair(last, last);
typedef typename thrust::iterator_traits<ForwardIterator>::value_type InputType;
typedef typename thrust::iterator_traits<ForwardIterator>::difference_type IndexType;
thrust::tuple< thrust::tuple<InputType,IndexType>, thrust::tuple<InputType,IndexType> > result =
thrust::transform_reduce
(exec,
thrust::make_zip_iterator(thrust::make_tuple(first, thrust::counting_iterator<IndexType>(0))),
thrust::make_zip_iterator(thrust::make_tuple(first, thrust::counting_iterator<IndexType>(0))) + (last - first),
detail::duplicate_tuple<InputType, IndexType>(),
detail::duplicate_tuple<InputType, IndexType>()(
thrust::tuple<InputType, IndexType>(thrust::detail::get_iterator_value(derived_cast(exec),first), 0)),
detail::minmax_element_reduction<InputType, IndexType, BinaryPredicate>(comp));
return thrust::make_pair(first + thrust::get<1>(thrust::get<0>(result)), first + thrust::get<1>(thrust::get<1>(result)));
} // end minmax_element()
} // end namespace generic
} // end namespace detail
} // end namespace system
THRUST_NAMESPACE_END
|
