Commit 61805e84 authored by incardon's avatar incardon

Adding SparseGridGpu iterators

parent 5056c725
......@@ -1091,28 +1091,6 @@ public:
const Box<dim,long int> & box_src,
const Box<dim,long int> & box_dst)
{
// sub-grid where to unpack
/* grid_key_dx_iterator_sub<dim> src(grid_src.getGrid(),box_src.getKP1(),box_src.getKP2());
grid_key_dx_iterator_sub<dim> dst(getGrid(),box_dst.getKP1(),box_dst.getKP2());
while (src.isNext())
{
auto key_src = src.get();
auto key_dst = dst.get();
get_o(key_dst) = grid_src.get_o(key_src);
++src;
++dst;
}*/
///////////////////////////////////////
// typedef typename std::remove_reference<decltype(grid_src)>::type grid_cp;
// typedef typename std::remove_reference<decltype(grid_src.getGrid())>::type grid_info_cp;
// fix box_dst
Box<dim,size_t> box_src_;
......@@ -1226,21 +1204,6 @@ public:
++sub_src;
++sub_dst;
}
/* grid_key_dx_iterator_sub<dim> sub_src(grid_src.getGrid(),box_src.getKP1(),box_src.getKP2());
grid_key_dx_iterator_sub<dim> sub_dst(this->getGrid(),box_dst.getKP1(),box_dst.getKP2());
// const auto & gs = loc_grid.get(i);
// auto & gd = loc_grid.get(sub_id_dst);
while (sub_src.isNext())
{
// write the object in the last element
object_s_di_op<op,decltype(grid_src.get_o(sub_src.get())),decltype(this->get_o(sub_dst.get())),OBJ_ENCAP,prp...>(grid_src.get_o(sub_src.get()),this->get_o(sub_dst.get()));
++sub_src;
++sub_dst;
}*/
}
/*! \brief Resize the grid
......
......@@ -11,6 +11,33 @@
#include <type_traits>
#include "util/tokernel_transformation.hpp"
/*! \brief this class is a functor for "for_each" algorithm
*
* This class is a functor for "for_each" algorithm. For each
* element of the boost::vector the operator() is called.
* Is mainly used to call hostToDevice for each properties
*
*/
template<typename aggrT_src, typename local_grids_type>
struct setBackground_impl
{
aggrT_src & bck;
local_grids_type loc_grid;
inline setBackground_impl(aggrT_src & bck, local_grids_type & loc_grid)
:bck(bck),loc_grid(loc_grid)
{};
//! It call the copy function for each property
template<typename T>
inline void operator()(T& t)
{
for (size_t i = 0 ; i < loc_grid.size() ; i++)
{loc_grid.get(i).template setBackgroundValue<T::value>(bck.template get<T::value>());}
}
};
/*! \brief this class is a functor for "for_each" algorithm
*
* This class is a functor for "for_each" algorithm. For each
......
......@@ -445,9 +445,10 @@ public:
* \return a point unsigned long int
*
*/
Point<dim,size_t> toPoint() const
template<typename typeT = size_t>
inline Point<dim,typeT> toPoint() const
{
Point<dim,size_t> p;
Point<dim,typeT> p;
for (size_t i = 0; i < dim ; i++)
{
......@@ -478,7 +479,7 @@ public:
* \return the index value
*
*/
__device__ __host__ mem_id get(size_t i) const
__device__ __host__ index_type get(index_type i) const
{
return k[i];
}
......@@ -491,7 +492,7 @@ public:
* \param id value to set
*
*/
__device__ __host__ void set_d(size_t i, mem_id id)
__device__ __host__ void set_d(index_type i, index_type id)
{
#if defined(SE_CLASS1) && !defined(__NVCC__)
......
......@@ -859,10 +859,10 @@ ite_gpu<dim> getGPUIterator_impl(const grid_sm<dim,T2> & g1, const grid_key_dx<d
{ig.thr.x = (key2.get(0) - key1.get(0) + 1);}
if (dim >= 2 && ig.wthr.y == 1)
{ig.wthr.y = key2.get(1) - key1.get(1) + 1;}
{ig.thr.y = key2.get(1) - key1.get(1) + 1;}
if (dim == 3 && ig.wthr.z == 1)
{ig.wthr.z = key2.get(2) - key1.get(2) + 1;}
{ig.thr.z = key2.get(2) - key1.get(2) + 1;}
for (size_t i = 0 ; i < dim ; i++)
{
......
......@@ -484,7 +484,7 @@ public:
{
#ifdef SE_CLASS1
if (initialized == false)
{std::cerr << "Error: " << __FILE__ << __LINE__ << " using unitialized iterator" << "\n";}
{std::cerr << "Error: " << __FILE__ << ":" << __LINE__ << " using unitialized iterator" << "\n";}
#endif
for (size_t i = 0 ; i < nsteps ; i++)
......@@ -506,7 +506,7 @@ public:
{
#ifdef SE_CLASS1
if (initialized == false)
{std::cerr << "Error: " << __FILE__ << __LINE__ << " using unitialized iterator" << "\n";}
{std::cerr << "Error: " << __FILE__ << ":" << __LINE__ << " using unitialized iterator" << "\n";}
#endif
if (this->gk.get(dim-1) <= gk_stop.get(dim-1))
......@@ -643,6 +643,17 @@ public:
{
this->stl_code.private_adjust(tot_add);
}
/* \brief Set the iterator in a way that isNext return false
*
*/
void invalidate()
{
this->gk.set_d(dim-1,1);
this->gk_stop.set_d(dim-1,0);
initialized = true;
}
};
......
......@@ -337,6 +337,18 @@ public:
return background;
}
/*! \brief Set the background value
*
* \tparam p property to set
*
*/
template<unsigned int p>
void setBackgroundValue(const typename boost::mpl::at<typename T::type,boost::mpl::int_<p>>::type & val)
{
background.template get<p>() = val;
}
/*! \brief assign operator
*
* \return itself
......
......@@ -586,7 +586,7 @@ public:
template<unsigned int p>
void setBackgroundValue(const typename boost::mpl::at<typename T::type,boost::mpl::int_<p>>::type & val)
{
return background.template get<p>() = val;
background.template get<p>() = val;
}
/*! \brief Get the background value
......
......@@ -30,6 +30,16 @@ class BlockMapGpu
private:
typedef BlockTypeOf<AggregateBlockT, 0> BlockT0;
#ifdef SE_CLASS1
//! Indicate if the setGPUInsertBuffer has been called
bool is_setGPUInsertBuffer = false;
//! Indicate if the initializeGPUInsertBuffer has been called
bool is_initializeGPUInsertBuffer = false;
#endif
protected:
const static unsigned char EXIST_BIT = 0;
typedef typename AggregateAppend<DataBlock<unsigned char, BlockT0::size>, AggregateBlockT>::type AggregateInternalT;
......@@ -42,7 +52,6 @@ protected:
public:
typedef AggregateBlockT AggregateType;
public:
BlockMapGpu() = default;
/*! \brief Get the background value
......@@ -133,12 +142,49 @@ public:
void hostToDevice();
void setGPUInsertBuffer(int nBlock, int nSlot);
/*! \Brief Before inser any element you have to call this function to initialize the insert buffer
*
* \param nBlock number of blocks the insert buffer has
* \param nSlot maximum number of insertion each thread block does
*
*/
void setGPUInsertBuffer(int nBlock, int nSlot)
{
// Prealloc the insert buffer on the underlying sparse vector
blockMap.setGPUInsertBuffer(nBlock, nSlot);
initializeGPUInsertBuffer();
void initializeGPUInsertBuffer();
#ifdef SE_CLASS1
is_setGPUInsertBuffer = true;
#endif
}
void initializeGPUInsertBuffer()
{
//todo: Test if it's enough to just initialize masks to 0, without any background value
// Initialize the blocks to background
auto & insertBuffer = blockMap.getGPUInsertBuffer();
typedef BlockTypeOf<AggregateInternalT, pMask> BlockType; // Here assuming that all block types in the aggregate have the same size!
constexpr unsigned int chunksPerBlock = threadBlockSize / BlockType::size; // Floor is good here...
BlockMapGpuKernels::initializeInsertBuffer<pMask, chunksPerBlock> <<< insertBuffer.size()/chunksPerBlock, chunksPerBlock*BlockType::size >>>(
insertBuffer.toKernel());
#ifdef SE_CLASS1
is_initializeGPUInsertBuffer = true;
#endif
}
template<typename ... v_reduce>
void flush(mgpu::ofp_context_t &context, flush_type opt = FLUSH_ON_HOST);
void flush(mgpu::ofp_context_t &context, flush_type opt = FLUSH_ON_HOST)
{
#ifdef SE_CLASS1
if (is_setGPUInsertBuffer == false || is_initializeGPUInsertBuffer == false)
{std::cout << __FILE__ << ":" << __LINE__ << " error setGPUInsertBuffer you must call before doing any insertion " << std::endl;}
#endif
blockMap.template flush<v_reduce ... >(context, opt);
}
template<unsigned int p>
void setBackgroundValue(ScalarTypeOf<AggregateBlockT, p> backgroundValue);
......@@ -234,33 +280,6 @@ void BlockMapGpu<AggregateBlockT, threadBlockSize, indexT, layout_base>::hostToD
blockMap.template hostToDevice<pMask>();
}
template<typename AggregateBlockT, unsigned int threadBlockSize, typename indexT, template<typename> class layout_base>
void BlockMapGpu<AggregateBlockT, threadBlockSize, indexT, layout_base>::setGPUInsertBuffer(int nBlock, int nSlot)
{
// Prealloc the insert buffer on the underlying sparse vector
blockMap.setGPUInsertBuffer(nBlock, nSlot);
initializeGPUInsertBuffer();
}
template<typename AggregateBlockT, unsigned int threadBlockSize, typename indexT, template<typename> class layout_base>
void BlockMapGpu<AggregateBlockT, threadBlockSize, indexT, layout_base>::initializeGPUInsertBuffer()
{
//todo: Test if it's enough to just initialize masks to 0, without any background value
// Initialize the blocks to background
auto & insertBuffer = blockMap.getGPUInsertBuffer();
typedef BlockTypeOf<AggregateInternalT, pMask> BlockType; // Here assuming that all block types in the aggregate have the same size!
constexpr unsigned int chunksPerBlock = threadBlockSize / BlockType::size; // Floor is good here...
BlockMapGpuKernels::initializeInsertBuffer<pMask, chunksPerBlock> <<< insertBuffer.size()/chunksPerBlock, chunksPerBlock*BlockType::size >>>(
insertBuffer.toKernel());
}
template<typename AggregateBlockT, unsigned int threadBlockSize, typename indexT, template<typename> class layout_base>
template<typename ... v_reduce>
void BlockMapGpu<AggregateBlockT, threadBlockSize, indexT, layout_base>::flush(mgpu::ofp_context_t &context, flush_type opt)
{
blockMap.template flush<v_reduce .../*, sBitwiseOr_<pMask>*/>(context, opt);
}
template<typename AggregateBlockT, unsigned int threadBlockSize, typename indexT, template<typename> class layout_base>
template<unsigned int p>
void BlockMapGpu<AggregateBlockT, threadBlockSize, indexT, layout_base>::setBackgroundValue(
......
/*
* SparseGridGpu_iterator.hpp
*
* Created on: Sep 4, 2019
* Author: i-bird
*/
#ifndef SPARSEGRIDGPU_ITERATOR_HPP_
#define SPARSEGRIDGPU_ITERATOR_HPP_
/*! \brief Element index contain a data chunk index and a point index
*
* \tparam SparseGridGpu type
*
*/
template<typename SparseGridGpu_type>
class sparse_grid_gpu_index
{
//! chunk position id
int cnk_pos_id;
//! data id
int data_id;
//! SparseGridGpu used to add functionalities
const SparseGridGpu_type & sparseGrid;
public:
/*! \brief Constructor from SparseGridGpu
*
*
*/
inline sparse_grid_gpu_index(const SparseGridGpu_type & sparseGrid, int cnk_pos_id, int data_id)
:sparseGrid(sparseGrid),cnk_pos_id(cnk_pos_id),data_id(data_id)
{}
/*! \brief Convert to a point this index
*
* \see toPointS
*
* \return a point unsigned long int
*
*/
inline Point<SparseGridGpu_type::dims,size_t> toPoint() const
{
auto indexCnk = sparseGrid.private_get_index_array().template get<0>(cnk_pos_id);
auto coord = sparseGrid.getCoord(indexCnk*sparseGrid.getBlockSize() + data_id);
Point<SparseGridGpu_type::dims,size_t> p;
for (size_t i = 0; i < SparseGridGpu_type::dims ; i++)
{
p.get(i) = coord.get(i);
}
return p;
}
/*! \brief Get chunk position id
*
* Return the position of the chunk in the chunks array \see SparseGridGpu \see private_get_data_array
*
* \return Get chunk position id
*
*/
int get_cnk_pos_id() const
{
return cnk_pos_id;
}
/*! \brief Get chunk local index (the returned index < getblockSize())
*
* \return Get chunk position id
*
*/
int get_data_id() const
{
return data_id;
}
};
template<unsigned int dim, typename SparseGridType>
class SparseGridGpu_iterator
{
//! actual chunk
unsigned int chunk;
//! actual point inside the chunk
unsigned int pnt;
grid_key_dx<dim> a_cnk;
//! original SparseGrid
const SparseGridType & sparseGrid;
//! array type for the indexes
typedef typename std::remove_reference<decltype(sparseGrid.private_get_index_array())>::type index_array_type;
//! array type for the data
typedef typename std::remove_reference<decltype(sparseGrid.private_get_data_array())>::type data_array_type;
//! vector of the chunk indexes
const decltype(sparseGrid.private_get_index_array()) & ids;
//! vector containing each chunks datas
const decltype(sparseGrid.private_get_data_array()) & data;
//Get the chunk type
typedef typename boost::mpl::at<typename data_array_type::value_type::type,boost::mpl::int_<0>>::type chunk_type;
//Get the chunk type
typedef boost::mpl::int_<boost::mpl::size<typename data_array_type::value_type::type>::type::value-1> pMask;
//! Select the first valid point chunk
void SelectValid()
{
while (pnt < chunk_type::size && data.template get<pMask::value>(chunk)[pnt] == 0)
{
pnt++;
}
while (pnt == chunk_type::size && chunk < ids.size())
{
chunk++;
pnt = 0;
while (pnt < chunk_type::size && data.template get<pMask::value>(chunk)[pnt] == 0)
{
pnt++;
}
}
}
public:
/*! \brief Constructor
*
* \param ids vector of ids
* \para data vector of chunk data
*
*/
inline SparseGridGpu_iterator(const SparseGridType & sparseGrid)
:chunk(0),
pnt(0),
sparseGrid(sparseGrid),
ids(sparseGrid.private_get_index_array()),
data(sparseGrid.private_get_data_array())
{
SelectValid();
}
/*! \brief Check if there is the next element
*
* Check if there is the next element
*
* \return true if there is the next, false otherwise
*
*/
bool isNext() const
{
return chunk < ids.size();
}
/*! \brief Get the next element
*
* Get the next element
*
* \return the next grid_key
*
*/
inline SparseGridGpu_iterator<dim,SparseGridType> & operator++()
{
++pnt;
if (pnt >= chunk_type::size)
{
++chunk;
pnt = 0;
}
SelectValid();
return *this;
}
/*! \brief return the actual point
*
* \return the index of the actual point
*
*/
inline sparse_grid_gpu_index<SparseGridType> get()
{
sparse_grid_gpu_index<SparseGridType> spgi(sparseGrid,chunk,pnt);
return spgi;
}
};
#endif /* SPARSEGRIDGPU_ITERATOR_HPP_ */
......@@ -8,39 +8,220 @@
#ifndef SPARSEGRIDGPU_ITERATOR_SUB_HPP_
#define SPARSEGRIDGPU_ITERATOR_SUB_HPP_
template<unsigned int dim/*, typename blockMap_type*/>
#include "SparseGridGpu_iterator.hpp"
template<unsigned int dim, typename SparseGridType>
class SparseGridGpu_iterator_sub
{
//! actual chunk
unsigned int chunk;
unsigned int pnt;
Box<dim,size_t> sub;
//! original SparseGrid
const SparseGridType * sparseGrid;
//! chunk coordinates
grid_key_dx<dim,int> chunk_coord;
//! array type for the indexes
typedef typename std::remove_reference<decltype(sparseGrid->private_get_index_array())>::type index_array_type;
//! array type for the data
typedef typename std::remove_reference<decltype(sparseGrid->private_get_data_array())>::type data_array_type;
//! vector of the chunk indexes
const typename std::remove_reference<decltype(sparseGrid->private_get_index_array())>::type * ids;
//! vector containing each chunks datas
const typename std::remove_reference<decltype(sparseGrid->private_get_data_array())>::type * data;
//Get the chunk type
typedef typename boost::mpl::at<typename data_array_type::value_type::type,boost::mpl::int_<0>>::type chunk_type;
//Get the chunk type
typedef boost::mpl::int_<boost::mpl::size<typename data_array_type::value_type::type>::type::value-1> pMask;
// subset in grid coordinates
Box<dim,int> sub_set;
// subset in local chunk coordinates
Box<dim,int> res;
// in chunk iterator
grid_key_dx_iterator_sub<dim> in_chunk_it;
// chunk size
grid_sm<dim,void> chunk_sz;
/*! \brief initializr the chunk interator
*
*
*/
void initialize_chunk_it()
{
// compute if the chunk intersect the start - stop box
chunk_coord = sparseGrid->getCoord(ids->template get<0>(chunk)*sparseGrid->getBlockSize());
Box<dim,int> box;
for (int i = 0 ; i < dim ; i++)
{
box.setLow(i,chunk_coord.get(i));
box.setHigh(i,chunk_coord.get(i) + sparseGrid->getBlockEdgeSize() - 1);
}
// blockMap_type & blockMap;
if (sub_set.Intersect(box,res) == true)
{
// remove the offset
for (int i = 0 ; i < dim ; i++)
{
res.setLow(i,res.getLow(i) - chunk_coord.get(i));
res.setHigh(i,res.getHigh(i) - chunk_coord.get(i));
}
in_chunk_it.reinitialize(grid_key_dx_iterator_sub<dim>(chunk_sz,res.getKP1(),res.getKP2()));
while (in_chunk_it.isNext() == true && data->template get<pMask::value>(chunk)[chunk_sz.LinId(in_chunk_it.get())] == 0)
{
++in_chunk_it;
}
}
}
//! Select the first valid point chunk
void SelectValid()
{
while (in_chunk_it.isNext() == true && data->template get<pMask::value>(chunk)[chunk_sz.LinId(in_chunk_it.get())] == 0)
{
++in_chunk_it;
}
while (in_chunk_it.isNext() == false && chunk < ids->size())
{
chunk++;
initialize_chunk_it();
}
}
/*! \brief Initialize chunk_sz member
*
*
*/
void initialize_chunk_sz()
{
size_t sz[dim];
for (int i = 0 ; i < dim ; i++)
{sz[i] = sparseGrid->getBlockEdgeSize();}
chunk_sz.setDimensions(sz);
}
public:
/*! \brief Default constructor
*
*/
inline SparseGridGpu_iterator_sub()
:chunk(0),
sparseGrid(NULL),
ids(NULL),
data(NULL)
{
initialize_chunk_sz();
}
/*! \brief Constructor
*
* \param sparseGrid original sparse grid
* \param start starting point
* \param stop stop point
*
*/
inline SparseGridGpu_iterator_sub(const SparseGridType & sparseGrid,const grid_key_dx<dim> & start,const grid_key_dx<dim> & stop)
:chunk(0),
sparseGrid(&sparseGrid),
ids(&sparseGrid.private_get_index_array()),
data(&sparseGrid.private_get_data_array())
{
for (int i = 0; i < dim ; i++)
{
sub_set.setLow(i,start.get(i));
sub_set.setHigh(i,stop.get(i));
}
initialize_chunk_sz();
in_chunk_it.invalidate();
initialize_chunk_it();
SelectValid();
}
/*! \brief Reinitialize the iterator
*
* it re-initialize the iterator with the passed grid_key_dx_iterator_sub
* the actual position of the grid_key_dx_iterator_sub is ignored
* \param it_sub subiterator
*
*/
inline void reinitialize(const SparseGridGpu_iterator_sub<dim,SparseGridType> & it_sub)
{
this->operator=(it_sub);
}
/*! \brief Check if there is the next element
*
* Check if there is the next element
*
* \param g_s_it grid_key_dx_iterator_sub
* \return true if there is the next, false otherwise
*
*/
inline void reinitialize(const SparseGridGpu_iterator_sub & g_s_it)
bool isNext() const
{
// Reinitialize the iterator
chunk = g_s_it.chunk;
pnt = g_s_it.pnt;
return chunk < ids->size();
}
/*! \brief Get the next element
*
* Get the next element
*
* \return the next grid_key
*
*/
inline SparseGridGpu_iterator_sub<dim,SparseGridType> & operator++()
{
++in_chunk_it;
};
SelectValid();
class SparseGridGpu_iterator
{
return *this;
}
/*! \brief return the actual point