From c55901f8ce2cdcf512b92e421035ce111902bf1c Mon Sep 17 00:00:00 2001
From: Stan Moore <stamoor@sandia.gov>
Date: Mon, 20 May 2024 14:28:31 -0600
Subject: [PATCH 1/5] Port hybrid topology (bond/angle/dihedral/improper)
 styles to Kokkos

---
 src/KOKKOS/Install.sh               |   2 +
 src/KOKKOS/bond_class2_kokkos.cpp   |  22 +--
 src/KOKKOS/bond_class2_kokkos.h     |   7 +-
 src/KOKKOS/bond_fene_kokkos.cpp     |  22 +--
 src/KOKKOS/bond_fene_kokkos.h       |   5 +-
 src/KOKKOS/bond_harmonic_kokkos.cpp |  25 ++--
 src/KOKKOS/bond_harmonic_kokkos.h   |   2 +-
 src/KOKKOS/bond_hybrid_kokkos.cpp   | 214 ++++++++++++++++++++++++++++
 src/KOKKOS/bond_hybrid_kokkos.h     |  58 ++++++++
 src/KOKKOS/kokkos_type.h            |   1 +
 src/KOKKOS/pair_hybrid_kokkos.cpp   |   4 +-
 src/bond.cpp                        |   3 +-
 src/bond.h                          |   2 +-
 src/bond_hybrid.cpp                 |  37 ++---
 src/bond_hybrid.h                   |   5 +-
 15 files changed, 344 insertions(+), 65 deletions(-)
 create mode 100644 src/KOKKOS/bond_hybrid_kokkos.cpp
 create mode 100644 src/KOKKOS/bond_hybrid_kokkos.h

diff --git a/src/KOKKOS/Install.sh b/src/KOKKOS/Install.sh
index 75949c35d8a..86acbe97b29 100755
--- a/src/KOKKOS/Install.sh
+++ b/src/KOKKOS/Install.sh
@@ -86,6 +86,8 @@ action bond_fene_kokkos.cpp bond_fene.cpp
 action bond_fene_kokkos.h bond_fene.h
 action bond_harmonic_kokkos.cpp bond_harmonic.cpp
 action bond_harmonic_kokkos.h bond_harmonic.h
+action bond_hybrid_kokkos.cpp bond_hybrid.cpp
+action bond_hybrid_kokkos.h bond_hybrid.h
 action comm_kokkos.cpp
 action comm_kokkos.h
 action comm_tiled_kokkos.cpp
diff --git a/src/KOKKOS/bond_class2_kokkos.cpp b/src/KOKKOS/bond_class2_kokkos.cpp
index eeecc385b02..2cf113ec009 100644
--- a/src/KOKKOS/bond_class2_kokkos.cpp
+++ b/src/KOKKOS/bond_class2_kokkos.cpp
@@ -34,6 +34,8 @@ using namespace LAMMPS_NS;
 template<class DeviceType>
 BondClass2Kokkos<DeviceType>::BondClass2Kokkos(LAMMPS *lmp) : BondClass2(lmp)
 {
+  kokkosable = 1;
+
   atomKK = (AtomKokkos *) atom;
   neighborKK = (NeighborKokkos *) neighbor;
   execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
@@ -122,12 +124,12 @@ void BondClass2Kokkos<DeviceType>::compute(int eflag_in, int vflag_in)
 
   if (eflag_atom) {
     k_eatom.template modify<DeviceType>();
-    k_eatom.template sync<LMPHostType>();
+    k_eatom.sync_host();
   }
 
   if (vflag_atom) {
     k_vatom.template modify<DeviceType>();
-    k_vatom.template sync<LMPHostType>();
+    k_vatom.sync_host();
   }
 
   copymode = 0;
@@ -227,13 +229,13 @@ void BondClass2Kokkos<DeviceType>::coeff(int narg, char **arg)
     k_r0.h_view[i] = r0[i];
   }
 
-  k_k2.template modify<LMPHostType>();
+  k_k2.modify_host();
   k_k2.template sync<DeviceType>();
-  k_k3.template modify<LMPHostType>();
+  k_k3.modify_host();
   k_k3.template sync<DeviceType>();
-  k_k4.template modify<LMPHostType>();
+  k_k4.modify_host();
   k_k4.template sync<DeviceType>();
-  k_r0.template modify<LMPHostType>();
+  k_r0.modify_host();
   k_r0.template sync<DeviceType>();
 }
 
@@ -264,13 +266,13 @@ void BondClass2Kokkos<DeviceType>::read_restart(FILE *fp)
     k_r0.h_view[i] = r0[i];
   }
 
