← Back to branch summary

~ubuntu-branches/ubuntu/utopic/libthrust/utopic

« back to all changes in this revision

Viewing changes to system/cuda/detail/temporary_indirect_permutation.h

  • Committer: Package Import Robot
  • Author(s): Andreas Beckmann
  • Date: 2013-07-10 12:57:33 UTC
  • mfrom: (1.1.4)
  • Revision ID: package-import@ubuntu.com-20130710125733-my19jic71sqsabaj
Tags: 1.7.0-1
http://bugs.debian.org/715362
* New upstream release.  (Closes: #715362)
* Update watch file.

Show diffs side-by-side

added added

removed removed

Lines of Context:
system/cuda/detail/temporary_indirect_permutation.h
 
1
/*
 
2
 *  Copyright 2008-2012 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/temporary_array.h>
 
20
#include <thrust/sequence.h>
 
21
#include <thrust/gather.h>
 
22
#include <thrust/detail/function.h>
 
23
 
 
24
namespace thrust
 
25
{
 
26
namespace system
 
27
{
 
28
namespace cuda
 
29
{
 
30
namespace detail
 
31
{
 
32
 
 
33
 
 
34
template<typename DerivedPolicy, typename RandomAccessIterator>
 
35
  struct temporary_indirect_permutation
 
36
{
 
37
  private:
 
38
    typedef unsigned int size_type;
 
39
    typedef thrust::detail::temporary_array<size_type, DerivedPolicy> array_type;
 
40
 
 
41
  public:
 
42
    temporary_indirect_permutation(DerivedPolicy &exec, RandomAccessIterator first, RandomAccessIterator last)
 
43
      : m_exec(exec),
 
44
        m_src_first(first),
 
45
        m_src_last(last),
 
46
        m_permutation(0, m_exec, last - first)
 
47
    {
 
48
      // generate sorted index sequence
 
49
      thrust::sequence(exec, m_permutation.begin(), m_permutation.end());
 
50
    }
 
51
 
 
52
    ~temporary_indirect_permutation()
 
53
    {
 
54
      // permute the source array using the indices
 
55
      typedef typename thrust::iterator_value<RandomAccessIterator>::type value_type;
 
56
      thrust::detail::temporary_array<value_type, DerivedPolicy> temp(m_exec, m_src_first, m_src_last);
 
57
      thrust::gather(m_exec, m_permutation.begin(), m_permutation.end(), temp.begin(), m_src_first);
 
58
    }
 
59
 
 
60
    typedef typename array_type::iterator iterator;
 
61
 
 
62
    iterator begin()
 
63
    {
 
64
      return m_permutation.begin();
 
65
    }
 
66
 
 
67
    iterator end()
 
68
    {
 
69
      return m_permutation.end();
 
70
    }
 
71
 
 
72
  private:
 
73
    DerivedPolicy &m_exec;
 
74
    RandomAccessIterator m_src_first, m_src_last;
 
75
    thrust::detail::temporary_array<size_type, DerivedPolicy> m_permutation;
 
76
};
 
77
 
 
78
 
 
79
template<typename DerivedPolicy, typename RandomAccessIterator>
 
80
  struct iterator_range_with_execution_policy
 
81
{
 
82
  iterator_range_with_execution_policy(DerivedPolicy &exec, RandomAccessIterator first, RandomAccessIterator last)
 
83
    : m_exec(exec), m_first(first), m_last(last)
 
84
  {}
 
85
 
 
86
  typedef RandomAccessIterator iterator;
 
87
 
 
88
  iterator begin()
 
89
  {
 
90
    return m_first;
 
91
  }
 
92
 
 
93
  iterator end()
 
94
  {
 
95
    return m_last;
 
96
  }
 
97
 
 
98
  DerivedPolicy &exec()
 
99
  {
 
100
    return m_exec;
 
101
  }
 
102
 
 
103
  DerivedPolicy &m_exec;
 
104
  RandomAccessIterator m_first, m_last;
 
105
};
 
106
 
 
107
 
 
108
template<typename Condition, typename DerivedPolicy, typename RandomAccessIterator>
 
109
  struct conditional_temporary_indirect_permutation
 
110
    : thrust::detail::eval_if<
 
111
        Condition::value,
 
112
        thrust::detail::identity_<temporary_indirect_permutation<DerivedPolicy, RandomAccessIterator> >,
 
113
        thrust::detail::identity_<iterator_range_with_execution_policy<DerivedPolicy, RandomAccessIterator> >
 
114
      >::type
 
115
{
 
116
  typedef typename thrust::detail::eval_if<
 
117
    Condition::value,
 
118
    thrust::detail::identity_<temporary_indirect_permutation<DerivedPolicy, RandomAccessIterator> >,
 
119
    thrust::detail::identity_<iterator_range_with_execution_policy<DerivedPolicy, RandomAccessIterator> >
 
120
  >::type super_t;
 
121
 
 
122
  conditional_temporary_indirect_permutation(DerivedPolicy &exec, RandomAccessIterator first, RandomAccessIterator last)
 
123
    : super_t(exec, first, last)
 
124
  {}
 
125
};
 
126
 
 
127
 
 
128
template<typename DerivedPolicy, typename RandomAccessIterator, typename Compare>
 
129
  struct temporary_indirect_ordering
 
130
    : temporary_indirect_permutation<DerivedPolicy,RandomAccessIterator>
 
131
{
 
132
  private:
 
133
    typedef temporary_indirect_permutation<DerivedPolicy,RandomAccessIterator> super_t;
 
134
 
 
135
  public:
 
136
    temporary_indirect_ordering(DerivedPolicy &exec, RandomAccessIterator first, RandomAccessIterator last, Compare comp)
 
137
      : super_t(exec, first, last),
 
138
        m_comp(first, comp)
 
139
    {}
 
140
 
 
141
    struct compare
 
142
    {
 
143
      RandomAccessIterator first;
 
144
 
 
145
      thrust::detail::host_device_function<
 
146
        Compare,
 
147
        bool
 
148
      > comp;
 
149
 
 
150
      compare(RandomAccessIterator first, Compare comp)
 
151
        : first(first), comp(comp)
 
152
      {}
 
153
 
 
154
      template<typename Integral>
 
155
      __host__ __device__
 
156
      bool operator()(Integral a, Integral b)
 
157
      {
 
158
        return comp(first[a], first[b]);
 
159
      }
 
160
    };
 
161
 
 
162
    compare comp() const
 
163
    {
 
164
      return m_comp;
 
165
    }
 
166
 
 
167
  private:
 
168
    compare m_comp;
 
169
};
 
170
 
 
171
 
 
172
template<typename DerivedPolicy, typename RandomAccessIterator, typename Compare>
 
173
  struct iterator_range_with_execution_policy_and_compare
 
174
    : iterator_range_with_execution_policy<DerivedPolicy, RandomAccessIterator>
 
175
{
 
176
  typedef iterator_range_with_execution_policy<DerivedPolicy, RandomAccessIterator> super_t;
 
177
 
 
178
  iterator_range_with_execution_policy_and_compare(DerivedPolicy &exec, RandomAccessIterator first, RandomAccessIterator last, Compare comp)
 
179
    : super_t(exec, first, last), m_comp(comp)
 
180
  {}
 
181
 
 
182
  typedef Compare compare;
 
183
 
 
184
  compare comp()
 
185
  {
 
186
    return m_comp;
 
187
  }
 
188
 
 
189
  Compare m_comp;
 
190
};
 
191
 
 
192
 
 
193
template<typename Condition, typename DerivedPolicy, typename RandomAccessIterator, typename Compare>
 
194
  struct conditional_temporary_indirect_ordering
 
195
    : thrust::detail::eval_if<
 
196
        Condition::value,
 
197
        thrust::detail::identity_<temporary_indirect_ordering<DerivedPolicy, RandomAccessIterator, Compare> >,
 
198
        thrust::detail::identity_<iterator_range_with_execution_policy_and_compare<DerivedPolicy, RandomAccessIterator, Compare> >
 
199
      >::type
 
200
{
 
201
  typedef typename thrust::detail::eval_if<
 
202
    Condition::value,
 
203
    thrust::detail::identity_<temporary_indirect_ordering<DerivedPolicy, RandomAccessIterator, Compare> >,
 
204
    thrust::detail::identity_<iterator_range_with_execution_policy_and_compare<DerivedPolicy, RandomAccessIterator, Compare> >
 
205
  >::type super_t;
 
206
 
 
207
  conditional_temporary_indirect_ordering(DerivedPolicy &exec, RandomAccessIterator first, RandomAccessIterator last, Compare comp)
 
208
    : super_t(exec, first, last, comp)
 
209
  {}
 
210
};
 
211
 
 
212
 
 
213
} // end detail
 
214
} // end cuda
 
215
} // end system
 
216
} // end thrust
 
217