diff --git a/doc/src/pair_born.txt b/doc/src/pair_born.txt index a016f77fa3..f867107426 100644 --- a/doc/src/pair_born.txt +++ b/doc/src/pair_born.txt @@ -12,12 +12,14 @@ pair_style born/omp command :h3 pair_style born/gpu command :h3 pair_style born/coul/long command :h3 pair_style born/coul/long/cs command :h3 +pair_style born/coul/long/cs/gpu command :h3 pair_style born/coul/long/gpu command :h3 pair_style born/coul/long/omp command :h3 pair_style born/coul/msm command :h3 pair_style born/coul/msm/omp command :h3 pair_style born/coul/wolf command :h3 pair_style born/coul/wolf/cs command :h3 +pair_style born/coul/wolf/cs/gpu command :h3 pair_style born/coul/wolf/gpu command :h3 pair_style born/coul/wolf/omp command :h3 pair_style born/coul/dsf command :h3 diff --git a/doc/src/pair_coul.txt b/doc/src/pair_coul.txt index 4cca5ee0d7..aa3a008bd3 100644 --- a/doc/src/pair_coul.txt +++ b/doc/src/pair_coul.txt @@ -20,6 +20,7 @@ pair_style coul/dsf/kk command :h3 pair_style coul/dsf/omp command :h3 pair_style coul/long command :h3 pair_style coul/long/cs command :h3 +pair_style coul/long/cs/gpu command :h3 pair_style coul/long/omp command :h3 pair_style coul/long/gpu command :h3 pair_style coul/long/kk command :h3 diff --git a/lib/gpu/Nvidia.makefile b/lib/gpu/Nvidia.makefile index 1b4cfc469f..0d7214260d 100644 --- a/lib/gpu/Nvidia.makefile +++ b/lib/gpu/Nvidia.makefile @@ -81,7 +81,8 @@ OBJS = $(OBJ_DIR)/lal_atom.o $(OBJ_DIR)/lal_ans.o \ $(OBJ_DIR)/lal_dipole_long_lj.o $(OBJ_DIR)/lal_dipole_long_lj_ext.o \ $(OBJ_DIR)/lal_lj_expand_coul_long.o $(OBJ_DIR)/lal_lj_expand_coul_long_ext.o \ $(OBJ_DIR)/lal_coul_long_cs.o $(OBJ_DIR)/lal_coul_long_cs_ext.o \ - $(OBJ_DIR)/lal_born_coul_long_cs.o $(OBJ_DIR)/lal_born_coul_long_cs_ext.o + $(OBJ_DIR)/lal_born_coul_long_cs.o $(OBJ_DIR)/lal_born_coul_long_cs_ext.o \ + $(OBJ_DIR)/lal_born_coul_wolf_cs.o $(OBJ_DIR)/lal_born_coul_wolf_cs_ext.o CBNS = $(OBJ_DIR)/device.cubin $(OBJ_DIR)/device_cubin.h \ $(OBJ_DIR)/atom.cubin $(OBJ_DIR)/atom_cubin.h \ @@ -141,7 +142,8 @@ CBNS = $(OBJ_DIR)/device.cubin $(OBJ_DIR)/device_cubin.h \ $(OBJ_DIR)/dipole_long_lj.cubin $(OBJ_DIR)/dipole_long_lj_cubin.h \ $(OBJ_DIR)/lj_expand_coul_long.cubin $(OBJ_DIR)/lj_expand_coul_long_cubin.h \ $(OBJ_DIR)/coul_long_cs.cubin $(OBJ_DIR)/coul_long_cs_cubin.h \ - $(OBJ_DIR)/born_coul_long_cs.cubin $(OBJ_DIR)/born_coul_long_cs_cubin.h + $(OBJ_DIR)/born_coul_long_cs.cubin $(OBJ_DIR)/born_coul_long_cs_cubin.h \ + $(OBJ_DIR)/born_coul_wolf_cs.cubin $(OBJ_DIR)/born_coul_wolf_cs_cubin.h all: $(OBJ_DIR) $(GPU_LIB) $(EXECS) @@ -847,10 +849,10 @@ $(OBJ_DIR)/coul_long_cs.cubin: lal_coul_long_cs.cu lal_precision.h lal_preproces $(OBJ_DIR)/coul_long_cs_cubin.h: $(OBJ_DIR)/coul_long_cs.cubin $(OBJ_DIR)/coul_long_cs.cubin $(BIN2C) -c -n coul_long_cs $(OBJ_DIR)/coul_long_cs.cubin > $(OBJ_DIR)/coul_long_cs_cubin.h -$(OBJ_DIR)/lal_coul_long_cs.o: $(ALL_H) lal_coul_long_cs.h lal_coul_long_cs.cpp $(OBJ_DIR)/coul_long_cs_cubin.h $(OBJ_DIR)/lal_base_charge.o +$(OBJ_DIR)/lal_coul_long_cs.o: $(ALL_H) lal_coul_long_cs.h lal_coul_long_cs.cpp $(OBJ_DIR)/coul_long_cs_cubin.h $(OBJ_DIR)/lal_base_charge.o $(OBJ_DIR)/lal_coul_long.o $(CUDR) -o $@ -c lal_coul_long_cs.cpp -I$(OBJ_DIR) -$(OBJ_DIR)/lal_coul_long_cs_ext.o: $(ALL_H) lal_coul_long_cs.h lal_coul_long_cs_ext.cpp lal_base_charge.h +$(OBJ_DIR)/lal_coul_long_cs_ext.o: $(ALL_H) lal_coul_long_cs.h lal_coul_long_cs_ext.cpp lal_coul_long.h $(CUDR) -o $@ -c lal_coul_long_cs_ext.cpp -I$(OBJ_DIR) $(OBJ_DIR)/born_coul_long_cs.cubin: lal_born_coul_long_cs.cu lal_precision.h lal_preprocessor.h @@ -859,12 +861,24 @@ $(OBJ_DIR)/born_coul_long_cs.cubin: lal_born_coul_long_cs.cu lal_precision.h lal $(OBJ_DIR)/born_coul_long_cs_cubin.h: $(OBJ_DIR)/born_coul_long_cs.cubin $(OBJ_DIR)/born_coul_long_cs.cubin $(BIN2C) -c -n born_coul_long_cs $(OBJ_DIR)/born_coul_long_cs.cubin > $(OBJ_DIR)/born_coul_long_cs_cubin.h -$(OBJ_DIR)/lal_born_coul_long_cs.o: $(ALL_H) lal_born_coul_long_cs.h lal_born_coul_long_cs.cpp $(OBJ_DIR)/born_coul_long_cs_cubin.h $(OBJ_DIR)/lal_base_charge.o +$(OBJ_DIR)/lal_born_coul_long_cs.o: $(ALL_H) lal_born_coul_long_cs.h lal_born_coul_long_cs.cpp $(OBJ_DIR)/born_coul_long_cs_cubin.h $(OBJ_DIR)/lal_base_charge.o $(OBJ_DIR)/lal_born_coul_long.o $(CUDR) -o $@ -c lal_born_coul_long_cs.cpp -I$(OBJ_DIR) -$(OBJ_DIR)/lal_born_coul_long_cs_ext.o: $(ALL_H) lal_born_coul_long_cs.h lal_born_coul_long_cs_ext.cpp lal_base_charge.h +$(OBJ_DIR)/lal_born_coul_long_cs_ext.o: $(ALL_H) lal_born_coul_long_cs.h lal_born_coul_long_cs_ext.cpp lal_born_coul_long.h $(CUDR) -o $@ -c lal_born_coul_long_cs_ext.cpp -I$(OBJ_DIR) +$(OBJ_DIR)/born_coul_wolf_cs.cubin: lal_born_coul_wolf_cs.cu lal_precision.h lal_preprocessor.h + $(CUDA) --cubin -DNV_KERNEL -o $@ lal_born_coul_wolf_cs.cu + +$(OBJ_DIR)/born_coul_wolf_cs_cubin.h: $(OBJ_DIR)/born_coul_wolf_cs.cubin $(OBJ_DIR)/born_coul_wolf_cs.cubin + $(BIN2C) -c -n born_coul_wolf_cs $(OBJ_DIR)/born_coul_wolf_cs.cubin > $(OBJ_DIR)/born_coul_wolf_cs_cubin.h + +$(OBJ_DIR)/lal_born_coul_wolf_cs.o: $(ALL_H) lal_born_coul_wolf_cs.h lal_born_coul_wolf_cs.cpp $(OBJ_DIR)/born_coul_wolf_cs_cubin.h $(OBJ_DIR)/lal_base_charge.o $(OBJ_DIR)/lal_born_coul_wolf.o + $(CUDR) -o $@ -c lal_born_coul_wolf_cs.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_born_coul_wolf_cs_ext.o: $(ALL_H) lal_born_coul_wolf_cs.h lal_born_coul_wolf_cs_ext.cpp lal_born_coul_wolf.h + $(CUDR) -o $@ -c lal_born_coul_wolf_cs_ext.cpp -I$(OBJ_DIR) + $(BIN_DIR)/nvc_get_devices: ./geryon/ucl_get_devices.cpp $(NVD_H) $(CUDR) -o $@ ./geryon/ucl_get_devices.cpp -DUCL_CUDADR $(CUDA_LIB) -lcuda diff --git a/lib/gpu/Opencl.makefile b/lib/gpu/Opencl.makefile index e6e84a0cf4..86991a1e98 100644 --- a/lib/gpu/Opencl.makefile +++ b/lib/gpu/Opencl.makefile @@ -70,7 +70,8 @@ OBJS = $(OBJ_DIR)/lal_atom.o $(OBJ_DIR)/lal_answer.o \ $(OBJ_DIR)/lal_dipole_long_lj.o $(OBJ_DIR)/lal_dipole_long_lj_ext.o \ $(OBJ_DIR)/lal_lj_expand_coul_long.o $(OBJ_DIR)/lal_lj_expand_coul_long_ext.o \ $(OBJ_DIR)/lal_coul_long_cs.o $(OBJ_DIR)/lal_coul_long_cs_ext.o \ - $(OBJ_DIR)/lal_born_coul_long_cs.o $(OBJ_DIR)/lal_born_coul_long_cs_ext.o + $(OBJ_DIR)/lal_born_coul_long_cs.o $(OBJ_DIR)/lal_born_coul_long_cs_ext.o \ + $(OBJ_DIR)/lal_born_coul_wolf_cs.o $(OBJ_DIR)/lal_born_coul_wolf_cs_ext.o KERS = $(OBJ_DIR)/device_cl.h $(OBJ_DIR)/atom_cl.h \ $(OBJ_DIR)/neighbor_cpu_cl.h $(OBJ_DIR)/pppm_cl.h \ @@ -101,7 +102,7 @@ KERS = $(OBJ_DIR)/device_cl.h $(OBJ_DIR)/atom_cl.h \ $(OBJ_DIR)/lj_cubic_cl.h $(OBJ_DIR)/vashishta_cl.h \ $(OBJ_DIR)/ufm_cl.h $(OBJ_DIR)/dipole_long_lj_cl.h \ $(OBJ_DIR)/lj_expand_coul_long_cl.h $(OBJ_DIR)/coul_long_cs_cl.h \ - $(OBJ_DIR)/born_coul_long_cs_cl.h + $(OBJ_DIR)/born_coul_long_cs_cl.h $(OBJ_DIR)/born_coul_wolf_cs_cl.h OCL_EXECS = $(BIN_DIR)/ocl_get_devices @@ -615,21 +616,30 @@ $(OBJ_DIR)/lal_lj_expand_coul_long_ext.o: $(ALL_H) lal_lj_expand_coul_long.h lal $(OBJ_DIR)/coul_long_cs_cl.h: lal_coul_long_cs.cu $(PRE1_H) $(BSH) ./geryon/file_to_cstr.sh coul_long_cs $(PRE1_H) lal_coul_long_cs.cu $(OBJ_DIR)/coul_long_cs_cl.h; -$(OBJ_DIR)/lal_coul_long_cs.o: $(ALL_H) lal_coul_long_cs.h lal_coul_long_cs.cpp $(OBJ_DIR)/coul_long_cs_cl.h $(OBJ_DIR)/coul_long_cs_cl.h $(OBJ_DIR)/lal_base_charge.o +$(OBJ_DIR)/lal_coul_long_cs.o: $(ALL_H) lal_coul_long_cs.h lal_coul_long_cs.cpp $(OBJ_DIR)/coul_long_cs_cl.h $(OBJ_DIR)/coul_long_cs_cl.h $(OBJ_DIR)/lal_base_charge.o $(OBJ_DIR)/lal_coul_long.o $(OCL) -o $@ -c lal_coul_long_cs.cpp -I$(OBJ_DIR) -$(OBJ_DIR)/lal_coul_long_cs_ext.o: $(ALL_H) lal_coul_long_cs.h lal_coul_long_cs_ext.cpp lal_base_charge.h +$(OBJ_DIR)/lal_coul_long_cs_ext.o: $(ALL_H) lal_coul_long_cs.h lal_coul_long_cs_ext.cpp lal_coul_long.h $(OCL) -o $@ -c lal_coul_long_cs_ext.cpp -I$(OBJ_DIR) $(OBJ_DIR)/born_coul_long_cs_cl.h: lal_born_coul_long_cs.cu $(PRE1_H) $(BSH) ./geryon/file_to_cstr.sh born_coul_long_cs $(PRE1_H) lal_born_coul_long_cs.cu $(OBJ_DIR)/born_coul_long_cs_cl.h; -$(OBJ_DIR)/lal_born_coul_long_cs.o: $(ALL_H) lal_born_coul_long_cs.h lal_born_coul_long_cs.cpp $(OBJ_DIR)/born_coul_long_cs_cl.h $(OBJ_DIR)/born_coul_long_cs_cl.h $(OBJ_DIR)/lal_base_charge.o +$(OBJ_DIR)/lal_born_coul_long_cs.o: $(ALL_H) lal_born_coul_long_cs.h lal_born_coul_long_cs.cpp $(OBJ_DIR)/born_coul_long_cs_cl.h $(OBJ_DIR)/born_coul_long_cs_cl.h $(OBJ_DIR)/lal_base_charge.o $(OBJ_DIR)/lal_born_coul_long.o $(OCL) -o $@ -c lal_born_coul_long_cs.cpp -I$(OBJ_DIR) -$(OBJ_DIR)/lal_born_coul_long_cs_ext.o: $(ALL_H) lal_born_coul_long_cs.h lal_born_coul_long_cs_ext.cpp lal_base_charge.h +$(OBJ_DIR)/lal_born_coul_long_cs_ext.o: $(ALL_H) lal_born_coul_long_cs.h lal_born_coul_long_cs_ext.cpp lal_born_coul_long.h $(OCL) -o $@ -c lal_born_coul_long_cs_ext.cpp -I$(OBJ_DIR) +$(OBJ_DIR)/born_coul_wolf_cs_cl.h: lal_born_coul_wolf_cs.cu $(PRE1_H) + $(BSH) ./geryon/file_to_cstr.sh born_coul_wolf_cs $(PRE1_H) lal_born_coul_wolf_cs.cu $(OBJ_DIR)/born_coul_wolf_cs_cl.h; + +$(OBJ_DIR)/lal_born_coul_wolf_cs.o: $(ALL_H) lal_born_coul_wolf_cs.h lal_born_coul_wolf_cs.cpp $(OBJ_DIR)/born_coul_wolf_cs_cl.h $(OBJ_DIR)/born_coul_wolf_cs_cl.h $(OBJ_DIR)/lal_base_charge.o $(OBJ_DIR)/lal_born_coul_wolf.o + $(OCL) -o $@ -c lal_born_coul_wolf_cs.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_born_coul_wolf_cs_ext.o: $(ALL_H) lal_born_coul_wolf_cs.h lal_born_coul_wolf_cs_ext.cpp lal_born_coul_wolf.h + $(OCL) -o $@ -c lal_born_coul_wolf_cs_ext.cpp -I$(OBJ_DIR) + $(BIN_DIR)/ocl_get_devices: ./geryon/ucl_get_devices.cpp $(OCL_H) $(OCL) -o $@ ./geryon/ucl_get_devices.cpp -DUCL_OPENCL $(OCL_LINK) diff --git a/lib/gpu/lal_born_coul_long.h b/lib/gpu/lal_born_coul_long.h index e0de27c71c..d236b7b57f 100644 --- a/lib/gpu/lal_born_coul_long.h +++ b/lib/gpu/lal_born_coul_long.h @@ -78,7 +78,7 @@ class BornCoulLong : public BaseCharge { numtyp _cut_coulsq, _qqrd2e, _g_ewald; - private: + protected: bool _allocated; void loop(const bool _eflag, const bool _vflag); }; diff --git a/lib/gpu/lal_born_coul_long_cs.cpp b/lib/gpu/lal_born_coul_long_cs.cpp index 99e71a75d4..e7fb946f14 100644 --- a/lib/gpu/lal_born_coul_long_cs.cpp +++ b/lib/gpu/lal_born_coul_long_cs.cpp @@ -10,7 +10,7 @@ __________________________________________________________________________ begin : - email : trung.nguyen@northwestern.edu + email : ndactrung@gmail.com ***************************************************************************/ #ifdef USE_OPENCL @@ -28,21 +28,6 @@ using namespace LAMMPS_AL; extern Device device; -template -BornCoulLongCST::BornCoulLongCS() : BaseCharge(), - _allocated(false) { -} - -template -BornCoulLongCST::~BornCoulLongCS() { - clear(); -} - -template -int BornCoulLongCST::bytes_per_atom(const int max_nbors) const { - return this->bytes_per_atom_atomic(max_nbors); -} - template int BornCoulLongCST::init(const int ntypes, double **host_cutsq, double **host_rhoinv, double **host_born1, double **host_born2, double **host_born3, @@ -63,13 +48,13 @@ int BornCoulLongCST::init(const int ntypes, double **host_cutsq, double **host_r // If atom type constants fit in shared memory use fast kernel int lj_types=ntypes; - shared_types=false; + this->shared_types=false; int max_shared_types=this->device->max_shared_types(); if (lj_types<=max_shared_types && this->_block_size>=max_shared_types) { lj_types=max_shared_types; - shared_types=true; + this->shared_types=true; } - _lj_types=lj_types; + this->_lj_types=lj_types; // Allocate a host write buffer for data initialization UCL_H_Vec host_write(lj_types*lj_types*32,*(this->ucl_device), @@ -78,98 +63,33 @@ int BornCoulLongCST::init(const int ntypes, double **host_cutsq, double **host_r for (int i=0; iucl_device),UCL_READ_ONLY); - this->atom->type_pack4(ntypes,lj_types,coeff1,host_write,host_rhoinv, + this->coeff1.alloc(lj_types*lj_types,*(this->ucl_device),UCL_READ_ONLY); + this->atom->type_pack4(ntypes,lj_types,this->coeff1,host_write,host_rhoinv, host_born1,host_born2,host_born3); - coeff2.alloc(lj_types*lj_types,*(this->ucl_device),UCL_READ_ONLY); - this->atom->type_pack4(ntypes,lj_types,coeff2,host_write,host_a,host_c, + this->coeff2.alloc(lj_types*lj_types,*(this->ucl_device),UCL_READ_ONLY); + this->atom->type_pack4(ntypes,lj_types,this->coeff2,host_write,host_a,host_c, host_d,host_offset); - cutsq_sigma.alloc(lj_types*lj_types,*(this->ucl_device),UCL_READ_ONLY); - this->atom->type_pack4(ntypes,lj_types,cutsq_sigma,host_write,host_cutsq, + this->cutsq_sigma.alloc(lj_types*lj_types,*(this->ucl_device),UCL_READ_ONLY); + this->atom->type_pack4(ntypes,lj_types,this->cutsq_sigma,host_write,host_cutsq, host_cut_ljsq,host_sigma); - sp_lj.alloc(8,*(this->ucl_device),UCL_READ_ONLY); + this->sp_lj.alloc(8,*(this->ucl_device),UCL_READ_ONLY); for (int i=0; i<4; i++) { host_write[i]=host_special_lj[i]; host_write[i+4]=host_special_coul[i]; } - ucl_copy(sp_lj,host_write,8,false); + ucl_copy(this->sp_lj,host_write,8,false); - _cut_coulsq=host_cut_coulsq; - _qqrd2e=qqrd2e; - _g_ewald=g_ewald; + this->_cut_coulsq=host_cut_coulsq; + this->_qqrd2e=qqrd2e; + this->_g_ewald=g_ewald; - _allocated=true; - this->_max_bytes=coeff1.row_bytes()+coeff2.row_bytes() - +cutsq_sigma.row_bytes()+sp_lj.row_bytes(); + this->_allocated=true; + this->_max_bytes=this->coeff1.row_bytes()+this->coeff2.row_bytes() + +this->cutsq_sigma.row_bytes()+this->sp_lj.row_bytes(); return 0; } -template -void BornCoulLongCST::clear() { - if (!_allocated) - return; - _allocated=false; - - coeff1.clear(); - coeff2.clear(); - cutsq_sigma.clear(); - sp_lj.clear(); - this->clear_atomic(); -} - -template -double BornCoulLongCST::host_memory_usage() const { - return this->host_memory_usage_atomic()+sizeof(BornCoulLongCS); -} - -// --------------------------------------------------------------------------- -// Calculate energies, forces, and torques -// --------------------------------------------------------------------------- -template -void BornCoulLongCST::loop(const bool _eflag, const bool _vflag) { - // Compute the block size and grid size to keep all cores busy - const int BX=this->block_size(); - int eflag, vflag; - if (_eflag) - eflag=1; - else - eflag=0; - - if (_vflag) - vflag=1; - else - vflag=0; - - int GX=static_cast(ceil(static_cast(this->ans->inum())/ - (BX/this->_threads_per_atom))); - - int ainum=this->ans->inum(); - int nbor_pitch=this->nbor->nbor_pitch(); - this->time_pair.start(); - if (shared_types) { - this->k_pair_fast.set_size(GX,BX); - this->k_pair_fast.run(&this->atom->x, &coeff1, &coeff2, &sp_lj, - &this->nbor->dev_nbor, - &this->_nbor_data->begin(), - &this->ans->force, - &this->ans->engv, &eflag, &vflag, - &ainum, &nbor_pitch, &this->atom->q, - &cutsq_sigma, &_cut_coulsq, &_qqrd2e, - &_g_ewald, &this->_threads_per_atom); - } else { - this->k_pair.set_size(GX,BX); - this->k_pair.run(&this->atom->x, &coeff1, &coeff2, &_lj_types, &sp_lj, - &this->nbor->dev_nbor, &this->_nbor_data->begin(), - &this->ans->force, &this->ans->engv, - &eflag, &vflag, &ainum, - &nbor_pitch, &this->atom->q, - &cutsq_sigma, &_cut_coulsq, - &_qqrd2e, &_g_ewald, &this->_threads_per_atom); - } - this->time_pair.stop(); -} - template class BornCoulLongCS; diff --git a/lib/gpu/lal_born_coul_long_cs.cu b/lib/gpu/lal_born_coul_long_cs.cu index a6821e1cf7..b3e79d9ec8 100644 --- a/lib/gpu/lal_born_coul_long_cs.cu +++ b/lib/gpu/lal_born_coul_long_cs.cu @@ -10,7 +10,7 @@ // __________________________________________________________________________ // // begin : June 2018 -// email : trung.nguyen@northwestern.edu +// email : ndactrung@gmail.com // ***************************************************************************/ #ifdef NV_KERNEL diff --git a/lib/gpu/lal_born_coul_long_cs.h b/lib/gpu/lal_born_coul_long_cs.h index 2b383dbea9..df0fa0fcd4 100644 --- a/lib/gpu/lal_born_coul_long_cs.h +++ b/lib/gpu/lal_born_coul_long_cs.h @@ -1,7 +1,7 @@ /*************************************************************************** - born_coul_long.h + born_coul_long_cs.h ------------------- - Trung Dac Nguyen (ORNL) + Trung Dac Nguyen (Northwestern) Class for acceleration of the born/coul/long/cs pair style. @@ -10,21 +10,21 @@ __________________________________________________________________________ begin : - email : trung.nguyen@northwestern.edu + email : ndactrung@gmail.com ***************************************************************************/ #ifndef LAL_BORN_COUL_LONG_CS_H #define LAL_BORN_COUL_LONG_CS_H -#include "lal_base_charge.h" +#include "lal_born_coul_long.h" namespace LAMMPS_AL { template -class BornCoulLongCS : public BaseCharge { +class BornCoulLongCS : public BornCoulLong { public: - BornCoulLongCS(); - ~BornCoulLongCS(); + BornCoulLongCS() {} + ~BornCoulLongCS() {} /// Clear any previous data and set up for a new LAMMPS run /** \param max_nbors initial number of rows in the neighbor matrix @@ -46,41 +46,6 @@ class BornCoulLongCS : public BaseCharge { const double gpu_split, FILE *screen, double **host_cut_ljsq, const double host_cut_coulsq, double *host_special_coul, const double qqrd2e, const double g_ewald); - - /// Clear all host and device data - /** \note This is called at the beginning of the init() routine **/ - void clear(); - - /// Returns memory usage on device per atom - int bytes_per_atom(const int max_nbors) const; - - /// Total host memory used by library for pair style - double host_memory_usage() const; - - // --------------------------- TYPE DATA -------------------------- - - /// coeff1.x = rhoinv, coeff1.y = born1, coeff1.z = born2, - /// coeff1.w = born3 - UCL_D_Vec coeff1; - /// coeff2.x = a, coeff2.y = c, coeff2.z = d, coeff2.w = offset - UCL_D_Vec coeff2; - /// cutsq_sigma.x = cutsq, cutsq_sigma.y = cutsq_lj, - /// cutsq_sigma.z = sigma - UCL_D_Vec cutsq_sigma; - /// Special LJ values [0-3] and Special Coul values [4-7] - UCL_D_Vec sp_lj; - - /// If atom type constants fit in shared memory, use fast kernels - bool shared_types; - - /// Number of atom types - int _lj_types; - - numtyp _cut_coulsq, _qqrd2e, _g_ewald; - - private: - bool _allocated; - void loop(const bool _eflag, const bool _vflag); }; } diff --git a/lib/gpu/lal_born_coul_long_cs_ext.cpp b/lib/gpu/lal_born_coul_long_cs_ext.cpp index 43363e217e..badc8b0808 100644 --- a/lib/gpu/lal_born_coul_long_cs_ext.cpp +++ b/lib/gpu/lal_born_coul_long_cs_ext.cpp @@ -10,7 +10,7 @@ __________________________________________________________________________ begin : - email : trung.nguyen@northwestern.edu + email : ndactrung@gmail.com ***************************************************************************/ #include diff --git a/lib/gpu/lal_born_coul_wolf.cpp b/lib/gpu/lal_born_coul_wolf.cpp index 7ebd7b744f..c44b841463 100644 --- a/lib/gpu/lal_born_coul_wolf.cpp +++ b/lib/gpu/lal_born_coul_wolf.cpp @@ -57,7 +57,7 @@ int BornCoulWolfT::init(const int ntypes, double **host_cutsq, double **host_rho const double alf, const double e_shift, const double f_shift) { int success; success=this->init_atomic(nlocal,nall,max_nbors,maxspecial,cell_size,gpu_split, - _screen,born_coul_wolf,"k_born_wolf"); + _screen,born_coul_wolf,"k_born_coul_wolf"); if (success!=0) return success; diff --git a/lib/gpu/lal_born_coul_wolf.cu b/lib/gpu/lal_born_coul_wolf.cu index 0dc7d08c63..2c2249feeb 100644 --- a/lib/gpu/lal_born_coul_wolf.cu +++ b/lib/gpu/lal_born_coul_wolf.cu @@ -31,7 +31,7 @@ texture q_tex; #define MY_PIS (acctyp)1.77245385090551602729 -__kernel void k_born_wolf(const __global numtyp4 *restrict x_, +__kernel void k_born_coul_wolf(const __global numtyp4 *restrict x_, const __global numtyp4 *restrict coeff1, const __global numtyp4 *restrict coeff2, const int lj_types, @@ -165,7 +165,7 @@ __kernel void k_born_wolf(const __global numtyp4 *restrict x_, } // if ii } -__kernel void k_born_wolf_fast(const __global numtyp4 *restrict x_, +__kernel void k_born_coul_wolf_fast(const __global numtyp4 *restrict x_, const __global numtyp4 *restrict coeff1_in, const __global numtyp4 *restrict coeff2_in, const __global numtyp *restrict sp_lj_in, diff --git a/lib/gpu/lal_born_coul_wolf.h b/lib/gpu/lal_born_coul_wolf.h index 4b2406b989..4b817ee0ce 100644 --- a/lib/gpu/lal_born_coul_wolf.h +++ b/lib/gpu/lal_born_coul_wolf.h @@ -13,8 +13,8 @@ email : nguyentd@ornl.gov ***************************************************************************/ -#ifndef LAL_BORN_COUL_LONG_H -#define LAL_BORN_COUL_LONG_H +#ifndef LAL_BORN_COUL_WOLF_H +#define LAL_BORN_COUL_WOLF_H #include "lal_base_charge.h" @@ -79,7 +79,7 @@ class BornCoulWolf : public BaseCharge { numtyp _cut_coulsq,_qqrd2e,_alf,_e_shift,_f_shift; - private: + protected: bool _allocated; void loop(const bool _eflag, const bool _vflag); }; diff --git a/lib/gpu/lal_coul_long.h b/lib/gpu/lal_coul_long.h index d12198fccc..ae25b91dc7 100644 --- a/lib/gpu/lal_coul_long.h +++ b/lib/gpu/lal_coul_long.h @@ -72,7 +72,7 @@ class CoulLong : public BaseCharge { numtyp _cut_coulsq, _qqrd2e, _g_ewald; - private: + protected: bool _allocated; void loop(const bool _eflag, const bool _vflag); }; diff --git a/lib/gpu/lal_coul_long_cs.cpp b/lib/gpu/lal_coul_long_cs.cpp index 7b88647fcc..7afa0ae5d2 100644 --- a/lib/gpu/lal_coul_long_cs.cpp +++ b/lib/gpu/lal_coul_long_cs.cpp @@ -10,7 +10,7 @@ __________________________________________________________________________ begin : June 2018 - email : trung.nguyen@northwestern.edu + email : ndactrung@gmail.com ***************************************************************************/ #if defined(USE_OPENCL) @@ -28,20 +28,6 @@ using namespace LAMMPS_AL; extern Device pair_gpu_device; -template -CoulLongCST::CoulLongCS() : BaseCharge(), _allocated(false) { -} - -template -CoulLongCST::~CoulLongCS() { - clear(); -} - -template -int CoulLongCST::bytes_per_atom(const int max_nbors) const { - return this->bytes_per_atom_atomic(max_nbors); -} - template int CoulLongCST::init(const int ntypes, double **host_scale, const int nlocal, const int nall, const int max_nbors, @@ -56,13 +42,13 @@ int CoulLongCST::init(const int ntypes, double **host_scale, return success; int lj_types=ntypes; - shared_types=false; + this->shared_types=false; int max_shared_types=this->device->max_shared_types(); if (lj_types<=max_shared_types && this->_block_size>=max_shared_types) { lj_types=max_shared_types; - shared_types=true; + this->shared_types=true; } - _lj_types=lj_types; + this->_lj_types=lj_types; // Allocate a host write buffer for data initialization UCL_H_Vec host_write(lj_types*lj_types*32,*(this->ucl_device), @@ -71,88 +57,22 @@ int CoulLongCST::init(const int ntypes, double **host_scale, for (int i=0; iucl_device),UCL_READ_ONLY); - this->atom->type_pack1(ntypes,lj_types,scale,host_write,host_scale); + this->scale.alloc(lj_types*lj_types,*(this->ucl_device),UCL_READ_ONLY); + this->atom->type_pack1(ntypes,lj_types,this->scale,host_write,host_scale); - sp_cl.alloc(4,*(this->ucl_device),UCL_READ_ONLY); + this->sp_cl.alloc(4,*(this->ucl_device),UCL_READ_ONLY); for (int i=0; i<4; i++) { host_write[i]=host_special_coul[i]; } - ucl_copy(sp_cl,host_write,4,false); + ucl_copy(this->sp_cl,host_write,4,false); - _cut_coulsq=host_cut_coulsq; - _qqrd2e=qqrd2e; - _g_ewald=g_ewald; + this->_cut_coulsq=host_cut_coulsq; + this->_qqrd2e=qqrd2e; + this->_g_ewald=g_ewald; - _allocated=true; - this->_max_bytes=scale.row_bytes()+sp_cl.row_bytes(); + this->_allocated=true; + this->_max_bytes=this->scale.row_bytes()+this->sp_cl.row_bytes(); return 0; } -template -void CoulLongCST::reinit(const int ntypes, double **host_scale) { - UCL_H_Vec hscale(_lj_types*_lj_types,*(this->ucl_device), - UCL_WRITE_ONLY); - this->atom->type_pack1(ntypes,_lj_types,scale,hscale,host_scale); -} - -template -void CoulLongCST::clear() { - if (!_allocated) - return; - _allocated=false; - - scale.clear(); - sp_cl.clear(); - this->clear_atomic(); -} - -template -double CoulLongCST::host_memory_usage() const { - return this->host_memory_usage_atomic()+sizeof(CoulLongCS); -} - -// --------------------------------------------------------------------------- -// Calculate energies, forces, and torques -// --------------------------------------------------------------------------- -template -void CoulLongCST::loop(const bool _eflag, const bool _vflag) { - // Compute the block size and grid size to keep all cores busy - const int BX=this->block_size(); - int eflag, vflag; - if (_eflag) - eflag=1; - else - eflag=0; - - if (_vflag) - vflag=1; - else - vflag=0; - - int GX=static_cast(ceil(static_cast(this->ans->inum())/ - (BX/this->_threads_per_atom))); - - int ainum=this->ans->inum(); - int nbor_pitch=this->nbor->nbor_pitch(); - this->time_pair.start(); - if (shared_types) { - this->k_pair_fast.set_size(GX,BX); - this->k_pair_fast.run(&this->atom->x, &scale, &sp_cl, - &this->nbor->dev_nbor, &this->_nbor_data->begin(), - &this->ans->force, &this->ans->engv, - &eflag, &vflag, &ainum, &nbor_pitch, - &this->atom->q, &_cut_coulsq, &_qqrd2e, &_g_ewald, - &this->_threads_per_atom); - } else { - this->k_pair.set_size(GX,BX); - this->k_pair.run(&this->atom->x, &scale, &_lj_types, &sp_cl, - &this->nbor->dev_nbor, &this->_nbor_data->begin(), - &this->ans->force, &this->ans->engv, &eflag, &vflag, - &ainum, &nbor_pitch, &this->atom->q, &_cut_coulsq, - &_qqrd2e, &_g_ewald, &this->_threads_per_atom); - } - this->time_pair.stop(); -} - template class CoulLongCS; diff --git a/lib/gpu/lal_coul_long_cs.cu b/lib/gpu/lal_coul_long_cs.cu index c038766129..3c34666131 100644 --- a/lib/gpu/lal_coul_long_cs.cu +++ b/lib/gpu/lal_coul_long_cs.cu @@ -10,7 +10,7 @@ // __________________________________________________________________________ // // begin : June 2018 -// email : trung.nguyen@northwestern.edu +// email : ndactrung@gmail.com // ***************************************************************************/ #ifdef NV_KERNEL diff --git a/lib/gpu/lal_coul_long_cs.h b/lib/gpu/lal_coul_long_cs.h index 7c7d6fbef1..3cfcb80fc8 100644 --- a/lib/gpu/lal_coul_long_cs.h +++ b/lib/gpu/lal_coul_long_cs.h @@ -10,21 +10,21 @@ __________________________________________________________________________ begin : June 2018 - email : trung.nguyen@northwestern.edu + email : ndactrung@gmail.com ***************************************************************************/ #ifndef LAL_COUL_LONG_CS_H #define LAL_COUL_LONG_CS_H -#include "lal_base_charge.h" +#include "lal_coul_long.h" namespace LAMMPS_AL { template -class CoulLongCS : public BaseCharge { +class CoulLongCS : public CoulLong { public: - CoulLongCS(); - ~CoulLongCS(); + CoulLongCS() {} + ~CoulLongCS() {} /// Clear any previous data and set up for a new LAMMPS run /** \param max_nbors initial number of rows in the neighbor matrix @@ -43,38 +43,6 @@ class CoulLongCS : public BaseCharge { const double gpu_split, FILE *screen, const double host_cut_coulsq, double *host_special_coul, const double qqrd2e, const double g_ewald); - - /// Send updated coeffs from host to device (to be compatible with fix adapt) - void reinit(const int ntypes, double **scale); - - /// Clear all host and device data - /** \note This is called at the beginning of the init() routine **/ - void clear(); - - /// Returns memory usage on device per atom - int bytes_per_atom(const int max_nbors) const; - - /// Total host memory used by library for pair style - double host_memory_usage() const; - - // --------------------------- TYPE DATA -------------------------- - - /// scale - UCL_D_Vec scale; - /// Special Coul values [0-3] - UCL_D_Vec sp_cl; - - /// If atom type constants fit in shared memory, use fast kernels - bool shared_types; - - /// Number of atom types - int _lj_types; - - numtyp _cut_coulsq, _qqrd2e, _g_ewald; - - private: - bool _allocated; - void loop(const bool _eflag, const bool _vflag); }; } diff --git a/lib/gpu/lal_coul_long_cs_ext.cpp b/lib/gpu/lal_coul_long_cs_ext.cpp index 25619a4921..ae57eb2038 100644 --- a/lib/gpu/lal_coul_long_cs_ext.cpp +++ b/lib/gpu/lal_coul_long_cs_ext.cpp @@ -10,7 +10,7 @@ __________________________________________________________________________ begin : June 2018 - email : trung.nguyen@northwestern.edu + email : ndactrung@gmail.com ***************************************************************************/ #include diff --git a/lib/gpu/lal_lj_expand_coul_long.cpp b/lib/gpu/lal_lj_expand_coul_long.cpp index 6a3da9b686..32a4f8120e 100644 --- a/lib/gpu/lal_lj_expand_coul_long.cpp +++ b/lib/gpu/lal_lj_expand_coul_long.cpp @@ -10,7 +10,7 @@ __________________________________________________________________________ begin : - email : trung.nguyen@northwestern.edu + email : ndactrung@gmail.com ***************************************************************************/ #if defined(USE_OPENCL) diff --git a/lib/gpu/lal_lj_expand_coul_long.cu b/lib/gpu/lal_lj_expand_coul_long.cu index bd8e8bfe2a..aa8f02be8c 100644 --- a/lib/gpu/lal_lj_expand_coul_long.cu +++ b/lib/gpu/lal_lj_expand_coul_long.cu @@ -10,7 +10,7 @@ // __________________________________________________________________________ // // begin : -// email : trung.nguyen@northwestern.edu +// email : ndactrung@gmail.com // ***************************************************************************/ #ifdef NV_KERNEL diff --git a/lib/gpu/lal_lj_expand_coul_long.h b/lib/gpu/lal_lj_expand_coul_long.h index c5c7a468a9..64fd9df1ab 100644 --- a/lib/gpu/lal_lj_expand_coul_long.h +++ b/lib/gpu/lal_lj_expand_coul_long.h @@ -10,7 +10,7 @@ __________________________________________________________________________ begin : - email : trung.nguyen@northwestern.edu + email : ndactrung@gmail.com ***************************************************************************/ #ifndef LAL_LJ_EXPAND_COUL_LONG_H diff --git a/lib/gpu/lal_lj_expand_coul_long_ext.cpp b/lib/gpu/lal_lj_expand_coul_long_ext.cpp index e4b53bbf20..3ff1bef701 100644 --- a/lib/gpu/lal_lj_expand_coul_long_ext.cpp +++ b/lib/gpu/lal_lj_expand_coul_long_ext.cpp @@ -1,5 +1,5 @@ /*************************************************************************** - lj_coul_long_ext.cpp + lj_expand_coul_long_ext.cpp ------------------------ Trung Nguyen (Northwestern) @@ -10,7 +10,7 @@ __________________________________________________________________________ begin : - email : trung.nguyen@northwestern.edu + email : ndactrung@gmail.com ***************************************************************************/ #include diff --git a/src/GPU/Install.sh b/src/GPU/Install.sh index 9b029d15cb..20429009db 100755 --- a/src/GPU/Install.sh +++ b/src/GPU/Install.sh @@ -39,6 +39,8 @@ action pair_born_coul_long_cs_gpu.cpp pair_born_coul_long_cs.cpp action pair_born_coul_long_cs_gpu.h pair_born_coul_long_cs.cpp action pair_born_coul_wolf_gpu.cpp action pair_born_coul_wolf_gpu.h +action pair_born_coul_wolf_cs_gpu.cpp pair_born_coul_wolf_cs.cpp +action pair_born_coul_wolf_cs_gpu.h pair_born_coul_wolf_cs.cpp action pair_born_gpu.cpp action pair_born_gpu.h action pair_buck_coul_cut_gpu.cpp pair_buck_coul_cut.cpp diff --git a/src/GPU/pair_born_coul_wolf_cs_gpu.cpp b/src/GPU/pair_born_coul_wolf_cs_gpu.cpp new file mode 100644 index 0000000000..90ddf825a9 --- /dev/null +++ b/src/GPU/pair_born_coul_wolf_cs_gpu.cpp @@ -0,0 +1,295 @@ +/* ---------------------------------------------------------------------- + LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator + http://lammps.sandia.gov, Sandia National Laboratories + Steve Plimpton, sjplimp@sandia.gov + + Copyright (2003) Sandia Corporation. Under the terms of Contract + DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains + certain rights in this software. This software is distributed under + the GNU General Public License. + + See the README file in the top-level LAMMPS directory. +------------------------------------------------------------------------- */ + +/* ---------------------------------------------------------------------- + Contributing authors: Trung Dac Nguyen (Northwestern) +------------------------------------------------------------------------- */ + +#include +#include +#include +#include +#include "pair_born_coul_wolf_cs_gpu.h" +#include "atom.h" +#include "atom_vec.h" +#include "comm.h" +#include "force.h" +#include "neighbor.h" +#include "neigh_list.h" +#include "integrate.h" +#include "math_const.h" +#include "memory.h" +#include "error.h" +#include "neigh_request.h" +#include "universe.h" +#include "update.h" +#include "domain.h" +#include "gpu_extra.h" + +using namespace LAMMPS_NS; +using namespace MathConst; + +#define EPSILON 1.0e-20 + +// External functions from cuda library for atom decomposition + +int borncwcs_gpu_init(const int ntypes, double **cutsq, double **host_rhoinv, + double **host_born1, double **host_born2, + double **host_born3, double **host_a, double **host_c, + double **host_d, double **sigma, double **offset, + double *special_lj, const int inum, + const int nall, const int max_nbors, const int maxspecial, + const double cell_size, int &gpu_mode, FILE *screen, + double **host_cut_ljsq, double host_cut_coulsq, + double *host_special_coul, const double qqrd2e, + const double alf, const double e_shift, const double f_shift); +void borncwcs_gpu_clear(); +int ** borncwcs_gpu_compute_n(const int ago, const int inum_full, const int nall, + double **host_x, int *host_type, double *sublo, + double *subhi, tagint *tag, int **nspecial, + tagint **special, const bool eflag, const bool vflag, + const bool eatom, const bool vatom, int &host_start, + int **ilist, int **jnum, const double cpu_time, + bool &success, double *host_q, double *boxlo, + double *prd); +void borncwcs_gpu_compute(const int ago, const int inum_full, const int nall, + double **host_x, int *host_type, int *ilist, int *numj, + int **firstneigh, const bool eflag, const bool vflag, + const bool eatom, const bool vatom, int &host_start, + const double cpu_time, bool &success, double *host_q, + const int nlocal, double *boxlo, double *prd); +double borncwcs_gpu_bytes(); + +/* ---------------------------------------------------------------------- */ + +PairBornCoulWolfCSGPU::PairBornCoulWolfCSGPU(LAMMPS *lmp) : PairBornCoulWolfCS(lmp), + gpu_mode(GPU_FORCE) +{ + respa_enable = 0; + reinitflag = 0; + cpu_time = 0.0; + GPU_EXTRA::gpu_ready(lmp->modify, lmp->error); +} + +/* ---------------------------------------------------------------------- + free all arrays +------------------------------------------------------------------------- */ + +PairBornCoulWolfCSGPU::~PairBornCoulWolfCSGPU() +{ + borncwcs_gpu_clear(); +} + +/* ---------------------------------------------------------------------- */ + +void PairBornCoulWolfCSGPU::compute(int eflag, int vflag) +{ + if (eflag || vflag) ev_setup(eflag,vflag); + else evflag = vflag_fdotr = 0; + + int nall = atom->nlocal + atom->nghost; + int inum, host_start; + + bool success = true; + int *ilist, *numneigh, **firstneigh; + if (gpu_mode != GPU_FORCE) { + inum = atom->nlocal; + firstneigh = borncwcs_gpu_compute_n(neighbor->ago, inum, nall, + atom->x, atom->type, domain->sublo, + domain->subhi, atom->tag, atom->nspecial, + atom->special, eflag, vflag, eflag_atom, + vflag_atom, host_start, + &ilist, &numneigh, cpu_time, success, + atom->q, domain->boxlo, domain->prd); + } else { + inum = list->inum; + ilist = list->ilist; + numneigh = list->numneigh; + firstneigh = list->firstneigh; + borncwcs_gpu_compute(neighbor->ago, inum, nall, atom->x, atom->type, + ilist, numneigh, firstneigh, eflag, vflag, eflag_atom, + vflag_atom, host_start, cpu_time, success, atom->q, + atom->nlocal, domain->boxlo, domain->prd); + } + if (!success) + error->one(FLERR,"Insufficient memory on accelerator"); + + if (host_startnewton_pair) + error->all(FLERR, + "Cannot use newton pair with born/coul/wolf/cs/gpu pair style"); + + // Repeat cutsq calculation because done after call to init_style + double maxcut = -1.0; + double cut; + for (int i = 1; i <= atom->ntypes; i++) { + for (int j = i; j <= atom->ntypes; j++) { + if (setflag[i][j] != 0 || (setflag[i][i] != 0 && setflag[j][j] != 0)) { + cut = init_one(i,j); + cut *= cut; + if (cut > maxcut) + maxcut = cut; + cutsq[i][j] = cutsq[j][i] = cut; + } else + cutsq[i][j] = cutsq[j][i] = 0.0; + } + } + double cell_size = sqrt(maxcut) + neighbor->skin; + + cut_coulsq = cut_coul * cut_coul; + + double e_shift = erfc(alf*cut_coul)/cut_coul; + double f_shift = -(e_shift+ 2.0*alf/MY_PIS * exp(-alf*alf*cut_coul*cut_coul)) / + cut_coul; + + int maxspecial=0; + if (atom->molecular) + maxspecial=atom->maxspecial; + int success = borncwcs_gpu_init(atom->ntypes+1, cutsq, rhoinv, + born1, born2, born3, a, c, d, sigma, offset, + force->special_lj, atom->nlocal, + atom->nlocal+atom->nghost, 300, maxspecial, + cell_size, gpu_mode, screen, cut_ljsq, + cut_coulsq, force->special_coul, force->qqrd2e, + alf, e_shift, f_shift); + GPU_EXTRA::check_flag(success,error,world); + + if (gpu_mode == GPU_FORCE) { + int irequest = neighbor->request(this,instance_me); + neighbor->requests[irequest]->half = 0; + neighbor->requests[irequest]->full = 1; + } +} + +/* ---------------------------------------------------------------------- */ + +double PairBornCoulWolfCSGPU::memory_usage() +{ + double bytes = Pair::memory_usage(); + return bytes + borncwcs_gpu_bytes(); +} + +/* ---------------------------------------------------------------------- */ + +void PairBornCoulWolfCSGPU::cpu_compute(int start, int inum, int eflag, int vflag, + int *ilist, int *numneigh, + int **firstneigh) { + int i,j,ii,jj,jnum,itype,jtype; + double xtmp,ytmp,ztmp,qtmp,delx,dely,delz,evdwl,ecoul,fpair; + double rsq,r2inv,r6inv,forcecoul,forceborn,factor_coul,factor_lj; + double erfcc,erfcd,v_sh,dvdrr,e_self,qisq; + double prefactor; + double r,rexp; + int *jlist; + + evdwl = ecoul = 0.0; + + double **x = atom->x; + double **f = atom->f; + double *q = atom->q; + int *type = atom->type; + int nlocal = atom->nlocal; + double *special_coul = force->special_coul; + double *special_lj = force->special_lj; + double qqrd2e = force->qqrd2e; + + double e_shift = erfc(alf*cut_coul)/cut_coul; + double f_shift = -(e_shift+ 2.0*alf/MY_PIS * exp(-alf*alf*cut_coul*cut_coul)) / + cut_coul; + + // loop over neighbors of my atoms + + for (ii = start; ii < inum; ii++) { + i = ilist[ii]; + qtmp = q[i]; + xtmp = x[i][0]; + ytmp = x[i][1]; + ztmp = x[i][2]; + itype = type[i]; + jlist = firstneigh[i]; + jnum = numneigh[i]; + + qisq = qtmp*qtmp; + e_self = -(e_shift/2.0 + alf/MY_PIS) * qisq*qqrd2e; + if (evflag) ev_tally(i,i,nlocal,0,0.0,e_self,0.0,0.0,0.0,0.0); + + for (jj = 0; jj < jnum; jj++) { + j = jlist[jj]; + factor_lj = special_lj[sbmask(j)]; + factor_coul = special_coul[sbmask(j)]; + j &= NEIGHMASK; + + delx = xtmp - x[j][0]; + dely = ytmp - x[j][1]; + delz = ztmp - x[j][2]; + rsq = delx*delx + dely*dely + delz*delz; + jtype = type[j]; + + if (rsq < cutsq[itype][jtype]) { + rsq += EPSILON; // Add EPISLON for case: r = 0; Interaction must be removed by special bond + r2inv = 1.0/rsq; + + if (rsq < cut_coulsq) { + r = sqrt(rsq); + prefactor = qqrd2e*qtmp*q[j]/r; + erfcc = erfc(alf*r); + erfcd = exp(-alf*alf*r*r); + v_sh = (erfcc - e_shift*r) * prefactor; + dvdrr = (erfcc/rsq + 2.0*alf/MY_PIS * erfcd/r) + f_shift; + forcecoul = dvdrr*rsq*prefactor; + if (factor_coul < 1.0) forcecoul -= (1.0-factor_coul)*prefactor; + } else forcecoul = 0.0; + + if (rsq < cut_ljsq[itype][jtype]) { + r6inv = r2inv*r2inv*r2inv; + r = sqrt(rsq); + rexp = exp((sigma[itype][jtype]-r)*rhoinv[itype][jtype]); + forceborn = born1[itype][jtype]*r*rexp - born2[itype][jtype]*r6inv + + born3[itype][jtype]*r2inv*r6inv; + } else forceborn = 0.0; + + fpair = (factor_coul*forcecoul + factor_lj*forceborn) * r2inv; + + f[i][0] += delx*fpair; + f[i][1] += dely*fpair; + f[i][2] += delz*fpair; + + if (eflag) { + if (rsq < cut_coulsq) { + ecoul = v_sh; + if (factor_coul < 1.0) ecoul -= (1.0-factor_coul)*prefactor; + } else ecoul = 0.0; + if (rsq < cut_ljsq[itype][jtype]) { + evdwl = a[itype][jtype]*rexp - c[itype][jtype]*r6inv + + d[itype][jtype]*r6inv*r2inv - offset[itype][jtype]; + evdwl *= factor_lj; + } else evdwl = 0.0; + } + + if (evflag) ev_tally_full(i,evdwl,ecoul,fpair,delx,dely,delz); + } + } + } +} diff --git a/src/GPU/pair_born_coul_wolf_cs_gpu.h b/src/GPU/pair_born_coul_wolf_cs_gpu.h new file mode 100644 index 0000000000..5f4f4ea35c --- /dev/null +++ b/src/GPU/pair_born_coul_wolf_cs_gpu.h @@ -0,0 +1,58 @@ +/* -*- c++ -*- ---------------------------------------------------------- + LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator + http://lammps.sandia.gov, Sandia National Laboratories + Steve Plimpton, sjplimp@sandia.gov + + Copyright (2003) Sandia Corporation. Under the terms of Contract + DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains + certain rights in this software. This software is distributed under + the GNU General Public License. + + See the README file in the top-level LAMMPS directory. +------------------------------------------------------------------------- */ + +#ifdef PAIR_CLASS + +PairStyle(born/coul/wolf/cs/gpu,PairBornCoulWolfCSGPU) + +#else + +#ifndef LMP_PAIR_BORN_COUL_WOLF_CS_GPU_H +#define LMP_PAIR_BORN_COUL_WOLF_CS_GPU_H + +#include "pair_born_coul_wolf_cs.h" + +namespace LAMMPS_NS { + +class PairBornCoulWolfCSGPU : public PairBornCoulWolfCS { + public: + PairBornCoulWolfCSGPU(LAMMPS *lmp); + ~PairBornCoulWolfCSGPU(); + void cpu_compute(int, int, int, int, int *, int *, int **); + void compute(int, int); + void init_style(); + double memory_usage(); + + enum { GPU_FORCE, GPU_NEIGH, GPU_HYB_NEIGH }; + + private: + int gpu_mode; + double cpu_time; +}; + +} +#endif +#endif + +/* ERROR/WARNING messages: + +E: Insufficient memory on accelerator + +There is insufficient memory on one of the devices specified for the gpu +package + +E: Cannot use newton pair with born/coul/wolf/cs/gpu pair style + +Self-explanatory. + +*/