30 #ifndef KOKKOS_EXPERIMENTAL_VIEW_SACADO_FAD_CONTIGUOUS_HPP 31 #define KOKKOS_EXPERIMENTAL_VIEW_SACADO_FAD_CONTIGUOUS_HPP 34 #if defined(HAVE_SACADO_KOKKOSCORE) 42 template <
typename ViewType,
typename Enabled =
void>
43 struct ThreadLocalScalarType {
44 typedef typename ViewType::non_const_value_type type;
47 template <
typename ViewType>
48 struct ViewScalarStride {
49 static const unsigned stride =
51 static const bool is_unit_stride =
62 template <
unsigned Size = 0>
64 static const unsigned PartitionSize = Size;
68 template<
typename iType0 ,
typename iType1 >
70 Partition(
const iType0 & i0 ,
const iType1 & i1 ) :
71 offset(i0), stride(i1) {
76 struct is_fad_partition {
77 static const bool value =
false;
80 template <
unsigned Str
ide>
81 struct is_fad_partition< Partition<Stride> > {
82 static const bool value =
true;
88 template <
typename T,
unsigned Str
ide = 0>
89 struct LocalScalarType {
92 template <
typename T,
unsigned Str
ide>
93 struct LocalScalarType<const
T, Stride> {
94 typedef typename LocalScalarType<T,Stride>::type lst;
95 typedef const lst type;
106 template <
typename T,
int N>
class StaticStorage;
107 template <
typename S>
class GeneralFad;
110 template <
typename T,
int N,
unsigned Str
ide>
111 struct LocalScalarType< Fad::Exp::GeneralFad< Fad::Exp::StaticStorage<T,N> >,
113 static const int Ns = (N+Stride-1) / Stride;
114 typedef Fad::Exp::GeneralFad< Fad::Exp::StaticStorage<T,Ns> > type;
116 #ifndef SACADO_NEW_FAD_DESIGN_IS_DEFAULT 118 template <
typename T,
int N>
class SLFad;
120 template <
typename T,
int N,
unsigned Str
ide>
121 struct LocalScalarType< Fad::
SLFad<T,N>, Stride > {
122 static const int Ns = (N+Stride-1) / Stride;
123 typedef Fad::SLFad<T,Ns> type;
133 template <
typename T,
int N>
class StaticFixedStorage;
134 template <
typename T,
int N>
class StaticStorage;
135 template <
typename S>
class GeneralFad;
138 template <
typename T,
int N,
unsigned Str
ide>
139 struct LocalScalarType< Fad::Exp::GeneralFad< Fad::Exp::StaticFixedStorage<T,N> >,
141 static const int Ns = (N+Stride-1) / Stride;
142 typedef typename std::conditional<
144 Fad::Exp::GeneralFad< Fad::Exp::StaticFixedStorage<T,Ns> > ,
145 Fad::Exp::GeneralFad< Fad::Exp::StaticStorage<T,Ns> >
149 #ifndef SACADO_NEW_FAD_DESIGN_IS_DEFAULT 151 template <
typename T,
int N>
class SFad;
153 template <
typename T,
int N,
unsigned Str
ide>
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;
160 template <
unsigned Str
ide,
typename T>
162 const T& partition_scalar(
const T& x) {
return x; }
167 #if defined(HAVE_SACADO_VIEW_SPEC) && !defined(SACADO_DISABLE_FAD_VIEW_SPEC) 170 #include "Kokkos_Core.hpp" 171 #include "impl/Kokkos_ViewMapping.hpp" 177 #if defined(__CUDA_ARCH__) 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;
186 #ifndef SACADO_VIEW_CUDA_HIERARCHICAL_DFAD 187 template <
unsigned Str
ide,
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());
200 const T*
dx = x.dx();
201 for (
int i=0; i<size; ++i)
202 xp.fastAccessDx(i) =
dx[offset+i*Stride];
207 template <
unsigned Str
ide,
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);
219 template <
unsigned Str
ide,
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);
232 #ifndef SACADO_NEW_FAD_DESIGN_IS_DEFAULT 234 template <
typename T>
class DFad;
235 template <
typename T,
int N>
class SLFad;
236 template <
typename T,
int N>
class SFad;
238 #ifndef SACADO_VIEW_CUDA_HIERARCHICAL_DFAD 239 template <
unsigned Str
ide,
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());
252 const T*
dx = x.dx();
253 for (
int i=0; i<size; ++i)
254 xp.fastAccessDx(i) =
dx[offset+i*Stride];
259 template <
unsigned Str
ide,
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);
271 template <
unsigned Str
ide,
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);
283 #endif // SACADO_NEW_FAD_DESIGN_IS_DEFAULT 285 #endif // defined(__CUDA_ARCH__) 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 )
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);
306 template <
typename ViewType>
307 struct ThreadLocalScalarType<
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;
317 #if defined (KOKKOS_ENABLE_CUDA) && defined(SACADO_VIEW_CUDA_HIERARCHICAL) 318 template<
class OutputView >
319 struct SacadoViewFill<
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 )
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;
335 const OutputView output ;
336 const_value_type input ;
339 void operator()(
const size_t i0 )
const 341 local_scalar_type input_stride = Sacado::partition_scalar<stride>(input);
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);
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 ;
363 void operator()(
const team_handle& team )
const 365 const size_t i0 = team.league_rank()*team.team_size() + team.team_rank();
366 if (i0 < output.extent(0))
370 SacadoViewFill(
const OutputView & arg_out , const_value_type & arg_in )
371 : output( arg_out ), input( arg_in )
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();
390 template<
class ... Args >
391 struct is_ViewSpecializeSacadoFadContiguous {
enum { value =
false }; };
393 template<
class D ,
class ... P ,
class ... Args >
394 struct is_ViewSpecializeSacadoFadContiguous<
Kokkos::
View<D,P...> , Args... > {
396 std::is_same<
typename Kokkos::ViewTraits<
D,P...>::specialize
397 , ViewSpecializeSacadoFadContiguous >::value
399 ( (
sizeof...(Args) == 0 ) ||
400 is_ViewSpecializeSacadoFadContiguous< Args... >::value ) };
414 constexpr
unsigned computeFadPartitionSize(
unsigned size,
unsigned stride)
417 ((size+stride-1)/stride) == (size/stride) ? ((size+stride-1)/stride) : 0;
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,
428 create_fad_array_layout(
const Layout& layout)
431 for (
int i=0; i<8; ++i)
432 dims[i] = layout.dimension[i];
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] );
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)
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];
452 if (static_dim > 0) {
453 dims[rank] = static_dim+1;
456 return Layout( dims[0], strides[0],
463 dims[7], strides[7] );
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)
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)
481 return Layout( dims[0], dims[1], dims[2], dims[3],
482 dims[4], dims[5], dims[6], dims[7] );
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)
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() )))))));
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)
506 return offset.dimension_0();
509 template<
class Traits >
510 class ViewMapping< Traits ,
511 typename
std::enable_if<
512 ( std::is_same< typename Traits::specialize
513 , ViewSpecializeSacadoFadContiguous >::value
515 ( std::is_same< typename Traits::array_layout
516 , Kokkos::LayoutLeft >::value
518 std::is_same< typename Traits::array_layout
519 , Kokkos::LayoutRight >::value
521 std::is_same< typename Traits::array_layout
522 , Kokkos::LayoutStride >::value
528 template< class ,
class ... >
friend class ViewMapping ;
529 template< class ,
class ... >
friend class Kokkos::View ;
531 typedef typename Traits::value_type fad_type ;
534 std::add_const< fad_value_type >::type const_fad_value_type ;
539 enum { PartitionedFadStride = Traits::array_layout::scalar_stride };
543 enum { PartitionedFadStaticDimension =
544 computeFadPartitionSize(FadStaticDimension,PartitionedFadStride) };
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;
550 typedef fad_type thread_local_scalar_type;
556 typedef fad_value_type * handle_type ;
558 typedef ViewArrayAnalysis< typename Traits::data_type > array_analysis ;
560 typedef ViewOffset<
typename Traits::dimension
561 ,
typename Traits::array_layout
566 static constexpr
bool is_layout_left =
567 std::is_same< typename Traits::array_layout, Kokkos::LayoutLeft>::value;
569 typename std::conditional<
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,
579 handle_type m_handle ;
580 offset_type m_offset ;
581 array_offset_type m_array_offset ;
582 sacado_size_type m_fad_size ;
585 unsigned m_original_fad_size ;
586 unsigned m_fad_stride ;
587 unsigned m_fad_index ;
594 enum { Rank = Traits::dimension::rank };
597 enum { Sacado_Rank = std::is_same< typename Traits::array_layout, Kokkos::LayoutLeft >::value ? 0 : Rank+1 };
600 template<
typename iType >
602 {
return m_offset.m_dim.extent(r); }
605 typename Traits::array_layout layout()
const 606 {
return m_offset.layout(); }
609 {
return m_offset.dimension_0(); }
611 {
return m_offset.dimension_1(); }
613 {
return m_offset.dimension_2(); }
615 {
return m_offset.dimension_3(); }
617 {
return m_offset.dimension_4(); }
619 {
return m_offset.dimension_5(); }
621 {
return m_offset.dimension_6(); }
623 {
return m_offset.dimension_7(); }
628 using is_regular = std::false_type ;
632 {
return m_offset.stride_0(); }
634 {
return m_offset.stride_1(); }
636 {
return m_offset.stride_2(); }
638 {
return m_offset.stride_3(); }
640 {
return m_offset.stride_4(); }
642 {
return m_offset.stride_5(); }
644 {
return m_offset.stride_6(); }
646 {
return m_offset.stride_7(); }
648 template<
typename iType >
650 { m_offset.stride(s); }
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; }
657 {
return m_fad_size.value+1; }
662 {
return m_fad_stride; }
667 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && defined(__CUDA_ARCH__) 681 typedef fad_value_type * pointer_type ;
685 {
return m_array_offset.span(); }
689 {
return m_array_offset.span_is_contiguous() && (m_fad_stride == 1); }
693 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && defined(__CUDA_ARCH__) 694 {
return m_handle + threadIdx.x; }
696 {
return m_handle + m_fad_index; }
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;
710 const unsigned index = m_fad_index;
711 const unsigned strd = m_fad_stride;
712 const unsigned size = m_fad_size.value;
714 return reference_type( m_handle + index
715 , m_handle + m_original_fad_size
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;
730 const unsigned index = m_fad_index;
731 const unsigned strd = m_fad_stride;
732 const unsigned size = m_fad_size.value;
734 return reference_type( beg + index
735 , beg + m_original_fad_size
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;
750 const unsigned index = m_fad_index;
751 const unsigned strd = m_fad_stride;
752 const unsigned size = m_fad_size.value;
754 return reference_type( beg + index
755 , beg + m_original_fad_size
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;
770 const unsigned index = m_fad_index;
771 const unsigned strd = m_fad_stride;
772 const unsigned size = m_fad_size.value;
774 return reference_type( beg + index
775 , beg + m_original_fad_size
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;
790 const unsigned index = m_fad_index;
791 const unsigned strd = m_fad_stride;
792 const unsigned size = m_fad_size.value;
794 return reference_type( beg + index
795 , beg + m_original_fad_size
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;
811 const unsigned index = m_fad_index;
812 const unsigned strd = m_fad_stride;
813 const unsigned size = m_fad_size.value;
815 return reference_type( beg + index
816 , beg + m_original_fad_size
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;
831 const unsigned index = m_fad_index;
832 const unsigned strd = m_fad_stride;
833 const unsigned size = m_fad_size.value;
835 return reference_type( beg + index
836 , beg + m_original_fad_size
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;
851 const unsigned index = m_fad_index;
852 const unsigned strd = m_fad_stride;
853 const unsigned size = m_fad_size.value;
855 return reference_type( beg + index
856 , beg + m_original_fad_size
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;
871 const unsigned index = m_fad_index;
872 const unsigned strd = m_fad_stride;
873 const unsigned size = m_fad_size.value;
875 return reference_type( beg + index
876 , beg + m_original_fad_size
880 template<
typename I0 ,
typename I1 ,
typename I2 ,
typename I3
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;
893 const unsigned index = m_fad_index;
894 const unsigned strd = m_fad_stride;
895 const unsigned size = m_fad_size.value;
897 return reference_type( beg + index
898 , beg + m_original_fad_size
902 template<
typename I0 ,
typename I1 ,
typename I2 ,
typename I3
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;
915 const unsigned index = m_fad_index;
916 const unsigned strd = m_fad_stride;
917 const unsigned size = m_fad_size.value;
919 return reference_type( beg + index
920 , beg + m_original_fad_size
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;
937 const unsigned index = m_fad_index;
938 const unsigned strd = m_fad_stride;
939 const unsigned size = m_fad_size.value;
941 return reference_type( beg + index
942 , beg + m_original_fad_size
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;
959 const unsigned index = m_fad_index;
960 const unsigned strd = m_fad_stride;
961 const unsigned size = m_fad_size.value;
963 return reference_type( beg + index
964 , beg + m_original_fad_size
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;
981 const unsigned index = m_fad_index;
982 const unsigned strd = m_fad_stride;
983 const unsigned size = m_fad_size.value;
985 return reference_type( beg + index
986 , beg + m_original_fad_size
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;
1003 const unsigned index = m_fad_index;
1004 const unsigned strd = m_fad_stride;
1005 const unsigned size = m_fad_size.value;
1007 return reference_type( beg + index
1008 , beg + m_original_fad_size
1016 static constexpr
size_t memory_span(
typename Traits::array_layout
const & layout )
1019 typedef std::integral_constant< unsigned , 0 > padding ;
1020 return array_offset_type(
1022 create_fad_array_layout<
unsigned(Rank),
unsigned(FadStaticDimension)>( layout ) ).span() *
sizeof(fad_value_type);
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) {}
1036 template<
class ... P >
1039 ( ViewCtorProp< P ... >
const & prop
1040 ,
typename Traits::array_layout
const & local_layout
1042 : m_handle( ( (ViewCtorProp<void,pointer_type> const &) prop ).value )
1043 , m_offset(
std::integral_constant< unsigned , 0 >()
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 )
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!");
1064 template<
class ... P >
1065 SharedAllocationRecord<> *
1066 allocate_shared( ViewCtorProp< P... >
const & prop
1067 ,
typename Traits::array_layout
const & local_layout )
1069 typedef ViewCtorProp< P... > ctor_prop ;
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 ;
1077 typedef std::integral_constant< unsigned , 0 > padding ;
1080 using CVTR =
typename Kokkos::Impl::CommonViewAllocProp<
typename Kokkos::Impl::ViewSpecializeSacadoFadContiguous
1081 ,
typename Traits::value_type>;
1083 enum { test_traits_check = Kokkos::Impl::check_has_common_view_alloc_prop< P... >::value };
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)
1091 m_offset = offset_type( padding(), internal_layout );
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 ;
1105 const size_t alloc_size = m_array_offset.span() *
sizeof(fad_value_type);
1108 record_type *
const record =
1109 record_type::allocate( ( (ViewCtorProp<void,memory_space>
const &) prop ).value
1110 , ( (ViewCtorProp<void,std::string>
const &) prop ).value
1117 m_handle = handle_type( reinterpret_cast< pointer_type >( record->data() ) );
1119 if ( ctor_prop::initialize ) {
1122 record->m_destroy = functor_type( ( (ViewCtorProp<void,execution_space>
const &) prop).value
1123 , (fad_value_type *) m_handle
1124 , m_array_offset.span()
1128 record->m_destroy.construct_shared_allocation();
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
1157 std::is_same< typename DstTraits::specialize
1158 , ViewSpecializeSacadoFadContiguous >::value
1161 std::is_same< typename SrcTraits::specialize
1162 , ViewSpecializeSacadoFadContiguous >::value
1167 enum { is_assignable =
true };
1169 typedef Kokkos::Impl::SharedAllocationTracker TrackType ;
1170 typedef ViewMapping< DstTraits , void > DstType ;
1171 typedef ViewMapping< SrcTraits , void > SrcFadType ;
1173 template<
class DstType >
1175 void assign( DstType & dst
1176 ,
const SrcFadType & src
1177 ,
const TrackType & )
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
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
1197 ,
"View of FAD requires LayoutLeft, LayoutRight, or LayoutStride" );
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" );
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" );
1214 ViewDimensionAssignable
1215 <
typename DstType::offset_type::dimension_type
1216 ,
typename SrcFadType::offset_type::dimension_type >::value ,
1217 "View assignment must have compatible dimensions" );
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" );
1225 typedef typename DstType::offset_type dst_offset_type ;
1226 typedef typename DstType::array_offset_type dst_array_offset_type ;
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 ;
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
1249 std::is_same< typename DstTraits::specialize
1250 , ViewSpecializeSacadoFad >::value
1253 std::is_same< typename SrcTraits::specialize
1254 , ViewSpecializeSacadoFadContiguous >::value
1257 std::is_same< typename DstTraits::array_layout
1258 , Kokkos::LayoutStride >::value
1263 enum { is_assignable =
true };
1265 typedef Kokkos::Impl::SharedAllocationTracker TrackType ;
1266 typedef ViewMapping< DstTraits , void > DstType ;
1267 typedef ViewMapping< SrcTraits , void > SrcFadType ;
1269 template<
class DstType >
1271 void assign( DstType & dst
1272 ,
const SrcFadType & src
1273 ,
const TrackType & )
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" );
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" );
1292 DstTraits::dimension::rank == SrcTraits::dimension::rank,
1293 "View assignment must have same rank" );
1295 typedef typename DstType::array_offset_type dst_offset_type ;
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;
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();
1322 if (std::is_same<
typename SrcTraits::array_layout
1323 , Kokkos::LayoutLeft >::value)
1325 const size_t N_fad = N[0];
1326 const size_t S_fad = S[0];
1327 for (
int i=0; i<7; ++i) {
1331 N[DstTraits::dimension::rank] = N_fad;
1332 S[DstTraits::dimension::rank] = S_fad;
1334 Kokkos::LayoutStride ls( N[0], S[0],
1342 dst.m_array_offset = dst_offset_type(std::integral_constant<unsigned,0>(), ls);
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
1357 std::is_same< typename DstTraits::specialize , void >::value
1360 std::is_same< typename SrcTraits::specialize
1361 , ViewSpecializeSacadoFadContiguous >::value
1366 enum { is_assignable =
true };
1368 typedef Kokkos::Impl::SharedAllocationTracker TrackType ;
1369 typedef ViewMapping< DstTraits , void > DstType ;
1370 typedef ViewMapping< SrcTraits , void > SrcFadType ;
1375 template <
class DstType,
class SrcFadType,
class Enable =
void >
1376 struct AssignOffset;
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 >
1383 static void assign( DstType & dst,
const SrcFadType & src )
1385 typedef typename SrcTraits::value_type TraitsValueType;
1392 typedef typename DstType::offset_type::array_layout DstLayoutType;
1394 typedef typename SrcFadType::array_offset_type::dimension_type SrcViewDimension;
1399 static constexpr
bool is_layout_left =
1400 std::is_same< DstLayoutType, Kokkos::LayoutLeft>::value;
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;
1407 typedef std::integral_constant< unsigned , 0 > padding ;
1409 typedef ViewOffset< SrcViewDimensionAppended, DstLayoutType > TmpOffsetType;
1411 auto src_layout = src.m_array_offset.layout();
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;
1419 TmpOffsetType offset_tmp( padding(), src_layout );
1420 dst.m_offset = offset_tmp;
1423 Kokkos::abort(
"Sacado error: Applying AssignOffset for case with nested Fads, but without nested Fads - something went wrong");
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 >
1432 static void assign( DstType & dst,
const SrcFadType & src )
1434 typedef typename DstType::offset_type dst_offset_type ;
1435 dst.m_offset = dst_offset_type( src.m_array_offset );
1439 template<
class DstType >
1441 void assign( DstType & dst
1442 ,
const SrcFadType & src
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
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
1464 ,
"View of FAD requires LayoutLeft, LayoutRight, or LayoutStride" );
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" );
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");
1477 AssignOffset< DstType, SrcFadType >::assign( dst, src );
1478 dst.m_handle =
reinterpret_cast< typename DstType::handle_type
>(src.m_handle) ;
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 };
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 };
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 };
1509 template<
class SrcTraits ,
class Arg0 ,
class ... Args >
1511 < typename
std::enable_if<(
1513 std::is_same< typename SrcTraits::specialize
1514 , ViewSpecializeSacadoFadContiguous >::value
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
1524 && !Sacado::Fad::is_fad_partition<Arg0>::value
1531 static_assert( SrcTraits::rank ==
sizeof...(Args)+1 ,
"" );
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)
1545 enum { rank = unsigned(R0) + unsigned(R1) + unsigned(R2) + unsigned(R3)
1546 + unsigned(R4) + unsigned(R5) + unsigned(R6) };
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 ))))))) };
1558 typedef typename std::conditional<
1564 ( rank <= 2 && R0 && std::is_same< typename SrcTraits::array_layout , Kokkos::LayoutLeft >::value )
1568 ( rank <= 2 && R0_rev && std::is_same< typename SrcTraits::array_layout , Kokkos::LayoutRight >::value )
1570 >::type array_layout ;
1572 typedef typename SrcTraits::value_type fad_type ;
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 ****** ,
1582 >::type >::type >::type >::type >::type >::type >::type
1587 typedef Kokkos::ViewTraits
1590 ,
typename SrcTraits::device_type
1591 ,
typename SrcTraits::memory_traits > traits_type ;
1593 typedef Kokkos::View
1596 ,
typename SrcTraits::device_type
1597 ,
typename SrcTraits::memory_traits > type ;
1601 static void assign( ViewMapping< traits_type , void > & dst
1602 , ViewMapping< SrcTraits , void >
const & src
1603 , Arg0 arg0 , Args ... args )
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 ;
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 ,
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 ,
1640 const SubviewExtents< SrcTraits::rank , rank >
1641 extents( src.m_offset.m_dim , arg0 , args... );
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;
1663 template<
class DataType,
class ...P,
unsigned Stride >
1666 ViewTraits<DataType,P...> ,
1667 Sacado::Fad::Partition<Stride> >
1671 enum { is_assignable =
true };
1673 typedef ViewTraits<DataType,P...> src_traits;
1674 typedef ViewMapping< src_traits , void > src_type ;
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;
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 ;
1686 void assign( dst_type & dst
1687 ,
const src_type & src
1688 ,
const Sacado::Fad::Partition<Stride> & part )
1690 if ( Stride != part.stride && Stride != 0 ) {
1691 Kokkos::abort(
"\n\n ****** Kokkos::View< Sacado::Fad ... > Invalid size in partitioned view assignment ******\n\n");
1693 if ( src.m_fad_stride != 1 ) {
1694 Kokkos::abort(
"\n\n ****** Kokkos::View< Sacado::Fad ... > Can't partition already partitioned view ******\n\n");
1697 dst.m_handle = src.m_handle ;
1698 dst.m_offset = src.m_offset ;
1699 dst.m_array_offset = src.m_array_offset ;
1704 (src.m_fad_size.value + part.stride-part.offset-1) / part.stride ;
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 ;
1715 #endif // defined(HAVE_SACADO_VIEW_SPEC) && !defined(SACADO_DISABLE_FAD_VIEW_SPEC) 1717 #endif // defined(HAVE_SACADO_KOKKOSCORE)
static const bool is_unit_stride
Base template specification for whether a type is a Fad type.
GeneralFad< StaticStorage< T, Num > > SLFad
static const unsigned stride
Base template specification for static size.
#define KOKKOS_INLINE_FUNCTION
GeneralFad< DynamicStorage< T > > DFad
#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.