diff --git a/src/Grid/Geometry/grid_smb.hpp b/src/Grid/Geometry/grid_smb.hpp
index da6a73f49f76049045f36011f6ab992eb7eeba0a..c6f6c377069a49b4c93e695f59fd075d0a67c89a 100644
--- a/src/Grid/Geometry/grid_smb.hpp
+++ b/src/Grid/Geometry/grid_smb.hpp
@@ -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)
diff --git a/src/Grid/grid_base_implementation.hpp b/src/Grid/grid_base_implementation.hpp
index 8409c6bd688e7466f1c5fcaa6976f20b2d113557..6143093d9b1e042ed035eb7d42ace5ec8847727d 100644
--- a/src/Grid/grid_base_implementation.hpp
+++ b/src/Grid/grid_base_implementation.hpp
@@ -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
diff --git a/src/Packer_Unpacker/Pack_selector.hpp b/src/Packer_Unpacker/Pack_selector.hpp
index f8ac005f1eac595b02aff1bf6db53eb93439b4ef..b0d3898c8db59a9b2280414ffc664654283d05ed 100644
--- a/src/Packer_Unpacker/Pack_selector.hpp
+++ b/src/Packer_Unpacker/Pack_selector.hpp
@@ -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
};
};
diff --git a/src/SparseGridGpu/BlockMapGpu.hpp b/src/SparseGridGpu/BlockMapGpu.hpp
index 088ca28f980342ff43159ae5ff729f8826dd20ef..b8c3a65b2eea953644fecb2019a8380600db56f1 100644
--- a/src/SparseGridGpu/BlockMapGpu.hpp
+++ b/src/SparseGridGpu/BlockMapGpu.hpp
@@ -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;
}
diff --git a/src/SparseGridGpu/SparseGridGpu.hpp b/src/SparseGridGpu/SparseGridGpu.hpp
index 606b29c6ebd38fd55080f134109a822f61903c02..f50c4c399a81fc4d93de8a65d8f2524e77719c54 100644
--- a/src/SparseGridGpu/SparseGridGpu.hpp
+++ b/src/SparseGridGpu/SparseGridGpu.hpp
@@ -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);
}
diff --git a/src/Vector/map_vector.hpp b/src/Vector/map_vector.hpp
index 29d3ccf7d11a452a24a67f6650f8592064e6a12d..953fdeabe0832996773bef89523ebaee614685f3 100644
--- a/src/Vector/map_vector.hpp
+++ b/src/Vector/map_vector.hpp
@@ -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
diff --git a/src/Vector/vector_pack_unpack.ipp b/src/Vector/vector_pack_unpack.ipp
index fcdab8b721827a18c190f45d4a957fbbd8a8ebae..f0c26245c9be06d549d2db3cd903176f4d7b7c63 100644
--- a/src/Vector/vector_pack_unpack.ipp
+++ b/src/Vector/vector_pack_unpack.ipp
@@ -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);
}
}
}