Skip to content

GitLab

Commit d5666536 authored by peastman's avatar peastman
Browse files

Fixed a bug in neighbor list construction

parent 9e10769a
Hide whitespace changes
Inline Side-by-side
Showing with 4 additions and 4 deletions
platforms/cuda/src/kernels/findInteractingBlocks.cu
View file @ d5666536
...@@ -256,7 +256,7 @@ extern "C" __global__ void findBlocksWithInteractions(real4 periodicBoxSize, rea ...@@ -256,7 +256,7 @@ extern "C" __global__ void findBlocksWithInteractions(real4 periodicBoxSize, rea
if (indexInWarp == 0) if (indexInWarp == 0)
tileStartIndex = atomicAdd(interactionCount, tilesToStore); tileStartIndex = atomicAdd(interactionCount, tilesToStore);
int newTileStartIndex = tileStartIndex; int newTileStartIndex = tileStartIndex;
if (newTileStartIndex+tilesToStore < maxTiles) { if (newTileStartIndex+tilesToStore <= maxTiles) {
if (indexInWarp < tilesToStore) if (indexInWarp < tilesToStore)
interactingTiles[newTileStartIndex+indexInWarp] = x; interactingTiles[newTileStartIndex+indexInWarp] = x;
for (int j = 0; j < tilesToStore; j++) for (int j = 0; j < tilesToStore; j++)
...@@ -275,7 +275,7 @@ extern "C" __global__ void findBlocksWithInteractions(real4 periodicBoxSize, rea ...@@ -275,7 +275,7 @@ extern "C" __global__ void findBlocksWithInteractions(real4 periodicBoxSize, rea
if (indexInWarp == 0) if (indexInWarp == 0)
tileStartIndex = atomicAdd(interactionCount, tilesToStore); tileStartIndex = atomicAdd(interactionCount, tilesToStore);
int newTileStartIndex = tileStartIndex; int newTileStartIndex = tileStartIndex;
if (newTileStartIndex+tilesToStore < maxTiles) { if (newTileStartIndex+tilesToStore <= maxTiles) {
if (indexInWarp < tilesToStore) if (indexInWarp < tilesToStore)
interactingTiles[newTileStartIndex+indexInWarp] = x; interactingTiles[newTileStartIndex+indexInWarp] = x;
for (int j = 0; j < tilesToStore; j++) for (int j = 0; j < tilesToStore; j++)
......
platforms/opencl/src/kernels/findInteractingBlocks.cl
View file @ d5666536
...@@ -216,7 +216,7 @@ __kernel void findBlocksWithInteractions(real4 periodicBoxSize, real4 invPeriodi ...@@ -216,7 +216,7 @@ __kernel void findBlocksWithInteractions(real4 periodicBoxSize, real4 invPeriodi
*tileStartIndex = atom_add(interactionCount, tilesToStore); *tileStartIndex = atom_add(interactionCount, tilesToStore);
SYNC_WARPS; SYNC_WARPS;
int newTileStartIndex = *tileStartIndex; int newTileStartIndex = *tileStartIndex;
if (newTileStartIndex+tilesToStore < maxTiles) { if (newTileStartIndex+tilesToStore <= maxTiles) {
if (indexInWarp < tilesToStore) if (indexInWarp < tilesToStore)
interactingTiles[newTileStartIndex+indexInWarp] = x; interactingTiles[newTileStartIndex+indexInWarp] = x;
for (int j = 0; j < tilesToStore; j++) for (int j = 0; j < tilesToStore; j++)
...@@ -237,7 +237,7 @@ __kernel void findBlocksWithInteractions(real4 periodicBoxSize, real4 invPeriodi ...@@ -237,7 +237,7 @@ __kernel void findBlocksWithInteractions(real4 periodicBoxSize, real4 invPeriodi
*tileStartIndex = atom_add(interactionCount, tilesToStore); *tileStartIndex = atom_add(interactionCount, tilesToStore);
SYNC_WARPS; SYNC_WARPS;
int newTileStartIndex = *tileStartIndex; int newTileStartIndex = *tileStartIndex;
if (newTileStartIndex+tilesToStore < maxTiles) { if (newTileStartIndex+tilesToStore <= maxTiles) {
if (indexInWarp < tilesToStore) if (indexInWarp < tilesToStore)
interactingTiles[newTileStartIndex+indexInWarp] = x; interactingTiles[newTileStartIndex+indexInWarp] = x;
for (int j = 0; j < tilesToStore; j++) for (int j = 0; j < tilesToStore; j++)
......
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