Commit e7010bdc authored by incardon's avatar incardon
Browse files

Fixing bug in cell-list populate

parent 6853316c
...@@ -49,8 +49,8 @@ struct populate_cell_list_no_sym_impl ...@@ -49,8 +49,8 @@ struct populate_cell_list_no_sym_impl
template<unsigned int dim, typename T, typename prop, typename Memory, template <typename> class layout_base , typename CellList> template<unsigned int dim, typename T, typename prop, typename Memory, template <typename> class layout_base , typename CellList>
static void populate(openfpm::vector<Point<dim,T>,Memory,typename layout_base<Point<dim,T>>::type,layout_base > & pos, static void populate(openfpm::vector<Point<dim,T>,Memory,typename layout_base<Point<dim,T>>::type,layout_base > & pos,
openfpm::vector<Point<dim,T>,Memory,typename layout_base<Point<dim,T>>::type,layout_base > & v_pos_out, openfpm::vector<Point<dim,T>,Memory,typename layout_base<Point<dim,T>>::type,layout_base > & v_pos_out,
openfpm::vector<prop,Memory,typename layout_base<prop>::type,layout_base > & v_prp_out,
openfpm::vector<prop,Memory,typename layout_base<prop>::type,layout_base > & v_prp, openfpm::vector<prop,Memory,typename layout_base<prop>::type,layout_base > & v_prp,
openfpm::vector<prop,Memory,typename layout_base<prop>::type,layout_base > & v_prp_out,
CellList & cli, CellList & cli,
size_t g_m) size_t g_m)
{ {
...@@ -69,8 +69,8 @@ struct populate_cell_list_no_sym_impl<true> ...@@ -69,8 +69,8 @@ struct populate_cell_list_no_sym_impl<true>
template<unsigned int dim, typename T, typename prop, typename Memory, template <typename> class layout_base , typename CellList> template<unsigned int dim, typename T, typename prop, typename Memory, template <typename> class layout_base , typename CellList>
static void populate(openfpm::vector<Point<dim,T>,Memory,typename layout_base<Point<dim,T>>::type,layout_base > & pos, static void populate(openfpm::vector<Point<dim,T>,Memory,typename layout_base<Point<dim,T>>::type,layout_base > & pos,
openfpm::vector<Point<dim,T>,Memory,typename layout_base<Point<dim,T>>::type,layout_base > & v_pos_out, openfpm::vector<Point<dim,T>,Memory,typename layout_base<Point<dim,T>>::type,layout_base > & v_pos_out,
openfpm::vector<prop,Memory,typename layout_base<prop>::type,layout_base > & v_prp_out,
openfpm::vector<prop,Memory,typename layout_base<prop>::type,layout_base > & v_prp, openfpm::vector<prop,Memory,typename layout_base<prop>::type,layout_base > & v_prp,
openfpm::vector<prop,Memory,typename layout_base<prop>::type,layout_base > & v_prp_out,
CellList & cli, CellList & cli,
size_t g_m) size_t g_m)
{ {
...@@ -130,12 +130,12 @@ struct populate_cell_list_sym_impl<true> ...@@ -130,12 +130,12 @@ struct populate_cell_list_sym_impl<true>
template<unsigned int dim, typename T, typename prop, typename Memory, template <typename> class layout_base , typename CellList> template<unsigned int dim, typename T, typename prop, typename Memory, template <typename> class layout_base , typename CellList>
void populate_cell_list_no_sym(openfpm::vector<Point<dim,T>,Memory,typename layout_base<Point<dim,T>>::type,layout_base > & pos, void populate_cell_list_no_sym(openfpm::vector<Point<dim,T>,Memory,typename layout_base<Point<dim,T>>::type,layout_base > & pos,
openfpm::vector<Point<dim,T>,Memory,typename layout_base<Point<dim,T>>::type,layout_base > & v_pos_out, openfpm::vector<Point<dim,T>,Memory,typename layout_base<Point<dim,T>>::type,layout_base > & v_pos_out,
openfpm::vector<prop,Memory,typename layout_base<prop>::type,layout_base > & v_prp_out,
openfpm::vector<prop,Memory,typename layout_base<prop>::type,layout_base > & v_prp, openfpm::vector<prop,Memory,typename layout_base<prop>::type,layout_base > & v_prp,
openfpm::vector<prop,Memory,typename layout_base<prop>::type,layout_base > & v_prp_out,
CellList & cli, CellList & cli,
size_t g_m) size_t g_m)
{ {
populate_cell_list_no_sym_impl<is_gpu_celllist<CellList>::value>::populate(pos,v_pos_out,v_prp_out,v_prp,cli,g_m); populate_cell_list_no_sym_impl<is_gpu_celllist<CellList>::value>::populate(pos,v_pos_out,v_prp,v_prp_out,cli,g_m);
} }
/*! \brief populate the Cell-list with particles symmetric case /*! \brief populate the Cell-list with particles symmetric case
...@@ -172,14 +172,14 @@ void populate_cell_list_sym(openfpm::vector<Point<dim,T>,Memory,typename layout_ ...@@ -172,14 +172,14 @@ void populate_cell_list_sym(openfpm::vector<Point<dim,T>,Memory,typename layout_
template<unsigned int dim, typename T, typename prop, typename Memory, template <typename> class layout_base, typename CellList> template<unsigned int dim, typename T, typename prop, typename Memory, template <typename> class layout_base, typename CellList>
void populate_cell_list(openfpm::vector<Point<dim,T>,Memory,typename layout_base<Point<dim,T>>::type,layout_base> & pos, void populate_cell_list(openfpm::vector<Point<dim,T>,Memory,typename layout_base<Point<dim,T>>::type,layout_base> & pos,
openfpm::vector<Point<dim,T>,Memory,typename layout_base<Point<dim,T>>::type,layout_base > & v_pos_out, openfpm::vector<Point<dim,T>,Memory,typename layout_base<Point<dim,T>>::type,layout_base > & v_pos_out,
openfpm::vector<prop,Memory,typename layout_base<prop>::type,layout_base > & v_prp_out,
openfpm::vector<prop,Memory,typename layout_base<prop>::type,layout_base > & v_prp, openfpm::vector<prop,Memory,typename layout_base<prop>::type,layout_base > & v_prp,
openfpm::vector<prop,Memory,typename layout_base<prop>::type,layout_base > & v_prp_out,
CellList & cli, CellList & cli,
size_t g_m, size_t g_m,
size_t opt) size_t opt)
{ {
if (opt == CL_NON_SYMMETRIC) if (opt == CL_NON_SYMMETRIC)
{populate_cell_list_no_sym(pos,v_pos_out,v_prp_out,v_prp,cli,g_m);} {populate_cell_list_no_sym(pos,v_pos_out,v_prp,v_prp_out,cli,g_m);}
else else
{populate_cell_list_sym(pos,cli,g_m);} {populate_cell_list_sym(pos,cli,g_m);}
} }
......
...@@ -81,9 +81,20 @@ public: ...@@ -81,9 +81,20 @@ 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)
{ {
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;
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 /*! \brief Copy constructor from temporal
* *
* *
* *
...@@ -99,6 +110,8 @@ public: ...@@ -99,6 +110,8 @@ public:
spacing_c = clg.spacing_c; spacing_c = clg.spacing_c;
div_c = clg.div_c; div_c = clg.div_c;
off = clg.off; off = clg.off;
g_m = clg.g_m;
n_dec = clg.n_dec;
} }
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)
...@@ -210,6 +223,7 @@ public: ...@@ -210,6 +223,7 @@ public:
sorted_to_not_sorted.toKernel(), sorted_to_not_sorted.toKernel(),
static_cast<cnt_type *>(cells.template getDeviceBuffer<0>())); static_cast<cnt_type *>(cells.template getDeviceBuffer<0>()));
#else #else
std::cout << "Error: " << __FILE__ << ":" << __LINE__ << " you are calling CellList_gpu.construct() this function is suppose must be compiled with NVCC compiler, but it look like has been compiled by the standard system compiler" << std::endl; std::cout << "Error: " << __FILE__ << ":" << __LINE__ << " you are calling CellList_gpu.construct() this function is suppose must be compiled with NVCC compiler, but it look like has been compiled by the standard system compiler" << std::endl;
......
...@@ -30,7 +30,7 @@ struct copy_fusion_vector ...@@ -30,7 +30,7 @@ struct copy_fusion_vector
* \param dst destination fusion vector * \param dst destination fusion vector
* *
*/ */
__device__ inline copy_fusion_vector(const bfv & src, bfv & dst) __device__ __host__ inline copy_fusion_vector(const bfv & src, bfv & dst)
:src(src),dst(dst){}; :src(src),dst(dst){};
#ifdef SE_CLASS1 #ifdef SE_CLASS1
...@@ -47,7 +47,7 @@ struct copy_fusion_vector ...@@ -47,7 +47,7 @@ struct copy_fusion_vector
//! It call the copy function for each property //! It call the copy function for each property
template<typename T> template<typename T>
__device__ inline void operator()(T& t) __device__ __host__ inline void operator()(T& t)
{ {
// This is the type of the object we have to copy // This is the type of the object we have to copy
typedef typename boost::fusion::result_of::at_c<bfv,T::value>::type copy_type; typedef typename boost::fusion::result_of::at_c<bfv,T::value>::type copy_type;
......
...@@ -241,8 +241,8 @@ MGPU_DEVICE inline c_type shfl_##dir##_op(c_type x, int offset, \ ...@@ -241,8 +241,8 @@ MGPU_DEVICE inline c_type shfl_##dir##_op(c_type x, int offset, \
".reg .u32 hi;" \ ".reg .u32 hi;" \
".reg .pred p;" \ ".reg .pred p;" \
"mov.b64 {lo, hi}, %1;" \ "mov.b64 {lo, hi}, %1;" \
"shfl."#dir".b32 lo|p, lo, %2, %3;" \ "shfl.sync."#dir".b32 lo|p, lo, %2, %3,0xFFFFFFFF;" \
"shfl."#dir".b32 hi , hi, %2, %3;" \ "shfl.sync."#dir".b32 hi , hi, %2, %3,0xFFFFFFFF;" \
"mov.b64 r0, {lo, hi};" \ "mov.b64 r0, {lo, hi};" \
"@p "#ptx_op"."#ptx_type" r0, r0, %4;" \ "@p "#ptx_op"."#ptx_type" r0, r0, %4;" \
"mov."#ptx_type" %0, r0; }" \ "mov."#ptx_type" %0, r0; }" \
......
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