-  k_k2.template modify<LMPHostType>();
+  k_k2.modify_host();
   k_k2.template sync<DeviceType>();
-  k_k3.template modify<LMPHostType>();
+  k_k3.modify_host();
   k_k3.template sync<DeviceType>();
-  k_k4.template modify<LMPHostType>();
+  k_k4.modify_host();
   k_k4.template sync<DeviceType>();
-  k_r0.template modify<LMPHostType>();
+  k_r0.modify_host();
   k_r0.template sync<DeviceType>();
 }
 
diff --git a/src/KOKKOS/bond_class2_kokkos.h b/src/KOKKOS/bond_class2_kokkos.h
index 0eb3c4d39a4..1056bf021dd 100644
--- a/src/KOKKOS/bond_class2_kokkos.h
+++ b/src/KOKKOS/bond_class2_kokkos.h
@@ -59,6 +59,10 @@ class BondClass2Kokkos : public BondClass2 {
       const F_FLOAT &ebond, const F_FLOAT &fbond, const F_FLOAT &delx,
                   const F_FLOAT &dely, const F_FLOAT &delz) const;
 
+  typedef typename KKDevice<DeviceType>::value KKDeviceType;
+  Kokkos::DualView<E_FLOAT*,Kokkos::LayoutRight,KKDeviceType> k_eatom;
+  Kokkos::DualView<F_FLOAT*[6],Kokkos::LayoutRight,KKDeviceType> k_vatom;
+
  protected:
 
   class NeighborKokkos *neighborKK;
@@ -67,9 +71,6 @@ class BondClass2Kokkos : public BondClass2 {
   typename Kokkos::View<double*[3],typename AT::t_f_array::array_layout,typename KKDevice<DeviceType>::value,Kokkos::MemoryTraits<Kokkos::Atomic> > f;
   typename AT::t_int_2d bondlist;
 
-  typedef typename KKDevice<DeviceType>::value KKDeviceType;
-  Kokkos::DualView<E_FLOAT*,Kokkos::LayoutRight,KKDeviceType> k_eatom;
-  Kokkos::DualView<F_FLOAT*[6],Kokkos::LayoutRight,KKDeviceType> k_vatom;
   Kokkos::View<E_FLOAT*,Kokkos::LayoutRight,KKDeviceType,Kokkos::MemoryTraits<Kokkos::Atomic> > d_eatom;
   Kokkos::View<F_FLOAT*[6],Kokkos::LayoutRight,KKDeviceType,Kokkos::MemoryTraits<Kokkos::Atomic> > d_vatom;
 
diff --git a/src/KOKKOS/bond_fene_kokkos.cpp b/src/KOKKOS/bond_fene_kokkos.cpp
index a6bbb1eddea..a056f35fde7 100644
--- a/src/KOKKOS/bond_fene_kokkos.cpp
+++ b/src/KOKKOS/bond_fene_kokkos.cpp
@@ -37,6 +37,8 @@ using MathConst::MY_CUBEROOT2;
 template<class DeviceType>
 BondFENEKokkos<DeviceType>::BondFENEKokkos(LAMMPS *lmp) : BondFENE(lmp)
 {
+  kokkosable = 1;
+
   atomKK = (AtomKokkos *) atom;
   neighborKK = (NeighborKokkos *) neighbor;
   execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
@@ -135,12 +137,12 @@ void BondFENEKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
 
   if (eflag_atom) {
     k_eatom.template modify<DeviceType>();
-    k_eatom.template sync<LMPHostType>();
+    k_eatom.sync_host();
   }
 
   if (vflag_atom) {
     k_vatom.template modify<DeviceType>();
-    k_vatom.template sync<LMPHostType>();
+    k_vatom.sync_host();
   }
 
   copymode = 0;
@@ -267,10 +269,10 @@ void BondFENEKokkos<DeviceType>::coeff(int narg, char **arg)
     k_sigma.h_view[i] = sigma[i];
   }
 
-  k_k.template modify<LMPHostType>();
-  k_r0.template modify<LMPHostType>();
-  k_epsilon.template modify<LMPHostType>();
-  k_sigma.template modify<LMPHostType>();
+  k_k.modify_host();
+  k_r0.modify_host();
+  k_epsilon.modify_host();
+  k_sigma.modify_host();
 }
 
 
@@ -291,10 +293,10 @@ void BondFENEKokkos<DeviceType>::read_restart(FILE *fp)
     k_sigma.h_view[i] = sigma[i];
   }
 
