diff --git a/doc/src/Commands_fix.rst b/doc/src/Commands_fix.rst index ff6e402a2e..0d3f69602e 100644 --- a/doc/src/Commands_fix.rst +++ b/doc/src/Commands_fix.rst @@ -65,7 +65,7 @@ OPT. * :doc:`drude ` * :doc:`drude/transform/direct ` * :doc:`drude/transform/inverse ` - * :doc:`dt/reset ` + * :doc:`dt/reset (k) ` * :doc:`edpd/source ` * :doc:`efield ` * :doc:`ehex ` @@ -250,7 +250,7 @@ OPT. * :doc:`tune/kspace ` * :doc:`vector ` * :doc:`viscosity ` - * :doc:`viscous ` + * :doc:`viscous (k) ` * :doc:`viscous/sphere ` * :doc:`wall/body/polygon ` * :doc:`wall/body/polyhedron ` diff --git a/doc/src/fix_dt_reset.rst b/doc/src/fix_dt_reset.rst index 368a3dcd70..e99bf37b75 100644 --- a/doc/src/fix_dt_reset.rst +++ b/doc/src/fix_dt_reset.rst @@ -3,6 +3,8 @@ fix dt/reset command ==================== +Accelerator Variants: *dt/reset/kk* + Syntax """""" @@ -86,6 +88,10 @@ allows dump files to be written at intervals specified by simulation time, rather than by timesteps. Simulation time is in time units; see the :doc:`units ` doc page for details. +---------- + +.. include:: accel_styles.rst + Restart, fix_modify, output, run start/stop, minimize info """"""""""""""""""""""""""""""""""""""""""""""""""""""""""" diff --git a/doc/src/fix_viscous.rst b/doc/src/fix_viscous.rst index 6c48cb5f12..723bc15085 100644 --- a/doc/src/fix_viscous.rst +++ b/doc/src/fix_viscous.rst @@ -3,6 +3,8 @@ fix viscous command =================== +Accelerator Variants: *viscous/kk* + Syntax """""" @@ -84,6 +86,10 @@ more easily be used as a thermostat. ---------- +.. include:: accel_styles.rst + +---------- + Restart, fix_modify, output, run start/stop, minimize info """"""""""""""""""""""""""""""""""""""""""""""""""""""""""" diff --git a/src/KOKKOS/Install.sh b/src/KOKKOS/Install.sh index b30461d5c6..d53e9273d5 100755 --- a/src/KOKKOS/Install.sh +++ b/src/KOKKOS/Install.sh @@ -117,6 +117,8 @@ action fix_acks2_reaxff_kokkos.cpp fix_acks2_reaxff.cpp action fix_acks2_reaxff_kokkos.h fix_acks2_reaxff.h action fix_deform_kokkos.cpp action fix_deform_kokkos.h +action fix_dt_reset_kokkos.cpp +action fix_dt_reset_kokkos.h action fix_enforce2d_kokkos.cpp action fix_enforce2d_kokkos.h action fix_eos_table_rx_kokkos.cpp fix_eos_table_rx.cpp @@ -165,6 +167,8 @@ action fix_wall_lj93_kokkos.cpp action fix_wall_lj93_kokkos.h action fix_wall_reflect_kokkos.cpp action fix_wall_reflect_kokkos.h +action fix_viscous_kokkos.cpp +action fix_viscous_kokkos.h action fix_dpd_energy_kokkos.cpp fix_dpd_energy.cpp action fix_dpd_energy_kokkos.h fix_dpd_energy.h action fix_rx_kokkos.cpp fix_rx.cpp diff --git a/src/KOKKOS/fix_dt_reset_kokkos.cpp b/src/KOKKOS/fix_dt_reset_kokkos.cpp new file mode 100644 index 0000000000..f3435e711e --- /dev/null +++ b/src/KOKKOS/fix_dt_reset_kokkos.cpp @@ -0,0 +1,195 @@ +// 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 "fix_dt_reset_kokkos.h" + +#include "atom_kokkos.h" +#include "atom_masks.h" +#include "error.h" +#include "force.h" +#include "input.h" +#include "integrate.h" +#include "kokkos_base.h" +#include "memory_kokkos.h" +#include "modify.h" +#include "output.h" +#include "pair.h" +#include "update.h" + +using namespace LAMMPS_NS; +using namespace FixConst; + +#define BIG 1.0e20 + +/* ---------------------------------------------------------------------- */ + +template +FixDtResetKokkos::FixDtResetKokkos(LAMMPS *lmp, int narg, char **arg) : + FixDtReset(lmp, narg, arg) +{ + kokkosable = 1; + atomKK = (AtomKokkos *) atom; + execution_space = ExecutionSpaceFromDevice::space; + datamask_read = EMPTY_MASK; + datamask_modify = EMPTY_MASK; +} + +/* ---------------------------------------------------------------------- */ + +template +void FixDtResetKokkos::init() +{ + FixDtReset::init(); + + k_emax = Kokkos::DualView("FixDtResetKokkos:gamma", 1); + + k_emax.h_view(0) = emax; + + + k_emax.template modify(); + k_emax.template sync(); + + if (utils::strmatch(update->integrate_style,"^respa")) + error->all(FLERR,"Cannot (yet) use respa with Kokkos"); +} + +/* ---------------------------------------------------------------------- */ + +template +void FixDtResetKokkos::end_of_step() +{ + atomKK->sync(execution_space, V_MASK | F_MASK | MASK_MASK | TYPE_MASK | RMASS_MASK); + + v = atomKK->k_v.view(); + f = atomKK->k_f.view(); + mask = atomKK->k_mask.view(); + type = atomKK->k_type.view(); + if (atomKK->rmass) + rmass = atomKK->k_rmass.view(); + else + mass = atomKK->k_mass.view(); + + int nlocal = atom->nlocal; + + double dt; + + copymode = 1; + if (atomKK->rmass) + Kokkos::parallel_reduce(Kokkos::RangePolicy(0,nlocal), *this, Kokkos::Min(dt)); + else + Kokkos::parallel_reduce(Kokkos::RangePolicy(0,nlocal), *this, Kokkos::Min(dt)); + copymode = 0; + + MPI_Allreduce(MPI_IN_PLACE, &dt, 1, MPI_DOUBLE, MPI_MIN, world); + + atomKK->modified(execution_space, F_MASK); + + if (minbound) dt = MAX(dt, tmin); + if (maxbound) dt = MIN(dt, tmax); + + // if timestep didn't change, just return + // else reset update->dt and other classes that depend on it + // rRESPA, pair style, fixes + + if (dt == update->dt) return; + + laststep = update->ntimestep; + + // calls to other classes that need to know timestep size changed + // similar logic is in Input::timestep() + + update->update_time(); + update->dt = dt; + update->dt_default = 0; + if (force->pair) force->pair->reset_dt(); + for (int i = 0; i < modify->nfix; i++) modify->fix[i]->reset_dt(); + output->reset_dt(); + +} + +/* ---------------------------------------------------------------------- */ + +template +KOKKOS_INLINE_FUNCTION +void FixDtResetKokkos::operator()(TagFixDtResetMass, const int &i, double &k_dt) const { + + double dtv, dtf, dte, dtsq; + double vsq, fsq, massinv; + double delx, dely, delz, delr; + + double emax = k_emax.d_view(0); + + if (mask[i] & groupbit) { + + massinv = 1.0 / mass[type[i]]; + vsq = v(i,0) * v(i,0) + v(i,1) * v(i,1) + v(i,2) * v(i,2); + fsq = f(i,0) * f(i,0) + f(i,1) * f(i,1) + f(i,2) * f(i,2); + dtv = dtf = dte = BIG; + if (vsq > 0.0) dtv = xmax / sqrt(vsq); + if (fsq > 0.0) dtf = sqrt(2.0 * xmax / (ftm2v * sqrt(fsq) * massinv)); + k_dt = MIN(dtv, dtf); + if ((emax > 0.0) && (fsq * vsq > 0.0)) { + dte = emax / sqrt(fsq * vsq) / sqrt(ftm2v * mvv2e); + k_dt = MIN(dt, dte); + } + dtsq = k_dt * k_dt; + delx = k_dt * v(i,0) + 0.5 * dtsq * massinv * f(i,0) * ftm2v; + dely = k_dt * v(i,1) + 0.5 * dtsq * massinv * f(i,1) * ftm2v; + delz = k_dt * v(i,2) + 0.5 * dtsq * massinv * f(i,2) * ftm2v; + delr = sqrt(delx * delx + dely * dely + delz * delz); + if (delr > xmax) k_dt *= xmax / delr; + } + } + +/* ---------------------------------------------------------------------- */ + +template +KOKKOS_INLINE_FUNCTION +void FixDtResetKokkos::operator()(TagFixDtResetRMass, const int &i, double &k_dt) const { + + double dtv, dtf, dte, dtsq; + double vsq, fsq, massinv; + double delx, dely, delz, delr; + + double emax = k_emax.d_view(0); + + if (mask[i] & groupbit) { + + massinv = 1.0 / rmass[i]; + vsq = v(i,0) * v(i,0) + v(i,1) * v(i,1) + v(i,2) * v(i,2); + fsq = f(i,0) * f(i,0) + f(i,1) * f(i,1) + f(i,2) * f(i,2); + dtv = dtf = dte = BIG; + if (vsq > 0.0) dtv = xmax / sqrt(vsq); + if (fsq > 0.0) dtf = sqrt(2.0 * xmax / (ftm2v * sqrt(fsq) * massinv)); + k_dt = MIN(dtv, dtf); + if ((emax > 0.0) && (fsq * vsq > 0.0)) { + dte = emax / sqrt(fsq * vsq) / sqrt(ftm2v * mvv2e); + k_dt = MIN(dt, dte); + } + dtsq = k_dt * k_dt; + delx = k_dt * v(i,0) + 0.5 * dtsq * massinv * f(i,0) * ftm2v; + dely = k_dt * v(i,1) + 0.5 * dtsq * massinv * f(i,1) * ftm2v; + delz = k_dt * v(i,2) + 0.5 * dtsq * massinv * f(i,2) * ftm2v; + delr = sqrt(delx * delx + dely * dely + delz * delz); + if (delr > xmax) k_dt *= xmax / delr; + } +} + +namespace LAMMPS_NS { +template class FixDtResetKokkos; +#ifdef LMP_KOKKOS_GPU +template class FixDtResetKokkos; +#endif +} + diff --git a/src/KOKKOS/fix_dt_reset_kokkos.h b/src/KOKKOS/fix_dt_reset_kokkos.h new file mode 100644 index 0000000000..a269b58137 --- /dev/null +++ b/src/KOKKOS/fix_dt_reset_kokkos.h @@ -0,0 +1,66 @@ +/* -*- 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 FIX_CLASS +// clang-format off +FixStyle(dt/reset/kk,FixDtResetKokkos); +FixStyle(dt/reset/kk/device,FixDtResetKokkos); +FixStyle(dt/reset/kk/host,FixDtResetKokkos); +// clang-format on +#else + +// clang-format off +#ifndef LMP_FIX_DT_RESET_KOKKOS_H +#define LMP_FIX_DT_RESET_KOKKOS_H + +#include "fix_dt_reset.h" +#include "kokkos_type.h" + +namespace LAMMPS_NS { + +struct TagFixDtResetMass{}; +struct TagFixDtResetRMass{}; + +template +class FixDtResetKokkos : public FixDtReset { + public: + typedef DeviceType device_type; + typedef ArrayTypes AT; + + FixDtResetKokkos(class LAMMPS *, int, char **); +// ~FixDtResetKokkos() override; + void init() override; + void end_of_step() override; + + KOKKOS_INLINE_FUNCTION + void operator()(TagFixDtResetMass, const int&, double&) const; + KOKKOS_INLINE_FUNCTION + void operator()(TagFixDtResetRMass, const int&, double&) const; + + private: + typename AT::t_v_array v; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread mask; + typename AT::t_int_1d_randomread type; + typename ArrayTypes::t_float_1d_randomread rmass; + typename ArrayTypes::t_float_1d_randomread mass; + + + Kokkos::DualView k_emax; +}; + +} + +#endif +#endif + diff --git a/src/KOKKOS/fix_viscous_kokkos.cpp b/src/KOKKOS/fix_viscous_kokkos.cpp new file mode 100644 index 0000000000..80ddff2fce --- /dev/null +++ b/src/KOKKOS/fix_viscous_kokkos.cpp @@ -0,0 +1,106 @@ +// 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 "fix_viscous_kokkos.h" + +#include "atom_kokkos.h" +#include "update.h" +#include "modify.h" +#include "input.h" +#include "memory_kokkos.h" +#include "error.h" +#include "atom_masks.h" +#include "kokkos_base.h" + +using namespace LAMMPS_NS; +using namespace FixConst; + +/* ---------------------------------------------------------------------- */ + +template +FixViscousKokkos::FixViscousKokkos(LAMMPS *lmp, int narg, char **arg) : + FixViscous(lmp, narg, arg) +{ + kokkosable = 1; + atomKK = (AtomKokkos *) atom; + execution_space = ExecutionSpaceFromDevice::space; + datamask_read = EMPTY_MASK; + datamask_modify = EMPTY_MASK; +} + +/* ---------------------------------------------------------------------- */ + +template +FixViscousKokkos::~FixViscousKokkos() +{ + if (copymode) return; +} + +/* ---------------------------------------------------------------------- */ + +template +void FixViscousKokkos::init() +{ + FixViscous::init(); + + k_gamma = Kokkos::DualView("FixViscousKokkos:gamma",atom->ntypes+1); + + for (int i = 1; i <= atom->ntypes; i++) k_gamma.h_view(i) = gamma[i]; + + k_gamma.template modify(); + k_gamma.template sync(); + + if (utils::strmatch(update->integrate_style,"^respa")) + error->all(FLERR,"Cannot (yet) use respa with Kokkos"); +} + +/* ---------------------------------------------------------------------- */ + +template +void FixViscousKokkos::post_force(int /*vflag*/) +{ + atomKK->sync(execution_space, V_MASK | F_MASK | MASK_MASK | TYPE_MASK); + + v = atomKK->k_v.view(); + f = atomKK->k_f.view(); + mask = atomKK->k_mask.view(); + type = atomKK->k_type.view(); + + int nlocal = atom->nlocal; + + copymode = 1; + Kokkos::parallel_for(Kokkos::RangePolicy(0,nlocal),*this); + copymode = 0; + + atomKK->modified(execution_space, F_MASK); +} + +template +KOKKOS_INLINE_FUNCTION +void FixViscousKokkos::operator()(TagFixViscous, const int &i) const { + if (mask[i] & groupbit) { + double drag = k_gamma.d_view(type[i]); + f(i,0) -= drag*v(i,0); + f(i,1) -= drag*v(i,1); + f(i,2) -= drag*v(i,2); + } +} + +namespace LAMMPS_NS { +template class FixViscousKokkos; +#ifdef LMP_KOKKOS_GPU +template class FixViscousKokkos; +#endif +} + diff --git a/src/KOKKOS/fix_viscous_kokkos.h b/src/KOKKOS/fix_viscous_kokkos.h new file mode 100644 index 0000000000..2594c0cf85 --- /dev/null +++ b/src/KOKKOS/fix_viscous_kokkos.h @@ -0,0 +1,60 @@ +/* -*- 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 FIX_CLASS +// clang-format off +FixStyle(viscous/kk,FixViscousKokkos); +FixStyle(viscous/kk/device,FixViscousKokkos); +FixStyle(viscous/kk/host,FixViscousKokkos); +// clang-format on +#else + +// clang-format off +#ifndef LMP_FIX_VISCOUS_KOKKOS_H +#define LMP_FIX_VISCOUS_KOKKOS_H + +#include "fix_viscous.h" +#include "kokkos_type.h" + +namespace LAMMPS_NS { + +struct TagFixViscous{}; + +template +class FixViscousKokkos : public FixViscous { + public: + typedef DeviceType device_type; + typedef ArrayTypes AT; + + FixViscousKokkos(class LAMMPS *, int, char **); + ~FixViscousKokkos() override; + void init() override; + void post_force(int) override; + + KOKKOS_INLINE_FUNCTION + void operator()(TagFixViscous, const int&) const; + + private: + typename AT::t_v_array v; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread mask; + typename AT::t_int_1d_randomread type; + + Kokkos::DualView k_gamma; +}; + +} + +#endif +#endif + diff --git a/src/fix_dt_reset.h b/src/fix_dt_reset.h index f0a0564685..00070dac4b 100644 --- a/src/fix_dt_reset.h +++ b/src/fix_dt_reset.h @@ -34,7 +34,7 @@ class FixDtReset : public Fix { void end_of_step() override; double compute_scalar() override; - private: + protected: bigint laststep; int minbound, maxbound; double tmin, tmax, xmax, emax; diff --git a/src/fix_viscous.cpp b/src/fix_viscous.cpp index 740dbd91fb..fb0b38f264 100644 --- a/src/fix_viscous.cpp +++ b/src/fix_viscous.cpp @@ -61,6 +61,8 @@ FixViscous::FixViscous(LAMMPS *lmp, int narg, char **arg) : FixViscous::~FixViscous() { + if (copymode) return; + delete [] gamma; }