...
 
Commits (2)
...@@ -217,6 +217,8 @@ install(FILES Vector/map_vector.hpp ...@@ -217,6 +217,8 @@ install(FILES Vector/map_vector.hpp
DESTINATION openfpm_data/include/Vector ) DESTINATION openfpm_data/include/Vector )
install(FILES Vector/cuda/map_vector_cuda_ker.cuh install(FILES Vector/cuda/map_vector_cuda_ker.cuh
Vector/cuda/map_vector_std_cuda_ker.cuh
Vector/cuda/map_vector_std_cuda.hpp
DESTINATION openfpm_data/include/Vector/cuda/ ) DESTINATION openfpm_data/include/Vector/cuda/ )
install(FILES util/multi_array_openfpm/array_openfpm.hpp install(FILES util/multi_array_openfpm/array_openfpm.hpp
......
...@@ -667,15 +667,16 @@ BOOST_AUTO_TEST_CASE (gpu_grid_test_se_class1) ...@@ -667,15 +667,16 @@ BOOST_AUTO_TEST_CASE (gpu_grid_test_se_class1)
BOOST_REQUIRE_EQUAL(dev_mem2[4],2); BOOST_REQUIRE_EQUAL(dev_mem2[4],2);
std::cout << "######### Testing error message #########" << std::endl; std::cout << "######### Testing error message #########" << std::endl;
dim3 wthr;
wthr.x = 32; ite_gpu<3> gr;
wthr.y = 1;
wthr.z = 1; gr.wthr.x = 32;
dim3 thr; gr.wthr.y = 1;
thr.x = 16; gr.wthr.z = 1;
thr.y = 1; gr.thr.x = 16;
thr.z = 1; gr.thr.y = 1;
CUDA_LAUNCH(test_se1_crash_gt2,wthr,thr,c3.toKernel(),c2.toKernel()); gr.thr.z = 1;
CUDA_LAUNCH(test_se1_crash_gt2,gr,c3.toKernel(),c2.toKernel());
std::cout << "######### End Testing error message #########" << std::endl; std::cout << "######### End Testing error message #########" << std::endl;
#endif #endif
......
...@@ -581,7 +581,7 @@ public: ...@@ -581,7 +581,7 @@ public:
* *
*/ */
inline grid_sm<N,T> & operator=(const grid_sm<N,T> & g) __device__ __host__ inline grid_sm<N,T> & operator=(const grid_sm<N,T> & g)
{ {
size_tot = g.size_tot; size_tot = g.size_tot;
......
...@@ -363,7 +363,7 @@ struct host_to_device_impl ...@@ -363,7 +363,7 @@ struct host_to_device_impl
kernel_type, kernel_type,
type_prp, type_prp,
layout_base, layout_base,
is_vector<typename mem_r_type::value_type>::value + 2*std::is_array<type_prp>::value + std::rank<type_prp>::value> (is_vector<typename mem_r_type::value_type>::value || is_vector_dist<typename mem_r_type::value_type>::value ) + 2*std::is_array<type_prp>::value + std::rank<type_prp>::value>
::template transform<Memory,mem_r_type>(static_cast<Memory *>(boost::fusion::at_c<boost::mpl::at<v_prp,boost::mpl::int_<T::value>>::type::value>(dst).mem), ::template transform<Memory,mem_r_type>(static_cast<Memory *>(boost::fusion::at_c<boost::mpl::at<v_prp,boost::mpl::int_<T::value>>::type::value>(dst).mem),
boost::fusion::at_c<boost::mpl::at<v_prp,boost::mpl::int_<T::value>>::type::value>(dst).mem_r, boost::fusion::at_c<boost::mpl::at<v_prp,boost::mpl::int_<T::value>>::type::value>(dst).mem_r,
start*sizeof(type_prp), start*sizeof(type_prp),
......
...@@ -434,6 +434,9 @@ public: ...@@ -434,6 +434,9 @@ public:
//! Type of the coordinate space (double float) //! Type of the coordinate space (double float)
typedef T stype; typedef T stype;
//!
typedef vector_pos_type internal_vector_pos_type;
/*! \brief Return the underlying grid information of the cell list /*! \brief Return the underlying grid information of the cell list
* *
* \return the grid infos * \return the grid infos
...@@ -969,14 +972,14 @@ public: ...@@ -969,14 +972,14 @@ public:
* \return An iterator across the neighborhood particles * \return An iterator across the neighborhood particles
* *
*/ */
template<unsigned int impl=NO_CHECK> inline CellNNIteratorRadius<dim,CellList<dim,T,Mem_type,transform>,impl> getNNIteratorRadius(size_t cell, T r_cut) template<unsigned int impl=NO_CHECK> inline CellNNIteratorRadius<dim,CellList<dim,T,Mem_type,transform,vector_pos_type>,impl> getNNIteratorRadius(size_t cell, T r_cut)
{ {
openfpm::vector<long int> & NNc = rcache[r_cut]; openfpm::vector<long int> & NNc = rcache[r_cut];
if (NNc.size() == 0) if (NNc.size() == 0)
{NNcalc_rad(r_cut,NNc,this->getCellBox(),this->getGrid());} {NNcalc_rad(r_cut,NNc,this->getCellBox(),this->getGrid());}
CellNNIteratorRadius<dim,CellList<dim,T,Mem_type,transform>,impl> cln(cell,NNc,*this); CellNNIteratorRadius<dim,CellList<dim,T,Mem_type,transform,vector_pos_type>,impl> cln(cell,NNc,*this);
return cln; return cln;
} }
......
...@@ -300,8 +300,8 @@ public: ...@@ -300,8 +300,8 @@ public:
getNNIteratorSym(size_t cell, getNNIteratorSym(size_t cell,
size_t pp, size_t pp,
size_t p, size_t p,
const openfpm::vector<Point<dim,typename CellBase::stype>> & pos, const typename CellBase::internal_vector_pos_type & pos,
const openfpm::vector<pos_v<dim,typename CellBase::stype>> & v) const openfpm::vector<pos_v<typename CellBase::internal_vector_pos_type>> & v)
{ {
CellNNIteratorSymM<dim,CellListM<dim,T,sh_byte,CellBase>,sh_byte,SYM,impl> cln(cell,pp,p,CellListM<dim,T,sh_byte,CellBase>::NNc_sym,*this,pos,v); CellNNIteratorSymM<dim,CellListM<dim,T,sh_byte,CellBase>,sh_byte,SYM,impl> cln(cell,pp,p,CellListM<dim,T,sh_byte,CellBase>::NNc_sym,*this,pos,v);
......
...@@ -265,9 +265,9 @@ template<unsigned int dim, typename T, typename CellS> void Test_cell_sM(SpaceBo ...@@ -265,9 +265,9 @@ template<unsigned int dim, typename T, typename CellS> void Test_cell_sM(SpaceBo
openfpm::vector<Point<dim,T>> phase1; openfpm::vector<Point<dim,T>> phase1;
openfpm::vector<Point<dim,T>> phase2; openfpm::vector<Point<dim,T>> phase2;
openfpm::vector<pos_v<dim,T>> phases; openfpm::vector<pos_v<openfpm::vector<Point<dim,T>>>> phases;
phases.add(pos_v<dim,T>(phase1)); phases.add(pos_v<openfpm::vector<Point<dim,T>>>(phase1));
phases.add(pos_v<dim,T>(phase2)); phases.add(pos_v<openfpm::vector<Point<dim,T>>>(phase2));
size_t id = 0; size_t id = 0;
......
...@@ -238,12 +238,12 @@ void populate_cell_list(openfpm::vector<Point<dim,T>,Memory,typename layout_base ...@@ -238,12 +238,12 @@ void populate_cell_list(openfpm::vector<Point<dim,T>,Memory,typename layout_base
* *
* *
*/ */
template<unsigned int dim, typename T> template<typename vector_pos_type>
struct pos_v struct pos_v
{ {
openfpm::vector<Point<dim,T>> & pos; vector_pos_type & pos;
pos_v(openfpm::vector<Point<dim,T>> & pos) pos_v(vector_pos_type & pos)
:pos(pos) :pos(pos)
{} {}
}; };
......
...@@ -39,10 +39,10 @@ class CellNNIteratorSymM : public CellNNIterator<dim,Cell,NNc_size,impl> ...@@ -39,10 +39,10 @@ class CellNNIteratorSymM : public CellNNIterator<dim,Cell,NNc_size,impl>
size_t p; size_t p;
//! Position of the particles p //! Position of the particles p
const openfpm::vector<Point<dim,typename Cell::stype>> & pos; const typename Cell::internal_vector_pos_type & pos;
//! Position of the particle p //! Position of the particle p
const openfpm::vector<pos_v<dim,typename Cell::stype>> & ps; const typename openfpm::vector<pos_v<typename Cell::internal_vector_pos_type>> & ps;
/*! Select the next valid element /*! Select the next valid element
* *
...@@ -91,8 +91,8 @@ public: ...@@ -91,8 +91,8 @@ public:
size_t p, size_t p,
const NNc_array<dim,NNc_size> & NNc, const NNc_array<dim,NNc_size> & NNc,
Cell & cl, Cell & cl,
const openfpm::vector<Point<dim,typename Cell::stype>> & pos, const typename Cell::internal_vector_pos_type & pos,
const openfpm::vector<pos_v<dim,typename Cell::stype>> & ps) const typename openfpm::vector<pos_v<typename Cell::internal_vector_pos_type>> & ps)
:CellNNIterator<dim,Cell,NNc_size,impl>(cell,NNc,cl),pp(pp),p(p),pos(pos),ps(ps) :CellNNIterator<dim,Cell,NNc_size,impl>(cell,NNc,cl),pp(pp),p(p),pos(pos),ps(ps)
{ {
selectValid(); selectValid();
......
...@@ -92,10 +92,10 @@ class CellNNIteratorSymM<dim,Cell,sh_byte,RUNTIME,impl> : public CellNNIterator< ...@@ -92,10 +92,10 @@ class CellNNIteratorSymM<dim,Cell,sh_byte,RUNTIME,impl> : public CellNNIterator<
//! index of the particle p //! index of the particle p
size_t p; size_t p;
const openfpm::vector<Point<dim,typename Cell::stype>> & pos; const typename Cell::internal_vector_pos_type & pos;
//! Position of the particles in the phases //! Position of the particles in the phases
const openfpm::vector<pos_v<dim,typename Cell::stype>> & ps; const openfpm::vector<pos_v<typename Cell::internal_vector_pos_type>> & ps;
/*! Select the next valid element /*! Select the next valid element
* *
...@@ -145,8 +145,8 @@ public: ...@@ -145,8 +145,8 @@ public:
const long int * NNc, const long int * NNc,
size_t NNc_size, size_t NNc_size,
Cell & cl, Cell & cl,
const openfpm::vector<Point<dim,typename Cell::stype>> & pos, const typename Cell::internal_vector_pos_type & pos,
const openfpm::vector<pos_v<dim,typename Cell::stype>> & ps) const openfpm::vector<pos_v<typename Cell::internal_vector_pos_type>> & ps)
:CellNNIterator<dim,Cell,RUNTIME,impl>(cell,NNc,NNc_size,cl),pp(pp),p(p),pos(pos),ps(ps) :CellNNIterator<dim,Cell,RUNTIME,impl>(cell,NNc,NNc_size,cl),pp(pp),p(p),pos(pos),ps(ps)
{} {}
......
...@@ -259,7 +259,8 @@ public: ...@@ -259,7 +259,8 @@ public:
* \return Return an iterator over the neighborhood particles * \return Return an iterator over the neighborhood particles
* *
*/ */
typename CellListType::SymNNIterator getNNIteratorCSRM(const openfpm::vector<Point<dim,typename CellListType::stype>> & pos ,const openfpm::vector<pos_v<dim,typename CellListType::stype>> & v) const typename CellListType::SymNNIterator getNNIteratorCSRM(const vector_pos_type & pos ,
const openfpm::vector<pos_v<vector_pos_type>> & v) const
{ {
if (dom_or_anom == 0) if (dom_or_anom == 0)
return typename CellListType::SymNNIterator(dom_cell.get(cid),CellListType::getV(*start),CellListType::getP(*start),NNc_sym.getPointer(),openfpm::math::pow(3,dim)/2+1,cli,pos,v); return typename CellListType::SymNNIterator(dom_cell.get(cid),CellListType::getV(*start),CellListType::getP(*start),NNc_sym.getPointer(),openfpm::math::pow(3,dim)/2+1,cli,pos,v);
......
...@@ -99,22 +99,8 @@ public: ...@@ -99,22 +99,8 @@ public:
* *
*/ */
CellList_gpu(const CellList_gpu<dim,T,Memory,transform,cnt_type,ids_type> & clg) CellList_gpu(const CellList_gpu<dim,T,Memory,transform,cnt_type,ids_type> & clg)
:CellDecomposer_sm<dim,T,transform>(clg)
{ {
cl_n = clg.cl_n; this->operator=(clg);
cells = clg.cells;
starts = clg.starts;
part_ids = clg.part_ids;
sorted_to_not_sorted = clg.sorted_to_not_sorted;
sorted_domain_particles_dg = clg.sorted_domain_particles_dg;
sorted_domain_particles_ids = clg.sorted_domain_particles_ids;
non_sorted_to_sorted = clg.non_sorted_to_sorted;
spacing_c = clg.spacing_c;
div_c = clg.div_c;
off = clg.off;
g_m = clg.g_m;
n_dec = clg.n_dec;
} }
/*! \brief Copy constructor from temporal /*! \brief Copy constructor from temporal
...@@ -123,24 +109,17 @@ public: ...@@ -123,24 +109,17 @@ public:
* *
*/ */
CellList_gpu(CellList_gpu<dim,T,Memory,transform,cnt_type,ids_type> && clg) CellList_gpu(CellList_gpu<dim,T,Memory,transform,cnt_type,ids_type> && clg)
:CellDecomposer_sm<dim,T,transform>(clg)
{ {
cl_n.swap(clg.cl_n); this->operator=(clg);
cells.swap(clg.cells);
starts.swap(clg.starts);
part_ids.swap(clg.part_ids);
sorted_to_not_sorted.swap(clg.sorted_to_not_sorted);
sorted_domain_particles_dg.swap(clg.sorted_domain_particles_dg);
sorted_domain_particles_ids.swap(clg.sorted_domain_particles_ids);
non_sorted_to_sorted.swap(clg.non_sorted_to_sorted);
spacing_c = clg.spacing_c;
div_c = clg.div_c;
off = clg.off;
g_m = clg.g_m;
n_dec = clg.n_dec;
} }
/*! \brief default constructor
*
*
*/
CellList_gpu()
{}
CellList_gpu(const Box<dim,T> & box, const size_t (&div)[dim], const size_t pad = 1) CellList_gpu(const Box<dim,T> & box, const size_t (&div)[dim], const size_t pad = 1)
{ {
Initialize(box,div,pad); Initialize(box,div,pad);
...@@ -489,6 +468,50 @@ public: ...@@ -489,6 +468,50 @@ public:
n_dec = clg.n_dec; n_dec = clg.n_dec;
clg.n_dec = n_dec_tmp; clg.n_dec = n_dec_tmp;
} }
CellList_gpu<dim,T,Memory,transform,cnt_type,ids_type> &
operator=(const CellList_gpu<dim,T,Memory,transform,cnt_type,ids_type> & clg)
{
*static_cast<CellDecomposer_sm<dim,T,transform> *>(this) = *static_cast<const CellDecomposer_sm<dim,T,transform> *>(&clg);
cl_n = clg.cl_n;
cells = clg.cells;
starts = clg.starts;
part_ids = clg.part_ids;
sorted_to_not_sorted = clg.sorted_to_not_sorted;
sorted_domain_particles_dg = clg.sorted_domain_particles_dg;
sorted_domain_particles_ids = clg.sorted_domain_particles_ids;
non_sorted_to_sorted = clg.non_sorted_to_sorted;
spacing_c = clg.spacing_c;
div_c = clg.div_c;
off = clg.off;
g_m = clg.g_m;
n_dec = clg.n_dec;
return *this;
}
CellList_gpu<dim,T,Memory,transform,cnt_type,ids_type> &
operator=(CellList_gpu<dim,T,Memory,transform,cnt_type,ids_type> && clg)
{
static_cast<CellDecomposer_sm<dim,T,transform> *>(this)->swap(*static_cast<CellDecomposer_sm<dim,T,transform> *>(&clg));
cl_n.swap(clg.cl_n);
cells.swap(clg.cells);
starts.swap(clg.starts);
part_ids.swap(clg.part_ids);
sorted_to_not_sorted.swap(clg.sorted_to_not_sorted);
sorted_domain_particles_dg.swap(clg.sorted_domain_particles_dg);
sorted_domain_particles_ids.swap(clg.sorted_domain_particles_ids);
non_sorted_to_sorted.swap(clg.non_sorted_to_sorted);
spacing_c = clg.spacing_c;
div_c = clg.div_c;
off = clg.off;
g_m = clg.g_m;
n_dec = clg.n_dec;
return *this;
}
}; };
......
...@@ -391,7 +391,7 @@ private: ...@@ -391,7 +391,7 @@ private:
Point<dim,T> xp = pos.template get<0>(i); Point<dim,T> xp = pos.template get<0>(i);
// Get the neighborhood of the particle // Get the neighborhood of the particle
NN_type NN = NNType<dim,T,CellListImpl,decltype(it),type,typename Mem_type::local_index_type>::get(it,pos,xp,i,cli,r_cut); auto NN = NNType<dim,T,CellListImpl,decltype(it),type,typename Mem_type::local_index_type>::get(it,pos,xp,i,cli,r_cut);
NNType<dim,T,CellListImpl,decltype(it),type,typename Mem_type::local_index_type>::add(i,dp); NNType<dim,T,CellListImpl,decltype(it),type,typename Mem_type::local_index_type>::add(i,dp);
while (NN.isNext()) while (NN.isNext())
...@@ -648,8 +648,8 @@ public: ...@@ -648,8 +648,8 @@ public:
*/ */
void Initialize(CellListImpl & cli, void Initialize(CellListImpl & cli,
T r_cut, T r_cut,
const openfpm::vector<Point<dim,T>> & pos, const vector_pos_type & pos,
const openfpm::vector<Point<dim,T>> & pos2, const vector_pos_type & pos2,
size_t g_m, size_t opt = VL_NON_SYMMETRIC) size_t g_m, size_t opt = VL_NON_SYMMETRIC)
{ {
Point<dim,T> spacing = cli.getCellBox().getP2(); Point<dim,T> spacing = cli.getCellBox().getP2();
......
...@@ -36,8 +36,8 @@ struct NNTypeM ...@@ -36,8 +36,8 @@ struct NNTypeM
* *
*/ */
static inline auto get(const PartIt & it, static inline auto get(const PartIt & it,
const openfpm::vector<Point<dim,T>> & pos, const typename CellListImpl::internal_vector_pos_type & pos,
const openfpm::vector<pos_v<dim,T>> & v, const openfpm::vector<pos_v<typename CellListImpl::internal_vector_pos_type>> & v,
Point<dim,T> & xp, Point<dim,T> & xp,
size_t pp, size_t pp,
size_t p, size_t p,
...@@ -74,8 +74,8 @@ struct NNTypeM<dim,T,CellListImpl,PartIt,VL_CRS_SYMMETRIC> ...@@ -74,8 +74,8 @@ struct NNTypeM<dim,T,CellListImpl,PartIt,VL_CRS_SYMMETRIC>
* *
*/ */
static inline auto get(const PartIt & it, static inline auto get(const PartIt & it,
const openfpm::vector<Point<dim,T>> & pos, const typename CellListImpl::internal_vector_pos_type & pos,
const openfpm::vector<pos_v<dim,T>> & v, const openfpm::vector<pos_v<typename CellListImpl::internal_vector_pos_type>> & v,
Point<dim,T> & xp, Point<dim,T> & xp,
size_t pp, size_t pp,
size_t p, size_t p,
...@@ -111,13 +111,13 @@ struct NNTypeM<dim,T,CellListImpl,PartIt,VL_SYMMETRIC> ...@@ -111,13 +111,13 @@ struct NNTypeM<dim,T,CellListImpl,PartIt,VL_SYMMETRIC>
* *
*/ */
static inline auto get(const PartIt & it, static inline auto get(const PartIt & it,
const openfpm::vector<Point<dim,T>> & pos, const typename CellListImpl::internal_vector_pos_type & pos,
const openfpm::vector<pos_v<dim,T>> & v, const openfpm::vector<pos_v<typename CellListImpl::internal_vector_pos_type>> & v,
Point<dim,T> & xp, Point<dim,T> & xp,
size_t pp, size_t pp,
size_t p, size_t p,
CellListImpl & cl, CellListImpl & cl,
T r_cut) -> decltype(cl.template getNNIteratorSym<NO_CHECK>(0,0,0,openfpm::vector<Point<dim,T>>(),openfpm::vector<pos_v<dim,T>>())) T r_cut) -> decltype(cl.template getNNIteratorSym<NO_CHECK>(0,0,0,typename CellListImpl::internal_vector_pos_type(),openfpm::vector<pos_v<typename CellListImpl::internal_vector_pos_type>>()))
{ {
return cl.template getNNIteratorSym<NO_CHECK>(cl.getCell(xp),pp,p,pos,v); return cl.template getNNIteratorSym<NO_CHECK>(cl.getCell(xp),pp,p,pos,v);
} }
...@@ -137,7 +137,8 @@ template<unsigned int dim, ...@@ -137,7 +137,8 @@ template<unsigned int dim,
unsigned int sh_byte , unsigned int sh_byte ,
typename CellListImpl=CellListM<dim,T,sh_byte>, typename CellListImpl=CellListM<dim,T,sh_byte>,
typename transform = shift<dim,T>, typename transform = shift<dim,T>,
typename VerletBase=VerletList<dim,T,Mem_fast<>,transform, size_t> > typename vector_pos_type = openfpm::vector<Point<dim,T>>,
typename VerletBase=VerletList<dim,T,Mem_fast<>,transform, vector_pos_type> >
class VerletListM : public VerletBase class VerletListM : public VerletBase
{ {
...@@ -159,8 +160,8 @@ class VerletListM : public VerletBase ...@@ -159,8 +160,8 @@ class VerletListM : public VerletBase
* \param opt options to create the verlet list like VL_SYMMETRIC or VL_NON_SYMMETRIC * \param opt options to create the verlet list like VL_SYMMETRIC or VL_NON_SYMMETRIC
* *
*/ */
inline void create(const openfpm::vector<Point<dim,T>> & pos, inline void create(const vector_pos_type & pos,
const openfpm::vector<pos_v<dim,T>> & pos2 , const openfpm::vector<pos_v<vector_pos_type>> & pos2 ,
const openfpm::vector<size_t> & dom, const openfpm::vector<size_t> & dom,
const openfpm::vector<subsub_lin<dim>> & anom, const openfpm::vector<subsub_lin<dim>> & anom,
size_t pp, size_t pp,
...@@ -197,8 +198,8 @@ class VerletListM : public VerletBase ...@@ -197,8 +198,8 @@ class VerletListM : public VerletBase
* \param opt options * \param opt options
* *
*/ */
template<typename NN_type, int type> inline void create_(const openfpm::vector<Point<dim,T>> & pos, template<typename NN_type, int type> inline void create_(const vector_pos_type & pos,
const openfpm::vector<pos_v<dim,T>> & pos2 , const openfpm::vector<pos_v<vector_pos_type>> & pos2 ,
const openfpm::vector<size_t> & dom, const openfpm::vector<size_t> & dom,
const openfpm::vector<subsub_lin<dim>> & anom, const openfpm::vector<subsub_lin<dim>> & anom,
size_t pp, size_t pp,
...@@ -209,7 +210,7 @@ class VerletListM : public VerletBase ...@@ -209,7 +210,7 @@ class VerletListM : public VerletBase
{ {
size_t end; size_t end;
auto it = PartItNN<type,dim,openfpm::vector<Point<dim,T>>,CellListImpl>::get(pos,dom,anom,cli,g_m,end); auto it = PartItNN<type,dim,vector_pos_type,CellListImpl>::get(pos,dom,anom,cli,g_m,end);
/* this->cl_n.resize(end); /* this->cl_n.resize(end);
this->cl_base.resize(end*this->slot); this->cl_base.resize(end*this->slot);
...@@ -263,7 +264,10 @@ public: ...@@ -263,7 +264,10 @@ public:
* \param opt options for the Verlet-list creation * \param opt options for the Verlet-list creation
* *
*/ */
void Initialize(CellListImpl & cli, size_t pp, T r_cut, const openfpm::vector<Point<dim,T>> & pos, const openfpm::vector<struct pos_v<dim,T>> & pos2, size_t g_m, size_t opt = VL_NON_SYMMETRIC) void Initialize(CellListImpl & cli, size_t pp,
T r_cut, const vector_pos_type & pos,
const openfpm::vector<struct pos_v<vector_pos_type>> & pos2, size_t g_m,
size_t opt = VL_NON_SYMMETRIC)
{ {
// this->cl_n.resize(g_m); // this->cl_n.resize(g_m);
// this->cl_base.resize(g_m*this->slot); // this->cl_base.resize(g_m*this->slot);
...@@ -328,9 +332,9 @@ public: ...@@ -328,9 +332,9 @@ public:
* \return an interator across the neighborhood particles * \return an interator across the neighborhood particles
* *
*/ */
template<unsigned int impl=NO_CHECK> inline VerletNNIteratorM<dim,VerletListM<dim,T,sh_byte,CellListImpl,transform,VerletBase>,sh_byte> getNNIterator(size_t part_id) template<unsigned int impl=NO_CHECK> inline VerletNNIteratorM<dim,VerletListM<dim,T,sh_byte,CellListImpl,transform,vector_pos_type,VerletBase>,sh_byte> getNNIterator(size_t part_id)
{ {
VerletNNIteratorM<dim,VerletListM<dim,T,sh_byte,CellListImpl,transform,VerletBase>,sh_byte> vln(part_id,*this); VerletNNIteratorM<dim,VerletListM<dim,T,sh_byte,CellListImpl,transform,vector_pos_type,VerletBase>,sh_byte> vln(part_id,*this);
return vln; return vln;
} }
......
...@@ -53,7 +53,11 @@ template<unsigned int dim, typename T> void create_particles_on_grid(grid_sm<dim ...@@ -53,7 +53,11 @@ template<unsigned int dim, typename T> void create_particles_on_grid(grid_sm<dim
* \param v vector of positions * \param v vector of positions
* *
*/ */
template<unsigned int dim, typename T> void create_particles_on_gridM(grid_sm<dim,void> & g_info, T r_cut, SpaceBox<dim,T> & box, openfpm::vector<pos_v<dim,T>> & v, CellListM<dim,T,2> & cl) template<unsigned int dim, typename T> void create_particles_on_gridM(grid_sm<dim,void> & g_info,
T r_cut,
SpaceBox<dim,T> & box,
openfpm::vector<pos_v<openfpm::vector<Point<dim,T>>>> & v,
CellListM<dim,T,2> & cl)
{ {
// Create a grid iterator // Create a grid iterator
grid_key_dx_iterator<dim> g_it(g_info); grid_key_dx_iterator<dim> g_it(g_info);
...@@ -304,10 +308,10 @@ template<unsigned int dim, typename T, typename VerS> void Verlet_list_sM(SpaceB ...@@ -304,10 +308,10 @@ template<unsigned int dim, typename T, typename VerS> void Verlet_list_sM(SpaceB
openfpm::vector<Point<dim,T>> ps2; openfpm::vector<Point<dim,T>> ps2;
openfpm::vector<Point<dim,T>> ps3; openfpm::vector<Point<dim,T>> ps3;
openfpm::vector<pos_v<dim,T>> pos2; openfpm::vector<pos_v<openfpm::vector<Point<dim,T>>>> pos2;
pos2.add(pos_v<dim,T>(ps1)); pos2.add(pos_v<openfpm::vector<Point<dim,T>>>(ps1));
pos2.add(pos_v<dim,T>(ps2)); pos2.add(pos_v<openfpm::vector<Point<dim,T>>>(ps2));
pos2.add(pos_v<dim,T>(ps3)); pos2.add(pos_v<openfpm::vector<Point<dim,T>>>(ps3));
CellListM<dim,T,2> cl; CellListM<dim,T,2> cl;
create_particles_on_gridM(ginfo,r_cut,box,pos2,cl); create_particles_on_gridM(ginfo,r_cut,box,pos2,cl);
...@@ -372,10 +376,10 @@ template<unsigned int dim, typename T, typename VerS> void Verlet_list_sM(SpaceB ...@@ -372,10 +376,10 @@ template<unsigned int dim, typename T, typename VerS> void Verlet_list_sM(SpaceB
openfpm::vector<Point<dim,T>> ps2; openfpm::vector<Point<dim,T>> ps2;
openfpm::vector<Point<dim,T>> ps3; openfpm::vector<Point<dim,T>> ps3;
openfpm::vector<pos_v<dim,T>> pos2; openfpm::vector<pos_v<openfpm::vector<Point<dim,T>>>> pos2;
pos2.add(pos_v<dim,T>(ps1)); pos2.add(pos_v<openfpm::vector<Point<dim,T>>>(ps1));
pos2.add(pos_v<dim,T>(ps2)); pos2.add(pos_v<openfpm::vector<Point<dim,T>>>(ps2));
pos2.add(pos_v<dim,T>(ps3)); pos2.add(pos_v<openfpm::vector<Point<dim,T>>>(ps3));
create_particles_on_gridM(ginfo,r_cut,box,pos2,cli); create_particles_on_gridM(ginfo,r_cut,box,pos2,cli);
......
...@@ -100,6 +100,88 @@ BOOST_AUTO_TEST_CASE( vector_cuda_to_kernel_recursive2 ) ...@@ -100,6 +100,88 @@ BOOST_AUTO_TEST_CASE( vector_cuda_to_kernel_recursive2 )
BOOST_REQUIRE_EQUAL(test,true); BOOST_REQUIRE_EQUAL(test,true);
} }
template<typename vv_rc,typename vector_output_type>
__global__ void kernel_recursive_check(vv_rc vvrc, vector_output_type vot)
{
int k = 0;
for (int i = 0 ; i < vvrc.size() ; i++)
{
for (int j = 0 ; j < vvrc.template get<1>(i).size() ; j++)
{
vot.template get<0>(k) = vvrc.template get<1>(i).template get<0>(j);
k++;
}
}
}
BOOST_AUTO_TEST_CASE( vector_cuda_to_kernel_recursive2_test_toKernel )
{
typedef openfpm::vector_gpu<aggregate<int,openfpm::vector_gpu<aggregate<long int>>>> test2_type;
typedef openfpm::vector_gpu<aggregate<int,openfpm::vector_gpu<aggregate<Box<2,float>>>>> test3_type;
test2_type tt2;
test3_type tt3;
tt2.add_no_device();
tt2.add_no_device();
tt2.add_no_device();
/* tt3.add();
tt3.add();
tt3.add();*/
tt2.template get<0>(0) = 80;
tt2.template get<1>(0).add();
tt2.template get<1>(0).template get<0>(0) = 500;
tt2.template get<0>(0) = 180;
tt2.template get<1>(0).add();
tt2.template get<1>(0).template get<0>(1) = 600;
tt2.template get<0>(0) = 280;;
tt2.template get<1>(0).add();
tt2.template get<1>(0).template get<0>(2) = 700;
tt2.template get<1>(0).template hostToDevice<0>();
tt2.template get<0>(1) = 10080;
tt2.template get<1>(1).add();
tt2.template get<1>(1).template get<0>(0) = 1500;
tt2.template get<0>(1) = 20080;
tt2.template get<1>(1).add();
tt2.template get<1>(1).template get<0>(1) = 1600;
tt2.template get<0>(1) = 30080;
tt2.template get<1>(1).add();
tt2.template get<1>(1).template get<0>(2) = 1700;
tt2.template get<1>(1).template hostToDevice<0>();
tt2.template get<0>(2) = 40080;
tt2.template get<1>(2).add();
tt2.template get<1>(2).template get<0>(0) = 2500;
tt2.template get<0>(2) = 50080;
tt2.template get<1>(2).add();
tt2.template get<1>(2).template get<0>(1) = 2600;
tt2.template get<0>(2) = 60080;
tt2.template get<1>(2).add();
tt2.template get<1>(2).template get<0>(2) = 2700;
tt2.template get<1>(2).template hostToDevice<0>();
tt2.template hostToDevice<1>();
openfpm::vector_gpu<aggregate<long int>> vg;
vg.resize(9);
kernel_recursive_check<<<1,1>>>(tt2.toKernel(),vg.toKernel());
vg.template deviceToHost<0>();
BOOST_REQUIRE_EQUAL(vg.template get<0>(0),500);
BOOST_REQUIRE_EQUAL(vg.template get<0>(1),600);
BOOST_REQUIRE_EQUAL(vg.template get<0>(2),700);
BOOST_REQUIRE_EQUAL(vg.template get<0>(3),1500);
BOOST_REQUIRE_EQUAL(vg.template get<0>(4),1600);
BOOST_REQUIRE_EQUAL(vg.template get<0>(5),1700);
BOOST_REQUIRE_EQUAL(vg.template get<0>(6),2500);
BOOST_REQUIRE_EQUAL(vg.template get<0>(7),2600);
BOOST_REQUIRE_EQUAL(vg.template get<0>(8),2700);
}
BOOST_AUTO_TEST_CASE( vector_cuda_to_cpu_operator_equal ) BOOST_AUTO_TEST_CASE( vector_cuda_to_cpu_operator_equal )
{ {
openfpm::vector_gpu<aggregate<int,int,double>> v1; openfpm::vector_gpu<aggregate<int,int,double>> v1;
......
...@@ -54,14 +54,14 @@ namespace openfpm ...@@ -54,14 +54,14 @@ namespace openfpm
typedef int yes_i_am_vector; typedef int yes_i_am_vector;
//! Type of the encapsulation memory parameter //! Type of the encapsulation memory parameter
typedef typename memory_traits_inte<T>::type layout_type; typedef typename memory_traits_inte<T_>::type layout_type;
//! Object container for T, it is the return type of get_o it return a object type trough //! Object container for T, it is the return type of get_o it return a object type trough
// you can access all the properties of T // you can access all the properties of T
typedef typename grid_cpu<1,T,CudaMemory,typename memory_traits_inte<T>::type>::container container; typedef typename grid_cpu<1,T_,CudaMemory,typename memory_traits_inte<T_>::type>::container container;
//! Type of the value the vector is storing //! Type of the value the vector is storing
typedef T value_type; typedef T_ value_type;
/*! \brief Return the size of the vector /*! \brief Return the size of the vector
* *
......
...@@ -178,6 +178,10 @@ namespace openfpm ...@@ -178,6 +178,10 @@ namespace openfpm
#include "map_vector_std.hpp" #include "map_vector_std.hpp"
#include "map_vector_std_ptr.hpp" #include "map_vector_std_ptr.hpp"
#ifdef CUDA_GPU
#include "cuda/map_vector_std_cuda.hpp"
#endif
/*! \brief Implementation of 1-D std::vector like structure /*! \brief Implementation of 1-D std::vector like structure
* *
* The layout is memory_traits_lin * The layout is memory_traits_lin
...@@ -424,6 +428,31 @@ namespace openfpm ...@@ -424,6 +428,31 @@ namespace openfpm
v_size++; v_size++;
} }
/*! \brief It insert a new emtpy object on the vector, eventually it reallocate the grid
*
* \warning It is not thread safe should not be used in multi-thread environment
* reallocation, work only on cpu
*
*/
void add_no_device()
{
#ifdef SE_CLASS2
check_valid(this,8);
#endif
//! Check if we have enough space
if (v_size >= base.size())
{
//! Resize the memory, double up the actual memory allocated for the vector
size_t sz[1];
non_zero_one(sz,2*base.size());
base.resize_no_device(sz);
}
//! increase the vector size
v_size++;
}
/*! \brief It insert a new object on the vector, eventually it reallocate the grid /*! \brief It insert a new object on the vector, eventually it reallocate the grid
* *
* \param v element to add * \param v element to add
...@@ -1794,7 +1823,7 @@ namespace openfpm ...@@ -1794,7 +1823,7 @@ namespace openfpm
template <typename T> using vector_std = vector<T, HeapMemory, typename memory_traits_lin<T>::type, memory_traits_lin, openfpm::grow_policy_double, STD_VECTOR>; template <typename T> using vector_std = vector<T, HeapMemory, typename memory_traits_lin<T>::type, memory_traits_lin, openfpm::grow_policy_double, STD_VECTOR>;
template<typename T> using vector_gpu = openfpm::vector<T,CudaMemory,typename memory_traits_inte<T>::type,memory_traits_inte>; template<typename T> using vector_gpu = openfpm::vector<T,CudaMemory,typename memory_traits_inte<T>::type,memory_traits_inte>;
template<typename T> using vector_gpu_single = openfpm::vector<T,CudaMemory,typename memory_traits_inte<T>::type,memory_traits_inte,openfpm::grow_policy_identity>; template<typename T> using vector_gpu_single = openfpm::vector<T,CudaMemory,typename memory_traits_inte<T>::type,memory_traits_inte,openfpm::grow_policy_identity>;
template<typename T> using vector_custd = vector<T, CudaMemory, typename memory_traits_inte<aggregate<T>>::type, memory_traits_inte, openfpm::grow_policy_double, STD_VECTOR>;
} }
#endif #endif
...@@ -46,4 +46,23 @@ template<typename T> ...@@ -46,4 +46,23 @@ template<typename T>
struct is_vector_native<T, typename Void< typename T::yes_i_am_vector_native>::type > : std::true_type struct is_vector_native<T, typename Void< typename T::yes_i_am_vector_native>::type > : std::true_type
{}; {};
///////////////////////////////////////////////////////////////////////////////////////////////////////////
template<typename T, typename Sfinae = void>
struct is_vector_dist: std::false_type {};
/*! \brief is_grid check if the type is a vector
*
* ### Example
*
* \snippet util.hpp Check is_vector
*
* return true if T is a vector
*
*/
template<typename T>
struct is_vector_dist<T, typename Void< typename T::yes_i_am_vector_dist>::type > : std::true_type
{};
#endif /* SRC_VECTOR_UTIL_HPP_ */ #endif /* SRC_VECTOR_UTIL_HPP_ */
...@@ -70,7 +70,7 @@ namespace openfpm ...@@ -70,7 +70,7 @@ namespace openfpm
// Definition of the box // Definition of the box
template<unsigned int dim , typename T> class Box; template<unsigned int dim , typename T> class Box;
template<template <typename> class layout_base, typename T, bool = is_vector_native<T>::value> template<template <typename> class layout_base, typename T, int = is_vector_native<T>::value + 2*is_vector_dist<T>::value >
struct toKernel_transform; struct toKernel_transform;
template<template <typename> class layout_base, typename T, typename ... args> template<template <typename> class layout_base, typename T, typename ... args>
...@@ -100,22 +100,24 @@ struct apply_transform ...@@ -100,22 +100,24 @@ struct apply_transform
typedef typename apply_trasform_impl<layout_base,T,typename T::type>::type type; typedef typename apply_trasform_impl<layout_base,T,typename T::type>::type type;
}; };
template<template <typename> class layout_base, typename T, bool > /////////////////////////////////////////////// TRANSFORMER NODE /////////////////////////////////////////////////
struct toKernel_transform
template<template <typename> class layout_base, typename T >
struct toKernel_transform<layout_base,T,0>
{ {
typedef T type; typedef T type;
}; };
template<template <typename> class layout_base, typename T> template<template <typename> class layout_base, typename T>
struct toKernel_transform<layout_base,T,true> struct toKernel_transform<layout_base,T,1>
{ {
typedef typename apply_transform<layout_base,typename T::value_type>::type aggr; typedef typename apply_transform<layout_base,typename T::value_type>::type aggr;
typedef openfpm::vector_gpu_ker<aggr,layout_base> type; typedef openfpm::vector_gpu_ker<aggr,layout_base> type;
}; };
/////////////////////////////////////////////////// KNOWN TYPE SPECIALIZATION ////////////////////////////////// /////////////////////////////////////////////////// KNOWN TYPE SPECIALIZATION TERMINATORS //////////////////////
template<template <typename> class layout_base,typename T, typename ... args> template<template <typename> class layout_base,typename T, typename ... args>
struct aggregate_or_known_type<layout_base,T,2,args ...> struct aggregate_or_known_type<layout_base,T,2,args ...>
......