ICODE-ADC/lammps
4
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

git-svn-id: svn://svn.icms.temple.edu/lammps-ro/trunk@13464 f3b2605a-c512-4ea7-a41b-209d697bcdaa

Browse Source
This commit is contained in:
sjplimp
2015-05-15 16:33:17 +00:00
parent 2284dcc452
commit c3fd41010d
23 changed files with 10777 additions and 20440 deletions
Download Patch File Download Diff File Expand all files Collapse all files

24
lib/gpu/lal_coul_dsf.cu
View File

@ -81,7 +81,7 @@ __kernel void k_coul_dsf(const __global numtyp4 *restrict x_,
int j=dev_packed[nbor];
numtyp factor_coul, r, prefactor, erfcc;
factor_coul = sp_lj[sbmask(j)];
factor_coul = (numtyp)1.0-sp_lj[sbmask(j)];
j &= NEIGHMASK;
numtyp4 jx; fetch4(jx,j,pos_tex); //x_[j];
@ -98,12 +98,12 @@ __kernel void k_coul_dsf(const __global numtyp4 *restrict x_,
r = ucl_sqrt(rsq);
fetch(prefactor,j,q_tex);
prefactor *= factor_coul * qqrd2e*qtmp/r;
prefactor *= qqrd2e*qtmp/r;
numtyp erfcd = ucl_exp(-alpha*alpha*rsq);
numtyp t = ucl_recip((numtyp)1.0 + EWALD_P*alpha*r);
erfcc = t * (A1+t*(A2+t*(A3+t*(A4+t*A5)))) * erfcd;
forcecoul = prefactor * (erfcc + (numtyp)2.0*alpha/MY_PIS*r*erfcd +
rsq*f_shift);
rsq*f_shift-factor_coul);
force = forcecoul * r2inv;
@ -112,10 +112,8 @@ __kernel void k_coul_dsf(const __global numtyp4 *restrict x_,
f.z+=delz*force;
if (eflag>0) {
if (rsq < cut_coulsq) {
numtyp e=prefactor*(erfcc-r*e_shift-rsq*f_shift);
e_coul += e;
}
numtyp e=prefactor*(erfcc-r*e_shift-rsq*f_shift-factor_coul);
e_coul += e;
}
if (vflag>0) {
virial[0] += delx*delx*force;
@ -182,7 +180,7 @@ __kernel void k_coul_dsf_fast(const __global numtyp4 *restrict x_,
int j=dev_packed[nbor];
numtyp factor_coul, r, prefactor, erfcc;
factor_coul = sp_lj[sbmask(j)];
factor_coul = (numtyp)1.0-sp_lj[sbmask(j)];
j &= NEIGHMASK;
numtyp4 jx; fetch4(jx,j,pos_tex); //x_[j];
@ -199,12 +197,12 @@ __kernel void k_coul_dsf_fast(const __global numtyp4 *restrict x_,
r = ucl_sqrt(rsq);
fetch(prefactor,j,q_tex);
prefactor *= factor_coul * qqrd2e*qtmp/r;
prefactor *= qqrd2e*qtmp/r;
numtyp erfcd = ucl_exp(-alpha*alpha*rsq);
numtyp t = ucl_recip((numtyp)1.0 + EWALD_P*alpha*r);
erfcc = t * (A1+t*(A2+t*(A3+t*(A4+t*A5)))) * erfcd;
forcecoul = prefactor * (erfcc + (numtyp)2.0*alpha/MY_PIS*r*erfcd +
rsq*f_shift);
rsq*f_shift-factor_coul);
force = forcecoul * r2inv;
@ -213,10 +211,8 @@ __kernel void k_coul_dsf_fast(const __global numtyp4 *restrict x_,
f.z+=delz*force;
if (eflag>0) {
if (rsq < cut_coulsq) {
numtyp e=prefactor*(erfcc-r*e_shift-rsq*f_shift);
e_coul += e;
}
numtyp e=prefactor*(erfcc-r*e_shift-rsq*f_shift-factor_coul);
e_coul += e;
}
if (vflag>0) {
virial[0] += delx*delx*force;

16
lib/gpu/lal_lj_dsf.cu
View File

@ -89,7 +89,7 @@ __kernel void k_lj_dsf(const __global numtyp4 *restrict x_,
numtyp factor_lj, factor_coul, r, prefactor, erfcc;
factor_lj = sp_lj[sbmask(j)];
factor_coul = sp_lj[sbmask(j)+4];
factor_coul = (numtyp)1.0-sp_lj[sbmask(j)+4];
j &= NEIGHMASK;
numtyp4 jx; fetch4(jx,j,pos_tex); //x_[j];
@ -115,12 +115,12 @@ __kernel void k_lj_dsf(const __global numtyp4 *restrict x_,
if (rsq < cut_coulsq) {
r = ucl_sqrt(rsq);
fetch(prefactor,j,q_tex);
prefactor *= factor_coul * qqrd2e*qtmp/r;
prefactor *= qqrd2e*qtmp/r;
numtyp erfcd = ucl_exp(-alpha*alpha*rsq);
numtyp t = ucl_recip((numtyp)1.0 + EWALD_P*alpha*r);
erfcc = t * (A1+t*(A2+t*(A3+t*(A4+t*A5)))) * erfcd;
forcecoul = prefactor * (erfcc + (numtyp)2.0*alpha/MY_PIS*r*erfcd +
rsq*f_shift);
rsq*f_shift-factor_coul);
} else
forcecoul = (numtyp)0.0;
@ -132,7 +132,7 @@ __kernel void k_lj_dsf(const __global numtyp4 *restrict x_,
if (eflag>0) {
if (rsq < cut_coulsq) {
numtyp e=prefactor*(erfcc-r*e_shift-rsq*f_shift);
numtyp e=prefactor*(erfcc-r*e_shift-rsq*f_shift-factor_coul);
e_coul += e;
}
if (rsq < lj1[mtype].z) {
@ -217,7 +217,7 @@ __kernel void k_lj_dsf_fast(const __global numtyp4 *restrict x_,
numtyp factor_lj, factor_coul, r, prefactor, erfcc;
factor_lj = sp_lj[sbmask(j)];
factor_coul = sp_lj[sbmask(j)+4];
factor_coul = (numtyp)1.0-sp_lj[sbmask(j)+4];
j &= NEIGHMASK;
numtyp4 jx; fetch4(jx,j,pos_tex); //x_[j];
@ -242,12 +242,12 @@ __kernel void k_lj_dsf_fast(const __global numtyp4 *restrict x_,
if (rsq < cut_coulsq) {
r = ucl_sqrt(rsq);
fetch(prefactor,j,q_tex);
prefactor *= factor_coul * qqrd2e*qtmp/r;
prefactor *= qqrd2e*qtmp/r;
numtyp erfcd = ucl_exp(-alpha*alpha*rsq);
numtyp t = ucl_recip((numtyp)1.0 + EWALD_P*alpha*r);
erfcc = t * (A1+t*(A2+t*(A3+t*(A4+t*A5)))) * erfcd;
forcecoul = prefactor * (erfcc + (numtyp)2.0*alpha/MY_PIS*r*erfcd +
rsq*f_shift);
rsq*f_shift-factor_coul);
} else
forcecoul = (numtyp)0.0;
@ -259,7 +259,7 @@ __kernel void k_lj_dsf_fast(const __global numtyp4 *restrict x_,
if (eflag>0) {
if (rsq < cut_coulsq) {
numtyp e=prefactor*(erfcc-r*e_shift-rsq*f_shift);
numtyp e=prefactor*(erfcc-r*e_shift-rsq*f_shift-factor_coul);
e_coul += e;
}
if (rsq < lj1[mtype].z) {