Speedup sparse-grid skip labelling

openfpm_data @ 776a4aa1

-Subproject commit e4ecadaf7f0def83285c4a8c06c54a9d158ad677
+Subproject commit 776a4aa1f036f83d16e767977f6bb2732b15f3a1

src/CMakeLists.txt

@@ -193,6 +193,9 @@ install(FILES Vector/se_class3_vector.hpp
 	      Vector/vector_dist_kernel.hpp
 	      DESTINATION openfpm_pdata/include/Vector )
 
+install(FILES util/common_pdata.hpp
+	      DESTINATION openfpm_pdata/include/util)
+
 install(FILES Vector/Iterators/vector_dist_iterator.hpp
 	      DESTINATION openfpm_pdata/include/Vector/Iterators/ )
 

src/Grid/grid_dist_id_comm.hpp

@@ -203,8 +203,15 @@ class grid_dist_id_comm
 											  bool use_bx_def,
 											  size_t opt)
 	{
-		for (size_t i = 0 ; i < loc_grid.size() ; i++)
-		{loc_grid.get(i).copyRemoveReset();}
+		rem_copy_opt opt_ = rem_copy_opt::NONE_OPT;
+		if (opt & SKIP_LABELLING)
+		{opt_ = rem_copy_opt::KEEP_GEOMETRY;}
+
+		if (opt_ != rem_copy_opt::KEEP_GEOMETRY)
+		{
+			for (size_t i = 0 ; i < loc_grid.size() ; i++)
+			{loc_grid.get(i).copyRemoveReset();}
+		}
 
 		grid_key_dx<dim> cnt[1];
 		cnt[0].zero();
@@ -259,11 +266,6 @@ class grid_dist_id_comm
 			}
 		}
 
-		rem_copy_opt opt_ = rem_copy_opt::NONE_OPT;
-		if (opt & SKIP_LABELLING == true)
-		{opt_ = rem_copy_opt::KEEP_GEOMETRY;}
-
-
 		for (size_t i = 0 ; i < loc_grid.size() ; i++)
 		{
 			loc_grid.get(i).template removeCopyToFinalize<prp ...>(v_cl.getmgpuContext(), rem_copy_opt::PHASE1 | opt_);

src/Grid/tests/sgrid_dist_id_gpu_unit_tests.cu

@@ -676,6 +676,13 @@ BOOST_AUTO_TEST_CASE( sgrid_gpu_test_skip_labelling )
 
 	gdist.template ghost_get<0,1>(RUN_ON_DEVICE | SKIP_LABELLING);
 
+	gdist.template conv2<0,1,0,1,1>({0,0,0},{(int)sz[0]-1,(int)sz[1]-1,(int)sz[2]-1},[] __device__ (float & u_out, float & v_out, CpBlockType & u, CpBlockType & v,int i, int j, int k){
+		u_out = 2*u(i,j,k);
+		v_out = 2*v(i,j,k);
+	});
+
+	gdist.template ghost_get<0,1>(RUN_ON_DEVICE | SKIP_LABELLING);
+
 	gdist.template conv2<0,1,2,3,1>({2,2,2},{(int)sz[0]-3,(int)sz[1]-3,(int)sz[2]-3},[] __device__ (float & u_out, float & v_out, CpBlockType & u, CpBlockType & v,int i, int j, int k){
 		u_out = u(i+1,j,k) - u(i-1,j,k) + u(i,j+1,k) - u(i,j-1,k) + u(i,j,k+1) - u(i,j,k-1);
 		v_out = v(i+1,j,k) - v(i-1,j,k) + v(i,j+1,k) - v(i,j-1,k) + v(i,j,k+1) - v(i,j,k-1);
@@ -704,7 +711,7 @@ BOOST_AUTO_TEST_CASE( sgrid_gpu_test_skip_labelling )
 		float sub1 = gdist.template get<2>(p);
 		float sub2 = gdist.template get<3>(p);
 
-		if (sub1 != 6.0*5 || sub2 != 6.0*5)
+		if (sub1 != 6.0*10.0 || sub2 != 6.0*10.0)
 		{
 			std::cout << sub1 << "  " << sub2 << std::endl;
 			std::cout << gdist.template get<0>(p_xp1) << "   " << gdist.template get<0>(p_xm1) << std::endl;