map_vector_cuda_ker.cuh

/*
 * map_vector_cuda.hpp
 *
 *  Created on: Jun 28, 2018
 *      Author: i-bird
 */

#ifndef MAP_VECTOR_CUDA_HPP_
#define MAP_VECTOR_CUDA_HPP_

#ifdef __NVCC__

template<typename vector_src_type, typename vector_dst_type, unsigned int ... args>
__global__ void merge_add_prp_device_impl(vector_src_type v_src, vector_dst_type v_dst, unsigned int old_sz)
{
    int i = threadIdx.x + blockIdx.x * blockDim.x;

    if (i >= v_src.size())
    {return;}

    // write the object in the last element
    object_s_di<decltype(v_src.get(i)),decltype(v_dst.get(old_sz+i)),OBJ_ENCAP,args...>(v_src.get(i),v_dst.get(old_sz+i));
}

template<typename vector_src_type, typename vector_dst_type>
__global__ void copy_two_vectors(vector_src_type v_dst, vector_dst_type v_src)
{
    int i = threadIdx.x + blockIdx.x * blockDim.x;

    if (i >= v_src.size())
    {return;}

    v_dst.get(i) = v_src.get(i);
}


template<template<typename,typename> class op,
         typename vector_src_type,
         typename vector_dst_type,
         typename vector_opart_type,
         unsigned int ... args>
__global__ void merge_add_prp_device_impl_src_dst_opar_offset(vector_src_type v_src, vector_dst_type v_dst, vector_opart_type opart, unsigned int start)
{
    int i = threadIdx.x + blockIdx.x * blockDim.x;

    if (i >= v_src.size())
    {return;}

    // write the object in the last element
    object_s_di_op<op,decltype(v_src.get(0)),decltype(v_dst.get(0)),OBJ_ENCAP,args...>(v_src.get(i),v_dst.get(opart.template get<1>(start + i)));
}

template<template<typename,typename> class op,
         typename vector_src_type,
         typename vector_dst_type,
         typename vector_opart_type,
         unsigned int ... args>
__global__ void merge_add_prp_device_impl_src_offset_dst_opar(vector_src_type v_src, vector_dst_type v_dst, vector_opart_type opart, unsigned int start)
{
    int i = threadIdx.x + blockIdx.x * blockDim.x;

    if (i >= opart.size())
    {return;}

    // write the object in the last element
    object_si_di_op<op,decltype(v_src.get(0)),decltype(v_dst.get(0)),OBJ_ENCAP,args...>(v_src.get(start + i),v_dst.get(opart.template get<0>(i)));
}

#endif


template<int prp>
__device__ void fill_vector_error_array_overflow(const void * sptr,int key)
{
#ifdef CUDA_GPU

    int * ptr = (int *)&global_cuda_error_array[0];

    ptr[0] = 1;
    ptr[1] = ((size_t)sptr) & 0xFFFFFFFF;
    ptr[2] = (((size_t)sptr) & 0xFFFFFFFF00000000) >> 32;
    ptr[3] = prp;
    ptr[4] = 1;

    for (int i = 0 ; i < 1 ; i++)
    {ptr[i+5] = key;}

#ifdef __NVCC__

    ptr[5+1] = blockIdx.x;
    ptr[6+1] = blockIdx.y;
    ptr[7+1] = blockIdx.z;

    ptr[8+1] = blockDim.x;
    ptr[9+1] = blockDim.y;
    ptr[10+1] = blockDim.z;

    ptr[11+1] = threadIdx.x;
    ptr[12+1] = threadIdx.y;
    ptr[13+1] = threadIdx.z;

#endif

#endif
}


namespace openfpm
{

    template<typename T, template <typename> class layout_base>
    struct vector_gpu_ker_ref;

    template<typename T, template <typename> class layout_base>
    struct vector_gpu_ker
    {
        typedef vector_gpu_ker<T,layout_base> self_type;

        typedef typename apply_transform<layout_base,T>::type T_;

