gmx_add_libgromacs_sources(
leapfrog_cuda_impl.cu
lincs_cuda.cu
- settle_cuda_impl.cu
+ settle_cuda.cu
update_constrain_cuda_impl.cu
)
endif()
*/
#include "gmxpre.h"
-#include "settle_cuda_impl.h"
+#include "settle_cuda.cuh"
#include <assert.h>
#include <stdio.h>
#include "gromacs/gpu_utils/gputraits.cuh"
#include "gromacs/gpu_utils/vectype_ops.cuh"
#include "gromacs/math/vec.h"
-#include "gromacs/mdlib/settle_cuda.h"
#include "gromacs/pbcutil/pbc.h"
#include "gromacs/pbcutil/pbc_aiuc_cuda.cuh"
__launch_bounds__(c_maxThreadsPerBlock)
__global__ void settle_kernel(const int numSettles,
const int3* __restrict__ gm_settles,
- const SettleCuda::SettleParameters pars,
+ const SettleParameters pars,
const float3* __restrict__ gm_x,
float3* __restrict__ gm_xprime,
const PbcAiuc pbcAiuc,
return kernelPtr;
}
-void SettleCuda::Impl::apply(const float3 *d_x,
- float3 *d_xp,
- const bool updateVelocities,
- float3 *d_v,
- const real invdt,
- const bool computeVirial,
- tensor virialScaled)
+void SettleCuda::apply(const float3 *d_x,
+ float3 *d_xp,
+ const bool updateVelocities,
+ float3 *d_v,
+ const real invdt,
+ const bool computeVirial,
+ tensor virialScaled)
{
ensureNoPendingCudaError("In CUDA version SETTLE");
return;
}
-void SettleCuda::Impl::copyApplyCopy(const int numAtoms,
- const rvec *h_x,
- rvec *h_xp,
- const bool updateVelocities,
- rvec *h_v,
- const real invdt,
- const bool computeVirial,
- tensor virialScaled)
-{
- float3 *d_x, *d_xp, *d_v;
-
- allocateDeviceBuffer(&d_x, numAtoms, nullptr);
- allocateDeviceBuffer(&d_xp, numAtoms, nullptr);
- allocateDeviceBuffer(&d_v, numAtoms, nullptr);
-
- copyToDeviceBuffer(&d_x, (float3*)h_x, 0, numAtoms, stream_, GpuApiCallBehavior::Sync, nullptr);
- copyToDeviceBuffer(&d_xp, (float3*)h_xp, 0, numAtoms, stream_, GpuApiCallBehavior::Sync, nullptr);
- if (updateVelocities)
- {
- copyToDeviceBuffer(&d_v, (float3*)h_v, 0, numAtoms, stream_, GpuApiCallBehavior::Sync, nullptr);
- }
- apply(d_x, d_xp,
- updateVelocities, d_v, invdt,
- computeVirial, virialScaled);
-
- copyFromDeviceBuffer((float3*)h_xp, &d_xp, 0, numAtoms, stream_, GpuApiCallBehavior::Sync, nullptr);
- if (updateVelocities)
- {
- copyFromDeviceBuffer((float3*)h_v, &d_v, 0, numAtoms, stream_, GpuApiCallBehavior::Sync, nullptr);
- }
-
- freeDeviceBuffer(&d_x);
- freeDeviceBuffer(&d_xp);
- freeDeviceBuffer(&d_v);
-}
-
-SettleCuda::Impl::Impl(const gmx_mtop_t &mtop)
+SettleCuda::SettleCuda(const gmx_mtop_t &mtop)
{
static_assert(sizeof(real) == sizeof(float),
"Real numbers should be in single precision in GPU code.");
}
-SettleCuda::Impl::Impl(const real mO, const real mH,
+SettleCuda::SettleCuda(const real mO, const real mH,
const real dOH, const real dHH)
{
static_assert(sizeof(real) == sizeof(float), "Real numbers should be in single precision in GPU code.");
}
-SettleCuda::Impl::~Impl()
+SettleCuda::~SettleCuda()
{
// Early exit if there is no settles
if (numSettles_ == 0)
}
}
-
-void SettleCuda::Impl::set(const t_idef &idef,
- const t_mdatoms gmx_unused &md)
+void SettleCuda::set(const t_idef &idef,
+ const t_mdatoms gmx_unused &md)
{
const int nral1 = 1 + NRAL(F_SETTLE);
t_ilist il_settle = idef.il[F_SETTLE];
}
-void SettleCuda::Impl::setPbc(const t_pbc *pbc)
-{
- setPbcAiuc(pbc->ndim_ePBC, pbc->box, &pbcAiuc_);
-}
-
-SettleCuda::SettleCuda(const gmx_mtop_t &mtop)
- : impl_(new Impl(mtop))
-{
-}
-
-SettleCuda::SettleCuda(const real mO, const real mH,
- const real dOH, const real dHH)
- : impl_(new Impl(mO, mH, dOH, dHH))
-{
-}
-
-SettleCuda::~SettleCuda() = default;
-
-void SettleCuda::copyApplyCopy(const int numAtoms,
- const rvec *h_x,
- rvec *h_xp,
- const bool updateVelocities,
- rvec *h_v,
- const real invdt,
- const bool computeVirial,
- tensor virialScaled)
-{
- impl_->copyApplyCopy(numAtoms, h_x, h_xp,
- updateVelocities, h_v, invdt,
- computeVirial, virialScaled);
-}
-
void SettleCuda::setPbc(const t_pbc *pbc)
{
- impl_->setPbc(pbc);
-}
-
-void SettleCuda::set(const t_idef &idef,
- const t_mdatoms &md)
-{
- impl_->set(idef, md);
+ setPbcAiuc(pbc->ndim_ePBC, pbc->box, &pbcAiuc_);
}
} // namespace gmx
*/
/*! \internal \file
*
- * \brief Declares CUDA implementation class for SETTLE
- *
- * This header file is needed to include from both the device-side
- * kernels file, and the host-side management code.
+ * \brief Declares class for CUDA implementation of SETTLE
*
* \author Artem Zhmurov <zhmurov@gmail.com>
*
* \ingroup module_mdlib
*/
-#ifndef GMX_MDLIB_SETTLE_CUDA_IMPL_H
-#define GMX_MDLIB_SETTLE_CUDA_IMPL_H
+#ifndef GMX_MDLIB_SETTLE_CUDA_CUH
+#define GMX_MDLIB_SETTLE_CUDA_CUH
+
+#include "gmxpre.h"
-#include "gromacs/mdlib/settle_cuda.h"
+#include "gromacs/math/functions.h"
+#include "gromacs/math/invertmatrix.h"
+#include "gromacs/math/vec.h"
#include "gromacs/mdtypes/mdatom.h"
-#include "gromacs/pbcutil/pbc_aiuc_cuda.cuh"
+#include "gromacs/pbcutil/pbc.h"
+#include "gromacs/pbcutil/pbc_aiuc.h"
#include "gromacs/topology/idef.h"
+#include "gromacs/topology/topology.h"
namespace gmx
{
* \todo Remove duplicates, check if recomputing makes more sense in some cases.
* \todo Move the projection parameters into separate structure.
*/
-struct SettleCuda::SettleParameters
+struct SettleParameters
{
//! Mass of oxygen atom
float mO;
* \param[in] dHH Target H-H bond length
*/
gmx_unused // Temporary solution to keep clang happy
-static void initSettleParameters(SettleCuda::SettleParameters *p,
+static void initSettleParameters(SettleParameters *p,
const real mO, const real mH,
const real dOH, const real dHH)
{
}
/*! \internal \brief Class with interfaces and data for CUDA version of SETTLE. */
-class SettleCuda::Impl
+class SettleCuda
{
public:
* target O-H and H-H distances. These values are also checked for
* consistency.
*/
- Impl(const gmx_mtop_t &mtop);
+ SettleCuda(const gmx_mtop_t &mtop);
/*! \brief Create SETTLE object
*
* \param[in] dOH Target distance for O-H bonds.
* \param[in] dHH Target for the distance between two hydrogen atoms.
*/
- Impl(const real mO, const real mH,
- const real dOH, const real dHH);
+ SettleCuda(const real mO, const real mH,
+ const real dOH, const real dHH);
- ~Impl();
+ ~SettleCuda();
/*! \brief Apply SETTLE.
*
const bool computeVirial,
tensor virialScaled);
- /*! \brief Apply SETTLE to the coordinates/velocities stored in CPU memory.
- *
- * This method should not be used in any code-path, where performance is of any value.
- * Only suitable for test and will be removed in future patch sets.
- * Allocates GPU memory, copies data from CPU, applies SETTLE to coordinates and,
- * if requested, to velocities, copies the results back, frees GPU memory.
- * Method uses this class data structures which should be filled with set() and setPbc()
- * methods.
- *
- * \todo Remove this method
- *
- * \param[in] numAtoms Number of atoms
- * \param[in] h_x Coordinates before timestep (in CPU memory)
- * \param[in,out] h_xp Coordinates after timestep (in CPU memory). The
- * resulting constrained coordinates will be saved here.
- * \param[in] updateVelocities If the velocities should be updated.
- * \param[in,out] h_v Velocities to update (in CPU memory, can be nullptr
- * if not updated)
- * \param[in] invdt Reciprocal timestep (to scale Lagrange
- * multipliers when velocities are updated)
- * \param[in] computeVirial If virial should be updated.
- * \param[in,out] virialScaled Scaled virial tensor to be updated.
- */
- void copyApplyCopy(const int numAtoms,
- const rvec *h_x,
- rvec *h_xp,
- const bool updateVelocities,
- rvec *h_v,
- const real invdt,
- const bool computeVirial,
- tensor virialScaled);
-
/*! \brief
* Update data-structures (e.g. after NB search step).
*
+++ /dev/null
-/*
- * This file is part of the GROMACS molecular simulation package.
- *
- * Copyright (c) 2019, 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.
- *
- * GROMACS is free software; you can redistribute it and/or
- * modify it under the terms of the GNU Lesser General Public License
- * as published by the Free Software Foundation; either version 2.1
- * of the License, or (at your option) any later version.
- *
- * GROMACS is distributed in the hope that it will be useful,
- * but WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
- * Lesser General Public License for more details.
- *
- * You should have received a copy of the GNU Lesser General Public
- * License along with GROMACS; if not, see
- * http://www.gnu.org/licenses, or write to the Free Software Foundation,
- * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
- *
- * If you want to redistribute modifications to GROMACS, please
- * consider that scientific software is very special. Version
- * control is crucial - bugs must be traceable. We will be happy to
- * consider code for inclusion in the official distribution, but
- * derived work must not be called official GROMACS. Details are found
- * in the README & COPYING files - if they are missing, get the
- * official version at http://www.gromacs.org.
- *
- * To help us fund GROMACS development, we humbly ask that you cite
- * the research papers on the package. Check out http://www.gromacs.org.
- */
-/*! \libinternal \file
- *
- * \brief Declaration of high-level functions of CUDA implementation of SETTLE.
- *
- * \todo This should only list interfaces needed for libgromacs clients.
- *
- * \author Artem Zhmurov <zhmurov@gmail.com>
- *
- * \ingroup module_mdlib
- * \inlibraryapi
- */
-#ifndef GMX_MDLIB_SETTLE_CUDA_H
-#define GMX_MDLIB_SETTLE_CUDA_H
-
-#include "gromacs/math/invertmatrix.h"
-#include "gromacs/mdtypes/mdatom.h"
-#include "gromacs/topology/idef.h"
-#include "gromacs/topology/topology.h"
-#include "gromacs/utility/classhelpers.h"
-
-namespace gmx
-{
-
-// TODO: Rename to SettleGpu
-class SettleCuda
-{
-
- public:
-
- /*! \brief Structure containing parameters for settles.
- *
- * Contains masses of atoms, distances between them and their pre-computed
- * derivatives (to avoid recomputing them for each water molecule).
- */
- struct SettleParameters;
-
- /*! \brief Create SETTLE object
- *
- * Extracts masses for oxygen and hydrogen as well as the O-H and H-H target distances
- * from the topology data (mtop), check their values for consistency and calls the
- * following constructor.
- *
- * \param[in] mtop Topology of the system to get the masses for O and H atoms and
- * target O-H and H-H distances. These values are also checked for
- * consistency.
- */
- SettleCuda(const gmx_mtop_t &mtop);
-
- /*! \brief Create SETTLE object
- *
- * \param[in] mO Mass of the oxygen atom.
- * \param[in] mH Mass of the hydrogen atom.
- * \param[in] dOH Target distance for O-H bonds.
- * \param[in] dHH Target for the distance between two hydrogen atoms.
- */
- SettleCuda(real mO, real mH,
- real dOH, real dHH);
-
- ~SettleCuda();
-
- /*! \brief Apply SETTLE to the coordinates/velocities stored in CPU memory.
- *
- * This method should not be used in any code-path, where performance is of any value.
- * Only suitable for test and will be removed in future patch sets.
- * Allocates GPU memory, copies data from CPU, applies SETTLE to coordinates and,
- * if requested, to velocities, copies the results back, frees GPU memory.
- * Method uses this class data structures which should be filled with set() and setPbc()
- * methods.
- *
- * \todo Remove this method
- *
- * \param[in] numAtoms Number of atoms
- * \param[in] h_x Coordinates before timestep (in CPU memory)
- * \param[in,out] h_xp Coordinates after timestep (in CPU memory). The
- * resulting constrained coordinates will be saved here.
- * \param[in] updateVelocities If the velocities should be updated.
- * \param[in,out] h_v Velocities to update (in CPU memory, can be nullptr
- * if not updated)
- * \param[in] invdt Reciprocal timestep (to scale Lagrange
- * multipliers when velocities are updated)
- * \param[in] computeVirial If virial should be updated.
- * \param[in,out] virialScaled Scaled virial tensor to be updated.
- */
- void copyApplyCopy(int numAtoms,
- const rvec *h_x,
- rvec *h_xp,
- bool updateVelocities,
- rvec *h_v,
- real invdt,
- bool computeVirial,
- tensor virialScaled);
-
- /*! \brief
- * Update data-structures (e.g. after NB search step).
- *
- * Updates the constraints data and copies it to the GPU. Should be
- * called if the particles were sorted, redistributed between domains, etc.
- * Does not recycle the data preparation routines from the CPU version.
- * All three atoms from single water molecule should be handled by the same GPU.
- *
- * SETTLEs atom ID's are taken from idef.il[F_SETTLE].iatoms.
- *
- * \param[in] idef System topology
- * \param[in] md Atoms data. Can be used to update masses if needed (not used now).
- */
- void set(const t_idef &idef,
- const t_mdatoms &md);
-
- /*! \brief
- * Update PBC data.
- *
- * \param[in] pbc The PBC data in t_pbc format.
- */
- void setPbc(const t_pbc *pbc);
-
- /*! \brief Class with hardware-specific interfaces and implementations.*/
- class Impl;
-
- private:
- gmx::PrivateImplPointer<Impl> impl_;
-
-
-
-};
-
-} // namespace gmx
-
-#endif
+++ /dev/null
-/*
- * This file is part of the GROMACS molecular simulation package.
- *
- * Copyright (c) 2019, 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.
- *
- * GROMACS is free software; you can redistribute it and/or
- * modify it under the terms of the GNU Lesser General Public License
- * as published by the Free Software Foundation; either version 2.1
- * of the License, or (at your option) any later version.
- *
- * GROMACS is distributed in the hope that it will be useful,
- * but WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
- * Lesser General Public License for more details.
- *
- * You should have received a copy of the GNU Lesser General Public
- * License along with GROMACS; if not, see
- * http://www.gnu.org/licenses, or write to the Free Software Foundation,
- * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
- *
- * If you want to redistribute modifications to GROMACS, please
- * consider that scientific software is very special. Version
- * control is crucial - bugs must be traceable. We will be happy to
- * consider code for inclusion in the official distribution, but
- * derived work must not be called official GROMACS. Details are found
- * in the README & COPYING files - if they are missing, get the
- * official version at http://www.gromacs.org.
- *
- * To help us fund GROMACS development, we humbly ask that you cite
- * the research papers on the package. Check out http://www.gromacs.org.
- */
-/*! \internal \file
- *
- * \brief Stub file to compile on systems without CUDA compiler.
- *
- * May be used to implement SETTLE CUDA interfaces for non-GPU builds.
- * Currently used to satisfy compiler on systems, where CUDA is not available.
- *
- * \author Artem Zhmurov <zhmurov@gmail.com>
- *
- * \ingroup module_mdlib
- */
-#include "gmxpre.h"
-
-#include "config.h"
-
-#include "gromacs/mdlib/settle_cuda.h"
-
-#if GMX_GPU != GMX_GPU_CUDA
-
-namespace gmx
-{
-
-class SettleCuda::Impl
-{
-};
-
-SettleCuda::SettleCuda(gmx_unused const gmx_mtop_t &mtop)
- : impl_(nullptr)
-{
- GMX_ASSERT(false, "A CPU stub for SETTLE was called insted of the correct implementation.");
-}
-
-SettleCuda::SettleCuda(gmx_unused const real mO, gmx_unused const real mH,
- gmx_unused const real dOH, gmx_unused const real dHH)
- : impl_(nullptr)
-{
- GMX_ASSERT(false, "A CPU stub for SETTLE was called insted of the correct implementation.");
-}
-
-SettleCuda::~SettleCuda() = default;
-
-void SettleCuda::copyApplyCopy(gmx_unused int numAtoms,
- gmx_unused const rvec *h_x,
- gmx_unused rvec *h_xp,
- gmx_unused bool updateVelocities,
- gmx_unused rvec *h_v,
- gmx_unused real invdt,
- gmx_unused bool computeVirial,
- gmx_unused tensor virialScaled)
-{
- GMX_ASSERT(false, "A CPU stub for SETTLE was called insted of the correct implementation.");
-}
-
-void SettleCuda::set(gmx_unused const t_idef &idef,
- gmx_unused const t_mdatoms &md)
-{
- GMX_ASSERT(false, "A CPU stub for SETTLE was called insted of the correct implementation.");
-}
-
-void SettleCuda::setPbc(gmx_unused const t_pbc *pbc)
-{
- GMX_ASSERT(false, "A CPU stub for SETTLE was called insted of the correct implementation.");
-}
-
-} // namespace gmx
-
-#endif /* GMX_GPU != GMX_GPU_CUDA */
if (GMX_USE_CUDA)
file(GLOB MDLIB_TEST_CUDA_SOURCES
constr_impl.cu
+ settle_runners.cu
)
endif()
#include "gromacs/math/paddedvector.h"
#include "gromacs/math/vec.h"
#include "gromacs/math/vectypes.h"
-#include "gromacs/mdlib/settle_cuda.h"
#include "gromacs/mdtypes/mdatom.h"
#include "gromacs/pbcutil/pbc.h"
#include "gromacs/topology/idef.h"
#include "gromacs/mdlib/tests/watersystem.h"
#include "testutils/testasserts.h"
+#include "settle_runners.h"
+
namespace gmx
{
{
public:
//! Updated water atom positions to constrain (DIM reals per atom)
- PaddedVector<gmx::RVec> updatedPositions_;
+ PaddedVector<RVec> updatedPositions_;
//! Water atom velocities to constrain (DIM reals per atom)
- PaddedVector<gmx::RVec> velocities_;
+ PaddedVector<RVec> velocities_;
//! No periodic boundary conditions
t_pbc pbcNone_;
//! Rectangular periodic box
const t_ilist ilist = { mtop.moltype[0].ilist[F_SETTLE].size(), 0, mtop.moltype[0].ilist[F_SETTLE].iatoms.data(), 0 };
// Copy the original positions from the array of doubles to a vector of reals
- PaddedVector<gmx::RVec> startingPositions;
+ PaddedVector<RVec> startingPositions;
startingPositions.resizeWithPadding(c_waterPositions.size());
std::copy(c_waterPositions.begin(), c_waterPositions.end(), startingPositions.begin());
#if GMX_GPU == GMX_GPU_CUDA
// Make a copy of all data-structures for GPU code testing
- PaddedVector<gmx::RVec> updatedPositionsGpu = updatedPositions_;
- PaddedVector<gmx::RVec> velocitiesGpu = velocities_;
- tensor virialGpu = {{0, 0, 0}, {0, 0, 0}, {0, 0, 0}};
-#endif
+ PaddedVector<RVec> updatedPositionsGpu = updatedPositions_;
+ PaddedVector<RVec> velocitiesGpu = velocities_;
+ tensor virialGpu = {{0, 0, 0}, {0, 0, 0}, {0, 0, 0}};
+#endif // GMX_GPU == GMX_GPU_CUDA
// Finally make the settle data structures
settledata *settled = settle_init(mtop);
// Verify the updated coordinates match the requirements
for (int i = 0; i < numSettles; ++i)
{
- const gmx::RVec &positionO = updatedPositions_[i*3 + 0];
- const gmx::RVec &positionH1 = updatedPositions_[i*3 + 1];
- const gmx::RVec &positionH2 = updatedPositions_[i*3 + 2];
+ const RVec &positionO = updatedPositions_[i*3 + 0];
+ const RVec &positionH1 = updatedPositions_[i*3 + 1];
+ const RVec &positionH2 = updatedPositions_[i*3 + 2];
EXPECT_REAL_EQ_TOL(dOH*dOH, distance2(positionO, positionH1), tolerance) << formatString("for water %d ", i) << testDescription;
EXPECT_REAL_EQ_TOL(dOH*dOH, distance2(positionO, positionH2), tolerance) << formatString("for water %d ", i) << testDescription;
if (canPerformGpuDetection())
{
// Run the CUDA code and check if it gives identical results to CPU code
- t_idef idef;
+ t_idef idef;
idef.il[F_SETTLE] = ilist;
- std::unique_ptr<SettleCuda> settleCuda = std::make_unique<SettleCuda>(mtop);
- settleCuda->setPbc(usePbc ? &pbcXyz_ : &pbcNone_);
- settleCuda->set(idef, mdatoms);
-
- settleCuda->copyApplyCopy(mdatoms.homenr,
- as_rvec_array(startingPositions.data()),
- as_rvec_array(updatedPositionsGpu.data()),
- useVelocities,
- as_rvec_array(velocitiesGpu.data()),
- reciprocalTimeStep,
- calcVirial,
- virialGpu);
+ applySettleCuda(mdatoms.homenr,
+ as_rvec_array(startingPositions.data()),
+ as_rvec_array(updatedPositionsGpu.data()),
+ useVelocities,
+ as_rvec_array(velocitiesGpu.data()),
+ reciprocalTimeStep,
+ calcVirial,
+ virialGpu,
+ usePbc ? &pbcXyz_ : &pbcNone_,
+ mtop,
+ idef,
+ mdatoms);
FloatingPointTolerance toleranceGpuCpu = absoluteTolerance(0.0001);
}
}
}
-#endif
+#endif // GMX_GPU == GMX_GPU_CUDA
}
// Scan the full Cartesian product of numbers of SETTLE interactions
--- /dev/null
+/*
+ * This file is part of the GROMACS molecular simulation package.
+ *
+ * Copyright (c) 2018,2019, 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.
+ *
+ * GROMACS is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public License
+ * as published by the Free Software Foundation; either version 2.1
+ * of the License, or (at your option) any later version.
+ *
+ * GROMACS is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with GROMACS; if not, see
+ * http://www.gnu.org/licenses, or write to the Free Software Foundation,
+ * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
+ *
+ * If you want to redistribute modifications to GROMACS, please
+ * consider that scientific software is very special. Version
+ * control is crucial - bugs must be traceable. We will be happy to
+ * consider code for inclusion in the official distribution, but
+ * derived work must not be called official GROMACS. Details are found
+ * in the README & COPYING files - if they are missing, get the
+ * official version at http://www.gromacs.org.
+ *
+ * To help us fund GROMACS development, we humbly ask that you cite
+ * the research papers on the package. Check out http://www.gromacs.org.
+ */
+/*! \internal \file
+ * \brief Implements the test runner for GPU version of SETTLE.
+ *
+ * \author Artem Zhmurov <zhmurov@gmail.com>
+ * \ingroup module_mdlib
+ */
+#include "gmxpre.h"
+
+#include "settle_runners.h"
+
+#include <assert.h>
+
+#include <cmath>
+
+#include <algorithm>
+#include <unordered_map>
+#include <vector>
+
+#include "gromacs/gpu_utils/devicebuffer.cuh"
+#include "gromacs/gpu_utils/gpu_utils.h"
+#include "gromacs/math/vec.h"
+#include "gromacs/math/vectypes.h"
+#include "gromacs/mdlib/settle_cuda.cuh"
+#include "gromacs/utility/unique_cptr.h"
+
+namespace gmx
+{
+namespace test
+{
+
+void applySettleCuda(const int numAtoms,
+ const rvec *h_x,
+ rvec *h_xp,
+ const bool updateVelocities,
+ rvec *h_v,
+ const real invdt,
+ const bool computeVirial,
+ tensor virialScaled,
+ const t_pbc *pbc,
+ const gmx_mtop_t &mtop,
+ const t_idef &idef,
+ const t_mdatoms &mdatoms)
+{
+ auto settleCuda = std::make_unique<SettleCuda>(mtop);
+ settleCuda->setPbc(pbc);
+ settleCuda->set(idef, mdatoms);
+
+ float3 *d_x, *d_xp, *d_v;
+
+ allocateDeviceBuffer(&d_x, numAtoms, nullptr);
+ allocateDeviceBuffer(&d_xp, numAtoms, nullptr);
+ allocateDeviceBuffer(&d_v, numAtoms, nullptr);
+
+ copyToDeviceBuffer(&d_x, (float3*)h_x, 0, numAtoms, nullptr, GpuApiCallBehavior::Sync, nullptr);
+ copyToDeviceBuffer(&d_xp, (float3*)h_xp, 0, numAtoms, nullptr, GpuApiCallBehavior::Sync, nullptr);
+ if (updateVelocities)
+ {
+ copyToDeviceBuffer(&d_v, (float3*)h_v, 0, numAtoms, nullptr, GpuApiCallBehavior::Sync, nullptr);
+ }
+ settleCuda->apply(d_x, d_xp,
+ updateVelocities, d_v, invdt,
+ computeVirial, virialScaled);
+
+ copyFromDeviceBuffer((float3*)h_xp, &d_xp, 0, numAtoms, nullptr, GpuApiCallBehavior::Sync, nullptr);
+ if (updateVelocities)
+ {
+ copyFromDeviceBuffer((float3*)h_v, &d_v, 0, numAtoms, nullptr, GpuApiCallBehavior::Sync, nullptr);
+ }
+
+ freeDeviceBuffer(&d_x);
+ freeDeviceBuffer(&d_xp);
+ freeDeviceBuffer(&d_v);
+}
+
+} // namespace test
+} // namespace gmx
--- /dev/null
+/*
+ * This file is part of the GROMACS molecular simulation package.
+ *
+ * Copyright (c) 2018,2019, 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.
+ *
+ * GROMACS is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public License
+ * as published by the Free Software Foundation; either version 2.1
+ * of the License, or (at your option) any later version.
+ *
+ * GROMACS is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with GROMACS; if not, see
+ * http://www.gnu.org/licenses, or write to the Free Software Foundation,
+ * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
+ *
+ * If you want to redistribute modifications to GROMACS, please
+ * consider that scientific software is very special. Version
+ * control is crucial - bugs must be traceable. We will be happy to
+ * consider code for inclusion in the official distribution, but
+ * derived work must not be called official GROMACS. Details are found
+ * in the README & COPYING files - if they are missing, get the
+ * official version at http://www.gromacs.org.
+ *
+ * To help us fund GROMACS development, we humbly ask that you cite
+ * the research papers on the package. Check out http://www.gromacs.org.
+ */
+/*! \internal \file
+ * \brief SETTLE tests header.
+ *
+ * Declares the functions that do the buffer management and apply
+ * SETTLE constraints ("test runners").
+ *
+ * \author Artem Zhmurov <zhmurov@gmail.com>
+ * \ingroup module_mdlib
+ */
+
+#ifndef GMX_MDLIB_TESTS_SETTLE_RUNNERS_H
+#define GMX_MDLIB_TESTS_SETTLE_RUNNERS_H
+
+#include "config.h"
+
+#include "gromacs/math/vectypes.h"
+
+struct t_pbc;
+struct gmx_mtop_t;
+struct t_idef;
+struct t_mdatoms;
+
+namespace gmx
+{
+namespace test
+{
+
+#if GMX_GPU == GMX_GPU_CUDA
+
+/*! \brief
+ * Initialize and apply SETTLE constraints on CUDA-enabled GPU.
+ *
+ * \param[in] numAtoms Number of atoms.
+ * \param[in] h_x Coordinates before timestep (in CPU memory).
+ * \param[in,out] h_xp Coordinates after timestep (in CPU memory). The
+ * resulting constrained coordinates will be saved here.
+ * \param[in] updateVelocities If the velocities should be updated.
+ * \param[in,out] h_v Velocities to update (in CPU memory, can be nullptr
+ * if not updated).
+ * \param[in] invdt Reciprocal timestep (to scale Lagrange
+ * multipliers when velocities are updated)
+ * \param[in] computeVirial If virial should be updated.
+ * \param[in,out] virialScaled Scaled virial tensor to be updated.
+ * \param[in] pbc Periodic boundary data.
+ * \param[in] mtop Topology of the system to get the masses for O and
+ * H atoms target O-H and H-H distances.
+ * \param[in] idef System topology.
+ * \param[in] mdatoms Atoms data.
+ */
+void applySettleCuda(int numAtoms,
+ const rvec *h_x,
+ rvec *h_xp,
+ bool updateVelocities,
+ rvec *h_v,
+ real invdt,
+ bool computeVirial,
+ tensor virialScaled,
+ const t_pbc *pbc,
+ const gmx_mtop_t &mtop,
+ const t_idef &idef,
+ const t_mdatoms &mdatoms);
+
+#endif // GMX_GPU == GMX_GPU_CUDA
+
+} // namespace test
+} // namespace gmx
+
+#endif // GMX_MDLIB_TESTS_SETTLE_RUNNERS_H
#include "gromacs/gpu_utils/vectype_ops.cuh"
#include "gromacs/math/vec.h"
#include "gromacs/mdlib/lincs_cuda.cuh"
+#include "gromacs/mdlib/settle_cuda.cuh"
#include "gromacs/mdlib/update_constrain_cuda.h"
#include "gromacs/pbcutil/pbc.h"
#include "gromacs/pbcutil/pbc_aiuc_cuda.cuh"
#include "leapfrog_cuda_impl.h"
-#include "settle_cuda_impl.h"
namespace gmx
{
GMX_RELEASE_ASSERT(numAtoms == mtop.natoms, "State and topology number of atoms should be the same.");
integrator_ = std::make_unique<LeapFrogCuda::Impl>();
lincsCuda_ = std::make_unique<LincsCuda>(ir.nLincsIter, ir.nProjOrder);
- settleCuda_ = std::make_unique<SettleCuda::Impl>(mtop);
+ settleCuda_ = std::make_unique<SettleCuda>(mtop);
}
#ifndef GMX_MDLIB_UPDATE_CONSTRAIN_CUDA_IMPL_H
#define GMX_MDLIB_UPDATE_CONSTRAIN_CUDA_IMPL_H
+#include "gmxpre.h"
+
#include "gromacs/mdlib/lincs_cuda.cuh"
+#include "gromacs/mdlib/settle_cuda.cuh"
#include "gromacs/mdlib/update_constrain_cuda.h"
#include "gromacs/mdtypes/inputrec.h"
#include "gromacs/pbcutil/pbc.h"
#include "gromacs/topology/idef.h"
#include "leapfrog_cuda_impl.h"
-#include "settle_cuda_impl.h"
namespace gmx
{
//! LINCS CUDA object to use for non-water constraints
std::unique_ptr<LincsCuda> lincsCuda_;
//! SETTLE CUDA object for water constrains
- std::unique_ptr<SettleCuda::Impl> settleCuda_;
+ std::unique_ptr<SettleCuda> settleCuda_;
};
biod.pnpi.spb.ru / alexxy / gromacs.git / commitdiff