This is to folow general naming conventions across the code.
Refs #2885, #2888.
Change-Id: Ifa7e3febeff1d958155ed02daa97d26e828e8381
endif()
if(GMX_USE_CUDA)
gmx_add_libgromacs_sources(
- lincs_cuda.cu
+ lincs_gpu.cu
settle_cuda.cu
leapfrog_gpu.cu
update_constrain_cuda_impl.cu
* using CUDA, including class initialization, data-structures management
* and GPU kernel.
*
- * \todo Reconsider naming, i.e. "cuda" suffics should be changed to "gpu".
- *
* \author Artem Zhmurov <zhmurov@gmail.com>
* \author Alan Gray <alang@nvidia.com>
*
*/
#include "gmxpre.h"
-#include "lincs_cuda.cuh"
+#include "lincs_gpu.cuh"
#include <assert.h>
#include <stdio.h>
*/
template<bool updateVelocities, bool computeVirial>
__launch_bounds__(c_maxThreadsPerBlock) __global__
- void lincs_kernel(LincsCudaKernelParameters kernelParams,
+ void lincs_kernel(LincsGpuKernelParameters kernelParams,
const float3* __restrict__ gm_x,
float3* gm_xp,
float3* gm_v,
return kernelPtr;
}
-void LincsCuda::apply(const float3* d_x,
- float3* d_xp,
- const bool updateVelocities,
- float3* d_v,
- const real invdt,
- const bool computeVirial,
- tensor virialScaled,
- const PbcAiuc pbcAiuc)
+void LincsGpu::apply(const float3* d_x,
+ float3* d_xp,
+ const bool updateVelocities,
+ float3* d_v,
+ const real invdt,
+ const bool computeVirial,
+ tensor virialScaled,
+ const PbcAiuc pbcAiuc)
{
ensureNoPendingCudaError("In CUDA version of LINCS");
return;
}
-LincsCuda::LincsCuda(int numIterations, int expansionOrder, CommandStream commandStream) :
+LincsGpu::LincsGpu(int numIterations, int expansionOrder, CommandStream commandStream) :
commandStream_(commandStream)
{
kernelParams_.numIterations = numIterations;
numAtomsAlloc_ = 0;
}
-LincsCuda::~LincsCuda()
+LincsGpu::~LincsGpu()
{
freeDeviceBuffer(&kernelParams_.d_virialScaled);
return numCoupledConstraints;
}
-bool LincsCuda::isNumCoupledConstraintsSupported(const gmx_mtop_t& mtop)
+bool LincsGpu::isNumCoupledConstraintsSupported(const gmx_mtop_t& mtop)
{
for (const gmx_moltype_t& molType : mtop.moltype)
{
return true;
}
-void LincsCuda::set(const t_idef& idef, const t_mdatoms& md)
+void LincsGpu::set(const t_idef& idef, const t_mdatoms& md)
{
int numAtoms = md.nr;
// List of constrained atoms (CPU memory)
/*
* This file is part of the GROMACS molecular simulation package.
*
- * Copyright (c) 2019, by the GROMACS development team, led by
+ * Copyright (c) 2019,2020, 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.
*/
/*! \libinternal \file
*
- * \brief Declares the class for CUDA implementation of LINCS.
+ * \brief Declares the class for GPU implementation of LINCS.
*
* \author Artem Zhmurov <zhmurov@gmail.com>
*
* \ingroup module_mdlib
* \inlibraryapi
*/
-#ifndef GMX_MDLIB_LINCS_CUDA_CUH
-#define GMX_MDLIB_LINCS_CUDA_CUH
+#ifndef GMX_MDLIB_LINCS_GPU_CUH
+#define GMX_MDLIB_LINCS_GPU_CUH
#include "gromacs/gpu_utils/gputraits.cuh"
#include "gromacs/mdlib/constr.h"
* to the GPU as a single structure.
*
*/
-struct LincsCudaKernelParameters
+struct LincsGpuKernelParameters
{
//! Periodic boundary data
PbcAiuc pbcAiuc;
float* d_massFactors;
};
-/*! \internal \brief Class with interfaces and data for CUDA version of LINCS. */
-class LincsCuda
+/*! \internal \brief Class with interfaces and data for GPU version of LINCS. */
+class LincsGpu
{
public:
* \param[in] expansionOrder Order of the matrix inversion algorithm.
* \param[in] commandStream Device command stream.
*/
- LincsCuda(int numIterations, int expansionOrder, CommandStream commandStream);
+ LincsGpu(int numIterations, int expansionOrder, CommandStream commandStream);
/*! \brief Destructor.*/
- ~LincsCuda();
+ ~LincsGpu();
/*! \brief Apply LINCS.
*
/*! \brief
* Returns whether the maximum number of coupled constraints is supported
- * by the CUDA LINCS code.
+ * by the GPU LINCS code.
*
* \param[in] mtop The molecular topology
*/
static bool isNumCoupledConstraintsSupported(const gmx_mtop_t& mtop);
private:
- //! CUDA stream
+ //! GPU stream
CommandStream commandStream_;
- //! Parameters and pointers, passed to the CUDA kernel
- LincsCudaKernelParameters kernelParams_;
+ //! Parameters and pointers, passed to the GPU kernel
+ LincsGpuKernelParameters kernelParams_;
//! Scaled virial tensor (6 floats: [XX, XY, XZ, YY, YZ, ZZ])
std::vector<float> h_virialScaled_;
} // namespace gmx
-#endif // GMX_MDLIB_LINCS_CUDA_CUH
+#endif // GMX_MDLIB_LINCS_GPU_CUH
runnersNames.emplace_back("LINCS");
if (GMX_GPU == GMX_GPU_CUDA && canComputeOnGpu())
{
- runnersNames.emplace_back("LINCS_CUDA");
+ runnersNames.emplace_back("LINCS_GPU");
}
return runnersNames;
}
algorithms_["SHAKE"] = applyShake;
// LINCS
algorithms_["LINCS"] = applyLincs;
- // LINCS using CUDA (will only be called if CUDA is available)
- algorithms_["LINCS_CUDA"] = applyLincsCuda;
+ // LINCS using GPU (will only be called if GPU is available)
+ algorithms_["LINCS_GPU"] = applyLincsGpu;
}
/*! \brief
/*
* This file is part of the GROMACS molecular simulation package.
*
- * Copyright (c) 2018,2019, by the GROMACS development team, led by
+ * Copyright (c) 2018,2019,2020, 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.
#if GMX_GPU != GMX_GPU_CUDA
/*! \brief
- * Stub for LINCS constraints on CUDA-enabled GPU to satisfy compiler.
+ * Stub for GPU version of LINCS constraints to satisfy compiler.
*
* \param[in] testData Test data structure.
* \param[in] pbc Periodic boundary data.
*/
-void applyLincsCuda(ConstraintsTestData gmx_unused* testData, t_pbc gmx_unused pbc)
+void applyLincsGpu(ConstraintsTestData gmx_unused* testData, t_pbc gmx_unused pbc)
{
FAIL() << "Dummy LINCS CUDA function was called instead of the real one.";
}
/*
* This file is part of the GROMACS molecular simulation package.
*
- * Copyright (c) 2018,2019, by the GROMACS development team, led by
+ * Copyright (c) 2018,2019,2020, 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.
#include "gromacs/gpu_utils/devicebuffer.cuh"
#include "gromacs/gpu_utils/gpu_utils.h"
-#include "gromacs/mdlib/lincs_cuda.cuh"
+#include "gromacs/mdlib/lincs_gpu.cuh"
#include "gromacs/pbcutil/pbc.h"
#include "gromacs/utility/unique_cptr.h"
{
/*! \brief
- * Initialize and apply LINCS constraints on CUDA-enabled GPU.
+ * Initialize and apply LINCS constraints on GPU.
*
* \param[in] testData Test data structure.
* \param[in] pbc Periodic boundary data.
*/
-void applyLincsCuda(ConstraintsTestData* testData, t_pbc pbc)
+void applyLincsGpu(ConstraintsTestData* testData, t_pbc pbc)
{
- auto lincsCuda =
- std::make_unique<LincsCuda>(testData->ir_.nLincsIter, testData->ir_.nProjOrder, nullptr);
+ auto lincsGpu =
+ std::make_unique<LincsGpu>(testData->ir_.nLincsIter, testData->ir_.nProjOrder, nullptr);
bool updateVelocities = true;
int numAtoms = testData->numAtoms_;
float3 *d_x, *d_xp, *d_v;
- lincsCuda->set(testData->idef_, testData->md_);
+ lincsGpu->set(testData->idef_, testData->md_);
PbcAiuc pbcAiuc;
setPbcAiuc(pbc.ndim_ePBC, pbc.box, &pbcAiuc);
copyToDeviceBuffer(&d_v, (float3*)(testData->v_.data()), 0, numAtoms, nullptr,
GpuApiCallBehavior::Sync, nullptr);
}
- lincsCuda->apply(d_x, d_xp, updateVelocities, d_v, testData->invdt_, testData->computeVirial_,
- testData->virialScaled_, pbcAiuc);
+ lincsGpu->apply(d_x, d_xp, updateVelocities, d_v, testData->invdt_, testData->computeVirial_,
+ testData->virialScaled_, pbcAiuc);
copyFromDeviceBuffer((float3*)(testData->xPrime_.data()), &d_xp, 0, numAtoms, nullptr,
GpuApiCallBehavior::Sync, nullptr);
/*
* This file is part of the GROMACS molecular simulation package.
*
- * Copyright (c) 2018,2019, by the GROMACS development team, led by
+ * Copyright (c) 2018,2019,2020, 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.
/*! \brief Apply LINCS constraints to the test data.
*/
void applyLincs(ConstraintsTestData* testData, t_pbc pbc);
-/*! \brief Apply CUDA version of LINCS constraints to the test data.
+/*! \brief Apply GPU version of LINCS constraints to the test data.
*
* All the data is copied to the GPU device, then LINCS is applied and
* the resulting coordinates are copied back.
*/
-void applyLincsCuda(ConstraintsTestData* testData, t_pbc pbc);
+void applyLincsGpu(ConstraintsTestData* testData, t_pbc pbc);
} // namespace test
} // namespace gmx
#include "gromacs/gpu_utils/gputraits.cuh"
#include "gromacs/gpu_utils/vectype_ops.cuh"
#include "gromacs/mdlib/leapfrog_gpu.cuh"
-#include "gromacs/mdlib/lincs_cuda.cuh"
+#include "gromacs/mdlib/lincs_gpu.cuh"
#include "gromacs/mdlib/settle_cuda.cuh"
#include "gromacs/mdlib/update_constrain_cuda.h"
// Constraints need both coordinates before (d_x_) and after (d_xp_) update. However, after constraints
// are applied, the d_x_ can be discarded. So we intentionally swap the d_x_ and d_xp_ here to avoid the
// d_xp_ -> d_x_ copy after constraints. Note that the integrate saves them in the wrong order as well.
- lincsCuda_->apply(d_xp_, d_x_, updateVelocities, d_v_, 1.0 / dt, computeVirial, virial, pbcAiuc_);
+ lincsGpu_->apply(d_xp_, d_x_, updateVelocities, d_v_, 1.0 / dt, computeVirial, virial, pbcAiuc_);
settleCuda_->apply(d_xp_, d_x_, updateVelocities, d_v_, 1.0 / dt, computeVirial, virial, pbcAiuc_);
// scaledVirial -> virial (methods above returns scaled values)
integrator_ = std::make_unique<LeapFrogGpu>(commandStream_);
- lincsCuda_ = std::make_unique<LincsCuda>(ir.nLincsIter, ir.nProjOrder, commandStream_);
+ lincsGpu_ = std::make_unique<LincsGpu>(ir.nLincsIter, ir.nProjOrder, commandStream_);
settleCuda_ = std::make_unique<SettleCuda>(mtop, commandStream_);
coordinateScalingKernelLaunchConfig_.blockSize[0] = c_threadsPerBlock;
// Integrator should also update something, but it does not even have a method yet
integrator_->set(md, numTempScaleValues, md.cTC);
- lincsCuda_->set(idef, md);
+ lincsGpu_->set(idef, md);
settleCuda_->set(idef, md);
coordinateScalingKernelLaunchConfig_.gridSize[0] =
bool UpdateConstrainCuda::isNumCoupledConstraintsSupported(const gmx_mtop_t& mtop)
{
- return LincsCuda::isNumCoupledConstraintsSupported(mtop);
+ return LincsGpu::isNumCoupledConstraintsSupported(mtop);
}
} // namespace gmx
#include "gromacs/gpu_utils/gpueventsynchronizer.cuh"
#include "gromacs/mdlib/leapfrog_gpu.cuh"
-#include "gromacs/mdlib/lincs_cuda.cuh"
+#include "gromacs/mdlib/lincs_gpu.cuh"
#include "gromacs/mdlib/settle_cuda.cuh"
#include "gromacs/mdlib/update_constrain_cuda.h"
#include "gromacs/mdtypes/inputrec.h"
//! Leap-Frog integrator
std::unique_ptr<LeapFrogGpu> integrator_;
//! LINCS CUDA object to use for non-water constraints
- std::unique_ptr<LincsCuda> lincsCuda_;
+ std::unique_ptr<LincsGpu> lincsGpu_;
//! SETTLE CUDA object for water constrains
std::unique_ptr<SettleCuda> settleCuda_;
biod.pnpi.spb.ru / alexxy / gromacs.git / commitdiff