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;
        }
 
        __host__ __device__ size_t size_local() const
        {
            return 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);
        }
        template <unsigned int p>
        __device__ __host__ inline auto getProp(unsigned int id) const -> decltype(base.template get<p>(grid_key_dx<1>(0)))
        {
            return this->get<p>(id);
        }
 
 
        template <unsigned int p, typename key_type>
        __device__ __host__ inline auto getProp(key_type id) const -> decltype(base.template get<p>(grid_key_dx<1>(0)))
        {
            return this->get<p>(id.getKey());
        }
 
        __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> getDomainIteratorGPU(size_t n_thr = default_kernel_wg_threads_) const
        {
            return getGPUIterator(n_thr);
        }
 
        //Stub for some expression
        void init() const {}
 
        __host__ __device__ auto value(unsigned int p) -> decltype(base.template get<0>(grid_key_dx<1>(0)))
        {
            return get<0>(p);
        }
        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__ __host__ vector_gpu_ker<T,layout_base> & getVector()
        {
            return *this;
        }
 
        __device__ __host__ const vector_gpu_ker<T,layout_base> & getVector() const
        {
            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();
        }
 
        __host__ __device__ size_t size_local() const
        {
            return 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);
        }
 
        vector_gpu_ker<T,layout_base> & getVector()
        {
         return *this;
        }
 
        const vector_gpu_ker<T,layout_base> & getVector() const
        {
            return *this;
        }
 
        __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_ */