-  k_k.template modify<LMPHostType>();
-  k_r0.template modify<LMPHostType>();
-  k_epsilon.template modify<LMPHostType>();
-  k_sigma.template modify<LMPHostType>();
+  k_k.modify_host();
+  k_r0.modify_host();
+  k_epsilon.modify_host();
+  k_sigma.modify_host();
 }
 
 /* ----------------------------------------------------------------------
diff --git a/src/KOKKOS/bond_fene_kokkos.h b/src/KOKKOS/bond_fene_kokkos.h
index 18f8d87b6aa..450490d6a9f 100644
--- a/src/KOKKOS/bond_fene_kokkos.h
+++ b/src/KOKKOS/bond_fene_kokkos.h
@@ -58,6 +58,9 @@ class BondFENEKokkos : public BondFENE {
       const F_FLOAT &ebond, const F_FLOAT &fbond, const F_FLOAT &delx,
                   const F_FLOAT &dely, const F_FLOAT &delz) const;
 
+  DAT::tdual_efloat_1d k_eatom;
+  DAT::tdual_virial_array k_vatom;
+
  protected:
 
   class NeighborKokkos *neighborKK;
@@ -66,8 +69,6 @@ class BondFENEKokkos : public BondFENE {
   typename ArrayTypes<DeviceType>::t_f_array f;
   typename ArrayTypes<DeviceType>::t_int_2d bondlist;
 
-  DAT::tdual_efloat_1d k_eatom;
-  DAT::tdual_virial_array k_vatom;
   typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
   typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
 
diff --git a/src/KOKKOS/bond_harmonic_kokkos.cpp b/src/KOKKOS/bond_harmonic_kokkos.cpp
index 7dcdf24c663..7e12400c9bc 100644
--- a/src/KOKKOS/bond_harmonic_kokkos.cpp
+++ b/src/KOKKOS/bond_harmonic_kokkos.cpp
@@ -34,6 +34,8 @@ using namespace LAMMPS_NS;
 template<class DeviceType>
 BondHarmonicKokkos<DeviceType>::BondHarmonicKokkos(LAMMPS *lmp) : BondHarmonic(lmp)
 {
+  kokkosable = 1;
+
   atomKK = (AtomKokkos *) atom;
   neighborKK = (NeighborKokkos *) neighbor;
   execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
@@ -65,23 +67,20 @@ void BondHarmonicKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
   // reallocate per-atom arrays if necessary
 
   if (eflag_atom) {
-    //if(k_eatom.extent(0)<maxeatom) { // won't work without adding zero functor
+    if (k_eatom.extent(0) < maxeatom) {
       memoryKK->destroy_kokkos(k_eatom,eatom);
       memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom");
       d_eatom = k_eatom.template view<KKDeviceType>();
-    //}
+    } else Kokkos::deep_copy(d_eatom,0.0);
   }
   if (vflag_atom) {
-    //if(k_vatom.extent(0)<maxvatom) { // won't work without adding zero functor
+    if (k_vatom.extent(0) < maxvatom) {
       memoryKK->destroy_kokkos(k_vatom,vatom);
       memoryKK->create_kokkos(k_vatom,vatom,maxvatom,"improper:vatom");
       d_vatom = k_vatom.template view<KKDeviceType>();
-    //}
+    } else Kokkos::deep_copy(d_vatom,0.0);
   }
 
-//  if (eflag || vflag) atomKK->modified(execution_space,datamask_modify);
-//  else atomKK->modified(execution_space,F_MASK);
-
   x = atomKK->k_x.template view<DeviceType>();
   f = atomKK->k_f.template view<DeviceType>();
   neighborKK->k_bondlist.template sync<DeviceType>();
@@ -122,12 +121,12 @@ void BondHarmonicKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
 
   if (eflag_atom) {
     k_eatom.template modify<DeviceType>();
-    k_eatom.template sync<LMPHostType>();
+    k_eatom.sync_host();
   }
 
   if (vflag_atom) {
     k_vatom.template modify<DeviceType>();
-    k_vatom.template sync<LMPHostType>();
+    k_vatom.sync_host();
   }
 
   copymode = 0;
@@ -214,8 +213,8 @@ void BondHarmonicKokkos<DeviceType>::coeff(int narg, char **arg)
     k_r0.h_view[i] = r0[i];
   }
 
-  k_k.template modify<LMPHostType>();
-  k_r0.template modify<LMPHostType>();
+  k_k.modify_host();
+  k_r0.modify_host();
   k_k.template sync<DeviceType>();
   k_r0.template sync<DeviceType>();
 }
@@ -241,8 +240,8 @@ void BondHarmonicKokkos<DeviceType>::read_restart(FILE *fp)
     k_r0.h_view[i] = r0[i];
   }
 
-  k_k.template modify<LMPHostType>();
-  k_r0.template modify<LMPHostType>();
+  k_k.modify_host();
+  k_r0.modify_host();
   k_k.template sync<DeviceType>();
   k_r0.template sync<DeviceType>();
 }
diff --git a/src/KOKKOS/bond_harmonic_kokkos.h b/src/KOKKOS/bond_harmonic_kokkos.h
index 837e669b98c..b613dbc3100 100644
--- a/src/KOKKOS/bond_harmonic_kokkos.h
+++ b/src/KOKKOS/bond_harmonic_kokkos.h
@@ -37,6 +37,7 @@ class BondHarmonicKokkos : public BondHarmonic {
  public:
   typedef DeviceType device_type;
   typedef EV_FLOAT value_type;
+  typedef ArrayTypes<DeviceType> AT;
 
   BondHarmonicKokkos(class LAMMPS *);
   ~BondHarmonicKokkos() override;
@@ -62,7 +63,6 @@ class BondHarmonicKokkos : public BondHarmonic {
 
   class NeighborKokkos *neighborKK;
 
-  typedef ArrayTypes<DeviceType> AT;
   typename AT::t_x_array_randomread x;
   typename Kokkos::View<double*[3],typename AT::t_f_array::array_layout,typename KKDevice<DeviceType>::value,Kokkos::MemoryTraits<Kokkos::Atomic> > f;
   typename AT::t_int_2d bondlist;
diff --git a/src/KOKKOS/bond_hybrid_kokkos.cpp b/src/KOKKOS/bond_hybrid_kokkos.cpp
new file mode 100644
index 00000000000..65208f8d8cf
--- /dev/null
+++ b/src/KOKKOS/bond_hybrid_kokkos.cpp
@@ -0,0 +1,214 @@
+/* ----------------------------------------------------------------------
+   LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
+   https://www.lammps.org/, Sandia National Laboratories
+   LAMMPS development team: developers@lammps.org
+
+   Copyright (2003) Sandia Corporation.  Under the terms of Contract
+   DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
+   certain rights in this software.  This software is distributed under
+   the GNU General Public License.
+
+   See the README file in the top-level LAMMPS directory.
+------------------------------------------------------------------------- */
+
+#include "bond_hybrid_kokkos.h"
+
+#include "atom_kokkos.h"
+#include "atom_masks.h"
+#include "comm.h"
+#include "error.h"
+#include "force.h"
+#include "kokkos.h"
+#include "memory_kokkos.h"
+#include "neighbor_kokkos.h"
+
+#include <cstring>
+
+using namespace LAMMPS_NS;
+
+#define EXTRA 1000
+
+/* ---------------------------------------------------------------------- */
+
+BondHybridKokkos::BondHybridKokkos(LAMMPS *lmp) : BondHybrid(lmp)
+{
+  kokkosable = 1;
+
+  atomKK = (AtomKokkos *) atom;
+  neighborKK = (NeighborKokkos *) neighbor;
+
+  execution_space = Device;
+  
+  datamask_read = EMPTY_MASK;
+  datamask_modify = EMPTY_MASK;
+}
+
+/* ---------------------------------------------------------------------- */
+
+BondHybridKokkos::~BondHybridKokkos()
+{
+  deallocate();
+}
+
+/* ---------------------------------------------------------------------- */
+
+void BondHybridKokkos::compute(int eflag, int vflag)
+{
+  // save ptrs to original bondlist
+
+  int nbondlist_orig = neighbor->nbondlist;
+  neighborKK->k_bondlist.sync_device();
+  auto k_bondlist_orig = neighborKK->k_bondlist;
+  auto d_bondlist_orig = k_bondlist_orig.d_view;
+  auto d_nbondlist = k_nbondlist.d_view;
+  auto h_nbondlist = k_nbondlist.h_view;
+
+  // if this is re-neighbor step, create sub-style bondlists
+  // nbondlist[] = length of each sub-style list
+  // realloc sub-style bondlist if necessary
+  // load sub-style bondlist with 3 values from original bondlist
+
+  if (neighbor->ago == 0) {
+    Kokkos::deep_copy(d_nbondlist,0);
+
+    k_map.sync_device();
+    auto d_map = k_map.d_view;
+
+    Kokkos::parallel_for(nbondlist_orig,LAMMPS_LAMBDA(int i) {
+      const int m = d_map[d_bondlist_orig(i,2)];
+      if (m >= 0) Kokkos::atomic_increment(&d_nbondlist[m]);
+    });
+
+    k_nbondlist.modify_device();
+    k_nbondlist.sync_host();
+
+    for (int m = 0; m < nstyles; m++)
+      if (h_nbondlist[m] > maxbond_all)
+        maxbond_all = h_nbondlist[m] + EXTRA;
+
+    if (k_bondlist.d_view.extent(1) < maxbond_all)
+      MemKK::realloc_kokkos(k_bondlist, "bond_hybrid:bondlist", nstyles, maxbond_all, 3);
+    auto d_bondlist = k_bondlist.d_view;
+
+    Kokkos::deep_copy(d_nbondlist,0);
+
+    Kokkos::parallel_for(nbondlist_orig,LAMMPS_LAMBDA(int i) {
+      const int m = d_map[d_bondlist_orig(i,2)];
+      if (m < 0) return;
+      const int n = Kokkos::atomic_fetch_add(&d_nbondlist[m],1);
+      d_bondlist(m,n,0) = d_bondlist_orig(i,0);
+      d_bondlist(m,n,1) = d_bondlist_orig(i,1);
+      d_bondlist(m,n,2) = d_bondlist_orig(i,2);
+    });
+  }
+
+  // call each sub-style's compute function
+  // set neighbor->bondlist to sub-style bondlist before call
+  // accumulate sub-style global/peratom energy/virial in hybrid
+
+  ev_init(eflag, vflag);
+
+  k_nbondlist.modify_device();
+  k_nbondlist.sync_host();
+
+  for (int m = 0; m < nstyles; m++) {
+    neighbor->nbondlist = h_nbondlist[m];
+    auto k_bondlist_m = Kokkos::subview(k_bondlist,m,Kokkos::ALL,Kokkos::ALL);
+    k_bondlist_m.modify_device();
+    neighborKK->k_bondlist = k_bondlist_m;
+
+    auto style = styles[m];
+    atomKK->sync(style->execution_space,style->datamask_read);
+    style->compute(eflag, vflag);
+    atomKK->modified(style->execution_space,style->datamask_modify);
+
+    if (eflag_global) energy += style->energy;
+    if (vflag_global)
+      for (int n = 0; n < 6; n++) virial[n] += style->virial[n];
+
+    if (eflag_atom) {
+      int n = atom->nlocal;
+      if (force->newton_bond) n += atom->nghost;
+      double *eatom_substyle = styles[m]->eatom;
+      for (int i = 0; i < n; i++) eatom[i] += eatom_substyle[i];
+    }
+    if (vflag_atom) {
+      int n = atom->nlocal;
+      if (force->newton_bond) n += atom->nghost;
+      double **vatom_substyle = styles[m]->vatom;
+      for (int i = 0; i < n; i++)
+        for (int j = 0; j < 6; j++) vatom[i][j] += vatom_substyle[i][j];
+    }
+  }
+
+  // restore ptrs to original bondlist
+
+  neighbor->nbondlist = nbondlist_orig;
+  neighborKK->k_bondlist = k_bondlist_orig;
+}
+
+/* ---------------------------------------------------------------------- */
+
+void BondHybridKokkos::allocate()
+{
+  allocated = 1;
+  int n = atom->nbondtypes;
+
+  memoryKK->create_kokkos(k_map, map, n + 1, "bond:map");
+  memory->create(setflag, n + 1, "bond:setflag");
+  for (int i = 1; i <= n; i++) setflag[i] = 0;
+
+  k_nbondlist = DAT::tdual_int_1d("bond:nbondlist", nstyles);
+}
+
+/* ---------------------------------------------------------------------- */
+
+void BondHybridKokkos::deallocate()
+{
+  if (!allocated) return;
+
+  allocated = 0;
+
+  memory->destroy(setflag);
+  memoryKK->destroy_kokkos(k_map,map);
+}
+
+/* ----------------------------------------------------------------------
+   set coeffs for one type
+---------------------------------------------------------------------- */
+
+void BondHybridKokkos::coeff(int narg, char **arg)
+{
+  BondHybrid::coeff(narg,arg);
+
+  k_map.modify_host();
+}
+
+/* ---------------------------------------------------------------------- */
+
+void BondHybridKokkos::init_style()
+{
+  BondHybrid::init_style();
+
+  for (int m = 0; m < nstyles; m++) {
+    if (!styles[m]->kokkosable)
+      error->all(FLERR,"Must use only Kokkos-enabled bond styles with bond_style hybrid/kk");
+
+    if (styles[m]->execution_space == Host)
+      lmp->kokkos->allow_overlap = 0;
+  }
+}
+
+/* ----------------------------------------------------------------------
+   memory usage
+------------------------------------------------------------------------- */
+
+double BondHybridKokkos::memory_usage()
+{
+  double bytes = (double) maxeatom * sizeof(double);
+  bytes += (double) maxvatom * 6 * sizeof(double);
+  for (int m = 0; m < nstyles; m++) bytes += (double) maxbond_all * 3 * sizeof(int);
+  for (int m = 0; m < nstyles; m++)
+    if (styles[m]) bytes += styles[m]->memory_usage();
+  return bytes;
+}
diff --git a/src/KOKKOS/bond_hybrid_kokkos.h b/src/KOKKOS/bond_hybrid_kokkos.h
new file mode 100644
index 00000000000..217beaca5f6
--- /dev/null
+++ b/src/KOKKOS/bond_hybrid_kokkos.h
@@ -0,0 +1,58 @@
+/* -*- c++ -*- ----------------------------------------------------------
+   LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
+   https://www.lammps.org/, Sandia National Laboratories
+   LAMMPS development team: developers@lammps.org
+
+   Copyright (2003) Sandia Corporation.  Under the terms of Contract
+   DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
+   certain rights in this software.  This software is distributed under
+   the GNU General Public License.
+
+   See the README file in the top-level LAMMPS directory.
+------------------------------------------------------------------------- */
+
+#ifdef BOND_CLASS
+// clang-format off
+BondStyle(hybrid/kk,BondHybridKokkos);
+BondStyle(hybrid/kk/device,BondHybridKokkos);
+BondStyle(hybrid/kk/host,BondHybridKokkos);
+// clang-format on
+#else
+
+// clang-format off
+#ifndef LMP_BOND_HYBRID_KOKKOS_H
+#define LMP_BOND_HYBRID_KOKKOS_H
+
+#include "bond_hybrid.h"
+#include "kokkos_type.h"
+
+namespace LAMMPS_NS {
+
+class BondHybridKokkos : public BondHybrid {
+  friend class Force;
+
+ public:
+  BondHybridKokkos(class LAMMPS *);
+  ~BondHybridKokkos() override;
+  void compute(int, int) override;
+  void coeff(int, char **) override;
+  void init_style() override;
+  double memory_usage() override;
+
+ private:
+  int maxbond_all;
+
+  class NeighborKokkos *neighborKK;
+
+  DAT::tdual_int_1d k_map;       // which style each bond type points to
+  DAT::tdual_int_1d k_nbondlist; // # of bonds in sub-style bondlists
+  DAT::tdual_int_3d k_bondlist;  // bondlist for each sub-style
+
+  void allocate() override;
+  void deallocate() override;
+};
+
+}    // namespace LAMMPS_NS
+
+#endif
+#endif
diff --git a/src/KOKKOS/kokkos_type.h b/src/KOKKOS/kokkos_type.h
index 6e3cb2a1d99..7f0eb5c1056 100644
--- a/src/KOKKOS/kokkos_type.h
+++ b/src/KOKKOS/kokkos_type.h
@@ -1408,6 +1408,7 @@ typedef SNAComplex<SNAreal> SNAcomplex;
 #endif
 
 #define LAMMPS_LAMBDA KOKKOS_LAMBDA
