Commit 7f2973d5 authored by incardon's avatar incardon

ghost local particles moving on with tests

parent e0df46eb
openfpm_data @ 42183947
Subproject commit c46dc01e03b1ea215239e8cb817ebff4c61a11db
Subproject commit 42183947ec6434fa7644185690b005c494b26676
......@@ -26,8 +26,7 @@ struct ID_operation
__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;
output.template get<1>(base + n) = pi;
output.template get<2>(base + n) = shift_act;
output.template get<1>(base + n) = (unsigned long int)pi + (((unsigned long int)shift_act) << 32);
}
};
......
......@@ -66,16 +66,10 @@ class ie_ghost
//! Cell-list that store the geometrical information of the internal ghost boxes
CellList<dim,T,Mem_fast<Memory,int>,shift<dim,T>> geo_cell;
//! Cell-list that store the geometrical information of the internal ghost boxes (on a processor based lavel)
CellList<dim,T,Mem_fast<Memory,int>,shift<dim,T>> geo_cell_proc;
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,int>,Memory,typename layout_base<aggregate<proc_boxes,int>>::type,layout_base> vb_int_proc;
//! shift vectors
openfpm::vector<Point<dim,T>> shifts;
openfpm::vector<Point<dim,T>,Memory,typename layout_base<Point<dim,T>>::type,layout_base> shifts;
//! Temporal buffers to return temporal information for ghost_processorID
openfpm::vector<std::pair<size_t,size_t>> ids_p;
......@@ -84,7 +78,7 @@ class ie_ghost
openfpm::vector<size_t> ids;
//! shift converter
shift_vect_converter<dim,T> sc_convert;
shift_vect_converter<dim,T,Memory,layout_base> sc_convert;
//! host to device transfer
bool host_dev_transfer = false;
......@@ -210,9 +204,6 @@ protected:
{
// Initialize the geo_cell structure
geo_cell.Initialize(domain,div,0);
// Initialize the geo_cell structure
geo_cell_proc.Initialize(domain,div,0);
}
/*! \brief Create the box_nn_processor_int (bx part) structure
......@@ -315,30 +306,6 @@ protected:
}
}
/*! \brief construct the vb_int_proc box
*
*
*/
void construct_vb_int_proc(const nn_prcs<dim,T> & nn_p)
{
vb_int_proc.resize_no_device(proc_int_box.size());
for (size_t i = 0 ; i < proc_int_box.size() ; i++)
{
vb_int_proc.template get<0>(i).resize(proc_int_box.get(i).ibx.size());
for (size_t j = 0 ; j < proc_int_box.get(i).ibx.size() ; j++)
{
for (size_t k = 0 ; k < dim ; k++)
{
vb_int_proc.template get<0>(i).template get<0>(j)[k] = proc_int_box.get(i).ibx.get(j).bx.getLow(k);
vb_int_proc.template get<0>(i).template get<1>(j)[k] = proc_int_box.get(i).ibx.get(j).bx.getHigh(k);
}
}
vb_int_proc.template get<1>(i) = nn_p.IDtoProc(i);
}
}
/*! \brief Create the box_nn_processor_int (nbx part) structure, the geo_cell list and proc_int_box
*
......@@ -470,24 +437,6 @@ protected:
geo_cell.addCell(cell,vb_int.size()-1);
// Check if p_id already exist at that cell
// and we add it only if does not exist
size_t nc = geo_cell_proc.getNelements(cell);
bool found = false;
for (size_t s = 0; s < nc ; s++)
{
if (geo_cell_proc.get(cell,s) == lc_proc)
{
found = true;
break;
}
}
if (found == false)
{geo_cell_proc.addCell(cell,lc_proc);}
++g_sub;
}
}
......@@ -495,8 +444,6 @@ protected:
}
}
construct_vb_int_proc(nn_p);
reorder_geo_cell();
}
......@@ -641,7 +588,7 @@ public:
* \return the shift vectors
*
*/
const openfpm::vector<Point<dim,T>> & getShiftVectors()
const openfpm::vector<Point<dim,T>,Memory,typename layout_base<Point<dim,T>>::type,layout_base> & getShiftVectors()
{
return shifts;
}
......@@ -1244,13 +1191,9 @@ public:
if (host_dev_transfer == false)
{
geo_cell.hostToDevice();
geo_cell_proc.hostToDevice();
vb_int_box.template hostToDevice<0,1>();
vb_int.template hostToDevice<0,1,2>();
for (size_t i = 0 ; i < vb_int_proc.size() ; i++)
{vb_int_proc.template get<0>(i). template hostToDevice<0,1>();}
vb_int_proc.template hostToDevice<0,1>();
shifts.template hostToDevice<0>();
host_dev_transfer = true;
}
......
......@@ -16,7 +16,7 @@
* handle such case
*
*/
template<unsigned int dim, typename T>
template<unsigned int dim, typename T, typename Memory, template<typename> class layout_base>
class shift_vect_converter
{
//! Indicate which indexes are non_periodic
......@@ -33,7 +33,8 @@ class shift_vect_converter
* \param domain box that describe the domain
*
*/
void generateShiftVectors_ld(const Box<dim,T> & domain, size_t (& bc)[dim], openfpm::vector<Point<dim,T>> & shifts)
void generateShiftVectors_ld(const Box<dim,T> & domain, size_t (& bc)[dim],
openfpm::vector<Point<dim,T>,Memory,typename layout_base<Point<dim,T>>::type,layout_base> & shifts)
{
shifts.resize(openfpm::math::pow(3,dim));
......@@ -69,7 +70,8 @@ class shift_vect_converter
* \param domain box that describe the domain
*
*/
void generateShiftVectors_hd(const Box<dim,T> & domain, size_t (& bc)[dim], openfpm::vector<Point<dim,T>> & shifts)
void generateShiftVectors_hd(const Box<dim,T> & domain, size_t (& bc)[dim],
openfpm::vector<Point<dim,T>,Memory,typename layout_base<Point<dim,T>>::type,layout_base> & shifts)
{
// get the indexes of the free degree of freedom
for (size_t i = 0 ; i < dim ; i++)
......@@ -123,7 +125,8 @@ public:
* \param domain box that describe the domain
*
*/
void generateShiftVectors(const Box<dim,T> & domain, size_t (& bc)[dim], openfpm::vector<Point<dim,T>> & shifts)
void generateShiftVectors(const Box<dim,T> & domain, size_t (& bc)[dim],
openfpm::vector<Point<dim,T>,Memory,typename layout_base<Point<dim,T>>::type,layout_base> & shifts)
{
if (dim < 10)
{generateShiftVectors_ld(domain,bc,shifts);}
......
......@@ -18,7 +18,7 @@ BOOST_AUTO_TEST_CASE( shift_vect_converter_tests_use )
{
{
Box<3,double> domain({0.0,0.0,0.0},{1.0,1.0,1.0});
shift_vect_converter<3,double> svc;
shift_vect_converter<3,double,HeapMemory,memory_traits_lin> svc;
size_t bc[3] = {PERIODIC,PERIODIC,PERIODIC};
openfpm::vector<Point<3,double>> sv;
......@@ -68,7 +68,7 @@ BOOST_AUTO_TEST_CASE( shift_vect_converter_tests_use )
bc[17] = PERIODIC;
bc[23] = PERIODIC;
shift_vect_converter<50,double> svc;
shift_vect_converter<50,double,HeapMemory,memory_traits_lin> svc;
svc.generateShiftVectors(domain,bc,sv);
......
......@@ -51,14 +51,14 @@ __global__ void process_id_proc_each_part(cartdec_gpu cdg, particles_type parts,
output.template get<0>(p) = p;
}
template<typename vector_type,typename vector_type_offs>
template<unsigned int prp_off, typename vector_type,typename vector_type_offs>
__global__ void find_buffer_offsets(vector_type vd, int * cnt, vector_type_offs offs)
{
int p = threadIdx.x + blockIdx.x * blockDim.x;
if (p >= vd.size() - 1) return;
if (vd.template get<1>(p) != vd.template get<1>(p+1))
if (vd.template get<prp_off>(p) != vd.template get<prp_off>(p+1))
{
int i = atomicAdd(cnt, 1);
offs.template get<0>(i) = p+1;
......@@ -78,5 +78,122 @@ __global__ void process_map_particles(vector_m_opart_type m_opart, vector_pos_t
process_map_device_particle<proc_without_prp_device>(i,offset,m_opart,m_pos,m_prp,v_pos,v_prp);
}
template<typename vector_g_opart_type, typename vector_prp_type_out, typename vector_prp_type_in, unsigned int ... prp>
__global__ void process_ghost_particles_prp(vector_g_opart_type g_opart, vector_prp_type_out m_prp,
vector_prp_type_in v_prp, unsigned int offset)
{
int i = threadIdx.x + blockIdx.x * blockDim.x;
if (i >= m_prp.size()) return;
process_ghost_device_particle_prp<vector_g_opart_type,vector_prp_type_out,vector_prp_type_in,prp...>(i,offset,g_opart,m_prp,v_prp);
}
template<unsigned int dim, typename vector_g_opart_type, typename vector_pos_type_out, typename vector_pos_type_in, typename vector_shift_type_in>
__global__ void process_ghost_particles_pos(vector_g_opart_type g_opart, vector_pos_type_out m_pos,
vector_pos_type_in v_pos, vector_shift_type_in shifts, unsigned int offset)
{
int i = threadIdx.x + blockIdx.x * blockDim.x;
if (i >= m_pos.size()) return;
unsigned long int psid = g_opart.template get<1>(i+offset);
unsigned int id = psid & 0xFFFFFFFF;
unsigned int shift_id = psid >> 32;
#pragma unroll
for (int j = 0; j < dim ; j++)
{
m_pos.template get<0>(i)[j] = v_pos.template get<0>(id)[j] - shifts.template get<0>(shift_id)[j];
}
}
template<bool with_pos, unsigned int dim, typename vector_g_opart_type, typename vector_pos_type,
typename vector_prp_type, typename vector_shift_type_in>
__global__ void process_ghost_particles_local(vector_g_opart_type g_opart, vector_pos_type v_pos, vector_prp_type v_prp,
vector_shift_type_in shifts, unsigned int base)
{
int i = threadIdx.x + blockIdx.x * blockDim.x;
if (i >= g_opart.size()) return;
unsigned int pid = g_opart.template get<0>(i);
unsigned int shift_id = g_opart.template get<1>(i);
if (with_pos == true)
{
#pragma unroll
for (int j = 0; j < dim ; j++)
{
v_pos.template get<0>(base+i)[j] = v_pos.template get<0>(pid)[j] - shifts.template get<0>(shift_id)[j];
}
}
v_prp.set(base+i,v_prp.get(pid));
}
template<unsigned int dim, typename St, typename vector_of_box, typename vector_type, typename output_type>
__global__ void num_shift_ghost_each_part(vector_of_box box_f, vector_type vd, output_type out)
{
int p = threadIdx.x + blockIdx.x * blockDim.x;
if (p >= vd.size()) return;
Point<dim,St> xp = vd.template get<0>(p);
unsigned int n = 0;
for (unsigned int i = 0 ; i < box_f.size() ; i++)
{
if (Box<dim,St>(box_f.get(i)).isInsideNP(xp) == true)
{n++;}
}
out.template get<0>(p) = n;
}
template<unsigned int dim, typename St,
typename vector_of_box,
typename vector_of_shifts,
typename vector_type_pos,
typename vector_type_prp,
typename start_type,
typename shifts_type,
typename output_type>
__global__ void shift_ghost_each_part(vector_of_box box_f, vector_of_shifts box_f_sv,
vector_type_pos v_pos, vector_type_prp v_prp,
start_type start, shifts_type shifts,
output_type output, unsigned int offset)
{
int p = threadIdx.x + blockIdx.x * blockDim.x;
if (p >= v_pos.size()) return;
Point<dim,St> xp = v_pos.template get<0>(p);
unsigned int base_o = start.template get<0>(p);
unsigned int base = base_o + offset;
unsigned int n = 0;
for (unsigned int i = 0 ; i < box_f.size() ; i++)
{
if (Box<dim,St>(box_f.get(i)).isInsideNP(xp) == true)
{
unsigned int shift_id = box_f_sv.template get<0>(i);
#pragma unroll
for (unsigned int j = 0 ; j < dim ; j++)
{
v_pos.template get<0>(base+n)[j] = xp.get(j) - shifts.template get<0>(shift_id)[j];
output.template get<0>(base_o+n) = p;
output.template get<1>(base_o+n) = shift_id;
}
n++;
}
}
}
#endif /* VECTOR_DIST_CUDA_FUNCS_CUH_ */
......@@ -494,9 +494,18 @@ BOOST_AUTO_TEST_CASE( vector_dist_map_on_gpu_test)
BOOST_REQUIRE_EQUAL(n_out,0);
BOOST_REQUIRE_EQUAL(l_cnt,vd.size_local());
vd.write("gpu_write_test");
vd.deviceToHostPos();
vd.deviceToHostProp<0,1,2>();
vd.write("gpu_write_before_test");
vd.ghost_get<0,1,2>(RUN_ON_DEVICE);
vd.deviceToHostPos();
vd.deviceToHostProp<0,1,2>();
vd.write("gpu_write_after_test");
}
BOOST_AUTO_TEST_SUITE_END()
......@@ -278,7 +278,7 @@ BOOST_AUTO_TEST_CASE( vector_dist_ghost )
BOOST_AUTO_TEST_CASE( vector_dist_ghost_inte )
{
typedef vector_dist<2,float, Point_test<float>,CartDecomposition<2,float>,HeapMemory,memory_traits_inte> vector;
typedef vector_dist_soa<2,float, Point_test<float>> vector;
Box<2,float> box({0.0,0.0},{1.0,1.0});
Test2D_ghost<vector>(box);
......
......@@ -2524,5 +2524,6 @@ public:
template<unsigned int dim, typename St, typename prop, typename Decomposition = CartDecomposition<dim,St,CudaMemory,memory_traits_inte>> using vector_dist_gpu = vector_dist<dim,St,prop,Decomposition,CudaMemory,memory_traits_inte>;
template<unsigned int dim, typename St, typename prop, typename Decomposition = CartDecomposition<dim,St,HeapMemory,memory_traits_inte>> using vector_dist_soa = vector_dist<dim,St,prop,Decomposition,HeapMemory,memory_traits_inte>;
#endif /* VECTOR_HPP_ */
This diff is collapsed.
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