diff --git a/doc/src/Commands_bond.rst b/doc/src/Commands_bond.rst index 0ebda6175b..2664b74076 100644 --- a/doc/src/Commands_bond.rst +++ b/doc/src/Commands_bond.rst @@ -73,7 +73,7 @@ OPT. * :doc:`none ` * :doc:`zero ` - * :doc:`hybrid ` + * :doc:`hybrid (k) ` * * * @@ -119,7 +119,7 @@ OPT. * :doc:`none ` * :doc:`zero ` - * :doc:`hybrid ` + * :doc:`hybrid (k) ` * * * @@ -157,7 +157,7 @@ OPT. * :doc:`none ` * :doc:`zero ` - * :doc:`hybrid ` + * :doc:`hybrid (k) ` * * * diff --git a/doc/src/angle_hybrid.rst b/doc/src/angle_hybrid.rst index cf8ea1a3fb..5eba42cd25 100644 --- a/doc/src/angle_hybrid.rst +++ b/doc/src/angle_hybrid.rst @@ -1,8 +1,11 @@ .. index:: angle_style hybrid +.. index:: angle_style hybrid/kk angle_style hybrid command ========================== +Accelerator Variants: *hybrid/kk* + Syntax """""" diff --git a/doc/src/dihedral_hybrid.rst b/doc/src/dihedral_hybrid.rst index 998a752275..b22fecef44 100644 --- a/doc/src/dihedral_hybrid.rst +++ b/doc/src/dihedral_hybrid.rst @@ -1,8 +1,11 @@ .. index:: dihedral_style hybrid +.. index:: dihedral_style hybrid/kk dihedral_style hybrid command ============================= +Accelerator Variants: *hybrid/kk* + Syntax """""" diff --git a/doc/src/improper_hybrid.rst b/doc/src/improper_hybrid.rst index a829d9989c..c2f80bdd25 100644 --- a/doc/src/improper_hybrid.rst +++ b/doc/src/improper_hybrid.rst @@ -1,8 +1,11 @@ .. index:: improper_style hybrid +.. index:: improper_style hybrid/kk improper_style hybrid command ============================= +Accelerator Variants: *hybrid/kk* + Syntax """""" diff --git a/src/KOKKOS/Install.sh b/src/KOKKOS/Install.sh index 7abdba17ea..94effc5d68 100755 --- a/src/KOKKOS/Install.sh +++ b/src/KOKKOS/Install.sh @@ -53,6 +53,8 @@ action angle_cosine_kokkos.cpp angle_cosine.cpp action angle_cosine_kokkos.h angle_cosine.h action angle_harmonic_kokkos.cpp angle_harmonic.cpp action angle_harmonic_kokkos.h angle_harmonic.h +action angle_hybrid_kokkos.cpp angle_hybrid.cpp +action angle_hybrid_kokkos.h angle_hybrid.h action angle_spica_kokkos.cpp angle_spica.cpp action angle_spica_kokkos.h angle_spica.h action atom_kokkos.cpp @@ -118,6 +120,8 @@ action dihedral_harmonic_kokkos.cpp dihedral_harmonic.cpp action dihedral_harmonic_kokkos.h dihedral_harmonic.h action dihedral_opls_kokkos.cpp dihedral_opls.cpp action dihedral_opls_kokkos.h dihedral_opls.h +action dihedral_hybrid_kokkos.cpp dihedral_hybrid.cpp +action dihedral_hybrid_kokkos.h dihedral_hybrid.h action domain_kokkos.cpp action domain_kokkos.h action dynamical_matrix_kokkos.cpp dynamical_matrix.cpp @@ -207,6 +211,8 @@ action improper_class2_kokkos.cpp improper_class2.cpp action improper_class2_kokkos.h improper_class2.h action improper_harmonic_kokkos.cpp improper_harmonic.cpp action improper_harmonic_kokkos.h improper_harmonic.h +action improper_hybrid_kokkos.cpp improper_hybrid.cpp +action improper_hybrid_kokkos.h improper_hybrid.h action kissfft_kokkos.h kissfft.h action kokkos_base_fft.h fft3d.h action kokkos_base.h diff --git a/src/KOKKOS/angle_charmm_kokkos.cpp b/src/KOKKOS/angle_charmm_kokkos.cpp index 666002686c..22d2b924dd 100644 --- a/src/KOKKOS/angle_charmm_kokkos.cpp +++ b/src/KOKKOS/angle_charmm_kokkos.cpp @@ -38,6 +38,7 @@ static constexpr double SMALL = 0.001; template AngleCharmmKokkos::AngleCharmmKokkos(LAMMPS *lmp) : AngleCharmm(lmp) { + kokkosable = 1; atomKK = (AtomKokkos *) atom; neighborKK = (NeighborKokkos *) neighbor; execution_space = ExecutionSpaceFromDevice::space; @@ -125,12 +126,12 @@ void AngleCharmmKokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { k_eatom.template modify(); - k_eatom.template sync(); + k_eatom.sync_host(); } if (vflag_atom) { k_vatom.template modify(); - k_vatom.template sync(); + k_vatom.sync_host(); } copymode = 0; @@ -284,10 +285,10 @@ void AngleCharmmKokkos::coeff(int narg, char **arg) k_r_ub.h_view[i] = r_ub[i]; } - k_k.template modify(); - k_theta0.template modify(); - k_k_ub.template modify(); - k_r_ub.template modify(); + k_k.modify_host(); + k_theta0.modify_host(); + k_k_ub.modify_host(); + k_r_ub.modify_host(); k_k.template sync(); k_theta0.template sync(); @@ -322,10 +323,10 @@ void AngleCharmmKokkos::read_restart(FILE *fp) k_r_ub.h_view[i] = r_ub[i]; } - k_k.template modify(); - k_theta0.template modify(); - k_k_ub.template modify(); - k_r_ub.template modify(); + k_k.modify_host(); + k_theta0.modify_host(); + k_k_ub.modify_host(); + k_r_ub.modify_host(); k_k.template sync(); k_theta0.template sync(); diff --git a/src/KOKKOS/angle_charmm_kokkos.h b/src/KOKKOS/angle_charmm_kokkos.h index 197f9160a0..2bb06725f9 100644 --- a/src/KOKKOS/angle_charmm_kokkos.h +++ b/src/KOKKOS/angle_charmm_kokkos.h @@ -58,19 +58,18 @@ class AngleCharmmKokkos : public AngleCharmm { const F_FLOAT &delx1, const F_FLOAT &dely1, const F_FLOAT &delz1, const F_FLOAT &delx2, const F_FLOAT &dely2, const F_FLOAT &delz2) const; + using KKDeviceType = typename KKDevice::value; + Kokkos::DualView k_eatom; + Kokkos::DualView k_vatom; + protected: class NeighborKokkos *neighborKK; typedef ArrayTypes AT; typename AT::t_x_array_randomread x; - - using KKDeviceType = typename KKDevice::value; typename Kokkos::View > f; typename AT::t_int_2d anglelist; - - Kokkos::DualView k_eatom; - Kokkos::DualView k_vatom; Kokkos::View> d_eatom; Kokkos::View> d_vatom; diff --git a/src/KOKKOS/angle_class2_kokkos.cpp b/src/KOKKOS/angle_class2_kokkos.cpp index e831ae2283..e9d4797e71 100644 --- a/src/KOKKOS/angle_class2_kokkos.cpp +++ b/src/KOKKOS/angle_class2_kokkos.cpp @@ -38,6 +38,7 @@ static constexpr double SMALL = 0.001; template AngleClass2Kokkos::AngleClass2Kokkos(LAMMPS *lmp) : AngleClass2(lmp) { + kokkosable = 1; atomKK = (AtomKokkos *) atom; neighborKK = (NeighborKokkos *) neighbor; execution_space = ExecutionSpaceFromDevice::space; @@ -141,12 +142,12 @@ void AngleClass2Kokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { k_eatom.template modify(); - k_eatom.template sync(); + k_eatom.sync_host(); } if (vflag_atom) { k_vatom.template modify(); - k_vatom.template sync(); + k_vatom.sync_host(); } copymode = 0; @@ -386,21 +387,21 @@ void AngleClass2Kokkos::coeff(int narg, char **arg) k_theta0.h_view[i] = theta0[i]; } - k_k2.template modify(); - k_k3.template modify(); - k_k4.template modify(); - k_bb_k.template modify(); - k_bb_r1.template modify(); - k_bb_r2.template modify(); - k_ba_k1.template modify(); - k_ba_k2.template modify(); - k_ba_r1.template modify(); - k_ba_r2.template modify(); - k_setflag.template modify(); - k_setflag_a.template modify(); - k_setflag_bb.template modify(); - k_setflag_ba.template modify(); - k_theta0.template modify(); + k_k2.modify_host(); + k_k3.modify_host(); + k_k4.modify_host(); + k_bb_k.modify_host(); + k_bb_r1.modify_host(); + k_bb_r2.modify_host(); + k_ba_k1.modify_host(); + k_ba_k2.modify_host(); + k_ba_r1.modify_host(); + k_ba_r2.modify_host(); + k_setflag.modify_host(); + k_setflag_a.modify_host(); + k_setflag_bb.modify_host(); + k_setflag_ba.modify_host(); + k_theta0.modify_host(); } /* ---------------------------------------------------------------------- @@ -465,21 +466,21 @@ void AngleClass2Kokkos::read_restart(FILE *fp) k_theta0.h_view[i] = theta0[i]; } - k_k2.template modify(); - k_k3.template modify(); - k_k4.template modify(); - k_bb_k.template modify(); - k_bb_r1.template modify(); - k_bb_r2.template modify(); - k_ba_k1.template modify(); - k_ba_k2.template modify(); - k_ba_r1.template modify(); - k_ba_r2.template modify(); - k_setflag.template modify(); - k_setflag_a.template modify(); - k_setflag_bb.template modify(); - k_setflag_ba.template modify(); - k_theta0.template modify(); + k_k2.modify_host(); + k_k3.modify_host(); + k_k4.modify_host(); + k_bb_k.modify_host(); + k_bb_r1.modify_host(); + k_bb_r2.modify_host(); + k_ba_k1.modify_host(); + k_ba_k2.modify_host(); + k_ba_r1.modify_host(); + k_ba_r2.modify_host(); + k_setflag.modify_host(); + k_setflag_a.modify_host(); + k_setflag_bb.modify_host(); + k_setflag_ba.modify_host(); + k_theta0.modify_host(); } /* ---------------------------------------------------------------------- diff --git a/src/KOKKOS/angle_class2_kokkos.h b/src/KOKKOS/angle_class2_kokkos.h index 81bed169bc..7ef9e9b652 100644 --- a/src/KOKKOS/angle_class2_kokkos.h +++ b/src/KOKKOS/angle_class2_kokkos.h @@ -36,8 +36,8 @@ class AngleClass2Kokkos : public AngleClass2 { public: typedef DeviceType device_type; - typedef ArrayTypes AT; typedef EV_FLOAT value_type; + typedef ArrayTypes AT; AngleClass2Kokkos(class LAMMPS *); ~AngleClass2Kokkos() override; @@ -60,6 +60,9 @@ class AngleClass2Kokkos : public AngleClass2 { const F_FLOAT &delx1, const F_FLOAT &dely1, const F_FLOAT &delz1, const F_FLOAT &delx2, const F_FLOAT &dely2, const F_FLOAT &delz2) const; + typename AT::tdual_efloat_1d k_eatom; + typename AT::tdual_virial_array k_vatom; + protected: class NeighborKokkos *neighborKK; @@ -67,9 +70,6 @@ class AngleClass2Kokkos : public AngleClass2 { typename AT::t_x_array_randomread x; typename AT::t_f_array f; typename AT::t_int_2d anglelist; - - typename AT::tdual_efloat_1d k_eatom; - typename AT::tdual_virial_array k_vatom; typename AT::t_efloat_1d d_eatom; typename AT::t_virial_array d_vatom; diff --git a/src/KOKKOS/angle_cosine_kokkos.cpp b/src/KOKKOS/angle_cosine_kokkos.cpp index 768dfd43ca..5d61213df8 100644 --- a/src/KOKKOS/angle_cosine_kokkos.cpp +++ b/src/KOKKOS/angle_cosine_kokkos.cpp @@ -36,6 +36,7 @@ using namespace MathConst; template AngleCosineKokkos::AngleCosineKokkos(LAMMPS *lmp) : AngleCosine(lmp) { + kokkosable = 1; atomKK = (AtomKokkos *) atom; neighborKK = (NeighborKokkos *) neighbor; execution_space = ExecutionSpaceFromDevice::space; @@ -124,12 +125,12 @@ void AngleCosineKokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { k_eatom.template modify(); - k_eatom.template sync(); + k_eatom.sync_host(); } if (vflag_atom) { k_vatom.template modify(); - k_vatom.template sync(); + k_vatom.sync_host(); } copymode = 0; @@ -254,7 +255,7 @@ void AngleCosineKokkos::coeff(int narg, char **arg) for (int i = 1; i <= n; i++) k_k.h_view[i] = k[i]; - k_k.template modify(); + k_k.modify_host(); } /* ---------------------------------------------------------------------- @@ -270,7 +271,7 @@ void AngleCosineKokkos::read_restart(FILE *fp) for (int i = 1; i <= n; i++) k_k.h_view[i] = k[i]; - k_k.template modify(); + k_k.modify_host(); } /* ---------------------------------------------------------------------- diff --git a/src/KOKKOS/angle_cosine_kokkos.h b/src/KOKKOS/angle_cosine_kokkos.h index 33b80f5a5f..3cfa5a73df 100644 --- a/src/KOKKOS/angle_cosine_kokkos.h +++ b/src/KOKKOS/angle_cosine_kokkos.h @@ -37,6 +37,7 @@ class AngleCosineKokkos : public AngleCosine { public: typedef DeviceType device_type; typedef EV_FLOAT value_type; + typedef ArrayTypes AT; AngleCosineKokkos(class LAMMPS *); ~AngleCosineKokkos() override; @@ -59,6 +60,9 @@ class AngleCosineKokkos : public AngleCosine { const F_FLOAT &delx1, const F_FLOAT &dely1, const F_FLOAT &delz1, const F_FLOAT &delx2, const F_FLOAT &dely2, const F_FLOAT &delz2) const; + typename AT::tdual_efloat_1d k_eatom; + typename AT::tdual_virial_array k_vatom; + protected: class NeighborKokkos *neighborKK; @@ -66,9 +70,6 @@ class AngleCosineKokkos : public AngleCosine { typename ArrayTypes::t_x_array_randomread x; typename ArrayTypes::t_f_array f; typename ArrayTypes::t_int_2d anglelist; - - typename ArrayTypes::tdual_efloat_1d k_eatom; - typename ArrayTypes::tdual_virial_array k_vatom; typename ArrayTypes::t_efloat_1d d_eatom; typename ArrayTypes::t_virial_array d_vatom; diff --git a/src/KOKKOS/angle_harmonic_kokkos.cpp b/src/KOKKOS/angle_harmonic_kokkos.cpp index d7be418326..2b3c283732 100644 --- a/src/KOKKOS/angle_harmonic_kokkos.cpp +++ b/src/KOKKOS/angle_harmonic_kokkos.cpp @@ -38,6 +38,7 @@ static constexpr double SMALL = 0.001; template AngleHarmonicKokkos::AngleHarmonicKokkos(LAMMPS *lmp) : AngleHarmonic(lmp) { + kokkosable = 1; atomKK = (AtomKokkos *) atom; neighborKK = (NeighborKokkos *) neighbor; execution_space = ExecutionSpaceFromDevice::space; @@ -71,14 +72,18 @@ void AngleHarmonicKokkos::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { + if(k_eatom.extent(0) < maxeatom) { memoryKK->destroy_kokkos(k_eatom,eatom); memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"angle:eatom"); d_eatom = k_eatom.template view(); + } else Kokkos::deep_copy(d_eatom,0.0); } if (vflag_atom) { + if(k_vatom.extent(0) < maxvatom) { memoryKK->destroy_kokkos(k_vatom,vatom); memoryKK->create_kokkos(k_vatom,vatom,maxvatom,"angle:vatom"); d_vatom = k_vatom.template view(); + } else Kokkos::deep_copy(d_vatom,0.0); } //atomKK->sync(execution_space,datamask_read); @@ -127,12 +132,12 @@ void AngleHarmonicKokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { k_eatom.template modify(); - k_eatom.template sync(); + k_eatom.sync_host(); } if (vflag_atom) { k_vatom.template modify(); - k_vatom.template sync(); + k_vatom.sync_host(); } copymode = 0; @@ -264,8 +269,8 @@ void AngleHarmonicKokkos::coeff(int narg, char **arg) k_theta0.h_view[i] = theta0[i]; } - k_k.template modify(); - k_theta0.template modify(); + k_k.modify_host(); + k_theta0.modify_host(); } /* ---------------------------------------------------------------------- @@ -283,8 +288,8 @@ void AngleHarmonicKokkos::read_restart(FILE *fp) k_theta0.h_view[i] = theta0[i]; } - k_k.template modify(); - k_theta0.template modify(); + k_k.modify_host(); + k_theta0.modify_host(); } /* ---------------------------------------------------------------------- diff --git a/src/KOKKOS/angle_harmonic_kokkos.h b/src/KOKKOS/angle_harmonic_kokkos.h index ce4b9d9976..4e427d2417 100644 --- a/src/KOKKOS/angle_harmonic_kokkos.h +++ b/src/KOKKOS/angle_harmonic_kokkos.h @@ -37,6 +37,7 @@ class AngleHarmonicKokkos : public AngleHarmonic { public: typedef DeviceType device_type; typedef EV_FLOAT value_type; + typedef ArrayTypes AT; AngleHarmonicKokkos(class LAMMPS *); ~AngleHarmonicKokkos() override; @@ -59,6 +60,9 @@ class AngleHarmonicKokkos : public AngleHarmonic { const F_FLOAT &delx1, const F_FLOAT &dely1, const F_FLOAT &delz1, const F_FLOAT &delx2, const F_FLOAT &dely2, const F_FLOAT &delz2) const; + typename AT::tdual_efloat_1d k_eatom; + typename AT::tdual_virial_array k_vatom; + protected: class NeighborKokkos *neighborKK; @@ -66,9 +70,6 @@ class AngleHarmonicKokkos : public AngleHarmonic { typename ArrayTypes::t_x_array_randomread x; typename ArrayTypes::t_f_array f; typename ArrayTypes::t_int_2d anglelist; - - typename ArrayTypes::tdual_efloat_1d k_eatom; - typename ArrayTypes::tdual_virial_array k_vatom; typename ArrayTypes::t_efloat_1d d_eatom; typename ArrayTypes::t_virial_array d_vatom; diff --git a/src/KOKKOS/angle_hybrid_kokkos.cpp b/src/KOKKOS/angle_hybrid_kokkos.cpp new file mode 100644 index 0000000000..06b2845545 --- /dev/null +++ b/src/KOKKOS/angle_hybrid_kokkos.cpp @@ -0,0 +1,224 @@ +/* ---------------------------------------------------------------------- + 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 "angle_hybrid_kokkos.h" + +#include "atom_kokkos.h" +#include "atom_masks.h" +#include "comm.h" +#include "error.h" +#include "force.h" +#include "kokkos.h" +#include "memory_kokkos.h" +#include "neighbor_kokkos.h" + +#include + +using namespace LAMMPS_NS; + +#define EXTRA 1000 + +/* ---------------------------------------------------------------------- */ + +AngleHybridKokkos::AngleHybridKokkos(LAMMPS *lmp) : AngleHybrid(lmp) +{ + kokkosable = 1; + + atomKK = (AtomKokkos *) atom; + neighborKK = (NeighborKokkos *) neighbor; + + execution_space = Device; + + datamask_read = EMPTY_MASK; + datamask_modify = EMPTY_MASK; +} + +/* ---------------------------------------------------------------------- */ + +AngleHybridKokkos::~AngleHybridKokkos() +{ + deallocate(); +} + +/* ---------------------------------------------------------------------- */ + +void AngleHybridKokkos::compute(int eflag, int vflag) +{ + // save ptrs to original anglelist + + int nanglelist_orig = neighbor->nanglelist; + neighborKK->k_anglelist.sync_device(); + auto k_anglelist_orig = neighborKK->k_anglelist; + auto d_anglelist_orig = k_anglelist_orig.d_view; + auto d_nanglelist = k_nanglelist.d_view; + auto h_nanglelist = k_nanglelist.h_view; + + // if this is re-neighbor step, create sub-style anglelists + // nanglelist[] = length of each sub-style list + // realloc sub-style anglelist if necessary + // load sub-style anglelist with 3 values from original anglelist + + if (neighbor->ago == 0) { + Kokkos::deep_copy(d_nanglelist,0); + + k_map.sync_device(); + auto d_map = k_map.d_view; + + Kokkos::parallel_for(nanglelist_orig,LAMMPS_LAMBDA(int i) { + const int m = d_map[d_anglelist_orig(i,3)]; + if (m >= 0) Kokkos::atomic_increment(&d_nanglelist[m]); + }); + + k_nanglelist.modify_device(); + k_nanglelist.sync_host(); + + maxangle_all = 0; + for (int m = 0; m < nstyles; m++) + if (h_nanglelist[m] > maxangle_all) + maxangle_all = h_nanglelist[m] + EXTRA; + + if (k_anglelist.d_view.extent(1) < maxangle_all) + MemKK::realloc_kokkos(k_anglelist, "angle_hybrid:anglelist", nstyles, maxangle_all, 4); + auto d_anglelist = k_anglelist.d_view; + + Kokkos::deep_copy(d_nanglelist,0); + + Kokkos::parallel_for(nanglelist_orig,LAMMPS_LAMBDA(int i) { + const int m = d_map[d_anglelist_orig(i,3)]; + if (m < 0) return; + const int n = Kokkos::atomic_fetch_add(&d_nanglelist[m],1); + d_anglelist(m,n,0) = d_anglelist_orig(i,0); + d_anglelist(m,n,1) = d_anglelist_orig(i,1); + d_anglelist(m,n,2) = d_anglelist_orig(i,2); + d_anglelist(m,n,3) = d_anglelist_orig(i,3); + }); + } + + // call each sub-style's compute function + // set neighbor->anglelist to sub-style anglelist before call + // accumulate sub-style global/peratom energy/virial in hybrid + + ev_init(eflag, vflag); + + k_nanglelist.modify_device(); + k_nanglelist.sync_host(); + + for (int m = 0; m < nstyles; m++) { + neighbor->nanglelist = h_nanglelist[m]; + auto k_anglelist_m = Kokkos::subview(k_anglelist,m,Kokkos::ALL,Kokkos::ALL); + k_anglelist_m.modify_device(); + neighborKK->k_anglelist = k_anglelist_m; + + auto style = styles[m]; + atomKK->sync(style->execution_space,style->datamask_read); + style->compute(eflag, vflag); + atomKK->modified(style->execution_space,style->datamask_modify); + + if (eflag_global) energy += style->energy; + if (vflag_global) + for (int n = 0; n < 6; n++) virial[n] += style->virial[n]; + + if (eflag_atom) { + int n = atom->nlocal; + if (force->newton_bond) n += atom->nghost; + double *eatom_substyle = styles[m]->eatom; + for (int i = 0; i < n; i++) eatom[i] += eatom_substyle[i]; + } + if (vflag_atom) { + int n = atom->nlocal; + if (force->newton_bond) n += atom->nghost; + double **vatom_substyle = styles[m]->vatom; + for (int i = 0; i < n; i++) + for (int j = 0; j < 6; j++) vatom[i][j] += vatom_substyle[i][j]; + } + if (cvflag_atom) { + int n = atom->nlocal; + if (force->newton_bond) n += atom->nghost; + double **cvatom_substyle = styles[m]->cvatom; + for (int i = 0; i < n; i++) + for (int j = 0; j < 9; j++) cvatom[i][j] += cvatom_substyle[i][j]; + } + } + + // restore ptrs to original anglelist + + neighbor->nanglelist = nanglelist_orig; + neighborKK->k_anglelist = k_anglelist_orig; +} + +/* ---------------------------------------------------------------------- */ + +void AngleHybridKokkos::allocate() +{ + allocated = 1; + int np1 = atom->nangletypes + 1; + + memoryKK->create_kokkos(k_map, map, np1, "angle:map"); + memory->create(setflag, np1, "angle:setflag"); + for (int i = 1; i < np1; i++) setflag[i] = 0; + + k_nanglelist = DAT::tdual_int_1d("angle:nanglelist", nstyles); +} + +/* ---------------------------------------------------------------------- */ + +void AngleHybridKokkos::deallocate() +{ + if (!allocated) return; + + allocated = 0; + + memory->destroy(setflag); + memoryKK->destroy_kokkos(k_map,map); +} + +/* ---------------------------------------------------------------------- + set coeffs for one type +---------------------------------------------------------------------- */ + +void AngleHybridKokkos::coeff(int narg, char **arg) +{ + AngleHybrid::coeff(narg,arg); + + k_map.modify_host(); +} + +/* ---------------------------------------------------------------------- */ + +void AngleHybridKokkos::init_style() +{ + AngleHybrid::init_style(); + + for (int m = 0; m < nstyles; m++) { + if (!styles[m]->kokkosable) + error->all(FLERR,"Must use only Kokkos-enabled angle styles with angle_style hybrid/kk"); + + if (styles[m]->execution_space == Host) + lmp->kokkos->allow_overlap = 0; + } +} + +/* ---------------------------------------------------------------------- + memory usage +------------------------------------------------------------------------- */ + +double AngleHybridKokkos::memory_usage() +{ + double bytes = (double) maxeatom * sizeof(double); + bytes += (double) maxvatom * 6 * sizeof(double); + bytes += (double) maxcvatom * 9 * sizeof(double); + for (int m = 0; m < nstyles; m++) bytes += (double) maxangle_all * 4 * sizeof(int); + for (int m = 0; m < nstyles; m++) + if (styles[m]) bytes += styles[m]->memory_usage(); + return bytes; +} diff --git a/src/KOKKOS/angle_hybrid_kokkos.h b/src/KOKKOS/angle_hybrid_kokkos.h new file mode 100644 index 0000000000..09b51958eb --- /dev/null +++ b/src/KOKKOS/angle_hybrid_kokkos.h @@ -0,0 +1,58 @@ +/* -*- 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 ANGLE_CLASS +// clang-format off +AngleStyle(hybrid/kk,AngleHybridKokkos); +AngleStyle(hybrid/kk/device,AngleHybridKokkos); +AngleStyle(hybrid/kk/host,AngleHybridKokkos); +// clang-format on +#else + +// clang-format off +#ifndef LMP_ANGLE_HYBRID_KOKKOS_H +#define LMP_ANGLE_HYBRID_KOKKOS_H + +#include "angle_hybrid.h" +#include "kokkos_type.h" + +namespace LAMMPS_NS { + +class AngleHybridKokkos : public AngleHybrid { + friend class Force; + + public: + AngleHybridKokkos(class LAMMPS *); + ~AngleHybridKokkos() override; + void compute(int, int) override; + void coeff(int, char **) override; + void init_style() override; + double memory_usage() override; + + private: + int maxangle_all; + + class NeighborKokkos *neighborKK; + + DAT::tdual_int_1d k_map; // which style each angle type points to + DAT::tdual_int_1d k_nanglelist; // # of angles in sub-style anglelists + DAT::tdual_int_3d k_anglelist; // anglelist for each sub-style + + void allocate() override; + void deallocate() override; +}; + +} // namespace LAMMPS_NS + +#endif +#endif diff --git a/src/KOKKOS/bond_hybrid_kokkos.cpp b/src/KOKKOS/bond_hybrid_kokkos.cpp index d63ebccac6..db247c7100 100644 --- a/src/KOKKOS/bond_hybrid_kokkos.cpp +++ b/src/KOKKOS/bond_hybrid_kokkos.cpp @@ -153,11 +153,11 @@ void BondHybridKokkos::compute(int eflag, int vflag) void BondHybridKokkos::allocate() { allocated = 1; - int n = atom->nbondtypes; + int np1 = atom->nbondtypes + 1; - memoryKK->create_kokkos(k_map, map, n + 1, "bond:map"); - memory->create(setflag, n + 1, "bond:setflag"); - for (int i = 1; i <= n; i++) setflag[i] = 0; + memoryKK->create_kokkos(k_map, map, np1, "bond:map"); + memory->create(setflag, np1, "bond:setflag"); + for (int i = 1; i < np1; i++) setflag[i] = 0; k_nbondlist = DAT::tdual_int_1d("bond:nbondlist", nstyles); } diff --git a/src/KOKKOS/dihedral_charmm_kokkos.cpp b/src/KOKKOS/dihedral_charmm_kokkos.cpp index b385ec7f01..a8939770df 100644 --- a/src/KOKKOS/dihedral_charmm_kokkos.cpp +++ b/src/KOKKOS/dihedral_charmm_kokkos.cpp @@ -40,6 +40,7 @@ static constexpr double TOLERANCE = 0.05; template DihedralCharmmKokkos::DihedralCharmmKokkos(LAMMPS *lmp) : DihedralCharmm(lmp) { + kokkosable = 1; atomKK = (AtomKokkos *) atom; neighborKK = (NeighborKokkos *) neighbor; execution_space = ExecutionSpaceFromDevice::space; @@ -115,7 +116,7 @@ void DihedralCharmmKokkos::compute(int eflag_in, int vflag_in) qqrd2e = force->qqrd2e; h_warning_flag() = 0; - k_warning_flag.template modify(); + k_warning_flag.modify_host(); k_warning_flag.template sync(); copymode = 1; @@ -141,7 +142,7 @@ void DihedralCharmmKokkos::compute(int eflag_in, int vflag_in) // error check k_warning_flag.template modify(); - k_warning_flag.template sync(); + k_warning_flag.sync_host(); if (h_warning_flag()) error->warning(FLERR,"Dihedral problem"); @@ -174,20 +175,20 @@ void DihedralCharmmKokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { k_eatom.template modify(); - k_eatom.template sync(); + k_eatom.sync_host(); k_eatom_pair.template modify(); - k_eatom_pair.template sync(); + k_eatom_pair.sync_host(); for (int i = 0; i < n; i++) force->pair->eatom[i] += k_eatom_pair.h_view(i); } if (vflag_atom) { k_vatom.template modify(); - k_vatom.template sync(); + k_vatom.sync_host(); k_vatom_pair.template modify(); - k_vatom_pair.template sync(); + k_vatom_pair.sync_host(); for (int i = 0; i < n; i++) { force->pair->vatom[i][0] += k_vatom_pair.h_view(i,0); force->pair->vatom[i][1] += k_vatom_pair.h_view(i,1); @@ -454,12 +455,12 @@ void DihedralCharmmKokkos::coeff(int narg, char **arg) k_weight.h_view[i] = weight[i]; } - k_k.template modify(); - k_multiplicity.template modify(); - k_shift.template modify(); - k_cos_shift.template modify(); - k_sin_shift.template modify(); - k_weight.template modify(); + k_k.modify_host(); + k_multiplicity.modify_host(); + k_shift.modify_host(); + k_cos_shift.modify_host(); + k_sin_shift.modify_host(); + k_weight.modify_host(); k_k.template sync(); k_multiplicity.template sync(); @@ -502,10 +503,10 @@ void DihedralCharmmKokkos::init_style() } } - k_lj14_1.template modify(); - k_lj14_2.template modify(); - k_lj14_3.template modify(); - k_lj14_4.template modify(); + k_lj14_1.modify_host(); + k_lj14_2.modify_host(); + k_lj14_3.modify_host(); + k_lj14_4.modify_host(); k_lj14_1.template sync(); k_lj14_2.template sync(); @@ -547,12 +548,12 @@ void DihedralCharmmKokkos::read_restart(FILE *fp) k_weight.h_view[i] = weight[i]; } - k_k.template modify(); - k_multiplicity.template modify(); - k_shift.template modify(); - k_cos_shift.template modify(); - k_sin_shift.template modify(); - k_weight.template modify(); + k_k.modify_host(); + k_multiplicity.modify_host(); + k_shift.modify_host(); + k_cos_shift.modify_host(); + k_sin_shift.modify_host(); + k_weight.modify_host(); k_k.template sync(); k_multiplicity.template sync(); diff --git a/src/KOKKOS/dihedral_charmm_kokkos.h b/src/KOKKOS/dihedral_charmm_kokkos.h index dea251473b..74510ed515 100644 --- a/src/KOKKOS/dihedral_charmm_kokkos.h +++ b/src/KOKKOS/dihedral_charmm_kokkos.h @@ -104,6 +104,10 @@ class DihedralCharmmKokkos : public DihedralCharmm { const F_FLOAT &evdwl, const F_FLOAT &ecoul, const F_FLOAT &fpair, const F_FLOAT &delx, const F_FLOAT &dely, const F_FLOAT &delz) const; + typedef typename KKDevice::value KKDeviceType; + Kokkos::DualView k_eatom; + Kokkos::DualView k_vatom; + protected: class NeighborKokkos *neighborKK; @@ -114,9 +118,6 @@ class DihedralCharmmKokkos : public DihedralCharmm { typename AT::t_f_array f; typename AT::t_int_2d dihedrallist; - typedef typename KKDevice::value KKDeviceType; - Kokkos::DualView k_eatom; - Kokkos::DualView k_vatom; Kokkos::View > d_eatom; Kokkos::View > d_vatom; diff --git a/src/KOKKOS/dihedral_charmmfsw_kokkos.cpp b/src/KOKKOS/dihedral_charmmfsw_kokkos.cpp index aeb9b022a7..3414a02ec4 100644 --- a/src/KOKKOS/dihedral_charmmfsw_kokkos.cpp +++ b/src/KOKKOS/dihedral_charmmfsw_kokkos.cpp @@ -47,6 +47,7 @@ static constexpr double TOLERANCE = 0.05; template DihedralCharmmfswKokkos::DihedralCharmmfswKokkos(LAMMPS *lmp) : DihedralCharmmfsw(lmp) { + kokkosable = 1; atomKK = (AtomKokkos *) atom; neighborKK = (NeighborKokkos *) neighbor; execution_space = ExecutionSpaceFromDevice::space; @@ -122,7 +123,7 @@ void DihedralCharmmfswKokkos::compute(int eflag_in, int vflag_in) qqrd2e = force->qqrd2e; h_warning_flag() = 0; - k_warning_flag.template modify(); + k_warning_flag.modify_host(); k_warning_flag.template sync(); copymode = 1; @@ -148,7 +149,7 @@ void DihedralCharmmfswKokkos::compute(int eflag_in, int vflag_in) // error check k_warning_flag.template modify(); - k_warning_flag.template sync(); + k_warning_flag.sync_host(); if (h_warning_flag()) error->warning(FLERR,"Dihedral problem"); @@ -181,20 +182,20 @@ void DihedralCharmmfswKokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { k_eatom.template modify(); - k_eatom.template sync(); + k_eatom.sync_host(); k_eatom_pair.template modify(); - k_eatom_pair.template sync(); + k_eatom_pair.sync_host(); for (int i = 0; i < n; i++) force->pair->eatom[i] += k_eatom_pair.h_view(i); } if (vflag_atom) { k_vatom.template modify(); - k_vatom.template sync(); + k_vatom.sync_host(); k_vatom_pair.template modify(); - k_vatom_pair.template sync(); + k_vatom_pair.sync_host(); for (int i = 0; i < n; i++) { force->pair->vatom[i][0] += k_vatom_pair.h_view(i,0); force->pair->vatom[i][1] += k_vatom_pair.h_view(i,1); @@ -471,12 +472,12 @@ void DihedralCharmmfswKokkos::coeff(int narg, char **arg) k_weight.h_view[i] = weight[i]; } - k_k.template modify(); - k_multiplicity.template modify(); - k_shift.template modify(); - k_cos_shift.template modify(); - k_sin_shift.template modify(); - k_weight.template modify(); + k_k.modify_host(); + k_multiplicity.modify_host(); + k_shift.modify_host(); + k_cos_shift.modify_host(); + k_sin_shift.modify_host(); + k_weight.modify_host(); k_k.template sync(); k_multiplicity.template sync(); @@ -519,10 +520,10 @@ void DihedralCharmmfswKokkos::init_style() } } - k_lj14_1.template modify(); - k_lj14_2.template modify(); - k_lj14_3.template modify(); - k_lj14_4.template modify(); + k_lj14_1.modify_host(); + k_lj14_2.modify_host(); + k_lj14_3.modify_host(); + k_lj14_4.modify_host(); k_lj14_1.template sync(); k_lj14_2.template sync(); @@ -564,12 +565,12 @@ void DihedralCharmmfswKokkos::read_restart(FILE *fp) k_weight.h_view[i] = weight[i]; } - k_k.template modify(); - k_multiplicity.template modify(); - k_shift.template modify(); - k_cos_shift.template modify(); - k_sin_shift.template modify(); - k_weight.template modify(); + k_k.modify_host(); + k_multiplicity.modify_host(); + k_shift.modify_host(); + k_cos_shift.modify_host(); + k_sin_shift.modify_host(); + k_weight.modify_host(); k_k.template sync(); k_multiplicity.template sync(); diff --git a/src/KOKKOS/dihedral_charmmfsw_kokkos.h b/src/KOKKOS/dihedral_charmmfsw_kokkos.h index b1c65ae477..845d2192d7 100644 --- a/src/KOKKOS/dihedral_charmmfsw_kokkos.h +++ b/src/KOKKOS/dihedral_charmmfsw_kokkos.h @@ -67,6 +67,10 @@ class DihedralCharmmfswKokkos : public DihedralCharmmfsw { const F_FLOAT &evdwl, const F_FLOAT &ecoul, const F_FLOAT &fpair, const F_FLOAT &delx, const F_FLOAT &dely, const F_FLOAT &delz) const; + typedef typename KKDevice::value KKDeviceType; + Kokkos::DualView k_eatom; + Kokkos::DualView k_vatom; + protected: class NeighborKokkos *neighborKK; @@ -76,10 +80,6 @@ class DihedralCharmmfswKokkos : public DihedralCharmmfsw { typename AT::t_ffloat_1d_randomread q; typename AT::t_f_array f; typename AT::t_int_2d dihedrallist; - - typedef typename KKDevice::value KKDeviceType; - Kokkos::DualView k_eatom; - Kokkos::DualView k_vatom; Kokkos::View > d_eatom; Kokkos::View > d_vatom; diff --git a/src/KOKKOS/dihedral_class2_kokkos.cpp b/src/KOKKOS/dihedral_class2_kokkos.cpp index 204a6d0d1a..83c03d856a 100644 --- a/src/KOKKOS/dihedral_class2_kokkos.cpp +++ b/src/KOKKOS/dihedral_class2_kokkos.cpp @@ -38,6 +38,7 @@ static constexpr double SMALL = 0.001; template DihedralClass2Kokkos::DihedralClass2Kokkos(LAMMPS *lmp) : DihedralClass2(lmp) { + kokkosable = 1; atomKK = (AtomKokkos *) atom; neighborKK = (NeighborKokkos *) neighbor; execution_space = ExecutionSpaceFromDevice::space; @@ -137,7 +138,7 @@ void DihedralClass2Kokkos::compute(int eflag_in, int vflag_in) newton_bond = force->newton_bond; h_warning_flag() = 0; - k_warning_flag.template modify(); + k_warning_flag.modify_host(); k_warning_flag.template sync(); copymode = 1; @@ -163,7 +164,7 @@ void DihedralClass2Kokkos::compute(int eflag_in, int vflag_in) // error check k_warning_flag.template modify(); - k_warning_flag.template sync(); + k_warning_flag.sync_host(); if (h_warning_flag()) error->warning(FLERR,"Dihedral problem"); @@ -179,12 +180,12 @@ void DihedralClass2Kokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { k_eatom.template modify(); - k_eatom.template sync(); + k_eatom.sync_host(); } if (vflag_atom) { k_vatom.template modify(); - k_vatom.template sync(); + k_vatom.sync_host(); } copymode = 0; @@ -786,44 +787,44 @@ void DihedralClass2Kokkos::coeff(int narg, char **arg) k_setflag_bb13t.h_view[i] = setflag_bb13t[i]; } - k_k1.template modify(); - k_k2.template modify(); - k_k3.template modify(); - k_phi1.template modify(); - k_phi2.template modify(); - k_phi3.template modify(); - k_mbt_f1.template modify(); - k_mbt_f2.template modify(); - k_mbt_f3.template modify(); - k_mbt_r0.template modify(); - k_ebt_f1_1.template modify(); - k_ebt_f2_1.template modify(); - k_ebt_f3_1.template modify(); - k_ebt_r0_1.template modify(); - k_ebt_f1_2.template modify(); - k_ebt_f2_2.template modify(); - k_ebt_f3_2.template modify(); - k_ebt_r0_2.template modify(); - k_at_f1_1.template modify(); - k_at_f2_1.template modify(); - k_at_f3_1.template modify(); - k_at_f1_2.template modify(); - k_at_f2_2.template modify(); - k_at_f3_2.template modify(); - k_at_theta0_1.template modify(); - k_at_theta0_2.template modify(); - k_aat_k.template modify(); - k_aat_theta0_1.template modify(); - k_aat_theta0_2.template modify(); - k_bb13t_k.template modify(); - k_bb13t_r10.template modify(); - k_bb13t_r30.template modify(); - k_setflag_d.template modify(); - k_setflag_mbt.template modify(); - k_setflag_ebt.template modify(); - k_setflag_at.template modify(); - k_setflag_aat.template modify(); - k_setflag_bb13t.template modify(); + k_k1.modify_host(); + k_k2.modify_host(); + k_k3.modify_host(); + k_phi1.modify_host(); + k_phi2.modify_host(); + k_phi3.modify_host(); + k_mbt_f1.modify_host(); + k_mbt_f2.modify_host(); + k_mbt_f3.modify_host(); + k_mbt_r0.modify_host(); + k_ebt_f1_1.modify_host(); + k_ebt_f2_1.modify_host(); + k_ebt_f3_1.modify_host(); + k_ebt_r0_1.modify_host(); + k_ebt_f1_2.modify_host(); + k_ebt_f2_2.modify_host(); + k_ebt_f3_2.modify_host(); + k_ebt_r0_2.modify_host(); + k_at_f1_1.modify_host(); + k_at_f2_1.modify_host(); + k_at_f3_1.modify_host(); + k_at_f1_2.modify_host(); + k_at_f2_2.modify_host(); + k_at_f3_2.modify_host(); + k_at_theta0_1.modify_host(); + k_at_theta0_2.modify_host(); + k_aat_k.modify_host(); + k_aat_theta0_1.modify_host(); + k_aat_theta0_2.modify_host(); + k_bb13t_k.modify_host(); + k_bb13t_r10.modify_host(); + k_bb13t_r30.modify_host(); + k_setflag_d.modify_host(); + k_setflag_mbt.modify_host(); + k_setflag_ebt.modify_host(); + k_setflag_at.modify_host(); + k_setflag_aat.modify_host(); + k_setflag_bb13t.modify_host(); } @@ -956,44 +957,44 @@ void DihedralClass2Kokkos::read_restart(FILE *fp) k_setflag_bb13t.h_view[i] = setflag_bb13t[i]; } - k_k1.template modify(); - k_k2.template modify(); - k_k3.template modify(); - k_phi1.template modify(); - k_phi2.template modify(); - k_phi3.template modify(); - k_mbt_f1.template modify(); - k_mbt_f2.template modify(); - k_mbt_f3.template modify(); - k_mbt_r0.template modify(); - k_ebt_f1_1.template modify(); - k_ebt_f2_1.template modify(); - k_ebt_f3_1.template modify(); - k_ebt_r0_1.template modify(); - k_ebt_f1_2.template modify(); - k_ebt_f2_2.template modify(); - k_ebt_f3_2.template modify(); - k_ebt_r0_2.template modify(); - k_at_f1_1.template modify(); - k_at_f2_1.template modify(); - k_at_f3_1.template modify(); - k_at_f1_2.template modify(); - k_at_f2_2.template modify(); - k_at_f3_2.template modify(); - k_at_theta0_1.template modify(); - k_at_theta0_2.template modify(); - k_aat_k.template modify(); - k_aat_theta0_1.template modify(); - k_aat_theta0_2.template modify(); - k_bb13t_k.template modify(); - k_bb13t_r10.template modify(); - k_bb13t_r30.template modify(); - k_setflag_d.template modify(); - k_setflag_mbt.template modify(); - k_setflag_ebt.template modify(); - k_setflag_at.template modify(); - k_setflag_aat.template modify(); - k_setflag_bb13t.template modify(); + k_k1.modify_host(); + k_k2.modify_host(); + k_k3.modify_host(); + k_phi1.modify_host(); + k_phi2.modify_host(); + k_phi3.modify_host(); + k_mbt_f1.modify_host(); + k_mbt_f2.modify_host(); + k_mbt_f3.modify_host(); + k_mbt_r0.modify_host(); + k_ebt_f1_1.modify_host(); + k_ebt_f2_1.modify_host(); + k_ebt_f3_1.modify_host(); + k_ebt_r0_1.modify_host(); + k_ebt_f1_2.modify_host(); + k_ebt_f2_2.modify_host(); + k_ebt_f3_2.modify_host(); + k_ebt_r0_2.modify_host(); + k_at_f1_1.modify_host(); + k_at_f2_1.modify_host(); + k_at_f3_1.modify_host(); + k_at_f1_2.modify_host(); + k_at_f2_2.modify_host(); + k_at_f3_2.modify_host(); + k_at_theta0_1.modify_host(); + k_at_theta0_2.modify_host(); + k_aat_k.modify_host(); + k_aat_theta0_1.modify_host(); + k_aat_theta0_2.modify_host(); + k_bb13t_k.modify_host(); + k_bb13t_r10.modify_host(); + k_bb13t_r30.modify_host(); + k_setflag_d.modify_host(); + k_setflag_mbt.modify_host(); + k_setflag_ebt.modify_host(); + k_setflag_at.modify_host(); + k_setflag_aat.modify_host(); + k_setflag_bb13t.modify_host(); } /* ---------------------------------------------------------------------- diff --git a/src/KOKKOS/dihedral_class2_kokkos.h b/src/KOKKOS/dihedral_class2_kokkos.h index 6e1dceb0cf..ddc79d21d9 100644 --- a/src/KOKKOS/dihedral_class2_kokkos.h +++ b/src/KOKKOS/dihedral_class2_kokkos.h @@ -60,6 +60,9 @@ class DihedralClass2Kokkos : public DihedralClass2 { const F_FLOAT &vb2x, const F_FLOAT &vb2y, const F_FLOAT &vb2z, const F_FLOAT &vb3x, const F_FLOAT &vb3y, const F_FLOAT &vb3z) const; + DAT::tdual_efloat_1d k_eatom; + DAT::tdual_virial_array k_vatom; + protected: class NeighborKokkos *neighborKK; @@ -67,9 +70,6 @@ class DihedralClass2Kokkos : public DihedralClass2 { typename AT::t_x_array_randomread x; typename AT::t_f_array f; typename AT::t_int_2d dihedrallist; - - DAT::tdual_efloat_1d k_eatom; - DAT::tdual_virial_array k_vatom; typename AT::t_efloat_1d d_eatom; typename AT::t_virial_array d_vatom; diff --git a/src/KOKKOS/dihedral_harmonic_kokkos.cpp b/src/KOKKOS/dihedral_harmonic_kokkos.cpp index 78860800be..05babd69b4 100644 --- a/src/KOKKOS/dihedral_harmonic_kokkos.cpp +++ b/src/KOKKOS/dihedral_harmonic_kokkos.cpp @@ -37,6 +37,7 @@ static constexpr double TOLERANCE = 0.05; template DihedralHarmonicKokkos::DihedralHarmonicKokkos(LAMMPS *lmp) : DihedralHarmonic(lmp) { + kokkosable = 1; atomKK = (AtomKokkos *) atom; neighborKK = (NeighborKokkos *) neighbor; execution_space = ExecutionSpaceFromDevice::space; @@ -74,14 +75,18 @@ void DihedralHarmonicKokkos::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { + if(k_eatom.extent(0) < maxeatom) { memoryKK->destroy_kokkos(k_eatom,eatom); memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"dihedral:eatom"); d_eatom = k_eatom.view(); + } else Kokkos::deep_copy(d_eatom,0.0); } if (vflag_atom) { + if(k_vatom.extent(0) < maxvatom) { memoryKK->destroy_kokkos(k_vatom,vatom); memoryKK->create_kokkos(k_vatom,vatom,maxvatom,"dihedral:vatom"); d_vatom = k_vatom.view(); + } else Kokkos::deep_copy(d_vatom,0.0); } k_k.template sync(); @@ -99,7 +104,7 @@ void DihedralHarmonicKokkos::compute(int eflag_in, int vflag_in) newton_bond = force->newton_bond; h_warning_flag() = 0; - k_warning_flag.template modify(); + k_warning_flag.modify_host(); k_warning_flag.template sync(); copymode = 1; @@ -125,7 +130,7 @@ void DihedralHarmonicKokkos::compute(int eflag_in, int vflag_in) // error check k_warning_flag.template modify(); - k_warning_flag.template sync(); + k_warning_flag.sync_host(); if (h_warning_flag()) error->warning(FLERR,"Dihedral problem"); @@ -141,12 +146,12 @@ void DihedralHarmonicKokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { k_eatom.template modify(); - k_eatom.template sync(); + k_eatom.sync_host(); } if (vflag_atom) { k_vatom.template modify(); - k_vatom.template sync(); + k_vatom.sync_host(); } copymode = 0; @@ -362,11 +367,11 @@ void DihedralHarmonicKokkos::coeff(int narg, char **arg) k_multiplicity.h_view[i] = multiplicity[i]; } - k_k.template modify(); - k_cos_shift.template modify(); - k_sin_shift.template modify(); - k_sign.template modify(); - k_multiplicity.template modify(); + k_k.modify_host(); + k_cos_shift.modify_host(); + k_sin_shift.modify_host(); + k_sign.modify_host(); + k_multiplicity.modify_host(); } /* ---------------------------------------------------------------------- @@ -387,11 +392,11 @@ void DihedralHarmonicKokkos::read_restart(FILE *fp) k_multiplicity.h_view[i] = multiplicity[i]; } - k_k.template modify(); - k_cos_shift.template modify(); - k_sin_shift.template modify(); - k_sign.template modify(); - k_multiplicity.template modify(); + k_k.modify_host(); + k_cos_shift.modify_host(); + k_sin_shift.modify_host(); + k_sign.modify_host(); + k_multiplicity.modify_host(); } /* ---------------------------------------------------------------------- diff --git a/src/KOKKOS/dihedral_harmonic_kokkos.h b/src/KOKKOS/dihedral_harmonic_kokkos.h index e73f19afd1..1ad62398f7 100644 --- a/src/KOKKOS/dihedral_harmonic_kokkos.h +++ b/src/KOKKOS/dihedral_harmonic_kokkos.h @@ -60,6 +60,9 @@ class DihedralHarmonicKokkos : public DihedralHarmonic { const F_FLOAT &vb2x, const F_FLOAT &vb2y, const F_FLOAT &vb2z, const F_FLOAT &vb3x, const F_FLOAT &vb3y, const F_FLOAT &vb3z) const; + DAT::tdual_efloat_1d k_eatom; + DAT::tdual_virial_array k_vatom; + protected: class NeighborKokkos *neighborKK; @@ -67,9 +70,6 @@ class DihedralHarmonicKokkos : public DihedralHarmonic { typename AT::t_x_array_randomread x; typename AT::t_f_array f; typename AT::t_int_2d dihedrallist; - - DAT::tdual_efloat_1d k_eatom; - DAT::tdual_virial_array k_vatom; typename ArrayTypes::t_efloat_1d d_eatom; typename ArrayTypes::t_virial_array d_vatom; diff --git a/src/KOKKOS/dihedral_hybrid_kokkos.cpp b/src/KOKKOS/dihedral_hybrid_kokkos.cpp new file mode 100644 index 0000000000..88dbeaf13b --- /dev/null +++ b/src/KOKKOS/dihedral_hybrid_kokkos.cpp @@ -0,0 +1,225 @@ +/* ---------------------------------------------------------------------- + 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 "dihedral_hybrid_kokkos.h" + +#include "atom_kokkos.h" +#include "atom_masks.h" +#include "comm.h" +#include "error.h" +#include "force.h" +#include "kokkos.h" +#include "memory_kokkos.h" +#include "neighbor_kokkos.h" + +#include + +using namespace LAMMPS_NS; + +#define EXTRA 1000 + +/* ---------------------------------------------------------------------- */ + +DihedralHybridKokkos::DihedralHybridKokkos(LAMMPS *lmp) : DihedralHybrid(lmp) +{ + kokkosable = 1; + + atomKK = (AtomKokkos *) atom; + neighborKK = (NeighborKokkos *) neighbor; + + execution_space = Device; + + datamask_read = EMPTY_MASK; + datamask_modify = EMPTY_MASK; +} + +/* ---------------------------------------------------------------------- */ + +DihedralHybridKokkos::~DihedralHybridKokkos() +{ + deallocate(); +} + +/* ---------------------------------------------------------------------- */ + +void DihedralHybridKokkos::compute(int eflag, int vflag) +{ + // save ptrs to original dihedrallist + + int ndihedrallist_orig = neighbor->ndihedrallist; + neighborKK->k_dihedrallist.sync_device(); + auto k_dihedrallist_orig = neighborKK->k_dihedrallist; + auto d_dihedrallist_orig = k_dihedrallist_orig.d_view; + auto d_ndihedrallist = k_ndihedrallist.d_view; + auto h_ndihedrallist = k_ndihedrallist.h_view; + + // if this is re-neighbor step, create sub-style dihedrallists + // ndihedrallist[] = length of each sub-style list + // realloc sub-style dihedrallist if necessary + // load sub-style dihedrallist with 3 values from original dihedrallist + + if (neighbor->ago == 0) { + Kokkos::deep_copy(d_ndihedrallist,0); + + k_map.sync_device(); + auto d_map = k_map.d_view; + + Kokkos::parallel_for(ndihedrallist_orig,LAMMPS_LAMBDA(int i) { + const int m = d_map[d_dihedrallist_orig(i,4)]; + if (m >= 0) Kokkos::atomic_increment(&d_ndihedrallist[m]); + }); + + k_ndihedrallist.modify_device(); + k_ndihedrallist.sync_host(); + + maxdihedral_all = 0; + for (int m = 0; m < nstyles; m++) + if (h_ndihedrallist[m] > maxdihedral_all) + maxdihedral_all = h_ndihedrallist[m] + EXTRA; + + if (k_dihedrallist.d_view.extent(1) < maxdihedral_all) + MemKK::realloc_kokkos(k_dihedrallist, "dihedral_hybrid:dihedrallist", nstyles, maxdihedral_all, 5); + auto d_dihedrallist = k_dihedrallist.d_view; + + Kokkos::deep_copy(d_ndihedrallist,0); + + Kokkos::parallel_for(ndihedrallist_orig,LAMMPS_LAMBDA(int i) { + const int m = d_map[d_dihedrallist_orig(i,4)]; + if (m < 0) return; + const int n = Kokkos::atomic_fetch_add(&d_ndihedrallist[m],1); + d_dihedrallist(m,n,0) = d_dihedrallist_orig(i,0); + d_dihedrallist(m,n,1) = d_dihedrallist_orig(i,1); + d_dihedrallist(m,n,2) = d_dihedrallist_orig(i,2); + d_dihedrallist(m,n,3) = d_dihedrallist_orig(i,3); + d_dihedrallist(m,n,4) = d_dihedrallist_orig(i,4); + }); + } + + // call each sub-style's compute function + // set neighbor->dihedrallist to sub-style dihedrallist before call + // accumulate sub-style global/peratom energy/virial in hybrid + + ev_init(eflag, vflag); + + k_ndihedrallist.modify_device(); + k_ndihedrallist.sync_host(); + + for (int m = 0; m < nstyles; m++) { + neighbor->ndihedrallist = h_ndihedrallist[m]; + auto k_dihedrallist_m = Kokkos::subview(k_dihedrallist,m,Kokkos::ALL,Kokkos::ALL); + k_dihedrallist_m.modify_device(); + neighborKK->k_dihedrallist = k_dihedrallist_m; + + auto style = styles[m]; + atomKK->sync(style->execution_space,style->datamask_read); + style->compute(eflag, vflag); + atomKK->modified(style->execution_space,style->datamask_modify); + + if (eflag_global) energy += style->energy; + if (vflag_global) + for (int n = 0; n < 6; n++) virial[n] += style->virial[n]; + + if (eflag_atom) { + int n = atom->nlocal; + if (force->newton_bond) n += atom->nghost; + double *eatom_substyle = styles[m]->eatom; + for (int i = 0; i < n; i++) eatom[i] += eatom_substyle[i]; + } + if (vflag_atom) { + int n = atom->nlocal; + if (force->newton_bond) n += atom->nghost; + double **vatom_substyle = styles[m]->vatom; + for (int i = 0; i < n; i++) + for (int j = 0; j < 6; j++) vatom[i][j] += vatom_substyle[i][j]; + } + if (cvflag_atom) { + int n = atom->nlocal; + if (force->newton_bond) n += atom->nghost; + double **cvatom_substyle = styles[m]->cvatom; + for (int i = 0; i < n; i++) + for (int j = 0; j < 9; j++) cvatom[i][j] += cvatom_substyle[i][j]; + } + } + + // restore ptrs to original dihedrallist + + neighbor->ndihedrallist = ndihedrallist_orig; + neighborKK->k_dihedrallist = k_dihedrallist_orig; +} + +/* ---------------------------------------------------------------------- */ + +void DihedralHybridKokkos::allocate() +{ + allocated = 1; + int np1 = atom->ndihedraltypes + 1; + + memoryKK->create_kokkos(k_map, map, np1, "dihedral:map"); + memory->create(setflag, np1, "dihedral:setflag"); + for (int i = 1; i < np1; i++) setflag[i] = 0; + + k_ndihedrallist = DAT::tdual_int_1d("dihedral:ndihedrallist", nstyles); +} + +/* ---------------------------------------------------------------------- */ + +void DihedralHybridKokkos::deallocate() +{ + if (!allocated) return; + + allocated = 0; + + memory->destroy(setflag); + memoryKK->destroy_kokkos(k_map,map); +} + +/* ---------------------------------------------------------------------- + set coeffs for one type +---------------------------------------------------------------------- */ + +void DihedralHybridKokkos::coeff(int narg, char **arg) +{ + DihedralHybrid::coeff(narg,arg); + + k_map.modify_host(); +} + +/* ---------------------------------------------------------------------- */ + +void DihedralHybridKokkos::init_style() +{ + DihedralHybrid::init_style(); + + for (int m = 0; m < nstyles; m++) { + if (!styles[m]->kokkosable) + error->all(FLERR,"Must use only Kokkos-enabled dihedral styles with dihedral_style hybrid/kk"); + + if (styles[m]->execution_space == Host) + lmp->kokkos->allow_overlap = 0; + } +} + +/* ---------------------------------------------------------------------- + memory usage +------------------------------------------------------------------------- */ + +double DihedralHybridKokkos::memory_usage() +{ + double bytes = (double) maxeatom * sizeof(double); + bytes += (double) maxvatom * 6 * sizeof(double); + bytes += (double) maxcvatom * 9 * sizeof(double); + for (int m = 0; m < nstyles; m++) bytes += (double) maxdihedral_all * 5 * sizeof(int); + for (int m = 0; m < nstyles; m++) + if (styles[m]) bytes += styles[m]->memory_usage(); + return bytes; +} diff --git a/src/KOKKOS/dihedral_hybrid_kokkos.h b/src/KOKKOS/dihedral_hybrid_kokkos.h new file mode 100644 index 0000000000..29a3d29689 --- /dev/null +++ b/src/KOKKOS/dihedral_hybrid_kokkos.h @@ -0,0 +1,58 @@ +/* -*- 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 BOND_CLASS +// clang-format off +DihedralStyle(hybrid/kk,DihedralHybridKokkos); +DihedralStyle(hybrid/kk/device,DihedralHybridKokkos); +DihedralStyle(hybrid/kk/host,DihedralHybridKokkos); +// clang-format on +#else + +// clang-format off +#ifndef LMP_BOND_HYBRID_KOKKOS_H +#define LMP_BOND_HYBRID_KOKKOS_H + +#include "dihedral_hybrid.h" +#include "kokkos_type.h" + +namespace LAMMPS_NS { + +class DihedralHybridKokkos : public DihedralHybrid { + friend class Force; + + public: + DihedralHybridKokkos(class LAMMPS *); + ~DihedralHybridKokkos() override; + void compute(int, int) override; + void coeff(int, char **) override; + void init_style() override; + double memory_usage() override; + + private: + int maxbond_all; + + class NeighborKokkos *neighborKK; + + DAT::tdual_int_1d k_map; // which style each bond type points to + DAT::tdual_int_1d k_nbondlist; // # of bonds in sub-style bondlists + DAT::tdual_int_3d k_bondlist; // bondlist for each sub-style + + void allocate() override; + void deallocate() override; +}; + +} // namespace LAMMPS_NS + +#endif +#endif diff --git a/src/KOKKOS/dihedral_opls_kokkos.cpp b/src/KOKKOS/dihedral_opls_kokkos.cpp index ce7502b25a..b45437b781 100644 --- a/src/KOKKOS/dihedral_opls_kokkos.cpp +++ b/src/KOKKOS/dihedral_opls_kokkos.cpp @@ -39,6 +39,7 @@ static constexpr double SMALLER = 0.00001; template DihedralOPLSKokkos::DihedralOPLSKokkos(LAMMPS *lmp) : DihedralOPLS(lmp) { + kokkosable = 1; atomKK = (AtomKokkos *) atom; neighborKK = (NeighborKokkos *) neighbor; execution_space = ExecutionSpaceFromDevice::space; @@ -100,7 +101,7 @@ void DihedralOPLSKokkos::compute(int eflag_in, int vflag_in) newton_bond = force->newton_bond; h_warning_flag() = 0; - k_warning_flag.template modify(); + k_warning_flag.modify_host(); k_warning_flag.template sync(); copymode = 1; @@ -126,7 +127,7 @@ void DihedralOPLSKokkos::compute(int eflag_in, int vflag_in) // error check k_warning_flag.template modify(); - k_warning_flag.template sync(); + k_warning_flag.sync_host(); if (h_warning_flag()) error->warning(FLERR,"Dihedral problem"); @@ -142,12 +143,12 @@ void DihedralOPLSKokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { k_eatom.template modify(); - k_eatom.template sync(); + k_eatom.sync_host(); } if (vflag_atom) { k_vatom.template modify(); - k_vatom.template sync(); + k_vatom.sync_host(); } copymode = 0; @@ -372,10 +373,10 @@ void DihedralOPLSKokkos::coeff(int narg, char **arg) k_k4.h_view[i] = k4[i]; } - k_k1.template modify(); - k_k2.template modify(); - k_k3.template modify(); - k_k4.template modify(); + k_k1.modify_host(); + k_k2.modify_host(); + k_k3.modify_host(); + k_k4.modify_host(); } /* ---------------------------------------------------------------------- @@ -395,10 +396,10 @@ void DihedralOPLSKokkos::read_restart(FILE *fp) k_k4.h_view[i] = k4[i]; } - k_k1.template modify(); - k_k2.template modify(); - k_k3.template modify(); - k_k4.template modify(); + k_k1.modify_host(); + k_k2.modify_host(); + k_k3.modify_host(); + k_k4.modify_host(); } /* ---------------------------------------------------------------------- diff --git a/src/KOKKOS/dihedral_opls_kokkos.h b/src/KOKKOS/dihedral_opls_kokkos.h index 886f4abcae..20c5083ad9 100644 --- a/src/KOKKOS/dihedral_opls_kokkos.h +++ b/src/KOKKOS/dihedral_opls_kokkos.h @@ -60,16 +60,15 @@ class DihedralOPLSKokkos : public DihedralOPLS { const F_FLOAT &vb2x, const F_FLOAT &vb2y, const F_FLOAT &vb2z, const F_FLOAT &vb3x, const F_FLOAT &vb3y, const F_FLOAT &vb3z) const; + DAT::tdual_efloat_1d k_eatom; + DAT::tdual_virial_array k_vatom; + protected: class NeighborKokkos *neighborKK; - typename AT::t_x_array_randomread x; typename AT::t_f_array f; typename AT::t_int_2d dihedrallist; - - DAT::tdual_efloat_1d k_eatom; - DAT::tdual_virial_array k_vatom; typename ArrayTypes::t_efloat_1d d_eatom; typename ArrayTypes::t_virial_array d_vatom; diff --git a/src/KOKKOS/improper_class2_kokkos.cpp b/src/KOKKOS/improper_class2_kokkos.cpp index 862ba2a52f..1aee9d2144 100644 --- a/src/KOKKOS/improper_class2_kokkos.cpp +++ b/src/KOKKOS/improper_class2_kokkos.cpp @@ -34,6 +34,7 @@ static constexpr double SMALL = 0.001; template ImproperClass2Kokkos::ImproperClass2Kokkos(LAMMPS *lmp) : ImproperClass2(lmp) { + kokkosable = 1; atomKK = (AtomKokkos *) atom; neighborKK = (NeighborKokkos *) neighbor; execution_space = ExecutionSpaceFromDevice::space; @@ -110,7 +111,7 @@ void ImproperClass2Kokkos::compute(int eflag_in, int vflag_in) newton_bond = force->newton_bond; h_warning_flag() = 0; - k_warning_flag.template modify(); + k_warning_flag.modify_host(); k_warning_flag.template sync(); copymode = 1; @@ -139,7 +140,7 @@ void ImproperClass2Kokkos::compute(int eflag_in, int vflag_in) // error check k_warning_flag.template modify(); - k_warning_flag.template sync(); + k_warning_flag.sync_host(); if (h_warning_flag()) error->warning(FLERR,"Improper problem"); @@ -171,12 +172,12 @@ void ImproperClass2Kokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { k_eatom.template modify(); - k_eatom.template sync(); + k_eatom.sync_host(); } if (vflag_atom) { k_vatom.template modify(); - k_vatom.template sync(); + k_vatom.sync_host(); } copymode = 0; @@ -918,17 +919,17 @@ void ImproperClass2Kokkos::coeff(int narg, char **arg) k_setflag_aa.h_view[i] = setflag_aa[i]; } - k_k0.template modify(); - k_chi0.template modify(); - k_aa_k1.template modify(); - k_aa_k2.template modify(); - k_aa_k3.template modify(); - k_aa_theta0_1.template modify(); - k_aa_theta0_2.template modify(); - k_aa_theta0_3 .template modify(); - k_setflag.template modify(); - k_setflag_i.template modify(); - k_setflag_aa.template modify(); + k_k0.modify_host(); + k_chi0.modify_host(); + k_aa_k1.modify_host(); + k_aa_k2.modify_host(); + k_aa_k3.modify_host(); + k_aa_theta0_1.modify_host(); + k_aa_theta0_2.modify_host(); + k_aa_theta0_3 .modify_host(); + k_setflag.modify_host(); + k_setflag_i.modify_host(); + k_setflag_aa.modify_host(); } /* ---------------------------------------------------------------------- @@ -979,17 +980,17 @@ void ImproperClass2Kokkos::read_restart(FILE *fp) k_setflag_aa.h_view[i] = setflag_aa[i]; } - k_k0.template modify(); - k_chi0.template modify(); - k_aa_k1.template modify(); - k_aa_k2.template modify(); - k_aa_k3.template modify(); - k_aa_theta0_1.template modify(); - k_aa_theta0_2.template modify(); - k_aa_theta0_3 .template modify(); - k_setflag.template modify(); - k_setflag_i.template modify(); - k_setflag_aa.template modify(); + k_k0.modify_host(); + k_chi0.modify_host(); + k_aa_k1.modify_host(); + k_aa_k2.modify_host(); + k_aa_k3.modify_host(); + k_aa_theta0_1.modify_host(); + k_aa_theta0_2.modify_host(); + k_aa_theta0_3 .modify_host(); + k_setflag.modify_host(); + k_setflag_i.modify_host(); + k_setflag_aa.modify_host(); } /* ---------------------------------------------------------------------- diff --git a/src/KOKKOS/improper_class2_kokkos.h b/src/KOKKOS/improper_class2_kokkos.h index da939b69b0..7a55c8a5f8 100644 --- a/src/KOKKOS/improper_class2_kokkos.h +++ b/src/KOKKOS/improper_class2_kokkos.h @@ -71,6 +71,9 @@ class ImproperClass2Kokkos : public ImproperClass2 { const F_FLOAT &vb2x, const F_FLOAT &vb2y, const F_FLOAT &vb2z, const F_FLOAT &vb3x, const F_FLOAT &vb3y, const F_FLOAT &vb3z) const; + DAT::tdual_efloat_1d k_eatom; + DAT::tdual_virial_array k_vatom; + protected: class NeighborKokkos *neighborKK; @@ -78,9 +81,6 @@ class ImproperClass2Kokkos : public ImproperClass2 { typename AT::t_x_array_randomread x; typename Kokkos::View::value,Kokkos::MemoryTraits > f; typename AT::t_int_2d improperlist; - - DAT::tdual_efloat_1d k_eatom; - DAT::tdual_virial_array k_vatom; typename AT::t_efloat_1d d_eatom; typename AT::t_virial_array d_vatom; diff --git a/src/KOKKOS/improper_harmonic_kokkos.cpp b/src/KOKKOS/improper_harmonic_kokkos.cpp index a075238f22..eafa7a08ec 100644 --- a/src/KOKKOS/improper_harmonic_kokkos.cpp +++ b/src/KOKKOS/improper_harmonic_kokkos.cpp @@ -36,6 +36,7 @@ static constexpr double SMALL = 0.001; template ImproperHarmonicKokkos::ImproperHarmonicKokkos(LAMMPS *lmp) : ImproperHarmonic(lmp) { + kokkosable = 1; atomKK = (AtomKokkos *) atom; neighborKK = (NeighborKokkos *) neighbor; execution_space = ExecutionSpaceFromDevice::space; @@ -73,18 +74,18 @@ void ImproperHarmonicKokkos::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - //if(k_eatom.extent(0)destroy_kokkos(k_eatom,eatom); memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); d_eatom = k_eatom.template view(); - //} + } else Kokkos::deep_copy(d_eatom,0.0); } if (vflag_atom) { - //if(k_vatom.extent(0)destroy_kokkos(k_vatom,vatom); memoryKK->create_kokkos(k_vatom,vatom,maxvatom,"improper:vatom"); d_vatom = k_vatom.template view(); - //} + } else Kokkos::deep_copy(d_vatom,0.0); } //atomKK->sync(execution_space,datamask_read); @@ -102,7 +103,7 @@ void ImproperHarmonicKokkos::compute(int eflag_in, int vflag_in) newton_bond = force->newton_bond; h_warning_flag() = 0; - k_warning_flag.template modify(); + k_warning_flag.modify_host(); k_warning_flag.template sync(); copymode = 1; @@ -128,7 +129,7 @@ void ImproperHarmonicKokkos::compute(int eflag_in, int vflag_in) // error check k_warning_flag.template modify(); - k_warning_flag.template sync(); + k_warning_flag.sync_host(); if (h_warning_flag()) error->warning(FLERR,"Dihedral problem"); @@ -144,12 +145,12 @@ void ImproperHarmonicKokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { k_eatom.template modify(); - k_eatom.template sync(); + k_eatom.sync_host(); } if (vflag_atom) { k_vatom.template modify(); - k_vatom.template sync(); + k_vatom.sync_host(); } copymode = 0; @@ -324,8 +325,8 @@ void ImproperHarmonicKokkos::coeff(int narg, char **arg) k_chi.h_view[i] = chi[i]; } - k_k.template modify(); - k_chi.template modify(); + k_k.modify_host(); + k_chi.modify_host(); } /* ---------------------------------------------------------------------- @@ -343,8 +344,8 @@ void ImproperHarmonicKokkos::read_restart(FILE *fp) k_chi.h_view[i] = chi[i]; } - k_k.template modify(); - k_chi.template modify(); + k_k.modify_host(); + k_chi.modify_host(); } /* ---------------------------------------------------------------------- diff --git a/src/KOKKOS/improper_harmonic_kokkos.h b/src/KOKKOS/improper_harmonic_kokkos.h index 8bd206aaf0..ad31447383 100644 --- a/src/KOKKOS/improper_harmonic_kokkos.h +++ b/src/KOKKOS/improper_harmonic_kokkos.h @@ -60,17 +60,17 @@ class ImproperHarmonicKokkos : public ImproperHarmonic { const F_FLOAT &vb2x, const F_FLOAT &vb2y, const F_FLOAT &vb2z, const F_FLOAT &vb3x, const F_FLOAT &vb3y, const F_FLOAT &vb3z) const; + typedef typename KKDevice::value KKDeviceType; + Kokkos::DualView k_eatom; + Kokkos::DualView k_vatom; + protected: class NeighborKokkos *neighborKK; - typedef typename KKDevice::value KKDeviceType; typename AT::t_x_array_randomread x; typename Kokkos::View > f; typename AT::t_int_2d improperlist; - - Kokkos::DualView k_eatom; - Kokkos::DualView k_vatom; Kokkos::View > d_eatom; Kokkos::View > d_vatom; diff --git a/src/KOKKOS/improper_hybrid_kokkos.cpp b/src/KOKKOS/improper_hybrid_kokkos.cpp new file mode 100644 index 0000000000..bfa55978cc --- /dev/null +++ b/src/KOKKOS/improper_hybrid_kokkos.cpp @@ -0,0 +1,226 @@ +/* ---------------------------------------------------------------------- + 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 "improper_hybrid_kokkos.h" + +#include "atom_kokkos.h" +#include "atom_masks.h" +#include "comm.h" +#include "error.h" +#include "force.h" +#include "kokkos.h" +#include "memory_kokkos.h" +#include "neighbor_kokkos.h" + +#include + +using namespace LAMMPS_NS; + +#define EXTRA 1000 + +/* ---------------------------------------------------------------------- */ + +ImproperHybridKokkos::ImproperHybridKokkos(LAMMPS *lmp) : ImproperHybrid(lmp) +{ + kokkosable = 1; + + atomKK = (AtomKokkos *) atom; + neighborKK = (NeighborKokkos *) neighbor; + + execution_space = Device; + + datamask_read = EMPTY_MASK; + datamask_modify = EMPTY_MASK; +} + +/* ---------------------------------------------------------------------- */ + +ImproperHybridKokkos::~ImproperHybridKokkos() +{ + deallocate(); +} + +/* ---------------------------------------------------------------------- */ + +void ImproperHybridKokkos::compute(int eflag, int vflag) +{ + + // save ptrs to original improperlist + + int nimproperlist_orig = neighbor->nimproperlist; + neighborKK->k_improperlist.sync_device(); + auto k_improperlist_orig = neighborKK->k_improperlist; + auto d_improperlist_orig = k_improperlist_orig.d_view; + auto d_nimproperlist = k_nimproperlist.d_view; + auto h_nimproperlist = k_nimproperlist.h_view; + + // if this is re-neighbor step, create sub-style improperlists + // nimproperlist[] = length of each sub-style list + // realloc sub-style improperlist if necessary + // load sub-style improperlist with 3 values from original improperlist + + if (neighbor->ago == 0) { + Kokkos::deep_copy(d_nimproperlist,0); + + k_map.sync_device(); + auto d_map = k_map.d_view; + + Kokkos::parallel_for(nimproperlist_orig,LAMMPS_LAMBDA(int i) { + const int m = d_map[d_improperlist_orig(i,4)]; + if (m >= 0) Kokkos::atomic_increment(&d_nimproperlist[m]); + }); + + k_nimproperlist.modify_device(); + k_nimproperlist.sync_host(); + + maximproper_all = 0; + for (int m = 0; m < nstyles; m++) + if (h_nimproperlist[m] > maximproper_all) + maximproper_all = h_nimproperlist[m] + EXTRA; + + if (k_improperlist.d_view.extent(1) < maximproper_all) + MemKK::realloc_kokkos(k_improperlist, "improper_hybrid:improperlist", nstyles, maximproper_all, 5); + auto d_improperlist = k_improperlist.d_view; + + Kokkos::deep_copy(d_nimproperlist,0); + + Kokkos::parallel_for(nimproperlist_orig,LAMMPS_LAMBDA(int i) { + const int m = d_map[d_improperlist_orig(i,4)]; + if (m < 0) return; + const int n = Kokkos::atomic_fetch_add(&d_nimproperlist[m],1); + d_improperlist(m,n,0) = d_improperlist_orig(i,0); + d_improperlist(m,n,1) = d_improperlist_orig(i,1); + d_improperlist(m,n,2) = d_improperlist_orig(i,2); + d_improperlist(m,n,3) = d_improperlist_orig(i,3); + d_improperlist(m,n,4) = d_improperlist_orig(i,4); + }); + } + + // call each sub-style's compute function + // set neighbor->improperlist to sub-style improperlist before call + // accumulate sub-style global/peratom energy/virial in hybrid + + ev_init(eflag, vflag); + + k_nimproperlist.modify_device(); + k_nimproperlist.sync_host(); + + for (int m = 0; m < nstyles; m++) { + neighbor->nimproperlist = h_nimproperlist[m]; + auto k_improperlist_m = Kokkos::subview(k_improperlist,m,Kokkos::ALL,Kokkos::ALL); + k_improperlist_m.modify_device(); + neighborKK->k_improperlist = k_improperlist_m; + + auto style = styles[m]; + atomKK->sync(style->execution_space,style->datamask_read); + style->compute(eflag, vflag); + atomKK->modified(style->execution_space,style->datamask_modify); + + if (eflag_global) energy += style->energy; + if (vflag_global) + for (int n = 0; n < 6; n++) virial[n] += style->virial[n]; + + if (eflag_atom) { + int n = atom->nlocal; + if (force->newton_bond) n += atom->nghost; + double *eatom_substyle = styles[m]->eatom; + for (int i = 0; i < n; i++) eatom[i] += eatom_substyle[i]; + } + if (vflag_atom) { + int n = atom->nlocal; + if (force->newton_bond) n += atom->nghost; + double **vatom_substyle = styles[m]->vatom; + for (int i = 0; i < n; i++) + for (int j = 0; j < 6; j++) vatom[i][j] += vatom_substyle[i][j]; + } + if (cvflag_atom) { + int n = atom->nlocal; + if (force->newton_bond) n += atom->nghost; + double **cvatom_substyle = styles[m]->cvatom; + for (int i = 0; i < n; i++) + for (int j = 0; j < 9; j++) cvatom[i][j] += cvatom_substyle[i][j]; + } + } + + // restore ptrs to original improperlist + + neighbor->nimproperlist = nimproperlist_orig; + neighborKK->k_improperlist = k_improperlist_orig; +} + +/* ---------------------------------------------------------------------- */ + +void ImproperHybridKokkos::allocate() +{ + allocated = 1; + int np1 = atom->nimpropertypes + 1; + + memoryKK->create_kokkos(k_map, map, np1, "improper:map"); + memory->create(setflag, np1, "improper:setflag"); + for (int i = 1; i < np1; i++) setflag[i] = 0; + + k_nimproperlist = DAT::tdual_int_1d("improper:nimproperlist", nstyles); +} + +/* ---------------------------------------------------------------------- */ + +void ImproperHybridKokkos::deallocate() +{ + if (!allocated) return; + + allocated = 0; + + memory->destroy(setflag); + memoryKK->destroy_kokkos(k_map,map); +} + +/* ---------------------------------------------------------------------- + set coeffs for one type +---------------------------------------------------------------------- */ + +void ImproperHybridKokkos::coeff(int narg, char **arg) +{ + ImproperHybrid::coeff(narg,arg); + + k_map.modify_host(); +} + +/* ---------------------------------------------------------------------- */ + +void ImproperHybridKokkos::init_style() +{ + ImproperHybrid::init_style(); + + for (int m = 0; m < nstyles; m++) { + if (!styles[m]->kokkosable) + error->all(FLERR,"Must use only Kokkos-enabled improper styles with improper_style hybrid/kk"); + + if (styles[m]->execution_space == Host) + lmp->kokkos->allow_overlap = 0; + } +} + +/* ---------------------------------------------------------------------- + memory usage +------------------------------------------------------------------------- */ + +double ImproperHybridKokkos::memory_usage() +{ + double bytes = (double) maxeatom * sizeof(double); + bytes += (double) maxvatom * 6 * sizeof(double); + bytes += (double) maxcvatom * 9 * sizeof(double); + for (int m = 0; m < nstyles; m++) bytes += (double) maximproper_all * 5 * sizeof(int); + for (int m = 0; m < nstyles; m++) + if (styles[m]) bytes += styles[m]->memory_usage(); + return bytes; +} diff --git a/src/KOKKOS/improper_hybrid_kokkos.h b/src/KOKKOS/improper_hybrid_kokkos.h new file mode 100644 index 0000000000..f2a80f6a0c --- /dev/null +++ b/src/KOKKOS/improper_hybrid_kokkos.h @@ -0,0 +1,58 @@ +/* -*- 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 BOND_CLASS +// clang-format off +ImproperStyle(hybrid/kk,ImproperHybridKokkos); +ImproperStyle(hybrid/kk/device,ImproperHybridKokkos); +ImproperStyle(hybrid/kk/host,ImproperHybridKokkos); +// clang-format on +#else + +// clang-format off +#ifndef LMP_BOND_HYBRID_KOKKOS_H +#define LMP_BOND_HYBRID_KOKKOS_H + +#include "improper_hybrid.h" +#include "kokkos_type.h" + +namespace LAMMPS_NS { + +class ImproperHybridKokkos : public ImproperHybrid { + friend class Force; + + public: + ImproperHybridKokkos(class LAMMPS *); + ~ImproperHybridKokkos() override; + void compute(int, int) override; + void coeff(int, char **) override; + void init_style() override; + double memory_usage() override; + + private: + int maximproper_all; + + class NeighborKokkos *neighborKK; + + DAT::tdual_int_1d k_map; // which style each improper type points to + DAT::tdual_int_1d k_nimproperlist; // # of impropers in sub-style improperlists + DAT::tdual_int_3d k_improperlist; // improperlist for each sub-style + + void allocate() override; + void deallocate() override; +}; + +} // namespace LAMMPS_NS + +#endif +#endif diff --git a/src/angle.cpp b/src/angle.cpp index 79893cc52f..ccb53dc84f 100644 --- a/src/angle.cpp +++ b/src/angle.cpp @@ -50,7 +50,7 @@ Angle::Angle(LAMMPS *_lmp) : Pointers(_lmp) datamask_read = ALL_MASK; datamask_modify = ALL_MASK; - copymode = 0; + copymode = kokkosable = 0; } /* ---------------------------------------------------------------------- */ diff --git a/src/angle.h b/src/angle.h index 542bad4911..759f1a1aa9 100644 --- a/src/angle.h +++ b/src/angle.h @@ -44,7 +44,7 @@ class Angle : protected Pointers { ExecutionSpace execution_space; unsigned int datamask_read, datamask_modify; - int copymode; + int copymode, kokkosable; Angle(class LAMMPS *); ~Angle() override; diff --git a/src/angle_hybrid.cpp b/src/angle_hybrid.cpp index 0c61970a1f..a015882a15 100644 --- a/src/angle_hybrid.cpp +++ b/src/angle_hybrid.cpp @@ -48,14 +48,7 @@ AngleHybrid::~AngleHybrid() delete[] keywords; } - if (allocated) { - memory->destroy(setflag); - memory->destroy(map); - delete[] nanglelist; - delete[] maxangle; - for (int i = 0; i < nstyles; i++) memory->destroy(anglelist[i]); - delete[] anglelist; - } + deallocate(); } /* ---------------------------------------------------------------------- */ @@ -171,6 +164,22 @@ void AngleHybrid::allocate() for (int m = 0; m < nstyles; m++) anglelist[m] = nullptr; } +/* ---------------------------------------------------------------------- */ + +void AngleHybrid::deallocate() +{ + if (!allocated) return; + + allocated = 0; + + memory->destroy(setflag); + memory->destroy(map); + delete[] nanglelist; + delete[] maxangle; + for (int i = 0; i < nstyles; i++) memory->destroy(anglelist[i]); + delete[] anglelist; +} + /* ---------------------------------------------------------------------- create one angle style for each arg in list ------------------------------------------------------------------------- */ @@ -190,15 +199,7 @@ void AngleHybrid::settings(int narg, char **arg) delete[] keywords; } - if (allocated) { - memory->destroy(setflag); - memory->destroy(map); - delete[] nanglelist; - delete[] maxangle; - for (i = 0; i < nstyles; i++) memory->destroy(anglelist[i]); - delete[] anglelist; - } - allocated = 0; + deallocate(); // allocate list of sub-styles @@ -367,7 +368,7 @@ void AngleHybrid::read_restart(FILE *fp) keywords[m] = new char[n]; if (me == 0) utils::sfread(FLERR, keywords[m], sizeof(char), n, fp, nullptr, error); MPI_Bcast(keywords[m], n, MPI_CHAR, 0, world); - styles[m] = force->new_angle(keywords[m], 0, dummy); + styles[m] = force->new_angle(keywords[m], 1, dummy); styles[m]->read_restart_settings(fp); } } diff --git a/src/angle_hybrid.h b/src/angle_hybrid.h index 474ce89673..a6da29245e 100644 --- a/src/angle_hybrid.h +++ b/src/angle_hybrid.h @@ -42,14 +42,14 @@ class AngleHybrid : public Angle { double single(int, int, int, int) override; double memory_usage() override; - private: + protected: int *map; // which style each angle type points to - int *nanglelist; // # of angles in sub-style anglelists int *maxangle; // max # of angles sub-style lists can store int ***anglelist; // anglelist for each sub-style - void allocate(); + virtual void allocate(); + virtual void deallocate(); }; } // namespace LAMMPS_NS diff --git a/src/dihedral.cpp b/src/dihedral.cpp index 3e995fc405..2f591b1fc1 100644 --- a/src/dihedral.cpp +++ b/src/dihedral.cpp @@ -48,7 +48,7 @@ Dihedral::Dihedral(LAMMPS *_lmp) : Pointers(_lmp) datamask_read = ALL_MASK; datamask_modify = ALL_MASK; - copymode = 0; + copymode = kokkosable = 0; } /* ---------------------------------------------------------------------- */ diff --git a/src/dihedral.h b/src/dihedral.h index cf3d3f7d9a..34210929cd 100644 --- a/src/dihedral.h +++ b/src/dihedral.h @@ -41,7 +41,7 @@ class Dihedral : protected Pointers { ExecutionSpace execution_space; unsigned int datamask_read, datamask_modify; - int copymode; + int copymode, kokkosable; Dihedral(class LAMMPS *); ~Dihedral() override; diff --git a/src/dihedral_hybrid.cpp b/src/dihedral_hybrid.cpp index 4ee0ffdad9..3671391f5d 100644 --- a/src/dihedral_hybrid.cpp +++ b/src/dihedral_hybrid.cpp @@ -48,14 +48,7 @@ DihedralHybrid::~DihedralHybrid() delete[] keywords; } - if (allocated) { - memory->destroy(setflag); - memory->destroy(map); - delete[] ndihedrallist; - delete[] maxdihedral; - for (int i = 0; i < nstyles; i++) memory->destroy(dihedrallist[i]); - delete[] dihedrallist; - } + deallocate(); } /* ---------------------------------------------------------------------- */ @@ -172,6 +165,20 @@ void DihedralHybrid::allocate() for (int m = 0; m < nstyles; m++) dihedrallist[m] = nullptr; } +void DihedralHybrid::deallocate() +{ + if (!allocated) return; + + allocated = 0; + + memory->destroy(setflag); + memory->destroy(map); + delete[] ndihedrallist; + delete[] maxdihedral; + for (int i = 0; i < nstyles; i++) memory->destroy(dihedrallist[i]); + delete[] dihedrallist; +} + /* ---------------------------------------------------------------------- create one dihedral style for each arg in list ------------------------------------------------------------------------- */ @@ -191,15 +198,7 @@ void DihedralHybrid::settings(int narg, char **arg) delete[] keywords; } - if (allocated) { - memory->destroy(setflag); - memory->destroy(map); - delete[] ndihedrallist; - delete[] maxdihedral; - for (i = 0; i < nstyles; i++) memory->destroy(dihedrallist[i]); - delete[] dihedrallist; - } - allocated = 0; + deallocate(); // allocate list of sub-styles @@ -365,7 +364,7 @@ void DihedralHybrid::read_restart(FILE *fp) keywords[m] = new char[n]; if (me == 0) utils::sfread(FLERR, keywords[m], sizeof(char), n, fp, nullptr, error); MPI_Bcast(keywords[m], n, MPI_CHAR, 0, world); - styles[m] = force->new_dihedral(keywords[m], 0, dummy); + styles[m] = force->new_dihedral(keywords[m], 1, dummy); styles[m]->read_restart_settings(fp); } } diff --git a/src/dihedral_hybrid.h b/src/dihedral_hybrid.h index b7d4013afe..debc8a9d8d 100644 --- a/src/dihedral_hybrid.h +++ b/src/dihedral_hybrid.h @@ -40,14 +40,15 @@ class DihedralHybrid : public Dihedral { void read_restart(FILE *) override; double memory_usage() override; - private: + protected: int *map; // which style each dihedral type points to int *ndihedrallist; // # of dihedrals in sub-style dihedrallists int *maxdihedral; // max # of dihedrals sub-style lists can store int ***dihedrallist; // dihedrallist for each sub-style - void allocate(); + virtual void allocate(); + virtual void deallocate(); }; } // namespace LAMMPS_NS diff --git a/src/improper.cpp b/src/improper.cpp index dd4b1b2b25..3476bcdb50 100644 --- a/src/improper.cpp +++ b/src/improper.cpp @@ -47,7 +47,7 @@ Improper::Improper(LAMMPS *_lmp) : Pointers(_lmp) datamask_read = ALL_MASK; datamask_modify = ALL_MASK; - copymode = 0; + copymode = kokkosable = 0; } /* ---------------------------------------------------------------------- */ diff --git a/src/improper.h b/src/improper.h index 22a5d09926..400e950967 100644 --- a/src/improper.h +++ b/src/improper.h @@ -46,7 +46,7 @@ class Improper : protected Pointers { ExecutionSpace execution_space; unsigned int datamask_read, datamask_modify; - int copymode; + int copymode, kokkosable; Improper(class LAMMPS *); ~Improper() override; diff --git a/src/improper_hybrid.cpp b/src/improper_hybrid.cpp index a847b7bc95..5337f062b4 100644 --- a/src/improper_hybrid.cpp +++ b/src/improper_hybrid.cpp @@ -48,14 +48,7 @@ ImproperHybrid::~ImproperHybrid() delete[] keywords; } - if (allocated) { - memory->destroy(setflag); - memory->destroy(map); - delete[] nimproperlist; - delete[] maximproper; - for (int i = 0; i < nstyles; i++) memory->destroy(improperlist[i]); - delete[] improperlist; - } + deallocate(); } /* ---------------------------------------------------------------------- */ @@ -172,6 +165,22 @@ void ImproperHybrid::allocate() for (int m = 0; m < nstyles; m++) improperlist[m] = nullptr; } +/* ---------------------------------------------------------------------- */ + +void ImproperHybrid::deallocate() +{ + if (!allocated) return; + + allocated = 0; + + memory->destroy(setflag); + memory->destroy(map); + delete[] nimproperlist; + delete[] maximproper; + for (int i = 0; i < nstyles; i++) memory->destroy(improperlist[i]); + delete[] improperlist; +} + /* ---------------------------------------------------------------------- create one improper style for each arg in list ------------------------------------------------------------------------- */ @@ -191,15 +200,7 @@ void ImproperHybrid::settings(int narg, char **arg) delete[] keywords; } - if (allocated) { - memory->destroy(setflag); - memory->destroy(map); - delete[] nimproperlist; - delete[] maximproper; - for (i = 0; i < nstyles; i++) memory->destroy(improperlist[i]); - delete[] improperlist; - } - allocated = 0; + deallocate(); // allocate list of sub-styles @@ -357,7 +358,7 @@ void ImproperHybrid::read_restart(FILE *fp) keywords[m] = new char[n]; if (me == 0) utils::sfread(FLERR, keywords[m], sizeof(char), n, fp, nullptr, error); MPI_Bcast(keywords[m], n, MPI_CHAR, 0, world); - styles[m] = force->new_improper(keywords[m], 0, dummy); + styles[m] = force->new_improper(keywords[m], 1, dummy); styles[m]->read_restart_settings(fp); } } diff --git a/src/improper_hybrid.h b/src/improper_hybrid.h index e7cb8383d4..89a4664da2 100644 --- a/src/improper_hybrid.h +++ b/src/improper_hybrid.h @@ -40,14 +40,15 @@ class ImproperHybrid : public Improper { void read_restart(FILE *) override; double memory_usage() override; - private: + protected: int *map; // which style each improper type points to int *nimproperlist; // # of impropers in sub-style improperlists int *maximproper; // max # of impropers sub-style lists can store int ***improperlist; // improperlist for each sub-style - void allocate(); + virtual void allocate(); + virtual void deallocate(); }; } // namespace LAMMPS_NS