+#define LAMMPS_CLASS_LAMBDA KOKKOS_CLASS_LAMBDA
 
 #if defined(KOKKOS_ENABLE_CUDA) || defined(KOKKOS_ENABLE_HIP)
 #define LAMMPS_DEVICE_FUNCTION __device__
diff --git a/src/KOKKOS/pair_hybrid_kokkos.cpp b/src/KOKKOS/pair_hybrid_kokkos.cpp
index 84d43bcec84..353998a9501 100644
--- a/src/KOKKOS/pair_hybrid_kokkos.cpp
+++ b/src/KOKKOS/pair_hybrid_kokkos.cpp
@@ -31,9 +31,7 @@ PairHybridKokkos::PairHybridKokkos(LAMMPS *lmp) : PairHybrid(lmp)
   kokkosable = 1;
   atomKK = (AtomKokkos *) atom;
 
- // prevent overlapping host/device computation, which isn't
- //  yet supported by pair_hybrid_kokkos
- execution_space = Device;
+  execution_space = Device;
 
   datamask_read = EMPTY_MASK;
   datamask_modify = EMPTY_MASK;
diff --git a/src/bond.cpp b/src/bond.cpp
index 943784427af..b414ddbdca1 100644
--- a/src/bond.cpp
+++ b/src/bond.cpp
@@ -63,8 +63,7 @@ Bond::Bond(LAMMPS *_lmp) : Pointers(_lmp)
   execution_space = Host;
   datamask_read = ALL_MASK;
   datamask_modify = ALL_MASK;
