Stokhos Package Browser (Single Doxygen Collection)  Version of the Day
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Groups Pages
Kokkos_Parallel_MP_Vector.hpp
Go to the documentation of this file.
1 // @HEADER
2 // ***********************************************************************
3 //
4 // Stokhos Package
5 // Copyright (2009) Sandia Corporation
6 //
7 // Under terms of Contract DE-AC04-94AL85000, there is a non-exclusive
8 // license for use of this work by or on behalf of the U.S. Government.
9 //
10 // Redistribution and use in source and binary forms, with or without
11 // modification, are permitted provided that the following conditions are
12 // met:
13 //
14 // 1. Redistributions of source code must retain the above copyright
15 // notice, this list of conditions and the following disclaimer.
16 //
17 // 2. Redistributions in binary form must reproduce the above copyright
18 // notice, this list of conditions and the following disclaimer in the
19 // documentation and/or other materials provided with the distribution.
20 //
21 // 3. Neither the name of the Corporation nor the names of the
22 // contributors may be used to endorse or promote products derived from
23 // this software without specific prior written permission.
24 //
25 // THIS SOFTWARE IS PROVIDED BY SANDIA CORPORATION "AS IS" AND ANY
26 // EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
27 // IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
28 // PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL SANDIA CORPORATION OR THE
29 // CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
30 // EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
31 // PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
32 // PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
33 // LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
34 // NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
35 // SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
36 //
37 // Questions? Contact Eric T. Phipps (etphipp@sandia.gov).
38 //
39 // ***********************************************************************
40 // @HEADER
41 
42 #ifndef KOKKOS_PARALLEL_MP_VECTOR_HPP
43 #define KOKKOS_PARALLEL_MP_VECTOR_HPP
44 
45 #include "Sacado_MP_Vector.hpp"
46 #include "Kokkos_Core.hpp"
47 
48 //----------------------------------------------------------------------------
49 // Kokkos execution policies useful for Sacado::MP::Vector scalar type
50 //----------------------------------------------------------------------------
51 
52 namespace Kokkos {
53 
57 template< class ExecSpace >
59 
61  typedef ExecSpace execution_space ;
62  typedef void work_tag ;
63 
64  size_t range;
65  size_t team;
66  size_t shared;
67 
68  MPVectorWorkConfig( const size_t range_,
69  const size_t team_,
70  const size_t shared_ = 0 ) :
71  range(range_), team(team_), shared(shared_) {}
72 };
73 
74 namespace Impl {
75 
76 #if defined( KOKKOS_ENABLE_THREADS )
77 // Specialization of ParallelFor<> for MPVectorWorkConfig and Threads
78 // The default implementation ignores the team size and uses the standard
79 // work-range implementation. In the future maybe we should try and use
80 // hyperthreads in a useful way. That would require:
81 // -- interpreting the team-size differently, rather as the sacado size
82 // -- determining the vector size of the architecture
83 // -- laying out the threads differently to use hyperthreads across the
84 // the sacado dimension
85 template< class FunctorType >
86 class ParallelFor< FunctorType , MPVectorWorkConfig< Threads > > :
87  public ParallelFor< FunctorType , Kokkos::RangePolicy< Threads > > {
88  typedef Kokkos::RangePolicy< Threads > Policy ;
89 public:
90  ParallelFor( const FunctorType & functor ,
91  const MPVectorWorkConfig< Threads > & work_config ) :
92  ParallelFor< FunctorType , Policy >( functor ,
93  Policy( 0, work_config.range ) ) {}
94 };
95 #endif
96 
97 #if defined( KOKKOS_ENABLE_OPENMP )
98 // Specialization of ParallelFor<> for MPVectorWorkConfig and OpenMP
99 // The default implementation ignores the team size and uses the standard
100 // work-range implementation. In the future maybe we should try and use
101 // hyperthreads in a useful way. That would require:
102 // -- interpreting the team-size differently, rather as the sacado size
103 // -- determining the vector size of the architecture
104 // -- laying out the threads differently to use hyperthreads across the
105 // the sacado dimension
106 template< class FunctorType >
107 class ParallelFor< FunctorType , MPVectorWorkConfig< OpenMP > > :
108  public ParallelFor< FunctorType , Kokkos::RangePolicy< OpenMP > > {
109  typedef Kokkos::RangePolicy< OpenMP > Policy ;
110 public:
111  ParallelFor( const FunctorType & functor ,
112  const MPVectorWorkConfig< OpenMP > & work_config ) :
113  ParallelFor< FunctorType , Policy >( functor ,
114  Policy( 0, work_config.range ) ) {}
115 };
116 #endif
117 
118 #if defined(KOKKOS_ENABLE_SERIAL)
119 // Specialization of ParallelFor<> for MPVectorWorkConfig and Serial
120 // The default implementation ignores the team size and uses the standard
121 // work-range implementation. In the future maybe we should try and use
122 // hyperthreads in a useful way. That would require:
123 // -- interpreting the team-size differently, rather as the sacado size
124 // -- determining the vector size of the architecture
125 // -- laying out the threads differently to use hyperthreads across the
126 // the sacado dimension
127 template< class FunctorType >
128 class ParallelFor< FunctorType , MPVectorWorkConfig< Serial > > :
129  public ParallelFor< FunctorType , Kokkos::RangePolicy< Serial > > {
130  typedef Kokkos::RangePolicy< Serial > Policy ;
131 public:
132  ParallelFor( const FunctorType & functor ,
133  const MPVectorWorkConfig< Serial > & work_config ) :
134  ParallelFor< FunctorType , Policy >( functor ,
135  Policy( 0, work_config.range ) ) {}
136 };
137 #endif // defined(KOKKOS_ENABLE_SERIAL)
138 
139 #if defined( KOKKOS_ENABLE_CUDA ) && defined( __CUDACC__ )
140 
141 // Specialization of ParallelFor<> for MPVectorWorkConfig on Cuda
142 // Here we use threadIdx.x for each entry in the specified team-size
143 template< class FunctorType >
144 class ParallelFor< FunctorType , MPVectorWorkConfig< Cuda > > {
145 public:
146 
147  const FunctorType m_functor ;
148  const MPVectorWorkConfig< Cuda > m_config;
149  const Cuda::size_type m_work ;
150 
151  inline
152  __device__
153  void operator()(void) const
154  {
155  const Cuda::size_type work_stride = blockDim.y * gridDim.x ;
156 
157  for ( Cuda::size_type iwork = threadIdx.y + blockDim.y * blockIdx.x ;
158  iwork < m_work ;
159  iwork += work_stride ) {
160  m_functor( iwork , threadIdx.x );
161  }
162  }
163 
164  ParallelFor( const FunctorType & functor ,
165  const MPVectorWorkConfig< Cuda > & work_config )
166  : m_functor( functor ) ,
167  m_config( work_config ) ,
168  m_work( work_config.range )
169  {
170  }
171 
172  inline
173  void execute() const
174  {
175  // To do: query number of registers used by functor and adjust
176  // nwarp accordingly to get maximum occupancy
177 
178  Cuda::size_type nwarp = 0;
179  if (m_config.team > CudaTraits::WarpSize) {
180  const Cuda::size_type warps_per_team =
181  ( m_config.team + CudaTraits::WarpSize-1 ) / CudaTraits::WarpSize;
182  nwarp = cuda_internal_maximum_warp_count() / warps_per_team;
183  }
184  else {
185  const Cuda::size_type teams_per_warp =
186  CudaTraits::WarpSize / m_config.team ;
187  nwarp = cuda_internal_maximum_warp_count() * teams_per_warp;
188  }
189  const dim3 block( m_config.team , nwarp , 1 );
190 
191  Cuda::size_type nblock =
192  std::min( (m_work + block.y - 1 ) / block.y ,
193  cuda_internal_maximum_grid_count() );
194  const dim3 grid( nblock , 1 , 1 );
195 
196  const Cuda::size_type shared = m_config.shared;
197  CudaParallelLaunch< ParallelFor >( *this , grid , block , shared );
198  }
199 };
200 
201 #endif
202 
203 } // namespace Impl
204 
205 } // namespace Kokkos
206 
207 //----------------------------------------------------------------------------
208 //----------------------------------------------------------------------------
209 
210 #endif /* #ifndef KOKKOS_ATOMIC_MP_VECTOR_HPP */
MPVectorWorkConfig(const size_t range_, const size_t team_, const size_t shared_=0)
KOKKOS_INLINE_FUNCTION PCE< Storage > min(const typename PCE< Storage >::value_type &a, const PCE< Storage > &b)
Team-based parallel work configuration for Sacado::MP::Vector.