Commit 5766014d authored by incardon's avatar incardon
Browse files

Fixing SKIP_LABELLING

parent c3fbb8da
Pipeline #1929 canceled with stages
openfpm_data @ 75b93fcc
Subproject commit 8c459e2608db17c0ad8fd2a27d29e4de9b01cc2c
Subproject commit 75b93fcc133165ba640eb0d1ab9c282d563f9c12
......@@ -1017,106 +1017,135 @@ public:
size_t req = 0;
// first we initialize the pack buffer on all internal grids
// Pack information
Pack_stat sts;
for (size_t i = 0 ; i < loc_grid.size() ; i++)
{loc_grid.get(i).packReset();}
// We check if skip labelling is possible in this condition
for (int i = 0 ; i < loc_grid.size() ; i++)
{opt &= (loc_grid.get(i).isSkipLabellingPossible())?(int)-1:~SKIP_LABELLING;}
// Calculating the size to pack all the data to send
for ( size_t i = 0 ; i < ig_box.size() ; i++ )
if (!(opt & SKIP_LABELLING))
{
// for each ghost box
for (size_t j = 0 ; j < ig_box.get(i).bid.size() ; j++)
// first we initialize the pack buffer on all internal grids
for (size_t i = 0 ; i < loc_grid.size() ; i++)
{loc_grid.get(i).packReset();}
// Calculating the size to pack all the data to send
for ( size_t i = 0 ; i < ig_box.size() ; i++ )
{
// And linked sub-domain
size_t sub_id = ig_box.get(i).bid.get(j).sub;
// Internal ghost box
Box<dim,long int> g_ig_box = ig_box.get(i).bid.get(j).box;
// for each ghost box
for (size_t j = 0 ; j < ig_box.get(i).bid.size() ; j++)
{
// And linked sub-domain
size_t sub_id = ig_box.get(i).bid.get(j).sub;
// Internal ghost box
Box<dim,long int> g_ig_box = ig_box.get(i).bid.get(j).box;
if (g_ig_box.isValid() == false)
{continue;}
if (g_ig_box.isValid() == false)
{continue;}
g_ig_box -= gdb_ext.get(sub_id).origin.template convertPoint<size_t>();
g_ig_box -= gdb_ext.get(sub_id).origin.template convertPoint<size_t>();
// Pack a size_t for the internal ghost id
Packer<size_t,Memory>::packRequest(req);
// Create a sub grid iterator spanning the internal ghost layer
auto sub_it = loc_grid.get(sub_id).getIterator(g_ig_box.getKP1(),g_ig_box.getKP2());
// Pack a size_t for the internal ghost id
Packer<size_t,Memory>::packRequest(req);
// Create a sub grid iterator spanning the internal ghost layer
auto sub_it = loc_grid.get(sub_id).getIterator(g_ig_box.getKP1(),g_ig_box.getKP2());
// get the size to pack
Packer<device_grid,Memory>::template packRequest<decltype(sub_it),prp...>(loc_grid.get(sub_id),sub_it,req);
// get the size to pack
Packer<device_grid,Memory>::template packRequest<decltype(sub_it),prp...>(loc_grid.get(sub_id),sub_it,req);
}
}
}
// Finalize calculation
for (size_t i = 0 ; i < loc_grid.size() ; i++)
{loc_grid.get(i).template packCalculate<prp ...>(req,v_cl.getmgpuContext());}
// Finalize calculation
for (size_t i = 0 ; i < loc_grid.size() ; i++)
{loc_grid.get(i).template packCalculate<prp ...>(req,v_cl.getmgpuContext());}
// resize the property buffer memory
g_send_prp_mem.resize(req);
// resize the property buffer memory
g_send_prp_mem.resize(req);
// Create an object of preallocated memory for properties
ExtPreAlloc<Memory> & prAlloc_prp = *(new ExtPreAlloc<Memory>(req,g_send_prp_mem));
// Create an object of preallocated memory for properties
ExtPreAlloc<Memory> & prAlloc_prp = *(new ExtPreAlloc<Memory>(req,g_send_prp_mem));
prAlloc_prp.incRef();
prAlloc_prp.incRef();
// Pack information
Pack_stat sts;
pointers.clear();
pointers2.clear();
pointers.clear();
pointers2.clear();
// Pack the information for each processor and send it
for ( size_t i = 0 ; i < ig_box.size() ; i++ )
{
// Pack the information for each processor and send it
for ( size_t i = 0 ; i < ig_box.size() ; i++ )
{
sts.mark();
sts.mark();
void * pointer;
void * pointer;
if (opt & RUN_ON_DEVICE)
{pointer = prAlloc_prp.getDevicePointerEnd();}
else
{pointer = prAlloc_prp.getPointerEnd();}
if (opt & RUN_ON_DEVICE)
{pointer = prAlloc_prp.getDevicePointerEnd();}
else
{pointer = prAlloc_prp.getPointerEnd();}
// for each ghost box
for (size_t j = 0 ; j < ig_box.get(i).bid.size() ; j++)
{
// we pack only if it is valid
if (ig_box.get(i).bid.get(j).box.isValid() == false)
continue;
// for each ghost box
for (size_t j = 0 ; j < ig_box.get(i).bid.size() ; j++)
{
// we pack only if it is valid
if (ig_box.get(i).bid.get(j).box.isValid() == false)
continue;
// And linked sub-domain
size_t sub_id = ig_box.get(i).bid.get(j).sub;
// Internal ghost box
Box<dim,size_t> g_ig_box = ig_box.get(i).bid.get(j).box;
g_ig_box -= gdb_ext.get(sub_id).origin.template convertPoint<size_t>();
// Ghost box global id
size_t g_id = ig_box.get(i).bid.get(j).g_id;
// Pack a size_t for the internal ghost id
Packer<size_t,Memory>::pack(prAlloc_prp,g_id,sts);
prAlloc_prp.hostToDevice(prAlloc_prp.getOffset(),prAlloc_prp.getOffsetEnd());
// Create a sub grid iterator spanning the internal ghost layer
auto sub_it = loc_grid.get(sub_id).getIterator(g_ig_box.getKP1(),g_ig_box.getKP2());
// and pack the internal ghost grid
Packer<device_grid,Memory>::template pack<decltype(sub_it),prp...>(prAlloc_prp,loc_grid.get(sub_id),sub_it,sts);
}
// send the request
// And linked sub-domain
size_t sub_id = ig_box.get(i).bid.get(j).sub;
// Internal ghost box
Box<dim,size_t> g_ig_box = ig_box.get(i).bid.get(j).box;
g_ig_box -= gdb_ext.get(sub_id).origin.template convertPoint<size_t>();
// Ghost box global id
size_t g_id = ig_box.get(i).bid.get(j).g_id;
void * pointer2;
// Pack a size_t for the internal ghost id
Packer<size_t,Memory>::pack(prAlloc_prp,g_id,sts);
prAlloc_prp.hostToDevice(prAlloc_prp.getOffset(),prAlloc_prp.getOffsetEnd());
// Create a sub grid iterator spanning the internal ghost layer
auto sub_it = loc_grid.get(sub_id).getIterator(g_ig_box.getKP1(),g_ig_box.getKP2());
// and pack the internal ghost grid
Packer<device_grid,Memory>::template pack<decltype(sub_it),prp...>(prAlloc_prp,loc_grid.get(sub_id),sub_it,sts);
}
// send the request
if (opt & RUN_ON_DEVICE)
{pointer2 = prAlloc_prp.getDevicePointerEnd();}
else
{pointer2 = prAlloc_prp.getPointerEnd();}
void * pointer2;
pointers.add(pointer);
pointers2.add(pointer2);
}
if (opt & RUN_ON_DEVICE)
{pointer2 = prAlloc_prp.getDevicePointerEnd();}
else
{pointer2 = prAlloc_prp.getPointerEnd();}
for (size_t i = 0 ; i < loc_grid.size() ; i++)
{
rem_copy_opt opt_ = rem_copy_opt::NONE_OPT;
if (opt & SKIP_LABELLING == true)
{opt_ = rem_copy_opt::KEEP_GEOMETRY;}
pointers.add(pointer);
pointers2.add(pointer2);
loc_grid.get(i).template packFinalize<prp ...>(prAlloc_prp,sts,opt_,true);
}
}
else
{
req = g_send_prp_mem.size();
for (size_t i = 0 ; i < loc_grid.size() ; i++)
{loc_grid.get(i).template packFinalize<prp ...>(prAlloc_prp,sts);}
// Create an object of preallocated memory for properties
ExtPreAlloc<Memory> & prAlloc_prp = *(new ExtPreAlloc<Memory>(req,g_send_prp_mem));
for (size_t i = 0 ; i < loc_grid.size() ; i++)
{
rem_copy_opt opt_ = rem_copy_opt::NONE_OPT;
if (opt & SKIP_LABELLING)
{opt_ = rem_copy_opt::KEEP_GEOMETRY;}
loc_grid.get(i).template packFinalize<prp ...>(prAlloc_prp,sts,opt_,true);
}
}
for ( size_t i = 0 ; i < ig_box.size() ; i++ )
{
......
......@@ -662,6 +662,13 @@ BOOST_AUTO_TEST_CASE( sgrid_gpu_test_skip_labelling )
typedef typename GetCpBlockType<decltype(gdist),0,1>::type CpBlockType;
gdist.template conv2<0,1,0,1,1>({0,0,0},{(int)sz[0]-1,(int)sz[1]-1,(int)sz[2]-1},[] __device__ (float & u_out, float & v_out, CpBlockType & u, CpBlockType & v,int i, int j, int k){
u_out = 1*u(i,j,k);
v_out = 1*v(i,j,k);
});
gdist.template ghost_get<0,1>(RUN_ON_DEVICE | SKIP_LABELLING);
gdist.template conv2<0,1,0,1,1>({0,0,0},{(int)sz[0]-1,(int)sz[1]-1,(int)sz[2]-1},[] __device__ (float & u_out, float & v_out, CpBlockType & u, CpBlockType & v,int i, int j, int k){
u_out = 5*u(i,j,k);
v_out = 5*v(i,j,k);
......
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