grid_dist_id work with SparseGridGPU

openfpm_data @ b0330228

-Subproject commit f982d0ac3f3ae0eefbd1ca9e39a700b8898f8ee9
+Subproject commit b0330228e8503af944dcc5a05fa4fb1bc3e3b594

src/Amr/grid_dist_amr_key.hpp

+/*
+ * grid_dist_amr_key.hpp
+ *
+ *  Created on: Sep 23, 2017
+ *      Author: i-bird
+ */
+
+#ifndef SRC_AMR_GRID_DIST_AMR_KEY_HPP_
+#define SRC_AMR_GRID_DIST_AMR_KEY_HPP_
+
+/*! \brief Amr grid distributed key
+ *
+ * \tparam dim dimensionality
+ *
+ */
+template<unsigned int dim>
+class grid_dist_amr_key
+{
+	//! actual level
+	size_t lvl;
+
+	//! actual position in the distributed grid
+	grid_dist_key_dx<dim> key;
+
+
+public:
+
+	/*! \constructor
+	 *
+	 * \param lvl level
+	 * \param key distributed grid key
+	 * \param offsets to move between levels
+	 *
+	 */
+	inline grid_dist_amr_key(size_t lvl,
+			          grid_dist_key_dx<dim> key)
+	:lvl(lvl),key(key)
+	{}
+
+	/*! \brief Return the grid key
+	 *
+	 * \return the distributed key
+	 *
+	 */
+	inline const grid_dist_key_dx<dim> & getKey() const
+	{
+		return key;
+	}
+
+	/*! \brief Return the grid key (as reference)
+	 *
+	 * \return the distributed key
+	 *
+	 */
+	inline grid_dist_key_dx<dim> & getKeyRef()
+	{
+		return key;
+	}
+
+
+	/*! \brief Return the level
+	 *
+	 * \return the level
+	 *
+	 */
+	inline size_t getLvl() const
+	{
+		return lvl;
+	}
+
+	/*! \brief Return the level
+	 *
+	 * \param lvl level to set
+	 *
+	 */
+	inline void setLvl(size_t lvl)
+	{
+		this->lvl = lvl;
+	}
+
+	/*! \brief Create a new key moving the old one
+	 *
+	 * \param s dimension id
+	 * \param s number of steps
+	 *
+	 * \return new key
+	 *
+	 */
+	inline grid_dist_amr_key<dim> moveSpace(size_t d,size_t s)
+	{
+		return grid_dist_amr_key<dim>(lvl,key.move(d,s));
+	}
+};
+
+
+
+
+#endif /* SRC_AMR_GRID_DIST_AMR_KEY_HPP_ */

src/Amr/grid_dist_amr_key_iterator.hpp

