diff --git a/src/Vector/cuda/vector_dist_gpu_unit_tests.cu b/src/Vector/cuda/vector_dist_gpu_unit_tests.cu index 437f1624e82dfe10907252d457db4da90829c4ac..8046a89a55d9fc1820eb92f5b12ace7149712fd4 100644 --- a/src/Vector/cuda/vector_dist_gpu_unit_tests.cu +++ b/src/Vector/cuda/vector_dist_gpu_unit_tests.cu @@ -1304,5 +1304,78 @@ BOOST_AUTO_TEST_CASE(vector_dist_compare_host_device) BOOST_REQUIRE_EQUAL(test,true); } +template<typename vector_dist_type> +__global__ void assign_to_ghost(vector_dist_type vds) +{ + int i = threadIdx.x + blockIdx.x * blockDim.x; + + if (i >= vds.size()) {return;} + + vds.template getProp<0>(i) = 1000.0 + i; + + vds.template getProp<1>(i)[0] = 2000.0 + i; + vds.template getProp<1>(i)[1] = 3000.0 + i; + vds.template getProp<1>(i)[2] = 4000.0 + i; + + vds.template getProp<2>(i)[0][0] = 12000.0 + i; + vds.template getProp<2>(i)[0][1] = 13000.0 + i; + vds.template getProp<2>(i)[0][2] = 14000.0 + i; + vds.template getProp<2>(i)[1][0] = 22000.0 + i; + vds.template getProp<2>(i)[1][1] = 23000.0 + i; + vds.template getProp<2>(i)[1][2] = 24000.0 + i; + vds.template getProp<2>(i)[2][0] = 32000.0 + i; + vds.template getProp<2>(i)[2][1] = 33000.0 + i; + vds.template getProp<2>(i)[2][2] = 34000.0 + i; + +} + +BOOST_AUTO_TEST_CASE(vector_dist_domain_and_ghost_test) +{ + Box<3,double> domain({0.0,0.0,0.0},{1.0,1.0,1.0}); + Ghost<3,double> g(0.1); + size_t bc[3] = {PERIODIC,PERIODIC,PERIODIC}; + + if (create_vcluster().size() >= 16) + {return;} + + vector_dist_gpu<3,double,aggregate<double,double[3],double[3][3]>> vdg(10000,domain,bc,g,DEC_GRAN(128)); + + auto ite = vdg.getDomainAndGhostIteratorGPU(); + + CUDA_LAUNCH(assign_to_ghost,ite,vdg.toKernel()); + + vdg.template deviceToHostProp<0,1,2>(); + + + auto it = vdg.getDomainAndGhostIterator(); + + bool check = true; + + while (it.isNext()) + { + auto k = it.get(); + + check &= vdg.template getProp<0>(k) == 1000.0 + k.getKey(); + + check &= vdg.template getProp<1>(k)[0] == 2000.0 + k.getKey(); + check &= vdg.template getProp<1>(k)[1] == 3000.0 + k.getKey(); + check &= vdg.template getProp<1>(k)[2] == 4000.0 + k.getKey(); + + check &= vdg.template getProp<2>(k)[0][0] == 12000.0 + k.getKey(); + check &= vdg.template getProp<2>(k)[0][1] == 13000.0 + k.getKey(); + check &= vdg.template getProp<2>(k)[0][2] == 14000.0 + k.getKey(); + check &= vdg.template getProp<2>(k)[1][0] == 22000.0 + k.getKey(); + check &= vdg.template getProp<2>(k)[1][1] == 23000.0 + k.getKey(); + check &= vdg.template getProp<2>(k)[1][2] == 24000.0 + k.getKey(); + check &= vdg.template getProp<2>(k)[2][0] == 32000.0 + k.getKey(); + check &= vdg.template getProp<2>(k)[2][1] == 33000.0 + k.getKey(); + check &= vdg.template getProp<2>(k)[2][2] == 34000.0 + k.getKey(); + + ++it; + } + + + BOOST_REQUIRE_EQUAL(check,true); +} BOOST_AUTO_TEST_SUITE_END() diff --git a/src/Vector/vector_dist.hpp b/src/Vector/vector_dist.hpp index 4c839d93aca6c6569b36877ab544b839bb8ea3a9..8d39a884982bff5ecdc81d784733fe21341a4878 100644 --- a/src/Vector/vector_dist.hpp +++ b/src/Vector/vector_dist.hpp @@ -1944,6 +1944,20 @@ public: return v_pos.getGPUIteratorTo(g_m,n_thr); } + /*! \brief Get an iterator that traverse the particles in the domain + * + * \return an iterator + * + */ + ite_gpu<1> getDomainAndGhostIteratorGPU(size_t n_thr = 1024) const + { +#ifdef SE_CLASS3 + se3.getIterator(); +#endif + + return v_pos.getGPUIteratorTo(v_pos.size(),n_thr); + } + /*! \brief Merge the properties calculated on the sorted vector on the original vector * * \parameter Cell-list from which has been constructed the sorted vector