memory_c.hpp

/*
 * memory_c.hpp
 *
 *  Created on: Aug 17, 2014
 *      Author: Pietro Incardona
 */

#include <boost/multi_array.hpp>
#include <boost/fusion/mpl.hpp>
#include <boost/fusion/include/mpl.hpp>
#include <boost/mpl/vector.hpp>
#include <array>
#include <boost/mpl/pop_front.hpp>
#include <boost/mpl/push_front.hpp>

//#include "util/boost/boost_multi_array_openfpm.hpp"
#include "util/multi_array_openfpm/multi_array_ref_openfpm.hpp"
#include "util/ct_array.hpp"
#include "memory_array.hpp"
#include "memory/memory.hpp"

#ifndef MEMORY_C_HPP_
#define MEMORY_C_HPP_

#define MEMORY_C_STANDARD 1
#define MEMORY_C_REDUCED  2

template<typename T, unsigned int impl = MEMORY_C_STANDARD, typename D = memory>
class memory_c
{

};

template<typename T, typename D>
class memory_c<T,MEMORY_C_STANDARD,D>
{
    bool manage_memory = true;

    public:

    typedef memory_c<T> type;

    typedef T& reference;

    typedef T vtype;

    memory * mem;

    memory_array<T> mem_r;

    void setMemory(memory & mem)
    {
        if (manage_memory)
        {
            if (this->mem != NULL)
            {
                this->mem->decRef();

                if (this->mem->ref() == 0 && &mem != this->mem)
                    delete(this->mem);
            }
            mem.incRef();
        }
        this->mem = &mem;
    }

    bool bind_ref(const memory_c<T,MEMORY_C_STANDARD,D> & ref)
    {
        mem = ref.mem;

        if (manage_memory)
        {mem->incRef();}

        mem_r = ref.mem_r;

        return true;
    }

    memory& getMemory()
    {
        return *this->mem;
    }

    void switchToDevicePtr()
    {
        if (mem != 0)
        {mem_r.set_pointer(mem->getDevicePointer());}
    }

    const memory& getMemory() const
    {
        return *this->mem;
    }

    bool allocate(const size_t sz, bool skip_initialization = false)
    {
        memory * mem = this->mem;

        mem->resize( sz*sizeof(T) );

        mem_r.initialize(mem->getPointer(),sz,mem->isInitialized() | skip_initialization);

        return true;
    }

    memory_c(bool manage_memory = true)
    :manage_memory(manage_memory),mem(NULL){}

    ~memory_c()
    {
        // deinitialixe mem_r
        if (manage_memory)
        {
            mem_r.deinit();
            if (mem != NULL)
            {
                mem->decRef();

                if (mem->ref() == 0)
                    delete(mem);
            }
        }
    }

    void disable_manage_memory()
    {
        manage_memory = false;
    }

    void swap(memory_c & mem_obj)
    {
        // Save on temporal

        void * mem_tmp = static_cast<void*>(mem);
        mem = mem_obj.mem;
        mem_obj.mem = static_cast<memory*>(mem_tmp);

        mem_obj.mem_r.swap(mem_r);
    }

    template<typename Mem_type>
    __host__ void swap_nomode(memory_c & mem_obj)
    {
        // It would be better a dynamic_cast, unfortunately nvcc
        // does not accept it. It seems that for some reason nvcc want to
        // produce device code out of this method. While it is true
        // that this method is called inside a generic __device__ __host__
        // function, tagging this method only __host__ does not stop
        // nvcc from the intention to produce device code.
        // The workaround is to use static_cast. Another workaround (to test)
        // could be create a duplicate for_each function tagged only __host__ ,
        // but I have to intention to duplicate code

        Mem_type * mem_tmp = static_cast<Mem_type*>(mem);
        mem_tmp->swap(*static_cast<Mem_type*>(mem_obj.mem));

        mem_obj.mem_r.swap(mem_r);
    }

};

template<typename T>
class multi_array
{
    typedef T type;
};


template<typename T, unsigned int N>
struct mult
{
    enum { value = mult<T,N-1>::value * boost::mpl::at<T,boost::mpl::int_<N>>::type::value };
};

template <typename T>
struct mult<T,1>
{
    enum { value = boost::mpl::at<T,boost::mpl::int_<1>>::type::value };
};

template<typename size_type, unsigned int dim>
static inline std::array<size_type ,dim> zero_dims()
{
    std::array<size_type ,dim> dimensions;

    // fill runtime, and the other dimensions
    for (size_t i = 0 ; i < dim ; i++)
    {dimensions[i] = 0;}

    return dimensions;
}

template<unsigned int dim, typename size_type>
struct array_ord
{
};

template<typename size_type>
struct array_ord<4,size_type>
{
    static constexpr size_type data[4] = {0,3,2,1};
};

template<typename size_type>
struct array_ord<3,size_type>
{
    static constexpr size_type data[3] = {0,2,1};
};

template<typename size_type>
struct array_ord<2,size_type>
{
    static constexpr size_type data[2] = {0,1};
};

template<unsigned int dim>
struct array_asc
{
};

