Commit 00b5aa4d authored by Peter Eastman's avatar Peter Eastman
Browse files

Bug fixes

parent cdec19f9
...@@ -208,7 +208,6 @@ void testErrorTolerance(NonbondedForce::NonbondedMethod method) { ...@@ -208,7 +208,6 @@ void testErrorTolerance(NonbondedForce::NonbondedMethod method) {
} }
force->setNonbondedMethod(method); force->setNonbondedMethod(method);
CudaPlatform platform; CudaPlatform platform;
VerletIntegrator integrator(0.01);
// For various values of the cutoff and error tolerance, see if the actual error is reasonable. // For various values of the cutoff and error tolerance, see if the actual error is reasonable.
...@@ -218,6 +217,7 @@ void testErrorTolerance(NonbondedForce::NonbondedMethod method) { ...@@ -218,6 +217,7 @@ void testErrorTolerance(NonbondedForce::NonbondedMethod method) {
double norm = 0.0; double norm = 0.0;
for (double tol = 5e-5; tol < 1e-3; tol *= 2.0) { for (double tol = 5e-5; tol < 1e-3; tol *= 2.0) {
force->setEwaldErrorTolerance(tol); force->setEwaldErrorTolerance(tol);
VerletIntegrator integrator(0.01);
Context context(system, integrator, platform); Context context(system, integrator, platform);
context.setPositions(positions); context.setPositions(positions);
State state = context.getState(State::Forces); State state = context.getState(State::Forces);
......
...@@ -84,7 +84,7 @@ __kernel void updateBsplines(__global float4* posq, __global float4* pmeBsplineT ...@@ -84,7 +84,7 @@ __kernel void updateBsplines(__global float4* posq, __global float4* pmeBsplineT
__kernel void gridSpreadCharge(__global float2* pmeAtomGridIndex, __global int* pmeAtomRange, __global float2* pmeGrid, __global float4* pmeBsplineTheta) { __kernel void gridSpreadCharge(__global float2* pmeAtomGridIndex, __global int* pmeAtomRange, __global float2* pmeGrid, __global float4* pmeBsplineTheta) {
unsigned int numGridPoints = GRID_SIZE_X*GRID_SIZE_Y*GRID_SIZE_Z; unsigned int numGridPoints = GRID_SIZE_X*GRID_SIZE_Y*GRID_SIZE_Z;
for (int gridIndex = get_global_id(0); gridIndex < numGridPoints; gridIndex += get_global_size(0)) { for (int gridIndex = get_global_id(0); gridIndex < numGridPoints; gridIndex += get_global_size(0)) {
int3 gridPoint; int4 gridPoint;
gridPoint.x = gridIndex/(GRID_SIZE_Y*GRID_SIZE_Z); gridPoint.x = gridIndex/(GRID_SIZE_Y*GRID_SIZE_Z);
int remainder = gridIndex-gridPoint.x*GRID_SIZE_Y*GRID_SIZE_Z; int remainder = gridIndex-gridPoint.x*GRID_SIZE_Y*GRID_SIZE_Z;
gridPoint.y = remainder/GRID_SIZE_Z; gridPoint.y = remainder/GRID_SIZE_Z;
...@@ -145,7 +145,7 @@ __kernel void reciprocalConvolution(__global float2* pmeGrid, __global float* en ...@@ -145,7 +145,7 @@ __kernel void reciprocalConvolution(__global float2* pmeGrid, __global float* en
__kernel void gridInterpolateForce(__global float4* posq, __global float4* forceBuffers, __global float4* pmeBsplineTheta, __global float4* pmeBsplineDTheta, __global float2* pmeGrid) { __kernel void gridInterpolateForce(__global float4* posq, __global float4* forceBuffers, __global float4* pmeBsplineTheta, __global float4* pmeBsplineDTheta, __global float2* pmeGrid) {
for (int atom = get_global_id(0); atom < NUM_ATOMS; atom += get_global_size(0)) { for (int atom = get_global_id(0); atom < NUM_ATOMS; atom += get_global_size(0)) {
float3 force = 0.0f; float4 force = 0.0f;
float4 pos = posq[atom]; float4 pos = posq[atom];
float4 t = (float4) ((pos.x/PERIODIC_BOX_SIZE_X+1.0f)*GRID_SIZE_X, float4 t = (float4) ((pos.x/PERIODIC_BOX_SIZE_X+1.0f)*GRID_SIZE_X,
(pos.y/PERIODIC_BOX_SIZE_Y+1.0f)*GRID_SIZE_Y, (pos.y/PERIODIC_BOX_SIZE_Y+1.0f)*GRID_SIZE_Y,
......
...@@ -114,7 +114,7 @@ __kernel void sortBuckets(__global TYPE* data, __global TYPE* buckets, int numBu ...@@ -114,7 +114,7 @@ __kernel void sortBuckets(__global TYPE* data, __global TYPE* buckets, int numBu
if (length <= get_local_size(0)) { if (length <= get_local_size(0)) {
// Load the data into local memory. // Load the data into local memory.
buffer[get_local_id(0)] = (get_local_id(0) < length ? buckets[startIndex+get_local_id(0)] : MAXFLOAT); buffer[get_local_id(0)] = (get_local_id(0) < length ? buckets[startIndex+get_local_id(0)] : (TYPE) MAXFLOAT);
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
// Perform a bitonic sort in local memory. // Perform a bitonic sort in local memory.
......
...@@ -207,7 +207,6 @@ void testErrorTolerance(NonbondedForce::NonbondedMethod method) { ...@@ -207,7 +207,6 @@ void testErrorTolerance(NonbondedForce::NonbondedMethod method) {
} }
force->setNonbondedMethod(method); force->setNonbondedMethod(method);
OpenCLPlatform platform; OpenCLPlatform platform;
VerletIntegrator integrator(0.01);
// For various values of the cutoff and error tolerance, see if the actual error is reasonable. // For various values of the cutoff and error tolerance, see if the actual error is reasonable.
...@@ -217,6 +216,7 @@ void testErrorTolerance(NonbondedForce::NonbondedMethod method) { ...@@ -217,6 +216,7 @@ void testErrorTolerance(NonbondedForce::NonbondedMethod method) {
double norm = 0.0; double norm = 0.0;
for (double tol = 5e-5; tol < 1e-3; tol *= 2.0) { for (double tol = 5e-5; tol < 1e-3; tol *= 2.0) {
force->setEwaldErrorTolerance(tol); force->setEwaldErrorTolerance(tol);
VerletIntegrator integrator(0.01);
Context context(system, integrator, platform); Context context(system, integrator, platform);
context.setPositions(positions); context.setPositions(positions);
State state = context.getState(State::Forces); State state = context.getState(State::Forces);
......
...@@ -320,7 +320,6 @@ void testErrorTolerance(NonbondedForce::NonbondedMethod method) { ...@@ -320,7 +320,6 @@ void testErrorTolerance(NonbondedForce::NonbondedMethod method) {
} }
force->setNonbondedMethod(method); force->setNonbondedMethod(method);
ReferencePlatform platform; ReferencePlatform platform;
VerletIntegrator integrator(0.01);
// For various values of the cutoff and error tolerance, see if the actual error is reasonable. // For various values of the cutoff and error tolerance, see if the actual error is reasonable.
...@@ -330,6 +329,7 @@ void testErrorTolerance(NonbondedForce::NonbondedMethod method) { ...@@ -330,6 +329,7 @@ void testErrorTolerance(NonbondedForce::NonbondedMethod method) {
double norm = 0.0; double norm = 0.0;
for (double tol = 5e-5; tol < 1e-3; tol *= 2.0) { for (double tol = 5e-5; tol < 1e-3; tol *= 2.0) {
force->setEwaldErrorTolerance(tol); force->setEwaldErrorTolerance(tol);
VerletIntegrator integrator(0.01);
Context context(system, integrator, platform); Context context(system, integrator, platform);
context.setPositions(positions); context.setPositions(positions);
State state = context.getState(State::Forces); State state = context.getState(State::Forces);
......
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