Added volatile keyword to fix crashes on some GPUs

platforms/cuda/src/kernels/customGBEnergyN2.cu

@@ -178,7 +178,7 @@ extern "C" __global__ void computeN2Energy(unsigned long long* __restrict__ forc
     int skipBase = 0;
     int currentSkipIndex = tbx;
     __shared__ int atomIndices[THREAD_BLOCK_SIZE];
-    __shared__ int skipTiles[THREAD_BLOCK_SIZE];
+    __shared__ volatile int skipTiles[THREAD_BLOCK_SIZE];
     skipTiles[threadIdx.x] = -1;
     
     while (pos < end) {

platforms/cuda/src/kernels/customGBValueN2.cu

@@ -156,7 +156,7 @@ extern "C" __global__ void computeN2Value(const real4* __restrict__ posq, const
     int skipBase = 0;
     int currentSkipIndex = tbx;
     __shared__ int atomIndices[THREAD_BLOCK_SIZE];
-    __shared__ int skipTiles[THREAD_BLOCK_SIZE];
+    __shared__ volatile int skipTiles[THREAD_BLOCK_SIZE];
     skipTiles[threadIdx.x] = -1;
     
     while (pos < end) {

platforms/cuda/src/kernels/gbsaObc1.cu

@@ -216,7 +216,7 @@ extern "C" __global__ void computeBornSum(unsigned long long* __restrict__ globa
     int skipBase = 0;
     int currentSkipIndex = tbx;
     __shared__ int atomIndices[FORCE_WORK_GROUP_SIZE];
-    __shared__ int skipTiles[FORCE_WORK_GROUP_SIZE];
+    __shared__ volatile int skipTiles[FORCE_WORK_GROUP_SIZE];
     skipTiles[threadIdx.x] = -1;
 
     while (pos < end) {
@@ -568,7 +568,7 @@ extern "C" __global__ void computeGBSAForce1(unsigned long long* __restrict__ fo
     int skipBase = 0;
     int currentSkipIndex = tbx;
     __shared__ int atomIndices[FORCE_WORK_GROUP_SIZE];
-    __shared__ int skipTiles[FORCE_WORK_GROUP_SIZE];
+    __shared__ volatile int skipTiles[FORCE_WORK_GROUP_SIZE];
     skipTiles[threadIdx.x] = -1;
 
     while (pos < end) {

platforms/cuda/src/kernels/nonbonded.cu

@@ -194,7 +194,7 @@ extern "C" __global__ void computeNonbonded(
     int skipBase = 0;
     int currentSkipIndex = tbx;
     __shared__ int atomIndices[THREAD_BLOCK_SIZE];
-    __shared__ int skipTiles[THREAD_BLOCK_SIZE];
+    __shared__ volatile int skipTiles[THREAD_BLOCK_SIZE];
     skipTiles[threadIdx.x] = -1;
     
     while (pos < end) {

platforms/opencl/src/kernels/customGBEnergyN2.cl

@@ -188,7 +188,7 @@ __kernel void computeN2Energy(
     int skipBase = 0;
     int currentSkipIndex = tbx;
     __local int atomIndices[FORCE_WORK_GROUP_SIZE];
-    __local int skipTiles[FORCE_WORK_GROUP_SIZE];
+    __local volatile int skipTiles[FORCE_WORK_GROUP_SIZE];
     skipTiles[get_local_id(0)] = -1;
 
     while (pos < end) {

platforms/opencl/src/kernels/customGBValueN2.cl

@@ -164,7 +164,7 @@ __kernel void computeN2Value(__global const real4* restrict posq, __local real4*
     int skipBase = 0;
     int currentSkipIndex = tbx;
     __local int atomIndices[FORCE_WORK_GROUP_SIZE];
-    __local int skipTiles[FORCE_WORK_GROUP_SIZE];
+    __local volatile int skipTiles[FORCE_WORK_GROUP_SIZE];
     skipTiles[get_local_id(0)] = -1;
 
     while (pos < end) {

platforms/opencl/src/kernels/gbsaObc.cl

@@ -176,7 +176,7 @@ __kernel void computeBornSum(
     int skipBase = 0;
     int currentSkipIndex = tbx;
     __local int atomIndices[FORCE_WORK_GROUP_SIZE];
-    __local int skipTiles[FORCE_WORK_GROUP_SIZE];
+    __local volatile int skipTiles[FORCE_WORK_GROUP_SIZE];
     skipTiles[get_local_id(0)] = -1;
 
     while (pos < end) {
@@ -550,7 +550,7 @@ __kernel void computeGBSAForce1(
     int skipBase = 0;
     int currentSkipIndex = tbx;
     __local int atomIndices[FORCE_WORK_GROUP_SIZE];
-    __local int skipTiles[FORCE_WORK_GROUP_SIZE];
+    __local volatile int skipTiles[FORCE_WORK_GROUP_SIZE];
     skipTiles[get_local_id(0)] = -1;
 
     while (pos < end) {

platforms/opencl/src/kernels/nonbonded.cl

@@ -202,7 +202,7 @@ __kernel void computeNonbonded(
     int skipBase = 0;
     int currentSkipIndex = tbx;
     __local int atomIndices[FORCE_WORK_GROUP_SIZE];
-    __local int skipTiles[FORCE_WORK_GROUP_SIZE];
+    __local volatile int skipTiles[FORCE_WORK_GROUP_SIZE];
     skipTiles[get_local_id(0)] = -1;
 
     while (pos < end) {

plugins/amoeba/platforms/cuda/src/kernels/amoebaGk.cu

@@ -760,7 +760,7 @@ extern "C" __global__ void computeEDiffForce(
     int end = startTileIndex+(warp+1)*numTiles/totalWarps;
     int skipBase = 0;
     int currentSkipIndex = tbx;
-    __shared__ int skipTiles[EDIFF_THREAD_BLOCK_SIZE];
+    __shared__ volatile int skipTiles[EDIFF_THREAD_BLOCK_SIZE];
     skipTiles[threadIdx.x] = -1;
 
     while (pos < end) {

plugins/amoeba/platforms/cuda/src/kernels/multipoleElectrostatics.cu

@@ -220,7 +220,7 @@ extern "C" __global__ void computeElectrostatics(
     int skipBase = 0;
     int currentSkipIndex = tbx;
     __shared__ int atomIndices[THREAD_BLOCK_SIZE];
-    __shared__ int skipTiles[THREAD_BLOCK_SIZE];
+    __shared__ volatile int skipTiles[THREAD_BLOCK_SIZE];
     skipTiles[threadIdx.x] = -1;
     
     while (pos < end) {

plugins/amoeba/platforms/cuda/src/kernels/multipoleFixedField.cu

@@ -559,7 +559,7 @@ extern "C" __global__ void computeFixedField(
     int skipBase = 0;
     int currentSkipIndex = tbx;
     __shared__ int atomIndices[THREAD_BLOCK_SIZE];
-    __shared__ int skipTiles[THREAD_BLOCK_SIZE];
+    __shared__ volatile int skipTiles[THREAD_BLOCK_SIZE];
     skipTiles[threadIdx.x] = -1;
 
     while (pos < end) {

plugins/amoeba/platforms/cuda/src/kernels/multipoleInducedField.cu

@@ -329,7 +329,7 @@ extern "C" __global__ void computeInducedField(
     int skipBase = 0;
     int currentSkipIndex = tbx;
     __shared__ int atomIndices[THREAD_BLOCK_SIZE];
-    __shared__ int skipTiles[THREAD_BLOCK_SIZE];
+    __shared__ volatile int skipTiles[THREAD_BLOCK_SIZE];
     skipTiles[threadIdx.x] = -1;
     
     while (pos < end) {

plugins/amoeba/platforms/cuda/src/kernels/pmeMultipoleElectrostatics.cu

@@ -302,7 +302,7 @@ extern "C" __global__ void computeElectrostatics(
     int skipBase = 0;
     int currentSkipIndex = tbx;
     __shared__ int atomIndices[THREAD_BLOCK_SIZE];
-    __shared__ int skipTiles[THREAD_BLOCK_SIZE];
+    __shared__ volatile int skipTiles[THREAD_BLOCK_SIZE];
     skipTiles[threadIdx.x] = -1;
     
     while (pos < end) {