8 #ifndef MEMORY_CONF_HPP_ 9 #define MEMORY_CONF_HPP_ 11 #include "util/variadic_to_vmpl.hpp" 12 #include "t_to_memory_c.hpp" 13 #include "Vector/vect_isel.hpp" 14 #include "Vector/util.hpp" 15 #include "util/tokernel_transformation.hpp" 16 #include "util/hostDevice_util_funcs.hpp" 18 constexpr
int SOA_layout_IA = 2;
19 constexpr
int SOA_layout = 1;
20 constexpr
int AOS_layout = 0;
39 template<
typename Seq>
62 template<
typename Seq>
86 typedef typename inter_memc<typename T::type>::type
type;
91 typedef boost::mpl::int_<SOA_layout_IA> type_value;
102 template<
unsigned int p,
typename data_type,
typename g1_type,
typename key_type>
103 __host__ __device__
static inline auto get(data_type & data_,
const g1_type & g1,
const key_type & v1) -> decltype(boost::fusion::at_c<p>(data_).mem_r.operator[](g1.LinId(v1)))
105 return boost::fusion::at_c<p>(data_).mem_r.operator[](g1.LinId(v1));
117 template<
unsigned int p,
typename data_type,
typename g1_type>
118 __host__ __device__
static inline auto get_lin(data_type & data_,
const g1_type & g1,
size_t lin_id) -> decltype(boost::fusion::at_c<p>(data_).mem_r.operator[](lin_id))
120 return boost::fusion::at_c<p>(data_).mem_r.operator[](lin_id);
132 template<
unsigned int p,
typename data_type,
typename g1_type,
typename key_type>
133 __host__ __device__
static inline auto get_c(
const data_type & data_,
const g1_type & g1,
const key_type & v1) -> decltype(boost::fusion::at_c<p>(data_).mem_r.operator[](g1.LinId(v1)))
135 return boost::fusion::at_c<p>(data_).mem_r.operator[](g1.LinId(v1));
147 template<
unsigned int p,
typename data_type,
typename g1_type>
148 __host__ __device__
static inline auto get_lin_c(
const data_type & data_,
const g1_type & g1,
size_t lin_id) -> decltype(boost::fusion::at_c<p>(data_).mem_r.operator[](lin_id))
150 return boost::fusion::at_c<p>(data_).mem_r.operator[](lin_id);
161 template<
typename S,
typename data_type,
unsigned int ... prp>
166 boost::mpl::for_each_ref< boost::mpl::range_c<
int,0,
sizeof...(prp)> >(dth);
177 template<
typename data_type,
unsigned int ... prp>
182 boost::mpl::for_each_ref< boost::mpl::range_c<
int,0,
sizeof...(prp)> >(dth);
213 template<
typename T,
bool is_agg>
247 typedef int yes_is_tlin;
249 typedef boost::mpl::int_<AOS_layout> type_value;
260 template<
unsigned int p,
typename data_type,
typename g1_type,
typename key_type>
261 __host__ __device__
static inline auto get(data_type & data_,
const g1_type & g1,
const key_type & v1) -> decltype(boost::fusion::at_c<p>(data_.mem_r.operator[](g1.LinId(v1)))) &
263 return boost::fusion::at_c<p>(data_.mem_r.operator[](g1.LinId(v1)));
275 template<
unsigned int p,
typename data_type,
typename g1_type>
276 __host__ __device__
static inline auto get_lin(data_type & data_,
const g1_type & g1,
const size_t lin_id) -> decltype(boost::fusion::at_c<p>(data_.mem_r.operator[](lin_id))) &
278 return boost::fusion::at_c<p>(data_.mem_r.operator[](lin_id));
290 template<
unsigned int p,
typename data_type,
typename g1_type,
typename key_type>
291 __host__ __device__
static inline auto get_c(
const data_type & data_,
const g1_type & g1,
const key_type & v1) -> decltype(boost::fusion::at_c<p>(data_.mem_r.operator[](g1.LinId(v1)))) &
293 return boost::fusion::at_c<p>(data_.mem_r.operator[](g1.LinId(v1)));
305 template<
unsigned int p,
typename data_type,
typename g1_type>
306 __host__ __device__
static inline auto get_lin_c(
const data_type & data_,
const g1_type & g1,
const size_t lin_id) -> decltype(boost::fusion::at_c<p>(data_.mem_r.operator[](lin_id))) &
308 return boost::fusion::at_c<p>(data_.mem_r.operator[](lin_id));
319 template<
typename S,
typename data_type,
unsigned int ... prp>
322 data_.mem->hostToDevice(start*
sizeof(T),(stop+1)*
sizeof(T));
333 template<
typename data_type,
unsigned int ... prp>
336 data_.mem->deviceToHost(start*
sizeof(T),(stop+1)*
sizeof(T));
343 template<
typename T,
typename Sfinae =
void>
361 template<
typename T,
typename Sfinae =
void>
389 template<typename T, unsigned int impl = is_vector<T>::value >
__host__ static __device__ auto get_lin(data_type &data_, const g1_type &g1, const size_t lin_id) -> decltype(boost::fusion::at_c< p >(data_.mem_r.operator[](lin_id))) &
Return a reference to the selected element.
Transform the boost::fusion::vector into memory specification (memory_traits)
__host__ static __device__ auto get_lin_c(const data_type &data_, const g1_type &g1, size_t lin_id) -> decltype(boost::fusion::at_c< p >(data_).mem_r.operator[](lin_id))
Return a reference to the selected element.
Transform the boost::fusion::vector into memory specification (memory_traits)
__host__ static __device__ auto get_c(const data_type &data_, const g1_type &g1, const key_type &v1) -> decltype(boost::fusion::at_c< p >(data_.mem_r.operator[](g1.LinId(v1)))) &
Return a reference to the selected element.
memory_traits_lin_type< T, openfpm::vect_isel< T >::value==OPENFPM_NATIVE >::type type
for each element in the vector interleave memory_c
static void deviceToHost(data_type &data_, size_t start, size_t stop)
Synchronize the memory buffer in the device with the memory in the host.
int yes_is_inte
indicate that it change the memory layout from the original
static void hostToDevice(data_type &data_, size_t start, size_t stop)
Copy the memory from host to device.
__host__ static __device__ auto get_lin_c(const data_type &data_, const g1_type &g1, const size_t lin_id) -> decltype(boost::fusion::at_c< p >(data_.mem_r.operator[](lin_id))) &
Return a reference to the selected element.
__host__ static __device__ auto get_c(const data_type &data_, const g1_type &g1, const key_type &v1) -> decltype(boost::fusion::at_c< p >(data_).mem_r.operator[](g1.LinId(v1)))
Return a reference to the selected element.
Transform the boost::fusion::vector into memory specification (memory_traits)
This class convert a boost::mpl::fusion/vector to a boost::mpl::fusion/vector with memory_c<....
static void hostToDevice(data_type &data_, size_t start, size_t stop)
Synchronize the memory buffer in the device with the memory in the host.
inter_memc< typename T::type >::type type
for each element in the vector interleave memory_c
static void deviceToHost(data_type &data_, size_t start, size_t stop)
Synchronize the memory buffer in the device with the memory in the host.
is_multiple_buffer_each_prp
this class is a functor for "for_each" algorithm
__host__ static __device__ auto get_lin(data_type &data_, const g1_type &g1, size_t lin_id) -> decltype(boost::fusion::at_c< p >(data_).mem_r.operator[](lin_id))
Return a reference to the selected element.
__host__ static __device__ auto get(data_type &data_, const g1_type &g1, const key_type &v1) -> decltype(boost::fusion::at_c< p >(data_.mem_r.operator[](g1.LinId(v1)))) &
Return a reference to the selected element.
KeyT const ValueT ValueT OffsetIteratorT OffsetIteratorT int
[in] The number of segments that comprise the sorting data
This class convert a boost::mpl::fusion/vector to a boost::mpl::fusion/vector with memory_c interleav...
__host__ static __device__ auto get(data_type &data_, const g1_type &g1, const key_type &v1) -> decltype(boost::fusion::at_c< p >(data_).mem_r.operator[](g1.LinId(v1)))
Return a reference to the selected element.
this class is a functor for "for_each" algorithm