+/*
+ * grid_amr_dist_key_iterator.hpp
+ *
+ *  Created on: Sep 22, 2017
+ *      Author: i-bird
+ */
+
+#ifndef SRC_AMR_GRID_DIST_AMR_KEY_ITERATOR_HPP_
+#define SRC_AMR_GRID_DIST_AMR_KEY_ITERATOR_HPP_
+
+#include "Vector/map_vector.hpp"
+#include "Grid/Iterators/grid_dist_id_iterator.hpp"
+#include "grid_dist_amr_key.hpp"
+
+template<unsigned int dim, typename device_grid, typename device_sub_it, typename it_type = grid_dist_iterator<dim,device_grid,device_sub_it,FREE>>
+class grid_dist_amr_key_iterator
+{
+	//! Array of grid iterators
+	openfpm::vector<it_type> & git;
+
+	//! actual it type
+	struct actual_it
+	{
+		it_type & it;
+	};
+
+	//! Actual distributed grid iterator
+	it_type * a_it;
+
+	//! iterator pointer
+	size_t g_c;
+
+
+
+	/*! \brief from g_c increment g_c until you find a valid grid
+	 *
+	 */
+	void selectValidGrid()
+	{
+		// When the grid has size 0 potentially all the other informations are garbage
+		while (g_c < git.size() && git.get(g_c).isNext() == false ) g_c++;
+
+		// get the next grid iterator
+		if (g_c < git.size())
+		{
+			a_it = &git.get(g_c);
+		}
+	}
+
+public:
+
+	/*! \brief Constructor
+	 *
+	 * \param git vector of iterator
+	 *
+	 */
+	grid_dist_amr_key_iterator(openfpm::vector<it_type> & git)
+	:git(git),g_c(0)
+	{
+		a_it = &git.get(0);
+
+		selectValidGrid();
+	}
+
+
+	//! Destructor
+	~grid_dist_amr_key_iterator()
+	{
+	}
+
+
+	/*! \brief Get the next element
+	 *
+	 * \return the next grid_key
+	 *
+	 */
+	inline grid_dist_amr_key_iterator<dim,device_grid,device_sub_it,it_type> & operator++()
+	{
+		++(*a_it);
+
+		// check if a_it is at the end
+
+		if (a_it->isNext() == true)
+		{return *this;}
+		else
+		{
+			// switch to the new iterator
+			g_c++;
+
+			selectValidGrid();
+		}
+
+		return *this;
+	}
+
+	/*! \brief Is there a next point
+	 *
+	 * \return true is there is a next point
+	 *
+	 */
+	inline bool isNext()
+	{
+		return g_c < git.size();
+	}
+
+	/*! \brief Return the actual AMR grid iterator point
+	 *
+	 *
+	 */
+	inline grid_dist_amr_key<dim> get()
+	{
+		return grid_dist_amr_key<dim>(g_c,a_it->get());
+	}
+
+	/*! \brief Return the actual global grid position in the AMR struct in global
+	 *         coordinates
+	 *
+	 *
+	 */
+	inline grid_key_dx<dim> getGKey()
+	{
+		return git.get(g_c).getGKey(a_it->get());
+	}
+
+	/*! \brief Return the level at which we are
+	 *
+	 *
+	 */
+	inline size_t getLvl() const
+	{
+		return g_c;
+	}
+};
+
+
+#endif /* SRC_AMR_GRID_DIST_AMR_KEY_ITERATOR_HPP_ */

src/Amr/tests/amr_base_unit_tests.cpp

