Commit 789d748e authored by incardon's avatar incardon

Fixing test with se-class1

parent f9e6bc7a
openfpm_data @ 483a48d1
Subproject commit a351d177e1a7e155ef2e0048d9d0e103a95faf83
Subproject commit 483a48d152b03c8422401ac242113e8a46172228
......@@ -293,12 +293,12 @@ class grid_dist_id : public grid_dist_id_comm<dim,St,T,Decomposition,Memory,devi
* rounding-off error can produce ghost bigger than the discrete selected
* one. This function adjust for this round-off error
*
* \param sub_domain the sub-domain
* \param sub_id sub-domain id
* \param sub_domain_other the other sub-domain
* \param ib internal ghost box to adjust
*
*/
void set_for_adjustment(const Box<dim,long int> & sub_domain,
void set_for_adjustment(size_t sub_id,
const Box<dim,St> & sub_domain_other,
const comb<dim> & cmb,
Box<dim,long int> & ib,
......@@ -307,6 +307,14 @@ class grid_dist_id : public grid_dist_id_comm<dim,St,T,Decomposition,Memory,devi
if (g.isInvalidGhost() == true || use_bx_def == true)
{return;}
Box<dim,long int> sub_domain;
if (use_bx_def == false)
{
sub_domain = gdb_ext.get(sub_id).Dbox;
sub_domain += gdb_ext.get(sub_id).origin;
}
// Convert from SpaceBox<dim,St> to SpaceBox<dim,long int>
Box<dim,long int> sub_domain_other_exp = cd_sm.convertDomainSpaceIntoGridUnits(sub_domain_other,dec.periodicity());
......@@ -368,10 +376,7 @@ class grid_dist_id : public grid_dist_id_comm<dim,St,T,Decomposition,Memory,devi
auto & n_box = dec.getNearSubdomains(dec.IDtoProc(i));
Box<dim,long int> sub = gdb_ext.get(sub_id).Dbox;
sub += gdb_ext.get(sub_id).origin;
set_for_adjustment(sub,
set_for_adjustment(sub_id,
n_box.get(r_sub),dec.getProcessorIGhostPos(i,j),
ib,ghost_int);
......@@ -685,10 +690,7 @@ class grid_dist_id : public grid_dist_id_comm<dim,St,T,Decomposition,Memory,devi
size_t sub_id = i;
size_t r_sub = dec.getLocalIGhostSub(i,j);
Box<dim,long int> sub = gdb_ext.get(sub_id).Dbox;
sub += gdb_ext.get(sub_id).origin;
set_for_adjustment(sub,dec.getSubDomain(r_sub),
set_for_adjustment(sub_id,dec.getSubDomain(r_sub),
dec.getLocalIGhostPos(i,j),ib,ghost_int);
// Check if ib is valid if not it mean that the internal ghost does not contain information so skip it
......
......@@ -224,6 +224,9 @@ class grid_dist_id_comm
if (bx_dst.isValid() == false)
continue;
Box<dim,size_t> bx_src = flip_box(loc_eg_box.get(sub_id_dst).bid.get(k).ebox,loc_eg_box.get(sub_id_dst).bid.get(k).cmb,ginfo);
bx_src -= gdb_ext.get(sub_id_src_gdb_ext).origin;
#ifdef SE_CLASS1
if (use_bx_def == false)
......@@ -237,9 +240,6 @@ class grid_dist_id_comm
#endif
Box<dim,size_t> bx_src = flip_box(loc_eg_box.get(sub_id_dst).bid.get(k).ebox,loc_eg_box.get(sub_id_dst).bid.get(k).cmb,ginfo);
bx_src -= gdb_ext.get(sub_id_src_gdb_ext).origin;
auto & gd = loc_grid.get(sub_id_dst_gdb_ext);
gd.remove(bx_dst);
gd.copy_to(loc_grid.get(sub_id_src_gdb_ext),bx_src,bx_dst);
......
#define BOOST_TEST_DYN_LINK
#define PRINT_STACKTRACE
#include "config.h"
#include <boost/test/unit_test.hpp>
#include "VCluster/VCluster.hpp"
#include <Vector/vector_dist.hpp>
......@@ -1600,5 +1600,42 @@ BOOST_AUTO_TEST_CASE(vector_dist_compare_host_device)
BOOST_REQUIRE_EQUAL(test,true);
}
template<typename vT>
__global__ void launch_overflow(vT vs, vT vs2)
{
vs2.template getProp<1>(57)[0];
}
BOOST_AUTO_TEST_CASE(vector_dist_overflow_se_class1)
{
Box<3,double> domain({0.0,0.0,0.0},{1.0,1.0,1.0});
Ghost<3,double> g(0.1);
size_t bc[3] = {PERIODIC,PERIODIC,PERIODIC};
if (create_vcluster().size() >= 16)
{return;}
std::cout << "****** TEST ERROR MESSAGE BEGIN ********" << std::endl;
vector_dist_gpu<3,double,aggregate<double,double[3],double[3][3]>> vdg(0,domain,bc,g,DEC_GRAN(128));
vector_dist_gpu<3,double,aggregate<double,double[3],double[3][3]>> vdg2(0,domain,bc,g,DEC_GRAN(128));
vdg.setCapacity(100);
ite_gpu<1> ite;
ite.wthr.x = 1;
ite.wthr.y = 1;
ite.wthr.z = 1;
ite.thr.x = 1;
ite.thr.y = 1;
ite.thr.z = 1;
CUDA_LAUNCH(launch_overflow,ite,vdg.toKernel(),vdg2.toKernel());
std::cout << "****** TEST ERROR MESSAGE END ********" << std::endl;
}
BOOST_AUTO_TEST_SUITE_END()
......@@ -1953,12 +1953,35 @@ BOOST_AUTO_TEST_CASE( vector_high_dimension )
BOOST_AUTO_TEST_CASE ( vector_of_cell_list_compile_test )
{
auto & v_cl = create_vcluster();
// set the seed
// create the random generator engine
std::srand(v_cl.getProcessUnitID());
std::default_random_engine eg;
std::uniform_real_distribution<float> ud(0.0f, 1.0f);
Box<3,double> domain({0.0,0.0,0.0},{1.0,1.0,1.0});
Ghost<3,double> g(0.1);
size_t bc[3] = {NON_PERIODIC,NON_PERIODIC,NON_PERIODIC};
vector_dist<3,double,aggregate<float,float[3]>> vd(100,domain,bc,g);
auto it = vd.getIterator();
while (it.isNext())
{
auto key = it.get();
vd.getPos(key)[0] = ud(eg);
vd.getPos(key)[1] = ud(eg);
vd.getPos(key)[2] = ud(eg);
++it;
}
vd.map();
std::vector<decltype(vd.getCellList(0.1))> vector_of_celllist;
typedef vector_dist<3,double,aggregate<float,float[3]>> my_particles;
......
......@@ -429,7 +429,7 @@ public:
//! dimensions of space
static const unsigned int dims = dim;
//!
//! yes I am vector dist
typedef int yes_i_am_vector_dist;
/*! \brief Operator= for distributed vector
......@@ -2489,6 +2489,17 @@ public:
h5l.load(filename,v_pos,v_prp,g_m);
}
/*! \brief Reserve space for the internal vectors
*
* \param
*
*/
void setCapacity(unsigned int ns)
{
v_pos.reserve(ns);
v_prp.reserve(ns);
}
/*! \brief Output particle position and properties
*
* \param out output filename
......
......@@ -78,6 +78,9 @@ public:
//! tag the type as a vector that run on kernel
typedef int vector_kernel;
//! Indicate this structure has a function to check the device pointer
typedef int yes_has_check_device_pointer;
vector_dist_ker(int g_m, const openfpm::vector_gpu_ker<Point<dim,St>,layout_base> & v_pos,
const openfpm::vector_gpu_ker<typename apply_transform<layout_base,prop>::type,layout_base> & v_prp)
:g_m(g_m),v_pos(v_pos),v_prp(v_prp)
......@@ -274,6 +277,40 @@ public:
return cv.check;
}
#ifdef SE_CLASS1
/*! \brief Check if the device pointer is owned by this structure
*
* \return a structure pointer check with information about the match
*
*/
pointer_check check_device_pointer(void * ptr)
{
pointer_check pc;
pc.match = false;
// we check the position vector and the property vector
pc = v_pos.check_device_pointer(ptr);
if (pc.match == true)
{
pc.match_str = std::string("Particle index overflow in position (v_pos): ") + "\n" + pc.match_str;
return pc;
}
pc = v_prp.check_device_pointer(ptr);
if (pc.match == true)
{
pc.match_str = std::string("Particle index overflow in properties (v_prp): ") + "\n" + pc.match_str;
return pc;
}
return pc;
}
#endif
};
// This is a tranformation node for vector_distributed for the algorithm toKernel_tranform
......
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