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

Added the API for the umutual kernel, needs work for storing the tdiptdip array

Browse Source
This commit is contained in:
Trung Nguyen
2021-09-09 17:22:09 -05:00
parent b654f293ee
commit a22923aee2
3 changed files with 41 additions and 23 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
lib/gpu/lal_amoeba.cu
View File

@ -1012,7 +1012,7 @@ __kernel void k_amoeba_umutual2b(const __global numtyp4 *restrict x_,
numtyp scalek = factor_uscale;
bcn[0] = bn[1] - ((numtyp)1.0-scalek*scale3)*rr3;
bcn[1] = bn[2] - ((numtyp)1.0-scalek*scale5)*rr5;
numtyp tdipdip[6];
numtyp tdipdip[6]; // the following tdipdip is incorrect!! needs work to store tdipdip
tdipdip[0] = -bcn[0] + bcn[1]*xr*xr;
tdipdip[1] = bcn[1]*xr*yr;
tdipdip[2] = bcn[1]*xr*zr;

54
lib/gpu/lal_amoeba_ext.cpp
View File

@ -105,6 +105,42 @@ void amoeba_gpu_clear() {
AMOEBAMF.clear();
}
int** amoeba_gpu_compute_udirect2b(const int ago, const int inum_full,
const int nall, double **host_x, int *host_type,
int *host_amtype, int *host_amgroup, double **host_rpole,
double **host_uind, double **host_uinp,
double *sublo, double *subhi, tagint *tag, int **nspecial,
tagint **special, int *nspecial15, tagint** special15,
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 **fieldp_ptr) {
return AMOEBAMF.compute_udirect2b(ago, inum_full, nall, host_x, host_type,
host_amtype, host_amgroup, host_rpole, host_uind, host_uinp,
sublo, subhi, tag, nspecial, special, nspecial15, special15,
eflag, vflag, eatom, vatom, host_start, ilist, jnum,
cpu_time, success, host_q, boxlo, prd, fieldp_ptr);
}
int** amoeba_gpu_compute_umutual2b(const int ago, const int inum_full,
const int nall, double **host_x, int *host_type,
int *host_amtype, int *host_amgroup, double **host_rpole,
double **host_uind, double **host_uinp,
double *sublo, double *subhi, tagint *tag, int **nspecial,
tagint **special, int *nspecial15, tagint** special15,
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 **fieldp_ptr) {
return AMOEBAMF.compute_umutual2b(ago, inum_full, nall, host_x, host_type,
host_amtype, host_amgroup, host_rpole, host_uind, host_uinp,
sublo, subhi, tag, nspecial, special, nspecial15, special15,
eflag, vflag, eatom, vatom, host_start, ilist, jnum,
cpu_time, success, host_q, boxlo, prd, fieldp_ptr);
}
int** amoeba_gpu_compute_polar_real(const int ago, const int inum_full,
const int nall, double **host_x, int *host_type,
int *host_amtype, int *host_amgroup,
@ -124,24 +160,6 @@ int** amoeba_gpu_compute_polar_real(const int ago, const int inum_full,
host_q, boxlo, prd, tep_ptr);
}
int** amoeba_gpu_compute_udirect2b(const int ago, const int inum_full,
const int nall, double **host_x, int *host_type,
int *host_amtype, int *host_amgroup, double **host_rpole,
double **host_uind, double **host_uinp,
double *sublo, double *subhi, tagint *tag, int **nspecial,
tagint **special, int *nspecial15, tagint** special15,
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 **fieldp_ptr) {
return AMOEBAMF.compute_udirect2b(ago, inum_full, nall, host_x, host_type,
host_amtype, host_amgroup, host_rpole, host_uind, host_uinp,
sublo, subhi, tag, nspecial, special, nspecial15, special15,
eflag, vflag, eatom, vatom, host_start, ilist, jnum,
cpu_time, success, host_q, boxlo, prd, fieldp_ptr);
}
double amoeba_gpu_bytes() {
return AMOEBAMF.host_memory_usage();
}

8
src/GPU/pair_amoeba_gpu.cpp
View File

@ -74,7 +74,7 @@ int ** amoeba_gpu_compute_udirect2b(const int ago, const int inum, const int nal
int **ilist, int **jnum, const double cpu_time,
bool &success, double *host_q, double *boxlo, double *prd,
void **fieldp_ptr);
/*
int ** amoeba_gpu_compute_umutual2b(const int ago, const int inum, const int nall,
double **host_x, int *host_type, int *host_amtype, int *host_amgroup,
double **host_rpole, double **host_uind, double **host_uinp,
@ -85,7 +85,7 @@ int ** amoeba_gpu_compute_umutual2b(const int ago, const int inum, const int nal
int **ilist, int **jnum, const double cpu_time,
bool &success, double *host_q, double *boxlo, double *prd,
void **fieldp_ptr);
*/
int ** amoeba_gpu_compute_polar_real(const int ago, const int inum, const int nall,
double **host_x, int *host_type, int *host_amtype, int *host_amgroup,
double **host_rpole, double **host_uind, double **host_uinp,
@ -1015,7 +1015,7 @@ void PairAmoebaGPU::umutual2b(double **field, double **fieldp)
PairAmoeba::umutual2b(field, fieldp);
return;
}
/*
int eflag=1, vflag=1;
int nall = atom->nlocal + atom->nghost;
int inum, host_start;
@ -1068,7 +1068,7 @@ void PairAmoebaGPU::umutual2b(double **field, double **fieldp)
fieldp[i][1] += fieldp_ptr[idx+1];
fieldp[i][2] += fieldp_ptr[idx+2];
}
*/
}
/* ---------------------------------------------------------------------- */