Skip to content

GitLab

Commit b8ba87a6 authored by incardon's avatar incardon
Browse files
Options

Test for local ghost comunication working

parent 7f2973d5
Hide whitespace changes
Inline Side-by-side
Showing with 154 additions and 1 deletion
src/Vector/cuda/vector_dist_cuda_func_test.cu
View file @ b8ba87a6
......@@ -152,6 +152,7 @@ BOOST_AUTO_TEST_CASE( vector_ghost_process_local_particles )
size_t old = v_pos.size();
v_pos.resize(v_pos.size() + tot);
v_prp.resize(v_prp.size() + tot);
shifts.template hostToDevice<0>();
openfpm::vector_gpu<aggregate<unsigned int,unsigned int>> o_part_loc2;
......@@ -168,6 +169,7 @@ BOOST_AUTO_TEST_CASE( vector_ghost_process_local_particles )
v_pos.deviceToHost<0>();
o_part_loc2.deviceToHost<0,1>();
v_prp.deviceToHost<0,1,2>();
size_t base = old;
size_t base_o = 0;
......@@ -185,6 +187,24 @@ BOOST_AUTO_TEST_CASE( vector_ghost_process_local_particles )
match &= o_part_loc2.template get<0>(base_o) == i;
match &= o_part_loc2.template get<1>(base_o) == j;
////// We check the properties
match &= v_prp.template get<0>(base) == v_prp.template get<0>(i);
match &= v_prp.template get<1>(base)[0] == v_prp.template get<1>(i)[0];
match &= v_prp.template get<1>(base)[1] == v_prp.template get<1>(i)[1];
match &= v_prp.template get<1>(base)[2] == v_prp.template get<1>(i)[2];
match &= v_prp.template get<2>(base)[0][0] == v_prp.template get<2>(i)[0][0];
match &= v_prp.template get<2>(base)[0][1] == v_prp.template get<2>(i)[0][1];
match &= v_prp.template get<2>(base)[0][2] == v_prp.template get<2>(i)[0][2];
match &= v_prp.template get<2>(base)[1][0] == v_prp.template get<2>(i)[1][0];
match &= v_prp.template get<2>(base)[1][1] == v_prp.template get<2>(i)[1][1];
match &= v_prp.template get<2>(base)[1][2] == v_prp.template get<2>(i)[1][2];
match &= v_prp.template get<2>(base)[2][0] == v_prp.template get<2>(i)[2][0];
match &= v_prp.template get<2>(base)[2][1] == v_prp.template get<2>(i)[2][1];
match &= v_prp.template get<2>(base)[2][2] == v_prp.template get<2>(i)[2][2];
base++;
base_o++;
}
......@@ -199,6 +219,24 @@ BOOST_AUTO_TEST_CASE( vector_ghost_process_local_particles )
match &= o_part_loc2.template get<0>(base_o) == i;
match &= o_part_loc2.template get<1>(base_o) == j;
////// We check the properties
match &= v_prp.template get<0>(base) == v_prp.template get<0>(i);
match &= v_prp.template get<1>(base)[0] == v_prp.template get<1>(i)[0];
match &= v_prp.template get<1>(base)[1] == v_prp.template get<1>(i)[1];
match &= v_prp.template get<1>(base)[2] == v_prp.template get<1>(i)[2];
match &= v_prp.template get<2>(base)[0][0] == v_prp.template get<2>(i)[0][0];
match &= v_prp.template get<2>(base)[0][1] == v_prp.template get<2>(i)[0][1];
match &= v_prp.template get<2>(base)[0][2] == v_prp.template get<2>(i)[0][2];
match &= v_prp.template get<2>(base)[1][0] == v_prp.template get<2>(i)[1][0];
match &= v_prp.template get<2>(base)[1][1] == v_prp.template get<2>(i)[1][1];
match &= v_prp.template get<2>(base)[1][2] == v_prp.template get<2>(i)[1][2];
match &= v_prp.template get<2>(base)[2][0] == v_prp.template get<2>(i)[2][0];
match &= v_prp.template get<2>(base)[2][1] == v_prp.template get<2>(i)[2][1];
match &= v_prp.template get<2>(base)[2][2] == v_prp.template get<2>(i)[2][2];
base++;
base_o++;
}
......@@ -213,6 +251,24 @@ BOOST_AUTO_TEST_CASE( vector_ghost_process_local_particles )
match &= o_part_loc2.template get<0>(base_o) == i;
match &= o_part_loc2.template get<1>(base_o) == j;
////// We check the properties
match &= v_prp.template get<0>(base) == v_prp.template get<0>(i);
match &= v_prp.template get<1>(base)[0] == v_prp.template get<1>(i)[0];
match &= v_prp.template get<1>(base)[1] == v_prp.template get<1>(i)[1];
match &= v_prp.template get<1>(base)[2] == v_prp.template get<1>(i)[2];
match &= v_prp.template get<2>(base)[0][0] == v_prp.template get<2>(i)[0][0];
match &= v_prp.template get<2>(base)[0][1] == v_prp.template get<2>(i)[0][1];
match &= v_prp.template get<2>(base)[0][2] == v_prp.template get<2>(i)[0][2];
match &= v_prp.template get<2>(base)[1][0] == v_prp.template get<2>(i)[1][0];
match &= v_prp.template get<2>(base)[1][1] == v_prp.template get<2>(i)[1][1];
match &= v_prp.template get<2>(base)[1][2] == v_prp.template get<2>(i)[1][2];
match &= v_prp.template get<2>(base)[2][0] == v_prp.template get<2>(i)[2][0];
match &= v_prp.template get<2>(base)[2][1] == v_prp.template get<2>(i)[2][1];
match &= v_prp.template get<2>(base)[2][2] == v_prp.template get<2>(i)[2][2];
base++;
base_o++;
}
......@@ -227,6 +283,24 @@ BOOST_AUTO_TEST_CASE( vector_ghost_process_local_particles )
match &= o_part_loc2.template get<0>(base_o) == i;
match &= o_part_loc2.template get<1>(base_o) == j;
////// We check the properties
match &= v_prp.template get<0>(base) == v_prp.template get<0>(i);
match &= v_prp.template get<1>(base)[0] == v_prp.template get<1>(i)[0];
match &= v_prp.template get<1>(base)[1] == v_prp.template get<1>(i)[1];
match &= v_prp.template get<1>(base)[2] == v_prp.template get<1>(i)[2];
match &= v_prp.template get<2>(base)[0][0] == v_prp.template get<2>(i)[0][0];
match &= v_prp.template get<2>(base)[0][1] == v_prp.template get<2>(i)[0][1];
match &= v_prp.template get<2>(base)[0][2] == v_prp.template get<2>(i)[0][2];
match &= v_prp.template get<2>(base)[1][0] == v_prp.template get<2>(i)[1][0];
match &= v_prp.template get<2>(base)[1][1] == v_prp.template get<2>(i)[1][1];
match &= v_prp.template get<2>(base)[1][2] == v_prp.template get<2>(i)[1][2];
match &= v_prp.template get<2>(base)[2][0] == v_prp.template get<2>(i)[2][0];
match &= v_prp.template get<2>(base)[2][1] == v_prp.template get<2>(i)[2][1];
match &= v_prp.template get<2>(base)[2][2] == v_prp.template get<2>(i)[2][2];
base++;
base_o++;
}
......@@ -234,6 +308,77 @@ BOOST_AUTO_TEST_CASE( vector_ghost_process_local_particles )
}
BOOST_REQUIRE_EQUAL(match,true);
////////////// Now we check that o_part_loc2 contain processble information
openfpm::vector_gpu<Point<3,float>> v_pos2;
openfpm::vector_gpu<prop> v_prp2;
v_pos2.resize(old);
v_prp2.resize(old);
for (size_t i = 0 ; i < old ; i++)
{
v_pos2.template get<0>(i)[0] = v_pos.template get<0>(i)[0];
v_pos2.template get<0>(i)[1] = v_pos.template get<0>(i)[1];
v_pos2.template get<0>(i)[2] = v_pos.template get<0>(i)[2];
v_prp2.template get<0>(i) = v_prp.template get<0>(i);
v_prp2.template get<1>(i)[0] = v_prp.template get<1>(i)[0];
v_prp2.template get<1>(i)[1] = v_prp.template get<1>(i)[1];
v_prp2.template get<1>(i)[2] = v_prp.template get<1>(i)[2];
v_prp2.template get<2>(i)[0][0] = v_prp.template get<2>(i)[0][0];
v_prp2.template get<2>(i)[0][1] = v_prp.template get<2>(i)[0][1];
v_prp2.template get<2>(i)[0][2] = v_prp.template get<2>(i)[0][2];
v_prp2.template get<2>(i)[1][0] = v_prp.template get<2>(i)[1][0];
v_prp2.template get<2>(i)[1][1] = v_prp.template get<2>(i)[1][1];
v_prp2.template get<2>(i)[1][2] = v_prp.template get<2>(i)[1][2];
v_prp2.template get<2>(i)[2][0] = v_prp.template get<2>(i)[2][0];
v_prp2.template get<2>(i)[2][1] = v_prp.template get<2>(i)[2][1];
v_prp2.template get<2>(i)[2][2] = v_prp.template get<2>(i)[2][2];
}
v_pos2.resize(v_pos.size());
v_prp2.resize(v_prp.size());
v_pos2.hostToDevice<0>();
v_prp2.hostToDevice<0,1,2>();
ite = o_part_loc2.getGPUIterator();
process_ghost_particles_local<true,3,decltype(o_part_loc2.toKernel()),decltype(v_pos2.toKernel()),decltype(v_prp2.toKernel()),decltype(shifts.toKernel())>
<<<ite.wthr,ite.thr>>>
(o_part_loc2.toKernel(),v_pos2.toKernel(),v_prp2.toKernel(),shifts.toKernel(),old);
v_pos2.template deviceToHost<0>();
v_prp2.template deviceToHost<0,1,2>();
for (size_t i = old ; i < v_pos.size() ; i++)
{
match &= v_pos.template get<0>(i)[0] == v_pos2.template get<0>(i)[0];
match &= v_pos.template get<0>(i)[1] == v_pos2.template get<0>(i)[1];
match &= v_pos.template get<0>(i)[2] == v_pos2.template get<0>(i)[2];
match &= v_prp2.template get<0>(i) == v_prp.template get<0>(i);
match &= v_prp2.template get<1>(i)[0] == v_prp.template get<1>(i)[0];
match &= v_prp2.template get<1>(i)[1] == v_prp.template get<1>(i)[1];
match &= v_prp2.template get<1>(i)[2] == v_prp.template get<1>(i)[2];
match &= v_prp2.template get<2>(i)[0][0] == v_prp.template get<2>(i)[0][0];
match &= v_prp2.template get<2>(i)[0][1] == v_prp.template get<2>(i)[0][1];
match &= v_prp2.template get<2>(i)[0][2] == v_prp.template get<2>(i)[0][2];
match &= v_prp2.template get<2>(i)[1][0] == v_prp.template get<2>(i)[1][0];
match &= v_prp2.template get<2>(i)[1][1] == v_prp.template get<2>(i)[1][1];
match &= v_prp2.template get<2>(i)[1][2] == v_prp.template get<2>(i)[1][2];
match &= v_prp2.template get<2>(i)[2][0] == v_prp.template get<2>(i)[2][0];
match &= v_prp2.template get<2>(i)[2][1] == v_prp.template get<2>(i)[2][1];
match &= v_prp2.template get<2>(i)[2][2] == v_prp.template get<2>(i)[2][2];
}
BOOST_REQUIRE_EQUAL(match,true);
}
BOOST_AUTO_TEST_CASE( vector_ghost_fill_send_buffer_test )
......
src/Vector/cuda/vector_dist_cuda_funcs.cuh
View file @ b8ba87a6
......@@ -191,6 +191,9 @@ __global__ void shift_ghost_each_part(vector_of_box box_f, vector_of_shifts box_
output.template get<0>(base_o+n) = p;
output.template get<1>(base_o+n) = shift_id;
}
v_prp.set(base+n,v_prp.get(p));
n++;
}
}
......
src/Vector/vector_dist_comm.hpp
View file @ b8ba87a6
......@@ -310,9 +310,14 @@ class vector_dist_comm
auto ite = o_part_loc.getGPUIterator();
size_t old = v_pos.size();
v_pos.resize(v_pos.size() + o_part_loc.size(),DATA_ON_DEVICE);
v_prp.resize(v_prp.size() + o_part_loc.size(),DATA_ON_DEVICE);
process_ghost_particles_local<true,dim,decltype(o_part_loc.toKernel()),decltype(v_pos.toKernel()),decltype(v_prp.toKernel()),decltype(shifts.toKernel())>
<<<ite.wthr,ite.thr>>>
(o_part_loc.toKernel(),v_pos.toKernel(),v_prp.toKernel(),shifts.toKernel(),v_pos.size());
(o_part_loc.toKernel(),v_pos.toKernel(),v_prp.toKernel(),shifts.toKernel(),old);
#else
std::cout << __FILE__ << ":" << __LINE__ << " error: to use the option RUN_ON_DEVICE you must compile with NVCC" << std::endl;
......
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