+/*
+ * amr_base_unit_test.cpp
+ *
+ *  Created on: Oct 5, 2017
+ *      Author: i-bird
+ */
+#define BOOST_TEST_DYN_LINK
+#include <boost/test/unit_test.hpp>
+
+#include "Grid/grid_dist_id.hpp"
+#include "Point_test.hpp"
+#include "Grid/tests/grid_dist_id_util_tests.hpp"
+
+BOOST_AUTO_TEST_SUITE( amr_grid_dist_id_test )
+
+
+BOOST_AUTO_TEST_CASE( grid_dist_id_amr )
+{
+	// Domain
+	Box<2,float> domain2({0.0,0.0},{1.0,1.0});
+
+	size_t sz[2] = {100,100};
+
+	// Ghost
+	Ghost<2,long int> g(1);
+
+	// periodicity
+	periodicity<2> pr = {{PERIODIC,PERIODIC}};
+
+	openfpm::vector<Box<2,long int>> C_draw;
+	C_draw.add(Box<2,long int>({20,20},{50,24}));
+	C_draw.add(Box<2,long int>({51,20},{60,24}));
+	C_draw.add(Box<2,long int>({61,20},{70,24}));
+	C_draw.add(Box<2,long int>({20,25},{24,66}));
+	C_draw.add(Box<2,long int>({15,67},{49,85}));
+	C_draw.add(Box<2,long int>({50,76},{70,81}));
+	C_draw.add(Box<2,long int>({30,25},{34,37}));
+	C_draw.add(Box<2,long int>({50,66},{70,70}));
+
+	size_t volume_key = 0;
+	for (size_t i = 0 ; i < C_draw.size() ; i++)
+	{
+		volume_key += Box<2,long int>(C_draw.get(i)).getVolumeKey();
+	}
+
+	// Distributed grid with id decomposition
+	grid_dist_id<2,float,Point_test<float>> g_dist(sz,domain2,g,pr,C_draw);
+
+	// fill with gkey
+
+	auto git = g_dist.getDomainIterator();
+	grid_sm<2,void> gs(sz);
+
+	size_t count = 0;
+
+	while (git.isNext())
+	{
+		auto key = git.get();
+		auto gkey = git.getGKey(key);
+
+		g_dist.template get<0>(key) = gs.LinId(gkey);
+
+		count++;
+
+		++git;
+	}
+
+	Vcluster<> & vcl = create_vcluster();
+
+	vcl.sum(count);
+	vcl.execute();
+
+	BOOST_REQUIRE_EQUAL(count,volume_key);
+
+	g_dist.ghost_get<0>();
+
+	// Check it is correct
+
+	bool check = true;
+	size_t check_count = 0;
+
+	auto git2 = g_dist.getDomainGhostIterator();
+	while (git2.isNext())
+	{
+		auto key = git2.get();
+		auto gkey = git2.getGKey(key);
+
+		float value = g_dist.template get<0>(key);
+
+		// check if the point is inside or outside the domain
+
+		for (size_t k = 0; k < C_draw.size() ; k++)
+		{
+			if (Box<2,long int>(C_draw.get(k)).isInside(gkey.toPoint()) == true)
+			{
+				check &= value == gs.LinId(gkey);
+
+				// get the gdb_ext
+				auto & gdb_ext = g_dist.getLocalGridsInfo();
+
+				for (size_t s = 0 ; s < gdb_ext.size() ; s++)
+				{
+					Box<2,long int> bx = gdb_ext.get(s).Dbox;
+					bx += gdb_ext.get(s).origin;
+					if (bx.isInside(gkey.toPoint()))
+					{
+						check_count++;
+						break;
+					}
+				}
+				break;
+			}
+		}
+
+		++git2;
+	}
+
+	vcl.sum(check_count);
+	vcl.execute();
+
+	BOOST_REQUIRE_EQUAL(check,true);
+	BOOST_REQUIRE(check_count >= volume_key);
+}
+
+BOOST_AUTO_TEST_CASE( amr_grid_dist_id_iterator_test_use_2D)
+{
+	// Domain
+	Box<2,float> domain({0.0,0.0},{1.0,1.0});
+
+#ifdef TEST_COVERAGE_MODE
+	long int k = 256*256*create_vcluster().getProcessingUnits();
+#else
+	long int k = 1024*1024*create_vcluster().getProcessingUnits();
+#endif
+	k = std::pow(k, 1/2.);
+
+	long int big_step = k / 30;
+	big_step = (big_step == 0)?1:big_step;
+	long int small_step = 21;
+
+	print_test( "AMR Testing 2D full grid k<=",k);
+
+	// 2D test
+	for ( ; k >= 2 ; k-= (k > 2*big_step)?big_step:small_step )
+	{
+		BOOST_TEST_CHECKPOINT( "AMR Testing 2D full grid k=" << k );
+
+		//! [Create and access a distributed grid]
+
+		// grid size
+		size_t sz[2];
+		sz[0] = k;
+		sz[1] = k;
+
+		// periodicity
+		periodicity<2> pr = {{PERIODIC,PERIODIC}};
+
+		// Ghost
+		Ghost<2,long int> g(1);
+
+		openfpm::vector<Box<2,long int>> bx_def;
+		bx_def.add(Box<2,long int>({0,0},{k-1,k-1}));
+
+		// Distributed grid with id decomposition
+		grid_dist_id<2, float, aggregate<double>> g_dist(sz,domain,g,pr,bx_def);
+
+		Test2D_core(g_dist,sz,k);
+	}
+}
+
+BOOST_AUTO_TEST_SUITE_END()

src/CMakeLists.txt

@@ -4,7 +4,9 @@ cmake_minimum_required(VERSION 3.8 FATAL_ERROR)
 ########################### Executables
 
 if(CUDA_FOUND)
