From f8b4e7ccb98648371fd782d8b6aa7e94238e7f97 Mon Sep 17 00:00:00 2001 From: Artem Zhmurov Date: Wed, 26 May 2021 12:05:03 +0300 Subject: [PATCH] Rename GpuBonded into ListedForcesGpu The listed forces name describes better what the class is doing and also follows the naming for other modules. --- src/gromacs/domdec/mdsetup.cpp | 2 +- src/gromacs/listed_forces/CMakeLists.txt | 8 +- src/gromacs/listed_forces/listed_forces.cpp | 2 +- .../{gpubonded.h => listed_forces_gpu.h} | 24 +++--- ...ed_impl.cpp => listed_forces_gpu_impl.cpp} | 50 ++++++------ ...nded_impl.cu => listed_forces_gpu_impl.cu} | 80 +++++++++---------- ...bonded_impl.h => listed_forces_gpu_impl.h} | 10 +-- ...rnels.cu => listed_forces_gpu_internal.cu} | 8 +- .../listed_forces/manage_threading.cpp | 2 +- src/gromacs/mdlib/forcerec.cpp | 2 +- src/gromacs/mdlib/sim_util.cpp | 38 ++++----- src/gromacs/mdrun/runner.cpp | 10 +-- src/gromacs/mdtypes/forcerec.h | 4 +- src/gromacs/nbnxm/gpu_common.h | 2 +- src/gromacs/nbnxm/gpu_common_utils.h | 2 +- src/gromacs/nbnxm/nbnxm.cpp | 4 +- src/gromacs/nbnxm/nbnxm.h | 5 +- src/gromacs/nbnxm/nbnxm_gpu.h | 14 ++-- src/gromacs/nbnxm/nbnxm_gpu_data_mgmt.cpp | 6 +- src/gromacs/pbcutil/pbc_aiuc_cuda.cuh | 2 +- src/gromacs/taskassignment/decidegpuusage.cpp | 6 +- 21 files changed, 143 insertions(+), 138 deletions(-) rename src/gromacs/listed_forces/{gpubonded.h => listed_forces_gpu.h} (92%) rename src/gromacs/listed_forces/{gpubonded_impl.cpp => listed_forces_gpu_impl.cpp} (73%) rename src/gromacs/listed_forces/{gpubonded_impl.cu => listed_forces_gpu_impl.cu} (84%) rename src/gromacs/listed_forces/{gpubonded_impl.h => listed_forces_gpu_impl.h} (97%) rename src/gromacs/listed_forces/{gpubondedkernels.cu => listed_forces_gpu_internal.cu} (99%) diff --git a/src/gromacs/domdec/mdsetup.cpp b/src/gromacs/domdec/mdsetup.cpp index ead9a61fd1..47e4203f82 100644 --- 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 c1de86e14b..a169778e4c 100644 --- 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 b0e40a9474..893a1d306c 100644 --- 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 41ad20d42a..5cdd8c786f 100644 --- 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 @@ -95,7 +95,7 @@ constexpr std::array 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 1fe74e0082..f9874822ab 100644 --- a/src/gromacs/listed_forces/gpubonded_impl.cpp +++ b/src/gromacs/listed_forces/listed_forces_gpu_impl.cpp @@ -48,7 +48,7 @@ #include #include -#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 /* nbnxnAtomOrder */, - const InteractionDefinitions& /* idef */, - void* /* xqDevice */, - DeviceBuffer /* forceDevice */, - DeviceBuffer /* fshiftDevice */) +void ListedForcesGpu::updateInteractionListsAndDeviceBuffers(ArrayRef /* nbnxnAtomOrder */, + const InteractionDefinitions& /* idef */, + void* /* xqDevice */, + DeviceBuffer /* forceDevice */, + DeviceBuffer /* 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 90c5449da9..d02f535623 100644 --- 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 nbnxnAtomOrder, - const InteractionDefinitions& idef, - void* d_xqPtr, - DeviceBuffer d_fPtr, - DeviceBuffer d_fShiftPtr) +void ListedForcesGpu::Impl::updateInteractionListsAndDeviceBuffers(ArrayRef nbnxnAtomOrder, + const InteractionDefinitions& idef, + void* d_xqPtr, + DeviceBuffer d_fPtr, + DeviceBuffer d_fShiftPtr) { // TODO wallcycle sub start haveInteractions_ = false; @@ -293,19 +293,19 @@ void GpuBonded::Impl::updateInteractionListsAndDeviceBuffers(ArrayRef // 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 nbnxnAtomOrder, - const InteractionDefinitions& idef, - void* d_xq, - DeviceBuffer d_f, - DeviceBuffer d_fShift) +void ListedForcesGpu::updateInteractionListsAndDeviceBuffers(ArrayRef nbnxnAtomOrder, + const InteractionDefinitions& idef, + void* d_xq, + DeviceBuffer d_f, + DeviceBuffer 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 5e66c52eea..902ceb24c4 100644 --- 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 407e447bdc..9b35ea6077 100644 --- 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 @@ -909,7 +909,7 @@ __global__ void exec_kernel_gpu(BondedCudaKernelParameters kernelParams) template -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 0eeac68641..26f9ce89ee 100644 --- a/src/gromacs/listed_forces/manage_threading.cpp +++ b/src/gromacs/listed_forces/manage_threading.cpp @@ -57,7 +57,7 @@ #include #include -#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 61659e59e4..8e8caec27a 100644 --- 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 ea4d010d7a..e73d83fdac 100644 --- 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 ec9c382e64..c4987f1cfc 100644 --- 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; std::unique_ptr vsite; - std::unique_ptr gpuBonded; + std::unique_ptr 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( + listedForcesGpu = std::make_unique( 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 218c7b26c9..8a1ebe14e5 100644 --- 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; /* 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 bf9c25a5e7..13aea52f4b 100644 --- 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 0d81e2f05a..2ede311e38 100644 --- 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 361483b84c..e0aae3761f 100644 --- 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 dfaaff41dc..545f85ef4a 100644 --- 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 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 1cf920d35c..2a49837c8f 100644 --- 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 6e7a5b799f..874f9a614f 100644 --- 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 eb12862bf4..405321490e 100644 --- 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 c18bba3ec4..c407f61d5d 100644 --- 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) { -- 2.22.0