Use multiple inheritance to remove accelerator_kokkos.h out of region.h and kspace.h

src/KOKKOS/fix_setforce_kokkos.cpp

@@ -26,6 +26,7 @@
 #include "error.h"
 #include "force.h"
 #include "atom_masks.h"
+#include "kokkos_base.h"
 
 using namespace LAMMPS_NS;
 using namespace FixConst;
@@ -90,7 +91,8 @@ void FixSetForceKokkos<DeviceType>::post_force(int vflag)
     region = domain->regions[iregion];
     region->prematch();
     DAT::tdual_int_1d k_match = DAT::tdual_int_1d("setforce:k_match",nlocal);
-    region->match_all_kokkos(groupbit,k_match);
+    KokkosBase* regionKKBase = (KokkosBase*) region;
+    regionKKBase->match_all_kokkos(groupbit,k_match);
     k_match.template sync<DeviceType>();
     d_match = k_match.template view<DeviceType>();
   }

src/KOKKOS/gridcomm_kokkos.cpp

@@ -17,6 +17,7 @@
 #include "kspace.h"
 #include "memory_kokkos.h"
 #include "error.h"
+#include "kokkos_base.h"
 
 using namespace LAMMPS_NS;
 
@@ -515,11 +516,13 @@ void GridCommKokkos<DeviceType>::forward_comm(KSpace *kspace, int which)
   k_packlist.sync<DeviceType>();
   k_unpacklist.sync<DeviceType>();
 
+  KokkosBase* kspaceKKBase = (KokkosBase*) kspace;
+
   for (int m = 0; m < nswap; m++) {
     if (swap[m].sendproc == me)
-      kspace->pack_forward_kokkos(which,k_buf2,swap[m].npack,k_packlist,m);
+      kspaceKKBase->pack_forward_kspace_kokkos(which,k_buf2,swap[m].npack,k_packlist,m);
     else
-      kspace->pack_forward_kokkos(which,k_buf1,swap[m].npack,k_packlist,m);
+      kspaceKKBase->pack_forward_kspace_kokkos(which,k_buf1,swap[m].npack,k_packlist,m);
 
     if (swap[m].sendproc != me) {
       MPI_Irecv(k_buf2.view<DeviceType>().ptr_on_device(),nforward*swap[m].nunpack,MPI_FFT_SCALAR,
@@ -529,7 +532,7 @@ void GridCommKokkos<DeviceType>::forward_comm(KSpace *kspace, int which)
       MPI_Wait(&request,MPI_STATUS_IGNORE);
     }
 
-    kspace->unpack_forward_kokkos(which,k_buf2,swap[m].nunpack,k_unpacklist,m);
+    kspaceKKBase->unpack_forward_kspace_kokkos(which,k_buf2,swap[m].nunpack,k_unpacklist,m);
   }
 }
 
@@ -544,11 +547,13 @@ void GridCommKokkos<DeviceType>::reverse_comm(KSpace *kspace, int which)
   k_packlist.sync<DeviceType>();
   k_unpacklist.sync<DeviceType>();
 
+  KokkosBase* kspaceKKBase = (KokkosBase*) kspace;
+
   for (int m = nswap-1; m >= 0; m--) {
     if (swap[m].recvproc == me)
-      kspace->pack_reverse_kokkos(which,k_buf2,swap[m].nunpack,k_unpacklist,m);
+      kspaceKKBase->pack_reverse_kspace_kokkos(which,k_buf2,swap[m].nunpack,k_unpacklist,m);
     else
-      kspace->pack_reverse_kokkos(which,k_buf1,swap[m].nunpack,k_unpacklist,m);
+      kspaceKKBase->pack_reverse_kspace_kokkos(which,k_buf1,swap[m].nunpack,k_unpacklist,m);
 
     if (swap[m].recvproc != me) {
       MPI_Irecv(k_buf2.view<DeviceType>().ptr_on_device(),nreverse*swap[m].npack,MPI_FFT_SCALAR,
@@ -558,7 +563,7 @@ void GridCommKokkos<DeviceType>::reverse_comm(KSpace *kspace, int which)
       MPI_Wait(&request,MPI_STATUS_IGNORE);
     }
 
-    kspace->unpack_reverse_kokkos(which,k_buf2,swap[m].npack,k_packlist,m);
+    kspaceKKBase->unpack_reverse_kspace_kokkos(which,k_buf2,swap[m].npack,k_packlist,m);
   }
 }
 

src/KOKKOS/kokkos_base.h

