Fixed bugs in CPU version of Custom GB

platforms/opencl/src/OpenCLKernels.cpp

@@ -2339,7 +2339,7 @@ double OpenCLCalcCustomGBForceKernel::execute(ContextImpl& context, bool include
         }
         int index = 0;
         pairValueKernel.setArg<cl::Buffer>(index++, cl.getPosq().getDeviceBuffer());
-        pairValueKernel.setArg(index++, nb.getForceThreadBlockSize()*sizeof(cl_float4), NULL);
+        pairValueKernel.setArg(index++, (deviceIsCpu ? OpenCLContext::TileSize : nb.getForceThreadBlockSize())*sizeof(cl_float4), NULL);
         pairValueKernel.setArg<cl::Buffer>(index++, cl.getNonbondedUtilities().getExclusions().getDeviceBuffer());
         pairValueKernel.setArg<cl::Buffer>(index++, cl.getNonbondedUtilities().getExclusionIndices().getDeviceBuffer());
         pairValueKernel.setArg<cl::Buffer>(index++, cl.getNonbondedUtilities().getExclusionRowIndices().getDeviceBuffer());
@@ -2361,7 +2361,7 @@ double OpenCLCalcCustomGBForceKernel::execute(ContextImpl& context, bool include
         for (int i = 0; i < (int) params->getBuffers().size(); i++) {
             const OpenCLNonbondedUtilities::ParameterInfo& buffer = params->getBuffers()[i];
             pairValueKernel.setArg<cl::Memory>(index++, buffer.getMemory());
-            pairValueKernel.setArg(index++, nb.getForceThreadBlockSize()*buffer.getSize(), NULL);
+            pairValueKernel.setArg(index++, (deviceIsCpu ? OpenCLContext::TileSize : nb.getForceThreadBlockSize())*buffer.getSize(), NULL);
         }
         if (tabulatedFunctionParams != NULL) {
             for (int i = 0; i < (int) tabulatedFunctions.size(); i++)

platforms/opencl/src/kernels/customGBEnergyN2_cpu.cl

@@ -80,6 +80,7 @@ __kernel void computeN2Energy(__global float4* restrict forceBuffers, __global f
 #endif
                 unsigned int atom1 = x*TILE_SIZE+tgx;
                 float4 force = 0.0f;
+                DECLARE_ATOM1_DERIVATIVES
                 float4 posq1 = posq[atom1];
                 LOAD_ATOM1_PARAMETERS
                 for (unsigned int j = 0; j < TILE_SIZE; j++) {
@@ -141,6 +142,7 @@ __kernel void computeN2Energy(__global float4* restrict forceBuffers, __global f
                     if ((flags2&(1<<tgx)) != 0) {
                         unsigned int atom1 = x*TILE_SIZE+tgx;
                         float value = 0.0f;
+                        DECLARE_ATOM1_DERIVATIVES
                         float4 posq1 = posq[atom1];
                         LOAD_ATOM1_PARAMETERS
                         for (unsigned int j = 0; j < TILE_SIZE; j++) {
@@ -188,6 +190,7 @@ __kernel void computeN2Energy(__global float4* restrict forceBuffers, __global f
                 for (unsigned int tgx = 0; tgx < TILE_SIZE; tgx++) {
                     unsigned int atom1 = x*TILE_SIZE+tgx;
                     float4 force = 0.0f;
+                    DECLARE_ATOM1_DERIVATIVES
                     float4 posq1 = posq[atom1];
                     LOAD_ATOM1_PARAMETERS
 #ifdef USE_EXCLUSIONS