diff --git a/core/src/Cabana_AoSoA.hpp b/core/src/Cabana_AoSoA.hpp
index 2387d47cd36caea2a638327255464f23859dcb6e..3da4eecef12539f615a78eeb1adb3d5074728099 100644
--- a/core/src/Cabana_AoSoA.hpp
+++ b/core/src/Cabana_AoSoA.hpp
@@ -2,36 +2,189 @@
#define CABANA_AOSOA_HPP
#include <Cabana_MemberDataTypes.hpp>
-#include <Cabana_MemoryPolicy.hpp>
#include <Cabana_SoA.hpp>
#include <Cabana_Index.hpp>
+#include <Cabana_InnerArraySize.hpp>
+#include <Kokkos_Core.hpp>
+#include <Kokkos_Core_fwd.hpp>
#include <Kokkos_Macros.hpp>
+#include <Kokkos_HostSpace.hpp>
+#include <Kokkos_MemoryTraits.hpp>
+#include <Kokkos_ExecPolicy.hpp>
#include <type_traits>
#include <memory>
#include <cmath>
#include <cstdlib>
+#include <string>
namespace Cabana
{
+//---------------------------------------------------------------------------//
+/*! \class AoSoATraits
+ \brief Traits class for accessing attributes of a AoSoA.
+
+ This is an implementation detail of AoSoA. It is only of interest
+ to developers implementing a new specialization of AoSoA.
+
+ Template argument options:
+ - AoSoA< DataTypes >
+ - AoSoA< DataTypes , Space >
+ - AoSoA< DataTypes , Space , MemoryTraits >
+ - AoSoA< DataTypes , ArraySize >
+ - AoSoA< DataTypes , ArraySize , Space >
+ - AoSoA< DataTypes , ArraySize , MemoryTraits >
+ - AoSoA< DataTypes , ArraySize , Space , MemoryTraits >
+ - AoSoA< DataTypes , MemoryTraits >
+
+ Note that this is effectively a reimplementation of Kokkos::ViewTraits for
+ the AoSoA with ArrayLayout replaced by ArraySize.
+*/
+template<class DataTypes , class ... Properties>
+class AoSoATraits ;
+
+// Void specialization.
+template<>
+class AoSoATraits<void>
+{
+ public:
+ using execution_space = void;
+ using memory_space = void;
+ using host_mirror_space = void;
+ using array_size = void;
+ using memory_traits = void;
+};
+
+// Extract the array size.
+template<class ArraySize, class ... Properties>
+class AoSoATraits<
+ typename std::enable_if<is_inner_array_size<ArraySize>::value>::type,
+ ArraySize, Properties...>
+{
+ public:
+ using execution_space = typename AoSoATraits<void,Properties...>::execution_space;
+ using memory_space = typename AoSoATraits<void,Properties...>::memory_space;
+ using host_mirror_space = typename AoSoATraits<void,Properties...>::host_mirror_space;
+ using array_size = ArraySize;
+ using memory_traits = typename AoSoATraits<void,Properties...>::memory_traits;
+};
+
+// Extract the space - either a Kokkos memory space or execution space. Can be
+// on or the other but not both.
+template<class Space, class ... Properties>
+class AoSoATraits<
+ typename std::enable_if<Kokkos::Impl::is_space<Space>::value>::type,
+ Space, Properties ...>
+{
+ public:
+ static_assert(
+ std::is_same<typename AoSoATraits<void,Properties...>::execution_space,void>::value &&
+ std::is_same<typename AoSoATraits<void,Properties...>::memory_space,void>::value &&
+ std::is_same<typename AoSoATraits<void,Properties...>::host_mirror_space,void>::value &&
+ std::is_same<typename AoSoATraits<void,Properties...>::array_size,void>::value
+ , "Only one AoSoA Execution or Memory Space template argument" );
+
+ using execution_space = typename Space::execution_space;
+ using memory_space = typename Space::memory_space;
+ using host_mirror_space = typename Kokkos::Impl::HostMirror<Space>::Space;
+ using array_size = ExecutionSpaceInnerArraySize<execution_space>;
+ using memory_traits = typename AoSoATraits<void,Properties...>::memory_traits;
+};
+
+// Extract the memory traits - this must be the last template parameter in the pack.
+template<class MemoryTraits, class ... Properties>
+class AoSoATraits<
+ typename std::enable_if<Kokkos::Impl::is_memory_traits<MemoryTraits>::value>::type,
+ MemoryTraits, Properties...>
+{
+ public:
+ static_assert( std::is_same<typename AoSoATraits<void,Properties...>::execution_space,void>::value &&
+ std::is_same<typename AoSoATraits<void,Properties...>::memory_space,void>::value &&
+ std::is_same<typename AoSoATraits<void,Properties...>::array_size,void>::value &&
+ std::is_same<typename AoSoATraits<void,Properties...>::memory_traits,void>::value
+ , "MemoryTrait is the final optional template argument for a AoSoA" );
+
+ using execution_space = void;
+ using memory_space = void;
+ using host_mirror_space = void;
+ using array_size = void;
+ using memory_traits = MemoryTraits;
+};
+
+// Set the traits for a given set of properties.
+template<class DataTypes, class ... Properties>
+class AoSoATraits
+{
+ private:
+
+ typedef AoSoATraits<void,Properties...> properties;
+
+ using ExecutionSpace =
+ typename
+ std::conditional<
+ !std::is_same<typename properties::execution_space,void>::value,
+ typename properties::execution_space,
+ Kokkos::DefaultExecutionSpace
+ >::type;
+
+ using MemorySpace =
+ typename std::conditional<
+ !std::is_same<typename properties::memory_space,void>::value,
+ typename properties::memory_space,
+ typename ExecutionSpace::memory_space
+ >::type;
+
+ using ArraySize =
+ typename std::conditional<
+ !std::is_same<typename properties::array_size,void>::value,
+ typename properties::array_size,
+ ExecutionSpaceInnerArraySize<ExecutionSpace>
+ >::type;
+
+ using HostMirrorSpace =
+ typename std::conditional<
+ !std::is_same<typename properties::host_mirror_space,void>::value,
+ typename properties::host_mirror_space,
+ typename Kokkos::Impl::HostMirror<ExecutionSpace>::Space
+ >::type;
+
+ using MemoryTraits =
+ typename std::conditional<
+ !std::is_same<typename properties::memory_traits,void>::value,
+ typename properties::memory_traits,
+ typename Kokkos::MemoryManaged
+ >::type;
+
+ public:
+
+ using data_types = DataTypes;
+ using execution_space = ExecutionSpace;
+ using memory_space = MemorySpace;
+ using device_type = Kokkos::Device<ExecutionSpace,MemorySpace>;
+ using memory_traits = MemoryTraits;
+ using host_mirror_space = HostMirrorSpace;
+ using size_type = typename memory_space::size_type;
+
+ static constexpr std::size_t array_size = ArraySize::value;
+};
+
//---------------------------------------------------------------------------//
// Forward declaration.
-template<typename DataTypes, typename Device, std::size_t ArraySize>
+template<typename DataTypes, typename ... Properties>
class AoSoA;
//---------------------------------------------------------------------------//
// Static type checker.
-template<typename >
-struct is_aosoa
- : public std::false_type {};
+template<class >
+struct is_aosoa : public std::false_type {};
-template<typename DataTypes, typename Device, std::size_t ArraySize>
-struct is_aosoa<AoSoA<DataTypes,Device,ArraySize> >
+template<class DataTypes, class ... Properties>
+struct is_aosoa<AoSoA<DataTypes,Properties...> >
: public std::true_type {};
-template<typename DataTypes, typename Device, std::size_t ArraySize>
-struct is_aosoa<const AoSoA<DataTypes,Device,ArraySize> >
+template<class DataTypes, class ... Properties>
+struct is_aosoa<const AoSoA<DataTypes,Properties...> >
: public std::true_type {};
//---------------------------------------------------------------------------//
@@ -39,22 +192,19 @@ struct is_aosoa<const AoSoA<DataTypes,Device,ArraySize> >
\class AoSoA
\brief Array-of-Structs-of-Arrays
*/
-template<typename Device, std::size_t ArraySize, typename... Types>
-class AoSoA<MemberDataTypes<Types...>,Device,ArraySize>
+template<class ... Types, class ... Properties>
+class AoSoA<MemberDataTypes<Types...>,Properties...>
{
public:
- // AoSoA type.
- using aosoa_type = AoSoA<MemberDataTypes<Types...>,Device,ArraySize>;
-
- // Device type.
- using device_type = Device;
+ // Traits.
+ using traits = AoSoATraits<MemberDataTypes<Types...>,Properties...>;
- // Memory policy.
- using memory_policy = MemoryPolicy<device_type>;
+ // AoSoA type.
+ using aosoa_type = AoSoA<MemberDataTypes<Types...>,Properties...>;
- // Inner array size (size of the arrays held by the structs).
- static constexpr std::size_t array_size = ArraySize;
+ // Array size.
+ static constexpr std::size_t array_size = traits::array_size;
// SoA type.
using soa_type = SoA<array_size,Types...>;
@@ -173,10 +323,10 @@ class AoSoA<MemberDataTypes<Types...>,Device,ArraySize>
// by inserting or erasing elements from it.
void resize( const std::size_t n )
{
+ reserve( n );
_size = n;
_num_soa = std::floor( n / array_size );
if ( 0 < n % array_size ) ++_num_soa;
- reserve( _size );
}
// Requests that the container capacity be at least enough to contain n
@@ -199,13 +349,16 @@ class AoSoA<MemberDataTypes<Types...>,Device,ArraySize>
if ( 0 < n % array_size ) ++num_soa_alloc;
_capacity = num_soa_alloc * array_size;
- soa_type* data_block;
- memory_policy::allocate( data_block, num_soa_alloc );
- std::shared_ptr<soa_type> sp(
- data_block, memory_policy::template deallocate<soa_type> );
+ std::shared_ptr<void> sp(
+ Kokkos::kokkos_malloc(num_soa_alloc * sizeof(soa_type)),
+ Kokkos::kokkos_free<typename traits::memory_space> );
if ( _managed_data != nullptr )
- memory_policy::copy( data_block, _managed_data.get(), _num_soa );
+ Kokkos::Impl::DeepCopy<
+ typename traits::memory_space,
+ typename traits::memory_space,
+ typename traits::execution_space>(
+ sp.get(), _managed_data.get(), _num_soa * sizeof(soa_type) );
std::swap( _managed_data, sp );
@@ -448,7 +601,8 @@ class AoSoA<MemberDataTypes<Types...>,Device,ArraySize>
{
static_assert( 0 <= N && N < number_of_members,
"Static loop out of bounds!" );
- soa_type* data_block = _managed_data.get();
+ soa_type* data_block =
+ std::static_pointer_cast<soa_type>(_managed_data).get();
_pointers[N] =
static_cast<void*>( getStructMember<N>(data_block[0]) );
static_assert( 0 ==
@@ -542,8 +696,9 @@ class AoSoA<MemberDataTypes<Types...>,Device,ArraySize>
// Structs-of-Arrays managed data. This shared pointer manages the block
// of memory owned by this class such that the copy constructor and
// assignment operator for this class perform a shallow and reference
- // counted copy of the data.
- std::shared_ptr<soa_type> _managed_data;
+ // counted copy of the data. The underlying pointer is to an array of
+ // soa_type objects.
+ std::shared_ptr<void> _managed_data;
// Pointers to the first element of each member.
void* _pointers[number_of_members];
diff --git a/core/src/Cabana_Cuda.hpp b/core/src/Cabana_Cuda.hpp
deleted file mode 100644
index 20633a91587ecab6a20e4ec8a42202a7a437d79f..0000000000000000000000000000000000000000
--- a/core/src/Cabana_Cuda.hpp
+++ /dev/null
@@ -1,96 +0,0 @@
-#ifndef CABANA_CUDA_HPP
-#define CABANA_CUDA_HPP
-
-#if defined( __NVCC__ )
-
-#include <type_traits>
-#include <cstdlib>
-
-#include <cuda.h>
-#include <cuda_runtime.h>
-
-namespace Cabana
-{
-//---------------------------------------------------------------------------//
-// Cuda tag.
-struct Cuda {};
-
-//---------------------------------------------------------------------------//
-/*!
- * \brief Memory policy for Cuda computations.
- */
-template<>
-struct MemoryPolicy<Cuda>
-{
- //! Allocate array of a number of objects of type T. This will only work
- //! if T is of trivial type (trivially copyable and contiguous).
- template<class T>
- static
- typename std::enable_if<std::is_trivial<T>::value,void>::type
- allocate( T*& ptr, const std::size_t num_t )
- {
- cudaMalloc( (void**) &ptr, num_t * sizeof(T) );
- }
-
- //! Dellocate an array.
- template<class T>
- static void deallocate( T* ptr )
- {
- cudaFree( ptr );
- }
-
- //! Copy from one address in the memory space to another in the same
- //! memory space.
- template<class T>
- static void copy( T* destination, const T* source, const std::size_t count )
- {
- cudaMemcpy(
- destination, source, count*sizeof(T), cudaMemcpyDeviceToDevice );
- }
-};
-
-//---------------------------------------------------------------------------//
-// CudaUVM tag.
-struct CudaUVM {};
-
-//---------------------------------------------------------------------------//
-/*!
- * \brief Memory policy for Cuda computations with unified-virtual-memory.
- */
-template<>
-struct MemoryPolicy<CudaUVM>
-{
- //! Allocate array of a number of objects of type T. This will only work
- //! if T is of trivial type (trivially copyable and contiguous).
- template<class T>
- static
- typename std::enable_if<std::is_trivial<T>::value,void>::type
- allocate( T*& ptr, const std::size_t num_t )
- {
- cudaMallocManaged( (void**) &ptr, num_t * sizeof(T) );
- }
-
- //! Dellocate an array.
- template<class T>
- static void deallocate( T* ptr )
- {
- cudaFree( ptr );
- }
-
- //! Copy from one address in the memory space to another in the same
- //! memory space.
- template<class T>
- static void copy( T* destination, const T* source, const std::size_t count )
- {
- cudaMemcpy(
- destination, source, count*sizeof(T), cudaMemcpyDeviceToDevice );
- }
-};
-
-//---------------------------------------------------------------------------//
-
-} // end namespace Cabana
-
-#endif // end defined( __NVCC__ )
-
-#endif // end CABANA_CUDA_HPP
diff --git a/core/src/Cabana_InnerArraySize.hpp b/core/src/Cabana_InnerArraySize.hpp
new file mode 100644
index 0000000000000000000000000000000000000000..7c3e8380912e5491f96f93a48837ae7062aa6770
--- /dev/null
+++ b/core/src/Cabana_InnerArraySize.hpp
@@ -0,0 +1,81 @@
+#ifndef CABANA_INNERARRAYSIZE_HPP
+#define CABANA_INNERARRAYSIZE_HPP
+
+#include <Kokkos_Core.hpp>
+
+#include <type_traits>
+
+namespace Cabana
+{
+//---------------------------------------------------------------------------//
+/*!
+ \class InnerArraySize
+ \brief Static inner array size definition. This is the size of the arrays in
+ the struct-of-arrays.
+*/
+template<std::size_t N>
+class InnerArraySize : public std::integral_constant<std::size_t,N> {};
+
+//---------------------------------------------------------------------------//
+// Static type checker.
+template<class >
+struct is_inner_array_size : public std::false_type {};
+
+template<std::size_t N>
+struct is_inner_array_size<InnerArraySize<N> > : public std::true_type {};
+
+template<std::size_t N>
+struct is_inner_array_size<const InnerArraySize<N> > : public std::true_type {};
+
+//---------------------------------------------------------------------------//
+/*!
+ \class ExecutionSpaceInnerArraySize
+ \brief Inner array sizes specific for execution spaces.
+
+ Default version has an inner array size of 1. Specializations will set this
+ specifically for the given space.
+*/
+template<class ExecutionSpace>
+class ExecutionSpaceInnerArraySize : public InnerArraySize<1> {};
+
+//---------------------------------------------------------------------------//
+// Serial specialization.
+#if defined( KOKKOS_ENABLE_SERIAL )
+
+#endif
+
+//---------------------------------------------------------------------------//
+// OpenMP specialization.
+#if defined( KOKKOS_ENABLE_OPENMP )
+
+#endif
+
+//---------------------------------------------------------------------------//
+// Qthread specialization
+#if defined( KOKKOS_ENABLE_QTHREADS )
+
+#endif
+
+//---------------------------------------------------------------------------//
+// Threads specialization.
+#if defined( KOKKOS_ENABLE_THREADS )
+
+#endif
+
+//---------------------------------------------------------------------------//
+// Cuda specialization.
+#if defined( KOKKOS_ENABLE_CUDA )
+
+#endif
+
+//---------------------------------------------------------------------------//
+// ROCm specialization.
+#if defined( KOKKOS_ENABLE_ROCM )
+
+#endif
+
+//---------------------------------------------------------------------------//
+
+} // end namespace Cabana
+
+#endif // end CABANA_INNERARRAYSIZE_HPP
diff --git a/core/src/Cabana_MemberSlice.hpp b/core/src/Cabana_MemberSlice.hpp
index ff786b80bc457db996c37d6d31c1059f1b0a4b3e..b376c9af65e3edb75430d2cc969f071c9e21c468 100644
--- a/core/src/Cabana_MemberSlice.hpp
+++ b/core/src/Cabana_MemberSlice.hpp
@@ -50,9 +50,6 @@ class MemberSlice
// AoSoA type this slice wraps.
using aosoa_type = AoSoA_t;
- // Device type.
- using device_type = typename aosoa_type::device_type;
-
// Inner array size.
static constexpr std::size_t array_size = aosoa_type::array_size;
diff --git a/core/src/Cabana_MemoryPolicy.hpp b/core/src/Cabana_MemoryPolicy.hpp
deleted file mode 100644
index c8a2eba6afe8adb3c53c178a99046e13ca807d8e..0000000000000000000000000000000000000000
--- a/core/src/Cabana_MemoryPolicy.hpp
+++ /dev/null
@@ -1,14 +0,0 @@
-#ifndef CABANA_MEMORYPOLICY_HPP
-#define CABANA_MEMORYPOLICY_HPP
-
-namespace Cabana
-{
-//---------------------------------------------------------------------------//
-template<typename >
-struct MemoryPolicy;
-
-//---------------------------------------------------------------------------//
-
-} // end namespace Cabana
-
-#endif // CABANA_MEMORYPOLICY_HPP
diff --git a/core/src/Cabana_Serial.hpp b/core/src/Cabana_Serial.hpp
deleted file mode 100644
index 2be1a15cc11ad47a4adaa7eab8225972022f5949..0000000000000000000000000000000000000000
--- a/core/src/Cabana_Serial.hpp
+++ /dev/null
@@ -1,50 +0,0 @@
-#ifndef CABANA_SERIAL_HPP
-#define CABANA_SERIAL_HPP
-
-#include <type_traits>
-#include <cstdlib>
-
-namespace Cabana
-{
-//---------------------------------------------------------------------------//
-// Serial tag.
-struct Serial {};
-
-//---------------------------------------------------------------------------//
-/*!
- * \brief Memory policy for serial computations.
- */
-template<>
-struct MemoryPolicy<Serial>
-{
- //! Allocate array of a number of objects of type T. This will only work
- //! if T is of trivial type (trivially copyable and contiguous).
- template<class T>
- static
- typename std::enable_if<std::is_trivial<T>::value,void>::type
- allocate( T*& ptr, const std::size_t n )
- {
- ptr = (T*) malloc( n * sizeof(T) );
- }
-
- //! Dellocate an array.
- template<class T>
- static void deallocate( T* ptr )
- {
- free( ptr );
- }
-
- //! Copy from one address in the memory space to another in the same
- //! memory space.
- template<class T>
- static void copy( T* destination, const T* source, const std::size_t count )
- {
- std::copy( source, source + count, destination );
- }
-};
-
-//---------------------------------------------------------------------------//
-
-} // end namespace Cabana
-
-#endif // end CABANA_SERIAL_HPP
diff --git a/core/unit_test/tstAoSoA.cpp b/core/unit_test/tstAoSoA.cpp
index 866809637308d01a9b274b7bda8cdd75364de993..4653c34bdadbb3f36eaa532da2ae4c5563675b93 100644
--- a/core/unit_test/tstAoSoA.cpp
+++ b/core/unit_test/tstAoSoA.cpp
@@ -1,5 +1,4 @@
#include <Cabana_AoSoA.hpp>
-#include <Cabana_Serial.hpp>
#include <boost/test/unit_test.hpp>
@@ -48,8 +47,8 @@ void checkDataMembers(
//---------------------------------------------------------------------------//
BOOST_AUTO_TEST_CASE( aosoa_serial_api_test )
{
- // Inner array size.
- const std::size_t array_size = 10;
+ // Manually set the inner array size.
+ using inner_array_size = Cabana::InnerArraySize<10>;
// Data dimensions.
const std::size_t dim_1 = 3;
@@ -67,7 +66,7 @@ BOOST_AUTO_TEST_CASE( aosoa_serial_api_test )
>;
// Declare the AoSoA type.
- using AoSoA_t = Cabana::AoSoA<DataTypes,Cabana::Serial,array_size>;
+ using AoSoA_t = Cabana::AoSoA<DataTypes,inner_array_size,Kokkos::Serial>;
// Make sure that it is actually an AoSoA.
BOOST_TEST( Cabana::is_aosoa<AoSoA_t>::value );
@@ -224,8 +223,8 @@ BOOST_AUTO_TEST_CASE( aosoa_serial_api_test )
//---------------------------------------------------------------------------//
BOOST_AUTO_TEST_CASE( aosoa_raw_data_test )
{
- // Inner array size.
- const std::size_t array_size = 103;
+ // Manually set the inner array size.
+ using inner_array_size = Cabana::InnerArraySize<103>;
// Multi dimensional member sizes.
const std::size_t dim_1 = 3;
@@ -241,7 +240,7 @@ BOOST_AUTO_TEST_CASE( aosoa_raw_data_test )
>;
// Declare the AoSoA type.
- using AoSoA_t = Cabana::AoSoA<DataTypes,Cabana::Serial,array_size>;
+ using AoSoA_t = Cabana::AoSoA<DataTypes,inner_array_size,Kokkos::Serial>;
// Create an AoSoA using the default constructor.
std::size_t num_data = 350;
diff --git a/core/unit_test/tstMemberSlice.cpp b/core/unit_test/tstMemberSlice.cpp
index c87a4f8aa4f85aa42c0c64b572505b052dd4dc22..1550426f13376bd61a8ff22b6378f936b235ee96 100644
--- a/core/unit_test/tstMemberSlice.cpp
+++ b/core/unit_test/tstMemberSlice.cpp
@@ -1,6 +1,5 @@
#include <Cabana_MemberSlice.hpp>
#include <Cabana_AoSoA.hpp>
-#include <Cabana_Serial.hpp>
#include <boost/test/unit_test.hpp>
@@ -49,8 +48,8 @@ void checkDataMembers(
//---------------------------------------------------------------------------//
BOOST_AUTO_TEST_CASE( slice_serial_api_test )
{
- // Inner array size.
- const std::size_t array_size = 10;
+ // Manually set the inner array size.
+ using inner_array_size = Cabana::InnerArraySize<10>;
// Data dimensions.
const std::size_t dim_1 = 3;
@@ -68,7 +67,7 @@ BOOST_AUTO_TEST_CASE( slice_serial_api_test )
>;
// Declare the AoSoA type.
- using AoSoA_t = Cabana::AoSoA<DataTypes,Cabana::Serial,array_size>;
+ using AoSoA_t = Cabana::AoSoA<DataTypes,inner_array_size,Kokkos::Serial>;
// Make sure that it is actually an AoSoA.
BOOST_TEST( Cabana::is_aosoa<AoSoA_t>::value );