From 882a72987be13f1b6e1829a1df45adce2fbb7979 Mon Sep 17 00:00:00 2001 From: Stan Gerald Moore Date: Fri, 17 Mar 2023 11:49:42 -0600 Subject: [PATCH] Port compute_erotate_sphere to Kokkos --- src/KOKKOS/Install.sh | 2 + src/KOKKOS/compute_erotate_sphere_kokkos.cpp | 91 ++++++++++++++++++++ src/KOKKOS/compute_erotate_sphere_kokkos.h | 50 +++++++++++ src/KOKKOS/compute_temp_kokkos.cpp | 1 - src/KOKKOS/compute_temp_kokkos.h | 11 ++- src/compute_erotate_sphere.h | 2 +- 6 files changed, 149 insertions(+), 8 deletions(-) create mode 100644 src/KOKKOS/compute_erotate_sphere_kokkos.cpp create mode 100644 src/KOKKOS/compute_erotate_sphere_kokkos.h diff --git a/src/KOKKOS/Install.sh b/src/KOKKOS/Install.sh index 9b972cdb46..9b390a23b2 100755 --- a/src/KOKKOS/Install.sh +++ b/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 diff --git a/src/KOKKOS/compute_erotate_sphere_kokkos.cpp b/src/KOKKOS/compute_erotate_sphere_kokkos.cpp new file mode 100644 index 0000000000..9fc477b3a0 --- /dev/null +++ b/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 +ComputeERotateSphereKokkos::ComputeERotateSphereKokkos(LAMMPS *lmp, int narg, char **arg) : + ComputeERotateSphere(lmp, narg, arg) +{ + kokkosable = 1; + atomKK = (AtomKokkos *) atom; + execution_space = ExecutionSpaceFromDevice::space; + + datamask_read = OMEGA_MASK | RADIUS_MASK | MASK_MASK | RMASS_MASK; + datamask_modify = EMPTY_MASK; +} + +/* ---------------------------------------------------------------------- */ + +template +double ComputeERotateSphereKokkos::compute_scalar() +{ + atomKK->sync(execution_space,datamask_read); + + invoked_scalar = update->ntimestep; + + omega = atomKK->k_omega.view(); + radius = atomKK->k_radius.view(); + rmass = atomKK->k_rmass.view(); + mask = atomKK->k_mask.view(); + 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(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; +#ifdef LMP_KOKKOS_GPU +template class ComputeERotateSphereKokkos; +#endif +} diff --git a/src/KOKKOS/compute_erotate_sphere_kokkos.h b/src/KOKKOS/compute_erotate_sphere_kokkos.h new file mode 100644 index 0000000000..2a8feb1fa3 --- /dev/null +++ b/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); +ComputeStyle(erotate/sphere/kk/device,ComputeERotateSphereKokkos); +ComputeStyle(erotate/sphere/kk/host,ComputeERotateSphereKokkos); +// 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 ComputeERotateSphereKokkos : public ComputeERotateSphere { + public: + typedef DeviceType device_type; + typedef ArrayTypes 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 diff --git a/src/KOKKOS/compute_temp_kokkos.cpp b/src/KOKKOS/compute_temp_kokkos.cpp index 159be08554..ebdd6971e0 100644 --- a/src/KOKKOS/compute_temp_kokkos.cpp +++ b/src/KOKKOS/compute_temp_kokkos.cpp @@ -159,4 +159,3 @@ template class ComputeTempKokkos; template class ComputeTempKokkos; #endif } - diff --git a/src/KOKKOS/compute_temp_kokkos.h b/src/KOKKOS/compute_temp_kokkos.h index 0bc56f13ba..80144acfc8 100644 --- a/src/KOKKOS/compute_temp_kokkos.h +++ b/src/KOKKOS/compute_temp_kokkos.h @@ -85,15 +85,14 @@ class ComputeTempKokkos : public ComputeTemp { void operator()(TagComputeTempVector, const int&, CTEMP&) const; protected: - typename ArrayTypes::t_v_array_randomread v; - typename ArrayTypes::t_float_1d_randomread rmass; - typename ArrayTypes::t_float_1d_randomread mass; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_int_1d_randomread mask; + typename AT::t_v_array_randomread v; + typename AT::t_float_1d_randomread rmass; + typename AT::t_float_1d_randomread mass; + typename AT::t_int_1d_randomread type; + typename AT::t_int_1d_randomread mask; }; } #endif #endif - diff --git a/src/compute_erotate_sphere.h b/src/compute_erotate_sphere.h index 06262b89ef..149ec9870d 100644 --- a/src/compute_erotate_sphere.h +++ b/src/compute_erotate_sphere.h @@ -31,7 +31,7 @@ class ComputeERotateSphere : public Compute { void init() override; double compute_scalar() override; - private: + protected: double pfactor; };