jianmu
/
lammps
forked from lijiext/lammps
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Cleaned up 3-body kernels, reverted some mistaken changes to vashishta

This commit is contained in:
Trung Nguyen 2017-07-23 00:08:55 -05:00
parent a71f5a0c20
commit aa60ef6ed8
4 changed files with 3 additions and 26 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
lib/gpu/lal_tersoff.cu
View File

@ -581,7 +581,6 @@ __kernel void k_tersoff_three_center(const __global numtyp4 *restrict x_,
delr1[2] = jx.z-ix.z; delr1[2] = jx.z-ix.z;
numtyp rsq1 = delr1[0]*delr1[0] + delr1[1]*delr1[1] + delr1[2]*delr1[2]; numtyp rsq1 = delr1[0]*delr1[0] + delr1[1]*delr1[1] + delr1[2]*delr1[2];
// if (rsq1 > cutsq[ijparam]) continue;
numtyp r1 = ucl_sqrt(rsq1); numtyp r1 = ucl_sqrt(rsq1);
numtyp r1inv = ucl_rsqrt(rsq1); numtyp r1inv = ucl_rsqrt(rsq1);
@ -770,8 +769,6 @@ __kernel void k_tersoff_three_end(const __global numtyp4 *restrict x_,
delr1[2] = jx.z-ix.z; delr1[2] = jx.z-ix.z;
numtyp rsq1 = delr1[0]*delr1[0] + delr1[1]*delr1[1] + delr1[2]*delr1[2]; numtyp rsq1 = delr1[0]*delr1[0] + delr1[1]*delr1[1] + delr1[2]*delr1[2];
// if (rsq1 > cutsq[ijparam]) continue;
numtyp mdelr1[3]; numtyp mdelr1[3];
mdelr1[0] = -delr1[0]; mdelr1[0] = -delr1[0];
mdelr1[1] = -delr1[1]; mdelr1[1] = -delr1[1];
@ -1017,8 +1014,6 @@ __kernel void k_tersoff_three_end_vatom(const __global numtyp4 *restrict x_,
delr1[2] = jx.z-ix.z; delr1[2] = jx.z-ix.z;
numtyp rsq1 = delr1[0]*delr1[0] + delr1[1]*delr1[1] + delr1[2]*delr1[2]; numtyp rsq1 = delr1[0]*delr1[0] + delr1[1]*delr1[1] + delr1[2]*delr1[2];
// if (rsq1 > cutsq[ijparam]) continue;
numtyp mdelr1[3]; numtyp mdelr1[3];
mdelr1[0] = -delr1[0]; mdelr1[0] = -delr1[0];
mdelr1[1] = -delr1[1]; mdelr1[1] = -delr1[1];

7
lib/gpu/lal_tersoff_mod.cu
View File

@ -308,8 +308,6 @@ __kernel void k_tersoff_mod_zeta(const __global numtyp4 *restrict x_,
delr1.z = jx.z-ix.z; delr1.z = jx.z-ix.z;
numtyp rsq1 = delr1.x*delr1.x+delr1.y*delr1.y+delr1.z*delr1.z; numtyp rsq1 = delr1.x*delr1.x+delr1.y*delr1.y+delr1.z*delr1.z;
// if (rsq1 > cutsq[ijparam]) continue;
// compute zeta_ij // compute zeta_ij
z = (acctyp)0; z = (acctyp)0;
@ -585,7 +583,6 @@ __kernel void k_tersoff_mod_three_center(const __global numtyp4 *restrict x_,
delr1[2] = jx.z-ix.z; delr1[2] = jx.z-ix.z;
numtyp rsq1 = delr1[0]*delr1[0] + delr1[1]*delr1[1] + delr1[2]*delr1[2]; numtyp rsq1 = delr1[0]*delr1[0] + delr1[1]*delr1[1] + delr1[2]*delr1[2];
// if (rsq1 > cutsq[ijparam]) continue;
numtyp r1 = ucl_sqrt(rsq1); numtyp r1 = ucl_sqrt(rsq1);
numtyp r1inv = ucl_rsqrt(rsq1); numtyp r1inv = ucl_rsqrt(rsq1);
@ -780,8 +777,6 @@ __kernel void k_tersoff_mod_three_end(const __global numtyp4 *restrict x_,
delr1[2] = jx.z-ix.z; delr1[2] = jx.z-ix.z;
numtyp rsq1 = delr1[0]*delr1[0] + delr1[1]*delr1[1] + delr1[2]*delr1[2]; numtyp rsq1 = delr1[0]*delr1[0] + delr1[1]*delr1[1] + delr1[2]*delr1[2];
// if (rsq1 > cutsq[ijparam]) continue;
numtyp mdelr1[3]; numtyp mdelr1[3];
mdelr1[0] = -delr1[0]; mdelr1[0] = -delr1[0];
mdelr1[1] = -delr1[1]; mdelr1[1] = -delr1[1];
@ -1036,8 +1031,6 @@ __kernel void k_tersoff_mod_three_end_vatom(const __global numtyp4 *restrict x_,
delr1[2] = jx.z-ix.z; delr1[2] = jx.z-ix.z;
numtyp rsq1 = delr1[0]*delr1[0] + delr1[1]*delr1[1] + delr1[2]*delr1[2]; numtyp rsq1 = delr1[0]*delr1[0] + delr1[1]*delr1[1] + delr1[2]*delr1[2];
// if (rsq1 > cutsq[ijparam]) continue;
numtyp mdelr1[3]; numtyp mdelr1[3];
mdelr1[0] = -delr1[0]; mdelr1[0] = -delr1[0];
mdelr1[1] = -delr1[1]; mdelr1[1] = -delr1[1];

7
lib/gpu/lal_tersoff_zbl.cu
View File

@ -314,8 +314,6 @@ __kernel void k_tersoff_zbl_zeta(const __global numtyp4 *restrict x_,
delr1.z = jx.z-ix.z; delr1.z = jx.z-ix.z;
numtyp rsq1 = delr1.x*delr1.x+delr1.y*delr1.y+delr1.z*delr1.z; numtyp rsq1 = delr1.x*delr1.x+delr1.y*delr1.y+delr1.z*delr1.z;
// if (rsq1 > cutsq[ijparam]) continue;
// compute zeta_ij // compute zeta_ij
z = (acctyp)0; z = (acctyp)0;
@ -601,7 +599,6 @@ __kernel void k_tersoff_zbl_three_center(const __global numtyp4 *restrict x_,
delr1[2] = jx.z-ix.z; delr1[2] = jx.z-ix.z;
numtyp rsq1 = delr1[0]*delr1[0] + delr1[1]*delr1[1] + delr1[2]*delr1[2]; numtyp rsq1 = delr1[0]*delr1[0] + delr1[1]*delr1[1] + delr1[2]*delr1[2];
// if (rsq1 > cutsq[ijparam]) continue;
numtyp r1 = ucl_sqrt(rsq1); numtyp r1 = ucl_sqrt(rsq1);
numtyp r1inv = ucl_rsqrt(rsq1); numtyp r1inv = ucl_rsqrt(rsq1);
@ -790,8 +787,6 @@ __kernel void k_tersoff_zbl_three_end(const __global numtyp4 *restrict x_,
delr1[2] = jx.z-ix.z; delr1[2] = jx.z-ix.z;
numtyp rsq1 = delr1[0]*delr1[0] + delr1[1]*delr1[1] + delr1[2]*delr1[2]; numtyp rsq1 = delr1[0]*delr1[0] + delr1[1]*delr1[1] + delr1[2]*delr1[2];
// if (rsq1 > cutsq[ijparam]) continue;
numtyp mdelr1[3]; numtyp mdelr1[3];
mdelr1[0] = -delr1[0]; mdelr1[0] = -delr1[0];
mdelr1[1] = -delr1[1]; mdelr1[1] = -delr1[1];
@ -1037,8 +1032,6 @@ __kernel void k_tersoff_zbl_three_end_vatom(const __global numtyp4 *restrict x_,
delr1[2] = jx.z-ix.z; delr1[2] = jx.z-ix.z;
numtyp rsq1 = delr1[0]*delr1[0] + delr1[1]*delr1[1] + delr1[2]*delr1[2]; numtyp rsq1 = delr1[0]*delr1[0] + delr1[1]*delr1[1] + delr1[2]*delr1[2];
// if (rsq1 > cutsq[ijparam]) continue;
numtyp mdelr1[3]; numtyp mdelr1[3];
mdelr1[0] = -delr1[0]; mdelr1[0] = -delr1[0];
mdelr1[1] = -delr1[1]; mdelr1[1] = -delr1[1];

10
lib/gpu/lal_vashishta.cu
View File

@ -474,9 +474,7 @@ __kernel void k_vashishta_three_center(const __global numtyp4 *restrict x_,
numtyp4 param4_ijparam; fetch4(param4_ijparam,ijparam,param4_tex); numtyp4 param4_ijparam; fetch4(param4_ijparam,ijparam,param4_tex);
param_r0sq_ij=param4_ijparam.x; param_r0sq_ij=param4_ijparam.x;
if (rsq1 > param_r0sq_ij) continue; // still keep this for neigh no and tpa > 1
// if (rsq1 > param_r0sq_ij) continue;
param_gamma_ij=param4_ijparam.y; param_gamma_ij=param4_ijparam.y;
param_r0_ij=param4_ijparam.w; param_r0_ij=param4_ijparam.w;
@ -617,8 +615,7 @@ __kernel void k_vashishta_three_end(const __global numtyp4 *restrict x_,
int ijparam=elem2param[itype*nelements*nelements+jtype*nelements+jtype]; int ijparam=elem2param[itype*nelements*nelements+jtype*nelements+jtype];
numtyp4 param4_ijparam; fetch4(param4_ijparam,ijparam,param4_tex); numtyp4 param4_ijparam; fetch4(param4_ijparam,ijparam,param4_tex);
param_r0sq_ij = param4_ijparam.x; param_r0sq_ij = param4_ijparam.x;
if (rsq1 > param_r0sq_ij) continue; // still keep this for neigh no and tpa > 1
// if (rsq1 > param_r0sq_ij) continue;
param_gamma_ij=param4_ijparam.y; param_gamma_ij=param4_ijparam.y;
param_r0_ij = param4_ijparam.w; param_r0_ij = param4_ijparam.w;
@ -773,8 +770,7 @@ __kernel void k_vashishta_three_end_vatom(const __global numtyp4 *restrict x_,
int ijparam=elem2param[itype*nelements*nelements+jtype*nelements+jtype]; int ijparam=elem2param[itype*nelements*nelements+jtype*nelements+jtype];
numtyp4 param4_ijparam; fetch4(param4_ijparam,ijparam,param4_tex); numtyp4 param4_ijparam; fetch4(param4_ijparam,ijparam,param4_tex);
param_r0sq_ij=param4_ijparam.x; param_r0sq_ij=param4_ijparam.x;
if (rsq1 > param_r0sq_ij) continue; // still keep this for neigh no and tpa > 1
// if (rsq1 > param_r0sq_ij) continue;
param_gamma_ij=param4_ijparam.y; param_gamma_ij=param4_ijparam.y;
param_r0_ij=param4_ijparam.w; param_r0_ij=param4_ijparam.w;