Rename GpuBonded into ListedForcesGpu

The listed forces name describes better what the class is
doing and also follows the naming for other modules.

diff --git a/src/gromacs/domdec/mdsetup.cpp b/src/gromacs/domdec/mdsetup.cpp
index ead9a61fd1c67f6c38f66624a9253ce24e5216e7..47e4203f82948069dd7131c1b06d43cef2b78359 100644 (file)
--- a/src/gromacs/domdec/mdsetup.cpp
+++ b/src/gromacs/domdec/mdsetup.cpp
@@ -137,7 +137,7 @@ void mdAlgorithmsSetupAtomData(const t_commrec*     cr,
 
     for (auto& listedForces : fr->listedForces)
     {
-        listedForces.setup(top->idef, fr->natoms_force, fr->gpuBonded != nullptr);
+        listedForces.setup(top->idef, fr->natoms_force, fr->listedForcesGpu != nullptr);
     }
 
     if (EEL_PME(fr->ic->eeltype) && (cr->duty & DUTY_PME))

diff --git a/src/gromacs/listed_forces/CMakeLists.txt b/src/gromacs/listed_forces/CMakeLists.txt
index c1de86e14b582b06edfbd3ef6ebb55cf97610291..a169778e4cd189d52aef3f2e80fe36dddf06f0af 100644 (file)
--- a/src/gromacs/listed_forces/CMakeLists.txt
+++ b/src/gromacs/listed_forces/CMakeLists.txt
@@ -1,7 +1,7 @@
 #
 # This file is part of the GROMACS molecular simulation package.
 #
-# Copyright (c) 2014,2015,2016,2018,2019,2020, by the GROMACS development team, led by
+# Copyright (c) 2014,2015,2016,2018,2019,2020,2021, by the GROMACS development team, led by
 # Mark Abraham, David van der Spoel, Berk Hess, and Erik Lindahl,
 # and including many others, as listed in the AUTHORS file in the
 # top-level source directory and at http://www.gromacs.org.
@@ -36,7 +36,7 @@ add_library(listed_forces INTERFACE)
 gmx_add_libgromacs_sources(
     bonded.cpp
     disre.cpp
-    gpubonded_impl.cpp
+    listed_forces_gpu_impl.cpp
     listed_forces.cpp
     listed_internal.cpp
     manage_threading.cpp
@@ -48,8 +48,8 @@ gmx_add_libgromacs_sources(
 
 if(GMX_GPU_CUDA)
     gmx_add_libgromacs_sources(
-       gpubonded_impl.cu
-       gpubondedkernels.cu
+       listed_forces_gpu_impl.cu
+       listed_forces_gpu_internal.cu
        )
 endif()
 

diff --git a/src/gromacs/listed_forces/listed_forces.cpp b/src/gromacs/listed_forces/listed_forces.cpp
index b0e40a9474656a3604a9175b6ef067af882851e3..893a1d306cd0ceb17b9e612a76ca8f4fcb2dc061 100644 (file)
--- a/src/gromacs/listed_forces/listed_forces.cpp
+++ b/src/gromacs/listed_forces/listed_forces.cpp
@@ -445,7 +445,7 @@ real calc_one_bond(int                           thread,
     const int nat1   = interaction_function[ftype].nratoms + 1;
     const int nbonds = iatoms.ssize() / nat1;
 
-    GMX_ASSERT(fr->gpuBonded != nullptr || workDivision.end(ftype) == iatoms.ssize(),
+    GMX_ASSERT(fr->listedForcesGpu != nullptr || workDivision.end(ftype) == iatoms.ssize(),
                "The thread division should match the topology");
 
     const int nb0 = workDivision.bound(ftype, thread);

diff --git a/src/gromacs/listed_forces/gpubonded.h b/src/gromacs/listed_forces/listed_forces_gpu.h
similarity index 92%
rename from src/gromacs/listed_forces/gpubonded.h
rename to src/gromacs/listed_forces/listed_forces_gpu.h
index 41ad20d42a58b04a1c754015da0fef7a1dc80117..5cdd8c786fc703696b2219ba34e24a5bd73955bb 100644 (file)
--- a/src/gromacs/listed_forces/gpubonded.h
+++ b/src/gromacs/listed_forces/listed_forces_gpu.h
@@ -46,8 +46,8 @@
  * \inlibraryapi
  * \ingroup module_listed_forces
  */
-#ifndef GMX_LISTED_FORCES_GPUBONDED_H
-#define GMX_LISTED_FORCES_GPUBONDED_H
+#ifndef GMX_LISTED_FORCES_LISTED_FORCES_GPU_H
+#define GMX_LISTED_FORCES_LISTED_FORCES_GPU_H
 
 #include <memory>
 
@@ -95,7 +95,7 @@ constexpr std::array<int, numFTypesOnGpu> fTypesOnGpu = { F_BONDS,  F_ANGLES, F_
  *
  * \throws std::bad_alloc when out of memory.
  */
-bool buildSupportsGpuBondeds(std::string* error);
+bool buildSupportsListedForcesGpu(std::string* error);
 
 /*! \brief Checks whether the input system allows to compute bonded interactions on a GPU.
  *
@@ -105,9 +105,9 @@ bool buildSupportsGpuBondeds(std::string* error);
  *
  * \returns true if PME can run on GPU with this input, false otherwise.
  */
-bool inputSupportsGpuBondeds(const t_inputrec& ir, const gmx_mtop_t& mtop, std::string* error);
+bool inputSupportsListedForcesGpu(const t_inputrec& ir, const gmx_mtop_t& mtop, std::string* error);
 
-class GpuBonded
+class ListedForcesGpu
 {
 public:
     /*! \brief Construct the manager with constant data and the stream to use.
@@ -120,13 +120,13 @@ public:
      * \param[in] wcycle                     The wallclock counter.
      *
      */
-    GpuBonded(const gmx_ffparams_t& ffparams,
-              float                 electrostaticsScaleFactor,
-              const DeviceContext&  deviceContext,
-              const DeviceStream&   deviceStream,
-              gmx_wallcycle*        wcycle);
+    ListedForcesGpu(const gmx_ffparams_t& ffparams,
+                    float                 electrostaticsScaleFactor,
+                    const DeviceContext&  deviceContext,
+                    const DeviceStream&   deviceStream,
+                    gmx_wallcycle*        wcycle);
     //! Destructor
-    ~GpuBonded();
+    ~ListedForcesGpu();
 
     /*! \brief Update lists of interactions from idef suitable for the GPU,
      * using the data structures prepared for PP work.
@@ -204,4 +204,4 @@ private:
 
 } // namespace gmx
 
-#endif
+#endif // GMX_LISTED_FORCES_LISTED_FORCES_GPU_H

diff --git a/src/gromacs/listed_forces/gpubonded_impl.cpp b/src/gromacs/listed_forces/listed_forces_gpu_impl.cpp
similarity index 73%
rename from src/gromacs/listed_forces/gpubonded_impl.cpp
rename to src/gromacs/listed_forces/listed_forces_gpu_impl.cpp
index 1fe74e0082e4a92ca18c8983371e9ffa073cd340..f9874822aba381d279f6100d54b74db0a85ea60e 100644 (file)
--- a/src/gromacs/listed_forces/gpubonded_impl.cpp
+++ b/src/gromacs/listed_forces/listed_forces_gpu_impl.cpp
@@ -48,7 +48,7 @@
 #include <algorithm>
 #include <string>
 
-#include "gromacs/listed_forces/gpubonded.h"
+#include "gromacs/listed_forces/listed_forces_gpu.h"
 #include "gromacs/mdtypes/inputrec.h"
 #include "gromacs/topology/topology.h"
 #include "gromacs/utility/message_string_collector.h"
@@ -93,7 +93,7 @@ static bool bondedInteractionsCanRunOnGpu(const gmx_mtop_t& mtop)
     return false;
 }
 
-bool buildSupportsGpuBondeds(std::string* error)
+bool buildSupportsListedForcesGpu(std::string* error)
 {
     MessageStringCollector errorReasons;
     // Before changing the prefix string, make sure that it is not searched for in regression tests.
@@ -110,7 +110,7 @@ bool buildSupportsGpuBondeds(std::string* error)
     return errorReasons.isEmpty();
 }
 
-bool inputSupportsGpuBondeds(const t_inputrec& ir, const gmx_mtop_t& mtop, std::string* error)
+bool inputSupportsListedForcesGpu(const t_inputrec& ir, const gmx_mtop_t& mtop, std::string* error)
 {
     MessageStringCollector errorReasons;
     // Before changing the prefix string, make sure that it is not searched for in regression tests.
@@ -135,52 +135,52 @@ bool inputSupportsGpuBondeds(const t_inputrec& ir, const gmx_mtop_t& mtop, std::
 
 #if !GMX_GPU_CUDA
 
-class GpuBonded::Impl
+class ListedForcesGpu::Impl
 {
 };
 
-GpuBonded::GpuBonded(const gmx_ffparams_t& /* ffparams */,
-                     const float /* electrostaticsScaleFactor */,
-                     const DeviceContext& /* deviceContext */,
-                     const DeviceStream& /* deviceStream */,
-                     gmx_wallcycle* /* wcycle */) :
+ListedForcesGpu::ListedForcesGpu(const gmx_ffparams_t& /* ffparams */,
+                                 const float /* electrostaticsScaleFactor */,
+                                 const DeviceContext& /* deviceContext */,
+                                 const DeviceStream& /* deviceStream */,
+                                 gmx_wallcycle* /* wcycle */) :
     impl_(nullptr)
 {
 }
 
-GpuBonded::~GpuBonded() = default;
+ListedForcesGpu::~ListedForcesGpu() = default;
 
-void GpuBonded::updateInteractionListsAndDeviceBuffers(ArrayRef<const int> /* nbnxnAtomOrder */,
-                                                       const InteractionDefinitions& /* idef */,
-                                                       void* /* xqDevice */,
-                                                       DeviceBuffer<RVec> /* forceDevice */,
-                                                       DeviceBuffer<RVec> /* fshiftDevice */)
+void ListedForcesGpu::updateInteractionListsAndDeviceBuffers(ArrayRef<const int> /* nbnxnAtomOrder */,
+                                                             const InteractionDefinitions& /* idef */,
+                                                             void* /* xqDevice */,
+                                                             DeviceBuffer<RVec> /* forceDevice */,
+                                                             DeviceBuffer<RVec> /* fshiftDevice */)
 {
 }
 
-void GpuBonded::setPbc(PbcType /* pbcType */, const matrix /* box */, bool /* canMoleculeSpanPbc */)
+void ListedForcesGpu::setPbc(PbcType /* pbcType */, const matrix /* box */, bool /* canMoleculeSpanPbc */)
 {
 }
 
-bool GpuBonded::haveInteractions() const
+bool ListedForcesGpu::haveInteractions() const
 {
     return !impl_;
 }
 
-void GpuBonded::launchKernel(const gmx::StepWorkload& /* stepWork */) {}
+void ListedForcesGpu::launchKernel(const gmx::StepWorkload& /* stepWork */) {}
 
-void GpuBonded::setPbcAndlaunchKernel(PbcType /* pbcType */,
-                                      const matrix /* box */,
-                                      bool /* canMoleculeSpanPbc */,
-                                      const gmx::StepWorkload& /* stepWork */)
+void ListedForcesGpu::setPbcAndlaunchKernel(PbcType /* pbcType */,
+                                            const matrix /* box */,
+                                            bool /* canMoleculeSpanPbc */,
+                                            const gmx::StepWorkload& /* stepWork */)
 {
 }
 
-void GpuBonded::launchEnergyTransfer() {}
+void ListedForcesGpu::launchEnergyTransfer() {}
 
-void GpuBonded::waitAccumulateEnergyTerms(gmx_enerdata_t* /* enerd */) {}
+void ListedForcesGpu::waitAccumulateEnergyTerms(gmx_enerdata_t* /* enerd */) {}
 
-void GpuBonded::clearEnergies() {}
+void ListedForcesGpu::clearEnergies() {}
 
 #endif // !GMX_GPU_CUDA
 

diff --git a/src/gromacs/listed_forces/gpubonded_impl.cu b/src/gromacs/listed_forces/listed_forces_gpu_impl.cu
similarity index 84%
rename from src/gromacs/listed_forces/gpubonded_impl.cu
rename to src/gromacs/listed_forces/listed_forces_gpu_impl.cu
index 90c5449da99e998db9958b0f469a976cee4c2436..d02f53562364c219f3fd1dd58f9babac2043afed 100644 (file)
--- a/src/gromacs/listed_forces/gpubonded_impl.cu
+++ b/src/gromacs/listed_forces/listed_forces_gpu_impl.cu
@@ -46,7 +46,7 @@
 
 #include "gmxpre.h"
 
-#include "gpubonded_impl.h"
+#include "listed_forces_gpu_impl.h"
 
 #include "gromacs/gpu_utils/cuda_arch_utils.cuh"
 #include "gromacs/gpu_utils/cudautils.cuh"
@@ -65,13 +65,13 @@ namespace gmx
 // Number of CUDA threads in a block
 constexpr static int c_threadsPerBlock = 256;
 
-// ---- GpuBonded::Impl
+// ---- ListedForcesGpu::Impl
 
-GpuBonded::Impl::Impl(const gmx_ffparams_t& ffparams,
-                      const float           electrostaticsScaleFactor,
-                      const DeviceContext&  deviceContext,
-                      const DeviceStream&   deviceStream,
-                      gmx_wallcycle*        wcycle) :
+ListedForcesGpu::Impl::Impl(const gmx_ffparams_t& ffparams,
+                            const float           electrostaticsScaleFactor,
+                            const DeviceContext&  deviceContext,
+                            const DeviceStream&   deviceStream,
+                            gmx_wallcycle*        wcycle) :
     deviceContext_(deviceContext), deviceStream_(deviceStream)
 {
     GMX_RELEASE_ASSERT(deviceStream.isValid(),
@@ -124,7 +124,7 @@ GpuBonded::Impl::Impl(const gmx_ffparams_t& ffparams,
             c_numShiftVectors * sizeof(float3) + (c_threadsPerBlock / warp_size) * 3 * sizeof(float);
 }
 
-GpuBonded::Impl::~Impl()
+ListedForcesGpu::Impl::~Impl()
 {
     for (int fType : fTypesOnGpu)
     {
@@ -188,7 +188,7 @@ static inline int roundUpToFactor(const int input, const int factor)
 // TODO Consider whether this function should be a factory method that
 // makes an object that is the only one capable of the device
 // operations needed for the lifetime of an interaction list. This
-// would be harder to misuse than GpuBonded, and exchange the problem
+// would be harder to misuse than ListedForcesGpu, and exchange the problem
 // of naming this method for the problem of what to name the
 // BondedDeviceInteractionListHandler type.
 
@@ -200,11 +200,11 @@ static inline int roundUpToFactor(const int input, const int factor)
  *
  * \todo Use DeviceBuffer for the d_xqPtr.
  */
-void GpuBonded::Impl::updateInteractionListsAndDeviceBuffers(ArrayRef<const int> nbnxnAtomOrder,
-                                                             const InteractionDefinitions& idef,
-                                                             void*                         d_xqPtr,
-                                                             DeviceBuffer<RVec>            d_fPtr,
-                                                             DeviceBuffer<RVec> d_fShiftPtr)
+void ListedForcesGpu::Impl::updateInteractionListsAndDeviceBuffers(ArrayRef<const int> nbnxnAtomOrder,
+                                                                   const InteractionDefinitions& idef,
+                                                                   void*              d_xqPtr,
+                                                                   DeviceBuffer<RVec> d_fPtr,
+                                                                   DeviceBuffer<RVec> d_fShiftPtr)
 {
     // TODO wallcycle sub start
     haveInteractions_ = false;
@@ -293,19 +293,19 @@ void GpuBonded::Impl::updateInteractionListsAndDeviceBuffers(ArrayRef<const int>
     // TODO wallcycle sub stop
 }
 
-void GpuBonded::Impl::setPbc(PbcType pbcType, const matrix box, bool canMoleculeSpanPbc)
+void ListedForcesGpu::Impl::setPbc(PbcType pbcType, const matrix box, bool canMoleculeSpanPbc)
 {
     PbcAiuc pbcAiuc;
     setPbcAiuc(canMoleculeSpanPbc ? numPbcDimensions(pbcType) : 0, box, &pbcAiuc);
     kernelParams_.pbcAiuc = pbcAiuc;
 }
 
-bool GpuBonded::Impl::haveInteractions() const
+bool ListedForcesGpu::Impl::haveInteractions() const
 {
     return haveInteractions_;
 }
 
-void GpuBonded::Impl::launchEnergyTransfer()
+void ListedForcesGpu::Impl::launchEnergyTransfer()
 {
     GMX_ASSERT(haveInteractions_,
                "No GPU bonded interactions, so no energies will be computed, so transfer should "
@@ -317,7 +317,7 @@ void GpuBonded::Impl::launchEnergyTransfer()
     wallcycle_sub_stop(wcycle_, WallCycleSubCounter::LaunchGpuBonded);
 }
 
-void GpuBonded::Impl::waitAccumulateEnergyTerms(gmx_enerdata_t* enerd)
+void ListedForcesGpu::Impl::waitAccumulateEnergyTerms(gmx_enerdata_t* enerd)
 {
     GMX_ASSERT(haveInteractions_,
                "No GPU bonded interactions, so no energies will be computed or transferred, so "
@@ -343,7 +343,7 @@ void GpuBonded::Impl::waitAccumulateEnergyTerms(gmx_enerdata_t* enerd)
     grppener->energyGroupPairTerms[NonBondedEnergyTerms::Coulomb14][0] += vTot_[F_COUL14];
 }
 
-void GpuBonded::Impl::clearEnergies()
+void ListedForcesGpu::Impl::clearEnergies()
 {
     wallcycle_start_nocount(wcycle_, WallCycleCounter::LaunchGpu);
     wallcycle_sub_start_nocount(wcycle_, WallCycleSubCounter::LaunchGpuBonded);
@@ -352,58 +352,58 @@ void GpuBonded::Impl::clearEnergies()
     wallcycle_stop(wcycle_, WallCycleCounter::LaunchGpu);
 }
 
-// ---- GpuBonded
+// ---- ListedForcesGpu
 
-GpuBonded::GpuBonded(const gmx_ffparams_t& ffparams,
-                     const float           electrostaticsScaleFactor,
-                     const DeviceContext&  deviceContext,
-                     const DeviceStream&   deviceStream,
-                     gmx_wallcycle*        wcycle) :
+ListedForcesGpu::ListedForcesGpu(const gmx_ffparams_t& ffparams,
+                                 const float           electrostaticsScaleFactor,
+                                 const DeviceContext&  deviceContext,
+                                 const DeviceStream&   deviceStream,
+                                 gmx_wallcycle*        wcycle) :
     impl_(new Impl(ffparams, electrostaticsScaleFactor, deviceContext, deviceStream, wcycle))
 {
 }
 
-GpuBonded::~GpuBonded() = default;
+ListedForcesGpu::~ListedForcesGpu() = default;
 
-void GpuBonded::updateInteractionListsAndDeviceBuffers(ArrayRef<const int>           nbnxnAtomOrder,
-                                                       const InteractionDefinitions& idef,
-                                                       void*                         d_xq,
-                                                       DeviceBuffer<RVec>            d_f,
-                                                       DeviceBuffer<RVec>            d_fShift)
+void ListedForcesGpu::updateInteractionListsAndDeviceBuffers(ArrayRef<const int> nbnxnAtomOrder,
+                                                             const InteractionDefinitions& idef,
+                                                             void*                         d_xq,
+                                                             DeviceBuffer<RVec>            d_f,
+                                                             DeviceBuffer<RVec>            d_fShift)
 {
     impl_->updateInteractionListsAndDeviceBuffers(nbnxnAtomOrder, idef, d_xq, d_f, d_fShift);
 }
 
-void GpuBonded::setPbc(PbcType pbcType, const matrix box, bool canMoleculeSpanPbc)
+void ListedForcesGpu::setPbc(PbcType pbcType, const matrix box, bool canMoleculeSpanPbc)
 {
     impl_->setPbc(pbcType, box, canMoleculeSpanPbc);
 }
 
-bool GpuBonded::haveInteractions() const
+bool ListedForcesGpu::haveInteractions() const
 {
     return impl_->haveInteractions();
 }
 
-void GpuBonded::setPbcAndlaunchKernel(PbcType                  pbcType,
-                                      const matrix             box,
-                                      bool                     canMoleculeSpanPbc,
-                                      const gmx::StepWorkload& stepWork)
+void ListedForcesGpu::setPbcAndlaunchKernel(PbcType                  pbcType,
+                                            const matrix             box,
+                                            bool                     canMoleculeSpanPbc,
+                                            const gmx::StepWorkload& stepWork)
 {
     setPbc(pbcType, box, canMoleculeSpanPbc);
     launchKernel(stepWork);
 }
 
-void GpuBonded::launchEnergyTransfer()
+void ListedForcesGpu::launchEnergyTransfer()
 {
     impl_->launchEnergyTransfer();
 }
 
-void GpuBonded::waitAccumulateEnergyTerms(gmx_enerdata_t* enerd)
+void ListedForcesGpu::waitAccumulateEnergyTerms(gmx_enerdata_t* enerd)
 {
     impl_->waitAccumulateEnergyTerms(enerd);
 }
 
-void GpuBonded::clearEnergies()
+void ListedForcesGpu::clearEnergies()
 {
     impl_->clearEnergies();
 }

diff --git a/src/gromacs/listed_forces/gpubonded_impl.h b/src/gromacs/listed_forces/listed_forces_gpu_impl.h
similarity index 97%
rename from src/gromacs/listed_forces/gpubonded_impl.h
rename to src/gromacs/listed_forces/listed_forces_gpu_impl.h
index 5e66c52eea3802e77153fbd0860703046010e8fc..902ceb24c400194c0185596874ccde83220d8de6 100644 (file)
--- a/src/gromacs/listed_forces/gpubonded_impl.h
+++ b/src/gromacs/listed_forces/listed_forces_gpu_impl.h
@@ -45,13 +45,13 @@
  *
  * \ingroup module_listed_forces
  */
-#ifndef GMX_LISTED_FORCES_GPUBONDED_IMPL_H
-#define GMX_LISTED_FORCES_GPUBONDED_IMPL_H
+#ifndef GMX_LISTED_FORCES_LISTED_FORCES_GPU_IMPL_H
+#define GMX_LISTED_FORCES_LISTED_FORCES_GPU_IMPL_H
 
 #include "gromacs/gpu_utils/device_context.h"
 #include "gromacs/gpu_utils/gputraits.cuh"
 #include "gromacs/gpu_utils/hostallocator.h"
-#include "gromacs/listed_forces/gpubonded.h"
+#include "gromacs/listed_forces/listed_forces_gpu.h"
 #include "gromacs/pbcutil/pbc_aiuc.h"
 
 struct gmx_ffparams_t;
@@ -122,7 +122,7 @@ struct BondedCudaKernelParameters
 };
 
 /*! \internal \brief Implements GPU bondeds */
-class GpuBonded::Impl
+class ListedForcesGpu::Impl
 {
 public:
     //! Constructor
@@ -211,4 +211,4 @@ private:
 
 } // namespace gmx
 
-#endif
+#endif // GMX_LISTED_FORCES_LISTED_FORCES_GPU_IMPL_H

diff --git a/src/gromacs/listed_forces/gpubondedkernels.cu b/src/gromacs/listed_forces/listed_forces_gpu_internal.cu
similarity index 99%
rename from src/gromacs/listed_forces/gpubondedkernels.cu
rename to src/gromacs/listed_forces/listed_forces_gpu_internal.cu
index 407e447bdc95c08369159a47b597a43e6ffb8658..9b35ea6077440dbde4350fe3939d686766a89e99 100644 (file)
--- a/src/gromacs/listed_forces/gpubondedkernels.cu
+++ b/src/gromacs/listed_forces/listed_forces_gpu_internal.cu
@@ -55,7 +55,7 @@
 #include "gromacs/gpu_utils/cudautils.cuh"
 #include "gromacs/gpu_utils/typecasts.cuh"
 #include "gromacs/gpu_utils/vectype_ops.cuh"
-#include "gromacs/listed_forces/gpubonded.h"
+#include "gromacs/listed_forces/listed_forces_gpu.h"
 #include "gromacs/math/units.h"
 #include "gromacs/mdlib/force_flags.h"
 #include "gromacs/mdtypes/interaction_const.h"
@@ -64,7 +64,7 @@
 #include "gromacs/timing/wallcycle.h"
 #include "gromacs/utility/gmxassert.h"
 
-#include "gpubonded_impl.h"
+#include "listed_forces_gpu_impl.h"
 
 #if defined(_MSVC)
 #    include <limits>
@@ -909,7 +909,7 @@ __global__ void exec_kernel_gpu(BondedCudaKernelParameters kernelParams)
 
 
 template<bool calcVir, bool calcEner>
-void GpuBonded::Impl::launchKernel()
+void ListedForcesGpu::Impl::launchKernel()
 {
     GMX_ASSERT(haveInteractions_,
                "Cannot launch bonded GPU kernels unless bonded GPU work was scheduled");
@@ -939,7 +939,7 @@ void GpuBonded::Impl::launchKernel()
     wallcycle_stop(wcycle_, WallCycleCounter::LaunchGpu);
 }
 
-void GpuBonded::launchKernel(const gmx::StepWorkload& stepWork)
+void ListedForcesGpu::launchKernel(const gmx::StepWorkload& stepWork)
 {
     if (stepWork.computeEnergy)
     {

diff --git a/src/gromacs/listed_forces/manage_threading.cpp b/src/gromacs/listed_forces/manage_threading.cpp
index 0eeac68641bfb4e790c32e29c2a67ee40763ca9b..26f9ce89eec7a96e552f0bba75e8731129462830 100644 (file)
--- a/src/gromacs/listed_forces/manage_threading.cpp
+++ b/src/gromacs/listed_forces/manage_threading.cpp
@@ -57,7 +57,7 @@
 #include <algorithm>
 #include <string>
 
-#include "gromacs/listed_forces/gpubonded.h"
+#include "gromacs/listed_forces/listed_forces_gpu.h"
 #include "gromacs/pbcutil/ishift.h"
 #include "gromacs/topology/ifunc.h"
 #include "gromacs/utility/exceptions.h"

diff --git a/src/gromacs/mdlib/forcerec.cpp b/src/gromacs/mdlib/forcerec.cpp
index 61659e59e4281cc950105e40ae39d04a0c42f4e6..8e8caec27a5c2afa16d81fb432a38501dcb70c41 100644 (file)
--- a/src/gromacs/mdlib/forcerec.cpp
+++ b/src/gromacs/mdlib/forcerec.cpp
@@ -57,7 +57,7 @@
 #include "gromacs/gmxlib/nonbonded/nonbonded.h"
 #include "gromacs/gpu_utils/gpu_utils.h"
 #include "gromacs/hardware/hw_info.h"
-#include "gromacs/listed_forces/gpubonded.h"
+#include "gromacs/listed_forces/listed_forces_gpu.h"
 #include "gromacs/listed_forces/listed_forces.h"
 #include "gromacs/listed_forces/pairs.h"
 #include "gromacs/math/functions.h"

diff --git a/src/gromacs/mdlib/sim_util.cpp b/src/gromacs/mdlib/sim_util.cpp
index ea4d010d7a187193fd7102b31f5e5987a1e33099..e73d83fdac694843f12671b81f8e6101ba7d0108 100644 (file)
--- a/src/gromacs/mdlib/sim_util.cpp
+++ b/src/gromacs/mdlib/sim_util.cpp
@@ -63,7 +63,7 @@
 #include "gromacs/gpu_utils/gpu_utils.h"
 #include "gromacs/imd/imd.h"
 #include "gromacs/listed_forces/disre.h"
-#include "gromacs/listed_forces/gpubonded.h"
+#include "gromacs/listed_forces/listed_forces_gpu.h"
 #include "gromacs/listed_forces/listed_forces.h"
 #include "gromacs/listed_forces/orires.h"
 #include "gromacs/math/arrayrefwithpadding.h"
@@ -925,7 +925,8 @@ static DomainLifetimeWorkload setupDomainLifetimeWorkload(const t_inputrec&
             domainWork.haveCpuBondedWork = true;
         }
     }
-    domainWork.haveGpuBondedWork = ((fr.gpuBonded != nullptr) && fr.gpuBonded->haveInteractions());
+    domainWork.haveGpuBondedWork =
+            ((fr.listedForcesGpu != nullptr) && fr.listedForcesGpu->haveInteractions());
     // Note that haveFreeEnergyWork is constant over the whole run
     domainWork.haveFreeEnergyWork =
             (fr.efep != FreeEnergyPerturbationType::No && mdatoms.nPerturbed != 0);
@@ -995,7 +996,7 @@ static StepWorkload setupStepWorkload(const int                     legacyFlags,
  *
  */
 static void launchGpuEndOfStepTasks(nonbonded_verlet_t*               nbv,
-                                    gmx::GpuBonded*                   gpuBonded,
+                                    gmx::ListedForcesGpu*             listedForcesGpu,
                                     gmx_pme_t*                        pmedata,
                                     gmx_enerdata_t*                   enerd,
                                     const gmx::MdrunScheduleWorkload& runScheduleWork,
@@ -1031,9 +1032,9 @@ static void launchGpuEndOfStepTasks(nonbonded_verlet_t*               nbv,
         // in principle this should be included in the DD balancing region,
         // but generally it is infrequent so we'll omit it for the sake of
         // simpler code
-        gpuBonded->waitAccumulateEnergyTerms(enerd);
+        listedForcesGpu->waitAccumulateEnergyTerms(enerd);
 
-        gpuBonded->clearEnergies();
+        listedForcesGpu->clearEnergies();
     }
 }
 
@@ -1414,14 +1415,14 @@ void do_force(FILE*                               fplog,
         /* initialize the GPU nbnxm atom data and bonded data structures */
         if (simulationWork.useGpuNonbonded)
         {
-            // Note: cycle counting only nononbondeds, gpuBonded counts internally
+            // Note: cycle counting only nononbondeds, GPU listed forces counts internally
             wallcycle_start_nocount(wcycle, WallCycleCounter::LaunchGpu);
             wallcycle_sub_start_nocount(wcycle, WallCycleSubCounter::LaunchGpuNonBonded);
             Nbnxm::gpu_init_atomdata(nbv->gpu_nbv, nbv->nbat.get());
             wallcycle_sub_stop(wcycle, WallCycleSubCounter::LaunchGpuNonBonded);
             wallcycle_stop(wcycle, WallCycleCounter::LaunchGpu);
 
-            if (fr->gpuBonded)
+            if (fr->listedForcesGpu)
             {
                 /* Now we put all atoms on the grid, we can assign bonded
                  * interactions to the GPU, where the grid order is
@@ -1431,11 +1432,12 @@ void do_force(FILE*                               fplog,
                 // TODO the xq, f, and fshift buffers are now shared
                 // resources, so they should be maintained by a
                 // higher-level object than the nb module.
-                fr->gpuBonded->updateInteractionListsAndDeviceBuffers(nbv->getGridIndices(),
-                                                                      top->idef,
-                                                                      Nbnxm::gpu_get_xq(nbv->gpu_nbv),
-                                                                      Nbnxm::gpu_get_f(nbv->gpu_nbv),
-                                                                      Nbnxm::gpu_get_fshift(nbv->gpu_nbv));
+                fr->listedForcesGpu->updateInteractionListsAndDeviceBuffers(
+                        nbv->getGridIndices(),
+                        top->idef,
+                        Nbnxm::gpu_get_xq(nbv->gpu_nbv),
+                        Nbnxm::gpu_get_f(nbv->gpu_nbv),
+                        Nbnxm::gpu_get_fshift(nbv->gpu_nbv));
             }
         }
 
@@ -1449,7 +1451,7 @@ void do_force(FILE*                               fplog,
         /* Note that with a GPU the launch overhead of the list transfer is not timed separately */
         nbv->constructPairlist(InteractionLocality::Local, top->excls, step, nrnb);
 
-        nbv->setupGpuShortRangeWork(fr->gpuBonded, InteractionLocality::Local);
+        nbv->setupGpuShortRangeWork(fr->listedForcesGpu, InteractionLocality::Local);
 
         wallcycle_sub_stop(wcycle, WallCycleSubCounter::NBSSearchLocal);
         wallcycle_stop(wcycle, WallCycleCounter::NS);
@@ -1503,7 +1505,7 @@ void do_force(FILE*                               fplog,
         // we can only launch the kernel after non-local coordinates have been received.
         if (domainWork.haveGpuBondedWork && !havePPDomainDecomposition(cr))
         {
-            fr->gpuBonded->setPbcAndlaunchKernel(fr->pbcType, box, fr->bMolPBC, stepWork);
+            fr->listedForcesGpu->setPbcAndlaunchKernel(fr->pbcType, box, fr->bMolPBC, stepWork);
         }
 
         /* launch local nonbonded work on GPU */
@@ -1538,7 +1540,7 @@ void do_force(FILE*                               fplog,
             /* Note that with a GPU the launch overhead of the list transfer is not timed separately */
             nbv->constructPairlist(InteractionLocality::NonLocal, top->excls, step, nrnb);
 
-            nbv->setupGpuShortRangeWork(fr->gpuBonded, InteractionLocality::NonLocal);
+            nbv->setupGpuShortRangeWork(fr->listedForcesGpu, InteractionLocality::NonLocal);
             wallcycle_sub_stop(wcycle, WallCycleSubCounter::NBSSearchNonLocal);
             wallcycle_stop(wcycle, WallCycleCounter::NS);
             // TODO refactor this GPU halo exchange re-initialisation
@@ -1606,7 +1608,7 @@ void do_force(FILE*                               fplog,
 
             if (domainWork.haveGpuBondedWork)
             {
-                fr->gpuBonded->setPbcAndlaunchKernel(fr->pbcType, box, fr->bMolPBC, stepWork);
+                fr->listedForcesGpu->setPbcAndlaunchKernel(fr->pbcType, box, fr->bMolPBC, stepWork);
             }
 
             /* launch non-local nonbonded tasks on GPU */
@@ -1633,7 +1635,7 @@ void do_force(FILE*                               fplog,
 
         if (domainWork.haveGpuBondedWork && stepWork.computeEnergy)
         {
-            fr->gpuBonded->launchEnergyTransfer();
+            fr->listedForcesGpu->launchEnergyTransfer();
         }
         wallcycle_stop(wcycle, WallCycleCounter::LaunchGpu);
     }
@@ -2279,7 +2281,7 @@ void do_force(FILE*                               fplog,
         }
     }
 
-    launchGpuEndOfStepTasks(nbv, fr->gpuBonded, fr->pmedata, enerd, *runScheduleWork, step, wcycle);
+    launchGpuEndOfStepTasks(nbv, fr->listedForcesGpu, fr->pmedata, enerd, *runScheduleWork, step, wcycle);
 
     if (DOMAINDECOMP(cr))
     {

diff --git a/src/gromacs/mdrun/runner.cpp b/src/gromacs/mdrun/runner.cpp
index ec9c382e64c4fa932eb30507a8f5b7bc9f544338..c4987f1cfc2e813ed5457824cf190d0c4777bfce 100644 (file)
--- a/src/gromacs/mdrun/runner.cpp
+++ b/src/gromacs/mdrun/runner.cpp
@@ -82,7 +82,7 @@
 #include "gromacs/hardware/printhardware.h"
 #include "gromacs/imd/imd.h"
 #include "gromacs/listed_forces/disre.h"
-#include "gromacs/listed_forces/gpubonded.h"
+#include "gromacs/listed_forces/listed_forces_gpu.h"
 #include "gromacs/listed_forces/listed_forces.h"
 #include "gromacs/listed_forces/orires.h"
 #include "gromacs/math/functions.h"
@@ -1599,7 +1599,7 @@ int Mdrunner::mdrunner()
     const bool               thisRankHasPmeGpuTask = gpuTaskAssignments.thisRankHasPmeGpuTask();
     std::unique_ptr<MDAtoms> mdAtoms;
     std::unique_ptr<VirtualSitesHandler> vsite;
-    std::unique_ptr<GpuBonded>           gpuBonded;
+    std::unique_ptr<ListedForcesGpu>     listedForcesGpu;
 
     t_nrnb nrnb;
     if (thisRankHasDuty(cr, DUTY_PP))
@@ -1664,13 +1664,13 @@ int Mdrunner::mdrunner()
             GMX_RELEASE_ASSERT(deviceStreamManager != nullptr,
                                "GPU device stream manager should be valid in order to use GPU "
                                "version of bonded forces.");
-            gpuBonded = std::make_unique<GpuBonded>(
+            listedForcesGpu = std::make_unique<ListedForcesGpu>(
                     mtop.ffparams,
                     fr->ic->epsfac * fr->fudgeQQ,
                     deviceStreamManager->context(),
                     deviceStreamManager->bondedStream(havePPDomainDecomposition(cr)),
                     wcycle.get());
-            fr->gpuBonded = gpuBonded.get();
+            fr->listedForcesGpu = listedForcesGpu.get();
         }
 
         /* Initialize the mdAtoms structure.
@@ -2060,7 +2060,7 @@ int Mdrunner::mdrunner()
     mdAtoms.reset(nullptr);
     globalState.reset(nullptr);
     mdModules_.reset(nullptr); // destruct force providers here as they might also use the GPU
-    gpuBonded.reset(nullptr);
+    listedForcesGpu.reset(nullptr);
     fr.reset(nullptr); // destruct forcerec before gpu
     // TODO convert to C++ so we can get rid of these frees
     sfree(disresdata);

diff --git a/src/gromacs/mdtypes/forcerec.h b/src/gromacs/mdtypes/forcerec.h
index 218c7b26c9a4dd9edb62135577e8d7fa31494d4d..8a1ebe14e554cb4796af1fd228eadf796b7c7c6e 100644 (file)
--- a/src/gromacs/mdtypes/forcerec.h
+++ b/src/gromacs/mdtypes/forcerec.h
@@ -65,7 +65,7 @@ struct interaction_const_t;
 namespace gmx
 {
 class DeviceStreamManager;
-class GpuBonded;
+class ListedForcesGpu;
 class GpuForceReduction;
 class ForceProviders;
 class StatePropagatorDataGpu;
@@ -297,7 +297,7 @@ struct t_forcerec
     std::vector<ListedForces> listedForces;
 
     /* TODO: Replace the pointer by an object once we got rid of C */
-    gmx::GpuBonded* gpuBonded = nullptr;
+    gmx::ListedForcesGpu* listedForcesGpu = nullptr;
 
     /* Ewald correction thread local virial and energy data */
     int                              nthread_ewc = 0;

diff --git a/src/gromacs/nbnxm/gpu_common.h b/src/gromacs/nbnxm/gpu_common.h
index bf9c25a5e70a16db2942bebf4822ee763a200fce..13aea52f4b3f093ac5069a89f0a2c1df4c9b44e6 100644 (file)
--- a/src/gromacs/nbnxm/gpu_common.h
+++ b/src/gromacs/nbnxm/gpu_common.h
@@ -73,7 +73,7 @@
 
 namespace gmx
 {
-class GpuBonded;
+class ListedForcesGpu;
 }
 
 namespace Nbnxm

diff --git a/src/gromacs/nbnxm/gpu_common_utils.h b/src/gromacs/nbnxm/gpu_common_utils.h
index 0d81e2f05a4dcb786aa91ec5752ea10db65b27f6..2ede311e38ef39ed6d9aec6a5f5487e013f6b722 100644 (file)
--- a/src/gromacs/nbnxm/gpu_common_utils.h
+++ b/src/gromacs/nbnxm/gpu_common_utils.h
@@ -42,7 +42,7 @@
 #ifndef GMX_NBNXM_GPU_COMMON_UTILS_H
 #define GMX_NBNXM_GPU_COMMON_UTILS_H
 
-#include "gromacs/listed_forces/gpubonded.h"
+#include "gromacs/listed_forces/listed_forces_gpu.h"
 #include "gromacs/mdtypes/locality.h"
 #include "gromacs/nbnxm/gpu_types_common.h"
 #include "gromacs/utility/exceptions.h"

diff --git a/src/gromacs/nbnxm/nbnxm.cpp b/src/gromacs/nbnxm/nbnxm.cpp
index 361483b84c1e529830742e6af993f73e0a27655a..e0aae3761fb68a88827323b7b9d03770ea63b7df 100644 (file)
--- a/src/gromacs/nbnxm/nbnxm.cpp
+++ b/src/gromacs/nbnxm/nbnxm.cpp
@@ -229,12 +229,12 @@ void nonbonded_verlet_t::changePairlistRadii(real rlistOuter, real rlistInner) c
     pairlistSets_->changePairlistRadii(rlistOuter, rlistInner);
 }
 
-void nonbonded_verlet_t::setupGpuShortRangeWork(const gmx::GpuBonded*          gpuBonded,
+void nonbonded_verlet_t::setupGpuShortRangeWork(const gmx::ListedForcesGpu*    listedForcesGpu,
                                                 const gmx::InteractionLocality iLocality) const
 {
     if (useGpu() && !emulateGpu())
     {
-        Nbnxm::setupGpuShortRangeWork(gpu_nbv, gpuBonded, iLocality);
+        Nbnxm::setupGpuShortRangeWork(gpu_nbv, listedForcesGpu, iLocality);
     }
 }
 

diff --git a/src/gromacs/nbnxm/nbnxm.h b/src/gromacs/nbnxm/nbnxm.h
index dfaaff41dc3a52a1ec818855d81dec37688cda9c..545f85ef4af624aa59d9ea9368b8087bee5332ae 100644 (file)
--- a/src/gromacs/nbnxm/nbnxm.h
+++ b/src/gromacs/nbnxm/nbnxm.h
@@ -146,7 +146,7 @@ namespace gmx
 {
 class DeviceStreamManager;
 class ForceWithShiftForces;
-class GpuBonded;
+class ListedForcesGpu;
 template<typename>
 class ListOfLists;
 class MDLogger;
@@ -415,7 +415,8 @@ public:
     void changePairlistRadii(real rlistOuter, real rlistInner) const;
 
     //! Set up internal flags that indicate what type of short-range work there is.
-    void setupGpuShortRangeWork(const gmx::GpuBonded* gpuBonded, gmx::InteractionLocality iLocality) const;
+    void setupGpuShortRangeWork(const gmx::ListedForcesGpu* listedForcesGpu,
+                                gmx::InteractionLocality    iLocality) const;
 
     // TODO: Make all data members private
     //! All data related to the pair lists

diff --git a/src/gromacs/nbnxm/nbnxm_gpu.h b/src/gromacs/nbnxm/nbnxm_gpu.h
index 1cf920d35cf85eb08856387a64073e775820e4cf..2a49837c8fc753ee1997330cd6aedbc11e476fb6 100644 (file)
--- a/src/gromacs/nbnxm/nbnxm_gpu.h
+++ b/src/gromacs/nbnxm/nbnxm_gpu.h
@@ -60,7 +60,7 @@ enum class GpuTaskCompletion;
 
 namespace gmx
 {
-class GpuBonded;
+class ListedForcesGpu;
 class StepWorkload;
 } // namespace gmx
 
@@ -282,14 +282,14 @@ void nbnxnInsertNonlocalGpuDependency(NbnxmGpu gmx_unused*                nb,
  * This function is expected to be called every time the work-distribution
  * can change (i.e. at search/domain decomposition steps).
  *
- * \param[inout]  nb         Pointer to the nonbonded GPU data structure
- * \param[in]     gpuBonded  Pointer to the GPU bonded data structure
- * \param[in]     iLocality  Interaction locality identifier
+ * \param[inout]  nb               Pointer to the nonbonded GPU data structure
+ * \param[in]     listedForcesGpu  Pointer to the GPU bonded data structure
+ * \param[in]     iLocality        Interaction locality identifier
  */
 GPU_FUNC_QUALIFIER
-void setupGpuShortRangeWork(NbnxmGpu gmx_unused* nb,
-                            const gmx::GpuBonded gmx_unused*    gpuBonded,
-                            gmx::InteractionLocality gmx_unused iLocality) GPU_FUNC_TERM;
+void setupGpuShortRangeWork(NbnxmGpu gmx_unused*       nb,
+                            const gmx::ListedForcesGpu gmx_unused* listedForcesGpu,
+                            gmx::InteractionLocality gmx_unused    iLocality) GPU_FUNC_TERM;
 
 /*! \brief Returns true if there is GPU short-range work for the given interaction locality.
  *

diff --git a/src/gromacs/nbnxm/nbnxm_gpu_data_mgmt.cpp b/src/gromacs/nbnxm/nbnxm_gpu_data_mgmt.cpp
index 6e7a5b799f109d9161cfc88e15191ea20932c4dc..874f9a614f190c4c38023bc3e3ec05ce80e5ba32 100644 (file)
--- a/src/gromacs/nbnxm/nbnxm_gpu_data_mgmt.cpp
+++ b/src/gromacs/nbnxm/nbnxm_gpu_data_mgmt.cpp
@@ -746,7 +746,9 @@ bool gpu_is_kernel_ewald_analytical(const NbnxmGpu* nb)
             || (nb->nbparam->elecType == ElecType::EwaldAnaTwin));
 }
 
-void setupGpuShortRangeWork(NbnxmGpu* nb, const gmx::GpuBonded* gpuBonded, const gmx::InteractionLocality iLocality)
+void setupGpuShortRangeWork(NbnxmGpu*                      nb,
+                            const gmx::ListedForcesGpu*    listedForcesGpu,
+                            const gmx::InteractionLocality iLocality)
 {
     GMX_ASSERT(nb, "Need a valid nbnxn_gpu object");
 
@@ -754,7 +756,7 @@ void setupGpuShortRangeWork(NbnxmGpu* nb, const gmx::GpuBonded* gpuBonded, const
     // interaction locality contains entries or if there is any
     // bonded work (as this is not split into local/nonlocal).
     nb->haveWork[iLocality] = ((nb->plist[iLocality]->nsci != 0)
-                               || (gpuBonded != nullptr && gpuBonded->haveInteractions()));
+                               || (listedForcesGpu != nullptr && listedForcesGpu->haveInteractions()));
 }
 
 bool haveGpuShortRangeWork(const NbnxmGpu* nb, const gmx::InteractionLocality interactionLocality)

diff --git a/src/gromacs/pbcutil/pbc_aiuc_cuda.cuh b/src/gromacs/pbcutil/pbc_aiuc_cuda.cuh
index eb12862bf4fff6d163b4a437a421ec6bd81c44ea..405321490e28e4c35397ae157b30d6d484fb4b9b 100644 (file)
--- a/src/gromacs/pbcutil/pbc_aiuc_cuda.cuh
+++ b/src/gromacs/pbcutil/pbc_aiuc_cuda.cuh
@@ -82,7 +82,7 @@ static inline __device__ int int3ToShiftIndex(int3 iv)
  * \todo This routine uses CUDA float4 types for input coordinates and
  *       returns in rvec data-type. Other than that, it does essentially
  *       the same thing as the version below, as well as SIMD and CPU
- *       versions. This routine is used in gpubonded module.
+ *       versions. This routine is used in GPU listed forces module.
  *       To avoid code duplication, these implementations should be
  *       unified. See Issue #2863:
  *       https://gitlab.com/gromacs/gromacs/-/issues/2863

diff --git a/src/gromacs/taskassignment/decidegpuusage.cpp b/src/gromacs/taskassignment/decidegpuusage.cpp
index c18bba3ec402ef5eb55df76454dc1ccfdc22b9c5..c407f61d5d2ff0678965d4ac6b152b1cb5ba095e 100644 (file)
--- a/src/gromacs/taskassignment/decidegpuusage.cpp
+++ b/src/gromacs/taskassignment/decidegpuusage.cpp
@@ -57,7 +57,7 @@
 #include "gromacs/hardware/detecthardware.h"
 #include "gromacs/hardware/hardwaretopology.h"
 #include "gromacs/hardware/hw_info.h"
-#include "gromacs/listed_forces/gpubonded.h"
+#include "gromacs/listed_forces/listed_forces_gpu.h"
 #include "gromacs/mdlib/gmx_omp_nthreads.h"
 #include "gromacs/mdlib/update_constrain_gpu.h"
 #include "gromacs/mdtypes/commrec.h"
@@ -476,7 +476,7 @@ bool decideWhetherToUseGpusForBonded(bool              useGpuForNonbonded,
 
     std::string errorMessage;
 
-    if (!buildSupportsGpuBondeds(&errorMessage))
+    if (!buildSupportsListedForcesGpu(&errorMessage))
     {
         if (bondedTarget == TaskTarget::Gpu)
         {
@@ -486,7 +486,7 @@ bool decideWhetherToUseGpusForBonded(bool              useGpuForNonbonded,
         return false;
     }
 
-    if (!inputSupportsGpuBondeds(inputrec, mtop, &errorMessage))
+    if (!inputSupportsListedForcesGpu(inputrec, mtop, &errorMessage))
     {
         if (bondedTarget == TaskTarget::Gpu)
         {