MemFast.hpp

/*
 * MemFast.hpp
 *
 *  Created on: Mar 22, 2015
 *      Author: Pietro Incardona
 */

#ifndef MEMFAST_HPP_
#define MEMFAST_HPP_

#include "config.h"
#include "Space/SpaceBox.hpp"
#include "util/mathutil.hpp"
#include "Space/Shape/HyperCube.hpp"
#include "NN/CellList/CellListIterator.hpp"
#include <unordered_map>
#include "util/common.hpp"
#include "Vector/map_vector.hpp"

template <typename Memory, template <typename> class layout_base,typename local_index>
class Mem_fast_ker
{
    local_index slot;

    openfpm::vector_gpu_ker<aggregate<local_index>,layout_base> cl_n;

    typedef openfpm::vector_gpu_ker<aggregate<local_index>,layout_base> base;

    base cl_base;

public:

    typedef local_index local_index_type;

    Mem_fast_ker(local_index slot,openfpm::vector_gpu_ker<aggregate<local_index>,layout_base> cl_n, openfpm::vector_gpu_ker<aggregate<local_index>,layout_base> cl_base)
    :slot(slot),cl_n(cl_n),cl_base(cl_base)
    {}

    inline __device__ int getNelements(int id) const
    {
        return (int)cl_n.template get<0>(id);
    }

    inline __device__ unsigned int get(unsigned int cell, unsigned int ele)
    {
        return cl_base.template get<0>(cell * slot + ele);
    }


    inline unsigned int get(unsigned int cell, unsigned int ele) const
    {
        return cl_base.template get<0>(cell * slot + ele);
    }
};

template <typename Memory = HeapMemory, typename local_index = size_t>
class Mem_fast
{
    local_index slot;

    openfpm::vector<aggregate<local_index>,Memory> cl_n;

    typedef typename openfpm::vector<aggregate<local_index>,Memory> base;

    base cl_base;

    inline void realloc()
    {
        // we do not have enough slots reallocate the basic structure with more
        // slots
        base cl_base_(2*slot * cl_n.size());

        // copy cl_base
        for (size_t i = 0 ; i < cl_n.size() ; i++)
        {
            for (local_index j = 0 ; j < cl_n.template get<0>(i) ; j++)
            {cl_base_.template get<0>(2*i*slot + j) = cl_base.template get<0>(slot * i + j);}
        }

        // Double the number of slots
        slot *= 2;

        // swap the memory
        cl_base.swap(cl_base_);
    }


public:

    typedef Mem_fast_ker<Memory,memory_traits_lin,local_index> toKernel_type;

    typedef local_index local_index_type;

    inline size_t size() const
    {
        return cl_n.size();
    }

    inline void destroy()
    {
        cl_n.swap(openfpm::vector<aggregate<local_index>,Memory>());
        cl_base.swap(base());
    }

    inline void init_to_zero(local_index slot, local_index tot_n_cell)
    {
        this->slot = slot;

        // create the array that store the number of particle on each cell and se it to 0

        cl_n.resize(tot_n_cell);
        cl_n.template fill<0>(0);

        // create the array that store the cell id

        cl_base.resize(tot_n_cell * slot);
    }

    inline void operator=(const Mem_fast<Memory,local_index> & mem)
    {
        slot = mem.slot;

        cl_n = mem.cl_n;
        cl_base = mem.cl_base;
    }

    inline void operator=(Mem_fast<Memory,local_index> && mem)
    {
        this->swap(mem);
    }

    template<typename Memory2>
    inline void copy_general(const Mem_fast<Memory2,local_index> & mem)
    {
        slot = mem.private_get_slot();

        cl_n = mem.private_get_cl_n();
        cl_base = mem.private_get_cl_base();
    }

    inline void addCell(local_index cell_id, local_index ele)
    {
        // Get the number of element the cell is storing

        local_index nl = getNelements(cell_id);

        if (nl + 1 >= slot)
        {
            realloc();
        }

        // we have enough slot to store another neighbor element

        cl_base.template get<0>(slot * cell_id + cl_n.template get<0>(cell_id)) = ele;
        cl_n.template get<0>(cell_id)++;
    }

    inline void add(local_index cell_id, local_index ele)
    {
        // add the element to the cell

        this->addCell(cell_id,ele);
    }

    inline auto get(local_index cell, local_index ele) -> decltype(cl_base.template get<0>(cell * slot + ele)) &
    {
        return cl_base.template get<0>(cell * slot + ele);
    }


    inline auto get(local_index cell, local_index ele) const -> decltype(cl_base.template get<0>(cell * slot + ele)) &
    {
        return cl_base.template get<0>(cell * slot + ele);
    }


    inline void remove(local_index cell, local_index ele)
    {
        cl_n.template get<0>(cell)--;
    }

    inline size_t getNelements(const local_index cell_id) const
    {
        return cl_n.template get<0>(cell_id);
    }


    inline void swap(Mem_fast<Memory,local_index> & mem)
    {
        cl_n.swap(mem.cl_n);
        cl_base.swap(mem.cl_base);

        size_t cl_slot_tmp = mem.slot;
        mem.slot = slot;
        slot = cl_slot_tmp;
    }

    inline void swap(Mem_fast<Memory,local_index> && mem)
    {
        slot = mem.slot;

        cl_n.swap(mem.cl_n);
        cl_base.swap(mem.cl_base);
    }

    inline void clear()
    {
        for (size_t i = 0 ; i < cl_n.size() ; i++)
        {cl_n.template get<0>(i) = 0;}
    }

    inline const local_index & getStartId(local_index cell_id) const
    {
        return cl_base.template get<0>(cell_id*slot);
    }

    inline const local_index & getStopId(local_index cell_id) const
    {
        return cl_base.template get<0>(cell_id*slot+cl_n.template get<0>(cell_id));
    }

    inline const local_index & get_lin(const local_index * part_id) const
    {
        return *part_id;
    }

public:

    typedef local_index loc_index;

    inline Mem_fast(local_index slot)
    :slot(slot)
    {}

    inline void set_slot(local_index slot)
    {
        this->slot = slot;
    }

#ifdef CUDA_GPU

    Mem_fast_ker<Memory,memory_traits_lin,local_index> toKernel()
    {
        Mem_fast_ker<Memory,memory_traits_lin,local_index> mfk(slot,cl_n.toKernel(),cl_base.toKernel());

        return mfk;
    }

    void hostToDevice()
    {
        cl_n.template hostToDevice<0>();
        cl_base.template hostToDevice<0>();
    }

#endif

    const openfpm::vector<aggregate<local_index>,Memory> & private_get_cl_n() const
    {
        return cl_n;
    }

    const int & private_get_slot() const
    {
        return slot;
    }

    const base & private_get_cl_base() const
    {
        return cl_base;
    }

};


#endif /* CELLLISTSTANDARD_HPP_ */