Commit 47f83be8 authored by incardon's avatar incardon

Compilable ie_ghost on gpu

parent 6e4ec106
......@@ -112,7 +112,9 @@ struct Box_sub_k
}
};
template<unsigned int dim,typename T>
template<unsigned int dim,typename T> using Box_map = aggregate<Box<dim,T>,long int>;
/*template<unsigned int dim,typename T>
struct Box_map
{
typedef boost::fusion::vector<Box<dim,T>,long int> type;
......@@ -125,7 +127,7 @@ struct Box_map
}
static const unsigned int max_prop = 2;
};
};*/
//! Case for local ghost box
template<unsigned int dim, typename T>
......
......@@ -14,8 +14,8 @@ constexpr unsigned int lc_proc_ = 0;
constexpr unsigned int proc_ = 1;
constexpr unsigned int shift_id_ = 2;
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)
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)
{
unsigned int cell = geo_cell.getCell(p);
unsigned int sz = geo_cell.getNelements(cell);
......@@ -26,9 +26,14 @@ __device__ __host__ inline unsigned int ghost_processorID_N_impl(const Point<dim
{
unsigned int bid = geo_cell.get(cell,i);
if (Box<dim,T>(vb_int_box.get(bid)).isInsideNP(p) == true)
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++;}
}
}
return n;
}
......@@ -49,28 +54,20 @@ class ie_ghost_gpu
CellList_cpu_ker<dim,T,Mem_fast_ker<Memory,memory_traits_lin,int>,shift<dim,T>> geo_cell;
//! internal ghost box
openfpm::vector_gpu_ker<Box<dim, T>,layout_base> vb_int_box;
//! internal ghost box
openfpm::vector_gpu_ker<aggregate<unsigned int,unsigned int,unsigned int>,layout_base> vb_int;
// maximum rank
unsigned int max_rank;
openfpm::vector_gpu_ker<aggregate<openfpm::vector_gpu_ker<Box<dim, T>,layout_base>,int>,layout_base> vb_int_proc;
public:
ie_ghost_gpu(CellList_cpu_ker<dim,T,Mem_fast_ker<Memory,memory_traits_lin,int>,shift<dim,T>> geo_cell,
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,
unsigned int max_rank)
:geo_cell(geo_cell),vb_int_box(vb_int_box),vb_int(vb_int),max_rank(max_rank)
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)
{
}
ie_ghost_gpu(const ie_ghost_gpu<dim,T,Memory,layout_base> & ieg)
:geo_cell(ieg.geo_cell),vb_int_box(ieg.vb_int_box),vb_int(ieg.vb_int),max_rank(ieg.max_rank)
:geo_cell(ieg.geo_cell),vb_int_proc(ieg.vb_int_proc)
{}
/*! \brief Get the cell from the particle position
......@@ -90,7 +87,7 @@ public:
*/
__device__ inline unsigned int ghost_processorID_N(const Point<dim,T> & p)
{
return ghost_processorID_N_impl(p,geo_cell,vb_int_box,vb_int);
return ghost_processorID_N_impl(p,geo_cell,vb_int_proc);
}
/*! \brief Get the number of processor a particle must sent
......@@ -109,15 +106,20 @@ public:
{
unsigned int bid = geo_cell.get(cell,i);
if (Box<dim,T>(vb_int_box.get(bid)).isInsideNP(p) == true)
unsigned int sz2 = vb_int_proc.template get<0>(bid).size();
for (int j = 0 ; j < sz2 ; j++)
{
output.template get<0>(base+n) = vb_int.template get<proc_>(bid);
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;
n++;
}
}
}
}
};
......
......@@ -52,7 +52,7 @@ class ie_ghost
typedef openfpm::vector<Box<dim,T>,Memory,typename layout_base<Box<dim,T>>::type,layout_base> proc_boxes;
//! internal ghost Boxes for each processor
openfpm::vector<aggregate<proc_boxes>,Memory,typename layout_base<aggregate<proc_boxes>>::type,layout_base> vb_int_proc;
openfpm::vector<aggregate<proc_boxes,int>,Memory,typename layout_base<aggregate<proc_boxes,int>>::type,layout_base> vb_int_proc;
//! shift vectors
openfpm::vector<Point<dim,T>> shifts;
......@@ -814,7 +814,7 @@ public:
*/
inline unsigned int ghost_processorID_N(const Point<dim,T> & p)
{
return ghost_processorID_N_impl(p,geo_cell,vb_int_box,vb_int);
return ghost_processorID_N_impl(p,geo_cell,vb_int_proc);
}
/*! \brief Given a position it return if the position belong to any neighborhood processor ghost
......@@ -1149,7 +1149,8 @@ public:
host_dev_transfer = true;
}
ie_ghost_gpu<dim,T,Memory,layout_base> igg(geo_cell.toKernel(),vb_int_box.toKernel(),vb_int.toKernel(),create_vcluster().size());
ie_ghost_gpu<dim,T,Memory,layout_base> igg(geo_cell.toKernel(),
vb_int_proc.toKernel());
return igg;
}
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment