Port compute_erotate_sphere to Kokkos

src/KOKKOS/Install.sh

@@ -94,6 +94,8 @@ action compute_ave_sphere_atom_kokkos.cpp compute_ave_sphere_atom.cpp
 action compute_ave_sphere_atom_kokkos.h compute_ave_sphere_atom.h
 action compute_coord_atom_kokkos.cpp
 action compute_coord_atom_kokkos.h
+action compute_erotate_sphere_kokkos.cpp
+action compute_erotate_sphere_kokkos.h
 action compute_orientorder_atom_kokkos.cpp
 action compute_orientorder_atom_kokkos.h
 action compute_temp_deform_kokkos.cpp

src/KOKKOS/compute_erotate_sphere_kokkos.cpp

@@ -0,0 +1,91 @@
+// clang-format off
+/* ----------------------------------------------------------------------
+   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 "compute_erotate_sphere_kokkos.h"
+
+#include "atom_kokkos.h"
+#include "atom_masks.h"
+#include "error.h"
+#include "force.h"
+#include "update.h"
+
+using namespace LAMMPS_NS;
+
+/* ---------------------------------------------------------------------- */
+
+template<class DeviceType>
+ComputeERotateSphereKokkos<DeviceType>::ComputeERotateSphereKokkos(LAMMPS *lmp, int narg, char **arg) :
+  ComputeERotateSphere(lmp, narg, arg)
+{
+  kokkosable = 1;
+  atomKK = (AtomKokkos *) atom;
+  execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
+
+  datamask_read = OMEGA_MASK | RADIUS_MASK | MASK_MASK | RMASS_MASK;
+  datamask_modify = EMPTY_MASK;
+}
+
+/* ---------------------------------------------------------------------- */
+
+template<class DeviceType>
+double ComputeERotateSphereKokkos<DeviceType>::compute_scalar()
+{
+  atomKK->sync(execution_space,datamask_read);
+
+  invoked_scalar = update->ntimestep;
+
+  omega = atomKK->k_omega.view<DeviceType>();
+  radius = atomKK->k_radius.view<DeviceType>();
+  rmass = atomKK->k_rmass.view<DeviceType>();
+  mask = atomKK->k_mask.view<DeviceType>();
+  int nlocal = atom->nlocal;
+
+  // sum rotational energy for each particle
+  // point particles will not contribute, due to radius = 0.0
+
+  double erotate = 0.0;
+
+  {
+    // local variables for lambda capture
+
+    auto l_omega = omega;
+    auto l_radius = radius;
+    auto l_rmass = rmass;
+    auto l_mask = mask;
+    auto l_groupbit = groupbit;
+
+    Kokkos::parallel_reduce(Kokkos::RangePolicy<DeviceType>(0,nlocal), LAMMPS_LAMBDA(int i, double &erotate) {
+      if (l_mask[i] & l_groupbit) {
+        auto omega0 = l_omega(i,0);
+        auto omega1 = l_omega(i,1);
+        auto omega2 = l_omega(i,2);
+        auto radius = l_radius(i);
+        erotate +=
+            (omega0 * omega0 + omega1 * omega1 + omega2 * omega2) *
+            radius * radius * l_rmass[i];
+      }
+    },erotate);
+  }
+
+  MPI_Allreduce(&erotate, &scalar, 1, MPI_DOUBLE, MPI_SUM, world);
+  scalar *= pfactor;
+  return scalar;
+}
+
+namespace LAMMPS_NS {
+template class ComputeERotateSphereKokkos<LMPDeviceType>;
+#ifdef LMP_KOKKOS_GPU
+template class ComputeERotateSphereKokkos<LMPHostType>;
+#endif
+}

src/KOKKOS/compute_erotate_sphere_kokkos.h

@@ -0,0 +1,50 @@
+/* -*- 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 COMPUTE_CLASS
+// clang-format off
+ComputeStyle(erotate/sphere/kk,ComputeERotateSphereKokkos<LMPDeviceType>);
+ComputeStyle(erotate/sphere/kk/device,ComputeERotateSphereKokkos<LMPDeviceType>);
+ComputeStyle(erotate/sphere/kk/host,ComputeERotateSphereKokkos<LMPHostType>);
+// clang-format on
+#else
+
+// clang-format off
+#ifndef LMP_COMPUTE_EROTATE_SPHERE_KOKKOS_H
+#define LMP_COMPUTE_EROTATE_SPHERE_KOKKOS_H
+
+#include "compute_erotate_sphere.h"
+#include "kokkos_type.h"
+
+namespace LAMMPS_NS {
+
+template<class DeviceType>
+class ComputeERotateSphereKokkos : public ComputeERotateSphere {
+ public:
+  typedef DeviceType device_type;
+  typedef ArrayTypes<DeviceType> AT;
+
+  ComputeERotateSphereKokkos(class LAMMPS *, int, char **);
+  double compute_scalar() override;
+
+ private:
+  typename AT::t_v_array_randomread omega;
+  typename AT::t_float_1d_randomread radius;
+  typename AT::t_float_1d_randomread rmass;
+  typename AT::t_int_1d_randomread mask;
+};
+
+}    // namespace LAMMPS_NS
+
+#endif
+#endif

src/KOKKOS/compute_temp_kokkos.cpp

@@ -159,4 +159,3 @@ template class ComputeTempKokkos<LMPDeviceType>;
 template class ComputeTempKokkos<LMPHostType>;
 #endif
 }
-

src/KOKKOS/compute_temp_kokkos.h

@@ -85,15 +85,14 @@ class ComputeTempKokkos : public ComputeTemp {
   void operator()(TagComputeTempVector<RMASS>, const int&, CTEMP&) const;
 
  protected:
-  typename ArrayTypes<DeviceType>::t_v_array_randomread v;
+  typename AT::t_v_array_randomread v;
-  typename ArrayTypes<DeviceType>::t_float_1d_randomread rmass;
+  typename AT::t_float_1d_randomread rmass;
-  typename ArrayTypes<DeviceType>::t_float_1d_randomread mass;
+  typename AT::t_float_1d_randomread mass;
-  typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
+  typename AT::t_int_1d_randomread type;
-  typename ArrayTypes<DeviceType>::t_int_1d_randomread mask;
+  typename AT::t_int_1d_randomread mask;
 };
 
 }
 
 #endif
 #endif
-

src/compute_erotate_sphere.h

@@ -31,7 +31,7 @@ class ComputeERotateSphere : public Compute {
   void init() override;
   double compute_scalar() override;
 
- private:
+ protected:
   double pfactor;
 };