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

major refactor for restart, data file handling. removal of dead code.

Browse Source
This commit is contained in:
Axel Kohlmeyer
2024-06-19 11:08:31 -04:00
parent 6ada6b7bf2
commit 45508baee5
7 changed files with 90 additions and 118 deletions
Download Patch File Download Diff File Expand all files Collapse all files

5
lib/gpu/lal_dpd_coul_slater_long.cpp
View File

@ -47,8 +47,7 @@ int DPDCoulSlaterLongT::init(const int ntypes,
double **host_cutsq, double **host_a0,
double **host_gamma, double **host_sigma,
double **host_cut_dpd, double **host_cut_dpdsq,
double **host_cut_slatersq, double **host_scale,
double *host_special_lj,
double **host_cut_slatersq, double *host_special_lj,
const bool tstat_only,
const int nlocal, const int nall,
const int max_nbors, const int maxspecial,
@ -115,7 +114,7 @@ int DPDCoulSlaterLongT::init(const int ntypes,
cutsq.alloc(lj_types*lj_types,*(this->ucl_device),UCL_READ_ONLY);
this->atom->type_pack4(ntypes,lj_types,cutsq,host_write,host_cutsq,
host_cut_dpdsq, host_scale, host_cut_slatersq);
host_cut_dpdsq,host_cut_slatersq);
double special_sqrt[4];
special_sqrt[0] = sqrt(host_special_lj[0]);

10
lib/gpu/lal_dpd_coul_slater_long.cu
View File

@ -289,8 +289,8 @@ __kernel void k_dpd_coul_slater_long(const __global numtyp4 *restrict x_,
// apply Slater electrostatic force if distance below Slater cutoff
// and the two species have a slater coeff
// cutsq[mtype].w -> Coulombic squared cutoff
if ( cutsq[mtype].w != 0.0 && rsq < cutsq[mtype].w){
// cutsq[mtype].z -> Coulombic squared cutoff
if ( cutsq[mtype].z != 0.0 && rsq < cutsq[mtype].z){
numtyp r2inv=ucl_recip(rsq);
numtyp _erfc;
numtyp grij = g_ewald * r;
@ -426,7 +426,7 @@ __kernel void k_dpd_coul_slater_long_fast(const __global numtyp4 *restrict x_,
int mtype=itype+jx.w;
/// cutsq.x = cutsq, cutsq.y = cut_dpdsq, cutsq.z = scale, cutsq.w = cut_slatersq
/// cutsq.x = cutsq, cutsq.y = cut_dpdsq, cutsq.z = cut_slatersq
if (rsq<cutsq[mtype].x) {
numtyp r=ucl_sqrt(rsq);
numtyp force_dpd = (numtyp)0.0;
@ -474,8 +474,8 @@ __kernel void k_dpd_coul_slater_long_fast(const __global numtyp4 *restrict x_,
// apply Slater electrostatic force if distance below Slater cutoff
// and the two species have a slater coeff
// cutsq[mtype].w -> Coulombic squared cutoff
if ( cutsq[mtype].w != 0.0 && rsq < cutsq[mtype].w){
// cutsq[mtype].z -> Coulombic squared cutoff
if ( cutsq[mtype].z != 0.0 && rsq < cutsq[mtype].z){
numtyp r2inv=ucl_recip(rsq);
numtyp _erfc;
numtyp grij = g_ewald * r;

15
lib/gpu/lal_dpd_coul_slater_long.h
View File

@ -37,13 +37,12 @@ class DPDCoulSlaterLong : public BaseDPD<numtyp, acctyp> {
* - -3 if there is an out of memory error
* - -4 if the GPU library was not compiled for GPU
* - -5 Double precision is not supported on card **/
int init(const int ntypes, double **host_cutsq, double **host_a0,
double **host_gamma, double **host_sigma,
double **host_cut_dpd, double **host_cut_dpdsq, double **host_cut_slatersq,
double **host_scale, double *host_special_lj, bool tstat_only, const int nlocal,
const int nall, const int max_nbors, const int maxspecial,
const double cell_size, const double gpu_split, FILE *screen, double *host_special_coul,
const double qqrd2e, const double g_ewald, const double lamda);
int init(const int ntypes, double **host_cutsq, double **host_a0, double **host_gamma,
double **host_sigma, double **host_cut_dpd, double **host_cut_dpdsq,
double **host_cut_slatersq, double *host_special_lj, bool tstat_only, const int nlocal,
const int nall, const int max_nbors, const int maxspecial, const double cell_size,
const double gpu_split, FILE *screen, double *host_special_coul, const double qqrd2e,
const double g_ewald, const double lamda);
/// Clear all host and device data
/** \note This is called at the beginning of the init() routine **/
@ -66,7 +65,7 @@ class DPDCoulSlaterLong : public BaseDPD<numtyp, acctyp> {
/// coeff.x = a0, coeff.y = gamma, coeff.z = sigma, coeff.w = cut_dpd
UCL_D_Vec<numtyp4> coeff;
/// cutsq.x = cutsq, cutsq.y = cut_dpdsq, cutsq.z = scale, cutsq.w = cut_slatersq
/// cutsq.x = cutsq, cutsq.y = cut_dpdsq, cutsq.w = cut_slatersq
UCL_D_Vec<numtyp4> cutsq;
/// Special LJ values

32
lib/gpu/lal_dpd_coul_slater_long_ext.cpp
View File

@ -27,12 +27,14 @@ static DPDCoulSlaterLong<PRECISION,ACC_PRECISION> DPDCMF;
// ---------------------------------------------------------------------------
// Allocate memory on host and device and copy constants to device
// ---------------------------------------------------------------------------
int dpd_coul_slater_long_gpu_init(const int ntypes, double **host_cutsq, double **host_a0, double **host_gamma,
double **host_sigma, double **host_cut_dpd, double **host_cut_dpdsq, double **host_cut_slatersq,
double **host_scale, 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_special_coul,
const double qqrd2e, const double g_ewald, const double lamda) {
int dpd_coul_slater_long_gpu_init(const int ntypes, double **host_cutsq, double **host_a0,
double **host_gamma, double **host_sigma, double **host_cut_dpd,
double **host_cut_dpdsq, double **host_cut_slatersq,
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_special_coul, const double qqrd2e,
const double g_ewald, const double lamda) {
DPDCMF.clear();
gpu_mode=DPDCMF.device->gpu_mode();
double gpu_split=DPDCMF.device->particle_split();
@ -55,11 +57,10 @@ int dpd_coul_slater_long_gpu_init(const int ntypes, double **host_cutsq, double
int init_ok=0;
if (world_me==0)
init_ok=DPDCMF.init(ntypes, host_cutsq, host_a0, host_gamma, host_sigma,
host_cut_dpd, host_cut_dpdsq, host_cut_slatersq,
host_scale, special_lj, false, inum, nall, max_nbors,
maxspecial, cell_size, gpu_split, screen,
host_special_coul,qqrd2e, g_ewald, lamda);
init_ok=DPDCMF.init(ntypes, host_cutsq, host_a0, host_gamma, host_sigma, host_cut_dpd,
host_cut_dpdsq, host_cut_slatersq, special_lj, false, inum, nall,
max_nbors, maxspecial, cell_size, gpu_split, screen, host_special_coul,
qqrd2e, g_ewald, lamda);
DPDCMF.device->world_barrier();
if (message)
@ -75,11 +76,10 @@ int dpd_coul_slater_long_gpu_init(const int ntypes, double **host_cutsq, double
fflush(screen);
}
if (gpu_rank==i && world_me!=0)
init_ok=DPDCMF.init(ntypes, host_cutsq, host_a0, host_gamma, host_sigma,
host_cut_dpd, host_cut_dpdsq, host_cut_slatersq,
host_scale, special_lj, false, inum, nall, max_nbors,
maxspecial, cell_size, gpu_split, screen,
host_special_coul,qqrd2e, g_ewald, lamda);
init_ok=DPDCMF.init(ntypes, host_cutsq, host_a0, host_gamma, host_sigma, host_cut_dpd,
host_cut_dpdsq, host_cut_slatersq, special_lj, false, inum, nall,
max_nbors, maxspecial, cell_size, gpu_split, screen, host_special_coul,
qqrd2e, g_ewald, lamda);
DPDCMF.device->serialize_init();
if (message)