Commit 001fd4b6 authored by incardon's avatar incardon
Browse files

Added load for sparsegrid GPU

parent e9ac5c0d
Pipeline #4196 failed with stages
in 12 minutes and 46 seconds
......@@ -93,6 +93,8 @@ public:
}
}
static bool noPointers() {return true;}
#ifdef __NVCC__
//Constructors from dim3 and uint3 objects
__host__ __device__ grid_smb(const dim3 blockDimensions)
......
......@@ -809,7 +809,7 @@ public:
* \param stop end point
*
*/
struct ite_gpu<dim> getGPUIterator(grid_key_dx<dim,long int> & key1, grid_key_dx<dim,long int> & key2, size_t n_thr = default_kernel_wg_threads_) const
struct ite_gpu<dim> getGPUIterator(const grid_key_dx<dim,long int> & key1, const grid_key_dx<dim,long int> & key2, size_t n_thr = default_kernel_wg_threads_) const
{
return getGPUIterator_impl<dim>(g1,key1,key2,n_thr);
}
......@@ -856,6 +856,19 @@ public:
}
/*! \brief Return the memory object
*
* Return the memory object
*
* \tparam p array to retrieve
*
*/
template<unsigned int p>
auto getMemory() -> decltype(boost::fusion::at_c<p>(data_).getMemory())
{
return boost::fusion::at_c<p>(data_).getMemory();
}
/*! \brief Set the object that provide memory from outside
*
* An external allocator is useful with allocator like PreAllocHeapMem
......
......@@ -43,7 +43,7 @@
*
*
*/
template <typename T, bool has_noPointers>
template <typename T, int has_noPointers>
struct Pack_selector_unknown_type_impl
{
enum
......@@ -53,7 +53,7 @@ struct Pack_selector_unknown_type_impl
};
template <typename T>
struct Pack_selector_unknown_type_impl<T,false>
struct Pack_selector_unknown_type_impl<T,0>
{
enum
{
......@@ -61,6 +61,28 @@ struct Pack_selector_unknown_type_impl<T,false>
};
};
/*! \brief Pack selector for unknown type
*
*
*/
template <typename T>
struct Pack_selector_unknown_type_impl<T,2>
{
enum
{
value = PACKER_ARRAY_CP_PRIMITIVE
};
};
template <typename T>
struct Pack_selector_unknown_type_impl<T,3>
{
enum
{
value = PACKER_ARRAY_CP_PRIMITIVE
};
};
/*! \brief Pack selector for unknown type
*
*
......@@ -70,7 +92,7 @@ struct Pack_selector_known_type_impl
{
enum
{
value = Pack_selector_unknown_type_impl<T, has_noPointers<T>::value >::value
value = Pack_selector_unknown_type_impl<T, has_noPointers<T>::value + 2*std::is_array<T>::value >::value
};
};
......
......@@ -14,6 +14,29 @@ using BlockTypeOf = typename std::remove_reference<typename boost::fusion::resul
template<typename AggregateT, unsigned int p>
using ScalarTypeOf = typename std::remove_reference<typename boost::fusion::result_of::at_c<typename AggregateT::type, p>::type>::type::scalarType;
template<typename T>
struct meta_copy_set_bck
{
template<typename destType>
inline static void set(destType & bP ,T & backgroundValue, int j)
{
bP[j] = backgroundValue;
}
};
template<unsigned int N, typename T>
struct meta_copy_set_bck<T[N]>
{
template<typename destType>
inline static void set(destType & bP ,T * backgroundValue, int j)
{
for (int i = 0 ; i < N ; i++)
{
bP[i][j] = backgroundValue[i];
}
}
};
template<typename AggregateBlockT, unsigned int threadBlockSize=128, typename indexT=long int, template<typename> class layout_base=memory_traits_inte>
class BlockMapGpu
{
......@@ -280,19 +303,21 @@ public:
* \tparam p property p
*
*/
template<unsigned int p>
void setBackgroundValue(ScalarTypeOf<AggregateBlockT, p> backgroundValue)
template<unsigned int p, typename TypeBck>
void setBackgroundValue(TypeBck backgroundValue)
{
// NOTE: Here we assume user only passes Blocks and not scalars in the templated aggregate type
typedef BlockTypeOf<AggregateInternalT, p> BlockT;
typedef typename std::remove_all_extents<BlockTypeOf<AggregateInternalT, p>>::type BlockT_noarr;
typedef BlockTypeOf<AggregateInternalT, pMask> BlockM;
BlockT bP;
BlockM bM;
for (unsigned int i = 0; i < BlockT::size; ++i)
for (unsigned int i = 0; i < BlockT_noarr::size; ++i)
{
bP[i] = backgroundValue;
meta_copy_set_bck<TypeBck>::set(bP,backgroundValue,i);
//meta_copy<TypeBck>::meta_copy_(backgroundValue,bP[][i]);
bM[i] = 0;
}
......
......@@ -2817,15 +2817,99 @@ public:
*
*/
template<unsigned int p>
void setBackgroundValue(ScalarTypeOf<AggregateBlockT, p> backgroundValue)
void setBackgroundValue(typename boost::mpl::at<typename AggregateT::type,boost::mpl::int_<p>>::type backgroundValue)
{
bck.template get<p>() = backgroundValue;
meta_copy<typename boost::mpl::at<typename AggregateT::type,boost::mpl::int_<p>>::type>::meta_copy_(backgroundValue,bck.template get<p>());
BMG::template setBackgroundValue<p>(backgroundValue);
BMG::template setBackgroundValue<p,typename boost::mpl::at<typename AggregateT::type,boost::mpl::int_<p>>::type>(backgroundValue);
}
/////////////////////////////////// DISTRIBUTED INTERFACE ///////////////////////
//Functions to check if the packing object is complex
static bool pack()
{
return false;
}
/*! \brief Asking to pack a SparseGrid GPU without GPU context pack the grid on CPU and host memory
*
*
*/
template<int ... prp> inline
void packRequest(size_t & req) const
{
// To fill
auto & indexBuffer = BlockMapGpu<AggregateInternalT, threadBlockSize, indexT, layout_base>::blockMap.getIndexBuffer();
auto & dataBuffer = BlockMapGpu<AggregateInternalT, threadBlockSize, indexT, layout_base>::blockMap.getDataBuffer();
indexBuffer.template packRequest<prp ...>(req);
dataBuffer.template packRequest<prp ...>(req);
Packer<decltype(gridGeometry),HeapMemory>::packRequest(req);
}
/*! \brief Pack the object into the memory
*
* \tparam prp properties to pack
*
* \param mem preallocated memory where to pack the objects
*
* \param sts pack statistic
*
*/
template<int ... prp> void pack(ExtPreAlloc<HeapMemory> & mem,
Pack_stat & sts) const
{
auto & indexBuffer = BlockMapGpu<AggregateInternalT, threadBlockSize, indexT, layout_base>::blockMap.getIndexBuffer();
auto & dataBuffer = BlockMapGpu<AggregateInternalT, threadBlockSize, indexT, layout_base>::blockMap.getDataBuffer();
// To fill
indexBuffer.template pack<prp ...>(mem,sts);
dataBuffer.template pack<prp ...>(mem,sts);
Packer<decltype(gridGeometry),HeapMemory>::pack(mem,gridGeometry,sts);
}
/*! \brief Unpack the object into the memory
*
* \tparam prp properties to pack
*
* \param mem preallocated memory where to pack the objects
*
* \param sts pack statistic
*
*/
template<int ... prp> void unpack(ExtPreAlloc<HeapMemory> & mem,
Unpack_stat & ps)
{
auto & indexBuffer = BlockMapGpu<AggregateInternalT, threadBlockSize, indexT, layout_base>::blockMap.getIndexBuffer();
auto & dataBuffer = BlockMapGpu<AggregateInternalT, threadBlockSize, indexT, layout_base>::blockMap.getDataBuffer();
// To fill
indexBuffer.template unpack<prp ...>(mem,ps);
dataBuffer.template unpack<prp ...>(mem,ps);
Unpacker<decltype(gridGeometry),HeapMemory>::unpack(mem,gridGeometry,ps);
}
/*! \brief Unpack the object into the memory
*
* \tparam prp properties to pack
*
* \param mem preallocated memory where to pack the objects
*
* \param sts pack statistic
*
*/
template<int ... prp> void unpack(ExtPreAlloc<CudaMemory> & mem,
Unpack_stat & ps)
{
if (mem.size() != 0)
{std::cout << __FILE__ << ":" << __LINE__ << " not implemented: " << std::endl;}
}
/*! \brief memory requested to pack this object
*
* \param req request
......@@ -3253,6 +3337,8 @@ public:
*/
void swap(self & gr)
{
gridGeometry.swap(gr.gridGeometry);
BMG::swap(gr);
}
......
......@@ -2071,6 +2071,19 @@ namespace openfpm
return 1;
}
/*! \brief Return the memory object
*
* Return the memory object
*
* \tparam p array to retrieve
*
*/
template<unsigned int p>
auto getMemory() -> decltype(base.template getMemory<p>())
{
return base.template getMemory<p>();
}
/*! \brief Set the memory of the base structure using an object
*
* \param mem Memory object to use for allocation
......
......@@ -339,7 +339,7 @@ template<int ... prp> inline void pack(ExtPreAlloc<HeapMemory> & mem, Pack_stat
* \param mem preallocated memory from where to unpack the vector
* \param ps unpack-stat info
*/
template<int ... prp> inline void unpack(ExtPreAlloc<HeapMemory> & mem, Unpack_stat & ps)
template<int ... prp, typename MemType> inline void unpack(ExtPreAlloc<MemType> & mem, Unpack_stat & ps)
{
//if all of the aggregate properties are simple (don't have "pack()" member)
if (has_pack_agg<T,prp...>::result::value == false)
......@@ -353,7 +353,7 @@ template<int ... prp> inline void unpack(ExtPreAlloc<HeapMemory> & mem, Unpack_s
{
//Unpack a size of a source vector
size_t u2 = 0;
Unpacker<size_t, HeapMemory>::unpack(mem,u2,ps);
Unpacker<size_t, MemType>::unpack(mem,u2,ps);
//Resize a destination vector
this->resize(u2);
......@@ -361,7 +361,7 @@ template<int ... prp> inline void unpack(ExtPreAlloc<HeapMemory> & mem, Unpack_s
for (size_t i = 0 ; i < this->size() ; i++)
{
//Call an unpacker in nested way
call_aggregateUnpack<decltype(this->get(i)),HeapMemory,prp ... >::call_unpack(this->get(i),mem,ps);
call_aggregateUnpack<decltype(this->get(i)),MemType,prp ... >::call_unpack(this->get(i),mem,ps);
}
}
}
......
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