             pos_frac.y *= (Scalar)Ny;
             pos_frac.z *= (Scalar)Nz;
 
-            float shift, shiftone, x0, y0, z0, dx, dy, dz;
+            Scalar shift, shiftone, x0, y0, z0, dx, dy, dz;
             int nlower, nupper, mx, my, mz, nxi, nyi, nzi; 
     
             nlower = -(order-1)/2;

    if (error != cudaSuccess)
        return error;

    // Update the force and virial with fixed exclusions. The memsets are merely commented and not removed
    // to avoid merge conflicts.
//    cudaMemset(d_force, 0, sizeof(float4)*N);
//    cudaMemset(d_virial, 0, 6*sizeof(float)*virial_pitch);

    gpu_fix_exclusions_kernel <<< grid, threads >>>  (d_force,
                                                      d_virial,
                                                      virial_pitch,

                wx = fast::pow(gpu_sinc(argx),order);

               for (iy = -nby; iy <= nby; iy++) {
                    qy = (j.y+(Scalar)(Ny*iy));
                    kn2 = b2 * qy;
                    argy = Scalar(0.5)*qy*kH.y;
                    wy = fast::pow(gpu_sinc(argy),order);
