* \todo Check if the force should be set to zero here.
* \todo This kernel can also accumulate incidental temperatures for each atom.
*
- * \tparam numTempScaleValues The number of different T-couple values.
- * \tparam velocityScaling Type of the Parrinello-Rahman velocity rescaling.
- * \param[in] numAtoms Total number of atoms.
- * \param[in,out] gm_x Coordinates to update upon integration.
- * \param[out] gm_xp A copy of the coordinates before the integration (for constraints).
- * \param[in,out] gm_v Velocities to update.
- * \param[in] gm_f Atomic forces.
- * \param[in] gm_inverseMasses Reciprocal masses.
- * \param[in] dt Timestep.
- * \param[in] gm_lambdas Temperature scaling factors (one per group)
- * \param[in] gm_tempScaleGroups Mapping of atoms into groups.
- * \param[in] dtPressureCouple Time step for pressure coupling
- * \param[in] velocityScalingMatrixDiagonal Diagonal elements of Parrinello-Rahman velocity scaling matrix
+ * \tparam numTempScaleValues The number of different T-couple values.
+ * \tparam velocityScaling Type of the Parrinello-Rahman velocity rescaling.
+ * \param[in] numAtoms Total number of atoms.
+ * \param[in,out] gm_x Coordinates to update upon integration.
+ * \param[out] gm_xp A copy of the coordinates before the integration (for constraints).
+ * \param[in,out] gm_v Velocities to update.
+ * \param[in] gm_f Atomic forces.
+ * \param[in] gm_inverseMasses Reciprocal masses.
+ * \param[in] dt Timestep.
+ * \param[in] gm_lambdas Temperature scaling factors (one per group)
+ * \param[in] gm_tempScaleGroups Mapping of atoms into groups.
+ * \param[in] dtPressureCouple Time step for pressure coupling
+ * \param[in] prVelocityScalingMatrixDiagonal Diagonal elements of Parrinello-Rahman velocity scaling matrix
*/
template<NumTempScaleValues numTempScaleValues, VelocityScalingType velocityScaling>
__launch_bounds__(c_maxThreadsPerBlock) __global__
const float dt,
const float* __restrict__ gm_lambdas,
const unsigned short* __restrict__ gm_tempScaleGroups,
- const float3 velocityScalingMatrixDiagonal);
+ const float3 prVelocityScalingMatrixDiagonal);
template<NumTempScaleValues numTempScaleValues, VelocityScalingType velocityScaling>
__launch_bounds__(c_maxThreadsPerBlock) __global__
const float dt,
const float* __restrict__ gm_lambdas,
const unsigned short* __restrict__ gm_tempScaleGroups,
- const float3 velocityScalingMatrixDiagonal)
+ const float3 prVelocityScalingMatrixDiagonal)
{
int threadIndex = blockIdx.x * blockDim.x + threadIdx.x;
if (threadIndex < numAtoms)
if (velocityScaling == VelocityScalingType::Diagonal)
{
- vp.x -= velocityScalingMatrixDiagonal.x * v.x;
- vp.y -= velocityScalingMatrixDiagonal.y * v.y;
- vp.z -= velocityScalingMatrixDiagonal.z * v.z;
+ vp.x -= prVelocityScalingMatrixDiagonal.x * v.x;
+ vp.y -= prVelocityScalingMatrixDiagonal.y * v.y;
+ vp.z -= prVelocityScalingMatrixDiagonal.z * v.z;
}
v = vp;
* Returns pointer to a CUDA kernel based on the number of temperature coupling groups.
* If zero is passed as an argument, it is assumed that no temperature coupling groups are used.
*
- * \param[in] numTempScaleValues Numer of temperature coupling groups in the system
- * \param[in] velocityScaling Type of the Parrinello-Rahman velocity scaling
+ * \param[in] numTempScaleValues Number of temperature coupling groups in the system
+ * \param[in] prVelocityScalingType Type of the Parrinello-Rahman velocity scaling
*
* \retrun Pointer to CUDA kernel
*/
-inline auto selectLeapFrogKernelPtr(int numTempScaleValues, VelocityScalingType velocityScaling)
+inline auto selectLeapFrogKernelPtr(int numTempScaleValues, VelocityScalingType prVelocityScalingType)
{
auto kernelPtr = leapfrog_kernel<NumTempScaleValues::None, VelocityScalingType::None>;
- if (velocityScaling == VelocityScalingType::None)
+ if (prVelocityScalingType == VelocityScalingType::None)
{
if (numTempScaleValues == 0)
{
"(zero for no coupling).");
}
}
- else if (velocityScaling == VelocityScalingType::Diagonal)
+ else if (prVelocityScalingType == VelocityScalingType::Diagonal)
{
if (numTempScaleValues == 0)
{
float3* d_v,
const float3* d_f,
const real dt,
- const bool doTempCouple,
+ const bool doTemperatureScaling,
gmx::ArrayRef<const t_grp_tcstat> tcstat,
- const bool doPressureCouple,
+ const bool doParrinelloRahman,
const float dtPressureCouple,
- const matrix velocityScalingMatrix)
+ const matrix prVelocityScalingMatrix)
{
ensureNoPendingCudaError("In CUDA version of Leap-Frog integrator");
auto kernelPtr = leapfrog_kernel<NumTempScaleValues::None, VelocityScalingType::None>;
- if (doTempCouple || doPressureCouple)
+ if (doTemperatureScaling || doParrinelloRahman)
{
- if (doTempCouple)
+ if (doTemperatureScaling)
{
GMX_ASSERT(numTempScaleValues_ == ssize(h_lambdas_),
"Number of temperature scaling factors changed since it was set for the "
copyToDeviceBuffer(&d_lambdas_, h_lambdas_.data(), 0, numTempScaleValues_,
commandStream_, GpuApiCallBehavior::Async, nullptr);
}
- VelocityScalingType velocityScaling = VelocityScalingType::None;
- if (doPressureCouple)
+ VelocityScalingType prVelocityScalingType = VelocityScalingType::None;
+ if (doParrinelloRahman)
{
- velocityScaling = VelocityScalingType::Diagonal;
- GMX_ASSERT(velocityScalingMatrix[YY][XX] == 0 && velocityScalingMatrix[ZZ][XX] == 0
- && velocityScalingMatrix[ZZ][YY] == 0 && velocityScalingMatrix[XX][YY] == 0
- && velocityScalingMatrix[XX][ZZ] == 0 && velocityScalingMatrix[YY][ZZ] == 0,
+ prVelocityScalingType = VelocityScalingType::Diagonal;
+ GMX_ASSERT(prVelocityScalingMatrix[YY][XX] == 0 && prVelocityScalingMatrix[ZZ][XX] == 0
+ && prVelocityScalingMatrix[ZZ][YY] == 0
+ && prVelocityScalingMatrix[XX][YY] == 0
+ && prVelocityScalingMatrix[XX][ZZ] == 0
+ && prVelocityScalingMatrix[YY][ZZ] == 0,
"Fully anisotropic Parrinello-Rahman pressure coupling is not yet supported "
"in GPU version of Leap-Frog integrator.");
- velocityScalingMatrixDiagonal_ =
- make_float3(dtPressureCouple * velocityScalingMatrix[XX][XX],
- dtPressureCouple * velocityScalingMatrix[YY][YY],
- dtPressureCouple * velocityScalingMatrix[ZZ][ZZ]);
+ prVelocityScalingMatrixDiagonal_ =
+ make_float3(dtPressureCouple * prVelocityScalingMatrix[XX][XX],
+ dtPressureCouple * prVelocityScalingMatrix[YY][YY],
+ dtPressureCouple * prVelocityScalingMatrix[ZZ][ZZ]);
}
- kernelPtr = selectLeapFrogKernelPtr(doTempCouple ? numTempScaleValues_ : 0, velocityScaling);
+ kernelPtr = selectLeapFrogKernelPtr(doTemperatureScaling ? numTempScaleValues_ : 0,
+ prVelocityScalingType);
}
const auto kernelArgs = prepareGpuKernelArguments(
kernelPtr, kernelLaunchConfig_, &numAtoms_, &d_x, &d_xp, &d_v, &d_f, &d_inverseMasses_,
- &dt, &d_lambdas_, &d_tempScaleGroups_, &velocityScalingMatrixDiagonal_);
+ &dt, &d_lambdas_, &d_tempScaleGroups_, &prVelocityScalingMatrixDiagonal_);
launchGpuKernel(kernelPtr, kernelLaunchConfig_, nullptr, "leapfrog_kernel", kernelArgs);
return;
* Integrates the equation of motion using Leap-Frog algorithm.
* Updates coordinates and velocities on the GPU. The current coordinates are saved for constraints.
*
- * \param[in,out] d_x Coordinates to update
- * \param[out] d_xp Place to save the values of initial coordinates coordinates to.
- * \param[in,out] d_v Velocities (will be updated).
- * \param[in] d_f Forces.
- * \param[in] dt Timestep.
- * \param[in] doTempCouple If the temperature coupling should be applied.
- * \param[in] tcstat Temperature coupling data.
- * \param[in] doPressureCouple If the temperature coupling should be applied.
- * \param[in] dtPressureCouple Period between pressure coupling steps
- * \param[in] velocityScalingMatrix Parrinello-Rahman velocity scaling matrix
+ * \param[in,out] d_x Coordinates to update
+ * \param[out] d_xp Place to save the values of initial coordinates coordinates to.
+ * \param[in,out] d_v Velocities (will be updated).
+ * \param[in] d_f Forces.
+ * \param[in] dt Timestep.
+ * \param[in] doTemperatureScaling If velocities should be scaled for temperature coupling.
+ * \param[in] tcstat Temperature coupling data.
+ * \param[in] doParrinelloRahman If current step is a Parrinello-Rahman pressure coupling step.
+ * \param[in] dtPressureCouple Period between pressure coupling steps
+ * \param[in] prVelocityScalingMatrix Parrinello-Rahman velocity scaling matrix
*/
void integrate(const float3* d_x,
float3* d_xp,
float3* d_v,
const float3* d_f,
const real dt,
- const bool doTempCouple,
+ const bool doTemperatureScaling,
gmx::ArrayRef<const t_grp_tcstat> tcstat,
- const bool doPressureCouple,
+ const bool doParrinelloRahman,
const float dtPressureCouple,
- const matrix velocityScalingMatrix);
+ const matrix prVelocityScalingMatrix);
/*! \brief Set the integrator
*
//! Maximum size of the temperature coupling groups array
int numTempScaleGroupsAlloc_ = -1;
- //! Vector with diagonal elements of the pressure coupling velocity rescale factors
- float3 velocityScalingMatrixDiagonal_;
+ //! Vector with diagonal elements of the Parrinello-Rahman pressure coupling velocity rescale factors
+ float3 prVelocityScalingMatrixDiagonal_;
};
} // namespace gmx
* This will extract temperature scaling factors from tcstat, transform them into the plain
* array and call the normal integrate method.
*
- * \param[in] fReadyOnDevice Event synchronizer indicating that the forces are
- * ready in the device memory.
- * \param[in] dt Timestep.
- * \param[in] updateVelocities If the velocities should be constrained.
- * \param[in] computeVirial If virial should be updated.
- * \param[out] virial Place to save virial tensor.
- * \param[in] doTempCouple If the temperature coupling should be performed.
- * \param[in] tcstat Temperature coupling data.
- * \param[in] doPressureCouple If the temperature coupling should be applied.
- * \param[in] dtPressureCouple Period between pressure coupling steps
- * \param[in] velocityScalingMatrix Parrinello-Rahman velocity scaling matrix
+ * \param[in] fReadyOnDevice Event synchronizer indicating that the forces are
+ * ready in the device memory.
+ * \param[in] dt Timestep.
+ * \param[in] updateVelocities If the velocities should be constrained.
+ * \param[in] computeVirial If virial should be updated.
+ * \param[out] virial Place to save virial tensor.
+ * \param[in] doTemperatureScaling If velocities should be scaled for temperature coupling.
+ * \param[in] tcstat Temperature coupling data.
+ * \param[in] doParrinelloRahman If current step is a Parrinello-Rahman pressure coupling step.
+ * \param[in] dtPressureCouple Period between pressure coupling steps.
+ * \param[in] prVelocityScalingMatrix Parrinello-Rahman velocity scaling matrix.
*/
void integrate(GpuEventSynchronizer* fReadyOnDevice,
real dt,
bool updateVelocities,
bool computeVirial,
tensor virial,
- bool doTempCouple,
+ bool doTemperatureScaling,
gmx::ArrayRef<const t_grp_tcstat> tcstat,
- bool doPressureCouple,
+ bool doParrinelloRahman,
float dtPressureCouple,
- const matrix velocityScalingMatrix);
+ const matrix prVelocityScalingMatrix);
/*! \brief Scale coordinates on the GPU for the pressure coupling.
*
{
};
-UpdateConstrainCuda::UpdateConstrainCuda(gmx_unused const t_inputrec& ir,
- gmx_unused const gmx_mtop_t& mtop,
- gmx_unused const void* commandStream,
- gmx_unused GpuEventSynchronizer* xUpdatedOnDevice) :
+UpdateConstrainCuda::UpdateConstrainCuda(const t_inputrec& /* ir */,
+ const gmx_mtop_t& /* mtop */,
+ const void* /* commandStream */,
+ GpuEventSynchronizer* /* xUpdatedOnDevice */) :
impl_(nullptr)
{
GMX_ASSERT(false,
UpdateConstrainCuda::~UpdateConstrainCuda() = default;
-void UpdateConstrainCuda::integrate(gmx_unused GpuEventSynchronizer* fReadyOnDevice,
- gmx_unused const real dt,
- gmx_unused const bool updateVelocities,
- gmx_unused const bool computeVirial,
- gmx_unused tensor virialScaled,
- gmx_unused const bool doTempCouple,
- gmx_unused gmx::ArrayRef<const t_grp_tcstat> tcstat,
- gmx_unused const bool doPressureCouple,
- gmx_unused const float dtPressureCouple,
- gmx_unused const matrix velocityScalingMatrix)
+void UpdateConstrainCuda::integrate(GpuEventSynchronizer* /* fReadyOnDevice */,
+ const real /* dt */,
+ const bool /* updateVelocities */,
+ const bool /* computeVirial */,
+ tensor /* virialScaled */,
+ const bool /* doTemperatureScaling */,
+ gmx::ArrayRef<const t_grp_tcstat> /* tcstat */,
+ const bool /* doParrinelloRahman */,
+ const float /* dtPressureCouple */,
+ const matrix /* prVelocityScalingMatrix*/)
{
GMX_ASSERT(false,
"A CPU stub for UpdateConstrain was called instead of the correct implementation.");
}
-void UpdateConstrainCuda::scaleCoordinates(gmx_unused const matrix scalingMatrix)
+void UpdateConstrainCuda::scaleCoordinates(const matrix /* scalingMatrix */)
{
GMX_ASSERT(false,
"A CPU stub for UpdateConstrain was called instead of the correct implementation.");
}
-void UpdateConstrainCuda::set(gmx_unused DeviceBuffer<float> d_x,
- gmx_unused DeviceBuffer<float> d_v,
- gmx_unused const DeviceBuffer<float> d_f,
- gmx_unused const t_idef& idef,
- gmx_unused const t_mdatoms& md,
- gmx_unused const int numTempScaleValues)
+void UpdateConstrainCuda::set(DeviceBuffer<float> /* d_x */,
+ DeviceBuffer<float> /* d_v */,
+ const DeviceBuffer<float> /* d_f */,
+ const t_idef& /* idef */,
+ const t_mdatoms& /* md */,
+ const int /* numTempScaleValues */)
{
GMX_ASSERT(false,
"A CPU stub for UpdateConstrain was called instead of the correct implementation.");
}
-void UpdateConstrainCuda::setPbc(gmx_unused const t_pbc* pbc)
+void UpdateConstrainCuda::setPbc(const t_pbc* /* pbc */)
{
GMX_ASSERT(false,
"A CPU stub for UpdateConstrain was called instead of the correct implementation.");
const bool updateVelocities,
const bool computeVirial,
tensor virial,
- const bool doTempCouple,
+ const bool doTemperatureScaling,
gmx::ArrayRef<const t_grp_tcstat> tcstat,
- const bool doPressureCouple,
+ const bool doParrinelloRahman,
const float dtPressureCouple,
- const matrix velocityScalingMatrix)
+ const matrix prVelocityScalingMatrix)
{
// Clearing virial matrix
// TODO There is no point in having separate virial matrix for constraints
// The integrate should save a copy of the current coordinates in d_xp_ and write updated once
// into d_x_. The d_xp_ is only needed by constraints.
- integrator_->integrate(d_x_, d_xp_, d_v_, d_f_, dt, doTempCouple, tcstat, doPressureCouple,
- dtPressureCouple, velocityScalingMatrix);
+ integrator_->integrate(d_x_, d_xp_, d_v_, d_f_, dt, doTemperatureScaling, tcstat,
+ doParrinelloRahman, dtPressureCouple, prVelocityScalingMatrix);
// 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.
const bool updateVelocities,
const bool computeVirial,
tensor virialScaled,
- const bool doTempCouple,
+ const bool doTemperatureScaling,
gmx::ArrayRef<const t_grp_tcstat> tcstat,
- const bool doPressureCouple,
+ const bool doParrinelloRahman,
const float dtPressureCouple,
- const matrix velocityScalingMatrix)
+ const matrix prVelocityScalingMatrix)
{
- impl_->integrate(fReadyOnDevice, dt, updateVelocities, computeVirial, virialScaled, doTempCouple,
- tcstat, doPressureCouple, dtPressureCouple, velocityScalingMatrix);
+ impl_->integrate(fReadyOnDevice, dt, updateVelocities, computeVirial, virialScaled, doTemperatureScaling,
+ tcstat, doParrinelloRahman, dtPressureCouple, prVelocityScalingMatrix);
}
void UpdateConstrainCuda::scaleCoordinates(const matrix scalingMatrix)
* 2. This is the temperature coupling step.
* Parameters virial/lambdas can be nullptr if computeVirial/doTempCouple are false.
*
- * \param[in] fReadyOnDevice Event synchronizer indicating that the forces are ready in
- * the device memory.
- * \param[in] dt Timestep.
- * \param[in] updateVelocities If the velocities should be constrained.
- * \param[in] computeVirial If virial should be updated.
- * \param[out] virial Place to save virial tensor.
- * \param[in] doTempCouple If the temperature coupling should be performed.
- * \param[in] tcstat Temperature coupling data.
- * \param[in] doPressureCouple If the temperature coupling should be applied.
- * \param[in] dtPressureCouple Period between pressure coupling steps
- * \param[in] velocityScalingMatrix Parrinello-Rahman velocity scaling matrix
+ * \param[in] fReadyOnDevice Event synchronizer indicating that the forces are ready in
+ * the device memory.
+ * \param[in] dt Timestep.
+ * \param[in] updateVelocities If the velocities should be constrained.
+ * \param[in] computeVirial If virial should be updated.
+ * \param[out] virial Place to save virial tensor.
+ * \param[in] doTemperatureScaling If velocities should be scaled for temperature coupling.
+ * \param[in] tcstat Temperature coupling data.
+ * \param[in] doParrinelloRahman If current step is a Parrinello-Rahman pressure coupling step.
+ * \param[in] dtPressureCouple Period between pressure coupling steps.
+ * \param[in] prVelocityScalingMatrix Parrinello-Rahman velocity scaling matrix.
*/
void integrate(GpuEventSynchronizer* fReadyOnDevice,
real dt,
bool updateVelocities,
bool computeVirial,
tensor virial,
- bool doTempCouple,
+ bool doTemperatureScaling,
gmx::ArrayRef<const t_grp_tcstat> tcstat,
- bool doPressureCouple,
+ bool doParrinelloRahman,
float dtPressureCouple,
- const matrix velocityScalingMatrix);
+ const matrix prVelocityScalingMatrix);
/*! \brief Scale coordinates on the GPU for the pressure coupling.
*
stateGpu->copyForcesToGpu(ArrayRef<RVec>(f), AtomLocality::Local);
}
- const bool doTempCouple =
+ const bool doTemperatureScaling =
(ir->etc != etcNO && do_per_step(step + ir->nsttcouple - 1, ir->nsttcouple));
// This applies Leap-Frog, LINCS and SETTLE in succession
integrator->integrate(stateGpu->getForcesReadyOnDeviceEvent(
AtomLocality::Local, runScheduleWork->stepWork.useGpuFBufferOps),
- ir->delta_t, true, bCalcVir, shake_vir, doTempCouple,
+ ir->delta_t, true, bCalcVir, shake_vir, doTemperatureScaling,
ekind->tcstat, doParrinelloRahman, ir->nstpcouple * ir->delta_t, M);
// Copy velocities D2H after update if:
biod.pnpi.spb.ru / alexxy / gromacs.git / commitdiff