Skip to content

GitLab

Commit 51b7f9e2 authored by Robert McGibbon's avatar Robert McGibbon
Browse files

Merge master

parents 85bfd73c be0387b6
Show whitespace changes
Inline Side-by-side
Showing with 315 additions and 204 deletions
platforms/cuda/src/CudaKernels.cpp
View file @ 51b7f9e2
......@@ -1458,16 +1458,24 @@ private:
class CudaCalcNonbondedForceKernel::SyncStreamPostComputation : public CudaContext::ForcePostComputation {
public:
SyncStreamPostComputation(CudaContext& cu, CUevent event, int forceGroup) : cu(cu), event(event), forceGroup(forceGroup) {
SyncStreamPostComputation(CudaContext& cu, CUevent event, CUfunction addEnergyKernel, CudaArray& pmeEnergyBuffer, int forceGroup) : cu(cu), event(event),
addEnergyKernel(addEnergyKernel), pmeEnergyBuffer(pmeEnergyBuffer), forceGroup(forceGroup) {
}
double computeForceAndEnergy(bool includeForces, bool includeEnergy, int groups) {
if ((groups&(1<<forceGroup)) != 0)
cuStreamWaitEvent(cu.getCurrentStream(), event, 0);
if (includeEnergy) {
int bufferSize = pmeEnergyBuffer.getSize();
void* args[] = {&pmeEnergyBuffer.getDevicePointer(), &cu.getEnergyBuffer().getDevicePointer(), &bufferSize};
cu.executeKernel(addEnergyKernel, args, bufferSize);
}
return 0.0;
}
private:
CudaContext& cu;
CUevent event;
CUfunction addEnergyKernel;
CudaArray& pmeEnergyBuffer;
int forceGroup;
};
......@@ -1493,6 +1501,8 @@ CudaCalcNonbondedForceKernel::~CudaCalcNonbondedForceKernel() {
delete pmeAtomRange;
if (pmeAtomGridIndex != NULL)
delete pmeAtomGridIndex;
if (pmeEnergyBuffer != NULL)
delete pmeEnergyBuffer;
if (sort != NULL)
delete sort;
if (fft != NULL)
......@@ -1681,6 +1691,9 @@ void CudaCalcNonbondedForceKernel::initialize(const System& system, const Nonbon
pmeBsplineModuliZ = new CudaArray(cu, gridSizeZ, elementSize, "pmeBsplineModuliZ");
pmeAtomRange = CudaArray::create<int>(cu, gridSizeX*gridSizeY*gridSizeZ+1, "pmeAtomRange");
pmeAtomGridIndex = CudaArray::create<int2>(cu, numParticles, "pmeAtomGridIndex");
int energyElementSize = (cu.getUseDoublePrecision() || cu.getUseMixedPrecision() ? sizeof(double) : sizeof(float));
pmeEnergyBuffer = new CudaArray(cu, cu.getNumThreadBlocks()*CudaContext::ThreadBlockSize, energyElementSize, "pmeEnergyBuffer");
cu.clearBuffer(*pmeEnergyBuffer);
sort = new CudaSort(cu, new SortTrait(), cu.getNumAtoms());
int cufftVersion;
cufftGetVersion(&cufftVersion);
......@@ -1714,7 +1727,7 @@ void CudaCalcNonbondedForceKernel::initialize(const System& system, const Nonbon
if (recipForceGroup < 0)
recipForceGroup = force.getForceGroup();
cu.addPreComputation(new SyncStreamPreComputation(cu, pmeStream, pmeSyncEvent, recipForceGroup));
cu.addPostComputation(new SyncStreamPostComputation(cu, pmeSyncEvent, recipForceGroup));
cu.addPostComputation(new SyncStreamPostComputation(cu, pmeSyncEvent, cu.getKernel(module, "addEnergy"), *pmeEnergyBuffer, recipForceGroup));
}
hasInitializedFFT = true;
......@@ -1889,7 +1902,7 @@ double CudaCalcNonbondedForceKernel::execute(ContextImpl& context, bool includeF
}
if (includeEnergy) {
void* computeEnergyArgs[] = {&reciprocalPmeGrid->getDevicePointer(), &cu.getEnergyBuffer().getDevicePointer(),
void* computeEnergyArgs[] = {&reciprocalPmeGrid->getDevicePointer(), usePmeStream ? &pmeEnergyBuffer->getDevicePointer() : &cu.getEnergyBuffer().getDevicePointer(),
&pmeBsplineModuliX->getDevicePointer(), &pmeBsplineModuliY->getDevicePointer(), &pmeBsplineModuliZ->getDevicePointer(),
cu.getPeriodicBoxSizePointer(), recipBoxVectorPointer[0], recipBoxVectorPointer[1], recipBoxVectorPointer[2]};
cu.executeKernel(pmeEvalEnergyKernel, computeEnergyArgs, cu.getNumAtoms());
......@@ -5674,7 +5687,6 @@ void CudaIntegrateVerletStepKernel::initialize(const System& system, const Verle
CUmodule module = cu.createModule(CudaKernelSources::verlet, defines, "");
kernel1 = cu.getKernel(module, "integrateVerletPart1");
kernel2 = cu.getKernel(module, "integrateVerletPart2");
prevStepSize = -1.0;
}
void CudaIntegrateVerletStepKernel::execute(ContextImpl& context, const VerletIntegrator& integrator) {
......@@ -5683,19 +5695,7 @@ void CudaIntegrateVerletStepKernel::execute(ContextImpl& context, const VerletIn
int numAtoms = cu.getNumAtoms();
int paddedNumAtoms = cu.getPaddedNumAtoms();
double dt = integrator.getStepSize();
if (dt != prevStepSize) {
if (cu.getUseDoublePrecision() || cu.getUseMixedPrecision()) {
vector<double2> stepSizeVec(1);
stepSizeVec[0] = make_double2(dt, dt);
cu.getIntegrationUtilities().getStepSize().upload(stepSizeVec);
}
else {
vector<float2> stepSizeVec(1);
stepSizeVec[0] = make_float2((float) dt, (float) dt);
cu.getIntegrationUtilities().getStepSize().upload(stepSizeVec);
}
prevStepSize = dt;
}
cu.getIntegrationUtilities().setNextStepSize(dt);
// Call the first integration kernel.
......@@ -5752,6 +5752,7 @@ void CudaIntegrateLangevinStepKernel::execute(ContextImpl& context, const Langev
double temperature = integrator.getTemperature();
double friction = integrator.getFriction();
double stepSize = integrator.getStepSize();
cu.getIntegrationUtilities().setNextStepSize(stepSize);
if (temperature != prevTemp || friction != prevFriction || stepSize != prevStepSize) {
// Calculate the integration parameters.
......@@ -5766,8 +5767,6 @@ void CudaIntegrateLangevinStepKernel::execute(ContextImpl& context, const Langev
p[1] = fscale;
p[2] = noisescale;
params->upload(p);
double2 ss = make_double2(0, stepSize);
integration.getStepSize().upload(&ss);
}
else {
vector<float> p(params->getSize());
......@@ -5775,8 +5774,6 @@ void CudaIntegrateLangevinStepKernel::execute(ContextImpl& context, const Langev
p[1] = (float) fscale;
p[2] = (float) noisescale;
params->upload(p);
float2 ss = make_float2(0, (float) stepSize);
integration.getStepSize().upload(&ss);
}
prevTemp = temperature;
prevFriction = friction;
......@@ -5929,20 +5926,13 @@ double CudaIntegrateVariableVerletStepKernel::execute(ContextImpl& context, cons
// Update the time and step count.
double dt, time;
double dt = cu.getIntegrationUtilities().getLastStepSize();
double time = cu.getTime()+dt;
if (useDouble) {
double2 stepSize;
cu.getIntegrationUtilities().getStepSize().download(&stepSize);
dt = stepSize.y;
time = cu.getTime()+dt;
if (dt == maxStepSize)
time = maxTime; // Avoid round-off error
}
else {
float2 stepSize;
cu.getIntegrationUtilities().getStepSize().download(&stepSize);
dt = stepSize.y;
time = cu.getTime()+dt;
if (dt == maxStepSizeFloat)
time = maxTime; // Avoid round-off error
}
......@@ -6023,20 +6013,13 @@ double CudaIntegrateVariableLangevinStepKernel::execute(ContextImpl& context, co
// Update the time and step count.
double dt, time;
double dt = cu.getIntegrationUtilities().getLastStepSize();
double time = cu.getTime()+dt;
if (useDouble) {
double2 stepSize;
cu.getIntegrationUtilities().getStepSize().download(&stepSize);
dt = stepSize.y;
time = cu.getTime()+dt;
if (dt == maxStepSize)
time = maxTime; // Avoid round-off error
}
else {
float2 stepSize;
cu.getIntegrationUtilities().getStepSize().download(&stepSize);
dt = stepSize.y;
time = cu.getTime()+dt;
if (dt == maxStepSizeFloat)
time = maxTime; // Avoid round-off error
}
......@@ -6139,7 +6122,6 @@ void CudaIntegrateCustomStepKernel::initialize(const System& system, const Custo
summedValue = new CudaArray(cu, 1, elementSize, "summedValue");
perDofValues = new CudaParameterSet(cu, integrator.getNumPerDofVariables(), 3*system.getNumParticles(), "perDofVariables", false, cu.getUseDoublePrecision() || cu.getUseMixedPrecision());
cu.addReorderListener(new ReorderListener(cu, *perDofValues, localPerDofValuesFloat, localPerDofValuesDouble, deviceValuesAreCurrent));
prevStepSize = -1.0;
SimTKOpenMMUtilities::setRandomNumberSeed(integrator.getRandomNumberSeed());
}
......@@ -6553,9 +6535,7 @@ void CudaIntegrateCustomStepKernel::prepareForComputation(ContextImpl& context,
}
localValuesAreCurrent = false;
double stepSize = integrator.getStepSize();
if (stepSize != prevStepSize) {
recordGlobalValue(stepSize, GlobalTarget(DT, dtVariableIndex));
}
for (int i = 0; i < (int) parameterNames.size(); i++) {
double value = context.getParameter(parameterNames[i]);
if (value != globalValuesDouble[parameterVariableIndex[i]]) {
......@@ -6760,17 +6740,10 @@ double CudaIntegrateCustomStepKernel::computeKineticEnergy(ContextImpl& context,
void CudaIntegrateCustomStepKernel::recordGlobalValue(double value, GlobalTarget target) {
switch (target.type) {
case DT:
globalValuesDouble[dtVariableIndex] = value;
if (value != globalValuesDouble[dtVariableIndex])
deviceGlobalsAreCurrent = false;
if (cu.getUseDoublePrecision() || cu.getUseMixedPrecision()) {
double size[] = {0, value};
cu.getIntegrationUtilities().getStepSize().upload(size);
}
else {
float size[] = {0, (float) value};
cu.getIntegrationUtilities().getStepSize().upload(size);
}
prevStepSize = value;
globalValuesDouble[dtVariableIndex] = value;
cu.getIntegrationUtilities().setNextStepSize(value);
break;
case VARIABLE:
case PARAMETER:
......
platforms/cuda/src/CudaPlatform.cpp
View file @ 51b7f9e2
......@@ -179,7 +179,7 @@ void CudaPlatform::contextDestroyed(ContextImpl& context) const {
}
CudaPlatform::PlatformData::PlatformData(ContextImpl* context, const System& system, const string& deviceIndexProperty, const string& blockingProperty, const string& precisionProperty,
const string& cpuPmeProperty, const string& compilerProperty, const string& tempProperty, const string& hostCompilerProperty) : context(context), removeCM(false), stepCount(0), computeForceCount(0), time(0.0) {
const string& cpuPmeProperty, const string& compilerProperty, const string& tempProperty, const string& hostCompilerProperty) : context(context), removeCM(false), stepCount(0), computeForceCount(0), time(0.0), hasInitializedContexts(false) {
bool blocking = (blockingProperty == "true");
vector<string> devices;
size_t searchPos = 0, nextPos;
......@@ -247,8 +247,11 @@ CudaPlatform::PlatformData::~PlatformData() {
}
void CudaPlatform::PlatformData::initializeContexts(const System& system) {
if (hasInitializedContexts)
return;
for (int i = 0; i < (int) contexts.size(); i++)
contexts[i]->initialize();
hasInitializedContexts = true;
}
void CudaPlatform::PlatformData::syncContexts() {
......
platforms/cuda/src/kernels/pme.cu
View file @ 51b7f9e2
......@@ -188,7 +188,7 @@ gridEvaluateEnergy(real2* __restrict__ halfcomplex_pmeGrid, mixed* __restrict__
energy += eterm*(grid.x*grid.x + grid.y*grid.y);
}
}
energyBuffer[blockIdx.x*blockDim.x+threadIdx.x] += 0.5f*energy;
energyBuffer[blockIdx.x*blockDim.x+threadIdx.x] = 0.5f*energy;
}
extern "C" __global__
......@@ -286,3 +286,9 @@ void addForces(const real4* __restrict__ forces, unsigned long long* __restrict_
forceBuffers[atom+2*PADDED_NUM_ATOMS] += static_cast<unsigned long long>((long long) (f.z*0x100000000));
}
}
extern "C" __global__
void addEnergy(const mixed* __restrict__ pmeEnergyBuffer, mixed* __restrict__ energyBuffer, int bufferSize) {
for (int i = blockIdx.x*blockDim.x+threadIdx.x; i < bufferSize; i += blockDim.x*gridDim.x)
energyBuffer[i] += pmeEnergyBuffer[i];
}
platforms/cuda/tests/TestCudaCompoundIntegrator.cpp 0 → 100644
View file @ 51b7f9e2
/* -------------------------------------------------------------------------- *
* 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) 2015 Stanford University and the Authors. *
* Authors: Peter Eastman *
* Contributors: *
* *
* Permission is hereby granted, free of charge, to any person obtaining a *
* copy of this software and associated documentation files (the "Software"), *
* to deal in the Software without restriction, including without limitation *
* the rights to use, copy, modify, merge, publish, distribute, sublicense, *
* and/or sell copies of the Software, and to permit persons to whom the *
* Software is furnished to do so, subject to the following conditions: *
* *
* The above copyright notice and this permission notice shall be included in *
* all copies or substantial portions of the Software. *
* *
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR *
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, *
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL *
* THE AUTHORS, CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, *
* DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR *
* OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE *
* USE OR OTHER DEALINGS IN THE SOFTWARE. *
* -------------------------------------------------------------------------- */
#include "CudaTests.h"
#include "TestCompoundIntegrator.h"
void runPlatformTests() {
}
platforms/opencl/include/OpenCLIntegrationUtilities.h
View file @ 51b7f9e2
......@@ -62,6 +62,14 @@ public:
OpenCLArray& getStepSize() {
return *stepSize;
}
/**
* Set the size to use for the next step.
*/
void setNextStepSize(double size);
/**
* Get the size that was used for the last step.
*/
double getLastStepSize();
/**
* Apply constraints to the atom positions.
*
......@@ -153,6 +161,7 @@ private:
int randomPos;
int lastSeed, numVsites;
bool hasInitializedPosConstraintKernels, hasInitializedVelConstraintKernels, ccmaUseDirectBuffer, hasOverlappingVsites;
mm_double2 lastStepSize;
struct ShakeCluster;
struct ConstraintOrderer;
};
......
platforms/opencl/include/OpenCLKernels.h
View file @ 51b7f9e2
......@@ -570,7 +570,7 @@ public:
OpenCLCalcNonbondedForceKernel(std::string name, const Platform& platform, OpenCLContext& cl, const System& system) : CalcNonbondedForceKernel(name, platform),
hasInitializedKernel(false), cl(cl), sigmaEpsilon(NULL), exceptionParams(NULL), cosSinSums(NULL), pmeGrid(NULL),
pmeGrid2(NULL), pmeBsplineModuliX(NULL), pmeBsplineModuliY(NULL), pmeBsplineModuliZ(NULL), pmeBsplineTheta(NULL),
pmeAtomRange(NULL), pmeAtomGridIndex(NULL), sort(NULL), fft(NULL), pmeio(NULL) {
pmeAtomRange(NULL), pmeAtomGridIndex(NULL), pmeEnergyBuffer(NULL), sort(NULL), fft(NULL), pmeio(NULL) {
}
~OpenCLCalcNonbondedForceKernel();
/**
......@@ -636,12 +636,14 @@ private:
OpenCLArray* pmeBsplineTheta;
OpenCLArray* pmeAtomRange;
OpenCLArray* pmeAtomGridIndex;
OpenCLArray* pmeEnergyBuffer;
OpenCLSort* sort;
cl::CommandQueue pmeQueue;
cl::Event pmeSyncEvent;
OpenCLFFT3D* fft;
Kernel cpuPme;
PmeIO* pmeio;
SyncQueuePostComputation* syncQueue;
cl::Kernel ewaldSumsKernel;
cl::Kernel ewaldForcesKernel;
cl::Kernel pmeGridIndexKernel;
......@@ -1103,7 +1105,6 @@ public:
double computeKineticEnergy(ContextImpl& context, const VerletIntegrator& integrator);
private:
OpenCLContext& cl;
double prevStepSize;
bool hasInitializedKernels;
cl::Kernel kernel1, kernel2;
};
......@@ -1342,7 +1343,7 @@ private:
void recordChangedParameters(ContextImpl& context);
bool evaluateCondition(int step);
OpenCLContext& cl;
double prevStepSize, energy;
double energy;
float energyFloat;
int numGlobalVariables;
bool hasInitializedKernels, deviceValuesAreCurrent, deviceGlobalsAreCurrent, modifiesParameters, keNeedsForce, hasAnyConstraints;
......
platforms/opencl/include/OpenCLPlatform.h
View file @ 51b7f9e2
......@@ -106,7 +106,7 @@ public:
ContextImpl* context;
std::vector<OpenCLContext*> contexts;
std::vector<double> contextEnergy;
bool removeCM, useCpuPme;
bool hasInitializedContexts, removeCM, useCpuPme;
int cmMotionFrequency;
int stepCount, computeForceCount;
double time;
......
platforms/opencl/src/OpenCLIntegrationUtilities.cpp
View file @ 51b7f9e2
......@@ -106,21 +106,21 @@ OpenCLIntegrationUtilities::OpenCLIntegrationUtilities(OpenCLContext& context, c
hasInitializedPosConstraintKernels(false), hasInitializedVelConstraintKernels(false), hasOverlappingVsites(false) {
// Create workspace arrays.
lastStepSize = mm_double2(0.0, 0.0);
if (context.getUseDoublePrecision() || context.getUseMixedPrecision()) {
posDelta = OpenCLArray::create<mm_double4>(context, context.getPaddedNumAtoms(), "posDelta");
vector<mm_double4> deltas(posDelta->getSize(), mm_double4(0.0, 0.0, 0.0, 0.0));
posDelta->upload(deltas);
stepSize = OpenCLArray::create<mm_double2>(context, 1, "stepSize");
vector<mm_double2> step(1, mm_double2(0.0, 0.0));
stepSize->upload(step);
stepSize->upload(&lastStepSize);
}
else {
posDelta = OpenCLArray::create<mm_float4>(context, context.getPaddedNumAtoms(), "posDelta");
vector<mm_float4> deltas(posDelta->getSize(), mm_float4(0.0f, 0.0f, 0.0f, 0.0f));
posDelta->upload(deltas);
stepSize = OpenCLArray::create<mm_float2>(context, 1, "stepSize");
vector<mm_float2> step(1, mm_float2(0.0f, 0.0f));
stepSize->upload(step);
mm_float2 lastStepSizeFloat = mm_float2(0.0f, 0.0f);
stepSize->upload(&lastStepSizeFloat);
}
// Create the time shift kernel for calculating kinetic energy.
......@@ -724,6 +724,29 @@ OpenCLIntegrationUtilities::~OpenCLIntegrationUtilities() {
delete vsiteLocalCoordsParams;
}
void OpenCLIntegrationUtilities::setNextStepSize(double size) {
if (size != lastStepSize.x || size != lastStepSize.y) {
lastStepSize = mm_double2(size, size);
if (context.getUseDoublePrecision() || context.getUseMixedPrecision())
stepSize->upload(&lastStepSize);
else {
mm_float2 lastStepSizeFloat = mm_float2((float) size, (float) size);
stepSize->upload(&lastStepSizeFloat);
}
}
}
double OpenCLIntegrationUtilities::getLastStepSize() {
if (context.getUseDoublePrecision() || context.getUseMixedPrecision())
stepSize->download(&lastStepSize);
else {
mm_float2 lastStepSizeFloat;
stepSize->download(&lastStepSizeFloat);
lastStepSize = mm_double2(lastStepSizeFloat.x, lastStepSizeFloat.y);
}
return lastStepSize.y;
}
void OpenCLIntegrationUtilities::applyConstraints(double tol) {
applyConstraints(false, tol);
}
......
platforms/opencl/src/OpenCLKernels.cpp
View file @ 51b7f9e2
......@@ -1454,7 +1454,14 @@ private:
class OpenCLCalcNonbondedForceKernel::SyncQueuePostComputation : public OpenCLContext::ForcePostComputation {
public:
SyncQueuePostComputation(OpenCLContext& cl, cl::Event& event, int forceGroup) : cl(cl), event(event), forceGroup(forceGroup) {
SyncQueuePostComputation(OpenCLContext& cl, cl::Event& event, OpenCLArray& pmeEnergyBuffer, int forceGroup) : cl(cl), event(event),
pmeEnergyBuffer(pmeEnergyBuffer), forceGroup(forceGroup) {
}
void setKernel(cl::Kernel kernel) {
addEnergyKernel = kernel;
addEnergyKernel.setArg<cl::Buffer>(0, pmeEnergyBuffer.getDeviceBuffer());
addEnergyKernel.setArg<cl::Buffer>(1, cl.getEnergyBuffer().getDeviceBuffer());
addEnergyKernel.setArg<cl_int>(2, pmeEnergyBuffer.getSize());
}
double computeForceAndEnergy(bool includeForces, bool includeEnergy, int groups) {
if ((groups&(1<<forceGroup)) != 0) {
......@@ -1463,11 +1470,15 @@ public:
event = cl::Event();
cl.getQueue().enqueueWaitForEvents(events);
}
if (includeEnergy)
cl.executeKernel(addEnergyKernel, pmeEnergyBuffer.getSize());
return 0.0;
}
private:
OpenCLContext& cl;
cl::Event& event;
cl::Kernel addEnergyKernel;
OpenCLArray& pmeEnergyBuffer;
int forceGroup;
};
......@@ -1494,6 +1505,8 @@ OpenCLCalcNonbondedForceKernel::~OpenCLCalcNonbondedForceKernel() {
delete pmeAtomRange;
if (pmeAtomGridIndex != NULL)
delete pmeAtomGridIndex;
if (pmeEnergyBuffer != NULL)
delete pmeEnergyBuffer;
if (sort != NULL)
delete sort;
if (fft != NULL)
......@@ -1663,6 +1676,9 @@ void OpenCLCalcNonbondedForceKernel::initialize(const System& system, const Nonb
pmeBsplineTheta = new OpenCLArray(cl, PmeOrder*numParticles, 4*elementSize, "pmeBsplineTheta");
pmeAtomRange = OpenCLArray::create<cl_int>(cl, gridSizeX*gridSizeY*gridSizeZ+1, "pmeAtomRange");
pmeAtomGridIndex = OpenCLArray::create<mm_int2>(cl, numParticles, "pmeAtomGridIndex");
int energyElementSize = (cl.getUseDoublePrecision() || cl.getUseMixedPrecision() ? sizeof(double) : sizeof(float));
pmeEnergyBuffer = new OpenCLArray(cl, cl.getNumThreadBlocks()*OpenCLContext::ThreadBlockSize, energyElementSize, "pmeEnergyBuffer");
cl.clearBuffer(*pmeEnergyBuffer);
sort = new OpenCLSort(cl, new SortTrait(), cl.getNumAtoms());
fft = new OpenCLFFT3D(cl, gridSizeX, gridSizeY, gridSizeZ, true);
string vendor = cl.getDevice().getInfo<CL_DEVICE_VENDOR>();
......@@ -1676,7 +1692,7 @@ void OpenCLCalcNonbondedForceKernel::initialize(const System& system, const Nonb
if (recipForceGroup < 0)
recipForceGroup = force.getForceGroup();
cl.addPreComputation(new SyncQueuePreComputation(cl, pmeQueue, recipForceGroup));
cl.addPostComputation(new SyncQueuePostComputation(cl, pmeSyncEvent, recipForceGroup));
cl.addPostComputation(syncQueue = new SyncQueuePostComputation(cl, pmeSyncEvent, *pmeEnergyBuffer, recipForceGroup));
}
// Initialize the b-spline moduli.
......@@ -1831,7 +1847,7 @@ double OpenCLCalcNonbondedForceKernel::execute(ContextImpl& context, bool includ
pmeConvolutionKernel.setArg<cl::Buffer>(2, pmeBsplineModuliY->getDeviceBuffer());
pmeConvolutionKernel.setArg<cl::Buffer>(3, pmeBsplineModuliZ->getDeviceBuffer());
pmeEvalEnergyKernel.setArg<cl::Buffer>(0, pmeGrid2->getDeviceBuffer());
pmeEvalEnergyKernel.setArg<cl::Buffer>(1, cl.getEnergyBuffer().getDeviceBuffer());
pmeEvalEnergyKernel.setArg<cl::Buffer>(1, usePmeQueue ? pmeEnergyBuffer->getDeviceBuffer() : cl.getEnergyBuffer().getDeviceBuffer());
pmeEvalEnergyKernel.setArg<cl::Buffer>(2, pmeBsplineModuliX->getDeviceBuffer());
pmeEvalEnergyKernel.setArg<cl::Buffer>(3, pmeBsplineModuliY->getDeviceBuffer());
pmeEvalEnergyKernel.setArg<cl::Buffer>(4, pmeBsplineModuliZ->getDeviceBuffer());
......@@ -1844,6 +1860,8 @@ double OpenCLCalcNonbondedForceKernel::execute(ContextImpl& context, bool includ
pmeFinishSpreadChargeKernel.setArg<cl::Buffer>(0, pmeGrid2->getDeviceBuffer());
pmeFinishSpreadChargeKernel.setArg<cl::Buffer>(1, pmeGrid->getDeviceBuffer());
}
if (usePmeQueue)
syncQueue->setKernel(cl::Kernel(program, "addEnergy"));
}
}
if (cosSinSums != NULL && includeReciprocal) {
......@@ -5866,7 +5884,6 @@ void OpenCLIntegrateVerletStepKernel::initialize(const System& system, const Ver
cl::Program program = cl.createProgram(OpenCLKernelSources::verlet, "");
kernel1 = cl::Kernel(program, "integrateVerletPart1");
kernel2 = cl::Kernel(program, "integrateVerletPart2");
prevStepSize = -1.0;
}
void OpenCLIntegrateVerletStepKernel::execute(ContextImpl& context, const VerletIntegrator& integrator) {
......@@ -5889,19 +5906,7 @@ void OpenCLIntegrateVerletStepKernel::execute(ContextImpl& context, const Verlet
kernel2.setArg<cl::Buffer>(4, cl.getVelm().getDeviceBuffer());
kernel2.setArg<cl::Buffer>(5, integration.getPosDelta().getDeviceBuffer());
}
if (dt != prevStepSize) {
if (cl.getUseDoublePrecision() || cl.getUseMixedPrecision()) {
vector<mm_double2> stepSizeVec(1);
stepSizeVec[0] = mm_double2(dt, dt);
cl.getIntegrationUtilities().getStepSize().upload(stepSizeVec);
}
else {
vector<mm_float2> stepSizeVec(1);
stepSizeVec[0] = mm_float2((cl_float) dt, (cl_float) dt);
cl.getIntegrationUtilities().getStepSize().upload(stepSizeVec);
}
prevStepSize = dt;
}
cl.getIntegrationUtilities().setNextStepSize(dt);
// Call the first integration kernel.
......@@ -5971,6 +5976,7 @@ void OpenCLIntegrateLangevinStepKernel::execute(ContextImpl& context, const Lang
double temperature = integrator.getTemperature();
double friction = integrator.getFriction();
double stepSize = integrator.getStepSize();
cl.getIntegrationUtilities().setNextStepSize(stepSize);
if (temperature != prevTemp || friction != prevFriction || stepSize != prevStepSize) {
// Calculate the integration parameters.
......@@ -5985,8 +5991,6 @@ void OpenCLIntegrateLangevinStepKernel::execute(ContextImpl& context, const Lang
p[1] = fscale;
p[2] = noisescale;
params->upload(p);
mm_double2 ss = mm_double2(0, stepSize);
integration.getStepSize().upload(&ss);
}
else {
vector<cl_float> p(params->getSize());
......@@ -5994,8 +5998,6 @@ void OpenCLIntegrateLangevinStepKernel::execute(ContextImpl& context, const Lang
p[1] = (cl_float) fscale;
p[2] = (cl_float) noisescale;
params->upload(p);
mm_float2 ss = mm_float2(0, (float) stepSize);
integration.getStepSize().upload(&ss);
}
prevTemp = temperature;
prevFriction = friction;
......@@ -6186,20 +6188,13 @@ double OpenCLIntegrateVariableVerletStepKernel::execute(ContextImpl& context, co
// Update the time and step count.
double dt, time;
double dt = cl.getIntegrationUtilities().getLastStepSize();
double time = cl.getTime()+dt;
if (useDouble) {
mm_double2 stepSize;
cl.getIntegrationUtilities().getStepSize().download(&stepSize);
dt = stepSize.y;
time = cl.getTime()+dt;
if (dt == maxStepSize)
time = maxTime; // Avoid round-off error
}
else {
mm_float2 stepSize;
cl.getIntegrationUtilities().getStepSize().download(&stepSize);
dt = stepSize.y;
time = cl.getTime()+dt;
if (dt == maxStepSizeFloat)
time = maxTime; // Avoid round-off error
}
......@@ -6300,20 +6295,13 @@ double OpenCLIntegrateVariableLangevinStepKernel::execute(ContextImpl& context,
// Update the time and step count.
double dt, time;
double dt = cl.getIntegrationUtilities().getLastStepSize();
double time = cl.getTime()+dt;
if (useDouble) {
mm_double2 stepSize;
cl.getIntegrationUtilities().getStepSize().download(&stepSize);
dt = stepSize.y;
time = cl.getTime()+dt;
if (dt == maxStepSize)
time = maxTime; // Avoid round-off error
}
else {
mm_float2 stepSize;
cl.getIntegrationUtilities().getStepSize().download(&stepSize);
dt = stepSize.y;
time = cl.getTime()+dt;
if (dt == maxStepSizeFloat)
time = maxTime; // Avoid round-off error
}
......@@ -6414,7 +6402,6 @@ void OpenCLIntegrateCustomStepKernel::initialize(const System& system, const Cus
summedValue = new OpenCLArray(cl, 1, elementSize, "summedValue");
perDofValues = new OpenCLParameterSet(cl, integrator.getNumPerDofVariables(), 3*system.getNumParticles(), "perDofVariables", false, cl.getUseDoublePrecision() || cl.getUseMixedPrecision());
cl.addReorderListener(new ReorderListener(cl, *perDofValues, localPerDofValuesFloat, localPerDofValuesDouble, deviceValuesAreCurrent));
prevStepSize = -1.0;
SimTKOpenMMUtilities::setRandomNumberSeed(integrator.getRandomNumberSeed());
}
......@@ -6825,9 +6812,7 @@ void OpenCLIntegrateCustomStepKernel::prepareForComputation(ContextImpl& context
}
localValuesAreCurrent = false;
double stepSize = integrator.getStepSize();
if (stepSize != prevStepSize) {
recordGlobalValue(stepSize, GlobalTarget(DT, dtVariableIndex));
}
for (int i = 0; i < (int) parameterNames.size(); i++) {
double value = context.getParameter(parameterNames[i]);
if (value != globalValuesDouble[parameterVariableIndex[i]]) {
......@@ -7032,17 +7017,10 @@ double OpenCLIntegrateCustomStepKernel::computeKineticEnergy(ContextImpl& contex
void OpenCLIntegrateCustomStepKernel::recordGlobalValue(double value, GlobalTarget target) {
switch (target.type) {
case DT:
globalValuesDouble[dtVariableIndex] = value;
if (value != globalValuesDouble[dtVariableIndex])
deviceGlobalsAreCurrent = false;
if (cl.getUseDoublePrecision() || cl.getUseMixedPrecision()) {
double size[] = {0, value};
cl.getIntegrationUtilities().getStepSize().upload(size);
}
else {
float size[] = {0, (float) value};
cl.getIntegrationUtilities().getStepSize().upload(size);
}
prevStepSize = value;
globalValuesDouble[dtVariableIndex] = value;
cl.getIntegrationUtilities().setNextStepSize(value);
break;
case VARIABLE:
case PARAMETER:
......
platforms/opencl/src/OpenCLPlatform.cpp
View file @ 51b7f9e2
......@@ -169,7 +169,7 @@ void OpenCLPlatform::contextDestroyed(ContextImpl& context) const {
}
OpenCLPlatform::PlatformData::PlatformData(const System& system, const string& platformPropValue, const string& deviceIndexProperty,
const string& precisionProperty, const string& cpuPmeProperty) : removeCM(false), stepCount(0), computeForceCount(0), time(0.0) {
const string& precisionProperty, const string& cpuPmeProperty) : removeCM(false), stepCount(0), computeForceCount(0), time(0.0), hasInitializedContexts(false) {
int platformIndex = -1;
if (platformPropValue.length() > 0)
stringstream(platformPropValue) >> platformIndex;
......@@ -227,8 +227,11 @@ OpenCLPlatform::PlatformData::~PlatformData() {
}
void OpenCLPlatform::PlatformData::initializeContexts(const System& system) {
if (hasInitializedContexts)
return;
for (int i = 0; i < (int) contexts.size(); i++)
contexts[i]->initialize();
hasInitializedContexts = true;
}
void OpenCLPlatform::PlatformData::syncContexts() {
......
platforms/opencl/src/kernels/pme.cl
View file @ 51b7f9e2
......@@ -362,7 +362,7 @@ __kernel void gridEvaluateEnergy(__global real2* restrict pmeGrid, __global mixe
energy += eterm*(grid.x*grid.x + grid.y*grid.y);
}
}
energyBuffer[get_global_id(0)] += 0.5f*energy;
energyBuffer[get_global_id(0)] = 0.5f*energy;
}
__kernel void gridInterpolateForce(__global const real4* restrict posq, __global real4* restrict forceBuffers, __global const real* restrict pmeGrid,
......@@ -445,3 +445,8 @@ __kernel void addForces(__global const real4* restrict forces, __global real4* r
for (int atom = get_global_id(0); atom < NUM_ATOMS; atom += get_global_size(0))
forceBuffers[atom] += forces[atom];
}
__kernel void addEnergy(__global const mixed* restrict pmeEnergyBuffer, __global mixed* restrict energyBuffer, int bufferSize) {
for (int i = get_global_id(0); i < bufferSize; i += get_global_size(0))
energyBuffer[i] += pmeEnergyBuffer[i];
}
platforms/opencl/tests/TestOpenCLCompoundIntegrator.cpp 0 → 100644
View file @ 51b7f9e2
/* -------------------------------------------------------------------------- *
* 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) 2015 Stanford University and the Authors. *
* Authors: Peter Eastman *
* Contributors: *
* *
* Permission is hereby granted, free of charge, to any person obtaining a *
* copy of this software and associated documentation files (the "Software"), *
* to deal in the Software without restriction, including without limitation *
* the rights to use, copy, modify, merge, publish, distribute, sublicense, *
* and/or sell copies of the Software, and to permit persons to whom the *
* Software is furnished to do so, subject to the following conditions: *
* *
* The above copyright notice and this permission notice shall be included in *
* all copies or substantial portions of the Software. *
* *
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR *
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, *
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL *
* THE AUTHORS, CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, *
* DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR *
* OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE *
* USE OR OTHER DEALINGS IN THE SOFTWARE. *
* -------------------------------------------------------------------------- */
#include "OpenCLTests.h"
#include "TestCompoundIntegrator.h"
void runPlatformTests() {
}
platforms/reference/tests/TestReferenceCompoundIntegrator.cpp 0 → 100644
View file @ 51b7f9e2
/* -------------------------------------------------------------------------- *
* 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) 2015 Stanford University and the Authors. *
* Authors: Peter Eastman *
* Contributors: *
* *
* Permission is hereby granted, free of charge, to any person obtaining a *
* copy of this software and associated documentation files (the "Software"), *
* to deal in the Software without restriction, including without limitation *
* the rights to use, copy, modify, merge, publish, distribute, sublicense, *
* and/or sell copies of the Software, and to permit persons to whom the *
* Software is furnished to do so, subject to the following conditions: *
* *
* The above copyright notice and this permission notice shall be included in *
* all copies or substantial portions of the Software. *
* *
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR *
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, *
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL *
* THE AUTHORS, CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, *
* DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR *
* OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE *
* USE OR OTHER DEALINGS IN THE SOFTWARE. *
* -------------------------------------------------------------------------- */
#include "ReferenceTests.h"
#include "TestCompoundIntegrator.h"
void runPlatformTests() {
}
plugins/amoeba/openmmapi/include/openmm/AmoebaAngleForce.h
View file @ 51b7f9e2
......@@ -128,7 +128,7 @@ public:
* @param particle2 the index of the second particle connected by the angle
* @param particle3 the index of the third particle connected by the angle
* @param length the equilibrium angle, measured in degrees
* @param quadratic k the quadratic force constant for the angle, measured in kJ/mol/radian^2
* @param quadraticK the quadratic force constant for the angle, measured in kJ/mol/radian^2
* @return the index of the angle that was added
*/
int addAngle(int particle1, int particle2, int particle3, double length, double quadraticK);
......@@ -137,11 +137,11 @@ public:
* Get the force field parameters for an angle term.
*
* @param index the index of the angle for which to get parameters
* @param particle1 the index of the first particle connected by the angle
* @param particle2 the index of the second particle connected by the angle
* @param particle3 the index of the third particle connected by the angle
* @param length the equilibrium angle, measured in degrees
* @param quadratic k the quadratic force constant for the angle, measured in kJ/mol/radian^2
* @param[out] particle1 the index of the first particle connected by the angle
* @param[out] particle2 the index of the second particle connected by the angle
* @param[out] particle3 the index of the third particle connected by the angle
* @param[out] length the equilibrium angle, measured in degrees
* @param[out] quadraticK the quadratic force constant for the angle, measured in kJ/mol/radian^2
*/
void getAngleParameters(int index, int& particle1, int& particle2, int& particle3, double& length, double& quadraticK) const;
......@@ -153,7 +153,7 @@ public:
* @param particle2 the index of the second particle connected by the angle
* @param particle3 the index of the third particle connected by the angle
* @param length the equilibrium angle, measured in degrees
* @param quadratic k the quadratic force constant for the angle, measured in kJ/mol/radian^2
* @param quadraticK the quadratic force constant for the angle, measured in kJ/mol/radian^2
*/
void setAngleParameters(int index, int particle1, int particle2, int particle3, double length, double quadraticK);
/**
......
plugins/amoeba/openmmapi/include/openmm/AmoebaBondForce.h
View file @ 51b7f9e2
......@@ -100,7 +100,7 @@ public:
* @param particle1 the index of the first particle connected by the bond
* @param particle2 the index of the second particle connected by the bond
* @param length the equilibrium length of the bond, measured in nm
* @param k the quadratic force constant for the bond
* @param quadraticK the quadratic force constant for the bond
* @return the index of the bond that was added
*/
......@@ -110,10 +110,10 @@ public:
* Get the force field parameters for a bond term.
*
* @param index the index of the bond for which to get parameters
* @param particle1 the index of the first particle connected by the bond
* @param particle2 the index of the second particle connected by the bond
* @param length the equilibrium length of the bond, measured in nm
* @param quadratic k the quadratic force constant for the bond
* @param[out] particle1 the index of the first particle connected by the bond
* @param[out] particle2 the index of the second particle connected by the bond
* @param[out] length the equilibrium length of the bond, measured in nm
* @param[out] quadraticK the quadratic force constant for the bond
*/
void getBondParameters(int index, int& particle1, int& particle2, double& length, double& quadraticK) const;
......@@ -125,7 +125,7 @@ public:
* @param particle1 the index of the first particle connected by the bond
* @param particle2 the index of the second particle connected by the bond
* @param length the equilibrium length of the bond, measured in nm
* @param k the quadratic force constant for the bond
* @param quadraticK the quadratic force constant for the bond
*/
void setBondParameters(int index, int particle1, int particle2, double length, double quadraticK);
/**
......
plugins/amoeba/openmmapi/include/openmm/AmoebaGeneralizedKirkwoodForce.h
View file @ 51b7f9e2
......@@ -80,9 +80,9 @@ public:
* Get the force field parameters for a particle.
*
* @param index the index of the particle for which to get parameters
* @param charge the charge of the particle, measured in units of the proton charge
* @param radius the atomic radius of the particle, measured in nm
* @param scalingFactor the scaling factor for the particle
* @param[out] charge the charge of the particle, measured in units of the proton charge
* @param[out] radius the atomic radius of the particle, measured in nm
* @param[out] scalingFactor the scaling factor for the particle
*/
void getParticleParameters(int index, double& charge, double& radius, double& scalingFactor) const;
......
plugins/amoeba/openmmapi/include/openmm/AmoebaInPlaneAngleForce.h
View file @ 51b7f9e2
......@@ -129,7 +129,7 @@ public:
* @param particle3 the index of the third particle connected by the angle
* @param particle4 the index of the fourth particle connected by the angle
* @param length the equilibrium angle, measured in radians
* @param quadratic k the quadratic force constant for the angle measured in kJ/mol/radian^2
* @param quadraticK the quadratic force constant for the angle measured in kJ/mol/radian^2
* @return the index of the angle that was added
*/
int addAngle(int particle1, int particle2, int particle3, int particle4, double length,
......@@ -139,12 +139,12 @@ public:
* Get the force field parameters for an angle term.
*
* @param index the index of the angle for which to get parameters
* @param particle1 the index of the first particle connected by the angle
* @param particle2 the index of the second particle connected by the angle
* @param particle3 the index of the third particle connected by the angle
* @param particle4 the index of the fourth particle connected by the angle
* @param length the equilibrium angle, measured in radians
* @param quadratic k the quadratic force constant for the angle measured in kJ/mol/radian^2
* @param[out] particle1 the index of the first particle connected by the angle
* @param[out] particle2 the index of the second particle connected by the angle
* @param[out] particle3 the index of the third particle connected by the angle
* @param[out] particle4 the index of the fourth particle connected by the angle
* @param[out] length the equilibrium angle, measured in radians
* @param[out] quadraticK the quadratic force constant for the angle measured in kJ/mol/radian^2
*/
void getAngleParameters(int index, int& particle1, int& particle2, int& particle3, int& particle4, double& length,
double& quadraticK) const;
......@@ -158,7 +158,7 @@ public:
* @param particle3 the index of the third particle connected by the angle
* @param particle4 the index of the fourth particle connected by the angle
* @param length the equilibrium angle, measured in radians
* @param quadratic k the quadratic force constant for the angle, measured in kJ/mol/radian^2
* @param quadraticK the quadratic force constant for the angle, measured in kJ/mol/radian^2
*/
void setAngleParameters(int index, int particle1, int particle2, int particle3, int particle4, double length, double quadraticK);
/**
......
plugins/amoeba/openmmapi/include/openmm/AmoebaMultipoleForce.h
View file @ 51b7f9e2
......@@ -148,7 +148,7 @@ public:
* Set the Ewald alpha parameter. If this is 0 (the default), a value is chosen automatically
* based on the Ewald error tolerance.
*
* @param Ewald alpha parameter
* @param aewald alpha parameter
*/
void setAEwald(double aewald);
......@@ -171,7 +171,7 @@ public:
* Set the PME grid dimensions. If Ewald alpha is 0 (the default), this is ignored and grid dimensions
* are chosen automatically based on the Ewald error tolerance.
*
* @param the PME grid dimensions
* @param gridDimension the PME grid dimensions
*/
void setPmeGridDimensions(const std::vector<int>& gridDimension);
......@@ -182,10 +182,10 @@ public:
* See the manual for details.
*
* @param context the Context for which to get the parameters
* @param alpha the separation parameter
* @param nx the number of grid points along the X axis
* @param ny the number of grid points along the Y axis
* @param nz the number of grid points along the Z axis
* @param[out] alpha the separation parameter
* @param[out] nx the number of grid points along the X axis
* @param[out] ny the number of grid points along the Y axis
* @param[out] nz the number of grid points along the Z axis
*/
void getPMEParametersInContext(const Context& context, double& alpha, int& nx, int& ny, int& nz) const;
......@@ -212,16 +212,16 @@ public:
* Get the multipole parameters for a particle.
*
* @param index the index of the atom for which to get parameters
* @param charge the particle's charge
* @param molecularDipole the particle's molecular dipole (vector of size 3)
* @param molecularQuadrupole the particle's molecular quadrupole (vector of size 9)
* @param axisType the particle's axis type
* @param multipoleAtomZ index of first atom used in constructing lab<->molecular frames
* @param multipoleAtomX index of second atom used in constructing lab<->molecular frames
* @param multipoleAtomY index of second atom used in constructing lab<->molecular frames
* @param thole Thole parameter
* @param dampingFactor dampingFactor parameter
* @param polarity polarity parameter
* @param[out] charge the particle's charge
* @param[out] molecularDipole the particle's molecular dipole (vector of size 3)
* @param[out] molecularQuadrupole the particle's molecular quadrupole (vector of size 9)
* @param[out] axisType the particle's axis type
* @param[out] multipoleAtomZ index of first atom used in constructing lab<->molecular frames
* @param[out] multipoleAtomX index of second atom used in constructing lab<->molecular frames
* @param[out] multipoleAtomY index of second atom used in constructing lab<->molecular frames
* @param[out] thole Thole parameter
* @param[out] dampingFactor dampingFactor parameter
* @param[out] polarity polarity parameter
*/
void getMultipoleParameters(int index, double& charge, std::vector<double>& molecularDipole, std::vector<double>& molecularQuadrupole,
int& axisType, int& multipoleAtomZ, int& multipoleAtomX, int& multipoleAtomY, double& thole, double& dampingFactor, double& polarity) const;
......@@ -237,6 +237,8 @@ public:
* @param multipoleAtomZ index of first atom used in constructing lab<->molecular frames
* @param multipoleAtomX index of second atom used in constructing lab<->molecular frames
* @param multipoleAtomY index of second atom used in constructing lab<->molecular frames
* @param thole thole parameter
* @param dampingFactor damping factor parameter
* @param polarity polarity parameter
*/
void setMultipoleParameters(int index, double charge, const std::vector<double>& molecularDipole, const std::vector<double>& molecularQuadrupole,
......@@ -256,7 +258,7 @@ public:
*
* @param index the index of the atom for which to set parameters
* @param typeId CovalentTypes type
* @param covalentAtoms output vector of covalent atoms associated w/ the specfied CovalentType
* @param[out] covalentAtoms output vector of covalent atoms associated w/ the specfied CovalentType
*/
void getCovalentMap(int index, CovalentType typeId, std::vector<int>& covalentAtoms) const;
......@@ -264,7 +266,7 @@ public:
* Get the CovalentMap for an atom
*
* @param index the index of the atom for which to set parameters
* @param covalentLists output vector of covalent lists of atoms
* @param[out] covalentLists output vector of covalent lists of atoms
*/
void getCovalentMaps(int index, std::vector < std::vector<int> >& covalentLists) const;
......@@ -278,7 +280,7 @@ public:
/**
* Set the max number of iterations to be used in calculating the mutual induced dipoles
*
* @param max number of iterations
* @param inputMutualInducedMaxIterations number of iterations
*/
void setMutualInducedMaxIterations(int inputMutualInducedMaxIterations);
......@@ -292,7 +294,7 @@ public:
/**
* Set the target epsilon to be used to test for convergence of iterative method used in calculating the mutual induced dipoles
*
* @param target epsilon
* @param inputMutualInducedTargetEpsilon target epsilon
*/
void setMutualInducedTargetEpsilon(double inputMutualInducedTargetEpsilon);
......@@ -319,7 +321,7 @@ public:
* Get the induced dipole moments of all particles.
*
* @param context the Context for which to get the induced dipoles
* @param dipoles the induced dipole moment of particle i is stored into the i'th element
* @param[out] dipoles the induced dipole moment of particle i is stored into the i'th element
*/
void getInducedDipoles(Context& context, std::vector<Vec3>& dipoles);
......@@ -328,7 +330,7 @@ public:
*
* @param inputGrid input grid points over which the potential is to be evaluated
* @param context context
* @param outputElectrostaticPotential output potential
* @param[out] outputElectrostaticPotential output potential
*/
void getElectrostaticPotential(const std::vector< Vec3 >& inputGrid,
......@@ -344,7 +346,7 @@ public:
* the quadrupole moment is still undefined and should be ignored.
*
* @param context context
* @param outputMultipoleMoments (charge,
* @param[out] outputMultipoleMoments (charge,
dipole_x, dipole_y, dipole_z,
quadrupole_xx, quadrupole_xy, quadrupole_xz,
quadrupole_yx, quadrupole_yy, quadrupole_yz,
......@@ -415,8 +417,8 @@ public:
MultipoleInfo(double charge, const std::vector<double>& inputMolecularDipole, const std::vector<double>& inputMolecularQuadrupole,
int axisType, int multipoleAtomZ, int multipoleAtomX, int multipoleAtomY, double thole, double dampingFactor, double polarity) :
charge(charge), axisType(axisType), multipoleAtomZ(multipoleAtomZ), multipoleAtomX(multipoleAtomX), multipoleAtomY(multipoleAtomY),
thole(thole), dampingFactor(dampingFactor), polarity(polarity) {
axisType(axisType), multipoleAtomZ(multipoleAtomZ), multipoleAtomX(multipoleAtomX), multipoleAtomY(multipoleAtomY),
charge(charge), thole(thole), dampingFactor(dampingFactor), polarity(polarity) {
covalentInfo.resize(CovalentEnd);
......
plugins/amoeba/openmmapi/include/openmm/AmoebaOutOfPlaneBendForce.h
View file @ 51b7f9e2
......@@ -134,11 +134,11 @@ public:
* Get the force field parameters for an out-of-plane bend term.
*
* @param index the index of the outOfPlaneBend for which to get parameters
* @param particle1 the index of the first particle connected by the outOfPlaneBend
* @param particle2 the index of the second particle connected by the outOfPlaneBend
* @param particle3 the index of the third particle connected by the outOfPlaneBend
* @param particle4 the index of the fourth particle connected by the outOfPlaneBend
* @param k the force constant for the out-of-plane bend
* @param[out] particle1 the index of the first particle connected by the outOfPlaneBend
* @param[out] particle2 the index of the second particle connected by the outOfPlaneBend
* @param[out] particle3 the index of the third particle connected by the outOfPlaneBend
* @param[out] particle4 the index of the fourth particle connected by the outOfPlaneBend
* @param[out] k the force constant for the out-of-plane bend
*/
void getOutOfPlaneBendParameters(int index, int& particle1, int& particle2, int& particle3, int& particle4, double& k) const;
......
plugins/amoeba/openmmapi/include/openmm/AmoebaPiTorsionForce.h
View file @ 51b7f9e2
......@@ -80,13 +80,13 @@ public:
* Get the force field parameters for a torsion term.
*
* @param index the index of the torsion for which to get parameters
* @param particle1 the index of the first particle connected by the torsion
* @param particle2 the index of the second particle connected by the torsion
* @param particle3 the index of the third particle connected by the torsion
* @param particle4 the index of the fourth particle connected by the torsion
* @param particle5 the index of the fifth particle connected by the torsion
* @param particle6 the index of the sixth particle connected by the torsion
* @param k the force constant for the torsion
* @param[out] particle1 the index of the first particle connected by the torsion
* @param[out] particle2 the index of the second particle connected by the torsion
* @param[out] particle3 the index of the third particle connected by the torsion
* @param[out] particle4 the index of the fourth particle connected by the torsion
* @param[out] particle5 the index of the fifth particle connected by the torsion
* @param[out] particle6 the index of the sixth particle connected by the torsion
* @param[out] k the force constant for the torsion
*/
void getPiTorsionParameters(int index, int& particle1, int& particle2, int& particle3, int& particle4, int& particle5, int& particle6, double& k) const;
......
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