Unverified Commit 79c0dda0 authored by Axel Kohlmeyer's avatar Axel Kohlmeyer Committed by GitHub
Browse files

Merge pull request #2187 from Vsevak/tip4p_pressure_fix 

Fix virial computation in GPU accelerated TIP4P
parents d55a0612 496bd55d

lib/gpu/lal_lj_tip4p_long.cu

+19 −27
Original line number Diff line number Diff line
@@ -110,10 +110,8 @@ __kernel void k_lj_tip4p_long_distrib(const __global numtyp4 *restrict x_, 
    const int typeO, const int typeH,

    const numtyp alpha,

    const __global numtyp *restrict q_, const __global acctyp4 *restrict ansO) {

  int tid, ii, offset;

  atom_info(t_per_atom,ii,tid,offset);

  int i = BLOCK_ID_X*(BLOCK_SIZE_X)+tid;



  int i = BLOCK_ID_X*(BLOCK_SIZE_X)+THREAD_ID_X;

  acctyp4 f;

  f.x=(acctyp)0; f.y=(acctyp)0; f.z=(acctyp)0;
@@ -122,6 +120,8 @@ __kernel void k_lj_tip4p_long_distrib(const __global numtyp4 *restrict x_, 
    int itype = ix.w;

    acctyp4 fM, vM;

    acctyp eM;

    // placement of the virial in engv depends on eflag value

    int engv_iter = eflag ? 2 : 0;

    if (itype == typeH) {

      int iO = hneigh[i*4];

      if (iO < inum) {
@@ -131,13 +131,13 @@ __kernel void k_lj_tip4p_long_distrib(const __global numtyp4 *restrict x_, 
        f.z += fM.z * (acctyp)0.5 * alpha;

        if (vflag > 0) {

          vM = ansO[inum  +iO];

          engv[inum*2 + i] += vM.x * (acctyp)0.5 * alpha;

          engv[inum*3 + i] += vM.y * (acctyp)0.5 * alpha;

          engv[inum*4 + i] += vM.z * (acctyp)0.5 * alpha;

          engv[inum*engv_iter + i] += vM.x * (acctyp)0.5 * alpha; engv_iter++;

          engv[inum*engv_iter + i] += vM.y * (acctyp)0.5 * alpha; engv_iter++;

          engv[inum*engv_iter + i] += vM.z * (acctyp)0.5 * alpha; engv_iter++;

          vM = ansO[inum*2+iO];

          engv[inum*5 + i] += vM.x * (acctyp)0.5 * alpha;

          engv[inum*6 + i] += vM.y * (acctyp)0.5 * alpha;

          engv[inum*7 + i] += vM.z * (acctyp)0.5 * alpha;

          engv[inum*engv_iter + i] += vM.x * (acctyp)0.5 * alpha; engv_iter++;

          engv[inum*engv_iter + i] += vM.y * (acctyp)0.5 * alpha; engv_iter++;

          engv[inum*engv_iter + i] += vM.z * (acctyp)0.5 * alpha;

        }

      }

    } else {
@@ -155,13 +155,13 @@ __kernel void k_lj_tip4p_long_distrib(const __global numtyp4 *restrict x_, 
      }

      if (vflag > 0) {

        vM = ansO[inum   + i];

        engv[inum*2 + i] += vM.x * (acctyp)(1 - alpha);

        engv[inum*3 + i] += vM.y * (acctyp)(1 - alpha);

        engv[inum*4 + i] += vM.z * (acctyp)(1 - alpha);

        engv[inum*engv_iter + i] += vM.x * (acctyp)(1 - alpha); engv_iter++;

        engv[inum*engv_iter + i] += vM.y * (acctyp)(1 - alpha); engv_iter++;

        engv[inum*engv_iter + i] += vM.z * (acctyp)(1 - alpha); engv_iter++;

        vM = ansO[inum*2 + i];

        engv[inum*5 + i] += vM.x * (acctyp)(1 - alpha);

        engv[inum*6 + i] += vM.y * (acctyp)(1 - alpha);

        engv[inum*7 + i] += vM.z * (acctyp)(1 - alpha);

        engv[inum*engv_iter + i] += vM.x * (acctyp)(1 - alpha); engv_iter++;

        engv[inum*engv_iter + i] += vM.y * (acctyp)(1 - alpha); engv_iter++;

        engv[inum*engv_iter + i] += vM.z * (acctyp)(1 - alpha);

      }

    }

    acctyp4 old=ans[i];
@@ -182,9 +182,8 @@ __kernel void k_lj_tip4p_reneigh(const __global numtyp4 *restrict x_, 
    const int typeO, const int typeH,

    const __global tagint *restrict tag, const __global int *restrict map,

    const __global int *restrict sametag) {

  int tid, ii, offset;

  atom_info(t_per_atom,ii,tid,offset);

  int i = BLOCK_ID_X*(BLOCK_SIZE_X)+tid;



  int i = BLOCK_ID_X*(BLOCK_SIZE_X)+THREAD_ID_X;



  if (i<nall) {

    numtyp4 ix; fetch4(ix,i,pos_tex); //x_[i];
@@ -237,13 +236,10 @@ __kernel void k_lj_tip4p_newsite(const __global numtyp4 *restrict x_, 
    __global numtyp4 *restrict m,

    const int typeO, const int typeH,

    const numtyp alpha, const __global numtyp *restrict q_) {

  int tid, ii, offset;

  atom_info(t_per_atom,ii,tid,offset);

  int i = BLOCK_ID_X*(BLOCK_SIZE_X)+tid;



  int i = BLOCK_ID_X*(BLOCK_SIZE_X)+THREAD_ID_X;



  if (i<nall) {

    int iO, iH1, iH2;

    iO  = i;

    numtyp4 ix; fetch4(ix,i,pos_tex); //x_[i];

    int itype = ix.w;

    if (itype == typeO) {
@@ -280,8 +276,6 @@ __kernel void k_lj_tip4p_long(const __global numtyp4 *restrict x_, 
  int tid, ii, offset;

  atom_info(t_per_atom,ii,tid,offset);



  const numtyp eq_zero = 1e-6;



  acctyp energy = (acctyp)0;

  acctyp e_coul = (acctyp)0;

  acctyp4 f, fO;
@@ -595,8 +589,6 @@ __kernel void k_lj_tip4p_long_fast(const __global numtyp4 *restrict x_, 
  int tid, ii, offset;

  atom_info(t_per_atom,ii,tid,offset);



  const numtyp eq_zero = 1e-6;



  __local numtyp4 lj1[MAX_SHARED_TYPES*MAX_SHARED_TYPES];

  __local numtyp4 lj3[MAX_SHARED_TYPES*MAX_SHARED_TYPES];

  __local numtyp sp_lj[8];

src/GPU/pair_lj_cut_tip4p_long_gpu.cpp

+1 −4
Original line number Diff line number Diff line
@@ -109,10 +109,7 @@ PairLJCutTIP4PLongGPU::~PairLJCutTIP4PLongGPU() 


void PairLJCutTIP4PLongGPU::compute(int eflag, int vflag)

{



  if (eflag || vflag) ev_setup(eflag,vflag);

  else evflag = vflag_fdotr = 0;



  ev_init(eflag,vflag);

  int nall = atom->nlocal + atom->nghost;

  int inum, host_start;

CRA Git | Maintained and supported by SUSTech CRA and CCSE