~ubuntu-branches/ubuntu/trusty/libthrust/trusty

« back to all changes in this revision

Viewing changes to detail/internal_functional.h

  • Committer: Bazaar Package Importer
  • Author(s): Andreas Beckmann
  • Date: 2011-05-28 09:32:48 UTC
  • Revision ID: james.westby@ubuntu.com-20110528093248-np3euv5sj7fw3nyv
Tags: upstream-1.4.0
ImportĀ upstreamĀ versionĀ 1.4.0

Show diffs side-by-side

added added

removed removed

Lines of Context:
 
1
/*
 
2
 *  Copyright 2008-2011 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
 
 
18
/*! \file internal_functional.inl
 
19
 *  \brief Non-public functionals used to implement algorithm internals.
 
20
 */
 
21
 
 
22
#pragma once
 
23
 
 
24
#include <thrust/tuple.h>
 
25
#include <thrust/iterator/iterator_traits.h>
 
26
#include <memory> // for ::new
 
27
 
 
28
namespace thrust
 
29
{
 
30
namespace detail
 
31
{
 
32
 
 
33
// unary_negate does not need to know argument_type
 
34
template <typename Predicate>
 
35
struct unary_negate
 
36
{
 
37
    typedef bool result_type;
 
38
 
 
39
    Predicate pred;
 
40
 
 
41
    __host__ __device__
 
42
    explicit unary_negate(const Predicate& pred) : pred(pred) {}
 
43
 
 
44
    template <typename T>
 
45
    __host__ __device__
 
46
    bool operator()(const T& x)
 
47
    {
 
48
        return !bool(pred(x));
 
49
    }
 
50
};
 
51
 
 
52
// binary_negate does not need to know first_argument_type or second_argument_type
 
53
template <typename Predicate>
 
54
struct binary_negate
 
55
{
 
56
    typedef bool result_type;
 
57
 
 
58
    Predicate pred;
 
59
 
 
60
    __host__ __device__
 
61
    explicit binary_negate(const Predicate& pred) : pred(pred) {}
 
62
 
 
63
    template <typename T1, typename T2>
 
64
        __host__ __device__
 
65
        bool operator()(const T1& x, const T2& y)
 
66
        {
 
67
            return !bool(pred(x,y));
 
68
        }
 
69
};
 
70
 
 
71
template<typename Predicate>
 
72
  __host__ __device__
 
73
  thrust::detail::unary_negate<Predicate> not1(const Predicate &pred)
 
74
{
 
75
    return thrust::detail::unary_negate<Predicate>(pred);
 
76
}
 
77
 
 
78
template<typename Predicate>
 
79
  __host__ __device__
 
80
  thrust::detail::binary_negate<Predicate> not2(const Predicate &pred)
 
81
{
 
82
    return thrust::detail::binary_negate<Predicate>(pred);
 
83
}
 
84
 
 
85
 
 
86
// convert a predicate to a 0 or 1 integral value
 
87
template <typename Predicate, typename IntegralType>
 
88
struct predicate_to_integral
 
89
{
 
90
    Predicate pred;
 
91
 
 
92
    __host__ __device__
 
93
    explicit predicate_to_integral(const Predicate& pred) : pred(pred) {}
 
94
 
 
95
    template <typename T>
 
96
        __host__ __device__
 
97
        bool operator()(const T& x)
 
98
        {
 
99
            return pred(x) ? IntegralType(1) : IntegralType(0);
 
100
        }
 
101
};
 
102
 
 
103
 
 
104
// note that detail::equal_to does not force conversion from T2 -> T1 as equal_to does
 
105
template <typename T1>
 
106
struct equal_to
 
107
{
 
108
    typedef bool result_type;
 
109
 
 
110
    template <typename T2>
 
111
        __host__ __device__
 
112
        bool operator()(const T1& lhs, const T2& rhs) const
 
113
        {
 
114
            return lhs == rhs;
 
115
        }
 
116
};
 
117
 
 
118
// note that equal_to_value does not force conversion from T2 -> T1 as equal_to does
 
119
template <typename T2>
 
120
struct equal_to_value
 
121
{
 
122
    T2 rhs;
 
123
 
 
124
    equal_to_value(const T2& rhs) : rhs(rhs) {}
 
125
 
 
126
    template <typename T1>
 
127
        __host__ __device__
 
128
        bool operator()(const T1& lhs) const
 
129
        {
 
130
            return lhs == rhs;
 
131
        }
 
132
};
 
133
 
 
134
template <typename Predicate>
 
135
struct tuple_binary_predicate
 
136
{
 
137
    typedef bool result_type;
 
138
 
 
139
    __host__ __device__
 
140
        tuple_binary_predicate(const Predicate& p) : pred(p) {}
 
141
 
 
142
    template<typename Tuple>
 
143
        __host__ __device__
 
144
        bool operator()(const Tuple& t) const
 
145
        { 
 
146
            return pred(thrust::get<0>(t), thrust::get<1>(t));
 
147
        }
 
148
 
 
149
    Predicate pred;
 
150
};
 
151
 
 
152
template <typename Predicate>
 
153
struct tuple_not_binary_predicate
 
154
{
 
155
    typedef bool result_type;
 
156
 
 
157
    __host__ __device__
 
158
        tuple_not_binary_predicate(const Predicate& p) : pred(p) {}
 
159
 
 
160
    template<typename Tuple>
 
161
        __host__ __device__
 
162
        bool operator()(const Tuple& t) const
 
163
        { 
 
164
            return !pred(thrust::get<0>(t), thrust::get<1>(t));
 
165
        }
 
166
 
 
167
    Predicate pred;
 
168
};
 
169
 
 
170
template<typename Generator>
 
171
  struct generate_functor
 
172
{
 
173
  typedef void result_type;
 
174
 
 
175
  __host__ __device__
 
176
  generate_functor(Generator g)
 
177
    : gen(g) {}
 
178
 
 
179
  // operator() does not take an lvalue reference because some iterators
 
180
  // produce temporary proxy references when dereferenced. for example,
 
181
  // consider the temporary tuple of references produced by zip_iterator.
 
182
  // such temporaries cannot bind to an lvalue reference.
 
183
  //
 
184
  // to WAR this, accept a const reference (which is bindable to a temporary),
 
185
  // and const_cast in the implementation.
 
186
  //
 
187
  // XXX change to an rvalue reference upon c++0x (which either a named variable
 
188
  //     or temporary can bind to)
 
189
  template<typename T>
 
190
  __host__ __device__
 
191
  void operator()(const T &x)
 
192
  {
 
193
    // we have to be naughty and const_cast this to get it to work
 
194
    T &lvalue = const_cast<T&>(x);
 
195
 
 
196
    // this assigns correctly whether x is a true reference or proxy
 
197
    lvalue = gen();
 
198
  }
 
199
 
 
200
  Generator gen;
 
201
};
 
202
 
 
203
 
 
204
template<typename ResultType, typename BinaryFunction>
 
205
  struct zipped_binary_op
 
206
{
 
207
  typedef ResultType result_type;
 
208
 
 
209
  __host__ __device__
 
210
  zipped_binary_op(BinaryFunction binary_op)
 
211
    : m_binary_op(binary_op) {}
 
212
 
 
213
  template<typename Tuple>
 
214
  __host__ __device__
 
215
  inline result_type operator()(Tuple t)
 
216
  {
 
217
    return m_binary_op(thrust::get<0>(t), thrust::get<1>(t));
 
218
  }
 
219
 
 
220
  BinaryFunction m_binary_op;
 
221
};
 
222
 
 
223
template<typename UnaryFunction>
 
224
  struct host_unary_transform_functor
 
225
{
 
226
  typedef void result_type;
 
227
 
 
228
  UnaryFunction f;
 
229
 
 
230
  host_unary_transform_functor(UnaryFunction f_)
 
231
    :f(f_) {}
 
232
 
 
233
  template<typename Tuple>
 
234
  __host__
 
235
  inline result_type operator()(Tuple t)
 
236
  {
 
237
    thrust::get<1>(t) = f(thrust::get<0>(t));
 
238
  }
 
239
};
 
240
 
 
241
template<typename UnaryFunction>
 
242
  struct device_unary_transform_functor
 
243
{
 
244
  typedef void result_type;
 
245
 
 
246
  UnaryFunction f;
 
247
 
 
248
  device_unary_transform_functor(UnaryFunction f_)
 
249
    :f(f_) {}
 
250
 
 
251
  // add __host__ to allow the omp backend compile with nvcc
 
252
  template<typename Tuple>
 
253
  __host__ __device__
 
254
  inline result_type operator()(Tuple t)
 
255
  {
 
256
    thrust::get<1>(t) = f(thrust::get<0>(t));
 
257
  }
 
258
};
 
259
 
 
260
 
 
261
template<typename Space, typename UnaryFunction>
 
262
  struct unary_transform_functor
 
263
    : thrust::detail::eval_if<
 
264
        thrust::detail::is_convertible<Space, thrust::host_space_tag>::value,
 
265
        thrust::detail::identity_<host_unary_transform_functor<UnaryFunction> >,
 
266
        thrust::detail::identity_<device_unary_transform_functor<UnaryFunction> >
 
267
      >
 
268
{};
 
269
 
 
270
 
 
271
template <typename BinaryFunction>
 
272
struct host_binary_transform_functor
 
273
{
 
274
  BinaryFunction f;
 
275
 
 
276
  host_binary_transform_functor(BinaryFunction f_)
 
277
    :f(f_)
 
278
  {}
 
279
 
 
280
  template <typename Tuple>
 
281
  __host__
 
282
  void operator()(Tuple t)
 
283
  { 
 
284
    thrust::get<2>(t) = f(thrust::get<0>(t), thrust::get<1>(t));
 
285
  }
 
286
}; // end binary_transform_functor
 
287
 
 
288
 
 
289
template <typename BinaryFunction>
 
290
struct device_binary_transform_functor
 
291
{
 
292
  BinaryFunction f;
 
293
 
 
294
  device_binary_transform_functor(BinaryFunction f_)
 
295
    :f(f_)
 
296
  {}
 
297
 
 
298
  // add __host__ to allow the omp backend compile with nvcc
 
299
  template <typename Tuple>
 
300
  __host__ __device__
 
301
  void operator()(Tuple t)
 
302
  { 
 
303
    thrust::get<2>(t) = f(thrust::get<0>(t), thrust::get<1>(t));
 
304
  }
 
305
}; // end binary_transform_functor
 
306
 
 
307
 
 
308
template<typename Space, typename BinaryFunction>
 
309
  struct binary_transform_functor
 
310
    : thrust::detail::eval_if<
 
311
        thrust::detail::is_convertible<Space, thrust::host_space_tag>::value,
 
312
        thrust::detail::identity_<host_binary_transform_functor<BinaryFunction> >,
 
313
        thrust::detail::identity_<device_binary_transform_functor<BinaryFunction> >
 
314
      >
 
315
{};
 
316
 
 
317
 
 
318
template <typename UnaryFunction, typename Predicate>
 
319
struct host_unary_transform_if_functor
 
320
{
 
321
  UnaryFunction unary_op;
 
322
  Predicate pred;
 
323
  
 
324
  host_unary_transform_if_functor(UnaryFunction _unary_op, Predicate _pred)
 
325
    : unary_op(_unary_op), pred(_pred) {} 
 
326
  
 
327
  template <typename Tuple>
 
328
  __host__
 
329
  void operator()(Tuple t)
 
330
  {
 
331
    if(pred(thrust::get<1>(t)))
 
332
      thrust::get<2>(t) = unary_op(thrust::get<0>(t));
 
333
  }
 
334
}; // end host_unary_transform_if_functor
 
335
 
 
336
 
 
337
template <typename UnaryFunction, typename Predicate>
 
338
struct device_unary_transform_if_functor
 
339
{
 
340
  UnaryFunction unary_op;
 
341
  Predicate pred;
 
342
  
 
343
  device_unary_transform_if_functor(UnaryFunction _unary_op, Predicate _pred)
 
344
    : unary_op(_unary_op), pred(_pred) {} 
 
345
  
 
346
  // add __host__ to allow the omp backend compile with nvcc
 
347
  template <typename Tuple>
 
348
  __host__ __device__
 
349
  void operator()(Tuple t)
 
350
  {
 
351
    if(pred(thrust::get<1>(t)))
 
352
      thrust::get<2>(t) = unary_op(thrust::get<0>(t));
 
353
  }
 
354
}; // end device_unary_transform_if_functor
 
355
 
 
356
 
 
357
template<typename Space, typename UnaryFunction, typename Predicate>
 
358
  struct unary_transform_if_functor
 
359
    : thrust::detail::eval_if<
 
360
        thrust::detail::is_convertible<Space, thrust::host_space_tag>::value,
 
361
        thrust::detail::identity_<host_unary_transform_if_functor<UnaryFunction,Predicate> >,
 
362
        thrust::detail::identity_<device_unary_transform_if_functor<UnaryFunction,Predicate> >
 
363
      >
 
364
{};
 
365
 
 
366
 
 
367
template <typename BinaryFunction, typename Predicate>
 
368
struct host_binary_transform_if_functor
 
369
{
 
370
  BinaryFunction binary_op;
 
371
  Predicate pred;
 
372
 
 
373
  host_binary_transform_if_functor(BinaryFunction _binary_op, Predicate _pred)
 
374
    : binary_op(_binary_op), pred(_pred) {} 
 
375
 
 
376
  template <typename Tuple>
 
377
  __host__
 
378
  void operator()(Tuple t)
 
379
  {
 
380
    if(pred(thrust::get<2>(t)))
 
381
      thrust::get<3>(t) = binary_op(thrust::get<0>(t), thrust::get<1>(t));
 
382
  }
 
383
}; // end host_binary_transform_if_functor
 
384
 
 
385
 
 
386
template <typename BinaryFunction, typename Predicate>
 
387
struct device_binary_transform_if_functor
 
388
{
 
389
  BinaryFunction binary_op;
 
390
  Predicate pred;
 
391
 
 
392
  device_binary_transform_if_functor(BinaryFunction _binary_op, Predicate _pred)
 
393
    : binary_op(_binary_op), pred(_pred) {} 
 
394
 
 
395
  // add __host__ to allow the omp backend compile with nvcc
 
396
  template <typename Tuple>
 
397
  __host__ __device__
 
398
  void operator()(Tuple t)
 
399
  {
 
400
    if(pred(thrust::get<2>(t)))
 
401
      thrust::get<3>(t) = binary_op(thrust::get<0>(t), thrust::get<1>(t));
 
402
  }
 
403
}; // end device_binary_transform_if_functor
 
404
 
 
405
 
 
406
template<typename Space, typename BinaryFunction, typename Predicate>
 
407
  struct binary_transform_if_functor
 
408
    : thrust::detail::eval_if<
 
409
        thrust::detail::is_convertible<Space, thrust::host_space_tag>::value,
 
410
        thrust::detail::identity_<host_binary_transform_if_functor<BinaryFunction,Predicate> >,
 
411
        thrust::detail::identity_<device_binary_transform_if_functor<BinaryFunction,Predicate> >
 
412
      >
 
413
{};
 
414
 
 
415
 
 
416
template<typename T>
 
417
  struct host_destroy_functor
 
418
{
 
419
  __host__
 
420
  void operator()(T &x) const
 
421
  {
 
422
    x.~T();
 
423
  } // end operator()()
 
424
}; // end host_destroy_functor
 
425
 
 
426
 
 
427
template<typename T>
 
428
  struct device_destroy_functor
 
429
{
 
430
  // add __host__ to allow the omp backend to compile with nvcc
 
431
  __host__ __device__
 
432
  void operator()(T &x) const
 
433
  {
 
434
    x.~T();
 
435
  } // end operator()()
 
436
}; // end device_destroy_functor
 
437
 
 
438
 
 
439
template<typename Space, typename T>
 
440
  struct destroy_functor
 
441
    : thrust::detail::eval_if<
 
442
        thrust::detail::is_convertible<Space, thrust::host_space_tag>::value,
 
443
        thrust::detail::identity_<host_destroy_functor<T> >,
 
444
        thrust::detail::identity_<device_destroy_functor<T> >
 
445
      >
 
446
{};
 
447
 
 
448
 
 
449
template <typename T>
 
450
struct fill_functor
 
451
{
 
452
  const T exemplar;
 
453
 
 
454
  fill_functor(const T& _exemplar) 
 
455
    : exemplar(_exemplar) {}
 
456
 
 
457
  __host__ __device__
 
458
  T operator()(void) const
 
459
  { 
 
460
    return exemplar;
 
461
  }
 
462
};
 
463
 
 
464
 
 
465
template<typename T>
 
466
  struct uninitialized_fill_functor
 
467
{
 
468
  T exemplar;
 
469
 
 
470
  uninitialized_fill_functor(T x):exemplar(x){}
 
471
 
 
472
  __host__ __device__
 
473
  void operator()(T &x)
 
474
  {
 
475
    ::new(static_cast<void*>(&x)) T(exemplar);
 
476
  } // end operator()()
 
477
}; // end uninitialized_fill_functor
 
478
 
 
479
 
 
480
// this predicate tests two two-element tuples
 
481
// we first use a Compare for the first element
 
482
// if the first elements are equivalent, we use
 
483
// < for the second elements
 
484
template<typename Compare>
 
485
  struct compare_first_less_second
 
486
{
 
487
  compare_first_less_second(Compare c)
 
488
    : comp(c) {}
 
489
 
 
490
  template<typename T1, typename T2>
 
491
  __host__ __device__
 
492
  bool operator()(T1 lhs, T2 rhs)
 
493
  {
 
494
    return comp(thrust::get<0>(lhs), thrust::get<0>(rhs)) || (!comp(thrust::get<0>(rhs), thrust::get<0>(lhs)) && thrust::get<1>(lhs) < thrust::get<1>(rhs));
 
495
  }
 
496
 
 
497
  Compare comp;
 
498
}; // end compare_first_less_second
 
499
 
 
500
 
 
501
} // end namespace detail
 
502
} // end namespace thrust
 
503