Fixed error with RMSDForce on GPUs without 64 bit atomics (#2963)

platforms/opencl/src/OpenCLContext.cpp

@@ -519,8 +519,7 @@ void OpenCLContext::initialize() {
     reduceForcesKernel.setArg<cl::Buffer>(1, forceBuffers.getDeviceBuffer());
     reduceForcesKernel.setArg<cl_int>(2, paddedNumAtoms);
     reduceForcesKernel.setArg<cl_int>(3, numForceBuffers);
-    if (supports64BitGlobalAtomics)
-        addAutoclearBuffer(longForceBuffer);
+    addAutoclearBuffer(longForceBuffer);
     addAutoclearBuffer(forceBuffers);
     addAutoclearBuffer(energyBuffer);
     int numEnergyParamDerivs = energyParamDerivNames.size();

platforms/opencl/src/kernels/utilities.cl

@@ -88,11 +88,7 @@ __kernel void reduceForces(__global long* restrict longBuffer, __global real4* r
     int totalSize = bufferSize*numBuffers;
     real scale = 1/(real) 0x100000000;
     for (int index = get_global_id(0); index < bufferSize; index += get_global_size(0)) {
-#ifdef SUPPORTS_64_BIT_ATOMICS
         real4 sum = (real4) (scale*longBuffer[index], scale*longBuffer[index+bufferSize], scale*longBuffer[index+2*bufferSize], 0);
-#else
-        real4 sum = (real4) 0;
-#endif
         for (int i = index; i < totalSize; i += bufferSize)
             sum += buffer[i];
         buffer[index] = sum;