 #define    epsilon0     8.8542e-12
 
 __global__ void spUpdateFieldYeeKernel(Real dt, Real3 dt_inv,
-                                       dim3 N, dim3 I, int rank,
+                                       dim3 N, dim3 I,
                                        Real const *Rho,
                                        Real const *Jx,
                                        Real const *Jy,

            ((By[s + I.x] - By[s]) * dt_inv.x - (Bx[s + I.y] - Bx[s]) * dt_inv.y) * speed_of_light2 -
            Jz[s] / epsilon0 * dt;

    Bx[s] /*= rank * 100 + s;//*/-= (Ez[s] - Ez[s - I.y]) * dt_inv.y - (Ey[s] - Ey[s - I.z]) * dt_inv.z;
    By[s] /*= rank * 100 + s;//*/-= (Ex[s] - Ex[s - I.z]) * dt_inv.z - (Ez[s] - Ez[s - I.x]) * dt_inv.x;
    Bz[s] /*= rank * 100 + s;//*/-= (Ey[s] - Ey[s - I.x]) * dt_inv.x - (Ex[s] - Ex[s - I.y]) * dt_inv.y;

}


    spFieldSubArray(fB, (void **) B);

    LOAD_KERNEL(spUpdateFieldYeeKernel, sizeType2Dim3(dims), 1,
                dt, real2Real3(dt_inv), sizeType2Dim3(dims), sizeType2Dim3(strides),
                (const Real *) rho,
                (const Real *) J[0],
                (const Real *) J[1],
