vector_dist_kernel.hpp

/*
 * vector_dist_gpu.hpp
 *
 *  Created on: Jul 28, 2018
 *      Author: i-bird
 */
 
#ifndef VECTOR_DIST_GPU_HPP_
#define VECTOR_DIST_GPU_HPP_
 
constexpr unsigned int POS_PROP = (unsigned int)-1;
 
#ifdef CUDA_GPU
 
#define GET_PARTICLE(vd) blockDim.x*blockIdx.x + threadIdx.x; if (blockDim.x*blockIdx.x + threadIdx.x >= static_cast<unsigned int>(vd.size_local())) {return;};
#define GET_PARTICLE_SORT(p,NN) if (blockDim.x*blockIdx.x + threadIdx.x >= NN.getGhostMarker()) {return;}\
    else{p = NN.getDomainSortIds().template get<0>(blockDim.x*blockIdx.x + threadIdx.x);}
 
 
#define GET_PARTICLE_BY_ID(p,ids) if (blockDim.x*blockIdx.x + threadIdx.x >= ids.size()) {return;}\
                              else{p = ids.template get<0>(blockDim.x*blockIdx.x + threadIdx.x);}
 
template<typename vector_dist_ker>
struct check_vector_dist_kernels
{
    const vector_dist_ker & o1;
    const vector_dist_ker & o2;
 
    bool check;
 
    inline check_vector_dist_kernels(const vector_dist_ker & o1, const vector_dist_ker & o2)
    :o1(o1),o2(o2),check(false)
    {};
 
    template<typename T>
    __device__ __host__ inline void operator()(T& t)
    {
        check &= o1.template getPointer<T::value>() == o2.template getPointer<T::value>();
    }
};
 
template<unsigned int dim,
         typename St,
         typename prop,
         template<typename> class layout_base = memory_traits_inte>
class vector_dist_ker
{
    int ghostMarker = 0;
 
    mutable openfpm::vector_gpu_ker<Point<dim,St>,layout_base> v_pos;
 
    mutable openfpm::vector_gpu_ker<typename apply_transform<layout_base,prop>::type,layout_base> v_prp;
 
public:
 
    typedef St stype;
 
    static const unsigned int dims = dim;
 
    typedef int vector_kernel;
 
    typedef int yes_has_check_device_pointer;
 
    typedef prop value_type;
 
    vector_dist_ker()
    :ghostMarker(0)
    {}
 
    vector_dist_ker(int ghostMarker, const openfpm::vector_gpu_ker<Point<dim,St>,layout_base> & v_pos,
                             const openfpm::vector_gpu_ker<typename apply_transform<layout_base,prop>::type,layout_base> & v_prp)
    :ghostMarker(ghostMarker),v_pos(v_pos),v_prp(v_prp)
    {}
 
    __device__ __host__ int size_local() const {return ghostMarker;}
 
    __device__ __host__ int size() const {return v_pos.size();}
 
    __device__ __host__ inline auto getPos(int vec_key) -> decltype(v_pos.template get<0>(vec_key))
    {
        return v_pos.template get<0>(vec_key);
    }
 
    __device__ __host__ inline auto getPos(const vect_dist_key_dx & vec_key) -> decltype(v_pos.template get<0>(vec_key.getKey()))
    {
        return v_pos.template get<0>(vec_key.getKey());
    }
 
    __device__ __host__ inline auto getPos(int vec_key) const -> decltype(v_pos.template get<0>(vec_key))
    {
        return v_pos.template get<0>(vec_key);
    }
 
    __device__ __host__ inline auto getPos(const vect_dist_key_dx & vec_key) const -> decltype(v_pos.template get<0>(vec_key.getKey()))
    {
        return v_pos.template get<0>(vec_key.getKey());
    }
 
    __device__ __host__ inline auto getPosOrig(int vec_key) -> decltype(v_pos.template get<0>(vec_key))
    {
        return v_pos.template get<0>(vec_key);
    }
 
