Skip to content

GitLab

Commit 05259f77 authored by Peter Eastman's avatar Peter Eastman
Browse files

Cleaned up constraint code and deleted obsolete constraint implementations

parent f62ba9d8
Show whitespace changes
Inline Side-by-side
Showing with 221 additions and 922 deletions
platforms/cuda/src/CudaKernels.cpp
View file @ 05259f77
...@@ -399,7 +399,7 @@ static void initializeIntegration(const System& system, CudaPlatform::PlatformDa ...@@ -399,7 +399,7 @@ static void initializeIntegration(const System& system, CudaPlatform::PlatformDa
invMass1[i] = 1.0f/mass[particle1Index]; invMass1[i] = 1.0f/mass[particle1Index];
invMass2[i] = 1.0f/mass[particle2Index]; invMass2[i] = 1.0f/mass[particle2Index];
} }
gpuSetConstraintParameters(gpu, particle1, particle2, distance, invMass1, invMass2, (float)integrator.getConstraintTolerance(), 4); gpuSetConstraintParameters(gpu, particle1, particle2, distance, invMass1, invMass2, (float)integrator.getConstraintTolerance());
// Finish initialization. // Finish initialization.
...@@ -438,7 +438,7 @@ void CudaIntegrateVerletStepKernel::execute(OpenMMContextImpl& context, const Ve ...@@ -438,7 +438,7 @@ void CudaIntegrateVerletStepKernel::execute(OpenMMContextImpl& context, const Ve
kVerletUpdatePart1(gpu); kVerletUpdatePart1(gpu);
kApplyFirstShake(gpu); kApplyFirstShake(gpu);
kApplyFirstSettle(gpu); kApplyFirstSettle(gpu);
kApplyFirstCShake(gpu); kApplyFirstCCMA(gpu);
if (data.removeCM) { if (data.removeCM) {
int step = (int) (context.getTime()/stepSize); int step = (int) (context.getTime()/stepSize);
if (step%data.cmMotionFrequency == 0) if (step%data.cmMotionFrequency == 0)
...@@ -477,7 +477,7 @@ void CudaIntegrateLangevinStepKernel::execute(OpenMMContextImpl& context, const ...@@ -477,7 +477,7 @@ void CudaIntegrateLangevinStepKernel::execute(OpenMMContextImpl& context, const
kUpdatePart1(gpu); kUpdatePart1(gpu);
kApplyFirstShake(gpu); kApplyFirstShake(gpu);
kApplyFirstSettle(gpu); kApplyFirstSettle(gpu);
kApplyFirstCShake(gpu); kApplyFirstCCMA(gpu);
if (data.removeCM) { if (data.removeCM) {
int step = (int) (context.getTime()/stepSize); int step = (int) (context.getTime()/stepSize);
if (step%data.cmMotionFrequency == 0) if (step%data.cmMotionFrequency == 0)
...@@ -486,7 +486,7 @@ void CudaIntegrateLangevinStepKernel::execute(OpenMMContextImpl& context, const ...@@ -486,7 +486,7 @@ void CudaIntegrateLangevinStepKernel::execute(OpenMMContextImpl& context, const
kUpdatePart2(gpu); kUpdatePart2(gpu);
kApplySecondShake(gpu); kApplySecondShake(gpu);
kApplySecondSettle(gpu); kApplySecondSettle(gpu);
kApplySecondCShake(gpu); kApplySecondCCMA(gpu);
} }
CudaIntegrateBrownianStepKernel::~CudaIntegrateBrownianStepKernel() { CudaIntegrateBrownianStepKernel::~CudaIntegrateBrownianStepKernel() {
...@@ -519,7 +519,7 @@ void CudaIntegrateBrownianStepKernel::execute(OpenMMContextImpl& context, const ...@@ -519,7 +519,7 @@ void CudaIntegrateBrownianStepKernel::execute(OpenMMContextImpl& context, const
kBrownianUpdatePart1(gpu); kBrownianUpdatePart1(gpu);
kApplyFirstShake(gpu); kApplyFirstShake(gpu);
kApplyFirstSettle(gpu); kApplyFirstSettle(gpu);
kApplyFirstCShake(gpu); kApplyFirstCCMA(gpu);
if (data.removeCM) { if (data.removeCM) {
int step = (int) (context.getTime()/stepSize); int step = (int) (context.getTime()/stepSize);
if (step%data.cmMotionFrequency == 0) if (step%data.cmMotionFrequency == 0)
......
platforms/cuda/src/kernels/cudaKernels.h
View file @ 05259f77
...@@ -42,14 +42,12 @@ extern void kCalculateAndersenThermostat(gpuContext gpu); ...@@ -42,14 +42,12 @@ extern void kCalculateAndersenThermostat(gpuContext gpu);
extern void kReduceBornSumAndForces(gpuContext gpu); extern void kReduceBornSumAndForces(gpuContext gpu);
extern void kUpdatePart1(gpuContext gpu); extern void kUpdatePart1(gpuContext gpu);
extern void kApplyFirstShake(gpuContext gpu); extern void kApplyFirstShake(gpuContext gpu);
extern void kApplyFirstCShake(gpuContext gpu); extern void kApplyFirstCCMA(gpuContext gpu);
extern void kApplyFirstSettle(gpuContext gpu); extern void kApplyFirstSettle(gpuContext gpu);
extern void kApplyFirstLincs(gpuContext gpu);
extern void kUpdatePart2(gpuContext gpu); extern void kUpdatePart2(gpuContext gpu);
extern void kApplySecondShake(gpuContext gpu); extern void kApplySecondShake(gpuContext gpu);
extern void kApplySecondCShake(gpuContext gpu); extern void kApplySecondCCMA(gpuContext gpu);
extern void kApplySecondSettle(gpuContext gpu); extern void kApplySecondSettle(gpuContext gpu);
extern void kApplySecondLincs(gpuContext gpu);
extern void kVerletUpdatePart1(gpuContext gpu); extern void kVerletUpdatePart1(gpuContext gpu);
extern void kVerletUpdatePart2(gpuContext gpu); extern void kVerletUpdatePart2(gpuContext gpu);
extern void kBrownianUpdatePart1(gpuContext gpu); extern void kBrownianUpdatePart1(gpuContext gpu);
...@@ -78,10 +76,8 @@ extern void SetUpdateShakeHSim(gpuContext gpu); ...@@ -78,10 +76,8 @@ extern void SetUpdateShakeHSim(gpuContext gpu);
extern void GetUpdateShakeHSim(gpuContext gpu); extern void GetUpdateShakeHSim(gpuContext gpu);
extern void SetSettleSim(gpuContext gpu); extern void SetSettleSim(gpuContext gpu);
extern void GetSettleSim(gpuContext gpu); extern void GetSettleSim(gpuContext gpu);
extern void SetCShakeSim(gpuContext gpu); extern void SetCCMASim(gpuContext gpu);
extern void GetCShakeSim(gpuContext gpu); extern void GetCCMASim(gpuContext gpu);
extern void SetLincsSim(gpuContext gpu);
extern void GetLincsSim(gpuContext gpu);
extern void SetVerletUpdateSim(gpuContext gpu); extern void SetVerletUpdateSim(gpuContext gpu);
extern void GetVerletUpdateSim(gpuContext gpu); extern void GetVerletUpdateSim(gpuContext gpu);
extern void SetBrownianUpdateSim(gpuContext gpu); extern void SetBrownianUpdateSim(gpuContext gpu);
......
platforms/cuda/src/kernels/cudatypes.h
View file @ 05259f77
...@@ -258,7 +258,7 @@ struct cudaGmxSimulation { ...@@ -258,7 +258,7 @@ struct cudaGmxSimulation {
unsigned int max_shake_threads_per_block; // Maximum threads per block in shake kernel calls unsigned int max_shake_threads_per_block; // Maximum threads per block in shake kernel calls
unsigned int shake_threads_per_block; // Threads per block in shake kernel calls unsigned int shake_threads_per_block; // Threads per block in shake kernel calls
unsigned int settle_threads_per_block; // Threads per block in SETTLE kernel calls unsigned int settle_threads_per_block; // Threads per block in SETTLE kernel calls
unsigned int lincs_threads_per_block; // Threads per block in LINCS kernel calls unsigned int ccma_threads_per_block; // Threads per block in CCMA kernel calls
unsigned int nonshake_threads_per_block; // Threads per block in nonshaking kernel call unsigned int nonshake_threads_per_block; // Threads per block in nonshaking kernel call
unsigned int max_localForces_threads_per_block; // Threads per block in local forces kernel calls unsigned int max_localForces_threads_per_block; // Threads per block in local forces kernel calls
unsigned int localForces_threads_per_block; // Threads per block in local forces kernel calls unsigned int localForces_threads_per_block; // Threads per block in local forces kernel calls
...@@ -360,7 +360,7 @@ struct cudaGmxSimulation { ...@@ -360,7 +360,7 @@ struct cudaGmxSimulation {
float inverseTotalMass; // Used in linear momentum removal float inverseTotalMass; // Used in linear momentum removal
unsigned int ShakeConstraints; // Total number of Shake constraints unsigned int ShakeConstraints; // Total number of Shake constraints
unsigned int settleConstraints; // Total number of Settle constraints unsigned int settleConstraints; // Total number of Settle constraints
unsigned int lincsConstraints; // Total number of LINCS constraints. unsigned int ccmaConstraints; // Total number of CCMA constraints.
unsigned int rigidClusters; // Total number of rigid clusters unsigned int rigidClusters; // Total number of rigid clusters
unsigned int maxRigidClusterSize; // The size of the largest rigid cluster unsigned int maxRigidClusterSize; // The size of the largest rigid cluster
unsigned int clusterShakeBlockSize; // The number of threads to process each rigid cluster unsigned int clusterShakeBlockSize; // The number of threads to process each rigid cluster
...@@ -368,7 +368,6 @@ struct cudaGmxSimulation { ...@@ -368,7 +368,6 @@ struct cudaGmxSimulation {
unsigned int maxShakeIterations; // Maximum shake iterations unsigned int maxShakeIterations; // Maximum shake iterations
unsigned int degreesOfFreedom; // Number of degrees of freedom in system unsigned int degreesOfFreedom; // Number of degrees of freedom in system
float shakeTolerance; // Shake tolerance float shakeTolerance; // Shake tolerance
unsigned int lincsTerms; // Number of terms in the matrix expansion for LINCS
float InvMassJ; // Shake inverse mass for hydrogens float InvMassJ; // Shake inverse mass for hydrogens
int* pNonShakeID; // Not Shaking atoms int* pNonShakeID; // Not Shaking atoms
int4* pShakeID; // Shake atoms and phase int4* pShakeID; // Shake atoms and phase
...@@ -385,23 +384,15 @@ struct cudaGmxSimulation { ...@@ -385,23 +384,15 @@ struct cudaGmxSimulation {
int* pAtomIndex; // The original index of each atom int* pAtomIndex; // The original index of each atom
float4* pGridBoundingBox; // The size of each grid cell float4* pGridBoundingBox; // The size of each grid cell
float4* pGridCenter; // The center of each grid cell float4* pGridCenter; // The center of each grid cell
int2* pLincsAtoms; // The atoms connected by each LINCS constraint int2* pCcmaAtoms; // The atoms connected by each CCMA constraint
float4* pLincsDistance; // The displacement vector (x, y, z) and constraint distance (w) for each LINCS constraint float4* pCcmaDistance; // The displacement vector (x, y, z) and constraint distance (w) for each CCMA constraint
int* pLincsConnections; // The indices of constraints that other constraints are connected to float* pCcmaDelta1; // Workspace for CCMA
int* pLincsNumConnections; // The number of other constraints that each constraint is linked to float* pCcmaDelta2; // Workspace for CCMA
float* pLincsS; // S matrix for LINCS int* pCcmaAtomConstraints; // The indices of constraints involving each atom
float* pLincsCoupling; // Coupling matrix for LINCS int* pCcmaNumAtomConstraints; // The number of constraints involving each atom
float* pLincsRhs1; // Workspace for LINCS
float* pLincsRhs2; // Workspace for LINCS
float* pLincsSolution; // Workspace for LINCS
int* pLincsAtomConstraints; // The indices of constraints involving each atom
int* pLincsNumAtomConstraints; // The number of constraints involving each atom
short* pSyncCounter; // Used for global thread synchronization short* pSyncCounter; // Used for global thread synchronization
unsigned int* pRequiredIterations; // Used by SHAKE to communicate whether iteration has converged unsigned int* pRequiredIterations; // Used by CCMA to communicate whether iteration has converged
float* pShakeReducedMass; // The reduced mass for each SHAKE constraint float* pCcmaReducedMass; // The reduced mass for each CCMA constraint
float* pRigidClusterMatrix; // The inverse constraint matrix for each rigid cluster
unsigned int* pRigidClusterConstraintIndex; // The index of each cluster in the stream containing cluster constraints.
unsigned int* pRigidClusterMatrixIndex; // The index of each cluster in the stream containing cluster matrices.
unsigned int* pConstraintMatrixColumn; // The column of each element in the constraint matrix. unsigned int* pConstraintMatrixColumn; // The column of each element in the constraint matrix.
float* pConstraintMatrixValue; // The value of each element in the constraint matrix. float* pConstraintMatrixValue; // The value of each element in the constraint matrix.
......
platforms/cuda/src/kernels/gpu.cpp
View file @ 05259f77
...@@ -48,12 +48,9 @@ using namespace std; ...@@ -48,12 +48,9 @@ using namespace std;
#include "cudaKernels.h" #include "cudaKernels.h"
#include "hilbert.h" #include "hilbert.h"
#include "openmm/OpenMMException.h" #include "openmm/OpenMMException.h"
#include "jama_svd.h"
#include "quern.h" #include "quern.h"
using OpenMM::OpenMMException; using OpenMM::OpenMMException;
using TNT::Array2D;
using JAMA::SVD;
struct ShakeCluster { struct ShakeCluster {
int centralID; int centralID;
...@@ -500,221 +497,9 @@ static void markShakeClusterInvalid(ShakeCluster& cluster, map<int, ShakeCluster ...@@ -500,221 +497,9 @@ static void markShakeClusterInvalid(ShakeCluster& cluster, map<int, ShakeCluster
} }
} }
static void findRigidClusters(gpuContext gpu, const vector<int>& firstAtom, const vector<int>& secondAtom, const vector<float>& invMass1, const vector<float>& invMass2, const vector<float>& distance, vector<int>& constraintIndices)
{
vector<map<int, int> > atomConstraints(firstAtom.size());
for (int i = 0; i < (int)constraintIndices.size(); i++) {
atomConstraints[firstAtom[i]][secondAtom[i]] = constraintIndices[i];
atomConstraints[secondAtom[i]][firstAtom[i]] = constraintIndices[i];
}
vector<vector<int> > rigidClusters;
int totalConstraints = 0;
int totalMatrixElements = 0;
for (int i = 0; i < (int)firstAtom.size(); i++) {
if (atomConstraints[i].size() < 2)
continue;
// Begin by looking for a triangle this atom is part of.
set<int> atoms;
atoms.insert(i);
for (map<int, int>::const_iterator atom1 = atomConstraints[i].begin(); atom1 != atomConstraints[i].end() && atoms.size() == 1; ++atom1) {
for (map<int, int>::const_iterator atom2 = atomConstraints[atom1->first].begin(); atom2 != atomConstraints[atom1->first].end(); ++atom2) {
if (atomConstraints[i].count(atom2->first) != 0) {
atoms.insert(atom1->first);
atoms.insert(atom2->first);
break;
}
}
}
if (atoms.size() == 1)
continue;
// We have three atoms that are part of a cluster, so look for other atoms we can add.
bool done = false;
while (!done) {
done = true;
for (set<int>::const_iterator atom1 = atoms.begin(); atom1 != atoms.end(); ++atom1) {
for (map<int, int>::const_iterator atom2 = atomConstraints[*atom1].begin(); atom2 != atomConstraints[*atom1].end(); ++atom2) {
if (atoms.find(atom2->first) != atoms.end())
continue; // This atom is already in the cluster.
// See if this atom is linked to three other atoms in the cluster.
int linkCount = 0;
for (map<int, int>::const_iterator atom3 = atomConstraints[atom2->first].begin(); atom3 != atomConstraints[atom2->first].end(); ++atom3)
if (atoms.find(atom3->first) != atoms.end())
linkCount++;
if (linkCount > 2) {
atoms.insert(atom2->first);
done = false;
}
}
}
}
// Record the cluster.
vector<int> constraints;
for (set<int>::const_iterator atom1 = atoms.begin(); atom1 != atoms.end(); ++atom1) {
for (map<int, int>::const_iterator atom2 = atomConstraints[*atom1].begin(); atom2 != atomConstraints[*atom1].end(); ++atom2) {
if (*atom1 < atom2->first && atoms.find(atom2->first) != atoms.end())
constraints.push_back(atom2->second);
}
}
rigidClusters.push_back(constraints);
totalConstraints += constraints.size();
totalMatrixElements += constraints.size()*constraints.size();
for (set<int>::const_iterator atom1 = atoms.begin(); atom1 != atoms.end(); ++atom1) {
for (map<int, int>::const_iterator atom2 = atomConstraints[*atom1].begin(); atom2 != atomConstraints[*atom1].end(); ++atom2)
atomConstraints[atom2->first].erase(*atom1);
atomConstraints[*atom1].clear();
}
}
// Reorder the constraints so those in a cluster are sequential.
// vector<int> constraintOrder(constraintIndices.size());
// vector<int> clusterStartIndex(rigidClusters.size());
// set<int> inCluster;
// int nextIndex = 0;
// for (int i = 0; i < (int) rigidClusters.size(); ++i) {
// clusterStartIndex[i] = nextIndex;
// for (int j = 0; j < (int) rigidClusters[i].size(); ++j) {
// int constraint = rigidClusters[i][j];
// constraintOrder[nextIndex++] = constraint;
// inCluster.insert(constraint);
// }
// }
// for (int i = 0; i < (int) constraintIndices.size(); ++i)
// if (inCluster.find(constraintIndices[i]) == inCluster.end())
// constraintOrder[nextIndex++] = constraintIndices[i];
// constraintIndices = constraintOrder;
// Build the CUDA streams.
CUDAStream<unsigned int>* psRigidClusterConstraintIndex = new CUDAStream<unsigned int>((int) rigidClusters.size()+1, 1, "RigidClusterConstraintIndex");
gpu->psRigidClusterConstraintIndex = psRigidClusterConstraintIndex;
gpu->sim.pRigidClusterConstraintIndex = psRigidClusterConstraintIndex->_pDevData;
CUDAStream<float>* psRigidClusterMatrix = new CUDAStream<float>(totalMatrixElements, 1, "RigidClusterMatrix");
gpu->psRigidClusterMatrix = psRigidClusterMatrix;
gpu->sim.pRigidClusterMatrix = psRigidClusterMatrix->_pDevData;
CUDAStream<unsigned int>* psRigidClusterMatrixIndex = new CUDAStream<unsigned int>((int) rigidClusters.size()+1, 1, "RigidClusterMatrixIndex");
gpu->psRigidClusterMatrixIndex = psRigidClusterMatrixIndex;
gpu->sim.pRigidClusterMatrixIndex = psRigidClusterMatrixIndex->_pDevData;
unsigned int constraintIndex = 0;
unsigned int maxClusterSize = 0;
for (unsigned int i = 0; i < rigidClusters.size(); i++) {
vector<int>& cluster = rigidClusters[i];
(*psRigidClusterConstraintIndex)[i] = constraintIndex;
constraintIndex += cluster.size();
if (cluster.size() > maxClusterSize)
maxClusterSize = cluster.size();
}
(*psRigidClusterConstraintIndex)[rigidClusters.size()] = constraintIndex;
gpu->sim.rigidClusters = rigidClusters.size();
gpu->sim.maxRigidClusterSize = maxClusterSize;
gpu->sim.clusterShakeBlockSize = 1;
while (gpu->sim.clusterShakeBlockSize < maxClusterSize)
gpu->sim.clusterShakeBlockSize *= 2;
psRigidClusterConstraintIndex->Upload();
// Build the inverse coupling matrix for each cluster.
unsigned int elementIndex = 0;
for (unsigned int i = 0; i < rigidClusters.size(); i++) {
// Compute the constraint coupling matrix for this cluster.
const vector<int>& cluster = rigidClusters[i];
unsigned int size = cluster.size();
Array2D<double> matrix(size, size);
for (int j = 0; j < (int)size; j++) {
for (int k = 0; k < (int)size; k++) {
if (j == k) {
matrix[j][j] = 1.0;
continue;
}
double scale;
int atomj0 = firstAtom[cluster[j]];
int atomj1 = secondAtom[cluster[j]];
int atomk0 = firstAtom[cluster[k]];
int atomk1 = secondAtom[cluster[k]];
int atoma, atomb;
if (atomj0 == atomk0) {
atoma = atomj1;
atomb = atomk1;
scale = invMass1[cluster[j]]/(invMass1[cluster[j]]+invMass2[cluster[j]]);
}
else if (atomj1 == atomk1) {
atoma = atomj0;
atomb = atomk0;
scale = invMass2[cluster[j]]/(invMass1[cluster[j]]+invMass2[cluster[j]]);
}
else if (atomj0 == atomk1) {
atoma = atomj1;
atomb = atomk0;
scale = invMass1[cluster[j]]/(invMass1[cluster[j]]+invMass2[cluster[j]]);
}
else if (atomj1 == atomk0) {
atoma = atomj0;
atomb = atomk1;
scale = invMass2[cluster[j]]/(invMass1[cluster[j]]+invMass2[cluster[j]]);
}
else {
matrix[j][k] = 0.0;
continue; // These constraints are not connected.
}
// Find the third constraint forming a triangle with these two.
for (int m = 0; m < size; m++) {
int other = cluster[m];
if ((firstAtom[other] == atoma && secondAtom[other] == atomb) || (firstAtom[other] == atomb && secondAtom[other] == atoma)) {
double d1 = distance[cluster[j]];
double d2 = distance[cluster[k]];
double d3 = distance[other];
matrix[j][k] = scale*(d1*d1+d2*d2-d3*d3)/(2.0*d1*d2);
break;
}
}
}
}
// Invert it using SVD.
Array2D<double> u, v;
Array1D<double> w;
SVD<double> svd(matrix);
svd.getU(u);
svd.getV(v);
svd.getSingularValues(w);
double singularValueCutoff = 0.01*w[0];
for (int j = 0; j < (int)size; j++)
w[j] = (w[j] < singularValueCutoff ? 0.0 : 1.0/w[j]);
for (int j = 0; j < (int)size; j++) {
for (int k = 0; k < (int)size; k++) {
matrix[j][k] = 0.0;
for (int m = 0; m < (int)size; m++)
matrix[j][k] += v[j][m]*w[m]*u[k][m];
}
}
// Record the inverted matrix.
(*gpu->psRigidClusterMatrixIndex)[i] = elementIndex;
for (int j = 0; j < (int)size; j++)
for (int k = 0; k < (int)size; k++)
(*gpu->psRigidClusterMatrix)[elementIndex++] = (float)(matrix[k][j]*distance[cluster[j]]/distance[cluster[k]]);
}
(*gpu->psRigidClusterMatrixIndex)[gpu->sim.rigidClusters] = elementIndex;
gpu->psRigidClusterMatrix->Upload();
gpu->psRigidClusterMatrixIndex->Upload();
}
extern "C" extern "C"
void gpuSetConstraintParameters(gpuContext gpu, const vector<int>& atom1, const vector<int>& atom2, const vector<float>& distance, void gpuSetConstraintParameters(gpuContext gpu, const vector<int>& atom1, const vector<int>& atom2, const vector<float>& distance,
const vector<float>& invMass1, const vector<float>& invMass2, float shakeTolerance, unsigned int lincsTerms) const vector<float>& invMass1, const vector<float>& invMass2, float constraintTolerance)
{ {
// Create a vector for recording which atoms are handled by SHAKE (or SETTLE). // Create a vector for recording which atoms are handled by SHAKE (or SETTLE).
...@@ -901,33 +686,29 @@ void gpuSetConstraintParameters(gpuContext gpu, const vector<int>& atom1, const ...@@ -901,33 +686,29 @@ void gpuSetConstraintParameters(gpuContext gpu, const vector<int>& atom1, const
} }
psShakeID->Upload(); psShakeID->Upload();
psShakeParameter->Upload(); psShakeParameter->Upload();
gpu->sim.shakeTolerance = shakeTolerance; gpu->sim.shakeTolerance = constraintTolerance;
gpu->sim.shake_threads_per_block = (gpu->sim.ShakeConstraints + gpu->sim.blocks - 1) / gpu->sim.blocks; gpu->sim.shake_threads_per_block = (gpu->sim.ShakeConstraints + gpu->sim.blocks - 1) / gpu->sim.blocks;
if (gpu->sim.shake_threads_per_block > gpu->sim.max_shake_threads_per_block) if (gpu->sim.shake_threads_per_block > gpu->sim.max_shake_threads_per_block)
gpu->sim.shake_threads_per_block = gpu->sim.max_shake_threads_per_block; gpu->sim.shake_threads_per_block = gpu->sim.max_shake_threads_per_block;
if (gpu->sim.shake_threads_per_block < 1) if (gpu->sim.shake_threads_per_block < 1)
gpu->sim.shake_threads_per_block = 1; gpu->sim.shake_threads_per_block = 1;
// Find connected constraints for LINCS. // Find connected constraints for CCMA.
vector<int> lincsConstraints; vector<int> ccmaConstraints;
for (unsigned i = 0; i < atom1.size(); i++) for (unsigned i = 0; i < atom1.size(); i++)
if (!isShakeAtom[atom1[i]]) if (!isShakeAtom[atom1[i]])
lincsConstraints.push_back(i); ccmaConstraints.push_back(i);
// Identify rigid clusters of atoms.
findRigidClusters(gpu, atom1, atom2, invMass1, invMass2, distance, lincsConstraints);
// Record the connections between constraints. // Record the connections between constraints.
int numLincs = (int) lincsConstraints.size(); int numCCMA = (int) ccmaConstraints.size();
vector<vector<int> > atomConstraints(gpu->natoms); vector<vector<int> > atomConstraints(gpu->natoms);
for (int i = 0; i < numLincs; i++) { for (int i = 0; i < numCCMA; i++) {
atomConstraints[atom1[lincsConstraints[i]]].push_back(i); atomConstraints[atom1[ccmaConstraints[i]]].push_back(i);
atomConstraints[atom2[lincsConstraints[i]]].push_back(i); atomConstraints[atom2[ccmaConstraints[i]]].push_back(i);
} }
vector<vector<int> > linkedConstraints(numLincs); vector<vector<int> > linkedConstraints(numCCMA);
for (unsigned atom = 0; atom < atomConstraints.size(); atom++) { for (unsigned atom = 0; atom < atomConstraints.size(); atom++) {
for (unsigned i = 0; i < atomConstraints[atom].size(); i++) for (unsigned i = 0; i < atomConstraints[atom].size(); i++)
for (unsigned j = 0; j < i; j++) { for (unsigned j = 0; j < i; j++) {
...@@ -949,17 +730,17 @@ void gpuSetConstraintParameters(gpuContext gpu, const vector<int>& atom1, const ...@@ -949,17 +730,17 @@ void gpuSetConstraintParameters(gpuContext gpu, const vector<int>& atom1, const
vector<vector<int> > atomAngles(gpu->natoms); vector<vector<int> > atomAngles(gpu->natoms);
for (int i = 0; i < gpu->sim.bond_angles; i++) for (int i = 0; i < gpu->sim.bond_angles; i++)
atomAngles[(*gpu->psBondAngleID1)[i].y].push_back(i); atomAngles[(*gpu->psBondAngleID1)[i].y].push_back(i);
vector<vector<pair<int, double> > > matrix(numLincs); vector<vector<pair<int, double> > > matrix(numCCMA);
if (numLincs > 0) { if (numCCMA > 0) {
for (int j = 0; j < numLincs; j++) { for (int j = 0; j < numCCMA; j++) {
for (int k = 0; k < numLincs; k++) { for (int k = 0; k < numCCMA; k++) {
if (j == k) { if (j == k) {
matrix[j].push_back(pair<int, double>(j, 1.0)); matrix[j].push_back(pair<int, double>(j, 1.0));
continue; continue;
} }
double scale; double scale;
int cj = lincsConstraints[j]; int cj = ccmaConstraints[j];
int ck = lincsConstraints[k]; int ck = ccmaConstraints[k];
int atomj0 = atom1[cj]; int atomj0 = atom1[cj];
int atomj1 = atom2[cj]; int atomj1 = atom2[cj];
int atomk0 = atom1[ck]; int atomk0 = atom1[ck];
...@@ -995,7 +776,7 @@ void gpuSetConstraintParameters(gpuContext gpu, const vector<int>& atom1, const ...@@ -995,7 +776,7 @@ void gpuSetConstraintParameters(gpuContext gpu, const vector<int>& atom1, const
// Look for a third constraint forming a triangle with these two. // Look for a third constraint forming a triangle with these two.
bool foundConstraint = false; bool foundConstraint = false;
for (int other = 0; other < numLincs; other++) { for (int other = 0; other < numCCMA; other++) {
if ((atom1[other] == atoma && atom2[other] == atomc) || (atom1[other] == atomc && atom2[other] == atoma)) { if ((atom1[other] == atoma && atom2[other] == atomc) || (atom1[other] == atomc && atom2[other] == atoma)) {
double d1 = distance[cj]; double d1 = distance[cj];
double d2 = distance[ck]; double d2 = distance[ck];
...@@ -1026,7 +807,7 @@ void gpuSetConstraintParameters(gpuContext gpu, const vector<int>& atom1, const ...@@ -1026,7 +807,7 @@ void gpuSetConstraintParameters(gpuContext gpu, const vector<int>& atom1, const
vector<int> matrixRowStart; vector<int> matrixRowStart;
vector<int> matrixColIndex; vector<int> matrixColIndex;
vector<double> matrixValue; vector<double> matrixValue;
for (int i = 0; i < numLincs; i++) { for (int i = 0; i < numCCMA; i++) {
matrixRowStart.push_back(matrixValue.size()); matrixRowStart.push_back(matrixValue.size());
for (int j = 0; j < (int) matrix[i].size(); j++) { for (int j = 0; j < (int) matrix[i].size(); j++) {
pair<int, double> element = matrix[i][j]; pair<int, double> element = matrix[i][j];
...@@ -1037,20 +818,20 @@ void gpuSetConstraintParameters(gpuContext gpu, const vector<int>& atom1, const ...@@ -1037,20 +818,20 @@ void gpuSetConstraintParameters(gpuContext gpu, const vector<int>& atom1, const
matrixRowStart.push_back(matrixValue.size()); matrixRowStart.push_back(matrixValue.size());
int *qRowStart, *qColIndex, *rRowStart, *rColIndex; int *qRowStart, *qColIndex, *rRowStart, *rColIndex;
double *qValue, *rValue; double *qValue, *rValue;
int result = QUERN_compute_qr(numLincs, numLincs, &matrixRowStart[0], &matrixColIndex[0], &matrixValue[0], NULL, int result = QUERN_compute_qr(numCCMA, numCCMA, &matrixRowStart[0], &matrixColIndex[0], &matrixValue[0], NULL,
&qRowStart, &qColIndex, &qValue, &rRowStart, &rColIndex, &rValue); &qRowStart, &qColIndex, &qValue, &rRowStart, &rColIndex, &rValue);
vector<double> rhs(numLincs); vector<double> rhs(numCCMA);
matrix.clear(); matrix.clear();
matrix.resize(numLincs); matrix.resize(numCCMA);
for (int i = 0; i < numLincs; i++) { for (int i = 0; i < numCCMA; i++) {
// Extract column i of the inverse matrix. // Extract column i of the inverse matrix.
for (int j = 0; j < numLincs; j++) for (int j = 0; j < numCCMA; j++)
rhs[j] = (i == j ? 1.0 : 0.0); rhs[j] = (i == j ? 1.0 : 0.0);
result = QUERN_multiply_with_q_transpose(numLincs, qRowStart, qColIndex, qValue, &rhs[0]); result = QUERN_multiply_with_q_transpose(numCCMA, qRowStart, qColIndex, qValue, &rhs[0]);
result = QUERN_solve_with_r(numLincs, rRowStart, rColIndex, rValue, &rhs[0], &rhs[0]); result = QUERN_solve_with_r(numCCMA, rRowStart, rColIndex, rValue, &rhs[0], &rhs[0]);
for (int j = 0; j < numLincs; j++) { for (int j = 0; j < numCCMA; j++) {
double value = rhs[j]*distance[lincsConstraints[i]]/distance[lincsConstraints[j]]; double value = rhs[j]*distance[ccmaConstraints[i]]/distance[ccmaConstraints[j]];
if (abs(value) > 0.1) if (abs(value) > 0.1)
matrix[j].push_back(pair<int, double>(i, value)); matrix[j].push_back(pair<int, double>(i, value));
} }
...@@ -1065,12 +846,12 @@ void gpuSetConstraintParameters(gpuContext gpu, const vector<int>& atom1, const ...@@ -1065,12 +846,12 @@ void gpuSetConstraintParameters(gpuContext gpu, const vector<int>& atom1, const
// Sort the constraints. // Sort the constraints.
vector<int> constraintOrder(numLincs); vector<int> constraintOrder(numCCMA);
for (int i = 0; i < numLincs; ++i) for (int i = 0; i < numCCMA; ++i)
constraintOrder[i] = i; constraintOrder[i] = i;
sort(constraintOrder.begin(), constraintOrder.end(), ConstraintOrderer(atom1, atom2)); sort(constraintOrder.begin(), constraintOrder.end(), ConstraintOrderer(atom1, atom2));
vector<int> inverseOrder(numLincs); vector<int> inverseOrder(numCCMA);
for (int i = 0; i < numLincs; ++i) for (int i = 0; i < numCCMA; ++i)
inverseOrder[constraintOrder[i]] = i; inverseOrder[constraintOrder[i]] = i;
for (int i = 0; i < matrix.size(); ++i) for (int i = 0; i < matrix.size(); ++i)
for (int j = 0; j < matrix[i].size(); ++j) for (int j = 0; j < matrix[i].size(); ++j)
...@@ -1078,100 +859,75 @@ void gpuSetConstraintParameters(gpuContext gpu, const vector<int>& atom1, const ...@@ -1078,100 +859,75 @@ void gpuSetConstraintParameters(gpuContext gpu, const vector<int>& atom1, const
// Fill in the CUDA streams. // Fill in the CUDA streams.
CUDAStream<int2>* psLincsAtoms = new CUDAStream<int2>(numLincs, 1, "LincsAtoms"); CUDAStream<int2>* psCcmaAtoms = new CUDAStream<int2>(numCCMA, 1, "CcmaAtoms");
gpu->psLincsAtoms = psLincsAtoms; gpu->psCcmaAtoms = psCcmaAtoms;
gpu->sim.pLincsAtoms = psLincsAtoms->_pDevData; gpu->sim.pCcmaAtoms = psCcmaAtoms->_pDevData;
CUDAStream<float4>* psLincsDistance = new CUDAStream<float4>(numLincs, 1, "LincsDistance"); CUDAStream<float4>* psCcmaDistance = new CUDAStream<float4>(numCCMA, 1, "CcmaDistance");
gpu->psLincsDistance = psLincsDistance; gpu->psCcmaDistance = psCcmaDistance;
gpu->sim.pLincsDistance = psLincsDistance->_pDevData; gpu->sim.pCcmaDistance = psCcmaDistance->_pDevData;
CUDAStream<int>* psLincsConnections = new CUDAStream<int>(numLincs*maxLinks, 1, "LincsConnections"); CUDAStream<int>* psCcmaAtomConstraints = new CUDAStream<int>(gpu->natoms*maxAtomConstraints, 1, "CcmaAtomConstraints");
gpu->psLincsConnections = psLincsConnections; gpu->psCcmaAtomConstraints = psCcmaAtomConstraints;
gpu->sim.pLincsConnections = psLincsConnections->_pDevData; gpu->sim.pCcmaAtomConstraints = psCcmaAtomConstraints->_pDevData;
CUDAStream<int>* psLincsNumConnections = new CUDAStream<int>(numLincs, 1, "LincsConnectionsIndex"); CUDAStream<int>* psCcmaNumAtomConstraints = new CUDAStream<int>(gpu->natoms, 1, "CcmaAtomConstraintsIndex");
gpu->psLincsNumConnections = psLincsNumConnections; gpu->psCcmaNumAtomConstraints = psCcmaNumAtomConstraints;
gpu->sim.pLincsNumConnections = psLincsNumConnections->_pDevData; gpu->sim.pCcmaNumAtomConstraints = psCcmaNumAtomConstraints->_pDevData;
CUDAStream<int>* psLincsAtomConstraints = new CUDAStream<int>(gpu->natoms*maxAtomConstraints, 1, "LincsAtomConstraints"); CUDAStream<float>* psCcmaDelta1 = new CUDAStream<float>(numCCMA, 1, "CcmaDelta1");
gpu->psLincsAtomConstraints = psLincsAtomConstraints; gpu->psCcmaDelta1 = psCcmaDelta1;
gpu->sim.pLincsAtomConstraints = psLincsAtomConstraints->_pDevData; gpu->sim.pCcmaDelta1 = psCcmaDelta1->_pDevData;
CUDAStream<int>* psLincsNumAtomConstraints = new CUDAStream<int>(gpu->natoms, 1, "LincsAtomConstraintsIndex"); CUDAStream<float>* psCcmaDelta2 = new CUDAStream<float>(numCCMA, 1, "CcmaDelta2");
gpu->psLincsNumAtomConstraints = psLincsNumAtomConstraints; gpu->psCcmaDelta2 = psCcmaDelta2;
gpu->sim.pLincsNumAtomConstraints = psLincsNumAtomConstraints->_pDevData; gpu->sim.pCcmaDelta2 = psCcmaDelta2->_pDevData;
CUDAStream<float>* psLincsS = new CUDAStream<float>(numLincs, 1, "LincsS");
gpu->psLincsS = psLincsS;
gpu->sim.pLincsS = psLincsS->_pDevData;
CUDAStream<float>* psLincsCoupling = new CUDAStream<float>(numLincs*maxLinks, 1, "LincsCoupling");
gpu->psLincsCoupling = psLincsCoupling;
gpu->sim.pLincsCoupling = psLincsCoupling->_pDevData;
CUDAStream<float>* psLincsRhs1 = new CUDAStream<float>(numLincs, 1, "LincsRhs1");
gpu->psLincsRhs1 = psLincsRhs1;
gpu->sim.pLincsRhs1 = psLincsRhs1->_pDevData;
CUDAStream<float>* psLincsRhs2 = new CUDAStream<float>(numLincs, 1, "LincsRhs2");
gpu->psLincsRhs2 = psLincsRhs2;
gpu->sim.pLincsRhs2 = psLincsRhs2->_pDevData;
CUDAStream<float>* psLincsSolution = new CUDAStream<float>(numLincs, 1, "LincsSolution");
gpu->psLincsSolution = psLincsSolution;
gpu->sim.pLincsSolution = psLincsSolution->_pDevData;
CUDAStream<short>* psSyncCounter = new CUDAStream<short>(3*gpu->sim.blocks, 1, "SyncCounter"); CUDAStream<short>* psSyncCounter = new CUDAStream<short>(3*gpu->sim.blocks, 1, "SyncCounter");
gpu->psSyncCounter = psSyncCounter; gpu->psSyncCounter = psSyncCounter;
gpu->sim.pSyncCounter = psSyncCounter->_pDevData; gpu->sim.pSyncCounter = psSyncCounter->_pDevData;
CUDAStream<unsigned int>* psRequiredIterations = new CUDAStream<unsigned int>(1, 1, "RequiredIterations"); CUDAStream<unsigned int>* psRequiredIterations = new CUDAStream<unsigned int>(1, 1, "RequiredIterations");
gpu->psRequiredIterations = psRequiredIterations; gpu->psRequiredIterations = psRequiredIterations;
gpu->sim.pRequiredIterations = psRequiredIterations->_pDevData; gpu->sim.pRequiredIterations = psRequiredIterations->_pDevData;
CUDAStream<float>* psShakeReducedMass = new CUDAStream<float>(numLincs, 1, "LincsSolution"); CUDAStream<float>* psCcmaReducedMass = new CUDAStream<float>(numCCMA, 1, "CcmaReducedMass");
gpu->psShakeReducedMass = psShakeReducedMass; gpu->psCcmaReducedMass = psCcmaReducedMass;
gpu->sim.pShakeReducedMass = psShakeReducedMass->_pDevData; gpu->sim.pCcmaReducedMass = psCcmaReducedMass->_pDevData;
CUDAStream<unsigned int>* psConstraintMatrixColumn = new CUDAStream<unsigned int>(numLincs*maxRowElements, 1, "ConstraintMatrixColumn"); CUDAStream<unsigned int>* psConstraintMatrixColumn = new CUDAStream<unsigned int>(numCCMA*maxRowElements, 1, "ConstraintMatrixColumn");
gpu->psConstraintMatrixColumn = psConstraintMatrixColumn; gpu->psConstraintMatrixColumn = psConstraintMatrixColumn;
gpu->sim.pConstraintMatrixColumn = psConstraintMatrixColumn->_pDevData; gpu->sim.pConstraintMatrixColumn = psConstraintMatrixColumn->_pDevData;
CUDAStream<float>* psConstraintMatrixValue = new CUDAStream<float>(numLincs*maxRowElements, 1, "ConstraintMatrixValue"); CUDAStream<float>* psConstraintMatrixValue = new CUDAStream<float>(numCCMA*maxRowElements, 1, "ConstraintMatrixValue");
gpu->psConstraintMatrixValue = psConstraintMatrixValue; gpu->psConstraintMatrixValue = psConstraintMatrixValue;
gpu->sim.pConstraintMatrixValue = psConstraintMatrixValue->_pDevData; gpu->sim.pConstraintMatrixValue = psConstraintMatrixValue->_pDevData;
gpu->sim.lincsConstraints = numLincs; gpu->sim.ccmaConstraints = numCCMA;
for (int i = 0; i < numLincs; i++) { for (int i = 0; i < numCCMA; i++) {
int index = constraintOrder[i]; int index = constraintOrder[i];
int c = lincsConstraints[index]; int c = ccmaConstraints[index];
(*psLincsAtoms)[i].x = atom1[c]; (*psCcmaAtoms)[i].x = atom1[c];
(*psLincsAtoms)[i].y = atom2[c]; (*psCcmaAtoms)[i].y = atom2[c];
(*psLincsDistance)[i].w = distance[c]; (*psCcmaDistance)[i].w = distance[c];
(*psLincsS)[i] = 1.0f/sqrt(invMass1[c]+invMass2[c]); (*psCcmaReducedMass)[i] = 0.5f/(invMass1[c]+invMass2[c]);
(*psShakeReducedMass)[i] = 0.5f/(invMass1[c]+invMass2[c]);
(*psLincsNumConnections)[i] = linkedConstraints[index].size();
for (unsigned int j = 0; j < linkedConstraints[index].size(); j++)
(*psLincsConnections)[i+j*numLincs] = linkedConstraints[index][j];
for (unsigned int j = 0; j < matrix[index].size(); j++) { for (unsigned int j = 0; j < matrix[index].size(); j++) {
(*psConstraintMatrixColumn)[i+j*numLincs] = matrix[index][j].first; (*psConstraintMatrixColumn)[i+j*numCCMA] = matrix[index][j].first;
(*psConstraintMatrixValue)[i+j*numLincs] = matrix[index][j].second; (*psConstraintMatrixValue)[i+j*numCCMA] = matrix[index][j].second;
} }
(*psConstraintMatrixColumn)[i+matrix[index].size()*numLincs] = numLincs; (*psConstraintMatrixColumn)[i+matrix[index].size()*numCCMA] = numCCMA;
} }
for (unsigned int i = 0; i < psSyncCounter->_length; i++) for (unsigned int i = 0; i < psSyncCounter->_length; i++)
(*psSyncCounter)[i] = -1; (*psSyncCounter)[i] = -1;
for (unsigned int i = 0; i < atomConstraints.size(); i++) { for (unsigned int i = 0; i < atomConstraints.size(); i++) {
(*psLincsNumAtomConstraints)[i] = atomConstraints[i].size(); (*psCcmaNumAtomConstraints)[i] = atomConstraints[i].size();
for (unsigned int j = 0; j < atomConstraints[i].size(); j++) { for (unsigned int j = 0; j < atomConstraints[i].size(); j++) {
bool forward = (atom1[lincsConstraints[atomConstraints[i][j]]] == i); bool forward = (atom1[ccmaConstraints[atomConstraints[i][j]]] == i);
(*psLincsAtomConstraints)[i+j*gpu->natoms] = (forward ? inverseOrder[atomConstraints[i][j]]+1 : -inverseOrder[atomConstraints[i][j]]-1); (*psCcmaAtomConstraints)[i+j*gpu->natoms] = (forward ? inverseOrder[atomConstraints[i][j]]+1 : -inverseOrder[atomConstraints[i][j]]-1);
} }
} }
psLincsAtoms->Upload(); psCcmaAtoms->Upload();
psLincsDistance->Upload(); psCcmaDistance->Upload();
psLincsS->Upload(); psCcmaReducedMass->Upload();
psShakeReducedMass->Upload(); psCcmaAtomConstraints->Upload();
psLincsConnections->Upload(); psCcmaNumAtomConstraints->Upload();
psLincsNumConnections->Upload();
psLincsAtomConstraints->Upload();
psLincsNumAtomConstraints->Upload();
psSyncCounter->Upload(); psSyncCounter->Upload();
psConstraintMatrixColumn->Upload(); psConstraintMatrixColumn->Upload();
psConstraintMatrixValue->Upload(); psConstraintMatrixValue->Upload();
gpu->sim.lincsTerms = lincsTerms; gpu->sim.ccma_threads_per_block = (gpu->sim.ccmaConstraints + gpu->sim.blocks - 1) / gpu->sim.blocks;
gpu->sim.lincs_threads_per_block = (gpu->sim.lincsConstraints + gpu->sim.blocks - 1) / gpu->sim.blocks; if (gpu->sim.ccma_threads_per_block > gpu->sim.threads_per_block)
if (gpu->sim.lincs_threads_per_block > gpu->sim.threads_per_block) gpu->sim.ccma_threads_per_block = gpu->sim.threads_per_block;
gpu->sim.lincs_threads_per_block = gpu->sim.threads_per_block; if (gpu->sim.ccma_threads_per_block < gpu->sim.blocks)
if (gpu->sim.lincs_threads_per_block < gpu->sim.blocks) gpu->sim.ccma_threads_per_block = gpu->sim.blocks;
gpu->sim.lincs_threads_per_block = gpu->sim.blocks;
if (gpu->sim.lincs_threads_per_block%gpu->sim.clusterShakeBlockSize != 0)
gpu->sim.lincs_threads_per_block += gpu->sim.clusterShakeBlockSize - gpu->sim.lincs_threads_per_block%gpu->sim.clusterShakeBlockSize;
// count number of atoms w/o constraint // count number of atoms w/o constraint
...@@ -1533,23 +1289,15 @@ void* gpuInit(int numAtoms) ...@@ -1533,23 +1289,15 @@ void* gpuInit(int numAtoms)
gpu->psInteractionCount = NULL; gpu->psInteractionCount = NULL;
gpu->psGridBoundingBox = NULL; gpu->psGridBoundingBox = NULL;
gpu->psGridCenter = NULL; gpu->psGridCenter = NULL;
gpu->psLincsAtoms = NULL; gpu->psCcmaAtoms = NULL;
gpu->psLincsDistance = NULL; gpu->psCcmaDistance = NULL;
gpu->psLincsConnections = NULL; gpu->psCcmaAtomConstraints = NULL;
gpu->psLincsNumConnections = NULL; gpu->psCcmaNumAtomConstraints = NULL;
gpu->psLincsAtomConstraints = NULL; gpu->psCcmaDelta1 = NULL;
gpu->psLincsNumAtomConstraints = NULL; gpu->psCcmaDelta2 = NULL;
gpu->psLincsS = NULL;
gpu->psLincsCoupling = NULL;
gpu->psLincsRhs1 = NULL;
gpu->psLincsRhs2 = NULL;
gpu->psLincsSolution = NULL;
gpu->psSyncCounter = NULL; gpu->psSyncCounter = NULL;
gpu->psRequiredIterations = NULL; gpu->psRequiredIterations = NULL;
gpu->psShakeReducedMass = NULL; gpu->psCcmaReducedMass = NULL;
gpu->psRigidClusterConstraintIndex = NULL;
gpu->psRigidClusterMatrix = NULL;
gpu->psRigidClusterMatrixIndex = NULL;
gpu->psConstraintMatrixColumn = NULL; gpu->psConstraintMatrixColumn = NULL;
gpu->psConstraintMatrixValue = NULL; gpu->psConstraintMatrixValue = NULL;
...@@ -1693,23 +1441,15 @@ void gpuShutDown(gpuContext gpu) ...@@ -1693,23 +1441,15 @@ void gpuShutDown(gpuContext gpu)
delete gpu->psAtomIndex; delete gpu->psAtomIndex;
delete gpu->psGridBoundingBox; delete gpu->psGridBoundingBox;
delete gpu->psGridCenter; delete gpu->psGridCenter;
delete gpu->psLincsAtoms; delete gpu->psCcmaAtoms;
delete gpu->psLincsDistance; delete gpu->psCcmaDistance;
delete gpu->psLincsConnections; delete gpu->psCcmaAtomConstraints;
delete gpu->psLincsNumConnections; delete gpu->psCcmaNumAtomConstraints;
delete gpu->psLincsAtomConstraints; delete gpu->psCcmaDelta1;
delete gpu->psLincsNumAtomConstraints; delete gpu->psCcmaDelta2;
delete gpu->psLincsS;
delete gpu->psLincsCoupling;
delete gpu->psLincsRhs1;
delete gpu->psLincsRhs2;
delete gpu->psLincsSolution;
delete gpu->psSyncCounter; delete gpu->psSyncCounter;
delete gpu->psRequiredIterations; delete gpu->psRequiredIterations;
delete gpu->psShakeReducedMass; delete gpu->psCcmaReducedMass;
delete gpu->psRigidClusterConstraintIndex;
delete gpu->psRigidClusterMatrix;
delete gpu->psRigidClusterMatrixIndex;
delete gpu->psConstraintMatrixColumn; delete gpu->psConstraintMatrixColumn;
delete gpu->psConstraintMatrixValue; delete gpu->psConstraintMatrixValue;
if (gpu->cudpp != 0) if (gpu->cudpp != 0)
...@@ -2021,8 +1761,7 @@ int gpuSetConstants(gpuContext gpu) ...@@ -2021,8 +1761,7 @@ int gpuSetConstants(gpuContext gpu)
SetVerletUpdateSim(gpu); SetVerletUpdateSim(gpu);
SetBrownianUpdateSim(gpu); SetBrownianUpdateSim(gpu);
SetSettleSim(gpu); SetSettleSim(gpu);
SetCShakeSim(gpu); SetCCMASim(gpu);
SetLincsSim(gpu);
SetRandomSim(gpu); SetRandomSim(gpu);
return 1; return 1;
} }
...@@ -2066,11 +1805,11 @@ static void findMoleculeGroups(gpuContext gpu) ...@@ -2066,11 +1805,11 @@ static void findMoleculeGroups(gpuContext gpu)
constraints.push_back(Constraint(atom1, atom3, distance12*distance12)); constraints.push_back(Constraint(atom1, atom3, distance12*distance12));
constraints.push_back(Constraint(atom2, atom3, distance23*distance23)); constraints.push_back(Constraint(atom2, atom3, distance23*distance23));
} }
for (int i = 0; i < (int)gpu->sim.lincsConstraints; i++) for (int i = 0; i < (int)gpu->sim.ccmaConstraints; i++)
{ {
int atom1 = (*gpu->psLincsAtoms)[i].x; int atom1 = (*gpu->psCcmaAtoms)[i].x;
int atom2 = (*gpu->psLincsAtoms)[i].y; int atom2 = (*gpu->psCcmaAtoms)[i].y;
float distance2 = (*gpu->psLincsDistance)[i].w; float distance2 = (*gpu->psCcmaDistance)[i].w;
constraints.push_back(Constraint(atom1, atom2, distance2)); constraints.push_back(Constraint(atom1, atom2, distance2));
} }
......
platforms/cuda/src/kernels/gputypes.h
View file @ 05259f77
...@@ -127,20 +127,15 @@ struct _gpuContext { ...@@ -127,20 +127,15 @@ struct _gpuContext {
CUDAStream<int>* psAtomIndex; // The original index of each atom CUDAStream<int>* psAtomIndex; // The original index of each atom
CUDAStream<float4>* psGridBoundingBox; // The size of each grid cell CUDAStream<float4>* psGridBoundingBox; // The size of each grid cell
CUDAStream<float4>* psGridCenter; // The center and radius for each grid cell CUDAStream<float4>* psGridCenter; // The center and radius for each grid cell
CUDAStream<int2>* psLincsAtoms; // The atoms connected by each LINCS constraint CUDAStream<int2>* psCcmaAtoms; // The atoms connected by each CCMA constraint
CUDAStream<float4>* psLincsDistance; // The displacement vector (x, y, z) and constraint distance (w) for each LINCS constraint CUDAStream<float4>* psCcmaDistance; // The displacement vector (x, y, z) and constraint distance (w) for each CCMA constraint
CUDAStream<int>* psLincsConnections; // The indices of constraints that other constraints are connected to CUDAStream<int>* psCcmaAtomConstraints; // The indices of constraints involving each atom
CUDAStream<int>* psLincsNumConnections; // The number of other constraints that each constraint is linked to CUDAStream<int>* psCcmaNumAtomConstraints; // The number of constraints involving each atom
CUDAStream<int>* psLincsAtomConstraints; // The indices of constraints involving each atom CUDAStream<float>* psCcmaDelta1; // Workspace for CCMA
CUDAStream<int>* psLincsNumAtomConstraints; // The number of constraints involving each atom CUDAStream<float>* psCcmaDelta2; // Workspace for CCMA
CUDAStream<float>* psLincsS; // S matrix for LINCS
CUDAStream<float>* psLincsCoupling; // Coupling matrix for LINCS
CUDAStream<float>* psLincsRhs1; // Workspace for LINCS
CUDAStream<float>* psLincsRhs2; // Workspace for LINCS
CUDAStream<float>* psLincsSolution; // Workspace for LINCS
CUDAStream<short>* psSyncCounter; // Used for global thread synchronization CUDAStream<short>* psSyncCounter; // Used for global thread synchronization
CUDAStream<unsigned int>* psRequiredIterations; // Used by SHAKE to communicate whether iteration has converged CUDAStream<unsigned int>* psRequiredIterations; // Used by SHAKE to communicate whether iteration has converged
CUDAStream<float>* psShakeReducedMass; // The reduced mass for each SHAKE constraint CUDAStream<float>* psCcmaReducedMass; // The reduced mass for each SHAKE constraint
CUDAStream<float>* psRigidClusterMatrix;// The inverse constraint matrix for each rigid cluster CUDAStream<float>* psRigidClusterMatrix;// The inverse constraint matrix for each rigid cluster
CUDAStream<unsigned int>* psRigidClusterConstraintIndex; // The index of each cluster in the stream containing cluster constraints. CUDAStream<unsigned int>* psRigidClusterConstraintIndex; // The index of each cluster in the stream containing cluster constraints.
CUDAStream<unsigned int>* psRigidClusterMatrixIndex; // The index of each cluster in the stream containing cluster matrices. CUDAStream<unsigned int>* psRigidClusterMatrixIndex; // The index of each cluster in the stream containing cluster matrices.
...@@ -192,7 +187,7 @@ void gpuSetObcParameters(gpuContext gpu, float innerDielectric, float solventDie ...@@ -192,7 +187,7 @@ void gpuSetObcParameters(gpuContext gpu, float innerDielectric, float solventDie
extern "C" extern "C"
void gpuSetConstraintParameters(gpuContext gpu, const std::vector<int>& atom1, const std::vector<int>& atom2, const std::vector<float>& distance, void gpuSetConstraintParameters(gpuContext gpu, const std::vector<int>& atom1, const std::vector<int>& atom2, const std::vector<float>& distance,
const std::vector<float>& invMass1, const std::vector<float>& invMass2, float shakeTolerance, unsigned int lincsTerms); const std::vector<float>& invMass1, const std::vector<float>& invMass2, float constraintTolerance);
extern "C" extern "C"
int gpuAllocateInitialBuffers(gpuContext gpu); int gpuAllocateInitialBuffers(gpuContext gpu);
......
platforms/cuda/src/kernels/kCShake.cu → platforms/cuda/src/kernels/kCCMA.cu
View file @ 05259f77
...@@ -34,14 +34,14 @@ using namespace std; ...@@ -34,14 +34,14 @@ using namespace std;
static __constant__ cudaGmxSimulation cSim; static __constant__ cudaGmxSimulation cSim;
void SetCShakeSim(gpuContext gpu) void SetCCMASim(gpuContext gpu)
{ {
cudaError_t status; cudaError_t status;
status = cudaMemcpyToSymbol(cSim, &gpu->sim, sizeof(cudaGmxSimulation)); status = cudaMemcpyToSymbol(cSim, &gpu->sim, sizeof(cudaGmxSimulation));
RTERROR(status, "cudaMemcpyToSymbol: SetSim copy to cSim failed"); RTERROR(status, "cudaMemcpyToSymbol: SetSim copy to cSim failed");
} }
void GetCShakeSim(gpuContext gpu) void GetCCMASim(gpuContext gpu)
{ {
cudaError_t status; cudaError_t status;
status = cudaMemcpyFromSymbol(&gpu->sim, cSim, sizeof(cudaGmxSimulation)); status = cudaMemcpyFromSymbol(&gpu->sim, cSim, sizeof(cudaGmxSimulation));
...@@ -66,7 +66,7 @@ __device__ void kSyncAllThreads_kernel(short* syncCounter, short newCount) ...@@ -66,7 +66,7 @@ __device__ void kSyncAllThreads_kernel(short* syncCounter, short newCount)
__syncthreads(); __syncthreads();
} }
__global__ void kApplyCShake_kernel(float4* atomPositions, bool addOldPosition) __global__ void kApplyCCMA_kernel(float4* atomPositions, bool addOldPosition)
{ {
// Initialize counters used for monitoring convergence and doing global thread synchronization. // Initialize counters used for monitoring convergence and doing global thread synchronization.
...@@ -82,16 +82,16 @@ __global__ void kApplyCShake_kernel(float4* atomPositions, bool addOldPosition) ...@@ -82,16 +82,16 @@ __global__ void kApplyCShake_kernel(float4* atomPositions, bool addOldPosition)
// Calculate the direction of each constraint. // Calculate the direction of each constraint.
unsigned int pos = threadIdx.x + blockIdx.x * blockDim.x; unsigned int pos = threadIdx.x + blockIdx.x * blockDim.x;
while (pos < cSim.lincsConstraints) while (pos < cSim.ccmaConstraints)
{ {
int2 atoms = cSim.pLincsAtoms[pos]; int2 atoms = cSim.pCcmaAtoms[pos];
float4 dir = cSim.pLincsDistance[pos]; float4 dir = cSim.pCcmaDistance[pos];
float4 oldPos1 = cSim.pOldPosq[atoms.x]; float4 oldPos1 = cSim.pOldPosq[atoms.x];
float4 oldPos2 = cSim.pOldPosq[atoms.y]; float4 oldPos2 = cSim.pOldPosq[atoms.y];
dir.x = oldPos1.x-oldPos2.x; dir.x = oldPos1.x-oldPos2.x;
dir.y = oldPos1.y-oldPos2.y; dir.y = oldPos1.y-oldPos2.y;
dir.z = oldPos1.z-oldPos2.z; dir.z = oldPos1.z-oldPos2.z;
cSim.pLincsDistance[pos] = dir; cSim.pCcmaDistance[pos] = dir;
pos += blockDim.x*gridDim.x; pos += blockDim.x*gridDim.x;
} }
__syncthreads(); __syncthreads();
...@@ -106,12 +106,12 @@ __global__ void kApplyCShake_kernel(float4* atomPositions, bool addOldPosition) ...@@ -106,12 +106,12 @@ __global__ void kApplyCShake_kernel(float4* atomPositions, bool addOldPosition)
// Calculate the constraint force for each constraint. // Calculate the constraint force for each constraint.
pos = threadIdx.x + blockIdx.x * blockDim.x; pos = threadIdx.x + blockIdx.x * blockDim.x;
while (pos < cSim.lincsConstraints) while (pos < cSim.ccmaConstraints)
{ {
int2 atoms = cSim.pLincsAtoms[pos]; int2 atoms = cSim.pCcmaAtoms[pos];
float4 delta1 = atomPositions[atoms.x]; float4 delta1 = atomPositions[atoms.x];
float4 delta2 = atomPositions[atoms.y]; float4 delta2 = atomPositions[atoms.y];
float4 dir = cSim.pLincsDistance[pos]; float4 dir = cSim.pCcmaDistance[pos];
float3 rp_ij = make_float3(delta1.x-delta2.x, delta1.y-delta2.y, delta1.z-delta2.z); float3 rp_ij = make_float3(delta1.x-delta2.x, delta1.y-delta2.y, delta1.z-delta2.z);
if (addOldPosition) if (addOldPosition)
{ {
...@@ -124,8 +124,8 @@ __global__ void kApplyCShake_kernel(float4* atomPositions, bool addOldPosition) ...@@ -124,8 +124,8 @@ __global__ void kApplyCShake_kernel(float4* atomPositions, bool addOldPosition)
float diff = dist2 - rp2; float diff = dist2 - rp2;
float rrpr = rp_ij.x*dir.x + rp_ij.y*dir.y + rp_ij.z*dir.z; float rrpr = rp_ij.x*dir.x + rp_ij.y*dir.y + rp_ij.z*dir.z;
float d_ij2 = dir.x*dir.x + dir.y*dir.y + dir.z*dir.z; float d_ij2 = dir.x*dir.x + dir.y*dir.y + dir.z*dir.z;
float reducedMass = cSim.pShakeReducedMass[pos]; float reducedMass = cSim.pCcmaReducedMass[pos];
cSim.pLincsRhs1[pos] = (rrpr > d_ij2*1e-6f ? reducedMass*diff/rrpr : 0.0f); cSim.pCcmaDelta1[pos] = (rrpr > d_ij2*1e-6f ? reducedMass*diff/rrpr : 0.0f);
if (requiredIterations == iteration && (rp2 < lowerTol*dist2 || rp2 > upperTol*dist2)) if (requiredIterations == iteration && (rp2 < lowerTol*dist2 || rp2 > upperTol*dist2))
requiredIterations = iteration+1; requiredIterations = iteration+1;
pos += blockDim.x * gridDim.x; pos += blockDim.x * gridDim.x;
...@@ -140,18 +140,18 @@ __global__ void kApplyCShake_kernel(float4* atomPositions, bool addOldPosition) ...@@ -140,18 +140,18 @@ __global__ void kApplyCShake_kernel(float4* atomPositions, bool addOldPosition)
// Multiply by the inverse constraint matrix. // Multiply by the inverse constraint matrix.
pos = threadIdx.x + blockIdx.x * blockDim.x; pos = threadIdx.x + blockIdx.x * blockDim.x;
while (pos < cSim.lincsConstraints) while (pos < cSim.ccmaConstraints)
{ {
float sum = 0.0f; float sum = 0.0f;
for (unsigned int i = 0; ; i++) for (unsigned int i = 0; ; i++)
{ {
unsigned int index = pos+i*cSim.lincsConstraints; unsigned int index = pos+i*cSim.ccmaConstraints;
unsigned int column = cSim.pConstraintMatrixColumn[index]; unsigned int column = cSim.pConstraintMatrixColumn[index];
if (column >= cSim.lincsConstraints) if (column >= cSim.ccmaConstraints)
break; break;
sum += cSim.pLincsRhs1[column]*cSim.pConstraintMatrixValue[index]; sum += cSim.pCcmaDelta1[column]*cSim.pConstraintMatrixValue[index];
} }
cSim.pLincsSolution[pos] = sum; cSim.pCcmaDelta2[pos] = sum;
pos += blockDim.x * gridDim.x; pos += blockDim.x * gridDim.x;
} }
kSyncAllThreads_kernel(&cSim.pSyncCounter[gridDim.x], iteration); kSyncAllThreads_kernel(&cSim.pSyncCounter[gridDim.x], iteration);
...@@ -164,16 +164,16 @@ __global__ void kApplyCShake_kernel(float4* atomPositions, bool addOldPosition) ...@@ -164,16 +164,16 @@ __global__ void kApplyCShake_kernel(float4* atomPositions, bool addOldPosition)
{ {
float4 atomPos = atomPositions[pos]; float4 atomPos = atomPositions[pos];
float invMass = cSim.pVelm4[pos].w; float invMass = cSim.pVelm4[pos].w;
int num = cSim.pLincsNumAtomConstraints[pos]; int num = cSim.pCcmaNumAtomConstraints[pos];
for (int i = 0; i < num; i++) for (int i = 0; i < num; i++)
{ {
int index = pos+i*cSim.atoms; int index = pos+i*cSim.atoms;
int constraint = cSim.pLincsAtomConstraints[index]; int constraint = cSim.pCcmaAtomConstraints[index];
bool forward = (constraint > 0); bool forward = (constraint > 0);
constraint = (forward ? constraint-1 : -constraint-1); constraint = (forward ? constraint-1 : -constraint-1);
float constraintForce = damping*invMass*cSim.pLincsSolution[constraint]; float constraintForce = damping*invMass*cSim.pCcmaDelta2[constraint];
constraintForce = (forward ? constraintForce : -constraintForce); constraintForce = (forward ? constraintForce : -constraintForce);
float4 dir = cSim.pLincsDistance[constraint]; float4 dir = cSim.pCcmaDistance[constraint];
atomPos.x += constraintForce*dir.x; atomPos.x += constraintForce*dir.x;
atomPos.y += constraintForce*dir.y; atomPos.y += constraintForce*dir.y;
atomPos.z += constraintForce*dir.z; atomPos.z += constraintForce*dir.z;
...@@ -192,22 +192,22 @@ __global__ void kApplyCShake_kernel(float4* atomPositions, bool addOldPosition) ...@@ -192,22 +192,22 @@ __global__ void kApplyCShake_kernel(float4* atomPositions, bool addOldPosition)
cSim.pSyncCounter[blockIdx.x] = -1; cSim.pSyncCounter[blockIdx.x] = -1;
} }
void kApplyFirstCShake(gpuContext gpu) void kApplyFirstCCMA(gpuContext gpu)
{ {
// printf("kApplyFirstCShake\n"); // printf("kApplyFirstCCMA\n");
if (gpu->sim.lincsConstraints > 0) if (gpu->sim.ccmaConstraints > 0)
{ {
kApplyCShake_kernel<<<gpu->sim.blocks, gpu->sim.lincs_threads_per_block>>>(gpu->sim.pPosqP, true); kApplyCCMA_kernel<<<gpu->sim.blocks, gpu->sim.ccma_threads_per_block>>>(gpu->sim.pPosqP, true);
LAUNCHERROR("kApplyCShake"); LAUNCHERROR("kApplyCCMA");
} }
} }
void kApplySecondCShake(gpuContext gpu) void kApplySecondCCMA(gpuContext gpu)
{ {
// printf("kApplySecondCShake\n"); // printf("kApplySecondCCMA\n");
if (gpu->sim.lincsConstraints > 0) if (gpu->sim.ccmaConstraints > 0)
{ {
kApplyCShake_kernel<<<gpu->sim.blocks, gpu->sim.lincs_threads_per_block>>>(gpu->sim.pPosq, false); kApplyCCMA_kernel<<<gpu->sim.blocks, gpu->sim.ccma_threads_per_block>>>(gpu->sim.pPosq, false);
LAUNCHERROR("kApplyCShake"); LAUNCHERROR("kApplyCCMA");
} }
} }
platforms/cuda/src/kernels/kLincs.cu deleted 100644 → 0
View file @ f62ba9d8
/* -------------------------------------------------------------------------- *
* OpenMM *
* -------------------------------------------------------------------------- *
* This is part of the OpenMM molecular simulation toolkit originating from *
* Simbios, the NIH National Center for Physics-Based Simulation of *
* Biological Structures at Stanford, funded under the NIH Roadmap for *
* Medical Research, grant U54 GM072970. See https://simtk.org. *
* *
* Portions copyright (c) 2009 Stanford University and the Authors. *
* Authors: Scott Le Grand, Peter Eastman *
* Contributors: *
* *
* This program 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 3 of the License, or *
* (at your option) any later version. *
* *
* This program 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 this program. If not, see <http://www.gnu.org/licenses/>. *
* -------------------------------------------------------------------------- */
#include <cuda.h>
#include <vector_functions.h>
using namespace std;
#include "gputypes.h"
static __constant__ cudaGmxSimulation cSim;
void SetLincsSim(gpuContext gpu)
{
cudaError_t status;
status = cudaMemcpyToSymbol(cSim, &gpu->sim, sizeof(cudaGmxSimulation));
RTERROR(status, "cudaMemcpyToSymbol: SetSim copy to cSim failed");
}
void GetLincsSim(gpuContext gpu)
{
cudaError_t status;
status = cudaMemcpyFromSymbol(&gpu->sim, cSim, sizeof(cudaGmxSimulation));
RTERROR(status, "cudaMemcpyFromSymbol: SetSim copy from cSim failed");
}
/**
* Synchronize all threads across all blocks.
*/
__device__ void kSyncAllThreads_kernel(short* syncCounter, short newCount)
{
// short* syncCounter = &cSim.pSyncCounter[newCount%2 == 0 ? 0 : gridDim.x];
__syncthreads();
if (threadIdx.x == 0)
syncCounter[blockIdx.x] = newCount;
if (threadIdx.x < gridDim.x)
{
volatile short* counter = &syncCounter[threadIdx.x];
do
{
} while (*counter != newCount);
}
__syncthreads();
}
__global__ void kSolveLincsMatrix_kernel(float4* atomPositions)
{
for (unsigned int iteration = 0; iteration < cSim.lincsTerms; iteration++) {
float* rhs1 = (iteration%2 == 0 ? cSim.pLincsRhs1 : cSim.pLincsRhs2);
float* rhs2 = (iteration%2 == 0 ? cSim.pLincsRhs2 : cSim.pLincsRhs1);
unsigned int pos = threadIdx.x + blockIdx.x * blockDim.x;
while (pos < cSim.lincsConstraints)
{
float rhs = 0.0f;
int num = cSim.pLincsNumConnections[pos];
for (int i = 0; i < num; i++)
{
int index = pos+i*cSim.lincsConstraints;
int otherConstraint = cSim.pLincsConnections[index];
rhs += cSim.pLincsCoupling[index]*rhs1[otherConstraint];
}
rhs2[pos] = rhs;
cSim.pLincsSolution[pos] += rhs;
pos += blockDim.x * gridDim.x;
}
kSyncAllThreads_kernel(&cSim.pSyncCounter[iteration%2 == 0 ? 0 : gridDim.x], iteration);
}
// Update the atom positions based on the solution to the matrix equations.
unsigned int pos = threadIdx.x + blockIdx.x * blockDim.x;
while (pos < cSim.atoms)
{
float4 atomPos = atomPositions[pos];
float invMass = cSim.pVelm4[pos].w;
int num = cSim.pLincsNumAtomConstraints[pos];
for (int i = 0; i < num; i++)
{
int index = pos+i*cSim.atoms;
int constraint = cSim.pLincsAtomConstraints[index];
bool forward = (constraint > 0);
constraint = (forward ? constraint-1 : -constraint-1);
float4 dir = cSim.pLincsDistance[constraint];
float c = invMass*cSim.pLincsS[constraint]*cSim.pLincsSolution[constraint];
c = (forward ? -c : c);
atomPos.x += c*dir.x;
atomPos.y += c*dir.y;
atomPos.z += c*dir.z;
}
atomPositions[pos] = atomPos;
pos += blockDim.x * gridDim.x;
}
}
__global__ void kApplyLincsPart1_kernel(float4* atomPositions, bool addOldPosition)
{
// Calculate the direction of each constraint, along with the initial RHS and solution vectors.
unsigned int pos = threadIdx.x + blockIdx.x * blockDim.x;
while (pos < cSim.lincsConstraints)
{
int2 atoms = cSim.pLincsAtoms[pos];
float4 delta1 = atomPositions[atoms.x];
float4 delta2 = atomPositions[atoms.y];
float4 dir = cSim.pLincsDistance[pos];
if (addOldPosition)
{
float4 oldPos1 = cSim.pOldPosq[atoms.x];
float4 oldPos2 = cSim.pOldPosq[atoms.y];
dir.x = (oldPos1.x-oldPos2.x)+(delta1.x-delta2.x);
dir.y = (oldPos1.y-oldPos2.y)+(delta1.y-delta2.y);
dir.z = (oldPos1.z-oldPos2.z)+(delta1.z-delta2.z);
}
else
{
dir.x = delta1.x-delta2.x;
dir.y = delta1.y-delta2.y;
dir.z = delta1.z-delta2.z;
}
float invLength = 1.0f/sqrt(dir.x*dir.x + dir.y*dir.y + dir.z*dir.z);
dir.x *= invLength;
dir.y *= invLength;
dir.z *= invLength;
cSim.pLincsDistance[pos] = dir;
float diff = cSim.pLincsS[pos]*(1.0f/invLength-dir.w);
cSim.pLincsRhs1[pos] = diff;
cSim.pLincsSolution[pos] = diff;
pos += blockDim.x * gridDim.x;
}
kSyncAllThreads_kernel(cSim.pSyncCounter, cSim.lincsTerms+1);
// Build the coupling matrix.
pos = threadIdx.x + blockIdx.x * blockDim.x;
while (pos < cSim.lincsConstraints)
{
float4 dir1 = cSim.pLincsDistance[pos];
int2 atoms1 = cSim.pLincsAtoms[pos];
int num = cSim.pLincsNumConnections[pos];
float s = cSim.pLincsS[pos];
float invMass = cSim.pVelm4[atoms1.x].w;
for (int i = 0; i < num; i++)
{
int index = pos+i*cSim.lincsConstraints;
int otherConstraint = cSim.pLincsConnections[index];
float4 dir2 = cSim.pLincsDistance[otherConstraint];
int2 atoms2 = cSim.pLincsAtoms[otherConstraint];
float signedMass = (atoms1.x == atoms2.x || atoms1.y == atoms2.y ? -invMass : cSim.pVelm4[atoms1.y].w);
cSim.pLincsCoupling[index] = signedMass*s*(dir1.x*dir2.x+dir1.y*dir2.y+dir1.z*dir2.z)*cSim.pLincsS[otherConstraint];
}
pos += blockDim.x * gridDim.x;
}
}
__global__ void kApplyLincsPart2_kernel(float4* atomPositions, bool addOldPosition)
{
// Correct for rotational lengthening.
unsigned int pos = threadIdx.x + blockIdx.x * blockDim.x;
while (pos < cSim.lincsConstraints)
{
int2 atoms = cSim.pLincsAtoms[pos];
float4 delta1 = atomPositions[atoms.x];
float4 delta2 = atomPositions[atoms.y];
float3 delta;
if (addOldPosition)
{
float4 oldPos1 = cSim.pOldPosq[atoms.x];
float4 oldPos2 = cSim.pOldPosq[atoms.y];
delta = make_float3((oldPos1.x-oldPos2.x)+(delta1.x-delta2.x),
(oldPos1.y-oldPos2.y)+(delta1.y-delta2.y),
(oldPos1.z-oldPos2.z)+(delta1.z-delta2.z));
}
else
{
delta = make_float3(delta1.x-delta2.x, delta1.y-delta2.y, delta1.z-delta2.z);
}
float distance = cSim.pLincsDistance[pos].w;
float p2 = 2.0f*distance*distance-(delta.x*delta.x+delta.y*delta.y+delta.z*delta.z);
p2 = (p2 < 0.0f ? 0.0f : p2);
float diff = cSim.pLincsS[pos]*(distance-sqrt(p2));
cSim.pLincsRhs1[pos] = diff;
cSim.pLincsSolution[pos] = diff;
pos += blockDim.x * gridDim.x;
}
}
static void kApplyLincs(gpuContext gpu, float4* atomPositions, bool addOldPosition)
{
kApplyLincsPart1_kernel<<<gpu->sim.blocks, gpu->sim.lincs_threads_per_block>>>(atomPositions, addOldPosition);
LAUNCHERROR("kApplyLincsPart1");
kSolveLincsMatrix_kernel<<<gpu->sim.blocks, gpu->sim.lincs_threads_per_block>>>(atomPositions);
LAUNCHERROR("kSolveLincsMatrix_kernel");
kApplyLincsPart2_kernel<<<gpu->sim.blocks, gpu->sim.lincs_threads_per_block>>>(atomPositions, addOldPosition);
LAUNCHERROR("kApplyLincsPart2");
kSolveLincsMatrix_kernel<<<gpu->sim.blocks, gpu->sim.lincs_threads_per_block>>>(atomPositions);
LAUNCHERROR("kSolveLincsMatrix_kernel");
}
void kApplyFirstLincs(gpuContext gpu)
{
// printf("kApplyFirstLincs\n");
if (gpu->sim.lincsConstraints > 0)
kApplyLincs(gpu, gpu->sim.pPosqP, true);
}
void kApplySecondLincs(gpuContext gpu)
{
// printf("kApplySecondLincs\n");
if (gpu->sim.lincsConstraints > 0)
kApplyLincs(gpu, gpu->sim.pPosq, false);
}
platforms/reference/src/ReferenceKernels.cpp
View file @ 05259f77
...@@ -38,14 +38,12 @@ ...@@ -38,14 +38,12 @@
#include "SimTKReference/ReferenceBondForce.h" #include "SimTKReference/ReferenceBondForce.h"
#include "SimTKReference/ReferenceBrownianDynamics.h" #include "SimTKReference/ReferenceBrownianDynamics.h"
#include "SimTKReference/ReferenceHarmonicBondIxn.h" #include "SimTKReference/ReferenceHarmonicBondIxn.h"
#include "SimTKReference/ReferenceLincsAlgorithm.h"
#include "SimTKReference/ReferenceLJCoulomb14.h" #include "SimTKReference/ReferenceLJCoulomb14.h"
#include "SimTKReference/ReferenceLJCoulombIxn.h" #include "SimTKReference/ReferenceLJCoulombIxn.h"
#include "SimTKReference/ReferenceProperDihedralBond.h" #include "SimTKReference/ReferenceProperDihedralBond.h"
#include "SimTKReference/ReferenceRbDihedralBond.h" #include "SimTKReference/ReferenceRbDihedralBond.h"
#include "SimTKReference/ReferenceStochasticDynamics.h" #include "SimTKReference/ReferenceStochasticDynamics.h"
#include "SimTKReference/ReferenceRigidShakeAlgorithm.h" #include "SimTKReference/ReferenceCCMAAlgorithm.h"
#include "SimTKReference/ReferenceShakeAlgorithm.h"
#include "SimTKReference/ReferenceVerletDynamics.h" #include "SimTKReference/ReferenceVerletDynamics.h"
#include "openmm/CMMotionRemover.h" #include "openmm/CMMotionRemover.h"
#include "openmm/System.h" #include "openmm/System.h"
...@@ -108,7 +106,7 @@ static void disposeRealArray(RealOpenMM** array, int size) { ...@@ -108,7 +106,7 @@ static void disposeRealArray(RealOpenMM** array, int size) {
} }
} }
static void findAnglesForShake(const System& system, vector<ReferenceRigidShakeAlgorithm::AngleInfo>& angles) { static void findAnglesForCCMA(const System& system, vector<ReferenceCCMAAlgorithm::AngleInfo>& angles) {
for (int i = 0; i < system.getNumForces(); i++) { for (int i = 0; i < system.getNumForces(); i++) {
const HarmonicAngleForce* force = dynamic_cast<const HarmonicAngleForce*>(&system.getForce(i)); const HarmonicAngleForce* force = dynamic_cast<const HarmonicAngleForce*>(&system.getForce(i));
if (force != NULL) { if (force != NULL) {
...@@ -116,7 +114,7 @@ static void findAnglesForShake(const System& system, vector<ReferenceRigidShakeA ...@@ -116,7 +114,7 @@ static void findAnglesForShake(const System& system, vector<ReferenceRigidShakeA
int atom1, atom2, atom3; int atom1, atom2, atom3;
double angle, k; double angle, k;
force->getAngleParameters(j, atom1, atom2, atom3, angle, k); force->getAngleParameters(j, atom1, atom2, atom3, angle, k);
angles.push_back(ReferenceRigidShakeAlgorithm::AngleInfo(atom1, atom2, atom3, angle)); angles.push_back(ReferenceCCMAAlgorithm::AngleInfo(atom1, atom2, atom3, angle));
} }
} }
} }
...@@ -623,9 +621,9 @@ void ReferenceIntegrateVerletStepKernel::execute(OpenMMContextImpl& context, con ...@@ -623,9 +621,9 @@ void ReferenceIntegrateVerletStepKernel::execute(OpenMMContextImpl& context, con
delete constraints; delete constraints;
} }
dynamics = new ReferenceVerletDynamics(context.getSystem().getNumParticles(), static_cast<RealOpenMM>(stepSize) ); dynamics = new ReferenceVerletDynamics(context.getSystem().getNumParticles(), static_cast<RealOpenMM>(stepSize) );
vector<ReferenceRigidShakeAlgorithm::AngleInfo> angles; vector<ReferenceCCMAAlgorithm::AngleInfo> angles;
findAnglesForShake(context.getSystem(), angles); findAnglesForCCMA(context.getSystem(), angles);
constraints = new ReferenceRigidShakeAlgorithm(context.getSystem().getNumParticles(), numConstraints, constraintIndices, constraintDistances, masses, angles, (RealOpenMM)integrator.getConstraintTolerance()); constraints = new ReferenceCCMAAlgorithm(context.getSystem().getNumParticles(), numConstraints, constraintIndices, constraintDistances, masses, angles, (RealOpenMM)integrator.getConstraintTolerance());
dynamics->setReferenceConstraintAlgorithm(constraints); dynamics->setReferenceConstraintAlgorithm(constraints);
prevStepSize = stepSize; prevStepSize = stepSize;
} }
...@@ -684,9 +682,9 @@ void ReferenceIntegrateLangevinStepKernel::execute(OpenMMContextImpl& context, c ...@@ -684,9 +682,9 @@ void ReferenceIntegrateLangevinStepKernel::execute(OpenMMContextImpl& context, c
static_cast<RealOpenMM>(stepSize), static_cast<RealOpenMM>(stepSize),
static_cast<RealOpenMM>(tau), static_cast<RealOpenMM>(tau),
static_cast<RealOpenMM>(temperature) ); static_cast<RealOpenMM>(temperature) );
vector<ReferenceRigidShakeAlgorithm::AngleInfo> angles; vector<ReferenceCCMAAlgorithm::AngleInfo> angles;
findAnglesForShake(context.getSystem(), angles); findAnglesForCCMA(context.getSystem(), angles);
constraints = new ReferenceRigidShakeAlgorithm(context.getSystem().getNumParticles(), numConstraints, constraintIndices, constraintDistances, masses, angles, (RealOpenMM)integrator.getConstraintTolerance()); constraints = new ReferenceCCMAAlgorithm(context.getSystem().getNumParticles(), numConstraints, constraintIndices, constraintDistances, masses, angles, (RealOpenMM)integrator.getConstraintTolerance());
dynamics->setReferenceConstraintAlgorithm(constraints); dynamics->setReferenceConstraintAlgorithm(constraints);
prevTemp = temperature; prevTemp = temperature;
prevFriction = friction; prevFriction = friction;
...@@ -746,9 +744,9 @@ void ReferenceIntegrateBrownianStepKernel::execute(OpenMMContextImpl& context, c ...@@ -746,9 +744,9 @@ void ReferenceIntegrateBrownianStepKernel::execute(OpenMMContextImpl& context, c
static_cast<RealOpenMM>(stepSize), static_cast<RealOpenMM>(stepSize),
static_cast<RealOpenMM>(friction), static_cast<RealOpenMM>(friction),
static_cast<RealOpenMM>(temperature) ); static_cast<RealOpenMM>(temperature) );
vector<ReferenceRigidShakeAlgorithm::AngleInfo> angles; vector<ReferenceCCMAAlgorithm::AngleInfo> angles;
findAnglesForShake(context.getSystem(), angles); findAnglesForCCMA(context.getSystem(), angles);
constraints = new ReferenceRigidShakeAlgorithm(context.getSystem().getNumParticles(), numConstraints, constraintIndices, constraintDistances, masses, angles, (RealOpenMM)integrator.getConstraintTolerance()); constraints = new ReferenceCCMAAlgorithm(context.getSystem().getNumParticles(), numConstraints, constraintIndices, constraintDistances, masses, angles, (RealOpenMM)integrator.getConstraintTolerance());
dynamics->setReferenceConstraintAlgorithm(constraints); dynamics->setReferenceConstraintAlgorithm(constraints);
prevTemp = temperature; prevTemp = temperature;
prevFriction = friction; prevFriction = friction;
......
platforms/reference/src/SimTKReference/ReferenceRigidShakeAlgorithm.cpp → platforms/reference/src/SimTKReference/ReferenceCCMAAlgorithm.cpp
View file @ 05259f77
...@@ -28,9 +28,8 @@ ...@@ -28,9 +28,8 @@
#include "../SimTKUtilities/SimTKOpenMMCommon.h" #include "../SimTKUtilities/SimTKOpenMMCommon.h"
#include "../SimTKUtilities/SimTKOpenMMUtilities.h" #include "../SimTKUtilities/SimTKOpenMMUtilities.h"
#include "../SimTKUtilities/SimTKOpenMMLog.h" #include "../SimTKUtilities/SimTKOpenMMLog.h"
#include "ReferenceRigidShakeAlgorithm.h" #include "ReferenceCCMAAlgorithm.h"
#include "ReferenceDynamics.h" #include "ReferenceDynamics.h"
#include "jama_svd.h"
#include "quern.h" #include "quern.h"
#include "openmm/Vec3.h" #include "openmm/Vec3.h"
#include <map> #include <map>
...@@ -40,22 +39,22 @@ using std::pair; ...@@ -40,22 +39,22 @@ using std::pair;
using std::vector; using std::vector;
using std::set; using std::set;
using OpenMM::Vec3; using OpenMM::Vec3;
using TNT::Array2D;
using JAMA::SVD;
/**--------------------------------------------------------------------------------------- /**---------------------------------------------------------------------------------------
ReferenceQShakeAlgorithm constructor ReferenceCCMAAlgorithm constructor
@param numberOfAtoms number of atoms @param numberOfAtoms number of atoms
@param numberOfConstraints number of constraints @param numberOfConstraints number of constraints
@param atomIndices block of atom indices @param atomIndices atom indices for contraints
@param shakeParameters Shake parameters @param distance distances for constraints
@param masses atom masses
@param angles angle force field terms
@param tolerance constraint tolerance @param tolerance constraint tolerance
--------------------------------------------------------------------------------------- */ --------------------------------------------------------------------------------------- */
ReferenceRigidShakeAlgorithm::ReferenceRigidShakeAlgorithm( int numberOfAtoms, ReferenceCCMAAlgorithm::ReferenceCCMAAlgorithm( int numberOfAtoms,
int numberOfConstraints, int numberOfConstraints,
int** atomIndices, int** atomIndices,
RealOpenMM* distance, RealOpenMM* distance,
...@@ -65,7 +64,7 @@ ReferenceRigidShakeAlgorithm::ReferenceRigidShakeAlgorithm( int numberOfAtoms, ...@@ -65,7 +64,7 @@ ReferenceRigidShakeAlgorithm::ReferenceRigidShakeAlgorithm( int numberOfAtoms,
// --------------------------------------------------------------------------------------- // ---------------------------------------------------------------------------------------
// static const char* methodName = "\nReferenceQShakeAlgorithm::ReferenceQShakeAlgorithm"; // static const char* methodName = "\nReferenceCCMAAlgorithm::ReferenceCCMAAlgorithm";
static const RealOpenMM zero = 0.0; static const RealOpenMM zero = 0.0;
static const RealOpenMM one = 1.0; static const RealOpenMM one = 1.0;
...@@ -201,181 +200,26 @@ ReferenceRigidShakeAlgorithm::ReferenceRigidShakeAlgorithm( int numberOfAtoms, ...@@ -201,181 +200,26 @@ ReferenceRigidShakeAlgorithm::ReferenceRigidShakeAlgorithm( int numberOfAtoms,
QUERN_solve_with_r(numberOfConstraints, rRowStart, rColIndex, rValue, &rhs[0], &rhs[0]); QUERN_solve_with_r(numberOfConstraints, rRowStart, rColIndex, rValue, &rhs[0], &rhs[0]);
for (int j = 0; j < numberOfConstraints; j++) { for (int j = 0; j < numberOfConstraints; j++) {
double value = rhs[j]*_distance[i]/_distance[j]; double value = rhs[j]*_distance[i]/_distance[j];
if (abs(value) > 0.02) if (FABS(value) > 0.02)
_matrix[j].push_back(pair<int, RealOpenMM>(i, (RealOpenMM) value)); _matrix[j].push_back(pair<int, RealOpenMM>(i, (RealOpenMM) value));
} }
} }
QUERN_free_result(qRowStart, qColIndex, qValue); QUERN_free_result(qRowStart, qColIndex, qValue);
QUERN_free_result(rRowStart, rColIndex, rValue); QUERN_free_result(rRowStart, rColIndex, rValue);
} }
// // Identify rigid clusters.
//
// vector<map<int, int> > atomConstraints(numberOfAtoms);
// for (int i = 0; i < numberOfConstraints; i++) {
// atomConstraints[atomIndices[i][0]][atomIndices[i][1]] = i;
// atomConstraints[atomIndices[i][1]][atomIndices[i][0]] = i;
// }
// for (int i = 0; i < numberOfAtoms; i++) {
// if (atomConstraints[i].size() < 2)
// continue;
//
// // Begin by looking for a triangle this atom is part of.
//
// set<int> atoms;
// atoms.insert(i);
// for (map<int, int>::const_iterator atom1 = atomConstraints[i].begin(); atom1 != atomConstraints[i].end() && atoms.size() == 1; ++atom1) {
// for (map<int, int>::const_iterator atom2 = atomConstraints[atom1->first].begin(); atom2 != atomConstraints[atom1->first].end(); ++atom2) {
// if (atomConstraints[i].count(atom2->first) != 0) {
// atoms.insert(atom1->first);
// atoms.insert(atom2->first);
// break;
// }
// }
// }
// if (atoms.size() == 1)
// continue;
//
// // We have three atoms that are part of a cluster, so look for other atoms we can add.
//
// bool done = false;
// while (!done) {
// done = true;
// for (set<int>::const_iterator atom1 = atoms.begin(); atom1 != atoms.end(); ++atom1) {
// for (map<int, int>::const_iterator atom2 = atomConstraints[*atom1].begin(); atom2 != atomConstraints[*atom1].end(); ++atom2) {
// if (atoms.find(atom2->first) != atoms.end())
// continue; // This atom is already in the cluster.
//
// // See if this atom is linked to three other atoms in the cluster.
//
// int linkCount = 0;
// for (map<int, int>::const_iterator atom3 = atomConstraints[atom2->first].begin(); atom3 != atomConstraints[atom2->first].end(); ++atom3)
// if (atoms.find(atom3->first) != atoms.end())
// linkCount++;
// if (linkCount > 2) {
// atoms.insert(atom2->first);
// done = false;
// }
// }
// }
// }
//
// // Record the cluster.
//
// vector<int> constraints;
// for (set<int>::const_iterator atom1 = atoms.begin(); atom1 != atoms.end(); ++atom1) {
// for (map<int, int>::const_iterator atom2 = atomConstraints[*atom1].begin(); atom2 != atomConstraints[*atom1].end(); ++atom2) {
// if (*atom1 < atom2->first && atoms.find(atom2->first) != atoms.end())
// constraints.push_back(atom2->second);
// }
// }
// _rigidClusters.push_back(constraints);
// for (set<int>::const_iterator atom1 = atoms.begin(); atom1 != atoms.end(); ++atom1) {
// for (map<int, int>::const_iterator atom2 = atomConstraints[*atom1].begin(); atom2 != atomConstraints[*atom1].end(); ++atom2)
// atomConstraints[atom2->first].erase(*atom1);
// atomConstraints[*atom1].clear();
// }
//
// // Compute the constraint coupling matrix for this cluster.
//
// const vector<int>& cluster = _rigidClusters[i];
// unsigned int size = cluster.size();
// Array2D<double> matrix(size, size);
// for (int j = 0; j < (int)size; j++) {
// for (int k = 0; k < (int)size; k++) {
// double scale;
// int atomj0 = _atomIndices[cluster[j]][0];
// int atomj1 = _atomIndices[cluster[j]][1];
// int atomk0 = _atomIndices[cluster[k]][0];
// int atomk1 = _atomIndices[cluster[k]][1];
// RealOpenMM invMass0 = one/masses[atomj0];
// RealOpenMM invMass1 = one/masses[atomj1];
// int atoma, atomb;
// if (atomj0 == atomk0) {
// atoma = atomj1;
// atomb = atomk1;
// scale = invMass0/(invMass0+invMass1);
// }
// else if (atomj1 == atomk1) {
// atoma = atomj0;
// atomb = atomk0;
// scale = invMass1/(invMass0+invMass1);
// }
// else if (atomj0 == atomk1) {
// atoma = atomj1;
// atomb = atomk0;
// scale = invMass0/(invMass0+invMass1);
// }
// else if (atomj1 == atomk0) {
// atoma = atomj0;
// atomb = atomk1;
// scale = invMass1/(invMass0+invMass1);
// }
// else {
// matrix[j][k] = 0.0;
// continue; // These constraints are not connected.
// }
//
// // Find the third constraint forming a triangle with these two.
//
// for (int m = 0; m < size; m++) {
// int other = cluster[m];
// if ((_atomIndices[other][0] == atoma && _atomIndices[other][1] == atomb) || (_atomIndices[other][0] == atomb && _atomIndices[other][1] == atoma)) {
// double d1 = _distance[cluster[j]];
// double d2 = _distance[cluster[k]];
// double d3 = _distance[other];
// matrix[j][k] = scale*(d1*d1+d2*d2-d3*d3)/(2.0*d1*d2);
// break;
// }
// }
// }
// matrix[j][j] = 1.0;
// }
//
// // Invert it using SVD.
//
// Array2D<double> u, v;
// Array1D<double> w;
// SVD<double> svd(matrix);
// svd.getU(u);
// svd.getV(v);
// svd.getSingularValues(w);
// double singularValueCutoff = 0.01*w[0];
// for (int j = 0; j < (int)size; j++)
// w[j] = (w[j] < singularValueCutoff ? 0.0 : 1.0/w[j]);
// for (int j = 0; j < (int)size; j++) {
// for (int k = 0; k < (int)size; k++) {
// matrix[j][k] = 0.0;
// for (int m = 0; m < (int)size; m++)
// matrix[j][k] += v[j][m]*w[m]*u[k][m];
// }
// }
//
// // Record the inverted matrix.
//
// _matrices.push_back(SimTKOpenMMUtilities::allocateTwoDRealOpenMMArray(constraints.size(), constraints.size(), NULL, false, 0.0, ""));
// for (int j = 0; j < (int)size; j++)
// for (int k = 0; k < (int)size; k++)
// _matrices[i][j][k] = (RealOpenMM)matrix[j][k]*_distance[cluster[k]]/_distance[cluster[j]];
// }
} }
/**--------------------------------------------------------------------------------------- /**---------------------------------------------------------------------------------------
ReferenceQShakeAlgorithm destructor ReferenceCCMAAlgorithm destructor
--------------------------------------------------------------------------------------- */ --------------------------------------------------------------------------------------- */
ReferenceRigidShakeAlgorithm::~ReferenceRigidShakeAlgorithm( ){ ReferenceCCMAAlgorithm::~ReferenceCCMAAlgorithm( ){
// --------------------------------------------------------------------------------------- // ---------------------------------------------------------------------------------------
// static const char* methodName = "\nReferenceQShakeAlgorithm::~ReferenceQShakeAlgorithm"; // static const char* methodName = "\nReferenceCCMAAlgorithm::~ReferenceCCMAAlgorithm";
// --------------------------------------------------------------------------------------- // ---------------------------------------------------------------------------------------
...@@ -398,11 +242,11 @@ ReferenceRigidShakeAlgorithm::~ReferenceRigidShakeAlgorithm( ){ ...@@ -398,11 +242,11 @@ ReferenceRigidShakeAlgorithm::~ReferenceRigidShakeAlgorithm( ){
--------------------------------------------------------------------------------------- */ --------------------------------------------------------------------------------------- */
int ReferenceRigidShakeAlgorithm::getNumberOfConstraints( void ) const { int ReferenceCCMAAlgorithm::getNumberOfConstraints( void ) const {
// --------------------------------------------------------------------------------------- // ---------------------------------------------------------------------------------------
// static const char* methodName = "\nReferenceQShakeAlgorithm::getNumberOfConstraints"; // static const char* methodName = "\nReferenceCCMAAlgorithm::getNumberOfConstraints";
// --------------------------------------------------------------------------------------- // ---------------------------------------------------------------------------------------
...@@ -417,11 +261,11 @@ int ReferenceRigidShakeAlgorithm::getNumberOfConstraints( void ) const { ...@@ -417,11 +261,11 @@ int ReferenceRigidShakeAlgorithm::getNumberOfConstraints( void ) const {
--------------------------------------------------------------------------------------- */ --------------------------------------------------------------------------------------- */
int ReferenceRigidShakeAlgorithm::getMaximumNumberOfIterations( void ) const { int ReferenceCCMAAlgorithm::getMaximumNumberOfIterations( void ) const {
// --------------------------------------------------------------------------------------- // ---------------------------------------------------------------------------------------
// static const char* methodName = "\nReferenceQShakeAlgorithm::getMaximumNumberOfIterations"; // static const char* methodName = "\nReferenceCCMAAlgorithm::getMaximumNumberOfIterations";
// --------------------------------------------------------------------------------------- // ---------------------------------------------------------------------------------------
...@@ -438,11 +282,11 @@ int ReferenceRigidShakeAlgorithm::getMaximumNumberOfIterations( void ) const { ...@@ -438,11 +282,11 @@ int ReferenceRigidShakeAlgorithm::getMaximumNumberOfIterations( void ) const {
--------------------------------------------------------------------------------------- */ --------------------------------------------------------------------------------------- */
int ReferenceRigidShakeAlgorithm::setMaximumNumberOfIterations( int maximumNumberOfIterations ){ int ReferenceCCMAAlgorithm::setMaximumNumberOfIterations( int maximumNumberOfIterations ){
// --------------------------------------------------------------------------------------- // ---------------------------------------------------------------------------------------
// static const char* methodName = "\nReferenceQShakeAlgorithm::setMaximumNumberOfIterations"; // static const char* methodName = "\nReferenceCCMAAlgorithm::setMaximumNumberOfIterations";
// --------------------------------------------------------------------------------------- // ---------------------------------------------------------------------------------------
...@@ -459,11 +303,11 @@ int ReferenceRigidShakeAlgorithm::setMaximumNumberOfIterations( int maximumNumbe ...@@ -459,11 +303,11 @@ int ReferenceRigidShakeAlgorithm::setMaximumNumberOfIterations( int maximumNumbe
--------------------------------------------------------------------------------------- */ --------------------------------------------------------------------------------------- */
RealOpenMM ReferenceRigidShakeAlgorithm::getTolerance( void ) const { RealOpenMM ReferenceCCMAAlgorithm::getTolerance( void ) const {
// --------------------------------------------------------------------------------------- // ---------------------------------------------------------------------------------------
// static const char* methodName = "\nReferenceQShakeAlgorithm::getTolerance"; // static const char* methodName = "\nReferenceCCMAAlgorithm::getTolerance";
// --------------------------------------------------------------------------------------- // ---------------------------------------------------------------------------------------
...@@ -480,12 +324,12 @@ RealOpenMM ReferenceRigidShakeAlgorithm::getTolerance( void ) const { ...@@ -480,12 +324,12 @@ RealOpenMM ReferenceRigidShakeAlgorithm::getTolerance( void ) const {
--------------------------------------------------------------------------------------- */ --------------------------------------------------------------------------------------- */
int ReferenceRigidShakeAlgorithm::setTolerance( RealOpenMM tolerance ){ int ReferenceCCMAAlgorithm::setTolerance( RealOpenMM tolerance ){
// --------------------------------------------------------------------------------------- // ---------------------------------------------------------------------------------------
// static const char* methodName = "\nReferenceQShakeAlgorithm::setTolerance"; // static const char* methodName = "\nReferenceCCMAAlgorithm::setTolerance";
// --------------------------------------------------------------------------------------- // ---------------------------------------------------------------------------------------
...@@ -497,7 +341,7 @@ int ReferenceRigidShakeAlgorithm::setTolerance( RealOpenMM tolerance ){ ...@@ -497,7 +341,7 @@ int ReferenceRigidShakeAlgorithm::setTolerance( RealOpenMM tolerance ){
/**--------------------------------------------------------------------------------------- /**---------------------------------------------------------------------------------------
Apply Shake algorithm Apply CCMA algorithm
@param numberOfAtoms number of atoms @param numberOfAtoms number of atoms
@param atomCoordinates atom coordinates @param atomCoordinates atom coordinates
...@@ -509,13 +353,13 @@ int ReferenceRigidShakeAlgorithm::setTolerance( RealOpenMM tolerance ){ ...@@ -509,13 +353,13 @@ int ReferenceRigidShakeAlgorithm::setTolerance( RealOpenMM tolerance ){
--------------------------------------------------------------------------------------- */ --------------------------------------------------------------------------------------- */
int ReferenceRigidShakeAlgorithm::apply( int numberOfAtoms, RealOpenMM** atomCoordinates, int ReferenceCCMAAlgorithm::apply( int numberOfAtoms, RealOpenMM** atomCoordinates,
RealOpenMM** atomCoordinatesP, RealOpenMM** atomCoordinatesP,
RealOpenMM* inverseMasses ){ RealOpenMM* inverseMasses ){
// --------------------------------------------------------------------------------------- // ---------------------------------------------------------------------------------------
static const char* methodName = "\nReferenceQShakeAlgorithm::apply"; static const char* methodName = "\nReferenceCCMAAlgorithm::apply";
static const RealOpenMM zero = 0.0; static const RealOpenMM zero = 0.0;
static const RealOpenMM one = 1.0; static const RealOpenMM one = 1.0;
...@@ -570,8 +414,8 @@ int ReferenceRigidShakeAlgorithm::apply( int numberOfAtoms, RealOpenMM** atomCoo ...@@ -570,8 +414,8 @@ int ReferenceRigidShakeAlgorithm::apply( int numberOfAtoms, RealOpenMM** atomCoo
int iterations = 0; int iterations = 0;
int numberConverged = 0; int numberConverged = 0;
vector<RealOpenMM> constraintForce(_numberOfConstraints); vector<RealOpenMM> constraintDelta(_numberOfConstraints);
vector<RealOpenMM> tempForce(_numberOfConstraints); vector<RealOpenMM> tempDelta(_numberOfConstraints);
while (iterations < getMaximumNumberOfIterations()) { while (iterations < getMaximumNumberOfIterations()) {
numberConverged = 0; numberConverged = 0;
for( int ii = 0; ii < _numberOfConstraints; ii++ ){ for( int ii = 0; ii < _numberOfConstraints; ii++ ){
...@@ -586,14 +430,14 @@ int ReferenceRigidShakeAlgorithm::apply( int numberOfAtoms, RealOpenMM** atomCoo ...@@ -586,14 +430,14 @@ int ReferenceRigidShakeAlgorithm::apply( int numberOfAtoms, RealOpenMM** atomCoo
RealOpenMM rp2 = DOT3( rp_ij, rp_ij ); RealOpenMM rp2 = DOT3( rp_ij, rp_ij );
RealOpenMM dist2= _distance[ii]*_distance[ii]; RealOpenMM dist2= _distance[ii]*_distance[ii];
RealOpenMM diff = dist2 - rp2; RealOpenMM diff = dist2 - rp2;
constraintForce[ii] = zero; constraintDelta[ii] = zero;
RealOpenMM rrpr = DOT3( rp_ij, r_ij[ii] ); RealOpenMM rrpr = DOT3( rp_ij, r_ij[ii] );
if( rrpr < d_ij2[ii]*epsilon6 ){ if( rrpr < d_ij2[ii]*epsilon6 ){
std::stringstream message; std::stringstream message;
message << iterations <<" "<<atomI<<" "<<atomJ<< " Error: sign of rrpr < 0?\n"; message << iterations <<" "<<atomI<<" "<<atomJ<< " Error: sign of rrpr < 0?\n";
SimTKOpenMMLog::printMessage( message ); SimTKOpenMMLog::printMessage( message );
} else { } else {
constraintForce[ii] = reducedMasses[ii]*diff/rrpr; constraintDelta[ii] = reducedMasses[ii]*diff/rrpr;
} }
if (rp2 >= lowerTol*dist2 && rp2 <= upperTol*dist2) { if (rp2 >= lowerTol*dist2 && rp2 <= upperTol*dist2) {
numberConverged++; numberConverged++;
...@@ -608,35 +452,18 @@ int ReferenceRigidShakeAlgorithm::apply( int numberOfAtoms, RealOpenMM** atomCoo ...@@ -608,35 +452,18 @@ int ReferenceRigidShakeAlgorithm::apply( int numberOfAtoms, RealOpenMM** atomCoo
RealOpenMM sum = 0.0; RealOpenMM sum = 0.0;
for (int j = 0; j < (int) _matrix[i].size(); j++) { for (int j = 0; j < (int) _matrix[i].size(); j++) {
pair<int, RealOpenMM> element = _matrix[i][j]; pair<int, RealOpenMM> element = _matrix[i][j];
sum += element.second*constraintForce[element.first]; sum += element.second*constraintDelta[element.first];
} }
tempForce[i] = sum; tempDelta[i] = sum;
} }
constraintForce = tempForce; constraintDelta = tempDelta;
} }
// for (unsigned int i = 0; i < _rigidClusters.size(); i++) {
// const vector<int>& cluster = _rigidClusters[i];
// RealOpenMM** matrix = _matrices[i];
// unsigned int size = cluster.size();
// if (size > tempForce.size())
// tempForce.resize(size);
// for (unsigned int j = 0; j < size; j++) {
// tempForce[j] = zero;
// for (unsigned int k = 0; k < size; k++)
// tempForce[j] += matrix[j][k]*constraintForce[cluster[k]];
// }
// for (unsigned int j = 0; j < size; j++)
// constraintForce[cluster[j]] = tempForce[j];
//
// }
RealOpenMM damping = one;//(RealOpenMM) (iterations < 2 ? 0.5 : 1.0);
for( int ii = 0; ii < _numberOfConstraints; ii++ ){ for( int ii = 0; ii < _numberOfConstraints; ii++ ){
int atomI = _atomIndices[ii][0]; int atomI = _atomIndices[ii][0];
int atomJ = _atomIndices[ii][1]; int atomJ = _atomIndices[ii][1];
for( int jj = 0; jj < 3; jj++ ){ for( int jj = 0; jj < 3; jj++ ){
RealOpenMM dr = constraintForce[ii]*r_ij[ii][jj]*damping; RealOpenMM dr = constraintDelta[ii]*r_ij[ii][jj];
atomCoordinatesP[atomI][jj] += inverseMasses[atomI]*dr; atomCoordinatesP[atomI][jj] += inverseMasses[atomI]*dr;
atomCoordinatesP[atomJ][jj] -= inverseMasses[atomJ]*dr; atomCoordinatesP[atomJ][jj] -= inverseMasses[atomJ]*dr;
} }
...@@ -655,7 +482,7 @@ RealOpenMM damping = one;//(RealOpenMM) (iterations < 2 ? 0.5 : 1.0); ...@@ -655,7 +482,7 @@ RealOpenMM damping = one;//(RealOpenMM) (iterations < 2 ? 0.5 : 1.0);
message << " FAILED"; message << " FAILED";
} }
message << "\n"; message << "\n";
int errors = reportShake( numberOfAtoms, atomCoordinatesP, message ); int errors = reportCCMA( numberOfAtoms, atomCoordinatesP, message );
if( !errors ){ if( !errors ){
message << "*** no errors recorded in explicit check ***"; message << "*** no errors recorded in explicit check ***";
} }
...@@ -679,12 +506,12 @@ RealOpenMM damping = one;//(RealOpenMM) (iterations < 2 ? 0.5 : 1.0); ...@@ -679,12 +506,12 @@ RealOpenMM damping = one;//(RealOpenMM) (iterations < 2 ? 0.5 : 1.0);
--------------------------------------------------------------------------------------- */ --------------------------------------------------------------------------------------- */
int ReferenceRigidShakeAlgorithm::reportShake( int numberOfAtoms, RealOpenMM** atomCoordinates, int ReferenceCCMAAlgorithm::reportCCMA( int numberOfAtoms, RealOpenMM** atomCoordinates,
std::stringstream& message ){ std::stringstream& message ){
// --------------------------------------------------------------------------------------- // ---------------------------------------------------------------------------------------
static const char* methodName = "\nReferenceQShakeAlgorithm::reportShake"; static const char* methodName = "\nReferenceCCMAAlgorithm::reportCCMA";
static const RealOpenMM zero = 0.0; static const RealOpenMM zero = 0.0;
static const RealOpenMM one = 1.0; static const RealOpenMM one = 1.0;
......
platforms/reference/src/SimTKReference/ReferenceRigidShakeAlgorithm.h → platforms/reference/src/SimTKReference/ReferenceCCMAAlgorithm.h
View file @ 05259f77
...@@ -22,8 +22,8 @@ ...@@ -22,8 +22,8 @@
* WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. * WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
*/ */
#ifndef __ReferenceQShakeAlgorithm_H__ #ifndef __ReferenceCCMAAlgorithm_H__
#define __ReferenceQShakeAlgorithm_H__ #define __ReferenceCCMAAlgorithm_H__
#include "ReferenceConstraintAlgorithm.h" #include "ReferenceConstraintAlgorithm.h"
#include <utility> #include <utility>
...@@ -32,7 +32,7 @@ ...@@ -32,7 +32,7 @@
// --------------------------------------------------------------------------------------- // ---------------------------------------------------------------------------------------
class ReferenceRigidShakeAlgorithm : public ReferenceConstraintAlgorithm { class ReferenceCCMAAlgorithm : public ReferenceConstraintAlgorithm {
protected: protected:
...@@ -48,8 +48,6 @@ class ReferenceRigidShakeAlgorithm : public ReferenceConstraintAlgorithm { ...@@ -48,8 +48,6 @@ class ReferenceRigidShakeAlgorithm : public ReferenceConstraintAlgorithm {
RealOpenMM* _distanceTolerance; RealOpenMM* _distanceTolerance;
RealOpenMM* _reducedMasses; RealOpenMM* _reducedMasses;
bool _hasInitializedMasses; bool _hasInitializedMasses;
// std::vector<std::vector<int> > _rigidClusters;
// std::vector<RealOpenMM**> _matrices;
std::vector<std::vector<std::pair<int, RealOpenMM> > > _matrix; std::vector<std::vector<std::pair<int, RealOpenMM> > > _matrix;
public: public:
...@@ -57,17 +55,19 @@ class ReferenceRigidShakeAlgorithm : public ReferenceConstraintAlgorithm { ...@@ -57,17 +55,19 @@ class ReferenceRigidShakeAlgorithm : public ReferenceConstraintAlgorithm {
/**--------------------------------------------------------------------------------------- /**---------------------------------------------------------------------------------------
ReferenceQShakeAlgorithm constructor ReferenceCCMAAlgorithm constructor
@param numberOfAtoms number of atoms @param numberOfAtoms number of atoms
@param numberOfConstraints number of constraints @param numberOfConstraints number of constraints
@param atomIndices atom indices for contraints @param atomIndices atom indices for contraints
@param distance distances for constraints @param distance distances for constraints
@param masses atom masses
@param angles angle force field terms
@param tolerance constraint tolerance @param tolerance constraint tolerance
--------------------------------------------------------------------------------------- */ --------------------------------------------------------------------------------------- */
ReferenceRigidShakeAlgorithm( int numberOfAtoms, int numberOfConstraints, int** atomIndices, RealOpenMM* distance, RealOpenMM* masses, std::vector<AngleInfo>& angles, RealOpenMM tolerance ); ReferenceCCMAAlgorithm( int numberOfAtoms, int numberOfConstraints, int** atomIndices, RealOpenMM* distance, RealOpenMM* masses, std::vector<AngleInfo>& angles, RealOpenMM tolerance );
/**--------------------------------------------------------------------------------------- /**---------------------------------------------------------------------------------------
...@@ -75,7 +75,7 @@ class ReferenceRigidShakeAlgorithm : public ReferenceConstraintAlgorithm { ...@@ -75,7 +75,7 @@ class ReferenceRigidShakeAlgorithm : public ReferenceConstraintAlgorithm {
--------------------------------------------------------------------------------------- */ --------------------------------------------------------------------------------------- */
~ReferenceRigidShakeAlgorithm( ); ~ReferenceCCMAAlgorithm( );
/**--------------------------------------------------------------------------------------- /**---------------------------------------------------------------------------------------
...@@ -132,19 +132,7 @@ class ReferenceRigidShakeAlgorithm : public ReferenceConstraintAlgorithm { ...@@ -132,19 +132,7 @@ class ReferenceRigidShakeAlgorithm : public ReferenceConstraintAlgorithm {
/**--------------------------------------------------------------------------------------- /**---------------------------------------------------------------------------------------
Print parameters Apply CCMA algorithm
@param message message
@return ReferenceQShakeAlgorithm::DefaultReturn
--------------------------------------------------------------------------------------- */
int printParameters( std::stringstream& message ) const;
/**---------------------------------------------------------------------------------------
Apply Shake algorithm
@param numberOfAtoms number of atoms @param numberOfAtoms number of atoms
@param atomCoordinates atom coordinates @param atomCoordinates atom coordinates
...@@ -171,10 +159,10 @@ class ReferenceRigidShakeAlgorithm : public ReferenceConstraintAlgorithm { ...@@ -171,10 +159,10 @@ class ReferenceRigidShakeAlgorithm : public ReferenceConstraintAlgorithm {
--------------------------------------------------------------------------------------- */ --------------------------------------------------------------------------------------- */
int reportShake( int numberOfAtoms, RealOpenMM** atomCoordinates, std::stringstream& message ); int reportCCMA( int numberOfAtoms, RealOpenMM** atomCoordinates, std::stringstream& message );
}; };
class ReferenceRigidShakeAlgorithm::AngleInfo class ReferenceCCMAAlgorithm::AngleInfo
{ {
public: public:
int atom1, atom2, atom3; int atom1, atom2, atom3;
...@@ -187,4 +175,4 @@ public: ...@@ -187,4 +175,4 @@ public:
// --------------------------------------------------------------------------------------- // ---------------------------------------------------------------------------------------
#endif // __ReferenceQShakeAlgorithm_H__ #endif // __ReferenceCCMAAlgorithm_H__
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment