diff --git a/lib/gpu/lal_dpd_coul_slater_long.cu b/lib/gpu/lal_dpd_coul_slater_long.cu index 4835b8777b..44931195ee 100644 --- a/lib/gpu/lal_dpd_coul_slater_long.cu +++ b/lib/gpu/lal_dpd_coul_slater_long.cu @@ -287,16 +287,16 @@ __kernel void k_dpd_coul_slater_long(const __global numtyp4 *restrict x_, // apply Slater electrostatic force if distance below Slater cutoff // and the two species have a slater coeff - // cutsq[mtype].z -> Coulombic squared cutoff - if ( cutsq[mtype].z != 0.0 && rsq < cutsq[mtype].z){ + // cutsq[mtype].z -> Slater cutoff + // extra[j].x -> q[j] ; particle j charge + if ( rsq < cutsq[mtype].z ){ numtyp r2inv=ucl_recip(rsq); numtyp _erfc; numtyp grij = g_ewald * r; numtyp expm2 = ucl_exp(-grij*grij); numtyp t = ucl_recip((numtyp)1.0 + EWALD_P*grij); _erfc = t * (A1+t*(A2+t*(A3+t*(A4+t*A5)))) * expm2; - numtyp prefactor = extra[j].x; - prefactor *= qqrd2e * cutsq[mtype].z * qtmp/r; + numtyp prefactor = qqrd2e * extra[j].x * qtmp / r; numtyp rlamdainv = r * lamdainv; numtyp exprlmdainv = ucl_exp((numtyp)-2.0*rlamdainv); numtyp slater_term = exprlmdainv*((numtyp)1.0 + ((numtyp)2.0*rlamdainv*((numtyp)1.0+rlamdainv))); @@ -306,9 +306,9 @@ __kernel void k_dpd_coul_slater_long(const __global numtyp4 *restrict x_, if (EVFLAG && eflag) { numtyp e_slater = ((numtyp)1.0 + rlamdainv)*exprlmdainv; - numtyp e = prefactor*(_erfc-e_slater); - if (factor_coul > (numtyp)0) e -= factor_coul*prefactor*((numtyp)1.0 - e_slater); - e_coul += e; + numtyp e_sf = prefactor*(_erfc-e_slater); + if (factor_coul > (numtyp)0) e_sf -= factor_coul*prefactor*((numtyp)1.0 - e_slater); + e_coul += e_sf; } } // if cut_coulsq @@ -471,16 +471,16 @@ __kernel void k_dpd_coul_slater_long_fast(const __global numtyp4 *restrict x_, // apply Slater electrostatic force if distance below Slater cutoff // and the two species have a slater coeff - // cutsq[mtype].z -> Coulombic squared cutoff - if ( cutsq[mtype].z != 0.0 && rsq < cutsq[mtype].z){ + // cutsq[mtype].z -> Slater cutoff + // extra[j].x -> q[j] ; particle j charge + if ( rsq < cutsq[mtype].z ){ numtyp r2inv=ucl_recip(rsq); numtyp _erfc; numtyp grij = g_ewald * r; numtyp expm2 = ucl_exp(-grij*grij); numtyp t = ucl_recip((numtyp)1.0 + EWALD_P*grij); _erfc = t * (A1+t*(A2+t*(A3+t*(A4+t*A5)))) * expm2; - numtyp prefactor = extra[j].x; - prefactor *= qqrd2e * cutsq[mtype].z * qtmp/r; + numtyp prefactor = qqrd2e * extra[j].x * qtmp / r; numtyp rlamdainv = r * lamdainv; numtyp exprlmdainv = ucl_exp((numtyp)-2.0*rlamdainv); numtyp slater_term = exprlmdainv*((numtyp)1.0 + ((numtyp)2.0*rlamdainv*((numtyp)1.0+rlamdainv))); diff --git a/lib/gpu/lal_dpd_coul_slater_long.h b/lib/gpu/lal_dpd_coul_slater_long.h index 1870255998..007e7a2e00 100644 --- a/lib/gpu/lal_dpd_coul_slater_long.h +++ b/lib/gpu/lal_dpd_coul_slater_long.h @@ -65,7 +65,7 @@ class DPDCoulSlaterLong : public BaseDPD { /// coeff.x = a0, coeff.y = gamma, coeff.z = sigma, coeff.w = cut_dpd UCL_D_Vec coeff; - /// cutsq.x = cutsq, cutsq.y = cut_dpdsq, cutsq.w = cut_slatersq + /// cutsq.x = cutsq, cutsq.y = cut_dpdsq, cutsq.z = cut_slatersq UCL_D_Vec cutsq; /// Special LJ values