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_ */