Commit b3e9006a authored by Tim Mattox's avatar Tim Mattox
Browse files

Merge branch 'master' into USER-DPD_kokkos

parents a1f4551a 997142a4

src/KOKKOS/angle_charmm_kokkos.cpp

+0 −3
Original line number Diff line number Diff line
@@ -84,9 +84,6 @@ void AngleCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_in) 
    }

  }



  if (eflag || vflag) atomKK->modified(execution_space,datamask_modify);

  else atomKK->modified(execution_space,F_MASK);



  x = atomKK->k_x.view<DeviceType>();

  f = atomKK->k_f.view<DeviceType>();

  neighborKK->k_anglelist.template sync<DeviceType>();

src/KOKKOS/bond_fene_kokkos.cpp

+0 −3
Original line number Diff line number Diff line
@@ -85,13 +85,10 @@ void BondFENEKokkos<DeviceType>::compute(int eflag_in, int vflag_in) 
    d_vatom = k_vatom.d_view;

  }



  atomKK->sync(execution_space,datamask_read);

  k_k.template sync<DeviceType>();

  k_r0.template sync<DeviceType>();

  k_epsilon.template sync<DeviceType>();

  k_sigma.template sync<DeviceType>();

  if (eflag || vflag) atomKK->modified(execution_space,datamask_modify);

  else atomKK->modified(execution_space,F_MASK);



  x = atomKK->k_x.view<DeviceType>();

  f = atomKK->k_f.view<DeviceType>();

src/KOKKOS/dihedral_charmm_kokkos.cpp

+0 −5
Original line number Diff line number Diff line
@@ -98,11 +98,6 @@ void DihedralCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_in) 
    }

  }





  //atomKK->sync(execution_space,datamask_read);

  if (eflag || vflag) atomKK->modified(execution_space,datamask_modify);

  else atomKK->modified(execution_space,F_MASK);



  x = atomKK->k_x.view<DeviceType>();

  f = atomKK->k_f.view<DeviceType>();

  q = atomKK->k_q.view<DeviceType>();

src/KOKKOS/dihedral_opls_kokkos.cpp

+0 −3
Original line number Diff line number Diff line
@@ -85,13 +85,10 @@ void DihedralOPLSKokkos<DeviceType>::compute(int eflag_in, int vflag_in) 
    d_vatom = k_vatom.d_view;

  }



  atomKK->sync(execution_space,datamask_read);

  k_k1.template sync<DeviceType>();

  k_k2.template sync<DeviceType>();

  k_k3.template sync<DeviceType>();

  k_k4.template sync<DeviceType>();

  if (eflag || vflag) atomKK->modified(execution_space,datamask_modify);

  else atomKK->modified(execution_space,F_MASK);



  x = atomKK->k_x.view<DeviceType>();

  f = atomKK->k_f.view<DeviceType>();

src/KOKKOS/npair_kokkos.cpp

