Commit c488e620 authored by Trung Nguyen's avatar Trung Nguyen
Browse files

Similar bug fixes to vashishta/gpu

parent adacc7f2
Loading
Loading
Loading
Loading
+2 −3
Original line number Diff line number Diff line
@@ -233,10 +233,9 @@ void VashishtaT::loop(const bool _eflag, const bool _vflag, const int evatom) {
                               (BX/this->_threads_per_atom)));

  this->k_short_nbor.set_size(GX,BX);
  this->k_short_nbor.run(&this->atom->x, &param4, &map,
                 &elem2param, &_nelements, &_nparams,
  this->k_short_nbor.run(&this->atom->x, &elem2param, &_nelements, &_nparams,
                 &this->nbor->dev_nbor, &this->_nbor_data->begin(),
                 &this->dev_short_nbor, &ainum,
                 &this->dev_short_nbor, &_cutshortsq, &ainum,
                 &nbor_pitch, &this->_threads_per_atom);

  // this->_nbor_data == nbor->dev_packed for gpu_nbor == 0 and tpa > 1
+2 −8
Original line number Diff line number Diff line
@@ -137,13 +137,12 @@ _texture( param5_tex,int4);
#endif

__kernel void k_vashishta_short_nbor(const __global numtyp4 *restrict x_,
                                     const __global numtyp4 *restrict param4,
                                     const __global int *restrict map,
                                     const __global int *restrict elem2param,
                                     const int nelements, const int nparams,
                                     const __global int * dev_nbor,
                                     const __global int * dev_packed,
                                     __global int * dev_short_nbor,
                                     const double _cutshortsq,
                                     const int inum, const int nbor_pitch,
                                     const int t_per_atom) {
  __local int n_stride;
@@ -157,8 +156,6 @@ __kernel void k_vashishta_short_nbor(const __global numtyp4 *restrict x_,
              n_stride,nbor_end,nbor);

    numtyp4 ix; fetch4(ix,i,pos_tex); //x_[i];
    int itype=ix.w;
    itype=map[itype];

    int ncount = 0;
    int m = nbor;
@@ -172,9 +169,6 @@ __kernel void k_vashishta_short_nbor(const __global numtyp4 *restrict x_,
      j &= NEIGHMASK;

      numtyp4 jx; fetch4(jx,j,pos_tex); //x_[j];
      int jtype=jx.w;
      jtype=map[jtype];
      int ijparam=elem2param[itype*nelements*nelements+jtype*nelements+jtype];

      // Compute r12
      numtyp delx = ix.x-jx.x;
@@ -182,7 +176,7 @@ __kernel void k_vashishta_short_nbor(const __global numtyp4 *restrict x_,
      numtyp delz = ix.z-jx.z;
      numtyp rsq = delx*delx+dely*dely+delz*delz;

      if (rsq<param4[ijparam].x) { //param4[ijparam].x = r0sq; //param4[ijparam].z=cutsq
      if (rsq<_cutshortsq) {
        dev_short_nbor[nbor_short] = nj;
        nbor_short += n_stride;
        ncount++;