ie_ghost_gpu.cuh 3.32 KB
Newer Older
incardon's avatar
incardon committed
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16
/*
 * ie_ghost_gpu.cuh
 *
 *  Created on: Aug 24, 2018
 *      Author: i-bird
 */

#ifndef IE_GHOST_GPU_CUH_
#define IE_GHOST_GPU_CUH_

#include "data_type/aggregate.hpp"

constexpr unsigned int lc_proc_ = 0;
constexpr unsigned int proc_ = 1;
constexpr unsigned int shift_id_ = 2;

incardon's avatar
incardon committed
17 18
template<unsigned int dim, typename T, typename cell_list_type, typename vb_int_box_type>
__device__ __host__ inline unsigned int ghost_processorID_N_impl(const Point<dim,T> & p, cell_list_type & geo_cell, vb_int_box_type & vb_int_proc)
incardon's avatar
incardon committed
19 20 21 22 23 24 25 26 27 28
{
	unsigned int cell = geo_cell.getCell(p);
	unsigned int sz = geo_cell.getNelements(cell);

	unsigned int n = 0;

	for (int i = 0 ; i < sz ; i++)
	{
		unsigned int bid = geo_cell.get(cell,i);

incardon's avatar
incardon committed
29 30 31 32 33 34 35
		unsigned int sz2 = vb_int_proc.template get<0>(bid).size();

		for (int j = 0 ; j < sz2 ; j++)
		{
			if (Box<dim,T>(vb_int_proc.template get<0>(bid).get(j)).isInsideNP(p) == true)
			{n++;}
		}
incardon's avatar
incardon committed
36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56
	}

	return n;
}

/*! \brief structure that store and compute the internal and external local ghost box. Version usable in kernel
 *
 * \tparam dim is the dimensionality of the physical domain we are going to decompose.
 * \tparam T type of the space we decompose, Real, Integer, Complex ...
 *
 * \see CartDecomposition
 *
 */
template<unsigned int dim, typename T, typename Memory, template<typename> class layout_base>
class ie_ghost_gpu
{

	//! Cell-list that store the geometrical information of the internal ghost boxes
	CellList_cpu_ker<dim,T,Mem_fast_ker<Memory,memory_traits_lin,int>,shift<dim,T>> geo_cell;

	//! internal ghost box
incardon's avatar
incardon committed
57
	openfpm::vector_gpu_ker<aggregate<openfpm::vector_gpu_ker<Box<dim, T>,layout_base>,int>,layout_base> vb_int_proc;
incardon's avatar
incardon committed
58 59 60 61 62

public:


	ie_ghost_gpu(CellList_cpu_ker<dim,T,Mem_fast_ker<Memory,memory_traits_lin,int>,shift<dim,T>> geo_cell,
incardon's avatar
incardon committed
63 64
				 openfpm::vector_gpu_ker<aggregate<openfpm::vector_gpu_ker<Box<dim, T>,layout_base>,int>,layout_base> vb_int_proc)
	:geo_cell(geo_cell),vb_int_proc(vb_int_proc)
incardon's avatar
incardon committed
65 66 67 68 69
	{

	}

	ie_ghost_gpu(const ie_ghost_gpu<dim,T,Memory,layout_base> & ieg)
incardon's avatar
incardon committed
70
	:geo_cell(ieg.geo_cell),vb_int_proc(ieg.vb_int_proc)
incardon's avatar
incardon committed
71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89
	{}

	/*! \brief Get the cell from the particle position
	 *
	 * \param p position of the particle
	 *
	 */
	__device__ inline unsigned int ghost_processorID_cell(const Point<dim,T> & p)
	{
		return geo_cell.getCell(p);
	}

	/*! \brief Get the number of processor a particle must sent
	 *
	 * \param p position of the particle
	 *
	 */
	__device__ inline unsigned int ghost_processorID_N(const Point<dim,T> & p)
	{
incardon's avatar
incardon committed
90
		return ghost_processorID_N_impl(p,geo_cell,vb_int_proc);
incardon's avatar
incardon committed
91 92 93 94 95 96 97 98 99 100 101 102 103 104 105 106 107 108
	}

	/*! \brief Get the number of processor a particle must sent
	 *
	 * \param p position of the particle
	 *
	 */
	template<typename output_type> __device__ inline void ghost_processor_ID(const Point<dim,T> & p, output_type & output, unsigned int base, unsigned int pi)
	{
		unsigned int cell = geo_cell.getCell(p);
		unsigned int sz = geo_cell.getNelements(cell);

		unsigned int n = 0;

		for (int i = 0 ; i < sz ; i++)
		{
			unsigned int bid = geo_cell.get(cell,i);

incardon's avatar
incardon committed
109 110 111
			unsigned int sz2 = vb_int_proc.template get<0>(bid).size();

			for (int j = 0 ; j < sz2 ; j++)
incardon's avatar
incardon committed
112
			{
incardon's avatar
incardon committed
113 114 115 116
				if (Box<dim,T>(vb_int_proc.template get<0>(bid).get(j)).isInsideNP(p) == true)
				{
					output.template get<0>(base+n) = vb_int_proc.template get<1>(bid);
					output.template get<1>(base+n) = pi;
incardon's avatar
incardon committed
117

incardon's avatar
incardon committed
118 119
					n++;
				}
incardon's avatar
incardon committed
120 121 122 123 124 125 126 127 128
			}
		}
	}

};



#endif /* IE_GHOST_GPU_CUH_ */