Sacado Package Browser (Single Doxygen Collection)  Version of the Day
KokkosExp_View_Fad_Contiguous.hpp
Go to the documentation of this file.
1 // @HEADER
2 // ***********************************************************************
3 //
4 // Sacado Package
5 // Copyright (2006) Sandia Corporation
6 //
7 // Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
8 // the U.S. Government retains certain rights in this software.
9 //
10 // This library is free software; you can redistribute it and/or modify
11 // it under the terms of the GNU Lesser General Public License as
12 // published by the Free Software Foundation; either version 2.1 of the
13 // License, or (at your option) any later version.
14 //
15 // This library is distributed in the hope that it will be useful, but
16 // WITHOUT ANY WARRANTY; without even the implied warranty of
17 // MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
18 // Lesser General Public License for more details.
19 //
20 // You should have received a copy of the GNU Lesser General Public
21 // License along with this library; if not, write to the Free Software
22 // Foundation, Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301
23 // USA
24 // Questions? Contact David M. Gay (dmgay@sandia.gov) or Eric T. Phipps
25 // (etphipp@sandia.gov).
26 //
27 // ***********************************************************************
28 // @HEADER
29 
30 #ifndef KOKKOS_EXPERIMENTAL_VIEW_SACADO_FAD_CONTIGUOUS_HPP
31 #define KOKKOS_EXPERIMENTAL_VIEW_SACADO_FAD_CONTIGUOUS_HPP
32 
33 #include "Sacado_ConfigDefs.h"
34 #if defined(HAVE_SACADO_KOKKOSCORE)
35 
37 
38 // Some default traits definitions that need to be defined even if the view
39 // specialization is disabled
40 namespace Kokkos {
41 
42 template <typename ViewType, typename Enabled = void>
43 struct ThreadLocalScalarType {
44  typedef typename ViewType::non_const_value_type type;
45 };
46 
47 template <typename ViewType>
48 struct ViewScalarStride {
49  static const unsigned stride =
51  static const bool is_unit_stride =
53 };
54 
55 } // namespace Kokkos
56 
57 namespace Sacado {
58 
59  namespace Fad {
60 
61  /* Define a partition of a View of Sacado Fad type */
62  template <unsigned Size = 0>
63  struct Partition {
64  static const unsigned PartitionSize = Size;
65  unsigned offset ;
66  unsigned stride ;
67 
68  template< typename iType0 , typename iType1 >
70  Partition( const iType0 & i0 , const iType1 & i1 ) :
71  offset(i0), stride(i1) {
72  }
73  };
74 
75  template <typename T>
76  struct is_fad_partition {
77  static const bool value = false;
78  };
79 
80  template <unsigned Stride>
81  struct is_fad_partition< Partition<Stride> > {
82  static const bool value = true;
83  };
84 
85  }
86 
87  // Type of local scalar type when partitioning a view
88  template <typename T, unsigned Stride = 0>
89  struct LocalScalarType {
90  typedef T type;
91  };
92  template <typename T, unsigned Stride>
93  struct LocalScalarType<const T, Stride> {
94  typedef typename LocalScalarType<T,Stride>::type lst;
95  typedef const lst type;
96  };
97 
98  // For DFad, the size is not part of the type, so the default implementation
99  // is sufficient
100 
101  // Type of local scalar type when partitioning a view
102  //
103  // For SLFad, divde the array size by the given stride
104  namespace Fad {
105  namespace Exp {
106  template <typename T, int N> class StaticStorage;
107  template <typename S> class GeneralFad;
108  }
109  }
110  template <typename T, int N, unsigned Stride>
111  struct LocalScalarType< Fad::Exp::GeneralFad< Fad::Exp::StaticStorage<T,N> >,
112  Stride > {
113  static const int Ns = (N+Stride-1) / Stride;
114  typedef Fad::Exp::GeneralFad< Fad::Exp::StaticStorage<T,Ns> > type;
115  };
116 #ifndef SACADO_NEW_FAD_DESIGN_IS_DEFAULT
117  namespace Fad {
118  template <typename T, int N> class SLFad;
119  }
120  template <typename T, int N, unsigned Stride>
121  struct LocalScalarType< Fad::SLFad<T,N>, Stride > {
122  static const int Ns = (N+Stride-1) / Stride;
123  typedef Fad::SLFad<T,Ns> type;
124  };
125 #endif
126 
127  // Type of local scalar type when partitioning a view
128  //
129  // For SFad, divde the array size by the given stride. If it divides evenly,
130  // use SFad, otherwise use SLFad
131  namespace Fad {
132  namespace Exp {
133  template <typename T, int N> class StaticFixedStorage;
134  template <typename T, int N> class StaticStorage;
135  template <typename S> class GeneralFad;
136  }
137  }
138  template <typename T, int N, unsigned Stride>
139  struct LocalScalarType< Fad::Exp::GeneralFad< Fad::Exp::StaticFixedStorage<T,N> >,
140  Stride > {
141  static const int Ns = (N+Stride-1) / Stride;
142  typedef typename std::conditional<
143  Ns == N/Stride ,
144  Fad::Exp::GeneralFad< Fad::Exp::StaticFixedStorage<T,Ns> > ,
145  Fad::Exp::GeneralFad< Fad::Exp::StaticStorage<T,Ns> >
146  >::type type;
147  };
148 
149 #ifndef SACADO_NEW_FAD_DESIGN_IS_DEFAULT
150  namespace Fad {
151  template <typename T, int N> class SFad;
152  }
153  template <typename T, int N, unsigned Stride>
154  struct LocalScalarType< Fad::SFad<T,N>, Stride > {
155  static const int Ns = (N+Stride-1) / Stride;
156  typedef typename std::conditional< Ns == N/Stride , Fad::SFad<T,Ns> , Fad::SLFad<T,Ns> >::type type;
157  };
158 #endif
159 
160  template <unsigned Stride, typename T>
162  const T& partition_scalar(const T& x) { return x; }
163 
164 } // namespace Sacado
165 
166 // Make sure the user really wants these View specializations
167 #if defined(HAVE_SACADO_VIEW_SPEC) && !defined(SACADO_DISABLE_FAD_VIEW_SPEC)
168 
169 #include "Sacado_Traits.hpp"
170 #include "Kokkos_Core.hpp"
171 #include "impl/Kokkos_ViewMapping.hpp"
172 
173 //----------------------------------------------------------------------------
174 
175 namespace Sacado {
176 
177 #if defined(__CUDA_ARCH__)
178  namespace Fad {
179  namespace Exp {
180  template <typename T, typename U> class DynamicStorage;
181  template <typename T, int N> class StaticFixedStorage;
182  template <typename T, int N> class StaticStorage;
183  template <typename S> class GeneralFad;
184  }
185  }
186 #ifndef SACADO_VIEW_CUDA_HIERARCHICAL_DFAD
187  template <unsigned Stride, typename T, typename U>
189  typename LocalScalarType< Fad::Exp::GeneralFad< Fad::Exp::DynamicStorage<T,U> >, Stride >::type
190  partition_scalar(const Fad::Exp::GeneralFad< Fad::Exp::DynamicStorage<T,U> >& x) {
191  typedef typename LocalScalarType< Fad::Exp::GeneralFad< Fad::Exp::DynamicStorage<T,U> >, Stride >::type ret_type;
192  const int size = (x.size()+blockDim.x-threadIdx.x-1) / blockDim.x;
193  const int offset = threadIdx.x;
194  ret_type xp(size, x.val());
195 
196  // Note: we can't use x.dx(offset+i*Stride) if
197  // SACADO_VIEW_CUDA_HIERARCHICAL_DFAD_STRIDED is defined because it already
198  // uses blockDim.x in its index calculation. This approach should work
199  // regardless
200  const T* dx = x.dx();
201  for (int i=0; i<size; ++i)
202  xp.fastAccessDx(i) = dx[offset+i*Stride];
203 
204  return xp;
205  }
206 #endif
207  template <unsigned Stride, typename T, int N>
209  typename LocalScalarType< Fad::Exp::GeneralFad< Fad::Exp::StaticStorage<T,N> >, Stride >::type
210  partition_scalar(const Fad::Exp::GeneralFad< Fad::Exp::StaticStorage<T,N> >& x) {
211  typedef typename LocalScalarType< Fad::Exp::GeneralFad< Fad::Exp::StaticStorage<T,N> >, Stride >::type ret_type;
212  const int size = (x.size()+blockDim.x-threadIdx.x-1) / blockDim.x;
213  const int offset = threadIdx.x;
214  ret_type xp(size, x.val());
215  for (int i=0; i<size; ++i)
216  xp.fastAccessDx(i) = x.fastAccessDx(offset+i*Stride);
217  return xp;
218  }
219  template <unsigned Stride, typename T, int N>
221  typename LocalScalarType< Fad::Exp::GeneralFad< Fad::Exp::StaticFixedStorage<T,N> >, Stride >::type
222  partition_scalar(const Fad::Exp::GeneralFad< Fad::Exp::StaticFixedStorage<T,N> >& x) {
223  typedef typename LocalScalarType< Fad::Exp::GeneralFad< Fad::Exp::StaticFixedStorage<T,N> >, Stride >::type ret_type;
224  const int size = (x.size()+blockDim.x-threadIdx.x-1) / blockDim.x;
225  const int offset = threadIdx.x;
226  ret_type xp(size, x.val());
227  for (int i=0; i<size; ++i)
228  xp.fastAccessDx(i) = x.fastAccessDx(offset+i*Stride);
229  return xp;
230  }
231 
232 #ifndef SACADO_NEW_FAD_DESIGN_IS_DEFAULT
233  namespace Fad {
234  template <typename T> class DFad;
235  template <typename T, int N> class SLFad;
236  template <typename T, int N> class SFad;
237  }
238 #ifndef SACADO_VIEW_CUDA_HIERARCHICAL_DFAD
239  template <unsigned Stride, typename T>
241  typename LocalScalarType< Fad::DFad<T>, Stride >::type
242  partition_scalar(const Fad::DFad<T>& x) {
243  typedef typename LocalScalarType< Fad::DFad<T>, Stride >::type ret_type;
244  const int size = (x.size()+blockDim.x-threadIdx.x-1) / blockDim.x;
245  const int offset = threadIdx.x;
246  ret_type xp(size, x.val());
247 
248  // Note: we can't use x.dx(offset+i*Stride) if
249  // SACADO_VIEW_CUDA_HIERARCHICAL_DFAD_STRIDED is defined because it already
250  // uses blockDim.x in its index calculation. This approach should work
251  // regardless
252  const T* dx = x.dx();
253  for (int i=0; i<size; ++i)
254  xp.fastAccessDx(i) = dx[offset+i*Stride];
255 
256  return xp;
257  }
258 #endif
259  template <unsigned Stride, typename T, int N>
261  typename LocalScalarType< Fad::SLFad<T,N>, Stride >::type
262  partition_scalar(const Fad::SLFad<T,N>& x) {
263  typedef typename LocalScalarType< Fad::SLFad<T,N>, Stride >::type ret_type;
264  const int size = (x.size()+blockDim.x-threadIdx.x-1) / blockDim.x;
265  const int offset = threadIdx.x;
266  ret_type xp(size, x.val());
267  for (int i=0; i<size; ++i)
268  xp.fastAccessDx(i) = x.fastAccessDx(offset+i*Stride);
269  return xp;
270  }
271  template <unsigned Stride, typename T, int N>
273  typename LocalScalarType< Fad::SFad<T,N>, Stride >::type
274  partition_scalar(const Fad::SFad<T,N>& x) {
275  typedef typename LocalScalarType< Fad::SFad<T,N>, Stride >::type ret_type;
276  const int size = (x.size()+blockDim.x-threadIdx.x-1) / blockDim.x;
277  const int offset = threadIdx.x;
278  ret_type xp(size, x.val());
279  for (int i=0; i<size; ++i)
280  xp.fastAccessDx(i) = x.fastAccessDx(offset+i*Stride);
281  return xp;
282  }
283 #endif // SACADO_NEW_FAD_DESIGN_IS_DEFAULT
284 
285 #endif // defined(__CUDA_ARCH__)
286 
287 } // namespace Sacado
288 
289 //----------------------------------------------------------------------------
290 
291 namespace Kokkos {
292 
293 template< unsigned Stride, typename D, typename ... P >
295 typename Kokkos::Impl::ViewMapping< void, typename Kokkos::ViewTraits<D,P...>, Sacado::Fad::Partition<Stride> >::type
296 partition( const Kokkos::View<D,P...> & src ,
297  const unsigned offset ,
298  const unsigned stride )
299 {
300  typedef Kokkos::ViewTraits<D,P...> traits;
301  typedef typename Kokkos::Impl::ViewMapping< void, traits, Sacado::Fad::Partition<Stride> >::type DstViewType;
302  const Sacado::Fad::Partition<Stride> part( offset , stride );
303  return DstViewType(src, part);
304 }
305 
306 template <typename ViewType>
307 struct ThreadLocalScalarType<
308  ViewType,
309  typename std::enable_if< is_view_fad_contiguous<ViewType>::value >::type > {
310  typedef typename ViewType::traits TraitsType;
311  typedef Impl::ViewMapping<TraitsType, void> MappingType;
312  typedef typename MappingType::thread_local_scalar_type type;
313 };
314 
315 namespace Impl {
316 
317 #if defined (KOKKOS_ENABLE_CUDA) && defined(SACADO_VIEW_CUDA_HIERARCHICAL)
318 template< class OutputView >
319 struct SacadoViewFill<
320  OutputView,
321  typename std::enable_if<
322  ( Kokkos::is_view_fad_contiguous<OutputView>::value &&
323  std::is_same<typename OutputView::execution_space, Kokkos::Cuda>::value &&
324  !Kokkos::ViewScalarStride<OutputView>::is_unit_stride )
325  >::type
326  >
327 {
328  typedef typename OutputView::const_value_type const_value_type ;
329  typedef typename OutputView::execution_space execution_space ;
330  typedef Kokkos::TeamPolicy< execution_space> team_policy;
331  typedef typename team_policy::member_type team_handle;
332  typedef typename Kokkos::ThreadLocalScalarType<OutputView>::type local_scalar_type;
333  static const unsigned stride = Kokkos::ViewScalarStride<OutputView>::stride;
334 
335  const OutputView output ;
336  const_value_type input ;
337 
339  void operator()( const size_t i0 ) const
340  {
341  local_scalar_type input_stride = Sacado::partition_scalar<stride>(input);
342 
343  const size_t n1 = output.extent(1);
344  const size_t n2 = output.extent(2);
345  const size_t n3 = output.extent(3);
346  const size_t n4 = output.extent(4);
347  const size_t n5 = output.extent(5);
348  const size_t n6 = output.extent(6);
349  const size_t n7 = output.extent(7);
350 
351  for ( size_t i1 = 0 ; i1 < n1 ; ++i1 ) {
352  for ( size_t i2 = 0 ; i2 < n2 ; ++i2 ) {
353  for ( size_t i3 = 0 ; i3 < n3 ; ++i3 ) {
354  for ( size_t i4 = 0 ; i4 < n4 ; ++i4 ) {
355  for ( size_t i5 = 0 ; i5 < n5 ; ++i5 ) {
356  for ( size_t i6 = 0 ; i6 < n6 ; ++i6 ) {
357  for ( size_t i7 = 0 ; i7 < n7 ; ++i7 ) {
358  output.access(i0,i1,i2,i3,i4,i5,i6,i7) = input_stride ;
359  }}}}}}}
360  }
361 
363  void operator()( const team_handle& team ) const
364  {
365  const size_t i0 = team.league_rank()*team.team_size() + team.team_rank();
366  if (i0 < output.extent(0))
367  (*this)(i0);
368  }
369 
370  SacadoViewFill( const OutputView & arg_out , const_value_type & arg_in )
371  : output( arg_out ), input( arg_in )
372  {
373  const size_t team_size = 256 / stride;
374  team_policy policy( (output.extent(0)+team_size-1)/team_size ,
375  team_size , stride );
376  Kokkos::parallel_for( policy, *this );
377  execution_space::fence();
378  }
379 };
380 #endif
381 
382 } // namespace Impl
383 } // namespace Kokkos
384 
385 //----------------------------------------------------------------------------
386 
387 namespace Kokkos {
388 namespace Impl {
389 
390 template< class ... Args >
391 struct is_ViewSpecializeSacadoFadContiguous { enum { value = false }; };
392 
393 template< class D , class ... P , class ... Args >
394 struct is_ViewSpecializeSacadoFadContiguous< Kokkos::View<D,P...> , Args... > {
395  enum { value =
396  std::is_same< typename Kokkos::ViewTraits<D,P...>::specialize
397  , ViewSpecializeSacadoFadContiguous >::value
398  &&
399  ( ( sizeof...(Args) == 0 ) ||
400  is_ViewSpecializeSacadoFadContiguous< Args... >::value ) };
401 };
402 
403 } // namespace Impl
404 } // namespace Kokkos
405 
406 //----------------------------------------------------------------------------
407 
408 namespace Kokkos {
409 namespace Impl {
410 
411 // Compute a partitioned fad size given a stride. Return 0 if the stride
412 // does not evenly divide the size
414 constexpr unsigned computeFadPartitionSize(unsigned size, unsigned stride)
415 {
416  return
417  ((size+stride-1)/stride) == (size/stride) ? ((size+stride-1)/stride) : 0;
418 }
419 
420 // Create new Layout with Fad dimension set to the last
421 // Note: we use enable_if<> here to handle both LayoutLeft and
422 // LayoutContiguous<LayoutLeft>
423 template <unsigned rank, unsigned static_dim, typename Layout>
425 typename std::enable_if< !std::is_same<Layout, LayoutLeft>::value &&
426  !std::is_same<Layout, LayoutStride>::value,
427  Layout>::type
428 create_fad_array_layout(const Layout& layout)
429 {
430  size_t dims[8];
431  for (int i=0; i<8; ++i)
432  dims[i] = layout.dimension[i];
433  if (static_dim > 0)
434  dims[rank] = static_dim+1;
435  return Layout( dims[0], dims[1], dims[2], dims[3],
436  dims[4], dims[5], dims[6], dims[7] );
437 }
438 
439 // Create new Layout with Fad dimension set to the last
440 // Note: we use enable_if<> here to handle both LayoutStride and
441 // LayoutContiguous<LayoutStride>
442 template <unsigned rank, unsigned static_dim, typename Layout>
444 typename std::enable_if< std::is_same<Layout, LayoutStride>::value, Layout>::type
445 create_fad_array_layout(const Layout& layout)
446 {
447  size_t dims[8], strides[8];
448  for (int i=0; i<8; ++i) {
449  dims[i] = layout.dimension[i];
450  strides[i] = layout.stride[i];
451  }
452  if (static_dim > 0) {
453  dims[rank] = static_dim+1;
454  strides[rank] = 1;
455  }
456  return Layout( dims[0], strides[0],
457  dims[1], strides[1],
458  dims[2], strides[2],
459  dims[3], strides[3],
460  dims[4], strides[4],
461  dims[5], strides[5],
462  dims[6], strides[6],
463  dims[7], strides[7] );
464 }
465 
466 // Create new LayoutLeft with Fad dimension shuffled to the first
467 // Note: we use enable_if<> here to handle both LayoutLeft and
468 // LayoutContiguous<LayoutLeft>
469  template <unsigned rank, unsigned static_dim, typename Layout>
471 typename std::enable_if< std::is_same<Layout, LayoutLeft>::value, Layout >::type
472 create_fad_array_layout(const Layout& layout)
473 {
474  size_t dims[8];
475  for (int i=0; i<8; ++i)
476  dims[i] = layout.dimension[i];
477  size_t fad_dim = static_dim == 0 ? dims[rank] : static_dim+1;
478  for (int i=rank; i>=1; --i)
479  dims[i] = dims[i-1];
480  dims[0] = fad_dim;
481  return Layout( dims[0], dims[1], dims[2], dims[3],
482  dims[4], dims[5], dims[6], dims[7] );
483 }
484 
485 template <unsigned Rank, typename Dimension, typename Layout>
487 typename std::enable_if< !std::is_same<Layout, LayoutLeft>::value, size_t>::type
488 getFadDimension(const ViewOffset<Dimension,Layout,void>& offset)
489 {
490  return
491  ( Rank == 0 ? offset.dimension_0() :
492  ( Rank == 1 ? offset.dimension_1() :
493  ( Rank == 2 ? offset.dimension_2() :
494  ( Rank == 3 ? offset.dimension_3() :
495  ( Rank == 4 ? offset.dimension_4() :
496  ( Rank == 5 ? offset.dimension_5() :
497  ( Rank == 6 ? offset.dimension_6() :
498  offset.dimension_7() )))))));
499 }
500 
501 template <unsigned Rank, typename Dimension, typename Layout>
503 typename std::enable_if< std::is_same<Layout, LayoutLeft>::value, size_t >::type
504 getFadDimension(const ViewOffset<Dimension,Layout,void>& offset)
505 {
506  return offset.dimension_0();
507 }
508 
509 template< class Traits >
510 class ViewMapping< Traits , /* View internal mapping */
511  typename std::enable_if<
512  ( std::is_same< typename Traits::specialize
513  , ViewSpecializeSacadoFadContiguous >::value
514  &&
515  ( std::is_same< typename Traits::array_layout
516  , Kokkos::LayoutLeft >::value
517  ||
518  std::is_same< typename Traits::array_layout
519  , Kokkos::LayoutRight >::value
520  ||
521  std::is_same< typename Traits::array_layout
522  , Kokkos::LayoutStride >::value
523  )
524  )>::type >
525 {
526 private:
527 
528  template< class , class ... > friend class ViewMapping ;
529  template< class , class ... > friend class Kokkos::View ;
530 
531  typedef typename Traits::value_type fad_type ;
532  typedef typename Sacado::ValueType< fad_type >::type fad_value_type ;
533  typedef typename
534  std::add_const< fad_value_type >::type const_fad_value_type ;
535 
536 public:
537 
538  enum { FadStaticDimension = Sacado::StaticSize< fad_type >::value };
539  enum { PartitionedFadStride = Traits::array_layout::scalar_stride };
540 
541  // The partitioned static size -- this will be 0 if ParitionedFadStride
542  // does not evenly divide FadStaticDimension
543  enum { PartitionedFadStaticDimension =
544  computeFadPartitionSize(FadStaticDimension,PartitionedFadStride) };
545 
546 #ifdef KOKKOS_ENABLE_CUDA
547  typedef typename Sacado::LocalScalarType< fad_type, unsigned(PartitionedFadStride) >::type strided_scalar_type;
548  typedef typename std::conditional< std::is_same<typename Traits::execution_space, Kokkos::Cuda>::value, strided_scalar_type, fad_type >::type thread_local_scalar_type;
549 #else
550  typedef fad_type thread_local_scalar_type;
551 #endif
552 
553 private:
555 
556  typedef fad_value_type * handle_type ;
557 
558  typedef ViewArrayAnalysis< typename Traits::data_type > array_analysis ;
559 
560  typedef ViewOffset< typename Traits::dimension
561  , typename Traits::array_layout
562  , void
563  > offset_type ;
564 
565  // Prepend/append the fad dimension for the internal offset mapping.
566  static constexpr bool is_layout_left =
567  std::is_same< typename Traits::array_layout, Kokkos::LayoutLeft>::value;
568  typedef ViewOffset<
569  typename std::conditional<
570  is_layout_left,
571  typename array_analysis::dimension::
572  template prepend<0>::type,
573  typename array_analysis::dimension::
574  template append<0>::type >::type,
575  typename Traits::array_layout,
576  void >
577  array_offset_type ;
578 
579  handle_type m_handle ;
580  offset_type m_offset ;
581  array_offset_type m_array_offset ;
582  sacado_size_type m_fad_size ;
583 
584  // These are for manual partitioning, and will likely be removed
585  unsigned m_original_fad_size ;
586  unsigned m_fad_stride ;
587  unsigned m_fad_index ;
588 
589 public:
590 
591  //----------------------------------------
592  // Domain dimensions
593 
594  enum { Rank = Traits::dimension::rank };
595 
596  // Rank corresponding to the sacado dimension
597  enum { Sacado_Rank = std::is_same< typename Traits::array_layout, Kokkos::LayoutLeft >::value ? 0 : Rank+1 };
598 
599  // Using the internal offset mapping so limit to public rank:
600  template< typename iType >
601  KOKKOS_INLINE_FUNCTION constexpr size_t extent( const iType & r ) const
602  { return m_offset.m_dim.extent(r); }
603 
604  KOKKOS_INLINE_FUNCTION constexpr
605  typename Traits::array_layout layout() const
606  { return m_offset.layout(); }
607 
608  KOKKOS_INLINE_FUNCTION constexpr size_t dimension_0() const
609  { return m_offset.dimension_0(); }
610  KOKKOS_INLINE_FUNCTION constexpr size_t dimension_1() const
611  { return m_offset.dimension_1(); }
612  KOKKOS_INLINE_FUNCTION constexpr size_t dimension_2() const
613  { return m_offset.dimension_2(); }
614  KOKKOS_INLINE_FUNCTION constexpr size_t dimension_3() const
615  { return m_offset.dimension_3(); }
616  KOKKOS_INLINE_FUNCTION constexpr size_t dimension_4() const
617  { return m_offset.dimension_4(); }
618  KOKKOS_INLINE_FUNCTION constexpr size_t dimension_5() const
619  { return m_offset.dimension_5(); }
620  KOKKOS_INLINE_FUNCTION constexpr size_t dimension_6() const
621  { return m_offset.dimension_6(); }
622  KOKKOS_INLINE_FUNCTION constexpr size_t dimension_7() const
623  { return m_offset.dimension_7(); }
624 
625  // Is a regular layout with uniform striding for each index.
626  // Since we allow for striding within the data type, we can't guarantee
627  // regular striding
628  using is_regular = std::false_type ;
629 
630  // FIXME: Adjust these for m_stride
631  KOKKOS_INLINE_FUNCTION constexpr size_t stride_0() const
632  { return m_offset.stride_0(); }
633  KOKKOS_INLINE_FUNCTION constexpr size_t stride_1() const
634  { return m_offset.stride_1(); }
635  KOKKOS_INLINE_FUNCTION constexpr size_t stride_2() const
636  { return m_offset.stride_2(); }
637  KOKKOS_INLINE_FUNCTION constexpr size_t stride_3() const
638  { return m_offset.stride_3(); }
639  KOKKOS_INLINE_FUNCTION constexpr size_t stride_4() const
640  { return m_offset.stride_4(); }
641  KOKKOS_INLINE_FUNCTION constexpr size_t stride_5() const
642  { return m_offset.stride_5(); }
643  KOKKOS_INLINE_FUNCTION constexpr size_t stride_6() const
644  { return m_offset.stride_6(); }
645  KOKKOS_INLINE_FUNCTION constexpr size_t stride_7() const
646  { return m_offset.stride_7(); }
647 
648  template< typename iType >
649  KOKKOS_INLINE_FUNCTION void stride( iType * const s ) const
650  { m_offset.stride(s); }
651 
652  // Size of sacado scalar dimension
653  KOKKOS_FORCEINLINE_FUNCTION constexpr unsigned dimension_scalar() const
654 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && defined(__CUDA_ARCH__)
655  { return PartitionedFadStaticDimension ? PartitionedFadStaticDimension+1 : (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x + 1; }
656 #else
657  { return m_fad_size.value+1; }
658 #endif
659 
660  // trode of sacado scalar dimension
661  KOKKOS_FORCEINLINE_FUNCTION constexpr unsigned stride_scalar() const
662  { return m_fad_stride; }
663 
664  //----------------------------------------
665  // Range of mapping
666 
667 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && defined(__CUDA_ARCH__)
668  // Return type of reference operators
669  // this only works if you are using a team-parallel operation on Cuda!
670  // typedef typename
671  // Sacado::ViewFadType< thread_local_scalar_type , PartitionedFadStaticDimension , (unsigned(ParitionedFadStride) > 1 ? PartitionedFadStride : 0) >::type reference_type ;
672  typedef typename
674 #else
675  // Return type of reference operators
676  typedef typename
678 #endif
679 
681  typedef fad_value_type * pointer_type ;
682 
684  KOKKOS_INLINE_FUNCTION constexpr size_t span() const
685  { return m_array_offset.span(); }
686 
688  KOKKOS_INLINE_FUNCTION constexpr bool span_is_contiguous() const
689  { return m_array_offset.span_is_contiguous() && (m_fad_stride == 1); }
690 
692  KOKKOS_INLINE_FUNCTION constexpr pointer_type data() const
693 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && defined(__CUDA_ARCH__)
694  { return m_handle + threadIdx.x; }
695 #else
696  { return m_handle + m_fad_index; }
697 #endif
698 
699  //----------------------------------------
700 
702  reference_type
703  reference() const
704  {
705 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && defined(__CUDA_ARCH__)
706  const unsigned index = threadIdx.x;
707  const unsigned strd = blockDim.x;
708  const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
709 #else
710  const unsigned index = m_fad_index;
711  const unsigned strd = m_fad_stride;
712  const unsigned size = m_fad_size.value;
713 #endif
714  return reference_type( m_handle + index
715  , m_handle + m_original_fad_size
716  , size
717  , strd ); }
718 
719  template< typename I0 >
721  typename std::enable_if< Kokkos::Impl::are_integral<I0>::value &&
722  is_layout_left, reference_type>::type
723  reference( const I0 & i0 ) const
724  { pointer_type beg = m_handle + m_array_offset(0,i0);
725 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && defined(__CUDA_ARCH__)
726  const unsigned index = threadIdx.x;
727  const unsigned strd = blockDim.x;
728  const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
729 #else
730  const unsigned index = m_fad_index;
731  const unsigned strd = m_fad_stride;
732  const unsigned size = m_fad_size.value;
733 #endif
734  return reference_type( beg + index
735  , beg + m_original_fad_size
736  , size
737  , strd ); }
738 
739  template< typename I0 >
741  typename std::enable_if< Kokkos::Impl::are_integral<I0>::value &&
742  !is_layout_left, reference_type>::type
743  reference( const I0 & i0 ) const
744  { pointer_type beg = m_handle + m_array_offset(i0,0);
745 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && defined(__CUDA_ARCH__)
746  const unsigned index = threadIdx.x;
747  const unsigned strd = blockDim.x;
748  const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
749 #else
750  const unsigned index = m_fad_index;
751  const unsigned strd = m_fad_stride;
752  const unsigned size = m_fad_size.value;
753 #endif
754  return reference_type( beg + index
755  , beg + m_original_fad_size
756  , size
757  , strd ); }
758 
759  template< typename I0 , typename I1 >
761  typename std::enable_if< Kokkos::Impl::are_integral<I0,I1>::value &&
762  is_layout_left, reference_type>::type
763  reference( const I0 & i0 , const I1 & i1 ) const
764  { pointer_type beg = m_handle + m_array_offset(0,i0,i1);
765 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && defined(__CUDA_ARCH__)
766  const unsigned index = threadIdx.x;
767  const unsigned strd = blockDim.x;
768  const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
769 #else
770  const unsigned index = m_fad_index;
771  const unsigned strd = m_fad_stride;
772  const unsigned size = m_fad_size.value;
773 #endif
774  return reference_type( beg + index
775  , beg + m_original_fad_size
776  , size
777  , strd ); }
778 
779  template< typename I0 , typename I1 >
781  typename std::enable_if< Kokkos::Impl::are_integral<I0,I1>::value &&
782  !is_layout_left, reference_type>::type
783  reference( const I0 & i0 , const I1 & i1 ) const
784  { pointer_type beg = m_handle + m_array_offset(i0,i1,0);
785 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && defined(__CUDA_ARCH__)
786  const unsigned index = threadIdx.x;
787  const unsigned strd = blockDim.x;
788  const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
789 #else
790  const unsigned index = m_fad_index;
791  const unsigned strd = m_fad_stride;
792  const unsigned size = m_fad_size.value;
793 #endif
794  return reference_type( beg + index
795  , beg + m_original_fad_size
796  , size
797  , strd ); }
798 
799 
800  template< typename I0 , typename I1 , typename I2 >
802  typename std::enable_if< Kokkos::Impl::are_integral<I0,I1,I2>::value &&
803  is_layout_left, reference_type>::type
804  reference( const I0 & i0 , const I1 & i1 , const I2 & i2 ) const
805  { pointer_type beg = m_handle + m_array_offset(0,i0,i1,i2);
806 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && defined(__CUDA_ARCH__)
807  const unsigned index = threadIdx.x;
808  const unsigned strd = blockDim.x;
809  const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
810 #else
811  const unsigned index = m_fad_index;
812  const unsigned strd = m_fad_stride;
813  const unsigned size = m_fad_size.value;
814 #endif
815  return reference_type( beg + index
816  , beg + m_original_fad_size
817  , size
818  , strd ); }
819 
820  template< typename I0 , typename I1 , typename I2 >
822  typename std::enable_if< Kokkos::Impl::are_integral<I0,I1,I2>::value &&
823  !is_layout_left, reference_type>::type
824  reference( const I0 & i0 , const I1 & i1 , const I2 & i2 ) const
825  { pointer_type beg = m_handle + m_array_offset(i0,i1,i2,0);
826 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && defined(__CUDA_ARCH__)
827  const unsigned index = threadIdx.x;
828  const unsigned strd = blockDim.x;
829  const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
830 #else
831  const unsigned index = m_fad_index;
832  const unsigned strd = m_fad_stride;
833  const unsigned size = m_fad_size.value;
834 #endif
835  return reference_type( beg + index
836  , beg + m_original_fad_size
837  , size
838  , strd ); }
839 
840  template< typename I0 , typename I1 , typename I2 , typename I3 >
842  typename std::enable_if< Kokkos::Impl::are_integral<I0,I1,I2,I3>::value &&
843  is_layout_left, reference_type>::type
844  reference( const I0 & i0 , const I1 & i1 , const I2 & i2 , const I3 & i3 ) const
845  { pointer_type beg = m_handle + m_array_offset(0,i0,i1,i2,i3);
846 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && defined(__CUDA_ARCH__)
847  const unsigned index = threadIdx.x;
848  const unsigned strd = blockDim.x;
849  const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
850 #else
851  const unsigned index = m_fad_index;
852  const unsigned strd = m_fad_stride;
853  const unsigned size = m_fad_size.value;
854 #endif
855  return reference_type( beg + index
856  , beg + m_original_fad_size
857  , size
858  , strd ); }
859 
860  template< typename I0 , typename I1 , typename I2 , typename I3 >
862  typename std::enable_if< Kokkos::Impl::are_integral<I0,I1,I2,I3>::value &&
863  !is_layout_left, reference_type>::type
864  reference( const I0 & i0 , const I1 & i1 , const I2 & i2 , const I3 & i3 ) const
865  { pointer_type beg = m_handle + m_array_offset(i0,i1,i2,i3,0);
866 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && defined(__CUDA_ARCH__)
867  const unsigned index = threadIdx.x;
868  const unsigned strd = blockDim.x;
869  const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
870 #else
871  const unsigned index = m_fad_index;
872  const unsigned strd = m_fad_stride;
873  const unsigned size = m_fad_size.value;
874 #endif
875  return reference_type( beg + index
876  , beg + m_original_fad_size
877  , size
878  , strd ); }
879 
880  template< typename I0 , typename I1 , typename I2 , typename I3
881  , typename I4 >
883  typename std::enable_if< Kokkos::Impl::are_integral<I0,I1,I2,I3,I4>::value &&
884  is_layout_left, reference_type>::type
885  reference( const I0 & i0 , const I1 & i1 , const I2 & i2 , const I3 & i3
886  , const I4 & i4 ) const
887  { pointer_type beg = m_handle + m_array_offset(0,i0,i1,i2,i3,i4);
888 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && defined(__CUDA_ARCH__)
889  const unsigned index = threadIdx.x;
890  const unsigned strd = blockDim.x;
891  const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
892 #else
893  const unsigned index = m_fad_index;
894  const unsigned strd = m_fad_stride;
895  const unsigned size = m_fad_size.value;
896 #endif
897  return reference_type( beg + index
898  , beg + m_original_fad_size
899  , size
900  , strd ); }
901 
902  template< typename I0 , typename I1 , typename I2 , typename I3
903  , typename I4 >
905  typename std::enable_if< Kokkos::Impl::are_integral<I0,I1,I2,I3,I4>::value &&
906  !is_layout_left, reference_type>::type
907  reference( const I0 & i0 , const I1 & i1 , const I2 & i2 , const I3 & i3
908  , const I4 & i4 ) const
909  { pointer_type beg = m_handle + m_array_offset(i0,i1,i2,i3,i4,0);
910 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && defined(__CUDA_ARCH__)
911  const unsigned index = threadIdx.x;
912  const unsigned strd = blockDim.x;
913  const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
914 #else
915  const unsigned index = m_fad_index;
916  const unsigned strd = m_fad_stride;
917  const unsigned size = m_fad_size.value;
918 #endif
919  return reference_type( beg + index
920  , beg + m_original_fad_size
921  , size
922  , strd ); }
923 
924  template< typename I0 , typename I1 , typename I2 , typename I3
925  , typename I4 , typename I5 >
927  typename std::enable_if< Kokkos::Impl::are_integral<I0,I1,I2,I3,I4,I5>::value &&
928  is_layout_left, reference_type>::type
929  reference( const I0 & i0 , const I1 & i1 , const I2 & i2 , const I3 & i3
930  , const I4 & i4 , const I5 & i5 ) const
931  { pointer_type beg = m_handle + m_array_offset(0,i0,i1,i2,i3,i4,i5);
932 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && defined(__CUDA_ARCH__)
933  const unsigned index = threadIdx.x;
934  const unsigned strd = blockDim.x;
935  const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
936 #else
937  const unsigned index = m_fad_index;
938  const unsigned strd = m_fad_stride;
939  const unsigned size = m_fad_size.value;
940 #endif
941  return reference_type( beg + index
942  , beg + m_original_fad_size
943  , size
944  , strd ); }
945 
946  template< typename I0 , typename I1 , typename I2 , typename I3
947  , typename I4 , typename I5 >
949  typename std::enable_if< Kokkos::Impl::are_integral<I0,I1,I2,I3,I4,I5>::value &&
950  !is_layout_left, reference_type>::type
951  reference( const I0 & i0 , const I1 & i1 , const I2 & i2 , const I3 & i3
952  , const I4 & i4 , const I5 & i5 ) const
953  { pointer_type beg = m_handle + m_array_offset(i0,i1,i2,i3,i4,i5,0);
954 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && defined(__CUDA_ARCH__)
955  const unsigned index = threadIdx.x;
956  const unsigned strd = blockDim.x;
957  const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
958 #else
959  const unsigned index = m_fad_index;
960  const unsigned strd = m_fad_stride;
961  const unsigned size = m_fad_size.value;
962 #endif
963  return reference_type( beg + index
964  , beg + m_original_fad_size
965  , size
966  , strd ); }
967 
968  template< typename I0 , typename I1 , typename I2 , typename I3
969  , typename I4 , typename I5 , typename I6 >
971  typename std::enable_if< Kokkos::Impl::are_integral<I0,I1,I2,I3,I4,I5,I6>::value &&
972  is_layout_left, reference_type>::type
973  reference( const I0 & i0 , const I1 & i1 , const I2 & i2 , const I3 & i3
974  , const I4 & i4 , const I5 & i5 , const I6 & i6 ) const
975  { pointer_type beg = m_handle + m_array_offset(0,i0,i1,i2,i3,i4,i5,i6);
976 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && defined(__CUDA_ARCH__)
977  const unsigned index = threadIdx.x;
978  const unsigned strd = blockDim.x;
979  const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
980 #else
981  const unsigned index = m_fad_index;
982  const unsigned strd = m_fad_stride;
983  const unsigned size = m_fad_size.value;
984 #endif
985  return reference_type( beg + index
986  , beg + m_original_fad_size
987  , size
988  , strd ); }
989 
990  template< typename I0 , typename I1 , typename I2 , typename I3
991  , typename I4 , typename I5 , typename I6 >
993  typename std::enable_if< Kokkos::Impl::are_integral<I0,I1,I2,I3,I4,I5,I6>::value &&
994  !is_layout_left, reference_type>::type
995  reference( const I0 & i0 , const I1 & i1 , const I2 & i2 , const I3 & i3
996  , const I4 & i4 , const I5 & i5 , const I6 & i6 ) const
997  { pointer_type beg = m_handle + m_array_offset(i0,i1,i2,i3,i4,i5,i6,0);
998 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && defined(__CUDA_ARCH__)
999  const unsigned index = threadIdx.x;
1000  const unsigned strd = blockDim.x;
1001  const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
1002 #else
1003  const unsigned index = m_fad_index;
1004  const unsigned strd = m_fad_stride;
1005  const unsigned size = m_fad_size.value;
1006 #endif
1007  return reference_type( beg + index
1008  , beg + m_original_fad_size
1009  , size
1010  , strd ); }
1011 
1012  //----------------------------------------
1013 
1016  static constexpr size_t memory_span( typename Traits::array_layout const & layout )
1017  {
1018  // Do not introduce padding...
1019  typedef std::integral_constant< unsigned , 0 > padding ;
1020  return array_offset_type(
1021  padding() ,
1022  create_fad_array_layout<unsigned(Rank), unsigned(FadStaticDimension)>( layout ) ).span() * sizeof(fad_value_type);
1023  }
1024 
1025  //----------------------------------------
1026 
1027  KOKKOS_INLINE_FUNCTION ~ViewMapping() = default ;
1028  KOKKOS_INLINE_FUNCTION ViewMapping() : m_handle(0) , m_offset() , m_array_offset() , m_fad_size(0) , m_original_fad_size(0) , m_fad_stride(1) , m_fad_index(0) {}
1029 
1030  KOKKOS_INLINE_FUNCTION ViewMapping( const ViewMapping & ) = default ;
1031  KOKKOS_INLINE_FUNCTION ViewMapping & operator = ( const ViewMapping & ) = default ;
1032 
1033  KOKKOS_INLINE_FUNCTION ViewMapping( ViewMapping && ) = default ;
1034  KOKKOS_INLINE_FUNCTION ViewMapping & operator = ( ViewMapping && ) = default ;
1035 
1036  template< class ... P >
1038  ViewMapping
1039  ( ViewCtorProp< P ... > const & prop
1040  , typename Traits::array_layout const & local_layout
1041  )
1042  : m_handle( ( (ViewCtorProp<void,pointer_type> const &) prop ).value )
1043  , m_offset( std::integral_constant< unsigned , 0 >()
1044  , local_layout )
1045  , m_array_offset(
1046  std::integral_constant< unsigned , 0 >()
1047  , create_fad_array_layout<unsigned(Rank), unsigned(FadStaticDimension)>( local_layout ) )
1048  , m_fad_size( getFadDimension<unsigned(Rank)>( m_array_offset ) - 1 )
1049  , m_original_fad_size( m_fad_size.value )
1050  , m_fad_stride( 1 )
1051  , m_fad_index( 0 )
1052  {
1053  const unsigned fad_dim =
1054  getFadDimension<unsigned(Rank)>( m_array_offset );
1055  if (unsigned(FadStaticDimension) == 0 && fad_dim == 0)
1056  Kokkos::abort("invalid fad dimension (0) supplied!");
1057  }
1058 
1059  //----------------------------------------
1060  /* Allocate and construct mapped array.
1061  * Allocate via shared allocation record and
1062  * return that record for allocation tracking.
1063  */
1064  template< class ... P >
1065  SharedAllocationRecord<> *
1066  allocate_shared( ViewCtorProp< P... > const & prop
1067  , typename Traits::array_layout const & local_layout )
1068  {
1069  typedef ViewCtorProp< P... > ctor_prop ;
1070 
1071  typedef typename ctor_prop::execution_space execution_space ;
1072  typedef typename Traits::memory_space memory_space ;
1073  typedef ViewValueFunctor< execution_space , fad_value_type > functor_type ;
1074  typedef SharedAllocationRecord< memory_space , functor_type > record_type ;
1075 
1076  // Disallow padding
1077  typedef std::integral_constant< unsigned , 0 > padding ;
1078 
1079  // Check if ViewCtorProp has CommonViewAllocProp - if so, retrieve the fad_size and append to layout
1080  using CVTR = typename Kokkos::Impl::CommonViewAllocProp< typename Kokkos::Impl::ViewSpecializeSacadoFadContiguous
1081  , typename Traits::value_type>;
1082 
1083  enum { test_traits_check = Kokkos::Impl::check_has_common_view_alloc_prop< P... >::value };
1084 
1085  typename Traits::array_layout internal_layout =
1086  (test_traits_check == true
1087  && ((Kokkos::Impl::ViewCtorProp<void, CVTR> const &)prop).value.is_view_type)
1088  ? Kokkos::Impl::appendFadToLayoutViewAllocHelper< Traits, ctor_prop >::returnNewLayoutPlusFad(prop, local_layout)
1089  : local_layout;
1090 
1091  m_offset = offset_type( padding(), internal_layout );
1092 
1093  m_array_offset =
1094  array_offset_type( padding() ,
1095  create_fad_array_layout<unsigned(Rank), unsigned(FadStaticDimension)>( internal_layout ) );
1096  const unsigned fad_dim =
1097  getFadDimension<unsigned(Rank)>( m_array_offset );
1098  if (unsigned(FadStaticDimension) == 0 && fad_dim == 0)
1099  Kokkos::abort("invalid fad dimension (0) supplied!");
1100  m_fad_size = fad_dim - 1 ;
1101  m_original_fad_size = m_fad_size.value ;
1102  m_fad_stride = 1;
1103  m_fad_index = 0;
1104 
1105  const size_t alloc_size = m_array_offset.span() * sizeof(fad_value_type);
1106 
1107  // Create shared memory tracking record with allocate memory from the memory space
1108  record_type * const record =
1109  record_type::allocate( ( (ViewCtorProp<void,memory_space> const &) prop ).value
1110  , ( (ViewCtorProp<void,std::string> const &) prop ).value
1111  , alloc_size );
1112 
1113  // Only set the the pointer and initialize if the allocation is non-zero.
1114  // May be zero if one of the dimensions is zero.
1115  if ( alloc_size ) {
1116 
1117  m_handle = handle_type( reinterpret_cast< pointer_type >( record->data() ) );
1118 
1119  if ( ctor_prop::initialize ) {
1120  // Assume destruction is only required when construction is requested.
1121  // The ViewValueFunctor has both value construction and destruction operators.
1122  record->m_destroy = functor_type( ( (ViewCtorProp<void,execution_space> const &) prop).value
1123  , (fad_value_type *) m_handle
1124  , m_array_offset.span()
1125  );
1126 
1127  // Construct values
1128  record->m_destroy.construct_shared_allocation();
1129  }
1130  }
1131 
1132  return record ;
1133  }
1134 
1135 };
1136 
1137 } // namespace Impl
1138 } // namespace Kokkos
1139 
1140 //----------------------------------------------------------------------------
1141 
1142 namespace Kokkos {
1143 namespace Impl {
1144 
1149 template< class DstTraits , class SrcTraits >
1150 class ViewMapping< DstTraits , SrcTraits ,
1151  typename std::enable_if<(
1152  Kokkos::Impl::MemorySpaceAccess
1153  < typename DstTraits::memory_space
1154  , typename SrcTraits::memory_space >::assignable
1155  &&
1156  // Destination view has FAD
1157  std::is_same< typename DstTraits::specialize
1158  , ViewSpecializeSacadoFadContiguous >::value
1159  &&
1160  // Source view has FAD
1161  std::is_same< typename SrcTraits::specialize
1162  , ViewSpecializeSacadoFadContiguous >::value
1163  )>::type >
1164 {
1165 public:
1166 
1167  enum { is_assignable = true };
1168 
1169  typedef Kokkos::Impl::SharedAllocationTracker TrackType ;
1170  typedef ViewMapping< DstTraits , void > DstType ;
1171  typedef ViewMapping< SrcTraits , void > SrcFadType ;
1172 
1173  template< class DstType >
1174  KOKKOS_INLINE_FUNCTION static
1175  void assign( DstType & dst
1176  , const SrcFadType & src
1177  , const TrackType & )
1178  {
1179  static_assert(
1180  (
1181  std::is_same< typename DstTraits::array_layout
1182  , Kokkos::LayoutLeft >::value ||
1183  std::is_same< typename DstTraits::array_layout
1184  , Kokkos::LayoutRight >::value ||
1185  std::is_same< typename DstTraits::array_layout
1186  , Kokkos::LayoutStride >::value
1187  )
1188  &&
1189  (
1190  std::is_same< typename SrcTraits::array_layout
1191  , Kokkos::LayoutLeft >::value ||
1192  std::is_same< typename SrcTraits::array_layout
1193  , Kokkos::LayoutRight >::value ||
1194  std::is_same< typename SrcTraits::array_layout
1195  , Kokkos::LayoutStride >::value
1196  )
1197  , "View of FAD requires LayoutLeft, LayoutRight, or LayoutStride" );
1198 
1199  static_assert(
1200  std::is_same< typename DstTraits::array_layout
1201  , typename SrcTraits::array_layout >::value ||
1202  std::is_same< typename DstTraits::array_layout
1203  , Kokkos::LayoutStride >::value ,
1204  "View assignment must have compatible layout" );
1205 
1206  static_assert(
1207  std::is_same< typename DstTraits::scalar_array_type
1208  , typename SrcTraits::scalar_array_type >::value ||
1209  std::is_same< typename DstTraits::scalar_array_type
1210  , typename SrcTraits::const_scalar_array_type >::value ,
1211  "View assignment must have same value type or const = non-const" );
1212 
1213  static_assert(
1214  ViewDimensionAssignable
1215  < typename DstType::offset_type::dimension_type
1216  , typename SrcFadType::offset_type::dimension_type >::value ,
1217  "View assignment must have compatible dimensions" );
1218 
1219  static_assert(
1220  ViewDimensionAssignable
1221  < typename DstType::array_offset_type::dimension_type
1222  , typename SrcFadType::array_offset_type::dimension_type >::value ,
1223  "View assignment must have compatible dimensions" );
1224 
1225  typedef typename DstType::offset_type dst_offset_type ;
1226  typedef typename DstType::array_offset_type dst_array_offset_type ;
1227 
1228  dst.m_handle = src.m_handle ;
1229  dst.m_offset = dst_offset_type( src.m_offset );
1230  dst.m_array_offset = dst_array_offset_type( src.m_array_offset );
1231  dst.m_fad_size = src.m_fad_size.value ;
1232  dst.m_original_fad_size = src.m_original_fad_size ;
1233  dst.m_fad_stride = src.m_fad_stride ;
1234  dst.m_fad_index = src.m_fad_index ;
1235  }
1236 };
1237 
1242 template< class DstTraits , class SrcTraits >
1243 class ViewMapping< DstTraits , SrcTraits ,
1244  typename std::enable_if<(
1245  std::is_same< typename DstTraits::memory_space
1246  , typename SrcTraits::memory_space >::value
1247  &&
1248  // Destination view has FAD
1249  std::is_same< typename DstTraits::specialize
1250  , ViewSpecializeSacadoFad >::value
1251  &&
1252  // Source view has FAD contiguous
1253  std::is_same< typename SrcTraits::specialize
1254  , ViewSpecializeSacadoFadContiguous >::value
1255  &&
1256  // Destination view is LayoutStride
1257  std::is_same< typename DstTraits::array_layout
1258  , Kokkos::LayoutStride >::value
1259  )>::type >
1260 {
1261 public:
1262 
1263  enum { is_assignable = true };
1264 
1265  typedef Kokkos::Impl::SharedAllocationTracker TrackType ;
1266  typedef ViewMapping< DstTraits , void > DstType ;
1267  typedef ViewMapping< SrcTraits , void > SrcFadType ;
1268 
1269  template< class DstType >
1270  KOKKOS_INLINE_FUNCTION static
1271  void assign( DstType & dst
1272  , const SrcFadType & src
1273  , const TrackType & )
1274  {
1275  static_assert(
1276  std::is_same< typename SrcTraits::array_layout
1277  , Kokkos::LayoutLeft >::value ||
1278  std::is_same< typename SrcTraits::array_layout
1279  , Kokkos::LayoutRight >::value ||
1280  std::is_same< typename SrcTraits::array_layout
1281  , Kokkos::LayoutStride >::value ,
1282  "View of FAD requires LayoutLeft, LayoutRight, or LayoutStride" );
1283 
1284  static_assert(
1285  std::is_same< typename DstTraits::value_type
1286  , typename SrcTraits::value_type >::value ||
1287  std::is_same< typename DstTraits::value_type
1288  , typename SrcTraits::const_value_type >::value ,
1289  "View assignment must have same value type or const = non-const" );
1290 
1291  static_assert(
1292  DstTraits::dimension::rank == SrcTraits::dimension::rank,
1293  "View assignment must have same rank" );
1294 
1295  typedef typename DstType::array_offset_type dst_offset_type ;
1296 
1297  dst.m_handle = src.m_handle ;
1298  dst.m_fad_size = src.m_fad_size.value ;
1299  dst.m_fad_stride = src.m_fad_stride ;
1300  dst.m_offset = src.m_offset;
1301 
1302  size_t N[8], S[8];
1303  N[0] = src.m_array_offset.dimension_0();
1304  N[1] = src.m_array_offset.dimension_1();
1305  N[2] = src.m_array_offset.dimension_2();
1306  N[3] = src.m_array_offset.dimension_3();
1307  N[4] = src.m_array_offset.dimension_4();
1308  N[5] = src.m_array_offset.dimension_5();
1309  N[6] = src.m_array_offset.dimension_6();
1310  N[7] = src.m_array_offset.dimension_7();
1311  S[0] = src.m_array_offset.stride_0();
1312  S[1] = src.m_array_offset.stride_1();
1313  S[2] = src.m_array_offset.stride_2();
1314  S[3] = src.m_array_offset.stride_3();
1315  S[4] = src.m_array_offset.stride_4();
1316  S[5] = src.m_array_offset.stride_5();
1317  S[6] = src.m_array_offset.stride_6();
1318  S[7] = src.m_array_offset.stride_7();
1319 
1320  // For LayoutLeft, we have to move the Sacado dimension from the first
1321  // to the last
1322  if (std::is_same< typename SrcTraits::array_layout
1323  , Kokkos::LayoutLeft >::value)
1324  {
1325  const size_t N_fad = N[0];
1326  const size_t S_fad = S[0];
1327  for (int i=0; i<7; ++i) {
1328  N[i] = N[i+1];
1329  S[i] = S[i+1];
1330  }
1331  N[DstTraits::dimension::rank] = N_fad;
1332  S[DstTraits::dimension::rank] = S_fad;
1333  }
1334  Kokkos::LayoutStride ls( N[0], S[0],
1335  N[1], S[1],
1336  N[2], S[2],
1337  N[3], S[3],
1338  N[4], S[4],
1339  N[5], S[5],
1340  N[6], S[6],
1341  N[7], S[7] );
1342  dst.m_array_offset = dst_offset_type(std::integral_constant<unsigned,0>(), ls);
1343  }
1344 };
1345 
1350 template< class DstTraits , class SrcTraits >
1351 class ViewMapping< DstTraits , SrcTraits ,
1352  typename std::enable_if<(
1353  std::is_same< typename DstTraits::memory_space
1354  , typename SrcTraits::memory_space >::value
1355  &&
1356  // Destination view has ordinary
1357  std::is_same< typename DstTraits::specialize , void >::value
1358  &&
1359  // Source view has FAD only
1360  std::is_same< typename SrcTraits::specialize
1361  , ViewSpecializeSacadoFadContiguous >::value
1362  )>::type >
1363 {
1364 public:
1365 
1366  enum { is_assignable = true };
1367 
1368  typedef Kokkos::Impl::SharedAllocationTracker TrackType ;
1369  typedef ViewMapping< DstTraits , void > DstType ;
1370  typedef ViewMapping< SrcTraits , void > SrcFadType ;
1371 
1372 
1373  // Helpers to assign, and generate if necessary, ViewOffset to the dst map
1374  // These are necessary to use Kokkos' deep_copy with nested fads
1375  template < class DstType, class SrcFadType, class Enable = void >
1376  struct AssignOffset;
1377 
1378  template < class DstType, class SrcFadType >
1379  struct AssignOffset< DstType, SrcFadType, typename std::enable_if< ((int)DstType::offset_type::dimension_type::rank != (int)SrcFadType::array_offset_type::dimension_type::rank) >::type >
1380  {
1381  // ViewOffset's Dimensions Ranks do not match
1383  static void assign( DstType & dst, const SrcFadType & src )
1384  {
1385  typedef typename SrcTraits::value_type TraitsValueType;
1386 
1389  )
1390  {
1391 
1392  typedef typename DstType::offset_type::array_layout DstLayoutType;
1393  //typedef typename ViewArrayLayoutSelector<typename DstType::offset_type::array_layout>::type DstLayoutType;
1394  typedef typename SrcFadType::array_offset_type::dimension_type SrcViewDimension;
1395 
1396  // This is the static dimension of the inner fad, missing from ViewDimension
1397  const size_t InnerStaticDim = Sacado::StaticSize< typename Sacado::ValueType< TraitsValueType >::type >::value;
1398 
1399  static constexpr bool is_layout_left =
1400  std::is_same< DstLayoutType, Kokkos::LayoutLeft>::value;
1401 
1402  typedef typename std::conditional< is_layout_left,
1403  typename SrcViewDimension:: template prepend< InnerStaticDim+1 >::type,
1404  typename SrcViewDimension:: template append < InnerStaticDim+1 >::type
1405  >::type SrcViewDimensionAppended;
1406 
1407  typedef std::integral_constant< unsigned , 0 > padding ;
1408 
1409  typedef ViewOffset< SrcViewDimensionAppended, DstLayoutType > TmpOffsetType;
1410 
1411  auto src_layout = src.m_array_offset.layout();
1412 
1413  if ( is_layout_left ) {
1414  auto prepend_layout = Kokkos::Impl::prependFadToLayout< DstLayoutType >::returnNewLayoutPlusFad(src_layout, InnerStaticDim+1);
1415  TmpOffsetType offset_tmp( padding(), prepend_layout );
1416  dst.m_offset = offset_tmp;
1417  }
1418  else {
1419  TmpOffsetType offset_tmp( padding(), src_layout );
1420  dst.m_offset = offset_tmp;
1421  }
1422  } else {
1423  Kokkos::abort("Sacado error: Applying AssignOffset for case with nested Fads, but without nested Fads - something went wrong");
1424  }
1425  }
1426  };
1427 
1428  template < class DstType, class SrcFadType >
1429  struct AssignOffset< DstType, SrcFadType, typename std::enable_if< ((int)DstType::offset_type::dimension_type::rank == (int)SrcFadType::array_offset_type::dimension_type::rank) >::type >
1430  {
1432  static void assign( DstType & dst, const SrcFadType & src )
1433  {
1434  typedef typename DstType::offset_type dst_offset_type ;
1435  dst.m_offset = dst_offset_type( src.m_array_offset );
1436  }
1437  };
1438 
1439  template< class DstType >
1440  KOKKOS_INLINE_FUNCTION static
1441  void assign( DstType & dst
1442  , const SrcFadType & src
1443  , const TrackType &
1444  )
1445  {
1446  static_assert(
1447  (
1448  std::is_same< typename DstTraits::array_layout
1449  , Kokkos::LayoutLeft >::value ||
1450  std::is_same< typename DstTraits::array_layout
1451  , Kokkos::LayoutRight >::value ||
1452  std::is_same< typename DstTraits::array_layout
1453  , Kokkos::LayoutStride >::value
1454  )
1455  &&
1456  (
1457  std::is_same< typename SrcTraits::array_layout
1458  , Kokkos::LayoutLeft >::value ||
1459  std::is_same< typename SrcTraits::array_layout
1460  , Kokkos::LayoutRight >::value ||
1461  std::is_same< typename SrcTraits::array_layout
1462  , Kokkos::LayoutStride >::value
1463  )
1464  , "View of FAD requires LayoutLeft, LayoutRight, or LayoutStride" );
1465 
1466  static_assert(
1467  std::is_same< typename DstTraits::array_layout
1468  , typename SrcTraits::array_layout >::value ||
1469  std::is_same< typename DstTraits::array_layout
1470  , Kokkos::LayoutStride >::value ,
1471  "View assignment must have compatible layout" );
1472 
1473  if ( src.m_fad_index != 0 || src.m_fad_stride != 1 ) {
1474  Kokkos::abort("\n\n ****** Kokkos::View< Sacado::Fad ... > Cannot assign to array with partitioned view ******\n\n");
1475  }
1476 
1477  AssignOffset< DstType, SrcFadType >::assign( dst, src );
1478  dst.m_handle = reinterpret_cast< typename DstType::handle_type >(src.m_handle) ;
1479  }
1480 };
1481 
1482 } // namespace Impl
1483 } // namespace Kokkos
1484 
1485 //----------------------------------------------------------------------------
1486 
1487 namespace Kokkos {
1488 namespace Impl {
1489 
1490 // Rules for subview arguments and layouts matching
1491 
1492 template<class LayoutDest, class LayoutSrc, int RankDest, int RankSrc, int CurrentArg, class ... SubViewArgs>
1493 struct SubviewLegalArgsCompileTime<Kokkos::LayoutContiguous<LayoutDest>,LayoutSrc,RankDest,RankSrc,CurrentArg,SubViewArgs...> {
1494  enum { value = SubviewLegalArgsCompileTime<LayoutDest,LayoutSrc,RankDest,RankSrc,CurrentArg,SubViewArgs...>::value };
1495 };
1496 
1497 template<class LayoutDest, class LayoutSrc, int RankDest, int RankSrc, int CurrentArg, class ... SubViewArgs>
1498 struct SubviewLegalArgsCompileTime<LayoutDest,Kokkos::LayoutContiguous<LayoutSrc>,RankDest,RankSrc,CurrentArg,SubViewArgs...> {
1499  enum { value = SubviewLegalArgsCompileTime<LayoutDest,LayoutSrc,RankDest,RankSrc,CurrentArg,SubViewArgs...>::value };
1500 };
1501 
1502 template<class LayoutDest, class LayoutSrc, int RankDest, int RankSrc, int CurrentArg, class ... SubViewArgs>
1503 struct SubviewLegalArgsCompileTime<Kokkos::LayoutContiguous<LayoutDest>,Kokkos::LayoutContiguous<LayoutSrc>,RankDest,RankSrc,CurrentArg,SubViewArgs...> {
1504  enum { value = SubviewLegalArgsCompileTime<LayoutDest,LayoutSrc,RankDest,RankSrc,CurrentArg,SubViewArgs...>::value };
1505 };
1506 
1507 // Subview mapping
1508 
1509 template< class SrcTraits , class Arg0 , class ... Args >
1510 struct ViewMapping
1511  < typename std::enable_if<(
1512  // Source view has FAD only
1513  std::is_same< typename SrcTraits::specialize
1514  , ViewSpecializeSacadoFadContiguous >::value
1515  &&
1516  (
1517  std::is_same< typename SrcTraits::array_layout
1518  , Kokkos::LayoutLeft >::value ||
1519  std::is_same< typename SrcTraits::array_layout
1520  , Kokkos::LayoutRight >::value ||
1521  std::is_same< typename SrcTraits::array_layout
1522  , Kokkos::LayoutStride >::value
1523  )
1524  && !Sacado::Fad::is_fad_partition<Arg0>::value
1525  )>::type
1526  , SrcTraits
1527  , Arg0, Args ... >
1528 {
1529 private:
1530 
1531  static_assert( SrcTraits::rank == sizeof...(Args)+1 , "" );
1532 
1533  enum
1534  { RZ = false
1535  , R0 = bool(is_integral_extent<0,Arg0,Args...>::value)
1536  , R1 = bool(is_integral_extent<1,Arg0,Args...>::value)
1537  , R2 = bool(is_integral_extent<2,Arg0,Args...>::value)
1538  , R3 = bool(is_integral_extent<3,Arg0,Args...>::value)
1539  , R4 = bool(is_integral_extent<4,Arg0,Args...>::value)
1540  , R5 = bool(is_integral_extent<5,Arg0,Args...>::value)
1541  , R6 = bool(is_integral_extent<6,Arg0,Args...>::value)
1542  };
1543 
1544  // Public rank
1545  enum { rank = unsigned(R0) + unsigned(R1) + unsigned(R2) + unsigned(R3)
1546  + unsigned(R4) + unsigned(R5) + unsigned(R6) };
1547 
1548  // Whether right-most non-FAD rank is a range.
1549  enum { R0_rev = ( 0 == SrcTraits::rank ? RZ : (
1550  1 == SrcTraits::rank ? R0 : (
1551  2 == SrcTraits::rank ? R1 : (
1552  3 == SrcTraits::rank ? R2 : (
1553  4 == SrcTraits::rank ? R3 : (
1554  5 == SrcTraits::rank ? R4 : (
1555  6 == SrcTraits::rank ? R5 : R6 ))))))) };
1556 
1557  // Subview's layout
1558  typedef typename std::conditional<
1559  ( /* Same array layout IF */
1560  ( rank == 0 ) /* output rank zero */
1561  ||
1562  // OutputRank 1 or 2, InputLayout Left, Interval 0
1563  // because single stride one or second index has a stride.
1564  ( rank <= 2 && R0 && std::is_same< typename SrcTraits::array_layout , Kokkos::LayoutLeft >::value )
1565  ||
1566  // OutputRank 1 or 2, InputLayout Right, Interval [InputRank-1]
1567  // because single stride one or second index has a stride.
1568  ( rank <= 2 && R0_rev && std::is_same< typename SrcTraits::array_layout , Kokkos::LayoutRight >::value )
1570  >::type array_layout ;
1571 
1572  typedef typename SrcTraits::value_type fad_type ;
1573 
1574  typedef typename std::conditional< rank == 0 , fad_type ,
1575  typename std::conditional< rank == 1 , fad_type * ,
1576  typename std::conditional< rank == 2 , fad_type ** ,
1577  typename std::conditional< rank == 3 , fad_type *** ,
1578  typename std::conditional< rank == 4 , fad_type **** ,
1579  typename std::conditional< rank == 5 , fad_type ***** ,
1580  typename std::conditional< rank == 6 , fad_type ****** ,
1581  fad_type *******
1582  >::type >::type >::type >::type >::type >::type >::type
1583  data_type ;
1584 
1585 public:
1586 
1587  typedef Kokkos::ViewTraits
1588  < data_type
1589  , array_layout
1590  , typename SrcTraits::device_type
1591  , typename SrcTraits::memory_traits > traits_type ;
1592 
1593  typedef Kokkos::View
1594  < data_type
1595  , array_layout
1596  , typename SrcTraits::device_type
1597  , typename SrcTraits::memory_traits > type ;
1598 
1599 
1601  static void assign( ViewMapping< traits_type , void > & dst
1602  , ViewMapping< SrcTraits , void > const & src
1603  , Arg0 arg0 , Args ... args )
1604  {
1605  typedef ViewMapping< traits_type , void > DstType ;
1606  typedef typename DstType::offset_type dst_offset_type ;
1607  typedef typename DstType::array_offset_type dst_array_offset_type ;
1608  typedef typename DstType::handle_type dst_handle_type ;
1609 
1610  size_t offset;
1611  if (std::is_same< typename SrcTraits::array_layout, LayoutLeft >::value) {
1612  const SubviewExtents< SrcTraits::rank + 1 , rank + 1 >
1613  array_extents( src.m_array_offset.m_dim , Kokkos::ALL() , arg0 , args... );
1614  offset = src.m_array_offset( array_extents.domain_offset(0)
1615  , array_extents.domain_offset(1)
1616  , array_extents.domain_offset(2)
1617  , array_extents.domain_offset(3)
1618  , array_extents.domain_offset(4)
1619  , array_extents.domain_offset(5)
1620  , array_extents.domain_offset(6)
1621  , array_extents.domain_offset(7) );
1622  dst.m_array_offset = dst_array_offset_type( src.m_array_offset ,
1623  array_extents );
1624  }
1625  else {
1626  const SubviewExtents< SrcTraits::rank + 1 , rank + 1 >
1627  array_extents( src.m_array_offset.m_dim , arg0 , args... , Kokkos::ALL() );
1628  offset = src.m_array_offset( array_extents.domain_offset(0)
1629  , array_extents.domain_offset(1)
1630  , array_extents.domain_offset(2)
1631  , array_extents.domain_offset(3)
1632  , array_extents.domain_offset(4)
1633  , array_extents.domain_offset(5)
1634  , array_extents.domain_offset(6)
1635  , array_extents.domain_offset(7) );
1636  dst.m_array_offset = dst_array_offset_type( src.m_array_offset ,
1637  array_extents );
1638  }
1639 
1640  const SubviewExtents< SrcTraits::rank , rank >
1641  extents( src.m_offset.m_dim , arg0 , args... );
1642 
1643  dst.m_offset = dst_offset_type( src.m_offset , extents );
1644  dst.m_handle = dst_handle_type( src.m_handle + offset );
1645  dst.m_fad_size = src.m_fad_size;
1646  dst.m_original_fad_size = src.m_original_fad_size;
1647  dst.m_fad_stride = src.m_fad_stride;
1648  dst.m_fad_index = src.m_fad_index;
1649  }
1650 
1651 };
1652 
1653 } // namespace Impl
1654 } // namespace Kokkos
1655 
1656 //---------------------------------------------------------------------------
1657 
1658 namespace Kokkos {
1659 namespace Impl {
1660 
1661 // Partition mapping
1662 
1663 template< class DataType, class ...P, unsigned Stride >
1664 class ViewMapping<
1665  void,
1666  ViewTraits<DataType,P...> ,
1667  Sacado::Fad::Partition<Stride> >
1668 {
1669 public:
1670 
1671  enum { is_assignable = true };
1672 
1673  typedef ViewTraits<DataType,P...> src_traits;
1674  typedef ViewMapping< src_traits , void > src_type ;
1675 
1676  typedef typename src_type::offset_type::dimension_type src_dimension;
1677  typedef typename src_traits::value_type fad_type;
1678  typedef typename Sacado::LocalScalarType<fad_type,Stride>::type strided_fad_type;
1679  typedef typename
1680  ViewDataType< strided_fad_type , src_dimension >::type strided_data_type;
1681  typedef ViewTraits<strided_data_type,P...> dst_traits;
1682  typedef View<strided_data_type,P...> type;
1683  typedef ViewMapping< dst_traits , void > dst_type ;
1684 
1685  KOKKOS_INLINE_FUNCTION static
1686  void assign( dst_type & dst
1687  , const src_type & src
1688  , const Sacado::Fad::Partition<Stride> & part )
1689  {
1690  if ( Stride != part.stride && Stride != 0 ) {
1691  Kokkos::abort("\n\n ****** Kokkos::View< Sacado::Fad ... > Invalid size in partitioned view assignment ******\n\n");
1692  }
1693  if ( src.m_fad_stride != 1 ) {
1694  Kokkos::abort("\n\n ****** Kokkos::View< Sacado::Fad ... > Can't partition already partitioned view ******\n\n");
1695  }
1696 
1697  dst.m_handle = src.m_handle ;
1698  dst.m_offset = src.m_offset ;
1699  dst.m_array_offset = src.m_array_offset ;
1700 
1701  // Assuming the alignment was choosen correctly for the partitioning,
1702  // each partition should get the same size. This allows the use of SFad.
1703  dst.m_fad_size =
1704  (src.m_fad_size.value + part.stride-part.offset-1) / part.stride ;
1705 
1706  dst.m_original_fad_size = src.m_original_fad_size ;
1707  dst.m_fad_stride = part.stride ;
1708  dst.m_fad_index = part.offset ;
1709  }
1710 };
1711 
1712 } // namespace Impl
1713 } // namespace Kokkos
1714 
1715 #endif // defined(HAVE_SACADO_VIEW_SPEC) && !defined(SACADO_DISABLE_FAD_VIEW_SPEC)
1716 
1717 #endif // defined(HAVE_SACADO_KOKKOSCORE)
1718 
1719 #endif /* #ifndef KOKKOS_EXPERIMENTAL_VIEW_SACADO_FAD_HPP */
View
expr expr dx(i)
Base template specification for whether a type is a Fad type.
GeneralFad< StaticStorage< T, Num > > SLFad
Base template specification for static size.
#define KOKKOS_INLINE_FUNCTION
#define T
Definition: Sacado_rad.hpp:573
#define D
Definition: Sacado_rad.hpp:577
GeneralFad< DynamicStorage< T > > DFad
expr expr expr bar false
#define KOKKOS_FORCEINLINE_FUNCTION
GeneralFad< StaticFixedStorage< T, Num > > SFad
Base template specification for testing whether type is statically sized.
Get view type for any Fad type.