Added missing synchronization

f1bbb8e7 · peastman · 8a42f4f6 · f1bbb8e7 · f1bbb8e7
Commit f1bbb8e7 authored Apr 25, 2018 by peastman
Showing with 2 additions and 0 deletions

platforms/cuda/src/kernels/customNonbondedGroups.cu platforms/cuda/src/kernels/customNonbondedGroups.cu +1 -0

platforms/opencl/src/kernels/customNonbondedGroups.cl platforms/opencl/src/kernels/customNonbondedGroups.cl +1 -0

No files found.
--- a/platforms/cuda/src/kernels/customNonbondedGroups.cu
+++ b/platforms/cuda/src/kernels/customNonbondedGroups.cu
@@ -156,6 +156,7 @@ extern "C" __global__  void buildNeighborList(int* __restrict__ rebuildNeighborL
        int rangeStop = rangeStart + reduceMax(rangeEnd-rangeStart);
        SYNC_WARPS;
        for (int j = rangeStart; j < rangeStop && !anyInteraction[local_warp]; j++) {
+            SYNC_WARPS;
            if (j < rangeEnd && tj < rangeEnd) {
                bool isExcluded = (((exclusions>>tj)&1) == 0);
                int localIndex = tbx+tj;

--- a/platforms/opencl/src/kernels/customNonbondedGroups.cl
+++ b/platforms/opencl/src/kernels/customNonbondedGroups.cl
@@ -198,6 +198,7 @@ __kernel void buildNeighborList(__global int* restrict rebuildNeighborList, __gl
        int rangeStop = rangeStart + reduceMax(rangeEnd-rangeStart, reductionBuffer);
        SYNC_WARPS;
        for (int j = rangeStart; j < rangeStop && !anyInteraction[local_warp]; j++) {
+            SYNC_WARPS;
            if (j < rangeEnd) {
                bool isExcluded = (((exclusions>>tj)&1) == 0);
                int localIndex = tbx+tj;