Commit 29256dfe authored by Yutong Zhao's avatar Yutong Zhao
Browse files

Fixes a race condition in buffer status in findBlocksWithInteractions()

parent cbc21173
...@@ -343,8 +343,8 @@ extern "C" __global__ void findBlocksWithInteractions(real4 periodicBoxSize, rea ...@@ -343,8 +343,8 @@ extern "C" __global__ void findBlocksWithInteractions(real4 periodicBoxSize, rea
valuesInBuffer = 0; valuesInBuffer = 0;
if (threadIdx.x == 0) if (threadIdx.x == 0)
bufferFull = false; bufferFull = false;
__syncthreads();
} }
__syncthreads();
} }
storeInteractionData(x, buffer, sum, temp, atoms, numAtoms, globalIndex, interactionCount, interactingTiles, interactingAtoms, periodicBoxSize, invPeriodicBoxSize, posq, posBuffer, blockCenterX, blockSizeX, maxTiles, true); storeInteractionData(x, buffer, sum, temp, atoms, numAtoms, globalIndex, interactionCount, interactingTiles, interactingAtoms, periodicBoxSize, invPeriodicBoxSize, posq, posBuffer, blockCenterX, blockSizeX, maxTiles, true);
} }
......
...@@ -320,8 +320,8 @@ __kernel void findBlocksWithInteractions(real4 periodicBoxSize, real4 invPeriodi ...@@ -320,8 +320,8 @@ __kernel void findBlocksWithInteractions(real4 periodicBoxSize, real4 invPeriodi
valuesInBuffer = 0; valuesInBuffer = 0;
if (get_local_id(0) == 0) if (get_local_id(0) == 0)
bufferFull = false; bufferFull = false;
barrier(CLK_LOCAL_MEM_FENCE);
} }
barrier(CLK_LOCAL_MEM_FENCE);
} }
storeInteractionData(x, buffer, sum, temp, atoms, &numAtoms, &globalIndex, interactionCount, interactingTiles, interactingAtoms, periodicBoxSize, invPeriodicBoxSize, posq, posBuffer, blockCenterX, blockSizeX, maxTiles, true); storeInteractionData(x, buffer, sum, temp, atoms, &numAtoms, &globalIndex, interactionCount, interactingTiles, interactingAtoms, periodicBoxSize, invPeriodicBoxSize, posq, posBuffer, blockCenterX, blockSizeX, maxTiles, true);
} }
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment