mirror of https://github.com/lammps/lammps.git
remove trailing whitespace
This commit is contained in:
parent
d53e468ce0
commit
f1aaad2b13
|
@ -109,7 +109,7 @@ __kernel void k_born_coul_long_cs(const __global numtyp4 *restrict x_,
|
||||||
numtyp rsq = delx*delx+dely*dely+delz*delz;
|
numtyp rsq = delx*delx+dely*dely+delz*delz;
|
||||||
|
|
||||||
int mtype=itype*lj_types+jtype;
|
int mtype=itype*lj_types+jtype;
|
||||||
if (rsq<cutsq_sigma[mtype].x) { // cutsq
|
if (rsq<cutsq_sigma[mtype].x) { // cutsq
|
||||||
numtyp forcecoul,forceborn,force,r6inv,prefactor,_erfc,rexp;
|
numtyp forcecoul,forceborn,force,r6inv,prefactor,_erfc,rexp;
|
||||||
|
|
||||||
rsq += EPSILON; // Add Epsilon for case: r = 0; Interaction must be removed by special bond;
|
rsq += EPSILON; // Add Epsilon for case: r = 0; Interaction must be removed by special bond;
|
||||||
|
@ -250,7 +250,7 @@ __kernel void k_born_coul_long_cs_fast(const __global numtyp4 *restrict x_,
|
||||||
numtyp delz = ix.z-jx.z;
|
numtyp delz = ix.z-jx.z;
|
||||||
numtyp rsq = delx*delx+dely*dely+delz*delz;
|
numtyp rsq = delx*delx+dely*dely+delz*delz;
|
||||||
|
|
||||||
if (rsq<cutsq_sigma[mtype].x) { // cutsq
|
if (rsq<cutsq_sigma[mtype].x) { // cutsq
|
||||||
numtyp forcecoul,forceborn,force,r6inv,prefactor,_erfc,rexp;
|
numtyp forcecoul,forceborn,force,r6inv,prefactor,_erfc,rexp;
|
||||||
|
|
||||||
rsq += EPSILON; // Add Epsilon for case: r = 0; Interaction must be removed by special bond;
|
rsq += EPSILON; // Add Epsilon for case: r = 0; Interaction must be removed by special bond;
|
||||||
|
|
|
@ -10,7 +10,7 @@
|
||||||
This file is part of the LAMMPS Accelerator Library (LAMMPS_AL)
|
This file is part of the LAMMPS Accelerator Library (LAMMPS_AL)
|
||||||
__________________________________________________________________________
|
__________________________________________________________________________
|
||||||
|
|
||||||
begin :
|
begin :
|
||||||
email : pl.rodolfo@gmail.com
|
email : pl.rodolfo@gmail.com
|
||||||
dekoning@ifi.unicamp.br
|
dekoning@ifi.unicamp.br
|
||||||
***************************************************************************/
|
***************************************************************************/
|
||||||
|
@ -31,10 +31,10 @@ __kernel void k_ufm(const __global numtyp4 *restrict x_,
|
||||||
const __global numtyp4 *restrict uf3,
|
const __global numtyp4 *restrict uf3,
|
||||||
const int lj_types,
|
const int lj_types,
|
||||||
const __global numtyp *restrict sp_lj,
|
const __global numtyp *restrict sp_lj,
|
||||||
const __global int * dev_nbor,
|
const __global int * dev_nbor,
|
||||||
const __global int * dev_packed,
|
const __global int * dev_packed,
|
||||||
__global acctyp4 *restrict ans,
|
__global acctyp4 *restrict ans,
|
||||||
__global acctyp *restrict engv,
|
__global acctyp *restrict engv,
|
||||||
const int eflag, const int vflag, const int inum,
|
const int eflag, const int vflag, const int inum,
|
||||||
const int nbor_pitch, const int t_per_atom) {
|
const int nbor_pitch, const int t_per_atom) {
|
||||||
int tid, ii, offset;
|
int tid, ii, offset;
|
||||||
|
@ -46,19 +46,19 @@ __kernel void k_ufm(const __global numtyp4 *restrict x_,
|
||||||
acctyp virial[6];
|
acctyp virial[6];
|
||||||
for (int i=0; i<6; i++)
|
for (int i=0; i<6; i++)
|
||||||
virial[i]=(acctyp)0;
|
virial[i]=(acctyp)0;
|
||||||
|
|
||||||
if (ii<inum) {
|
if (ii<inum) {
|
||||||
int i, numj, nbor, nbor_end;
|
int i, numj, nbor, nbor_end;
|
||||||
__local int n_stride;
|
__local int n_stride;
|
||||||
nbor_info(dev_nbor,dev_packed,nbor_pitch,t_per_atom,ii,offset,i,numj,
|
nbor_info(dev_nbor,dev_packed,nbor_pitch,t_per_atom,ii,offset,i,numj,
|
||||||
n_stride,nbor_end,nbor);
|
n_stride,nbor_end,nbor);
|
||||||
|
|
||||||
numtyp4 ix; fetch4(ix,i,pos_tex); //x_[i];
|
numtyp4 ix; fetch4(ix,i,pos_tex); //x_[i];
|
||||||
int itype=ix.w;
|
int itype=ix.w;
|
||||||
|
|
||||||
numtyp factor_lj;
|
numtyp factor_lj;
|
||||||
for ( ; nbor<nbor_end; nbor+=n_stride) {
|
for ( ; nbor<nbor_end; nbor+=n_stride) {
|
||||||
|
|
||||||
int j=dev_packed[nbor];
|
int j=dev_packed[nbor];
|
||||||
factor_lj = sp_lj[sbmask(j)];
|
factor_lj = sp_lj[sbmask(j)];
|
||||||
j &= NEIGHMASK;
|
j &= NEIGHMASK;
|
||||||
|
@ -71,12 +71,12 @@ __kernel void k_ufm(const __global numtyp4 *restrict x_,
|
||||||
numtyp dely = ix.y-jx.y;
|
numtyp dely = ix.y-jx.y;
|
||||||
numtyp delz = ix.z-jx.z;
|
numtyp delz = ix.z-jx.z;
|
||||||
numtyp rsq = delx*delx+dely*dely+delz*delz;
|
numtyp rsq = delx*delx+dely*dely+delz*delz;
|
||||||
|
|
||||||
int mtype=itype*lj_types+jtype;
|
int mtype=itype*lj_types+jtype;
|
||||||
if (rsq<uf1[mtype].z) {
|
if (rsq<uf1[mtype].z) {
|
||||||
numtyp expuf = exp(- rsq * uf1[mtype].y);
|
numtyp expuf = exp(- rsq * uf1[mtype].y);
|
||||||
numtyp force = factor_lj * uf1[mtype].x * expuf / (1.0 - expuf);
|
numtyp force = factor_lj * uf1[mtype].x * expuf / (1.0 - expuf);
|
||||||
|
|
||||||
f.x += delx*force;
|
f.x += delx*force;
|
||||||
f.y += dely*force;
|
f.y += dely*force;
|
||||||
f.z += delz*force;
|
f.z += delz*force;
|
||||||
|
@ -102,17 +102,17 @@ __kernel void k_ufm(const __global numtyp4 *restrict x_,
|
||||||
|
|
||||||
__kernel void k_ufm_fast(const __global numtyp4 *restrict x_,
|
__kernel void k_ufm_fast(const __global numtyp4 *restrict x_,
|
||||||
const __global numtyp4 *restrict uf1_in,
|
const __global numtyp4 *restrict uf1_in,
|
||||||
const __global numtyp4 *restrict uf3_in,
|
const __global numtyp4 *restrict uf3_in,
|
||||||
const __global numtyp *restrict sp_lj_in,
|
const __global numtyp *restrict sp_lj_in,
|
||||||
const __global int * dev_nbor,
|
const __global int * dev_nbor,
|
||||||
const __global int * dev_packed,
|
const __global int * dev_packed,
|
||||||
__global acctyp4 *restrict ans,
|
__global acctyp4 *restrict ans,
|
||||||
__global acctyp *restrict engv,
|
__global acctyp *restrict engv,
|
||||||
const int eflag, const int vflag, const int inum,
|
const int eflag, const int vflag, const int inum,
|
||||||
const int nbor_pitch, const int t_per_atom) {
|
const int nbor_pitch, const int t_per_atom) {
|
||||||
int tid, ii, offset;
|
int tid, ii, offset;
|
||||||
atom_info(t_per_atom,ii,tid,offset);
|
atom_info(t_per_atom,ii,tid,offset);
|
||||||
|
|
||||||
__local numtyp4 uf1[MAX_SHARED_TYPES*MAX_SHARED_TYPES];
|
__local numtyp4 uf1[MAX_SHARED_TYPES*MAX_SHARED_TYPES];
|
||||||
__local numtyp4 uf3[MAX_SHARED_TYPES*MAX_SHARED_TYPES];
|
__local numtyp4 uf3[MAX_SHARED_TYPES*MAX_SHARED_TYPES];
|
||||||
__local numtyp sp_lj[4];
|
__local numtyp sp_lj[4];
|
||||||
|
@ -123,7 +123,7 @@ __kernel void k_ufm_fast(const __global numtyp4 *restrict x_,
|
||||||
if (eflag>0)
|
if (eflag>0)
|
||||||
uf3[tid]=uf3_in[tid];
|
uf3[tid]=uf3_in[tid];
|
||||||
}
|
}
|
||||||
|
|
||||||
acctyp energy=(acctyp)0;
|
acctyp energy=(acctyp)0;
|
||||||
acctyp4 f;
|
acctyp4 f;
|
||||||
f.x=(acctyp)0; f.y=(acctyp)0; f.z=(acctyp)0;
|
f.x=(acctyp)0; f.y=(acctyp)0; f.z=(acctyp)0;
|
||||||
|
@ -132,7 +132,7 @@ __kernel void k_ufm_fast(const __global numtyp4 *restrict x_,
|
||||||
virial[i]=(acctyp)0;
|
virial[i]=(acctyp)0;
|
||||||
|
|
||||||
__syncthreads();
|
__syncthreads();
|
||||||
|
|
||||||
if (ii<inum) {
|
if (ii<inum) {
|
||||||
int i, numj, nbor, nbor_end;
|
int i, numj, nbor, nbor_end;
|
||||||
__local int n_stride;
|
__local int n_stride;
|
||||||
|
@ -145,7 +145,7 @@ __kernel void k_ufm_fast(const __global numtyp4 *restrict x_,
|
||||||
|
|
||||||
numtyp factor_lj;
|
numtyp factor_lj;
|
||||||
for ( ; nbor<nbor_end; nbor+=n_stride) {
|
for ( ; nbor<nbor_end; nbor+=n_stride) {
|
||||||
|
|
||||||
int j=dev_packed[nbor];
|
int j=dev_packed[nbor];
|
||||||
factor_lj = sp_lj[sbmask(j)];
|
factor_lj = sp_lj[sbmask(j)];
|
||||||
j &= NEIGHMASK;
|
j &= NEIGHMASK;
|
||||||
|
@ -158,11 +158,11 @@ __kernel void k_ufm_fast(const __global numtyp4 *restrict x_,
|
||||||
numtyp dely = ix.y-jx.y;
|
numtyp dely = ix.y-jx.y;
|
||||||
numtyp delz = ix.z-jx.z;
|
numtyp delz = ix.z-jx.z;
|
||||||
numtyp rsq = delx*delx+dely*dely+delz*delz;
|
numtyp rsq = delx*delx+dely*dely+delz*delz;
|
||||||
|
|
||||||
if (rsq<uf1[mtype].z) {
|
if (rsq<uf1[mtype].z) {
|
||||||
numtyp expuf = exp(- rsq * uf1[mtype].y);
|
numtyp expuf = exp(- rsq * uf1[mtype].y);
|
||||||
numtyp force = factor_lj * uf1[mtype].x * expuf / (1.0 - expuf);
|
numtyp force = factor_lj * uf1[mtype].x * expuf / (1.0 - expuf);
|
||||||
|
|
||||||
f.x += delx*force;
|
f.x += delx*force;
|
||||||
f.y += dely*force;
|
f.y += dely*force;
|
||||||
f.z += delz*force;
|
f.z += delz*force;
|
||||||
|
|
|
@ -290,7 +290,7 @@ __kernel void k_vashishta(const __global numtyp4 *restrict x_,
|
||||||
|
|
||||||
if (eflag>0)
|
if (eflag>0)
|
||||||
energy += (param3_bigh*reta+vc2-vc3-param3_bigw*r6inv-r*param3_dvrc+param3_c0);
|
energy += (param3_bigh*reta+vc2-vc3-param3_bigw*r6inv-r*param3_dvrc+param3_c0);
|
||||||
|
|
||||||
if (vflag>0) {
|
if (vflag>0) {
|
||||||
virial[0] += delx*delx*force;
|
virial[0] += delx*delx*force;
|
||||||
virial[1] += dely*dely*force;
|
virial[1] += dely*dely*force;
|
||||||
|
@ -471,13 +471,13 @@ __kernel void k_vashishta_three_center(const __global numtyp4 *restrict x_,
|
||||||
numtyp rsq1 = delr1x*delr1x+delr1y*delr1y+delr1z*delr1z;
|
numtyp rsq1 = delr1x*delr1x+delr1y*delr1y+delr1z*delr1z;
|
||||||
|
|
||||||
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; // still keep this for neigh no and tpa > 1
|
||||||
param_gamma_ij=param4_ijparam.y;
|
param_gamma_ij=param4_ijparam.y;
|
||||||
param_r0_ij=param4_ijparam.w;
|
param_r0_ij=param4_ijparam.w;
|
||||||
|
|
||||||
int nbor_k,k_end;
|
int nbor_k,k_end;
|
||||||
if (dev_packed==dev_nbor) {
|
if (dev_packed==dev_nbor) {
|
||||||
nbor_k=nborj_start-offset_j+offset_k;
|
nbor_k=nborj_start-offset_j+offset_k;
|
||||||
|
@ -619,7 +619,7 @@ __kernel void k_vashishta_three_end(const __global numtyp4 *restrict x_,
|
||||||
|
|
||||||
param_gamma_ij=param4_ijparam.y;
|
param_gamma_ij=param4_ijparam.y;
|
||||||
param_r0_ij = param4_ijparam.w;
|
param_r0_ij = param4_ijparam.w;
|
||||||
|
|
||||||
int nbor_k,numk;
|
int nbor_k,numk;
|
||||||
if (dev_nbor==dev_packed) {
|
if (dev_nbor==dev_packed) {
|
||||||
if (gpu_nbor) nbor_k=j+nbor_pitch;
|
if (gpu_nbor) nbor_k=j+nbor_pitch;
|
||||||
|
@ -665,14 +665,14 @@ __kernel void k_vashishta_three_end(const __global numtyp4 *restrict x_,
|
||||||
if (rsq2 < param_r0sq_ik) {
|
if (rsq2 < param_r0sq_ik) {
|
||||||
param_gamma_ik=param4_ikparam.y;
|
param_gamma_ik=param4_ikparam.y;
|
||||||
param_r0_ik=param4_ikparam.w;
|
param_r0_ik=param4_ikparam.w;
|
||||||
|
|
||||||
int ijkparam=elem2param[jtype*nelements*nelements+itype*nelements+ktype]; //jik
|
int ijkparam=elem2param[jtype*nelements*nelements+itype*nelements+ktype]; //jik
|
||||||
numtyp4 param5_ijkparam; fetch4(param5_ijkparam,ijkparam,param5_tex);
|
numtyp4 param5_ijkparam; fetch4(param5_ijkparam,ijkparam,param5_tex);
|
||||||
param_bigc_ijk=param5_ijkparam.x;
|
param_bigc_ijk=param5_ijkparam.x;
|
||||||
param_costheta_ijk=param5_ijkparam.y;
|
param_costheta_ijk=param5_ijkparam.y;
|
||||||
param_bigb_ijk=param5_ijkparam.z;
|
param_bigb_ijk=param5_ijkparam.z;
|
||||||
param_big2b_ijk=param5_ijkparam.w;
|
param_big2b_ijk=param5_ijkparam.w;
|
||||||
|
|
||||||
numtyp fjx, fjy, fjz;
|
numtyp fjx, fjy, fjz;
|
||||||
//if (evatom==0) {
|
//if (evatom==0) {
|
||||||
threebody_half(delr1x,delr1y,delr1z);
|
threebody_half(delr1x,delr1y,delr1z);
|
||||||
|
@ -774,7 +774,7 @@ __kernel void k_vashishta_three_end_vatom(const __global numtyp4 *restrict x_,
|
||||||
|
|
||||||
param_gamma_ij=param4_ijparam.y;
|
param_gamma_ij=param4_ijparam.y;
|
||||||
param_r0_ij=param4_ijparam.w;
|
param_r0_ij=param4_ijparam.w;
|
||||||
|
|
||||||
int nbor_k,numk;
|
int nbor_k,numk;
|
||||||
if (dev_nbor==dev_packed) {
|
if (dev_nbor==dev_packed) {
|
||||||
if (gpu_nbor) nbor_k=j+nbor_pitch;
|
if (gpu_nbor) nbor_k=j+nbor_pitch;
|
||||||
|
@ -827,7 +827,7 @@ __kernel void k_vashishta_three_end_vatom(const __global numtyp4 *restrict x_,
|
||||||
param_costheta_ijk=param5_ijkparam.y;
|
param_costheta_ijk=param5_ijkparam.y;
|
||||||
param_bigb_ijk=param5_ijkparam.z;
|
param_bigb_ijk=param5_ijkparam.z;
|
||||||
param_big2b_ijk=param5_ijkparam.w;
|
param_big2b_ijk=param5_ijkparam.w;
|
||||||
|
|
||||||
numtyp fjx, fjy, fjz, fkx, fky, fkz;
|
numtyp fjx, fjy, fjz, fkx, fky, fkz;
|
||||||
threebody(delr1x,delr1y,delr1z,eflag,energy);
|
threebody(delr1x,delr1y,delr1z,eflag,energy);
|
||||||
|
|
||||||
|
|
Loading…
Reference in New Issue