memory_conf.hpp

/*
 * memory_conf.hpp
 *
 *  Created on: Aug 28, 2014
 *      Author: Pietro Incardona
 */

#ifndef MEMORY_CONF_HPP_
#define MEMORY_CONF_HPP_

#include "util/variadic_to_vmpl.hpp"
#include "t_to_memory_c.hpp"
#include "Vector/vect_isel.hpp"
#include "Vector/util.hpp"
#include "util/tokernel_transformation.hpp"
#include "util/hostDevice_util_funcs.hpp"

constexpr int SOA_layout_IA = 2;
constexpr int SOA_layout = 1;
constexpr int AOS_layout = 0;

template<typename Seq>
struct inter_memc
{
    typedef typename v_transform<t_to_memory_c,Seq>::type type;
};

template<typename Seq>
struct inter_memc_red
{
    typedef typename v_transform<t_to_memory_c_red,Seq>::type type;
};


template<typename T>
struct memory_traits_inte
{
    typedef typename inter_memc<typename T::type>::type type;

    typedef int yes_is_inte;

    typedef boost::mpl::int_<SOA_layout_IA> type_value;

    template<unsigned int p, typename data_type, typename g1_type, typename key_type>
    __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)))
    {
        return boost::fusion::at_c<p>(data_).mem_r.operator[](g1.LinId(v1));
    }

    template<unsigned int p, typename data_type, typename g1_type>
    __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))
    {
        return boost::fusion::at_c<p>(data_).mem_r.operator[](lin_id);
    }

    template<unsigned int p, typename data_type, typename g1_type, typename key_type>
    __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)))
    {
        return boost::fusion::at_c<p>(data_).mem_r.operator[](g1.LinId(v1));
    }

    template<unsigned int p, typename data_type, typename g1_type>
    __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))
    {
        return boost::fusion::at_c<p>(data_).mem_r.operator[](lin_id);
    }

    template<typename S, typename data_type, unsigned int ... prp> 
        static void hostToDevice(data_type & data_, size_t start, size_t stop)
    {
        host_to_device_impl<T,memory_traits_inte,S, prp ...> dth(data_,start,stop);

        boost::mpl::for_each_ref< boost::mpl::range_c<int,0,sizeof...(prp)> >(dth);
    }

    template<typename data_type, unsigned int ... prp> 
        static void deviceToHost(data_type & data_, size_t start, size_t stop)
    {
        device_to_host_start_stop_impl<T, memory_traits_inte, prp ...> dth(data_,start,stop);

        boost::mpl::for_each_ref< boost::mpl::range_c<int,0,sizeof...(prp)> >(dth);
    }
};

/*template<typename T>
struct memory_traits_inte_red
{
    typedef typename inter_memc_red<typename T::type>::type type;

    typedef int yes_is_inte;
};*/

template<typename T, bool is_agg>
struct memory_traits_lin_type
{
    typedef memory_c<typename T::type> type;
};

template<typename T>
struct memory_traits_lin_type<T,false>
{
    typedef void type;
};

template<typename T>
struct memory_traits_lin
{
    typedef typename memory_traits_lin_type<T,openfpm::vect_isel<T>::value == OPENFPM_NATIVE>::type type;

    typedef int yes_is_tlin;

    typedef boost::mpl::int_<AOS_layout> type_value;

    template<unsigned int p, typename data_type, typename g1_type, typename key_type>
    __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)))) &
    {
        return boost::fusion::at_c<p>(data_.mem_r.operator[](g1.LinId(v1)));
    }

    template<unsigned int p, typename data_type, typename g1_type>
    __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))) &
    {
        return boost::fusion::at_c<p>(data_.mem_r.operator[](lin_id));
    }

    template<unsigned int p, typename data_type, typename g1_type, typename key_type>
    __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)))) &
    {
        return boost::fusion::at_c<p>(data_.mem_r.operator[](g1.LinId(v1)));
    }

    template<unsigned int p, typename data_type, typename g1_type>
    __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))) &
    {
        return boost::fusion::at_c<p>(data_.mem_r.operator[](lin_id));
    }

    template<typename S, typename data_type, unsigned int ... prp> 
        static void hostToDevice(data_type & data_, size_t start, size_t stop)
    {
        data_.mem->hostToDevice(start*sizeof(T),(stop+1)*sizeof(T));
    }

    template<typename data_type, unsigned int ... prp> 
        static void deviceToHost(data_type & data_, size_t start, size_t stop)
    {
        data_.mem->deviceToHost(start*sizeof(T),(stop+1)*sizeof(T));
    }
};



template<typename T, typename Sfinae = void>
struct is_layout_mlin: std::false_type {};


template<typename T>
struct is_layout_mlin<T, typename Void< typename T::yes_is_tlin>::type> : std::true_type
{};


template<typename T, typename Sfinae = void>
struct is_layout_inte: std::false_type {};


template<typename T>
struct is_layout_inte<T, typename Void< typename T::yes_is_inte>::type> : std::true_type
{};

template<typename T, unsigned int impl = is_vector<T>::value >
struct is_multiple_buffer_each_prp: std::false_type
{};

template<typename T>
struct is_multiple_buffer_each_prp<T,true>: is_layout_inte<typename T::layout_base_>
{};

#endif