-
-  copymode = 0;
+  copymode = kokkosable = 0;
 }
 
 /* ---------------------------------------------------------------------- */
diff --git a/src/bond.h b/src/bond.h
index be6d7375976..f52204662c4 100644
--- a/src/bond.h
+++ b/src/bond.h
@@ -49,7 +49,7 @@ class Bond : protected Pointers {
 
   ExecutionSpace execution_space;
   unsigned int datamask_read, datamask_modify;
-  int copymode;
+  int copymode, kokkosable;
 
   Bond(class LAMMPS *);
   ~Bond() override;
diff --git a/src/bond_hybrid.cpp b/src/bond_hybrid.cpp
index 6e5ae8d5e7f..2fb5882e9a0 100644
--- a/src/bond_hybrid.cpp
+++ b/src/bond_hybrid.cpp
@@ -51,14 +51,7 @@ BondHybrid::~BondHybrid()
 
   delete[] svector;
 
-  if (allocated) {
-    memory->destroy(setflag);
-    memory->destroy(map);
-    delete[] nbondlist;
-    delete[] maxbond;
-    for (int i = 0; i < nstyles; i++) memory->destroy(bondlist[i]);
-    delete[] bondlist;
-  }
+  deallocate();
 }
 
 /* ---------------------------------------------------------------------- */