    __device__ __host__ inline auto getPosOrig(const vect_dist_key_dx & vec_key) -> decltype(v_pos.template get<0>(vec_key.getKey()))
    {
        return v_pos.template get<0>(vec_key.getKey());
    }
 
    __device__ __host__ inline auto getPosOrig(int vec_key) const -> decltype(v_pos.template get<0>(vec_key))
    {
        return v_pos.template get<0>(vec_key);
    }
 
    __device__ __host__ inline auto getPosOrig(const vect_dist_key_dx & vec_key) const -> decltype(v_pos.template get<0>(vec_key.getKey()))
    {
        return v_pos.template get<0>(vec_key.getKey());
    }
 
    template<unsigned int id> __device__ __host__ inline auto getProp(int vec_key) -> decltype(v_prp.template get<id>(vec_key))
    {
        return v_prp.template get<id>(vec_key);
    }
 
    template<unsigned int id> __device__ __host__ inline auto getProp(const vect_dist_key_dx & vec_key) -> decltype(v_prp.template get<id>(vec_key.getKey()))
    {
        return v_prp.template get<id>(vec_key.getKey());
    }
 
    template<unsigned int id> __device__  __host__ inline auto getProp(int vec_key) const -> decltype(v_prp.template get<id>(vec_key))
    {
        return v_prp.template get<id>(vec_key);
    }
 
    template<unsigned int id> __device__ __host__  inline auto getProp(const vect_dist_key_dx & vec_key) const -> decltype(v_prp.template get<id>(vec_key.getKey()))
    {
        return v_prp.template get<id>(vec_key.getKey());
    }
 
    __device__ openfpm::vector_gpu_ker<Point<dim,St>,memory_traits_inte> & getPosVector()
    {
        return v_pos;
    }
 
    __device__ openfpm::vector_gpu_ker<prop,memory_traits_inte> & getPropVector()
    {
        return v_prp;
    }
 
    __host__ vector_dist_iterator getDomainIterator() const
    {
        std::cout << __FILE__ << ":" << __LINE__ << " error getDomainIterator used on a vector_dist_ker object is not allowed" << std::endl;
 
        return vector_dist_iterator(0,0);
    }
 
    __host__ ite_gpu<1> getDomainIteratorGPU(size_t n_thr = default_kernel_wg_threads_) const
    {
        return v_pos.getGPUIteratorTo(ghostMarker,n_thr);
    }
 
    __host__ bool operator==(const vector_dist_ker & v)
    {
        if (v.v_pos.template getPointer<0>() != v_pos.template getPointer<0>())
        {return false;}
 
        check_vector_dist_kernels<openfpm::vector_gpu_ker<prop,memory_traits_inte>> cv(this->v_prp,v.v_prp);
 
        cv.check = true;
 
        // Do the same for the properties
        boost::mpl::for_each_ref< boost::mpl::range_c<int,0,prop::max_prop> >(cv);
 
        return cv.check;
    }
 
    __host__ bool isSubset() const
    {
        return false;
    }
 
#ifdef SE_CLASS1
 
        pointer_check check_device_pointer(void * ptr)
        {
            pointer_check pc;
 
            pc.match = false;
 
            // we check the position vector and the property vector
            pc = v_pos.check_device_pointer(ptr);
 
            if (pc.match == true)
            {
                pc.match_str = std::string("Particle index overflow in position (v_pos): ") + "\n" + pc.match_str;
                return pc;
            }
 
            pc = v_prp.check_device_pointer(ptr);
            if (pc.match == true)
            {
                pc.match_str = std::string("Particle index overflow in properties (v_prp): ") + "\n" + pc.match_str;
                return pc;
            }
 
            return pc;
        }
 
#endif
};
 
// This is a tranformation node for vector_distributed for the algorithm toKernel_tranform
template<template <typename> class layout_base, typename T>
struct toKernel_transform<layout_base,T,2>
{
    typedef typename apply_transform<layout_base,typename T::value_type>::type aggr;
 
    typedef vector_dist_ker<T::dims,typename T::stype,aggr,layout_base> type;
};
 
#endif
 
#endif /* VECTOR_DIST_GPU_HPP_ */