Skip to content

GitLab

Commit 95dc4dc1 authored by Peter Eastman's avatar Peter Eastman
Browse files

Fixed errors

parent 122c84d2
Hide whitespace changes
Inline Side-by-side
Showing with 17 additions and 17 deletions
platforms/cuda/include/CudaKernels.h
View file @ 95dc4dc1
...@@ -816,6 +816,7 @@ private: ...@@ -816,6 +816,7 @@ private:
int maxTiles; int maxTiles;
CudaContext& cu; CudaContext& cu;
ForceInfo* info; ForceInfo* info;
CudaArray charges;
CudaArray params; CudaArray params;
CudaArray bornSum; CudaArray bornSum;
CudaArray bornRadii; CudaArray bornRadii;
......
platforms/cuda/src/CudaKernels.cpp
View file @ 95dc4dc1
...@@ -1610,6 +1610,8 @@ void CudaCalcNonbondedForceKernel::initialize(const System& system, const Nonbon ...@@ -1610,6 +1610,8 @@ void CudaCalcNonbondedForceKernel::initialize(const System& system, const Nonbon
bool useCutoff = (nonbondedMethod != NoCutoff); bool useCutoff = (nonbondedMethod != NoCutoff);
bool usePeriodic = (nonbondedMethod != NoCutoff && nonbondedMethod != CutoffNonPeriodic); bool usePeriodic = (nonbondedMethod != NoCutoff && nonbondedMethod != CutoffNonPeriodic);
doLJPME = (nonbondedMethod == LJPME && hasLJ); doLJPME = (nonbondedMethod == LJPME && hasLJ);
if (hasCoulomb)
usePosqCharges = cu.requestPosqCharges();
map<string, string> defines; map<string, string> defines;
defines["HAS_COULOMB"] = (hasCoulomb ? "1" : "0"); defines["HAS_COULOMB"] = (hasCoulomb ? "1" : "0");
defines["HAS_LENNARD_JONES"] = (hasLJ ? "1" : "0"); defines["HAS_LENNARD_JONES"] = (hasLJ ? "1" : "0");
...@@ -1933,7 +1935,6 @@ void CudaCalcNonbondedForceKernel::initialize(const System& system, const Nonbon ...@@ -1933,7 +1935,6 @@ void CudaCalcNonbondedForceKernel::initialize(const System& system, const Nonbon
string source = cu.replaceStrings(CudaKernelSources::coulombLennardJones, defines); string source = cu.replaceStrings(CudaKernelSources::coulombLennardJones, defines);
charges.initialize(cu, cu.getPaddedNumAtoms(), cu.getUseDoublePrecision() ? sizeof(double) : sizeof(float), "charges"); charges.initialize(cu, cu.getPaddedNumAtoms(), cu.getUseDoublePrecision() ? sizeof(double) : sizeof(float), "charges");
if (hasCoulomb) { if (hasCoulomb) {
usePosqCharges = cu.requestPosqCharges();
map<string, string> replacements; map<string, string> replacements;
if (usePosqCharges) { if (usePosqCharges) {
cu.setCharges(chargeVec); cu.setCharges(chargeVec);
...@@ -1965,9 +1966,6 @@ void CudaCalcNonbondedForceKernel::initialize(const System& system, const Nonbon ...@@ -1965,9 +1966,6 @@ void CudaCalcNonbondedForceKernel::initialize(const System& system, const Nonbon
cu.getNonbondedUtilities().addParameter(CudaNonbondedUtilities::ParameterInfo(prefix+"sigmaEpsilon", "float", 2, sizeof(float2), sigmaEpsilon.getDevicePointer())); cu.getNonbondedUtilities().addParameter(CudaNonbondedUtilities::ParameterInfo(prefix+"sigmaEpsilon", "float", 2, sizeof(float2), sigmaEpsilon.getDevicePointer()));
} }
cu.getNonbondedUtilities().addInteraction(useCutoff, usePeriodic, true, force.getCutoffDistance(), exclusionList, source, force.getForceGroup(), true); cu.getNonbondedUtilities().addInteraction(useCutoff, usePeriodic, true, force.getCutoffDistance(), exclusionList, source, force.getForceGroup(), true);
if (hasLJ)
cu.getNonbondedUtilities().addParameter(CudaNonbondedUtilities::ParameterInfo("sigmaEpsilon", "float", 2,
sizeof(float2), sigmaEpsilon.getDevicePointer()));
// Initialize the exceptions. // Initialize the exceptions.
...@@ -2873,7 +2871,7 @@ void CudaCalcGBSAOBCForceKernel::initialize(const System& system, const GBSAOBCF ...@@ -2873,7 +2871,7 @@ void CudaCalcGBSAOBCForceKernel::initialize(const System& system, const GBSAOBCF
string prefix = "obc"+cu.intToString(forceIndex)+"_"; string prefix = "obc"+cu.intToString(forceIndex)+"_";
CudaNonbondedUtilities& nb = cu.getNonbondedUtilities(); CudaNonbondedUtilities& nb = cu.getNonbondedUtilities();
params.initialize<float2>(cu, cu.getPaddedNumAtoms(), "gbsaObcParams"); params.initialize<float2>(cu, cu.getPaddedNumAtoms(), "gbsaObcParams");
int elementSize = (cu.getUseDoublePrecision() ? sizeof(cl_double) : sizeof(cl_float)); int elementSize = (cu.getUseDoublePrecision() ? sizeof(double) : sizeof(float));
charges.initialize(cu, cu.getPaddedNumAtoms(), elementSize, "gbsaObcCharges"); charges.initialize(cu, cu.getPaddedNumAtoms(), elementSize, "gbsaObcCharges");
bornRadii.initialize(cu, cu.getPaddedNumAtoms(), elementSize, "bornRadii"); bornRadii.initialize(cu, cu.getPaddedNumAtoms(), elementSize, "bornRadii");
obcChain.initialize(cu, cu.getPaddedNumAtoms(), elementSize, "obcChain"); obcChain.initialize(cu, cu.getPaddedNumAtoms(), elementSize, "obcChain");
......
platforms/cuda/src/kernels/gbsaObc1.cu
View file @ 95dc4dc1
...@@ -441,8 +441,8 @@ extern "C" __global__ void computeGBSAForce1(unsigned long long* __restrict__ fo ...@@ -441,8 +441,8 @@ extern "C" __global__ void computeGBSAForce1(unsigned long long* __restrict__ fo
localData[threadIdx.x].bornRadius = bornRadius1; localData[threadIdx.x].bornRadius = bornRadius1;
for (unsigned int j = 0; j < TILE_SIZE; j++) { for (unsigned int j = 0; j < TILE_SIZE; j++) {
if (atom1 < NUM_ATOMS && y*TILE_SIZE+j < NUM_ATOMS) { if (atom1 < NUM_ATOMS && y*TILE_SIZE+j < NUM_ATOMS) {
real4 pos2 = make_real3(localData[tbx+j].x, localData[tbx+j].y, localData[tbx+j].z); real3 pos2 = make_real3(localData[tbx+j].x, localData[tbx+j].y, localData[tbx+j].z);
real charge1 = localData[tbx+j].q; real charge2 = localData[tbx+j].q;
real3 delta = make_real3(pos2.x-posq1.x, pos2.y-posq1.y, pos2.z-posq1.z); real3 delta = make_real3(pos2.x-posq1.x, pos2.y-posq1.y, pos2.z-posq1.z);
#ifdef USE_PERIODIC #ifdef USE_PERIODIC
APPLY_PERIODIC_TO_DELTA(delta) APPLY_PERIODIC_TO_DELTA(delta)
...@@ -498,8 +498,8 @@ extern "C" __global__ void computeGBSAForce1(unsigned long long* __restrict__ fo ...@@ -498,8 +498,8 @@ extern "C" __global__ void computeGBSAForce1(unsigned long long* __restrict__ fo
unsigned int tj = tgx; unsigned int tj = tgx;
for (j = 0; j < TILE_SIZE; j++) { for (j = 0; j < TILE_SIZE; j++) {
if (atom1 < NUM_ATOMS && y*TILE_SIZE+tj < NUM_ATOMS) { if (atom1 < NUM_ATOMS && y*TILE_SIZE+tj < NUM_ATOMS) {
real4 pos2 = make_real3(localData[tbx+tj].x, localData[tbx+tj].y, localData[tbx+tj].z); real3 pos2 = make_real3(localData[tbx+tj].x, localData[tbx+tj].y, localData[tbx+tj].z);
real charge1 = localData[tbx+tj].q; real charge2 = localData[tbx+tj].q;
real3 delta = make_real3(pos2.x-posq1.x, pos2.y-posq1.y, pos2.z-posq1.z); real3 delta = make_real3(pos2.x-posq1.x, pos2.y-posq1.y, pos2.z-posq1.z);
#ifdef USE_PERIODIC #ifdef USE_PERIODIC
APPLY_PERIODIC_TO_DELTA(delta) APPLY_PERIODIC_TO_DELTA(delta)
...@@ -654,8 +654,8 @@ extern "C" __global__ void computeGBSAForce1(unsigned long long* __restrict__ fo ...@@ -654,8 +654,8 @@ extern "C" __global__ void computeGBSAForce1(unsigned long long* __restrict__ fo
for (j = 0; j < TILE_SIZE; j++) { for (j = 0; j < TILE_SIZE; j++) {
int atom2 = atomIndices[tbx+tj]; int atom2 = atomIndices[tbx+tj];
if (atom1 < NUM_ATOMS && atom2 < NUM_ATOMS) { if (atom1 < NUM_ATOMS && atom2 < NUM_ATOMS) {
real4 pos2 = make_real3(localData[tbx+tj].x, localData[tbx+tj].y, localData[tbx+tj].z); real3 pos2 = make_real3(localData[tbx+tj].x, localData[tbx+tj].y, localData[tbx+tj].z);
real charge1 = localData[tbx+tj].q; real charge2 = localData[tbx+tj].q;
real3 delta = make_real3(pos2.x-posq1.x, pos2.y-posq1.y, pos2.z-posq1.z); real3 delta = make_real3(pos2.x-posq1.x, pos2.y-posq1.y, pos2.z-posq1.z);
real r2 = delta.x*delta.x + delta.y*delta.y + delta.z*delta.z; real r2 = delta.x*delta.x + delta.y*delta.y + delta.z*delta.z;
if (r2 < CUTOFF_SQUARED) { if (r2 < CUTOFF_SQUARED) {
...@@ -700,8 +700,8 @@ extern "C" __global__ void computeGBSAForce1(unsigned long long* __restrict__ fo ...@@ -700,8 +700,8 @@ extern "C" __global__ void computeGBSAForce1(unsigned long long* __restrict__ fo
for (j = 0; j < TILE_SIZE; j++) { for (j = 0; j < TILE_SIZE; j++) {
int atom2 = atomIndices[tbx+tj]; int atom2 = atomIndices[tbx+tj];
if (atom1 < NUM_ATOMS && atom2 < NUM_ATOMS) { if (atom1 < NUM_ATOMS && atom2 < NUM_ATOMS) {
real4 pos2 = make_real3(localData[tbx+tj].x, localData[tbx+tj].y, localData[tbx+tj].z); real3 pos2 = make_real3(localData[tbx+tj].x, localData[tbx+tj].y, localData[tbx+tj].z);
real charge1 = localData[tbx+tj].q; real charge2 = localData[tbx+tj].q;
real3 delta = make_real3(pos2.x-posq1.x, pos2.y-posq1.y, pos2.z-posq1.z); real3 delta = make_real3(pos2.x-posq1.x, pos2.y-posq1.y, pos2.z-posq1.z);
#ifdef USE_PERIODIC #ifdef USE_PERIODIC
APPLY_PERIODIC_TO_DELTA(delta) APPLY_PERIODIC_TO_DELTA(delta)
......
platforms/cuda/src/kernels/gbsaObc2.cu
View file @ 95dc4dc1
...@@ -16,8 +16,8 @@ ...@@ -16,8 +16,8 @@
real t2I = (l_ij2I-u_ij2I); real t2I = (l_ij2I-u_ij2I);
real term1 = (0.5f*(0.25f+OBC_PARAMS2.y*OBC_PARAMS2.y*invRSquaredOver4)*t2J + t1J*invRSquaredOver4)*invR; real term1 = (0.5f*(0.25f+OBC_PARAMS2.y*OBC_PARAMS2.y*invRSquaredOver4)*t2J + t1J*invRSquaredOver4)*invR;
real term2 = (0.5f*(0.25f+OBC_PARAMS1.y*OBC_PARAMS1.y*invRSquaredOver4)*t2I + t1I*invRSquaredOver4)*invR; real term2 = (0.5f*(0.25f+OBC_PARAMS1.y*OBC_PARAMS1.y*invRSquaredOver4)*t2I + t1I*invRSquaredOver4)*invR;
real tempdEdR = (OBC_PARAMS1.x < rScaledRadiusJ ? bornForce1*term1/0x100000000 : 0); real tempdEdR = (OBC_PARAMS1.x < rScaledRadiusJ ? BORN_FORCE1*term1/0x100000000 : 0);
tempdEdR += (OBC_PARAMS2.x < rScaledRadiusI ? bornForce2*term2/0x100000000 : 0); tempdEdR += (OBC_PARAMS2.x < rScaledRadiusI ? BORN_FORCE2*term2/0x100000000 : 0);
#ifdef USE_CUTOFF #ifdef USE_CUTOFF
unsigned int includeInteraction = (atom1 < NUM_ATOMS && atom2 < NUM_ATOMS && atom1 != atom2 && r2 < CUTOFF_SQUARED); unsigned int includeInteraction = (atom1 < NUM_ATOMS && atom2 < NUM_ATOMS && atom1 != atom2 && r2 < CUTOFF_SQUARED);
#else #else
......
platforms/opencl/src/OpenCLKernels.cpp
View file @ 95dc4dc1
...@@ -1602,6 +1602,8 @@ void OpenCLCalcNonbondedForceKernel::initialize(const System& system, const Nonb ...@@ -1602,6 +1602,8 @@ void OpenCLCalcNonbondedForceKernel::initialize(const System& system, const Nonb
bool useCutoff = (nonbondedMethod != NoCutoff); bool useCutoff = (nonbondedMethod != NoCutoff);
bool usePeriodic = (nonbondedMethod != NoCutoff && nonbondedMethod != CutoffNonPeriodic); bool usePeriodic = (nonbondedMethod != NoCutoff && nonbondedMethod != CutoffNonPeriodic);
doLJPME = (nonbondedMethod == LJPME && hasLJ); doLJPME = (nonbondedMethod == LJPME && hasLJ);
if (hasCoulomb)
usePosqCharges = cl.requestPosqCharges();
map<string, string> defines; map<string, string> defines;
defines["HAS_COULOMB"] = (hasCoulomb ? "1" : "0"); defines["HAS_COULOMB"] = (hasCoulomb ? "1" : "0");
defines["HAS_LENNARD_JONES"] = (hasLJ ? "1" : "0"); defines["HAS_LENNARD_JONES"] = (hasLJ ? "1" : "0");
...@@ -1864,7 +1866,6 @@ void OpenCLCalcNonbondedForceKernel::initialize(const System& system, const Nonb ...@@ -1864,7 +1866,6 @@ void OpenCLCalcNonbondedForceKernel::initialize(const System& system, const Nonb
string source = cl.replaceStrings(OpenCLKernelSources::coulombLennardJones, defines); string source = cl.replaceStrings(OpenCLKernelSources::coulombLennardJones, defines);
charges.initialize(cl, cl.getPaddedNumAtoms(), cl.getUseDoublePrecision() ? sizeof(double) : sizeof(float), "charges"); charges.initialize(cl, cl.getPaddedNumAtoms(), cl.getUseDoublePrecision() ? sizeof(double) : sizeof(float), "charges");
if (hasCoulomb) { if (hasCoulomb) {
usePosqCharges = cl.requestPosqCharges();
map<string, string> replacements; map<string, string> replacements;
if (usePosqCharges) { if (usePosqCharges) {
cl.setCharges(chargeVec); cl.setCharges(chargeVec);
......
tests/TestNonbondedForce.h
View file @ 95dc4dc1
...@@ -6,7 +6,7 @@ ...@@ -6,7 +6,7 @@
* Biological Structures at Stanford, funded under the NIH Roadmap for * * Biological Structures at Stanford, funded under the NIH Roadmap for *
* Medical Research, grant U54 GM072970. See https://simtk.org. * * Medical Research, grant U54 GM072970. See https://simtk.org. *
* * * *
* Portions copyright (c) 2008-2015 Stanford University and the Authors. * * Portions copyright (c) 2008-2018 Stanford University and the Authors. *
* Authors: Peter Eastman * * Authors: Peter Eastman *
* Contributors: * * Contributors: *
* * * *
......
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