+55 −23
Original line number Diff line number Diff line
@@ -24,8 +24,8 @@ namespace LAMMPS_NS { 


/* ---------------------------------------------------------------------- */



template<class DeviceType, int HALF_NEIGH, int GHOST>

NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::NPairKokkos(LAMMPS *lmp) : NPair(lmp) {

template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>

NPairKokkos<DeviceType,HALF_NEIGH,GHOST,TRI>::NPairKokkos(LAMMPS *lmp) : NPair(lmp) {



}
@@ -33,8 +33,8 @@ NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::NPairKokkos(LAMMPS *lmp) : NPair(lmp) 
   copy needed info from Neighbor class to this build class

   ------------------------------------------------------------------------- */



template<class DeviceType, int HALF_NEIGH, int GHOST>

void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::copy_neighbor_info()

template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>

void NPairKokkos<DeviceType,HALF_NEIGH,GHOST,TRI>::copy_neighbor_info()

{

  NPair::copy_neighbor_info();
@@ -62,8 +62,8 @@ void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::copy_neighbor_info() 
 copy per-atom and per-bin vectors from NBin class to this build class

 ------------------------------------------------------------------------- */



template<class DeviceType, int HALF_NEIGH, int GHOST>

void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::copy_bin_info()

template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>

void NPairKokkos<DeviceType,HALF_NEIGH,GHOST,TRI>::copy_bin_info()

{

  NPair::copy_bin_info();
@@ -78,8 +78,8 @@ void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::copy_bin_info() 
 copy needed info from NStencil class to this build class

 ------------------------------------------------------------------------- */



template<class DeviceType, int HALF_NEIGH, int GHOST>

void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::copy_stencil_info()

template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>

void NPairKokkos<DeviceType,HALF_NEIGH,GHOST,TRI>::copy_stencil_info()

{

  NPair::copy_stencil_info();
@@ -106,8 +106,8 @@ void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::copy_stencil_info() 


/* ---------------------------------------------------------------------- */



template<class DeviceType, int HALF_NEIGH, int GHOST>

void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::build(NeighList *list_)

template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>

void NPairKokkos<DeviceType,HALF_NEIGH,GHOST,TRI>::build(NeighList *list_)

{

  NeighListKokkos<DeviceType>* list = (NeighListKokkos<DeviceType>*) list_;

  const int nlocal = includegroup?atom->nfirst:atom->nlocal;
@@ -196,14 +196,14 @@ if (GHOST) { 
  Kokkos::parallel_for(nall, f);

} else {

  if (newton_pair) {

    NPairKokkosBuildFunctor<DeviceType,HALF_NEIGH,1> f(data,atoms_per_bin * 5 * sizeof(X_FLOAT) * factor);

    NPairKokkosBuildFunctor<DeviceType,TRI?0:HALF_NEIGH,1,TRI> f(data,atoms_per_bin * 5 * sizeof(X_FLOAT) * factor);

#ifdef KOKKOS_HAVE_CUDA

    Kokkos::parallel_for(config, f);

#else

    Kokkos::parallel_for(nall, f);

#endif

  } else {

    NPairKokkosBuildFunctor<DeviceType,HALF_NEIGH,0> f(data,atoms_per_bin * 5 * sizeof(X_FLOAT) * factor);

    NPairKokkosBuildFunctor<DeviceType,HALF_NEIGH,0,0> f(data,atoms_per_bin * 5 * sizeof(X_FLOAT) * factor);

#ifdef KOKKOS_HAVE_CUDA

    Kokkos::parallel_for(config, f);

#else
@@ -293,7 +293,7 @@ int NeighborKokkosExecute<DeviceType>::exclusion(const int &i,const int &j, 


/* ---------------------------------------------------------------------- */



template<class DeviceType> template<int HalfNeigh,int Newton>

template<class DeviceType> template<int HalfNeigh,int Newton,int Tri>

void NeighborKokkosExecute<DeviceType>::

   build_Item(const int &i) const

{
@@ -374,6 +374,16 @@ void NeighborKokkosExecute<DeviceType>:: 


        if(HalfNeigh && !Newton && (j < i)) continue;

        if(!HalfNeigh && j==i) continue;

        if(Tri) {

          if (x(j,2) < ztmp) continue;

          if (x(j,2) == ztmp) {

            if (x(j,1) < ytmp) continue;

            if (x(j,1) == ytmp) {

              if (x(j,0) < xtmp) continue;

              if (x(j,0) == xtmp && j <= i) continue;

            }

          }

        }

        if(exclude && exclusion(i,j,itype,jtype)) continue;



        const X_FLOAT delx = xtmp - x(j, 0);
@@ -428,7 +438,7 @@ extern __shared__ X_FLOAT sharedmem[]; 


/* ---------------------------------------------------------------------- */



template<class DeviceType> template<int HalfNeigh,int Newton>

template<class DeviceType> template<int HalfNeigh,int Newton,int Tri>

__device__ inline

void NeighborKokkosExecute<DeviceType>::build_ItemCuda(typename Kokkos::TeamPolicy<DeviceType>::member_type dev) const

{
@@ -491,6 +501,16 @@ void NeighborKokkosExecute<DeviceType>::build_ItemCuda(typename Kokkos::TeamPoli 
            ((j >= nlocal) && ((x(j, 2) < ztmp) || (x(j, 2) == ztmp && x(j, 1) < ytmp) ||

              (x(j, 2) == ztmp && x(j, 1)  == ytmp && x(j, 0) < xtmp)))))

        ) continue;

        if(Tri) {

          if (x(j,2) < ztmp) continue;

          if (x(j,2) == ztmp) {

            if (x(j,1) < ytmp) continue;

            if (x(j,1) == ytmp) {

              if (x(j,0) < xtmp) continue;

              if (x(j,0) == xtmp && j <= i) continue;

            }

          }

        }

      if(exclude && exclusion(i,j,itype,jtype)) continue;

      const X_FLOAT delx = xtmp - other_x[m];

      const X_FLOAT dely = ytmp - other_x[m + atoms_per_bin];
@@ -558,6 +578,16 @@ void NeighborKokkosExecute<DeviceType>::build_ItemCuda(typename Kokkos::TeamPoli 
        //if(HalfNeigh && (j < i))  continue;

        if(HalfNeigh && !Newton && (j < i)) continue;

        if(!HalfNeigh && j==i) continue;

        if(Tri) {

          if (x(j,2) < ztmp) continue;

          if (x(j,2) == ztmp) {

            if (x(j,1) < ytmp) continue;

            if (x(j,1) == ytmp) {

              if (x(j,0) < xtmp) continue;

              if (x(j,0) == xtmp && j <= i) continue;

            }

          }

        }

        if(exclude && exclusion(i,j,itype,jtype)) continue;



        const X_FLOAT delx = xtmp - other_x[m];
@@ -736,14 +766,16 @@ void NeighborKokkosExecute<DeviceType>:: 
}



namespace LAMMPS_NS {

template class NPairKokkos<LMPDeviceType,0,0>;

template class NPairKokkos<LMPDeviceType,0,1>;

template class NPairKokkos<LMPDeviceType,1,0>;

template class NPairKokkos<LMPDeviceType,1,1>;

template class NPairKokkos<LMPDeviceType,0,0,0>;

template class NPairKokkos<LMPDeviceType,0,1,0>;

template class NPairKokkos<LMPDeviceType,1,0,0>;

template class NPairKokkos<LMPDeviceType,1,1,0>;

template class NPairKokkos<LMPDeviceType,1,0,1>;

#ifdef KOKKOS_HAVE_CUDA

template class NPairKokkos<LMPHostType,0,0>;

template class NPairKokkos<LMPHostType,0,1>;

template class NPairKokkos<LMPHostType,1,0>;

template class NPairKokkos<LMPHostType,1,1>;

template class NPairKokkos<LMPHostType,0,0,0>;

template class NPairKokkos<LMPHostType,0,1,0>;

template class NPairKokkos<LMPHostType,1,0,0>;

template class NPairKokkos<LMPHostType,1,1,0>;

template class NPairKokkos<LMPHostType,1,0,1>;

#endif

}

CRA Git | Maintained and supported by SUSTech CRA and CCSE