[UPDATE] Add initial granular/kk KOKKOS port with shared evaluator path

doc/src/Commands_pair.rst

@@ -117,7 +117,7 @@ parenthesis: g = GPU, i = INTEL, k = KOKKOS, o = OPENMP, t = OPT.
    * :doc:`gran/hertz/history (o) <pair_gran>`
    * :doc:`gran/hooke (o) <pair_gran>`
    * :doc:`gran/hooke/history (ko) <pair_gran>`
-   * :doc:`granular <pair_granular>`
+   * :doc:`granular (k) <pair_granular>`
    * :doc:`granular/superellipsoid <pair_granular_superellipsoid>`
    * :doc:`gw (o) <pair_gw>`
    * :doc:`gw/zbl (o) <pair_gw>`

doc/src/pair_granular.rst

@@ -1,8 +1,11 @@
 .. index:: pair_style granular
+.. index:: pair_style granular/kk
 
 pair_style granular command
 ===========================
 
+Accelerator Variants: *granular/kk*
+
 Syntax
 """"""
 
@@ -1054,13 +1057,49 @@ These extra quantities can be accessed by the :doc:`compute pair/local
 
 ----------
 
+.. include:: accel_styles.rst
+
+----------
+
 Restrictions
 """"""""""""
 
 This pair style is part of the GRANULAR package.  It is only enabled if
 LAMMPS was built with that package.  See the :doc:`Build package
 <Build_package>` page for more info.
 
+The *granular/kk* style also requires the KOKKOS package.  Its model
+support matrix is:
+
+.. list-table::
+   :header-rows: 1
+
+   * - Category
+     - Supported models
+     - Unsupported models
+   * - Normal
+     - *hooke*, *hertz*, *hertz/material*
+     - *mdr* and other normal models
+   * - Damping
+     - *none*, *velocity*, *mass_velocity*, *viscoelastic*, *tsuji*,
+       *coeff_restitution*
+     - *mdr* and other damping models
+   * - Tangential
+     - *none*, *linear_nohistory*, *linear_history*, *mindlin*,
+       *mindlin/force*
+     - *mindlin_rescale*, *mindlin_rescale/force*, and other tangential models
+   * - Rolling
+     - *none*, *sds*
+     - other rolling models
+   * - Twisting
+     - *none*, *sds*, *marshall*
+     - other twisting models
+   * - Heat
+     - *none*, *radius*, *area*
+     - other heat models
+
+Models requiring nondefault contact-history transfer are unsupported.
+
 This pair style requires that atoms store per-particle radius, torque,
 and angular velocity (omega) as defined by the :doc:`atom_style sphere
 <atom_style>`.

src/GRANULAR/granular_model.cpp

@@ -249,10 +249,8 @@ void GranularModel::init()
   for (int i = 0; i < NSUBMODELS; i++)
     if (!sub_models[i]) construct_sub_model("none", (SubModelType) i);
 
-  // Must have valid normal, damping, and tangential models
+  // Must have a valid normal model
   if (normal_model->name == "none") error->all(FLERR, "Must specify normal granular model");
-  if (damping_model->name == "none") error->all(FLERR, "Must specify damping granular model");
-  if (tangential_model->name == "none") error->all(FLERR, "Must specify tangential granular model");
 
   // Twisting, rolling, and heat are optional
   twisting_defined = rolling_defined = heat_defined = 1;

src/GRANULAR/pair_granular.cpp

@@ -95,6 +95,8 @@ id_history = utils::strdup(std::string("NEIGH_HISTORY_GRANULAR") + std::to_strin
 
 PairGranular::~PairGranular()
 {
+  if (copymode) return;
+
   delete[] svector;
 
   if (!fix_history) modify->delete_fix(id_dummy);
@@ -492,7 +494,8 @@ void PairGranular::init_style()
   // this is so its order in the fix list is preserved
 
   if (use_history && fix_history == nullptr) {
-    fix_history = dynamic_cast<FixNeighHistory *>(modify->replace_fix(id_dummy, fmt::format("{} all NEIGH_HISTORY {}", id_history, size_history),1));
+    fix_history = dynamic_cast<FixNeighHistory *>(modify->replace_fix(
+        id_dummy, fmt::format("{} all {} {}", id_history, history_fix_style(), size_history), 1));
     fix_history->pair = this;
   } else if (use_history) {
     fix_history = dynamic_cast<FixNeighHistory *>(modify->get_fix_by_id(id_history));
@@ -559,6 +562,13 @@ void PairGranular::init_style()
   MPI_Allreduce(&onerad_frozen[1],&maxrad_frozen[1],atom->ntypes,MPI_DOUBLE,MPI_MAX,world);
 }
 
+/* ---------------------------------------------------------------------- */
+
+const char *PairGranular::history_fix_style() const
+{
+  return "NEIGH_HISTORY";
+}
+
 /* ----------------------------------------------------------------------
    init for one type pair i,j and corresponding j,i
 ------------------------------------------------------------------------- */

src/GRANULAR/pair_granular.h

@@ -72,8 +72,9 @@ class PairGranular : public Pair {
   double *mass_rigid;      // rigid mass for owned+ghost atoms
   int nmax;                // allocated size of mass_rigid
 
-  void allocate();
+  virtual void allocate();
   void prune_models();
+  virtual const char *history_fix_style() const;
 
   int size_history;
   int heat_flag;

src/KOKKOS/Install.sh

@@ -416,6 +416,8 @@ action pair_gauss_kokkos.cpp pair_gauss.cpp
 action pair_gauss_kokkos.h pair_gauss.h
 action pair_gran_hooke_history_kokkos.cpp pair_gran_hooke_history.cpp
 action pair_gran_hooke_history_kokkos.h pair_gran_hooke_history.h
+action pair_granular_kokkos.cpp pair_granular.cpp
+action pair_granular_kokkos.h pair_granular.h
 action pair_gauss_cut_kokkos.cpp pair_gauss_cut.cpp
 action pair_gauss_cut_kokkos.h pair_gauss_cut.h
 action pair_hybrid_kokkos.cpp