Fixed illegal memory access in nonbonded kernel (#2837)

platforms/cuda/src/CudaNonbondedUtilities.cpp

@@ -651,6 +651,21 @@ CUfunction CudaNonbondedUtilities::createInteractionKernel(const string& source,
         }
     }
     replacements["LOAD_ATOM2_PARAMETERS"] = load2j.str();
+    
+    stringstream clearLocal;
+    for (int i = 0; i < (int) params.size(); i++) {
+        if (useShuffle)
+            clearLocal<<"shfl";
+        else
+            clearLocal<<"localData[atom2].";
+        clearLocal<<params[i].getName()<<" = ";
+        if (params[i].getNumComponents() == 1)
+            clearLocal<<"0;\n";
+        else
+            clearLocal<<"make_"<<params[i].getType()<<"(0);\n";
+    }
+    replacements["CLEAR_LOCAL_PARAMETERS"] = clearLocal.str();
+
     stringstream initDerivs;
     for (int i = 0; i < energyParameterDerivatives.size(); i++)
         initDerivs<<"mixed energyParamDeriv"<<i<<" = 0;\n";

platforms/cuda/src/kernels/nonbonded.cu

@@ -419,6 +419,7 @@ extern "C" __global__ void computeNonbonded(
                 localData[threadIdx.x].y = 0;
                 localData[threadIdx.x].z = 0;
 #endif
+                CLEAR_LOCAL_PARAMETERS
             }
 #ifdef USE_PERIODIC
             if (singlePeriodicCopy) {
@@ -607,11 +608,11 @@ extern "C" __global__ void computeNonbonded(
         real4 posq2 = posq[atom2];
         LOAD_ATOM1_PARAMETERS
         int j = atom2;
-atom2 = threadIdx.x;
+        atom2 = threadIdx.x;
         DECLARE_LOCAL_PARAMETERS
         LOAD_LOCAL_PARAMETERS_FROM_GLOBAL
         LOAD_ATOM2_PARAMETERS
-atom2 = pair.y;
+        atom2 = pair.y;
         real3 delta = make_real3(posq2.x-posq1.x, posq2.y-posq1.y, posq2.z-posq1.z);
 #ifdef USE_PERIODIC
         APPLY_PERIODIC_TO_DELTA(delta)

platforms/opencl/src/OpenCLNonbondedUtilities.cpp

@@ -640,6 +640,15 @@ cl::Kernel OpenCLNonbondedUtilities::createInteractionKernel(const string& sourc
         }
     }
     replacements["LOAD_ATOM2_PARAMETERS"] = load2j.str();
+    stringstream clearLocal;
+    for (int i = 0; i < (int) params.size(); i++) {
+        if (params[i].getNumComponents() == 1)
+            clearLocal<<"localData[localAtomIndex]."<<params[i].getName()<<" = 0;\n";
+        else
+            for (int j = 0; j < params[i].getNumComponents(); ++j)
+                clearLocal<<"localData[localAtomIndex]."<<params[i].getName()<<"_"<<suffixes[j]<<" = 0;\n";
+    }
+    replacements["CLEAR_LOCAL_PARAMETERS"] = clearLocal.str();
     stringstream initDerivs;
     for (int i = 0; i < energyParameterDerivatives.size(); i++)
         initDerivs<<"mixed energyParamDeriv"<<i<<" = 0;\n";

platforms/opencl/src/kernels/nonbonded.cl

@@ -288,6 +288,7 @@ __kernel void computeNonbonded(
                 localData[localAtomIndex].x = 0;
                 localData[localAtomIndex].y = 0;
                 localData[localAtomIndex].z = 0;
+                CLEAR_LOCAL_PARAMETERS
             }
             SYNC_WARPS;
 #ifdef USE_PERIODIC