Merge branch 'python' into fix-4143

doc/sphinx/installation.rst

@@ -383,8 +383,6 @@ Fluid dynamics and fluid structure interaction
 
 -  ``EK_DEBUG``
 
--  ``EK_DOUBLE_PREC``
-
 
 .. _Interaction features:
 

doc/tutorials/active_matter/active_matter.ipynb

@@ -4,7 +4,7 @@
    "cell_type": "markdown",
    "metadata": {},
    "source": [
-    "# Active Matter Tutorial\n",
+    "# Active Matter\n",
     "## Table of Contents\n",
     "1. [Introduction](#Introduction)\n",
     "2. [Active particles](#Active-particles)\n",

doc/tutorials/charged_system/CMakeLists.txt

@@ -1,3 +1,3 @@
-configure_tutorial_target(TARGET tutorial_crg DEPENDS charged_system-1.ipynb)
+configure_tutorial_target(TARGET tutorial_crg DEPENDS charged_system.ipynb)
 
-nb_export(TARGET tutorial_crg SUFFIX "1" FILE "charged_system-1.ipynb" HTML_RUN)
+nb_export(TARGET tutorial_crg SUFFIX "1" FILE "charged_system.ipynb" HTML_RUN)

doc/tutorials/charged_system/charged_system-1.ipynb → doc/tutorials/charged_system/charged_system.ipynb

@@ -4,7 +4,7 @@
    "cell_type": "markdown",
    "metadata": {},
    "source": [
-    "# Tutorial 2: Charged Systems: Counterion Condensation\n",
+    "# A Charged System: Counterion Condensation\n",
     "\n",
     "## Table of contents\n",
     "* [Introduction](#Introduction)\n",

doc/tutorials/constant_pH/constant_pH.ipynb

@@ -4,7 +4,7 @@
    "cell_type": "markdown",
    "metadata": {},
    "source": [
-    "# Tutorial 12: The Reaction-Ensemble Method"
+    "# The Reaction-Ensemble Method"
    ]
   },
   {

doc/tutorials/electrokinetics/electrokinetics.ipynb

@@ -4,7 +4,7 @@
    "cell_type": "markdown",
    "metadata": {},
    "source": [
-    "# Tutorial 7: Electrokinetics"
+    "# Electrokinetics"
    ]
   },
   {

doc/tutorials/ferrofluid/ferrofluid_part1.ipynb

@@ -4,7 +4,7 @@
    "cell_type": "markdown",
    "metadata": {},
    "source": [
-    "# Ferrofluid - Part I"
+    "# Ferrofluid - Part 1"
    ]
   },
   {

doc/tutorials/ferrofluid/ferrofluid_part2.ipynb

@@ -4,7 +4,7 @@
    "cell_type": "markdown",
    "metadata": {},
    "source": [
-    "# Ferrofluid - Part II"
+    "# Ferrofluid - Part 2"
    ]
   },
   {

doc/tutorials/ferrofluid/ferrofluid_part3.ipynb

@@ -4,7 +4,7 @@
    "cell_type": "markdown",
    "metadata": {},
    "source": [
-    "# Ferrofluid - Part III "
+    "# Ferrofluid - Part 3 "
    ]
   },
   {

doc/tutorials/lattice_boltzmann/lattice_boltzmann_part1.ipynb

@@ -4,7 +4,7 @@
    "cell_type": "markdown",
    "metadata": {},
    "source": [
-    "# Tutorial 4 : The lattice-Boltzmann method in ESPResSo - Part 1\n",
+    "# The Lattice-Boltzmann Method in ESPResSo - Part 1\n",
     "\n",
     "#### Before you start:\n",
     "\n",

doc/tutorials/lattice_boltzmann/lattice_boltzmann_part2.ipynb

@@ -4,7 +4,7 @@
    "cell_type": "markdown",
    "metadata": {},
    "source": [
-    "# Tutorial 4 : The lattice-Boltzmann method in ESPResSo - Part 2"
+    "# The Lattice-Boltzmann Method in ESPResSo - Part 2"
    ]
   },
   {

doc/tutorials/lattice_boltzmann/lattice_boltzmann_part3.ipynb

@@ -4,7 +4,7 @@
    "cell_type": "markdown",
    "metadata": {},
    "source": [
-    "# Tutorial 4 : The lattice-Boltzmann method in ESPResSo - Part 3"
+    "# The Lattice-Boltzmann Method in ESPResSo - Part 3"
    ]
   },
   {

doc/tutorials/lattice_boltzmann/lattice_boltzmann_part4.ipynb

@@ -4,7 +4,7 @@
    "cell_type": "markdown",
    "metadata": {},
    "source": [
-    "# Tutorial 4 : The lattice-Boltzmann method in ESPResSo - Part 4"
+    "# The Lattice-Boltzmann Method in ESPResSo - Part 4"
    ]
   },
   {

doc/tutorials/lennard_jones/lennard_jones.ipynb

@@ -4,7 +4,7 @@
    "cell_type": "markdown",
    "metadata": {},
    "source": [
-    "# Tutorial 1: Lennard-Jones Liquid"
+    "# Introductory Tutorial: Lennard-Jones Liquid"
    ]
   },
   {

doc/tutorials/raspberry_electrophoresis/raspberry_electrophoresis.ipynb

@@ -4,7 +4,7 @@
    "cell_type": "markdown",
    "metadata": {},
    "source": [
-    "# Tutorial 5: Raspberry Electrophoresis"
+    "# Raspberry Electrophoresis"
    ]
   },
   {

doc/tutorials/visualization/visualization.ipynb

@@ -4,7 +4,7 @@
    "cell_type": "markdown",
    "metadata": {},
    "source": [
-    "# Tutorial 8: Visualization"
+    "# Visualization"
    ]
   },
   {

src/config/features.def

@@ -62,7 +62,6 @@ ELECTROKINETICS                 requires CUDA
 EK_BOUNDARIES                   implies ELECTROKINETICS, LB_BOUNDARIES_GPU, EXTERNAL_FORCES, ELECTROSTATICS
 EK_BOUNDARIES                   requires CUDA
 EK_DEBUG                        requires ELECTROKINETICS
-EK_DOUBLE_PREC                  requires ELECTROKINETICS
 
 /* Interaction features */
 TABULATED

src/core/CellStructure.cpp

@@ -154,6 +154,9 @@ static unsigned map_data_parts(unsigned data_parts) {
          | ((DATA_PART_POSITION & data_parts) ? GHOSTTRANS_POSITION : 0u)
          | ((DATA_PART_MOMENTUM & data_parts) ? GHOSTTRANS_MOMENTUM : 0u)
          | ((DATA_PART_FORCE & data_parts) ? GHOSTTRANS_FORCE : 0u)
+#ifdef BOND_CONSTRAINT
+         | ((DATA_PART_RATTLE & data_parts) ? GHOSTTRANS_RATTLE : 0u)
+#endif
          | ((DATA_PART_BONDS & data_parts) ? GHOSTTRANS_BONDS : 0u);
   /* clang-format on */
 }
@@ -166,6 +169,12 @@ void CellStructure::ghosts_reduce_forces() {
   ghost_communicator(decomposition().collect_ghost_force_comm(),
                      GHOSTTRANS_FORCE);
 }
+#ifdef BOND_CONSTRAINT
+void CellStructure::ghosts_reduce_rattle_correction() {
+  ghost_communicator(decomposition().collect_ghost_force_comm(),
+                     GHOSTTRANS_RATTLE);
+}
+#endif
 
 Utils::Span<Cell *> CellStructure::local_cells() {
   return decomposition().local_cells();

src/core/CellStructure.hpp

@@ -66,7 +66,10 @@ enum DataPart : unsigned {
   DATA_PART_POSITION = 2u,   /**< Particle::r */
   DATA_PART_MOMENTUM = 8u,   /**< Particle::m */
   DATA_PART_FORCE = 16u,     /**< Particle::f */
-  DATA_PART_BONDS = 32u      /**< Particle::bonds */
+#ifdef BOND_CONSTRAINT
+  DATA_PART_RATTLE = 32u, /**< Particle::rattle */
+#endif
+  DATA_PART_BONDS = 64u /**< Particle::bonds */
 };
 } // namespace Cells
 
@@ -369,6 +372,12 @@ struct CellStructure {
    * @brief Add forces from ghost particles to real particles.
    */
   void ghosts_reduce_forces();
+#ifdef BOND_CONSTRAINT
+  /**
+   * @brief Add rattle corrections from ghost particles to real particles.
+   */
+  void ghosts_reduce_rattle_correction();
+#endif
 
 private:
   /**

src/core/Particle.hpp

@@ -268,7 +268,7 @@ struct ParticlePosition {
 
 #ifdef BOND_CONSTRAINT
   /** particle position at the previous time step (RATTLE algorithm) */
-  Utils::Vector3d p_old = {0., 0., 0.};
+  Utils::Vector3d p_last_timestep = {0., 0., 0.};
 #endif
 
   template <class Archive> void serialize(Archive &ar, long int /* version */) {
@@ -277,7 +277,7 @@ struct ParticlePosition {
     ar &quat;
 #endif
 #ifdef BOND_CONSTRAINT
-    ar &p_old;
+    ar &p_last_timestep;
 #endif
   }
 };
@@ -365,6 +365,26 @@ struct ParticleLocal {
   }
 };
 
+#ifdef BOND_CONSTRAINT
+struct ParticleRattle {
+  /** position/velocity correction */
+  Utils::Vector3d correction = {0, 0, 0};
+
+  friend ParticleRattle operator+(ParticleRattle const &lhs,
+                                  ParticleRattle const &rhs) {
+    return {lhs.correction + rhs.correction};
+  }
+
+  ParticleRattle &operator+=(ParticleRattle const &rhs) {
+    return *this = *this + rhs;
+  }
+
+  template <class Archive> void serialize(Archive &ar, long int /* version */) {
+    ar &correction;
+  }
+};
+#endif
+
 /** Struct holding all information for one particle. */
 struct Particle { // NOLINT(bugprone-exception-escape)
   int &identity() { return p.identity; }
@@ -391,6 +411,10 @@ struct Particle { // NOLINT(bugprone-exception-escape)
   ParticleForce f;
   ///
   ParticleLocal l;
+#ifdef BOND_CONSTRAINT
+  ///
+  ParticleRattle rattle;
+#endif
 
 private:
   BondList bl;

src/core/actor/DipolarBarnesHut.hpp

@@ -33,9 +33,6 @@
 
 #include <memory>
 
-// This needs to be done in the .cu file too
-typedef float dds_float;
-
 class DipolarBarnesHut : public Actor {
 public:
   DipolarBarnesHut(SystemInterface &s, float epssq, float itolsq) {

src/core/actor/DipolarBarnesHut_cuda.cu

@@ -36,12 +36,8 @@
 
 #include <cuda.h>
 
-typedef float dds_float;
-
 #define IND (blockDim.x * blockIdx.x + threadIdx.x)
 
-using namespace std;
-
 // Method performance/accuracy parameters
 __constant__ float epssqd[1], itolsqd[1];
 // blkcntd is a factual blocks' count.
@@ -57,7 +53,7 @@ __device__ __constant__ volatile BHData bhpara[1];
 // The "half-convolution" multi-thread reduction.
 // The thread with a lower index will operate longer and
 // final result (here: the sum) is flowing down towards zero thread.
-__device__ void dds_sumReduction_BH(dds_float *input, dds_float *sum) {
+__device__ void dds_sumReduction_BH(float *input, float *sum) {
   auto tid = static_cast<int>(threadIdx.x);
   for (auto i = static_cast<int>(blockDim.x); i > 1; i /= 2) {
     __syncthreads();
@@ -690,7 +686,7 @@ __global__ __launch_bounds__(THREADS4, FACTOR4) void sortKernel() {
 /******************************************************************************/
 
 __global__ __launch_bounds__(THREADS5, FACTOR5) void forceCalculationKernel(
-    dds_float pf, float *force, float *torque) {
+    float pf, float *force, float *torque) {
   int i, j, k, l, n, depth, base, sbase, diff, t;
   float tmp;
   // dr is a distance between particles.
@@ -902,7 +898,7 @@ __global__ __launch_bounds__(THREADS5, FACTOR5) void forceCalculationKernel(
 /******************************************************************************/
 
 __global__ __launch_bounds__(THREADS5, FACTOR5) void energyCalculationKernel(
-    dds_float pf, dds_float *energySum) {
+    float pf, float *energySum) {
   // NOTE: the algorithm of this kernel is almost identical to
   // forceCalculationKernel. See comments there.
 
@@ -912,8 +908,8 @@ __global__ __launch_bounds__(THREADS5, FACTOR5) void energyCalculationKernel(
   __shared__ int pos[MAXDEPTH * THREADS5 / WARPSIZE],
       node[MAXDEPTH * THREADS5 / WARPSIZE];
   __shared__ float dq[MAXDEPTH * THREADS5 / WARPSIZE];
-  dds_float sum = 0.0;
-  extern __shared__ dds_float res[];
+  float sum = 0.0;
+  extern __shared__ float res[];
 
   float b, d1, dd5;
 
@@ -1120,7 +1116,7 @@ void sortBH(int blocks) {
 }
 
 // Force calculation.
-int forceBH(BHData *bh_data, dds_float k, float *f, float *torque) {
+int forceBH(BHData *bh_data, float k, float *f, float *torque) {
   int error_code = 0;
   dim3 grid(1, 1, 1);
   dim3 block(1, 1, 1);
@@ -1137,26 +1133,26 @@ int forceBH(BHData *bh_data, dds_float k, float *f, float *torque) {
 }
 
 // Energy calculation.
-int energyBH(BHData *bh_data, dds_float k, float *E) {
+int energyBH(BHData *bh_data, float k, float *E) {
   int error_code = 0;
   dim3 grid(1, 1, 1);
   dim3 block(1, 1, 1);
 
   grid.x = bh_data->blocks * FACTOR5;
   block.x = THREADS5;
 
-  dds_float *energySum;
-  cuda_safe_mem(cudaMalloc(&energySum, (int)(sizeof(dds_float) * grid.x)));
+  float *energySum;
+  cuda_safe_mem(cudaMalloc(&energySum, (int)(sizeof(float) * grid.x)));
   // cleanup the memory for the energy sum
-  cuda_safe_mem(cudaMemset(energySum, 0, (int)(sizeof(dds_float) * grid.x)));
+  cuda_safe_mem(cudaMemset(energySum, 0, (int)(sizeof(float) * grid.x)));
 
   KERNELCALL_shared(energyCalculationKernel, grid, block,
-                    block.x * sizeof(dds_float), k, energySum);
+                    block.x * sizeof(float), k, energySum);
   cuda_safe_mem(cudaDeviceSynchronize());
 
   // Sum the results of all blocks
   // One energy part per block in the prev kernel
-  thrust::device_ptr<dds_float> t(energySum);
+  thrust::device_ptr<float> t(energySum);
   float x = thrust::reduce(t, t + grid.x);
   cuda_safe_mem(cudaMemcpy(E, &x, sizeof(float), cudaMemcpyHostToDevice));
 

src/core/actor/DipolarBarnesHut_cuda.cuh

@@ -25,8 +25,6 @@
 
 #ifdef DIPOLAR_BARNES_HUT
 
-typedef float dds_float;
-
 typedef struct {
   /// CUDA blocks
   int blocks;
@@ -117,10 +115,10 @@ void summarizeBH(int blocks);
 void sortBH(int blocks);
 
 /// Barnes-Hut force calculation.
-int forceBH(BHData *bh_data, dds_float k, float *f, float *torque);
+int forceBH(BHData *bh_data, float k, float *f, float *torque);
 
 /// Barnes-Hut energy calculation.
-int energyBH(BHData *bh_data, dds_float k, float *E);
+int energyBH(BHData *bh_data, float k, float *E);
 
 #endif // DIPOLAR_BARNES_HUT
 #endif /* DIPOLARBARNESHUT_CUH_ */