From ebfe7f68af4214a9efa78fb5b9f2b817d5a40bd2 Mon Sep 17 00:00:00 2001 From: Trung Nguyen Date: Sun, 5 Jul 2020 00:20:28 -0500 Subject: [PATCH] More cleanup --- lib/gpu/lal_tersoff.cpp | 8 ++++---- lib/gpu/lal_tersoff.cu | 4 +--- lib/gpu/lal_tersoff_mod.cpp | 8 ++++---- lib/gpu/lal_tersoff_mod.cu | 4 +--- lib/gpu/lal_tersoff_zbl.cpp | 8 ++++---- lib/gpu/lal_tersoff_zbl.cu | 4 +--- lib/gpu/lal_vashishta.cpp | 12 ++++++------ lib/gpu/lal_vashishta.cu | 4 +--- 8 files changed, 22 insertions(+), 30 deletions(-) diff --git a/lib/gpu/lal_tersoff.cpp b/lib/gpu/lal_tersoff.cpp index 6574e7dfa6..63691a2047 100644 --- a/lib/gpu/lal_tersoff.cpp +++ b/lib/gpu/lal_tersoff.cpp @@ -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; diff --git a/lib/gpu/lal_tersoff.cu b/lib/gpu/lal_tersoff.cu index 604ba7e64b..7d3ede8dfc 100644 --- a/lib/gpu/lal_tersoff.cu +++ b/lib/gpu/lal_tersoff.cu @@ -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; diff --git a/lib/gpu/lal_tersoff_mod.cpp b/lib/gpu/lal_tersoff_mod.cpp index fa6ae67dfc..2b56991cc6 100644 --- a/lib/gpu/lal_tersoff_mod.cpp +++ b/lib/gpu/lal_tersoff_mod.cpp @@ -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; diff --git a/lib/gpu/lal_tersoff_mod.cu b/lib/gpu/lal_tersoff_mod.cu index ff8625b99d..bb5aabca44 100644 --- a/lib/gpu/lal_tersoff_mod.cu +++ b/lib/gpu/lal_tersoff_mod.cu @@ -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; diff --git a/lib/gpu/lal_tersoff_zbl.cpp b/lib/gpu/lal_tersoff_zbl.cpp index 052167cabd..7d254d568d 100644 --- a/lib/gpu/lal_tersoff_zbl.cpp +++ b/lib/gpu/lal_tersoff_zbl.cpp @@ -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; diff --git a/lib/gpu/lal_tersoff_zbl.cu b/lib/gpu/lal_tersoff_zbl.cu index 1a75cb09e7..cd8d453b3c 100644 --- a/lib/gpu/lal_tersoff_zbl.cu +++ b/lib/gpu/lal_tersoff_zbl.cu @@ -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; diff --git a/lib/gpu/lal_vashishta.cpp b/lib/gpu/lal_vashishta.cpp index e926772ab9..4af8a0f71c 100644 --- a/lib/gpu/lal_vashishta.cpp +++ b/lib/gpu/lal_vashishta.cpp @@ -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 diff --git a/lib/gpu/lal_vashishta.cu b/lib/gpu/lal_vashishta.cu index 9899ba632c..da15aaf09a 100644 --- a/lib/gpu/lal_vashishta.cu +++ b/lib/gpu/lal_vashishta.cu @@ -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;