        unsigned int v_size;

        mutable grid_gpu_ker<1,T_,layout_base,grid_sm<1,void>> base;

        __device__ __host__ inline bool check_bound(size_t v1) const
        {
            return v1 < size();
        }

    public:

        typedef int yes_i_am_vector;

        typedef typename layout_base<T_>::type layout_type;

        // you can access all the properties of T
        typedef typename grid_base<1,T_,CudaMemory,typename layout_base<T_>::type>::container container;

        typedef T_ value_type;

        typedef int yes_has_check_device_pointer;

        __device__ __host__ unsigned int size() const
        {
            return v_size;
        }

        __device__ __host__ unsigned int capacity() const
        {
            return base.size();
        }

        template <unsigned int p>
        __device__ __host__ inline auto get(unsigned int id) const -> decltype(base.template get<p>(grid_key_dx<1>(0)))
        {
#ifdef SE_CLASS1
            if (check_bound(id) == false)
            {fill_vector_error_array_overflow<p>(this->getPointer<p>(),id);}
#endif
            grid_key_dx<1> key(id);

            return base.template get<p>(key);
        }

        __device__ __host__ inline auto get(unsigned int id) -> decltype(base.get_o(grid_key_dx<1>(id)))
        {
#ifdef SE_CLASS1
            if (check_bound(id) == false)
            {fill_vector_error_array_overflow<-1>(this->template getPointer<0>(),id);}
#endif

            grid_key_dx<1> key(id);

            return base.get_o(key);
        }

        inline __device__ __host__ auto get(unsigned int id) const -> const decltype(base.get_o(grid_key_dx<1>(id)))
        {
#ifdef SE_CLASS1
            if (check_bound(id) == false)
            {fill_vector_error_array_overflow<-1>(this->getPointer<0>(),id);}
#endif

            grid_key_dx<1> key(id);

            return base.get_o(key);
        }

        inline __device__ __host__ auto get_o(unsigned int id) const -> decltype(base.get_o(grid_key_dx<1>(id)))
        {
#ifdef SE_CLASS1
            if (check_bound(id) == false)
            {fill_vector_error_array_overflow<-1>(this->template getPointer<0>(),id);}
#endif

            grid_key_dx<1> key(id);

            return base.get_o(key);
        }

        inline __device__ __host__ auto get_o(unsigned int id) -> decltype(base.get_o(grid_key_dx<1>(id)))
        {
#ifdef SE_CLASS1
            if (check_bound(id) == false)
            {fill_vector_error_array_overflow<-1>(this->template getPointer<0>(),id);}
#endif

            grid_key_dx<1> key(id);

            return base.get_o(key);
        }

        inline auto last() const -> decltype(base.get_o(grid_key_dx<1>(0)))
        {
            grid_key_dx<1> key(size()-1);

            return base.get_o(key);
        }

        template <unsigned int p>
        __device__ __host__ inline auto get(unsigned int id) -> decltype(base.template get<p>(grid_key_dx<1>(0)))
        {
#ifdef SE_CLASS1
            if (check_bound(id) == false)
            {fill_vector_error_array_overflow<p>(this->template getPointer<p>(),id);}
#endif

            grid_key_dx<1> key(id);

            return base.template get<p>(key);
        }

        inline auto last() -> decltype(base.get_o(grid_key_dx<1>(0)))
        {
            grid_key_dx<1> key(size()-1);

            return base.get_o(key);
        }

        vector_gpu_ker()
        :v_size(0)
        {}

        vector_gpu_ker(int v_size, const grid_gpu_ker<1,T_,layout_base,grid_sm<1,void>> & cpy)
        :v_size(v_size),base(cpy)
        {}

        vector_gpu_ker(const vector_gpu_ker_ref<T,layout_base> & vref)
        {
            this->operator=(vref.vref);
        }

