vector_dist_operators_apply_kernel.hpp

/*
 * vector_dist_operators_apply_kernel.hpp
 *
 *  Created on: Jun 19, 2016
 *      Author: i-bird
 */
 
#ifndef OPENFPM_NUMERICS_SRC_OPERATORS_VECTOR_VECTOR_DIST_OPERATORS_APPLY_KERNEL_HPP_
#define OPENFPM_NUMERICS_SRC_OPERATORS_VECTOR_VECTOR_DIST_OPERATORS_APPLY_KERNEL_HPP_
 
 
#define DEFINE_INTERACTION_3D(name) struct name \
{\
\
    Point<3,double> value(const Point<3,double> & xp, const Point<3,double> xq)\
    {
 
#define END_INTERACTION }\
                        };
 
template<typename ObjType, typename Sfinae = void>
struct is_expression: std::false_type {};
 
template<typename ObjType>
struct is_expression<ObjType, typename Void<typename ObjType::is_expression>::type> : std::true_type
{};
 
template<typename exp, bool is_exp = is_expression<exp>::value>
struct apply_kernel_rtype
{
    typedef typename exp::return_type rtype;
};
 
template<typename exp>
struct apply_kernel_rtype<exp,false>
{
    typedef exp rtype;
};
 
 
template<typename rtype>
struct set_zero
{
    __device__ __host__ static rtype create()
    {
        return 0.0;
    }
};
 
template<unsigned int dim, typename T>
struct set_zero<Point<dim,T>>
{
    __device__ __host__ static Point<dim,T> create()
    {
        Point<dim,T> ret;
 
        ret.zero();
        return ret;
    }
};
 
template<typename T, typename vector, typename exp,typename NN_type, typename Kernel, typename rtype, bool is_exp=is_expression<T>::value>
struct apply_kernel_is_number_or_expression
{
    inline __host__ __device__ static typename std::remove_reference<rtype>::type apply(const vector & vd, NN_type & cl, const exp & v_exp, const vect_dist_key_dx & key, Kernel & lker)
    {
        // accumulator
        typename std::remove_reference<rtype>::type pse = set_zero<typename std::remove_reference<rtype>::type>::create();
 
        // position of particle p
        Point<vector::dims,typename vector::stype> p = vd.getPos(key);
 
        // property of the particle x
        rtype prp_p = v_exp.value(key);
 
        // Get the neighborhood of the particle
        auto NN = cl.getNNIteratorBox(cl.getCell(p));
        while(NN.isNext())
        {
            auto nnp = NN.get();
 
            // Calculate contribution given by the kernel value at position p,
            // given by the Near particle, exclude itself
            if (nnp != key.getKey())
            {
                vect_dist_key_dx nnp_k;
                nnp_k.setKey(nnp);
 
                // property of the particle x
                rtype prp_q = v_exp.value(nnp_k);
 
                // position of the particle q
                Point<vector::dims,typename vector::stype> q = vd.getPos(nnp);
 
                pse += lker.value(p,q,prp_p,prp_q);
            }
 
            // Next particle
            ++NN;
        }
 
        return pse;
    }
};
 
 
template<typename vector, typename exp,typename NN_type, typename Kernel, typename rtype>
struct apply_kernel_is_number_or_expression_sim
{
    __device__ __host__ inline static typename std::remove_reference<rtype>::type apply(const vector & vd, NN_type & cl, const vect_dist_key_dx & key, Kernel & lker)
    {
        // accumulator
        typename std::remove_reference<rtype>::type pse = set_zero<typename std::remove_reference<rtype>::type>::create();
 
        // position of particle p
        Point<vector::dims,typename vector::stype> p = vd.getPos(key);
 
        // Get the neighborhood of the particle
        auto NN = cl.getNNIteratorBox(cl.getCell(p));
        while(NN.isNext())
        {
            auto nnp = NN.get();
 
            // Calculate contribution given by the kernel value at position p,
            // given by the Near particle, exclude itself
            if (nnp != key.getKey())
            {
                // position of the particle q
                Point<vector::dims,typename vector::stype> q = vd.getPos(nnp);
 
                pse += lker.value(p,q);
            }
 
            // Next particle
            ++NN;
        }
 
        return pse;
    }
};
 
 
 
template<typename T, typename vector, typename exp,typename NN_type, typename Kernel, typename rtype, bool is_exp=is_expression<T>::value>
struct apply_kernel_is_number_or_expression_gen
{
    __device__ __host__ inline static typename std::remove_reference<rtype>::type apply(const vector & vd, NN_type & cl, const exp & v_exp, const vect_dist_key_dx & key, Kernel & lker)
    {
        // accumulator
        typename std::remove_reference<rtype>::type pse = set_zero<typename std::remove_reference<rtype>::type>::create();
 
        // property of the particle x
        rtype prp_p = v_exp.value(key);
 
        // position of particle p
        Point<vector::dims,typename vector::stype> p = vd.getPos(key);
 
        // Get the neighborhood of the particle
        auto NN = cl.getNNIteratorBox(cl.getCell(p));
        while(NN.isNext())
        {
            auto nnp = NN.get();
 
            // Calculate contribution given by the kernel value at position p,
            // given by the Near particle, exclude itself
            if (nnp != key.getKey())
            {
                vect_dist_key_dx nnp_k;
                nnp_k.setKey(nnp);
 
                // property of the particle x
                rtype prp_q = v_exp.value(nnp_k);
 
                pse += lker.value(key.getKey(),nnp,prp_p,prp_q,vd);
            }
 
            // Next particle
            ++NN;
        }
 
        return pse;
    }
};
 
template<typename T, bool mm>
struct mutable_or_not
{
    T obj;
 
    mutable_or_not(const T & obj)
    :obj(obj)
    {}
};
 
template<typename T>
struct mutable_or_not<T,true>
{
    mutable T obj;
 
    mutable_or_not(const T & obj)
    :obj(obj)
    {}
};
 
template<typename T, bool is_reference = std::is_reference<T>::value>
struct add_const_reference
{
    typedef const T type;
};
 
template<typename T>
struct add_const_reference<T,true>
{
    typedef typename std::remove_reference<T>::type const & type;
};
 
template <typename exp1,typename vector_type>
class vector_dist_expression_op<exp1,vector_type,VECT_APPLYKER_IN>
{
    typedef typename boost::mpl::at<vector_type,boost::mpl::int_<0>>::type NN;
    typedef typename boost::mpl::at<vector_type,boost::mpl::int_<1>>::type Kernel;
    typedef typename boost::mpl::at<vector_type,boost::mpl::int_<2>>::type vector_orig;
 
    typedef typename std::remove_reference<NN>::type NN_nr;
    typedef typename std::remove_reference<Kernel>::type Kernel_nr;
    typedef typename std::remove_reference<vector_orig>::type vector_orig_nr;
 
    const exp1 o1;
 
    mutable_or_not<NN,is_gpu_celllist<NN>::type::value || is_gpu_ker_celllist<NN>::type::value> cl;
 
    mutable_or_not<Kernel,is_gpu_celllist<NN>::type::value || is_gpu_ker_celllist<NN>::type::value> ker;
 
    typename add_const_reference<vector_orig>::type vd;
 
    typedef typename apply_kernel_rtype<decltype(o1.value(vect_dist_key_dx()))>::rtype rtype;
 
public:
 
    typedef typename has_vector_kernel<vector_orig_nr>::type is_ker;
 
    typedef vector_orig_nr vtype;
 
    typedef NN_nr NN_type;
 
    inline NN_nr * getNN() const
    {
        return &cl.obj;
    }
 
    inline const vtype & getVector() const
    {
        return vd;
    }
 
    inline void init() const
    {
        o1.init();
    }
 
    vector_dist_expression_op(const exp1 & o1, const NN& cl, Kernel & ker, const vector_orig_nr & vd)
    :o1(o1),cl(cl),ker(ker),vd(vd)
    {}
 
    __device__ __host__ inline typename std::remove_reference<rtype>::type value(const vect_dist_key_dx & key) const
    {
        return apply_kernel_is_number_or_expression<decltype(o1.value(key)),
                                                    vector_orig_nr,
                                                    exp1,
                                                    NN_nr,
                                                    Kernel_nr,
                                                    rtype>::apply(vd,cl.obj,o1,key,ker.obj);
    }
};
 
template <typename exp1,typename vector_type>
class vector_dist_expression_op<exp1,vector_type,VECT_APPLYKER_IN_SIM>
{
    typedef typename boost::mpl::at<vector_type,boost::mpl::int_<0>>::type NN;
    typedef typename boost::mpl::at<vector_type,boost::mpl::int_<1>>::type Kernel;
    typedef typename boost::mpl::at<vector_type,boost::mpl::int_<2>>::type vector_orig;
 
    typedef typename std::remove_reference<NN>::type NN_nr;
    typedef typename std::remove_reference<Kernel>::type Kernel_nr;
    typedef typename std::remove_reference<vector_orig>::type vector_orig_nr;
 
    mutable_or_not<NN,is_gpu_celllist<NN>::type::value || is_gpu_ker_celllist<NN>::type::value> cl;
 
    mutable_or_not<Kernel,is_gpu_celllist<NN>::type::value || is_gpu_ker_celllist<NN>::type::value> ker;
 
    const vector_orig vd;
 
    typedef typename apply_kernel_rtype<decltype(std::declval<Kernel_nr>().value(Point<vector_orig_nr::dims,typename vector_orig_nr::stype>(0.0), Point<vector_orig_nr::dims,typename vector_orig_nr::stype>(0.0) ) )>::rtype rtype;
 
 
public:
 
    typedef typename has_vector_kernel<vector_orig_nr>::type is_ker;
 
    typedef vector_orig_nr vtype;
 
    typedef NN_nr NN_type;
 
    inline NN_nr * getNN() const
    {
        return &cl.obj;
    }
 
    inline const vtype & getVector() const
    {
        return vd;
    }
 
    inline void init() const
    {
    }
 
    vector_dist_expression_op(const NN & cl, Kernel & ker, const vector_orig & vd)
    :cl(cl),ker(ker),vd(vd)
    {}
 
    __device__ __host__ inline typename std::remove_reference<rtype>::type value(const vect_dist_key_dx & key) const
    {
        return apply_kernel_is_number_or_expression_sim<vector_orig_nr,
                                                        exp1,
                                                        NN_nr,
                                                        Kernel_nr,
                                                        rtype>::apply(vd,cl.obj,key,ker.obj);
    }
};
 
template <typename exp1,typename vector_type>
class vector_dist_expression_op<exp1,vector_type,VECT_APPLYKER_IN_GEN>
{
    typedef typename boost::mpl::at<vector_type,boost::mpl::int_<0>>::type NN;
    typedef typename boost::mpl::at<vector_type,boost::mpl::int_<1>>::type Kernel;
 
    typedef typename boost::mpl::at<vector_type,boost::mpl::int_<2>>::type vector_orig;
 
    typedef typename std::remove_reference<NN>::type NN_nr;
    typedef typename std::remove_reference<Kernel>::type Kernel_nr;
    typedef typename std::remove_reference<vector_orig>::type vector_orig_nr;
 
    const exp1 o1;
 
    mutable_or_not<NN,is_gpu_celllist<NN>::type::value || is_gpu_ker_celllist<NN>::type::value> cl;
 
    mutable_or_not<Kernel,is_gpu_celllist<NN>::type::value || is_gpu_ker_celllist<NN>::type::value> ker;
 
    const vector_orig vd;
 
    typedef typename apply_kernel_rtype<decltype(o1.value(vect_dist_key_dx()))>::rtype rtype;
 
public:
 
    typedef typename has_vector_kernel<vector_orig_nr>::type is_ker;
 
    typedef vector_orig_nr vtype;
 
    typedef NN_nr NN_type;
 
    inline NN_nr * getNN() const
    {
        return &cl.obj;
    }
 
    inline const vtype & getVector() const
    {
        return vd;
    }
 
    inline void init() const
    {
        o1.init();
    }
 
    vector_dist_expression_op(const exp1 & o1, const NN & cl, Kernel & ker, const vector_orig & vd)
    :o1(o1),cl(cl),ker(ker),vd(vd)
    {}
 
    __device__ __host__ inline typename std::remove_reference<rtype>::type value(const vect_dist_key_dx & key) const
    {
        return apply_kernel_is_number_or_expression_gen<
            decltype(o1.value(key)),
            vector_orig_nr,
            exp1,
            NN_nr,
            Kernel_nr,
            rtype
            >::apply(vd,cl.obj,o1,key,ker.obj);
    }
};
 
template<typename cl_type, int impl=2*is_gpu_celllist<cl_type>::type::value + is_gpu_ker_celllist<cl_type>::type::value >
struct cl_selector_impl
{
    typedef cl_type& ctype;
 
    static ctype & get(cl_type & cl)
    {
        return cl;
    }
};
 
template<typename cl_type>
struct cl_selector_impl<cl_type,2>
{
    typedef decltype(std::declval<cl_type>().toKernel()) ctype;
 
    static ctype get(cl_type & cl)
    {
        return cl.toKernel();
    }
};
 
template<typename cl_type>
struct cl_selector_impl<cl_type,1>
{
    typedef cl_type ctype;
 
    static ctype & get(cl_type & cl)
    {
        return cl;
    }
};
 
template<typename kl_type, typename cl_type, int impl=2*is_gpu_celllist<cl_type>::type::value + is_gpu_ker_celllist<cl_type>::type::value >
struct kl_selector_impl
{
    typedef kl_type& ktype;
};
 
template<typename kl_type, typename cl_type>
struct kl_selector_impl<kl_type,cl_type,2>
{
    typedef kl_type ktype;
};
 
template<typename kl_type, typename cl_type>
struct kl_selector_impl<kl_type,cl_type,1>
{
    typedef kl_type ktype;
};
 
template<bool uwk, typename T>
struct unroll_with_to_kernel
{
    typedef T type;
 
    static T & get(T & v)
    {
        return v;
    }
};
 
template<typename T>
struct unroll_with_to_kernel<true,T>
{
    typedef decltype(std::declval<T>().toKernel()) type;
 
    static type get(T & v)
    {
        return v.toKernel();
    }
};
 
template<typename vl_type, typename cl_type, int impl=2*is_gpu_celllist<cl_type>::type::value + is_gpu_ker_celllist<cl_type>::type::value>
struct vl_selector_impl
{
    typedef vl_type& vtype;
 
    static vtype & get(vl_type & v)
    {
        return v;
    }
};
 
template<typename vl_type, typename cl_type>
struct vl_selector_impl<vl_type,cl_type,2>
{
    typedef decltype(std::declval<vl_type>().toKernel()) vtype;
 
    static vtype get(vl_type & v)
    {
        return v.toKernel();
    }
};
 
template<typename vl_type, typename cl_type>
struct vl_selector_impl<vl_type,cl_type,1>
{
    typedef typename unroll_with_to_kernel<has_toKernel<vl_type>::type::value,vl_type>::type vtype;
 
    static vtype get(vl_type & v)
    {
        return unroll_with_to_kernel<has_toKernel<vl_type>::type::value,vl_type>::get(v);
    }
};
 
 
 
/* \brief Apply kernel expression
 *
 * \param va vector expression one
 * \param vb vector expression two
 *
 * \return an object that encapsulate the expression
 *
 */
template<typename exp1, typename exp2, unsigned int op1, typename NN, typename Kernel, typename vector_type>
inline vector_dist_expression_op<vector_dist_expression_op<exp1,exp2,op1>,
    boost::mpl::vector<typename cl_selector_impl<NN>::ctype,
    typename kl_selector_impl<Kernel,NN>::ktype,
    typename vl_selector_impl<vector_type,NN>::vtype>,
    VECT_APPLYKER_IN>
applyKernel_in(const vector_dist_expression_op<exp1,exp2,op1> & va, vector_type & vd, NN & cl, Kernel & ker)
{
    vector_dist_expression_op<
        vector_dist_expression_op<exp1,exp2,op1>,
        boost::mpl::vector<typename cl_selector_impl<NN>::ctype,
        typename kl_selector_impl<Kernel,NN>::ktype,
        typename vl_selector_impl<vector_type,NN>::vtype>,
        VECT_APPLYKER_IN
        > exp_sum(va,
            cl_selector_impl<NN>::get(cl), ker,
            vl_selector_impl<vector_type,NN>::get(vd)
        );
 
    return exp_sum;
}
 
 
 
/* \brief Apply kernel expression
 *
 * \param va vector expression one
 * \param vb vector expression two
 *
 * \return an object that encapsulate the expression
 *
 */
template<typename exp1, typename exp2, unsigned int op1, typename NN, typename Kernel, typename vector_type>
inline vector_dist_expression_op<vector_dist_expression_op<exp1,exp2,op1>,
    boost::mpl::vector<typename cl_selector_impl<NN>::ctype,
    typename kl_selector_impl<Kernel,NN>::ktype,
    typename vl_selector_impl<vector_type,NN>::vtype>,
    VECT_APPLYKER_IN_GEN>
applyKernel_in_gen(const vector_dist_expression_op<exp1,exp2,op1> & va, vector_type & vd, NN & cl, Kernel & ker)
{
    vector_dist_expression_op<
        vector_dist_expression_op<exp1,exp2,op1>,
        boost::mpl::vector<typename cl_selector_impl<NN>::ctype,
        typename kl_selector_impl<Kernel,NN>::ktype,
        typename vl_selector_impl<vector_type,NN>::vtype>,
        VECT_APPLYKER_IN_GEN
        > exp_sum(va,
            cl_selector_impl<NN>::get(cl),
            ker,
            vl_selector_impl<vector_type,NN>::get(vd)
        );
 
    return exp_sum;
}
 
 
/* \brief Apply kernel expression
 *
 * \param va vector expression one
 * \param vb vector expression two
 *
 * \return an object that encapsulate the expression
 *
 */
template<unsigned int prp, typename NN, typename Kernel, typename vector_type, typename vector_type_ker>
inline vector_dist_expression_op<vector_dist_expression<prp,vector_type_ker>,
    boost::mpl::vector<typename cl_selector_impl<NN>::ctype,
    typename kl_selector_impl<Kernel,NN>::ktype,
    typename vl_selector_impl<vector_type,NN>::vtype>,
    VECT_APPLYKER_IN>
applyKernel_in(const vector_dist_expression<prp,vector_type_ker> & va, vector_type & vd, NN & cl, Kernel & ker)
{
    vector_dist_expression_op<
        vector_dist_expression<prp,vector_type_ker>,
        boost::mpl::vector<typename cl_selector_impl<NN>::ctype,
        typename kl_selector_impl<Kernel,NN>::ktype,
        typename vl_selector_impl<vector_type,NN>::vtype>,
        VECT_APPLYKER_IN
        > exp_sum(va,
            cl_selector_impl<NN>::get(cl),
            ker,
            vl_selector_impl<vector_type,NN>::get(vd)
        );
 
    return exp_sum;
}
 
/* \brief Apply kernel expression
 *
 * \param va vector expression one
 * \param vb vector expression two
 *
 * \return an object that encapsulate the expression
 *
 */
template<unsigned int prp, typename NN, typename Kernel, typename vector_type, typename vector_type_ker>
inline vector_dist_expression_op<vector_dist_expression<prp,vector_type_ker>,
    boost::mpl::vector<typename cl_selector_impl<NN>::ctype,
    typename kl_selector_impl<Kernel,NN>::ktype,
    typename vl_selector_impl<vector_type,NN>::vtype>,
    VECT_APPLYKER_IN_GEN>
applyKernel_in_gen(const vector_dist_expression<prp,vector_type_ker> & va, vector_type & vd, NN & cl, Kernel & ker)
{
    vector_dist_expression_op<
        vector_dist_expression<prp,vector_type_ker>,
        boost::mpl::vector<typename cl_selector_impl<NN>::ctype,
        typename kl_selector_impl<Kernel,NN>::ktype,
        typename vl_selector_impl<vector_type,NN>::vtype>,
        VECT_APPLYKER_IN_GEN
        > exp_sum(va,
            cl_selector_impl<NN>::get(cl),
            ker,
            vl_selector_impl<vector_type,NN>::get(vd)
        );
 
    return exp_sum;
}
 
/* \brief Apply kernel expression
 *
 * \param va vector expression one
 * \param vb vector expression two
 *
 * \return an object that encapsulate the expression
 *
 */
template<typename NN, typename Kernel, typename vector_type>
inline vector_dist_expression_op<void,
    boost::mpl::vector<typename cl_selector_impl<NN>::ctype,
    typename kl_selector_impl<Kernel,NN>::ktype,
    typename vl_selector_impl<vector_type,NN>::vtype>,
    VECT_APPLYKER_IN_SIM>
applyKernel_in_sim(vector_type & vd, NN & cl, Kernel & ker)
{
    vector_dist_expression_op<
        void,
        boost::mpl::vector<typename cl_selector_impl<NN>::ctype,
        typename kl_selector_impl<Kernel,NN>::ktype,
        typename vl_selector_impl<vector_type,NN>::vtype>,
        VECT_APPLYKER_IN_SIM
        > exp_sum(cl_selector_impl<NN>::get(cl),
            ker,
            vl_selector_impl<vector_type,NN>::get(vd)
        );
 
    return exp_sum;
}
 
#endif /* OPENFPM_NUMERICS_SRC_OPERATORS_VECTOR_VECTOR_DIST_OPERATORS_APPLY_KERNEL_HPP_ */