Commit 804ad631 authored by incardon's avatar incardon

merging with release 2.0.0

parents 47732113 33a5ec1e
......@@ -34,7 +34,7 @@ mkdir openfpm_data/src/config
git clone https://git.mpi-cbg.de/openfpm/openfpm_devices.git openfpm_devices
cd openfpm_devices
git checkout GPU_test
git checkout master
cd ..
cd "$workspace/openfpm_data"
......
......@@ -226,6 +226,9 @@ do
se_class3)
conf_options="$conf_options -DSE_CLASS3=ON"
;;
test_coverage)
conf_options="$conf_options -DTEST_COVERAGE=ON"
;;
gpu)
if [ x"$CXX" == x"" ]; then
conf_options="$conf_options"
......@@ -466,6 +469,9 @@ do
boost)
conf_options="$conf_options -DBOOST_ROOT=$ac_optarg"
;;
mpivendor)
conf_options="$conf_options -DMPI_VENDOR=$ac_optarg"
;;
*) ac_unrecognized_opts="$ac_unrecognized_opts$ac_unrecognized_sep--with-$ac_useropt_orig"
ac_unrecognized_sep=', ';;
esac
......
......@@ -12,12 +12,18 @@ add_executable(mem_map ../../openfpm_devices/src/Memleak_check.cpp main.cpp util
if ( CMAKE_COMPILER_IS_GNUCC )
target_compile_options(mem_map PRIVATE "-Wno-deprecated-declarations")
if (TEST_COVERAGE)
target_compile_options(mem_map PRIVATE $<$<COMPILE_LANGUAGE:CXX>: -fprofile-arcs -ftest-coverage>)
endif()
endif()
###########################
if (CUDA_FOUND)
target_compile_options(mem_map PUBLIC $<$<COMPILE_LANGUAGE:CUDA>:-Xcudafe "--display_error_number --diag_suppress=611 --diag_suppress=2885 --diag_suppress=2886 --diag_suppress=2887 --diag_suppress=2888 --diag_suppress=186 --diag_suppress=111" --expt-extended-lambda>)
if (TEST_COVERAGE)
target_compile_options(mem_map PRIVATE $<$<COMPILE_LANGUAGE:CUDA>: -Xcompiler "-fprofile-arcs -ftest-coverage">)
endif()
endif()
target_include_directories (mem_map PUBLIC ${CUDA_INCLUDE_DIRS})
......@@ -31,12 +37,19 @@ target_include_directories (mem_map PUBLIC ${Boost_INCLUDE_DIRS})
target_link_libraries(mem_map ${Boost_LIBRARIES})
target_link_libraries(mem_map -L${LIBHILBERT_LIBRARY_DIRS} ${LIBHILBERT_LIBRARIES})
if (TEST_COVERAGE)
target_link_libraries(mem_map -lgcov)
endif()
# Request that particles be built with -std=c++11
# As this is a public compile feature anything that links to particles
# will also build with -std=c++11
target_compile_features(mem_map PUBLIC cxx_std_11)
target_link_libraries(mem_map ${MPI_C_LIBRARIES})
target_link_libraries(mem_map m)
if (NOT APPLE)
target_link_libraries(mem_map rt)
endif ()
install(FILES Grid/comb.hpp
Grid/copy_grid_fast.hpp
......@@ -61,7 +74,8 @@ install(FILES Grid/cuda/cuda_grid_gpu_funcs.cuh
install(FILES data_type/aggregate.hpp
DESTINATION openfpm_data/include/data_type )
install(FILES Graph/map_graph.hpp
install(FILES Graph/CartesianGraphFactory.hpp
Graph/map_graph.hpp
DESTINATION openfpm_data/include/Graph)
install(FILES Point_test.hpp
......@@ -203,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
......@@ -222,6 +238,7 @@ install(FILES util/cuda/scan_cuda.cuh
DESTINATION openfpm_data/include/util/cuda )
install(FILES util/cuda/moderngpu/context.hxx
util/cuda/moderngpu/context_reduced.hxx
util/cuda/moderngpu/cpp11.hxx
util/cuda/moderngpu/cta_load_balance.hxx
util/cuda/moderngpu/cta_merge.hxx
......
......@@ -230,7 +230,7 @@ public:
* \tparam Graph Graph
* \tparam pos Array of properties
*/
template<int i, typename p, typename Graph, int ... pos>
template<unsigned int dim, int i, typename p, typename Graph, int ... pos>
struct fill_prop_by_type
{
......@@ -254,8 +254,8 @@ struct fill_prop_by_type
* \tparam Graph Graph
* \tparam pos Array of properties
*/
template<typename p, typename Graph, int ... pos>
struct fill_prop_by_type<0, p, Graph, pos...>
template<unsigned int dim, typename p, typename Graph, int ... pos>
struct fill_prop_by_type<dim, 0, p, Graph, pos...>
{
enum
{
......@@ -286,7 +286,7 @@ public:
* \return the constructed graph
*
*/
static Graph construct(const size_t (& sz)[dim], Box<dim,T> dom, const size_t(& bc)[dim])
static Graph construct(const size_t (& sz)[dim], Box<dim,T> & dom, const size_t(& bc)[dim])
{
// Calculate the size of the hyper-cubes on each dimension
T szd[dim];
......@@ -335,7 +335,7 @@ public:
// vertex spatial properties functor
fill_prop<dim, lin_id, T, decltype(gp.vertex(g.LinId(key))), typename to_boost_vmpl<pos...>::type, fill_prop_by_type<sizeof...(pos), p, Graph, pos...>::value> flp(obj, szd, key, g, dom);
fill_prop<dim, lin_id, T, decltype(gp.vertex(g.LinId(key))), typename to_boost_vmpl<pos...>::type, fill_prop_by_type<dim,sizeof...(pos), p, Graph, pos...>::value> flp(obj, szd, key, g, dom);
// fill properties
......@@ -407,7 +407,7 @@ public:
* \return the constructed graph
*
*/
static Graph construct(const size_t ( & sz)[dim], Box<dim,T> dom, const size_t(& bc)[dim])
static Graph construct(const size_t ( & sz)[dim], Box<dim,T> & dom, const size_t(& bc)[dim])
{
// Calculate the size of the hyper-cubes on each dimension
......@@ -456,7 +456,11 @@ public:
// vertex spatial properties functor
<<<<<<< HEAD
fill_prop<dim, lin_id, T, decltype(gp.vertex(g.LinId(key))), typename to_boost_vmpl<pos...>::type, fill_prop_by_type<sizeof...(pos), p, Graph, pos...>::value> flp(obj, szd, key, g, dom);
=======
fill_prop<dim, lin_id, T, decltype(gp.vertex(g.LinId(key))), typename to_boost_vmpl<pos...>::type, fill_prop_by_type<dim,sizeof...(pos), p, Graph, pos...>::value> flp(obj, szd, key, g, dom);
>>>>>>> 33a5ec1ebc3fc47dd5def0a60f7c352ba810d247
// fill properties
......@@ -537,7 +541,11 @@ public:
*
*/
template<int se, int id_prp, typename T, unsigned int dim_c, int ... pos>
<<<<<<< HEAD
static Graph construct(const size_t (&sz)[dim], Box<dim, T> dom, const size_t (& bc)[dim])
=======
static Graph construct(const size_t (&sz)[dim], Box<dim, T> & dom, const size_t (& bc)[dim])
>>>>>>> 33a5ec1ebc3fc47dd5def0a60f7c352ba810d247
{
return Graph_constructor_impl<dim, id_prp, Graph, se, T, dim_c, pos...>::construct(sz, dom, bc);
}
......
......@@ -487,7 +487,7 @@ public:
*
*/
template <unsigned int p, typename r_type=decltype(boost::fusion::at_c<p>(data_c))>
inline r_type get()
__device__ __host__ inline r_type get()
{
#ifdef SE_CLASS2
check_valid(&boost::fusion::at_c<p>(data_c),sizeof(typename boost::mpl::at<type,boost::mpl::int_<p>>::type));
......@@ -501,7 +501,7 @@ public:
*
*/
template <unsigned int p, typename r_type=decltype(boost::fusion::at_c<p>(data_c))>
inline const r_type get() const
__device__ __host__ inline const r_type get() const
{
#ifdef SE_CLASS2
check_valid(&boost::fusion::at_c<p>(data_c),sizeof(typename boost::mpl::at<type,boost::mpl::int_<p>>::type));
......@@ -516,7 +516,8 @@ public:
* \param ele value to set
*
*/
template <unsigned int p> inline void set(decltype(boost::fusion::at_c<p>(data_c)) & ele)
template <unsigned int p> inline
__device__ __host__ void set(decltype(boost::fusion::at_c<p>(data_c)) & ele)
{
#ifdef SE_CLASS2
check_valid(&boost::fusion::at_c<p>(data_c),sizeof(typename boost::mpl::at<type,boost::mpl::int_<p>>::type));
......@@ -534,7 +535,8 @@ public:
* \return itself
*
*/
template<unsigned int dim2> inline encapc<dim,T,Mem> & set(const encapc<dim2,T,Mem> & ec)
template<unsigned int dim2>
__device__ __host__ inline encapc<dim,T,Mem> & set(const encapc<dim2,T,Mem> & ec)
{
copy_cpu_encap_encap<encapc<dim2,T,Mem>,encapc<dim,T,Mem>> cp(ec,*this);
......
......@@ -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
......
......@@ -902,7 +902,9 @@ public:
check_init();
check_bound(v1);
#endif
return mem_geto<dim,T,layout_base<T>,decltype(this->data_),decltype(this->g1),decltype(v1)>::get(const_cast<decltype(this->data_) &>(data_),g1,v1);
return mem_geto<dim,T,layout_base<T>,decltype(this->data_),decltype(this->g1),decltype(v1)>
::get(const_cast<typename std::add_lvalue_reference<decltype(this->data_)>::type>(data_),
g1,v1);
}
/*! \brief Get the of the selected element as a boost::fusion::vector
......@@ -948,7 +950,8 @@ public:
check_init();
check_bound(v1);
#endif
return mem_geto<dim,T,layout_base<T>,decltype(this->data_),decltype(this->g1),decltype(v1)>::get_lin(const_cast<decltype(this->data_) &>(data_),v1);
return mem_geto<dim,T,layout_base<T>,decltype(this->data_),decltype(this->g1),decltype(v1)>
::get_lin(const_cast<typename std::add_lvalue_reference<decltype(this->data_)>::type>(data_),v1);
}
/*! \brief Fill the memory with the selected byte
......
......@@ -463,7 +463,7 @@ public:
mem_id lid = gk.get(0);
for (mem_id i = 1 ; i < N ; i++)
{
/* coverity[dead_error_begin */
/* coverity[dead_error_line] */
lid += gk.get(i) * sz_s[i-1];
}
......@@ -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;
......
......@@ -752,7 +752,8 @@ BOOST_AUTO_TEST_CASE(copy_encap_vector_fusion_test)
g.template get<2>(key)[2][1] = 12.0;
g.template get<2>(key)[2][2] = 13.0;
copy_encap_vector_fusion<decltype(g.get_o(key)),typename aggregate<float,float[3],float[3][3]>::type> cp(g.get_o(key),tmp);
auto ge = g.get_o(key);
copy_encap_vector_fusion<decltype(g.get_o(key)),typename aggregate<float,float[3],float[3][3]>::type> cp(ge,tmp);
boost::mpl::for_each_ref< boost::mpl::range_c<int,0,aggregate<float,float[3],float[3][3]>::max_prop> >(cp);
g.get_o(key1) = tmp;
......
......@@ -640,7 +640,7 @@ public:
*/
inline void private_adjust(size_t tot_add)
{
this->stl_code.template private_adjust(tot_add);
this->stl_code.private_adjust(tot_add);
}
};
......
......@@ -353,17 +353,23 @@ struct host_to_device_impl
template<typename T>
inline void operator()(T& t) const
{
typedef decltype(boost::fusion::at_c<boost::mpl::at<v_prp,boost::mpl::int_<T::value>>::type::value>(dst).mem_r) mem_r_type;
typedef typename boost::mpl::at<v_prp,boost::mpl::int_<T::value>>::type ele_type;
typedef typename boost::mpl::at<typename T_type::type,boost::mpl::int_<T::value>>::type type_prp;
typedef decltype(boost::fusion::at_c<ele_type::value>(dst).mem_r) mem_r_type;
typedef typename boost::mpl::at<typename T_type::type,ele_type>::type type_prp;
typedef typename toKernel_transform<layout_base,typename mem_r_type::value_type>::type kernel_type;
typedef boost::mpl::int_<(is_vector<typename mem_r_type::value_type>::value ||
is_vector_dist<typename mem_r_type::value_type>::value ||
is_gpu_celllist<typename mem_r_type::value_type>::value) + 2*std::is_array<type_prp>::value + std::rank<type_prp>::value> crh_cond;
call_recursive_host_device_if_vector<typename mem_r_type::value_type,
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>
crh_cond::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),
......@@ -374,7 +380,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>
0>
::call(boost::fusion::at_c<boost::mpl::at<v_prp,boost::mpl::int_<T::value>>::type::value>(dst).mem_r,start,stop);
}
};
......
......@@ -1487,6 +1487,10 @@ public:
CellDecomposer_sm(const SpaceBox<dim,T> & box, const size_t (&div)[dim], Matrix<dim,T> & mat, const size_t pad)
:t(Matrix<dim,T>::identity(),box.getP1()),box(box),gr_cell()
{
// set div_wp to zero
for (size_t i = 0 ; i < dim ; i++)
{div_wp[i] = 0;}
Initialize(pad);
}
......
......@@ -426,7 +426,7 @@ public:
//! Type of internal memory structure
typedef Mem_type Mem_type_type;
typedef CellNNIteratorSym<dim,CellList<dim,T,Mem_type,transform,vector_pos_type>,RUNTIME,vector_pos_type,NO_CHECK> SymNNIterator;
typedef CellNNIteratorSym<dim,CellList<dim,T,Mem_type,transform,vector_pos_type>,vector_pos_type,RUNTIME,NO_CHECK> SymNNIterator;
//! Object type that the structure store
typedef typename Mem_type::local_index_type value_type;
......@@ -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
......@@ -891,7 +894,10 @@ public:
Mem_type::swap(static_cast<Mem_type &>(cl));
static_cast<CellDecomposer_sm<dim,T,transform> &>(*this) = static_cast<const CellDecomposer_sm<dim,T,transform> &>(cl);
static_cast<CellDecomposer_sm<dim,T,transform> &>(*this).swap(static_cast<CellDecomposer_sm<dim,T,transform> &>(cl));
n_dec = cl.n_dec;
from_cd = cl.from_cd;
}
/*! \brief Get the Cell iterator
......@@ -969,14 +975,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;
}
......@@ -1004,7 +1010,7 @@ public:
*
*/
template<unsigned int impl>
inline CellNNIteratorSym<dim,CellList<dim,T,Mem_type,transform,vector_pos_type>,(unsigned int)SYM,vector_pos_type,impl>
inline CellNNIteratorSym<dim,CellList<dim,T,Mem_type,transform,vector_pos_type>,vector_pos_type,(unsigned int)SYM,impl>
getNNIteratorSym(size_t cell, size_t p, const vector_pos_type & v)
{
#ifdef SE_CLASS1
......@@ -1012,7 +1018,7 @@ public:
{std::cerr << __FILE__ << ":" << __LINE__ << " Warning when you try to get a symmetric neighborhood iterator, you must construct the Cell-list in a symmetric way" << std::endl;}
#endif
CellNNIteratorSym<dim,CellList<dim,T,Mem_type,transform,vector_pos_type>,SYM,vector_pos_type,impl> cln(cell,p,NNc_sym,*this,v);
CellNNIteratorSym<dim,CellList<dim,T,Mem_type,transform,vector_pos_type>,vector_pos_type,SYM,impl> cln(cell,p,NNc_sym,*this,v);
return cln;
}
......@@ -1038,16 +1044,16 @@ public:
* \return An aiterator across the neighborhood particles
*
*/
template<unsigned int impl>
inline CellNNIteratorSymMP<dim,CellList<dim,T,Mem_type,transform>,(unsigned int)SYM,impl>
getNNIteratorSymMP(size_t cell, size_t p, const openfpm::vector<Point<dim,T>> & v_p1, const openfpm::vector<Point<dim,T>> & v_p2)
template<unsigned int impl, typename vector_pos_type2>
inline CellNNIteratorSymMP<dim,CellList<dim,T,Mem_type,transform,vector_pos_type>,vector_pos_type2,(unsigned int)SYM,impl>
getNNIteratorSymMP(size_t cell, size_t p, const vector_pos_type2 & v_p1, const vector_pos_type2 & v_p2)
{
#ifdef SE_CLASS1
if (from_cd == false)
std::cerr << __FILE__ << ":" << __LINE__ << " Warning when you try to get a symmetric neighborhood iterator, you must construct the Cell-list in a symmetric way" << std::endl;
#endif
CellNNIteratorSymMP<dim,CellList<dim,T,Mem_type,transform>,SYM,impl> cln(cell,p,NNc_sym,*this,v_p1,v_p2);
CellNNIteratorSymMP<dim,CellList<dim,T,Mem_type,transform,vector_pos_type>,vector_pos_type2,SYM,impl> cln(cell,p,NNc_sym,*this,v_p1,v_p2);
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;
......
......@@ -22,18 +22,6 @@ enum cl_construct_opt
#endif
#include "util/cuda/ofp_context.hxx"
/*! \brief Check this is a gpu or cpu type cell-list
*
*/
template<typename T, typename Sfinae = void>
struct is_gpu_celllist: std::false_type {};
template<typename T>
struct is_gpu_celllist<T, typename Void<typename T::yes_is_gpu_celllist>::type> : std::true_type
{};
/*! \brief populate the Cell-list with particles non symmetric case on GPU
*
......@@ -238,12 +226,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)
{}
};
......
......@@ -163,7 +163,7 @@ public:
* \tparam impl implementation specific options NO_CHECK do not do check on access, SAFE do check on access
*
*/
template<unsigned int dim, typename Cell,int NNc_size, typename vector_pos_type, unsigned int impl>
template<unsigned int dim, typename Cell, typename vector_pos_type,int NNc_size, unsigned int impl>
class CellNNIteratorSym : public CellNNIterator<dim,Cell,NNc_size,impl>
{
//! index of the particle p
......@@ -226,7 +226,7 @@ public:
* \return itself
*
*/
inline CellNNIteratorSym<dim,Cell,NNc_size,vector_pos_type,impl> & operator++()
inline CellNNIteratorSym<dim,Cell,vector_pos_type,NNc_size,impl> & operator++()
{
this->start_id++;
......@@ -252,16 +252,17 @@ public:
* \tparam impl implementation specific options NO_CHECK do not do check on access, SAFE do check on access
*
*/
template<unsigned int dim, typename Cell,int NNc_size, unsigned int impl> class CellNNIteratorSymMP : public CellNNIterator<dim,Cell,NNc_size,impl>
template<unsigned int dim, typename Cell, typename vector_pos_type , int NNc_size, unsigned int impl>
class CellNNIteratorSymMP : public CellNNIterator<dim,Cell,NNc_size,impl>
{
//! index of the particle p
size_t p;
//! Phase vector for particle p
const openfpm::vector<Point<dim,typename Cell::stype>> & v_p1;
const vector_pos_type & v_p1;
//! Phase vector for particle q
const openfpm::vector<Point<dim,typename Cell::stype>> & v_p2;
const vector_pos_type & v_p2;
/*! Select the next valid element
*
......@@ -309,8 +310,8 @@ public:
size_t p,
const NNc_array<dim,NNc_size> &NNc,
Cell & cl,
const openfpm::vector<Point<dim,typename Cell::stype>> & v_p1,
const openfpm::vector<Point<dim,typename Cell::stype>> & v_p2)
const vector_pos_type & v_p1,
const vector_pos_type & v_p2)
:CellNNIterator<dim,Cell,NNc_size,impl>(cell,NNc,cl),p(p),v_p1(v_p1),v_p2(v_p2)
{
selectValid();
......@@ -322,7 +323,7 @@ public:
* \return itself
*
*/
inline CellNNIteratorSymMP<dim,Cell,NNc_size,impl> & operator++()
inline CellNNIteratorSymMP<dim,Cell,vector_pos_type,NNc_size,impl> & operator++()
{
this->start_id++;
......
......@@ -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();
......
......@@ -158,7 +158,7 @@ public:
*
*/
template<unsigned int dim, typename Cell,typename vector_pos_type, unsigned int impl>
class CellNNIteratorSym<dim,Cell,RUNTIME,vector_pos_type,impl> : public CellNNIterator<dim,Cell,RUNTIME,impl>
class CellNNIteratorSym<dim,Cell,vector_pos_type,RUNTIME,impl> : public CellNNIterator<dim,Cell,RUNTIME,impl>
{
//! index of the particle p
size_t p;
......@@ -228,7 +228,7 @@ public:
* \return itself
*
*/
inline CellNNIteratorSym<dim,Cell,RUNTIME,vector_pos_type,impl> & operator++()
inline CellNNIteratorSym<dim,Cell,vector_pos_type,RUNTIME,impl> & operator++()
{
this->start_id++;
......
......@@ -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);
......
......@@ -27,7 +27,12 @@
constexpr int count = 0;
constexpr int start = 1;
template<unsigned int dim, typename T, typename Memory, typename transform = no_transform_only<dim,T>, typename cnt_type = unsigned int, typename ids_type = int>
template<unsigned int dim,
typename T,
typename Memory,
typename transform = no_transform_only<dim,T>,
typename cnt_type = unsigned int,
typename ids_type = int>
class CellList_gpu : public CellDecomposer_sm<dim,T,transform>
{
typedef openfpm::vector<aggregate<cnt_type>,Memory,typename memory_traits_inte<aggregate<cnt_type>>::type,memory_traits_inte> vector_cnt_type;
......@@ -93,28 +98,29 @@ public:
//! Indicate that this cell list is a gpu type cell-list
typedef int yes_is_gpu_celllist;
//! the type of the space
typedef T stype;
//! dimensions of space
static const unsigned int dims = dim;
//! count type
typedef cnt_type cnt_type_;
//! id type
typedef ids_type ids_type_;
//! transform type
typedef transform transform_;
/*! \brief Copy constructor
*
*
*
*/
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 +129,17 @@ public:
*
*/
CellList_gpu(CellList_gpu<dim,T,Memory,transform,cnt_type,ids_type> && clg)
:CellDecomposer_sm<dim,T,transform>(clg)