Fixed illegal read in kernel (#2843)

050a1aa1 · peastman · GitHub · 817fd3bc · 050a1aa1 · 050a1aa1
Unverified Commit 050a1aa1 authored Sep 16, 2020 by peastman Committed by GitHub Sep 16, 2020
Showing with 4 additions and 2 deletions

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

platforms/opencl/src/kernels/findInteractingBlocks.cl platforms/opencl/src/kernels/findInteractingBlocks.cl +2 -1

No files found.
--- a/platforms/cuda/src/kernels/findInteractingBlocks.cu
+++ b/platforms/cuda/src/kernels/findInteractingBlocks.cu
@@ -344,7 +344,8 @@ extern "C" __global__ void findBlocksWithInteractions(real4 periodicBoxSize, rea
                            for (int j = 0; j < tilesToStore; j++)
                                interactingAtoms[(newTileStartIndex+j)*TILE_SIZE+indexInWarp] = buffer[indexInWarp+j*TILE_SIZE];
                        }
-                        buffer[indexInWarp] = buffer[indexInWarp+TILE_SIZE*tilesToStore];
+                        if (indexInWarp+TILE_SIZE*tilesToStore < BUFFER_SIZE)
+                            buffer[indexInWarp] = buffer[indexInWarp+TILE_SIZE*tilesToStore];
                        neighborsInBuffer -= TILE_SIZE*tilesToStore;
                    }
                }

--- a/platforms/opencl/src/kernels/findInteractingBlocks.cl
+++ b/platforms/opencl/src/kernels/findInteractingBlocks.cl
@@ -244,7 +244,8 @@ __kernel void findBlocksWithInteractions(real4 periodicBoxSize, real4 invPeriodi
                            for (int j = 0; j < tilesToStore; j++)
                                interactingAtoms[(newTileStartIndex+j)*TILE_SIZE+indexInWarp] = buffer[indexInWarp+j*TILE_SIZE];
                        }
-                        buffer[indexInWarp] = buffer[indexInWarp+TILE_SIZE*tilesToStore];
+                        if (indexInWarp+TILE_SIZE*tilesToStore < BUFFER_SIZE)
+                            buffer[indexInWarp] = buffer[indexInWarp+TILE_SIZE*tilesToStore];
                        neighborsInBuffer -= TILE_SIZE*tilesToStore;
                   }
                }