30 #ifndef KOKKOS_EXPERIMENTAL_VIEW_SACADO_FAD_CONTIGUOUS_HPP
31 #define KOKKOS_EXPERIMENTAL_VIEW_SACADO_FAD_CONTIGUOUS_HPP
34 #if defined(HAVE_SACADO_KOKKOS)
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 =
50 Impl::LayoutScalarStride< typename ViewType::array_layout>::stride;
51 static const bool is_unit_stride =
52 Impl::LayoutScalarStride< typename ViewType::array_layout>::is_unit_stride;
62 template <
unsigned Size = 0>
64 static const unsigned PartitionSize = Size;
68 template<
typename iType0 ,
typename iType1 >
69 KOKKOS_INLINE_FUNCTION
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>
161 KOKKOS_INLINE_FUNCTION
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__) || defined(__HIP_DEVICE_COMPILE__)
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>
188 KOKKOS_INLINE_FUNCTION
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>
208 KOKKOS_INLINE_FUNCTION
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>
220 KOKKOS_INLINE_FUNCTION
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>
240 KOKKOS_INLINE_FUNCTION
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>
260 KOKKOS_INLINE_FUNCTION
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>
272 KOKKOS_INLINE_FUNCTION
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__) || defined(__HIP_DEVICE_COMPILE__)
293 template<
unsigned Stride,
typename D,
typename ... P >
294 KOKKOS_INLINE_FUNCTION
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, typename TraitsType::specialize> 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_impl_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 ;
338 KOKKOS_INLINE_FUNCTION
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);
350 for (
size_t i1 = 0 ; i1 < n1 ; ++i1 ) {
351 for (
size_t i2 = 0 ; i2 < n2 ; ++i2 ) {
352 for (
size_t i3 = 0 ; i3 < n3 ; ++i3 ) {
353 for (
size_t i4 = 0 ; i4 < n4 ; ++i4 ) {
354 for (
size_t i5 = 0 ; i5 < n5 ; ++i5 ) {
355 for (
size_t i6 = 0 ; i6 < n6 ; ++i6 ) {
356 output.access(i0,i1,i2,i3,i4,i5,i6) = input_stride ;
360 KOKKOS_INLINE_FUNCTION
361 void operator()(
const team_impl_handle& team )
const
363 const size_t i0 = team.league_rank()*team.team_size() + team.team_rank();
364 if (i0 < output.extent(0))
368 SacadoViewFill(
const OutputView & arg_out , const_value_type & arg_in )
369 : output( arg_out ), input( arg_in )
371 const size_t team_size = 256 / stride;
372 team_policy policy( (output.extent(0)+team_size-1)/team_size ,
373 team_size , stride );
374 Kokkos::parallel_for( policy, *
this );
379 #if defined (KOKKOS_ENABLE_HIP) && defined(SACADO_VIEW_CUDA_HIERARCHICAL)
380 template<
class OutputView >
381 struct SacadoViewFill<
383 typename std::enable_if<
384 ( Kokkos::is_view_fad_contiguous<OutputView>::value &&
385 std::is_same<typename OutputView::execution_space, Kokkos::HIP>::value &&
386 !Kokkos::ViewScalarStride<OutputView>::is_unit_stride )
390 typedef typename OutputView::const_value_type const_value_type ;
391 typedef typename OutputView::execution_space execution_space ;
392 typedef Kokkos::TeamPolicy< execution_space> team_policy;
393 typedef typename team_policy::member_type team_impl_handle;
394 typedef typename Kokkos::ThreadLocalScalarType<OutputView>::type local_scalar_type;
395 static const unsigned stride = Kokkos::ViewScalarStride<OutputView>::stride;
397 const OutputView output ;
398 const_value_type input ;
400 KOKKOS_INLINE_FUNCTION
401 void operator()(
const size_t i0 )
const
403 local_scalar_type input_stride = Sacado::partition_scalar<stride>(input);
405 const size_t n1 = output.extent(1);
406 const size_t n2 = output.extent(2);
407 const size_t n3 = output.extent(3);
408 const size_t n4 = output.extent(4);
409 const size_t n5 = output.extent(5);
410 const size_t n6 = output.extent(6);
411 const size_t n7 = output.extent(7);
413 for (
size_t i1 = 0 ; i1 < n1 ; ++i1 ) {
414 for (
size_t i2 = 0 ; i2 < n2 ; ++i2 ) {
415 for (
size_t i3 = 0 ; i3 < n3 ; ++i3 ) {
416 for (
size_t i4 = 0 ; i4 < n4 ; ++i4 ) {
417 for (
size_t i5 = 0 ; i5 < n5 ; ++i5 ) {
418 for (
size_t i6 = 0 ; i6 < n6 ; ++i6 ) {
419 output.access(i0,i1,i2,i3,i4,i5,i6) = input_stride ;
423 KOKKOS_INLINE_FUNCTION
424 void operator()(
const team_impl_handle& team )
const
426 const size_t i0 = team.league_rank()*team.team_size() + team.team_rank();
427 if (i0 < output.extent(0))
431 SacadoViewFill(
const OutputView & arg_out , const_value_type & arg_in )
432 : output( arg_out ), input( arg_in )
434 const size_t team_size = 256 / stride;
435 team_policy policy( (output.extent(0)+team_size-1)/team_size ,
436 team_size , stride );
437 Kokkos::parallel_for( policy, *
this );
450 template<
class ... Args >
451 struct is_ViewSpecializeSacadoFadContiguous {
enum {
value =
false }; };
453 template<
class D ,
class ... P ,
class ... Args >
454 struct is_ViewSpecializeSacadoFadContiguous< Kokkos::
View<D,P...> , Args... > {
456 std::is_same<
typename Kokkos::ViewTraits<
D,P...>::specialize
457 , ViewSpecializeSacadoFadContiguous >
::value
459 ( (
sizeof...(Args) == 0 ) ||
460 is_ViewSpecializeSacadoFadContiguous< Args... >
::value ) };
473 KOKKOS_INLINE_FUNCTION
474 constexpr
unsigned computeFadPartitionSize(
unsigned size,
unsigned stride)
477 ((size+stride-1)/stride) == (size/stride) ? ((size+stride-1)/stride) : 0;
483 template <
unsigned rank,
unsigned static_dim,
typename Layout>
484 KOKKOS_INLINE_FUNCTION
488 create_fad_array_layout(
const Layout& layout)
491 for (
int i=0;
i<8; ++
i)
492 dims[
i] = layout.dimension[
i];
494 dims[rank] = static_dim+1;
495 return Layout( dims[0], dims[1], dims[2], dims[3],
496 dims[4], dims[5], dims[6], dims[7] );
502 template <
unsigned rank,
unsigned static_dim,
typename Layout>
503 KOKKOS_INLINE_FUNCTION
505 create_fad_array_layout(
const Layout& layout)
507 size_t dims[8], strides[8];
508 for (
int i=0;
i<8; ++
i) {
509 dims[
i] = layout.dimension[
i];
510 strides[
i] = layout.stride[
i];
512 if (static_dim > 0) {
513 dims[rank] = static_dim+1;
516 return Layout( dims[0], strides[0],
523 dims[7], strides[7] );
529 template <
unsigned rank,
unsigned static_dim,
typename Layout>
530 KOKKOS_INLINE_FUNCTION
532 create_fad_array_layout(
const Layout& layout)
535 for (
int i=0;
i<8; ++
i)
536 dims[
i] = layout.dimension[
i];
537 size_t fad_dim = static_dim == 0 ? dims[rank] : static_dim+1;
538 for (
int i=rank;
i>=1; --
i)
541 return Layout( dims[0], dims[1], dims[2], dims[3],
542 dims[4], dims[5], dims[6], dims[7] );
545 template <
unsigned Rank,
typename Dimension,
typename Layout>
546 KOKKOS_INLINE_FUNCTION
548 getFadDimension(
const ViewOffset<Dimension,Layout,void>& offset)
551 ( Rank == 0 ? offset.dimension_0() :
552 ( Rank == 1 ? offset.dimension_1() :
553 ( Rank == 2 ? offset.dimension_2() :
554 ( Rank == 3 ? offset.dimension_3() :
555 ( Rank == 4 ? offset.dimension_4() :
556 ( Rank == 5 ? offset.dimension_5() :
557 ( Rank == 6 ? offset.dimension_6() :
558 offset.dimension_7() )))))));
561 template <
unsigned Rank,
typename Dimension,
typename Layout>
562 KOKKOS_INLINE_FUNCTION
564 getFadDimension(
const ViewOffset<Dimension,Layout,void>& offset)
566 return offset.dimension_0();
569 template<
class Traits >
570 class ViewMapping< Traits ,
571 typename std::enable_if<
572 ( std::is_same< typename Traits::specialize
573 , ViewSpecializeSacadoFadContiguous >::value
575 ( std::is_same< typename Traits::array_layout
576 , Kokkos::LayoutLeft >::value
578 std::is_same< typename Traits::array_layout
579 , Kokkos::LayoutRight >::value
581 std::is_same< typename Traits::array_layout
582 , Kokkos::LayoutStride >::value
585 , typename Traits::specialize
590 template< class ,
class ... >
friend class ViewMapping ;
591 template< class ,
class ... >
friend class Kokkos::View ;
593 typedef typename Traits::value_type fad_type ;
596 std::add_const< fad_value_type >::type const_fad_value_type ;
599 enum { is_assignable_data_type =
true };
602 enum { PartitionedFadStride = Traits::array_layout::scalar_stride };
606 enum { PartitionedFadStaticDimension =
607 computeFadPartitionSize(FadStaticDimension,PartitionedFadStride) };
609 #ifdef KOKKOS_ENABLE_CUDA
610 typedef typename Sacado::LocalScalarType< fad_type, unsigned(PartitionedFadStride) >::type strided_scalar_type;
612 #elif defined(KOKKOS_ENABLE_HIP)
613 typedef typename Sacado::LocalScalarType< fad_type, unsigned(PartitionedFadStride) >::type strided_scalar_type;
616 typedef fad_type thread_local_scalar_type;
622 typedef fad_value_type * handle_type ;
624 typedef ViewArrayAnalysis< typename Traits::data_type > array_analysis ;
626 typedef ViewOffset<
typename Traits::dimension
627 ,
typename Traits::array_layout
632 static constexpr
bool is_layout_left =
635 typename std::conditional<
637 typename array_analysis::dimension::
638 template prepend<0>::type,
639 typename array_analysis::dimension::
640 template append<0>::type >::type,
641 typename Traits::array_layout,
645 handle_type m_impl_handle ;
646 offset_type m_impl_offset ;
647 array_offset_type m_array_offset ;
648 sacado_size_type m_fad_size ;
651 unsigned m_original_fad_size ;
652 unsigned m_fad_stride ;
653 unsigned m_fad_index ;
660 enum { Rank = Traits::dimension::rank };
666 template<
typename iType >
667 KOKKOS_INLINE_FUNCTION constexpr
size_t extent(
const iType & r )
const
668 {
return m_impl_offset.m_dim.extent(r); }
670 KOKKOS_INLINE_FUNCTION constexpr
671 typename Traits::array_layout layout()
const
672 {
return m_impl_offset.layout(); }
674 KOKKOS_INLINE_FUNCTION constexpr
size_t dimension_0()
const
675 {
return m_impl_offset.dimension_0(); }
676 KOKKOS_INLINE_FUNCTION constexpr
size_t dimension_1()
const
677 {
return m_impl_offset.dimension_1(); }
678 KOKKOS_INLINE_FUNCTION constexpr
size_t dimension_2()
const
679 {
return m_impl_offset.dimension_2(); }
680 KOKKOS_INLINE_FUNCTION constexpr
size_t dimension_3()
const
681 {
return m_impl_offset.dimension_3(); }
682 KOKKOS_INLINE_FUNCTION constexpr
size_t dimension_4()
const
683 {
return m_impl_offset.dimension_4(); }
684 KOKKOS_INLINE_FUNCTION constexpr
size_t dimension_5()
const
685 {
return m_impl_offset.dimension_5(); }
686 KOKKOS_INLINE_FUNCTION constexpr
size_t dimension_6()
const
687 {
return m_impl_offset.dimension_6(); }
688 KOKKOS_INLINE_FUNCTION constexpr
size_t dimension_7()
const
689 {
return m_impl_offset.dimension_7(); }
694 using is_regular = std::false_type ;
697 KOKKOS_INLINE_FUNCTION constexpr
size_t stride_0()
const
698 {
return m_impl_offset.stride_0(); }
699 KOKKOS_INLINE_FUNCTION constexpr
size_t stride_1()
const
700 {
return m_impl_offset.stride_1(); }
701 KOKKOS_INLINE_FUNCTION constexpr
size_t stride_2()
const
702 {
return m_impl_offset.stride_2(); }
703 KOKKOS_INLINE_FUNCTION constexpr
size_t stride_3()
const
704 {
return m_impl_offset.stride_3(); }
705 KOKKOS_INLINE_FUNCTION constexpr
size_t stride_4()
const
706 {
return m_impl_offset.stride_4(); }
707 KOKKOS_INLINE_FUNCTION constexpr
size_t stride_5()
const
708 {
return m_impl_offset.stride_5(); }
709 KOKKOS_INLINE_FUNCTION constexpr
size_t stride_6()
const
710 {
return m_impl_offset.stride_6(); }
711 KOKKOS_INLINE_FUNCTION constexpr
size_t stride_7()
const
712 {
return m_impl_offset.stride_7(); }
714 template<
typename iType >
715 KOKKOS_INLINE_FUNCTION
void stride( iType *
const s )
const
716 { m_impl_offset.stride(s); }
719 KOKKOS_FORCEINLINE_FUNCTION constexpr
unsigned dimension_scalar() const
720 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && ( defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) )
721 {
return PartitionedFadStaticDimension ? PartitionedFadStaticDimension+1 : (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x + 1; }
723 {
return m_fad_size.value+1; }
727 KOKKOS_FORCEINLINE_FUNCTION constexpr
unsigned stride_scalar()
const
728 {
return m_fad_stride; }
733 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && ( defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) )
747 typedef fad_value_type * pointer_type ;
750 KOKKOS_INLINE_FUNCTION constexpr
size_t span()
const
751 {
return m_array_offset.span(); }
754 KOKKOS_INLINE_FUNCTION constexpr
bool span_is_contiguous()
const
755 {
return m_array_offset.span_is_contiguous() && (m_fad_stride == 1); }
758 KOKKOS_INLINE_FUNCTION constexpr pointer_type data() const
759 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && (defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) )
760 {
return m_impl_handle + threadIdx.x; }
762 {
return m_impl_handle + m_fad_index; }
767 KOKKOS_FORCEINLINE_FUNCTION
771 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && ( defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) )
772 const unsigned index = threadIdx.x;
773 const unsigned strd = blockDim.x;
774 const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
776 const unsigned index = m_fad_index;
777 const unsigned strd = m_fad_stride;
778 const unsigned size = m_fad_size.value;
780 return reference_type( m_impl_handle + index
781 , m_impl_handle + m_original_fad_size
785 template<
typename I0 >
786 KOKKOS_FORCEINLINE_FUNCTION
788 is_layout_left, reference_type>::type
789 reference(
const I0 & i0 )
const
790 { pointer_type beg = m_impl_handle + m_array_offset(0,i0);
791 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && ( defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) )
792 const unsigned index = threadIdx.x;
793 const unsigned strd = blockDim.x;
794 const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
796 const unsigned index = m_fad_index;
797 const unsigned strd = m_fad_stride;
798 const unsigned size = m_fad_size.value;
800 return reference_type( beg + index
801 , beg + m_original_fad_size
805 template<
typename I0 >
806 KOKKOS_FORCEINLINE_FUNCTION
808 !is_layout_left, reference_type>::type
809 reference(
const I0 & i0 )
const
810 { pointer_type beg = m_impl_handle + m_array_offset(i0,0);
811 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && ( defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) )
812 const unsigned index = threadIdx.x;
813 const unsigned strd = blockDim.x;
814 const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
816 const unsigned index = m_fad_index;
817 const unsigned strd = m_fad_stride;
818 const unsigned size = m_fad_size.value;
820 return reference_type( beg + index
821 , beg + m_original_fad_size
825 template<
typename I0 ,
typename I1 >
826 KOKKOS_FORCEINLINE_FUNCTION
828 is_layout_left, reference_type>::type
829 reference(
const I0 & i0 ,
const I1 & i1 )
const
830 { pointer_type beg = m_impl_handle + m_array_offset(0,i0,i1);
831 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && (defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) )
832 const unsigned index = threadIdx.x;
833 const unsigned strd = blockDim.x;
834 const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
836 const unsigned index = m_fad_index;
837 const unsigned strd = m_fad_stride;
838 const unsigned size = m_fad_size.value;
840 return reference_type( beg + index
841 , beg + m_original_fad_size
845 template<
typename I0 ,
typename I1 >
846 KOKKOS_FORCEINLINE_FUNCTION
848 !is_layout_left, reference_type>::type
849 reference(
const I0 & i0 ,
const I1 & i1 )
const
850 { pointer_type beg = m_impl_handle + m_array_offset(i0,i1,0);
851 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && ( defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) )
852 const unsigned index = threadIdx.x;
853 const unsigned strd = blockDim.x;
854 const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
856 const unsigned index = m_fad_index;
857 const unsigned strd = m_fad_stride;
858 const unsigned size = m_fad_size.value;
860 return reference_type( beg + index
861 , beg + m_original_fad_size
866 template<
typename I0 ,
typename I1 ,
typename I2 >
867 KOKKOS_FORCEINLINE_FUNCTION
869 is_layout_left, reference_type>::type
870 reference(
const I0 & i0 ,
const I1 & i1 ,
const I2 & i2 )
const
871 { pointer_type beg = m_impl_handle + m_array_offset(0,i0,i1,i2);
872 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && ( defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) )
873 const unsigned index = threadIdx.x;
874 const unsigned strd = blockDim.x;
875 const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
877 const unsigned index = m_fad_index;
878 const unsigned strd = m_fad_stride;
879 const unsigned size = m_fad_size.value;
881 return reference_type( beg + index
882 , beg + m_original_fad_size
886 template<
typename I0 ,
typename I1 ,
typename I2 >
887 KOKKOS_FORCEINLINE_FUNCTION
889 !is_layout_left, reference_type>::type
890 reference(
const I0 & i0 ,
const I1 & i1 ,
const I2 & i2 )
const
891 { pointer_type beg = m_impl_handle + m_array_offset(i0,i1,i2,0);
892 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && ( defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) )
893 const unsigned index = threadIdx.x;
894 const unsigned strd = blockDim.x;
895 const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
897 const unsigned index = m_fad_index;
898 const unsigned strd = m_fad_stride;
899 const unsigned size = m_fad_size.value;
901 return reference_type( beg + index
902 , beg + m_original_fad_size
906 template<
typename I0 ,
typename I1 ,
typename I2 ,
typename I3 >
907 KOKKOS_FORCEINLINE_FUNCTION
909 is_layout_left, reference_type>::type
910 reference(
const I0 & i0 ,
const I1 & i1 ,
const I2 & i2 ,
const I3 & i3 )
const
911 { pointer_type beg = m_impl_handle + m_array_offset(0,i0,i1,i2,i3);
912 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && ( defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) )
913 const unsigned index = threadIdx.x;
914 const unsigned strd = blockDim.x;
915 const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
917 const unsigned index = m_fad_index;
918 const unsigned strd = m_fad_stride;
919 const unsigned size = m_fad_size.value;
921 return reference_type( beg + index
922 , beg + m_original_fad_size
926 template<
typename I0 ,
typename I1 ,
typename I2 ,
typename I3 >
927 KOKKOS_FORCEINLINE_FUNCTION
929 !is_layout_left, reference_type>::type
930 reference(
const I0 & i0 ,
const I1 & i1 ,
const I2 & i2 ,
const I3 & i3 )
const
931 { pointer_type beg = m_impl_handle + m_array_offset(i0,i1,i2,i3,0);
932 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && ( defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) )
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
948 KOKKOS_FORCEINLINE_FUNCTION
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
953 { pointer_type beg = m_impl_handle + m_array_offset(0,i0,i1,i2,i3,i4);
954 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && ( defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) )
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
970 KOKKOS_FORCEINLINE_FUNCTION
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
975 { pointer_type beg = m_impl_handle + m_array_offset(i0,i1,i2,i3,i4,0);
976 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && ( defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) )
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 >
992 KOKKOS_FORCEINLINE_FUNCTION
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
997 { pointer_type beg = m_impl_handle + m_array_offset(0,i0,i1,i2,i3,i4,i5);
998 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && ( defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) )
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
1012 template<
typename I0 ,
typename I1 ,
typename I2 ,
typename I3
1013 ,
typename I4 ,
typename I5 >
1014 KOKKOS_FORCEINLINE_FUNCTION
1016 !is_layout_left, reference_type>::type
1017 reference(
const I0 & i0 ,
const I1 & i1 ,
const I2 & i2 ,
const I3 & i3
1018 ,
const I4 & i4 ,
const I5 & i5 )
const
1019 { pointer_type beg = m_impl_handle + m_array_offset(i0,i1,i2,i3,i4,i5,0);
1020 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && ( defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) )
1021 const unsigned index = threadIdx.x;
1022 const unsigned strd = blockDim.x;
1023 const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
1025 const unsigned index = m_fad_index;
1026 const unsigned strd = m_fad_stride;
1027 const unsigned size = m_fad_size.value;
1029 return reference_type( beg + index
1030 , beg + m_original_fad_size
1034 template<
typename I0 ,
typename I1 ,
typename I2 ,
typename I3
1035 ,
typename I4 ,
typename I5 ,
typename I6 >
1036 KOKKOS_FORCEINLINE_FUNCTION
1038 is_layout_left, reference_type>::type
1039 reference(
const I0 & i0 ,
const I1 & i1 ,
const I2 & i2 ,
const I3 & i3
1040 ,
const I4 & i4 ,
const I5 & i5 ,
const I6 & i6 )
const
1041 { pointer_type beg = m_impl_handle + m_array_offset(0,i0,i1,i2,i3,i4,i5,i6);
1042 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && ( defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) )
1043 const unsigned index = threadIdx.x;
1044 const unsigned strd = blockDim.x;
1045 const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
1047 const unsigned index = m_fad_index;
1048 const unsigned strd = m_fad_stride;
1049 const unsigned size = m_fad_size.value;
1051 return reference_type( beg + index
1052 , beg + m_original_fad_size
1056 template<
typename I0 ,
typename I1 ,
typename I2 ,
typename I3
1057 ,
typename I4 ,
typename I5 ,
typename I6 >
1058 KOKKOS_FORCEINLINE_FUNCTION
1060 !is_layout_left, reference_type>::type
1061 reference(
const I0 & i0 ,
const I1 & i1 ,
const I2 & i2 ,
const I3 & i3
1062 ,
const I4 & i4 ,
const I5 & i5 ,
const I6 & i6 )
const
1063 { pointer_type beg = m_impl_handle + m_array_offset(i0,i1,i2,i3,i4,i5,i6,0);
1064 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) && ( defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) )
1065 const unsigned index = threadIdx.x;
1066 const unsigned strd = blockDim.x;
1067 const unsigned size = (m_fad_size.value+blockDim.x-threadIdx.x-1) / blockDim.x;
1069 const unsigned index = m_fad_index;
1070 const unsigned strd = m_fad_stride;
1071 const unsigned size = m_fad_size.value;
1073 return reference_type( beg + index
1074 , beg + m_original_fad_size
1081 KOKKOS_INLINE_FUNCTION
1082 static constexpr
size_t memory_span(
typename Traits::array_layout
const & layout )
1085 typedef std::integral_constant< unsigned , 0 > padding ;
1086 return array_offset_type(
1088 create_fad_array_layout<
unsigned(Rank),
unsigned(FadStaticDimension)>( layout ) ).span() *
sizeof(fad_value_type);
1093 KOKKOS_DEFAULTED_FUNCTION ~ViewMapping() = default ;
1094 KOKKOS_INLINE_FUNCTION ViewMapping() : m_impl_handle(0) , m_impl_offset() , m_array_offset() , m_fad_size(0) , m_original_fad_size(0) , m_fad_stride(1) , m_fad_index(0) {}
1096 KOKKOS_DEFAULTED_FUNCTION ViewMapping(
const ViewMapping & ) = default ;
1097 KOKKOS_DEFAULTED_FUNCTION ViewMapping & operator = (
const ViewMapping & ) = default ;
1099 KOKKOS_DEFAULTED_FUNCTION ViewMapping( ViewMapping && ) = default ;
1100 KOKKOS_DEFAULTED_FUNCTION ViewMapping & operator = ( ViewMapping && ) = default ;
1102 template<
class ... P >
1103 KOKKOS_INLINE_FUNCTION
1105 ( ViewCtorProp< P ... >
const & prop
1106 ,
typename Traits::array_layout
const & local_layout
1108 : m_impl_handle( ( (ViewCtorProp<
void,pointer_type> const &) prop ).
value )
1109 , m_impl_offset( std::integral_constant< unsigned , 0 >()
1112 std::integral_constant< unsigned , 0 >()
1113 , create_fad_array_layout<unsigned(Rank), unsigned(FadStaticDimension)>( local_layout ) )
1114 , m_fad_size( getFadDimension<unsigned(Rank)>( m_array_offset ) - 1 )
1115 , m_original_fad_size( m_fad_size.
value )
1120 getFadDimension<unsigned(Rank)>( m_array_offset );
1121 if (
unsigned(FadStaticDimension) == 0 && fad_dim == 0)
1122 Kokkos::abort(
"invalid fad dimension (0) supplied!");
1130 template<
class ... P >
1131 SharedAllocationRecord<> *
1132 allocate_shared( ViewCtorProp< P... >
const & prop
1133 ,
typename Traits::array_layout
const & local_layout
1134 ,
bool execution_space_specified)
1136 typedef ViewCtorProp< P... > ctor_prop ;
1138 typedef typename ctor_prop::execution_space execution_space ;
1139 typedef typename Traits::memory_space memory_space ;
1140 typedef ViewValueFunctor< execution_space , fad_value_type > functor_type ;
1141 typedef SharedAllocationRecord< memory_space , functor_type > record_type ;
1144 typedef std::integral_constant< unsigned , 0 > padding ;
1147 enum { test_traits_check = Kokkos::Impl::check_has_common_view_alloc_prop< P... >
::value };
1149 typename Traits::array_layout internal_layout =
1150 (test_traits_check ==
true)
1151 ? Kokkos::Impl::appendFadToLayoutViewAllocHelper< Traits, P... >::returnNewLayoutPlusFad(prop, local_layout)
1154 m_impl_offset = offset_type( padding(), internal_layout );
1157 array_offset_type( padding() ,
1158 create_fad_array_layout<
unsigned(Rank),
unsigned(FadStaticDimension)>( internal_layout ) );
1159 const unsigned fad_dim =
1160 getFadDimension<unsigned(Rank)>( m_array_offset );
1161 if (
unsigned(FadStaticDimension) == 0 && fad_dim == 0)
1162 Kokkos::abort(
"invalid fad dimension (0) supplied!");
1163 m_fad_size = fad_dim - 1 ;
1164 m_original_fad_size = m_fad_size.value ;
1168 const size_t alloc_size = m_array_offset.span() *
sizeof(fad_value_type);
1171 record_type *
const record =
1172 record_type::allocate( ( (ViewCtorProp<void,memory_space>
const &) prop ).
value
1173 , ( (ViewCtorProp<void,std::string>
const &) prop ).value
1180 m_impl_handle = handle_type( reinterpret_cast< pointer_type >( record->data() ) );
1182 if ( ctor_prop::initialize ) {
1185 if (execution_space_specified)
1186 record->m_destroy = functor_type( ( (ViewCtorProp<void,execution_space>
const &) prop).value
1187 , (fad_value_type *) m_impl_handle
1188 , m_array_offset.span()
1189 , record->get_label()
1192 record->m_destroy = functor_type((fad_value_type *) m_impl_handle
1193 , m_array_offset.span()
1194 , record->get_label()
1198 record->m_destroy.construct_shared_allocation();
1219 template<
class DstTraits ,
class SrcTraits >
1220 class ViewMapping< DstTraits , SrcTraits ,
1221 typename std::enable_if<(
1222 Kokkos::Impl::MemorySpaceAccess
1223 < typename DstTraits::memory_space
1224 , typename SrcTraits::memory_space >::assignable
1227 std::is_same< typename DstTraits::specialize
1228 , ViewSpecializeSacadoFadContiguous >::value
1231 std::is_same< typename SrcTraits::specialize
1232 , ViewSpecializeSacadoFadContiguous >::value
1234 , typename DstTraits::specialize
1239 enum { is_assignable =
true };
1240 enum { is_assignable_data_type =
true };
1242 typedef Kokkos::Impl::SharedAllocationTracker TrackType ;
1243 typedef ViewMapping< DstTraits , typename DstTraits::specialize > DstType ;
1244 typedef ViewMapping< SrcTraits , typename SrcTraits::specialize > SrcFadType ;
1246 template<
class DstType >
1247 KOKKOS_INLINE_FUNCTION
static
1248 void assign( DstType & dst
1249 ,
const SrcFadType & src
1250 ,
const TrackType & )
1254 std::is_same<
typename DstTraits::array_layout
1255 , Kokkos::LayoutLeft >::value ||
1256 std::is_same<
typename DstTraits::array_layout
1257 , Kokkos::LayoutRight >::value ||
1258 std::is_same<
typename DstTraits::array_layout
1259 , Kokkos::LayoutStride >::value
1263 std::is_same<
typename SrcTraits::array_layout
1264 , Kokkos::LayoutLeft >::value ||
1265 std::is_same<
typename SrcTraits::array_layout
1266 , Kokkos::LayoutRight >::value ||
1267 std::is_same<
typename SrcTraits::array_layout
1268 , Kokkos::LayoutStride >::value
1270 ,
"View of FAD requires LayoutLeft, LayoutRight, or LayoutStride" );
1273 std::is_same<
typename DstTraits::array_layout
1274 ,
typename SrcTraits::array_layout >::value ||
1275 std::is_same<
typename DstTraits::array_layout
1276 , Kokkos::LayoutStride >::value ,
1277 "View assignment must have compatible layout" );
1280 std::is_same<
typename DstTraits::value_type
1281 ,
typename SrcTraits::value_type >::value ||
1282 std::is_same<
typename DstTraits::value_type
1283 ,
typename SrcTraits::const_value_type >::value ,
1284 "View assignment must have same value type or const = non-const" );
1287 ViewDimensionAssignable
1288 <
typename DstType::offset_type::dimension_type
1289 ,
typename SrcFadType::offset_type::dimension_type >::value ,
1290 "View assignment must have compatible dimensions" );
1293 ViewDimensionAssignable
1294 <
typename DstType::array_offset_type::dimension_type
1295 ,
typename SrcFadType::array_offset_type::dimension_type >::value ,
1296 "View assignment must have compatible dimensions" );
1298 typedef typename DstType::offset_type dst_offset_type ;
1299 typedef typename DstType::array_offset_type dst_array_offset_type ;
1301 dst.m_impl_handle = src.m_impl_handle ;
1302 dst.m_impl_offset = dst_offset_type( src.m_impl_offset );
1303 dst.m_array_offset = dst_array_offset_type( src.m_array_offset );
1304 dst.m_fad_size = src.m_fad_size.value ;
1305 dst.m_original_fad_size = src.m_original_fad_size ;
1306 dst.m_fad_stride = src.m_fad_stride ;
1307 dst.m_fad_index = src.m_fad_index ;
1315 template<
class DstTraits ,
class SrcTraits >
1316 class ViewMapping< DstTraits , SrcTraits ,
1317 typename std::enable_if<(
1318 std::is_same< typename DstTraits::memory_space
1319 , typename SrcTraits::memory_space >::value
1322 std::is_same< typename DstTraits::specialize
1323 , ViewSpecializeSacadoFad >::value
1326 std::is_same< typename SrcTraits::specialize
1327 , ViewSpecializeSacadoFadContiguous >::value
1330 std::is_same< typename DstTraits::array_layout
1331 , Kokkos::LayoutStride >::value
1333 , typename DstTraits::specialize
1338 enum { is_assignable =
true };
1339 enum { is_assignable_data_type =
true };
1341 typedef Kokkos::Impl::SharedAllocationTracker TrackType ;
1342 typedef ViewMapping< DstTraits , typename DstTraits::specialize > DstType ;
1343 typedef ViewMapping< SrcTraits , typename SrcTraits::specialize > SrcFadType ;
1345 template<
class DstType >
1346 KOKKOS_INLINE_FUNCTION
static
1347 void assign( DstType & dst
1348 ,
const SrcFadType & src
1349 ,
const TrackType & )
1352 std::is_same<
typename SrcTraits::array_layout
1353 , Kokkos::LayoutLeft >::value ||
1354 std::is_same<
typename SrcTraits::array_layout
1355 , Kokkos::LayoutRight >::value ||
1356 std::is_same<
typename SrcTraits::array_layout
1357 , Kokkos::LayoutStride >::value ,
1358 "View of FAD requires LayoutLeft, LayoutRight, or LayoutStride" );
1361 std::is_same<
typename DstTraits::value_type
1362 ,
typename SrcTraits::value_type >::value ||
1363 std::is_same<
typename DstTraits::value_type
1364 ,
typename SrcTraits::const_value_type >::value ,
1365 "View assignment must have same value type or const = non-const" );
1368 DstTraits::dimension::rank == SrcTraits::dimension::rank,
1369 "View assignment must have same rank" );
1371 typedef typename DstType::array_offset_type dst_offset_type ;
1373 dst.m_impl_handle = src.m_impl_handle ;
1374 dst.m_fad_size = src.m_fad_size.value ;
1375 dst.m_fad_stride = src.m_fad_stride ;
1376 dst.m_impl_offset = src.m_impl_offset;
1379 N[0] = src.m_array_offset.dimension_0();
1380 N[1] = src.m_array_offset.dimension_1();
1381 N[2] = src.m_array_offset.dimension_2();
1382 N[3] = src.m_array_offset.dimension_3();
1383 N[4] = src.m_array_offset.dimension_4();
1384 N[5] = src.m_array_offset.dimension_5();
1385 N[6] = src.m_array_offset.dimension_6();
1386 N[7] = src.m_array_offset.dimension_7();
1387 S[0] = src.m_array_offset.stride_0();
1388 S[1] = src.m_array_offset.stride_1();
1389 S[2] = src.m_array_offset.stride_2();
1390 S[3] = src.m_array_offset.stride_3();
1391 S[4] = src.m_array_offset.stride_4();
1392 S[5] = src.m_array_offset.stride_5();
1393 S[6] = src.m_array_offset.stride_6();
1394 S[7] = src.m_array_offset.stride_7();
1398 if (std::is_same<
typename SrcTraits::array_layout
1399 , Kokkos::LayoutLeft >::value)
1401 const size_t N_fad = N[0];
1402 const size_t S_fad = S[0];
1403 for (
int i=0;
i<7; ++
i) {
1407 N[DstTraits::dimension::rank] = N_fad;
1408 S[DstTraits::dimension::rank] = S_fad;
1410 Kokkos::LayoutStride ls( N[0], S[0],
1418 dst.m_array_offset = dst_offset_type(std::integral_constant<unsigned,0>(), ls);
1426 template<
class DstTraits ,
class SrcTraits >
1427 class ViewMapping< DstTraits , SrcTraits ,
1428 typename std::enable_if<(
1429 std::is_same< typename DstTraits::memory_space
1430 , typename SrcTraits::memory_space >::value
1433 std::is_same< typename DstTraits::specialize , void >::value
1436 std::is_same< typename SrcTraits::specialize
1437 , ViewSpecializeSacadoFadContiguous >::value
1439 , typename DstTraits::specialize
1444 enum { is_assignable =
true };
1445 enum { is_assignable_data_type =
true };
1447 typedef Kokkos::Impl::SharedAllocationTracker TrackType ;
1448 typedef ViewMapping< DstTraits , typename DstTraits::specialize > DstType ;
1449 typedef ViewMapping< SrcTraits , typename SrcTraits::specialize > SrcFadType ;
1454 template <
class DstType,
class SrcFadType,
class Enable =
void >
1455 struct AssignOffset;
1457 template <
class DstType,
class SrcFadType >
1458 struct AssignOffset< DstType, SrcFadType, typename std::enable_if< ((int)DstType::offset_type::dimension_type::rank != (int)SrcFadType::array_offset_type::dimension_type::rank) >::type >
1461 KOKKOS_INLINE_FUNCTION
1462 static void assign( DstType & dst,
const SrcFadType & src )
1464 typedef typename SrcTraits::value_type TraitsValueType;
1471 typedef typename DstType::offset_type::array_layout DstLayoutType;
1473 typedef typename SrcFadType::array_offset_type::dimension_type SrcViewDimension;
1478 static constexpr
bool is_layout_left =
1481 typedef typename std::conditional< is_layout_left,
1482 typename SrcViewDimension:: template prepend< InnerStaticDim+1 >::type,
1483 typename SrcViewDimension:: template append < InnerStaticDim+1 >::type
1484 >::type SrcViewDimensionAppended;
1486 typedef std::integral_constant< unsigned , 0 > padding ;
1488 typedef ViewOffset< SrcViewDimensionAppended, DstLayoutType > TmpOffsetType;
1490 auto src_layout = src.m_array_offset.layout();
1492 if ( is_layout_left ) {
1493 auto prepend_layout = Kokkos::Impl::prependFadToLayout< DstLayoutType >::returnNewLayoutPlusFad(src_layout, InnerStaticDim+1);
1494 TmpOffsetType offset_tmp( padding(), prepend_layout );
1495 dst.m_impl_offset = offset_tmp;
1498 TmpOffsetType offset_tmp( padding(), src_layout );
1499 dst.m_impl_offset = offset_tmp;
1502 Kokkos::abort(
"Sacado error: Applying AssignOffset for case with nested Fads, but without nested Fads - something went wrong");
1507 template <
class DstType,
class SrcFadType >
1508 struct AssignOffset< DstType, SrcFadType, typename std::enable_if< ((int)DstType::offset_type::dimension_type::rank == (int)SrcFadType::array_offset_type::dimension_type::rank) >::type >
1510 KOKKOS_INLINE_FUNCTION
1511 static void assign( DstType & dst,
const SrcFadType & src )
1513 typedef typename DstType::offset_type dst_offset_type ;
1514 dst.m_impl_offset = dst_offset_type( src.m_array_offset );
1518 template<
class DstType >
1519 KOKKOS_INLINE_FUNCTION
static
1520 void assign( DstType & dst
1521 ,
const SrcFadType & src
1527 std::is_same<
typename DstTraits::array_layout
1528 , Kokkos::LayoutLeft >::value ||
1529 std::is_same<
typename DstTraits::array_layout
1530 , Kokkos::LayoutRight >::value ||
1531 std::is_same<
typename DstTraits::array_layout
1532 , Kokkos::LayoutStride >::value
1536 std::is_same<
typename SrcTraits::array_layout
1537 , Kokkos::LayoutLeft >::value ||
1538 std::is_same<
typename SrcTraits::array_layout
1539 , Kokkos::LayoutRight >::value ||
1540 std::is_same<
typename SrcTraits::array_layout
1541 , Kokkos::LayoutStride >::value
1543 ,
"View of FAD requires LayoutLeft, LayoutRight, or LayoutStride" );
1546 std::is_same<
typename DstTraits::array_layout
1547 ,
typename SrcTraits::array_layout >::value ||
1548 std::is_same<
typename DstTraits::array_layout
1549 , Kokkos::LayoutStride >::value ,
1550 "View assignment must have compatible layout" );
1552 if ( src.m_fad_index != 0 || src.m_fad_stride != 1 ) {
1553 Kokkos::abort(
"\n\n ****** Kokkos::View< Sacado::Fad ... > Cannot assign to array with partitioned view ******\n\n");
1556 AssignOffset< DstType, SrcFadType >::assign( dst, src );
1557 dst.m_impl_handle =
reinterpret_cast< typename DstType::handle_type
>(src.m_impl_handle) ;
1571 template<
class LayoutDest,
class LayoutSrc,
int RankDest,
int RankSrc,
int CurrentArg,
class ... SubViewArgs>
1572 struct SubviewLegalArgsCompileTime<Kokkos::LayoutContiguous<LayoutDest>,LayoutSrc,RankDest,RankSrc,CurrentArg,SubViewArgs...> {
1573 enum { value = SubviewLegalArgsCompileTime<LayoutDest,LayoutSrc,RankDest,RankSrc,CurrentArg,SubViewArgs...>::value };
1576 template<
class LayoutDest,
class LayoutSrc,
int RankDest,
int RankSrc,
int CurrentArg,
class ... SubViewArgs>
1577 struct SubviewLegalArgsCompileTime<LayoutDest,Kokkos::LayoutContiguous<LayoutSrc>,RankDest,RankSrc,CurrentArg,SubViewArgs...> {
1578 enum { value = SubviewLegalArgsCompileTime<LayoutDest,LayoutSrc,RankDest,RankSrc,CurrentArg,SubViewArgs...>::value };
1581 template<
class LayoutDest,
class LayoutSrc,
int RankDest,
int RankSrc,
int CurrentArg,
class ... SubViewArgs>
1582 struct SubviewLegalArgsCompileTime<Kokkos::LayoutContiguous<LayoutDest>,
Kokkos::LayoutContiguous<LayoutSrc>,RankDest,RankSrc,CurrentArg,SubViewArgs...> {
1583 enum { value = SubviewLegalArgsCompileTime<LayoutDest,LayoutSrc,RankDest,RankSrc,CurrentArg,SubViewArgs...>::value };
1588 template<
class SrcTraits ,
class Arg0 ,
class ... Args >
1590 < typename std::enable_if<(
1592 std::is_same< typename SrcTraits::specialize
1593 , ViewSpecializeSacadoFadContiguous >::value
1596 std::is_same< typename SrcTraits::array_layout
1597 , Kokkos::LayoutLeft >::value ||
1598 std::is_same< typename SrcTraits::array_layout
1599 , Kokkos::LayoutRight >::value ||
1600 std::is_same< typename SrcTraits::array_layout
1601 , Kokkos::LayoutStride >::value
1603 && !Sacado::Fad::is_fad_partition<Arg0>::value
1611 static_assert( SrcTraits::rank ==
sizeof...(Args)+1 ,
"" );
1625 enum { rank = unsigned(R0) + unsigned(R1) + unsigned(R2) + unsigned(R3)
1626 + unsigned(R4) + unsigned(R5) + unsigned(R6) };
1629 enum { R0_rev = ( 0 == SrcTraits::rank ? RZ : (
1630 1 == SrcTraits::rank ? R0 : (
1631 2 == SrcTraits::rank ? R1 : (
1632 3 == SrcTraits::rank ? R2 : (
1633 4 == SrcTraits::rank ? R3 : (
1634 5 == SrcTraits::rank ? R4 : (
1635 6 == SrcTraits::rank ? R5 : R6 ))))))) };
1638 typedef typename std::conditional<
1650 >::type array_layout ;
1652 typedef typename SrcTraits::value_type fad_type ;
1654 typedef typename std::conditional< rank == 0 , fad_type ,
1655 typename std::conditional< rank == 1 , fad_type * ,
1656 typename std::conditional< rank == 2 , fad_type ** ,
1657 typename std::conditional< rank == 3 , fad_type *** ,
1658 typename std::conditional< rank == 4 , fad_type **** ,
1659 typename std::conditional< rank == 5 , fad_type ***** ,
1660 typename std::conditional< rank == 6 , fad_type ****** ,
1662 >::type >::type >::type >::type >::type >::type >::type
1667 typedef Kokkos::ViewTraits
1670 ,
typename SrcTraits::device_type
1671 ,
typename SrcTraits::memory_traits > traits_type ;
1673 typedef Kokkos::View
1676 ,
typename SrcTraits::device_type
1677 ,
typename SrcTraits::memory_traits > type ;
1680 KOKKOS_INLINE_FUNCTION
1681 static void assign( ViewMapping< traits_type , typename traits_type::specialize > & dst
1682 , ViewMapping< SrcTraits , typename SrcTraits::specialize >
const & src
1683 , Arg0 arg0 , Args ... args )
1685 typedef ViewMapping< traits_type , typename traits_type::specialize > DstType ;
1686 typedef typename DstType::offset_type dst_offset_type ;
1687 typedef typename DstType::array_offset_type dst_array_offset_type ;
1688 typedef typename DstType::handle_type dst_handle_type ;
1692 const SubviewExtents< SrcTraits::rank + 1 , rank + 1 >
1693 array_extents( src.m_array_offset.m_dim , Kokkos::ALL() , arg0 , args... );
1694 offset = src.m_array_offset( array_extents.domain_offset(0)
1695 , array_extents.domain_offset(1)
1696 , array_extents.domain_offset(2)
1697 , array_extents.domain_offset(3)
1698 , array_extents.domain_offset(4)
1699 , array_extents.domain_offset(5)
1700 , array_extents.domain_offset(6)
1701 , array_extents.domain_offset(7) );
1702 dst_array_offset_type dst_array_offset( src.m_array_offset ,
1710 Kokkos::LayoutStride ls(
1711 dst_array_offset.m_dim.N0, dst_array_offset.m_stride.S0,
1712 dst_array_offset.m_dim.N1, dst_array_offset.m_stride.S1,
1713 dst_array_offset.m_dim.N2, dst_array_offset.m_stride.S2,
1714 dst_array_offset.m_dim.N3, dst_array_offset.m_stride.S3,
1715 dst_array_offset.m_dim.N4, dst_array_offset.m_stride.S4,
1716 dst_array_offset.m_dim.N5, dst_array_offset.m_stride.S5,
1717 dst_array_offset.m_dim.N6, dst_array_offset.m_stride.S6,
1718 dst_array_offset.m_dim.N7, dst_array_offset.m_stride.S7);
1719 auto t1 = ls.dimension[0];
1720 for (
unsigned i=0;
i<rank; ++
i)
1721 ls.dimension[
i] = ls.dimension[
i+1];
1722 ls.dimension[rank] = t1;
1723 auto t2 = ls.stride[0];
1724 for (
unsigned i=0;
i<rank; ++
i)
1725 ls.stride[
i] = ls.stride[
i+1];
1726 ls.stride[rank] = t2;
1727 dst.m_array_offset = dst_array_offset_type(std::integral_constant<unsigned, 0>(), ls);
1730 dst.m_array_offset = dst_array_offset;
1733 const SubviewExtents< SrcTraits::rank + 1 , rank + 1 >
1734 array_extents( src.m_array_offset.m_dim , arg0 , args... , Kokkos::ALL() );
1735 offset = src.m_array_offset( array_extents.domain_offset(0)
1736 , array_extents.domain_offset(1)
1737 , array_extents.domain_offset(2)
1738 , array_extents.domain_offset(3)
1739 , array_extents.domain_offset(4)
1740 , array_extents.domain_offset(5)
1741 , array_extents.domain_offset(6)
1742 , array_extents.domain_offset(7) );
1743 dst.m_array_offset = dst_array_offset_type( src.m_array_offset ,
1747 const SubviewExtents< SrcTraits::rank , rank >
1748 extents( src.m_impl_offset.m_dim , arg0 , args... );
1750 dst.m_impl_offset = dst_offset_type( src.m_impl_offset , extents );
1751 dst.m_impl_handle = dst_handle_type( src.m_impl_handle + offset );
1752 dst.m_fad_size = src.m_fad_size;
1753 dst.m_original_fad_size = src.m_original_fad_size;
1754 dst.m_fad_stride = src.m_fad_stride;
1755 dst.m_fad_index = src.m_fad_index;
1770 template<
class DataType,
class ...P,
unsigned Stride >
1773 ViewTraits<DataType,P...> ,
1774 Sacado::Fad::Partition<Stride>
1779 enum { is_assignable =
true };
1780 enum { is_assignable_data_type =
true };
1782 typedef ViewTraits<DataType,P...> src_traits;
1783 typedef ViewMapping< src_traits , typename src_traits::specialize > src_type ;
1785 typedef typename src_type::offset_type::dimension_type src_dimension;
1786 typedef typename src_traits::value_type fad_type;
1787 typedef typename Sacado::LocalScalarType<fad_type,Stride>::type strided_fad_type;
1789 ViewDataType< strided_fad_type , src_dimension >::type strided_data_type;
1790 typedef ViewTraits<strided_data_type,P...> dst_traits;
1791 typedef View<strided_data_type,P...> type;
1792 typedef ViewMapping< dst_traits , typename dst_traits::specialize > dst_type ;
1794 KOKKOS_INLINE_FUNCTION
static
1795 void assign( dst_type & dst
1796 ,
const src_type & src
1797 ,
const Sacado::Fad::Partition<Stride> & part )
1799 if ( Stride != part.stride && Stride != 0 ) {
1800 Kokkos::abort(
"\n\n ****** Kokkos::View< Sacado::Fad ... > Invalid size in partitioned view assignment ******\n\n");
1802 if ( src.m_fad_stride != 1 ) {
1803 Kokkos::abort(
"\n\n ****** Kokkos::View< Sacado::Fad ... > Can't partition already partitioned view ******\n\n");
1806 dst.m_impl_handle = src.m_impl_handle ;
1807 dst.m_impl_offset = src.m_impl_offset ;
1808 dst.m_array_offset = src.m_array_offset ;
1813 (src.m_fad_size.value + part.stride-part.offset-1) / part.stride ;
1815 dst.m_original_fad_size = src.m_original_fad_size ;
1816 dst.m_fad_stride = part.stride ;
1817 dst.m_fad_index = part.offset ;
1824 #endif // defined(HAVE_SACADO_VIEW_SPEC) && !defined(SACADO_DISABLE_FAD_VIEW_SPEC)
1826 #endif // defined(HAVE_SACADO_KOKKOS)
Base template specification for whether a type is a Fad type.
GeneralFad< StaticStorage< T, Num > > SLFad
Base template specification for static size.
GeneralFad< DynamicStorage< T > > DFad
GeneralFad< StaticFixedStorage< T, Num > > SFad
Base template specification for testing whether type is statically sized.
Get view type for any Fad type.