ie_ghost_gpu.cuh 4.93 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
19
20
21
22
23
24
25
26
27
28
template<typename output_type>
struct ID_operation
{
	output_type & output;

	__device__ __host__ ID_operation(output_type & output)
	:output(output)
	{}

	__device__ __host__ inline void op(unsigned int base, unsigned int n, unsigned int proc_act, unsigned int shift_act, unsigned int pi)
	{
		output.template get<0>(base + n) = proc_act;
29
		output.template get<1>(base + n) = (unsigned long int)pi + (((unsigned long int)shift_act) << 32);
incardon's avatar
incardon committed
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
	}
};

struct N_operation
{
	__device__ __host__ inline void op(unsigned int base, unsigned int n, unsigned int proc_act, unsigned int shift_act, unsigned int pi)
	{
	}
};

template<unsigned int dim, typename T, typename cell_list_type, typename vb_int_box_type, typename vb_int_type, typename operation>
__device__ __host__ inline unsigned int ghost_processorID_general_impl(const Point<dim,T> & p,
																 unsigned int base,
																 unsigned int pi,
																 cell_list_type & geo_cell,
																 vb_int_box_type & vb_int_box,
																 vb_int_type & vb_int,
																 operation & op)
incardon's avatar
incardon committed
48
49
50
51
52
53
{
	unsigned int cell = geo_cell.getCell(p);
	unsigned int sz = geo_cell.getNelements(cell);

	unsigned int n = 0;

incardon's avatar
incardon committed
54
55
56
	bool switch_prc = false;

	if (sz != 0)
incardon's avatar
incardon committed
57
	{
incardon's avatar
incardon committed
58
59
60
61
62
63
64
65
66
67
		int i = 0;
		unsigned int bid = geo_cell.get(cell,0);
		unsigned int proc_prev = vb_int.template get<proc_>(bid);
		unsigned int shift_prev = vb_int.template get<shift_id_>(bid);
		unsigned int proc_act;
		unsigned int shift_act;

		if (Box<dim,T>(vb_int_box.get(bid)).isInsideNP(p) == true)
		{
			op.op(base,n,proc_prev,shift_prev,pi);
incardon's avatar
incardon committed
68

incardon's avatar
incardon committed
69
70
71
72
73
			switch_prc = true;
			n++;
		}

		i++;
incardon's avatar
incardon committed
74

incardon's avatar
incardon committed
75
		for ( ; i < sz ; i++)
incardon's avatar
incardon committed
76
		{
incardon's avatar
incardon committed
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
			unsigned int bid = geo_cell.get(cell,i);
			proc_act = vb_int.template get<proc_>(bid);
			shift_act = vb_int.template get<shift_id_>(bid);

			switch_prc = (proc_act == proc_prev && shift_act == shift_prev) & switch_prc;

			if (Box<dim,T>(vb_int_box.get(bid)).isInsideNP(p) == true && switch_prc == false)
			{
				op.op(base,n,proc_act,shift_act,pi);

				switch_prc = true;
				n++;
			}
			proc_prev = proc_act;
			shift_prev = shift_act;
incardon's avatar
incardon committed
92
		}
incardon's avatar
incardon committed
93
94
95
96
97
	}

	return n;
}

incardon's avatar
incardon committed
98
99
100
101
102
103
104
105
106
107
108
template<unsigned int dim, typename T, typename cell_list_type, typename vb_int_box_type, typename vb_int_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_box,
																 vb_int_type & vb_int)
{
	N_operation op;

	return ghost_processorID_general_impl(p,0,0,geo_cell,vb_int_box,vb_int,op);
}

incardon's avatar
incardon committed
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
/*! \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
125
126
127
128
	openfpm::vector_gpu_ker<Box<dim, T>,layout_base> vb_int_box;

	//! internal ghost box processor infos
	openfpm::vector_gpu_ker<aggregate<unsigned int,unsigned int,unsigned int>,layout_base> vb_int;
incardon's avatar
incardon committed
129
130
131
132
133

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
134
135
136
				 openfpm::vector_gpu_ker<Box<dim, T>,layout_base> vb_int_box,
				 openfpm::vector_gpu_ker<aggregate<unsigned int,unsigned int,unsigned int>,layout_base> vb_int)
	:geo_cell(geo_cell),vb_int_box(vb_int_box),vb_int(vb_int)
incardon's avatar
incardon committed
137
138
139
140
141
	{

	}

	ie_ghost_gpu(const ie_ghost_gpu<dim,T,Memory,layout_base> & ieg)
incardon's avatar
incardon committed
142
	:geo_cell(ieg.geo_cell),vb_int_box(ieg.vb_int_box),vb_int(ieg.vb_int)
incardon's avatar
incardon committed
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
	{}

	/*! \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
162
		return ghost_processorID_N_impl(p,geo_cell,vb_int_box,vb_int);
incardon's avatar
incardon committed
163
164
165
166
167
168
169
170
171
	}

	/*! \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)
	{
incardon's avatar
incardon committed
172
		ID_operation<output_type> op(output);
incardon's avatar
incardon committed
173

incardon's avatar
incardon committed
174
		ghost_processorID_general_impl(p,base,pi,geo_cell,vb_int_box,vb_int,op);
incardon's avatar
incardon committed
175
176
177
178
179
180
181
	}

};



#endif /* IE_GHOST_GPU_CUH_ */