Fixed issue that caused inefficient sorting when a block contained only one atom (#5215)

* Fixed issue that caused inefficient sorting when a block contained only one atom

* Add the fix to OpenCL and HIP

platforms/cuda/src/kernels/findInteractingBlocks.cu

@@ -109,7 +109,8 @@ extern "C" __global__ void computeSortKeys(const real4* __restrict__ blockBoundi
         sizeRange = blockSizeRange[0];
         for (int i = 1; i < numSizes; i++) {
             real2 size = blockSizeRange[i];
-            sizeRange.x = min(sizeRange.x, size.x);
+            if (size.x > 0)
+                sizeRange.x = min(sizeRange.x, size.x);
             sizeRange.y = max(sizeRange.y, size.y);
         }
         sizeRange.x = LOG(sizeRange.x);

platforms/hip/src/kernels/findInteractingBlocks.hip

@@ -115,7 +115,8 @@ extern "C" __global__ void findBlockBounds(int numAtoms, real4 periodicBoxSize,
 
         // Record the range of sizes.
         real totalSize = blockSize.x+blockSize.y+blockSize.z;
-        atomicMin(&blockSizeRange->x, totalSize);
+        if (totalSize > 0)
+            atomicMin(&blockSizeRange->x, totalSize);
         atomicMax(&blockSizeRange->y, totalSize);
     }
     if (blockIdx.x == 0 && threadIdx.x == 0)

platforms/opencl/src/kernels/findInteractingBlocks.cl

@@ -75,7 +75,8 @@ __kernel void computeSortKeys(__global const real4* restrict blockBoundingBox, _
         sizeRange = blockSizeRange[0];
         for (int i = 1; i < numSizes; i++) {
             real2 size = blockSizeRange[i];
-            sizeRange.x = min(sizeRange.x, size.x);
+            if (size.x > 0)
+                sizeRange.x = min(sizeRange.x, size.x);
             sizeRange.y = max(sizeRange.y, size.y);
         }
         sizeRange.x = LOG(sizeRange.x);