@@ -21,10 +21,21 @@ namespace LAMMPS_NS {
 class KokkosBase {
  public:
   KokkosBase() {}
+
+  //Kspace
+  virtual void pack_forward_kspace_kokkos(int, DAT::tdual_FFT_SCALAR_1d &, int, DAT::tdual_int_2d &, int) {};
+  virtual void unpack_forward_kspace_kokkos(int, DAT::tdual_FFT_SCALAR_1d &, int, DAT::tdual_int_2d &, int) {};
+  virtual void pack_reverse_kspace_kokkos(int, DAT::tdual_FFT_SCALAR_1d &, int, DAT::tdual_int_2d &, int) {};
+  virtual void unpack_reverse_kspace_kokkos(int, DAT::tdual_FFT_SCALAR_1d &, int, DAT::tdual_int_2d &, int) {};
+
+  // Pair
   virtual int pack_forward_comm_kokkos(int, DAT::tdual_int_2d, 
                                        int, DAT::tdual_xfloat_1d &, 
                                        int, int *) {return 0;};
   virtual void unpack_forward_comm_kokkos(int, int, DAT::tdual_xfloat_1d &) {}
+
+  // Region
+  virtual void match_all_kokkos(int, DAT::tdual_int_1d) {}
 };
 
 }

src/KOKKOS/pppm_kokkos.cpp

@@ -35,6 +35,7 @@
 #include "memory_kokkos.h"
 #include "error.h"
 #include "atom_masks.h"
+#include "kokkos.h"
 
 #include "math_const.h"
 #include "math_special_kokkos.h"
