Fixed RPMD error created by changes to computing and storing charges

plugins/rpmd/platforms/cuda/src/kernels/rpmd.cu

@@ -194,7 +194,9 @@ extern "C" __global__ void copyDataToContext(mixed4* srcVel, mixed4* dstVel, mix
         int index = base+order[particle];
         dstVel[particle] = srcVel[index];
         mixed4 posq = srcPos[index];
-        dstPos[particle] = make_real4(posq.x, posq.y, posq.z, posq.w);
+        dstPos[particle].x = posq.x;
+        dstPos[particle].y = posq.y;
+        dstPos[particle].z = posq.z;
     }
 }
 
@@ -211,7 +213,9 @@ extern "C" __global__ void copyDataFromContext(long long* srcForce, long long* d
         dstForce[base*3+index+PADDED_NUM_ATOMS*2] = srcForce[particle+PADDED_NUM_ATOMS*2];
         dstVel[base+index] = srcVel[particle];
         real4 posq = srcPos[particle];
-        dstPos[base+index] = make_mixed4(posq.x, posq.y, posq.z, posq.w);
+        dstPos[base+index].x = posq.x;
+        dstPos[base+index].y = posq.y;
+        dstPos[base+index].z = posq.z;
 
     }
 }
@@ -224,7 +228,10 @@ extern "C" __global__ void applyCellTranslations(mixed4* posq, real4* movedPos,
         int index = order[particle];
         real4 p = movedPos[particle];
         mixed4 delta = make_mixed4(p.x, p.y, p.z, p.w)-posq[movedCopy*PADDED_NUM_ATOMS+index];
-        for (int copy = 0; copy < NUM_COPIES; copy++)
-            posq[copy*PADDED_NUM_ATOMS+index] += delta;
+        for (int copy = 0; copy < NUM_COPIES; copy++) {
+            posq[copy*PADDED_NUM_ATOMS+index].x += delta.x;
+            posq[copy*PADDED_NUM_ATOMS+index].y += delta.y;
+            posq[copy*PADDED_NUM_ATOMS+index].z += delta.z;
+        }
     }
 }

plugins/rpmd/platforms/opencl/src/kernels/rpmd.cl

@@ -185,7 +185,7 @@ __kernel void copyDataToContext(__global mixed4* srcVel, __global mixed4* dstVel
     for (int particle = get_global_id(0); particle < NUM_ATOMS; particle += get_global_size(0)) {
         int index = base+order[particle];
         dstVel[particle] = srcVel[index];
-        dstPos[particle] = convert_real4(srcPos[index]);
+        dstPos[particle].xyz = convert_real4(srcPos[index]).xyz;
     }
 }
 
@@ -199,7 +199,7 @@ __kernel void copyDataFromContext(__global real4* srcForce, __global real4* dstF
         int index = base+order[particle];
         dstForce[index] = srcForce[particle];
         dstVel[index] = srcVel[particle];
-        dstPos[index] = convert_mixed4(srcPos[particle]);
+        dstPos[index].xyz = convert_mixed4(srcPos[particle]).xyz;
     }
 }
 
@@ -211,6 +211,6 @@ __kernel void applyCellTranslations(__global mixed4* posq, __global real4* moved
         int index = order[particle];
         mixed4 delta = convert_mixed4(movedPos[particle])-posq[movedCopy*PADDED_NUM_ATOMS+index];
         for (int copy = 0; copy < NUM_COPIES; copy++)
-            posq[copy*PADDED_NUM_ATOMS+index] += delta;
+            posq[copy*PADDED_NUM_ATOMS+index].xyz += delta.xyz;
     }
 }