Skip to content

GitLab

"platforms/hip/src/HipProgram.cpp" did not exist on "abb8cb4bb2d199e526b91f5307b18e6a5168e08e"
Commit b97ff259 authored by Peter Eastman's avatar Peter Eastman
Browse files

All platforms now handle nonbonded exceptions in a consistent way (they ignore... 

All platforms now handle nonbonded exceptions in a consistent way (they ignore cutoffs and periodic boundary conditions).  Also started OpenCL implementation of CustomHbondForce.
parent ca63e6d8
Expand all Hide whitespace changes
Inline Side-by-side
Showing with 639 additions and 329 deletions
platforms/cuda/src/kernels/kCalculateCDLJForces.h
View file @ b97ff259
......@@ -213,7 +213,7 @@ void METHOD_NAME(kCalculateCDLJ, Forces_kernel)(unsigned int* workUnit)
dEdR *= invR * invR;
#ifdef USE_CUTOFF
#ifdef USE_EWALD
if ((!(excl & 0x1) && !needCorrection) || r2 > cSim.nonbondedCutoffSqr)
if (!needCorrection && (!(excl & 0x1) || r2 > cSim.nonbondedCutoffSqr))
#else
if (!(excl & 0x1) || r2 > cSim.nonbondedCutoffSqr)
#endif
......@@ -501,7 +501,7 @@ void METHOD_NAME(kCalculateCDLJ, Forces_kernel)(unsigned int* workUnit)
dEdR *= invR * invR;
#ifdef USE_CUTOFF
#ifdef USE_EWALD
if ((!(excl & 0x1) && !needCorrection) || r2 > cSim.nonbondedCutoffSqr)
if (!needCorrection && (!(excl & 0x1) || r2 > cSim.nonbondedCutoffSqr))
#else
if (!(excl & 0x1) || r2 > cSim.nonbondedCutoffSqr)
#endif
......
platforms/cuda/src/kernels/kCalculateLocalForces.cu
View file @ b97ff259
......@@ -123,14 +123,14 @@ void GetCalculateLocalForcesSim(gpuContext gpu)
}
__global__
#if (__CUDA_ARCH__ >= 200)
__launch_bounds__(GF1XX_LOCALFORCES_THREADS_PER_BLOCK, 1)
#elif (__CUDA_ARCH__ >= 130)
__launch_bounds__(GT2XX_LOCALFORCES_THREADS_PER_BLOCK, 1)
#else
__launch_bounds__(G8X_LOCALFORCES_THREADS_PER_BLOCK, 1)
#endif
__global__
#if (__CUDA_ARCH__ >= 200)
__launch_bounds__(GF1XX_LOCALFORCES_THREADS_PER_BLOCK, 1)
#elif (__CUDA_ARCH__ >= 130)
__launch_bounds__(GT2XX_LOCALFORCES_THREADS_PER_BLOCK, 1)
#else
__launch_bounds__(G8X_LOCALFORCES_THREADS_PER_BLOCK, 1)
#endif
void kCalculateLocalForces_kernel()
{
unsigned int pos = blockIdx.x * blockDim.x + threadIdx.x;
......@@ -461,166 +461,50 @@ void kCalculateLocalForces_kernel()
pos += blockDim.x * gridDim.x;
}
if (cSim.nonbondedMethod == NO_CUTOFF || cSim.nonbondedMethod == EWALD || cSim.nonbondedMethod == PARTICLE_MESH_EWALD)
while (pos < cSim.LJ14_offset)
{
while (pos < cSim.LJ14_offset)
unsigned int pos1 = pos - cSim.rb_dihedral_offset;
if (pos1 < cSim.LJ14s)
{
unsigned int pos1 = pos - cSim.rb_dihedral_offset;
if (pos1 < cSim.LJ14s)
{
int4 atom = cSim.pLJ14ID[pos1];
float4 LJ14 = cSim.pLJ14Parameter[pos1];
float4 a1 = cSim.pPosq[atom.x];
float4 a2 = cSim.pPosq[atom.y];
float3 d;
d.x = a1.x - a2.x;
d.y = a1.y - a2.y;
d.z = a1.z - a2.z;
float r2 = DOT3(d, d);
float inverseR = 1.0f / sqrt(r2);
float sig2 = inverseR * LJ14.y;
sig2 *= sig2;
float sig6 = sig2 * sig2 * sig2;
float dEdR = LJ14.x * (12.0f * sig6 - 6.0f) * sig6;
/* E */
energy += LJ14.x * (sig6 - 1.0f) * sig6;
energy += LJ14.z * inverseR;
dEdR += LJ14.z * inverseR;
dEdR *= inverseR * inverseR;
unsigned int offsetA = atom.x + atom.z * cSim.stride;
unsigned int offsetB = atom.y + atom.w * cSim.stride;
float4 forceA = cSim.pForce4[offsetA];
float4 forceB = cSim.pForce4[offsetB];
d.x *= dEdR;
d.y *= dEdR;
d.z *= dEdR;
forceA.x += d.x;
forceA.y += d.y;
forceA.z += d.z;
forceB.x -= d.x;
forceB.y -= d.y;
forceB.z -= d.z;
cSim.pForce4[offsetA] = forceA;
cSim.pForce4[offsetB] = forceB;
}
pos += blockDim.x * gridDim.x;
}
}
else if (cSim.nonbondedMethod == CUTOFF)
{
float LJ14_energy;
while (pos < cSim.LJ14_offset)
{
unsigned int pos1 = pos - cSim.rb_dihedral_offset;
if (pos1 < cSim.LJ14s)
{
int4 atom = cSim.pLJ14ID[pos1];
float4 LJ14 = cSim.pLJ14Parameter[pos1];
float4 a1 = cSim.pPosq[atom.x];
float4 a2 = cSim.pPosq[atom.y];
float3 d;
d.x = a1.x - a2.x;
d.y = a1.y - a2.y;
d.z = a1.z - a2.z;
float r2 = DOT3(d, d);
float inverseR = 1.0f / sqrt(r2);
float sig2 = inverseR * LJ14.y;
sig2 *= sig2;
float sig6 = sig2 * sig2 * sig2;
float dEdR = LJ14.x * (12.0f * sig6 - 6.0f) * sig6;
/* E */
LJ14_energy = LJ14.x * (sig6 - 1.0f) * sig6;
LJ14_energy += LJ14.z * (inverseR + cSim.reactionFieldK * r2 - cSim.reactionFieldC);
dEdR += LJ14.z * (inverseR - 2.0f * cSim.reactionFieldK * r2);
dEdR *= inverseR * inverseR;
if (r2 > cSim.nonbondedCutoffSqr)
{
dEdR = 0.0f;
/* E */
LJ14_energy = 0.0f;
}
/* E */
energy += LJ14_energy;
unsigned int offsetA = atom.x + atom.z * cSim.stride;
unsigned int offsetB = atom.y + atom.w * cSim.stride;
float4 forceA = cSim.pForce4[offsetA];
float4 forceB = cSim.pForce4[offsetB];
d.x *= dEdR;
d.y *= dEdR;
d.z *= dEdR;
forceA.x += d.x;
forceA.y += d.y;
forceA.z += d.z;
forceB.x -= d.x;
forceB.y -= d.y;
forceB.z -= d.z;
cSim.pForce4[offsetA] = forceA;
cSim.pForce4[offsetB] = forceB;
}
pos += blockDim.x * gridDim.x;
}
}
else if (cSim.nonbondedMethod == PERIODIC)
{
float LJ14_energy;
while (pos < cSim.LJ14_offset)
{
unsigned int pos1 = pos - cSim.rb_dihedral_offset;
if (pos1 < cSim.LJ14s)
{
int4 atom = cSim.pLJ14ID[pos1];
float4 LJ14 = cSim.pLJ14Parameter[pos1];
float4 a1 = cSim.pPosq[atom.x];
float4 a2 = cSim.pPosq[atom.y];
float3 d;
d.x = a1.x - a2.x;
d.y = a1.y - a2.y;
d.z = a1.z - a2.z;
d.x -= floor(d.x/cSim.periodicBoxSizeX+0.5f)*cSim.periodicBoxSizeX;
d.y -= floor(d.y/cSim.periodicBoxSizeY+0.5f)*cSim.periodicBoxSizeY;
d.z -= floor(d.z/cSim.periodicBoxSizeZ+0.5f)*cSim.periodicBoxSizeZ;
float r2 = DOT3(d, d);
float inverseR = 1.0f / sqrt(r2);
float sig2 = inverseR * LJ14.y;
sig2 *= sig2;
float sig6 = sig2 * sig2 * sig2;
float dEdR = LJ14.x * (12.0f * sig6 - 6.0f) * sig6;
/* E */
LJ14_energy = LJ14.x * (sig6 - 1.0f) * sig6;
LJ14_energy += LJ14.z * (inverseR + cSim.reactionFieldK * r2 - cSim.reactionFieldC);
dEdR += LJ14.z * (inverseR - 2.0f * cSim.reactionFieldK * r2);
dEdR *= inverseR * inverseR;
if (r2 > cSim.nonbondedCutoffSqr)
{
dEdR = 0.0f;
/* E */
LJ14_energy = 0.0f;
}
/* E */
energy += LJ14_energy;
unsigned int offsetA = atom.x + atom.z * cSim.stride;
unsigned int offsetB = atom.y + atom.w * cSim.stride;
float4 forceA = cSim.pForce4[offsetA];
float4 forceB = cSim.pForce4[offsetB];
d.x *= dEdR;
d.y *= dEdR;
d.z *= dEdR;
forceA.x += d.x;
forceA.y += d.y;
forceA.z += d.z;
forceB.x -= d.x;
forceB.y -= d.y;
forceB.z -= d.z;
cSim.pForce4[offsetA] = forceA;
cSim.pForce4[offsetB] = forceB;
}
pos += blockDim.x * gridDim.x;
int4 atom = cSim.pLJ14ID[pos1];
float4 LJ14 = cSim.pLJ14Parameter[pos1];
float4 a1 = cSim.pPosq[atom.x];
float4 a2 = cSim.pPosq[atom.y];
float3 d;
d.x = a1.x - a2.x;
d.y = a1.y - a2.y;
d.z = a1.z - a2.z;
float r2 = DOT3(d, d);
float inverseR = 1.0f / sqrt(r2);
float sig2 = inverseR * LJ14.y;
sig2 *= sig2;
float sig6 = sig2 * sig2 * sig2;
float dEdR = LJ14.x * (12.0f * sig6 - 6.0f) * sig6;
/* E */
energy += LJ14.x * (sig6 - 1.0f) * sig6;
energy += LJ14.z * inverseR;
dEdR += LJ14.z * inverseR;
dEdR *= inverseR * inverseR;
unsigned int offsetA = atom.x + atom.z * cSim.stride;
unsigned int offsetB = atom.y + atom.w * cSim.stride;
float4 forceA = cSim.pForce4[offsetA];
float4 forceB = cSim.pForce4[offsetB];
d.x *= dEdR;
d.y *= dEdR;
d.z *= dEdR;
forceA.x += d.x;
forceA.y += d.y;
forceA.z += d.z;
forceB.x -= d.x;
forceB.y -= d.y;
forceB.z -= d.z;
cSim.pForce4[offsetA] = forceA;
cSim.pForce4[offsetB] = forceB;
}
pos += blockDim.x * gridDim.x;
}
cSim.pEnergy[blockIdx.x * blockDim.x + threadIdx.x] += energy;
}
......
platforms/cuda/tests/TestCudaEwald.cpp
View file @ b97ff259
......@@ -79,17 +79,20 @@ void testEwaldPME(bool includeExceptions) {
nonbonded->addParticle(1.0, 1.0,0.0);
for (int i = 0; i < numParticles/2; i++)
nonbonded->addParticle(-1.0, 1.0,0.0);
if (includeExceptions) {
// Add some exclusions.
for (int i = 0; i < numParticles-1; i++)
nonbonded->addException(i, i+1, 0.0, 1.0, 0.0);
}
system.setPeriodicBoxVectors(Vec3(boxSize, 0, 0), Vec3(0, boxSize, 0), Vec3(0, 0, boxSize));
system.addForce(nonbonded);
vector<Vec3> positions(numParticles);
#include "nacl_amorph.dat"
if (includeExceptions) {
// Add some exclusions.
for (int i = 0; i < numParticles-1; i++) {
Vec3 delta = positions[i]-positions[i+1];
if (sqrt(delta.dot(delta)) < 0.5*cutoff)
nonbonded->addException(i, i+1, i%2 == 0 ? 0.0 : 0.5, 1.0, 0.0);
}
}
// (1) Check whether the Reference and Cuda platforms agree when using Ewald Method
......
platforms/cuda/tests/TestCudaNonbondedForce.cpp
View file @ b97ff259
......@@ -312,10 +312,8 @@ void testCutoff14() {
context.setPositions(positions);
state = context.getState(State::Forces | State::Energy);
const vector<Vec3>& forces2 = state.getForces();
const double krf = (1.0/(cutoff*cutoff*cutoff))*(eps-1.0)/(2.0*eps+1.0);
const double crf = (1.0/cutoff)*(3.0*eps)/(2.0*eps+1.0);
force = ONE_4PI_EPS0*q*q*(1.0/(r*r)-2.0*krf*r);
energy = ONE_4PI_EPS0*q*q*(1.0/r+krf*r*r-crf);
force = ONE_4PI_EPS0*q*q/(r*r);
energy = ONE_4PI_EPS0*q*q/r;
if (i == 3) {
force /= 1.2;
energy /= 1.2;
......
platforms/opencl/src/OpenCLKernelFactory.cpp
View file @ b97ff259
......@@ -62,6 +62,8 @@ KernelImpl* OpenCLKernelFactory::createKernelImpl(std::string name, const Platfo
return new OpenCLCalcCustomGBForceKernel(name, platform, cl, context.getSystem());
if (name == CalcCustomExternalForceKernel::Name())
return new OpenCLCalcCustomExternalForceKernel(name, platform, cl, context.getSystem());
if (name == CalcCustomHbondForceKernel::Name())
return new OpenCLCalcCustomHbondForceKernel(name, platform, cl, context.getSystem());
if (name == IntegrateVerletStepKernel::Name())
return new OpenCLIntegrateVerletStepKernel(name, platform, cl);
if (name == IntegrateLangevinStepKernel::Name())
......
platforms/opencl/src/OpenCLKernels.cpp
View file @ b97ff259
This diff is collapsed.
platforms/opencl/src/OpenCLKernels.h
View file @ b97ff259
......@@ -488,7 +488,7 @@ private:
cl::Kernel pmeConvolutionKernel;
cl::Kernel pmeInterpolateForceKernel;
std::map<std::string, std::string> pmeDefines;
double cutoffSquared, ewaldSelfEnergy;
double ewaldSelfEnergy;
static const int PmeOrder = 5;
};
......@@ -664,6 +664,55 @@ private:
cl::Kernel kernel;
};
/**
* This kernel is invoked by CustomHbondForce to calculate the forces acting on the system.
*/
class OpenCLCalcCustomHbondForceKernel : public CalcCustomHbondForceKernel {
public:
OpenCLCalcCustomHbondForceKernel(std::string name, const Platform& platform, OpenCLContext& cl, System& system) : CalcCustomHbondForceKernel(name, platform),
hasInitializedKernel(false), cl(cl), donorParams(NULL), acceptorParams(NULL), donors(NULL), acceptors(NULL),
donorBufferIndices(NULL), acceptorBufferIndices(NULL), globals(NULL), tabulatedFunctionParams(NULL), system(system) {
}
~OpenCLCalcCustomHbondForceKernel();
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the CustomHbondForce this kernel will be used for
*/
void initialize(const System& system, const CustomHbondForce& force);
/**
* Execute the kernel to calculate the forces.
*
* @param context the context in which to execute this kernel
*/
void executeForces(ContextImpl& context);
/**
* Execute the kernel to calculate the energy.
*
* @param context the context in which to execute this kernel
* @return the potential energy due to the CustomHbondForce
*/
double executeEnergy(ContextImpl& context);
private:
int numDonors, numAcceptors;
bool hasInitializedKernel;
OpenCLContext& cl;
OpenCLParameterSet* donorParams;
OpenCLParameterSet* acceptorParams;
OpenCLArray<cl_float>* globals;
OpenCLArray<mm_int4>* donors;
OpenCLArray<mm_int4>* acceptors;
OpenCLArray<mm_int4>* donorBufferIndices;
OpenCLArray<mm_int4>* acceptorBufferIndices;
OpenCLArray<mm_float4>* tabulatedFunctionParams;
std::vector<std::string> globalParamNames;
std::vector<cl_float> globalParamValues;
std::vector<OpenCLArray<mm_float4>*> tabulatedFunctions;
System& system;
cl::Kernel kernel;
};
/**
* This kernel is invoked by VerletIntegrator to take one time step.
*/
......
platforms/opencl/src/OpenCLPlatform.cpp
View file @ b97ff259
......@@ -58,6 +58,7 @@ OpenCLPlatform::OpenCLPlatform() {
registerKernelFactory(CalcGBSAOBCForceKernel::Name(), factory);
registerKernelFactory(CalcCustomGBForceKernel::Name(), factory);
registerKernelFactory(CalcCustomExternalForceKernel::Name(), factory);
registerKernelFactory(CalcCustomHbondForceKernel::Name(), factory);
registerKernelFactory(IntegrateVerletStepKernel::Name(), factory);
registerKernelFactory(IntegrateLangevinStepKernel::Name(), factory);
registerKernelFactory(IntegrateBrownianStepKernel::Name(), factory);
......
platforms/opencl/src/kernels/coulombLennardJones.cl
View file @ b97ff259
#if USE_EWALD
if (r2 < CUTOFF_SQUARED) {
bool needCorrection = isExcluded && atom1 != atom2 && atom1 < NUM_ATOMS && atom2 < NUM_ATOMS;
if (!isExcluded || needCorrection) {
const float prefactor = 138.935456f*posq1.w*posq2.w*invR;
float alphaR = EWALD_ALPHA*r;
float erfcAlphaR = erfc(alphaR);
float tempForce;
if (needCorrection) {
// Subtract off the part of this interaction that was included in the reciprocal space contribution.
bool needCorrection = isExcluded && atom1 != atom2 && atom1 < NUM_ATOMS && atom2 < NUM_ATOMS;
if (!isExcluded || needCorrection) {
const float prefactor = 138.935456f*posq1.w*posq2.w*invR;
float alphaR = EWALD_ALPHA*r;
float erfcAlphaR = erfc(alphaR);
float tempForce = 0.0f;
if (needCorrection) {
// Subtract off the part of this interaction that was included in the reciprocal space contribution.
tempForce = -prefactor*((1.0f-erfcAlphaR)-alphaR*exp(-alphaR*alphaR)*TWO_OVER_SQRT_PI);
tempEnergy += -prefactor*(1.0f-erfcAlphaR);
}
else {
tempForce = -prefactor*((1.0f-erfcAlphaR)-alphaR*exp(-alphaR*alphaR)*TWO_OVER_SQRT_PI);
tempEnergy += -prefactor*(1.0f-erfcAlphaR);
}
else if (r2 < CUTOFF_SQUARED) {
#if HAS_LENNARD_JONES
float sig = sigmaEpsilon1.x + sigmaEpsilon2.x;
float sig2 = invR*sig;
sig2 *= sig2;
float sig6 = sig2*sig2*sig2;
float eps = sigmaEpsilon1.y*sigmaEpsilon2.y;
tempForce = eps*(12.0f*sig6 - 6.0f)*sig6 + prefactor*(erfcAlphaR+alphaR*exp(-alphaR*alphaR)*TWO_OVER_SQRT_PI);
tempEnergy += eps*(sig6 - 1.0f)*sig6 + prefactor*erfcAlphaR;
float sig = sigmaEpsilon1.x + sigmaEpsilon2.x;
float sig2 = invR*sig;
sig2 *= sig2;
float sig6 = sig2*sig2*sig2;
float eps = sigmaEpsilon1.y*sigmaEpsilon2.y;
tempForce = eps*(12.0f*sig6 - 6.0f)*sig6 + prefactor*(erfcAlphaR+alphaR*exp(-alphaR*alphaR)*TWO_OVER_SQRT_PI);
tempEnergy += eps*(sig6 - 1.0f)*sig6 + prefactor*erfcAlphaR;
#else
tempForce = prefactor*(erfcAlphaR+alphaR*exp(-alphaR*alphaR)*TWO_OVER_SQRT_PI);
tempEnergy += prefactor*erfcAlphaR;
tempForce = prefactor*(erfcAlphaR+alphaR*exp(-alphaR*alphaR)*TWO_OVER_SQRT_PI);
tempEnergy += prefactor*erfcAlphaR;
#endif
}
dEdR += tempForce*invR*invR;
}
dEdR += tempForce*invR*invR;
}
#else
#ifdef USE_CUTOFF
......
platforms/opencl/src/kernels/customHbondForce.cl 0 → 100644
View file @ b97ff259
/**
* Compute the difference between two vectors, setting the fourth component to the squared magnitude.
*/
float4 delta(float4 vec1, float4 vec2) {
float4 result = (float4) (vec1.x-vec2.x, vec1.y-vec2.y, vec1.z-vec2.z, 0.0f);
result.w = result.x*result.x + result.y*result.y + result.z*result.z;
return result;
}
/**
* Compute the difference between two vectors, taking periodic boundary conditions into account
* and setting the fourth component to the squared magnitude.
*/
float4 deltaPeriodic(float4 vec1, float4 vec2) {
float4 result = (float4) (vec1.x-vec2.x, vec1.y-vec2.y, vec1.z-vec2.z, 0.0f);
#ifdef USE_PERIODIC
result.x -= floor(result.x/PERIODIC_BOX_SIZE_X+0.5f)*PERIODIC_BOX_SIZE_X;
result.y -= floor(result.y/PERIODIC_BOX_SIZE_Y+0.5f)*PERIODIC_BOX_SIZE_Y;
result.z -= floor(result.z/PERIODIC_BOX_SIZE_Z+0.5f)*PERIODIC_BOX_SIZE_Z;
#endif
result.w = result.x*result.x + result.y*result.y + result.z*result.z;
return result;
}
/**
* Compute the angle between two vectors. The w component of each vector should contain the squared magnitude.
*/
float computeAngle(float4 vec1, float4 vec2) {
float dot = vec1.x*vec2.x + vec1.y*vec2.y + vec1.z*vec2.z;
float cosine = dot/sqrt(vec1.w*vec2.w);
float angle;
if (cosine > 0.99f || cosine < -0.99f) {
// We're close to the singularity in acos(), so take the cross product and use asin() instead.
float4 cross_prod = cross(vec1, vec2);
float scale = vec1.w*vec2.w;
angle = asin(sqrt(dot(cross_prod, cross_prod)/scale));
if (cosine < 0.0f)
angle = M_PI-angle;
}
else
angle = acos(cosine);
return angle;
}
/**
* Compute hbond interactions.
*/
__kernel void computeHbonds(__global float4* forceBuffers, __global float* energyBuffer, __global float4* posq, /*__global unsigned int* exclusions,
__global unsigned int* exclusionIndices, */__global int4* donorAtoms, __global int4* acceptorAtoms, __global int4* donorBufferIndices, __global int4* acceptorBufferIndices, __local float4* posBuffer, __local float4* deltaBuffer
PARAMETER_ARGUMENTS) {
float energy = 0.0f;
unsigned int tgx = get_local_id(0) & (get_local_size(0)-1);
unsigned int tbx = get_local_id(0) - tgx;
float4 f1 = 0;
float4 f2 = 0;
for (int donorIndex = get_global_id(0); donorIndex < NUM_DONORS; donorIndex += get_global_size(0)) {
// Load information about the donor this thread will compute forces on.
int4 atoms = donorAtoms[donorIndex];
float4 d1 = posq[atoms.x];
float4 d2 = posq[atoms.y];
float4 d3 = posq[atoms.z];
float4 deltaD1D2 = delta(d1, d2);
for (int acceptorStart = 0; acceptorStart < NUM_ACCEPTORS; acceptorStart += get_local_size(0)) {
// Load the next block of acceptors into local memory.
int blockSize = min((int) get_local_size(0), NUM_ACCEPTORS-acceptorStart);
if (tgx < blockSize) {
int4 atoms2 = acceptorAtoms[acceptorStart+tgx];
float4 pos1 = posq[atoms2.x];
float4 pos2 = posq[atoms2.y];
float4 pos3 = posq[atoms2.z];
posBuffer[get_local_id(0)] = pos1;
deltaBuffer[get_local_id(0)] = delta(pos2, pos1);
}
barrier(CLK_LOCAL_MEM_FENCE);
for (int index = 0; index < blockSize; index++) {
// Compute the interaction between a donor and an acceptor.
float4 a1 = posBuffer[index];
float4 deltaD1A1 = deltaPeriodic(d1, a1);
#ifdef USE_CUTOFF
if (deltaD1A1.w < CUTOFF_SQUARED) {
#endif
// Compute variables the force can depend on.
float r = sqrt(deltaD1A1.w);
float4 deltaA2A1 = deltaBuffer[index];
float theta = computeAngle(deltaD1A1, deltaD1D2);
float psi = computeAngle(deltaD1A1, deltaA2A1);
float4 cross1 = cross(deltaA2A1, deltaD1A1);
float4 cross2 = cross(deltaD1A1, deltaD1D2);
cross1.w = cross1.x*cross1.x + cross1.y*cross1.y + cross1.z*cross1.z;
cross2.w = cross2.x*cross2.x + cross2.y*cross2.y + cross2.z*cross2.z;
float chi = computeAngle(cross1, cross2);
chi = (dot(deltaA2A1, cross2) < 0 ? -chi : chi);
COMPUTE_FORCE
#ifdef INCLUDE_R
// Apply forces based on r.
f1.xyz -= (dEdR/r)*deltaD1A1.xyz;
#endif
#ifdef INCLUDE_THETA
// Apply forces based on theta.
float4 thetaCross = cross(deltaD1D2, deltaD1A1);
float lengthThetaCross = max(length(thetaCross), 1e-6f);
float4 deltaCross0 = cross(deltaD1D2, thetaCross)*dEdTheta/(deltaD1D2.w*lengthThetaCross);
float4 deltaCross2 = -cross(deltaD1A1, thetaCross)*dEdTheta/(deltaD1A1.w*lengthThetaCross);
float4 deltaCross1 = -(deltaCross0+deltaCross2);
f1.xyz += deltaCross1.xyz;
f2.xyz += deltaCross0.xyz;
#endif
#ifdef INCLUDE_PSI
// Apply forces based on psi.
float4 psiCross = cross(deltaA2A1, deltaD1A1);
float lengthPsiCross = max(length(psiCross), 1e-6f);
deltaCross0 = cross(deltaD1A1, psiCross)*dEdPsi/(deltaD1A1.w*lengthPsiCross);
// float4 deltaCross2 = -cross(deltaA2A1, psiCross)*dEdPsi/(deltaA2A1.w*lengthPsiCross);
// float4 deltaCross1 = -(deltaCross0+deltaCross2);
f1.xyz += deltaCross0.xyz;
#endif
#ifdef INCLUDE_CHI
// Apply forces based on chi.
float4 ff;
ff.x = (-dEdChi*r)/cross1.w;
ff.y = (deltaA2A1.x*deltaD1A1.x + deltaA2A1.y*deltaD1A1.y + deltaA2A1.z*deltaD1A1.z)/deltaD1A1.w;
ff.z = (deltaD1D2.x*deltaD1A1.x + deltaD1D2.y*deltaD1A1.y + deltaD1D2.z*deltaD1A1.z)/deltaD1A1.w;
ff.w = (dEdChi*r)/cross2.w;
float4 internalF0 = ff.x*cross1;
float4 internalF3 = ff.w*cross2;
float4 s = ff.y*internalF0 - ff.z*internalF3;
f1.xyz -= s.xyz+internalF3.xyz;
f2.xyz += internalF3.xyz;
#endif
#ifdef USE_CUTOFF
}
#endif
}
}
// Write results
int4 bufferIndices = donorBufferIndices[donorIndex];
unsigned int offset1 = atoms.x+bufferIndices.x*PADDED_NUM_ATOMS;
unsigned int offset2 = atoms.y+bufferIndices.y*PADDED_NUM_ATOMS;
float4 force1 = forceBuffers[offset1];
float4 force2 = forceBuffers[offset2];
force1.xyz += f1.xyz;
force2.xyz += f2.xyz;
forceBuffers[offset1] = force1;
forceBuffers[offset2] = force2;
}
energyBuffer[get_global_id(0)] += energy;
}
platforms/opencl/src/kernels/nonbondedExceptions.cl
View file @ b97ff259
......@@ -2,21 +2,16 @@
* Compute nonbonded exceptions.
*/
__kernel void computeNonbondedExceptions(int numAtoms, int numExceptions, float cutoffSquared, float4 periodicBoxSize, __global float4* forceBuffers, __global float* energyBuffer,
__kernel void computeNonbondedExceptions(__global float4* forceBuffers, __global float* energyBuffer,
__global float4* posq, __global float4* params, __global int4* indices) {
int index = get_global_id(0);
float energy = 0.0f;
while (index < numExceptions) {
while (index < NUM_EXCEPTIONS) {
// Look up the data for this bonds.
int4 atoms = indices[index];
float4 exceptionParams = params[index];
float4 delta = posq[atoms.y]-posq[atoms.x];
#ifdef USE_PERIODIC
delta.x -= floor(delta.x/periodicBoxSize.x+0.5f)*periodicBoxSize.x;
delta.y -= floor(delta.y/periodicBoxSize.y+0.5f)*periodicBoxSize.y;
delta.z -= floor(delta.z/periodicBoxSize.z+0.5f)*periodicBoxSize.z;
#endif
// Compute the force.
......@@ -27,27 +22,16 @@ __kernel void computeNonbondedExceptions(int numAtoms, int numExceptions, float
float sig6 = sig2*sig2*sig2;
float dEdR = exceptionParams.z*(12.0f*sig6-6.0f)*sig6;
float tempEnergy = exceptionParams.z*(sig6-1.0f)*sig6;
#ifdef USE_CUTOFF
dEdR += exceptionParams.x*(invR-2.0f*REACTION_FIELD_K*r2);
tempEnergy += exceptionParams.x*(invR+REACTION_FIELD_K*r2-REACTION_FIELD_C);
#else
dEdR += exceptionParams.x*invR;
tempEnergy += exceptionParams.x*invR;
#endif
dEdR *= invR*invR;
#ifdef USE_CUTOFF
if (r2 > cutoffSquared) {
dEdR = 0.0f;
tempEnergy = 0.0f;
}
#endif
tempEnergy += exceptionParams.x*invR;
energy += tempEnergy;
delta.xyz *= dEdR;
// Record the force on each of the two atoms.
unsigned int offsetA = atoms.x+atoms.z*numAtoms;
unsigned int offsetB = atoms.y+atoms.w*numAtoms;
unsigned int offsetA = atoms.x+atoms.z*NUM_ATOMS;
unsigned int offsetB = atoms.y+atoms.w*NUM_ATOMS;
float4 forceA = forceBuffers[offsetA];
float4 forceB = forceBuffers[offsetB];
forceA.xyz -= delta.xyz;
......
platforms/opencl/tests/TestOpenCLEwald.cpp
View file @ b97ff259
......@@ -78,17 +78,20 @@ void testEwaldPME(bool includeExceptions) {
nonbonded->addParticle(1.0, 1.0,0.0);
for (int i = 0; i < numParticles/2; i++)
nonbonded->addParticle(-1.0, 1.0,0.0);
if (includeExceptions) {
// Add some exclusions.
for (int i = 0; i < numParticles-1; i++)
nonbonded->addException(i, i+1, 0.0, 1.0, 0.0);
}
system.setPeriodicBoxVectors(Vec3(boxSize, 0, 0), Vec3(0, boxSize, 0), Vec3(0, 0, boxSize));
system.addForce(nonbonded);
vector<Vec3> positions(numParticles);
#include "nacl_amorph.dat"
if (includeExceptions) {
// Add some exclusions.
for (int i = 0; i < numParticles-1; i++) {
Vec3 delta = positions[i]-positions[i+1];
if (sqrt(delta.dot(delta)) < 0.5*cutoff)
nonbonded->addException(i, i+1, i%2 == 0 ? 0.0 : 0.5, 1.0, 0.0);
}
}
// (1) Check whether the Reference and OpenCL platforms agree when using Ewald Method
......
platforms/opencl/tests/TestOpenCLNonbondedForce.cpp
View file @ b97ff259
......@@ -307,10 +307,8 @@ void testCutoff14() {
context.setPositions(positions);
state = context.getState(State::Forces | State::Energy);
const vector<Vec3>& forces2 = state.getForces();
const double krf = (1.0/(cutoff*cutoff*cutoff))*(eps-1.0)/(2.0*eps+1.0);
const double crf = (1.0/cutoff)*(3.0*eps)/(2.0*eps+1.0);
force = ONE_4PI_EPS0*q*q*(1.0/(r*r)-2.0*krf*r);
energy = ONE_4PI_EPS0*q*q*(1.0/r+krf*r*r-crf);
force = ONE_4PI_EPS0*q*q/(r*r);
energy = ONE_4PI_EPS0*q*q/r;
if (i == 3) {
force /= 1.2;
energy /= 1.2;
......
platforms/reference/src/ReferenceKernels.cpp
View file @ b97ff259
......@@ -694,7 +694,7 @@ void ReferenceCalcNonbondedForceKernel::executeForces(ContextImpl& context) {
computeNeighborListVoxelHash(*neighborList, numParticles, posData, exclusions, (periodic || ewald || pme) ? periodicBoxSize : NULL, nonbondedCutoff, 0.0);
clj.setUseCutoff(nonbondedCutoff, *neighborList, rfDielectric);
}
if (periodic||ewald||pme)
if (periodic || ewald || pme)
clj.setPeriodic(periodicBoxSize);
if (ewald)
clj.setUseEwald(ewaldAlpha, kmax[0], kmax[1], kmax[2]);
......@@ -703,8 +703,6 @@ void ReferenceCalcNonbondedForceKernel::executeForces(ContextImpl& context) {
clj.calculatePairIxn(numParticles, posData, particleParamArray, exclusionArray, 0, forceData, 0, 0);
ReferenceBondForce refBondForce;
ReferenceLJCoulomb14 nonbonded14;
if (nonbondedMethod == CutoffNonPeriodic || nonbondedMethod == CutoffPeriodic)
nonbonded14.setUseCutoff(nonbondedCutoff, rfDielectric);
refBondForce.calculateForce(num14, bonded14IndexArray, posData, bonded14ParamArray, forceData, 0, 0, 0, nonbonded14);
}
......@@ -729,8 +727,6 @@ double ReferenceCalcNonbondedForceKernel::executeEnergy(ContextImpl& context) {
clj.calculatePairIxn(numParticles, posData, particleParamArray, exclusionArray, 0, forceData, 0, &energy);
ReferenceBondForce refBondForce;
ReferenceLJCoulomb14 nonbonded14;
if (nonbondedMethod == CutoffNonPeriodic || nonbondedMethod == CutoffPeriodic)
nonbonded14.setUseCutoff(nonbondedCutoff, rfDielectric);
RealOpenMM* energyArray = new RealOpenMM[num14];
for (int i = 0; i < num14; ++i)
energyArray[i] = 0;
......
platforms/reference/src/SimTKReference/ReferenceLJCoulomb14.cpp
View file @ b97ff259
......@@ -37,7 +37,7 @@
--------------------------------------------------------------------------------------- */
ReferenceLJCoulomb14::ReferenceLJCoulomb14( ) : cutoff(false) {
ReferenceLJCoulomb14::ReferenceLJCoulomb14( ) {
// ---------------------------------------------------------------------------------------
......@@ -63,27 +63,6 @@ ReferenceLJCoulomb14::~ReferenceLJCoulomb14( ){
}
/**---------------------------------------------------------------------------------------
Set the force to use a cutoff.
@param distance the cutoff distance
@param solventDielectric the dielectric constant of the bulk solvent
@return ReferenceForce::DefaultReturn
--------------------------------------------------------------------------------------- */
int ReferenceLJCoulomb14::setUseCutoff( RealOpenMM distance, RealOpenMM solventDielectric ) {
cutoff = true;
cutoffDistance = distance;
krf = pow(cutoffDistance, -3.0)*(solventDielectric-1.0)/(2.0*solventDielectric+1.0);
crf = (1.0f/cutoffDistance)*(3.0*solventDielectric)/(2.0*solventDielectric+1.0);
return ReferenceForce::DefaultReturn;
}
/**---------------------------------------------------------------------------------------
Calculate LJ 1-4 ixn
......@@ -147,8 +126,6 @@ int ReferenceLJCoulomb14::calculateBondIxn( int* atomIndices, RealOpenMM** atomC
int atomBIndex = atomIndices[1];
ReferenceForce::getDeltaR( atomCoordinates[atomBIndex], atomCoordinates[atomAIndex], deltaR[0] );
if (cutoff && deltaR[0][ReferenceForce::RIndex] > cutoffDistance)
return ReferenceForce::DefaultReturn;
RealOpenMM r2 = deltaR[0][ReferenceForce::R2Index];
RealOpenMM inverseR = one/(deltaR[0][ReferenceForce::RIndex]);
RealOpenMM sig2 = inverseR*parameters[0];
......@@ -156,10 +133,7 @@ int ReferenceLJCoulomb14::calculateBondIxn( int* atomIndices, RealOpenMM** atomC
RealOpenMM sig6 = sig2*sig2*sig2;
RealOpenMM dEdR = parameters[1]*( twelve*sig6 - six )*sig6;
if (cutoff)
dEdR += (RealOpenMM) (ONE_4PI_EPS0*parameters[2]*(inverseR-2.0f*krf*r2));
else
dEdR += (RealOpenMM) (ONE_4PI_EPS0*parameters[2]*inverseR);
dEdR += (RealOpenMM) (ONE_4PI_EPS0*parameters[2]*inverseR);
dEdR *= inverseR*inverseR;
// accumulate forces
......@@ -171,10 +145,7 @@ int ReferenceLJCoulomb14::calculateBondIxn( int* atomIndices, RealOpenMM** atomC
}
RealOpenMM energy = parameters[1]*( sig6 - one )*sig6;
if (cutoff)
energy += (RealOpenMM) (ONE_4PI_EPS0*parameters[2]*(inverseR+krf*r2-crf));
else
energy += (RealOpenMM) (ONE_4PI_EPS0*parameters[2]*inverseR);
energy += (RealOpenMM) (ONE_4PI_EPS0*parameters[2]*inverseR);
// accumulate energies
......
platforms/reference/src/SimTKReference/ReferenceLJCoulomb14.h
View file @ b97ff259
......@@ -31,12 +31,6 @@
class ReferenceLJCoulomb14 : public ReferenceBondIxn {
private:
bool cutoff;
RealOpenMM cutoffDistance;
RealOpenMM krf, crf;
public:
/**---------------------------------------------------------------------------------------
......@@ -55,19 +49,6 @@ class ReferenceLJCoulomb14 : public ReferenceBondIxn {
~ReferenceLJCoulomb14( );
/**---------------------------------------------------------------------------------------
Set the force to use a cutoff.
@param distance the cutoff distance
@param solventDielectric the dielectric constant of the bulk solvent
@return ReferenceForce::DefaultReturn
--------------------------------------------------------------------------------------- */
int setUseCutoff( RealOpenMM distance, RealOpenMM solventDielectric );
/**---------------------------------------------------------------------------------------
Calculate Ryckaert-Bellemans bond ixn
......
platforms/reference/src/SimTKReference/ReferenceLJCoulombIxn.cpp
View file @ b97ff259
......@@ -417,10 +417,8 @@ int ReferenceLJCoulombIxn::calculateEwaldIxn( int numberOfAtoms, RealOpenMM** at
int jj = exclusions[i][j];
RealOpenMM deltaR[2][ReferenceForce::LastDeltaRIndex];
ReferenceForce::getDeltaRPeriodic( atomCoordinates[jj], atomCoordinates[ii], periodicBoxSize, deltaR[0] );
ReferenceForce::getDeltaR( atomCoordinates[jj], atomCoordinates[ii], deltaR[0] );
RealOpenMM r = deltaR[0][ReferenceForce::RIndex];
if (r >= cutoffDistance)
continue;
RealOpenMM inverseR = one/(deltaR[0][ReferenceForce::RIndex]);
RealOpenMM alphaR = alphaEwald * r;
RealOpenMM dEdR = (RealOpenMM) (ONE_4PI_EPS0 * atomParameters[ii][QIndex] * atomParameters[jj][QIndex] * inverseR * inverseR * inverseR);
......
platforms/reference/tests/TestReferenceNonbondedForce.cpp
View file @ b97ff259
......@@ -302,10 +302,8 @@ void testCutoff14() {
context.setPositions(positions);
state = context.getState(State::Forces | State::Energy);
const vector<Vec3>& forces2 = state.getForces();
const double krf = (1.0/(cutoff*cutoff*cutoff))*(eps-1.0)/(2.0*eps+1.0);
const double crf = (1.0/cutoff)*(3.0*eps)/(2.0*eps+1.0);
force = ONE_4PI_EPS0*q*q*(1.0/(r*r)-2.0*krf*r);
energy = ONE_4PI_EPS0*q*q*(1.0/r+krf*r*r-crf);
force = ONE_4PI_EPS0*q*q/(r*r);
energy = ONE_4PI_EPS0*q*q/r;
if (i == 3) {
force /= 1.2;
energy /= 1.2;
......
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