Fixed thread synchronization problem in LINCS kernel

8a83a2fb · Peter Eastman · 306d99e8 · 8a83a2fb
Commit 8a83a2fb authored Apr 08, 2009 by Peter Eastman
Show whitespace changes
Inline Side-by-side

Showing with 2 additions and 1 deletion

platforms/cuda/src/kernels/kLincs.cu platforms/cuda/src/kernels/kLincs.cu +2 -1

No files found.
--- a/platforms/cuda/src/kernels/kLincs.cu
+++ b/platforms/cuda/src/kernels/kLincs.cu
@@ -184,6 +184,7 @@ __global__ void kApplyLincs_kernel(float4* atomPositions, bool addOldPosition)

    kSolveMatrix_kernel(cSim.pSyncCounter+1);
    kUpdateAtomPositions_kernel(atomPositions);
+    kSyncAllThreads_kernel(cSim.pSyncCounter+cSim.lincsTerms+1);

    // Correct for rotational lengthening.

@@ -217,7 +218,7 @@ __global__ void kApplyLincs_kernel(float4* atomPositions, bool addOldPosition)

    // Solve the matrix equation and update the atom positions.

-    kSolveMatrix_kernel(cSim.pSyncCounter+cSim.lincsTerms+1);
+    kSolveMatrix_kernel(cSim.pSyncCounter+cSim.lincsTerms+2);
    kUpdateAtomPositions_kernel(atomPositions);
 }