Commit 2f4d1bae authored by incardon's avatar incardon
Browse files

Merge branch 'master' of git.mpi-cbg.de:openfpm/openfpm_pdata

parents 0235f075 37977db6
openfpm_data @ 9234d279
Subproject commit 1e67b751bf63de73c2d23da4793d71d1f662eaa6
Subproject commit 9234d279fdd0735a93b01389e3a47a298e320f52
openfpm_vcluster @ 2a56c313
Subproject commit 96b8c774751903afedeabdd4d79dae43fce85510
Subproject commit 2a56c313040d0fd816d396c3b75c651a894b5577
......@@ -13,7 +13,8 @@ if(CUDA_FOUND OR CUDA_ON_CPU OR HIP_FOUND)
Decomposition/cuda/decomposition_cuda_tests.cu
Vector/cuda/vector_dist_gpu_unit_tests.cu
Decomposition/cuda/Domain_icells_cart_unit_test.cu
Amr/tests/amr_base_gpu_unit_tests.cu)
Amr/tests/amr_base_gpu_unit_tests.cu
Grid/tests/grid_dist_id_unit_test.cu)
endif()
if(CMAKE_CXX_COMPILER_ID STREQUAL "Intel")
......
......@@ -20,6 +20,8 @@
#include "SparseGridGpu/encap_num.hpp"
#endif
#include "Grid/cuda/grid_dist_id_kernels.cuh"
template<unsigned int dim>
struct launch_insert_sparse_lambda_call
{
......@@ -183,6 +185,56 @@ struct launch_insert_sparse
}
};
template<unsigned int dim>
struct launch_set_dense
{
template<typename grid_type, typename ite_type, typename lambda_f2>
__device__ void operator()(grid_type & grid, ite_type itg, lambda_f2 f2)
{
#ifdef __NVCC__
printf("grid on GPU Dimension %d not implemented, yet\n",(int)dim);
#endif
}
};
template<>
struct launch_set_dense<2>
{
template<typename grid_type, typename ite_type, typename lambda_f2>
__device__ void operator()(grid_type & grid, ite_type itg, lambda_f2 f2)
{
#ifdef __NVCC__
GRID_ID_2_GLOBAL(itg);
auto obj = grid.get_o(key);
f2(obj,keyg.get(0),keyg.get(1));
#endif
}
};
template<>
struct launch_set_dense<3>
{
template<typename grid_type, typename ite_type, typename lambda_f2>
__device__ void operator()(grid_type & grid, ite_type itg, lambda_f2 f2)
{
#ifdef __NVCC__
GRID_ID_3_GLOBAL(itg);
auto obj = grid.get_o(key);
f2(obj,keyg.get(0),keyg.get(1),keyg.get(2));
#endif
}
};
template<bool is_free>
struct selvg
{
......
......@@ -77,28 +77,28 @@ class grid_dist_id_iterator_gpu
* \return itself
*
*/
grid_dist_id_iterator_gpu<Decomposition,deviceGrids> & operator=(const grid_dist_id_iterator_gpu<Decomposition,deviceGrids> & tmp)
{
g_c = tmp.g_c;
gdb_ext = tmp.gdb_ext;
// grid_dist_id_iterator_gpu<Decomposition,deviceGrids> & operator=(const grid_dist_id_iterator_gpu<Decomposition,deviceGrids> & tmp)
// {
// g_c = tmp.g_c;
// gdb_ext = tmp.gdb_ext;
start = tmp.start;
stop = tmp.stop;
loc_grids = tmp.loc_grids;
// start = tmp.start;
// stop = tmp.stop;
// loc_grids = tmp.loc_grids;
return *this;
}
// return *this;
// }
/*! \brief Copy constructor
*
* \param tmp iterator to copy
*
*/
grid_dist_id_iterator_gpu(const grid_dist_id_iterator_gpu<Decomposition,deviceGrids> & tmp)
:loc_grids(tmp.loc_grids)
{
this->operator=(tmp);
}
// grid_dist_id_iterator_gpu(const grid_dist_id_iterator_gpu<Decomposition,deviceGrids> & tmp)
// :loc_grids(tmp.loc_grids)
// {
// this->operator=(tmp);
// }
/*! \brief Constructor of the distributed grid iterator
*
......@@ -269,6 +269,7 @@ class grid_dist_id_iterator_gpu
}
}
/*! \brief Get the starting point of the sub-grid we are iterating
*
* \return the starting point
......
......@@ -8,6 +8,8 @@
#ifndef GRID_DIST_ID_KERNELS_CUH_
#define GRID_DIST_ID_KERNELS_CUH_
#include "config.h"
#ifdef CUDA_GPU
template<unsigned int dim>
......
......@@ -1768,6 +1768,32 @@ public:
#ifdef __NVCC__
/*! \brief set existing points in the grid
*
* \param f2 lambda function to set points
*/
template<typename lambda_t2>
void setPoints(lambda_t2 f2)
{
auto it = getGridIteratorGPU();
it.template launch<1>(launch_set_dense<dim>(),f2);
}
/*! \brief set point existing in the grid between start and stop
*
* \param start point
* \param stop point
* \param f2 lambda function to set points
*/
template<typename lambda_t2>
void setPoints(grid_key_dx<dim> k1, grid_key_dx<dim> k2, lambda_t2 f2)
{
auto it = getGridIteratorGPU(k1,k2);
it.template launch<0>(launch_set_dense<dim>(),f2);
}
/*! \brief Insert point in the grid
*
* \param f1 lambda function to insert point
......@@ -1779,7 +1805,7 @@ public:
auto it = getGridIteratorGPU();
it.setGPUInsertBuffer(1);
it.template launch<1>(launch_insert_sparse(),f1,f2);
it.template launch<0>(launch_insert_sparse(),f1,f2);
}
/*! \brief Insert point in the grid between start and stop
......@@ -3306,6 +3332,9 @@ using grid_dist_id_devg = grid_dist_id<dim,St,T,Decomposition,Memory,devg>;
template<unsigned int dim, typename St, typename T, typename Memory = CudaMemory, typename Decomposition = CartDecomposition<dim,St,CudaMemory,memory_traits_inte> >
using sgrid_dist_id_gpu = grid_dist_id<dim,St,T,Decomposition,Memory,SparseGridGpu<dim,T>>;
template<unsigned int dim, typename St, typename T, typename Memory = CudaMemory, typename Decomposition = CartDecomposition<dim,St,CudaMemory,memory_traits_inte> >
using grid_dist_id_gpu = grid_dist_id<dim,St,T,Decomposition,Memory,grid_gpu<dim,T>>;
template<unsigned int dim, typename St, typename T, typename Memory = CudaMemory, typename Decomposition = CartDecomposition<dim,St,CudaMemory,memory_traits_inte> >
using sgrid_dist_sid_gpu = grid_dist_id<dim,St,T,Decomposition,Memory,SparseGridGpu<dim,T,default_edge<dim>::type::value,default_edge<dim>::tb::value,int>>;
#endif
......
#define BOOST_TEST_DYN_LINK
#include <boost/test/unit_test.hpp>
#include "Point_test.hpp"
#include "Grid/grid_dist_id.hpp"
#include "data_type/aggregate.hpp"
extern void print_test_v(std::string test, size_t sz);
BOOST_AUTO_TEST_SUITE( grid_dist_id_test )
BOOST_AUTO_TEST_CASE( grid_dist_id_gpu_test )
{
// Test grid periodic
/* Box<3,float> domain({-1.0,-1.0,-1.0},{1.0,1.0,1.0});
Vcluster<> & v_cl = create_vcluster();
if ( v_cl.getProcessingUnits() > 32 )
{return;}
// grid size
size_t sz[3];
sz[0] = 32;
sz[1] = 32;
sz[2] = 32;
// Ghost
Ghost<3,long int> g(1);
// periodicity
periodicity<3> pr = {{PERIODIC,PERIODIC,PERIODIC}};
// Distributed grid with id decomposition
grid_dist_id_gpu<3, float, aggregate<float, float>> g_dist(sz,domain,g,pr);
Box<3,size_t> box({1,1,1},{30,30,30});
auto it = g_dist.getGridIterator(box.getKP1(),box.getKP2());
float c = 5.0;
typedef typename GetSetBlockType<decltype(g_dist)>::type BlockT;
g_dist.setPoints(box.getKP1(),box.getKP2(),
[c] __device__ (BlockT & data, int i, int j, int k)
{
data.template get<0>() = c + i*i + j*j + k*k;
}
);
*/
}
BOOST_AUTO_TEST_SUITE_END()
......@@ -259,7 +259,9 @@ template<unsigned int dim,
typename prop,
typename Decomposition = CartDecomposition<dim,St>,
typename Memory = HeapMemory,
template<typename> class layout_base = memory_traits_lin>
template<typename> class layout_base = memory_traits_lin,
typename vector_dist_pos = openfpm::vector<Point<dim, St>,Memory,layout_base>,
typename vector_dist_prop = openfpm::vector<prop,Memory,layout_base> >
class vector_dist : public vector_dist_comm<dim,St,prop,Decomposition,Memory,layout_base>,
#ifdef CUDA_GPU
private vector_dist_ker_list<vector_dist_ker<dim,St,prop,layout_base>>
......@@ -280,17 +282,17 @@ private:
//! Particle position vector, (It has 2 elements) the first has real particles assigned to a processor
//! the second element contain unassigned particles
openfpm::vector<Point<dim, St>,Memory,layout_base> v_pos;
vector_dist_pos v_pos;
//! Particle properties vector, (It has 2 elements) the first has real particles assigned to a processor
//! the second element contain unassigned particles
openfpm::vector<prop,Memory,layout_base> v_prp;
vector_dist_prop v_prp;
//! reordered v_pos buffer
openfpm::vector<prop,Memory,layout_base> v_prp_out;
vector_dist_prop v_prp_out;
//! reordered v_prp buffer
openfpm::vector<Point<dim, St>,Memory,layout_base> v_pos_out;
vector_dist_pos v_pos_out;
//! option used to create this vector
size_t opt = 0;
......@@ -2606,8 +2608,8 @@ public:
if ((opt & 0x0FFF0000) == CSV_WRITER)
{
// CSVWriter test
CSVWriter<openfpm::vector<Point<dim, St>,Memory,layout_base>,
openfpm::vector<prop,Memory,layout_base> > csv_writer;
CSVWriter<vector_dist_pos,
vector_dist_prop > csv_writer;
std::string output = std::to_string(out + "_" + std::to_string(v_cl.getProcessUnitID()) + std::to_string(".csv"));
......@@ -2622,8 +2624,8 @@ public:
ft = file_type::BINARY;
// VTKWriter for a set of points
VTKWriter<boost::mpl::pair<openfpm::vector<Point<dim, St>,Memory,layout_base>,
openfpm::vector<prop,Memory,layout_base>>,
VTKWriter<boost::mpl::pair<vector_dist_pos,
vector_dist_prop>,
VECTOR_POINTS> vtk_writer;
vtk_writer.add(v_pos,v_prp,g_m);
......@@ -2699,8 +2701,8 @@ public:
if ((opt & 0x0FFF0000) == CSV_WRITER)
{
// CSVWriter test
CSVWriter<openfpm::vector<Point<dim, St>,Memory,layout_base>,
openfpm::vector<prop,Memory,layout_base> > csv_writer;
CSVWriter<vector_dist_pos,
vector_dist_prop > csv_writer;
std::string output = std::to_string(out + "_" + std::to_string(v_cl.getProcessUnitID()) + "_" + std::to_string(iteration) + std::to_string(".csv"));
......@@ -2715,8 +2717,8 @@ public:
ft = file_type::BINARY;
// VTKWriter for a set of points
VTKWriter<boost::mpl::pair<openfpm::vector<Point<dim, St>,Memory,layout_base>,
openfpm::vector<prop,Memory,layout_base>>, VECTOR_POINTS> vtk_writer;
VTKWriter<boost::mpl::pair<vector_dist_pos,
vector_dist_prop>, VECTOR_POINTS> vtk_writer;
vtk_writer.add(v_pos,v_prp,g_m);
std::string output = std::to_string(out + "_" + std::to_string(v_cl.getProcessUnitID()) + "_" + std::to_string(iteration) + std::to_string(".vtk"));
......@@ -2785,7 +2787,7 @@ public:
* \return the particle position vector
*
*/
const openfpm::vector<Point<dim, St>,Memory,layout_base> & getPosVector() const
const vector_dist_pos & getPosVector() const
{
return v_pos;
}
......@@ -2795,7 +2797,7 @@ public:
* \return the particle position vector
*
*/
openfpm::vector<Point<dim, St>,Memory,layout_base> & getPosVector()
vector_dist_pos & getPosVector()
{
return v_pos;
}
......@@ -2805,7 +2807,7 @@ public:
* \return the particle property vector
*
*/
const openfpm::vector<prop,Memory,layout_base> & getPropVector() const
const vector_dist_prop & getPropVector() const
{
return v_prp;
}
......@@ -2815,7 +2817,7 @@ public:
* \return the particle property vector
*
*/
openfpm::vector<prop,Memory,layout_base> & getPropVector()
vector_dist_prop & getPropVector()
{
return v_prp;
}
......@@ -2825,7 +2827,7 @@ public:
* \return the particle position vector
*
*/
const openfpm::vector<Point<dim, St>,Memory,layout_base> & getPosVectorSort() const
const vector_dist_pos & getPosVectorSort() const
{
return v_pos_out;
}
......@@ -2835,7 +2837,7 @@ public:
* \return the particle position vector
*
*/
openfpm::vector<Point<dim, St>,Memory,layout_base> & getPosVectorSort()
vector_dist_pos & getPosVectorSort()
{
return v_pos_out;
}
......@@ -2845,7 +2847,7 @@ public:
* \return the particle property vector
*
*/
const openfpm::vector<prop,Memory,layout_base> & getPropVectorSort() const
const vector_dist_prop & getPropVectorSort() const
{
return v_prp_out;
}
......@@ -2855,7 +2857,7 @@ public:
* \return the particle property vector
*
*/
openfpm::vector<prop,Memory,layout_base> & getPropVectorSort()
vector_dist_prop & getPropVectorSort()
{
return v_prp_out;
}
......
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