Fixed error simulating very large systems (see bug 896)

platforms/cuda/src/kernels/gpu.cpp

@@ -1839,6 +1839,8 @@ int gpuBuildThreadBlockWorkList(gpuContext gpu)
     gpu->sim.bf2WorkUnitsPerBlockRemainder  = cells - gpu->sim.bornForce2_blocks * gpu->sim.bf2WorkUnitsPerBlock;
     gpu->sim.interaction_threads_per_block = 64;
     gpu->sim.interaction_blocks = (gpu->sim.workUnits + gpu->sim.interaction_threads_per_block - 1) / gpu->sim.interaction_threads_per_block;
+    if (gpu->sim.interaction_blocks > 8*gpu->sim.blocks)
+        gpu->sim.interaction_blocks = 8*gpu->sim.blocks;
 
     // Decrease thread count for extra small molecules to spread computation
     // across entire chip

platforms/cuda/src/kernels/kFindInteractingBlocks.h

@@ -79,7 +79,7 @@ __global__ void METHOD_NAME(kFindBlockBounds, _kernel)()
 __global__ void METHOD_NAME(kFindBlocksWithInteractions, _kernel)()
 {
     unsigned int pos = blockIdx.x * blockDim.x + threadIdx.x;
-    if (pos < cSim.workUnits)
+    while (pos < cSim.workUnits)
     {
         // Extract cell coordinates from appropriate work unit
 
@@ -105,6 +105,7 @@ __global__ void METHOD_NAME(kFindBlocksWithInteractions, _kernel)()
         dy = max(0.0f, abs(dy)-boxSizea.y-boxSizeb.y);
         dz = max(0.0f, abs(dz)-boxSizea.z-boxSizeb.z);
         cSim.pInteractionFlag[pos] = (dx*dx+dy*dy+dz*dz > cSim.nonbondedCutoffSqr ? 0 : 1);
+        pos += gridDim.x*blockDim.x;
     }
 }