Skip to content

GitLab

Commit 4c17f597 authored by Peter Eastman's avatar Peter Eastman
Browse files

Fixed bug accessing parameters for CustomHbondForce

parent 111844a2
Hide whitespace changes
Inline Side-by-side
Showing with 36 additions and 8 deletions
platforms/cuda/src/CudaKernels.cpp
View file @ 4c17f597
...@@ -4521,12 +4521,12 @@ void CudaCalcCustomHbondForceKernel::initialize(const System& system, const Cust ...@@ -4521,12 +4521,12 @@ void CudaCalcCustomHbondForceKernel::initialize(const System& system, const Cust
for (int i = 0; i < (int) donorParams->getBuffers().size(); i++) { for (int i = 0; i < (int) donorParams->getBuffers().size(); i++) {
CudaNonbondedUtilities::ParameterInfo& buffer = donorParams->getBuffers()[i]; CudaNonbondedUtilities::ParameterInfo& buffer = donorParams->getBuffers()[i];
extraArgs << ", const "+buffer.getType()+"* __restrict__ donor"+buffer.getName(); extraArgs << ", const "+buffer.getType()+"* __restrict__ donor"+buffer.getName();
addDonorAndAcceptorCode(computeDonor, computeAcceptor, buffer.getType()+" donorParams"+cu.intToString(i+1)+" = donor"+buffer.getName()+"[index];\n"); addDonorAndAcceptorCode(computeDonor, computeAcceptor, buffer.getType()+" donorParams"+cu.intToString(i+1)+" = donor"+buffer.getName()+"[donorIndex];\n");
} }
for (int i = 0; i < (int) acceptorParams->getBuffers().size(); i++) { for (int i = 0; i < (int) acceptorParams->getBuffers().size(); i++) {
CudaNonbondedUtilities::ParameterInfo& buffer = acceptorParams->getBuffers()[i]; CudaNonbondedUtilities::ParameterInfo& buffer = acceptorParams->getBuffers()[i];
extraArgs << ", const "+buffer.getType()+"* __restrict__ acceptor"+buffer.getName(); extraArgs << ", const "+buffer.getType()+"* __restrict__ acceptor"+buffer.getName();
addDonorAndAcceptorCode(computeDonor, computeAcceptor, buffer.getType()+" acceptorParams"+cu.intToString(i+1)+" = acceptor"+buffer.getName()+"[index];\n"); addDonorAndAcceptorCode(computeDonor, computeAcceptor, buffer.getType()+" acceptorParams"+cu.intToString(i+1)+" = acceptor"+buffer.getName()+"[acceptorIndex];\n");
} }
// Now evaluate the expressions. // Now evaluate the expressions.
......
platforms/cuda/src/kernels/customHbondForce.cu
View file @ 4c17f597
...@@ -105,8 +105,8 @@ extern "C" __global__ void computeDonorForces(unsigned long long* __restrict__ f ...@@ -105,8 +105,8 @@ extern "C" __global__ void computeDonorForces(unsigned long long* __restrict__ f
__syncthreads(); __syncthreads();
if (donorIndex < NUM_DONORS) { if (donorIndex < NUM_DONORS) {
for (int index = 0; index < blockSize; index++) { for (int index = 0; index < blockSize; index++) {
#ifdef USE_EXCLUSIONS
int acceptorIndex = acceptorStart+index; int acceptorIndex = acceptorStart+index;
#ifdef USE_EXCLUSIONS
if (acceptorIndex == exclusionIndices.x || acceptorIndex == exclusionIndices.y || acceptorIndex == exclusionIndices.z || acceptorIndex == exclusionIndices.w) if (acceptorIndex == exclusionIndices.x || acceptorIndex == exclusionIndices.y || acceptorIndex == exclusionIndices.z || acceptorIndex == exclusionIndices.w)
continue; continue;
#endif #endif
...@@ -193,8 +193,8 @@ extern "C" __global__ void computeAcceptorForces(unsigned long long* __restrict_ ...@@ -193,8 +193,8 @@ extern "C" __global__ void computeAcceptorForces(unsigned long long* __restrict_
__syncthreads(); __syncthreads();
if (acceptorIndex < NUM_ACCEPTORS) { if (acceptorIndex < NUM_ACCEPTORS) {
for (int index = 0; index < blockSize; index++) { for (int index = 0; index < blockSize; index++) {
#ifdef USE_EXCLUSIONS
int donorIndex = donorStart+index; int donorIndex = donorStart+index;
#ifdef USE_EXCLUSIONS
if (donorIndex == exclusionIndices.x || donorIndex == exclusionIndices.y || donorIndex == exclusionIndices.z || donorIndex == exclusionIndices.w) if (donorIndex == exclusionIndices.x || donorIndex == exclusionIndices.y || donorIndex == exclusionIndices.z || donorIndex == exclusionIndices.w)
continue; continue;
#endif #endif
......
platforms/opencl/src/OpenCLKernels.cpp
View file @ 4c17f597
...@@ -4798,12 +4798,12 @@ void OpenCLCalcCustomHbondForceKernel::initialize(const System& system, const Cu ...@@ -4798,12 +4798,12 @@ void OpenCLCalcCustomHbondForceKernel::initialize(const System& system, const Cu
for (int i = 0; i < (int) donorParams->getBuffers().size(); i++) { for (int i = 0; i < (int) donorParams->getBuffers().size(); i++) {
const OpenCLNonbondedUtilities::ParameterInfo& buffer = donorParams->getBuffers()[i]; const OpenCLNonbondedUtilities::ParameterInfo& buffer = donorParams->getBuffers()[i];
extraArgs << ", __global const "+buffer.getType()+"* restrict donor"+buffer.getName(); extraArgs << ", __global const "+buffer.getType()+"* restrict donor"+buffer.getName();
addDonorAndAcceptorCode(computeDonor, computeAcceptor, buffer.getType()+" donorParams"+cl.intToString(i+1)+" = donor"+buffer.getName()+"[index];\n"); addDonorAndAcceptorCode(computeDonor, computeAcceptor, buffer.getType()+" donorParams"+cl.intToString(i+1)+" = donor"+buffer.getName()+"[donorIndex];\n");
} }
for (int i = 0; i < (int) acceptorParams->getBuffers().size(); i++) { for (int i = 0; i < (int) acceptorParams->getBuffers().size(); i++) {
const OpenCLNonbondedUtilities::ParameterInfo& buffer = acceptorParams->getBuffers()[i]; const OpenCLNonbondedUtilities::ParameterInfo& buffer = acceptorParams->getBuffers()[i];
extraArgs << ", __global const "+buffer.getType()+"* restrict acceptor"+buffer.getName(); extraArgs << ", __global const "+buffer.getType()+"* restrict acceptor"+buffer.getName();
addDonorAndAcceptorCode(computeDonor, computeAcceptor, buffer.getType()+" acceptorParams"+cl.intToString(i+1)+" = acceptor"+buffer.getName()+"[index];\n"); addDonorAndAcceptorCode(computeDonor, computeAcceptor, buffer.getType()+" acceptorParams"+cl.intToString(i+1)+" = acceptor"+buffer.getName()+"[acceptorIndex];\n");
} }
// Now evaluate the expressions. // Now evaluate the expressions.
......
platforms/opencl/src/kernels/customHbondForce.cl
View file @ 4c17f597
...@@ -91,8 +91,8 @@ __kernel void computeDonorForces(__global real4* restrict forceBuffers, __global ...@@ -91,8 +91,8 @@ __kernel void computeDonorForces(__global real4* restrict forceBuffers, __global
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
if (donorIndex < NUM_DONORS) { if (donorIndex < NUM_DONORS) {
for (int index = 0; index < blockSize; index++) { for (int index = 0; index < blockSize; index++) {
#ifdef USE_EXCLUSIONS
int acceptorIndex = acceptorStart+index; int acceptorIndex = acceptorStart+index;
#ifdef USE_EXCLUSIONS
if (acceptorIndex == exclusionIndices.x || acceptorIndex == exclusionIndices.y || acceptorIndex == exclusionIndices.z || acceptorIndex == exclusionIndices.w) if (acceptorIndex == exclusionIndices.x || acceptorIndex == exclusionIndices.y || acceptorIndex == exclusionIndices.z || acceptorIndex == exclusionIndices.w)
continue; continue;
#endif #endif
...@@ -179,8 +179,8 @@ __kernel void computeAcceptorForces(__global real4* restrict forceBuffers, __glo ...@@ -179,8 +179,8 @@ __kernel void computeAcceptorForces(__global real4* restrict forceBuffers, __glo
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
if (acceptorIndex < NUM_ACCEPTORS) { if (acceptorIndex < NUM_ACCEPTORS) {
for (int index = 0; index < blockSize; index++) { for (int index = 0; index < blockSize; index++) {
#ifdef USE_EXCLUSIONS
int donorIndex = donorStart+index; int donorIndex = donorStart+index;
#ifdef USE_EXCLUSIONS
if (donorIndex == exclusionIndices.x || donorIndex == exclusionIndices.y || donorIndex == exclusionIndices.z || donorIndex == exclusionIndices.w) if (donorIndex == exclusionIndices.x || donorIndex == exclusionIndices.y || donorIndex == exclusionIndices.z || donorIndex == exclusionIndices.w)
continue; continue;
#endif #endif
......
tests/TestCustomHbondForce.h
View file @ 4c17f597
...@@ -244,6 +244,33 @@ void testIllegalVariable() { ...@@ -244,6 +244,33 @@ void testIllegalVariable() {
ASSERT(threwException); ASSERT(threwException);
} }
void testParameters() {
System system;
system.addParticle(1.0);
system.addParticle(1.0);
system.addParticle(1.0);
VerletIntegrator integrator(0.01);
CustomHbondForce* custom = new CustomHbondForce("(2*d+a)*distance(d1,a1)");
custom->addPerDonorParameter("d");
custom->addPerAcceptorParameter("a");
custom->addDonor(1, 0, -1, vector<double>({1.5}));
custom->addDonor(2, 0, -1, vector<double>({1.8}));
custom->addAcceptor(0, 1, -1, vector<double>({2.1}));
system.addForce(custom);
Context context(system, integrator, platform);
vector<Vec3> positions(3);
positions[0] = Vec3(0, 0, 0);
positions[1] = Vec3(0, 2, 0);
positions[2] = Vec3(2, 0, 0);
context.setPositions(positions);
State state = context.getState(State::Forces | State::Energy);
const vector<Vec3>& forces = state.getForces();
ASSERT_EQUAL_VEC(Vec3((2*1.8+2.1), (2*1.5+2.1), 0), forces[0], TOL);
ASSERT_EQUAL_VEC(Vec3(0, -(2*1.5+2.1), 0), forces[1], TOL);
ASSERT_EQUAL_VEC(Vec3(-(2*1.8+2.1), 0, 0), forces[2], TOL);
ASSERT_EQUAL_TOL(2*(2*1.8+2.1)+2*(2*1.5+2.1), state.getPotentialEnergy(), TOL);
}
void runPlatformTests(); void runPlatformTests();
int main(int argc, char* argv[]) { int main(int argc, char* argv[]) {
...@@ -254,6 +281,7 @@ int main(int argc, char* argv[]) { ...@@ -254,6 +281,7 @@ int main(int argc, char* argv[]) {
testCutoff(); testCutoff();
testCustomFunctions(); testCustomFunctions();
testIllegalVariable(); testIllegalVariable();
testParameters();
runPlatformTests(); runPlatformTests();
} }
catch(const exception& e) { catch(const exception& e) {
......
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