@@ -166,6 +159,22 @@ void BondHybrid::allocate()
   for (int m = 0; m < nstyles; m++) bondlist[m] = nullptr;
 }
 
+/* ---------------------------------------------------------------------- */
+
+void BondHybrid::deallocate()
+{ 
+  if (!allocated) return;
+  
+  allocated = 0;
+  
+  memory->destroy(setflag);
+  memory->destroy(map);
+  delete[] nbondlist;
+  delete[] maxbond;
+  for (int i = 0; i < nstyles; i++) memory->destroy(bondlist[i]);
+  delete[] bondlist;
+}
+
 /* ----------------------------------------------------------------------
    create one bond style for each arg in list
 ------------------------------------------------------------------------- */
@@ -186,15 +195,7 @@ void BondHybrid::settings(int narg, char **arg)
     has_quartic = -1;
   }
 
-  if (allocated) {
-    memory->destroy(setflag);
-    memory->destroy(map);
-    delete[] nbondlist;
-    delete[] maxbond;
-    for (i = 0; i < nstyles; i++) memory->destroy(bondlist[i]);
-    delete[] bondlist;
-  }
-  allocated = 0;
+  deallocate();
 
   // allocate list of sub-styles
 
@@ -400,7 +401,7 @@ void BondHybrid::read_restart(FILE *fp)
     keywords[m] = new char[n];
     if (me == 0) utils::sfread(FLERR, keywords[m], sizeof(char), n, fp, nullptr, error);
     MPI_Bcast(keywords[m], n, MPI_CHAR, 0, world);
-    styles[m] = force->new_bond(keywords[m], 0, dummy);
+    styles[m] = force->new_bond(keywords[m], 1, dummy);
     styles[m]->read_restart_settings(fp);
   }
 }
diff --git a/src/bond_hybrid.h b/src/bond_hybrid.h
index df1437c038c..d2496b8297e 100644
--- a/src/bond_hybrid.h
+++ b/src/bond_hybrid.h
@@ -44,14 +44,15 @@ class BondHybrid : public Bond {
   double single(int, double, int, int, double &) override;
   double memory_usage() override;
 
- private:
+ protected:
   int *map;           // which style each bond type points to
   int has_quartic;    // which style, if any is a quartic bond style
   int *nbondlist;     // # of bonds in sub-style bondlists
   int *maxbond;       // max # of bonds sub-style lists can store
   int ***bondlist;    // bondlist for each sub-style
 
-  void allocate();
+  virtual void allocate();
+  virtual void deallocate();
   void flags();
 
   virtual void init_svector();

From ea3bd6043f5820c4ab89897156489b7a34e6c03a Mon Sep 17 00:00:00 2001
From: Stan Moore <stamoor@sandia.gov>
Date: Mon, 20 May 2024 14:36:01 -0600
Subject: [PATCH 2/5] whitespace

---
 src/KOKKOS/bond_hybrid_kokkos.cpp | 2 +-
 src/bond_hybrid.cpp               | 6 +++---
 2 files changed, 4 insertions(+), 4 deletions(-)

diff --git a/src/KOKKOS/bond_hybrid_kokkos.cpp b/src/KOKKOS/bond_hybrid_kokkos.cpp
index 65208f8d8cf..6bf4786a347 100644
--- a/src/KOKKOS/bond_hybrid_kokkos.cpp
+++ b/src/KOKKOS/bond_hybrid_kokkos.cpp
@@ -38,7 +38,7 @@ BondHybridKokkos::BondHybridKokkos(LAMMPS *lmp) : BondHybrid(lmp)
   neighborKK = (NeighborKokkos *) neighbor;
 
   execution_space = Device;
-  
+
   datamask_read = EMPTY_MASK;
   datamask_modify = EMPTY_MASK;
 }
