30#ifndef KOKKOS_EXPERIMENTAL_VIEW_SACADO_FAD_CONTIGUOUS_HPP
31#define KOKKOS_EXPERIMENTAL_VIEW_SACADO_FAD_CONTIGUOUS_HPP
34#if defined(HAVE_SACADO_KOKKOSCORE)
42template <
typename ViewType,
typename Enabled =
void>
43struct ThreadLocalScalarType {
44 typedef typename ViewType::non_const_value_type type;
47template <
typename ViewType>
48struct 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>
68 template<
typename iType0 ,
typename iType1 >
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);
293template<
unsigned Stride,
typename D,
typename ... P >
294KOKKOS_INLINE_FUNCTION
296partition(
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;
303 return DstViewType(src, part);
306template <
typename ViewType>
307struct 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)
318template<
class OutputView >
319struct 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)
380template<
class OutputView >
381struct SacadoViewFill<
383 typename
std::enable_if<
384 ( Kokkos::is_view_fad_contiguous<OutputView>::value &&
385 std::is_same<typename OutputView::execution_space, Kokkos::Experimental::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 );
450template<
class ... Args >
451struct is_ViewSpecializeSacadoFadContiguous {
enum {
value =
false }; };
453template<
class D ,
class ... P ,
class ... Args >
454struct 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 ) };
473KOKKOS_INLINE_FUNCTION
474constexpr unsigned computeFadPartitionSize(
unsigned size,
unsigned stride)
477 ((size+stride-1)/stride) == (size/stride) ? ((size+stride-1)/stride) : 0;
483template <
unsigned rank,
unsigned static_dim,
typename Layout>
484KOKKOS_INLINE_FUNCTION
485typename std::enable_if< !std::is_same<Layout, LayoutLeft>::value &&
486 !std::is_same<Layout, LayoutStride>::value,
488create_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] );
502template <
unsigned rank,
unsigned static_dim,
typename Layout>
503KOKKOS_INLINE_FUNCTION
504typename std::enable_if< std::is_same<Layout, LayoutStride>::value, Layout>::type
505create_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>
530KOKKOS_INLINE_FUNCTION
531typename std::enable_if< std::is_same<Layout, LayoutLeft>::value, Layout >::type
532create_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] );
545template <
unsigned Rank,
typename Dimension,
typename Layout>
546KOKKOS_INLINE_FUNCTION
547typename std::enable_if< !std::is_same<Layout, LayoutLeft>::value,
size_t>::type
548getFadDimension(
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() )))))));
561template <
unsigned Rank,
typename Dimension,
typename Layout>
562KOKKOS_INLINE_FUNCTION
563typename std::enable_if< std::is_same<Layout, LayoutLeft>::value,
size_t >::type
564getFadDimension(
const ViewOffset<Dimension,Layout,void>& offset)
566 return offset.dimension_0();
569template<
class Traits >
570class 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;
611 typedef typename std::conditional< std::is_same<typename Traits::execution_space, Kokkos::Cuda>::value, strided_scalar_type, fad_type >::type thread_local_scalar_type;
612#elif defined(KOKKOS_ENABLE_HIP)
613 typedef typename Sacado::LocalScalarType< fad_type, unsigned(PartitionedFadStride) >::type strided_scalar_type;
614 typedef typename std::conditional< std::is_same<typename Traits::execution_space, Kokkos::Experimental::HIP>::value, strided_scalar_type, fad_type >::type thread_local_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 =
633 std::is_same< typename Traits::array_layout, Kokkos::LayoutLeft>::value;
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 };
663 enum { Sacado_Rank = std::is_same< typename Traits::array_layout, Kokkos::LayoutLeft >::value ? 0 : Rank+1 };
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
787 typename std::enable_if< Kokkos::Impl::are_integral<I0>::value &&
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
807 typename std::enable_if< Kokkos::Impl::are_integral<I0>::value &&
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
827 typename std::enable_if< Kokkos::Impl::are_integral<I0,I1>::value &&
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
847 typename std::enable_if< Kokkos::Impl::are_integral<I0,I1>::value &&
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
868 typename std::enable_if< Kokkos::Impl::are_integral<I0,I1,I2>::value &&
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
888 typename std::enable_if< Kokkos::Impl::are_integral<I0,I1,I2>::value &&
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
908 typename std::enable_if< Kokkos::Impl::are_integral<I0,I1,I2,I3>::value &&
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
928 typename std::enable_if< Kokkos::Impl::are_integral<I0,I1,I2,I3>::value &&
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
949 typename std::enable_if< Kokkos::Impl::are_integral<I0,I1,I2,I3,I4>::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
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
971 typename std::enable_if< Kokkos::Impl::are_integral<I0,I1,I2,I3,I4>::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
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
993 typename std::enable_if< Kokkos::Impl::are_integral<I0,I1,I2,I3,I4,I5>::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
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
1015 typename std::enable_if< Kokkos::Impl::are_integral<I0,I1,I2,I3,I4,I5>::value &&
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
1037 typename std::enable_if< Kokkos::Impl::are_integral<I0,I1,I2,I3,I4,I5,I6>::value &&
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
1059 typename std::enable_if< Kokkos::Impl::are_integral<I0,I1,I2,I3,I4,I5,I6>::value &&
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 ) );
1160 getFadDimension<unsigned(Rank)>( m_array_offset );
1161 if (
unsigned(FadStaticDimension) == 0 &&
fad_dim == 0)
1162 Kokkos::abort(
"invalid fad dimension (0) supplied!");
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();
1219template<
class DstTraits ,
class SrcTraits >
1220class 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 ;
1315template<
class DstTraits ,
class SrcTraits >
1316class 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);
1426template<
class DstTraits ,
class SrcTraits >
1427class 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 =
1479 std::is_same< DstLayoutType, Kokkos::LayoutLeft>::value;
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) ;
1571template<
class LayoutDest,
class LayoutSrc,
int RankDest,
int RankSrc,
int CurrentArg,
class ... SubViewArgs>
1572struct SubviewLegalArgsCompileTime<
Kokkos::LayoutContiguous<LayoutDest>,LayoutSrc,RankDest,RankSrc,CurrentArg,SubViewArgs...> {
1573 enum {
value = SubviewLegalArgsCompileTime<LayoutDest,LayoutSrc,RankDest,RankSrc,CurrentArg,SubViewArgs...>
::value };
1576template<
class LayoutDest,
class LayoutSrc,
int RankDest,
int RankSrc,
int CurrentArg,
class ... SubViewArgs>
1577struct SubviewLegalArgsCompileTime<LayoutDest,
Kokkos::LayoutContiguous<LayoutSrc>,RankDest,RankSrc,CurrentArg,SubViewArgs...> {
1578 enum {
value = SubviewLegalArgsCompileTime<LayoutDest,LayoutSrc,RankDest,RankSrc,CurrentArg,SubViewArgs...>
::value };
1581template<
class LayoutDest,
class LayoutSrc,
int RankDest,
int RankSrc,
int CurrentArg,
class ... SubViewArgs>
1582struct SubviewLegalArgsCompileTime<
Kokkos::LayoutContiguous<LayoutDest>,
Kokkos::LayoutContiguous<LayoutSrc>,RankDest,RankSrc,CurrentArg,SubViewArgs...> {
1583 enum {
value = SubviewLegalArgsCompileTime<LayoutDest,LayoutSrc,RankDest,RankSrc,CurrentArg,SubViewArgs...>
::value };
1588template<
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 ,
"" );
1615 , R0 = bool(is_integral_extent<0,Arg0,Args...>::value)
1616 , R1 = bool(is_integral_extent<1,Arg0,Args...>::value)
1617 , R2 = bool(is_integral_extent<2,Arg0,Args...>::value)
1618 , R3 = bool(is_integral_extent<3,Arg0,Args...>::value)
1619 , R4 = bool(is_integral_extent<4,Arg0,Args...>::value)
1620 , R5 = bool(is_integral_extent<5,Arg0,Args...>::value)
1621 , R6 = bool(is_integral_extent<6,Arg0,Args...>::value)
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<
1644 ( rank <= 2 && R0 && std::is_same< typename SrcTraits::array_layout , Kokkos::LayoutLeft >::value )
1648 ( rank <= 2 && R0_rev && std::is_same< typename SrcTraits::array_layout , Kokkos::LayoutRight >::value )
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 ;
1691 if (std::is_same< typename SrcTraits::array_layout, LayoutLeft >::value) {
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.m_array_offset = dst_array_offset_type( src.m_array_offset ,
1706 const SubviewExtents< SrcTraits::rank + 1 , rank + 1 >
1707 array_extents( src.m_array_offset.m_dim , arg0 , args... , Kokkos::ALL() );
1708 offset = src.m_array_offset( array_extents.domain_offset(0)
1709 , array_extents.domain_offset(1)
1710 , array_extents.domain_offset(2)
1711 , array_extents.domain_offset(3)
1712 , array_extents.domain_offset(4)
1713 , array_extents.domain_offset(5)
1714 , array_extents.domain_offset(6)
1715 , array_extents.domain_offset(7) );
1716 dst.m_array_offset = dst_array_offset_type( src.m_array_offset ,
1720 const SubviewExtents< SrcTraits::rank , rank >
1721 extents( src.m_impl_offset.m_dim , arg0 , args... );
1723 dst.m_impl_offset = dst_offset_type( src.m_impl_offset , extents );
1724 dst.m_impl_handle = dst_handle_type( src.m_impl_handle + offset );
1725 dst.m_fad_size = src.m_fad_size;
1726 dst.m_original_fad_size = src.m_original_fad_size;
1727 dst.m_fad_stride = src.m_fad_stride;
1728 dst.m_fad_index = src.m_fad_index;
1743template<
class DataType,
class ...P,
unsigned Stride >
1746 ViewTraits<DataType,P...> ,
1747 Sacado::Fad::Partition<Stride>
1752 enum { is_assignable =
true };
1753 enum { is_assignable_data_type =
true };
1755 typedef ViewTraits<DataType,P...> src_traits;
1756 typedef ViewMapping< src_traits , typename src_traits::specialize > src_type ;
1758 typedef typename src_type::offset_type::dimension_type src_dimension;
1759 typedef typename src_traits::value_type fad_type;
1760 typedef typename Sacado::LocalScalarType<fad_type,Stride>::type strided_fad_type;
1762 ViewDataType< strided_fad_type , src_dimension >::type strided_data_type;
1763 typedef ViewTraits<strided_data_type,P...> dst_traits;
1764 typedef View<strided_data_type,P...> type;
1765 typedef ViewMapping< dst_traits , typename dst_traits::specialize > dst_type ;
1767 KOKKOS_INLINE_FUNCTION
static
1768 void assign( dst_type & dst
1769 ,
const src_type & src
1772 if ( Stride != part.stride && Stride != 0 ) {
1773 Kokkos::abort(
"\n\n ****** Kokkos::View< Sacado::Fad ... > Invalid size in partitioned view assignment ******\n\n");
1775 if ( src.m_fad_stride != 1 ) {
1776 Kokkos::abort(
"\n\n ****** Kokkos::View< Sacado::Fad ... > Can't partition already partitioned view ******\n\n");
1779 dst.m_impl_handle = src.m_impl_handle ;
1780 dst.m_impl_offset = src.m_impl_offset ;
1781 dst.m_array_offset = src.m_array_offset ;
1786 (src.m_fad_size.value + part.stride-part.offset-1) / part.stride ;
1788 dst.m_original_fad_size = src.m_original_fad_size ;
1789 dst.m_fad_stride = part.stride ;
1790 dst.m_fad_index = part.offset ;
Fad specializations for Teuchos::BLAS wrappers.
BLAS(bool use_default_impl=true, bool use_dynamic=true, OrdinalType static_workspace_size=0)
Default constructor.
GeneralFad< DynamicStorage< T > > DFad
GeneralFad< StaticStorage< T, Num > > SLFad
GeneralFad< StaticFixedStorage< T, Num > > SFad
Base template specification for whether a type is a Fad type.
Base template specification for testing whether type is statically sized.
Base template specification for static size.
Get view type for any Fad type.