        inline void constructor_impl(int v_size, const grid_gpu_ker<1,T_,layout_base,grid_sm<1,void>> & cpy)
        {
            this->v_size = v_size;
            base.constructor_impl(cpy);
        }

        inline void constructor_impl(int v_size, const grid_gpu_ker_ref<1,T_,layout_base,grid_sm<1,void>> & cpy)
        {
            this->v_size = v_size;
            base.constructor_impl(cpy);
        }

        __device__ void set(int id, const container & obj)
        {
#ifdef SE_CLASS1
            if (check_bound(id) == false)
            {fill_vector_error_array_overflow<-1>(this->template getPointer<0>(),id);}
#endif

            base.set(id,obj);
        }

        template<unsigned int p> __device__ __host__ void * getPointer()
        {
            return base.template getPointer<p>();
        }

        template<unsigned int p> __device__ __host__ const void * getPointer() const
        {
            return base.template getPointer<p>();
        }

        template <typename encap_S, unsigned int ...args> void set_o(unsigned int i, const encap_S & obj)
        {
#ifdef SE_CLASS1
            if (check_bound(i) == false)
            {fill_vector_error_array_overflow<-1>(this->template getPointer<0>(),i);}
#endif

            // write the object in the last element
            object_s_di<encap_S,decltype(get(i)),OBJ_ENCAP,args...>(obj,get(i));
        }

        __device__ void set(unsigned int id, const vector_gpu_ker<T_,layout_base> & v, unsigned int src)
        {
#ifdef SE_CLASS1
            if (check_bound(id) == false)
            {fill_vector_error_array_overflow<-1>(this->template getPointer<0>(),id);}
#endif

            base.set(id,v.base,src);
        }

        template<unsigned int ... prp>
        __device__ void set(unsigned int id, const vector_gpu_ker<T_,layout_base> & v, unsigned int src)
        {
#ifdef SE_CLASS1
            if (check_bound(id) == false)
            {fill_vector_error_array_overflow<-1>(this->template getPointer<0>(),id);}
#endif

            base.template set<prp...>(id,v.base,src);
        }

        __host__ ite_gpu<1> getGPUIterator(size_t n_thr = default_kernel_wg_threads_) const
        {
            grid_key_dx<1> start(0);
            grid_key_dx<1> stop(size()-1);

            return base.getGPUIterator(start,stop,n_thr);
        }

        ite_gpu<1> getGPUIteratorTo(size_t stop, size_t n_thr = default_kernel_wg_threads_) const
        {
            grid_key_dx<1> start(0);
            grid_key_dx<1> stop_(stop);

            return base.getGPUIterator(start,stop_,n_thr);
        }

        __host__ vector_gpu_ker<T,layout_base> & operator=(const vector_gpu_ker<T,layout_base> & v)
        {
            v_size = v.v_size;
            base = v.base;

            return *this;
        }

        __device__ grid_gpu_ker<1,T_,layout_base, grid_sm<1,void>> & getBase()
        {
            return base;
        }

        void * internal_get_size_pointer()  {return &v_size;}

        void print_size()
        {
#ifndef DISABLE_ALL_RTTI
            std::cout << "the size of: " << demangle(typeid(self_type).name()) << " is " << sizeof(self_type) << std::endl;
            std::cout << "    " << demangle(typeid(decltype(v_size)).name()) << ":" << sizeof(decltype(v_size)) << std::endl;
            std::cout << "    " << demangle(typeid(decltype(base)).name()) << ":" << sizeof(decltype(base)) << std::endl;
#endif
        }

#ifdef SE_CLASS1

        pointer_check check_device_pointer(void * ptr)
        {
            pointer_check pc;
            pc.match = false;

            check_device_ptr<self_type> ptr_chk(ptr,*this);

            boost::mpl::for_each_ref<boost::mpl::range_c<int,0,T::max_prop>>(ptr_chk);

            if (ptr_chk.result == true)
            {
                pc.match = true;
                pc.match_str += std::string("Property: ") + std::to_string(ptr_chk.prp) + "\n";
            }

            return pc;
        }

#endif
    };

