Fix for errors on Pascal GPUs

5044c12d · Peter Eastman · d4465cb4 · 5044c12d
Commit 5044c12d authored Jan 11, 2017 by Peter Eastman
Show whitespace changes
Inline Side-by-side

Showing with 1 addition and 1 deletion

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

No files found.
--- a/platforms/opencl/src/kernels/findInteractingBlocks.cl
+++ b/platforms/opencl/src/kernels/findInteractingBlocks.cl
@@ -99,7 +99,7 @@ __kernel void findBlocksWithInteractions(real4 periodicBoxSize, real4 invPeriodi
    __local real3 posBuffer[GROUP_SIZE];
    __local volatile int workgroupTileIndex[GROUP_SIZE/32];
    __local bool includeBlockFlags[GROUP_SIZE];
-    __local short2 atomCountBuffer[GROUP_SIZE];
+    __local volatile short2 atomCountBuffer[GROUP_SIZE];
    __local int* buffer = workgroupBuffer+BUFFER_SIZE*(warpStart/32);
    __local int* exclusionsForX = warpExclusions+MAX_EXCLUSIONS*(warpStart/32);
    __local volatile int* tileStartIndex = workgroupTileIndex+(warpStart/32);