utiltools/X86MAC64/cuda/include/thrust/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