    template<typename T, template <typename> class layout_base>
    struct vector_gpu_ker_ref
    {
        typedef vector_gpu_ker<T,layout_base> self_type;

        typedef typename apply_transform<layout_base,T>::type T_;

        vector_gpu_ker<T,layout_base> & vref;

    public:

        typedef int yes_i_am_vector;

        typedef typename layout_base<T_>::type layout_type;

        typedef typename grid_base<1,T_,CudaMemory,typename layout_base<T_>::type>::container container;

        typedef T_ value_type;

        typedef int yes_has_check_device_pointer;

        __device__ __host__ unsigned int size() const
        {
            return vref.size();
        }

        __device__ __host__ unsigned int capacity() const
        {
            return vref.capacity;
        }

        template <unsigned int p>
        __device__ __host__ inline auto get(unsigned int id) const -> decltype(vref.template get<p>(id))
        {
            return vref.template get<p>(id);
        }

        __device__ __host__ inline auto get(unsigned int id) -> decltype(vref.get(id))
        {
            return vref.get(id);
        }

        inline __device__ __host__ auto get(unsigned int id) const -> decltype(vref.get(id))
        {
            return vref.get(id);
        }

        inline __device__ __host__ auto get_o(unsigned int id) const -> decltype(vref.get_o(id))
        {
            return vref.get_o(id);
        }

        inline __device__ __host__ auto get_o(unsigned int id) -> decltype(vref.get_o(id))
        {
            return vref.get_o(id);
        }

        inline auto last() const -> decltype(vref.last())
        {
            return vref.last();
        }

        template <unsigned int p>
        __device__ __host__ inline auto get(unsigned int id) -> decltype(vref.template get<p>(id))
        {
            return vref.template get<p>(id);
        }

        inline auto last() -> decltype(vref.last())
        {
            return vref.last();
        }

        vector_gpu_ker_ref(vector_gpu_ker<T,layout_base> & vref)
        :vref(vref)
        {}

        __device__ void set(int id, const container & obj)
        {
            vref.set(id,obj);
        }

        template<unsigned int p> __device__ __host__ void * getPointer()
        {
            return vref.template getPointer<p>();
        }

        template<unsigned int p> __device__ __host__ const void * getPointer() const
        {
            return vref.template getPointer<p>();
        }

        template <typename encap_S, unsigned int ...args> void set_o(unsigned int i, const encap_S & obj)
        {
            vref.set(i,obj);
        }

        __device__ void set(unsigned int id, const vector_gpu_ker<T_,layout_base> & v, unsigned int src)
        {
            vref.set(id,v,src);
        }

        template<unsigned int ... prp>
        __device__ void set(unsigned int id, const vector_gpu_ker<T_,layout_base> & v, unsigned int src)
        {
            vref.template set<prp ...>(id,v,src);
        }

        __host__ ite_gpu<1> getGPUIterator(size_t n_thr = default_kernel_wg_threads_) const
        {
            return vref.getGPUIterator(n_thr);
        }

        ite_gpu<1> getGPUIteratorTo(size_t stop, size_t n_thr = default_kernel_wg_threads_) const
        {
            return vref.getGPUItertatorTo(stop,n_thr);
        }

        __host__ vector_gpu_ker_ref<T,layout_base> & operator=(const vector_gpu_ker<T,layout_base> & v)
        {
            vref.operator=(v);
            return this;
        }

        __device__ grid_gpu_ker<1,T_,layout_base, grid_sm<1,void>> & getBase()
        {
            return vref.getBase();
        }

        pointer_check check_device_pointer(void * ptr)
        {
            return vref.check_device_pointer(ptr);
        }

        void * internal_get_size_pointer()  {return &vref.internal_get_size_pointer();}

        void print_size()
        {
            return vref.print_size();
        }
    };

}

#endif /* MAP_VECTOR_CUDA_HPP_ */