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

Fixed issue with accessing type2frho array in eam energy kernels

Browse Source
This commit is contained in:
Trung Nguyen
2021-05-04 23:12:55 -05:00
committed by Axel Kohlmeyer
parent 147e561dca
commit b3083f1982
1 changed files with 9 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files

20
lib/gpu/lal_eam.cu
View File

@ -225,7 +225,7 @@ __kernel void k_energy(const __global numtyp4 *restrict x_,
const numtyp rdr, const numtyp rdrho,
const numtyp rhomax, const int nrho,
const int nr, const int t_per_atom) {
int tid, ii, offset, i, itype;
int tid, ii, offset, i, itype, tfrho;
atom_info(t_per_atom,ii,tid,offset);
int n_stride;
@ -242,6 +242,7 @@ __kernel void k_energy(const __global numtyp4 *restrict x_,
numtyp4 ix; fetch4(ix,i,pos_tex); //x_[i];
itype=ix.w;
tfrho=type2frho[itype];
for ( ; nbor<nbor_end; nbor+=n_stride) {
int j=dev_packed[nbor];
@ -270,7 +271,6 @@ __kernel void k_energy(const __global numtyp4 *restrict x_,
}
} // for nbor
} // if ii
const numtyp tfrho=type2frho[itype];
store_energy_fp(rho,energy,ii,inum,tid,t_per_atom,offset,
eflag,vflag,engv,rdrho,nrho,i,rhomax,tfrho);
}
@ -291,7 +291,7 @@ __kernel void k_energy_fast(const __global numtyp4 *restrict x_,
const numtyp rdrho, const numtyp rhomax,
const int nrho, const int nr,
const int t_per_atom) {
int tid, ii, offset, i, itype;
int tid, ii, offset, i, itype, tfrho;
atom_info(t_per_atom,ii,tid,offset);
#ifndef ONETYPE
@ -305,9 +305,9 @@ __kernel void k_energy_fast(const __global numtyp4 *restrict x_,
}
__syncthreads();
#else
const numtyp type2rhor_z2rx=
const int type2rhor_z2rx=
type2rhor_z2r_in[ONETYPE*MAX_SHARED_TYPES+ONETYPE].x;
const numtyp tfrho=type2frho_in[ONETYPE];
tfrho=type2frho_in[ONETYPE];
#endif
int n_stride;
@ -325,6 +325,7 @@ __kernel void k_energy_fast(const __global numtyp4 *restrict x_,
numtyp4 ix; fetch4(ix,i,pos_tex); //x_[i];
#ifndef ONETYPE
itype=ix.w;
tfrho=type2frho[itype];
#endif
for ( ; nbor<nbor_end; nbor+=n_stride) {
@ -347,7 +348,7 @@ __kernel void k_energy_fast(const __global numtyp4 *restrict x_,
p = MIN(p,(numtyp)1.0);
#ifndef ONETYPE
int jtype=fast_mul((int)MAX_SHARED_TYPES,jx.w);
int jtype = fast_mul((int)MAX_SHARED_TYPES,jx.w);
int mtype = jtype+itype;
int index = type2rhor_z2r[mtype].x*(nr+1)+m;
#else
@ -358,9 +359,6 @@ __kernel void k_energy_fast(const __global numtyp4 *restrict x_,
}
} // for nbor
} // if ii
#ifndef ONETYPE
const numtyp tfrho=type2frho[itype];
#endif
store_energy_fp(rho,energy,ii,inum,tid,t_per_atom,offset,
eflag,vflag,engv,rdrho,nrho,i,rhomax,tfrho);
}
@ -498,8 +496,8 @@ __kernel void k_eam_fast(const __global numtyp4 *x_,
__syncthreads();
#else
const int oi=ONETYPE*MAX_SHARED_TYPES+ONETYPE;
const numtyp type2rhor_z2rx=type2rhor_z2r_in[oi].x;
const numtyp type2rhor_z2ry=type2rhor_z2r_in[oi].y;
const int type2rhor_z2rx=type2rhor_z2r_in[oi].x;
const int type2rhor_z2ry=type2rhor_z2r_in[oi].y;
#endif
int n_stride;