...
 
Commits (2)
......@@ -217,6 +217,8 @@ install(FILES Vector/map_vector.hpp
DESTINATION openfpm_data/include/Vector )
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/ )
install(FILES util/multi_array_openfpm/array_openfpm.hpp
......
......@@ -667,15 +667,16 @@ BOOST_AUTO_TEST_CASE (gpu_grid_test_se_class1)
BOOST_REQUIRE_EQUAL(dev_mem2[4],2);
std::cout << "######### Testing error message #########" << std::endl;
dim3 wthr;
wthr.x = 32;
wthr.y = 1;
wthr.z = 1;
dim3 thr;
thr.x = 16;
thr.y = 1;
thr.z = 1;
CUDA_LAUNCH(test_se1_crash_gt2,wthr,thr,c3.toKernel(),c2.toKernel());
ite_gpu<3> gr;
gr.wthr.x = 32;
gr.wthr.y = 1;
gr.wthr.z = 1;
gr.thr.x = 16;
gr.thr.y = 1;
gr.thr.z = 1;
CUDA_LAUNCH(test_se1_crash_gt2,gr,c3.toKernel(),c2.toKernel());
std::cout << "######### End Testing error message #########" << std::endl;
#endif
......
......@@ -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;
......
......@@ -363,7 +363,7 @@ struct host_to_device_impl
kernel_type,
type_prp,
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),
boost::fusion::at_c<boost::mpl::at<v_prp,boost::mpl::int_<T::value>>::type::value>(dst).mem_r,
start*sizeof(type_prp),
......
......@@ -434,6 +434,9 @@ public:
//! Type of the coordinate space (double float)
typedef T stype;
//!
typedef vector_pos_type internal_vector_pos_type;
/*! \brief Return the underlying grid information of the cell list
*
* \return the grid infos
......@@ -969,14 +972,14 @@ public:
* \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];
if (NNc.size() == 0)
{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;
}
......
......@@ -300,8 +300,8 @@ public:
getNNIteratorSym(size_t cell,
size_t pp,
size_t p,
const openfpm::vector<Point<dim,typename CellBase::stype>> & pos,
const openfpm::vector<pos_v<dim,typename CellBase::stype>> & v)
const typename CellBase::internal_vector_pos_type & pos,
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);
......
......@@ -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>> phase2;
openfpm::vector<pos_v<dim,T>> phases;
phases.add(pos_v<dim,T>(phase1));
phases.add(pos_v<dim,T>(phase2));
openfpm::vector<pos_v<openfpm::vector<Point<dim,T>>>> phases;
phases.add(pos_v<openfpm::vector<Point<dim,T>>>(phase1));
phases.add(pos_v<openfpm::vector<Point<dim,T>>>(phase2));
size_t id = 0;
......
......@@ -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
{
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)
{}
};
......
......@@ -39,10 +39,10 @@ class CellNNIteratorSymM : public CellNNIterator<dim,Cell,NNc_size,impl>
size_t 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
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
*
......@@ -91,8 +91,8 @@ public:
size_t p,
const NNc_array<dim,NNc_size> & NNc,
Cell & cl,
const openfpm::vector<Point<dim,typename Cell::stype>> & pos,
const openfpm::vector<pos_v<dim,typename Cell::stype>> & ps)
const typename Cell::internal_vector_pos_type & pos,
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)
{
selectValid();
......
......@@ -92,10 +92,10 @@ class CellNNIteratorSymM<dim,Cell,sh_byte,RUNTIME,impl> : public CellNNIterator<
//! index of the particle 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
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
*
......@@ -145,8 +145,8 @@ public:
const long int * NNc,
size_t NNc_size,
Cell & cl,
const openfpm::vector<Point<dim,typename Cell::stype>> & pos,
const openfpm::vector<pos_v<dim,typename Cell::stype>> & ps)
const typename Cell::internal_vector_pos_type & pos,
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)
{}
......
......@@ -259,7 +259,8 @@ public:
* \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)
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:
*
*/
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;
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;
this->operator=(clg);
}
/*! \brief Copy constructor from temporal
......@@ -123,24 +109,17 @@ public:
*
*/
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);
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;
this->operator=(clg);
}
/*! \brief default constructor
*
*
*/
CellList_gpu()
{}
CellList_gpu(const Box<dim,T> & box, const size_t (&div)[dim], const size_t pad = 1)
{
Initialize(box,div,pad);
......@@ -489,6 +468,50 @@ public:
n_dec = clg.n_dec;
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:
Point<dim,T> xp = pos.template get<0>(i);
// 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);
while (NN.isNext())
......@@ -648,8 +648,8 @@ public:
*/
void Initialize(CellListImpl & cli,
T r_cut,
const openfpm::vector<Point<dim,T>> & pos,
const openfpm::vector<Point<dim,T>> & pos2,
const vector_pos_type & pos,
const vector_pos_type & pos2,
size_t g_m, size_t opt = VL_NON_SYMMETRIC)
{
Point<dim,T> spacing = cli.getCellBox().getP2();
......
......@@ -36,8 +36,8 @@ struct NNTypeM
*
*/
static inline auto get(const PartIt & it,
const openfpm::vector<Point<dim,T>> & pos,
const openfpm::vector<pos_v<dim,T>> & v,
const typename CellListImpl::internal_vector_pos_type & pos,
const openfpm::vector<pos_v<typename CellListImpl::internal_vector_pos_type>> & v,
Point<dim,T> & xp,
size_t pp,
size_t p,
......@@ -74,8 +74,8 @@ struct NNTypeM<dim,T,CellListImpl,PartIt,VL_CRS_SYMMETRIC>
*
*/
static inline auto get(const PartIt & it,
const openfpm::vector<Point<dim,T>> & pos,
const openfpm::vector<pos_v<dim,T>> & v,
const typename CellListImpl::internal_vector_pos_type & pos,
const openfpm::vector<pos_v<typename CellListImpl::internal_vector_pos_type>> & v,
Point<dim,T> & xp,
size_t pp,
size_t p,
......@@ -111,13 +111,13 @@ struct NNTypeM<dim,T,CellListImpl,PartIt,VL_SYMMETRIC>
*
*/
static inline auto get(const PartIt & it,
const openfpm::vector<Point<dim,T>> & pos,
const openfpm::vector<pos_v<dim,T>> & v,
const typename CellListImpl::internal_vector_pos_type & pos,
const openfpm::vector<pos_v<typename CellListImpl::internal_vector_pos_type>> & v,
Point<dim,T> & xp,
size_t pp,
size_t p,
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);
}
......@@ -137,7 +137,8 @@ template<unsigned int dim,
unsigned int sh_byte ,
typename CellListImpl=CellListM<dim,T,sh_byte>,
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
{
......@@ -159,8 +160,8 @@ class VerletListM : public VerletBase
* \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,
const openfpm::vector<pos_v<dim,T>> & pos2 ,
inline void create(const vector_pos_type & pos,
const openfpm::vector<pos_v<vector_pos_type>> & pos2 ,
const openfpm::vector<size_t> & dom,
const openfpm::vector<subsub_lin<dim>> & anom,
size_t pp,
......@@ -197,8 +198,8 @@ class VerletListM : public VerletBase
* \param opt options
*
*/
template<typename NN_type, int type> inline void create_(const openfpm::vector<Point<dim,T>> & pos,
const openfpm::vector<pos_v<dim,T>> & pos2 ,
template<typename NN_type, int type> inline void create_(const vector_pos_type & pos,
const openfpm::vector<pos_v<vector_pos_type>> & pos2 ,
const openfpm::vector<size_t> & dom,
const openfpm::vector<subsub_lin<dim>> & anom,
size_t pp,
......@@ -209,7 +210,7 @@ class VerletListM : public VerletBase
{
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_base.resize(end*this->slot);
......@@ -263,7 +264,10 @@ public:
* \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_base.resize(g_m*this->slot);
......@@ -328,9 +332,9 @@ public:
* \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;
}
......
......@@ -53,7 +53,11 @@ template<unsigned int dim, typename T> void create_particles_on_grid(grid_sm<dim
* \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
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
openfpm::vector<Point<dim,T>> ps2;
openfpm::vector<Point<dim,T>> ps3;
openfpm::vector<pos_v<dim,T>> pos2;
pos2.add(pos_v<dim,T>(ps1));
pos2.add(pos_v<dim,T>(ps2));
pos2.add(pos_v<dim,T>(ps3));
openfpm::vector<pos_v<openfpm::vector<Point<dim,T>>>> pos2;
pos2.add(pos_v<openfpm::vector<Point<dim,T>>>(ps1));
pos2.add(pos_v<openfpm::vector<Point<dim,T>>>(ps2));
pos2.add(pos_v<openfpm::vector<Point<dim,T>>>(ps3));
CellListM<dim,T,2> 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
openfpm::vector<Point<dim,T>> ps2;
openfpm::vector<Point<dim,T>> ps3;
openfpm::vector<pos_v<dim,T>> pos2;
pos2.add(pos_v<dim,T>(ps1));
pos2.add(pos_v<dim,T>(ps2));
pos2.add(pos_v<dim,T>(ps3));
openfpm::vector<pos_v<openfpm::vector<Point<dim,T>>>> pos2;
pos2.add(pos_v<openfpm::vector<Point<dim,T>>>(ps1));
pos2.add(pos_v<openfpm::vector<Point<dim,T>>>(ps2));
pos2.add(pos_v<openfpm::vector<Point<dim,T>>>(ps3));
create_particles_on_gridM(ginfo,r_cut,box,pos2,cli);
......
......@@ -100,6 +100,88 @@ BOOST_AUTO_TEST_CASE( vector_cuda_to_kernel_recursive2 )
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 )
{
openfpm::vector_gpu<aggregate<int,int,double>> v1;
......
......@@ -54,14 +54,14 @@ namespace openfpm
typedef int yes_i_am_vector;
//! 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
// 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
typedef T value_type;
typedef T_ value_type;
/*! \brief Return the size of the vector
*
......
......@@ -178,6 +178,10 @@ namespace openfpm
#include "map_vector_std.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
*
* The layout is memory_traits_lin
......@@ -424,6 +428,31 @@ namespace openfpm
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
*
* \param v element to add
......@@ -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_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_custd = vector<T, CudaMemory, typename memory_traits_inte<aggregate<T>>::type, memory_traits_inte, openfpm::grow_policy_double, STD_VECTOR>;
}
#endif
......@@ -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
{};
///////////////////////////////////////////////////////////////////////////////////////////////////////////
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_ */
......@@ -70,7 +70,7 @@ namespace openfpm
// Definition of the 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;
template<template <typename> class layout_base, typename T, typename ... args>
......@@ -100,22 +100,24 @@ struct apply_transform
typedef typename apply_trasform_impl<layout_base,T,typename T::type>::type type;
};
template<template <typename> class layout_base, typename T, bool >
struct toKernel_transform
/////////////////////////////////////////////// TRANSFORMER NODE /////////////////////////////////////////////////
template<template <typename> class layout_base, typename T >
struct toKernel_transform<layout_base,T,0>
{
typedef T type;
};
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 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>
struct aggregate_or_known_type<layout_base,T,2,args ...>
......