Fixed round-off error that occasionally caused errors in the neighbor list

platforms/opencl/src/kernels/customGBEnergyN2_cpu.cl

@@ -40,8 +40,8 @@ __kernel void computeN2Energy(__global float4* forceBuffers, __global float* ene
         {
             y = (unsigned int) floor(NUM_BLOCKS+0.5f-sqrt((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
             x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
-            if (x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
-                y++;
+            if (x < y || x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
+                y += (x < y ? -1 : 1);
                 x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
             }
         }

platforms/opencl/src/kernels/customGBEnergyN2_default.cl

@@ -44,8 +44,8 @@ void computeN2Energy(__global float4* forceBuffers, __global float* energyBuffer
         {
             y = (unsigned int) floor(NUM_BLOCKS+0.5f-sqrt((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
             x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
-            if (x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
-                y++;
+            if (x < y || x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
+                y += (x < y ? -1 : 1);
                 x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
             }
         }

platforms/opencl/src/kernels/customGBEnergyN2_nvidia.cl

@@ -45,8 +45,8 @@ void computeN2Energy(__global float4* forceBuffers, __global float* energyBuffer
         {
             y = (unsigned int) floor(NUM_BLOCKS+0.5f-sqrt((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
             x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
-            if (x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
-                y++;
+            if (x < y || x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
+                y += (x < y ? -1 : 1);
                 x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
             }
         }

platforms/opencl/src/kernels/customGBValueN2_cpu.cl

@@ -37,8 +37,8 @@ __kernel void computeN2Value(__global float4* posq, __local float4* local_posq,
         {
             y = (unsigned int) floor(NUM_BLOCKS+0.5f-sqrt((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
             x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
-            if (x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
-                y++;
+            if (x < y || x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
+                y += (x < y ? -1 : 1);
                 x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
             }
         }

platforms/opencl/src/kernels/customGBValueN2_default.cl

@@ -41,8 +41,8 @@ void computeN2Value(__global float4* posq, __local float4* local_posq, __global
         {
             y = (unsigned int) floor(NUM_BLOCKS+0.5f-sqrt((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
             x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
-            if (x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
-                y++;
+            if (x < y || x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
+                y += (x < y ? -1 : 1);
                 x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
             }
         }

platforms/opencl/src/kernels/customGBValueN2_nvidia.cl

@@ -43,8 +43,8 @@ void computeN2Value(__global float4* posq, __local float4* local_posq, __global
         {
             y = (unsigned int) floor(NUM_BLOCKS+0.5f-sqrt((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
             x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
-            if (x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
-                y++;
+            if (x < y || x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
+                y += (x < y ? -1 : 1);
                 x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
             }
         }

platforms/opencl/src/kernels/findInteractingBlocks.cl

@@ -180,8 +180,8 @@ __kernel void findBlocksWithInteractions(float cutoffSquared, float4 periodicBox
         if (index < numTiles) {
             unsigned int y = (unsigned int) floor(NUM_BLOCKS+0.5f-sqrt((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*index));
             unsigned int x = (index-y*NUM_BLOCKS+y*(y+1)/2);
-            if (x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
-                y++;
+            if (x < y || x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
+                y += (x < y ? -1 : 1);
                 x = (index-y*NUM_BLOCKS+y*(y+1)/2);
             }
 

platforms/opencl/src/kernels/findInteractingBlocks_cpu.cl

@@ -134,8 +134,8 @@ __kernel void findBlocksWithInteractions(float cutoffSquared, float4 periodicBox
 
         unsigned int y = (unsigned int) floor(NUM_BLOCKS+0.5f-sqrt((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*index));
         unsigned int x = (index-y*NUM_BLOCKS+y*(y+1)/2);
-        if (x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
-            y++;
+        if (x < y || x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
+            y += (x < y ? -1 : 1);
             x = (index-y*NUM_BLOCKS+y*(y+1)/2);
         }
 

platforms/opencl/src/kernels/gbsaObc_cpu.cl

@@ -44,8 +44,8 @@ __kernel void computeBornSum(__global float* global_bornSum, __global float4* po
         {
             y = (unsigned int) floor(NUM_BLOCKS+0.5f-sqrt((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
             x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
-            if (x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
-                y++;
+            if (x < y || x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
+                y += (x < y ? -1 : 1);
                 x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
             }
         }
@@ -222,8 +222,8 @@ __kernel void computeGBSAForce1(__global float4* forceBuffers, __global float* e
         {
             y = (unsigned int) floor(NUM_BLOCKS+0.5f-sqrt((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
             x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
-            if (x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
-                y++;
+            if (x < y || x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
+                y += (x < y ? -1 : 1);
                 x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
             }
         }

platforms/opencl/src/kernels/gbsaObc_default.cl

@@ -45,8 +45,8 @@ void computeBornSum(__global float* global_bornSum, __global float4* posq, __glo
         {
             y = (unsigned int) floor(NUM_BLOCKS+0.5f-sqrt((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
             x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
-            if (x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
-                y++;
+            if (x < y || x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
+                y += (x < y ? -1 : 1);
                 x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
             }
         }
@@ -232,8 +232,8 @@ void computeGBSAForce1(__global float4* forceBuffers, __global float* energyBuff
         {
             y = (unsigned int) floor(NUM_BLOCKS+0.5f-sqrt((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
             x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
-            if (x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
-                y++;
+            if (x < y || x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
+                y += (x < y ? -1 : 1);
                 x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
             }
         }

platforms/opencl/src/kernels/gbsaObc_nvidia.cl

@@ -47,8 +47,8 @@ void computeBornSum(__global float* global_bornSum, __global float4* posq, __glo
         {
             y = (unsigned int) floor(NUM_BLOCKS+0.5f-sqrt((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
             x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
-            if (x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
-                y++;
+            if (x < y || x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
+                y += (x < y ? -1 : 1);
                 x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
             }
         }
@@ -300,8 +300,8 @@ void computeGBSAForce1(__global float4* forceBuffers, __global float* energyBuff
         {
             y = (unsigned int) floor(NUM_BLOCKS+0.5f-sqrt((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
             x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
-            if (x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
-                y++;
+            if (x < y || x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
+                y += (x < y ? -1 : 1);
                 x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
             }
         }

platforms/opencl/src/kernels/nonbonded_cpu.cl

@@ -44,8 +44,8 @@ __kernel void computeNonbonded(__global float4* forceBuffers, __global float* en
         {
             y = (unsigned int) floor(NUM_BLOCKS+0.5f-sqrt((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
             x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
-            if (x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
-                y++;
+            if (x < y || x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
+                y += (x < y ? -1 : 1);
                 x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
             }
         }

platforms/opencl/src/kernels/nonbonded_default.cl

@@ -47,8 +47,8 @@ void computeNonbonded(__global float4* forceBuffers, __global float* energyBuffe
         {
             y = (unsigned int) floor(NUM_BLOCKS+0.5f-sqrt((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
             x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
-            if (x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
-                y++;
+            if (x < y || x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
+                y += (x < y ? -1 : 1);
                 x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
             }
         }

platforms/opencl/src/kernels/nonbonded_nvidia.cl

@@ -49,8 +49,8 @@ void computeNonbonded(__global float4* forceBuffers, __global float* energyBuffe
         {
             y = (unsigned int) floor(NUM_BLOCKS+0.5f-sqrt((NUM_BLOCKS+0.5f)*(NUM_BLOCKS+0.5f)-2*pos));
             x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
-            if (x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
-                y++;
+            if (x < y || x >= NUM_BLOCKS) { // Occasionally happens due to roundoff error.
+                y += (x < y ? -1 : 1);
                 x = (pos-y*NUM_BLOCKS+y*(y+1)/2);
             }
         }