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

More cleanup

This commit is contained in:
Trung Nguyen 2020-07-05 00:20:28 -05:00
parent c488e6208d
commit ebfe7f68af
8 changed files with 22 additions and 30 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
lib/gpu/lal_tersoff.cpp
View File

@ -250,10 +250,10 @@ void TersoffT::loop(const bool _eflag, const bool _vflag, const int evatom) {
(BX/this->_threads_per_atom)));
this->k_short_nbor.set_size(GX,BX);
this->k_short_nbor.run(&this->atom->x, &elem2param, &_nelements, &_nparams,
&this->nbor->dev_nbor, &this->_nbor_data->begin(),
&this->dev_short_nbor, &_cutshortsq, &ainum,
&nbor_pitch, &this->_threads_per_atom);
this->k_short_nbor.run(&this->atom->x, &this->nbor->dev_nbor,
&this->_nbor_data->begin(),
&this->dev_short_nbor, &_cutshortsq, &ainum,
&nbor_pitch, &this->_threads_per_atom);
// re-allocate zetaij if necessary
int nall = this->_nall;

4
lib/gpu/lal_tersoff.cu
View File

@ -165,12 +165,10 @@ _texture( ts5_tex,int4);
#endif
__kernel void k_tersoff_short_nbor(const __global numtyp4 *restrict x_,
const __global int *restrict elem2param,
const int nelements, const int nparams,
const __global int * dev_nbor,
const __global int * dev_packed,
__global int * dev_short_nbor,
const double _cutshortsq,
const numtyp _cutshortsq,
const int inum, const int nbor_pitch,
const int t_per_atom) {
__local int n_stride;

8
lib/gpu/lal_tersoff_mod.cpp
View File

@ -250,10 +250,10 @@ void TersoffMT::loop(const bool _eflag, const bool _vflag, const int evatom) {
(BX/this->_threads_per_atom)));
this->k_short_nbor.set_size(GX,BX);
this->k_short_nbor.run(&this->atom->x, &elem2param, &_nelements, &_nparams,
&this->nbor->dev_nbor, &this->_nbor_data->begin(),
&this->dev_short_nbor, &_cutshortsq, &ainum,
&nbor_pitch, &this->_threads_per_atom);
this->k_short_nbor.run(&this->atom->x, &this->nbor->dev_nbor,
&this->_nbor_data->begin(),
&this->dev_short_nbor, &_cutshortsq, &ainum,
&nbor_pitch, &this->_threads_per_atom);
// re-allocate zetaij if necessary
int nall = this->_nall;

4
lib/gpu/lal_tersoff_mod.cu
View File

@ -165,12 +165,10 @@ _texture( ts5_tex,int4);
#endif
__kernel void k_tersoff_mod_short_nbor(const __global numtyp4 *restrict x_,
const __global int *restrict elem2param,
const int nelements, const int nparams,
const __global int * dev_nbor,
const __global int * dev_packed,
__global int * dev_short_nbor,
const double _cutshortsq,
const numtyp _cutshortsq,
const int inum, const int nbor_pitch,
const int t_per_atom) {
__local int n_stride;

8
lib/gpu/lal_tersoff_zbl.cpp
View File

@ -275,10 +275,10 @@ void TersoffZT::loop(const bool _eflag, const bool _vflag, const int evatom) {
(BX/this->_threads_per_atom)));
this->k_short_nbor.set_size(GX,BX);
this->k_short_nbor.run(&this->atom->x, &elem2param, &_nelements, &_nparams,
&this->nbor->dev_nbor, &this->_nbor_data->begin(),
&this->dev_short_nbor, &_cutshortsq, &ainum,
&nbor_pitch, &this->_threads_per_atom);
this->k_short_nbor.run(&this->atom->x, &this->nbor->dev_nbor,
&this->_nbor_data->begin(),
&this->dev_short_nbor, &_cutshortsq, &ainum,
&nbor_pitch, &this->_threads_per_atom);
// re-allocate zetaij if necessary
int nall = this->_nall;

4
lib/gpu/lal_tersoff_zbl.cu
View File

@ -168,12 +168,10 @@ _texture( ts6_tex,int4);
#endif
__kernel void k_tersoff_zbl_short_nbor(const __global numtyp4 *restrict x_,
const __global int *restrict elem2param,
const int nelements, const int nparams,
const __global int * dev_nbor,
const __global int * dev_packed,
__global int * dev_short_nbor,
const double _cutshortsq,
const numtyp _cutshortsq,
const int inum, const int nbor_pitch,
const int t_per_atom) {
__local int n_stride;

12
lib/gpu/lal_vashishta.cpp
View File

@ -180,8 +180,8 @@ int VashishtaT::init(const int ntypes, const int nlocal, const int nall, const i
ucl_copy(map,dview_map,false);
_allocated=true;
this->_max_bytes=param1.row_bytes()+param2.row_bytes()+param3.row_bytes()+param4.row_bytes()+param5.row_bytes()+
map.row_bytes()+elem2param.row_bytes();
this->_max_bytes=param1.row_bytes()+param2.row_bytes()+param3.row_bytes()+
param4.row_bytes()+param5.row_bytes()+map.row_bytes()+elem2param.row_bytes();
return 0;
}
@ -233,10 +233,10 @@ void VashishtaT::loop(const bool _eflag, const bool _vflag, const int evatom) {
(BX/this->_threads_per_atom)));
this->k_short_nbor.set_size(GX,BX);
this->k_short_nbor.run(&this->atom->x, &elem2param, &_nelements, &_nparams,
&this->nbor->dev_nbor, &this->_nbor_data->begin(),
&this->dev_short_nbor, &_cutshortsq, &ainum,
&nbor_pitch, &this->_threads_per_atom);
this->k_short_nbor.run(&this->atom->x, &this->nbor->dev_nbor,
&this->_nbor_data->begin(),
&this->dev_short_nbor, &_cutshortsq, &ainum,
&nbor_pitch, &this->_threads_per_atom);
// this->_nbor_data == nbor->dev_packed for gpu_nbor == 0 and tpa > 1
// this->_nbor_data == nbor->dev_nbor for gpu_nbor == 1 or tpa == 1

4
lib/gpu/lal_vashishta.cu
View File

@ -137,12 +137,10 @@ _texture( param5_tex,int4);
#endif
__kernel void k_vashishta_short_nbor(const __global numtyp4 *restrict x_,
const __global int *restrict elem2param,
const int nelements, const int nparams,
const __global int * dev_nbor,
const __global int * dev_packed,
__global int * dev_short_nbor,
const double _cutshortsq,
const numtyp _cutshortsq,
const int inum, const int nbor_pitch,
const int t_per_atom) {
__local int n_stride;