util/X86LINUX64/cuda-6.5/include/thrust/iterator/detail/zip_iterator_base.h

   1 /*
   2  *  Copyright 2008-2013 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/iterator/iterator_traits.h>
  20 #include <thrust/iterator/iterator_facade.h>
  21 #include <thrust/iterator/iterator_categories.h>
  22 #include <thrust/iterator/detail/minimum_category.h>
  23 #include <thrust/iterator/detail/minimum_system.h>
  24 #include <thrust/tuple.h>
  25 #include <thrust/detail/tuple_meta_transform.h>
  26 #include <thrust/detail/tuple_transform.h>
  27 #include <thrust/detail/type_traits.h>
  28 #include <thrust/iterator/detail/tuple_of_iterator_references.h>
  29
  30 namespace thrust
  31 {
  32
  33 // forward declare zip_iterator for zip_iterator_base
  34 template<typename IteratorTuple> class zip_iterator;
  35
  36 namespace detail
  37 {
  38
  39
  40 // Functors to be used with tuple algorithms
  41 //
  42 template<typename DiffType>
  43 class advance_iterator
  44 {
  45 public:
  46   inline __host__ __device__
  47   advance_iterator(DiffType step) : m_step(step) {}
  48
  49   template<typename Iterator>
  50   inline __host__ __device__
  51   void operator()(Iterator& it) const
  52   { it += m_step; }
  53
  54 private:
  55   DiffType m_step;
  56 }; // end advance_iterator
  57
  58
  59 struct increment_iterator
  60 {
  61   template<typename Iterator>
  62   inline __host__ __device__
  63   void operator()(Iterator& it)
  64   { ++it; }
  65 }; // end increment_iterator
  66
  67
  68 struct decrement_iterator
  69 {
  70   template<typename Iterator>
  71   inline __host__ __device__
  72   void operator()(Iterator& it)
  73   { --it; }
  74 }; // end decrement_iterator
  75
  76
  77 struct dereference_iterator
  78 {
  79   template<typename Iterator>
  80   struct apply
  81   {
  82     typedef typename
  83       iterator_traits<Iterator>::reference
  84     type;
  85   }; // end apply
  86
  87   // XXX silence warnings of the form "calling a __host__ function from a __host__ __device__ function is not allowed
  88   __thrust_hd_warning_disable__
  89   template<typename Iterator>
  90   __host__ __device__
  91     typename apply<Iterator>::type operator()(Iterator const& it)
  92   {
  93     return *it;
  94   }
  95 }; // end dereference_iterator
  96
  97
  98 // The namespace tuple_impl_specific provides two meta-
  99 // algorithms and two algorithms for tuples.
 100 namespace tuple_impl_specific
 101 {
 102
 103 // define apply1 for tuple_meta_transform_impl
 104 template<typename UnaryMetaFunctionClass, class Arg>
 105   struct apply1
 106     : UnaryMetaFunctionClass::template apply<Arg>
 107 {
 108 }; // end apply1
 109
 110
 111 // define apply2 for tuple_meta_accumulate_impl
 112 template<typename UnaryMetaFunctionClass, class Arg1, class Arg2>
 113   struct apply2
 114     : UnaryMetaFunctionClass::template apply<Arg1,Arg2>
 115 {
 116 }; // end apply2
 117
 118
 119 // Meta-accumulate algorithm for tuples. Note: The template
 120 // parameter StartType corresponds to the initial value in
 121 // ordinary accumulation.
 122 //
 123 template<class Tuple, class BinaryMetaFun, class StartType>
 124   struct tuple_meta_accumulate;
 125
 126 template<
 127     typename Tuple
 128   , class BinaryMetaFun
 129   , typename StartType
 130 >
 131   struct tuple_meta_accumulate_impl
 132 {
 133    typedef typename apply2<
 134        BinaryMetaFun
 135      , typename Tuple::head_type
 136      , typename tuple_meta_accumulate<
 137            typename Tuple::tail_type
 138          , BinaryMetaFun
 139          , StartType
 140        >::type
 141    >::type type;
 142 };
 143
 144
 145 template<
 146     typename Tuple
 147   , class BinaryMetaFun
 148   , typename StartType
 149 >
 150 struct tuple_meta_accumulate
 151   : thrust::detail::eval_if<
 152         thrust::detail::is_same<Tuple, thrust::null_type>::value
 153       , thrust::detail::identity_<StartType>
 154       , tuple_meta_accumulate_impl<
 155             Tuple
 156           , BinaryMetaFun
 157           , StartType
 158         >
 159     > // end eval_if
 160 {
 161 }; // end tuple_meta_accumulate
 162
 163
 164 // transform algorithm for tuples. The template parameter Fun
 165 // must be a unary functor which is also a unary metafunction
 166 // class that computes its return type based on its argument
 167 // type. For example:
 168 //
 169 // struct to_ptr
 170 // {
 171 //     template <class Arg>
 172 //     struct apply
 173 //     {
 174 //          typedef Arg* type;
 175 //     }
 176 //
 177 //     template <class Arg>
 178 //     Arg* operator()(Arg x);
 179 // };
 180
 181
 182
 183 // for_each algorithm for tuples.
 184 //
 185 template<typename Fun, typename System>
 186 inline __host__ __device__
 187 Fun tuple_for_each(thrust::null_type, Fun f, System *)
 188 {
 189   return f;
 190 } // end tuple_for_each()
 191
 192
 193 template<typename Tuple, typename Fun, typename System>
 194 inline __host__ __device__
 195 Fun tuple_for_each(Tuple& t, Fun f, System *dispatch_tag)
 196 {
 197   f( t.get_head() );
 198   return tuple_for_each(t.get_tail(), f, dispatch_tag);
 199 } // end tuple_for_each()
 200
 201
 202 template<typename Tuple, typename Fun>
 203 inline __host__ __device__
 204 Fun tuple_for_each(Tuple& t, Fun f, thrust::host_system_tag *dispatch_tag)
 205 {
 206 // XXX this path is required in order to accomodate pure host iterators
 207 //     (such as std::vector::iterator) in a zip_iterator
 208 #ifndef __CUDA_ARCH__
 209   f( t.get_head() );
 210   return tuple_for_each(t.get_tail(), f, dispatch_tag);
 211 #else
 212   // this code will never be called
 213   return f;
 214 #endif
 215 } // end tuple_for_each()
 216
 217
 218 // Equality of tuples. NOTE: "==" for tuples currently (7/2003)
 219 // has problems under some compilers, so I just do my own.
 220 // No point in bringing in a bunch of #ifdefs here. This is
 221 // going to go away with the next tuple implementation anyway.
 222 //
 223 __host__ __device__
 224 inline bool tuple_equal(thrust::null_type, thrust::null_type)
 225 { return true; }
 226
 227
 228 template<typename Tuple1, typename Tuple2>
 229 __host__ __device__
 230 bool tuple_equal(Tuple1 const& t1, Tuple2 const& t2)
 231 {
 232   return t1.get_head() == t2.get_head() &&
 233   tuple_equal(t1.get_tail(), t2.get_tail());
 234 } // end tuple_equal()
 235
 236 } // end end tuple_impl_specific
 237
 238
 239 // Metafunction to obtain the type of the tuple whose element types
 240 // are the value_types of an iterator tupel.
 241 //
 242 template<typename IteratorTuple>
 243   struct tuple_of_value_types
 244     : tuple_meta_transform<
 245           IteratorTuple,
 246           iterator_value
 247         >
 248 {
 249 }; // end tuple_of_value_types
 250
 251
 252 struct minimum_category_lambda
 253 {
 254   template<typename T1, typename T2>
 255     struct apply : minimum_category<T1,T2>
 256   {};
 257 };
 258
 259
 260
 261 // Metafunction to obtain the minimal traversal tag in a tuple
 262 // of iterators.
 263 //
 264 template<typename IteratorTuple>
 265 struct minimum_traversal_category_in_iterator_tuple
 266 {
 267   typedef typename tuple_meta_transform<
 268       IteratorTuple
 269     , thrust::iterator_traversal
 270   >::type tuple_of_traversal_tags;
 271
 272   typedef typename tuple_impl_specific::tuple_meta_accumulate<
 273       tuple_of_traversal_tags
 274     , minimum_category_lambda
 275     , thrust::random_access_traversal_tag
 276   >::type type;
 277 };
 278
 279
 280 struct minimum_system_lambda
 281 {
 282   template<typename T1, typename T2>
 283     struct apply : minimum_system<T1,T2>
 284   {};
 285 };
 286
 287
 288
 289 // Metafunction to obtain the minimal system tag in a tuple
 290 // of iterators.
 291 template<typename IteratorTuple>
 292 struct minimum_system_in_iterator_tuple
 293 {
 294   typedef typename thrust::detail::tuple_meta_transform<
 295     IteratorTuple,
 296     thrust::iterator_system
 297   >::type tuple_of_system_tags;
 298
 299   typedef typename tuple_impl_specific::tuple_meta_accumulate<
 300     tuple_of_system_tags,
 301     minimum_system_lambda,
 302     thrust::any_system_tag
 303   >::type type;
 304 };
 305
 306 namespace zip_iterator_base_ns
 307 {
 308
 309
 310 template<int i, typename Tuple>
 311   struct tuple_elements_helper
 312     : eval_if<
 313         (i < tuple_size<Tuple>::value),
 314         tuple_element<i,Tuple>,
 315         identity_<thrust::null_type>
 316       >
 317 {};
 318
 319
 320 template<typename Tuple>
 321   struct tuple_elements
 322 {
 323   typedef typename tuple_elements_helper<0,Tuple>::type T0;
 324   typedef typename tuple_elements_helper<1,Tuple>::type T1;
 325   typedef typename tuple_elements_helper<2,Tuple>::type T2;
 326   typedef typename tuple_elements_helper<3,Tuple>::type T3;
 327   typedef typename tuple_elements_helper<4,Tuple>::type T4;
 328   typedef typename tuple_elements_helper<5,Tuple>::type T5;
 329   typedef typename tuple_elements_helper<6,Tuple>::type T6;
 330   typedef typename tuple_elements_helper<7,Tuple>::type T7;
 331   typedef typename tuple_elements_helper<8,Tuple>::type T8;
 332   typedef typename tuple_elements_helper<9,Tuple>::type T9;
 333 };
 334
 335
 336 template<typename IteratorTuple>
 337   struct tuple_of_iterator_references
 338 {
 339   // get a thrust::tuple of the iterators' references
 340   typedef typename tuple_meta_transform<
 341     IteratorTuple,
 342     iterator_reference
 343   >::type tuple_of_references;
 344
 345   // get at the individual tuple element types by name
 346   typedef tuple_elements<tuple_of_references> elements;
 347
 348   // map thrust::tuple<T...> to tuple_of_iterator_references<T...>
 349   typedef thrust::detail::tuple_of_iterator_references<
 350     typename elements::T0,
 351     typename elements::T1,
 352     typename elements::T2,
 353     typename elements::T3,
 354     typename elements::T4,
 355     typename elements::T5,
 356     typename elements::T6,
 357     typename elements::T7,
 358     typename elements::T8,
 359     typename elements::T9
 360   > type;
 361 };
 362
 363
 364 } // end zip_iterator_base_ns
 365
 366 ///////////////////////////////////////////////////////////////////
 367 //
 368 // Class zip_iterator_base
 369 //
 370 // Builds and exposes the iterator facade type from which the zip
 371 // iterator will be derived.
 372 //
 373 template<typename IteratorTuple>
 374   struct zip_iterator_base
 375 {
 376  //private:
 377     // reference type is the type of the tuple obtained from the
 378     // iterators' reference types.
 379     typedef typename zip_iterator_base_ns::tuple_of_iterator_references<IteratorTuple>::type reference;
 380
 381     // Boost's Value type is the same as reference type.
 382     //typedef reference value_type;
 383     typedef typename tuple_of_value_types<IteratorTuple>::type value_type;
 384
 385     // Difference type is the first iterator's difference type
 386     typedef typename thrust::iterator_traits<
 387       typename thrust::tuple_element<0, IteratorTuple>::type
 388     >::difference_type difference_type;
 389
 390     // Iterator system is the minimum system tag in the
 391     // iterator tuple
 392     typedef typename
 393     minimum_system_in_iterator_tuple<IteratorTuple>::type system;
 394
 395     // Traversal category is the minimum traversal category in the
 396     // iterator tuple
 397     typedef typename
 398     minimum_traversal_category_in_iterator_tuple<IteratorTuple>::type traversal_category;
 399
 400  public:
 401
 402     // The iterator facade type from which the zip iterator will
 403     // be derived.
 404     typedef thrust::iterator_facade<
 405         zip_iterator<IteratorTuple>,
 406         value_type,
 407         system,
 408         traversal_category,
 409         reference,
 410         difference_type
 411     > type;
 412 }; // end zip_iterator_base
 413
 414 } // end detail
 415
 416 } // end thrust
 417
 418