util/X86LINUX64/cuda-6.5/include/thrust/system/detail/internal/decompose.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/detail/config.h>
  20
  21 namespace thrust
  22 {
  23 namespace system
  24 {
  25 namespace detail
  26 {
  27 namespace internal
  28 {
  29
  30   template <typename IndexType>
  31     class index_range
  32     {
  33       public:
  34         typedef IndexType index_type;
  35
  36         __host__ __device__
  37           index_range(index_type begin, index_type end) : m_begin(begin), m_end(end) {}
  38
  39         __host__ __device__
  40           index_type begin(void) const { return m_begin; }
  41
  42         __host__ __device__
  43           index_type end(void)   const { return m_end; }
  44
  45         __host__ __device__
  46           index_type size(void)  const { return m_end - m_begin; }
  47
  48       private:
  49         index_type m_begin;
  50         index_type m_end;
  51     };
  52
  53   template <typename IndexType>
  54     class uniform_decomposition
  55     {
  56       public:
  57         typedef IndexType               index_type;
  58         typedef index_range<index_type> range_type;
  59
  60         uniform_decomposition(index_type N, index_type granularity, index_type max_intervals)
  61           : m_N(N),
  62             m_intervals((N + granularity - 1) / granularity),
  63             m_threshold(0),
  64             m_small_interval(granularity),
  65             m_large_interval(0)
  66         {
  67           if(m_intervals > max_intervals)
  68           {
  69             m_small_interval = granularity * (m_intervals / max_intervals);
  70             m_large_interval = m_small_interval + granularity;
  71             m_threshold      = m_intervals % max_intervals;
  72             m_intervals      = max_intervals;
  73           }
  74         }
  75
  76         __host__ __device__
  77           index_range<index_type> operator[](const index_type& i) const
  78           {
  79             if (i < m_threshold)
  80             {
  81               index_type begin = m_large_interval * i;
  82               index_type end   = begin + m_large_interval;
  83               return range_type(begin, end);
  84             }
  85             else
  86             {
  87               index_type begin = m_large_interval * m_threshold + m_small_interval * (i - m_threshold);
  88               index_type end   = (begin + m_small_interval < m_N) ? begin + m_small_interval : m_N;
  89               return range_type(begin, end);
  90             }
  91           }
  92
  93         __host__ __device__
  94           index_type size(void) const
  95           {
  96             return m_intervals;
  97           }
  98
  99       private:
 100
 101         index_type m_N;
 102         index_type m_intervals;
 103         index_type m_threshold;
 104         index_type m_small_interval;
 105         index_type m_large_interval;
 106     };
 107
 108
 109 } // end namespace internal
 110 } // end namespace detail
 111 } // end namespace system
 112 } // end namespace thrust
 113