map_grid.hpp

#ifndef MAP_HPP_
#define MAP_HPP_
 
 
#include "config.h"
#include "util/cuda_launch.hpp"
#include "util/object_util.hpp"
#include "Grid/util.hpp"
#include "Vector/vect_isel.hpp"
#include "Vector/util.hpp"
#include "Vector/map_vector_grow_p.hpp"
#include "memory/ExtPreAlloc.hpp"
#include "util/util_debug.hpp"
#include "util/Pack_stat.hpp"
#include <boost/fusion/include/mpl.hpp>
#include <boost/fusion/sequence/intrinsic/at_c.hpp>
#include <boost/fusion/include/at_c.hpp>
#include <boost/fusion/include/for_each.hpp>
#include <boost/fusion/container/vector.hpp>
#include <boost/fusion/include/vector.hpp>
#include <boost/fusion/container/vector/vector_fwd.hpp>
#include <boost/fusion/include/vector_fwd.hpp>
#include <boost/type_traits.hpp>
#include <boost/fusion/include/for_each.hpp>
#include <boost/mpl/range_c.hpp>
#include <boost/mpl/for_each.hpp>
#include "memory_ly/memory_conf.hpp"
#include "util/copy_compare/meta_copy.hpp"
#include "util/for_each_ref.hpp"
#include "util.hpp"
#include <utility>
#ifdef CUDA_GPU
#include "memory/CudaMemory.cuh"
#endif
#include "grid_sm.hpp"
#include "grid_zm.hpp"
#include "memory_ly/Encap.hpp"
#include "memory_ly/memory_array.hpp"
#include "memory_ly/memory_c.hpp"
#include <vector>
#include "se_grid.hpp"
#include "memory/HeapMemory.hpp"
#include "memory/PtrMemory.hpp"
#include "grid_common.hpp"
#include "util/se_util.hpp"
#include "iterators/grid_key_dx_iterator.hpp"
#include "iterators/grid_key_dx_iterator_sub.hpp"
#include "iterators/grid_key_dx_iterator_sp.hpp"
#include "iterators/grid_key_dx_iterator_sub_bc.hpp"
#include "Packer_Unpacker/Packer_util.hpp"
#include "Packer_Unpacker/has_pack_agg.hpp"
#include "cuda/cuda_grid_gpu_funcs.cuh"
#include "grid_base_implementation.hpp"
#include "util/for_each_ref.hpp"
#include "Geometry/grid_smb.hpp"
#include "Geometry/grid_zmb.hpp"
 
#ifndef CUDA_GPU
typedef HeapMemory CudaMemory;
#endif
 
template<typename SGridGpu>
struct GetSetBlockType
{
    typedef typename SGridGpu::device_grid_type::container type;
};
 
template<unsigned int dim, typename T, typename S=HeapMemory, typename layout = typename memory_traits_lin<T>::type, typename linearizer = grid_sm<dim,void>  >
class grid_base
{
};
 
template<unsigned int dim, typename T, typename S, typename linearizer>
class grid_base<dim,T,S,typename memory_traits_lin<T>::type, linearizer> : public grid_base_impl<dim,T,S, memory_traits_lin,linearizer>
{
    typedef typename apply_transform<memory_traits_lin,T>::type T_;
 
    T background;
 
public:
 
    typedef typename memory_traits_lin<T>::type layout;
 
    // you can access all the properties of T
    typedef typename grid_base_impl<dim,T,S, memory_traits_lin>::container container;
 
    typedef void grow_policy;
 
    typedef grid_key_dx<dim> base_key;
 
    typedef grid_key_dx_iterator_sub<dim> sub_grid_iterator_type;
 
    typedef typename grid_base_impl<dim,T,S, memory_traits_lin>::linearizer_type linearizer_type;
 
    inline grid_base() THROW
    :grid_base_impl<dim,T,S,memory_traits_lin, linearizer>()
    {}
 
    inline grid_base(const grid_base<dim,T,S,typename memory_traits_lin<T>::type> & g) THROW
    :grid_base_impl<dim,T,S,memory_traits_lin, linearizer>(g)
    {
    }
 
    inline grid_base(const size_t & sz) THROW
    :grid_base_impl<dim,T,S,memory_traits_lin,linearizer>(sz)
    {
    }
 
    inline grid_base(const size_t (& sz)[dim]) THROW
    :grid_base_impl<dim,T,S,memory_traits_lin,linearizer>(sz)
    {
    }
 
    template<typename pointers_type, 
             typename headers_type, 
             typename result_type, 
             unsigned int ... prp >
    static void unpack_headers(pointers_type & pointers, headers_type & headers, result_type & result, int n_slot)
    {}
 
