diff --git a/runtime/domain/ParticleDataHanding.hpp b/runtime/domain/ParticleDataHanding.hpp
index b36cd0e0db416a8e38d828efd29dd5e6efb5d657..6ee2710135327c582dc4e56d60d04fcba57078db 100644
--- a/runtime/domain/ParticleDataHanding.hpp
+++ b/runtime/domain/ParticleDataHanding.hpp
@@ -25,7 +25,7 @@ inline bool operator==(const ParticleDeleter& lhs, const ParticleDeleter& rhs) {
class ParticleDataHandling : public blockforest::BlockDataHandling<internal::ParticleDeleter>{
public:
- ParticleDataHandling() {}
+ ParticleDataHandling(PairsSimulation *_ps) { ps = _ps; }
virtual ~ParticleDataHandling() {}
virtual internal::ParticleDeleter *initialize(IBlock *const block) WALBERLA_OVERRIDE {
@@ -71,6 +71,8 @@ public:
}
private:
+ PairsSimulation *ps;
+
void serializeImpl(Block *const block, const BlockDataID&, mpi::SendBuffer& buffer, const uint_t child, bool check_child) {
auto ptr = buffer.allocate<uint_t>();
double aabb_check[6];
@@ -95,7 +97,7 @@ private:
}
for(auto& p: ps->getNonVolatileProperties()) {
- ps->copyPropertyToHost(p, ReadOnly);
+ ps->copyPropertyToHost(p, WriteAfterRead);
}
auto position = ps->getPropertyByName("position");
@@ -167,6 +169,8 @@ private:
void deserializeImpl(IBlock *const, const BlockDataID&, mpi::RecvBuffer& buffer) {
int nlocal = ps->getNumberOfLocalParticles();
+ real_t real_tmp;
+ int int_tmp;
uint_t nrecv;
buffer >> nrecv;
@@ -184,28 +188,33 @@ private:
constexpr int nelems = 3;
for(int e = 0; e < nelems; e++) {
- buffer >> vector_ptr(nlocal + i, e);
+ buffer >> real_tmp;
+ vector_ptr(nlocal + i, e) = real_tmp;
}
} else if(prop_type == Prop_Matrix) {
auto matrix_ptr = ps->getAsMatrixProperty(prop);
constexpr int nelems = 9;
for(int e = 0; e < nelems; e++) {
- buffer >> matrix_ptr(nlocal + i, e);
+ buffer >> real_tmp;
+ matrix_ptr(nlocal + i, e) = real_tmp;
}
} else if(prop_type == Prop_Quaternion) {
auto quat_ptr = ps->getAsQuaternionProperty(prop);
constexpr int nelems = 4;
for(int e = 0; e < nelems; e++) {
- buffer >> quat_ptr(nlocal + i, e);
+ buffer >> real_tmp;
+ quat_ptr(nlocal + i, e) = real_tmp;
}
} else if(prop_type == Prop_Integer) {
auto int_ptr = ps->getAsIntegerProperty(prop);
- buffer >> int_ptr(nlocal + i);
+ buffer >> int_tmp;
+ int_ptr(nlocal + i) = int_tmp;
} else if(prop_type == Prop_Real) {
auto float_ptr = ps->getAsFloatProperty(prop);
- buffer >> float_ptr(nlocal + i);
+ buffer >> real_tmp;
+ float_ptr(nlocal + i) = real_tmp;
} else {
std::cerr << "deserializeImpl(): Invalid property type!" << std::endl;
return 0;
@@ -213,10 +222,6 @@ private:
}
}
- for(auto& p: ps->getNonVolatileProperties()) {
- ps->clearDeviceFlags(p);
- }
-
ps->setNumberOfLocalParticles(nlocal + nrecv);
}
};
diff --git a/runtime/domain/block_forest.cpp b/runtime/domain/block_forest.cpp
index 02a8fca8204d551e912c4112f99deb2ed71f724c..54aba1db9f82dd0429cc84390b8ac99b4499c5f0 100644
--- a/runtime/domain/block_forest.cpp
+++ b/runtime/domain/block_forest.cpp
@@ -308,7 +308,7 @@ void BlockForest::initializeWorkloadBalancer() {
forest->setRefreshPhantomBlockMigrationPreparationFunction(prepFunc);
}
- forest->addBlockData(make_shared<ParticleDataHandling>(), "Interface");
+ forest->addBlockData(make_shared<ParticleDataHandling>(ps), "Interface");
}
void BlockForest::finalize() {
diff --git a/runtime/domain/block_forest.hpp b/runtime/domain/block_forest.hpp
index 6facb7d9e4d7f1e2a72935a8bdab67e060c87771..22b7fe123a018cc6c3419bb3c7626dbd224d6153 100644
--- a/runtime/domain/block_forest.hpp
+++ b/runtime/domain/block_forest.hpp
@@ -1,5 +1,16 @@
#include <map>
//---
+#include <blockforest/BlockForest.h>
+#include <blockforest/Initialization.h>
+#include <blockforest/loadbalancing/DynamicCurve.h>
+#include <blockforest/loadbalancing/DynamicDiffusive.h>
+#include <blockforest/loadbalancing/DynamicParMetis.h>
+#include <blockforest/loadbalancing/InfoCollection.h>
+#include <blockforest/loadbalancing/PODPhantomData.h>
+#include <pe/amr/level_determination/MinMaxLevelDetermination.h>
+#include <pe/amr/weight_assignment/MetisAssignmentFunctor.h>
+#include <pe/amr/weight_assignment/WeightAssignmentFunctor.h>
+//---
#include "../pairs_common.hpp"
#include "domain_partitioning.hpp"
diff --git a/runtime/interfaces/md.hpp b/runtime/interfaces/md.hpp
new file mode 100644
index 0000000000000000000000000000000000000000..c8bb6b29da34073cb58e82a8b54dc5d0f832f3ce
--- /dev/null
+++ b/runtime/interfaces/md.hpp
@@ -0,0 +1,27 @@
+#include "../pairs.hpp"
+
+namespace pairs_host_interface {
+
+int get_uid(int *uid, int i) { return uid[i]; }
+int get_shape(int *shape, int i) { return shape[i]; }
+int get_flags(int *flags, int i) { return flags[i]; }
+double get_position(double *position, int i, int j, int capacity) { return position[i * 3 + j]; }
+double get_mass(double *mass, int i) { return mass[i]; }
+double get_linear_velocity(double *linear_velocity, int i, int j, int capacity) { return linear_velocity[i * 3 + j]; }
+double get_force(double *force, int i, int j, int capacity) { return force[i * 3 + j]; }
+int get_type(int *type, int i) { return type[i]; }
+
+}
+
+namespace pairs_cuda_interface {
+
+__inline__ __device__ int get_uid(int *uid, int i) { return uid[i]; }
+__inline__ __device__ int get_shape(int *shape, int i) { return shape[i]; }
+__inline__ __device__ int get_flags(int *flags, int i) { return flags[i]; }
+__inline__ __device__ double get_position(double *position, int i, int j, int capacity) { return position[i * 3 + j]; }
+__inline__ __device__ double get_mass(double *mass, int i) { return mass[i]; }
+__inline__ __device__ double get_linear_velocity(double *linear_velocity, int i, int j, int capacity) { return linear_velocity[i * 3 + j]; }
+__inline__ __device__ double get_force(double *force, int i, int j, int capacity) { return force[i * 3 + j]; }
+__inline__ __device__ int get_type(int *type, int i) { return type[i]; }
+
+}
diff --git a/runtime/pairs.hpp b/runtime/pairs.hpp
index 93f458392e5f2c28a50e04ad3c9d9996e984cfdf..f47e7148007024a50e964c6de58298b1b8fd0bf7 100644
--- a/runtime/pairs.hpp
+++ b/runtime/pairs.hpp
@@ -70,7 +70,9 @@ public:
return RuntimeVar<T>(h_ptr);
}
+ void setNumberOfLocalParticles(int n) { *nlocal = n; }
const int getNumberOfLocalParticles() { return *nlocal; }
+ void setNumberOfGhostParticles(int n) { *nghost = n; }
const int getNumberOfGhostParticles() { return *nghost; }
// Arrays
diff --git a/src/pairs/sim/domain_partitioning.py b/src/pairs/sim/domain_partitioning.py
index dae2d28fc0a518741ecfd5797e240ed9877adabd..93822feacca8a32fb2e4a0a093ce2cda2cb87bde 100644
--- a/src/pairs/sim/domain_partitioning.py
+++ b/src/pairs/sim/domain_partitioning.py
@@ -75,12 +75,13 @@ class BlockForest:
self.sim = sim
self.nranks = sim.add_var('nranks', Types.Int32)
self.nranks_capacity = sim.add_var('nranks_capacity', Types.Int32)
+ self.ntotal_aabbs = sim.add_var('ntotal_aabbs', Types.Int32)
self.aabb_capacity = sim.add_var('aabb_capacity', Types.Int32)
- self.ranks = sim.add_static_array('ranks', [self.nranks_capacity], Types.Int32)
- self.naabbs = sim.add_static_array('naabbs', [self.nranks_capacity], Types.Int32)
- self.aabb_offsets = sim.add_static_array('aabb_offsets', [self.nranks_capacity], Types.Int32)
- self.aabbs = sim.add_static_array('aabbs', [self.aabb_capacity, 6], Types.Real)
- self.subdom = sim.add_static_array('subdom', [sim.ndims() * 2], Types.Real)
+ self.ranks = sim.add_array('ranks', [self.nranks_capacity], Types.Int32)
+ self.naabbs = sim.add_array('naabbs', [self.nranks_capacity], Types.Int32)
+ self.aabb_offsets = sim.add_array('aabb_offsets', [self.nranks_capacity], Types.Int32)
+ self.aabbs = sim.add_array('aabbs', [self.aabb_capacity, 6], Types.Real)
+ self.subdom = sim.add_array('subdom', [sim.ndims() * 2], Types.Real)
def min(self, dim):
return self.subdom[dim * 2 + 0]
@@ -99,7 +100,7 @@ class BlockForest:
Call_Void(self.sim, "pairs->initDomain", [param for delim in grid_array for param in delim])
Assign(self.sim, self.nranks, Call_Int(self.sim, "pairs->getNumberOfNeighborRanks", []))
- Assign(self.sim, self.naabbs, Call_Int(self.sim, "pairs->getNumberOfNeighborAABBs", []))
+ Assign(self.sim, self.ntotal_aabbs, Call_Int(self.sim, "pairs->getNumberOfNeighborAABBs", []))
for _ in Filter(self.sim, self.nranks_capacity < self.nranks):
Assign(self.sim, self.nranks_capacity, self.nranks + 10)
@@ -107,15 +108,14 @@ class BlockForest:
self.naabbs.realloc()
self.aabb_offsets.realloc()
- for _ in Filter(self.sim, self.aabb_capacity < self.naabbs):
- Assign(self.sim, self.aabb_capacity, self.naabbs + 20)
- self.pbc.realloc()
+ for _ in Filter(self.sim, self.aabb_capacity < self.ntotal_aabbs):
+ Assign(self.sim, self.aabb_capacity, self.ntotal_aabbs + 20)
self.aabbs.realloc()
Call_Void(self.sim, "pairs->copyRuntimeArray", ['ranks', self.ranks, self.nranks])
Call_Void(self.sim, "pairs->copyRuntimeArray", ['naabbs', self.naabbs, self.nranks])
Call_Void(self.sim, "pairs->copyRuntimeArray", ['aabb_offsets', self.aabb_offsets, self.nranks])
- Call_Void(self.sim, "pairs->copyRuntimeArray", ['aabbs', self.aabbs, self.naabbs * 6])
+ Call_Void(self.sim, "pairs->copyRuntimeArray", ['aabbs', self.aabbs, self.ntotal_aabbs * 6])
Call_Void(self.sim, "pairs->copyRuntimeArray", ['subdom', self.subdom, self.sim.ndims() * 2])
def ghost_particles(self, step, position, offset=0.0):