Merge pull request #1223 from peastman/gberrors

Fixed some test failures on OpenCL with CPU device

platforms/opencl/src/kernels/customGBEnergyN2_cpu.cl

@@ -279,6 +279,7 @@ __kernel void computeN2Energy(
                     real4 force = 0;
                     DECLARE_ATOM1_DERIVATIVES
                     real4 posq1 = posq[atom1];
+                    APPLY_PERIODIC_TO_POS_WITH_CENTER(posq1, blockCenterX)
                     LOAD_ATOM1_PARAMETERS
                     for (unsigned int j = 0; j < TILE_SIZE; j++) {
                         real4 posq2 = local_posq[j];

platforms/opencl/src/kernels/customGBValueN2_cpu.cl

@@ -245,6 +245,7 @@ __kernel void computeN2Value(__global const real4* restrict posq, __local real4*
                     unsigned int atom1 = x*TILE_SIZE+tgx;
                     real value = 0;
                     real4 posq1 = posq[atom1];
+                    APPLY_PERIODIC_TO_POS_WITH_CENTER(posq1, blockCenterX)
                     LOAD_ATOM1_PARAMETERS
                     for (unsigned int j = 0; j < TILE_SIZE; j++) {
                         real4 posq2 = local_posq[j];

tests/TestGBSAOBCForce.h

@@ -255,7 +255,7 @@ void testForce(int numParticles, NonbondedForce::NonbondedMethod method, GBSAOBC
         State state2 = context.getState(State::Energy);
         context.setPositions(positions3);
         State state3 = context.getState(State::Energy);
-        ASSERT_EQUAL_TOL(norm, (state2.getPotentialEnergy()-state3.getPotentialEnergy())/delta, 1e-2)
+        ASSERT_EQUAL_TOL(state2.getPotentialEnergy(), state3.getPotentialEnergy()+norm*delta, 1e-5)
     }
 }