← Back to branch summary

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

« back to all changes in this revision

Viewing changes to detail/backend/cuda/set_intersection.inl

  • Committer: Package Import Robot
  • Author(s): Andreas Beckmann
  • Date: 2011-12-02 01:48:24 UTC
  • mfrom: (1.1.1)
  • Revision ID: package-import@ubuntu.com-20111202014824-bpfczhbx39usefge
Tags: 1.5.0-1
* New upstream release.
* debian/copyright:
  - Update to dep5.mdwn?revision=202.
  - Update copyright entries for added/moved files.

Show diffs side-by-side

added added

removed removed

Lines of Context:
detail/backend/cuda/set_intersection.inl
 
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
#include <thrust/detail/config.h>
 
18
 
 
19
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
 
20
 
 
21
#include <thrust/iterator/iterator_traits.h>
 
22
#include <thrust/pair.h>
 
23
#include <thrust/detail/backend/cuda/block/set_intersection.h>
 
24
#include <thrust/detail/backend/cuda/detail/split_for_set_operation.h>
 
25
#include <thrust/detail/backend/cuda/detail/set_operation.h>
 
26
 
 
27
namespace thrust
 
28
{
 
29
namespace detail
 
30
{
 
31
namespace backend
 
32
{
 
33
namespace cuda
 
34
{
 
35
 
 
36
namespace set_intersection_detail
 
37
{
 
38
 
 
39
struct block_convergent_set_intersection_functor
 
40
{
 
41
  __host__ __device__ __forceinline__
 
42
  static size_t get_min_size_of_result_in_number_of_elements(size_t size_of_range1,
 
43
                                                             size_t size_of_range2)
 
44
  {
 
45
    // set_intersection could result in zero output
 
46
    return 0u;
 
47
  }
 
48
 
 
49
  __host__ __device__ __forceinline__
 
50
  static size_t get_max_size_of_result_in_number_of_elements(size_t size_of_range1,
 
51
                                                             size_t size_of_range2)
 
52
  {
 
53
    // set_intersection could output all of range1
 
54
    return size_of_range1;
 
55
  }
 
56
 
 
57
  __host__ __device__ __forceinline__
 
58
  static unsigned int get_temporary_array_size_in_number_of_bytes(unsigned int block_size)
 
59
  {
 
60
    return block_size * sizeof(int);
 
61
  }
 
62
 
 
63
  // operator() simply calls the block-wise function
 
64
  template<typename RandomAccessIterator1,
 
65
           typename RandomAccessIterator2,
 
66
           typename RandomAccessIterator3,
 
67
           typename StrictWeakOrdering>
 
68
  __device__ __forceinline__
 
69
    RandomAccessIterator3 operator()(RandomAccessIterator1 first1,
 
70
                                     RandomAccessIterator1 last1,
 
71
                                     RandomAccessIterator2 first2,
 
72
                                     RandomAccessIterator2 last2,
 
73
                                     void *temporary,
 
74
                                     RandomAccessIterator3 result,
 
75
                                     StrictWeakOrdering comp)
 
76
  {
 
77
    return block::set_intersection(first1,last1,first2,last2,reinterpret_cast<int*>(temporary),result,comp);
 
78
  } // end operator()()
 
79
}; // end block_convergent_set_intersection_functor
 
80
 
 
81
} // end namespace set_intersection_detail
 
82
 
 
83
 
 
84
template<typename RandomAccessIterator1,
 
85
         typename RandomAccessIterator2, 
 
86
         typename RandomAccessIterator3,
 
87
         typename Compare>
 
88
RandomAccessIterator3 set_intersection(RandomAccessIterator1 first1,
 
89
                                       RandomAccessIterator1 last1,
 
90
                                       RandomAccessIterator2 first2,
 
91
                                       RandomAccessIterator2 last2,
 
92
                                       RandomAccessIterator3 result,
 
93
                                       Compare comp)
 
94
{
 
95
  typedef typename thrust::iterator_difference<RandomAccessIterator1>::type difference1;
 
96
  typedef typename thrust::iterator_difference<RandomAccessIterator2>::type difference2;
 
97
 
 
98
  const difference1 num_elements1 = last1 - first1;
 
99
  const difference2 num_elements2 = last2 - first2;
 
100
 
 
101
  // check for trivial problem
 
102
  if(num_elements1 == 0 || num_elements2 == 0)
 
103
    return result;
 
104
 
 
105
  return detail::set_operation(first1, last1,
 
106
                               first2, last2,
 
107
                               result,
 
108
                               comp,
 
109
                               detail::split_for_set_operation(),
 
110
                               set_intersection_detail::block_convergent_set_intersection_functor());
 
111
} // end set_intersection
 
112
 
 
113
} // end namespace cuda
 
114
} // end namespace backend
 
115
} // end namespace detail
 
116
} // end namespace thrust
 
117
 
 
118
#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
 
119