template<>
struct array_asc<4>
{
    static constexpr bool data[4] = {true,true,true,true};
};

template<>
struct array_asc<3>
{
    static constexpr bool data[3] = {true,true,true};
};

template<>
struct array_asc<2>
{
    static constexpr bool data[2] = {true,true};
};


template<typename T, typename D>
class memory_c<multi_array<T>, MEMORY_C_STANDARD, D>
{
    bool manage_memory = true;

    template<size_t index,size_t N> struct ascending
    {
        enum { value = true };
    };

    typedef typename boost::mpl::push_front<typename boost::mpl::pop_front<T>::type,boost::mpl::int_<-1>>::type Tv;

    template<int S> using int_ = boost::mpl::int_<S>;

    typedef typename boost::mpl::size<T> size_p;

    template< typename S, unsigned int n> using at = boost::mpl::at<S,boost::mpl::int_<n> >;

    typedef typename at<T,0>::type base;

    typedef typename openfpm::multi_array_ref_openfpm<base,size_p::value,Tv>::size_type size_type;

    public:

    void setMemory(memory & mem)
    {
        if (manage_memory)
        {
            if (this->mem != NULL)
            {
                this->mem->decRef();

                if (this->mem->ref() == 0 && &mem != this->mem)
                    delete(this->mem);
            }
            mem.incRef();
        }
        this->mem = &mem;
    }

    memory& getMemory()
    {
        return *this->mem;
    }

    void switchToDevicePtr()
    {
        if (mem != NULL)
        {mem_r.set_pointer(mem->getDevicePointer());}
    }


    bool bind_ref(const memory_c<multi_array<T>, MEMORY_C_STANDARD, D> & ref)
    {
        mem = ref.mem;
        if (manage_memory)
        {mem->incRef();}

        mem_r = ref.mem_r;

        return true;
    }

    bool allocate(const size_t sz, bool skip_initialization = false)
    {
        memory * mem = this->mem;

        mem->resize( sz*mult<T,size_p::value-1>::value*sizeof(base) );

        openfpm::multi_array_ref_openfpm<base,size_p::value,Tv> tmp(static_cast<base *>(mem->getPointer()),
                                                                                        sz,
                                                                                        openfpm::general_storage_order<size_p::value>(openfpm::ofp_storage_order()));

        mem_r.swap(tmp);

        return true;
    }

    typedef boost::multi_array<base,size_p::value> type;

    D * mem;

    openfpm::multi_array_ref_openfpm<base,boost::mpl::size<T>::value,Tv> mem_r;

    memory_c(bool manage_memory = true)
    :manage_memory(manage_memory),mem(NULL),
     mem_r(static_cast<base *>(NULL),0,openfpm::ofp_storage_order())
    {}

    ~memory_c()
    {
        // Indicate that this memory_c does not manage the memory
        if (manage_memory)
        {
            if (mem != NULL)
            {
                mem->decRef();
                if (mem->ref() == 0)
                    delete(mem);
            }
        }
    }

    void disable_manage_memory()
    {
        manage_memory = false;
    }

    void set_mem(memory & mem)
    {
        this->mem = &mem;
    }

    void swap(memory_c & mem_obj)
    {
        // Save on temporal

        void * mem_tmp = static_cast<void*>(mem);
        mem = mem_obj.mem;
        mem_obj.mem = static_cast<D*>(mem_tmp);

        mem_r.swap(mem_obj.mem_r);
    }


    template<typename Mem_type>
    __host__ void swap_nomode(memory_c & mem_obj)
    {
        // It would be better a dynamic_cast, unfortunately nvcc
        // does not accept it. It seems that for some reason nvcc want to
        // produce device code out of this method. While it is true
        // that this function is called inside a generic __device__ __host__
        // function, tagging this method only __host__ does not stop
        // nvcc from the intention to produce device code.
        // The workaround is to use static_cast. Another workaround (to test)
        // could be create a duplicate for_each function tagged only __host__ ,
        // but I have to intention to duplicate code

        Mem_type * mem_tmp = static_cast<Mem_type*>(mem);
        mem_tmp->swap(*static_cast<Mem_type*>(mem_obj.mem));

        mem_obj.mem_r.swap(mem_r);
    }
};

template<typename T>
struct array_to_vmpl
{
};


template<typename T,size_t N1>
struct array_to_vmpl<T[N1]>
{
    typedef boost::mpl::vector<T,boost::mpl::int_<N1>> prim_vmpl;
};


template<typename T,size_t N1,size_t N2>
struct array_to_vmpl<T[N1][N2]>
{
    typedef boost::mpl::vector<T,boost::mpl::int_<N1>,
                                 boost::mpl::int_<N2>> prim_vmpl;
};

template<typename T>
struct to_memory_multi_array_ref_view
{
    // first we convert the type into a boost vector containing the primitive, followed by array dimension
    typedef typename array_to_vmpl<T>::prim_vmpl prim_vmpl;

    // Than we operate at compile-time the same operation memory_c<multi_array does
    typedef typename boost::mpl::push_front<typename boost::mpl::pop_front<prim_vmpl>::type,boost::mpl::int_<-1>>::type vmpl;


};

#endif