Bug fixes to OpenCL version of CustomManyParticleForce

platforms/cuda/src/CudaKernels.cpp

@@ -5009,7 +5009,8 @@ double CudaCalcCustomManyParticleForceKernel::execute(ContextImpl& context, bool
             cu.executeKernel(startIndicesKernel, &startIndicesArgs[0], 256, 256, 256*sizeof(int));
             cu.executeKernel(copyPairsKernel, &copyPairsArgs[0], maxNeighborPairs);
         }
-        cu.executeKernel(forceKernel, &forceArgs[0], cu.getNumAtoms()*forceWorkgroupSize, forceWorkgroupSize);
+        int maxThreads = min(cu.getNumAtoms()*forceWorkgroupSize, cu.getEnergyBuffer().getSize());
+        cu.executeKernel(forceKernel, &forceArgs[0], maxThreads, forceWorkgroupSize);
         if (nonbondedMethod != NoCutoff) {
             // Make sure there was enough memory for the neighbor list.
 

platforms/opencl/src/OpenCLKernels.cpp

@@ -5073,7 +5073,7 @@ double OpenCLCalcCustomManyParticleForceKernel::execute(ContextImpl& context, bo
         // Set arguments for the force kernel.
         
         int index = 0;
-        forceKernel.setArg<cl::Buffer>(index++, cl.getForce().getDeviceBuffer());
+        forceKernel.setArg<cl::Buffer>(index++, cl.getLongForceBuffer().getDeviceBuffer());
         forceKernel.setArg<cl::Buffer>(index++, cl.getEnergyBuffer().getDeviceBuffer());
         forceKernel.setArg<cl::Buffer>(index++, cl.getPosq().getDeviceBuffer());
         setPeriodicBoxSizeArg(cl, forceKernel, index++);
@@ -5175,7 +5175,8 @@ double OpenCLCalcCustomManyParticleForceKernel::execute(ContextImpl& context, bo
             cl.executeKernel(startIndicesKernel, 256, 256);
             cl.executeKernel(copyPairsKernel, maxNeighborPairs);
         }
-        cl.executeKernel(forceKernel, cl.getNumAtoms()*forceWorkgroupSize, forceWorkgroupSize);
+        int maxThreads = min(cl.getNumAtoms()*forceWorkgroupSize, cl.getEnergyBuffer().getSize());
+        cl.executeKernel(forceKernel, maxThreads, forceWorkgroupSize);
         if (nonbondedMethod != NoCutoff) {
             // Make sure there was enough memory for the neighbor list.
 

platforms/opencl/src/kernels/customManyParticle.cl

+#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
 #pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
 
 /**
@@ -82,7 +83,7 @@ __kernel void computeInteraction(
         , __global int* restrict particleTypes, __global int* restrict orderIndex, __global int* restrict particleOrder
 #endif
 #ifdef USE_EXCLUSIONS
-        , int* __global restrict exclusions, __global int* restrict exclusionStartIndex
+        , __global int* restrict exclusions, __global int* restrict exclusionStartIndex
 #endif
         PARAMETER_ARGUMENTS) {
     real energy = 0.0f;