Skip to content

GitLab

Commit 3031c4da authored by peastman's avatar peastman
Browse files

Fixed bug on Intel GPUs

parent ae466134
Show whitespace changes
Inline Side-by-side
Showing with 2 additions and 0 deletions
platforms/opencl/src/kernels/findInteractingBlocks.cl
View file @ 3031c4da
...@@ -483,6 +483,7 @@ __kernel void findBlocksWithInteractions(real4 periodicBoxSize, real4 invPeriodi ...@@ -483,6 +483,7 @@ __kernel void findBlocksWithInteractions(real4 periodicBoxSize, real4 invPeriodi
int tilesToStore = neighborsInBuffer/TILE_SIZE; int tilesToStore = neighborsInBuffer/TILE_SIZE;
if (indexInWarp == 0) if (indexInWarp == 0)
*tileStartIndex = atom_add(interactionCount, tilesToStore); *tileStartIndex = atom_add(interactionCount, tilesToStore);
SYNC_WARPS;
int newTileStartIndex = *tileStartIndex; int newTileStartIndex = *tileStartIndex;
if (newTileStartIndex+tilesToStore < maxTiles) { if (newTileStartIndex+tilesToStore < maxTiles) {
if (indexInWarp < tilesToStore) if (indexInWarp < tilesToStore)
...@@ -503,6 +504,7 @@ __kernel void findBlocksWithInteractions(real4 periodicBoxSize, real4 invPeriodi ...@@ -503,6 +504,7 @@ __kernel void findBlocksWithInteractions(real4 periodicBoxSize, real4 invPeriodi
int tilesToStore = (neighborsInBuffer+TILE_SIZE-1)/TILE_SIZE; int tilesToStore = (neighborsInBuffer+TILE_SIZE-1)/TILE_SIZE;
if (indexInWarp == 0) if (indexInWarp == 0)
*tileStartIndex = atom_add(interactionCount, tilesToStore); *tileStartIndex = atom_add(interactionCount, tilesToStore);
SYNC_WARPS;
int newTileStartIndex = *tileStartIndex; int newTileStartIndex = *tileStartIndex;
if (newTileStartIndex+tilesToStore < maxTiles) { if (newTileStartIndex+tilesToStore < maxTiles) {
if (indexInWarp < tilesToStore) if (indexInWarp < tilesToStore)
......
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