-	set(CUDA_SOURCES Vector/cuda/vector_dist_gpu_MP_tests.cu 
+	set(CUDA_SOURCES 
+		Grid/tests/sgrid_dist_id_gpu_unit_tests.cu
+		Vector/cuda/vector_dist_gpu_MP_tests.cu 
 					 Vector/cuda/vector_dist_cuda_func_test.cu
 					 Decomposition/cuda/decomposition_cuda_tests.cu
 					 Vector/cuda/vector_dist_gpu_unit_tests.cu

src/Decomposition/Domain_icells_cart.hpp

@@ -156,7 +156,7 @@ class domain_icell_calculator
 		openfpm::vector_sparse_gpu<aggregate<unsigned int>> vs;
 		openfpm::vector_sparse_gpu<aggregate<unsigned int>> vsi;
 
-		vs.template getBackground<0>() = 0;
+		vs.template setBackground<0>(0);
 
 		// insert Domain cells
 
@@ -182,7 +182,7 @@ class domain_icell_calculator
 
 			CUDA_LAUNCH((insert_icell<dim>),ite,vsi.toKernel(),cld,ite.start,p2);
 
-			vsi.template flush<>(v_cl.getmgpuContext(),flust_type::FLUSH_ON_DEVICE);
+			vsi.template flush<>(v_cl.getmgpuContext(),flush_type::FLUSH_ON_DEVICE);
 		}
 
 		// calculate the number of kernel launch
@@ -209,8 +209,8 @@ class domain_icell_calculator
 
 			CUDA_LAUNCH(insert_remove_icell<dim>,ite,vs.toKernel(),vsi.toKernel(),cld,ite.start,p2);
 
-			vs.template flush<>(v_cl.getmgpuContext(),flust_type::FLUSH_ON_DEVICE);
-			vsi.flush_remove(v_cl.getmgpuContext(),flust_type::FLUSH_ON_DEVICE);
+			vs.template flush<>(v_cl.getmgpuContext(),flush_type::FLUSH_ON_DEVICE);
+			vsi.flush_remove(v_cl.getmgpuContext(),flush_type::FLUSH_ON_DEVICE);
 		}
 
 		vs.swapIndexVector(icells);

src/Grid/Iterators/grid_dist_id_iterator_dec.hpp

@@ -155,6 +155,56 @@ class grid_dist_id_iterator_dec
 	{
 	}
 
+	/*! \brief Return true if we point to a valid grid
+	 *
+	 * \return true if valid grid
+	 *
+	 */
+	inline bool isNextGrid()
+	{
+		return g_c < gdb_ext.size();
+	}
+
+	/*! \brief Return the index of the grid in which we are iterating
+	 *
+	 *
+	 */
+	inline size_t getGridId()
+	{
+		return g_c;
+	}
+
+	/*! \brief next grid
+	 *
+	 *
+	 */
+	inline void nextGrid()
+	{
+		g_c++;
+		selectValidGrid();
+	}
+
+	/*! \brief Return the actual pointed grid
+	 *
+	 * \return the grid index
+	 *
+	 */
+	inline Box<Decomposition::dims,size_t> getGridBox()
+	{
+		Box<Decomposition::dims,size_t> bx;
+
+		auto start = a_it.getStart();
+		auto stop = a_it.getStop();
+
+		for (int i = 0 ; i < Decomposition::dims ; i++)
+		{
+			bx.setHigh(i,stop.get(i));
+			bx.setLow(i,start.get(i));
+		}
+
+		return bx;
+	}
+
 	/*! \brief Get the next element
 	 *
 	 * \return the next grid_key

src/Grid/cuda/grid_dist_id_kernels.cuh

+/*
+ * grid_dist_id_kernels.cuh
+ *
+ *  Created on: Jun 25, 2019
+ *      Author: i-bird
+ */
+
+#ifndef GRID_DIST_ID_KERNELS_CUH_
+#define GRID_DIST_ID_KERNELS_CUH_
+
+
+template<typename grid_type, typename func_t,typename ... args_t>
+__global__ void grid_apply_functor(grid_type g, ite_gpu<grid_type::d> ite, func_t f, args_t ... args)
+{
+	f(g,ite,args...);
+}
+
+
+
+#endif /* GRID_DIST_ID_KERNELS_CUH_ */

src/Grid/grid_dist_id.hpp

@@ -23,8 +23,11 @@
 #include "grid_dist_id_comm.hpp"
 #include "HDF5_wr/HDF5_wr.hpp"
 #include "SparseGrid/SparseGrid.hpp"
+#ifdef __NVCC__
+#include "SparseGridGpu/SparseGridGpu.hpp"
+#include "cuda/grid_dist_id_kernels.cuh"
+#endif
 
-template <typename> struct Debug;
 
 //! Internal ghost box sent to construct external ghost box into the other processors
 template<unsigned int dim>
@@ -115,6 +118,9 @@ class grid_dist_id : public grid_dist_id_comm<dim,St,T,Decomposition,Memory,devi
 	//! Structure that divide the space into cells
 	CellDecomposer_sm<dim,St,shift<dim,St>> cd_sm;
 
+	//! size of the insert pool on gpu
+	size_t gpu_insert_pool_size;
+
 	//! Communicator class
 	Vcluster<> & v_cl;
 
@@ -553,11 +559,11 @@ class grid_dist_id : public grid_dist_id_comm<dim,St,T,Decomposition,Memory,devi
 						tmp.GDbox = box_int_recv.get(i).get(j).bx;
 						tmp.GDbox -= tmp.GDbox.getP1();
 						tmp.origin = output.getP1();
-						for (size_t i = 0 ; i < dim ; i++)
+						for (size_t s = 0 ; s < dim ; s++)
 						{
 							// we set an invalid box, there is no-domain
-							tmp.Dbox.setLow(i,0);
-							tmp.Dbox.setHigh(i,-1);
+							tmp.Dbox.setLow(s,0);
+							tmp.Dbox.setHigh(s,-1);
 						}
 						tmp.k = -1;
 						gdb_ext.add(tmp);
@@ -1185,6 +1191,19 @@ public:
 		{meta_copy<T>::meta_copy_(bv,loc_grid.get(i).getBackgroundValue());}
 	}
 
+	/*! \brief set the background value
+	 *
+	 * You can use this function make sense in case of sparse in case of dense
+	 * it does nothing
+	 *
+	 */
+	template<unsigned int p>
+	void setBackgroundValue(const typename boost::mpl::at<typename T::type,boost::mpl::int_<p>>::type & bv)
+	{
+		for (size_t i = 0 ; i < loc_grid.size() ; i++)
+		{loc_grid.get(i).template setBackgroundValue<p>(bv);}
+	}
+
 	/*! \brief Return the local total number of points inserted in the grid
 	 *
 	 * in case of dense grid it return the number of local points, in case of
@@ -1992,6 +2011,16 @@ public:
 		return loc_grid.get(v1.getSub()).remove_no_flush(v1.getKey());
 	}
 
+
+	template<typename ... v_reduce>
+	void flush(flush_type opt = flush_type::FLUSH_ON_HOST)
+	{
+		for (size_t i = 0 ; i < loc_grid.size() ; i++)
+		{
+			loc_grid.get(i).template flush<v_reduce ...>(v_cl.getmgpuContext(),opt);
+		}
+	}
+
 	/*! \brief remove an element in the grid
 	 *
 	 * In case of dense grid this function print a warning, in case of sparse
@@ -2654,6 +2683,73 @@ public:
 		return this->ig_box;
 	}
 
+#ifdef __NVCC__
+
+	/*! \brief Set the size of the gpu insert buffer pool
+	 *
+	 * \param size of the insert pool
+	 *
+	 */
+	void setInsertBuffer(size_t n_pool)
+	{
+		gpu_insert_pool_size = n_pool;
+	}
+
+	template<typename func_t,typename it_t, typename ... args_t>
+	void iterateGridGPU(it_t & it, args_t ... args)
+	{
+		while(it.isNextGrid())
+		{
+			Box<dim,size_t> b = it.getGridBox();
+
+			size_t i = it.getGridId();
+
+			auto ite = loc_grid.get(i).getGridGPUIterator(b.getKP1(),b.getKP2());
+
+			loc_grid.get(i).setGPUInsertBuffer(ite.nblocks(),gpu_insert_pool_size);
+			loc_grid.get(i).initializeGPUInsertBuffer();
+
+			grid_apply_functor<<<ite.wthr,ite.thr>>>(loc_grid.get(i).toKernel(),ite,func_t(),args...);
+
+			it.nextGrid();
+		}
+	}