diff --git a/src/bond_hybrid.cpp b/src/bond_hybrid.cpp
index 2fb5882e9a0..7df4f314556 100644
--- a/src/bond_hybrid.cpp
+++ b/src/bond_hybrid.cpp
@@ -162,11 +162,11 @@ void BondHybrid::allocate()
 /* ---------------------------------------------------------------------- */
 
 void BondHybrid::deallocate()
-{ 
+{
   if (!allocated) return;
-  
+
   allocated = 0;
-  
+
   memory->destroy(setflag);
   memory->destroy(map);
   delete[] nbondlist;

From d512326f0606c181595fb7767b8e73eb9aca6f07 Mon Sep 17 00:00:00 2001
From: Stan Moore <stamoor@sandia.gov>
Date: Tue, 4 Jun 2024 15:06:34 -0600
Subject: [PATCH 3/5] Fix uninitialized variable

---
 src/KOKKOS/bond_hybrid_kokkos.cpp | 1 +
 1 file changed, 1 insertion(+)

diff --git a/src/KOKKOS/bond_hybrid_kokkos.cpp b/src/KOKKOS/bond_hybrid_kokkos.cpp
index 6bf4786a347..d63ebccac6d 100644
--- a/src/KOKKOS/bond_hybrid_kokkos.cpp
+++ b/src/KOKKOS/bond_hybrid_kokkos.cpp
@@ -82,6 +82,7 @@ void BondHybridKokkos::compute(int eflag, int vflag)
     k_nbondlist.modify_device();
     k_nbondlist.sync_host();
 
+    maxbond_all = 0;
     for (int m = 0; m < nstyles; m++)
       if (h_nbondlist[m] > maxbond_all)
         maxbond_all = h_nbondlist[m] + EXTRA;

From ecdc1bc336e42ec382dee06b04a9f154d3778ce0 Mon Sep 17 00:00:00 2001
From: Stan Gerald Moore <stamoor@sandia.gov>
Date: Wed, 5 Jun 2024 08:29:33 -0600
Subject: [PATCH 4/5] Update docs

---
 doc/src/Commands_bond.rst | 2 +-
 doc/src/bond_hybrid.rst   | 8 +++++++-
 2 files changed, 8 insertions(+), 2 deletions(-)

diff --git a/doc/src/Commands_bond.rst b/doc/src/Commands_bond.rst
index 0c389df3994..fcd725d787a 100644
--- a/doc/src/Commands_bond.rst
+++ b/doc/src/Commands_bond.rst
@@ -72,7 +72,7 @@ OPT.
 
    * :doc:`none <angle_none>`
    * :doc:`zero <angle_zero>`
-   * :doc:`hybrid <angle_hybrid>`
+   * :doc:`hybrid (k) <angle_hybrid>`
    *
    *
    *
diff --git a/doc/src/bond_hybrid.rst b/doc/src/bond_hybrid.rst
index 1ffd4a1c987..49bfefdcc84 100644
--- a/doc/src/bond_hybrid.rst
+++ b/doc/src/bond_hybrid.rst
@@ -3,6 +3,8 @@
 bond_style hybrid command
 =========================
 
+Accelerator Variants: *hybrid/kk*
+
 Syntax
 """"""
 
@@ -15,7 +17,7 @@ Syntax
 Examples
 """"""""
 
-.. code-block: LAMMPS
+.. code-block:: LAMMPS
 
    bond_style hybrid harmonic fene
    bond_coeff 1 harmonic 80.0 1.2
@@ -60,6 +62,10 @@ bond types.
 
 ----------
 
+.. include:: accel_styles.rst
+
+----------
+
 Restrictions
 """"""""""""
 

From c1c3f21b36521ddbbc7dabf4226633e925825c66 Mon Sep 17 00:00:00 2001
From: Stan Gerald Moore <stamoor@sandia.gov>
Date: Wed, 5 Jun 2024 09:05:13 -0600
Subject: [PATCH 5/5] Need to update index as well

---
 doc/src/bond_hybrid.rst | 1 +
 1 file changed, 1 insertion(+)

diff --git a/doc/src/bond_hybrid.rst b/doc/src/bond_hybrid.rst
index 49bfefdcc84..ea90ed2b578 100644
--- a/doc/src/bond_hybrid.rst
+++ b/doc/src/bond_hybrid.rst
@@ -1,4 +1,5 @@
 .. index:: bond_style hybrid
+.. index:: bond_style hybrid/kk
 
 bond_style hybrid command
 =========================