    template<unsigned int ... prp, typename S2, typename header_type, typename ite_type, typename context_type>
    void unpack_with_headers(ExtPreAlloc<S2> & mem,
                ite_type & sub_it,
                header_type & headers,
                int ih,
                Unpack_stat & ps,
                context_type &context,
                rem_copy_opt opt = rem_copy_opt::NONE_OPT)
    {}
 
#if defined(__HIP__)
 
    __device__ grid_base<dim,T,S> & operator=(const grid_base<dim,T,S> & g)
    {
        printf("Error grid_base operator= is not defined in device code\n");
 
        return *this;
    }
 
#endif
 
    __host__ grid_base<dim,T,S> & operator=(const grid_base<dim,T,S> & g)
    {
        (static_cast<grid_base_impl<dim,T,S, memory_traits_lin> *>(this))->swap(g.duplicate());
 
        meta_copy<T>::meta_copy_(g.background,background);
 
        return *this;
    }
 
    grid_base<dim,T,S,typename memory_traits_lin<T>::type> & operator=(grid_base<dim,T,S,typename memory_traits_lin<T>::type> && g)
    {
        (static_cast<grid_base_impl<dim,T,S, memory_traits_lin> *>(this))->swap(g);
 
        meta_copy<T>::meta_copy_(g.background,background);
 
        return *this;
    }
 
    static bool noPointers()
    {
        return false;
    }
 
    template<unsigned int id> void * getDeviceBuffer()
    {
        return this->data_.mem->getDevicePointer();
    }
 
    template <typename stencil = no_stencil>
    static grid_key_dx_iterator_sub<dim, stencil> type_of_subiterator()
    {
        return grid_key_dx_iterator_sub<dim, stencil>();
    }
 
    static constexpr bool isCompressed()
    {
        return false;
    }
 
    static grid_key_dx_iterator<dim> type_of_iterator()
    {
        return grid_key_dx_iterator<dim>();
    }
 
    void convert_key(grid_key_dx<dim> & key_out, const grid_key_dx<dim> & key_in) const
    {
        for (size_t i = 0 ; i < dim ; i++)
        {key_out.set_d(i,key_in.get(i));}
    }
 
    T & getBackgroundValue()
    {
        return background;
    }
 
    T & getBackgroundValueAggr()
    {
        return background;
    }
 
    template<unsigned int p>
    void setBackgroundValue(const typename boost::mpl::at<typename T::type,boost::mpl::int_<p>>::type & val)
    {
        meta_copy<typename boost::mpl::at<typename T::type,boost::mpl::int_<p>>::type>::meta_copy_(val,background.template get<p>());
    }
 
 
    grid_base<dim,T,S,typename memory_traits_lin<T>::type> & operator=(const grid_base_impl<dim,T,S, memory_traits_lin> & base)
    {
        grid_base_impl<dim,T,S, memory_traits_inte>::operator=(base);
 
        return *this;
    }
 
    grid_base<dim,T,S,typename memory_traits_lin<T>::type> & operator=(grid_base_impl<dim,T,S, memory_traits_lin> && base)
    {
        grid_base_impl<dim,T,S, memory_traits_lin>::operator=((grid_base_impl<dim,T,S, memory_traits_lin> &&)base);
 
        return *this;
    }
};
 
 
 
template<typename T_type, unsigned int ... prp>
struct switch_copy_host_to_device
{
    typename memory_traits_inte<T_type>::type & dst;
 
    typedef typename to_boost_vmpl<prp...>::type v_prp;
 
    inline switch_copy_host_to_device(typename memory_traits_inte<T_type>::type & dst)
    :dst(dst)
    {
    };
 
 
    template<typename T>
    inline void operator()(T& t) const
    {
        boost::fusion::at_c<boost::mpl::at<v_prp,boost::mpl::int_<T::value>>::type::value>(dst).switchToDevicePtr();
    }
};
 
 
template<typename T_type, template<typename> class layout_base , typename Memory>
struct deconstruct_impl
{
    typename memory_traits_inte<T_type>::type & dst;
 
    inline deconstruct_impl(typename memory_traits_inte<T_type>::type & dst)
    :dst(dst)
    {};
 
 
    template<typename T>
    inline void operator()(T& t) const
    {
        typedef decltype(boost::fusion::at_c<T::value>(dst).mem_r) mem_r_type;
 
        typedef typename boost::mpl::at<typename T_type::type,T>::type type_prp;
 
        typedef typename toKernel_transform<layout_base,typename mem_r_type::value_type>::type kernel_type;
 
        typedef boost::mpl::int_<(is_vector<typename mem_r_type::value_type>::value ||
                                  is_vector_dist<typename mem_r_type::value_type>::value ||
                                  is_gpu_celllist<typename mem_r_type::value_type>::value) + 2*std::is_array<type_prp>::value + std::rank<type_prp>::value> crh_cond;
 
        call_recursive_destructor_if_vector<typename mem_r_type::value_type,
                                             kernel_type,
                                             type_prp,
                                             layout_base,
                                             crh_cond::value>
        ::template destruct<Memory,mem_r_type>(static_cast<Memory *>(boost::fusion::at_c<T::value>(dst).mem),
                                     boost::fusion::at_c<T::value>(dst).mem_r);
    }
};
 
 
 
