Skip to content

GitLab

Commit ecbbf442 authored by Peter Eastman's avatar Peter Eastman
Browse files

Continuing to implement new CUDA platform: constraints, LangevinIntegrator,... 

Continuing to implement new CUDA platform: constraints, LangevinIntegrator, BrownianIntegrator, VariableLangevinIntegrator, VariableVerletIntegrator
parent 03cc3523
Expand all Show whitespace changes
Inline Side-by-side
Showing with 1663 additions and 738 deletions
platforms/cuda2/src/CudaArray.cpp
View file @ ecbbf442
......@@ -32,8 +32,8 @@
using namespace OpenMM;
CudaArray::CudaArray(int size, int elementSize, const std::string& name) :
size(size), elementSize(elementSize), name(name), ownsMemory(true) {
CudaArray::CudaArray(CudaContext& context, int size, int elementSize, const std::string& name) :
context(context), size(size), elementSize(elementSize), name(name), ownsMemory(true) {
CUresult result = cuMemAlloc(&pointer, size*elementSize);
if (result != CUDA_SUCCESS) {
std::stringstream str;
......@@ -43,7 +43,7 @@ CudaArray::CudaArray(int size, int elementSize, const std::string& name) :
}
CudaArray::~CudaArray() {
if (ownsMemory) {
if (ownsMemory && context.getContextIsValid()) {
CUresult result = cuMemFree(pointer);
if (result != CUDA_SUCCESS) {
std::stringstream str;
......
platforms/cuda2/src/CudaArray.h
View file @ ecbbf442
......@@ -35,6 +35,8 @@
namespace OpenMM {
class CudaContext;
/**
* This class encapsulates a block of CUDA device memory. It provides a simplified API
* for working with it and for copying data to and from device memory.
......@@ -46,21 +48,23 @@ public:
* Create a CudaArray object. The object is allocated on the heap with the "new" operator.
* The template argument is the data type of each array element.
*
* @param context the context for which to create the array
* @param size the number of elements in the array
* @param name the name of the array
*/
template <class T>
static CudaArray* create(int size, const std::string& name) {
return new CudaArray(size, sizeof(T), name);
static CudaArray* create(CudaContext& context, int size, const std::string& name) {
return new CudaArray(context, size, sizeof(T), name);
}
/**
* Create a CudaArray object.
*
* @param context the context for which to create the array
* @param size the number of elements in the array
* @param elementSize the size of each element in bytes
* @param name the name of the array
*/
CudaArray(int size, int elementSize, const std::string& name);
CudaArray(CudaContext& context, int size, int elementSize, const std::string& name);
~CudaArray();
/**
* Get the number of elements in the array.
......@@ -123,6 +127,7 @@ public:
*/
void download(void* data, bool blocking = true) const;
private:
CudaContext& context;
CUdeviceptr pointer;
int size, elementSize;
bool ownsMemory;
......
platforms/cuda2/src/CudaBondedUtilities.cpp
View file @ ecbbf442
......@@ -86,7 +86,7 @@ void CudaBondedUtilities::initialize(const System& system) {
for (int atom = 0; atom < width; atom++)
indexVec[bond*width+atom] = forceAtoms[i][bond][startAtom+atom];
}
CudaArray* indices = new CudaArray(numBonds, 4*width, "bondedIndices");
CudaArray* indices = new CudaArray(context, numBonds, 4*width, "bondedIndices");
indices->upload(&indexVec[0]);
atomIndices[i].push_back(indices);
startAtom += width;
......
platforms/cuda2/src/CudaContext.cpp
View file @ ecbbf442
......@@ -141,23 +141,23 @@ CudaContext::CudaContext(const System& system, int deviceIndex, bool useBlocking
nonbonded = new CudaNonbondedUtilities(*this);
int numEnergyBuffers = max(numThreadBlocks*ThreadBlockSize, nonbonded->getNumEnergyBuffers());
if (useDoublePrecision) {
posq = CudaArray::create<double4>(paddedNumAtoms, "posq");
velm = CudaArray::create<double4>(paddedNumAtoms, "velm");
posq = CudaArray::create<double4>(*this, paddedNumAtoms, "posq");
velm = CudaArray::create<double4>(*this, paddedNumAtoms, "velm");
compilationDefines["USE_DOUBLE_PRECISION"] = "1";
compilationDefines["make_real2"] = "make_double2";
compilationDefines["make_real3"] = "make_double3";
compilationDefines["make_real4"] = "make_double4";
energyBuffer = CudaArray::create<double>(numEnergyBuffers, "energyBuffer");
energyBuffer = CudaArray::create<double>(*this, numEnergyBuffers, "energyBuffer");
int pinnedBufferSize = max(paddedNumAtoms*4, numEnergyBuffers);
CHECK_RESULT(cuMemHostAlloc(&pinnedBuffer, pinnedBufferSize*sizeof(double), 0));
}
else {
posq = CudaArray::create<float4>(paddedNumAtoms, "posq");
velm = CudaArray::create<float4>(paddedNumAtoms, "velm");
posq = CudaArray::create<float4>(*this, paddedNumAtoms, "posq");
velm = CudaArray::create<float4>(*this, paddedNumAtoms, "velm");
compilationDefines["make_real2"] = "make_float2";
compilationDefines["make_real3"] = "make_float3";
compilationDefines["make_real4"] = "make_float4";
energyBuffer = CudaArray::create<float>(numEnergyBuffers, "energyBuffer");
energyBuffer = CudaArray::create<float>(*this, numEnergyBuffers, "energyBuffer");
int pinnedBufferSize = max(paddedNumAtoms*6, numEnergyBuffers);
CHECK_RESULT(cuMemHostAlloc(&pinnedBuffer, pinnedBufferSize*sizeof(float), 0));
}
......@@ -198,7 +198,7 @@ CudaContext::CudaContext(const System& system, int deviceIndex, bool useBlocking
}
CudaContext::~CudaContext() {
cuCtxSetCurrent(context);
setAsCurrent();
for (int i = 0; i < (int) forces.size(); i++)
delete forces[i];
for (int i = 0; i < (int) reorderListeners.size(); i++)
......@@ -226,6 +226,7 @@ CudaContext::~CudaContext() {
string errorMessage = "Error deleting Context";
if (contextIsValid)
CHECK_RESULT(cuCtxDestroy(context));
contextIsValid = false;
}
void CudaContext::initialize() {
......@@ -240,10 +241,10 @@ void CudaContext::initialize() {
}
velm->upload(pinnedBuffer);
bonded->initialize(system);
force = CudaArray::create<long long>(paddedNumAtoms*3, "force");
force = CudaArray::create<long long>(*this, paddedNumAtoms*3, "force");
addAutoclearBuffer(force->getDevicePointer(), force->getSize()*force->getElementSize());
addAutoclearBuffer(energyBuffer->getDevicePointer(), energyBuffer->getSize()*energyBuffer->getElementSize());
atomIndexDevice = CudaArray::create<int>(paddedNumAtoms, "atomIndex");
atomIndexDevice = CudaArray::create<int>(*this, paddedNumAtoms, "atomIndex");
atomIndex.resize(paddedNumAtoms);
for (int i = 0; i < paddedNumAtoms; ++i)
atomIndex[i] = i;
......@@ -257,6 +258,11 @@ void CudaContext::addForce(CudaForceInfo* force) {
forces.push_back(force);
}
void CudaContext::setAsCurrent() {
if (contextIsValid)
cuCtxSetCurrent(context);
}
string CudaContext::replaceStrings(const string& input, const std::map<std::string, std::string>& replacements) const {
string result = input;
for (map<string, string>::const_iterator iter = replacements.begin(); iter != replacements.end(); iter++) {
......
platforms/cuda2/src/CudaContext.h
View file @ ecbbf442
......@@ -88,6 +88,17 @@ public:
CUcontext getContext() {
return context;
}
/**
* Get whether the CUcontext associated with this object is currently a valid contex.
*/
bool getContextIsValid() const {
return contextIsValid;
}
/**
* Set the CUcontext associated with this object to be the current context. If the context is not
* valid, this returns without doing anything.
*/
void setAsCurrent();
/**
* Get the CUdevice associated with this object.
*/
......
platforms/cuda2/src/CudaIntegrationUtilities.h
View file @ ecbbf442
......@@ -135,6 +135,7 @@ private:
CudaArray* ccmaDelta2;
CudaArray* ccmaConverged;
int* ccmaConvergedMemory;
CUevent ccmaEvent;
CudaArray* vsite2AvgAtoms;
CudaArray* vsite2AvgWeights;
CudaArray* vsite3AvgAtoms;
......@@ -143,7 +144,6 @@ private:
CudaArray* vsiteOutOfPlaneWeights;
int randomPos;
int lastSeed, numVsites;
bool hasInitializedPosConstraintKernels, hasInitializedVelConstraintKernels;
struct ShakeCluster;
struct ConstraintOrderer;
};
......
platforms/cuda2/src/CudaKernelFactory.cpp
View file @ ecbbf442
......@@ -108,14 +108,14 @@ KernelImpl* CudaKernelFactory::createKernelImpl(std::string name, const Platform
return new CudaCalcCustomCompoundBondForceKernel(name, platform, cu, context.getSystem());
if (name == IntegrateVerletStepKernel::Name())
return new CudaIntegrateVerletStepKernel(name, platform, cu);
// if (name == IntegrateLangevinStepKernel::Name())
// return new CudaIntegrateLangevinStepKernel(name, platform, cu);
// if (name == IntegrateBrownianStepKernel::Name())
// return new CudaIntegrateBrownianStepKernel(name, platform, cu);
// if (name == IntegrateVariableVerletStepKernel::Name())
// return new CudaIntegrateVariableVerletStepKernel(name, platform, cu);
// if (name == IntegrateVariableLangevinStepKernel::Name())
// return new CudaIntegrateVariableLangevinStepKernel(name, platform, cu);
if (name == IntegrateLangevinStepKernel::Name())
return new CudaIntegrateLangevinStepKernel(name, platform, cu);
if (name == IntegrateBrownianStepKernel::Name())
return new CudaIntegrateBrownianStepKernel(name, platform, cu);
if (name == IntegrateVariableVerletStepKernel::Name())
return new CudaIntegrateVariableVerletStepKernel(name, platform, cu);
if (name == IntegrateVariableLangevinStepKernel::Name())
return new CudaIntegrateVariableLangevinStepKernel(name, platform, cu);
// if (name == IntegrateCustomStepKernel::Name())
// return new CudaIntegrateCustomStepKernel(name, platform, cu);
// if (name == ApplyAndersenThermostatKernel::Name())
......
platforms/cuda2/src/CudaKernels.cpp
View file @ ecbbf442
This diff is collapsed.
platforms/cuda2/src/CudaKernels.h
View file @ ecbbf442
......@@ -942,133 +942,126 @@ private:
CUfunction kernel1, kernel2;
};
///**
// * This kernel is invoked by LangevinIntegrator to take one time step.
// */
//class CudaIntegrateLangevinStepKernel : public IntegrateLangevinStepKernel {
//public:
// CudaIntegrateLangevinStepKernel(std::string name, const Platform& platform, CudaContext& cu) : IntegrateLangevinStepKernel(name, platform), cu(cu),
// hasInitializedKernels(false), params(NULL) {
// }
// ~CudaIntegrateLangevinStepKernel();
// /**
// * Initialize the kernel, setting up the particle masses.
// *
// * @param system the System this kernel will be applied to
// * @param integrator the LangevinIntegrator this kernel will be used for
// */
// void initialize(const System& system, const LangevinIntegrator& integrator);
// /**
// * Execute the kernel.
// *
// * @param context the context in which to execute this kernel
// * @param integrator the LangevinIntegrator this kernel is being used for
// */
// void execute(ContextImpl& context, const LangevinIntegrator& integrator);
//private:
// CudaContext& cu;
// double prevTemp, prevFriction, prevStepSize;
// bool hasInitializedKernels;
// CudaArray<cl_float>* params;
// CUfunction kernel1, kernel2;
//};
//
///**
// * This kernel is invoked by BrownianIntegrator to take one time step.
// */
//class CudaIntegrateBrownianStepKernel : public IntegrateBrownianStepKernel {
//public:
// CudaIntegrateBrownianStepKernel(std::string name, const Platform& platform, CudaContext& cu) : IntegrateBrownianStepKernel(name, platform), cu(cu),
// hasInitializedKernels(false), prevTemp(-1), prevFriction(-1), prevStepSize(-1) {
// }
// ~CudaIntegrateBrownianStepKernel();
// /**
// * Initialize the kernel.
// *
// * @param system the System this kernel will be applied to
// * @param integrator the BrownianIntegrator this kernel will be used for
// */
// void initialize(const System& system, const BrownianIntegrator& integrator);
// /**
// * Execute the kernel.
// *
// * @param context the context in which to execute this kernel
// * @param integrator the BrownianIntegrator this kernel is being used for
// */
// void execute(ContextImpl& context, const BrownianIntegrator& integrator);
//private:
// CudaContext& cu;
// double prevTemp, prevFriction, prevStepSize;
// bool hasInitializedKernels;
// CUfunction kernel1, kernel2;
//};
//
///**
// * This kernel is invoked by VariableVerletIntegrator to take one time step.
// */
//class CudaIntegrateVariableVerletStepKernel : public IntegrateVariableVerletStepKernel {
//public:
// CudaIntegrateVariableVerletStepKernel(std::string name, const Platform& platform, CudaContext& cu) : IntegrateVariableVerletStepKernel(name, platform), cu(cu),
// hasInitializedKernels(false) {
// }
// ~CudaIntegrateVariableVerletStepKernel();
// /**
// * Initialize the kernel.
// *
// * @param system the System this kernel will be applied to
// * @param integrator the VerletIntegrator this kernel will be used for
// */
// void initialize(const System& system, const VariableVerletIntegrator& integrator);
// /**
// * Execute the kernel.
// *
// * @param context the context in which to execute this kernel
// * @param integrator the VerletIntegrator this kernel is being used for
// * @param maxTime the maximum time beyond which the simulation should not be advanced
// * @return the size of the step that was taken
// */
// double execute(ContextImpl& context, const VariableVerletIntegrator& integrator, double maxTime);
//private:
// CudaContext& cu;
// bool hasInitializedKernels;
// int blockSize;
// CUfunction kernel1, kernel2, selectSizeKernel;
//};
//
///**
// * This kernel is invoked by VariableLangevinIntegrator to take one time step.
// */
//class CudaIntegrateVariableLangevinStepKernel : public IntegrateVariableLangevinStepKernel {
//public:
// CudaIntegrateVariableLangevinStepKernel(std::string name, const Platform& platform, CudaContext& cu) : IntegrateVariableLangevinStepKernel(name, platform), cu(cu),
// hasInitializedKernels(false), params(NULL) {
// }
// ~CudaIntegrateVariableLangevinStepKernel();
// /**
// * Initialize the kernel, setting up the particle masses.
// *
// * @param system the System this kernel will be applied to
// * @param integrator the VariableLangevinIntegrator this kernel will be used for
// */
// void initialize(const System& system, const VariableLangevinIntegrator& integrator);
// /**
// * Execute the kernel.
// *
// * @param context the context in which to execute this kernel
// * @param integrator the VariableLangevinIntegrator this kernel is being used for
// * @param maxTime the maximum time beyond which the simulation should not be advanced
// * @return the size of the step that was taken
// */
// double execute(ContextImpl& context, const VariableLangevinIntegrator& integrator, double maxTime);
//private:
// CudaContext& cu;
// bool hasInitializedKernels;
// int blockSize;
// CudaArray<cl_float>* params;
// CUfunction kernel1, kernel2, selectSizeKernel;
// double prevTemp, prevFriction, prevErrorTol;
//};
//
/**
* This kernel is invoked by LangevinIntegrator to take one time step.
*/
class CudaIntegrateLangevinStepKernel : public IntegrateLangevinStepKernel {
public:
CudaIntegrateLangevinStepKernel(std::string name, const Platform& platform, CudaContext& cu) : IntegrateLangevinStepKernel(name, platform), cu(cu), params(NULL) {
}
~CudaIntegrateLangevinStepKernel();
/**
* Initialize the kernel, setting up the particle masses.
*
* @param system the System this kernel will be applied to
* @param integrator the LangevinIntegrator this kernel will be used for
*/
void initialize(const System& system, const LangevinIntegrator& integrator);
/**
* Execute the kernel.
*
* @param context the context in which to execute this kernel
* @param integrator the LangevinIntegrator this kernel is being used for
*/
void execute(ContextImpl& context, const LangevinIntegrator& integrator);
private:
CudaContext& cu;
double prevTemp, prevFriction, prevStepSize;
CudaArray* params;
CUfunction kernel1, kernel2;
};
/**
* This kernel is invoked by BrownianIntegrator to take one time step.
*/
class CudaIntegrateBrownianStepKernel : public IntegrateBrownianStepKernel {
public:
CudaIntegrateBrownianStepKernel(std::string name, const Platform& platform, CudaContext& cu) : IntegrateBrownianStepKernel(name, platform), cu(cu) {
}
~CudaIntegrateBrownianStepKernel();
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param integrator the BrownianIntegrator this kernel will be used for
*/
void initialize(const System& system, const BrownianIntegrator& integrator);
/**
* Execute the kernel.
*
* @param context the context in which to execute this kernel
* @param integrator the BrownianIntegrator this kernel is being used for
*/
void execute(ContextImpl& context, const BrownianIntegrator& integrator);
private:
CudaContext& cu;
double prevTemp, prevFriction, prevStepSize;
CUfunction kernel1, kernel2;
};
/**
* This kernel is invoked by VariableVerletIntegrator to take one time step.
*/
class CudaIntegrateVariableVerletStepKernel : public IntegrateVariableVerletStepKernel {
public:
CudaIntegrateVariableVerletStepKernel(std::string name, const Platform& platform, CudaContext& cu) : IntegrateVariableVerletStepKernel(name, platform), cu(cu) {
}
~CudaIntegrateVariableVerletStepKernel();
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param integrator the VerletIntegrator this kernel will be used for
*/
void initialize(const System& system, const VariableVerletIntegrator& integrator);
/**
* Execute the kernel.
*
* @param context the context in which to execute this kernel
* @param integrator the VerletIntegrator this kernel is being used for
* @param maxTime the maximum time beyond which the simulation should not be advanced
* @return the size of the step that was taken
*/
double execute(ContextImpl& context, const VariableVerletIntegrator& integrator, double maxTime);
private:
CudaContext& cu;
int blockSize;
CUfunction kernel1, kernel2, selectSizeKernel;
};
/**
* This kernel is invoked by VariableLangevinIntegrator to take one time step.
*/
class CudaIntegrateVariableLangevinStepKernel : public IntegrateVariableLangevinStepKernel {
public:
CudaIntegrateVariableLangevinStepKernel(std::string name, const Platform& platform, CudaContext& cu) : IntegrateVariableLangevinStepKernel(name, platform),
cu(cu), params(NULL) {
}
~CudaIntegrateVariableLangevinStepKernel();
/**
* Initialize the kernel, setting up the particle masses.
*
* @param system the System this kernel will be applied to
* @param integrator the VariableLangevinIntegrator this kernel will be used for
*/
void initialize(const System& system, const VariableLangevinIntegrator& integrator);
/**
* Execute the kernel.
*
* @param context the context in which to execute this kernel
* @param integrator the VariableLangevinIntegrator this kernel is being used for
* @param maxTime the maximum time beyond which the simulation should not be advanced
* @return the size of the step that was taken
*/
double execute(ContextImpl& context, const VariableLangevinIntegrator& integrator, double maxTime);
private:
CudaContext& cu;
int blockSize;
CudaArray* params;
CUfunction kernel1, kernel2, selectSizeKernel;
double prevTemp, prevFriction, prevErrorTol;
};
///**
// * This kernel is invoked by CustomIntegrator to take one time step.
// */
......
platforms/cuda2/src/CudaNonbondedUtilities.cpp
View file @ ecbbf442
......@@ -169,14 +169,14 @@ void CudaNonbondedUtilities::initialize(const System& system) {
exclusionIndicesVec.push_back(iter->second);
}
exclusionRowIndicesVec[++currentRow] = exclusionIndicesVec.size();
exclusionIndices = CudaArray::create<unsigned int>(exclusionIndicesVec.size(), "exclusionIndices");
exclusionRowIndices = CudaArray::create<unsigned int>(exclusionRowIndicesVec.size(), "exclusionRowIndices");
exclusionIndices = CudaArray::create<unsigned int>(context, exclusionIndicesVec.size(), "exclusionIndices");
exclusionRowIndices = CudaArray::create<unsigned int>(context, exclusionRowIndicesVec.size(), "exclusionRowIndices");
exclusionIndices->upload(exclusionIndicesVec);
exclusionRowIndices->upload(exclusionRowIndicesVec);
// Record the exclusion data.
exclusions = CudaArray::create<unsigned int>(tilesWithExclusions.size()*CudaContext::TileSize, "exclusions");
exclusions = CudaArray::create<unsigned int>(context, tilesWithExclusions.size()*CudaContext::TileSize, "exclusions");
vector<unsigned int> exclusionVec(exclusions->getSize());
for (int i = 0; i < exclusions->getSize(); ++i)
exclusionVec[i] = 0xFFFFFFFF;
......@@ -231,11 +231,11 @@ void CudaNonbondedUtilities::initialize(const System& system) {
maxTiles = numTiles;
if (maxTiles < 1)
maxTiles = 1;
interactingTiles = CudaArray::create<ushort2>(maxTiles, "interactingTiles");
interactionFlags = CudaArray::create<unsigned int>(maxTiles, "interactionFlags");
interactionCount = CudaArray::create<unsigned int>(1, "interactionCount");
blockCenter = CudaArray::create<float4>(numAtomBlocks, "blockCenter");
blockBoundingBox = CudaArray::create<float4>(numAtomBlocks, "blockBoundingBox");
interactingTiles = CudaArray::create<ushort2>(context, maxTiles, "interactingTiles");
interactionFlags = CudaArray::create<unsigned int>(context, maxTiles, "interactionFlags");
interactionCount = CudaArray::create<unsigned int>(context, 1, "interactionCount");
blockCenter = CudaArray::create<float4>(context, numAtomBlocks, "blockCenter");
blockBoundingBox = CudaArray::create<float4>(context, numAtomBlocks, "blockBoundingBox");
CHECK_RESULT(cuMemHostAlloc((void**) &pinnedInteractionCount, sizeof(unsigned int), 0));
pinnedInteractionCount[0] = 0;
interactionCount->upload(pinnedInteractionCount);
......@@ -330,11 +330,11 @@ void CudaNonbondedUtilities::updateNeighborListSize() {
if (maxTiles > numTiles)
maxTiles = numTiles;
delete interactingTiles;
interactingTiles = CudaArray::create<ushort2>(maxTiles, "interactingTiles");
interactingTiles = CudaArray::create<ushort2>(context, maxTiles, "interactingTiles");
forceArgs[8] = &interactingTiles->getDevicePointer();
findInteractingBlocksArgs[5] = &interactingTiles->getDevicePointer();
delete interactionFlags;
interactionFlags = CudaArray::create<unsigned int>(maxTiles, "interactionFlags");
interactionFlags = CudaArray::create<unsigned int>(context, maxTiles, "interactionFlags");
forceArgs[13] = &interactionFlags->getDevicePointer();
findInteractingBlocksArgs[6] = &interactionFlags->getDevicePointer();
findInteractionsWithinBlocksArgs[3] = &interactingTiles->getDevicePointer();
......
platforms/cuda2/src/CudaSort.cpp
View file @ ecbbf442
......@@ -73,11 +73,11 @@ CudaSort::CudaSort(CudaContext& context, SortTrait* trait, unsigned int length)
// Create workspace arrays.
dataRange = new CudaArray(2, trait->getKeySize(), "sortDataRange");
bucketOffset = CudaArray::create<uint1>(numBuckets, "bucketOffset");
bucketOfElement = CudaArray::create<uint1>(length, "bucketOfElement");
offsetInBucket = CudaArray::create<uint1>(length, "offsetInBucket");
buckets = new CudaArray(length, trait->getDataSize(), "buckets");
dataRange = new CudaArray(context, 2, trait->getKeySize(), "sortDataRange");
bucketOffset = CudaArray::create<uint1>(context, numBuckets, "bucketOffset");
bucketOfElement = CudaArray::create<uint1>(context, length, "bucketOfElement");
offsetInBucket = CudaArray::create<uint1>(context, length, "offsetInBucket");
buckets = new CudaArray(context, length, trait->getDataSize(), "buckets");
}
CudaSort::~CudaSort() {
......
platforms/cuda2/src/kernels/brownian.cu 0 → 100644
View file @ ecbbf442
/**
* Perform the first step of Brownian integration.
*/
extern "C" __global__ void integrateBrownianPart1(real tauDeltaT, real noiseAmplitude, const long long* __restrict__ force,
real4* __restrict__ posDelta, const real4* __restrict__ velm, const float4* __restrict__ random, unsigned int randomIndex) {
randomIndex += blockIdx.x*blockDim.x+threadIdx.x;
const real fscale = tauDeltaT/(real) 0xFFFFFFFF;
for (int index = blockIdx.x*blockDim.x+threadIdx.x; index < NUM_ATOMS; index += blockDim.x*gridDim.x) {
real invMass = velm[index].w;
if (invMass != 0) {
posDelta[index].x = fscale*invMass*force[index] + noiseAmplitude*SQRT(invMass)*random[randomIndex].x;
posDelta[index].y = fscale*invMass*force[index+PADDED_NUM_ATOMS] + noiseAmplitude*SQRT(invMass)*random[randomIndex].y;
posDelta[index].z = fscale*invMass*force[index+PADDED_NUM_ATOMS*2] + noiseAmplitude*SQRT(invMass)*random[randomIndex].z;
}
randomIndex += blockDim.x*gridDim.x;
}
}
/**
* Perform the second step of Brownian integration.
*/
extern "C" __global__ void integrateBrownianPart2(real deltaT, real4* posq, real4* velm, const real4* __restrict__ posDelta) {
const real oneOverDeltaT = RECIP(deltaT);
for (int index = blockIdx.x*blockDim.x+threadIdx.x; index < NUM_ATOMS; index += blockDim.x*gridDim.x) {
if (velm[index].w != 0) {
real4 delta = posDelta[index];
velm[index].x = oneOverDeltaT*delta.x;
velm[index].y = oneOverDeltaT*delta.y;
velm[index].z = oneOverDeltaT*delta.z;
posq[index].x = posq[index].x + delta.x;
posq[index].y = posq[index].y + delta.y;
posq[index].z = posq[index].z + delta.z;
}
}
}
platforms/cuda2/src/kernels/ccma.cu 0 → 100644
View file @ ecbbf442
/**
* Compute the direction each constraint is pointing in. This is called once at the beginning of constraint evaluation.
*/
extern "C" __global__ void computeConstraintDirections(const int2* __restrict__ constraintAtoms, real4* __restrict__ constraintDistance, const real4* __restrict__ atomPositions) {
for (int index = blockIdx.x*blockDim.x+threadIdx.x; index < NUM_CONSTRAINTS; index += blockDim.x*gridDim.x) {
// Compute the direction for this constraint.
int2 atoms = constraintAtoms[index];
real4 dir = constraintDistance[index];
real4 oldPos1 = atomPositions[atoms.x];
real4 oldPos2 = atomPositions[atoms.y];
dir.x = oldPos1.x-oldPos2.x;
dir.y = oldPos1.y-oldPos2.y;
dir.z = oldPos1.z-oldPos2.z;
constraintDistance[index] = dir;
}
}
/**
* Compute the force applied by each constraint.
*/
extern "C" __global__ void computeConstraintForce(const int2* __restrict__ constraintAtoms, const real4* __restrict__ constraintDistance, const real4* __restrict__ atomPositions,
const real* __restrict__ reducedMass, real* __restrict__ delta1, int* __restrict__ converged, float tol, int iteration) {
__shared__ int groupConverged;
if (converged[1-iteration%2]) {
if (blockIdx.x == 0 && threadIdx.x == 0)
converged[iteration%2] = 1;
return; // The constraint iteration has already converged.
}
if (threadIdx.x == 0)
groupConverged = 1;
__syncthreads();
real lowerTol = 1.0f-2.0f*tol+tol*tol;
real upperTol = 1.0f+2.0f*tol+tol*tol;
for (int index = blockIdx.x*blockDim.x+threadIdx.x; index < NUM_CONSTRAINTS; index += blockDim.x*gridDim.x) {
// Compute the force due to this constraint.
int2 atoms = constraintAtoms[index];
real4 dir = constraintDistance[index];
real4 rp_ij = atomPositions[atoms.x]-atomPositions[atoms.y];
#ifndef CONSTRAIN_VELOCITIES
rp_ij.x += dir.x;
rp_ij.y += dir.y;
rp_ij.z += dir.z;
#endif
real rrpr = rp_ij.x*dir.x + rp_ij.y*dir.y + rp_ij.z*dir.z;
real d_ij2 = dir.x*dir.x + dir.y*dir.y + dir.z*dir.z;
#ifdef CONSTRAIN_VELOCITIES
delta1[index] = -2.0f*reducedMass[index]*rrpr/d_ij2;
// See whether it has converged.
if (groupConverged && fabs(delta1[index]) > tol) {
groupConverged = 0;
converged[iteration%2] = 0;
}
#else
real rp2 = rp_ij.x*rp_ij.x + rp_ij.y*rp_ij.y + rp_ij.z*rp_ij.z;
real dist2 = dir.w*dir.w;
real diff = dist2 - rp2;
delta1[index] = (rrpr > d_ij2*1e-6f ? reducedMass[index]*diff/rrpr : 0.0f);
// See whether it has converged.
if (groupConverged && (rp2 < lowerTol*dist2 || rp2 > upperTol*dist2)) {
groupConverged = 0;
converged[iteration%2] = 0;
}
#endif
}
}
/**
* Multiply the vector of constraint forces by the constraint matrix.
*/
extern "C" __global__ void multiplyByConstraintMatrix(const real* __restrict__ delta1, real* __restrict__ delta2, const int* __restrict__ constraintMatrixColumn,
const real* __restrict__ constraintMatrixValue, const int* __restrict__ converged, int iteration) {
if (converged[iteration%2])
return; // The constraint iteration has already converged.
// Multiply by the inverse constraint matrix.
for (int index = blockIdx.x*blockDim.x+threadIdx.x; index < NUM_CONSTRAINTS; index += blockDim.x*gridDim.x) {
real sum = 0.0f;
for (int i = 0; ; i++) {
int element = index+i*NUM_CONSTRAINTS;
int column = constraintMatrixColumn[element];
if (column >= NUM_CONSTRAINTS)
break;
sum += delta1[column]*constraintMatrixValue[element];
}
delta2[index] = sum;
}
}
/**
* Update the atom positions based on constraint forces.
*/
extern "C" __global__ void updateAtomPositions(const int* __restrict__ numAtomConstraints, const int* __restrict__ atomConstraints, const real4* __restrict__ constraintDistance,
real4* __restrict__ atomPositions, const real4* __restrict__ velm, const real* __restrict__ delta1, const real* __restrict__ delta2, int* __restrict__ converged, int iteration) {
if (blockIdx.x == 0 && threadIdx.x == 0)
converged[1-iteration%2] = 1;
if (converged[iteration%2])
return; // The constraint iteration has already converged.
real damping = (iteration < 2 ? 0.5f : 1.0f);
for (int index = blockIdx.x*blockDim.x+threadIdx.x; index < NUM_ATOMS; index += blockDim.x*gridDim.x) {
// Compute the new position of this atom.
real4 atomPos = atomPositions[index];
real invMass = velm[index].w;
int num = numAtomConstraints[index];
for (int i = 0; i < num; i++) {
int constraint = atomConstraints[index+i*NUM_ATOMS];
bool forward = (constraint > 0);
constraint = (forward ? constraint-1 : -constraint-1);
real constraintForce = damping*invMass*delta2[constraint];
constraintForce = (forward ? constraintForce : -constraintForce);
real4 dir = constraintDistance[constraint];
atomPos.x += constraintForce*dir.x;
atomPos.y += constraintForce*dir.y;
atomPos.z += constraintForce*dir.z;
}
atomPositions[index] = atomPos;
}
}
platforms/cuda2/src/kernels/constraints.cu 0 → 100644
View file @ ecbbf442
extern "C" __global__ void applyPositionDeltas(real4* __restrict__ posq, real4* __restrict__ posDelta) {
for (unsigned int index = blockIdx.x*blockDim.x+threadIdx.x; index < NUM_ATOMS; index += blockDim.x*gridDim.x) {
real4 position = posq[index];
position.x += posDelta[index].x;
position.y += posDelta[index].y;
position.z += posDelta[index].z;
posq[index] = position;
}
}
platforms/cuda2/src/kernels/langevin.cu 0 → 100644
View file @ ecbbf442
enum {VelScale, ForceScale, NoiseScale, MaxParams};
/**
* Perform the first step of Langevin integration.
*/
extern "C" __global__ void integrateLangevinPart1(real4* __restrict__ velm, const long long* __restrict__ force, real4* __restrict__ posDelta,
const real* __restrict__ paramBuffer, const real2* __restrict__ dt, const float4* __restrict__ random, unsigned int randomIndex) {
real vscale = paramBuffer[VelScale];
real fscale = paramBuffer[ForceScale]/(real) 0xFFFFFFFF;
real noisescale = paramBuffer[NoiseScale];
real stepSize = dt[0].y;
int index = blockIdx.x*blockDim.x+threadIdx.x;
randomIndex += index;
while (index < NUM_ATOMS) {
real4 velocity = velm[index];
if (velocity.w != 0) {
real sqrtInvMass = SQRT(velocity.w);
velocity.x = vscale*velocity.x + fscale*velocity.w*force[index] + noisescale*sqrtInvMass*random[randomIndex].x;
velocity.y = vscale*velocity.y + fscale*velocity.w*force[index+PADDED_NUM_ATOMS] + noisescale*sqrtInvMass*random[randomIndex].y;
velocity.z = vscale*velocity.z + fscale*velocity.w*force[index+PADDED_NUM_ATOMS*2] + noisescale*sqrtInvMass*random[randomIndex].z;
velm[index] = velocity;
posDelta[index] = make_real4(stepSize*velocity.x, stepSize*velocity.y, stepSize*velocity.z, 0);
}
randomIndex += blockDim.x*gridDim.x;
index += blockDim.x*gridDim.x;
}
}
/**
* Perform the second step of Langevin integration.
*/
extern "C" __global__ void integrateLangevinPart2(real4* __restrict__ posq, const real4* __restrict__ posDelta, real4* __restrict__ velm, const real2* __restrict__ dt) {
double invStepSize = 1.0/dt[0].y;
int index = blockIdx.x*blockDim.x+threadIdx.x;
while (index < NUM_ATOMS) {
real4 vel = velm[index];
if (vel.w != 0) {
real4 pos = posq[index];
real4 delta = posDelta[index];
pos.x += delta.x;
pos.y += delta.y;
pos.z += delta.z;
vel.x = (real) invStepSize*delta.x;
vel.y = (real) invStepSize*delta.y;
vel.z = (real) invStepSize*delta.z;
posq[index] = pos;
velm[index] = vel;
}
index += blockDim.x*gridDim.x;
}
}
/**
* Select the step size to use for the next step.
*/
extern "C" __global__ void selectLangevinStepSize(real maxStepSize, real errorTol, real tau, real kT, real2* __restrict__ dt,
const real4* __restrict__ velm, const long long* __restrict__ force, real* __restrict__ paramBuffer) {
// Calculate the error.
extern __shared__ real params[];
real* error = &params[MaxParams];
real err = 0;
unsigned int index = threadIdx.x;
const real scale = RECIP((real) 0xFFFFFFFF);
while (index < NUM_ATOMS) {
real3 f = make_real3(scale*force[index], scale*force[index+PADDED_NUM_ATOMS], scale*force[index+PADDED_NUM_ATOMS*2]);
real invMass = velm[index].w;
err += (f.x*f.x + f.y*f.y + f.z*f.z)*invMass;
index += blockDim.x*gridDim.x;
}
error[threadIdx.x] = err;
__syncthreads();
// Sum the errors from all threads.
for (unsigned int offset = 1; offset < blockDim.x; offset *= 2) {
if (threadIdx.x+offset < blockDim.x && (threadIdx.x&(2*offset-1)) == 0)
error[threadIdx.x] += error[threadIdx.x+offset];
__syncthreads();
}
if (blockIdx.x*blockDim.x+threadIdx.x == 0) {
// Select the new step size.
real totalError = sqrt(error[0]/(NUM_ATOMS*3));
real newStepSize = sqrt(errorTol/totalError);
real oldStepSize = dt[0].y;
if (oldStepSize > 0.0f)
newStepSize = min(newStepSize, oldStepSize*2.0f); // For safety, limit how quickly dt can increase.
if (newStepSize > oldStepSize && newStepSize < 1.1f*oldStepSize)
newStepSize = oldStepSize; // Keeping dt constant between steps improves the behavior of the integrator.
if (newStepSize > maxStepSize)
newStepSize = maxStepSize;
dt[0].y = newStepSize;
// Recalculate the integration parameters.
real vscale = exp(-newStepSize/tau);
real fscale = (1-vscale)*tau;
real noisescale = sqrt(2*kT/tau)*sqrt(0.5f*(1-vscale*vscale)*tau);
params[VelScale] = vscale;
params[ForceScale] = fscale;
params[NoiseScale] = noisescale;
}
__syncthreads();
if (threadIdx.x < MaxParams)
paramBuffer[threadIdx.x] = params[threadIdx.x];
}
platforms/cuda2/src/kernels/settle.cu 0 → 100644
View file @ ecbbf442
/**
* Enforce constraints on SETTLE clusters
*/
extern "C" __global__ void applySettle(int numClusters, float tol, const real4* __restrict__ oldPos, real4* __restrict__ posDelta, const real4* __restrict__ velm, const int4* __restrict__ clusterAtoms, const float2* __restrict__ clusterParams) {
int index = blockIdx.x*blockDim.x+threadIdx.x;
while (index < numClusters) {
// Load the data for this cluster.
int4 atoms = clusterAtoms[index];
float2 params = clusterParams[index];
real4 apos0 = oldPos[atoms.x];
real4 xp0 = posDelta[atoms.x];
real4 apos1 = oldPos[atoms.y];
real4 xp1 = posDelta[atoms.y];
real4 apos2 = oldPos[atoms.z];
real4 xp2 = posDelta[atoms.z];
real m0 = RECIP(velm[atoms.x].w);
real m1 = RECIP(velm[atoms.y].w);
real m2 = RECIP(velm[atoms.z].w);
// Apply the SETTLE algorithm.
real xb0 = apos1.x-apos0.x;
real yb0 = apos1.y-apos0.y;
real zb0 = apos1.z-apos0.z;
real xc0 = apos2.x-apos0.x;
real yc0 = apos2.y-apos0.y;
real zc0 = apos2.z-apos0.z;
real invTotalMass = RECIP(m0+m1+m2);
real xcom = (xp0.x*m0 + (xb0+xp1.x)*m1 + (xc0+xp2.x)*m2) * invTotalMass;
real ycom = (xp0.y*m0 + (yb0+xp1.y)*m1 + (yc0+xp2.y)*m2) * invTotalMass;
real zcom = (xp0.z*m0 + (zb0+xp1.z)*m1 + (zc0+xp2.z)*m2) * invTotalMass;
real xa1 = xp0.x - xcom;
real ya1 = xp0.y - ycom;
real za1 = xp0.z - zcom;
real xb1 = xb0 + xp1.x - xcom;
real yb1 = yb0 + xp1.y - ycom;
real zb1 = zb0 + xp1.z - zcom;
real xc1 = xc0 + xp2.x - xcom;
real yc1 = yc0 + xp2.y - ycom;
real zc1 = zc0 + xp2.z - zcom;
real xaksZd = yb0*zc0 - zb0*yc0;
real yaksZd = zb0*xc0 - xb0*zc0;
real zaksZd = xb0*yc0 - yb0*xc0;
real xaksXd = ya1*zaksZd - za1*yaksZd;
real yaksXd = za1*xaksZd - xa1*zaksZd;
real zaksXd = xa1*yaksZd - ya1*xaksZd;
real xaksYd = yaksZd*zaksXd - zaksZd*yaksXd;
real yaksYd = zaksZd*xaksXd - xaksZd*zaksXd;
real zaksYd = xaksZd*yaksXd - yaksZd*xaksXd;
real axlng = SQRT(xaksXd*xaksXd + yaksXd*yaksXd + zaksXd*zaksXd);
real aylng = SQRT(xaksYd*xaksYd + yaksYd*yaksYd + zaksYd*zaksYd);
real azlng = SQRT(xaksZd*xaksZd + yaksZd*yaksZd + zaksZd*zaksZd);
real trns11 = xaksXd / axlng;
real trns21 = yaksXd / axlng;
real trns31 = zaksXd / axlng;
real trns12 = xaksYd / aylng;
real trns22 = yaksYd / aylng;
real trns32 = zaksYd / aylng;
real trns13 = xaksZd / azlng;
real trns23 = yaksZd / azlng;
real trns33 = zaksZd / azlng;
real xb0d = trns11*xb0 + trns21*yb0 + trns31*zb0;
real yb0d = trns12*xb0 + trns22*yb0 + trns32*zb0;
real xc0d = trns11*xc0 + trns21*yc0 + trns31*zc0;
real yc0d = trns12*xc0 + trns22*yc0 + trns32*zc0;
real za1d = trns13*xa1 + trns23*ya1 + trns33*za1;
real xb1d = trns11*xb1 + trns21*yb1 + trns31*zb1;
real yb1d = trns12*xb1 + trns22*yb1 + trns32*zb1;
real zb1d = trns13*xb1 + trns23*yb1 + trns33*zb1;
real xc1d = trns11*xc1 + trns21*yc1 + trns31*zc1;
real yc1d = trns12*xc1 + trns22*yc1 + trns32*zc1;
real zc1d = trns13*xc1 + trns23*yc1 + trns33*zc1;
// --- Step2 A2' ---
float rc = 0.5f*params.y;
float rb = SQRT(params.x*params.x-rc*rc);
real ra = rb*(m1+m2)*invTotalMass;
rb -= ra;
real sinphi = za1d/ra;
real cosphi = SQRT(1-sinphi*sinphi);
real sinpsi = (zb1d-zc1d) / (2*rc*cosphi);
real cospsi = SQRT(1-sinpsi*sinpsi);
real ya2d = ra*cosphi;
real xb2d = - rc*cospsi;
real yb2d = - rb*cosphi - rc*sinpsi*sinphi;
real yc2d = - rb*cosphi + rc*sinpsi*sinphi;
real xb2d2 = xb2d*xb2d;
real hh2 = 4.0f*xb2d2 + (yb2d-yc2d)*(yb2d-yc2d) + (zb1d-zc1d)*(zb1d-zc1d);
real deltx = 2.0f*xb2d + SQRT(4.0f*xb2d2 - hh2 + params.y*params.y);
xb2d -= deltx*0.5f;
// --- Step3 al,be,ga ---
real alpha = (xb2d*(xb0d-xc0d) + yb0d*yb2d + yc0d*yc2d);
real beta = (xb2d*(yc0d-yb0d) + xb0d*yb2d + xc0d*yc2d);
real gamma = xb0d*yb1d - xb1d*yb0d + xc0d*yc1d - xc1d*yc0d;
real al2be2 = alpha*alpha + beta*beta;
real sintheta = (alpha*gamma - beta*SQRT(al2be2 - gamma*gamma)) / al2be2;
// --- Step4 A3' ---
real costheta = SQRT(1-sintheta*sintheta);
real xa3d = - ya2d*sintheta;
real ya3d = ya2d*costheta;
real za3d = za1d;
real xb3d = xb2d*costheta - yb2d*sintheta;
real yb3d = xb2d*sintheta + yb2d*costheta;
real zb3d = zb1d;
real xc3d = - xb2d*costheta - yc2d*sintheta;
real yc3d = - xb2d*sintheta + yc2d*costheta;
real zc3d = zc1d;
// --- Step5 A3 ---
real xa3 = trns11*xa3d + trns12*ya3d + trns13*za3d;
real ya3 = trns21*xa3d + trns22*ya3d + trns23*za3d;
real za3 = trns31*xa3d + trns32*ya3d + trns33*za3d;
real xb3 = trns11*xb3d + trns12*yb3d + trns13*zb3d;
real yb3 = trns21*xb3d + trns22*yb3d + trns23*zb3d;
real zb3 = trns31*xb3d + trns32*yb3d + trns33*zb3d;
real xc3 = trns11*xc3d + trns12*yc3d + trns13*zc3d;
real yc3 = trns21*xc3d + trns22*yc3d + trns23*zc3d;
real zc3 = trns31*xc3d + trns32*yc3d + trns33*zc3d;
xp0.x = xcom + xa3;
xp0.y = ycom + ya3;
xp0.z = zcom + za3;
xp1.x = xcom + xb3 - xb0;
xp1.y = ycom + yb3 - yb0;
xp1.z = zcom + zb3 - zb0;
xp2.x = xcom + xc3 - xc0;
xp2.y = ycom + yc3 - yc0;
xp2.z = zcom + zc3 - zc0;
// Record the new positions.
posDelta[atoms.x] = xp0;
posDelta[atoms.y] = xp1;
posDelta[atoms.z] = xp2;
index += blockDim.x*gridDim.x;
}
}
/**
* Enforce velocity constraints on SETTLE clusters
*/
extern "C" __global__ void constrainVelocities(int numClusters, float tol, const real4* __restrict__ oldPos, real4* __restrict__ posDelta, real4* __restrict__ velm, const int4* __restrict__ clusterAtoms, const float2* __restrict__ clusterParams) {
for (int index = blockIdx.x*blockDim.x+threadIdx.x; index < numClusters; index += blockDim.x*gridDim.x) {
// Load the data for this cluster.
int4 atoms = clusterAtoms[index];
real4 apos0 = oldPos[atoms.x];
real4 apos1 = oldPos[atoms.y];
real4 apos2 = oldPos[atoms.z];
real4 v0 = velm[atoms.x];
real4 v1 = velm[atoms.y];
real4 v2 = velm[atoms.z];
// Compute intermediate quantities: the atom masses, the bond directions, the relative velocities,
// and the angle cosines and sines.
real mA = RECIP(v0.w);
real mB = RECIP(v1.w);
real mC = RECIP(v2.w);
real3 eAB = make_real3(apos1.x-apos0.x, apos1.y-apos0.y, apos1.z-apos0.z);
real3 eBC = make_real3(apos2.x-apos1.x, apos2.y-apos1.y, apos2.z-apos1.z);
real3 eCA = make_real3(apos0.x-apos2.x, apos0.y-apos2.y, apos0.z-apos2.z);
eAB *= RECIP(SQRT(eAB.x*eAB.x + eAB.y*eAB.y + eAB.z*eAB.z));
eBC *= RECIP(SQRT(eBC.x*eBC.x + eBC.y*eBC.y + eBC.z*eBC.z));
eCA *= RECIP(SQRT(eCA.x*eCA.x + eCA.y*eCA.y + eCA.z*eCA.z));
real vAB = (v1.x-v0.x)*eAB.x + (v1.y-v0.y)*eAB.y + (v1.z-v0.z)*eAB.z;
real vBC = (v2.x-v1.x)*eBC.x + (v2.y-v1.y)*eBC.y + (v2.z-v1.z)*eBC.z;
real vCA = (v0.x-v2.x)*eCA.x + (v0.y-v2.y)*eCA.y + (v0.z-v2.z)*eCA.z;
real cA = -(eAB.x*eCA.x + eAB.y*eCA.y + eAB.z*eCA.z);
real cB = -(eAB.x*eBC.x + eAB.y*eBC.y + eAB.z*eBC.z);
real cC = -(eBC.x*eCA.x + eBC.y*eCA.y + eBC.z*eCA.z);
real s2A = 1-cA*cA;
real s2B = 1-cB*cB;
real s2C = 1-cC*cC;
// Solve the equations. These are different from those in the SETTLE paper (JCC 13(8), pp. 952-962, 1992), because
// in going from equations B1 to B2, they make the assumption that mB=mC (but don't bother to mention they're
// making that assumption). We allow all three atoms to have different masses.
real mABCinv = RECIP(mA*mB*mC);
real denom = (((s2A*mB+s2B*mA)*mC+(s2A*mB*mB+2*(cA*cB*cC+1)*mA*mB+s2B*mA*mA))*mC+s2C*mA*mB*(mA+mB))*mABCinv;
real tab = ((cB*cC*mA-cA*mB-cA*mC)*vCA + (cA*cC*mB-cB*mC-cB*mA)*vBC + (s2C*mA*mA*mB*mB*mABCinv+(mA+mB+mC))*vAB)/denom;
real tbc = ((cA*cB*mC-cC*mB-cC*mA)*vCA + (s2A*mB*mB*mC*mC*mABCinv+(mA+mB+mC))*vBC + (cA*cC*mB-cB*mA-cB*mC)*vAB)/denom;
real tca = ((s2B*mA*mA*mC*mC*mABCinv+(mA+mB+mC))*vCA + (cA*cB*mC-cC*mB-cC*mA)*vBC + (cB*cC*mA-cA*mB-cA*mC)*vAB)/denom;
v0.x += (tab*eAB.x - tca*eCA.x)*v0.w;
v0.y += (tab*eAB.y - tca*eCA.y)*v0.w;
v0.z += (tab*eAB.z - tca*eCA.z)*v0.w;
v1.x += (tbc*eBC.x - tab*eAB.x)*v1.w;
v1.y += (tbc*eBC.y - tab*eAB.y)*v1.w;
v1.z += (tbc*eBC.z - tab*eAB.z)*v1.w;
v2.x += (tca*eCA.x - tbc*eBC.x)*v2.w;
v2.y += (tca*eCA.y - tbc*eBC.y)*v2.w;
v2.z += (tca*eCA.z - tbc*eBC.z)*v2.w;
velm[atoms.x] = v0;
velm[atoms.y] = v1;
velm[atoms.z] = v2;
}
}
\ No newline at end of file
platforms/cuda2/src/kernels/shakeHydrogens.cu 0 → 100644
View file @ ecbbf442
/**
* Enforce constraints on SHAKE clusters
*/
extern "C" __global__ void applyShakeToHydrogens(int numClusters, real tol, const real4* __restrict__ oldPos, real4* __restrict__ posDelta, const int4* __restrict__ clusterAtoms, const float4* __restrict__ clusterParams) {
int index = blockIdx.x*blockDim.x+threadIdx.x;
while (index < numClusters) {
// Load the data for this cluster.
int4 atoms = clusterAtoms[index];
float4 params = clusterParams[index];
real4 pos = oldPos[atoms.x];
real4 xpi = posDelta[atoms.x];
real4 pos1 = oldPos[atoms.y];
real4 xpj1 = posDelta[atoms.y];
real4 pos2 = make_real4(0);
real4 xpj2 = make_real4(0);
real invMassCentral = params.x;
real avgMass = params.y;
float d2 = params.z;
float invMassPeripheral = params.w;
if (atoms.z != -1) {
pos2 = oldPos[atoms.z];
xpj2 = posDelta[atoms.z];
}
real4 pos3 = make_real4(0);
real4 xpj3 = make_real4(0);
if (atoms.w != -1) {
pos3 = oldPos[atoms.w];
xpj3 = posDelta[atoms.w];
}
// Precompute quantities.
real3 rij1 = make_real3(pos.x-pos1.x, pos.y-pos1.y, pos.z-pos1.z);
real3 rij2 = make_real3(pos.x-pos2.x, pos.y-pos2.y, pos.z-pos2.z);
real3 rij3 = make_real3(pos.x-pos3.x, pos.y-pos3.y, pos.z-pos3.z);
real rij1sq = rij1.x*rij1.x + rij1.y*rij1.y + rij1.z*rij1.z;
real rij2sq = rij2.x*rij2.x + rij2.y*rij2.y + rij2.z*rij2.z;
real rij3sq = rij3.x*rij3.x + rij3.y*rij3.y + rij3.z*rij3.z;
real ld1 = d2-rij1sq;
real ld2 = d2-rij2sq;
real ld3 = d2-rij3sq;
// Iterate until convergence.
bool converged = false;
int iteration = 0;
while (iteration < 15 && !converged) {
converged = true;
#ifdef CONSTRAIN_VELOCITIES
real3 rpij = make_real3(xpi.x-xpj1.x, xpi.y-xpj1.y, xpi.z-xpj1.z);
real rrpr = rpij.x*rij1.x + rpij.y*rij1.y + rpij.z*rij1.z;
real delta = -2.0f*avgMass*rrpr/rij1sq;
real3 dr = rij1*delta;
xpi.x += dr.x*invMassCentral;
xpi.y += dr.y*invMassCentral;
xpi.z += dr.z*invMassCentral;
xpj1.x -= dr.x*invMassPeripheral;
xpj1.y -= dr.y*invMassPeripheral;
xpj1.z -= dr.z*invMassPeripheral;
if (fabs(delta) > tol)
converged = false;
if (atoms.z != -1) {
rpij = make_real3(xpi.x-xpj2.x, xpi.y-xpj2.y, xpi.z-xpj2.z);
rrpr = rpij.x*rij2.x + rpij.y*rij2.y + rpij.z*rij2.z;
delta = -2.0f*avgMass*rrpr/rij2sq;
dr = rij2*delta;
xpi.x += dr.x*invMassCentral;
xpi.y += dr.y*invMassCentral;
xpi.z += dr.z*invMassCentral;
xpj2.x -= dr.x*invMassPeripheral;
xpj2.y -= dr.y*invMassPeripheral;
xpj2.z -= dr.z*invMassPeripheral;
if (fabs(delta) > tol)
converged = false;
}
if (atoms.w != -1) {
rpij = make_real3(xpi.x-xpj3.x, xpi.y-xpj3.y, xpi.z-xpj3.z);
rrpr = rpij.x*rij3.x + rpij.y*rij3.y + rpij.z*rij3.z;
delta = -2.0f*avgMass*rrpr/rij3sq;
dr = rij3*delta;
xpi.x += dr.x*invMassCentral;
xpi.y += dr.y*invMassCentral;
xpi.z += dr.z*invMassCentral;
xpj3.x -= dr.x*invMassPeripheral;
xpj3.y -= dr.y*invMassPeripheral;
xpj3.z -= dr.z*invMassPeripheral;
if (fabs(delta) > tol)
converged = false;
}
#else
real3 rpij = make_real3(xpi.x-xpj1.x, xpi.y-xpj1.y, xpi.z-xpj1.z);
real rpsqij = rpij.x*rpij.x + rpij.y*rpij.y + rpij.z*rpij.z;
real rrpr = rij1.x*rpij.x + rij1.y*rpij.y + rij1.z*rpij.z;
real diff = fabs(ld1-2.0f*rrpr-rpsqij) / (d2*tol);
if (diff >= 1.0f) {
real acor = (ld1-2.0f*rrpr-rpsqij)*avgMass / (rrpr+rij1sq);
real3 dr = rij1*acor;
xpi.x += dr.x*invMassCentral;
xpi.y += dr.y*invMassCentral;
xpi.z += dr.z*invMassCentral;
xpj1.x -= dr.x*invMassPeripheral;
xpj1.y -= dr.y*invMassPeripheral;
xpj1.z -= dr.z*invMassPeripheral;
converged = false;
}
if (atoms.z != -1) {
rpij = make_real3(xpi.x-xpj2.x, xpi.y-xpj2.y, xpi.z-xpj2.z);
rpsqij = rpij.x*rpij.x + rpij.y*rpij.y + rpij.z*rpij.z;
rrpr = rij2.x*rpij.x + rij2.y*rpij.y + rij2.z*rpij.z;
diff = fabs(ld2-2.0f*rrpr-rpsqij) / (d2*tol);
if (diff >= 1.0f) {
real acor = (ld2 - 2.0f*rrpr - rpsqij)*avgMass / (rrpr + rij2sq);
real3 dr = rij2*acor;
xpi.x += dr.x*invMassCentral;
xpi.y += dr.y*invMassCentral;
xpi.z += dr.z*invMassCentral;
xpj2.x -= dr.x*invMassPeripheral;
xpj2.y -= dr.y*invMassPeripheral;
xpj2.z -= dr.z*invMassPeripheral;
converged = false;
}
}
if (atoms.w != -1) {
rpij = make_real3(xpi.x-xpj3.x, xpi.y-xpj3.y, xpi.z-xpj3.z);
rpsqij = rpij.x*rpij.x + rpij.y*rpij.y + rpij.z*rpij.z;
rrpr = rij3.x*rpij.x + rij3.y*rpij.y + rij3.z*rpij.z;
diff = fabs(ld3 - 2.0f*rrpr - rpsqij) / (d2*tol);
if (diff >= 1.0f) {
real acor = (ld3-2.0f*rrpr-rpsqij)*avgMass / (rrpr+rij3sq);
real3 dr = rij3*acor;
xpi.x += dr.x*invMassCentral;
xpi.y += dr.y*invMassCentral;
xpi.z += dr.z*invMassCentral;
xpj3.x -= dr.x*invMassPeripheral;
xpj3.y -= dr.y*invMassPeripheral;
xpj3.z -= dr.z*invMassPeripheral;
converged = false;
}
}
#endif
iteration++;
}
// Record the new positions.
posDelta[atoms.x] = xpi;
posDelta[atoms.y] = xpj1;
if (atoms.z != -1)
posDelta[atoms.z] = xpj2;
if (atoms.w != -1)
posDelta[atoms.w] = xpj3;
index += blockDim.x*gridDim.x;
}
}
platforms/cuda2/src/kernels/verlet.cu
View file @ ecbbf442
......@@ -51,36 +51,38 @@ extern "C" __global__ void integrateVerletPart2(real2* __restrict__ dt, real4* _
/**
* Select the step size to use for the next step.
*/
//
//extern "C" __global__ void selectVerletStepSize(real maxStepSize, real errorTol, real2* __restrict__ dt, const real4* __restrict__ velm, const real4* __restrict__ force, __local real* __restrict__ error) {
// // Calculate the error.
//
// real err = 0.0f;
// for (int index = threadIdx.x; index < NUM_ATOMS; index += blockDim.x*gridDim.x) {
// real4 f = force[index];
// real invMass = velm[index].w;
// err += (f.x*f.x + f.y*f.y + f.z*f.z)*invMass;
// }
// error[threadIdx.x] = err;
// __syncthreads;
//
// // Sum the errors from all threads.
//
// for (unsigned int offset = 1; offset < get_local_size(0); offset *= 2) {
// if (threadIdx.x+offset < get_local_size(0) && (threadIdx.x&(2*offset-1)) == 0)
// error[threadIdx.x] += error[threadIdx.x+offset];
// __syncthreads;
// }
// if (threadIdx.x == 0) {
// real totalError = sqrt(error[0]/(NUM_ATOMS*3));
// real newStepSize = sqrt(errorTol/totalError);
// real oldStepSize = dt[0].y;
// if (oldStepSize > 0.0f)
// newStepSize = min(newStepSize, oldStepSize*2.0f); // For safety, limit how quickly dt can increase.
// if (newStepSize > oldStepSize && newStepSize < 1.1f*oldStepSize)
// newStepSize = oldStepSize; // Keeping dt constant between steps improves the behavior of the integrator.
// if (newStepSize > maxStepSize)
// newStepSize = maxStepSize;
// dt[0].y = newStepSize;
// }
//}
extern "C" __global__ void selectVerletStepSize(real maxStepSize, real errorTol, real2* __restrict__ dt, const real4* __restrict__ velm, const long long* __restrict__ force) {
// Calculate the error.
extern __shared__ real error[];
real err = 0.0f;
const real scale = RECIP((real) 0xFFFFFFFF);
for (int index = threadIdx.x; index < NUM_ATOMS; index += blockDim.x*gridDim.x) {
real3 f = make_real3(scale*force[index], scale*force[index+PADDED_NUM_ATOMS], scale*force[index+PADDED_NUM_ATOMS*2]);
real invMass = velm[index].w;
err += (f.x*f.x + f.y*f.y + f.z*f.z)*invMass;
}
error[threadIdx.x] = err;
__syncthreads();
// Sum the errors from all threads.
for (unsigned int offset = 1; offset < blockDim.x; offset *= 2) {
if (threadIdx.x+offset < blockDim.x && (threadIdx.x&(2*offset-1)) == 0)
error[threadIdx.x] += error[threadIdx.x+offset];
__syncthreads();
}
if (threadIdx.x == 0) {
real totalError = SQRT(error[0]/(NUM_ATOMS*3));
real newStepSize = SQRT(errorTol/totalError);
real oldStepSize = dt[0].y;
if (oldStepSize > 0.0f)
newStepSize = min(newStepSize, oldStepSize*2.0f); // For safety, limit how quickly dt can increase.
if (newStepSize > oldStepSize && newStepSize < 1.1f*oldStepSize)
newStepSize = oldStepSize; // Keeping dt constant between steps improves the behavior of the integrator.
if (newStepSize > maxStepSize)
newStepSize = maxStepSize;
dt[0].y = newStepSize;
}
}
platforms/cuda2/tests/TestCudaBrownianIntegrator.cpp 0 → 100644
View file @ ecbbf442
/* -------------------------------------------------------------------------- *
* 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) 2008-2012 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 "openmm/System.h"
/**
* This tests the CUDA implementation of BrownianIntegrator.
*/
#include "openmm/internal/AssertionUtilities.h"
#include "openmm/Context.h"
#include "CudaPlatform.h"
#include "openmm/HarmonicBondForce.h"
#include "openmm/NonbondedForce.h"
#include "openmm/System.h"
#include "openmm/BrownianIntegrator.h"
#include "../src/SimTKUtilities/SimTKOpenMMRealType.h"
#include "sfmt/SFMT.h"
#include <iostream>
#include <vector>
using namespace OpenMM;
using namespace std;
const double TOL = 1e-5;
void testSingleBond() {
CudaPlatform platform;
System system;
system.addParticle(2.0);
system.addParticle(2.0);
double dt = 0.01;
BrownianIntegrator integrator(0, 0.1, dt);
HarmonicBondForce* forceField = new HarmonicBondForce();
forceField->addBond(0, 1, 1.5, 1);
system.addForce(forceField);
Context context(system, integrator, platform);
vector<Vec3> positions(2);
positions[0] = Vec3(-1, 0, 0);
positions[1] = Vec3(1, 0, 0);
context.setPositions(positions);
// This is simply an overdamped harmonic oscillator, so compare it to the analytical solution.
double rate = 2*1.0/(0.1*2.0);
for (int i = 0; i < 1000; ++i) {
State state = context.getState(State::Positions | State::Velocities);
double time = state.getTime();
double expectedDist = 1.5+0.5*std::exp(-rate*time);
ASSERT_EQUAL_VEC(Vec3(-0.5*expectedDist, 0, 0), state.getPositions()[0], 0.02);
ASSERT_EQUAL_VEC(Vec3(0.5*expectedDist, 0, 0), state.getPositions()[1], 0.02);
if (i > 0) {
double expectedSpeed = -0.5*rate*std::exp(-rate*(time-0.5*dt));
ASSERT_EQUAL_VEC(Vec3(-0.5*expectedSpeed, 0, 0), state.getVelocities()[0], 0.11);
ASSERT_EQUAL_VEC(Vec3(0.5*expectedSpeed, 0, 0), state.getVelocities()[1], 0.11);
}
integrator.step(1);
}
}
void testTemperature() {
const int numParticles = 8;
const int numBonds = numParticles-1;
const double temp = 10.0;
CudaPlatform platform;
System system;
BrownianIntegrator integrator(temp, 2.0, 0.01);
HarmonicBondForce* forceField = new HarmonicBondForce();
for (int i = 0; i < numParticles; ++i)
system.addParticle(2.0);
for (int i = 0; i < numBonds; ++i)
forceField->addBond(i, i+1, 1.0, 5.0);
system.addForce(forceField);
Context context(system, integrator, platform);
vector<Vec3> positions(numParticles);
for (int i = 0; i < numParticles; ++i)
positions[i] = Vec3(i, 0, 0);
context.setPositions(positions);
// Let it equilibrate.
integrator.step(10000);
// Now run it for a while and see if the temperature is correct.
double pe = 0.0;
const int steps = 50000;
for (int i = 0; i < steps; ++i) {
State state = context.getState(State::Energy);
pe += state.getPotentialEnergy();
integrator.step(1);
}
pe /= steps;
double expected = 0.5*numBonds*BOLTZ*temp;
ASSERT_USUALLY_EQUAL_TOL(expected, pe, 0.1*expected);
}
void testConstraints() {
const int numParticles = 8;
const int numConstraints = 5;
const double temp = 20.0;
CudaPlatform platform;
System system;
BrownianIntegrator integrator(temp, 2.0, 0.001);
integrator.setConstraintTolerance(1e-5);
NonbondedForce* forceField = new NonbondedForce();
for (int i = 0; i < numParticles; ++i) {
system.addParticle(10.0);
forceField->addParticle((i%2 == 0 ? 0.2 : -0.2), 0.5, 5.0);
}
system.addConstraint(0, 1, 1.0);
system.addConstraint(1, 2, 1.0);
system.addConstraint(2, 3, 1.0);
system.addConstraint(4, 5, 1.0);
system.addConstraint(6, 7, 1.0);
system.addForce(forceField);
Context context(system, integrator, platform);
vector<Vec3> positions(numParticles);
vector<Vec3> velocities(numParticles);
OpenMM_SFMT::SFMT sfmt;
init_gen_rand(0, sfmt);
for (int i = 0; i < numParticles; ++i) {
positions[i] = Vec3(i, 0, 0);
velocities[i] = Vec3(genrand_real2(sfmt)-0.5, genrand_real2(sfmt)-0.5, genrand_real2(sfmt)-0.5);
}
context.setPositions(positions);
context.setVelocities(velocities);
// Simulate it and see whether the constraints remain satisfied.
for (int i = 0; i < 1000; ++i) {
State state = context.getState(State::Positions);
for (int j = 0; j < numConstraints; ++j) {
int particle1, particle2;
double distance;
system.getConstraintParameters(j, particle1, particle2, distance);
Vec3 p1 = state.getPositions()[particle1];
Vec3 p2 = state.getPositions()[particle2];
double dist = std::sqrt((p1[0]-p2[0])*(p1[0]-p2[0])+(p1[1]-p2[1])*(p1[1]-p2[1])+(p1[2]-p2[2])*(p1[2]-p2[2]));
ASSERT_EQUAL_TOL(distance, dist, 1e-4);
}
integrator.step(1);
}
}
void testRandomSeed() {
const int numParticles = 8;
const double temp = 100.0;
const double collisionFreq = 10.0;
CudaPlatform platform;
System system;
BrownianIntegrator integrator(temp, 2.0, 0.001);
NonbondedForce* forceField = new NonbondedForce();
for (int i = 0; i < numParticles; ++i) {
system.addParticle(2.0);
forceField->addParticle((i%2 == 0 ? 1.0 : -1.0), 1.0, 5.0);
}
system.addForce(forceField);
vector<Vec3> positions(numParticles);
vector<Vec3> velocities(numParticles);
for (int i = 0; i < numParticles; ++i) {
positions[i] = Vec3((i%2 == 0 ? 2 : -2), (i%4 < 2 ? 2 : -2), (i < 4 ? 2 : -2));
velocities[i] = Vec3(0, 0, 0);
}
// Try twice with the same random seed.
integrator.setRandomNumberSeed(5);
Context context(system, integrator, platform);
context.setPositions(positions);
context.setVelocities(velocities);
integrator.step(10);
State state1 = context.getState(State::Positions);
context.reinitialize();
context.setPositions(positions);
context.setVelocities(velocities);
integrator.step(10);
State state2 = context.getState(State::Positions);
// Try twice with a different random seed.
integrator.setRandomNumberSeed(10);
context.reinitialize();
context.setPositions(positions);
context.setVelocities(velocities);
integrator.step(10);
State state3 = context.getState(State::Positions);
context.reinitialize();
context.setPositions(positions);
context.setVelocities(velocities);
integrator.step(10);
State state4 = context.getState(State::Positions);
// Compare the results.
for (int i = 0; i < numParticles; i++) {
for (int j = 0; j < 3; j++) {
ASSERT(state1.getPositions()[i][j] == state2.getPositions()[i][j]);
ASSERT(state3.getPositions()[i][j] == state4.getPositions()[i][j]);
ASSERT(state1.getPositions()[i][j] != state3.getPositions()[i][j]);
}
}
}
int main() {
try {
testSingleBond();
testTemperature();
testConstraints();
testRandomSeed();
}
catch(const exception& e) {
cout << "exception: " << e.what() << endl;
return 1;
}
cout << "Done" << endl;
return 0;
}
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