forked from lijiext/lammps
1438 lines
43 KiB
Plaintext
1438 lines
43 KiB
Plaintext
/* ----------------------------------------------------------------------
|
|
LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
|
|
|
|
Original Version:
|
|
http://lammps.sandia.gov, Sandia National Laboratories
|
|
Steve Plimpton, sjplimp@sandia.gov
|
|
|
|
See the README file in the top-level LAMMPS directory.
|
|
|
|
-----------------------------------------------------------------------
|
|
|
|
USER-CUDA Package and associated modifications:
|
|
https://sourceforge.net/projects/lammpscuda/
|
|
|
|
Christian Trott, christian.trott@tu-ilmenau.de
|
|
Lars Winterfeld, lars.winterfeld@tu-ilmenau.de
|
|
Theoretical Physics II, University of Technology Ilmenau, Germany
|
|
|
|
See the README file in the USER-CUDA directory.
|
|
|
|
This software is distributed under the GNU General Public License.
|
|
------------------------------------------------------------------------- */
|
|
#define EWALD_F 1.12837917
|
|
#define EWALD_P 0.3275911
|
|
#define A1 0.254829592
|
|
#define A2 -0.284496736
|
|
#define A3 1.421413741
|
|
#define A4 -1.453152027
|
|
#define A5 1.061405429
|
|
|
|
|
|
template <const PAIR_FORCES pair_type, const COUL_FORCES coul_type, const unsigned int extended_data>
|
|
__global__ void Pair_Kernel_TpA(int eflag, int vflag, int eflag_atom, int vflag_atom)
|
|
{
|
|
ENERGY_FLOAT evdwl = ENERGY_F(0.0);
|
|
ENERGY_FLOAT ecoul = ENERGY_F(0.0);
|
|
|
|
ENERGY_FLOAT* sharedE;
|
|
ENERGY_FLOAT* sharedECoul;
|
|
ENERGY_FLOAT* sharedV = &sharedmem[threadIdx.x];
|
|
|
|
if(eflag || eflag_atom) {
|
|
sharedE = &sharedmem[threadIdx.x];
|
|
sharedE[0] = ENERGY_F(0.0);
|
|
sharedV += blockDim.x;
|
|
|
|
if(coul_type != COUL_NONE) {
|
|
sharedECoul = sharedE + blockDim.x;
|
|
sharedECoul[0] = ENERGY_F(0.0);
|
|
sharedV += blockDim.x;
|
|
}
|
|
}
|
|
|
|
if(vflag || vflag_atom) {
|
|
sharedV[0 * blockDim.x] = ENERGY_F(0.0);
|
|
sharedV[1 * blockDim.x] = ENERGY_F(0.0);
|
|
sharedV[2 * blockDim.x] = ENERGY_F(0.0);
|
|
sharedV[3 * blockDim.x] = ENERGY_F(0.0);
|
|
sharedV[4 * blockDim.x] = ENERGY_F(0.0);
|
|
sharedV[5 * blockDim.x] = ENERGY_F(0.0);
|
|
}
|
|
|
|
int ii = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
|
|
|
|
X_FLOAT xtmp, ytmp, ztmp;
|
|
X_FLOAT4 myxtype;
|
|
F_FLOAT fxtmp, fytmp, fztmp, fpair;
|
|
F_FLOAT delx, dely, delz;
|
|
F_FLOAT factor_lj, factor_coul;
|
|
F_FLOAT qtmp;
|
|
int itype, i, j;
|
|
int jnum = 0;
|
|
int* jlist;
|
|
|
|
if(ii < _inum) {
|
|
i = _ilist[ii];
|
|
|
|
myxtype = fetchXType(i);
|
|
xtmp = myxtype.x;
|
|
ytmp = myxtype.y;
|
|
ztmp = myxtype.z;
|
|
itype = static_cast <int>(myxtype.w);
|
|
|
|
|
|
fxtmp = F_F(0.0);
|
|
fytmp = F_F(0.0);
|
|
fztmp = F_F(0.0);
|
|
|
|
if(coul_type != COUL_NONE)
|
|
qtmp = fetchQ(i);
|
|
|
|
jnum = _numneigh[i];
|
|
jlist = &_neighbors[i];
|
|
}
|
|
|
|
__syncthreads();
|
|
|
|
for(int jj = 0; jj < jnum; jj++) {
|
|
if(ii < _inum)
|
|
if(jj < jnum) {
|
|
fpair = F_F(0.0);
|
|
j = jlist[jj * _nlocal];
|
|
factor_lj = _special_lj[sbmask(j)];
|
|
|
|
if(coul_type != COUL_NONE)
|
|
factor_coul = _special_coul[sbmask(j)];
|
|
|
|
j &= NEIGHMASK;
|
|
|
|
myxtype = fetchXType(j);
|
|
delx = xtmp - myxtype.x;
|
|
dely = ytmp - myxtype.y;
|
|
delz = ztmp - myxtype.z;
|
|
int jtype = static_cast <int>(myxtype.w);
|
|
|
|
|
|
const F_FLOAT rsq = delx * delx + dely * dely + delz * delz;
|
|
|
|
bool in_cutoff = rsq < (_cutsq_global > X_F(0.0) ? _cutsq_global : _cutsq[itype * _cuda_ntypes + jtype]);
|
|
|
|
if(in_cutoff) {
|
|
switch(pair_type) {
|
|
case PAIR_BORN:
|
|
fpair += PairBornCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_BUCK:
|
|
fpair += PairBuckCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_CG_CMM:
|
|
fpair += PairLJSDKCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_CHARMM:
|
|
fpair += PairLJCharmmCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_CLASS2:
|
|
fpair += PairLJClass2Cuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_CUT:
|
|
fpair += PairLJCutCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_EXPAND:
|
|
fpair += PairLJExpandCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_GROMACS:
|
|
fpair += PairLJGromacsCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_SMOOTH:
|
|
fpair += PairLJSmoothCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ96_CUT:
|
|
fpair += PairLJ96CutCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_MORSE_R6:
|
|
fpair += PairMorseR6Cuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_MORSE:
|
|
fpair += PairMorseCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
}
|
|
}
|
|
|
|
if(coul_type != COUL_NONE) {
|
|
const F_FLOAT qiqj = qtmp * fetchQ(j);
|
|
|
|
if(qiqj * qiqj > 1e-8) {
|
|
const bool in_coul_cutoff =
|
|
rsq < (_cut_coulsq_global > X_F(0.0) ? _cut_coulsq_global : _cut_coulsq[itype * _cuda_ntypes + jtype]);
|
|
|
|
if(in_coul_cutoff) {
|
|
switch(coul_type) {
|
|
case COUL_CHARMM:
|
|
fpair += CoulCharmmCuda_Eval(rsq, factor_coul, eflag, ecoul, qiqj);
|
|
break;
|
|
|
|
case COUL_CHARMM_IMPLICIT:
|
|
fpair += CoulCharmmImplicitCuda_Eval(rsq, factor_coul, eflag, ecoul, qiqj);
|
|
break;
|
|
|
|
case COUL_CUT: {
|
|
const F_FLOAT forcecoul = factor_coul * _qqrd2e * qiqj * _RSQRT_(rsq);
|
|
|
|
if(eflag) {
|
|
ecoul += forcecoul;
|
|
}
|
|
|
|
fpair += forcecoul * (F_F(1.0) / rsq);
|
|
}
|
|
break;
|
|
|
|
case COUL_DEBYE: {
|
|
const F_FLOAT r2inv = F_F(1.0) / rsq;
|
|
const X_FLOAT r = _RSQRT_(r2inv);
|
|
const X_FLOAT rinv = F_F(1.0) / r;
|
|
const F_FLOAT screening = _EXP_(-_kappa * r);
|
|
F_FLOAT forcecoul = factor_coul * _qqrd2e * qiqj * screening ;
|
|
|
|
if(eflag) {
|
|
ecoul += forcecoul * rinv;
|
|
}
|
|
|
|
forcecoul *= (_kappa + rinv);
|
|
fpair += forcecoul * r2inv;
|
|
}
|
|
break;
|
|
|
|
case COUL_GROMACS:
|
|
fpair += CoulGromacsCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_coul, eflag, ecoul, qiqj);
|
|
break;
|
|
|
|
case COUL_LONG: {
|
|
const F_FLOAT r2inv = F_F(1.0) / rsq;
|
|
const F_FLOAT r = _RSQRT_(r2inv);
|
|
const F_FLOAT grij = _g_ewald * r;
|
|
const F_FLOAT expm2 = _EXP_(-grij * grij);
|
|
const F_FLOAT t = F_F(1.0) / (F_F(1.0) + EWALD_P * grij);
|
|
const F_FLOAT erfc = t * (A1 + t * (A2 + t * (A3 + t * (A4 + t * A5)))) * expm2;
|
|
const F_FLOAT prefactor = _qqrd2e * qiqj * (F_F(1.0) / r);
|
|
F_FLOAT forcecoul = prefactor * (erfc + EWALD_F * grij * expm2);
|
|
|
|
if(factor_coul < 1.0) forcecoul -= (1.0 - factor_coul) * prefactor;
|
|
|
|
if(eflag) {
|
|
ecoul += prefactor * erfc;
|
|
|
|
if(factor_coul < 1.0) ecoul -= (1.0 - factor_coul) * prefactor;
|
|
}
|
|
|
|
fpair += forcecoul * r2inv;
|
|
}
|
|
break;
|
|
}
|
|
}
|
|
|
|
in_cutoff = in_cutoff || in_coul_cutoff;
|
|
}
|
|
}
|
|
|
|
|
|
if(in_cutoff) {
|
|
F_FLOAT dxfp, dyfp, dzfp;
|
|
fxtmp += dxfp = delx * fpair;
|
|
fytmp += dyfp = dely * fpair;
|
|
fztmp += dzfp = delz * fpair;
|
|
|
|
if(vflag) {
|
|
sharedV[0 * blockDim.x] += delx * dxfp;
|
|
sharedV[1 * blockDim.x] += dely * dyfp;
|
|
sharedV[2 * blockDim.x] += delz * dzfp;
|
|
sharedV[3 * blockDim.x] += delx * dyfp;
|
|
sharedV[4 * blockDim.x] += delx * dzfp;
|
|
sharedV[5 * blockDim.x] += dely * dzfp;
|
|
}
|
|
}
|
|
}
|
|
}
|
|
|
|
__syncthreads();
|
|
|
|
if(ii < _inum) {
|
|
F_FLOAT* my_f;
|
|
|
|
if(_collect_forces_later) {
|
|
ENERGY_FLOAT* buffer = (ENERGY_FLOAT*) _buffer;
|
|
|
|
if(eflag) {
|
|
buffer = &buffer[1 * gridDim.x * gridDim.y];
|
|
|
|
if(coul_type != COUL_NONE)
|
|
buffer = &buffer[1 * gridDim.x * gridDim.y];
|
|
}
|
|
|
|
if(vflag) {
|
|
buffer = &buffer[6 * gridDim.x * gridDim.y];
|
|
}
|
|
|
|
my_f = (F_FLOAT*) buffer;
|
|
my_f += i;
|
|
*my_f = fxtmp;
|
|
my_f += _nmax;
|
|
*my_f = fytmp;
|
|
my_f += _nmax;
|
|
*my_f = fztmp;
|
|
} else {
|
|
my_f = _f + i;
|
|
*my_f += fxtmp;
|
|
my_f += _nmax;
|
|
*my_f += fytmp;
|
|
my_f += _nmax;
|
|
*my_f += fztmp;
|
|
}
|
|
}
|
|
|
|
__syncthreads();
|
|
|
|
if(eflag) {
|
|
sharedE[0] = evdwl;
|
|
|
|
if(coul_type != COUL_NONE)
|
|
sharedECoul[0] = ecoul;
|
|
}
|
|
|
|
if(eflag_atom && i < _nlocal) {
|
|
if(coul_type != COUL_NONE)
|
|
_eatom[i] += evdwl + ecoul;
|
|
else
|
|
_eatom[i] += evdwl;
|
|
}
|
|
|
|
if(vflag_atom && i < _nlocal) {
|
|
_vatom[i] += ENERGY_F(0.5) * sharedV[0 * blockDim.x];
|
|
_vatom[i + _nmax] += ENERGY_F(0.5) * sharedV[1 * blockDim.x];
|
|
_vatom[i + 2 * _nmax] += ENERGY_F(0.5) * sharedV[2 * blockDim.x];
|
|
_vatom[i + 3 * _nmax] += ENERGY_F(0.5) * sharedV[3 * blockDim.x];
|
|
_vatom[i + 4 * _nmax] += ENERGY_F(0.5) * sharedV[4 * blockDim.x];
|
|
_vatom[i + 5 * _nmax] += ENERGY_F(0.5) * sharedV[5 * blockDim.x];
|
|
}
|
|
|
|
if(vflag || eflag) PairVirialCompute_A_Kernel(eflag, vflag, coul_type != COUL_NONE ? 1 : 0);
|
|
}
|
|
|
|
template <const PAIR_FORCES pair_type, const COUL_FORCES coul_type, const unsigned int extended_data>
|
|
__global__ void Pair_Kernel_BpA(int eflag, int vflag, int eflag_atom, int vflag_atom)
|
|
{
|
|
int ii = (blockIdx.x * gridDim.y + blockIdx.y);
|
|
|
|
if(ii >= _inum)
|
|
return;
|
|
|
|
ENERGY_FLOAT evdwl = ENERGY_F(0.0);
|
|
ENERGY_FLOAT ecoul = ENERGY_F(0.0);
|
|
F_FLOAT3* sharedVirial1;
|
|
F_FLOAT3* sharedVirial2;
|
|
F_FLOAT* sharedEnergy;
|
|
F_FLOAT* sharedEnergyCoul;
|
|
|
|
F_FLOAT3* sharedForce = (F_FLOAT3*) &sharedmem[0];
|
|
|
|
if(vflag) {
|
|
sharedVirial1 = &sharedForce[64];
|
|
sharedVirial2 = &sharedVirial1[64];
|
|
} else {
|
|
sharedVirial1 = &sharedForce[0];
|
|
sharedVirial2 = &sharedVirial1[0];
|
|
}
|
|
|
|
if(eflag) {
|
|
if(vflag || vflag_atom)
|
|
sharedEnergy = (F_FLOAT*) &sharedVirial2[64];
|
|
else
|
|
sharedEnergy = (F_FLOAT*) &sharedForce[64];
|
|
|
|
if(coul_type != COUL_NONE)
|
|
sharedEnergyCoul = (F_FLOAT*) &sharedEnergy[64];
|
|
|
|
}
|
|
|
|
F_FLOAT3 partialForce = { F_F(0.0), F_F(0.0), F_F(0.0) };
|
|
F_FLOAT3 partialVirial1 = { F_F(0.0), F_F(0.0), F_F(0.0) };
|
|
F_FLOAT3 partialVirial2 = { F_F(0.0), F_F(0.0), F_F(0.0) };
|
|
|
|
X_FLOAT xtmp, ytmp, ztmp;
|
|
X_FLOAT4 myxtype;
|
|
F_FLOAT delx, dely, delz;
|
|
F_FLOAT factor_lj, factor_coul;
|
|
F_FLOAT fpair;
|
|
F_FLOAT qtmp;
|
|
int itype, jnum, i, j;
|
|
int* jlist;
|
|
|
|
i = _ilist[ii];
|
|
|
|
myxtype = fetchXType(i);
|
|
|
|
xtmp = myxtype.x;
|
|
ytmp = myxtype.y;
|
|
ztmp = myxtype.z;
|
|
itype = static_cast <int>(myxtype.w);
|
|
|
|
if(coul_type != COUL_NONE)
|
|
qtmp = fetchQ(i);
|
|
|
|
jnum = _numneigh[i];
|
|
|
|
jlist = &_neighbors[i * _maxneighbors];
|
|
__syncthreads();
|
|
|
|
for(int jj = threadIdx.x; jj < jnum + blockDim.x; jj += blockDim.x) {
|
|
if(jj < jnum) {
|
|
fpair = F_F(0.0);
|
|
j = jlist[jj];
|
|
factor_lj = _special_lj[sbmask(j)];
|
|
|
|
if(coul_type != COUL_NONE)
|
|
factor_coul = _special_coul[sbmask(j)];
|
|
|
|
j &= NEIGHMASK;
|
|
|
|
myxtype = fetchXType(j);
|
|
|
|
delx = xtmp - myxtype.x;
|
|
dely = ytmp - myxtype.y;
|
|
delz = ztmp - myxtype.z;
|
|
int jtype = static_cast <int>(myxtype.w);
|
|
|
|
const F_FLOAT rsq = delx * delx + dely * dely + delz * delz;
|
|
|
|
bool in_cutoff = rsq < (_cutsq_global > X_F(0.0) ? _cutsq_global : _cutsq[itype * _cuda_ntypes + jtype]);
|
|
bool in_coul_cutoff;
|
|
|
|
if(in_cutoff) {
|
|
switch(pair_type) {
|
|
case PAIR_BORN:
|
|
fpair += PairBornCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_BUCK:
|
|
fpair += PairBuckCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_CG_CMM:
|
|
fpair += PairLJSDKCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_CHARMM:
|
|
fpair += PairLJCharmmCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_CLASS2:
|
|
fpair += PairLJClass2Cuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_CUT:
|
|
fpair += PairLJCutCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_EXPAND:
|
|
fpair += PairLJExpandCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_GROMACS:
|
|
fpair += PairLJGromacsCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_SMOOTH:
|
|
fpair += PairLJSmoothCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ96_CUT:
|
|
fpair += PairLJ96CutCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_MORSE_R6:
|
|
fpair += PairMorseR6Cuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_MORSE:
|
|
fpair += PairMorseCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
}
|
|
}
|
|
|
|
if(coul_type != COUL_NONE) {
|
|
const F_FLOAT qiqj = qtmp * fetchQ(j);
|
|
|
|
if(qiqj * qiqj > (1e-8f)) {
|
|
in_coul_cutoff =
|
|
rsq < (_cut_coulsq_global > X_F(0.0) ? _cut_coulsq_global : _cut_coulsq[itype * _cuda_ntypes + jtype]);
|
|
|
|
if(in_coul_cutoff) {
|
|
switch(coul_type) {
|
|
case COUL_CHARMM:
|
|
fpair += CoulCharmmCuda_Eval(rsq, factor_coul, eflag, ecoul, qiqj);
|
|
break;
|
|
|
|
case COUL_CHARMM_IMPLICIT:
|
|
fpair += CoulCharmmImplicitCuda_Eval(rsq, factor_coul, eflag, ecoul, qiqj);
|
|
break;
|
|
|
|
case COUL_GROMACS:
|
|
fpair += CoulGromacsCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_coul, eflag, ecoul, qiqj);
|
|
break;
|
|
|
|
case COUL_LONG: {
|
|
const F_FLOAT r2inv = F_F(1.0) / rsq;
|
|
const F_FLOAT r = _RSQRT_(r2inv);
|
|
const F_FLOAT grij = _g_ewald * r;
|
|
const F_FLOAT expm2 = _EXP_(-grij * grij);
|
|
const F_FLOAT t = F_F(1.0) / (F_F(1.0) + EWALD_P * grij);
|
|
const F_FLOAT erfc = t * (A1 + t * (A2 + t * (A3 + t * (A4 + t * A5)))) * expm2;
|
|
const F_FLOAT prefactor = _qqrd2e * qiqj * (F_F(1.0) / r);
|
|
F_FLOAT forcecoul = prefactor * (erfc + EWALD_F * grij * expm2);
|
|
|
|
if(factor_coul < 1.0) forcecoul -= (1.0 - factor_coul) * prefactor;
|
|
|
|
if(eflag) {
|
|
ecoul += prefactor * erfc;
|
|
|
|
if(factor_coul < 1.0) ecoul -= (1.0 - factor_coul) * prefactor;
|
|
}
|
|
|
|
fpair += forcecoul * r2inv;
|
|
}
|
|
break;
|
|
|
|
case COUL_DEBYE: {
|
|
const F_FLOAT r2inv = F_F(1.0) / rsq;
|
|
const X_FLOAT r = _RSQRT_(r2inv);
|
|
const X_FLOAT rinv = F_F(1.0) / r;
|
|
const F_FLOAT screening = _EXP_(-_kappa * r);
|
|
F_FLOAT forcecoul = factor_coul * _qqrd2e * qiqj * screening ;
|
|
|
|
if(eflag) {
|
|
ecoul += forcecoul * rinv;
|
|
}
|
|
|
|
forcecoul *= (_kappa + rinv);
|
|
fpair += forcecoul * r2inv;
|
|
}
|
|
break;
|
|
|
|
case COUL_CUT: {
|
|
const F_FLOAT forcecoul = factor_coul * _qqrd2e * qiqj * _RSQRT_(rsq);
|
|
|
|
if(eflag) {
|
|
ecoul += forcecoul;
|
|
}
|
|
|
|
fpair += forcecoul * (F_F(1.0) / rsq);
|
|
}
|
|
break;
|
|
|
|
|
|
}
|
|
}
|
|
}
|
|
}
|
|
|
|
|
|
|
|
if(in_cutoff || in_coul_cutoff) {
|
|
F_FLOAT dxfp, dyfp, dzfp;
|
|
partialForce.x += dxfp = delx * fpair;
|
|
partialForce.y += dyfp = dely * fpair;
|
|
partialForce.z += dzfp = delz * fpair;
|
|
|
|
if(vflag) {
|
|
partialVirial1.x += delx * dxfp;
|
|
partialVirial1.y += dely * dyfp;
|
|
partialVirial1.z += delz * dzfp;
|
|
partialVirial2.x += delx * dyfp;
|
|
partialVirial2.y += delx * dzfp;
|
|
partialVirial2.z += dely * dzfp;
|
|
}
|
|
}
|
|
}
|
|
}
|
|
|
|
if(eflag) {
|
|
sharedEnergy[threadIdx.x] = evdwl;
|
|
|
|
if(coul_type != COUL_NONE)
|
|
sharedEnergyCoul[threadIdx.x] = ecoul;
|
|
}
|
|
|
|
sharedForce[threadIdx.x] = partialForce;
|
|
|
|
if(vflag) {
|
|
sharedVirial1[threadIdx.x] = partialVirial1;
|
|
sharedVirial2[threadIdx.x] = partialVirial2;
|
|
}
|
|
|
|
__syncthreads();
|
|
|
|
|
|
for(unsigned int s = blockDim.x >> 1; s > 0; s >>= 1) {
|
|
|
|
if(threadIdx.x < s) {
|
|
sharedForce[ threadIdx.x ].x += sharedForce[ threadIdx.x + s ].x;
|
|
sharedForce[ threadIdx.x ].y += sharedForce[ threadIdx.x + s ].y;
|
|
sharedForce[ threadIdx.x ].z += sharedForce[ threadIdx.x + s ].z;
|
|
|
|
if(vflag) {
|
|
sharedVirial1[ threadIdx.x ].x += sharedVirial1[ threadIdx.x + s ].x;
|
|
sharedVirial1[ threadIdx.x ].y += sharedVirial1[ threadIdx.x + s ].y;
|
|
sharedVirial1[ threadIdx.x ].z += sharedVirial1[ threadIdx.x + s ].z;
|
|
|
|
sharedVirial2[ threadIdx.x ].x += sharedVirial2[ threadIdx.x + s ].x;
|
|
sharedVirial2[ threadIdx.x ].y += sharedVirial2[ threadIdx.x + s ].y;
|
|
sharedVirial2[ threadIdx.x ].z += sharedVirial2[ threadIdx.x + s ].z;
|
|
}
|
|
|
|
if(eflag) {
|
|
sharedEnergy[ threadIdx.x ] += sharedEnergy[ threadIdx.x + s ];
|
|
|
|
if(coul_type != COUL_NONE)
|
|
sharedEnergyCoul[ threadIdx.x ] += sharedEnergyCoul[ threadIdx.x + s ];
|
|
}
|
|
}
|
|
|
|
__syncthreads();
|
|
}
|
|
|
|
if(threadIdx.x == 0) {
|
|
|
|
ENERGY_FLOAT* buffer = (ENERGY_FLOAT*) _buffer;
|
|
|
|
if(eflag) {
|
|
ENERGY_FLOAT tmp_evdwl;
|
|
buffer[blockIdx.x * gridDim.y + blockIdx.y + 0 * gridDim.x * gridDim.y] = tmp_evdwl = ENERGY_F(0.5) * sharedEnergy[0];
|
|
|
|
if(eflag_atom)
|
|
_eatom[i] = tmp_evdwl;
|
|
|
|
buffer = &buffer[gridDim.x * gridDim.y];
|
|
|
|
if(coul_type != COUL_NONE) {
|
|
buffer[blockIdx.x * gridDim.y + blockIdx.y + 0 * gridDim.x * gridDim.y] = tmp_evdwl = ENERGY_F(0.5) * sharedEnergyCoul[0];
|
|
|
|
if(eflag_atom)
|
|
_eatom[i] += tmp_evdwl;
|
|
|
|
buffer = &buffer[gridDim.x * gridDim.y];
|
|
}
|
|
}
|
|
|
|
if(vflag) {
|
|
ENERGY_FLOAT tmp;
|
|
buffer[blockIdx.x * gridDim.y + blockIdx.y + 0 * gridDim.x * gridDim.y] = tmp = ENERGY_F(0.5) * sharedVirial1[0].x;
|
|
|
|
if(vflag_atom) _vatom[i + 0 * _nmax] = tmp;
|
|
|
|
buffer[blockIdx.x * gridDim.y + blockIdx.y + 1 * gridDim.x * gridDim.y] = tmp = ENERGY_F(0.5) * sharedVirial1[0].y;
|
|
|
|
if(vflag_atom) _vatom[i + 1 * _nmax] = tmp;
|
|
|
|
buffer[blockIdx.x * gridDim.y + blockIdx.y + 2 * gridDim.x * gridDim.y] = tmp = ENERGY_F(0.5) * sharedVirial1[0].z;
|
|
|
|
if(vflag_atom) _vatom[i + 2 * _nmax] = tmp;
|
|
|
|
buffer[blockIdx.x * gridDim.y + blockIdx.y + 3 * gridDim.x * gridDim.y] = tmp = ENERGY_F(0.5) * sharedVirial2[0].x;
|
|
|
|
if(vflag_atom) _vatom[i + 3 * _nmax] = tmp;
|
|
|
|
buffer[blockIdx.x * gridDim.y + blockIdx.y + 4 * gridDim.x * gridDim.y] = tmp = ENERGY_F(0.5) * sharedVirial2[0].y;
|
|
|
|
if(vflag_atom) _vatom[i + 4 * _nmax] = tmp;
|
|
|
|
buffer[blockIdx.x * gridDim.y + blockIdx.y + 5 * gridDim.x * gridDim.y] = tmp = ENERGY_F(0.5) * sharedVirial2[0].z;
|
|
|
|
if(vflag_atom) _vatom[i + 5 * _nmax] = tmp;
|
|
|
|
buffer = &buffer[6 * gridDim.x * gridDim.y];
|
|
}
|
|
|
|
F_FLOAT* my_f;
|
|
|
|
if(_collect_forces_later) {
|
|
my_f = (F_FLOAT*) buffer;
|
|
my_f += i;
|
|
*my_f = sharedForce[0].x;
|
|
my_f += _nmax;
|
|
*my_f = sharedForce[0].y;
|
|
my_f += _nmax;
|
|
*my_f = sharedForce[0].z;
|
|
} else {
|
|
my_f = _f + i;
|
|
*my_f += sharedForce[0].x;
|
|
my_f += _nmax;
|
|
*my_f += sharedForce[0].y;
|
|
my_f += _nmax;
|
|
*my_f += sharedForce[0].z;
|
|
}
|
|
}
|
|
}
|
|
|
|
|
|
template <const PAIR_FORCES pair_type, const COUL_FORCES coul_type, const unsigned int extended_data>
|
|
__global__ void Pair_Kernel_TpA_opt(int eflag, int vflag, int eflag_atom, int vflag_atom, int comm_phase)
|
|
{
|
|
ENERGY_FLOAT evdwl = ENERGY_F(0.0);
|
|
ENERGY_FLOAT ecoul = ENERGY_F(0.0);
|
|
|
|
ENERGY_FLOAT* sharedE;
|
|
ENERGY_FLOAT* sharedECoul;
|
|
ENERGY_FLOAT* sharedV = &sharedmem[threadIdx.x];
|
|
|
|
if(eflag || eflag_atom) {
|
|
sharedE = &sharedmem[threadIdx.x];
|
|
sharedE[0] = ENERGY_F(0.0);
|
|
sharedV += blockDim.x;
|
|
|
|
if(coul_type != COUL_NONE) {
|
|
sharedECoul = sharedE + blockDim.x;
|
|
sharedECoul[0] = ENERGY_F(0.0);
|
|
sharedV += blockDim.x;
|
|
}
|
|
}
|
|
|
|
if(vflag || vflag_atom) {
|
|
sharedV[0 * blockDim.x] = ENERGY_F(0.0);
|
|
sharedV[1 * blockDim.x] = ENERGY_F(0.0);
|
|
sharedV[2 * blockDim.x] = ENERGY_F(0.0);
|
|
sharedV[3 * blockDim.x] = ENERGY_F(0.0);
|
|
sharedV[4 * blockDim.x] = ENERGY_F(0.0);
|
|
sharedV[5 * blockDim.x] = ENERGY_F(0.0);
|
|
}
|
|
|
|
int ii = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
|
|
|
|
X_FLOAT xtmp, ytmp, ztmp;
|
|
X_FLOAT4 myxtype;
|
|
F_FLOAT fxtmp, fytmp, fztmp, fpair;
|
|
F_FLOAT delx, dely, delz;
|
|
F_FLOAT factor_lj, factor_coul;
|
|
F_FLOAT qtmp;
|
|
int itype, i, j;
|
|
int jnum = 0;
|
|
int* jlist;
|
|
|
|
if(ii < (comm_phase < 2 ? _inum : _inum_border[0])) {
|
|
i = comm_phase < 2 ? _ilist[ii] : _ilist_border[ii] ;
|
|
|
|
myxtype = fetchXType(i);
|
|
myxtype = _x_type[i];
|
|
xtmp = myxtype.x;
|
|
ytmp = myxtype.y;
|
|
ztmp = myxtype.z;
|
|
itype = static_cast <int>(myxtype.w);
|
|
|
|
|
|
fxtmp = F_F(0.0);
|
|
fytmp = F_F(0.0);
|
|
fztmp = F_F(0.0);
|
|
|
|
if(coul_type != COUL_NONE)
|
|
qtmp = fetchQ(i);
|
|
|
|
jnum = comm_phase == 0 ? _numneigh[i] : (comm_phase == 1 ? _numneigh_inner[i] : _numneigh_border[ii]);
|
|
|
|
|
|
jlist = comm_phase == 0 ? &_neighbors[i] : (comm_phase == 1 ? &_neighbors_inner[i] : &_neighbors_border[ii]);
|
|
}
|
|
|
|
__syncthreads();
|
|
|
|
for(int jj = 0; jj < jnum; jj++) {
|
|
if(ii < (comm_phase < 2 ? _inum : _inum_border[0]))
|
|
if(jj < jnum) {
|
|
fpair = F_F(0.0);
|
|
j = jlist[jj * _nlocal];
|
|
|
|
factor_lj = j < _nall ? F_F(1.0) : _special_lj[j / _nall];
|
|
|
|
if(coul_type != COUL_NONE)
|
|
factor_coul = j < _nall ? F_F(1.0) : _special_coul[j / _nall];
|
|
|
|
j = j < _nall ? j : j % _nall;
|
|
|
|
myxtype = fetchXType(j);
|
|
delx = xtmp - myxtype.x;
|
|
dely = ytmp - myxtype.y;
|
|
delz = ztmp - myxtype.z;
|
|
int jtype = static_cast <int>(myxtype.w);
|
|
|
|
|
|
const F_FLOAT rsq = delx * delx + dely * dely + delz * delz;
|
|
|
|
bool in_cutoff = rsq < (_cutsq_global > X_F(0.0) ? _cutsq_global : _cutsq[itype * _cuda_ntypes + jtype]);
|
|
|
|
if(in_cutoff) {
|
|
switch(pair_type) {
|
|
case PAIR_BORN:
|
|
fpair += PairBornCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_BUCK:
|
|
fpair += PairBuckCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_CG_CMM:
|
|
fpair += PairLJSDKCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_CHARMM:
|
|
fpair += PairLJCharmmCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_CLASS2:
|
|
fpair += PairLJClass2Cuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_CUT:
|
|
fpair += PairLJCutCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_EXPAND:
|
|
fpair += PairLJExpandCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_GROMACS:
|
|
fpair += PairLJGromacsCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_SMOOTH:
|
|
fpair += PairLJSmoothCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ96_CUT:
|
|
fpair += PairLJ96CutCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_MORSE_R6:
|
|
fpair += PairMorseR6Cuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_MORSE:
|
|
fpair += PairMorseCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
}
|
|
}
|
|
|
|
if(coul_type != COUL_NONE) {
|
|
const F_FLOAT qiqj = qtmp * fetchQ(j);
|
|
|
|
if(qiqj * qiqj > 1e-8) {
|
|
const bool in_coul_cutoff =
|
|
rsq < (_cut_coulsq_global > X_F(0.0) ? _cut_coulsq_global : _cut_coulsq[itype * _cuda_ntypes + jtype]);
|
|
|
|
if(in_coul_cutoff) {
|
|
switch(coul_type) {
|
|
case COUL_CHARMM:
|
|
fpair += CoulCharmmCuda_Eval(rsq, factor_coul, eflag, ecoul, qiqj);
|
|
break;
|
|
|
|
case COUL_CHARMM_IMPLICIT:
|
|
fpair += CoulCharmmImplicitCuda_Eval(rsq, factor_coul, eflag, ecoul, qiqj);
|
|
break;
|
|
|
|
case COUL_CUT: {
|
|
const F_FLOAT forcecoul = factor_coul * _qqrd2e * qiqj * _RSQRT_(rsq);
|
|
|
|
if(eflag) {
|
|
ecoul += forcecoul;
|
|
}
|
|
|
|
fpair += forcecoul * (F_F(1.0) / rsq);
|
|
}
|
|
break;
|
|
|
|
case COUL_DEBYE: {
|
|
const F_FLOAT r2inv = F_F(1.0) / rsq;
|
|
const X_FLOAT r = _RSQRT_(r2inv);
|
|
const X_FLOAT rinv = F_F(1.0) / r;
|
|
const F_FLOAT screening = _EXP_(-_kappa * r);
|
|
F_FLOAT forcecoul = factor_coul * _qqrd2e * qiqj * screening ;
|
|
|
|
if(eflag) {
|
|
ecoul += forcecoul * rinv;
|
|
}
|
|
|
|
forcecoul *= (_kappa + rinv);
|
|
fpair += forcecoul * r2inv;
|
|
}
|
|
break;
|
|
|
|
case COUL_GROMACS:
|
|
fpair += CoulGromacsCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_coul, eflag, ecoul, qiqj);
|
|
break;
|
|
|
|
case COUL_LONG: {
|
|
const F_FLOAT r2inv = F_F(1.0) / rsq;
|
|
const F_FLOAT r = _RSQRT_(r2inv);
|
|
const F_FLOAT grij = _g_ewald * r;
|
|
const F_FLOAT expm2 = _EXP_(-grij * grij);
|
|
const F_FLOAT t = F_F(1.0) / (F_F(1.0) + EWALD_P * grij);
|
|
const F_FLOAT erfc = t * (A1 + t * (A2 + t * (A3 + t * (A4 + t * A5)))) * expm2;
|
|
const F_FLOAT prefactor = _qqrd2e * qiqj * (F_F(1.0) / r);
|
|
F_FLOAT forcecoul = prefactor * (erfc + EWALD_F * grij * expm2);
|
|
|
|
if(factor_coul < 1.0) forcecoul -= (1.0 - factor_coul) * prefactor;
|
|
|
|
if(eflag) {
|
|
ecoul += prefactor * erfc;
|
|
|
|
if(factor_coul < 1.0) ecoul -= (1.0 - factor_coul) * prefactor;
|
|
}
|
|
|
|
fpair += forcecoul * r2inv;
|
|
}
|
|
break;
|
|
|
|
}
|
|
}
|
|
|
|
in_cutoff = in_cutoff || in_coul_cutoff;
|
|
}
|
|
}
|
|
|
|
|
|
if(in_cutoff) {
|
|
F_FLOAT dxfp, dyfp, dzfp;
|
|
fxtmp += dxfp = delx * fpair;
|
|
fytmp += dyfp = dely * fpair;
|
|
fztmp += dzfp = delz * fpair;
|
|
|
|
if(vflag) {
|
|
sharedV[0 * blockDim.x] += delx * dxfp;
|
|
sharedV[1 * blockDim.x] += dely * dyfp;
|
|
sharedV[2 * blockDim.x] += delz * dzfp;
|
|
sharedV[3 * blockDim.x] += delx * dyfp;
|
|
sharedV[4 * blockDim.x] += delx * dzfp;
|
|
sharedV[5 * blockDim.x] += dely * dzfp;
|
|
}
|
|
}
|
|
}
|
|
}
|
|
|
|
__syncthreads();
|
|
|
|
if(ii < (comm_phase < 2 ? _inum : _inum_border[0])) {
|
|
F_FLOAT* my_f;
|
|
|
|
if(_collect_forces_later) {
|
|
ENERGY_FLOAT* buffer = (ENERGY_FLOAT*) _buffer;
|
|
|
|
if(eflag) {
|
|
buffer = &buffer[1 * gridDim.x * gridDim.y];
|
|
|
|
if(coul_type != COUL_NONE)
|
|
buffer = &buffer[1 * gridDim.x * gridDim.y];
|
|
}
|
|
|
|
if(vflag) {
|
|
buffer = &buffer[6 * gridDim.x * gridDim.y];
|
|
}
|
|
|
|
my_f = (F_FLOAT*) buffer;
|
|
my_f += i;
|
|
*my_f = fxtmp;
|
|
my_f += _nmax;
|
|
*my_f = fytmp;
|
|
my_f += _nmax;
|
|
*my_f = fztmp;
|
|
} else {
|
|
my_f = _f + i;
|
|
*my_f += fxtmp;
|
|
my_f += _nmax;
|
|
*my_f += fytmp;
|
|
my_f += _nmax;
|
|
*my_f += fztmp;
|
|
}
|
|
}
|
|
|
|
__syncthreads();
|
|
|
|
if(eflag) {
|
|
sharedE[0] = evdwl;
|
|
|
|
if(coul_type != COUL_NONE)
|
|
sharedECoul[0] = ecoul;
|
|
}
|
|
|
|
if(eflag_atom && i < _nlocal) {
|
|
if(coul_type != COUL_NONE)
|
|
_eatom[i] += evdwl + ecoul;
|
|
else
|
|
_eatom[i] += evdwl;
|
|
}
|
|
|
|
if(vflag_atom && i < _nlocal) {
|
|
_vatom[i] += ENERGY_F(0.5) * sharedV[0 * blockDim.x];
|
|
_vatom[i + _nmax] += ENERGY_F(0.5) * sharedV[1 * blockDim.x];
|
|
_vatom[i + 2 * _nmax] += ENERGY_F(0.5) * sharedV[2 * blockDim.x];
|
|
_vatom[i + 3 * _nmax] += ENERGY_F(0.5) * sharedV[3 * blockDim.x];
|
|
_vatom[i + 4 * _nmax] += ENERGY_F(0.5) * sharedV[4 * blockDim.x];
|
|
_vatom[i + 5 * _nmax] += ENERGY_F(0.5) * sharedV[5 * blockDim.x];
|
|
}
|
|
|
|
if(vflag || eflag) PairVirialCompute_A_Kernel(eflag, vflag, coul_type != COUL_NONE ? 1 : 0);
|
|
}
|
|
|
|
template <const PAIR_FORCES pair_type, const COUL_FORCES coul_type, const unsigned int extended_data>
|
|
__global__ void Pair_Kernel_BpA_opt(int eflag, int vflag, int eflag_atom, int vflag_atom, int comm_phase)
|
|
{
|
|
int ii = (blockIdx.x * gridDim.y + blockIdx.y);
|
|
|
|
if(ii >= (comm_phase < 2 ? _inum : _inum_border[0]))
|
|
return;
|
|
|
|
ENERGY_FLOAT evdwl = ENERGY_F(0.0);
|
|
ENERGY_FLOAT ecoul = ENERGY_F(0.0);
|
|
F_FLOAT3* sharedVirial1;
|
|
F_FLOAT3* sharedVirial2;
|
|
F_FLOAT* sharedEnergy;
|
|
F_FLOAT* sharedEnergyCoul;
|
|
|
|
F_FLOAT3* sharedForce = (F_FLOAT3*) &sharedmem[0];
|
|
|
|
if(vflag) {
|
|
sharedVirial1 = &sharedForce[64];
|
|
sharedVirial2 = &sharedVirial1[64];
|
|
} else {
|
|
sharedVirial1 = &sharedForce[0];
|
|
sharedVirial2 = &sharedVirial1[0];
|
|
}
|
|
|
|
if(eflag) {
|
|
if(vflag || vflag_atom)
|
|
sharedEnergy = (F_FLOAT*) &sharedVirial2[64];
|
|
else
|
|
sharedEnergy = (F_FLOAT*) &sharedForce[64];
|
|
|
|
if(coul_type != COUL_NONE)
|
|
sharedEnergyCoul = (F_FLOAT*) &sharedEnergy[64];
|
|
|
|
}
|
|
|
|
F_FLOAT3 partialForce = { F_F(0.0), F_F(0.0), F_F(0.0) };
|
|
F_FLOAT3 partialVirial1 = { F_F(0.0), F_F(0.0), F_F(0.0) };
|
|
F_FLOAT3 partialVirial2 = { F_F(0.0), F_F(0.0), F_F(0.0) };
|
|
|
|
X_FLOAT xtmp, ytmp, ztmp;
|
|
X_FLOAT4 myxtype;
|
|
F_FLOAT delx, dely, delz;
|
|
F_FLOAT factor_lj, factor_coul;
|
|
F_FLOAT fpair;
|
|
F_FLOAT qtmp;
|
|
int itype, jnum, i, j;
|
|
int* jlist;
|
|
|
|
i = comm_phase < 2 ? _ilist[ii] : _ilist_border[ii];
|
|
|
|
myxtype = fetchXType(i);
|
|
|
|
xtmp = myxtype.x;
|
|
ytmp = myxtype.y;
|
|
ztmp = myxtype.z;
|
|
itype = static_cast <int>(myxtype.w);
|
|
|
|
if(coul_type != COUL_NONE)
|
|
qtmp = fetchQ(i);
|
|
|
|
jnum = comm_phase == 0 ? _numneigh[i] : (comm_phase == 1 ? _numneigh_inner[i] : _numneigh_border[ii]);
|
|
|
|
jlist = comm_phase == 0 ? &_neighbors[i * _maxneighbors] : (comm_phase == 1 ? &_neighbors_inner[i * _maxneighbors] : &_neighbors_border[ii * _maxneighbors]);
|
|
__syncthreads();
|
|
|
|
for(int jj = threadIdx.x; jj < jnum + blockDim.x; jj += blockDim.x) {
|
|
if(jj < jnum) {
|
|
fpair = F_F(0.0);
|
|
j = jlist[jj];
|
|
factor_lj = j < _nall ? F_F(1.0) : _special_lj[j / _nall];
|
|
|
|
if(coul_type != COUL_NONE)
|
|
factor_coul = j < _nall ? F_F(1.0) : _special_coul[j / _nall];
|
|
|
|
j = j < _nall ? j : j % _nall;
|
|
|
|
myxtype = fetchXType(j);
|
|
|
|
delx = xtmp - myxtype.x;
|
|
dely = ytmp - myxtype.y;
|
|
delz = ztmp - myxtype.z;
|
|
int jtype = static_cast <int>(myxtype.w);
|
|
|
|
const F_FLOAT rsq = delx * delx + dely * dely + delz * delz;
|
|
|
|
bool in_cutoff = rsq < (_cutsq_global > X_F(0.0) ? _cutsq_global : _cutsq[itype * _cuda_ntypes + jtype]);
|
|
bool in_coul_cutoff;
|
|
|
|
if(in_cutoff) {
|
|
switch(pair_type) {
|
|
case PAIR_BORN:
|
|
fpair += PairBornCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_BUCK:
|
|
fpair += PairBuckCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_CG_CMM:
|
|
fpair += PairLJSDKCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_CHARMM:
|
|
fpair += PairLJCharmmCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_CLASS2:
|
|
fpair += PairLJClass2Cuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_CUT:
|
|
fpair += PairLJCutCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_EXPAND:
|
|
fpair += PairLJExpandCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_GROMACS:
|
|
fpair += PairLJGromacsCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ_SMOOTH:
|
|
fpair += PairLJSmoothCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_LJ96_CUT:
|
|
fpair += PairLJ96CutCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_MORSE_R6:
|
|
fpair += PairMorseR6Cuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
|
|
case PAIR_MORSE:
|
|
fpair += PairMorseCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_lj, eflag, evdwl);
|
|
break;
|
|
}
|
|
}
|
|
|
|
if(coul_type != COUL_NONE) {
|
|
const F_FLOAT qiqj = qtmp * fetchQ(j);
|
|
|
|
if(qiqj * qiqj > (1e-8f)) {
|
|
in_coul_cutoff =
|
|
rsq < (_cut_coulsq_global > X_F(0.0) ? _cut_coulsq_global : _cut_coulsq[itype * _cuda_ntypes + jtype]);
|
|
|
|
if(in_coul_cutoff) {
|
|
switch(coul_type) {
|
|
case COUL_CHARMM:
|
|
fpair += CoulCharmmCuda_Eval(rsq, factor_coul, eflag, ecoul, qiqj);
|
|
break;
|
|
|
|
case COUL_CHARMM_IMPLICIT:
|
|
fpair += CoulCharmmImplicitCuda_Eval(rsq, factor_coul, eflag, ecoul, qiqj);
|
|
break;
|
|
|
|
case COUL_GROMACS:
|
|
fpair += CoulGromacsCuda_Eval(rsq, itype * _cuda_ntypes + jtype, factor_coul, eflag, ecoul, qiqj);
|
|
break;
|
|
|
|
case COUL_LONG: {
|
|
const F_FLOAT r2inv = F_F(1.0) / rsq;
|
|
const F_FLOAT r = _RSQRT_(r2inv);
|
|
const F_FLOAT grij = _g_ewald * r;
|
|
const F_FLOAT expm2 = _EXP_(-grij * grij);
|
|
const F_FLOAT t = F_F(1.0) / (F_F(1.0) + EWALD_P * grij);
|
|
const F_FLOAT erfc = t * (A1 + t * (A2 + t * (A3 + t * (A4 + t * A5)))) * expm2;
|
|
const F_FLOAT prefactor = _qqrd2e * qiqj * (F_F(1.0) / r);
|
|
F_FLOAT forcecoul = prefactor * (erfc + EWALD_F * grij * expm2);
|
|
|
|
if(factor_coul < 1.0) forcecoul -= (1.0 - factor_coul) * prefactor;
|
|
|
|
if(eflag) {
|
|
ecoul += prefactor * erfc;
|
|
|
|
if(factor_coul < 1.0) ecoul -= (1.0 - factor_coul) * prefactor;
|
|
}
|
|
|
|
fpair += forcecoul * r2inv;
|
|
}
|
|
break;
|
|
|
|
case COUL_DEBYE: {
|
|
const F_FLOAT r2inv = F_F(1.0) / rsq;
|
|
const X_FLOAT r = _RSQRT_(r2inv);
|
|
const X_FLOAT rinv = F_F(1.0) / r;
|
|
const F_FLOAT screening = _EXP_(-_kappa * r);
|
|
F_FLOAT forcecoul = factor_coul * _qqrd2e * qiqj * screening ;
|
|
|
|
if(eflag) {
|
|
ecoul += forcecoul * rinv;
|
|
}
|
|
|
|
forcecoul *= (_kappa + rinv);
|
|
fpair += forcecoul * r2inv;
|
|
}
|
|
break;
|
|
|
|
case COUL_CUT: {
|
|
const F_FLOAT forcecoul = factor_coul * _qqrd2e * qiqj * _RSQRT_(rsq);
|
|
|
|
if(eflag) {
|
|
ecoul += forcecoul;
|
|
}
|
|
|
|
fpair += forcecoul * (F_F(1.0) / rsq);
|
|
}
|
|
break;
|
|
|
|
|
|
}
|
|
}
|
|
}
|
|
}
|
|
|
|
|
|
|
|
if(in_cutoff || in_coul_cutoff) {
|
|
F_FLOAT dxfp, dyfp, dzfp;
|
|
partialForce.x += dxfp = delx * fpair;
|
|
partialForce.y += dyfp = dely * fpair;
|
|
partialForce.z += dzfp = delz * fpair;
|
|
|
|
if(vflag) {
|
|
partialVirial1.x += delx * dxfp;
|
|
partialVirial1.y += dely * dyfp;
|
|
partialVirial1.z += delz * dzfp;
|
|
partialVirial2.x += delx * dyfp;
|
|
partialVirial2.y += delx * dzfp;
|
|
partialVirial2.z += dely * dzfp;
|
|
}
|
|
}
|
|
}
|
|
}
|
|
|
|
if(eflag) {
|
|
sharedEnergy[threadIdx.x] = evdwl;
|
|
|
|
if(coul_type != COUL_NONE)
|
|
sharedEnergyCoul[threadIdx.x] = ecoul;
|
|
}
|
|
|
|
sharedForce[threadIdx.x] = partialForce;
|
|
|
|
if(vflag) {
|
|
sharedVirial1[threadIdx.x] = partialVirial1;
|
|
sharedVirial2[threadIdx.x] = partialVirial2;
|
|
}
|
|
|
|
__syncthreads();
|
|
|
|
|
|
for(unsigned int s = blockDim.x >> 1; s > 0; s >>= 1) {
|
|
|
|
if(threadIdx.x < s) {
|
|
sharedForce[ threadIdx.x ].x += sharedForce[ threadIdx.x + s ].x;
|
|
sharedForce[ threadIdx.x ].y += sharedForce[ threadIdx.x + s ].y;
|
|
sharedForce[ threadIdx.x ].z += sharedForce[ threadIdx.x + s ].z;
|
|
|
|
if(vflag) {
|
|
sharedVirial1[ threadIdx.x ].x += sharedVirial1[ threadIdx.x + s ].x;
|
|
sharedVirial1[ threadIdx.x ].y += sharedVirial1[ threadIdx.x + s ].y;
|
|
sharedVirial1[ threadIdx.x ].z += sharedVirial1[ threadIdx.x + s ].z;
|
|
|
|
sharedVirial2[ threadIdx.x ].x += sharedVirial2[ threadIdx.x + s ].x;
|
|
sharedVirial2[ threadIdx.x ].y += sharedVirial2[ threadIdx.x + s ].y;
|
|
sharedVirial2[ threadIdx.x ].z += sharedVirial2[ threadIdx.x + s ].z;
|
|
}
|
|
|
|
if(eflag) {
|
|
sharedEnergy[ threadIdx.x ] += sharedEnergy[ threadIdx.x + s ];
|
|
|
|
if(coul_type != COUL_NONE)
|
|
sharedEnergyCoul[ threadIdx.x ] += sharedEnergyCoul[ threadIdx.x + s ];
|
|
}
|
|
}
|
|
|
|
__syncthreads();
|
|
}
|
|
|
|
if(threadIdx.x == 0) {
|
|
|
|
ENERGY_FLOAT* buffer = (ENERGY_FLOAT*) _buffer;
|
|
|
|
if(eflag) {
|
|
ENERGY_FLOAT tmp_evdwl;
|
|
buffer[blockIdx.x * gridDim.y + blockIdx.y + 0 * gridDim.x * gridDim.y] = tmp_evdwl = ENERGY_F(0.5) * sharedEnergy[0];
|
|
|
|
if(eflag_atom)
|
|
_eatom[i] = tmp_evdwl;
|
|
|
|
buffer = &buffer[gridDim.x * gridDim.y];
|
|
|
|
if(coul_type != COUL_NONE) {
|
|
buffer[blockIdx.x * gridDim.y + blockIdx.y + 0 * gridDim.x * gridDim.y] = tmp_evdwl = ENERGY_F(0.5) * sharedEnergyCoul[0];
|
|
|
|
if(eflag_atom)
|
|
_eatom[i] += tmp_evdwl;
|
|
|
|
buffer = &buffer[gridDim.x * gridDim.y];
|
|
}
|
|
}
|
|
|
|
if(vflag) {
|
|
ENERGY_FLOAT tmp;
|
|
buffer[blockIdx.x * gridDim.y + blockIdx.y + 0 * gridDim.x * gridDim.y] = tmp = ENERGY_F(0.5) * sharedVirial1[0].x;
|
|
|
|
if(vflag_atom) _vatom[i + 0 * _nmax] = tmp;
|
|
|
|
buffer[blockIdx.x * gridDim.y + blockIdx.y + 1 * gridDim.x * gridDim.y] = tmp = ENERGY_F(0.5) * sharedVirial1[0].y;
|
|
|
|
if(vflag_atom) _vatom[i + 1 * _nmax] = tmp;
|
|
|
|
buffer[blockIdx.x * gridDim.y + blockIdx.y + 2 * gridDim.x * gridDim.y] = tmp = ENERGY_F(0.5) * sharedVirial1[0].z;
|
|
|
|
if(vflag_atom) _vatom[i + 2 * _nmax] = tmp;
|
|
|
|
buffer[blockIdx.x * gridDim.y + blockIdx.y + 3 * gridDim.x * gridDim.y] = tmp = ENERGY_F(0.5) * sharedVirial2[0].x;
|
|
|
|
if(vflag_atom) _vatom[i + 3 * _nmax] = tmp;
|
|
|
|
buffer[blockIdx.x * gridDim.y + blockIdx.y + 4 * gridDim.x * gridDim.y] = tmp = ENERGY_F(0.5) * sharedVirial2[0].y;
|
|
|
|
if(vflag_atom) _vatom[i + 4 * _nmax] = tmp;
|
|
|
|
buffer[blockIdx.x * gridDim.y + blockIdx.y + 5 * gridDim.x * gridDim.y] = tmp = ENERGY_F(0.5) * sharedVirial2[0].z;
|
|
|
|
if(vflag_atom) _vatom[i + 5 * _nmax] = tmp;
|
|
|
|
buffer = &buffer[6 * gridDim.x * gridDim.y];
|
|
}
|
|
|
|
F_FLOAT* my_f;
|
|
|
|
if(_collect_forces_later) {
|
|
my_f = (F_FLOAT*) buffer;
|
|
my_f += i;
|
|
*my_f = sharedForce[0].x;
|
|
my_f += _nmax;
|
|
*my_f = sharedForce[0].y;
|
|
my_f += _nmax;
|
|
*my_f = sharedForce[0].z;
|
|
} else {
|
|
my_f = _f + i;
|
|
*my_f += sharedForce[0].x;
|
|
my_f += _nmax;
|
|
*my_f += sharedForce[0].y;
|
|
my_f += _nmax;
|
|
*my_f += sharedForce[0].z;
|
|
}
|
|
}
|
|
}
|
|
|
|
__global__ void Pair_GenerateXType_Kernel()
|
|
{
|
|
int i = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
|
|
|
|
if(i < _nall) {
|
|
X_FLOAT4 xtype;
|
|
xtype.x = _x[i];
|
|
xtype.y = _x[i + _nmax];
|
|
xtype.z = _x[i + 2 * _nmax];
|
|
xtype.w = _type[i];
|
|
_x_type[i] = xtype;
|
|
}
|
|
|
|
}
|
|
|
|
__global__ void Pair_GenerateVRadius_Kernel()
|
|
{
|
|
int i = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
|
|
|
|
if(i < _nall) {
|
|
V_FLOAT4 vradius;
|
|
vradius.x = _v[i];
|
|
vradius.y = _v[i + _nmax];
|
|
vradius.z = _v[i + 2 * _nmax];
|
|
vradius.w = _radius[i];
|
|
_v_radius[i] = vradius;
|
|
}
|
|
}
|
|
|
|
__global__ void Pair_GenerateOmegaRmass_Kernel()
|
|
{
|
|
int i = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
|
|
|
|
if(i < _nall) {
|
|
V_FLOAT4 omegarmass;
|
|
omegarmass.x = _omega[i];
|
|
omegarmass.y = _omega[i + _nmax];
|
|
omegarmass.z = _omega[i + 2 * _nmax];
|
|
omegarmass.w = _rmass[i];
|
|
_omega_rmass[i] = omegarmass;
|
|
}
|
|
}
|
|
|
|
__global__ void Pair_RevertXType_Kernel()
|
|
{
|
|
int i = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
|
|
|
|
if(i < _nall) {
|
|
X_FLOAT4 xtype = _x_type[i];
|
|
_x[i] = xtype.x;
|
|
_x[i + _nmax] = xtype.y;
|
|
_x[i + 2 * _nmax] = xtype.z;
|
|
_type[i] = static_cast <int>(xtype.w);
|
|
}
|
|
|
|
}
|
|
|
|
__global__ void Pair_BuildXHold_Kernel()
|
|
{
|
|
int i = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
|
|
|
|
if(i < _nall) {
|
|
X_FLOAT4 xtype = _x_type[i];
|
|
_xhold[i] = xtype.x;
|
|
_xhold[i + _nmax] = xtype.y;
|
|
_xhold[i + 2 * _nmax] = xtype.z;
|
|
}
|
|
|
|
}
|
|
|
|
__global__ void Pair_CollectForces_Kernel(int nperblock, int n)
|
|
{
|
|
int i = (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x;
|
|
|
|
if(i >= _nlocal) return;
|
|
|
|
ENERGY_FLOAT* buf = (ENERGY_FLOAT*) _buffer;
|
|
|
|
F_FLOAT* buf_f = (F_FLOAT*) &buf[nperblock * n];
|
|
F_FLOAT* my_f = _f + i;
|
|
buf_f += i;
|
|
*my_f += * buf_f;
|
|
my_f += _nmax;
|
|
buf_f += _nmax;
|
|
*my_f += * buf_f;
|
|
my_f += _nmax;
|
|
buf_f += _nmax;
|
|
*my_f += * buf_f;
|
|
my_f += _nmax;
|
|
}
|