Refactor wall potential, Lennard-Jones 9-3 potential

doc/index.rst

@@ -56,6 +56,7 @@ Contents
     module-azplugins-external
     module-azplugins-flow
     module-azplugins-pair
+    module-azplugins-wall
 
 .. toctree::
     :maxdepth: 1

doc/module-azplugins-wall.rst

@@ -0,0 +1,23 @@
+.. Copyright (c) 2018-2020, Michael P. Howard
+.. Copyright (c) 2021-2025, Auburn University
+.. Part of azplugins, released under the BSD 3-Clause License.
+
+azplugins.wall
+--------------
+
+.. rubric:: Overview
+
+.. py:currentmodule:: hoomd.azplugins.wall
+
+.. autosummary::
+    :nosignatures:
+
+    LJ93
+
+.. rubric:: Details
+
+.. automodule:: hoomd.azplugins.wall
+    :synopsis: Wall potentials.
+    :members: LJ93
+    :no-inherited-members:
+    :show-inheritance:

src/CMakeLists.txt

@@ -32,6 +32,7 @@ set(python_files
     external.py
     flow.py
     pair.py
+    wall.py
     )
 
 # TODO: Add names of all bond evaluators
@@ -48,6 +49,11 @@ set(_pair_evaluators
     PerturbedLennardJones
     )
 
+# TODO: Add names of all wall evaluators
+set(_wall_evaluators
+    LJ93
+    )
+
 # TODO: Add names of all dpd evaluators
 set(_dpd_evaluators
     GeneralWeight
@@ -96,6 +102,25 @@ foreach(_evaluator ${_pair_evaluators})
     endif()
 endforeach()
 
+# process wall potentials
+foreach(_evaluator ${_wall_evaluators})
+    configure_file(export_PotentialExternalWall.cc.inc
+                   export_PotentialExternalWall${_evaluator}.cc
+                   @ONLY)
+    set(_${COMPONENT_NAME}_sources ${_${COMPONENT_NAME}_sources} export_PotentialExternalWall${_evaluator}.cc)
+
+    if (ENABLE_HIP)
+        configure_file(export_PotentialExternalWallGPU.cc.inc
+                       export_PotentialExternalWall${_evaluator}GPU.cc
+                       @ONLY)
+        configure_file(PotentialExternalWallGPUKernel.cu.inc
+                       PotentialExternalWall${_evaluator}GPUKernel.cu
+                       @ONLY)
+        set(_${COMPONENT_NAME}_sources ${_${COMPONENT_NAME}_sources} export_PotentialExternalWall${_evaluator}GPU.cc)
+        set(_${COMPONENT_NAME}_cu_sources ${_${COMPONENT_NAME}_cu_sources} PotentialExternalWall${_evaluator}GPUKernel.cu)
+    endif()
+endforeach()
+
 # process DPD potentials
 foreach(_evaluator ${_dpd_evaluators})
     configure_file(export_PotentialPairDPDThermo.cc.inc

src/PotentialExternalWallGPUKernel.cu.inc

@@ -0,0 +1,33 @@
+// Copyright (c) 2018-2020, Michael P. Howard
+// Copyright (c) 2021-2025, Auburn University
+// Part of azplugins, released under the BSD 3-Clause License.
+
+// Adapted from hoomd/md/export_PotentialExternalWallGPUKernel.cu.inc of HOOMD-blue.
+// Copyright (c) 2009-2024 The Regents of the University of Michigan.
+// Part of HOOMD-blue, released under the BSD 3-Clause License.
+
+// See CMakeLists.txt for the source of these variables to be processed by CMake's
+// configure_file().
+
+// clang-format off
+#include "hoomd/md/PotentialExternalGPU.cuh"
+#include "hoomd/md/EvaluatorWalls.h"
+#include "WallEvaluator@_evaluator@.h"
+
+#define EVALUATOR_CLASS hoomd::md::EvaluatorWalls<hoomd::azplugins::detail::WallEvaluator@_evaluator@>
+// clang-format on
+
+namespace hoomd
+    {
+namespace md
+    {
+namespace kernel
+    {
+template __attribute__((visibility("default"))) hipError_t
+gpu_compute_potential_external_forces<EVALUATOR_CLASS>(
+    const external_potential_args_t& external_potential_args,
+    const typename EVALUATOR_CLASS::param_type* d_params,
+    const typename EVALUATOR_CLASS::field_type* d_field);
+    } // end namespace kernel
+    } // end namespace md
+    } // end namespace hoomd

src/WallEvaluatorLJ93.h

@@ -10,91 +10,94 @@
 #ifndef AZPLUGINS_WALL_EVALUATOR_LJ_93_H_
 #define AZPLUGINS_WALL_EVALUATOR_LJ_93_H_
 
-#ifndef NVCC
-#include <string>
-#endif
-
-#include "hoomd/HOOMDMath.h"
+#include "PairEvaluator.h"
 
-#ifdef NVCC
+#ifdef __HIPCC__
 #define DEVICE __device__
 #else
 #define DEVICE
 #endif
 
+namespace hoomd
+    {
 namespace azplugins
     {
 namespace detail
     {
-//! Evaluates the Lennard-Jones 9-3 wall force
+//! Define the paramter type used by this wall potential evaluator
+struct WallParametersLJ93 : public PairParameters
+    {
+#ifndef __HIPCC__
+    WallParametersLJ93() : sigma_3(0), A(0) { }
+
+    WallParametersLJ93(pybind11::dict v, bool managed = false)
+        {
+        auto sigma(v["sigma"].cast<Scalar>());
+        A = v["A"].cast<Scalar>();
+
+        sigma_3 = sigma * sigma * sigma;
+        }
+
+    pybind11::dict asDict()
+        {
+        pybind11::dict v;
+        v["sigma"] = std::cbrt(sigma_3);
+        v["A"] = A;
+        return v;
+        }
+#endif // __HIPCC__
+
+    Scalar sigma_3; //!< Lennard-Jones sigma
+    Scalar A;       //!< Hamaker constant
+    }
+
+#if HOOMD_LONGREAL_SIZE == 32
+    __attribute__((aligned(8)));
+#else
+    __attribute__((aligned(16)));
+#endif
+
+//! Class for evaluating the Lennard-Jones 9-3 wall force
 /*!
  * WallEvaluatorLJ93 computes the Lennard-Jones 9-3 wall potential, which is derived from
  * integrating the standard Lennard-Jones potential between a point particle and a half plane:
  *
- * \f[ V(r) = \varepsilon \left( \frac{2}{15}\left(\frac{\sigma}{r}\right)^9 -
- * \left(\frac{\sigma}{r}\right)^3 \right) \f]
+ * \f[ V(r) = A \left[ \frac{2}{15}\left(\frac{\sigma}{r}\right)^9 -
+ * \left(\frac{\sigma}{r}\right)^3 \right] \f]
  *
- * where \f$\sigma\f$ is the diameter of Lennard-Jones particles in the wall, and \f$ \varepsilon
- * \f$ is the effective Hamaker constant \f$ \varepsilon = (2/3) \pi \varepsilon_{\rm LJ} \rho_{\rm
- * w} \sigma^3 \f$ with \f$\varepsilon_{\rm LJ}\f$ the energy of interaction and \f$\rho_{\rm w}\f$
- * the density of particles in the wall. Evaluation of this energy is simplified into the following
+ * where \f$\sigma\f$ is the diameter of Lennard-Jones particles in the wall, and \f$ A \f$ is the
+ * Hamaker constant \f$ A = (2/3) \pi \varepsilon_{\rm LJ} \rho_{\rm w} \sigma^3 \f$ with
+ * \f$\varepsilon_{\rm LJ}\f$ the energy of interaction and \f$\rho_{\rm w}\f$ the density of
+ * particles in the wall. Evaluation of this energy is simplified into the following
  * parameters:
  *
- * - \verbatim lj1 = (2.0/15.0) * epsilon * pow(sigma,9.0) \endverbatim
- * - \verbatim lj2 = epsilon * pow(sigma,3.0) \endverbatim
+ * - \verbatim lj1 = (2.0/15.0) * A * pow(sigma,9.0) \endverbatim
+ * - \verbatim lj2 = A * pow(sigma,3.0) \endverbatim
  *
  * The force acting on the particle is then
- * \f[ F(r)/r = \frac{\varepsilon}{r^2} \left ( \frac{6}{5}\left(\frac{\sigma}{r}\right)^9 - 3
+ * \f[ F(r)/r = \frac{A}{r^2} \left ( \frac{6}{5}\left(\frac{\sigma}{r}\right)^9 - 3
  * \left(\frac{\sigma}{r}\right)^3 \right) \f]
  */
-class WallEvaluatorLJ93
+class WallEvaluatorLJ93 : public PairEvaluator
     {
     public:
-    //! Define the parameter type used by this wall potential evaluator
-    typedef Scalar2 param_type;
+    typedef WallParametersLJ93 param_type;
 
     //! Constructor
     /*!
-     * \param _rsq Squared distance between particles
+     * \param _rsq Sqaured distance between particles
      * \param _rcutsq Cutoff radius squared
-     * \param _params Pair potential parameters, given by typedef above
+     * \param _params Wall potential paramters, given by typedef above
      *
      * The functor initializes its members from \a _params.
      */
     DEVICE WallEvaluatorLJ93(Scalar _rsq, Scalar _rcutsq, const param_type& _params)
-        : rsq(_rsq), rcutsq(_rcutsq), lj1(_params.x), lj2(_params.y)
-        {
-        }
-
-    //! LJ 9-3 doesn't use diameter
-    DEVICE static bool needsDiameter()
+        : PairEvaluator(_rsq, _rcutsq)
         {
-        return false;
+        lj1 = (Scalar(2.0) / Scalar(15.0)) * _params.A * _params.sigma_3 * _params.sigma_3
+              * _params.sigma_3;
+        lj2 = _params.A * _params.sigma_3;
         }
-    //! Accept the optional diameter values
-    /*!
-     * \param di Diameter of particle
-     * \param dj Dummy diameter
-     *
-     * \note The way HOOMD computes wall forces by recycling evaluators requires that we give
-     *       a second diameter, even though this is meaningless for the potential.
-     */
-    DEVICE void setDiameter(Scalar di, Scalar dj) { }
-
-    //! LJ 9-3 doesn't use charge
-    DEVICE static bool needsCharge()
-        {
-        return false;
-        }
-    //! Accept the optional charge values
-    /*!
-     * \param qi Charge of particle
-     * \param qj Dummy charge
-     *
-     * \note The way HOOMD computes wall forces by recycling evaluators requires that we give
-     *       a second charge, even though this is meaningless for the potential.
-     */
-    DEVICE void setCharge(Scalar qi, Scalar qj) { }
 
     //! Evaluate the force and energy
     /*!
@@ -135,22 +138,22 @@ class WallEvaluatorLJ93
             return false;
         }
 
-#ifndef NVCC
+#ifndef __HIPCC__
     //! Return the name of this potential
     static std::string getName()
         {
         return std::string("lj93");
         }
 #endif
 
-    protected:
-    Scalar rsq;    //!< Stored rsq from the constructor
-    Scalar rcutsq; //!< Stored rcutsq from the constructor
-    Scalar lj1;    //!< lj1 parameter extracted from the params passed to the constructor
-    Scalar lj2;    //!< lj2 parameter extracted from the params passed to the constructor
+    private:
+    Scalar lj1;
+    Scalar lj2;
     };
+
     } // end namespace detail
     } // end namespace azplugins
+    } // end namespace hoomd
 
 #undef DEVICE
 #endif // AZPLUGINS_WALL_EVALUATOR_LJ_93_H_