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

Handle inconsistent J molecules in tip4p/gpu

Browse Source
This commit is contained in:
Vsevak
2022-05-28 00:24:24 +03:00
parent bd4bbbddbe
commit d8f8a3a36a
3 changed files with 38 additions and 13 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
lib/gpu/geryon/hip_device.h
View File

@ -394,7 +394,7 @@ UCL_Device::~UCL_Device() {
clear();
}
int UCL_Device::set_platform(const int) {
int UCL_Device::set_platform(const int pid) {
clear();
#ifdef UCL_DEBUG
assert(pid<num_platforms());

42
lib/gpu/lal_lj_tip4p_long.cu
View File

@ -217,26 +217,43 @@ __kernel void k_lj_tip4p_reneigh(const __global numtyp4 *restrict x_,
// set iH1,iH2 to closest image to O
iH1 = closest_image(i, iH1, sametag, x_);
iH2 = closest_image(i, iH2, sametag, x_);
//printf("%d %f %f %f\n", (iH1 < 0 || iH2 < 0) ? 0:1, ix.x, ix.y, ix.z );
/*
if (iH1 < 0) {
printf("i=%d\ttag[i]=%d\tmap[tag[i]-1]=%d\tmap[tag[i]-2]=%d\n", i, tag[i], map[tag[i]+1], map[tag[i]+2]);
}
if (iH2 < 0) {
printf("i=%d\ttag[i]=%d\tmap[tag[i]-1]=%d\tmap[tag[i]-2]=%d\n", i, tag[i], map[tag[i]+1], map[tag[i]+2]);
}*/
hneigh[i*4 ] = iH1;
hneigh[i*4+1] = iH2;
hneigh[i*4+2] = -1;
}
}
if (itype == typeH) {
if (i<inum && itype == typeH) {
if (hneigh[i*4+2] != -1) {
int iI, iH;
iI = atom_mapping(map,tag[i] - 1);
numtyp4 iIx; fetch4(iIx,iI,pos_tex); //x_[iI];
iO = closest_image(i,iI,sametag, x_);
//printf("%d %f %f %f\n", (iI < 0) ? 2:3, ix.x, ix.y, ix.z );
/*
// printf("iI = %d iO closest = %d\n",iI, iO);
if (iI < 0) {
printf("i=%d\ttag[i]=%d\tmap[tag[i]-1]=%d\tmap[tag[i]-2]=%d\n", i, tag[i], map[tag[i]-1],map[tag[i]-2]);
}*/
numtyp4 iIx; fetch4(iIx,iO,pos_tex); //x_[iI];
if ((int)iIx.w == typeH) {
iO = atom_mapping(map,tag[i] - 2);
iO = closest_image(i, iO, sametag, x_);
iH1 = closest_image(i, iI, sametag, x_);
iH2 = i;
//iH1 = closest_image(i, iI, sametag, x_);
//iH2 = i;
} else { //if ((int)iIx.w == typeO)
iH = atom_mapping(map, tag[i] + 1);
iO = closest_image(i,iI,sametag, x_);
iH1 = i;
iH2 = closest_image(i,iH,sametag, x_);
//iH = atom_mapping(map, tag[i] + 1);
//iO = closest_image(i,iI,sametag, x_);
//iH1 = i;
//iH2 = closest_image(i,iH,sametag, x_);
}
hneigh[i*4+0] = iO;
hneigh[i*4+1] += -1;
@ -268,7 +285,14 @@ __kernel void k_lj_tip4p_newsite(const __global numtyp4 *restrict x_,
iH2 = hneigh[i*4+1];
iO = i;
numtyp qO; fetch(qO,iO,q_tex);
compute_newsite(iO,iH1,iH2, &m[iO], qO, alpha, x_);
if (iH1>=0 && iH2>=0) {
compute_newsite(iO,iH1,iH2, &m[iO], qO, alpha, x_);
} else {
m[iO] = ix;
m[iO].w = qO;
hneigh[i*4] = iO;
hneigh[i*4+1] = iO;
}
}
}
}

7
src/GPU/pair_lj_cut_tip4p_long_gpu.cpp
View File

@ -186,8 +186,9 @@ void PairLJCutTIP4PLongGPU::init_style()
alpha = qdist / (cos(0.5 * theta) * blen);
cut_coulsq = cut_coul * cut_coul;
double cut_coulsqplus = (cut_coul + qdist + blen) * (cut_coul + qdist + blen);
if (maxcut < cut_coulsqplus) { cell_size = (cut_coul + qdist + blen) + neighbor->skin; }
double cut_coulplus = cut_coul + qdist + blen;
double cut_coulsqplus = cut_coulplus*cut_coulplus;
if (sqrt(maxcut) < cut_coulplus+blen) { cell_size = (cut_coulplus + blen) + neighbor->skin; }
if (comm->cutghostuser < cell_size) {
if (comm->me == 0)
error->warning(FLERR,
@ -205,7 +206,7 @@ void PairLJCutTIP4PLongGPU::init_style()
GPU_EXTRA::check_flag(success, error, world);
if (gpu_mode == GPU_FORCE) {
auto req = neighbor->add_request(this, NeighConst::REQ_FULL);
req->set_cutoff(cut_coul + qdist + blen + neighbor->skin);
req->set_cutoff(cut_coulplus + neighbor->skin);
}
}