Refactor Kokkos AtomVec

src/DPD-REACT/atom_vec_dpd.h

@@ -24,7 +24,7 @@ AtomStyle(dpd,AtomVecDPD);
 
 namespace LAMMPS_NS {
 
-class AtomVecDPD : public AtomVec {
+class AtomVecDPD : virtual public AtomVec {
  public:
   AtomVecDPD(class LAMMPS *);
 

src/KOKKOS/Install.sh

@@ -64,8 +64,8 @@ action atom_vec_bond_kokkos.cpp atom_vec_bond.cpp
 action atom_vec_bond_kokkos.h atom_vec_bond.h
 action atom_vec_charge_kokkos.cpp
 action atom_vec_charge_kokkos.h
-action atom_vec_spin_kokkos.cpp
-action atom_vec_spin_kokkos.h
+action atom_vec_spin_kokkos.cpp atom_vec_spin.cpp
+action atom_vec_spin_kokkos.h atom_vec_spin.h
 action atom_vec_dpd_kokkos.cpp atom_vec_dpd.cpp
 action atom_vec_dpd_kokkos.h atom_vec_dpd.h
 action atom_vec_full_kokkos.cpp atom_vec_full.cpp

src/KOKKOS/atom_kokkos.cpp

@@ -104,23 +104,23 @@ AtomKokkos::~AtomKokkos()
 
 void AtomKokkos::sync(const ExecutionSpace space, unsigned int mask)
 {
-  if (space == Device && lmp->kokkos->auto_sync) ((AtomVecKokkos *) avec)->modified(Host, mask);
+  if (space == Device && lmp->kokkos->auto_sync) avecKK->modified(Host, mask);
 
-  ((AtomVecKokkos *) avec)->sync(space, mask);
+  avecKK->sync(space, mask);
 }
 
 /* ---------------------------------------------------------------------- */
 
 void AtomKokkos::modified(const ExecutionSpace space, unsigned int mask)
 {
-  ((AtomVecKokkos *) avec)->modified(space, mask);
+  avecKK->modified(space, mask);
 
-  if (space == Device && lmp->kokkos->auto_sync) ((AtomVecKokkos *) avec)->sync(Host, mask);
+  if (space == Device && lmp->kokkos->auto_sync) avecKK->sync(Host, mask);
 }
 
 void AtomKokkos::sync_overlapping_device(const ExecutionSpace space, unsigned int mask)
 {
-  ((AtomVecKokkos *) avec)->sync_overlapping_device(space, mask);
+  avecKK->sync_overlapping_device(space, mask);
 }
 /* ---------------------------------------------------------------------- */
 
@@ -380,5 +380,8 @@ AtomVec *AtomKokkos::new_avec(const std::string &style, int trysuffix, int &sfla
 {
   AtomVec *avec = Atom::new_avec(style, trysuffix, sflag);
   if (!avec->kokkosable) error->all(FLERR, "KOKKOS package requires a kokkos enabled atom_style");
+
+  avecKK = dynamic_cast<AtomVecKokkos*>(avec);
+
   return avec;
 }

src/KOKKOS/atom_kokkos.h

@@ -78,6 +78,8 @@ class AtomKokkos : public Atom {
   DAT::tdual_int_scalar k_error_flag;
   dual_hash_type k_map_hash;
 
+  class AtomVecKokkos* avecKK;
+
   // map lookup function inlined for efficiency
   // return -1 if no map defined
 

src/KOKKOS/atom_vec_angle_kokkos.cpp

@@ -27,27 +27,10 @@ using namespace LAMMPS_NS;
 
 /* ---------------------------------------------------------------------- */
 
-AtomVecAngleKokkos::AtomVecAngleKokkos(LAMMPS *lmp) : AtomVecKokkos(lmp)
+AtomVecAngleKokkos::AtomVecAngleKokkos(LAMMPS *lmp) : AtomVec(lmp),
+AtomVecKokkos(lmp), AtomVecAngle(lmp)
 {
-  molecular = Atom::MOLECULAR;
-  bonds_allow = angles_allow = 1;
-  mass_type = PER_TYPE;
 
-  comm_x_only = comm_f_only = 1;
-  size_forward = 3;
-  size_reverse = 3;
-  size_border = 7;
-  size_velocity = 3;
-  size_data_atom = 6;
-  size_data_vel = 4;
-  xcol_data = 4;
-
-  atom->molecule_flag = 1;
-
-  k_count = DAT::tdual_int_1d("atom::k_count",1);
-  atomKK = (AtomKokkos *) atom;
-  commKK = (CommKokkos *) comm;
-  buffer = nullptr;
 }
 
 /* ----------------------------------------------------------------------
@@ -172,52 +155,6 @@ void AtomVecAngleKokkos::grow_pointers()
   h_angle_atom3 = atomKK->k_angle_atom3.h_view;
 }
 
-/* ----------------------------------------------------------------------
-   copy atom I info to atom J
-------------------------------------------------------------------------- */
-
-void AtomVecAngleKokkos::copy(int i, int j, int delflag)
-{
-  int k;
-
-  h_tag[j] = h_tag[i];
-  h_type[j] = h_type[i];
-  mask[j] = mask[i];
-  h_image[j] = h_image[i];
-  h_x(j,0) = h_x(i,0);
-  h_x(j,1) = h_x(i,1);
-  h_x(j,2) = h_x(i,2);
-  h_v(j,0) = h_v(i,0);
-  h_v(j,1) = h_v(i,1);
-  h_v(j,2) = h_v(i,2);
-
-  h_molecule(j) = h_molecule(i);
-
-  h_num_bond(j) = h_num_bond(i);
-  for (k = 0; k < h_num_bond(j); k++) {
-    h_bond_type(j,k) = h_bond_type(i,k);
-    h_bond_atom(j,k) = h_bond_atom(i,k);
-  }
-
-  h_nspecial(j,0) = h_nspecial(i,0);
-  h_nspecial(j,1) = h_nspecial(i,1);
-  h_nspecial(j,2) = h_nspecial(i,2);
-  for (k = 0; k < h_nspecial(j,2); k++)
-    h_special(j,k) = h_special(i,k);
-
-  h_num_angle(j) = h_num_angle(i);
-  for (k = 0; k < h_num_angle(j); k++) {
-    h_angle_type(j,k) = h_angle_type(i,k);
-    h_angle_atom1(j,k) = h_angle_atom1(i,k);
-    h_angle_atom2(j,k) = h_angle_atom2(i,k);
-    h_angle_atom3(j,k) = h_angle_atom3(i,k);
-  }
-
-  if (atom->nextra_grow)
-    for (int iextra = 0; iextra < atom->nextra_grow; iextra++)
-      modify->fix[atom->extra_grow[iextra]]->copy_arrays(i,j,delflag);
-}
-
 /* ---------------------------------------------------------------------- */
 
 template<class DeviceType,int PBC_FLAG,int TRICLINIC>
@@ -508,171 +445,6 @@ void AtomVecAngleKokkos::unpack_comm_kokkos(const int &n, const int &first,
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecAngleKokkos::pack_comm(int n, int *list, double *buf,
-                                 int pbc_flag, int *pbc)
-{
-  int i,j,m;
-  double dx,dy,dz;
-
-  m = 0;
-  if (pbc_flag == 0) {
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0);
-      buf[m++] = h_x(j,1);
-      buf[m++] = h_x(j,2);
-    }
-  } else {
-    if (domain->triclinic == 0) {
-      dx = pbc[0]*domain->xprd;
-      dy = pbc[1]*domain->yprd;
-      dz = pbc[2]*domain->zprd;
-    } else {
-      dx = pbc[0]*domain->xprd + pbc[5]*domain->xy + pbc[4]*domain->xz;
-      dy = pbc[1]*domain->yprd + pbc[3]*domain->yz;
-      dz = pbc[2]*domain->zprd;
-    }
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0) + dx;
-      buf[m++] = h_x(j,1) + dy;
-      buf[m++] = h_x(j,2) + dz;
-    }
-  }
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecAngleKokkos::pack_comm_vel(int n, int *list, double *buf,
-                                     int pbc_flag, int *pbc)
-{
-  int i,j,m;
-  double dx,dy,dz,dvx,dvy,dvz;
-
-  m = 0;
-  if (pbc_flag == 0) {
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0);
-      buf[m++] = h_x(j,1);
-      buf[m++] = h_x(j,2);
-      buf[m++] = h_v(j,0);
-      buf[m++] = h_v(j,1);
-      buf[m++] = h_v(j,2);
-    }
-  } else {
-    if (domain->triclinic == 0) {
-      dx = pbc[0]*domain->xprd;
-      dy = pbc[1]*domain->yprd;
-      dz = pbc[2]*domain->zprd;
-    } else {
-      dx = pbc[0]*domain->xprd + pbc[5]*domain->xy + pbc[4]*domain->xz;
-      dy = pbc[1]*domain->yprd + pbc[3]*domain->yz;
-      dz = pbc[2]*domain->zprd;
-    }
-    if (!deform_vremap) {
-      for (i = 0; i < n; i++) {
-        j = list[i];
-        buf[m++] = h_x(j,0) + dx;
-        buf[m++] = h_x(j,1) + dy;
-        buf[m++] = h_x(j,2) + dz;
-        buf[m++] = h_v(j,0);
-        buf[m++] = h_v(j,1);
-        buf[m++] = h_v(j,2);
-      }
-    } else {
-      dvx = pbc[0]*h_rate[0] + pbc[5]*h_rate[5] + pbc[4]*h_rate[4];
-      dvy = pbc[1]*h_rate[1] + pbc[3]*h_rate[3];
-      dvz = pbc[2]*h_rate[2];
-      for (i = 0; i < n; i++) {
-        j = list[i];
-        buf[m++] = h_x(j,0) + dx;
-        buf[m++] = h_x(j,1) + dy;
-        buf[m++] = h_x(j,2) + dz;
-        if (mask[i] & deform_groupbit) {
-          buf[m++] = h_v(j,0) + dvx;
-          buf[m++] = h_v(j,1) + dvy;
-          buf[m++] = h_v(j,2) + dvz;
-        } else {
-          buf[m++] = h_v(j,0);
-          buf[m++] = h_v(j,1);
-          buf[m++] = h_v(j,2);
-        }
-      }
-    }
-  }
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-void AtomVecAngleKokkos::unpack_comm(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  for (i = first; i < last; i++) {
-    h_x(i,0) = buf[m++];
-    h_x(i,1) = buf[m++];
-    h_x(i,2) = buf[m++];
-  }
-}
-
-/* ---------------------------------------------------------------------- */
-
-void AtomVecAngleKokkos::unpack_comm_vel(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  for (i = first; i < last; i++) {
-    h_x(i,0) = buf[m++];
-    h_x(i,1) = buf[m++];
-    h_x(i,2) = buf[m++];
-    h_v(i,0) = buf[m++];
-    h_v(i,1) = buf[m++];
-    h_v(i,2) = buf[m++];
-  }
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecAngleKokkos::pack_reverse(int n, int first, double *buf)
-{
-  if (n > 0)
-    atomKK->sync(Host,F_MASK);
-
-  int m = 0;
-  const int last = first + n;
-  for (int i = first; i < last; i++) {
-    buf[m++] = h_f(i,0);
-    buf[m++] = h_f(i,1);
-    buf[m++] = h_f(i,2);
-  }
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-void AtomVecAngleKokkos::unpack_reverse(int n, int *list, double *buf)
-{
-  if (n > 0)
-    atomKK->modified(Host,F_MASK);
-
-  int m = 0;
-  for (int i = 0; i < n; i++) {
-    const int j = list[i];
-    h_f(j,0) += buf[m++];
-    h_f(j,1) += buf[m++];
-    h_f(j,2) += buf[m++];
-  }
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType,int PBC_FLAG>
 struct AtomVecAngleKokkos_PackBorder {
   typedef DeviceType device_type;
@@ -774,149 +546,6 @@ int AtomVecAngleKokkos::pack_border_kokkos(int n, DAT::tdual_int_2d k_sendlist,
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecAngleKokkos::pack_border(int n, int *list, double *buf,
-                                   int pbc_flag, int *pbc)
-{
-  int i,j,m;
-  double dx,dy,dz;
-
-  m = 0;
-  if (pbc_flag == 0) {
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0);
-      buf[m++] = h_x(j,1);
-      buf[m++] = h_x(j,2);
-      buf[m++] = ubuf(h_tag(j)).d;
-      buf[m++] = ubuf(h_type(j)).d;
-      buf[m++] = ubuf(h_mask(j)).d;
-      buf[m++] = ubuf(h_molecule(j)).d;
-    }
-  } else {
-    if (domain->triclinic == 0) {
-      dx = pbc[0]*domain->xprd;
-      dy = pbc[1]*domain->yprd;
-      dz = pbc[2]*domain->zprd;
-    } else {
-      dx = pbc[0];
-      dy = pbc[1];
-      dz = pbc[2];
-    }
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0) + dx;
-      buf[m++] = h_x(j,1) + dy;
-      buf[m++] = h_x(j,2) + dz;
-      buf[m++] = ubuf(h_tag(j)).d;
-      buf[m++] = ubuf(h_type(j)).d;
-      buf[m++] = ubuf(h_mask(j)).d;
-      buf[m++] = ubuf(h_molecule(j)).d;
-    }
-  }
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->pack_border(n,list,&buf[m]);
-
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecAngleKokkos::pack_border_vel(int n, int *list, double *buf,
-                                       int pbc_flag, int *pbc)
-{
-  int i,j,m;
-  double dx,dy,dz,dvx,dvy,dvz;
-
-  m = 0;
-  if (pbc_flag == 0) {
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0);
-      buf[m++] = h_x(j,1);
-      buf[m++] = h_x(j,2);
-      buf[m++] = ubuf(h_tag(j)).d;
-      buf[m++] = ubuf(h_type(j)).d;
-      buf[m++] = ubuf(h_mask(j)).d;
-      buf[m++] = ubuf(h_molecule(j)).d;
-      buf[m++] = h_v(j,0);
-      buf[m++] = h_v(j,1);
-      buf[m++] = h_v(j,2);
-    }
-  } else {
-    if (domain->triclinic == 0) {
-      dx = pbc[0]*domain->xprd;
-      dy = pbc[1]*domain->yprd;
-      dz = pbc[2]*domain->zprd;
-    } else {
-      dx = pbc[0];
-      dy = pbc[1];
-      dz = pbc[2];
-    }
-    if (!deform_vremap) {
-      for (i = 0; i < n; i++) {
-        j = list[i];
-        buf[m++] = h_x(j,0) + dx;
-        buf[m++] = h_x(j,1) + dy;
-        buf[m++] = h_x(j,2) + dz;
-        buf[m++] = ubuf(h_tag(j)).d;
-        buf[m++] = ubuf(h_type(j)).d;
-        buf[m++] = ubuf(h_mask(j)).d;
-        buf[m++] = ubuf(h_molecule(j)).d;
-        buf[m++] = h_v(j,0);
-        buf[m++] = h_v(j,1);
-        buf[m++] = h_v(j,2);
-      }
-    } else {
-      dvx = pbc[0]*h_rate[0] + pbc[5]*h_rate[5] + pbc[4]*h_rate[4];
-      dvy = pbc[1]*h_rate[1] + pbc[3]*h_rate[3];
-      dvz = pbc[2]*h_rate[2];
-      for (i = 0; i < n; i++) {
-        j = list[i];
-        buf[m++] = h_x(j,0) + dx;
-        buf[m++] = h_x(j,1) + dy;
-        buf[m++] = h_x(j,2) + dz;
-        buf[m++] = ubuf(h_tag(j)).d;
-        buf[m++] = ubuf(h_type(j)).d;
-        buf[m++] = ubuf(h_mask(j)).d;
-        buf[m++] = ubuf(h_molecule(j)).d;
-        if (mask[i] & deform_groupbit) {
-          buf[m++] = h_v(j,0) + dvx;
-          buf[m++] = h_v(j,1) + dvy;
-          buf[m++] = h_v(j,2) + dvz;
-        } else {
-          buf[m++] = h_v(j,0);
-          buf[m++] = h_v(j,1);
-          buf[m++] = h_v(j,2);
-        }
-      }
-    }
-  }
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->pack_border(n,list,&buf[m]);
-
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecAngleKokkos::pack_border_hybrid(int n, int *list, double *buf)
-{
-  int i,j,m;
-
-  m = 0;
-  for (i = 0; i < n; i++) {
-    j = list[i];
-    buf[m++] = h_molecule(j);
-  }
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType>
 struct AtomVecAngleKokkos_UnpackBorder {
   typedef DeviceType device_type;
@@ -977,78 +606,6 @@ void AtomVecAngleKokkos::unpack_border_kokkos(const int &n, const int &first,
 
 /* ---------------------------------------------------------------------- */
 
-void AtomVecAngleKokkos::unpack_border(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  while (last > nmax) grow(0);
-
-  for (i = first; i < last; i++) {
-    h_x(i,0) = buf[m++];
-    h_x(i,1) = buf[m++];
-    h_x(i,2) = buf[m++];
-    h_tag(i) =  (tagint)  ubuf(buf[m++]).i;
-    h_type(i) = (int) ubuf(buf[m++]).i;
-    h_mask(i) = (int) ubuf(buf[m++]).i;
-    h_molecule(i) = (tagint) ubuf(buf[m++]).i;
-  }
-
-  atomKK->modified(Host,X_MASK|TAG_MASK|TYPE_MASK|MASK_MASK|MOLECULE_MASK);
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->
-        unpack_border(n,first,&buf[m]);
-}
-
-/* ---------------------------------------------------------------------- */
-
-void AtomVecAngleKokkos::unpack_border_vel(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  while (last > nmax) grow(0);
-
-  for (i = first; i < last; i++) {
-    h_x(i,0) = buf[m++];
-    h_x(i,1) = buf[m++];
-    h_x(i,2) = buf[m++];
-    h_tag(i) =  (tagint)  ubuf(buf[m++]).i;
-    h_type(i) = (int) ubuf(buf[m++]).i;
-    h_mask(i) = (int) ubuf(buf[m++]).i;
-    h_molecule(i) = (tagint) ubuf(buf[m++]).i;
-    h_v(i,0) = buf[m++];
-    h_v(i,1) = buf[m++];
-    h_v(i,2) = buf[m++];
-  }
-
-  atomKK->modified(Host,X_MASK|TAG_MASK|TYPE_MASK|MASK_MASK|MOLECULE_MASK|V_MASK);
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->
-        unpack_border(n,first,&buf[m]);
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecAngleKokkos::unpack_border_hybrid(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  for (i = first; i < last; i++)
-    h_molecule(i) = (tagint) ubuf(buf[m++]).i;
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType>
 struct AtomVecAngleKokkos_PackExchangeFunctor {
   typedef DeviceType device_type;
@@ -1245,50 +802,6 @@ int AtomVecAngleKokkos::pack_exchange_kokkos(const int &nsend,DAT::tdual_xfloat_
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecAngleKokkos::pack_exchange(int i, double *buf)
-{
-  int k;
-  int m = 1;
-  buf[m++] = h_x(i,0);
-  buf[m++] = h_x(i,1);
-  buf[m++] = h_x(i,2);
-  buf[m++] = h_v(i,0);
-  buf[m++] = h_v(i,1);
-  buf[m++] = h_v(i,2);
-  buf[m++] = ubuf(h_tag(i)).d;
-  buf[m++] = ubuf(h_type(i)).d;
-  buf[m++] = ubuf(h_mask(i)).d;
-  buf[m++] = ubuf(h_image(i)).d;
-  buf[m++] = ubuf(h_molecule(i)).d;
-
-  buf[m++] = ubuf(h_num_bond(i)).d;
-  for (k = 0; k < h_num_bond(i); k++) {
-    buf[m++] = ubuf(h_bond_type(i,k)).d;
-    buf[m++] = ubuf(h_bond_atom(i,k)).d;
-  }
-  buf[m++] = ubuf(h_num_angle(i)).d;
-  for (k = 0; k < h_num_angle(i); k++) {
-    buf[m++] = ubuf(h_angle_type(i,k)).d;
-    buf[m++] = ubuf(h_angle_atom1(i,k)).d;
-    buf[m++] = ubuf(h_angle_atom2(i,k)).d;
-    buf[m++] = ubuf(h_angle_atom3(i,k)).d;
-  }
-  buf[m++] = ubuf(h_nspecial(i,0)).d;
-  buf[m++] = ubuf(h_nspecial(i,1)).d;
-  buf[m++] = ubuf(h_nspecial(i,2)).d;
-  for (k = 0; k < h_nspecial(i,2); k++)
-    buf[m++] = ubuf(h_special(i,k)).d;
-
-  if (atom->nextra_grow)
-    for (int iextra = 0; iextra < atom->nextra_grow; iextra++)
-      m += modify->fix[atom->extra_grow[iextra]]->pack_exchange(i,&buf[m]);
-
-  buf[0] = m;
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType>
 struct AtomVecAngleKokkos_UnpackExchangeFunctor {
   typedef DeviceType device_type;
@@ -1416,361 +929,6 @@ int AtomVecAngleKokkos::unpack_exchange_kokkos(DAT::tdual_xfloat_2d &k_buf,int n
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecAngleKokkos::unpack_exchange(double *buf)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) grow(0);
-  atomKK->modified(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-           MASK_MASK | IMAGE_MASK | MOLECULE_MASK | BOND_MASK |
-           ANGLE_MASK | SPECIAL_MASK);
-
-  int k;
-  int m = 1;
-  h_x(nlocal,0) = buf[m++];
-  h_x(nlocal,1) = buf[m++];
-  h_x(nlocal,2) = buf[m++];
-  h_v(nlocal,0) = buf[m++];
-  h_v(nlocal,1) = buf[m++];
-  h_v(nlocal,2) = buf[m++];
-  h_tag(nlocal) = (tagint) ubuf(buf[m++]).i;
-  h_type(nlocal) = (int) ubuf(buf[m++]).i;
-  h_mask(nlocal) = (int) ubuf(buf[m++]).i;
-  h_image(nlocal) = (imageint) ubuf(buf[m++]).i;
-  h_molecule(nlocal) = (tagint) ubuf(buf[m++]).i;
-
-  h_num_bond(nlocal) = (int) ubuf(buf[m++]).i;
-  for (k = 0; k < h_num_bond(nlocal); k++) {
-    h_bond_type(nlocal,k) = (int) ubuf(buf[m++]).i;
-    h_bond_atom(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-  }
-  h_num_angle(nlocal) = (int) ubuf(buf[m++]).i;
-  for (k = 0; k < h_num_angle(nlocal); k++) {
-    h_angle_type(nlocal,k) = (int) ubuf(buf[m++]).i;
-    h_angle_atom1(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-    h_angle_atom2(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-    h_angle_atom3(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-  }
-  h_nspecial(nlocal,0) = (int) ubuf(buf[m++]).i;
-  h_nspecial(nlocal,1) = (int) ubuf(buf[m++]).i;
-  h_nspecial(nlocal,2) = (int) ubuf(buf[m++]).i;
-  for (k = 0; k < h_nspecial(nlocal,2); k++)
-   h_special(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-
-  if (atom->nextra_grow)
-    for (int iextra = 0; iextra < atom->nextra_grow; iextra++)
-      m += modify->fix[atom->extra_grow[iextra]]->
-        unpack_exchange(nlocal,&buf[m]);
-
-  atom->nlocal++;
-  return m;
-}
-
-/* ----------------------------------------------------------------------
-   size of restart data for all atoms owned by this proc
-   include extra data stored by fixes
-------------------------------------------------------------------------- */
-
-int AtomVecAngleKokkos::size_restart()
-{
-  int i;
-
-  int nlocal = atom->nlocal;
-  int n = 0;
-  for (i = 0; i < nlocal; i++)
-    n += 14 + 2*h_num_bond(i) + 4*h_num_angle(i);
-
-  if (atom->nextra_restart)
-    for (int iextra = 0; iextra < atom->nextra_restart; iextra++)
-      for (i = 0; i < nlocal; i++)
-        n += modify->fix[atom->extra_restart[iextra]]->size_restart(i);
-
-  return n;
-}
-
-/* ----------------------------------------------------------------------
-   pack atom I's data for restart file including extra quantities
-   xyz must be 1st 3 values, so that read_restart can test on them
-   molecular types may be negative, but write as positive
-------------------------------------------------------------------------- */
-
-int AtomVecAngleKokkos::pack_restart(int i, double *buf)
-{
-  atomKK->sync(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-            MASK_MASK | IMAGE_MASK | MOLECULE_MASK | BOND_MASK |
-            ANGLE_MASK | SPECIAL_MASK);
-
-  int m = 1;
-  buf[m++] = h_x(i,0);
-  buf[m++] = h_x(i,1);
-  buf[m++] = h_x(i,2);
-  buf[m++] = ubuf(h_tag(i)).d;
-  buf[m++] = ubuf(h_type(i)).d;
-  buf[m++] = ubuf(h_mask(i)).d;
-  buf[m++] = ubuf(h_image(i)).d;
-  buf[m++] = h_v(i,0);
-  buf[m++] = h_v(i,1);
-  buf[m++] = h_v(i,2);
-
-  buf[m++] = ubuf(h_molecule(i)).d;
-
-  buf[m++] = ubuf(h_num_bond(i)).d;
-  for (int k = 0; k < h_num_bond(i); k++) {
-    buf[m++] = ubuf(MAX(h_bond_type(i,k),-h_bond_type(i,k))).d;
-    buf[m++] = ubuf(h_bond_atom(i,k)).d;
-  }
-
-  buf[m++] = ubuf(h_num_angle(i)).d;
-  for (int k = 0; k < h_num_angle(i); k++) {
-    buf[m++] = ubuf(MAX(h_angle_type(i,k),-h_angle_type(i,k))).d;
-    buf[m++] = ubuf(h_angle_atom1(i,k)).d;
-    buf[m++] = ubuf(h_angle_atom2(i,k)).d;
-    buf[m++] = ubuf(h_angle_atom3(i,k)).d;
-  }
-
-  if (atom->nextra_restart)
-    for (int iextra = 0; iextra < atom->nextra_restart; iextra++)
-      m += modify->fix[atom->extra_restart[iextra]]->pack_restart(i,&buf[m]);
-
-  buf[0] = m;
-  return m;
-}
-
-/* ----------------------------------------------------------------------
-   unpack data for one atom from restart file including extra quantities
-------------------------------------------------------------------------- */
-
-int AtomVecAngleKokkos::unpack_restart(double *buf)
-{
-  int k;
-
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) {
-    grow(0);
-    if (atom->nextra_store)
-      memory->grow(atom->extra,nmax,atom->nextra_store,"atom:extra");
-  }
-  atomKK->modified(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-                MASK_MASK | IMAGE_MASK | MOLECULE_MASK | BOND_MASK |
-                ANGLE_MASK | SPECIAL_MASK);
-
-  int m = 1;
-  h_x(nlocal,0) = buf[m++];
-  h_x(nlocal,1) = buf[m++];
-  h_x(nlocal,2) = buf[m++];
-  h_tag(nlocal) = (tagint) ubuf(buf[m++]).i;
-  h_type(nlocal) = (int) ubuf(buf[m++]).i;
-  h_mask(nlocal) = (int) ubuf(buf[m++]).i;
-  h_image(nlocal) = (imageint) ubuf(buf[m++]).i;
-  h_v(nlocal,0) = buf[m++];
-  h_v(nlocal,1) = buf[m++];
-  h_v(nlocal,2) = buf[m++];
-
-  h_molecule(nlocal) = (tagint) ubuf(buf[m++]).i;
-
-  h_num_bond(nlocal) = (int) ubuf(buf[m++]).i;
-  for (k = 0; k < h_num_bond(nlocal); k++) {
-    h_bond_type(nlocal,k) = (int) ubuf(buf[m++]).i;
-    h_bond_atom(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-  }
-
-  h_num_angle(nlocal) = (int) ubuf(buf[m++]).i;
-  for (k = 0; k < h_num_angle(nlocal); k++) {
-    h_angle_type(nlocal,k) = (int) ubuf(buf[m++]).i;
-    h_angle_atom1(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-    h_angle_atom2(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-    h_angle_atom3(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-  }
-
-  h_nspecial(nlocal,0) = h_nspecial(nlocal,1) = h_nspecial(nlocal,2) = 0;
-
-  double **extra = atom->extra;
-  if (atom->nextra_store) {
-    int size = static_cast<int> (buf[0]) - m;
-    for (int i = 0; i < size; i++) extra[nlocal][i] = buf[m++];
-  }
-
-  atom->nlocal++;
-  return m;
-}
-
-/* ----------------------------------------------------------------------
-   create one atom of itype at coord
-   set other values to defaults
-------------------------------------------------------------------------- */
-
-void AtomVecAngleKokkos::create_atom(int itype, double *coord)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) {
-    atomKK->modified(Host,ALL_MASK);
-    grow(0);
-  }
-  atomKK->modified(Host,ALL_MASK);
-
-  tag[nlocal] = 0;
-  type[nlocal] = itype;
-  h_x(nlocal,0) = coord[0];
-  h_x(nlocal,1) = coord[1];
-  h_x(nlocal,2) = coord[2];
-  h_mask(nlocal) = 1;
-  h_image(nlocal) = ((imageint) IMGMAX << IMG2BITS) |
-    ((imageint) IMGMAX << IMGBITS) | IMGMAX;
-  h_v(nlocal,0) = 0.0;
-  h_v(nlocal,1) = 0.0;
-  h_v(nlocal,2) = 0.0;
-
-  h_molecule(nlocal) = 0;
-  h_num_bond(nlocal) = 0;
-  h_num_angle(nlocal) = 0;
-  h_nspecial(nlocal,0) = h_nspecial(nlocal,1) = h_nspecial(nlocal,2) = 0;
-
-  atom->nlocal++;
-}
-
-/* ----------------------------------------------------------------------
-   unpack one line from Atoms section of data file
-   initialize other atom quantities
-------------------------------------------------------------------------- */
-
-void AtomVecAngleKokkos::data_atom(double *coord, imageint imagetmp,
-                                  const std::vector<std::string> &values, std::string &extract)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) grow(0);
-  atomKK->modified(Host,ALL_MASK);
-
-  h_tag(nlocal) = utils::inumeric(FLERR,values[0],true,lmp);
-  h_molecule(nlocal) = utils::inumeric(FLERR,values[1],true,lmp);
-  h_type(nlocal) = utils::inumeric(FLERR,values[2],true,lmp);
-  extract = values[2];
-  if (h_type(nlocal) <= 0 || h_type(nlocal) > atom->ntypes)
-    error->one(FLERR,"Invalid atom type in Atoms section of data file");
-
-  h_x(nlocal,0) = coord[0];
-  h_x(nlocal,1) = coord[1];
-  h_x(nlocal,2) = coord[2];
-
-  h_image(nlocal) = imagetmp;
-
-  h_mask(nlocal) = 1;
-  h_v(nlocal,0) = 0.0;
-  h_v(nlocal,1) = 0.0;
-  h_v(nlocal,2) = 0.0;
-  h_num_bond(nlocal) = 0;
-  h_num_angle(nlocal) = 0;
-
-  atom->nlocal++;
-}
-
-/* ----------------------------------------------------------------------
-   unpack hybrid quantities from one line in Atoms section of data file
-   initialize other atom quantities for this sub-style
-------------------------------------------------------------------------- */
-
-int AtomVecAngleKokkos::data_atom_hybrid(int nlocal, const std::vector<std::string> &values,
-                                         int offset)
-{
-  h_molecule(nlocal) = utils::inumeric(FLERR,values[offset],true,lmp);
-  h_num_bond(nlocal) = 0;
-  h_num_angle(nlocal) = 0;
-  return 1;
-}
-
-/* ----------------------------------------------------------------------
-   pack atom info for data file including 3 image flags
-------------------------------------------------------------------------- */
-
-void AtomVecAngleKokkos::pack_data(double **buf)
-{
-  int nlocal = atom->nlocal;
-  for (int i = 0; i < nlocal; i++) {
-    buf[i][0] = h_tag(i);
-    buf[i][1] = h_molecule(i);
-    buf[i][2] = h_type(i);
-    buf[i][3] = h_x(i,0);
-    buf[i][4] = h_x(i,1);
-    buf[i][5] = h_x(i,2);
-    buf[i][6] = (h_image[i] & IMGMASK) - IMGMAX;
-    buf[i][7] = (h_image[i] >> IMGBITS & IMGMASK) - IMGMAX;
-    buf[i][8] = (h_image[i] >> IMG2BITS) - IMGMAX;
-  }
-}
-
-/* ----------------------------------------------------------------------
-   pack hybrid atom info for data file
-------------------------------------------------------------------------- */
-
-int AtomVecAngleKokkos::pack_data_hybrid(int i, double *buf)
-{
-  buf[0] = h_molecule(i);
-  return 1;
-}
-
-/* ----------------------------------------------------------------------
-   write atom info to data file including 3 image flags
-------------------------------------------------------------------------- */
-
-void AtomVecAngleKokkos::write_data(FILE *fp, int n, double **buf)
-{
-  for (int i = 0; i < n; i++)
-    fprintf(fp,"%d %d %d %-1.16e %-1.16e %-1.16e %d %d %d\n",
-            (int) buf[i][0],(int) buf[i][1], (int) buf[i][2],
-            buf[i][3],buf[i][4],buf[i][5],
-            (int) buf[i][6],(int) buf[i][7],(int) buf[i][8]);
-}
-
-/* ----------------------------------------------------------------------
-   write hybrid atom info to data file
-------------------------------------------------------------------------- */
-
-int AtomVecAngleKokkos::write_data_hybrid(FILE *fp, double *buf)
-{
-  fprintf(fp," " TAGINT_FORMAT, (tagint) (buf[0]));
-  return 1;
-}
-
-/* ----------------------------------------------------------------------
-   return # of bytes of allocated memory
-------------------------------------------------------------------------- */
-
-double AtomVecAngleKokkos::memory_usage()
-{
-  double bytes = 0;
-
-  if (atom->memcheck("tag")) bytes += memory->usage(tag,nmax);
-  if (atom->memcheck("type")) bytes += memory->usage(type,nmax);
-  if (atom->memcheck("mask")) bytes += memory->usage(mask,nmax);
-  if (atom->memcheck("image")) bytes += memory->usage(image,nmax);
-  if (atom->memcheck("x")) bytes += memory->usage(x,nmax,3);
-  if (atom->memcheck("v")) bytes += memory->usage(v,nmax,3);
-  if (atom->memcheck("f")) bytes += memory->usage(f,nmax*commKK->nthreads,3);
-
-  if (atom->memcheck("molecule")) bytes += memory->usage(molecule,nmax);
-  if (atom->memcheck("nspecial")) bytes += memory->usage(nspecial,nmax,3);
-  if (atom->memcheck("special"))
-    bytes += memory->usage(special,nmax,atom->maxspecial);
-
-  if (atom->memcheck("num_bond")) bytes += memory->usage(num_bond,nmax);
-  if (atom->memcheck("bond_type"))
-    bytes += memory->usage(bond_type,nmax,atom->bond_per_atom);
-  if (atom->memcheck("bond_atom"))
-    bytes += memory->usage(bond_atom,nmax,atom->bond_per_atom);
-
-  if (atom->memcheck("num_angle")) bytes += memory->usage(num_angle,nmax);
-  if (atom->memcheck("angle_type"))
-    bytes += memory->usage(angle_type,nmax,atom->angle_per_atom);
-  if (atom->memcheck("angle_atom1"))
-    bytes += memory->usage(angle_atom1,nmax,atom->angle_per_atom);
-  if (atom->memcheck("angle_atom2"))
-    bytes += memory->usage(angle_atom2,nmax,atom->angle_per_atom);
-  if (atom->memcheck("angle_atom3"))
-    bytes += memory->usage(angle_atom3,nmax,atom->angle_per_atom);
-
-  return bytes;
-}
-
-/* ---------------------------------------------------------------------- */
-
 void AtomVecAngleKokkos::sync(ExecutionSpace space, unsigned int mask)
 {
   if (space == Device) {
@@ -1826,6 +984,8 @@ void AtomVecAngleKokkos::sync(ExecutionSpace space, unsigned int mask)
   }
 }
 
+/* ---------------------------------------------------------------------- */
+
 void AtomVecAngleKokkos::sync_overlapping_device(ExecutionSpace space, unsigned int mask)
 {
   if (space == Device) {
@@ -1973,4 +1133,3 @@ void AtomVecAngleKokkos::modified(ExecutionSpace space, unsigned int mask)
     }
   }
 }
-

src/KOKKOS/atom_vec_angle_kokkos.h

@@ -24,41 +24,15 @@ AtomStyle(angle/kk/host,AtomVecAngleKokkos);
 #define LMP_ATOM_VEC_ANGLE_KOKKOS_H
 
 #include "atom_vec_kokkos.h"
+#include "atom_vec_angle.h"
 
 namespace LAMMPS_NS {
 
-class AtomVecAngleKokkos : public AtomVecKokkos {
+class AtomVecAngleKokkos : public AtomVecKokkos, public AtomVecAngle {
  public:
   AtomVecAngleKokkos(class LAMMPS *);
 
   void grow(int) override;
-  void copy(int, int, int) override;
-  int pack_comm(int, int *, double *, int, int *) override;
-  int pack_comm_vel(int, int *, double *, int, int *) override;
-  void unpack_comm(int, int, double *) override;
-  void unpack_comm_vel(int, int, double *) override;
-  int pack_reverse(int, int, double *) override;
-  void unpack_reverse(int, int *, double *) override;
-  int pack_border(int, int *, double *, int, int *) override;
-  int pack_border_vel(int, int *, double *, int, int *) override;
-  int pack_border_hybrid(int, int *, double *) override;
-  void unpack_border(int, int, double *) override;
-  void unpack_border_vel(int, int, double *) override;
-  int unpack_border_hybrid(int, int, double *) override;
-  int pack_exchange(int, double *) override;
-  int unpack_exchange(double *) override;
-  int size_restart() override;
-  int pack_restart(int, double *) override;
-  int unpack_restart(double *) override;
-  void create_atom(int, double *) override;
-  void data_atom(double *, imageint, const std::vector<std::string> &, std::string &) override;
-  int data_atom_hybrid(int, const std::vector<std::string> &, int) override;
-  void pack_data(double **) override;
-  int pack_data_hybrid(int, double *) override;
-  void write_data(FILE *, int, double **) override;
-  int write_data_hybrid(FILE *, double *) override;
-  double memory_usage() override;
-
   void grow_pointers() override;
   int pack_comm_kokkos(const int &n, const DAT::tdual_int_2d &k_sendlist,
                        const int & iswap,
@@ -142,9 +116,6 @@ class AtomVecAngleKokkos : public AtomVecKokkos {
   HAT::t_int_1d h_num_angle;
   HAT::t_int_2d h_angle_type;
   HAT::t_tagint_2d h_angle_atom1,h_angle_atom2,h_angle_atom3;
-
-  DAT::tdual_int_1d k_count;
-
 };
 
 }

src/KOKKOS/atom_vec_atomic_kokkos.cpp

@@ -27,23 +27,10 @@ using namespace LAMMPS_NS;
 
 /* ---------------------------------------------------------------------- */
 
-AtomVecAtomicKokkos::AtomVecAtomicKokkos(LAMMPS *lmp) : AtomVecKokkos(lmp)
+AtomVecAtomicKokkos::AtomVecAtomicKokkos(LAMMPS *lmp) : AtomVec(lmp),
+AtomVecKokkos(lmp), AtomVecAtomic(lmp)
 {
-  molecular = Atom::ATOMIC;
-  mass_type = PER_TYPE;
 
-  comm_x_only = comm_f_only = 1;
-  size_forward = 3;
-  size_reverse = 3;
-  size_border = 6;
-  size_velocity = 3;
-  size_data_atom = 5;
-  size_data_vel = 4;
-  xcol_data = 3;
-
-  k_count = DAT::tdual_int_1d("atom::k_count",1);
-  atomKK = (AtomKokkos *) atom;
-  commKK = (CommKokkos *) comm;
 }
 
 /* ----------------------------------------------------------------------
@@ -113,28 +100,6 @@ void AtomVecAtomicKokkos::grow_pointers()
   h_f = atomKK->k_f.h_view;
 }
 
-/* ----------------------------------------------------------------------
-   copy atom I info to atom J
-------------------------------------------------------------------------- */
-
-void AtomVecAtomicKokkos::copy(int i, int j, int delflag)
-{
-  h_tag[j] = h_tag[i];
-  h_type[j] = h_type[i];
-  mask[j] = mask[i];
-  h_image[j] = h_image[i];
-  h_x(j,0) = h_x(i,0);
-  h_x(j,1) = h_x(i,1);
-  h_x(j,2) = h_x(i,2);
-  h_v(j,0) = h_v(i,0);
-  h_v(j,1) = h_v(i,1);
-  h_v(j,2) = h_v(i,2);
-
-  if (atom->nextra_grow)
-    for (int iextra = 0; iextra < atom->nextra_grow; iextra++)
-      modify->fix[atom->extra_grow[iextra]]->copy_arrays(i,j,delflag);
-}
-
 /* ---------------------------------------------------------------------- */
 
 template<class DeviceType,int PBC_FLAG>
@@ -232,130 +197,6 @@ int AtomVecAtomicKokkos::pack_border_kokkos(int n, DAT::tdual_int_2d k_sendlist,
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecAtomicKokkos::pack_border(int n, int *list, double *buf,
-                               int pbc_flag, int *pbc)
-{
-  int i,j,m;
-  double dx,dy,dz;
-
-  m = 0;
-  if (pbc_flag == 0) {
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0);
-      buf[m++] = h_x(j,1);
-      buf[m++] = h_x(j,2);
-      buf[m++] = ubuf(h_tag(j)).d;
-      buf[m++] = ubuf(h_type(j)).d;
-      buf[m++] = ubuf(h_mask(j)).d;
-    }
-  } else {
-    if (domain->triclinic == 0) {
-      dx = pbc[0]*domain->xprd;
-      dy = pbc[1]*domain->yprd;
-      dz = pbc[2]*domain->zprd;
-    } else {
-      dx = pbc[0];
-      dy = pbc[1];
-      dz = pbc[2];
-    }
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0) + dx;
-      buf[m++] = h_x(j,1) + dy;
-      buf[m++] = h_x(j,2) + dz;
-      buf[m++] = ubuf(h_tag(j)).d;
-      buf[m++] = ubuf(h_type(j)).d;
-      buf[m++] = ubuf(h_mask(j)).d;
-    }
-  }
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->pack_border(n,list,&buf[m]);
-
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecAtomicKokkos::pack_border_vel(int n, int *list, double *buf,
-                                   int pbc_flag, int *pbc)
-{
-  int i,j,m;
-  double dx,dy,dz,dvx,dvy,dvz;
-
-  m = 0;
-  if (pbc_flag == 0) {
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0);
-      buf[m++] = h_x(j,1);
-      buf[m++] = h_x(j,2);
-      buf[m++] = ubuf(h_tag(j)).d;
-      buf[m++] = ubuf(h_type(j)).d;
-      buf[m++] = ubuf(h_mask(j)).d;
-      buf[m++] = h_v(j,0);
-      buf[m++] = h_v(j,1);
-      buf[m++] = h_v(j,2);
-    }
-  } else {
-    if (domain->triclinic == 0) {
-      dx = pbc[0]*domain->xprd;
-      dy = pbc[1]*domain->yprd;
-      dz = pbc[2]*domain->zprd;
-    } else {
-      dx = pbc[0];
-      dy = pbc[1];
-      dz = pbc[2];
-    }
-    if (!deform_vremap) {
-      for (i = 0; i < n; i++) {
-        j = list[i];
-        buf[m++] = h_x(j,0) + dx;
-        buf[m++] = h_x(j,1) + dy;
-        buf[m++] = h_x(j,2) + dz;
-        buf[m++] = ubuf(h_tag(j)).d;
-        buf[m++] = ubuf(h_type(j)).d;
-        buf[m++] = ubuf(h_mask(j)).d;
-        buf[m++] = h_v(j,0);
-        buf[m++] = h_v(j,1);
-        buf[m++] = h_v(j,2);
-      }
-    } else {
-      dvx = pbc[0]*h_rate[0] + pbc[5]*h_rate[5] + pbc[4]*h_rate[4];
-      dvy = pbc[1]*h_rate[1] + pbc[3]*h_rate[3];
-      dvz = pbc[2]*h_rate[2];
-      for (i = 0; i < n; i++) {
-        j = list[i];
-        buf[m++] = h_x(j,0) + dx;
-        buf[m++] = h_x(j,1) + dy;
-        buf[m++] = h_x(j,2) + dz;
-        buf[m++] = ubuf(h_tag(j)).d;
-        buf[m++] = ubuf(h_type(j)).d;
-        buf[m++] = ubuf(h_mask(j)).d;
-        if (mask[i] & deform_groupbit) {
-          buf[m++] = h_v(j,0) + dvx;
-          buf[m++] = h_v(j,1) + dvy;
-          buf[m++] = h_v(j,2) + dvz;
-        } else {
-          buf[m++] = h_v(j,0);
-          buf[m++] = h_v(j,1);
-          buf[m++] = h_v(j,2);
-        }
-      }
-    }
-  }
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->pack_border(n,list,&buf[m]);
-
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType>
 struct AtomVecAtomicKokkos_UnpackBorder {
   typedef DeviceType device_type;
@@ -408,63 +249,6 @@ void AtomVecAtomicKokkos::unpack_border_kokkos(const int &n, const int &first,
 
 /* ---------------------------------------------------------------------- */
 
-void AtomVecAtomicKokkos::unpack_border(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  while (last > nmax) grow(0);
-
-  for (i = first; i < last; i++) {
-    h_x(i,0) = buf[m++];
-    h_x(i,1) = buf[m++];
-    h_x(i,2) = buf[m++];
-    h_tag(i) =  (tagint)  ubuf(buf[m++]).i;
-    h_type(i) = (int) ubuf(buf[m++]).i;
-    h_mask(i) = (int) ubuf(buf[m++]).i;
-  }
-
-  atomKK->modified(Host,X_MASK|TAG_MASK|TYPE_MASK|MASK_MASK);
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->
-        unpack_border(n,first,&buf[m]);
-}
-
-/* ---------------------------------------------------------------------- */
-
-void AtomVecAtomicKokkos::unpack_border_vel(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  while (last > nmax) grow(0);
-
-  for (i = first; i < last; i++) {
-    h_x(i,0) = buf[m++];
-    h_x(i,1) = buf[m++];
-    h_x(i,2) = buf[m++];
-    h_tag(i) =  (tagint)  ubuf(buf[m++]).i;
-    h_type(i) = (int) ubuf(buf[m++]).i;
-    h_mask(i) = (int) ubuf(buf[m++]).i;
-    h_v(i,0) = buf[m++];
-    h_v(i,1) = buf[m++];
-    h_v(i,2) = buf[m++];
-  }
-
-  atomKK->modified(Host,X_MASK|TAG_MASK|TYPE_MASK|MASK_MASK|V_MASK);
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->
-        unpack_border(n,first,&buf[m]);
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType>
 struct AtomVecAtomicKokkos_PackExchangeFunctor {
   typedef DeviceType device_type;
@@ -568,30 +352,6 @@ int AtomVecAtomicKokkos::pack_exchange_kokkos(const int &nsend,DAT::tdual_xfloat
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecAtomicKokkos::pack_exchange(int i, double *buf)
-{
-  int m = 1;
-  buf[m++] = h_x(i,0);
-  buf[m++] = h_x(i,1);
-  buf[m++] = h_x(i,2);
-  buf[m++] = h_v(i,0);
-  buf[m++] = h_v(i,1);
-  buf[m++] = h_v(i,2);
-  buf[m++] = ubuf(h_tag(i)).d;
-  buf[m++] = ubuf(h_type(i)).d;
-  buf[m++] = ubuf(h_mask(i)).d;
-  buf[m++] = ubuf(h_image(i)).d;
-
-  if (atom->nextra_grow)
-    for (int iextra = 0; iextra < atom->nextra_grow; iextra++)
-      m += modify->fix[atom->extra_grow[iextra]]->pack_exchange(i,&buf[m]);
-
-  buf[0] = m;
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType>
 struct AtomVecAtomicKokkos_UnpackExchangeFunctor {
   typedef DeviceType device_type;
@@ -671,238 +431,6 @@ int AtomVecAtomicKokkos::unpack_exchange_kokkos(DAT::tdual_xfloat_2d &k_buf,int
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecAtomicKokkos::unpack_exchange(double *buf)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) grow(0);
-  atomKK->modified(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-           MASK_MASK | IMAGE_MASK);
-
-  int m = 1;
-  h_x(nlocal,0) = buf[m++];
-  h_x(nlocal,1) = buf[m++];
-  h_x(nlocal,2) = buf[m++];
-  h_v(nlocal,0) = buf[m++];
-  h_v(nlocal,1) = buf[m++];
-  h_v(nlocal,2) = buf[m++];
-  h_tag(nlocal) = (tagint) ubuf(buf[m++]).i;
-  h_type(nlocal) = (int) ubuf(buf[m++]).i;
-  h_mask(nlocal) = (int) ubuf(buf[m++]).i;
-  h_image(nlocal) = (imageint) ubuf(buf[m++]).i;
-
-  if (atom->nextra_grow)
-    for (int iextra = 0; iextra < atom->nextra_grow; iextra++)
-      m += modify->fix[atom->extra_grow[iextra]]->
-        unpack_exchange(nlocal,&buf[m]);
-
-  atom->nlocal++;
-  return m;
-}
-
-/* ----------------------------------------------------------------------
-   size of restart data for all atoms owned by this proc
-   include extra data stored by fixes
-------------------------------------------------------------------------- */
-
-int AtomVecAtomicKokkos::size_restart()
-{
-  int i;
-
-  int nlocal = atom->nlocal;
-  int n = 11 * nlocal;
-
-  if (atom->nextra_restart)
-    for (int iextra = 0; iextra < atom->nextra_restart; iextra++)
-      for (i = 0; i < nlocal; i++)
-        n += modify->fix[atom->extra_restart[iextra]]->size_restart(i);
-
-  return n;
-}
-
-/* ----------------------------------------------------------------------
-   pack atom I's data for restart file including extra quantities
-   xyz must be 1st 3 values, so that read_restart can test on them
-   molecular types may be negative, but write as positive
-------------------------------------------------------------------------- */
-
-int AtomVecAtomicKokkos::pack_restart(int i, double *buf)
-{
-  atomKK->sync(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-            MASK_MASK | IMAGE_MASK );
-
-  int m = 1;
-  buf[m++] = h_x(i,0);
-  buf[m++] = h_x(i,1);
-  buf[m++] = h_x(i,2);
-  buf[m++] = ubuf(h_tag(i)).d;
-  buf[m++] = ubuf(h_type(i)).d;
-  buf[m++] = ubuf(h_mask(i)).d;
-  buf[m++] = ubuf(h_image(i)).d;
-  buf[m++] = h_v(i,0);
-  buf[m++] = h_v(i,1);
-  buf[m++] = h_v(i,2);
-
-  if (atom->nextra_restart)
-    for (int iextra = 0; iextra < atom->nextra_restart; iextra++)
-      m += modify->fix[atom->extra_restart[iextra]]->pack_restart(i,&buf[m]);
-
-  buf[0] = m;
-  return m;
-}
-
-/* ----------------------------------------------------------------------
-   unpack data for one atom from restart file including extra quantities
-------------------------------------------------------------------------- */
-
-int AtomVecAtomicKokkos::unpack_restart(double *buf)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) {
-    grow(0);
-    if (atom->nextra_store)
-      memory->grow(atom->extra,nmax,atom->nextra_store,"atom:extra");
-  }
-  atomKK->modified(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-                MASK_MASK | IMAGE_MASK );
-
-  int m = 1;
-  h_x(nlocal,0) = buf[m++];
-  h_x(nlocal,1) = buf[m++];
-  h_x(nlocal,2) = buf[m++];
-  h_tag(nlocal) = (tagint) ubuf(buf[m++]).i;
-  h_type(nlocal) = (int) ubuf(buf[m++]).i;
-  h_mask(nlocal) = (int) ubuf(buf[m++]).i;
-  h_image(nlocal) = (imageint) ubuf(buf[m++]).i;
-  h_v(nlocal,0) = buf[m++];
-  h_v(nlocal,1) = buf[m++];
-  h_v(nlocal,2) = buf[m++];
-
-  double **extra = atom->extra;
-  if (atom->nextra_store) {
-    int size = static_cast<int> (buf[0]) - m;
-    for (int i = 0; i < size; i++) extra[nlocal][i] = buf[m++];
-  }
-
-  atom->nlocal++;
-  return m;
-}
-
-/* ----------------------------------------------------------------------
-   create one atom of itype at coord
-   set other values to defaults
-------------------------------------------------------------------------- */
-
-void AtomVecAtomicKokkos::create_atom(int itype, double *coord)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) {
-    //if(nlocal>2) printf("typeA: %i %i\n",type[0],type[1]);
-    atomKK->modified(Host,ALL_MASK);
-    grow(0);
-    //if(nlocal>2) printf("typeB: %i %i\n",type[0],type[1]);
-  }
-  atomKK->modified(Host,ALL_MASK);
-
-  tag[nlocal] = 0;
-  type[nlocal] = itype;
-  h_x(nlocal,0) = coord[0];
-  h_x(nlocal,1) = coord[1];
-  h_x(nlocal,2) = coord[2];
-  h_mask[nlocal] = 1;
-  h_image[nlocal] = ((tagint) IMGMAX << IMG2BITS) |
-    ((tagint) IMGMAX << IMGBITS) | IMGMAX;
-  h_v(nlocal,0) = 0.0;
-  h_v(nlocal,1) = 0.0;
-  h_v(nlocal,2) = 0.0;
-
-  atom->nlocal++;
-}
-
-/* ----------------------------------------------------------------------
-   unpack one line from Atoms section of data file
-   initialize other atom quantities
-------------------------------------------------------------------------- */
-
-void AtomVecAtomicKokkos::data_atom(double *coord, imageint imagetmp,
-                                    const std::vector<std::string> &values, std::string &extract)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) grow(0);
-
-  h_tag[nlocal] = utils::inumeric(FLERR,values[0],true,lmp);
-  h_type[nlocal] = utils::inumeric(FLERR,values[1],true,lmp);
-  extract = values[1];
-  if (type[nlocal] <= 0 || type[nlocal] > atom->ntypes)
-    error->one(FLERR,"Invalid atom type in Atoms section of data file");
-
-  h_x(nlocal,0) = coord[0];
-  h_x(nlocal,1) = coord[1];
-  h_x(nlocal,2) = coord[2];
-
-  h_image[nlocal] = imagetmp;
-
-  h_mask[nlocal] = 1;
-  h_v(nlocal,0) = 0.0;
-  h_v(nlocal,1) = 0.0;
-  h_v(nlocal,2) = 0.0;
-
-  atomKK->modified(Host,ALL_MASK);
-
-  atom->nlocal++;
-}
-
-/* ----------------------------------------------------------------------
-   pack atom info for data file including 3 image flags
-------------------------------------------------------------------------- */
-
-void AtomVecAtomicKokkos::pack_data(double **buf)
-{
-  int nlocal = atom->nlocal;
-  for (int i = 0; i < nlocal; i++) {
-    buf[i][0] = h_tag[i];
-    buf[i][1] = h_type[i];
-    buf[i][2] = h_x(i,0);
-    buf[i][3] = h_x(i,1);
-    buf[i][4] = h_x(i,2);
-    buf[i][5] = (h_image[i] & IMGMASK) - IMGMAX;
-    buf[i][6] = (h_image[i] >> IMGBITS & IMGMASK) - IMGMAX;
-    buf[i][7] = (h_image[i] >> IMG2BITS) - IMGMAX;
-  }
-}
-
-/* ----------------------------------------------------------------------
-   write atom info to data file including 3 image flags
-------------------------------------------------------------------------- */
-
-void AtomVecAtomicKokkos::write_data(FILE *fp, int n, double **buf)
-{
-  for (int i = 0; i < n; i++)
-    fprintf(fp,"%d %d %-1.16e %-1.16e %-1.16e %d %d %d\n",
-            (int) buf[i][0],(int) buf[i][1],buf[i][2],buf[i][3],buf[i][4],
-            (int) buf[i][5],(int) buf[i][6],(int) buf[i][7]);
-}
-
-/* ----------------------------------------------------------------------
-   return # of bytes of allocated memory
-------------------------------------------------------------------------- */
-
-double AtomVecAtomicKokkos::memory_usage()
-{
-  double bytes = 0;
-
-  if (atom->memcheck("tag")) bytes += memory->usage(tag,nmax);
-  if (atom->memcheck("type")) bytes += memory->usage(type,nmax);
-  if (atom->memcheck("mask")) bytes += memory->usage(mask,nmax);
-  if (atom->memcheck("image")) bytes += memory->usage(image,nmax);
-  if (atom->memcheck("x")) bytes += memory->usage(x,nmax,3);
-  if (atom->memcheck("v")) bytes += memory->usage(v,nmax,3);
-  if (atom->memcheck("f")) bytes += memory->usage(f,nmax*commKK->nthreads,3);
-
-  return bytes;
-}
-
-/* ---------------------------------------------------------------------- */
-
 void AtomVecAtomicKokkos::sync(ExecutionSpace space, unsigned int mask)
 {
   if (space == Device) {
@@ -983,4 +511,3 @@ void AtomVecAtomicKokkos::modified(ExecutionSpace space, unsigned int mask)
     if (mask & IMAGE_MASK) atomKK->k_image.modify<LMPHostType>();
   }
 }
-

src/KOKKOS/atom_vec_atomic_kokkos.h

@@ -24,31 +24,16 @@ AtomStyle(atomic/kk/host,AtomVecAtomicKokkos);
 #define LMP_ATOM_VEC_ATOMIC_KOKKOS_H
 
 #include "atom_vec_kokkos.h"
+#include "atom_vec_atomic.h"
 #include "kokkos_type.h"
 
 namespace LAMMPS_NS {
 
-class AtomVecAtomicKokkos : public AtomVecKokkos {
+class AtomVecAtomicKokkos : public AtomVecKokkos, public AtomVecAtomic {
  public:
   AtomVecAtomicKokkos(class LAMMPS *);
 
   void grow(int) override;
-  void copy(int, int, int) override;
-  int pack_border(int, int *, double *, int, int *) override;
-  int pack_border_vel(int, int *, double *, int, int *) override;
-  void unpack_border(int, int, double *) override;
-  void unpack_border_vel(int, int, double *) override;
-  int pack_exchange(int, double *) override;
-  int unpack_exchange(double *) override;
-  int size_restart() override;
-  int pack_restart(int, double *) override;
-  int unpack_restart(double *) override;
-  void create_atom(int, double *) override;
-  void data_atom(double *, imageint, const std::vector<std::string> &, std::string &) override;
-  void pack_data(double **) override;
-  void write_data(FILE *, int, double **) override;
-  double memory_usage() override;
-
   void grow_pointers() override;
   int pack_border_kokkos(int n, DAT::tdual_int_2d k_sendlist,
                          DAT::tdual_xfloat_2d buf,int iswap,
@@ -70,11 +55,6 @@ class AtomVecAtomicKokkos : public AtomVecKokkos {
   void sync_overlapping_device(ExecutionSpace space, unsigned int mask) override;
 
  protected:
-  tagint *tag;
-  imageint *image;
-  int *type,*mask;
-  double **x,**v,**f;
-
   DAT::t_tagint_1d d_tag;
   HAT::t_tagint_1d h_tag;
   DAT::t_imageint_1d d_image;
@@ -85,8 +65,6 @@ class AtomVecAtomicKokkos : public AtomVecKokkos {
   DAT::t_x_array d_x;
   DAT::t_v_array d_v;
   DAT::t_f_array d_f;
-
-  DAT::tdual_int_1d k_count;
 };
 
 }

src/KOKKOS/atom_vec_bond_kokkos.cpp

@@ -27,26 +27,10 @@ using namespace LAMMPS_NS;
 
 /* ---------------------------------------------------------------------- */
 
-AtomVecBondKokkos::AtomVecBondKokkos(LAMMPS *lmp) : AtomVecKokkos(lmp)
+AtomVecBondKokkos::AtomVecBondKokkos(LAMMPS *lmp) : AtomVec(lmp),
+AtomVecKokkos(lmp), AtomVecBond(lmp)
 {
-  molecular = Atom::MOLECULAR;
-  bonds_allow = 1;
-  mass_type = PER_TYPE;
 
-  comm_x_only = comm_f_only = 1;
-  size_forward = 3;
-  size_reverse = 3;
-  size_border = 7;
-  size_velocity = 3;
-  size_data_atom = 6;
-  size_data_vel = 4;
-  xcol_data = 4;
-
-  atom->molecule_flag = 1;
-
-  k_count = DAT::tdual_int_1d("atom::k_count",1);
-  atomKK = (AtomKokkos *) atom;
-  commKK = (CommKokkos *) comm;
 }
 
 /* ----------------------------------------------------------------------
@@ -142,43 +126,6 @@ void AtomVecBondKokkos::grow_pointers()
   h_bond_atom = atomKK->k_bond_atom.h_view;
 }
 
-/* ----------------------------------------------------------------------
-   copy atom I info to atom J
-------------------------------------------------------------------------- */
-
-void AtomVecBondKokkos::copy(int i, int j, int delflag)
-{
-  int k;
-
-  h_tag[j] = h_tag[i];
-  h_type[j] = h_type[i];
-  mask[j] = mask[i];
-  h_image[j] = h_image[i];
-  h_x(j,0) = h_x(i,0);
-  h_x(j,1) = h_x(i,1);
-  h_x(j,2) = h_x(i,2);
-  h_v(j,0) = h_v(i,0);
-  h_v(j,1) = h_v(i,1);
-  h_v(j,2) = h_v(i,2);
-
-  h_molecule(j) = h_molecule(i);
-
-  h_num_bond(j) = h_num_bond(i);
-  for (k = 0; k < h_num_bond(j); k++) {
-    h_bond_type(j,k) = h_bond_type(i,k);
-    h_bond_atom(j,k) = h_bond_atom(i,k);
-  }
-
-  h_nspecial(j,0) = h_nspecial(i,0);
-  h_nspecial(j,1) = h_nspecial(i,1);
-  h_nspecial(j,2) = h_nspecial(i,2);
-  for (k = 0; k < h_nspecial(j,2); k++) h_special(j,k) = h_special(i,k);
-
-  if (atom->nextra_grow)
-    for (int iextra = 0; iextra < atom->nextra_grow; iextra++)
-      modify->fix[atom->extra_grow[iextra]]->copy_arrays(i,j,delflag);
-}
-
 /* ---------------------------------------------------------------------- */
 
 template<class DeviceType,int PBC_FLAG>
@@ -282,149 +229,6 @@ int AtomVecBondKokkos::pack_border_kokkos(int n, DAT::tdual_int_2d k_sendlist,
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecBondKokkos::pack_border(int n, int *list, double *buf,
-                                   int pbc_flag, int *pbc)
-{
-  int i,j,m;
-  double dx,dy,dz;
-
-  m = 0;
-  if (pbc_flag == 0) {
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0);
-      buf[m++] = h_x(j,1);
-      buf[m++] = h_x(j,2);
-      buf[m++] = ubuf(h_tag(j)).d;
-      buf[m++] = ubuf(h_type(j)).d;
-      buf[m++] = ubuf(h_mask(j)).d;
-      buf[m++] = ubuf(h_molecule(j)).d;
-    }
-  } else {
-    if (domain->triclinic == 0) {
-      dx = pbc[0]*domain->xprd;
-      dy = pbc[1]*domain->yprd;
-      dz = pbc[2]*domain->zprd;
-    } else {
-      dx = pbc[0];
-      dy = pbc[1];
-      dz = pbc[2];
-    }
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0) + dx;
-      buf[m++] = h_x(j,1) + dy;
-      buf[m++] = h_x(j,2) + dz;
-      buf[m++] = ubuf(h_tag(j)).d;
-      buf[m++] = ubuf(h_type(j)).d;
-      buf[m++] = ubuf(h_mask(j)).d;
-      buf[m++] = ubuf(h_molecule(j)).d;
-    }
-  }
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->pack_border(n,list,&buf[m]);
-
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecBondKokkos::pack_border_vel(int n, int *list, double *buf,
-                                       int pbc_flag, int *pbc)
-{
-  int i,j,m;
-  double dx,dy,dz,dvx,dvy,dvz;
-
-  m = 0;
-  if (pbc_flag == 0) {
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0);
-      buf[m++] = h_x(j,1);
-      buf[m++] = h_x(j,2);
-      buf[m++] = ubuf(h_tag(j)).d;
-      buf[m++] = ubuf(h_type(j)).d;
-      buf[m++] = ubuf(h_mask(j)).d;
-      buf[m++] = ubuf(h_molecule(j)).d;
-      buf[m++] = h_v(j,0);
-      buf[m++] = h_v(j,1);
-      buf[m++] = h_v(j,2);
-    }
-  } else {
-    if (domain->triclinic == 0) {
-      dx = pbc[0]*domain->xprd;
-      dy = pbc[1]*domain->yprd;
-      dz = pbc[2]*domain->zprd;
-    } else {
-      dx = pbc[0];
-      dy = pbc[1];
-      dz = pbc[2];
-    }
-    if (!deform_vremap) {
-      for (i = 0; i < n; i++) {
-        j = list[i];
-        buf[m++] = h_x(j,0) + dx;
-        buf[m++] = h_x(j,1) + dy;
-        buf[m++] = h_x(j,2) + dz;
-        buf[m++] = ubuf(h_tag(j)).d;
-        buf[m++] = ubuf(h_type(j)).d;
-        buf[m++] = ubuf(h_mask(j)).d;
-        buf[m++] = ubuf(h_molecule(j)).d;
-        buf[m++] = h_v(j,0);
-        buf[m++] = h_v(j,1);
-        buf[m++] = h_v(j,2);
-      }
-    } else {
-      dvx = pbc[0]*h_rate[0] + pbc[5]*h_rate[5] + pbc[4]*h_rate[4];
-      dvy = pbc[1]*h_rate[1] + pbc[3]*h_rate[3];
-      dvz = pbc[2]*h_rate[2];
-      for (i = 0; i < n; i++) {
-        j = list[i];
-        buf[m++] = h_x(j,0) + dx;
-        buf[m++] = h_x(j,1) + dy;
-        buf[m++] = h_x(j,2) + dz;
-        buf[m++] = ubuf(h_tag(j)).d;
-        buf[m++] = ubuf(h_type(j)).d;
-        buf[m++] = ubuf(h_mask(j)).d;
-        buf[m++] = ubuf(h_molecule(j)).d;
-        if (mask[i] & deform_groupbit) {
-          buf[m++] = h_v(j,0) + dvx;
-          buf[m++] = h_v(j,1) + dvy;
-          buf[m++] = h_v(j,2) + dvz;
-        } else {
-          buf[m++] = h_v(j,0);
-          buf[m++] = h_v(j,1);
-          buf[m++] = h_v(j,2);
-        }
-      }
-    }
-  }
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->pack_border(n,list,&buf[m]);
-
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecBondKokkos::pack_border_hybrid(int n, int *list, double *buf)
-{
-  int i,j,m;
-
-  m = 0;
-  for (i = 0; i < n; i++) {
-    j = list[i];
-    buf[m++] = h_molecule(j);
-  }
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType>
 struct AtomVecBondKokkos_UnpackBorder {
   typedef DeviceType device_type;
@@ -485,78 +289,6 @@ void AtomVecBondKokkos::unpack_border_kokkos(const int &n, const int &first,
 
 /* ---------------------------------------------------------------------- */
 
-void AtomVecBondKokkos::unpack_border(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  while (last > nmax) grow(0);
-
-  for (i = first; i < last; i++) {
-    h_x(i,0) = buf[m++];
-    h_x(i,1) = buf[m++];
-    h_x(i,2) = buf[m++];
-    h_tag(i) =  (tagint)  ubuf(buf[m++]).i;
-    h_type(i) = (int) ubuf(buf[m++]).i;
-    h_mask(i) = (int) ubuf(buf[m++]).i;
-    h_molecule(i) = (tagint) ubuf(buf[m++]).i;
-  }
-
-  atomKK->modified(Host,X_MASK|TAG_MASK|TYPE_MASK|MASK_MASK|MOLECULE_MASK);
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->
-        unpack_border(n,first,&buf[m]);
-}
-
-/* ---------------------------------------------------------------------- */
-
-void AtomVecBondKokkos::unpack_border_vel(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  while (last > nmax) grow(0);
-
-  for (i = first; i < last; i++) {
-    h_x(i,0) = buf[m++];
-    h_x(i,1) = buf[m++];
-    h_x(i,2) = buf[m++];
-    h_tag(i) =  (tagint)  ubuf(buf[m++]).i;
-    h_type(i) = (int) ubuf(buf[m++]).i;
-    h_mask(i) = (int) ubuf(buf[m++]).i;
-    h_molecule(i) = (tagint) ubuf(buf[m++]).i;
-    h_v(i,0) = buf[m++];
-    h_v(i,1) = buf[m++];
-    h_v(i,2) = buf[m++];
-  }
-
-  atomKK->modified(Host,X_MASK|TAG_MASK|TYPE_MASK|MASK_MASK|MOLECULE_MASK|V_MASK);
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->
-        unpack_border(n,first,&buf[m]);
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecBondKokkos::unpack_border_hybrid(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  for (i = first; i < last; i++)
-    h_molecule(i) = (tagint) ubuf(buf[m++]).i;
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType>
 struct AtomVecBondKokkos_PackExchangeFunctor {
   typedef DeviceType device_type;
@@ -721,43 +453,6 @@ int AtomVecBondKokkos::pack_exchange_kokkos(const int &nsend,DAT::tdual_xfloat_2
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecBondKokkos::pack_exchange(int i, double *buf)
-{
-  int k;
-  int m = 1;
-  buf[m++] = h_x(i,0);
-  buf[m++] = h_x(i,1);
-  buf[m++] = h_x(i,2);
-  buf[m++] = h_v(i,0);
-  buf[m++] = h_v(i,1);
-  buf[m++] = h_v(i,2);
-  buf[m++] = ubuf(h_tag(i)).d;
-  buf[m++] = ubuf(h_type(i)).d;
-  buf[m++] = ubuf(h_mask(i)).d;
-  buf[m++] = ubuf(h_image(i)).d;
-  buf[m++] = ubuf(h_molecule(i)).d;
-
-  buf[m++] = ubuf(h_num_bond(i)).d;
-  for (k = 0; k < h_num_bond(i); k++) {
-    buf[m++] = ubuf(h_bond_type(i,k)).d;
-    buf[m++] = ubuf(h_bond_atom(i,k)).d;
-  }
-  buf[m++] = ubuf(h_nspecial(i,0)).d;
-  buf[m++] = ubuf(h_nspecial(i,1)).d;
-  buf[m++] = ubuf(h_nspecial(i,2)).d;
-  for (k = 0; k < h_nspecial(i,2); k++)
-    buf[m++] = ubuf(h_special(i,k)).d;
-
-  if (atom->nextra_grow)
-    for (int iextra = 0; iextra < atom->nextra_grow; iextra++)
-      m += modify->fix[atom->extra_grow[iextra]]->pack_exchange(i,&buf[m]);
-
-  buf[0] = m;
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType>
 struct AtomVecBondKokkos_UnpackExchangeFunctor {
   typedef DeviceType device_type;
@@ -870,321 +565,6 @@ int AtomVecBondKokkos::unpack_exchange_kokkos(DAT::tdual_xfloat_2d &k_buf,int nr
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecBondKokkos::unpack_exchange(double *buf)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) grow(0);
-  atomKK->modified(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-           MASK_MASK | IMAGE_MASK | MOLECULE_MASK | BOND_MASK | SPECIAL_MASK);
-
-  int k;
-  int m = 1;
-  h_x(nlocal,0) = buf[m++];
-  h_x(nlocal,1) = buf[m++];
-  h_x(nlocal,2) = buf[m++];
-  h_v(nlocal,0) = buf[m++];
-  h_v(nlocal,1) = buf[m++];
-  h_v(nlocal,2) = buf[m++];
-  h_tag(nlocal) = (tagint) ubuf(buf[m++]).i;
-  h_type(nlocal) = (int) ubuf(buf[m++]).i;
-  h_mask(nlocal) = (int) ubuf(buf[m++]).i;
-  h_image(nlocal) = (imageint) ubuf(buf[m++]).i;
-  h_molecule(nlocal) = (tagint) ubuf(buf[m++]).i;
-
-  h_num_bond(nlocal) = (int) ubuf(buf[m++]).i;
-  for (k = 0; k < h_num_bond(nlocal); k++) {
-    h_bond_type(nlocal,k) = (int) ubuf(buf[m++]).i;
-    h_bond_atom(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-  }
-
-  h_nspecial(nlocal,0) = (int) ubuf(buf[m++]).i;
-  h_nspecial(nlocal,1) = (int) ubuf(buf[m++]).i;
-  h_nspecial(nlocal,2) = (int) ubuf(buf[m++]).i;
-  for (k = 0; k < h_nspecial(nlocal,2); k++)
-   h_special(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-
-  if (atom->nextra_grow)
-    for (int iextra = 0; iextra < atom->nextra_grow; iextra++)
-      m += modify->fix[atom->extra_grow[iextra]]->
-        unpack_exchange(nlocal,&buf[m]);
-
-  atom->nlocal++;
-  return m;
-}
-
-/* ----------------------------------------------------------------------
-   size of restart data for all atoms owned by this proc
-   include extra data stored by fixes
-------------------------------------------------------------------------- */
-
-int AtomVecBondKokkos::size_restart()
-{
-  int i;
-
-  int nlocal = atom->nlocal;
-  int n = 0;
-  for (i = 0; i < nlocal; i++)
-    n += 13 + 2*h_num_bond[i];
-
-  if (atom->nextra_restart)
-    for (int iextra = 0; iextra < atom->nextra_restart; iextra++)
-      for (i = 0; i < nlocal; i++)
-        n += modify->fix[atom->extra_restart[iextra]]->size_restart(i);
-
-  return n;
-}
-
-/* ----------------------------------------------------------------------
-   pack atom I's data for restart file including extra quantities
-   xyz must be 1st 3 values, so that read_restart can test on them
-   molecular types may be negative, but write as positive
-------------------------------------------------------------------------- */
-
-int AtomVecBondKokkos::pack_restart(int i, double *buf)
-{
-  atomKK->sync(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-            MASK_MASK | IMAGE_MASK | MOLECULE_MASK | BOND_MASK | SPECIAL_MASK);
-  int m = 1;
-  buf[m++] = h_x(i,0);
-  buf[m++] = h_x(i,1);
-  buf[m++] = h_x(i,2);
-  buf[m++] = ubuf(h_tag(i)).d;
-  buf[m++] = ubuf(h_type(i)).d;
-  buf[m++] = ubuf(h_mask(i)).d;
-  buf[m++] = ubuf(h_image(i)).d;
-  buf[m++] = h_v(i,0);
-  buf[m++] = h_v(i,1);
-  buf[m++] = h_v(i,2);
-
-  buf[m++] = ubuf(h_molecule(i)).d;
-
-  buf[m++] = ubuf(h_num_bond(i)).d;
-  for (int k = 0; k < h_num_bond(i); k++) {
-    buf[m++] = ubuf(MAX(h_bond_type(i,k),-h_bond_type(i,k))).d;
-    buf[m++] = ubuf(h_bond_atom(i,k)).d;
-  }
-
-  if (atom->nextra_restart)
-    for (int iextra = 0; iextra < atom->nextra_restart; iextra++)
-      m += modify->fix[atom->extra_restart[iextra]]->pack_restart(i,&buf[m]);
-
-  buf[0] = m;
-  return m;
-}
-
-/* ----------------------------------------------------------------------
-   unpack data for one atom from restart file including extra quantities
-------------------------------------------------------------------------- */
-
-int AtomVecBondKokkos::unpack_restart(double *buf)
-{
-  int k;
-
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) {
-    grow(0);
-    if (atom->nextra_store)
-      memory->grow(atom->extra,nmax,atom->nextra_store,"atom:extra");
-  }
-  atomKK->modified(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-           MASK_MASK | IMAGE_MASK | MOLECULE_MASK | BOND_MASK | SPECIAL_MASK);
-  int m = 1;
-  h_x(nlocal,0) = buf[m++];
-  h_x(nlocal,1) = buf[m++];
-  h_x(nlocal,2) = buf[m++];
-  h_tag(nlocal) = (tagint) ubuf(buf[m++]).i;
-  h_type(nlocal) = (int) ubuf(buf[m++]).i;
-  h_mask(nlocal) = (int) ubuf(buf[m++]).i;
-  h_image(nlocal) = (imageint) ubuf(buf[m++]).i;
-  h_v(nlocal,0) = buf[m++];
-  h_v(nlocal,1) = buf[m++];
-  h_v(nlocal,2) = buf[m++];
-
-  h_molecule(nlocal) = (tagint) ubuf(buf[m++]).i;
-
-  h_num_bond(nlocal) = (int) ubuf(buf[m++]).i;
-  for (k = 0; k < h_num_bond(nlocal); k++) {
-    h_bond_type(nlocal,k) = (int) ubuf(buf[m++]).i;
-    h_bond_atom(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-  }
-
-  h_nspecial(nlocal,0) = h_nspecial(nlocal,1) = h_nspecial(nlocal,2) = 0;
-
-  double **extra = atom->extra;
-  if (atom->nextra_store) {
-    int size = static_cast<int> (buf[0]) - m;
-    for (int i = 0; i < size; i++) extra[nlocal][i] = buf[m++];
-  }
-
-  atom->nlocal++;
-  return m;
-}
-
-/* ----------------------------------------------------------------------
-   create one atom of itype at coord
-   set other values to defaults
-------------------------------------------------------------------------- */
-
-void AtomVecBondKokkos::create_atom(int itype, double *coord)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) {
-    atomKK->modified(Host,ALL_MASK);
-    grow(0);
-  }
-  atomKK->modified(Host,ALL_MASK);
-
-  tag[nlocal] = 0;
-  type[nlocal] = itype;
-  h_x(nlocal,0) = coord[0];
-  h_x(nlocal,1) = coord[1];
-  h_x(nlocal,2) = coord[2];
-  h_mask(nlocal) = 1;
-  h_image(nlocal) = ((imageint) IMGMAX << IMG2BITS) |
-    ((imageint) IMGMAX << IMGBITS) | IMGMAX;
-  h_v(nlocal,0) = 0.0;
-  h_v(nlocal,1) = 0.0;
-  h_v(nlocal,2) = 0.0;
-
-  h_molecule(nlocal) = 0;
-  h_num_bond(nlocal) = 0;
-  h_nspecial(nlocal,0) = h_nspecial(nlocal,1) = h_nspecial(nlocal,2) = 0;
-
-  atom->nlocal++;
-}
-
-/* ----------------------------------------------------------------------
-   unpack one line from Atoms section of data file
-   initialize other atom quantities
-------------------------------------------------------------------------- */
-
-void AtomVecBondKokkos::data_atom(double *coord, imageint imagetmp,
-                                  const std::vector<std::string> &values, std::string &extract)
-{
-  int nlocal = atomKK->nlocal;
-  if (nlocal == nmax) grow(0);
-  atomKK->modified(Host,ALL_MASK);
-
-  h_tag(nlocal) = utils::inumeric(FLERR,values[0],true,lmp);
-  h_molecule(nlocal) = utils::inumeric(FLERR,values[1],true,lmp);
-  h_type(nlocal) = utils::inumeric(FLERR,values[2],true,lmp);
-  extract = values[2];
-  if (h_type(nlocal) <= 0 || h_type(nlocal) > atom->ntypes)
-    error->one(FLERR,"Invalid atom type in Atoms section of data file");
-
-  h_x(nlocal,0) = coord[0];
-  h_x(nlocal,1) = coord[1];
-  h_x(nlocal,2) = coord[2];
-
-  h_image(nlocal) = imagetmp;
-
-  h_mask(nlocal) = 1;
-  h_v(nlocal,0) = 0.0;
-  h_v(nlocal,1) = 0.0;
-  h_v(nlocal,2) = 0.0;
-  h_num_bond(nlocal) = 0;
-
-  atomKK->nlocal++;
-}
-
-/* ----------------------------------------------------------------------
-   unpack hybrid quantities from one line in Atoms section of data file
-   initialize other atom quantities for this sub-style
-------------------------------------------------------------------------- */
-
-int AtomVecBondKokkos::data_atom_hybrid(int nlocal, const std::vector<std::string> &values,
-                                        int offset)
-{
-  h_molecule(nlocal) = utils::inumeric(FLERR,values[offset],true,lmp);
-  h_num_bond(nlocal) = 0;
-  return 1;
-}
-
-/* ----------------------------------------------------------------------
-   pack atom info for data file including 3 image flags
-------------------------------------------------------------------------- */
-
-void AtomVecBondKokkos::pack_data(double **buf)
-{
-  int nlocal = atom->nlocal;
-  for (int i = 0; i < nlocal; i++) {
-    buf[i][0] = h_tag(i);
-    buf[i][1] = h_molecule(i);
-    buf[i][2] = h_type(i);
-    buf[i][3] = h_x(i,0);
-    buf[i][4] = h_x(i,1);
-    buf[i][5] = h_x(i,2);
-    buf[i][6] = (h_image[i] & IMGMASK) - IMGMAX;
-    buf[i][7] = (h_image[i] >> IMGBITS & IMGMASK) - IMGMAX;
-    buf[i][8] = (h_image[i] >> IMG2BITS) - IMGMAX;
-  }
-}
-
-/* ----------------------------------------------------------------------
-   pack hybrid atom info for data file
-------------------------------------------------------------------------- */
-
-int AtomVecBondKokkos::pack_data_hybrid(int i, double *buf)
-{
-  buf[0] = h_molecule(i);
-  return 1;
-}
-
-/* ----------------------------------------------------------------------
-   write atom info to data file including 3 image flags
-------------------------------------------------------------------------- */
-
-void AtomVecBondKokkos::write_data(FILE *fp, int n, double **buf)
-{
-  for (int i = 0; i < n; i++)
-    fprintf(fp,"%d %d %d %-1.16e %-1.16e %-1.16e %d %d %d\n",
-            (int) buf[i][0],(int) buf[i][1], (int) buf[i][2],
-            buf[i][3],buf[i][4],buf[i][5],
-            (int) buf[i][6],(int) buf[i][7],(int) buf[i][8]);
-}
-
-/* ----------------------------------------------------------------------
-   write hybrid atom info to data file
-------------------------------------------------------------------------- */
-
-int AtomVecBondKokkos::write_data_hybrid(FILE *fp, double *buf)
-{
-  fprintf(fp," " TAGINT_FORMAT, (tagint) (buf[0]));
-  return 1;
-}
-
-/* ----------------------------------------------------------------------
-   return # of bytes of allocated memory
-------------------------------------------------------------------------- */
-
-double AtomVecBondKokkos::memory_usage()
-{
-  double bytes = 0;
-
-  if (atom->memcheck("tag")) bytes += memory->usage(tag,nmax);
-  if (atom->memcheck("type")) bytes += memory->usage(type,nmax);
-  if (atom->memcheck("mask")) bytes += memory->usage(mask,nmax);
-  if (atom->memcheck("image")) bytes += memory->usage(image,nmax);
-  if (atom->memcheck("x")) bytes += memory->usage(x,nmax,3);
-  if (atom->memcheck("v")) bytes += memory->usage(v,nmax,3);
-  if (atom->memcheck("f")) bytes += memory->usage(f,nmax*commKK->nthreads,3);
-
-  if (atom->memcheck("molecule")) bytes += memory->usage(molecule,nmax);
-  if (atom->memcheck("nspecial")) bytes += memory->usage(nspecial,nmax,3);
-  if (atom->memcheck("special"))
-    bytes += memory->usage(special,nmax,atom->maxspecial);
-
-  if (atom->memcheck("num_bond")) bytes += memory->usage(num_bond,nmax);
-  if (atom->memcheck("bond_type"))
-    bytes += memory->usage(bond_type,nmax,atom->bond_per_atom);
-  if (atom->memcheck("bond_atom"))
-    bytes += memory->usage(bond_atom,nmax,atom->bond_per_atom);
-
-  return bytes;
-}
-
-/* ---------------------------------------------------------------------- */
-
 void AtomVecBondKokkos::sync(ExecutionSpace space, unsigned int mask)
 {
   if (space == Device) {
@@ -1337,4 +717,3 @@ void AtomVecBondKokkos::modified(ExecutionSpace space, unsigned int mask)
     }
   }
 }
-

src/KOKKOS/atom_vec_bond_kokkos.h

@@ -24,35 +24,15 @@ AtomStyle(bond/kk/host,AtomVecBondKokkos);
 #define LMP_ATOM_VEC_BOND_KOKKOS_H
 
 #include "atom_vec_kokkos.h"
+#include "atom_vec_bond.h"
 
 namespace LAMMPS_NS {
 
-class AtomVecBondKokkos : public AtomVecKokkos {
+class AtomVecBondKokkos : public AtomVecKokkos, public AtomVecBond {
  public:
   AtomVecBondKokkos(class LAMMPS *);
 
   void grow(int) override;
-  void copy(int, int, int) override;
-  int pack_border(int, int *, double *, int, int *) override;
-  int pack_border_vel(int, int *, double *, int, int *) override;
-  int pack_border_hybrid(int, int *, double *) override;
-  void unpack_border(int, int, double *) override;
-  void unpack_border_vel(int, int, double *) override;
-  int unpack_border_hybrid(int, int, double *) override;
-  int pack_exchange(int, double *) override;
-  int unpack_exchange(double *) override;
-  int size_restart() override;
-  int pack_restart(int, double *) override;
-  int unpack_restart(double *) override;
-  void create_atom(int, double *) override;
-  void data_atom(double *, imageint, const std::vector<std::string> &, std::string &) override;
-  int data_atom_hybrid(int, const std::vector<std::string> &, int) override;
-  void pack_data(double **) override;
-  int pack_data_hybrid(int, double *) override;
-  void write_data(FILE *, int, double **) override;
-  int write_data_hybrid(FILE *, double *) override;
-  double memory_usage() override;
-
   void grow_pointers() override;
   int pack_border_kokkos(int n, DAT::tdual_int_2d k_sendlist,
                          DAT::tdual_xfloat_2d buf,int iswap,
@@ -112,9 +92,6 @@ class AtomVecBondKokkos : public AtomVecKokkos {
   HAT::t_int_1d h_num_bond;
   HAT::t_int_2d h_bond_type;
   HAT::t_tagint_2d h_bond_atom;
-
-  DAT::tdual_int_1d k_count;
-
 };
 
 }

src/KOKKOS/atom_vec_charge_kokkos.cpp

@@ -27,25 +27,9 @@ using namespace LAMMPS_NS;
 
 /* ---------------------------------------------------------------------- */
 
-AtomVecChargeKokkos::AtomVecChargeKokkos(LAMMPS *lmp) : AtomVecKokkos(lmp)
+AtomVecChargeKokkos::AtomVecChargeKokkos(LAMMPS *lmp) : AtomVec(lmp),
+AtomVecKokkos(lmp), AtomVecCharge(lmp)
 {
-  molecular = Atom::ATOMIC;
-  mass_type = PER_TYPE;
-
-  comm_x_only = comm_f_only = 1;
-  size_forward = 3;
-  size_reverse = 3;
-  size_border = 7;
-  size_velocity = 3;
-  size_data_atom = 6;
-  size_data_vel = 4;
-  xcol_data = 4;
-
-  atom->q_flag = 1;
-
-  k_count = DAT::tdual_int_1d("atom::k_count",1);
-  atomKK = (AtomKokkos *) atom;
-  commKK = (CommKokkos *) comm;
 
 }
 
@@ -123,30 +107,6 @@ void AtomVecChargeKokkos::grow_pointers()
 
 }
 
-/* ----------------------------------------------------------------------
-   copy atom I info to atom J
-------------------------------------------------------------------------- */
-
-void AtomVecChargeKokkos::copy(int i, int j, int delflag)
-{
-  h_tag[j] = h_tag[i];
-  h_type[j] = h_type[i];
-  mask[j] = mask[i];
-  h_image[j] = h_image[i];
-  h_x(j,0) = h_x(i,0);
-  h_x(j,1) = h_x(i,1);
-  h_x(j,2) = h_x(i,2);
-  h_v(j,0) = h_v(i,0);
-  h_v(j,1) = h_v(i,1);
-  h_v(j,2) = h_v(i,2);
-
-  h_q[j] = h_q[i];
-
-  if (atom->nextra_grow)
-    for (int iextra = 0; iextra < atom->nextra_grow; iextra++)
-      modify->fix[atom->extra_grow[iextra]]->copy_arrays(i,j,delflag);
-}
-
 /* ---------------------------------------------------------------------- */
 
 template<class DeviceType,int PBC_FLAG,int TRICLINIC>
@@ -297,150 +257,6 @@ int AtomVecChargeKokkos::pack_border_kokkos(int n, DAT::tdual_int_2d k_sendlist,
   return n*size_border;
 }
 
-/* ---------------------------------------------------------------------- */
-
-int AtomVecChargeKokkos::pack_border(int n, int *list, double *buf,
-                               int pbc_flag, int *pbc)
-{
-  int i,j,m;
-  double dx,dy,dz;
-
-  m = 0;
-  if (pbc_flag == 0) {
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0);
-      buf[m++] = h_x(j,1);
-      buf[m++] = h_x(j,2);
-      buf[m++] = ubuf(h_tag(j)).d;
-      buf[m++] = ubuf(h_type(j)).d;
-      buf[m++] = ubuf(h_mask(j)).d;
-      buf[m++] = h_q(j);
-    }
-  } else {
-    if (domain->triclinic == 0) {
-      dx = pbc[0]*domain->xprd;
-      dy = pbc[1]*domain->yprd;
-      dz = pbc[2]*domain->zprd;
-    } else {
-      dx = pbc[0];
-      dy = pbc[1];
-      dz = pbc[2];
-    }
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0) + dx;
-      buf[m++] = h_x(j,1) + dy;
-      buf[m++] = h_x(j,2) + dz;
-      buf[m++] = ubuf(h_tag(j)).d;
-      buf[m++] = ubuf(h_type(j)).d;
-      buf[m++] = ubuf(h_mask(j)).d;
-      buf[m++] = h_q(j);
-    }
-  }
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->pack_border(n,list,&buf[m]);
-
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecChargeKokkos::pack_border_vel(int n, int *list, double *buf,
-                                   int pbc_flag, int *pbc)
-{
-  int i,j,m;
-  double dx,dy,dz,dvx,dvy,dvz;
-
-  m = 0;
-  if (pbc_flag == 0) {
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0);
-      buf[m++] = h_x(j,1);
-      buf[m++] = h_x(j,2);
-      buf[m++] = ubuf(h_tag(j)).d;
-      buf[m++] = ubuf(h_type(j)).d;
-      buf[m++] = ubuf(h_mask(j)).d;
-      buf[m++] = h_q[j];
-      buf[m++] = h_v(j,0);
-      buf[m++] = h_v(j,1);
-      buf[m++] = h_v(j,2);
-    }
-  } else {
-    if (domain->triclinic == 0) {
-      dx = pbc[0]*domain->xprd;
-      dy = pbc[1]*domain->yprd;
-      dz = pbc[2]*domain->zprd;
-    } else {
-      dx = pbc[0];
-      dy = pbc[1];
-      dz = pbc[2];
-    }
-    if (!deform_vremap) {
-      for (i = 0; i < n; i++) {
-        j = list[i];
-        buf[m++] = h_x(j,0) + dx;
-        buf[m++] = h_x(j,1) + dy;
-        buf[m++] = h_x(j,2) + dz;
-        buf[m++] = ubuf(h_tag(j)).d;
-        buf[m++] = ubuf(h_type(j)).d;
-        buf[m++] = ubuf(h_mask(j)).d;
-        buf[m++] = h_q[j];
-        buf[m++] = h_v(j,0);
-        buf[m++] = h_v(j,1);
-        buf[m++] = h_v(j,2);
-      }
-    } else {
-      dvx = pbc[0]*h_rate[0] + pbc[5]*h_rate[5] + pbc[4]*h_rate[4];
-      dvy = pbc[1]*h_rate[1] + pbc[3]*h_rate[3];
-      dvz = pbc[2]*h_rate[2];
-      for (i = 0; i < n; i++) {
-        j = list[i];
-        buf[m++] = h_x(j,0) + dx;
-        buf[m++] = h_x(j,1) + dy;
-        buf[m++] = h_x(j,2) + dz;
-        buf[m++] = ubuf(h_tag(j)).d;
-        buf[m++] = ubuf(h_type(j)).d;
-        buf[m++] = ubuf(h_mask(j)).d;
-        buf[m++] = h_q[j];
-        if (mask[i] & deform_groupbit) {
-          buf[m++] = h_v(j,0) + dvx;
-          buf[m++] = h_v(j,1) + dvy;
-          buf[m++] = h_v(j,2) + dvz;
-        } else {
-          buf[m++] = h_v(j,0);
-          buf[m++] = h_v(j,1);
-          buf[m++] = h_v(j,2);
-        }
-      }
-    }
-  }
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->pack_border(n,list,&buf[m]);
-
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecChargeKokkos::pack_border_hybrid(int n, int *list, double *buf)
-{
-  int i,j,m;
-
-  m = 0;
-  for (i = 0; i < n; i++) {
-    j = list[i];
-    buf[m++] = h_q[j];
-  }
-  return m;
-}
-
-
 /* ---------------------------------------------------------------------- */
 
 template<class DeviceType>
@@ -500,78 +316,6 @@ void AtomVecChargeKokkos::unpack_border_kokkos(const int &n, const int &first,
 
 /* ---------------------------------------------------------------------- */
 
-void AtomVecChargeKokkos::unpack_border(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  while (last > nmax) grow(0);
-
-  for (i = first; i < last; i++) {
-    h_x(i,0) = buf[m++];
-    h_x(i,1) = buf[m++];
-    h_x(i,2) = buf[m++];
-    h_tag(i) =  (tagint)  ubuf(buf[m++]).i;
-    h_type(i) = (int) ubuf(buf[m++]).i;
-    h_mask(i) = (int) ubuf(buf[m++]).i;
-    h_q[i] = buf[m++];
-  }
-
-  atomKK->modified(Host,X_MASK|TAG_MASK|TYPE_MASK|MASK_MASK|Q_MASK);
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->
-        unpack_border(n,first,&buf[m]);
-}
-
-/* ---------------------------------------------------------------------- */
-
-void AtomVecChargeKokkos::unpack_border_vel(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  while (last > nmax) grow(0);
-
-  for (i = first; i < last; i++) {
-    h_x(i,0) = buf[m++];
-    h_x(i,1) = buf[m++];
-    h_x(i,2) = buf[m++];
-    h_tag(i) =  (tagint)  ubuf(buf[m++]).i;
-    h_type(i) = (int) ubuf(buf[m++]).i;
-    h_mask(i) = (int) ubuf(buf[m++]).i;
-    h_q[i] = buf[m++];
-    h_v(i,0) = buf[m++];
-    h_v(i,1) = buf[m++];
-    h_v(i,2) = buf[m++];
-  }
-
-  atomKK->modified(Host,X_MASK|TAG_MASK|TYPE_MASK|MASK_MASK|Q_MASK|V_MASK);
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->
-        unpack_border(n,first,&buf[m]);
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecChargeKokkos::unpack_border_hybrid(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  for (i = first; i < last; i++)
-    h_q[i] = buf[m++];
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType>
 struct AtomVecChargeKokkos_PackExchangeFunctor {
   typedef DeviceType device_type;
@@ -688,31 +432,6 @@ int AtomVecChargeKokkos::pack_exchange_kokkos(const int &nsend,DAT::tdual_xfloat
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecChargeKokkos::pack_exchange(int i, double *buf)
-{
-  int m = 1;
-  buf[m++] = h_x(i,0);
-  buf[m++] = h_x(i,1);
-  buf[m++] = h_x(i,2);
-  buf[m++] = h_v(i,0);
-  buf[m++] = h_v(i,1);
-  buf[m++] = h_v(i,2);
-  buf[m++] = ubuf(h_tag(i)).d;
-  buf[m++] = ubuf(h_type(i)).d;
-  buf[m++] = ubuf(h_mask(i)).d;
-  buf[m++] = ubuf(h_image(i)).d;
-  buf[m++] = h_q[i];
-
-  if (atom->nextra_grow)
-    for (int iextra = 0; iextra < atom->nextra_grow; iextra++)
-      m += modify->fix[atom->extra_grow[iextra]]->pack_exchange(i,&buf[m]);
-
-  buf[0] = m;
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType>
 struct AtomVecChargeKokkos_UnpackExchangeFunctor {
   typedef DeviceType device_type;
@@ -797,281 +516,6 @@ int AtomVecChargeKokkos::unpack_exchange_kokkos(DAT::tdual_xfloat_2d &k_buf,int
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecChargeKokkos::unpack_exchange(double *buf)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) grow(0);
-  atomKK->modified(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-           MASK_MASK | IMAGE_MASK | Q_MASK);
-
-  int m = 1;
-  h_x(nlocal,0) = buf[m++];
-  h_x(nlocal,1) = buf[m++];
-  h_x(nlocal,2) = buf[m++];
-  h_v(nlocal,0) = buf[m++];
-  h_v(nlocal,1) = buf[m++];
-  h_v(nlocal,2) = buf[m++];
-  h_tag(nlocal) = (tagint) ubuf(buf[m++]).i;
-  h_type(nlocal) = (int) ubuf(buf[m++]).i;
-  h_mask(nlocal) = (int) ubuf(buf[m++]).i;
-  h_image(nlocal) = (imageint) ubuf(buf[m++]).i;
-  h_q[nlocal] = buf[m++];
-
-  if (atom->nextra_grow)
-    for (int iextra = 0; iextra < atom->nextra_grow; iextra++)
-      m += modify->fix[atom->extra_grow[iextra]]->
-        unpack_exchange(nlocal,&buf[m]);
-
-  atom->nlocal++;
-  return m;
-}
-
-/* ----------------------------------------------------------------------
-   size of restart data for all atoms owned by this proc
-   include extra data stored by fixes
-------------------------------------------------------------------------- */
-
-int AtomVecChargeKokkos::size_restart()
-{
-  int i;
-
-  int nlocal = atom->nlocal;
-  int n = 12 * nlocal;
-
-  if (atom->nextra_restart)
-    for (int iextra = 0; iextra < atom->nextra_restart; iextra++)
-      for (i = 0; i < nlocal; i++)
-        n += modify->fix[atom->extra_restart[iextra]]->size_restart(i);
-
-  return n;
-}
-
-/* ----------------------------------------------------------------------
-   pack atom I's data for restart file including extra quantities
-   xyz must be 1st 3 values, so that read_restart can test on them
-   molecular types may be negative, but write as positive
-------------------------------------------------------------------------- */
-
-int AtomVecChargeKokkos::pack_restart(int i, double *buf)
-{
-  atomKK->sync(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-            MASK_MASK | IMAGE_MASK | Q_MASK);
-
-  int m = 1;
-  buf[m++] = h_x(i,0);
-  buf[m++] = h_x(i,1);
-  buf[m++] = h_x(i,2);
-  buf[m++] = ubuf(h_tag(i)).d;
-  buf[m++] = ubuf(h_type(i)).d;
-  buf[m++] = ubuf(h_mask(i)).d;
-  buf[m++] = ubuf(h_image(i)).d;
-  buf[m++] = h_v(i,0);
-  buf[m++] = h_v(i,1);
-  buf[m++] = h_v(i,2);
-
-  buf[m++] = h_q[i];
-
-  if (atom->nextra_restart)
-    for (int iextra = 0; iextra < atom->nextra_restart; iextra++)
-      m += modify->fix[atom->extra_restart[iextra]]->pack_restart(i,&buf[m]);
-
-  buf[0] = m;
-  return m;
-}
-
-/* ----------------------------------------------------------------------
-   unpack data for one atom from restart file including extra quantities
-------------------------------------------------------------------------- */
-
-int AtomVecChargeKokkos::unpack_restart(double *buf)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) {
-    grow(0);
-    if (atom->nextra_store)
-      memory->grow(atom->extra,nmax,atom->nextra_store,"atom:extra");
-  }
-
-  atomKK->modified(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-           MASK_MASK | IMAGE_MASK | Q_MASK);
-
-  int m = 1;
-  h_x(nlocal,0) = buf[m++];
-  h_x(nlocal,1) = buf[m++];
-  h_x(nlocal,2) = buf[m++];
-  h_tag(nlocal) = (tagint) ubuf(buf[m++]).i;
-  h_type(nlocal) = (int) ubuf(buf[m++]).i;
-  h_mask(nlocal) = (int) ubuf(buf[m++]).i;
-  h_image(nlocal) = (imageint) ubuf(buf[m++]).i;
-  h_v(nlocal,0) = buf[m++];
-  h_v(nlocal,1) = buf[m++];
-  h_v(nlocal,2) = buf[m++];
-
-  h_q[nlocal] = buf[m++];
-
-  double **extra = atom->extra;
-  if (atom->nextra_store) {
-    int size = static_cast<int> (buf[0]) - m;
-    for (int i = 0; i < size; i++) extra[nlocal][i] = buf[m++];
-  }
-
-  atom->nlocal++;
-  return m;
-}
-
-/* ----------------------------------------------------------------------
-   create one atom of itype at coord
-   set other values to defaults
-------------------------------------------------------------------------- */
-
-void AtomVecChargeKokkos::create_atom(int itype, double *coord)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) {
-    atomKK->modified(Host,ALL_MASK);
-    grow(0);
-  }
-  atomKK->sync(Host,ALL_MASK);
-  atomKK->modified(Host,ALL_MASK);
-
-  tag[nlocal] = 0;
-  type[nlocal] = itype;
-  h_x(nlocal,0) = coord[0];
-  h_x(nlocal,1) = coord[1];
-  h_x(nlocal,2) = coord[2];
-  h_mask[nlocal] = 1;
-  h_image[nlocal] = ((imageint) IMGMAX << IMG2BITS) |
-    ((imageint) IMGMAX << IMGBITS) | IMGMAX;
-  h_v(nlocal,0) = 0.0;
-  h_v(nlocal,1) = 0.0;
-  h_v(nlocal,2) = 0.0;
-
-  h_q[nlocal] = 0.0;
-
-  atom->nlocal++;
-}
-
-/* ----------------------------------------------------------------------
-   unpack one line from Atoms section of data file
-   initialize other atom quantities
-------------------------------------------------------------------------- */
-
-void AtomVecChargeKokkos::data_atom(double *coord, imageint imagetmp,
-                                    const std::vector<std::string> &values, std::string &extract)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) grow(0);
-
-  h_tag[nlocal] = utils::inumeric(FLERR,values[0],true,lmp);
-  h_type[nlocal] = utils::inumeric(FLERR,values[1],true,lmp);
-  extract = values[1];
-  if (type[nlocal] <= 0 || type[nlocal] > atom->ntypes)
-    error->one(FLERR,"Invalid atom type in Atoms section of data file");
-
-  h_q[nlocal] = utils::numeric(FLERR,values[2],true,lmp);
-
-  h_x(nlocal,0) = coord[0];
-  h_x(nlocal,1) = coord[1];
-  h_x(nlocal,2) = coord[2];
-
-  h_image[nlocal] = imagetmp;
-
-  h_mask[nlocal] = 1;
-  h_v(nlocal,0) = 0.0;
-  h_v(nlocal,1) = 0.0;
-  h_v(nlocal,2) = 0.0;
-
-  atomKK->modified(Host,ALL_MASK);
-
-  atom->nlocal++;
-}
-/* ----------------------------------------------------------------------
-   unpack hybrid quantities from one line in Atoms section of data file
-   initialize other atom quantities for this sub-style
-------------------------------------------------------------------------- */
-
-int AtomVecChargeKokkos::data_atom_hybrid(int nlocal, const std::vector<std::string> &values,
-                                          int offset)
-{
-  h_q[nlocal] = utils::numeric(FLERR,values[offset],true,lmp);
-
-  return 1;
-}
-/* ----------------------------------------------------------------------
-   pack atom info for data file including 3 image flags
-------------------------------------------------------------------------- */
-
-void AtomVecChargeKokkos::pack_data(double **buf)
-{
-  int nlocal = atom->nlocal;
-  for (int i = 0; i < nlocal; i++) {
-    buf[i][0] = h_tag[i];
-    buf[i][1] = h_type[i];
-    buf[i][2] = h_q[i];
-    buf[i][3] = h_x(i,0);
-    buf[i][4] = h_x(i,1);
-    buf[i][5] = h_x(i,2);
-    buf[i][6] = (h_image[i] & IMGMASK) - IMGMAX;
-    buf[i][7] = (h_image[i] >> IMGBITS & IMGMASK) - IMGMAX;
-    buf[i][8] = (h_image[i] >> IMG2BITS) - IMGMAX;
-  }
-}
-
-/* ----------------------------------------------------------------------
-   pack hybrid atom info for data file
-------------------------------------------------------------------------- */
-
-int AtomVecChargeKokkos::pack_data_hybrid(int i, double *buf)
-{
-  buf[0] = h_q[i];
-  return 1;
-}
-
-/* ----------------------------------------------------------------------
-   write atom info to data file including 3 image flags
-------------------------------------------------------------------------- */
-
-void AtomVecChargeKokkos::write_data(FILE *fp, int n, double **buf)
-{
-  for (int i = 0; i < n; i++)
-    fprintf(fp,"%d %d %-1.16e %-1.16e %-1.16e %-1.16e %d %d %d\n",
-            (int) buf[i][0],(int) buf[i][1],buf[i][2],buf[i][3],buf[i][4],buf[i][5],
-            (int) buf[i][6],(int) buf[i][7],(int) buf[i][8]);
-}
-
-/* ----------------------------------------------------------------------
-   write hybrid atom info to data file
-------------------------------------------------------------------------- */
-
-int AtomVecChargeKokkos::write_data_hybrid(FILE *fp, double *buf)
-{
-  fprintf(fp," %-1.16e",buf[0]);
-  return 1;
-}
-
-/* ----------------------------------------------------------------------
-   return # of bytes of allocated memory
-------------------------------------------------------------------------- */
-
-double AtomVecChargeKokkos::memory_usage()
-{
-  double bytes = 0;
-
-  if (atom->memcheck("tag")) bytes += memory->usage(tag,nmax);
-  if (atom->memcheck("type")) bytes += memory->usage(type,nmax);
-  if (atom->memcheck("mask")) bytes += memory->usage(mask,nmax);
-  if (atom->memcheck("image")) bytes += memory->usage(image,nmax);
-  if (atom->memcheck("x")) bytes += memory->usage(x,nmax,3);
-  if (atom->memcheck("v")) bytes += memory->usage(v,nmax,3);
-  if (atom->memcheck("f")) bytes += memory->usage(f,nmax*commKK->nthreads,3);
-
-  if (atom->memcheck("q")) bytes += memory->usage(q,nmax);
-
-  return bytes;
-}
-
-/* ---------------------------------------------------------------------- */
-
 void AtomVecChargeKokkos::sync(ExecutionSpace space, unsigned int mask)
 {
   if (space == Device) {
@@ -1120,6 +564,8 @@ void AtomVecChargeKokkos::modified(ExecutionSpace space, unsigned int mask)
   }
 }
 
+/* ---------------------------------------------------------------------- */
+
 void AtomVecChargeKokkos::sync_overlapping_device(ExecutionSpace space, unsigned int mask)
 {
   if (space == Device) {
@@ -1158,4 +604,3 @@ void AtomVecChargeKokkos::sync_overlapping_device(ExecutionSpace space, unsigned
       perform_async_copy<DAT::tdual_float_1d>(atomKK->k_q,space);
   }
 }
-

src/KOKKOS/atom_vec_charge_kokkos.h

@@ -24,36 +24,16 @@ AtomStyle(charge/kk/host,AtomVecChargeKokkos);
 #define LMP_ATOM_VEC_CHARGE_KOKKOS_H
 
 #include "atom_vec_kokkos.h"
+#include "atom_vec_charge.h"
 #include "kokkos_type.h"
 
 namespace LAMMPS_NS {
 
-class AtomVecChargeKokkos : public AtomVecKokkos {
+class AtomVecChargeKokkos : public AtomVecKokkos, public AtomVecCharge {
  public:
   AtomVecChargeKokkos(class LAMMPS *);
 
   void grow(int) override;
-  void copy(int, int, int) override;
-  int pack_border(int, int *, double *, int, int *) override;
-  int pack_border_vel(int, int *, double *, int, int *) override;
-  int pack_border_hybrid(int, int *, double *) override;
-  void unpack_border(int, int, double *) override;
-  void unpack_border_vel(int, int, double *) override;
-  int unpack_border_hybrid(int, int, double *) override;
-  int pack_exchange(int, double *) override;
-  int unpack_exchange(double *) override;
-  int size_restart() override;
-  int pack_restart(int, double *) override;
-  int unpack_restart(double *) override;
-  void create_atom(int, double *) override;
-  void data_atom(double *, imageint, const std::vector<std::string> &, std::string &) override;
-  int data_atom_hybrid(int , const std::vector<std::string> &, int) override;
-  void pack_data(double **) override;
-  int pack_data_hybrid(int, double *) override;
-  void write_data(FILE *, int, double **) override;
-  int write_data_hybrid(FILE *, double *) override;
-  double memory_usage() override;
-
   void grow_pointers() override;
   int pack_border_kokkos(int n, DAT::tdual_int_2d k_sendlist,
                          DAT::tdual_xfloat_2d buf,int iswap,
@@ -98,8 +78,6 @@ class AtomVecChargeKokkos : public AtomVecKokkos {
   DAT::t_float_1d d_q;
 
   HAT::t_float_1d h_q;
-
-  DAT::tdual_int_1d k_count;
 };
 
 }

src/KOKKOS/atom_vec_dpd_kokkos.cpp

@@ -27,27 +27,9 @@ using namespace LAMMPS_NS;
 
 /* ---------------------------------------------------------------------- */
 
-AtomVecDPDKokkos::AtomVecDPDKokkos(LAMMPS *lmp) : AtomVecKokkos(lmp)
+AtomVecDPDKokkos::AtomVecDPDKokkos(LAMMPS *lmp) : AtomVec(lmp),
+AtomVecKokkos(lmp), AtomVecDPD(lmp)
 {
-  molecular = Atom::ATOMIC;
-  mass_type = PER_TYPE;
-
-  comm_x_only = comm_f_only = 0;
-  size_forward = 7;
-  size_reverse = 3;
-  size_border = 12;
-  size_velocity = 3;
-  size_data_atom = 6;
-  size_data_vel = 4;
-  xcol_data = 4;
-
-  atom->rho_flag = 1;
-  atom->dpd_flag = 1;
-
-  k_count = DAT::tdual_int_1d("atom::k_count",1);
-  atomKK = (AtomKokkos *) atom;
-  commKK = (CommKokkos *) comm;
-
   no_comm_vel_flag = 1;
 }
 
@@ -153,44 +135,6 @@ void AtomVecDPDKokkos::grow_pointers()
   h_duChem = atomKK->k_duChem.h_view;
 }
 
-/* ----------------------------------------------------------------------
-   copy atom I info to atom J
-------------------------------------------------------------------------- */
-
-void AtomVecDPDKokkos::copy(int i, int j, int delflag)
-{
-  atomKK->sync(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-            MASK_MASK | IMAGE_MASK | DPDTHETA_MASK |
-            UCG_MASK | UCGNEW_MASK |
-            UCOND_MASK | UMECH_MASK | UCHEM_MASK | DVECTOR_MASK);
-
-  h_tag[j] = h_tag[i];
-  h_type[j] = h_type[i];
-  mask[j] = mask[i];
-  h_image[j] = h_image[i];
-  h_x(j,0) = h_x(i,0);
-  h_x(j,1) = h_x(i,1);
-  h_x(j,2) = h_x(i,2);
-  h_v(j,0) = h_v(i,0);
-  h_v(j,1) = h_v(i,1);
-  h_v(j,2) = h_v(i,2);
-  h_dpdTheta[j] = h_dpdTheta[i];
-  h_uCond[j] = h_uCond[i];
-  h_uMech[j] = h_uMech[i];
-  h_uChem[j] = h_uChem[i];
-  h_uCG[j] = h_uCG[i];
-  h_uCGnew[j] = h_uCGnew[i];
-
-  if (atom->nextra_grow)
-    for (int iextra = 0; iextra < atom->nextra_grow; iextra++)
-      modify->fix[atom->extra_grow[iextra]]->copy_arrays(i,j,delflag);
-
-  atomKK->modified(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-                MASK_MASK | IMAGE_MASK | DPDTHETA_MASK |
-                UCG_MASK | UCGNEW_MASK |
-                UCOND_MASK | UMECH_MASK | UCHEM_MASK | DVECTOR_MASK);
-}
-
 /* ---------------------------------------------------------------------- */
 
 template<class DeviceType,int PBC_FLAG,int TRICLINIC>
@@ -547,209 +491,6 @@ void AtomVecDPDKokkos::unpack_comm_kokkos(const int &n, const int &first,
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecDPDKokkos::pack_comm(int n, int *list, double *buf,
-                             int pbc_flag, int *pbc)
-{
-  int i,j,m;
-  double dx,dy,dz;
-
-  atomKK->sync(Host,X_MASK|DPDTHETA_MASK|UCOND_MASK|UMECH_MASK|UCHEM_MASK);
-
-  m = 0;
-  if (pbc_flag == 0) {
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0);
-      buf[m++] = h_x(j,1);
-      buf[m++] = h_x(j,2);
-      buf[m++] = h_dpdTheta[j];
-      buf[m++] = h_uCond[j];
-      buf[m++] = h_uMech[j];
-      buf[m++] = h_uChem[j];
-    }
-  } else {
-    if (domain->triclinic == 0) {
-      dx = pbc[0]*domain->xprd;
-      dy = pbc[1]*domain->yprd;
-      dz = pbc[2]*domain->zprd;
-    } else {
-      dx = pbc[0]*domain->xprd + pbc[5]*domain->xy + pbc[4]*domain->xz;
-      dy = pbc[1]*domain->yprd + pbc[3]*domain->yz;
-      dz = pbc[2]*domain->zprd;
-    }
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0) + dx;
-      buf[m++] = h_x(j,1) + dy;
-      buf[m++] = h_x(j,2) + dz;
-      buf[m++] = h_dpdTheta[j];
-      buf[m++] = h_uCond[j];
-      buf[m++] = h_uMech[j];
-      buf[m++] = h_uChem[j];
-    }
-  }
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecDPDKokkos::pack_comm_vel(int n, int *list, double *buf,
-                                 int pbc_flag, int *pbc)
-{
-  int i,j,m;
-  double dx,dy,dz,dvx,dvy,dvz;
-
-  atomKK->sync(Host,X_MASK|V_MASK|DPDTHETA_MASK|UCOND_MASK|UMECH_MASK|UCHEM_MASK);
-
-  m = 0;
-  if (pbc_flag == 0) {
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0);
-      buf[m++] = h_x(j,1);
-      buf[m++] = h_x(j,2);
-      buf[m++] = h_v(j,0);
-      buf[m++] = h_v(j,1);
-      buf[m++] = h_v(j,2);
-      buf[m++] = h_dpdTheta[j];
-      buf[m++] = h_uCond[j];
-      buf[m++] = h_uMech[j];
-      buf[m++] = h_uChem[j];
-    }
-  } else {
-    if (domain->triclinic == 0) {
-      dx = pbc[0]*domain->xprd;
-      dy = pbc[1]*domain->yprd;
-      dz = pbc[2]*domain->zprd;
-    } else {
-      dx = pbc[0]*domain->xprd + pbc[5]*domain->xy + pbc[4]*domain->xz;
-      dy = pbc[1]*domain->yprd + pbc[3]*domain->yz;
-      dz = pbc[2]*domain->zprd;
-    }
-    if (!deform_vremap) {
-      for (i = 0; i < n; i++) {
-        j = list[i];
-        buf[m++] = h_x(j,0) + dx;
-        buf[m++] = h_x(j,1) + dy;
-        buf[m++] = h_x(j,2) + dz;
-        buf[m++] = h_v(j,0);
-        buf[m++] = h_v(j,1);
-        buf[m++] = h_v(j,2);
-        buf[m++] = h_dpdTheta[j];
-        buf[m++] = h_uCond[j];
-        buf[m++] = h_uMech[j];
-        buf[m++] = h_uChem[j];
-      }
-    } else {
-      dvx = pbc[0]*h_rate[0] + pbc[5]*h_rate[5] + pbc[4]*h_rate[4];
-      dvy = pbc[1]*h_rate[1] + pbc[3]*h_rate[3];
-      dvz = pbc[2]*h_rate[2];
-      for (i = 0; i < n; i++) {
-        j = list[i];
-        buf[m++] = h_x(j,0) + dx;
-        buf[m++] = h_x(j,1) + dy;
-        buf[m++] = h_x(j,2) + dz;
-        if (mask[i] & deform_groupbit) {
-          buf[m++] = h_v(j,0) + dvx;
-          buf[m++] = h_v(j,1) + dvy;
-          buf[m++] = h_v(j,2) + dvz;
-        } else {
-          buf[m++] = h_v(j,0);
-          buf[m++] = h_v(j,1);
-          buf[m++] = h_v(j,2);
-        }
-        buf[m++] = h_dpdTheta(j);
-        buf[m++] = h_uCond(j);
-        buf[m++] = h_uMech(j);
-        buf[m++] = h_uChem(j);
-      }
-    }
-  }
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-void AtomVecDPDKokkos::unpack_comm(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  for (i = first; i < last; i++) {
-    h_x(i,0) = buf[m++];
-    h_x(i,1) = buf[m++];
-    h_x(i,2) = buf[m++];
-    h_dpdTheta[i] = buf[m++];
-    h_uCond[i] = buf[m++];
-    h_uMech[i] = buf[m++];
-    h_uChem[i] = buf[m++];
-  }
-
-  atomKK->modified(Host,X_MASK|DPDTHETA_MASK|UCOND_MASK|UMECH_MASK|UCHEM_MASK);
-}
-
-/* ---------------------------------------------------------------------- */
-
-void AtomVecDPDKokkos::unpack_comm_vel(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  for (i = first; i < last; i++) {
-    h_x(i,0) = buf[m++];
-    h_x(i,1) = buf[m++];
-    h_x(i,2) = buf[m++];
-    h_v(i,0) = buf[m++];
-    h_v(i,1) = buf[m++];
-    h_v(i,2) = buf[m++];
-    h_dpdTheta[i] = buf[m++];
-    h_uCond[i] = buf[m++];
-    h_uMech[i] = buf[m++];
-    h_uChem[i] = buf[m++];
-  }
-
-  atomKK->modified(Host,X_MASK|V_MASK|DPDTHETA_MASK|UCOND_MASK|UMECH_MASK|UCHEM_MASK);
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecDPDKokkos::pack_reverse(int n, int first, double *buf)
-{
-  if (n > 0)
-    atomKK->sync(Host,F_MASK);
-
-  int m = 0;
-  const int last = first + n;
-  for (int i = first; i < last; i++) {
-    buf[m++] = h_f(i,0);
-    buf[m++] = h_f(i,1);
-    buf[m++] = h_f(i,2);
-  }
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-void AtomVecDPDKokkos::unpack_reverse(int n, int *list, double *buf)
-{
-  if (n > 0) {
-    atomKK->sync(Host,F_MASK);
-    atomKK->modified(Host,F_MASK);
-  }
-
-  int m = 0;
-  for (int i = 0; i < n; i++) {
-    const int j = list[i];
-    h_f(j,0) += buf[m++];
-    h_f(j,1) += buf[m++];
-    h_f(j,2) += buf[m++];
-  }
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType,int PBC_FLAG>
 struct AtomVecDPDKokkos_PackBorder {
   typedef DeviceType device_type;
@@ -871,206 +612,6 @@ int AtomVecDPDKokkos::pack_border_kokkos(int n, DAT::tdual_int_2d k_sendlist, DA
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecDPDKokkos::pack_border(int n, int *list, double *buf,
-                               int pbc_flag, int *pbc)
-{
-  int i,j,m;
-  double dx,dy,dz;
-
-  atomKK->sync(Host,ALL_MASK);
-
-  m = 0;
-  if (pbc_flag == 0) {
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0);
-      buf[m++] = h_x(j,1);
-      buf[m++] = h_x(j,2);
-      buf[m++] = ubuf(h_tag(j)).d;
-      buf[m++] = ubuf(h_type(j)).d;
-      buf[m++] = ubuf(h_mask(j)).d;
-      buf[m++] = h_dpdTheta(j);
-      buf[m++] = h_uCond(j);
-      buf[m++] = h_uMech(j);
-      buf[m++] = h_uChem(j);
-      buf[m++] = h_uCG(j);
-      buf[m++] = h_uCGnew(j);
-    }
-  } else {
-    if (domain->triclinic == 0) {
-      dx = pbc[0]*domain->xprd;
-      dy = pbc[1]*domain->yprd;
-      dz = pbc[2]*domain->zprd;
-    } else {
-      dx = pbc[0];
-      dy = pbc[1];
-      dz = pbc[2];
-    }
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0) + dx;
-      buf[m++] = h_x(j,1) + dy;
-      buf[m++] = h_x(j,2) + dz;
-      buf[m++] = ubuf(h_tag(j)).d;
-      buf[m++] = ubuf(h_type(j)).d;
-      buf[m++] = ubuf(h_mask(j)).d;
-      buf[m++] = h_dpdTheta(j);
-      buf[m++] = h_uCond(j);
-      buf[m++] = h_uMech(j);
-      buf[m++] = h_uChem(j);
-      buf[m++] = h_uCG(j);
-      buf[m++] = h_uCGnew(j);
-    }
-  }
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->pack_border(n,list,&buf[m]);
-
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecDPDKokkos::pack_border_vel(int n, int *list, double *buf,
-                                   int pbc_flag, int *pbc)
-{
-  int i,j,m;
-  double dx,dy,dz,dvx,dvy,dvz;
-
-  atomKK->sync(Host,ALL_MASK);
-
-  m = 0;
-  if (pbc_flag == 0) {
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0);
-      buf[m++] = h_x(j,1);
-      buf[m++] = h_x(j,2);
-      buf[m++] = ubuf(h_tag(j)).d;
-      buf[m++] = ubuf(h_type(j)).d;
-      buf[m++] = ubuf(h_mask(j)).d;
-      buf[m++] = h_v(j,0);
-      buf[m++] = h_v(j,1);
-      buf[m++] = h_v(j,2);
-      buf[m++] = h_dpdTheta(j);
-      buf[m++] = h_uCond(j);
-      buf[m++] = h_uMech(j);
-      buf[m++] = h_uChem(j);
-      buf[m++] = h_uCG(j);
-      buf[m++] = h_uCGnew(j);
-    }
-  } else {
-    if (domain->triclinic == 0) {
-      dx = pbc[0]*domain->xprd;
-      dy = pbc[1]*domain->yprd;
-      dz = pbc[2]*domain->zprd;
-    } else {
-      dx = pbc[0];
-      dy = pbc[1];
-      dz = pbc[2];
-    }
-    if (!deform_vremap) {
-      for (i = 0; i < n; i++) {
-        j = list[i];
-        buf[m++] = h_x(j,0) + dx;
-        buf[m++] = h_x(j,1) + dy;
-        buf[m++] = h_x(j,2) + dz;
-        buf[m++] = ubuf(h_tag(j)).d;
-        buf[m++] = ubuf(h_type(j)).d;
-        buf[m++] = ubuf(h_mask(j)).d;
-        buf[m++] = h_v(j,0);
-        buf[m++] = h_v(j,1);
-        buf[m++] = h_v(j,2);
-        buf[m++] = h_dpdTheta(j);
-        buf[m++] = h_uCond(j);
-        buf[m++] = h_uMech(j);
-        buf[m++] = h_uChem(j);
-        buf[m++] = h_uCG(j);
-        buf[m++] = h_uCGnew(j);
-      }
-    } else {
-      dvx = pbc[0]*h_rate[0] + pbc[5]*h_rate[5] + pbc[4]*h_rate[4];
-      dvy = pbc[1]*h_rate[1] + pbc[3]*h_rate[3];
-      dvz = pbc[2]*h_rate[2];
-      for (i = 0; i < n; i++) {
-        j = list[i];
-        buf[m++] = h_x(j,0) + dx;
-        buf[m++] = h_x(j,1) + dy;
-        buf[m++] = h_x(j,2) + dz;
-        buf[m++] = ubuf(h_tag(j)).d;
-        buf[m++] = ubuf(h_type(j)).d;
-        buf[m++] = ubuf(h_mask(j)).d;
-        if (mask[i] & deform_groupbit) {
-          buf[m++] = h_v(j,0) + dvx;
-          buf[m++] = h_v(j,1) + dvy;
-          buf[m++] = h_v(j,2) + dvz;
-        } else {
-          buf[m++] = h_v(j,0);
-          buf[m++] = h_v(j,1);
-          buf[m++] = h_v(j,2);
-        }
-        buf[m++] = h_dpdTheta(j);
-        buf[m++] = h_uCond(j);
-        buf[m++] = h_uMech(j);
-        buf[m++] = h_uChem(j);
-        buf[m++] = h_uCG(j);
-        buf[m++] = h_uCGnew(j);
-      }
-    }
-  }
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->pack_border(n,list,&buf[m]);
-
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecDPDKokkos::pack_comm_hybrid(int n, int *list, double *buf)
-{
-  int i,j,m;
-
-  atomKK->sync(Host,DPDTHETA_MASK | UCOND_MASK |
-            UMECH_MASK | UCHEM_MASK);
-
-  m = 0;
-  for (i = 0; i < n; i++) {
-    j = list[i];
-    buf[m++] = h_dpdTheta[j];
-    buf[m++] = h_uCond[j];
-    buf[m++] = h_uMech[j];
-    buf[m++] = h_uChem[j];
-  }
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecDPDKokkos::pack_border_hybrid(int n, int *list, double *buf)
-{
-  int i,j,m;
-
-  atomKK->sync(Host,DPDTHETA_MASK | UCOND_MASK |
-            UMECH_MASK | UCHEM_MASK | UCG_MASK | UCGNEW_MASK);
-
-  m = 0;
-  for (i = 0; i < n; i++) {
-    j = list[i];
-    buf[m++] = h_dpdTheta[j];
-    buf[m++] = h_uCond[j];
-    buf[m++] = h_uMech[j];
-    buf[m++] = h_uChem[j];
-    buf[m++] = h_uCG[j];
-    buf[m++] = h_uCGnew[j];
-  }
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType>
 struct AtomVecDPDKokkos_UnpackBorder {
   typedef DeviceType device_type;
@@ -1152,123 +693,6 @@ void AtomVecDPDKokkos::unpack_border_kokkos(const int &n, const int &first,
 
 /* ---------------------------------------------------------------------- */
 
-void AtomVecDPDKokkos::unpack_border(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  while (last > nmax) grow(0);
-
-  for (i = first; i < last; i++) {
-    h_x(i,0) = buf[m++];
-    h_x(i,1) = buf[m++];
-    h_x(i,2) = buf[m++];
-    h_tag(i) =  (tagint)  ubuf(buf[m++]).i;
-    h_type(i) = (int) ubuf(buf[m++]).i;
-    h_mask(i) = (int) ubuf(buf[m++]).i;
-    h_dpdTheta(i) = buf[m++];
-    h_uCond(i) = buf[m++];
-    h_uMech(i) = buf[m++];
-    h_uChem(i) = buf[m++];
-    h_uCG(i) = buf[m++];
-    h_uCGnew(i) = buf[m++];
-  }
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->
-        unpack_border(n,first,&buf[m]);
-
-  atomKK->modified(Host,X_MASK|TAG_MASK|TYPE_MASK|MASK_MASK|
-                DPDTHETA_MASK|UCOND_MASK|UMECH_MASK|UCHEM_MASK|
-                UCG_MASK|UCGNEW_MASK|DVECTOR_MASK);
-}
-
-/* ---------------------------------------------------------------------- */
-
-void AtomVecDPDKokkos::unpack_border_vel(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  while (last > nmax) grow(0);
-
-  for (i = first; i < last; i++) {
-    h_x(i,0) = buf[m++];
-    h_x(i,1) = buf[m++];
-    h_x(i,2) = buf[m++];
-    h_tag(i) =  (tagint)  ubuf(buf[m++]).i;
-    h_type(i) = (int) ubuf(buf[m++]).i;
-    h_mask(i) = (int) ubuf(buf[m++]).i;
-    h_v(i,0) = buf[m++];
-    h_v(i,1) = buf[m++];
-    h_v(i,2) = buf[m++];
-    h_dpdTheta(i) = buf[m++];
-    h_uCond(i) = buf[m++];
-    h_uMech(i) = buf[m++];
-    h_uChem(i) = buf[m++];
-    h_uCG(i) = buf[m++];
-    h_uCGnew(i) = buf[m++];
-  }
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->
-        unpack_border(n,first,&buf[m]);
-
-  atomKK->modified(Host,X_MASK|V_MASK|TAG_MASK|TYPE_MASK|MASK_MASK|
-                DPDTHETA_MASK|UCOND_MASK|UMECH_MASK|UCHEM_MASK|
-                UCG_MASK|UCGNEW_MASK|DVECTOR_MASK);
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecDPDKokkos::unpack_comm_hybrid(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  for (i = first; i < last; i++) {
-    h_dpdTheta(i) = buf[m++];
-    h_uCond(i) = buf[m++];
-    h_uMech(i) = buf[m++];
-    h_uChem(i) = buf[m++];
-  }
-
-  atomKK->modified(Host,DPDTHETA_MASK | UCOND_MASK |
-                UMECH_MASK | UCHEM_MASK );
-
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecDPDKokkos::unpack_border_hybrid(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  for (i = first; i < last; i++) {
-    h_dpdTheta(i) = buf[m++];
-    h_uCond(i) = buf[m++];
-    h_uMech(i) = buf[m++];
-    h_uChem(i) = buf[m++];
-    h_uCG(i) = buf[m++];
-    h_uCGnew(i) = buf[m++];
-  }
-
-  atomKK->modified(Host,DPDTHETA_MASK | UCOND_MASK |
-                UMECH_MASK | UCHEM_MASK | UCG_MASK | UCGNEW_MASK);
-
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType>
 struct AtomVecDPDKokkos_PackExchangeFunctor {
   typedef DeviceType device_type;
@@ -1401,41 +825,6 @@ int AtomVecDPDKokkos::pack_exchange_kokkos(const int &nsend,DAT::tdual_xfloat_2d
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecDPDKokkos::pack_exchange(int i, double *buf)
-{
-  atomKK->sync(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-            MASK_MASK | IMAGE_MASK| DPDTHETA_MASK | UCOND_MASK |
-            UMECH_MASK | UCHEM_MASK | UCG_MASK | UCGNEW_MASK |
-            DVECTOR_MASK);
-
-  int m = 1;
-  buf[m++] = h_x(i,0);
-  buf[m++] = h_x(i,1);
-  buf[m++] = h_x(i,2);
-  buf[m++] = h_v(i,0);
-  buf[m++] = h_v(i,1);
-  buf[m++] = h_v(i,2);
-  buf[m++] = ubuf(h_tag(i)).d;
-  buf[m++] = ubuf(h_type(i)).d;
-  buf[m++] = ubuf(h_mask(i)).d;
-  buf[m++] = ubuf(h_image(i)).d;
-  buf[m++] = h_dpdTheta[i];
-  buf[m++] = h_uCond[i];
-  buf[m++] = h_uMech[i];
-  buf[m++] = h_uChem[i];
-  buf[m++] = h_uCG[i];
-  buf[m++] = h_uCGnew[i];
-
-  if (atom->nextra_grow)
-    for (int iextra = 0; iextra < atom->nextra_grow; iextra++)
-      m += modify->fix[atom->extra_grow[iextra]]->pack_exchange(i,&buf[m]);
-
-  buf[0] = m;
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType>
 struct AtomVecDPDKokkos_UnpackExchangeFunctor {
   typedef DeviceType device_type;
@@ -1531,334 +920,6 @@ int AtomVecDPDKokkos::unpack_exchange_kokkos(DAT::tdual_xfloat_2d &k_buf,int nre
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecDPDKokkos::unpack_exchange(double *buf)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) grow(0);
-
-  int m = 1;
-  h_x(nlocal,0) = buf[m++];
-  h_x(nlocal,1) = buf[m++];
-  h_x(nlocal,2) = buf[m++];
-  h_v(nlocal,0) = buf[m++];
-  h_v(nlocal,1) = buf[m++];
-  h_v(nlocal,2) = buf[m++];
-  h_tag(nlocal) = (tagint) ubuf(buf[m++]).i;
-  h_type(nlocal) = (int) ubuf(buf[m++]).i;
-  h_mask(nlocal) = (int) ubuf(buf[m++]).i;
-  h_image(nlocal) = (imageint) ubuf(buf[m++]).i;
-  h_dpdTheta[nlocal] = buf[m++];
-  h_uCond[nlocal] = buf[m++];
-  h_uMech[nlocal] = buf[m++];
-  h_uChem[nlocal] = buf[m++];
-  h_uCG[nlocal] = buf[m++];
-  h_uCGnew[nlocal] = buf[m++];
-
-  if (atom->nextra_grow)
-    for (int iextra = 0; iextra < atom->nextra_grow; iextra++)
-      m += modify->fix[atom->extra_grow[iextra]]->
-        unpack_exchange(nlocal,&buf[m]);
-
-  atomKK->modified(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-           MASK_MASK | IMAGE_MASK| DPDTHETA_MASK | UCOND_MASK |
-           UMECH_MASK | UCHEM_MASK | UCG_MASK | UCGNEW_MASK |
-           DVECTOR_MASK);
-
-  atom->nlocal++;
-  return m;
-}
-
-/* ----------------------------------------------------------------------
-   size of restart data for all atoms owned by this proc
-   include extra data stored by fixes
-------------------------------------------------------------------------- */
-
-int AtomVecDPDKokkos::size_restart()
-{
-  int i;
-
-  int nlocal = atom->nlocal;
-  int n = 15 * nlocal; // 11 + dpdTheta + uCond + uMech + uChem
-
-  if (atom->nextra_restart)
-    for (int iextra = 0; iextra < atom->nextra_restart; iextra++)
-      for (i = 0; i < nlocal; i++)
-        n += modify->fix[atom->extra_restart[iextra]]->size_restart(i);
-
-  return n;
-}
-
-/* ----------------------------------------------------------------------
-   pack atom I's data for restart file including extra quantities
-   xyz must be 1st 3 values, so that read_restart can test on them
-   molecular types may be negative, but write as positive
-------------------------------------------------------------------------- */
-
-int AtomVecDPDKokkos::pack_restart(int i, double *buf)
-{
-  atomKK->sync(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-            MASK_MASK | IMAGE_MASK | DPDTHETA_MASK |
-            UCOND_MASK | UMECH_MASK | UCHEM_MASK | DVECTOR_MASK);
-
-  int m = 1;
-  buf[m++] = h_x(i,0);
-  buf[m++] = h_x(i,1);
-  buf[m++] = h_x(i,2);
-  buf[m++] = ubuf(h_tag(i)).d;
-  buf[m++] = ubuf(h_type(i)).d;
-  buf[m++] = ubuf(h_mask(i)).d;
-  buf[m++] = ubuf(h_image(i)).d;
-  buf[m++] = h_v(i,0);
-  buf[m++] = h_v(i,1);
-  buf[m++] = h_v(i,2);
-  buf[m++] = h_dpdTheta[i];
-  buf[m++] = h_uCond[i];
-  buf[m++] = h_uMech[i];
-  buf[m++] = h_uChem[i];
-
-  if (atom->nextra_restart)
-    for (int iextra = 0; iextra < atom->nextra_restart; iextra++)
-      m += modify->fix[atom->extra_restart[iextra]]->pack_restart(i,&buf[m]);
-
-  buf[0] = m;
-  return m;
-}
-
-/* ----------------------------------------------------------------------
-   unpack data for one atom from restart file including extra quantities
-------------------------------------------------------------------------- */
-
-int AtomVecDPDKokkos::unpack_restart(double *buf)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) {
-    grow(0);
-    if (atom->nextra_store)
-      memory->grow(atom->extra,nmax,atom->nextra_store,"atom:extra");
-  }
-
-  int m = 1;
-  h_x(nlocal,0) = buf[m++];
-  h_x(nlocal,1) = buf[m++];
-  h_x(nlocal,2) = buf[m++];
-  h_tag(nlocal) = (tagint) ubuf(buf[m++]).i;
-  h_type(nlocal) = (int) ubuf(buf[m++]).i;
-  h_mask(nlocal) = (int) ubuf(buf[m++]).i;
-  h_image(nlocal) = (imageint) ubuf(buf[m++]).i;
-  h_v(nlocal,0) = buf[m++];
-  h_v(nlocal,1) = buf[m++];
-  h_v(nlocal,2) = buf[m++];
-  h_dpdTheta[nlocal] = buf[m++];
-  h_uCond[nlocal] = buf[m++];
-  h_uMech[nlocal] = buf[m++];
-  h_uChem[nlocal] = buf[m++];
-  h_uCG[nlocal] = 0.0;
-  h_uCGnew[nlocal] = 0.0;
-
-  double **extra = atom->extra;
-  if (atom->nextra_store) {
-    int size = static_cast<int> (buf[0]) - m;
-    for (int i = 0; i < size; i++) extra[nlocal][i] = buf[m++];
-  }
-
-  atomKK->modified(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-                MASK_MASK | IMAGE_MASK | DPDTHETA_MASK |
-                UCG_MASK | UCGNEW_MASK |
-                UCOND_MASK | UMECH_MASK | UCHEM_MASK | DVECTOR_MASK);
-
-  atom->nlocal++;
-  return m;
-}
-
-/* ----------------------------------------------------------------------
-   create one atom of itype at coord
-   set other values to defaults
-------------------------------------------------------------------------- */
-
-void AtomVecDPDKokkos::create_atom(int itype, double *coord)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) {
-    //if(nlocal>2) printf("typeA: %i %i\n",type[0],type[1]);
-    atomKK->modified(Host,ALL_MASK);
-    grow(0);
-    //if(nlocal>2) printf("typeB: %i %i\n",type[0],type[1]);
-  }
-  atomKK->modified(Host,ALL_MASK);
-
-  tag[nlocal] = 0;
-  type[nlocal] = itype;
-  h_x(nlocal,0) = coord[0];
-  h_x(nlocal,1) = coord[1];
-  h_x(nlocal,2) = coord[2];
-  h_mask[nlocal] = 1;
-  h_image[nlocal] = ((tagint) IMGMAX << IMG2BITS) |
-    ((tagint) IMGMAX << IMGBITS) | IMGMAX;
-  h_v(nlocal,0) = 0.0;
-  h_v(nlocal,1) = 0.0;
-  h_v(nlocal,2) = 0.0;
-  h_rho[nlocal] = 0.0;
-  h_dpdTheta[nlocal] = 0.0;
-  h_uCond[nlocal] = 0.0;
-  h_uMech[nlocal] = 0.0;
-  h_uChem[nlocal] = 0.0;
-  h_uCG[nlocal] = 0.0;
-  h_uCGnew[nlocal] = 0.0;
-  h_duChem[nlocal] = 0.0;
-
-  //atomKK->modified(Host,TAG_MASK|TYPE_MASK|DPDTHETA_MASK|X_MASK|IMAGE_MASK|
-  //                      MASK_MASK|V_MASK|DPDRHO_MASK|UCOND_MASK|UMECH_MASK|
-  //                      UCHEM_MASK|UCG_MASK|UCGNEW_MASK);
-
-  atom->nlocal++;
-}
-
-/* ----------------------------------------------------------------------
-   unpack one line from Atoms section of data file
-   initialize other atom quantities
-------------------------------------------------------------------------- */
-
-void AtomVecDPDKokkos::data_atom(double *coord, imageint imagetmp,
-                                 const std::vector<std::string> &values, std::string &extract)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) grow(0);
-
-  h_tag[nlocal] = utils::tnumeric(FLERR,values[0],true,lmp);
-  h_type[nlocal] = utils::inumeric(FLERR,values[1],true,lmp);
-  extract = values[1];
-  if (type[nlocal] <= 0 || type[nlocal] > atom->ntypes)
-    error->one(FLERR,"Invalid atom type in Atoms section of data file");
-
-  h_dpdTheta[nlocal] = utils::numeric(FLERR,values[2],true,lmp);
-  if (h_dpdTheta[nlocal] <= 0)
-    error->one(FLERR,"Internal temperature in Atoms section of date file must be > zero");
-
-  h_x(nlocal,0) = coord[0];
-  h_x(nlocal,1) = coord[1];
-  h_x(nlocal,2) = coord[2];
-
-  h_image[nlocal] = imagetmp;
-
-  h_mask[nlocal] = 1;
-  h_v(nlocal,0) = 0.0;
-  h_v(nlocal,1) = 0.0;
-  h_v(nlocal,2) = 0.0;
-
-  h_rho[nlocal] = 0.0;
-  h_uCond[nlocal] = 0.0;
-  h_uMech[nlocal] = 0.0;
-  h_uChem[nlocal] = 0.0;
-  h_uCG[nlocal] = 0.0;
-  h_uCGnew[nlocal] = 0.0;
-
-  atomKK->modified(Host,ALL_MASK);
-
-  atom->nlocal++;
-}
-
-/* ----------------------------------------------------------------------
-   unpack hybrid quantities from one line in Atoms section of data file
-   initialize other atom quantities for this sub-style
-------------------------------------------------------------------------- */
-
-int AtomVecDPDKokkos::data_atom_hybrid(int nlocal, const std::vector<std::string> &values,
-                                       int offset)
-{
-  h_dpdTheta(nlocal) = utils::numeric(FLERR,values[offset],true,lmp);
-
-  atomKK->modified(Host,DPDTHETA_MASK);
-
-  return 1;
-}
-
-/* ----------------------------------------------------------------------
-   pack atom info for data file including 3 image flags
-------------------------------------------------------------------------- */
-
-void AtomVecDPDKokkos::pack_data(double **buf)
-{
-  atomKK->sync(Host,TAG_MASK|TYPE_MASK|DPDTHETA_MASK|X_MASK|IMAGE_MASK);
-
-  int nlocal = atom->nlocal;
-  for (int i = 0; i < nlocal; i++) {
-    buf[i][0] = ubuf(h_tag(i)).d;
-    buf[i][1] = ubuf(h_type(i)).d;
-    buf[i][2] = h_dpdTheta(i);
-    buf[i][3] = h_x(i,0);
-    buf[i][4] = h_x(i,1);
-    buf[i][5] = h_x(i,2);
-    buf[i][6] = (h_image[i] & IMGMASK) - IMGMAX;
-    buf[i][7] = (h_image[i] >> IMGBITS & IMGMASK) - IMGMAX;
-    buf[i][8] = (h_image[i] >> IMG2BITS) - IMGMAX;
-  }
-}
-
-/* ----------------------------------------------------------------------
-   pack hybrid atom info for data file
-------------------------------------------------------------------------- */
-
-int AtomVecDPDKokkos::pack_data_hybrid(int i, double *buf)
-{
-  atomKK->sync(Host,DPDTHETA_MASK);
-
-  buf[0] = h_dpdTheta(i);
-  return 1;
-}
-
-/* ----------------------------------------------------------------------
-   write atom info to data file including 3 image flags
-------------------------------------------------------------------------- */
-
-void AtomVecDPDKokkos::write_data(FILE *fp, int n, double **buf)
-{
-  for (int i = 0; i < n; i++)
-    fprintf(fp,TAGINT_FORMAT " %d %-1.16e %-1.16e %-1.16e %-1.16e %d %d %d\n",
-            (tagint) ubuf(buf[i][0]).i,(int) ubuf(buf[i][1]).i,
-            buf[i][2],buf[i][3],buf[i][4],buf[i][5],
-            (int) ubuf(buf[i][6]).i,(int) ubuf(buf[i][7]).i,
-            (int) ubuf(buf[i][8]).i);
-}
-
-/* ----------------------------------------------------------------------
-   write hybrid atom info to data file
-------------------------------------------------------------------------- */
-
-int AtomVecDPDKokkos::write_data_hybrid(FILE *fp, double *buf)
-{
-  fprintf(fp," %-1.16e",buf[0]);
-  return 1;
-}
-
-/* ----------------------------------------------------------------------
-   return # of bytes of allocated memory
-------------------------------------------------------------------------- */
-
-double AtomVecDPDKokkos::memory_usage()
-{
-  double bytes = 0;
-
-  if (atom->memcheck("tag")) bytes += memory->usage(tag,nmax);
-  if (atom->memcheck("type")) bytes += memory->usage(type,nmax);
-  if (atom->memcheck("mask")) bytes += memory->usage(mask,nmax);
-  if (atom->memcheck("image")) bytes += memory->usage(image,nmax);
-  if (atom->memcheck("x")) bytes += memory->usage(x,nmax,3);
-  if (atom->memcheck("v")) bytes += memory->usage(v,nmax,3);
-  if (atom->memcheck("f")) bytes += memory->usage(f,nmax*commKK->nthreads,3);
-  if (atom->memcheck("rho")) bytes += memory->usage(rho,nmax);
-  if (atom->memcheck("dpdTheta")) bytes += memory->usage(dpdTheta,nmax);
-  if (atom->memcheck("uCond")) bytes += memory->usage(uCond,nmax);
-  if (atom->memcheck("uMech")) bytes += memory->usage(uMech,nmax);
-  if (atom->memcheck("uChem")) bytes += memory->usage(uChem,nmax);
-  if (atom->memcheck("uCG")) bytes += memory->usage(uCG,nmax);
-  if (atom->memcheck("uCGnew")) bytes += memory->usage(uCGnew,nmax);
-  if (atom->memcheck("duChem")) bytes += memory->usage(duChem,nmax);
-
-  return bytes;
-}
-
-/* ---------------------------------------------------------------------- */
-
 void AtomVecDPDKokkos::sync(ExecutionSpace space, unsigned int mask)
 {
   if (space == Device) {
@@ -2011,4 +1072,3 @@ void AtomVecDPDKokkos::modified(ExecutionSpace space, unsigned int mask)
     if (mask & DVECTOR_MASK) atomKK->k_dvector.modify<LMPHostType>();
   }
 }
-

src/KOKKOS/atom_vec_dpd_kokkos.h

@@ -24,44 +24,16 @@ AtomStyle(dpd/kk/host,AtomVecDPDKokkos);
 #define LMP_ATOM_VEC_DPD_KOKKOS_H
 
 #include "atom_vec_kokkos.h"
+#include "atom_vec_dpd.h"
 #include "kokkos_type.h"
 
 namespace LAMMPS_NS {
 
-class AtomVecDPDKokkos : public AtomVecKokkos {
+class AtomVecDPDKokkos : public AtomVecKokkos, public AtomVecDPD {
  public:
   AtomVecDPDKokkos(class LAMMPS *);
 
   void grow(int) override;
-  void copy(int, int, int) override;
-  int pack_comm(int, int *, double *, int, int *) override;
-  int pack_comm_vel(int, int *, double *, int, int *) override;
-  int pack_comm_hybrid(int, int *, double *) override;
-  void unpack_comm(int, int, double *) override;
-  void unpack_comm_vel(int, int, double *) override;
-  int unpack_comm_hybrid(int, int, double *) override;
-  int pack_reverse(int, int, double *) override;
-  void unpack_reverse(int, int *, double *) override;
-  int pack_border(int, int *, double *, int, int *) override;
-  int pack_border_vel(int, int *, double *, int, int *) override;
-  int pack_border_hybrid(int, int *, double *) override;
-  void unpack_border(int, int, double *) override;
-  void unpack_border_vel(int, int, double *) override;
-  int unpack_border_hybrid(int, int, double *) override;
-  int pack_exchange(int, double *) override;
-  int unpack_exchange(double *) override;
-  int size_restart() override;
-  int pack_restart(int, double *) override;
-  int unpack_restart(double *) override;
-  void create_atom(int, double *) override;
-  void data_atom(double *, imageint, const std::vector<std::string> &, std::string &) override;
-  int data_atom_hybrid(int, const std::vector<std::string> &, int) override;
-  void pack_data(double **) override;
-  int pack_data_hybrid(int, double *) override;
-  void write_data(FILE *, int, double **) override;
-  int write_data_hybrid(FILE *, double *) override;
-  double memory_usage() override;
-
   void grow_pointers() override;
   int pack_comm_kokkos(const int &n, const DAT::tdual_int_2d &k_sendlist,
                        const int & iswap,
@@ -112,8 +84,6 @@ class AtomVecDPDKokkos : public AtomVecKokkos {
   DAT::t_x_array d_x;
   DAT::t_v_array d_v;
   DAT::t_f_array d_f;
-
-  DAT::tdual_int_1d k_count;
 };
 
 }

src/KOKKOS/atom_vec_full_kokkos.cpp

@@ -27,26 +27,10 @@ using namespace LAMMPS_NS;
 
 /* ---------------------------------------------------------------------- */
 
-AtomVecFullKokkos::AtomVecFullKokkos(LAMMPS *lmp) : AtomVecKokkos(lmp)
+AtomVecFullKokkos::AtomVecFullKokkos(LAMMPS *lmp) : AtomVec(lmp),
+AtomVecKokkos(lmp), AtomVecFull(lmp)
 {
-  molecular = Atom::MOLECULAR;
-  bonds_allow = angles_allow = dihedrals_allow = impropers_allow = 1;
-  mass_type = PER_TYPE;
 
-  comm_x_only = comm_f_only = 1;
-  size_forward = 3;
-  size_reverse = 3;
-  size_border = 8;
-  size_velocity = 3;
-  size_data_atom = 7;
-  size_data_vel = 4;
-  xcol_data = 5;
-
-  atom->molecule_flag = atom->q_flag = 1;
-
-  k_count = DAT::tdual_int_1d("atom::k_count",1);
-  atomKK = (AtomKokkos *) atom;
-  commKK = (CommKokkos *) comm;
 }
 
 /* ----------------------------------------------------------------------
@@ -241,71 +225,6 @@ void AtomVecFullKokkos::grow_pointers()
   h_improper_atom4 = atomKK->k_improper_atom4.h_view;
 }
 
-/* ----------------------------------------------------------------------
-   copy atom I info to atom J
-------------------------------------------------------------------------- */
-
-void AtomVecFullKokkos::copy(int i, int j, int delflag)
-{
-  int k;
-
-  h_tag[j] = h_tag[i];
-  h_type[j] = h_type[i];
-  mask[j] = mask[i];
-  h_image[j] = h_image[i];
-  h_x(j,0) = h_x(i,0);
-  h_x(j,1) = h_x(i,1);
-  h_x(j,2) = h_x(i,2);
-  h_v(j,0) = h_v(i,0);
-  h_v(j,1) = h_v(i,1);
-  h_v(j,2) = h_v(i,2);
-
-  h_q[j] = h_q[i];
-  h_molecule(j) = h_molecule(i);
-
-  h_num_bond(j) = h_num_bond(i);
-  for (k = 0; k < h_num_bond(j); k++) {
-    h_bond_type(j,k) = h_bond_type(i,k);
-    h_bond_atom(j,k) = h_bond_atom(i,k);
-  }
-
-  h_nspecial(j,0) = h_nspecial(i,0);
-  h_nspecial(j,1) = h_nspecial(i,1);
-  h_nspecial(j,2) = h_nspecial(i,2);
-  for (k = 0; k < h_nspecial(j,2); k++)
-    h_special(j,k) = h_special(i,k);
-
-  h_num_angle(j) = h_num_angle(i);
-  for (k = 0; k < h_num_angle(j); k++) {
-    h_angle_type(j,k) = h_angle_type(i,k);
-    h_angle_atom1(j,k) = h_angle_atom1(i,k);
-    h_angle_atom2(j,k) = h_angle_atom2(i,k);
-    h_angle_atom3(j,k) = h_angle_atom3(i,k);
-  }
-
-  h_num_dihedral(j) = h_num_dihedral(i);
-  for (k = 0; k < h_num_dihedral(j); k++) {
-    h_dihedral_type(j,k) = h_dihedral_type(i,k);
-    h_dihedral_atom1(j,k) = h_dihedral_atom1(i,k);
-    h_dihedral_atom2(j,k) = h_dihedral_atom2(i,k);
-    h_dihedral_atom3(j,k) = h_dihedral_atom3(i,k);
-    h_dihedral_atom4(j,k) = h_dihedral_atom4(i,k);
-  }
-
-  h_num_improper(j) = h_num_improper(i);
-  for (k = 0; k < h_num_improper(j); k++) {
-    h_improper_type(j,k) = h_improper_type(i,k);
-    h_improper_atom1(j,k) = h_improper_atom1(i,k);
-    h_improper_atom2(j,k) = h_improper_atom2(i,k);
-    h_improper_atom3(j,k) = h_improper_atom3(i,k);
-    h_improper_atom4(j,k) = h_improper_atom4(i,k);
-  }
-
-  if (atom->nextra_grow)
-    for (int iextra = 0; iextra < atom->nextra_grow; iextra++)
-      modify->fix[atom->extra_grow[iextra]]->copy_arrays(i,j,delflag);
-}
-
 /* ---------------------------------------------------------------------- */
 
 template<class DeviceType,int PBC_FLAG>
@@ -413,155 +332,6 @@ int AtomVecFullKokkos::pack_border_kokkos(int n, DAT::tdual_int_2d k_sendlist,
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecFullKokkos::pack_border(int n, int *list, double *buf,
-                                        int pbc_flag, int *pbc)
-{
-  int i,j,m;
-  double dx,dy,dz;
-
-  m = 0;
-  if (pbc_flag == 0) {
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0);
-      buf[m++] = h_x(j,1);
-      buf[m++] = h_x(j,2);
-      buf[m++] = ubuf(h_tag(j)).d;
-      buf[m++] = ubuf(h_type(j)).d;
-      buf[m++] = ubuf(h_mask(j)).d;
-      buf[m++] = h_q(j);
-      buf[m++] = ubuf(h_molecule(j)).d;
-    }
-  } else {
-    if (domain->triclinic == 0) {
-      dx = pbc[0]*domain->xprd;
-      dy = pbc[1]*domain->yprd;
-      dz = pbc[2]*domain->zprd;
-    } else {
-      dx = pbc[0];
-      dy = pbc[1];
-      dz = pbc[2];
-    }
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0) + dx;
-      buf[m++] = h_x(j,1) + dy;
-      buf[m++] = h_x(j,2) + dz;
-      buf[m++] = ubuf(h_tag(j)).d;
-      buf[m++] = ubuf(h_type(j)).d;
-      buf[m++] = ubuf(h_mask(j)).d;
-      buf[m++] = h_q(j);
-      buf[m++] = ubuf(h_molecule(j)).d;
-    }
-  }
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->pack_border(n,list,&buf[m]);
-
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecFullKokkos::pack_border_vel(int n, int *list, double *buf,
-                                            int pbc_flag, int *pbc)
-{
-  int i,j,m;
-  double dx,dy,dz,dvx,dvy,dvz;
-
-  m = 0;
-  if (pbc_flag == 0) {
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0);
-      buf[m++] = h_x(j,1);
-      buf[m++] = h_x(j,2);
-      buf[m++] = ubuf(h_tag(j)).d;
-      buf[m++] = ubuf(h_type(j)).d;
-      buf[m++] = ubuf(h_mask(j)).d;
-      buf[m++] = h_q(j);
-      buf[m++] = ubuf(h_molecule(j)).d;
-      buf[m++] = h_v(j,0);
-      buf[m++] = h_v(j,1);
-      buf[m++] = h_v(j,2);
-    }
-  } else {
-    if (domain->triclinic == 0) {
-      dx = pbc[0]*domain->xprd;
-      dy = pbc[1]*domain->yprd;
-      dz = pbc[2]*domain->zprd;
-    } else {
-      dx = pbc[0];
-      dy = pbc[1];
-      dz = pbc[2];
-    }
-    if (!deform_vremap) {
-      for (i = 0; i < n; i++) {
-        j = list[i];
-        buf[m++] = h_x(j,0) + dx;
-        buf[m++] = h_x(j,1) + dy;
-        buf[m++] = h_x(j,2) + dz;
-        buf[m++] = ubuf(h_tag(j)).d;
-        buf[m++] = ubuf(h_type(j)).d;
-        buf[m++] = ubuf(h_mask(j)).d;
-        buf[m++] = h_q(j);
-        buf[m++] = ubuf(h_molecule(j)).d;
-        buf[m++] = h_v(j,0);
-        buf[m++] = h_v(j,1);
-        buf[m++] = h_v(j,2);
-      }
-    } else {
-      dvx = pbc[0]*h_rate[0] + pbc[5]*h_rate[5] + pbc[4]*h_rate[4];
-      dvy = pbc[1]*h_rate[1] + pbc[3]*h_rate[3];
-      dvz = pbc[2]*h_rate[2];
-      for (i = 0; i < n; i++) {
-        j = list[i];
-        buf[m++] = h_x(j,0) + dx;
-        buf[m++] = h_x(j,1) + dy;
-        buf[m++] = h_x(j,2) + dz;
-        buf[m++] = ubuf(h_tag(j)).d;
-        buf[m++] = ubuf(h_type(j)).d;
-        buf[m++] = ubuf(h_mask(j)).d;
-        buf[m++] = h_q(j);
-        buf[m++] = ubuf(h_molecule(j)).d;
-        if (mask[i] & deform_groupbit) {
-          buf[m++] = h_v(j,0) + dvx;
-          buf[m++] = h_v(j,1) + dvy;
-          buf[m++] = h_v(j,2) + dvz;
-        } else {
-          buf[m++] = h_v(j,0);
-          buf[m++] = h_v(j,1);
-          buf[m++] = h_v(j,2);
-        }
-      }
-    }
-  }
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->pack_border(n,list,&buf[m]);
-
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecFullKokkos::pack_border_hybrid(int n, int *list, double *buf)
-{
-  int i,j,m;
-
-  m = 0;
-  for (i = 0; i < n; i++) {
-    j = list[i];
-    buf[m++] = h_q(j);
-    buf[m++] = ubuf(h_molecule(j)).d;
-  }
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType>
 struct AtomVecFullKokkos_UnpackBorder {
   typedef DeviceType device_type;
@@ -628,82 +398,6 @@ void AtomVecFullKokkos::unpack_border_kokkos(const int &n, const int &first,
 
 /* ---------------------------------------------------------------------- */
 
-void AtomVecFullKokkos::unpack_border(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  while (last > nmax) grow(0);
-
-  for (i = first; i < last; i++) {
-    h_x(i,0) = buf[m++];
-    h_x(i,1) = buf[m++];
-    h_x(i,2) = buf[m++];
-    h_tag(i) =  (tagint)  ubuf(buf[m++]).i;
-    h_type(i) = (int) ubuf(buf[m++]).i;
-    h_mask(i) = (int) ubuf(buf[m++]).i;
-    h_q(i) = buf[m++];
-    h_molecule(i) = (tagint) ubuf(buf[m++]).i;
-  }
-
-  atomKK->modified(Host,X_MASK|TAG_MASK|TYPE_MASK|MASK_MASK|Q_MASK|MOLECULE_MASK);
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->
-        unpack_border(n,first,&buf[m]);
-}
-
-/* ---------------------------------------------------------------------- */
-
-void AtomVecFullKokkos::unpack_border_vel(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  while (last > nmax) grow(0);
-
-  for (i = first; i < last; i++) {
-    h_x(i,0) = buf[m++];
-    h_x(i,1) = buf[m++];
-    h_x(i,2) = buf[m++];
-    h_tag(i) =  (tagint)  ubuf(buf[m++]).i;
-    h_type(i) = (int) ubuf(buf[m++]).i;
-    h_mask(i) = (int) ubuf(buf[m++]).i;
-    h_q(i) = buf[m++];
-    h_molecule(i) = (tagint) ubuf(buf[m++]).i;
-    h_v(i,0) = buf[m++];
-    h_v(i,1) = buf[m++];
-    h_v(i,2) = buf[m++];
-  }
-
-  atomKK->modified(Host,X_MASK|TAG_MASK|TYPE_MASK|MASK_MASK|Q_MASK|MOLECULE_MASK|V_MASK);
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->
-        unpack_border(n,first,&buf[m]);
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecFullKokkos::unpack_border_hybrid(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  for (i = first; i < last; i++) {
-    h_q(i) = buf[m++];
-    h_molecule(i) = (tagint) ubuf(buf[m++]).i;
-  }
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType>
 struct AtomVecFullKokkos_PackExchangeFunctor {
   typedef DeviceType device_type;
@@ -983,66 +677,6 @@ int AtomVecFullKokkos::pack_exchange_kokkos(const int &nsend,DAT::tdual_xfloat_2
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecFullKokkos::pack_exchange(int i, double *buf)
-{
-  int k;
-  int m = 1;
-  buf[m++] = h_x(i,0);
-  buf[m++] = h_x(i,1);
-  buf[m++] = h_x(i,2);
-  buf[m++] = h_v(i,0);
-  buf[m++] = h_v(i,1);
-  buf[m++] = h_v(i,2);
-  buf[m++] = ubuf(h_tag(i)).d;
-  buf[m++] = ubuf(h_type(i)).d;
-  buf[m++] = ubuf(h_mask(i)).d;
-  buf[m++] = ubuf(h_image(i)).d;
-  buf[m++] = h_q(i);
-  buf[m++] = ubuf(h_molecule(i)).d;
-  buf[m++] = ubuf(h_num_bond(i)).d;
-  for (k = 0; k < h_num_bond(i); k++) {
-    buf[m++] = ubuf(h_bond_type(i,k)).d;
-    buf[m++] = ubuf(h_bond_atom(i,k)).d;
-  }
-  buf[m++] = ubuf(h_num_angle(i)).d;
-  for (k = 0; k < h_num_angle(i); k++) {
-    buf[m++] = ubuf(h_angle_type(i,k)).d;
-    buf[m++] = ubuf(h_angle_atom1(i,k)).d;
-    buf[m++] = ubuf(h_angle_atom2(i,k)).d;
-    buf[m++] = ubuf(h_angle_atom3(i,k)).d;
-  }
-  buf[m++] = ubuf(h_num_dihedral(i)).d;
-  for (k = 0; k < h_num_dihedral(i); k++) {
-    buf[m++] = ubuf(h_dihedral_type(i,k)).d;
-    buf[m++] = ubuf(h_dihedral_atom1(i,k)).d;
-    buf[m++] = ubuf(h_dihedral_atom2(i,k)).d;
-    buf[m++] = ubuf(h_dihedral_atom3(i,k)).d;
-    buf[m++] = ubuf(h_dihedral_atom4(i,k)).d;
-  }
-  buf[m++] = ubuf(h_num_improper(i)).d;
-  for (k = 0; k < h_num_improper(i); k++) {
-    buf[m++] = ubuf(h_improper_type(i,k)).d;
-    buf[m++] = ubuf(h_improper_atom1(i,k)).d;
-    buf[m++] = ubuf(h_improper_atom2(i,k)).d;
-    buf[m++] = ubuf(h_improper_atom3(i,k)).d;
-    buf[m++] = ubuf(h_improper_atom4(i,k)).d;
-  }
-  buf[m++] = ubuf(h_nspecial(i,0)).d;
-  buf[m++] = ubuf(h_nspecial(i,1)).d;
-  buf[m++] = ubuf(h_nspecial(i,2)).d;
-  for (k = 0; k < h_nspecial(i,2); k++)
-    buf[m++] = ubuf(h_special(i,k)).d;
-
-  if (atom->nextra_grow)
-    for (int iextra = 0; iextra < atom->nextra_grow; iextra++)
-      m += modify->fix[atom->extra_grow[iextra]]->pack_exchange(i,&buf[m]);
-
-  buf[0] = m;
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType>
 struct AtomVecFullKokkos_UnpackExchangeFunctor {
   typedef DeviceType device_type;
@@ -1211,457 +845,6 @@ int AtomVecFullKokkos::unpack_exchange_kokkos(DAT::tdual_xfloat_2d &k_buf,int nr
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecFullKokkos::unpack_exchange(double *buf)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) grow(0);
-  atomKK->modified(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-           MASK_MASK | IMAGE_MASK | Q_MASK | MOLECULE_MASK | BOND_MASK |
-           ANGLE_MASK | DIHEDRAL_MASK | IMPROPER_MASK | SPECIAL_MASK);
-
-  int k;
-  int m = 1;
-  h_x(nlocal,0) = buf[m++];
-  h_x(nlocal,1) = buf[m++];
-  h_x(nlocal,2) = buf[m++];
-  h_v(nlocal,0) = buf[m++];
-  h_v(nlocal,1) = buf[m++];
-  h_v(nlocal,2) = buf[m++];
-  h_tag(nlocal) = (tagint) ubuf(buf[m++]).i;
-  h_type(nlocal) = (int) ubuf(buf[m++]).i;
-  h_mask(nlocal) = (int) ubuf(buf[m++]).i;
-  h_image(nlocal) = (imageint) ubuf(buf[m++]).i;
-  h_q(nlocal) = buf[m++];
-  h_molecule(nlocal) = (tagint) ubuf(buf[m++]).i;
-
-  h_num_bond(nlocal) = (int) ubuf(buf[m++]).i;
-  for (k = 0; k < h_num_bond(nlocal); k++) {
-    h_bond_type(nlocal,k) = (int) ubuf(buf[m++]).i;
-    h_bond_atom(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-  }
-  h_num_angle(nlocal) = (int) ubuf(buf[m++]).i;
-  for (k = 0; k < h_num_angle(nlocal); k++) {
-    h_angle_type(nlocal,k) = (int) ubuf(buf[m++]).i;
-    h_angle_atom1(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-    h_angle_atom2(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-    h_angle_atom3(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-  }
-  h_num_dihedral(nlocal) = (int) ubuf(buf[m++]).i;
-  for (k = 0; k < h_num_dihedral(nlocal); k++) {
-    h_dihedral_type(nlocal,k) = (int) ubuf(buf[m++]).i;
-    h_dihedral_atom1(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-    h_dihedral_atom2(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-    h_dihedral_atom3(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-    h_dihedral_atom4(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-  }
-  h_num_improper(nlocal) = (int) ubuf(buf[m++]).i;
-  for (k = 0; k < h_num_improper(nlocal); k++) {
-    h_improper_type(nlocal,k) = (int) ubuf(buf[m++]).i;
-    h_improper_atom1(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-    h_improper_atom2(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-    h_improper_atom3(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-    h_improper_atom4(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-  }
-  h_nspecial(nlocal,0) = (int) ubuf(buf[m++]).i;
-  h_nspecial(nlocal,1) = (int) ubuf(buf[m++]).i;
-  h_nspecial(nlocal,2) = (int) ubuf(buf[m++]).i;
-  for (k = 0; k < h_nspecial(nlocal,2); k++)
-         h_special(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-
-  if (atom->nextra_grow)
-    for (int iextra = 0; iextra < atom->nextra_grow; iextra++)
-      m += modify->fix[atom->extra_grow[iextra]]->
-        unpack_exchange(nlocal,&buf[m]);
-
-  atom->nlocal++;
-  return m;
-}
-
-/* ----------------------------------------------------------------------
-   size of restart data for all atoms owned by this proc
-   include extra data stored by fixes
-------------------------------------------------------------------------- */
-
-int AtomVecFullKokkos::size_restart()
-{
-  int i;
-
-  int nlocal = atom->nlocal;
-  int n = 0;
-  for (i = 0; i < nlocal; i++)
-    n += 17 + 2*num_bond[i] + 4*num_angle[i] +
-      5*num_dihedral[i] + 5*num_improper[i];
-
-  if (atom->nextra_restart)
-    for (int iextra = 0; iextra < atom->nextra_restart; iextra++)
-      for (i = 0; i < nlocal; i++)
-        n += modify->fix[atom->extra_restart[iextra]]->size_restart(i);
-
-  return n;
-}
-
-/* ----------------------------------------------------------------------
-   pack atom I's data for restart file including extra quantities
-   xyz must be 1st 3 values, so that read_restart can test on them
-   molecular types may be negative, but write as positive
-------------------------------------------------------------------------- */
-
-int AtomVecFullKokkos::pack_restart(int i, double *buf)
-{
-  atomKK->sync(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-       MASK_MASK | IMAGE_MASK | Q_MASK | MOLECULE_MASK | BOND_MASK |
-       ANGLE_MASK | DIHEDRAL_MASK | IMPROPER_MASK | SPECIAL_MASK);
-
-  int m = 1;
-  buf[m++] = h_x(i,0);
-  buf[m++] = h_x(i,1);
-  buf[m++] = h_x(i,2);
-  buf[m++] = ubuf(h_tag(i)).d;
-  buf[m++] = ubuf(h_type(i)).d;
-  buf[m++] = ubuf(h_mask(i)).d;
-  buf[m++] = ubuf(h_image(i)).d;
-  buf[m++] = h_v(i,0);
-  buf[m++] = h_v(i,1);
-  buf[m++] = h_v(i,2);
-
-  buf[m++] = h_q(i);
-  buf[m++] = ubuf(h_molecule(i)).d;
-
-  buf[m++] = ubuf(h_num_bond(i)).d;
-  for (int k = 0; k < h_num_bond(i); k++) {
-    buf[m++] = ubuf(MAX(h_bond_type(i,k),-h_bond_type(i,k))).d;
-    buf[m++] = ubuf(h_bond_atom(i,k)).d;
-  }
-
-  buf[m++] = ubuf(h_num_angle(i)).d;
-  for (int k = 0; k < h_num_angle(i); k++) {
-    buf[m++] = ubuf(MAX(h_angle_type(i,k),-h_angle_type(i,k))).d;
-    buf[m++] = ubuf(h_angle_atom1(i,k)).d;
-    buf[m++] = ubuf(h_angle_atom2(i,k)).d;
-    buf[m++] = ubuf(h_angle_atom3(i,k)).d;
-  }
-
-  buf[m++] = ubuf(h_num_dihedral(i)).d;
-  for (int k = 0; k < h_num_dihedral(i); k++) {
-    buf[m++] = ubuf(MAX(h_dihedral_type(i,k),-h_dihedral_type(i,k))).d;
-    buf[m++] = ubuf(h_dihedral_atom1(i,k)).d;
-    buf[m++] = ubuf(h_dihedral_atom2(i,k)).d;
-    buf[m++] = ubuf(h_dihedral_atom3(i,k)).d;
-    buf[m++] = ubuf(h_dihedral_atom4(i,k)).d;
-  }
-
-  buf[m++] = ubuf(h_num_improper(i)).d;
-  for (int k = 0; k < h_num_improper(i); k++) {
-    buf[m++] = ubuf(MAX(h_improper_type(i,k),-h_improper_type(i,k))).d;
-    buf[m++] = ubuf(h_improper_atom1(i,k)).d;
-    buf[m++] = ubuf(h_improper_atom2(i,k)).d;
-    buf[m++] = ubuf(h_improper_atom3(i,k)).d;
-    buf[m++] = ubuf(h_improper_atom4(i,k)).d;
-  }
-
-  if (atom->nextra_restart)
-    for (int iextra = 0; iextra < atom->nextra_restart; iextra++)
-      m += modify->fix[atom->extra_restart[iextra]]->pack_restart(i,&buf[m]);
-
-  buf[0] = m;
-  return m;
-}
-
-/* ----------------------------------------------------------------------
-   unpack data for one atom from restart file including extra quantities
-------------------------------------------------------------------------- */
-
-int AtomVecFullKokkos::unpack_restart(double *buf)
-{
-  int k;
-
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) {
-    grow(0);
-    if (atom->nextra_store)
-      memory->grow(atom->extra,nmax,atom->nextra_store,"atom:extra");
-  }
-  atomKK->sync(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-       MASK_MASK | IMAGE_MASK | Q_MASK | MOLECULE_MASK | BOND_MASK |
-       ANGLE_MASK | DIHEDRAL_MASK | IMPROPER_MASK | SPECIAL_MASK);
-  atomKK->modified(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-           MASK_MASK | IMAGE_MASK | Q_MASK | MOLECULE_MASK | BOND_MASK |
-           ANGLE_MASK | DIHEDRAL_MASK | IMPROPER_MASK | SPECIAL_MASK);
-
-  int m = 1;
-  h_x(nlocal,0) = buf[m++];
-  h_x(nlocal,1) = buf[m++];
-  h_x(nlocal,2) = buf[m++];
-  h_tag(nlocal) = (tagint) ubuf(buf[m++]).i;
-  h_type(nlocal) = (int) ubuf(buf[m++]).i;
-  h_mask(nlocal) = (int) ubuf(buf[m++]).i;
-  h_image(nlocal) = (imageint) ubuf(buf[m++]).i;
-  h_v(nlocal,0) = buf[m++];
-  h_v(nlocal,1) = buf[m++];
-  h_v(nlocal,2) = buf[m++];
-
-  h_q(nlocal) = buf[m++];
-  h_molecule(nlocal) = (tagint) ubuf(buf[m++]).i;
-
-  h_num_bond(nlocal) = (int) ubuf(buf[m++]).i;
-  for (k = 0; k < h_num_bond(nlocal); k++) {
-    h_bond_type(nlocal,k) = (int) ubuf(buf[m++]).i;
-    h_bond_atom(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-  }
-
-  h_num_angle(nlocal) = (int) ubuf(buf[m++]).i;
-  for (k = 0; k < h_num_angle(nlocal); k++) {
-    h_angle_type(nlocal,k) = (int) ubuf(buf[m++]).i;
-    h_angle_atom1(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-    h_angle_atom2(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-    h_angle_atom3(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-  }
-
-  h_num_dihedral(nlocal) = (int) ubuf(buf[m++]).i;
-  for (k = 0; k < h_num_dihedral(nlocal); k++) {
-    h_dihedral_type(nlocal,k) = (int) ubuf(buf[m++]).i;
-    h_dihedral_atom1(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-    h_dihedral_atom2(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-    h_dihedral_atom3(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-    h_dihedral_atom4(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-  }
-
-  h_num_improper(nlocal) = (int) ubuf(buf[m++]).i;
-  for (k = 0; k < h_num_improper(nlocal); k++) {
-    h_improper_type(nlocal,k) = (int) ubuf(buf[m++]).i;
-    h_improper_atom1(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-    h_improper_atom2(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-    h_improper_atom3(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-    h_improper_atom4(nlocal,k) = (tagint) ubuf(buf[m++]).i;
-  }
-
-  h_nspecial(nlocal,0) = h_nspecial(nlocal,1) = h_nspecial(nlocal,2) = 0;
-
-  double **extra = atom->extra;
-  if (atom->nextra_store) {
-    int size = static_cast<int> (buf[0]) - m;
-    for (int i = 0; i < size; i++) extra[nlocal][i] = buf[m++];
-  }
-
-  atom->nlocal++;
-  return m;
-}
-
-/* ----------------------------------------------------------------------
-   create one atom of itype at coord
-   set other values to defaults
-------------------------------------------------------------------------- */
-
-void AtomVecFullKokkos::create_atom(int itype, double *coord)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) {
-    atomKK->modified(Host,ALL_MASK);
-    grow(0);
-  }
-  atomKK->sync(Host,ALL_MASK);
-  atomKK->modified(Host,ALL_MASK);
-
-  tag[nlocal] = 0;
-  type[nlocal] = itype;
-  h_x(nlocal,0) = coord[0];
-  h_x(nlocal,1) = coord[1];
-  h_x(nlocal,2) = coord[2];
-  h_mask(nlocal) = 1;
-  h_image(nlocal) = ((imageint) IMGMAX << IMG2BITS) |
-    ((imageint) IMGMAX << IMGBITS) | IMGMAX;
-  h_v(nlocal,0) = 0.0;
-  h_v(nlocal,1) = 0.0;
-  h_v(nlocal,2) = 0.0;
-
-  h_q(nlocal) = 0.0;
-  h_molecule(nlocal) = 0;
-  h_num_bond(nlocal) = 0;
-  h_num_angle(nlocal) = 0;
-  h_num_dihedral(nlocal) = 0;
-  h_num_improper(nlocal) = 0;
-  h_nspecial(nlocal,0) = h_nspecial(nlocal,1) = h_nspecial(nlocal,2) = 0;
-
-  atom->nlocal++;
-}
-
-/* ----------------------------------------------------------------------
-   unpack one line from Atoms section of data file
-   initialize other atom quantities
-------------------------------------------------------------------------- */
-
-void AtomVecFullKokkos::data_atom(double *coord, imageint imagetmp,
-                                  const std::vector<std::string> &values, std::string &extract)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) grow(0);
-  atomKK->modified(Host,ALL_MASK);
-
-  h_tag(nlocal) = utils::inumeric(FLERR,values[0],true,lmp);
-  h_molecule(nlocal) = utils::inumeric(FLERR,values[1],true,lmp);
-  h_type(nlocal) = utils::inumeric(FLERR,values[2],true,lmp);
-  extract = values[2];
-  if (h_type(nlocal) <= 0 || h_type(nlocal) > atom->ntypes)
-    error->one(FLERR,"Invalid atom type in Atoms section of data file");
-
-  h_q(nlocal) = utils::numeric(FLERR,values[3],true,lmp);
-
-  h_x(nlocal,0) = coord[0];
-  h_x(nlocal,1) = coord[1];
-  h_x(nlocal,2) = coord[2];
-
-  h_image(nlocal) = imagetmp;
-
-  h_mask(nlocal) = 1;
-  h_v(nlocal,0) = 0.0;
-  h_v(nlocal,1) = 0.0;
-  h_v(nlocal,2) = 0.0;
-  h_num_bond(nlocal) = 0;
-  h_num_angle(nlocal) = 0;
-  h_num_dihedral(nlocal) = 0;
-  h_num_improper(nlocal) = 0;
-
-  atom->nlocal++;
-}
-
-/* ----------------------------------------------------------------------
-   unpack hybrid quantities from one line in Atoms section of data file
-   initialize other atom quantities for this sub-style
-------------------------------------------------------------------------- */
-
-int AtomVecFullKokkos::data_atom_hybrid(int nlocal, const std::vector<std::string> &values,
-                                        int offset)
-{
-  h_molecule(nlocal) = utils::inumeric(FLERR,values[offset],true,lmp);
-  h_q(nlocal) = utils::numeric(FLERR,values[offset+1],true,lmp);
-  h_num_bond(nlocal) = 0;
-  h_num_angle(nlocal) = 0;
-  h_num_dihedral(nlocal) = 0;
-  h_num_improper(nlocal) = 0;
-  return 2;
-}
-
-/* ----------------------------------------------------------------------
-   pack atom info for data file including 3 image flags
-------------------------------------------------------------------------- */
-
-void AtomVecFullKokkos::pack_data(double **buf)
-{
-  int nlocal = atom->nlocal;
-  for (int i = 0; i < nlocal; i++) {
-    buf[i][0] = h_tag(i);
-    buf[i][1] = h_molecule(i);
-    buf[i][2] = h_type(i);
-    buf[i][3] = h_q(i);
-    buf[i][4] = h_x(i,0);
-    buf[i][5] = h_x(i,1);
-    buf[i][6] = h_x(i,2);
-    buf[i][7] = (h_image[i] & IMGMASK) - IMGMAX;
-    buf[i][8] = (h_image[i] >> IMGBITS & IMGMASK) - IMGMAX;
-    buf[i][9] = (h_image[i] >> IMG2BITS) - IMGMAX;
-  }
-}
-
-/* ----------------------------------------------------------------------
-   pack hybrid atom info for data file
-------------------------------------------------------------------------- */
-
-int AtomVecFullKokkos::pack_data_hybrid(int i, double *buf)
-{
-  buf[0] = h_molecule(i);
-  buf[1] = h_q(i);
-  return 2;
-}
-
-/* ----------------------------------------------------------------------
-   write atom info to data file including 3 image flags
-------------------------------------------------------------------------- */
-
-void AtomVecFullKokkos::write_data(FILE *fp, int n, double **buf)
-{
-  for (int i = 0; i < n; i++)
-    fprintf(fp,"%d %d %d %-1.16e %-1.16e %-1.16e %-1.16e %d %d %d\n",
-            (int) buf[i][0],(int) buf[i][1], (int) buf[i][2], buf[i][3],
-            buf[i][4],buf[i][5],buf[i][6],
-            (int) buf[i][7],(int) buf[i][8],(int) buf[i][9]);
-}
-
-/* ----------------------------------------------------------------------
-   write hybrid atom info to data file
-------------------------------------------------------------------------- */
-
-int AtomVecFullKokkos::write_data_hybrid(FILE *fp, double *buf)
-{
-  fprintf(fp," " TAGINT_FORMAT " %-1.16e",(tagint) ubuf(buf[0]).i,buf[1]);
-  return 2;
-}
-
-/* ----------------------------------------------------------------------
-   return # of bytes of allocated memory
-------------------------------------------------------------------------- */
-
-double AtomVecFullKokkos::memory_usage()
-{
-  double bytes = 0;
-
-  if (atom->memcheck("tag")) bytes += memory->usage(tag,nmax);
-  if (atom->memcheck("type")) bytes += memory->usage(type,nmax);
-  if (atom->memcheck("mask")) bytes += memory->usage(mask,nmax);
-  if (atom->memcheck("image")) bytes += memory->usage(image,nmax);
-  if (atom->memcheck("x")) bytes += memory->usage(x,nmax,3);
-  if (atom->memcheck("v")) bytes += memory->usage(v,nmax,3);
-  if (atom->memcheck("f")) bytes += memory->usage(f,nmax*commKK->nthreads,3);
-
-  if (atom->memcheck("q")) bytes += memory->usage(q,nmax);
-  if (atom->memcheck("molecule")) bytes += memory->usage(molecule,nmax);
-  if (atom->memcheck("nspecial")) bytes += memory->usage(nspecial,nmax,3);
-  if (atom->memcheck("special"))
-    bytes += memory->usage(special,nmax,atom->maxspecial);
-
-  if (atom->memcheck("num_bond")) bytes += memory->usage(num_bond,nmax);
-  if (atom->memcheck("bond_type"))
-    bytes += memory->usage(bond_type,nmax,atom->bond_per_atom);
-  if (atom->memcheck("bond_atom"))
-    bytes += memory->usage(bond_atom,nmax,atom->bond_per_atom);
-
-  if (atom->memcheck("num_angle")) bytes += memory->usage(num_angle,nmax);
-  if (atom->memcheck("angle_type"))
-    bytes += memory->usage(angle_type,nmax,atom->angle_per_atom);
-  if (atom->memcheck("angle_atom1"))
-    bytes += memory->usage(angle_atom1,nmax,atom->angle_per_atom);
-  if (atom->memcheck("angle_atom2"))
-    bytes += memory->usage(angle_atom2,nmax,atom->angle_per_atom);
-  if (atom->memcheck("angle_atom3"))
-    bytes += memory->usage(angle_atom3,nmax,atom->angle_per_atom);
-
-  if (atom->memcheck("num_dihedral")) bytes += memory->usage(num_dihedral,nmax);
-  if (atom->memcheck("dihedral_type"))
-    bytes += memory->usage(dihedral_type,nmax,atom->dihedral_per_atom);
-  if (atom->memcheck("dihedral_atom1"))
-    bytes += memory->usage(dihedral_atom1,nmax,atom->dihedral_per_atom);
-  if (atom->memcheck("dihedral_atom2"))
-    bytes += memory->usage(dihedral_atom2,nmax,atom->dihedral_per_atom);
-  if (atom->memcheck("dihedral_atom3"))
-    bytes += memory->usage(dihedral_atom3,nmax,atom->dihedral_per_atom);
-  if (atom->memcheck("dihedral_atom4"))
-    bytes += memory->usage(dihedral_atom4,nmax,atom->dihedral_per_atom);
-  if (atom->memcheck("num_improper")) bytes += memory->usage(num_improper,nmax);
-  if (atom->memcheck("improper_type"))
-    bytes += memory->usage(improper_type,nmax,atom->improper_per_atom);
-  if (atom->memcheck("improper_atom1"))
-    bytes += memory->usage(improper_atom1,nmax,atom->improper_per_atom);
-  if (atom->memcheck("improper_atom2"))
-    bytes += memory->usage(improper_atom2,nmax,atom->improper_per_atom);
-  if (atom->memcheck("improper_atom3"))
-    bytes += memory->usage(improper_atom3,nmax,atom->improper_per_atom);
-  if (atom->memcheck("improper_atom4"))
-    bytes += memory->usage(improper_atom4,nmax,atom->improper_per_atom);
-
-  return bytes;
-}
-
-/* ---------------------------------------------------------------------- */
-
 void AtomVecFullKokkos::sync(ExecutionSpace space, unsigned int mask)
 {
   if (space == Device) {
@@ -1992,4 +1175,3 @@ void AtomVecFullKokkos::modified(ExecutionSpace space, unsigned int mask)
     }
   }
 }
-

src/KOKKOS/atom_vec_full_kokkos.h

@@ -24,35 +24,15 @@ AtomStyle(full/kk/host,AtomVecFullKokkos);
 #define LMP_ATOM_VEC_FULL_KOKKOS_H
 
 #include "atom_vec_kokkos.h"
+#include "atom_vec_full.h"
 
 namespace LAMMPS_NS {
 
-class AtomVecFullKokkos : public AtomVecKokkos {
+class AtomVecFullKokkos : public AtomVecKokkos, public AtomVecFull {
  public:
   AtomVecFullKokkos(class LAMMPS *);
 
   void grow(int) override;
-  void copy(int, int, int) override;
-  int pack_border(int, int *, double *, int, int *) override;
-  int pack_border_vel(int, int *, double *, int, int *) override;
-  int pack_border_hybrid(int, int *, double *) override;
-  void unpack_border(int, int, double *) override;
-  void unpack_border_vel(int, int, double *) override;
-  int unpack_border_hybrid(int, int, double *) override;
-  int pack_exchange(int, double *) override;
-  int unpack_exchange(double *) override;
-  int size_restart() override;
-  int pack_restart(int, double *) override;
-  int unpack_restart(double *) override;
-  void create_atom(int, double *) override;
-  void data_atom(double *, imageint, const std::vector<std::string> &, std::string &) override;
-  int data_atom_hybrid(int, const std::vector<std::string> &, int) override;
-  void pack_data(double **) override;
-  int pack_data_hybrid(int, double *) override;
-  void write_data(FILE *, int, double **) override;
-  int write_data_hybrid(FILE *, double *) override;
-  double memory_usage() override;
-
   void grow_pointers() override;
   int pack_border_kokkos(int n, DAT::tdual_int_2d k_sendlist,
                          DAT::tdual_xfloat_2d buf,int iswap,
@@ -154,9 +134,6 @@ class AtomVecFullKokkos : public AtomVecKokkos {
   HAT::t_int_2d h_improper_type;
   HAT::t_tagint_2d h_improper_atom1,h_improper_atom2,
     h_improper_atom3,h_improper_atom4;
-
-  HAT::tdual_int_1d k_count;
-
 };
 
 }

src/KOKKOS/atom_vec_hybrid_kokkos.h

@@ -22,11 +22,12 @@ AtomStyle(hybrid/kk,AtomVecHybridKokkos);
 #define LMP_ATOM_VEC_HYBRID_KOKKOS_H
 
 #include "atom_vec_kokkos.h"
+#include "atom_vec_hybrid.h"
 #include "kokkos_type.h"
 
 namespace LAMMPS_NS {
 
-class AtomVecHybridKokkos : public AtomVecKokkos {
+class AtomVecHybridKokkos : public AtomVecKokkos, public AtomVecHybrid {
  public:
   int nstyles;
   class AtomVec **styles;
@@ -34,39 +35,6 @@ class AtomVecHybridKokkos : public AtomVecKokkos {
 
   AtomVecHybridKokkos(class LAMMPS *);
   ~AtomVecHybridKokkos() override;
-  void process_args(int, char **) override;
-  void init() override;
-  void grow(int) override;
-  void grow_pointers() override;
-  void copy(int, int, int) override;
-  void clear_bonus() override;
-  void force_clear(int, size_t) override;
-  int pack_comm(int, int *, double *, int, int *) override;
-  int pack_comm_vel(int, int *, double *, int, int *) override;
-  void unpack_comm(int, int, double *) override;
-  void unpack_comm_vel(int, int, double *) override;
-  int pack_reverse(int, int, double *) override;
-  void unpack_reverse(int, int *, double *) override;
-  int pack_border(int, int *, double *, int, int *) override;
-  int pack_border_vel(int, int *, double *, int, int *) override;
-  void unpack_border(int, int, double *) override;
-  void unpack_border_vel(int, int, double *) override;
-  int pack_exchange(int, double *) override;
-  int unpack_exchange(double *) override;
-  int size_restart() override;
-  int pack_restart(int, double *) override;
-  int unpack_restart(double *) override;
-  void create_atom(int, double *) override;
-  void data_atom(double *, imageint, const std::vector<std::string> &, std::string &) override;
-  int data_atom_hybrid(int, const std::vector<std::string> &, int) override {return 0;}
-  void data_vel(int, const std::vector<std::string> &) override;
-  void pack_data(double **) override;
-  void write_data(FILE *, int, double **) override;
-  void pack_vel(double **) override;
-  void write_vel(FILE *, int, double **) override;
-  int property_atom(const std::string &) override;
-  void pack_property_atom(int, double *, int, int) override;
-  double memory_usage() override;
 
   int pack_comm_kokkos(const int &n, const DAT::tdual_int_2d &k_sendlist,
                        const int & iswap,

src/KOKKOS/atom_vec_kokkos.cpp

@@ -32,22 +32,10 @@ AtomVecKokkos::AtomVecKokkos(LAMMPS *lmp) : AtomVec(lmp)
 
   no_comm_vel_flag = 0;
   no_border_vel_flag = 1;
-}
 
-/* ----------------------------------------------------------------------
-   roundup N so it is a multiple of DELTA
-   error if N exceeds 32-bit int, since will be used as arg to grow()
-   overload needed because Kokkos uses a smaller DELTA than in atom_vec.cpp
-   and an exponential instead of a linear growth
-------------------------------------------------------------------------- */
-
-bigint AtomVecKokkos::roundup(bigint n)
-{
-  auto DELTA = LMP_KOKKOS_AV_DELTA;
-  if (n % DELTA) n = n/DELTA * DELTA + DELTA;
-  if (n > MAXSMALLINT)
-    error->one(FLERR,"Too many atoms created on one or more procs");
-  return n;
+  k_count = DAT::tdual_int_1d("atom::k_count",1);
+  atomKK = (AtomKokkos *) atom;
+  commKK = (CommKokkos *) comm;
 }
 
 /* ---------------------------------------------------------------------- */
@@ -708,138 +696,6 @@ void AtomVecKokkos::unpack_comm_vel_kokkos(const int &n, const int &first,
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecKokkos::pack_comm(int n, int *list, double *buf,
-                             int pbc_flag, int *pbc)
-{
-  int i,j,m;
-  double dx,dy,dz;
-
-  m = 0;
-  if (pbc_flag == 0) {
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0);
-      buf[m++] = h_x(j,1);
-      buf[m++] = h_x(j,2);
-    }
-  } else {
-    if (domain->triclinic == 0) {
-      dx = pbc[0]*domain->xprd;
-      dy = pbc[1]*domain->yprd;
-      dz = pbc[2]*domain->zprd;
-    } else {
-      dx = pbc[0]*domain->xprd + pbc[5]*domain->xy + pbc[4]*domain->xz;
-      dy = pbc[1]*domain->yprd + pbc[3]*domain->yz;
-      dz = pbc[2]*domain->zprd;
-    }
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0) + dx;
-      buf[m++] = h_x(j,1) + dy;
-      buf[m++] = h_x(j,2) + dz;
-    }
-  }
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecKokkos::pack_comm_vel(int n, int *list, double *buf,
-                                 int pbc_flag, int *pbc)
-{
-  int i,j,m;
-  double dx,dy,dz,dvx,dvy,dvz;
-
-  m = 0;
-  if (pbc_flag == 0) {
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0);
-      buf[m++] = h_x(j,1);
-      buf[m++] = h_x(j,2);
-      buf[m++] = h_v(j,0);
-      buf[m++] = h_v(j,1);
-      buf[m++] = h_v(j,2);
-    }
-  } else {
-    if (domain->triclinic == 0) {
-      dx = pbc[0]*domain->xprd;
-      dy = pbc[1]*domain->yprd;
-      dz = pbc[2]*domain->zprd;
-    } else {
-      dx = pbc[0]*domain->xprd + pbc[5]*domain->xy + pbc[4]*domain->xz;
-      dy = pbc[1]*domain->yprd + pbc[3]*domain->yz;
-      dz = pbc[2]*domain->zprd;
-    }
-    if (!deform_vremap) {
-      for (i = 0; i < n; i++) {
-        j = list[i];
-        buf[m++] = h_x(j,0) + dx;
-        buf[m++] = h_x(j,1) + dy;
-        buf[m++] = h_x(j,2) + dz;
-        buf[m++] = h_v(j,0);
-        buf[m++] = h_v(j,1);
-        buf[m++] = h_v(j,2);
-      }
-    } else {
-      dvx = pbc[0]*h_rate[0] + pbc[5]*h_rate[5] + pbc[4]*h_rate[4];
-      dvy = pbc[1]*h_rate[1] + pbc[3]*h_rate[3];
-      dvz = pbc[2]*h_rate[2];
-      for (i = 0; i < n; i++) {
-        j = list[i];
-        buf[m++] = h_x(j,0) + dx;
-        buf[m++] = h_x(j,1) + dy;
-        buf[m++] = h_x(j,2) + dz;
-        if (atom->mask[i] & deform_groupbit) {
-          buf[m++] = h_v(j,0) + dvx;
-          buf[m++] = h_v(j,1) + dvy;
-          buf[m++] = h_v(j,2) + dvz;
-        } else {
-          buf[m++] = h_v(j,0);
-          buf[m++] = h_v(j,1);
-          buf[m++] = h_v(j,2);
-        }
-      }
-    }
-  }
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-void AtomVecKokkos::unpack_comm(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  for (i = first; i < last; i++) {
-    h_x(i,0) = buf[m++];
-    h_x(i,1) = buf[m++];
-    h_x(i,2) = buf[m++];
-  }
-}
-
-/* ---------------------------------------------------------------------- */
-
-void AtomVecKokkos::unpack_comm_vel(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  for (i = first; i < last; i++) {
-    h_x(i,0) = buf[m++];
-    h_x(i,1) = buf[m++];
-    h_x(i,2) = buf[m++];
-    h_v(i,0) = buf[m++];
-    h_v(i,1) = buf[m++];
-    h_v(i,2) = buf[m++];
-  }
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType>
 struct AtomVecKokkos_PackReverse {
   typedef DeviceType device_type;
@@ -977,84 +833,3 @@ void AtomVecKokkos::unpack_reverse_kokkos(const int &n,
     atomKK->modified(Device,F_MASK);
   }
 }
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecKokkos::pack_reverse(int n, int first, double *buf)
-{
-  if (n > 0)
-    atomKK->sync(Host,F_MASK);
-
-  int m = 0;
-  const int last = first + n;
-  for (int i = first; i < last; i++) {
-    buf[m++] = h_f(i,0);
-    buf[m++] = h_f(i,1);
-    buf[m++] = h_f(i,2);
-  }
-
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-void AtomVecKokkos::unpack_reverse(int n, int *list, double *buf)
-{
-  int m = 0;
-  for (int i = 0; i < n; i++) {
-    const int j = list[i];
-    h_f(j,0) += buf[m++];
-    h_f(j,1) += buf[m++];
-    h_f(j,2) += buf[m++];
-  }
-
-  if (n > 0)
-    atomKK->modified(Host,F_MASK);
-}
-
-/* ----------------------------------------------------------------------
- *    unpack one line from Velocities section of data file
- *    ------------------------------------------------------------------------- */
-
-void AtomVecKokkos::data_vel(int m, const std::vector<std::string> &values)
-{
-  double **v = atom->v;
-  int ivalue = 1;
-  v[m][0] = utils::numeric(FLERR,values[ivalue++],true,lmp);
-  v[m][1] = utils::numeric(FLERR,values[ivalue++],true,lmp);
-  v[m][2] = utils::numeric(FLERR,values[ivalue++],true,lmp);
-
-  atomKK->modified(Host,V_MASK);
-}
-
-/* ----------------------------------------------------------------------
- *    pack velocity info for data file
- *    ------------------------------------------------------------------------- */
-
-void AtomVecKokkos::pack_vel(double **buf)
-{
-  double **v = atom->v;
-  tagint *tag = atom->tag;
-  int nlocal = atom->nlocal;
-
-  atomKK->sync(Host,V_MASK|TAG_MASK);
-
-  for (int i = 0; i < nlocal; i++) {
-    buf[i][0] = ubuf(tag[i]).d;
-    buf[i][1] = v[i][0];
-    buf[i][2] = v[i][1];
-    buf[i][3] = v[i][2];
-  }
-}
-
-/* ----------------------------------------------------------------------
- *    write velocity info to data file
- *    ------------------------------------------------------------------------- */
-
-void AtomVecKokkos::write_vel(FILE *fp, int n, double **buf)
-{
-  for (int i = 0; i < n; i++)
-    fprintf(fp,TAGINT_FORMAT " %-1.16e %-1.16e %-1.16e\n",
-            (tagint) ubuf(buf[i][0]).i,buf[i][1],buf[i][2],buf[i][3]);
-}
-

src/KOKKOS/atom_vec_kokkos.h

@@ -33,21 +33,10 @@ union d_ubuf {
   d_ubuf(int arg) : i(arg) {}
 };
 
-class AtomVecKokkos : public AtomVec {
+class AtomVecKokkos : virtual public AtomVec {
  public:
   AtomVecKokkos(class LAMMPS *);
 
-  bigint roundup(bigint) override;
-  int pack_comm(int, int *, double *, int, int *) override;
-  int pack_comm_vel(int, int *, double *, int, int *) override;
-  void unpack_comm(int, int, double *) override;
-  void unpack_comm_vel(int, int, double *) override;
-  int pack_reverse(int, int, double *) override;
-  void unpack_reverse(int, int *, double *) override;
-  void data_vel(int, const std::vector<std::string> &) override;
-  void pack_vel(double **) override;
-  void write_vel(FILE *, int, double **) override;
-
   virtual void sync(ExecutionSpace space, unsigned int mask) = 0;
   virtual void modified(ExecutionSpace space, unsigned int mask) = 0;
   virtual void sync_overlapping_device(ExecutionSpace space, unsigned int mask) = 0;
@@ -130,7 +119,6 @@ class AtomVecKokkos : public AtomVec {
   int no_comm_vel_flag,no_border_vel_flag;
 
  protected:
-
   HAT::t_x_array h_x;
   HAT::t_v_array h_v;
   HAT::t_f_array h_f;
@@ -139,6 +127,8 @@ class AtomVecKokkos : public AtomVec {
   size_t buffer_size;
   void* buffer;
 
+  DAT::tdual_int_1d k_count;
+
   #ifdef LMP_KOKKOS_GPU
   template<class ViewType>
   Kokkos::View<typename ViewType::data_type,
@@ -203,4 +193,3 @@ class AtomVecKokkos : public AtomVec {
 }
 
 #endif
-

src/KOKKOS/atom_vec_molecular_kokkos.h

@@ -24,41 +24,15 @@ AtomStyle(molecular/kk/host,AtomVecMolecularKokkos);
 #define LMP_ATOM_VEC_MOLECULAR_KOKKOS_H
 
 #include "atom_vec_kokkos.h"
+#include "atom_vec_molecular.h"
 
 namespace LAMMPS_NS {
 
-class AtomVecMolecularKokkos : public AtomVecKokkos {
+class AtomVecMolecularKokkos : public AtomVecKokkos, public AtomVecMolecular {
  public:
   AtomVecMolecularKokkos(class LAMMPS *);
 
   void grow(int) override;
-  void copy(int, int, int) override;
-  int pack_comm(int, int *, double *, int, int *) override;
-  int pack_comm_vel(int, int *, double *, int, int *) override;
-  void unpack_comm(int, int, double *) override;
-  void unpack_comm_vel(int, int, double *) override;
-  int pack_reverse(int, int, double *) override;
-  void unpack_reverse(int, int *, double *) override;
-  int pack_border(int, int *, double *, int, int *) override;
-  int pack_border_vel(int, int *, double *, int, int *) override;
-  int pack_border_hybrid(int, int *, double *) override;
-  void unpack_border(int, int, double *) override;
-  void unpack_border_vel(int, int, double *) override;
-  int unpack_border_hybrid(int, int, double *) override;
-  int pack_exchange(int, double *) override;
-  int unpack_exchange(double *) override;
-  int size_restart() override;
-  int pack_restart(int, double *) override;
-  int unpack_restart(double *) override;
-  void create_atom(int, double *) override;
-  void data_atom(double *, imageint, const std::vector<std::string> &, std::string &) override;
-  int data_atom_hybrid(int, const std::vector<std::string> &, int) override;
-  void pack_data(double **) override;
-  int pack_data_hybrid(int, double *) override;
-  void write_data(FILE *, int, double **) override;
-  int write_data_hybrid(FILE *, double *) override;
-  double memory_usage() override;
-
   void grow_pointers() override;
   int pack_comm_kokkos(const int &n, const DAT::tdual_int_2d &k_sendlist,
                        const int & iswap,
@@ -167,9 +141,6 @@ class AtomVecMolecularKokkos : public AtomVecKokkos {
   HAT::t_int_2d h_improper_type;
   HAT::t_tagint_2d h_improper_atom1,h_improper_atom2,
     h_improper_atom3,h_improper_atom4;
-
-  HAT::tdual_int_1d k_count;
-
 };
 
 }

src/KOKKOS/atom_vec_sphere_kokkos.h

@@ -24,53 +24,17 @@ AtomStyle(sphere/kk/host,AtomVecSphereKokkos);
 #define LMP_ATOM_VEC_SPHERE_KOKKOS_H
 
 #include "atom_vec_kokkos.h"
+#include "atom_vec_sphere.h"
 #include "kokkos_type.h"
 
 namespace LAMMPS_NS {
 
-class AtomVecSphereKokkos : public AtomVecKokkos {
+class AtomVecSphereKokkos : public AtomVecKokkos, public AtomVecSphere {
  public:
   AtomVecSphereKokkos(class LAMMPS *);
 
-  void init() override;
   void grow(int) override;
   void grow_pointers() override;
-  void copy(int, int, int) override;
-  int pack_comm(int, int *, double *, int, int *) override;
-  int pack_comm_vel(int, int *, double *, int, int *) override;
-  int pack_comm_hybrid(int, int *, double *) override;
-  void unpack_comm(int, int, double *) override;
-  void unpack_comm_vel(int, int, double *) override;
-  int unpack_comm_hybrid(int, int, double *) override;
-  int pack_reverse(int, int, double *) override;
-  int pack_reverse_hybrid(int, int, double *) override;
-  void unpack_reverse(int, int *, double *) override;
-  int unpack_reverse_hybrid(int, int *, double *) override;
-  int pack_border(int, int *, double *, int, int *) override;
-  int pack_border_vel(int, int *, double *, int, int *) override;
-  int pack_border_hybrid(int, int *, double *) override;
-  void unpack_border(int, int, double *) override;
-  void unpack_border_vel(int, int, double *) override;
-  int unpack_border_hybrid(int, int, double *) override;
-  int pack_exchange(int, double *) override;
-  int unpack_exchange(double *) override;
-  int size_restart() override;
-  int pack_restart(int, double *) override;
-  int unpack_restart(double *) override;
-  void create_atom(int, double *) override;
-  void data_atom(double *, imageint, const std::vector<std::string> &, std::string &) override;
-  int data_atom_hybrid(int, const std::vector<std::string> &, int) override;
-  void data_vel(int, const std::vector<std::string> &) override;
-  int data_vel_hybrid(int, const std::vector<std::string> &, int) override;
-  void pack_data(double **) override;
-  int pack_data_hybrid(int, double *) override;
-  void write_data(FILE *, int, double **) override;
-  int write_data_hybrid(FILE *, double *) override;
-  void pack_vel(double **) override;
-  int pack_vel_hybrid(int, double *) override;
-  void write_vel(FILE *, int, double **) override;
-  int write_vel_hybrid(FILE *, double *) override;
-  double memory_usage() override;
 
   int pack_comm_kokkos(const int &n, const DAT::tdual_int_2d &k_sendlist,
                        const int & iswap,
@@ -139,8 +103,6 @@ class AtomVecSphereKokkos : public AtomVecKokkos {
   HAT::t_v_array h_omega;
   DAT::t_f_array d_torque;
   HAT::t_f_array h_torque;
-
-  DAT::tdual_int_1d k_count;
 };
 
 }

src/KOKKOS/atom_vec_spin_kokkos.cpp

@@ -41,26 +41,10 @@ using namespace LAMMPS_NS;
 
 /* ---------------------------------------------------------------------- */
 
-AtomVecSpinKokkos::AtomVecSpinKokkos(LAMMPS *lmp) : AtomVecKokkos(lmp)
+AtomVecSpinKokkos::AtomVecSpinKokkos(LAMMPS *lmp) : AtomVec(lmp),
+AtomVecKokkos(lmp), AtomVecSpin(lmp)
 {
-  molecular = 0;
-  mass_type = 1;
-  forceclearflag = 1;
 
-  comm_x_only = comm_f_only = 0;
-  size_forward = 7;
-  size_reverse = 9;
-  size_border = 10;
-  size_velocity = 3;
-  size_data_atom = 9;
-  size_data_vel = 4;
-  xcol_data = 4;
-
-  atom->sp_flag = 1;
-
-  k_count = DAT::tdual_int_1d("atom::k_count",1);
-  atomKK = (AtomKokkos *) atom;
-  commKK = (CommKokkos *) comm;
 }
 
 /* ----------------------------------------------------------------------
@@ -147,33 +131,6 @@ void AtomVecSpinKokkos::grow_pointers()
   h_fm_long = atomKK->k_fm_long.h_view;
 }
 
-/* ----------------------------------------------------------------------
-   copy atom I info to atom J
-------------------------------------------------------------------------- */
-
-void AtomVecSpinKokkos::copy(int i, int j, int delflag)
-{
-  h_tag[j] = h_tag[i];
-  h_type[j] = h_type[i];
-  mask[j] = mask[i];
-  h_image[j] = h_image[i];
-  h_x(j,0) = h_x(i,0);
-  h_x(j,1) = h_x(i,1);
-  h_x(j,2) = h_x(i,2);
-  h_v(j,0) = h_v(i,0);
-  h_v(j,1) = h_v(i,1);
-  h_v(j,2) = h_v(i,2);
-
-  h_sp(j,0) = h_sp(i,0);
-  h_sp(j,1) = h_sp(i,1);
-  h_sp(j,2) = h_sp(i,2);
-  h_sp(j,3) = h_sp(i,3);
-
-  if (atom->nextra_grow)
-    for (int iextra = 0; iextra < atom->nextra_grow; iextra++)
-      modify->fix[atom->extra_grow[iextra]]->copy_arrays(i,j,delflag);
-}
-
 /* ---------------------------------------------------------------------- */
 
 template<class DeviceType,int PBC_FLAG,int TRICLINIC>
@@ -347,167 +304,6 @@ int AtomVecSpinKokkos::pack_border_kokkos(int n, DAT::tdual_int_2d k_sendlist, D
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecSpinKokkos::pack_border(int n, int *list, double *buf,
-                               int pbc_flag, int *pbc)
-{
-  int i,j,m;
-  double dx,dy,dz;
-
-  m = 0;
-  if (pbc_flag == 0) {
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0);
-      buf[m++] = h_x(j,1);
-      buf[m++] = h_x(j,2);
-      buf[m++] = ubuf(h_tag(j)).d;
-      buf[m++] = ubuf(h_type(j)).d;
-      buf[m++] = ubuf(h_mask(j)).d;
-      buf[m++] = h_sp(j,0);
-      buf[m++] = h_sp(j,1);
-      buf[m++] = h_sp(j,2);
-      buf[m++] = h_sp(j,3);
-    }
-  } else {
-    if (domain->triclinic == 0) {
-      dx = pbc[0]*domain->xprd;
-      dy = pbc[1]*domain->yprd;
-      dz = pbc[2]*domain->zprd;
-    } else {
-      dx = pbc[0];
-      dy = pbc[1];
-      dz = pbc[2];
-    }
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0) + dx;
-      buf[m++] = h_x(j,1) + dy;
-      buf[m++] = h_x(j,2) + dz;
-      buf[m++] = ubuf(h_tag(j)).d;
-      buf[m++] = ubuf(h_type(j)).d;
-      buf[m++] = ubuf(h_mask(j)).d;
-      buf[m++] = h_sp(j,0);
-      buf[m++] = h_sp(j,1);
-      buf[m++] = h_sp(j,2);
-      buf[m++] = h_sp(j,3);
-    }
-  }
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->pack_border(n,list,&buf[m]);
-
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecSpinKokkos::pack_border_vel(int n, int *list, double *buf,
-                                   int pbc_flag, int *pbc)
-{
-  int i,j,m;
-  double dx,dy,dz,dvx,dvy,dvz;
-
-  m = 0;
-  if (pbc_flag == 0) {
-    for (i = 0; i < n; i++) {
-      j = list[i];
-      buf[m++] = h_x(j,0);
-      buf[m++] = h_x(j,1);
-      buf[m++] = h_x(j,2);
-      buf[m++] = ubuf(h_tag(j)).d;
-      buf[m++] = ubuf(h_type(j)).d;
-      buf[m++] = ubuf(h_mask(j)).d;
-      buf[m++] = h_sp(j,0);
-      buf[m++] = h_sp(j,1);
-      buf[m++] = h_sp(j,2);
-      buf[m++] = h_sp(j,3);
-      buf[m++] = h_v(j,0);
-      buf[m++] = h_v(j,1);
-      buf[m++] = h_v(j,2);
-    }
-  } else {
-    if (domain->triclinic == 0) {
-      dx = pbc[0]*domain->xprd;
-      dy = pbc[1]*domain->yprd;
-      dz = pbc[2]*domain->zprd;
-    } else {
-      dx = pbc[0];
-      dy = pbc[1];
-      dz = pbc[2];
-    }
-    if (!deform_vremap) {
-      for (i = 0; i < n; i++) {
-        j = list[i];
-        buf[m++] = h_x(j,0) + dx;
-        buf[m++] = h_x(j,1) + dy;
-        buf[m++] = h_x(j,2) + dz;
-        buf[m++] = ubuf(h_tag(j)).d;
-        buf[m++] = ubuf(h_type(j)).d;
-        buf[m++] = ubuf(h_mask(j)).d;
-        buf[m++] = h_sp(j,0);
-        buf[m++] = h_sp(j,1);
-        buf[m++] = h_sp(j,2);
-        buf[m++] = h_sp(j,3);
-        buf[m++] = h_v(j,0);
-        buf[m++] = h_v(j,1);
-        buf[m++] = h_v(j,2);
-      }
-    } else {
-      dvx = pbc[0]*h_rate[0] + pbc[5]*h_rate[5] + pbc[4]*h_rate[4];
-      dvy = pbc[1]*h_rate[1] + pbc[3]*h_rate[3];
-      dvz = pbc[2]*h_rate[2];
-      for (i = 0; i < n; i++) {
-        j = list[i];
-        buf[m++] = h_x(j,0) + dx;
-        buf[m++] = h_x(j,1) + dy;
-        buf[m++] = h_x(j,2) + dz;
-        buf[m++] = ubuf(h_tag(j)).d;
-        buf[m++] = ubuf(h_type(j)).d;
-        buf[m++] = ubuf(h_mask(j)).d;
-        buf[m++] = h_sp(j,0);
-        buf[m++] = h_sp(j,1);
-        buf[m++] = h_sp(j,2);
-        buf[m++] = h_sp(j,3);
-        if (mask[i] & deform_groupbit) {
-          buf[m++] = h_v(j,0) + dvx;
-          buf[m++] = h_v(j,1) + dvy;
-          buf[m++] = h_v(j,2) + dvz;
-        } else {
-          buf[m++] = h_v(j,0);
-          buf[m++] = h_v(j,1);
-          buf[m++] = h_v(j,2);
-        }
-      }
-    }
-  }
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->pack_border(n,list,&buf[m]);
-
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecSpinKokkos::pack_border_hybrid(int n, int *list, double *buf)
-{
-  int i,j,m;
-
-  m = 0;
-  for (i = 0; i < n; i++) {
-    j = list[i];
-    buf[m++] = h_sp(j,0);
-    buf[m++] = h_sp(j,1);
-    buf[m++] = h_sp(j,2);
-    buf[m++] = h_sp(j,3);
-  }
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType>
 struct AtomVecSpinKokkos_UnpackBorder {
   typedef DeviceType device_type;
@@ -568,87 +364,6 @@ void AtomVecSpinKokkos::unpack_border_kokkos(const int &n, const int &first,
 
 /* ---------------------------------------------------------------------- */
 
-void AtomVecSpinKokkos::unpack_border(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-
-  for (i = first; i < last; i++) {
-    if (i == nmax) {
-      grow(0);
-    }
-    atomKK->modified(Host,X_MASK|TAG_MASK|TYPE_MASK|MASK_MASK|SP_MASK);
-    h_x(i,0) = buf[m++];
-    h_x(i,1) = buf[m++];
-    h_x(i,2) = buf[m++];
-    h_tag(i) =  (tagint)  ubuf(buf[m++]).i;
-    h_type(i) = (int) ubuf(buf[m++]).i;
-    h_mask(i) = (int) ubuf(buf[m++]).i;
-    h_sp(i,0) = buf[m++];
-    h_sp(i,1) = buf[m++];
-    h_sp(i,2) = buf[m++];
-    h_sp(i,3) = buf[m++];
-  }
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->
-        unpack_border(n,first,&buf[m]);
-}
-
-/* ---------------------------------------------------------------------- */
-
-void AtomVecSpinKokkos::unpack_border_vel(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  for (i = first; i < last; i++) {
-    if (i == nmax) grow(0);
-    atomKK->modified(Host,X_MASK|V_MASK|TAG_MASK|TYPE_MASK|MASK_MASK|SP_MASK);
-    h_x(i,0) = buf[m++];
-    h_x(i,1) = buf[m++];
-    h_x(i,2) = buf[m++];
-    h_tag(i) =  (tagint)  ubuf(buf[m++]).i;
-    h_type(i) = (int) ubuf(buf[m++]).i;
-    h_mask(i) = (int) ubuf(buf[m++]).i;
-    h_sp(i,0) = buf[m++];
-    h_sp(i,1) = buf[m++];
-    h_sp(i,2) = buf[m++];
-    h_sp(i,3) = buf[m++];
-    h_v(i,0) = buf[m++];
-    h_v(i,1) = buf[m++];
-    h_v(i,2) = buf[m++];
-  }
-
-  if (atom->nextra_border)
-    for (int iextra = 0; iextra < atom->nextra_border; iextra++)
-      m += modify->fix[atom->extra_border[iextra]]->
-        unpack_border(n,first,&buf[m]);
-}
-
-/* ---------------------------------------------------------------------- */
-
-int AtomVecSpinKokkos::unpack_border_hybrid(int n, int first, double *buf)
-{
-  int i,m,last;
-
-  m = 0;
-  last = first + n;
-  for (i = first; i < last; i++) {
-    h_sp(i,0) = buf[m++];
-    h_sp(i,1) = buf[m++];
-    h_sp(i,2) = buf[m++];
-    h_sp(i,3) = buf[m++];
-  }
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType>
 struct AtomVecSpinKokkos_PackExchangeFunctor {
   typedef DeviceType device_type;
@@ -771,34 +486,6 @@ int AtomVecSpinKokkos::pack_exchange_kokkos(const int &nsend,DAT::tdual_xfloat_2
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecSpinKokkos::pack_exchange(int i, double *buf)
-{
-  int m = 1;
-  buf[m++] = h_x(i,0);
-  buf[m++] = h_x(i,1);
-  buf[m++] = h_x(i,2);
-  buf[m++] = h_v(i,0);
-  buf[m++] = h_v(i,1);
-  buf[m++] = h_v(i,2);
-  buf[m++] = ubuf(h_tag(i)).d;
-  buf[m++] = ubuf(h_type(i)).d;
-  buf[m++] = ubuf(h_mask(i)).d;
-  buf[m++] = ubuf(h_image(i)).d;
-  buf[m++] = h_sp(i,0);
-  buf[m++] = h_sp(i,1);
-  buf[m++] = h_sp(i,2);
-  buf[m++] = h_sp(i,3);
-
-  if (atom->nextra_grow)
-    for (int iextra = 0; iextra < atom->nextra_grow; iextra++)
-      m += modify->fix[atom->extra_grow[iextra]]->pack_exchange(i,&buf[m]);
-
-  buf[0] = m;
-  return m;
-}
-
-/* ---------------------------------------------------------------------- */
-
 template<class DeviceType>
 struct AtomVecSpinKokkos_UnpackExchangeFunctor {
   typedef DeviceType device_type;
@@ -886,321 +573,6 @@ int AtomVecSpinKokkos::unpack_exchange_kokkos(DAT::tdual_xfloat_2d &k_buf,int nr
 
 /* ---------------------------------------------------------------------- */
 
-int AtomVecSpinKokkos::unpack_exchange(double *buf)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) grow(0);
-  atomKK->modified(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-           MASK_MASK | IMAGE_MASK | SP_MASK);
-
-  int m = 1;
-  h_x(nlocal,0) = buf[m++];
-  h_x(nlocal,1) = buf[m++];
-  h_x(nlocal,2) = buf[m++];
-  h_v(nlocal,0) = buf[m++];
-  h_v(nlocal,1) = buf[m++];
-  h_v(nlocal,2) = buf[m++];
-  h_tag(nlocal) = (tagint) ubuf(buf[m++]).i;
-  h_type(nlocal) = (int) ubuf(buf[m++]).i;
-  h_mask(nlocal) = (int) ubuf(buf[m++]).i;
-  h_image(nlocal) = (imageint) ubuf(buf[m++]).i;
-  h_sp(nlocal,0) = buf[m++];
-  h_sp(nlocal,1) = buf[m++];
-  h_sp(nlocal,2) = buf[m++];
-  h_sp(nlocal,3) = buf[m++];
-
-  if (atom->nextra_grow)
-    for (int iextra = 0; iextra < atom->nextra_grow; iextra++)
-      m += modify->fix[atom->extra_grow[iextra]]->
-        unpack_exchange(nlocal,&buf[m]);
-
-  atom->nlocal++;
-  return m;
-}
-
-/* ----------------------------------------------------------------------
-   size of restart data for all atoms owned by this proc
-   include extra data stored by fixes
-------------------------------------------------------------------------- */
-
-int AtomVecSpinKokkos::size_restart()
-{
-  int i;
-
-  int nlocal = atom->nlocal;
-  int n = 15 * nlocal;
-
-  if (atom->nextra_restart)
-    for (int iextra = 0; iextra < atom->nextra_restart; iextra++)
-      for (i = 0; i < nlocal; i++)
-        n += modify->fix[atom->extra_restart[iextra]]->size_restart(i);
-
-  return n;
-}
-
-/* ----------------------------------------------------------------------
-   pack atom I's data for restart file including extra quantities
-   xyz must be 1st 3 values, so that read_restart can test on them
-   molecular types may be negative, but write as positive
-------------------------------------------------------------------------- */
-
-int AtomVecSpinKokkos::pack_restart(int i, double *buf)
-{
-  atomKK->sync(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-            MASK_MASK | IMAGE_MASK | SP_MASK);
-
-  int m = 1;
-  buf[m++] = h_x(i,0);
-  buf[m++] = h_x(i,1);
-  buf[m++] = h_x(i,2);
-  buf[m++] = ubuf(h_tag(i)).d;
-  buf[m++] = ubuf(h_type(i)).d;
-  buf[m++] = ubuf(h_mask(i)).d;
-  buf[m++] = ubuf(h_image(i)).d;
-  buf[m++] = h_v(i,0);
-  buf[m++] = h_v(i,1);
-  buf[m++] = h_v(i,2);
-
-  buf[m++] = h_sp(i,0);
-  buf[m++] = h_sp(i,1);
-  buf[m++] = h_sp(i,2);
-  buf[m++] = h_sp(i,3);
-
-  if (atom->nextra_restart)
-    for (int iextra = 0; iextra < atom->nextra_restart; iextra++)
-      m += modify->fix[atom->extra_restart[iextra]]->pack_restart(i,&buf[m]);
-
-  buf[0] = m;
-  return m;
-}
-
-/* ----------------------------------------------------------------------
-   unpack data for one atom from restart file including extra quantities
-------------------------------------------------------------------------- */
-
-int AtomVecSpinKokkos::unpack_restart(double *buf)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) {
-    grow(0);
-    if (atom->nextra_store)
-      memory->grow(atom->extra,nmax,atom->nextra_store,"atom:extra");
-  }
-
-  atomKK->modified(Host,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
-           MASK_MASK | IMAGE_MASK | SP_MASK);
-
-  int m = 1;
-  h_x(nlocal,0) = buf[m++];
-  h_x(nlocal,1) = buf[m++];
-  h_x(nlocal,2) = buf[m++];
-  h_tag(nlocal) = (tagint) ubuf(buf[m++]).i;
-  h_type(nlocal) = (int) ubuf(buf[m++]).i;
-  h_mask(nlocal) = (int) ubuf(buf[m++]).i;
-  h_image(nlocal) = (imageint) ubuf(buf[m++]).i;
-  h_v(nlocal,0) = buf[m++];
-  h_v(nlocal,1) = buf[m++];
-  h_v(nlocal,2) = buf[m++];
-
-  h_sp(nlocal,0) = buf[m++];
-  h_sp(nlocal,1) = buf[m++];
-  h_sp(nlocal,2) = buf[m++];
-  h_sp(nlocal,3) = buf[m++];
-
-  double **extra = atom->extra;
-  if (atom->nextra_store) {
-    int size = static_cast<int> (buf[0]) - m;
-    for (int i = 0; i < size; i++) extra[nlocal][i] = buf[m++];
-  }
-
-  atom->nlocal++;
-  return m;
-}
-
-/* ----------------------------------------------------------------------
-   create one atom of itype at coord
-   set other values to defaults
-------------------------------------------------------------------------- */
-
-void AtomVecSpinKokkos::create_atom(int itype, double *coord)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) {
-    atomKK->modified(Host,ALL_MASK);
-    grow(0);
-  }
-  atomKK->sync(Host,ALL_MASK);
-  atomKK->modified(Host,ALL_MASK);
-
-  h_tag[nlocal] = 0;
-  h_type[nlocal] = itype;
-  h_x(nlocal,0) = coord[0];
-  h_x(nlocal,1) = coord[1];
-  h_x(nlocal,2) = coord[2];
-  h_mask[nlocal] = 1;
-  h_image[nlocal] = ((imageint) IMGMAX << IMG2BITS) |
-    ((imageint) IMGMAX << IMGBITS) | IMGMAX;
-  h_v(nlocal,0) = 0.0;
-  h_v(nlocal,1) = 0.0;
-  h_v(nlocal,2) = 0.0;
-
-  h_sp(nlocal,0) = 0.0;
-  h_sp(nlocal,1) = 0.0;
-  h_sp(nlocal,2) = 0.0;
-  h_sp(nlocal,3) = 0.0;
-
-  atom->nlocal++;
-}
-
-/* ----------------------------------------------------------------------
-   unpack one line from Atoms section of data file
-   initialize other atom quantities
-------------------------------------------------------------------------- */
-
-void AtomVecSpinKokkos::data_atom(double *coord, imageint imagetmp,
-                                  const std::vector<std::string> &values, std::string &extract)
-{
-  int nlocal = atom->nlocal;
-  if (nlocal == nmax) grow(0);
-
-  h_tag[nlocal] = utils::inumeric(FLERR,values[0],true,lmp);
-  h_type[nlocal] = utils::inumeric(FLERR,values[1],true,lmp);
-  extract = values[1];
-  if (type[nlocal] <= 0 || type[nlocal] > atom->ntypes)
-    error->one(FLERR,"Invalid atom type in Atoms section of data file");
-
-  h_sp(nlocal,3) = utils::numeric(FLERR,values[2],true,lmp);
-  h_sp(nlocal,0) = utils::numeric(FLERR,values[6],true,lmp);
-  h_sp(nlocal,1) = utils::numeric(FLERR,values[7],true,lmp);
-  h_sp(nlocal,2) = utils::numeric(FLERR,values[8],true,lmp);
-  double inorm = 1.0/sqrt(sp[nlocal][0]*sp[nlocal][0] +
-                          sp[nlocal][1]*sp[nlocal][1] +
-                          sp[nlocal][2]*sp[nlocal][2]);
-  h_sp(nlocal,0) *= inorm;
-  h_sp(nlocal,1) *= inorm;
-  h_sp(nlocal,2) *= inorm;
-
-  h_x(nlocal,0) = coord[0];
-  h_x(nlocal,1) = coord[1];
-  h_x(nlocal,2) = coord[2];
-
-  h_image[nlocal] = imagetmp;
-
-  h_mask[nlocal] = 1;
-  h_v(nlocal,0) = 0.0;
-  h_v(nlocal,1) = 0.0;
-  h_v(nlocal,2) = 0.0;
-
-  atomKK->modified(Host,ALL_MASK);
-
-  atom->nlocal++;
-}
-
-/* ----------------------------------------------------------------------
-   unpack hybrid quantities from one line in Atoms section of data file
-   initialize other atom quantities for this sub-style
-------------------------------------------------------------------------- */
-
-int AtomVecSpinKokkos::data_atom_hybrid(int nlocal, const std::vector<std::string> &values,
-                                        int offset)
-{
-  h_sp(nlocal,3) = utils::numeric(FLERR,values[offset],true,lmp);
-  h_sp(nlocal,0) = utils::numeric(FLERR,values[offset+1],true,lmp);
-  h_sp(nlocal,1) = utils::numeric(FLERR,values[offset+2],true,lmp);
-  h_sp(nlocal,2) = utils::numeric(FLERR,values[offset+3],true,lmp);
-  double inorm = 1.0/sqrt(sp[nlocal][0]*sp[nlocal][0] +
-                          sp[nlocal][1]*sp[nlocal][1] +
-                          sp[nlocal][2]*sp[nlocal][2]);
-  sp[nlocal][0] *= inorm;
-  sp[nlocal][1] *= inorm;
-  sp[nlocal][2] *= inorm;
-
-  return 4;
-}
-
-/* ----------------------------------------------------------------------
-   pack atom info for data file including 3 image flags
-------------------------------------------------------------------------- */
-
-void AtomVecSpinKokkos::pack_data(double **buf)
-{
-  int nlocal = atom->nlocal;
-  for (int i = 0; i < nlocal; i++) {
-    buf[i][0] = h_tag[i];
-    buf[i][1] = h_type[i];
-    buf[i][2] = h_sp(i,0);
-    buf[i][3] = h_x(i,0);
-    buf[i][4] = h_x(i,1);
-    buf[i][5] = h_x(i,2);
-    buf[i][2] = h_sp(i,1);
-    buf[i][2] = h_sp(i,2);
-    buf[i][2] = h_sp(i,3);
-    buf[i][6] = (h_image[i] & IMGMASK) - IMGMAX;
-    buf[i][7] = (h_image[i] >> IMGBITS & IMGMASK) - IMGMAX;
-    buf[i][8] = (h_image[i] >> IMG2BITS) - IMGMAX;
-  }
-}
-
-/* ----------------------------------------------------------------------
-   pack hybrid atom info for data file
-------------------------------------------------------------------------- */
-
-int AtomVecSpinKokkos::pack_data_hybrid(int i, double *buf)
-{
-  buf[0] = h_sp(i,3);
-  buf[1] = h_sp(i,0);
-  buf[2] = h_sp(i,1);
-  buf[3] = h_sp(i,2);
-  return 4;
-}
-
-/* ----------------------------------------------------------------------
-   write atom info to data file including 3 image flags
-------------------------------------------------------------------------- */
-
-void AtomVecSpinKokkos::write_data(FILE *fp, int n, double **buf)
-{
-  for (int i = 0; i < n; i++)
-    fprintf(fp,"%d %d %-1.16e %-1.16e %-1.16e %-1.16e %d %d %d\n",
-            (int) buf[i][0],(int) buf[i][1],buf[i][2],buf[i][3],buf[i][4],
-            buf[i][5],(int) buf[i][6],(int) buf[i][7],(int) buf[i][8]);
-}
-
-/* ----------------------------------------------------------------------
-   write hybrid atom info to data file
-------------------------------------------------------------------------- */
-
-int AtomVecSpinKokkos::write_data_hybrid(FILE *fp, double *buf)
-{
-  fprintf(fp," %-1.16e %-1.16e %-1.16e %-1.16e",buf[0],buf[1],buf[2],buf[3]);
-  return 4;
-}
-
-/* ----------------------------------------------------------------------
-   return # of bytes of allocated memory
-------------------------------------------------------------------------- */
-
-double AtomVecSpinKokkos::memory_usage()
-{
-  bigint bytes = 0;
-
-  if (atom->memcheck("tag")) bytes += memory->usage(tag,nmax);
-  if (atom->memcheck("type")) bytes += memory->usage(type,nmax);
-  if (atom->memcheck("mask")) bytes += memory->usage(mask,nmax);
-  if (atom->memcheck("image")) bytes += memory->usage(image,nmax);
-  if (atom->memcheck("x")) bytes += memory->usage(x,nmax,3);
-  if (atom->memcheck("v")) bytes += memory->usage(v,nmax,3);
-  if (atom->memcheck("f")) bytes += memory->usage(f,nmax*commKK->nthreads,3);
-
-  if (atom->memcheck("sp")) bytes += memory->usage(sp,nmax,4);
-  if (atom->memcheck("fm")) bytes += memory->usage(fm,nmax*comm->nthreads,3);
-  if (atom->memcheck("fm_long")) bytes += memory->usage(fm_long,nmax*comm->nthreads,3);
-
-  return bytes;
-}
-
-/* ---------------------------------------------------------------------- */
-
 void AtomVecSpinKokkos::sync(ExecutionSpace space, unsigned int mask)
 {
   if (space == Device) {
@@ -1303,14 +675,3 @@ void AtomVecSpinKokkos::sync_overlapping_device(ExecutionSpace space, unsigned i
       perform_async_copy<DAT::tdual_f_array>(atomKK->k_fm_long,space);
   }
 }
-
-/* ----------------------------------------------------------------------
-   clear all forces (mech and mag)
-------------------------------------------------------------------------- */
-
-void AtomVecSpinKokkos::force_clear(int n, size_t nbytes)
-{
-  memset(&f[n][0],0,3*nbytes);
-  memset(&fm[n][0],0,3*nbytes);
-  memset(&fm_long[n][0],0,3*nbytes);
-}

src/KOKKOS/atom_vec_spin_kokkos.h

@@ -24,39 +24,15 @@ AtomStyle(spin/kk/host,AtomVecSpinKokkos);
 #define LMP_ATOM_VEC_SPIN_KOKKOS_H
 
 #include "atom_vec_kokkos.h"
+#include "atom_vec_spin.h"
 #include "kokkos_type.h"
 
 namespace LAMMPS_NS {
 
-class AtomVecSpinKokkos : public AtomVecKokkos {
+class AtomVecSpinKokkos : public AtomVecKokkos, public AtomVecSpin {
  public:
   AtomVecSpinKokkos(class LAMMPS *);
   void grow(int) override;
-  void copy(int, int, int) override;
-  int pack_border(int, int *, double *, int, int *) override;
-  int pack_border_vel(int, int *, double *, int, int *) override;
-  int pack_border_hybrid(int, int *, double *) override;
-  void unpack_border(int, int, double *) override;
-  void unpack_border_vel(int, int, double *) override;
-  int unpack_border_hybrid(int, int, double *) override;
-  int pack_exchange(int, double *) override;
-  int unpack_exchange(double *) override;
-  int size_restart() override;
-  int pack_restart(int, double *) override;
-  int unpack_restart(double *) override;
-  void create_atom(int, double *) override;
-  void data_atom(double *, imageint, const std::vector<std::string> &, std::string &) override;
-  int data_atom_hybrid(int, const std::vector<std::string> &, int) override;
-  void pack_data(double **) override;
-  int pack_data_hybrid(int, double *) override;
-  void write_data(FILE *, int, double **) override;
-  int write_data_hybrid(FILE *, double *) override;
-  double memory_usage() override;
-
-  // clear magnetic and mechanic forces
-
-  void force_clear(int, size_t) override;
-
   void grow_pointers() override;
   // input lists to be checked
   int pack_border_kokkos(int n, DAT::tdual_int_2d k_sendlist,
@@ -110,8 +86,6 @@ class AtomVecSpinKokkos : public AtomVecKokkos {
   HAT::t_sp_array h_sp;
   HAT::t_fm_array h_fm;
   HAT::t_fm_long_array h_fm_long;
-
-  DAT::tdual_int_1d k_count;
 };
 
 }

src/KOKKOS/comm_kokkos.cpp

@@ -146,8 +146,12 @@ void CommKokkos::init()
   if (!comm_f_only) // not all Kokkos atom_vec styles have reverse pack/unpack routines yet
     reverse_comm_classic = true;
 
-  if (ghost_velocity && ((AtomVecKokkos*)atom->avec)->no_comm_vel_flag) // not all Kokkos atom_vec styles have comm vel pack/unpack routines yet
+  atomKK->avecKK = dynamic_cast<AtomVecKokkos*>(atom->avec);
+
+  if (ghost_velocity && atomKK->avecKK->no_comm_vel_flag) // not all Kokkos atom_vec styles have comm vel pack/unpack routines yet
     forward_comm_classic = true;
+
+
 }
 
 /* ----------------------------------------------------------------------
@@ -186,7 +190,6 @@ void CommKokkos::forward_comm_device(int)
 {
   int n;
   MPI_Request request;
-  AtomVecKokkos *avec = (AtomVecKokkos *) atom->avec;
   double *buf;
 
   // exchange data with another proc
@@ -200,7 +203,7 @@ void CommKokkos::forward_comm_device(int)
     k_swap.sync<DeviceType>();
     k_swap2.sync<DeviceType>();
     k_pbc.sync<DeviceType>();
-    n = avec->pack_comm_self_fused(totalsend,k_sendlist,k_sendnum_scan,
+    n = atomKK->avecKK->pack_comm_self_fused(totalsend,k_sendlist,k_sendnum_scan,
                     k_firstrecv,k_pbc_flag,k_pbc,k_g2l);
   } else {
 
@@ -213,7 +216,7 @@ void CommKokkos::forward_comm_device(int)
             MPI_Irecv(buf,size_forward_recv[iswap],MPI_DOUBLE,
                       recvproc[iswap],0,world,&request);
           }
-          n = avec->pack_comm_kokkos(sendnum[iswap],k_sendlist,
+          n = atomKK->avecKK->pack_comm_kokkos(sendnum[iswap],k_sendlist,
                                      iswap,k_buf_send,pbc_flag[iswap],pbc[iswap]);
           DeviceType().fence();
           if (n) {
@@ -232,7 +235,7 @@ void CommKokkos::forward_comm_device(int)
                       size_forward_recv[iswap],MPI_DOUBLE,
                       recvproc[iswap],0,world,&request);
           }
-          n = avec->pack_comm_vel_kokkos(sendnum[iswap],k_sendlist,iswap,
+          n = atomKK->avecKK->pack_comm_vel_kokkos(sendnum[iswap],k_sendlist,iswap,
                                          k_buf_send,pbc_flag[iswap],pbc[iswap]);
           DeviceType().fence();
           if (n) {
@@ -240,34 +243,34 @@ void CommKokkos::forward_comm_device(int)
                      MPI_DOUBLE,sendproc[iswap],0,world);
           }
           if (size_forward_recv[iswap]) MPI_Wait(&request,MPI_STATUS_IGNORE);
-          avec->unpack_comm_vel_kokkos(recvnum[iswap],firstrecv[iswap],k_buf_recv);
+          atomKK->avecKK->unpack_comm_vel_kokkos(recvnum[iswap],firstrecv[iswap],k_buf_recv);
           DeviceType().fence();
         } else {
           if (size_forward_recv[iswap])
             MPI_Irecv(k_buf_recv.view<DeviceType>().data(),
                       size_forward_recv[iswap],MPI_DOUBLE,
                       recvproc[iswap],0,world,&request);
-          n = avec->pack_comm_kokkos(sendnum[iswap],k_sendlist,iswap,
+          n = atomKK->avecKK->pack_comm_kokkos(sendnum[iswap],k_sendlist,iswap,
                                      k_buf_send,pbc_flag[iswap],pbc[iswap]);
           DeviceType().fence();
           if (n)
             MPI_Send(k_buf_send.view<DeviceType>().data(),n,
                      MPI_DOUBLE,sendproc[iswap],0,world);
           if (size_forward_recv[iswap]) MPI_Wait(&request,MPI_STATUS_IGNORE);
-          avec->unpack_comm_kokkos(recvnum[iswap],firstrecv[iswap],k_buf_recv);
+          atomKK->avecKK->unpack_comm_kokkos(recvnum[iswap],firstrecv[iswap],k_buf_recv);
           DeviceType().fence();
         }
       } else {
         if (!ghost_velocity) {
           if (sendnum[iswap])
-            n = avec->pack_comm_self(sendnum[iswap],k_sendlist,iswap,
+            n = atomKK->avecKK->pack_comm_self(sendnum[iswap],k_sendlist,iswap,
                                      firstrecv[iswap],pbc_flag[iswap],pbc[iswap]);
           DeviceType().fence();
         } else {
-          n = avec->pack_comm_vel_kokkos(sendnum[iswap],k_sendlist,iswap,
+          n = atomKK->avecKK->pack_comm_vel_kokkos(sendnum[iswap],k_sendlist,iswap,
                                          k_buf_send,pbc_flag[iswap],pbc[iswap]);
           DeviceType().fence();
-          avec->unpack_comm_vel_kokkos(recvnum[iswap],firstrecv[iswap],k_buf_send);
+          atomKK->avecKK->unpack_comm_vel_kokkos(recvnum[iswap],firstrecv[iswap],k_buf_send);
           DeviceType().fence();
         }
       }
@@ -310,7 +313,6 @@ void CommKokkos::reverse_comm_device()
 {
   int n;
   MPI_Request request;
-  AtomVecKokkos *avec = (AtomVecKokkos *) atom->avec;
   double *buf;
 
   // exchange data with another proc
@@ -343,19 +345,19 @@ void CommKokkos::reverse_comm_device()
           MPI_Irecv(k_buf_recv.view<DeviceType>().data(),
                     size_reverse_recv[iswap],MPI_DOUBLE,
                     sendproc[iswap],0,world,&request);
-        n = avec->pack_reverse_kokkos(recvnum[iswap],firstrecv[iswap],k_buf_send);
+        n = atomKK->avecKK->pack_reverse_kokkos(recvnum[iswap],firstrecv[iswap],k_buf_send);
         DeviceType().fence();
         if (n)
           MPI_Send(k_buf_send.view<DeviceType>().data(),n,
                    MPI_DOUBLE,recvproc[iswap],0,world);
         if (size_reverse_recv[iswap]) MPI_Wait(&request,MPI_STATUS_IGNORE);
       }
-      avec->unpack_reverse_kokkos(sendnum[iswap],k_sendlist,iswap,
+      atomKK->avecKK->unpack_reverse_kokkos(sendnum[iswap],k_sendlist,iswap,
                                 k_buf_recv);
       DeviceType().fence();
     } else {
       if (sendnum[iswap])
-        n = avec->unpack_reverse_self(sendnum[iswap],k_sendlist,iswap,
+        n = atomKK->avecKK->unpack_reverse_self(sendnum[iswap],k_sendlist,iswap,
                                  firstrecv[iswap]);
     }
   }
@@ -717,7 +719,6 @@ void CommKokkos::exchange_device()
   double **x;
   double *sublo,*subhi;
   MPI_Request request;
-  AtomVecKokkos *avec = (AtomVecKokkos *) atom->avec;
 
   // clear global->local map for owned and ghost atoms
   // b/c atoms migrate to new procs in exchange() and
@@ -800,7 +801,7 @@ void CommKokkos::exchange_device()
         nsend = k_count.h_view();
         if (nsend > maxsend) grow_send_kokkos(nsend,1);
         nsend =
-          avec->pack_exchange_kokkos(k_count.h_view(),k_buf_send,
+          atomKK->avecKK->pack_exchange_kokkos(k_count.h_view(),k_buf_send,
                                      k_exchange_sendlist,k_exchange_copylist,
                                      ExecutionSpaceFromDevice<DeviceType>::space,
                                      dim,lo,hi);
@@ -809,8 +810,8 @@ void CommKokkos::exchange_device()
         while (i < nlocal) {
           if (x[i][dim] < lo || x[i][dim] >= hi) {
             if (nsend > maxsend) grow_send_kokkos(nsend,1);
-            nsend += avec->pack_exchange(i,&buf_send[nsend]);
-            avec->copy(nlocal-1,i,1);
+            nsend += atomKK->avecKK->pack_exchange(i,&buf_send[nsend]);
+            atomKK->avecKK->copy(nlocal-1,i,1);
             nlocal--;
           } else i++;
         }
@@ -825,7 +826,7 @@ void CommKokkos::exchange_device()
       if (procgrid[dim] == 1) {
         nrecv = nsend;
         if (nrecv) {
-          atom->nlocal=avec->
+          atom->nlocal=atomKK->avecKK->
             unpack_exchange_kokkos(k_buf_send,nrecv,atom->nlocal,dim,lo,hi,
                                    ExecutionSpaceFromDevice<DeviceType>::space);
           DeviceType().fence();
@@ -858,7 +859,7 @@ void CommKokkos::exchange_device()
         }
 
         if (nrecv) {
-          atom->nlocal = avec->
+          atom->nlocal = atomKK->avecKK->
             unpack_exchange_kokkos(k_buf_recv,nrecv,atom->nlocal,dim,lo,hi,
                                    ExecutionSpaceFromDevice<DeviceType>::space);
           DeviceType().fence();
@@ -896,7 +897,7 @@ void CommKokkos::borders()
     static int print = 1;
 
     if (mode != Comm::SINGLE || bordergroup ||
-         (ghost_velocity && ((AtomVecKokkos*)atom->avec)->no_border_vel_flag)) {
+         (ghost_velocity && atomKK->avecKK->no_border_vel_flag)) {
       if (print && comm->me==0) {
         error->warning(FLERR,"Required border comm not yet implemented in Kokkos communication, "
                       "switching to classic exchange/border communication");
@@ -982,7 +983,6 @@ void CommKokkos::borders_device() {
   double **x;
   double *mlo,*mhi;
   MPI_Request request;
-  AtomVecKokkos *avec = (AtomVecKokkos *) atom->avec;
 
   ExecutionSpace exec_space = ExecutionSpaceFromDevice<DeviceType>::space;
   atomKK->sync(exec_space,ALL_MASK);
@@ -1122,13 +1122,13 @@ void CommKokkos::borders_device() {
       if (nsend*size_border > maxsend)
         grow_send_kokkos(nsend*size_border,0);
       if (ghost_velocity) {
-        n = avec->
+        n = atomKK->avecKK->
           pack_border_vel_kokkos(nsend,k_sendlist,k_buf_send,iswap,
                                  pbc_flag[iswap],pbc[iswap],exec_space);
         DeviceType().fence();
       }
       else {
-        n = avec->
+        n = atomKK->avecKK->
           pack_border_kokkos(nsend,k_sendlist,k_buf_send,iswap,
                              pbc_flag[iswap],pbc[iswap],exec_space);
         DeviceType().fence();
@@ -1157,21 +1157,21 @@ void CommKokkos::borders_device() {
 
       if (ghost_velocity) {
         if (sendproc[iswap] != me) {
-          avec->unpack_border_vel_kokkos(nrecv,atom->nlocal+atom->nghost,
+          atomKK->avecKK->unpack_border_vel_kokkos(nrecv,atom->nlocal+atom->nghost,
                                          k_buf_recv,exec_space);
           DeviceType().fence();
         } else {
-          avec->unpack_border_vel_kokkos(nrecv,atom->nlocal+atom->nghost,
+          atomKK->avecKK->unpack_border_vel_kokkos(nrecv,atom->nlocal+atom->nghost,
                                          k_buf_send,exec_space);
           DeviceType().fence();
         }
       } else {
         if (sendproc[iswap] != me) {
-          avec->unpack_border_kokkos(nrecv,atom->nlocal+atom->nghost,
+          atomKK->avecKK->unpack_border_kokkos(nrecv,atom->nlocal+atom->nghost,
                                      k_buf_recv,exec_space);
           DeviceType().fence();
         } else {
-          avec->unpack_border_kokkos(nrecv,atom->nlocal+atom->nghost,
+          atomKK->avecKK->unpack_border_kokkos(nrecv,atom->nlocal+atom->nghost,
                                      k_buf_send,exec_space);
           DeviceType().fence();
         }
@@ -1299,7 +1299,7 @@ void CommKokkos::grow_recv(int n)
 void CommKokkos::grow_send_kokkos(int n, int flag, ExecutionSpace space)
 {
   maxsend = static_cast<int> (BUFFACTOR * n);
-  int maxsend_border = (maxsend+BUFEXTRA+5)/atom->avec->size_border + 2;
+  int maxsend_border = (maxsend+BUFEXTRA+5)/atomKK->avecKK->size_border + 2;
   if (flag) {
     if (space == Device)
       k_buf_send.modify<LMPDeviceType>();
@@ -1308,9 +1308,9 @@ void CommKokkos::grow_send_kokkos(int n, int flag, ExecutionSpace space)
 
     if (ghost_velocity)
       k_buf_send.resize(maxsend_border,
-                        atom->avec->size_border + atom->avec->size_velocity);
+                        atomKK->avecKK->size_border + atomKK->avecKK->size_velocity);
     else
-      k_buf_send.resize(maxsend_border,atom->avec->size_border);
+      k_buf_send.resize(maxsend_border,atomKK->avecKK->size_border);
     buf_send = k_buf_send.view<LMPHostType>().data();
   }
   else {
@@ -1318,10 +1318,10 @@ void CommKokkos::grow_send_kokkos(int n, int flag, ExecutionSpace space)
       k_buf_send = DAT::
         tdual_xfloat_2d("comm:k_buf_send",
                         maxsend_border,
-                        atom->avec->size_border + atom->avec->size_velocity);
+                        atomKK->avecKK->size_border + atomKK->avecKK->size_velocity);
     else
       k_buf_send = DAT::
-        tdual_xfloat_2d("comm:k_buf_send",maxsend_border,atom->avec->size_border);
+        tdual_xfloat_2d("comm:k_buf_send",maxsend_border,atomKK->avecKK->size_border);
     buf_send = k_buf_send.view<LMPHostType>().data();
   }
 }
@@ -1333,9 +1333,9 @@ void CommKokkos::grow_send_kokkos(int n, int flag, ExecutionSpace space)
 void CommKokkos::grow_recv_kokkos(int n, ExecutionSpace /*space*/)
 {
   maxrecv = static_cast<int> (BUFFACTOR * n);
-  int maxrecv_border = (maxrecv+BUFEXTRA+5)/atom->avec->size_border + 2;
+  int maxrecv_border = (maxrecv+BUFEXTRA+5)/atomKK->avecKK->size_border + 2;
   k_buf_recv = DAT::
-    tdual_xfloat_2d("comm:k_buf_recv",maxrecv_border,atom->avec->size_border);
+    tdual_xfloat_2d("comm:k_buf_recv",maxrecv_border,atomKK->avecKK->size_border);
   buf_recv = k_buf_recv.view<LMPHostType>().data();
 }
 

src/KOKKOS/comm_kokkos.h

@@ -96,6 +96,8 @@ class CommKokkos : public CommBrick {
   void grow_list(int, int) override;
   void grow_swap(int) override;
   void copy_swap_info();
+
+  class AtomVecKokkos* avec_kk;
 };
 
 }

src/MOLECULE/atom_vec_angle.h

@@ -24,7 +24,7 @@ AtomStyle(angle,AtomVecAngle);
 
 namespace LAMMPS_NS {
 
-class AtomVecAngle : public AtomVec {
+class AtomVecAngle : virtual public AtomVec {
  public:
   AtomVecAngle(class LAMMPS *);
   ~AtomVecAngle() override;

src/MOLECULE/atom_vec_bond.h

@@ -24,7 +24,7 @@ AtomStyle(bond,AtomVecBond);
 
 namespace LAMMPS_NS {
 
-class AtomVecBond : public AtomVec {
+class AtomVecBond : virtual public AtomVec {
  public:
   AtomVecBond(class LAMMPS *);
   ~AtomVecBond() override;

src/MOLECULE/atom_vec_full.h

@@ -24,7 +24,7 @@ AtomStyle(full,AtomVecFull);
 
 namespace LAMMPS_NS {
 
-class AtomVecFull : public AtomVec {
+class AtomVecFull : virtual public AtomVec {
  public:
   AtomVecFull(class LAMMPS *);
   ~AtomVecFull() override;

src/MOLECULE/atom_vec_molecular.h

@@ -24,7 +24,7 @@ AtomStyle(molecular,AtomVecMolecular);
 
 namespace LAMMPS_NS {
 
-class AtomVecMolecular : public AtomVec {
+class AtomVecMolecular : virtual public AtomVec {
  public:
   AtomVecMolecular(class LAMMPS *);
   ~AtomVecMolecular() override;

src/SPIN/atom_vec_spin.h

@@ -24,7 +24,7 @@ AtomStyle(spin,AtomVecSpin);
 
 namespace LAMMPS_NS {
 
-class AtomVecSpin : public AtomVec {
+class AtomVecSpin : virtual public AtomVec {
  public:
   AtomVecSpin(class LAMMPS *);
 

src/atom_vec.h

@@ -157,21 +157,6 @@ class AtomVec : protected Pointers {
   virtual double memory_usage();
   virtual double memory_usage_bonus() { return 0; }
 
-  // old hybrid functions, needed by Kokkos package
-
-  virtual int pack_comm_hybrid(int, int *, double *) { return 0; }
-  virtual int unpack_comm_hybrid(int, int, double *) { return 0; }
-  virtual int pack_reverse_hybrid(int, int, double *) { return 0; }
-  virtual int unpack_reverse_hybrid(int, int *, double *) { return 0; }
-  virtual int pack_border_hybrid(int, int *, double *) { return 0; }
-  virtual int unpack_border_hybrid(int, int, double *) { return 0; }
-  virtual int data_atom_hybrid(int, const std::vector<std::string> &, int) { return 0; }
-  virtual int data_vel_hybrid(int, const std::vector<std::string> &, int) { return 0; }
-  virtual int pack_data_hybrid(int, double *) { return 0; }
-  virtual int write_data_hybrid(FILE *, double *) { return 0; }
-  virtual int pack_vel_hybrid(int, double *) { return 0; }
-  virtual int write_vel_hybrid(FILE *, double *) { return 0; }
-
  protected:
   int nmax;             // local copy of atom->nmax
   int deform_vremap;    // local copy of domain properties

src/atom_vec_atomic.h

@@ -24,7 +24,7 @@ AtomStyle(atomic,AtomVecAtomic);
 
 namespace LAMMPS_NS {
 
-class AtomVecAtomic : public AtomVec {
+class AtomVecAtomic : virtual public AtomVec {
  public:
   AtomVecAtomic(class LAMMPS *);
 };

src/atom_vec_charge.h

@@ -24,7 +24,7 @@ AtomStyle(charge,AtomVecCharge);
 
 namespace LAMMPS_NS {
 
-class AtomVecCharge : public AtomVec {
+class AtomVecCharge : virtual public AtomVec {
  public:
   AtomVecCharge(class LAMMPS *);
 };

src/atom_vec_hybrid.h

@@ -24,7 +24,7 @@ AtomStyle(hybrid,AtomVecHybrid);
 
 namespace LAMMPS_NS {
 
-class AtomVecHybrid : public AtomVec {
+class AtomVecHybrid : virtual public AtomVec {
  public:
   int nstyles;
   class AtomVec **styles;

src/atom_vec_sphere.h

@@ -24,7 +24,7 @@ AtomStyle(sphere,AtomVecSphere);
 
 namespace LAMMPS_NS {
 
-class AtomVecSphere : public AtomVec {
+class AtomVecSphere : virtual public AtomVec {
  public:
   AtomVecSphere(class LAMMPS *);
   void process_args(int, char **) override;