@@ -2631,7 +2632,7 @@ void PPPMKokkos<DeviceType>::operator()(TagPPPM_fieldforce_peratom, const int &i
 ------------------------------------------------------------------------- */
 
 template<class DeviceType>
-void PPPMKokkos<DeviceType>::pack_forward_kokkos(int flag, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &k_buf, int nlist, DAT::tdual_int_2d &k_list, int index)
+void PPPMKokkos<DeviceType>::pack_forward_kspace_kokkos(int flag, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &k_buf, int nlist, DAT::tdual_int_2d &k_list, int index)
 {
   typename AT::t_int_2d_um d_list = k_list.view<DeviceType>();
   d_list_index = Kokkos::subview(d_list,index,Kokkos::ALL());
@@ -2687,7 +2688,7 @@ void PPPMKokkos<DeviceType>::operator()(TagPPPM_pack_forward2, const int &i) con
 ------------------------------------------------------------------------- */
 
 template<class DeviceType>
-void PPPMKokkos<DeviceType>::unpack_forward_kokkos(int flag, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &k_buf, int nlist, DAT::tdual_int_2d &k_list, int index)
+void PPPMKokkos<DeviceType>::unpack_forward_kspace_kokkos(int flag, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &k_buf, int nlist, DAT::tdual_int_2d &k_list, int index)
 {
   typename AT::t_int_2d_um d_list = k_list.view<DeviceType>();
   d_list_index = Kokkos::subview(d_list,index,Kokkos::ALL());
@@ -2744,7 +2745,7 @@ void PPPMKokkos<DeviceType>::operator()(TagPPPM_unpack_forward2, const int &i) c
 ------------------------------------------------------------------------- */
 
 template<class DeviceType>
-void PPPMKokkos<DeviceType>::pack_reverse_kokkos(int flag, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &k_buf, int nlist, DAT::tdual_int_2d &k_list, int index)
+void PPPMKokkos<DeviceType>::pack_reverse_kspace_kokkos(int flag, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &k_buf, int nlist, DAT::tdual_int_2d &k_list, int index)
 {
   typename AT::t_int_2d_um d_list = k_list.view<DeviceType>();
   d_list_index = Kokkos::subview(d_list,index,Kokkos::ALL());
@@ -2774,7 +2775,7 @@ void PPPMKokkos<DeviceType>::operator()(TagPPPM_pack_reverse, const int &i) cons
 ------------------------------------------------------------------------- */
 
 template<class DeviceType>
-void PPPMKokkos<DeviceType>::unpack_reverse_kokkos(int flag, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &k_buf, int nlist, DAT::tdual_int_2d &k_list, int index)
+void PPPMKokkos<DeviceType>::unpack_reverse_kspace_kokkos(int flag, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &k_buf, int nlist, DAT::tdual_int_2d &k_list, int index)
 {
   typename AT::t_int_2d_um d_list = k_list.view<DeviceType>();
   d_list_index = Kokkos::subview(d_list,index,Kokkos::ALL());

src/KOKKOS/pppm_kokkos.h

@@ -24,6 +24,7 @@ KSpaceStyle(pppm/kk/host,PPPMKokkos<LMPHostType>)
 
 #include "pppm.h"
 #include "gridcomm_kokkos.h"
+#include "kokkos_base.h"
 #include "kokkos_type.h"
 
 namespace LAMMPS_NS {
@@ -86,7 +87,7 @@ struct TagPPPM_slabcorr4{};
 struct TagPPPM_timing_zero{};
 
 template<class DeviceType>
-class PPPMKokkos : public PPPM {
+class PPPMKokkos : public PPPM, public KokkosBase {
  public:
   typedef DeviceType device_type;
   typedef ArrayTypes<DeviceType> AT;
@@ -379,10 +380,10 @@ class PPPMKokkos : public PPPM {
 
   // grid communication
 
-  virtual void pack_forward_kokkos(int, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &, int, DAT::tdual_int_2d &, int);
+  void pack_forward_kspace_kokkos(int, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &, int, DAT::tdual_int_2d &, int);
-  virtual void unpack_forward_kokkos(int, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &, int, DAT::tdual_int_2d &, int);
+  void unpack_forward_kspace_kokkos(int, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &, int, DAT::tdual_int_2d &, int);
-  virtual void pack_reverse_kokkos(int, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &, int, DAT::tdual_int_2d &, int);
+  void pack_reverse_kspace_kokkos(int, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &, int, DAT::tdual_int_2d &, int);
-  virtual void unpack_reverse_kokkos(int, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &, int, DAT::tdual_int_2d &, int);
+  void unpack_reverse_kspace_kokkos(int, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &, int, DAT::tdual_int_2d &, int);
 
   // triclinic
 

src/KOKKOS/region_block_kokkos.h

@@ -23,6 +23,7 @@ RegionStyle(block/kk/host,RegBlockKokkos<LMPHostType>)
 #define LMP_REGION_BLOCK_KOKKOS_H
 
 #include "region_block.h"
+#include "kokkos_base.h"
 #include "kokkos_type.h"
 
 namespace LAMMPS_NS {
@@ -30,7 +31,7 @@ namespace LAMMPS_NS {
 struct TagRegBlockMatchAll{};
 
 template<class DeviceType>
-class RegBlockKokkos : public RegBlock {
+class RegBlockKokkos : public RegBlock, public KokkosBase {
   friend class FixPour;
 
  public:

src/KOKKOS/verlet_kokkos.cpp

@@ -34,6 +34,7 @@
 #include "timer.h"
 #include "memory_kokkos.h"
 #include "error.h"
+#include "kokkos.h"
 
 #include <ctime>
 

src/delete_atoms.cpp

@@ -28,6 +28,7 @@
 #include "random_mars.h"
 #include "memory.h"
 #include "error.h"
+#include "modify.h"
 
 #include <map>
 

src/kspace.h

@@ -15,7 +15,6 @@
 #define LMP_KSPACE_H
 
 #include "pointers.h"
-#include "accelerator_kokkos.h"
 
 #ifdef FFT_SINGLE
 typedef float FFT_SCALAR;
@@ -124,11 +123,6 @@ class KSpace : protected Pointers {
   virtual void pack_reverse(int, FFT_SCALAR *, int, int *) {};
   virtual void unpack_reverse(int, FFT_SCALAR *, int, int *) {};
 
-  virtual void pack_forward_kokkos(int, DAT::tdual_FFT_SCALAR_1d &, int, DAT::tdual_int_2d &, int) {};
-  virtual void unpack_forward_kokkos(int, DAT::tdual_FFT_SCALAR_1d &, int, DAT::tdual_int_2d &, int) {};
-  virtual void pack_reverse_kokkos(int, DAT::tdual_FFT_SCALAR_1d &, int, DAT::tdual_int_2d &, int) {};
-  virtual void unpack_reverse_kokkos(int, DAT::tdual_FFT_SCALAR_1d &, int, DAT::tdual_int_2d &, int) {};
-
   virtual int timing(int, double &, double &) {return 0;}
   virtual int timing_1d(int, double &) {return 0;}
   virtual int timing_3d(int, double &) {return 0;}

src/region.cpp

@@ -142,15 +142,6 @@ int Region::match(double x, double y, double z)
   return !(inside(x,y,z) ^ interior);
 }
 
-/* ----------------------------------------------------------------------
-   generate error if Kokkos function defaults to base class
-------------------------------------------------------------------------- */
-
-void Region::match_all_kokkos(int, DAT::tdual_int_1d)
-{
-  error->all(FLERR,"Can only use Kokkos supported regions with Kokkos package");
-}
-
 /* ----------------------------------------------------------------------
    generate list of contact points for interior or exterior regions
    if region has variable shape, invoke shape_update() once per timestep

src/region.h

@@ -15,7 +15,6 @@
 #define LMP_REGION_H
 
 #include "pointers.h"
-#include "accelerator_kokkos.h"
 
 namespace LAMMPS_NS {
 
@@ -97,10 +96,6 @@ class Region : protected Pointers {
   virtual void set_velocity_shape() {}
   virtual void velocity_contact_shape(double*, double*) {}
 
-  // Kokkos function, implemented by each Kokkos region
-
-  virtual void match_all_kokkos(int, DAT::tdual_int_1d);
-
  protected:
   void add_contact(int, double *, double, double, double);
   void options(int, char **);

src/set.cpp

@@ -36,6 +36,7 @@
 #include "math_const.h"
 #include "memory.h"
 #include "error.h"
+#include "modify.h"
 
 using namespace LAMMPS_NS;
 using namespace MathConst;