Bug fixes

platforms/cuda/src/CudaNonbondedUtilities.cpp

@@ -266,6 +266,7 @@ void CudaNonbondedUtilities::initialize(const System& system) {
         blockSorter = new CudaSort(context, new BlockSortTrait(context.getUseDoublePrecision()), numAtomBlocks);
         vector<unsigned int> count(2, 0);
         interactionCount.upload(count);
+        rebuildNeighborList.upload(count);
     }
 
     // Record arguments for kernels.

platforms/cuda/src/kernels/customNonbondedGroups.cu

@@ -151,6 +151,7 @@ extern "C" __global__  void buildNeighborList(int* __restrict__ rebuildNeighborL
             SYNC_WARPS;
         }
         if (anyInteraction[local_warp]) {
+            SYNC_WARPS;
             if (tgx == 0)
                 tileIndex[local_warp] = atomicAdd(numGroupTiles, 1);
             SYNC_WARPS;

platforms/opencl/src/OpenCLNonbondedUtilities.cpp

@@ -296,6 +296,7 @@ void OpenCLNonbondedUtilities::initialize(const System& system) {
         blockSorter = new OpenCLSort(context, new BlockSortTrait(context.getUseDoublePrecision()), numAtomBlocks);
         vector<cl_uint> count(1, 0);
         interactionCount.upload(count);
+        rebuildNeighborList.upload(count);
     }
 }
 

platforms/opencl/src/kernels/customNonbondedGroups.cl

@@ -144,7 +144,7 @@ __kernel void prepareToBuildNeighborList(__global int* restrict rebuildNeighborL
  * Filter the list of tiles to include only ones that have interactions within the
  * padded cutoff.
  */
-__kernel void buildNeighborList(__global int* restrict rebuildNeighborList, __global int* numGroupTiles,
+__kernel void buildNeighborList(__global int* restrict rebuildNeighborList, __global int* restrict numGroupTiles,
         __global const real4* restrict posq, __global const int4* restrict groupData, __global int4* restrict filteredGroupData,
         real4 periodicBoxSize, real4 invPeriodicBoxSize, real4 periodicBoxVecX, real4 periodicBoxVecY, real4 periodicBoxVecZ) {
     
@@ -193,8 +193,9 @@ __kernel void buildNeighborList(__global int* restrict rebuildNeighborList, __gl
             SYNC_WARPS;
         }
         if (anyInteraction[local_warp]) {
+            SYNC_WARPS;
             if (tgx == 0)
-                tileIndex[local_warp] = atom_add(numGroupTiles, 1);
+                tileIndex[local_warp] = atomic_add(numGroupTiles, 1);
             SYNC_WARPS;
             filteredGroupData[TILE_SIZE*tileIndex[local_warp]+tgx] = atomData;
         }