struct dim3_
{
    unsigned int x;
 
    unsigned int y;
 
    unsigned int z;
};
 
template<unsigned int dim>
struct device_grid
{
    dim3_ threads;
 
    dim3_ grids;
};
 
 
template<unsigned int dim, typename T, typename S, typename linearizer>
class grid_base<dim,T,S,typename memory_traits_inte<T>::type,linearizer> : public grid_base_impl<dim,T,S, memory_traits_inte,linearizer>
{
    typedef typename apply_transform<memory_traits_inte,T>::type T_;
 
    T background;
 
public:
 
    typedef typename memory_traits_inte<T>::type layout;
 
    // you can access all the properties of T
    typedef typename grid_base_impl<dim,T,S, memory_traits_inte,linearizer>::container container;
 
    typedef typename grid_base_impl<dim,T,S, memory_traits_inte,linearizer>::linearizer_type linearizer_type;
 
    inline grid_base() THROW
    :grid_base_impl<dim,T,S,memory_traits_inte,linearizer>()
    {
    }
 
    inline grid_base(const grid_base & g) THROW
    :grid_base_impl<dim,T,S,memory_traits_inte,linearizer>(g)
    {
    }
 
    inline grid_base(grid_base && g) THROW
    :grid_base_impl<dim,T,S,memory_traits_inte,linearizer>(g)
    {
    }
 
    inline grid_base(const size_t & sz) THROW
    :grid_base_impl<dim,T,S,memory_traits_inte,linearizer>(sz)
    {
    }
 
    inline grid_base(const size_t (& sz)[dim]) THROW
    :grid_base_impl<dim,T,S,memory_traits_inte,linearizer>(sz)
    {
    }
 
    static void unpack_headers()
    {}
 
    template<unsigned int id> void fill(unsigned char c)
    {
        boost::fusion::at_c<id>(this->data_).mem->fill(c);
    }
 
    template<unsigned int id> void * getDeviceBuffer()
    {
        return boost::fusion::at_c<id>(this->data_).mem->getDevicePointer();
    }
 
    template <typename stencil = no_stencil>
    static grid_key_dx_iterator_sub<dim, stencil> type_of_subiterator()
    {
        return grid_key_dx_iterator_sub<dim, stencil>();
    }
 
    static constexpr bool isCompressed()
    {
        return false;
    }
 
    static grid_key_dx_iterator<dim> type_of_iterator()
    {
        return grid_key_dx_iterator<dim>();
    }
 
    void convert_key(grid_key_dx<dim> & key_out, const grid_key_dx<dim> & key_in) const
    {
        for (size_t i = 0 ; i < dim ; i++)
        {key_out.set_d(i,key_in.get(i));}
    }
 
    T & getBackgroundValue()
    {
        return background;
    }
 
    T & getBackgroundValueAggr()
    {
        return background;
    }
 
    grid_base<dim,T,S,typename memory_traits_inte<T>::type,linearizer> & operator=(const grid_base_impl<dim,T,S, memory_traits_inte,linearizer> & base)
    {
        grid_base_impl<dim,T,S, memory_traits_inte,linearizer>::operator=(base);
 
        return *this;
    }
 
    grid_base<dim,T,S,typename memory_traits_inte<T>::type,linearizer> & operator=(grid_base_impl<dim,T,S, memory_traits_inte,linearizer> && base)
    {
        grid_base_impl<dim,T,S, memory_traits_inte,linearizer>::operator=(base);
 
        return *this;
    }
 
    ~grid_base()
    {
        deconstruct_impl<T,memory_traits_inte,S> dth(this->data_);
 
        boost::mpl::for_each_ref< boost::mpl::range_c<int,0,T::max_prop> >(dth);
    }
};
 
template <unsigned int dim, typename T, typename linearizer = grid_sm<dim,void> > using grid_gpu = grid_base<dim,T,CudaMemory,typename memory_traits_inte<T>::type,linearizer>;
 
template <unsigned int dim, typename T, typename linearizer = grid_sm<dim,void> > using grid_cpu = grid_base<dim,T,HeapMemory,typename memory_traits_lin<T>::type,linearizer>;
 
 
#endif