Fix compilation error in OpenCL kernels on some platforms (#3857)

platforms/opencl/src/OpenCLBondedUtilities.cpp

@@ -106,7 +106,7 @@ void OpenCLBondedUtilities::initialize(const System& system) {
     stringstream s;
     for (int i = 0; i < (int) prefixCode.size(); i++)
         s<<prefixCode[i];
-    s<<"__kernel void computeBondedForces(__global long* restrict forceBuffers, __global mixed* restrict energyBuffer, __global const real4* restrict posq, int groups, real4 periodicBoxSize, real4 invPeriodicBoxSize, real4 periodicBoxVecX, real4 periodicBoxVecY, real4 periodicBoxVecZ";
+    s<<"__kernel void computeBondedForces(__global unsigned long* restrict forceBuffers, __global mixed* restrict energyBuffer, __global const real4* restrict posq, int groups, real4 periodicBoxSize, real4 invPeriodicBoxSize, real4 periodicBoxVecX, real4 periodicBoxVecY, real4 periodicBoxVecZ";
     for (int force = 0; force < numForces; force++) {
         string indexType = "uint"+(indexWidth[force] == 1 ? "" : context.intToString(indexWidth[force]));
         s<<", __global const "<<indexType<<"* restrict atomIndices"<<force;

platforms/opencl/src/kernels/common.cl

@@ -7,7 +7,7 @@
 #ifdef cl_khr_int64_base_atomics
 #pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
 #else
-void atom_add(volatile __global unsigned long* p, long unsigned val) {
+__attribute__((overloadable)) unsigned long atom_add(volatile __global unsigned long* p, unsigned long val) {
     volatile __global unsigned int* word = (volatile __global unsigned int*) p;
 #ifdef __ENDIAN_LITTLE__
     int lowIndex = 0;
@@ -21,6 +21,7 @@ void atom_add(volatile __global unsigned long* p, long unsigned val) {
     upper += carry;
     if (upper != 0)
         atomic_add(&word[1-lowIndex], upper);
+    return 0;
 }
 #endif
 

platforms/opencl/src/kernels/nonbonded.cl

@@ -14,7 +14,7 @@ typedef struct {
  * Compute nonbonded interactions.
  */
 __kernel void computeNonbonded(
-        __global long* restrict forceBuffers,
+        __global unsigned long* restrict forceBuffers,
         __global mixed* restrict energyBuffer, __global const real4* restrict posq, __global const unsigned int* restrict exclusions,
         __global const int2* restrict exclusionTiles, unsigned int startTileIndex, unsigned long numTileIndices
 #ifdef USE_CUTOFF