Merge pull request #4167 from stanmoore1/kk_hybrid_topo
Port hybrid bond topology styles to Kokkos
This commit is contained in:
@ -72,7 +72,7 @@ OPT.
|
|||||||
|
|
||||||
* :doc:`none <angle_none>`
|
* :doc:`none <angle_none>`
|
||||||
* :doc:`zero <angle_zero>`
|
* :doc:`zero <angle_zero>`
|
||||||
* :doc:`hybrid <angle_hybrid>`
|
* :doc:`hybrid (k) <angle_hybrid>`
|
||||||
*
|
*
|
||||||
*
|
*
|
||||||
*
|
*
|
||||||
|
|||||||
@ -1,8 +1,11 @@
|
|||||||
.. index:: bond_style hybrid
|
.. index:: bond_style hybrid
|
||||||
|
.. index:: bond_style hybrid/kk
|
||||||
|
|
||||||
bond_style hybrid command
|
bond_style hybrid command
|
||||||
=========================
|
=========================
|
||||||
|
|
||||||
|
Accelerator Variants: *hybrid/kk*
|
||||||
|
|
||||||
Syntax
|
Syntax
|
||||||
""""""
|
""""""
|
||||||
|
|
||||||
@ -15,7 +18,7 @@ Syntax
|
|||||||
Examples
|
Examples
|
||||||
""""""""
|
""""""""
|
||||||
|
|
||||||
.. code-block: LAMMPS
|
.. code-block:: LAMMPS
|
||||||
|
|
||||||
bond_style hybrid harmonic fene
|
bond_style hybrid harmonic fene
|
||||||
bond_coeff 1 harmonic 80.0 1.2
|
bond_coeff 1 harmonic 80.0 1.2
|
||||||
@ -60,6 +63,10 @@ bond types.
|
|||||||
|
|
||||||
----------
|
----------
|
||||||
|
|
||||||
|
.. include:: accel_styles.rst
|
||||||
|
|
||||||
|
----------
|
||||||
|
|
||||||
Restrictions
|
Restrictions
|
||||||
""""""""""""
|
""""""""""""
|
||||||
|
|
||||||
|
|||||||
@ -86,6 +86,8 @@ action bond_fene_kokkos.cpp bond_fene.cpp
|
|||||||
action bond_fene_kokkos.h bond_fene.h
|
action bond_fene_kokkos.h bond_fene.h
|
||||||
action bond_harmonic_kokkos.cpp bond_harmonic.cpp
|
action bond_harmonic_kokkos.cpp bond_harmonic.cpp
|
||||||
action bond_harmonic_kokkos.h bond_harmonic.h
|
action bond_harmonic_kokkos.h bond_harmonic.h
|
||||||
|
action bond_hybrid_kokkos.cpp bond_hybrid.cpp
|
||||||
|
action bond_hybrid_kokkos.h bond_hybrid.h
|
||||||
action comm_kokkos.cpp
|
action comm_kokkos.cpp
|
||||||
action comm_kokkos.h
|
action comm_kokkos.h
|
||||||
action comm_tiled_kokkos.cpp
|
action comm_tiled_kokkos.cpp
|
||||||
|
|||||||
@ -34,6 +34,8 @@ using namespace LAMMPS_NS;
|
|||||||
template<class DeviceType>
|
template<class DeviceType>
|
||||||
BondClass2Kokkos<DeviceType>::BondClass2Kokkos(LAMMPS *lmp) : BondClass2(lmp)
|
BondClass2Kokkos<DeviceType>::BondClass2Kokkos(LAMMPS *lmp) : BondClass2(lmp)
|
||||||
{
|
{
|
||||||
|
kokkosable = 1;
|
||||||
|
|
||||||
atomKK = (AtomKokkos *) atom;
|
atomKK = (AtomKokkos *) atom;
|
||||||
neighborKK = (NeighborKokkos *) neighbor;
|
neighborKK = (NeighborKokkos *) neighbor;
|
||||||
execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
|
execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
|
||||||
@ -122,12 +124,12 @@ void BondClass2Kokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
|||||||
|
|
||||||
if (eflag_atom) {
|
if (eflag_atom) {
|
||||||
k_eatom.template modify<DeviceType>();
|
k_eatom.template modify<DeviceType>();
|
||||||
k_eatom.template sync<LMPHostType>();
|
k_eatom.sync_host();
|
||||||
}
|
}
|
||||||
|
|
||||||
if (vflag_atom) {
|
if (vflag_atom) {
|
||||||
k_vatom.template modify<DeviceType>();
|
k_vatom.template modify<DeviceType>();
|
||||||
k_vatom.template sync<LMPHostType>();
|
k_vatom.sync_host();
|
||||||
}
|
}
|
||||||
|
|
||||||
copymode = 0;
|
copymode = 0;
|
||||||
@ -227,13 +229,13 @@ void BondClass2Kokkos<DeviceType>::coeff(int narg, char **arg)
|
|||||||
k_r0.h_view[i] = r0[i];
|
k_r0.h_view[i] = r0[i];
|
||||||
}
|
}
|
||||||
|
|
||||||
k_k2.template modify<LMPHostType>();
|
k_k2.modify_host();
|
||||||
k_k2.template sync<DeviceType>();
|
k_k2.template sync<DeviceType>();
|
||||||
k_k3.template modify<LMPHostType>();
|
k_k3.modify_host();
|
||||||
k_k3.template sync<DeviceType>();
|
k_k3.template sync<DeviceType>();
|
||||||
k_k4.template modify<LMPHostType>();
|
k_k4.modify_host();
|
||||||
k_k4.template sync<DeviceType>();
|
k_k4.template sync<DeviceType>();
|
||||||
k_r0.template modify<LMPHostType>();
|
k_r0.modify_host();
|
||||||
k_r0.template sync<DeviceType>();
|
k_r0.template sync<DeviceType>();
|
||||||
}
|
}
|
||||||
|
|
||||||
@ -264,13 +266,13 @@ void BondClass2Kokkos<DeviceType>::read_restart(FILE *fp)
|
|||||||
k_r0.h_view[i] = r0[i];
|
k_r0.h_view[i] = r0[i];
|
||||||
}
|
}
|
||||||
|
|
||||||
k_k2.template modify<LMPHostType>();
|
k_k2.modify_host();
|
||||||
k_k2.template sync<DeviceType>();
|
k_k2.template sync<DeviceType>();
|
||||||
k_k3.template modify<LMPHostType>();
|
k_k3.modify_host();
|
||||||
k_k3.template sync<DeviceType>();
|
k_k3.template sync<DeviceType>();
|
||||||
k_k4.template modify<LMPHostType>();
|
k_k4.modify_host();
|
||||||
k_k4.template sync<DeviceType>();
|
k_k4.template sync<DeviceType>();
|
||||||
k_r0.template modify<LMPHostType>();
|
k_r0.modify_host();
|
||||||
k_r0.template sync<DeviceType>();
|
k_r0.template sync<DeviceType>();
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|||||||
@ -59,6 +59,10 @@ class BondClass2Kokkos : public BondClass2 {
|
|||||||
const F_FLOAT &ebond, const F_FLOAT &fbond, const F_FLOAT &delx,
|
const F_FLOAT &ebond, const F_FLOAT &fbond, const F_FLOAT &delx,
|
||||||
const F_FLOAT &dely, const F_FLOAT &delz) const;
|
const F_FLOAT &dely, const F_FLOAT &delz) const;
|
||||||
|
|
||||||
|
typedef typename KKDevice<DeviceType>::value KKDeviceType;
|
||||||
|
Kokkos::DualView<E_FLOAT*,Kokkos::LayoutRight,KKDeviceType> k_eatom;
|
||||||
|
Kokkos::DualView<F_FLOAT*[6],Kokkos::LayoutRight,KKDeviceType> k_vatom;
|
||||||
|
|
||||||
protected:
|
protected:
|
||||||
|
|
||||||
class NeighborKokkos *neighborKK;
|
class NeighborKokkos *neighborKK;
|
||||||
@ -67,9 +71,6 @@ class BondClass2Kokkos : public BondClass2 {
|
|||||||
typename Kokkos::View<double*[3],typename AT::t_f_array::array_layout,typename KKDevice<DeviceType>::value,Kokkos::MemoryTraits<Kokkos::Atomic> > f;
|
typename Kokkos::View<double*[3],typename AT::t_f_array::array_layout,typename KKDevice<DeviceType>::value,Kokkos::MemoryTraits<Kokkos::Atomic> > f;
|
||||||
typename AT::t_int_2d bondlist;
|
typename AT::t_int_2d bondlist;
|
||||||
|
|
||||||
typedef typename KKDevice<DeviceType>::value KKDeviceType;
|
|
||||||
Kokkos::DualView<E_FLOAT*,Kokkos::LayoutRight,KKDeviceType> k_eatom;
|
|
||||||
Kokkos::DualView<F_FLOAT*[6],Kokkos::LayoutRight,KKDeviceType> k_vatom;
|
|
||||||
Kokkos::View<E_FLOAT*,Kokkos::LayoutRight,KKDeviceType,Kokkos::MemoryTraits<Kokkos::Atomic> > d_eatom;
|
Kokkos::View<E_FLOAT*,Kokkos::LayoutRight,KKDeviceType,Kokkos::MemoryTraits<Kokkos::Atomic> > d_eatom;
|
||||||
Kokkos::View<F_FLOAT*[6],Kokkos::LayoutRight,KKDeviceType,Kokkos::MemoryTraits<Kokkos::Atomic> > d_vatom;
|
Kokkos::View<F_FLOAT*[6],Kokkos::LayoutRight,KKDeviceType,Kokkos::MemoryTraits<Kokkos::Atomic> > d_vatom;
|
||||||
|
|
||||||
|
|||||||
@ -37,6 +37,8 @@ using MathConst::MY_CUBEROOT2;
|
|||||||
template<class DeviceType>
|
template<class DeviceType>
|
||||||
BondFENEKokkos<DeviceType>::BondFENEKokkos(LAMMPS *lmp) : BondFENE(lmp)
|
BondFENEKokkos<DeviceType>::BondFENEKokkos(LAMMPS *lmp) : BondFENE(lmp)
|
||||||
{
|
{
|
||||||
|
kokkosable = 1;
|
||||||
|
|
||||||
atomKK = (AtomKokkos *) atom;
|
atomKK = (AtomKokkos *) atom;
|
||||||
neighborKK = (NeighborKokkos *) neighbor;
|
neighborKK = (NeighborKokkos *) neighbor;
|
||||||
execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
|
execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
|
||||||
@ -135,12 +137,12 @@ void BondFENEKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
|||||||
|
|
||||||
if (eflag_atom) {
|
if (eflag_atom) {
|
||||||
k_eatom.template modify<DeviceType>();
|
k_eatom.template modify<DeviceType>();
|
||||||
k_eatom.template sync<LMPHostType>();
|
k_eatom.sync_host();
|
||||||
}
|
}
|
||||||
|
|
||||||
if (vflag_atom) {
|
if (vflag_atom) {
|
||||||
k_vatom.template modify<DeviceType>();
|
k_vatom.template modify<DeviceType>();
|
||||||
k_vatom.template sync<LMPHostType>();
|
k_vatom.sync_host();
|
||||||
}
|
}
|
||||||
|
|
||||||
copymode = 0;
|
copymode = 0;
|
||||||
@ -267,10 +269,10 @@ void BondFENEKokkos<DeviceType>::coeff(int narg, char **arg)
|
|||||||
k_sigma.h_view[i] = sigma[i];
|
k_sigma.h_view[i] = sigma[i];
|
||||||
}
|
}
|
||||||
|
|
||||||
k_k.template modify<LMPHostType>();
|
k_k.modify_host();
|
||||||
k_r0.template modify<LMPHostType>();
|
k_r0.modify_host();
|
||||||
k_epsilon.template modify<LMPHostType>();
|
k_epsilon.modify_host();
|
||||||
k_sigma.template modify<LMPHostType>();
|
k_sigma.modify_host();
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
@ -291,10 +293,10 @@ void BondFENEKokkos<DeviceType>::read_restart(FILE *fp)
|
|||||||
k_sigma.h_view[i] = sigma[i];
|
k_sigma.h_view[i] = sigma[i];
|
||||||
}
|
}
|
||||||
|
|
||||||
k_k.template modify<LMPHostType>();
|
k_k.modify_host();
|
||||||
k_r0.template modify<LMPHostType>();
|
k_r0.modify_host();
|
||||||
k_epsilon.template modify<LMPHostType>();
|
k_epsilon.modify_host();
|
||||||
k_sigma.template modify<LMPHostType>();
|
k_sigma.modify_host();
|
||||||
}
|
}
|
||||||
|
|
||||||
/* ----------------------------------------------------------------------
|
/* ----------------------------------------------------------------------
|
||||||
|
|||||||
@ -58,6 +58,9 @@ class BondFENEKokkos : public BondFENE {
|
|||||||
const F_FLOAT &ebond, const F_FLOAT &fbond, const F_FLOAT &delx,
|
const F_FLOAT &ebond, const F_FLOAT &fbond, const F_FLOAT &delx,
|
||||||
const F_FLOAT &dely, const F_FLOAT &delz) const;
|
const F_FLOAT &dely, const F_FLOAT &delz) const;
|
||||||
|
|
||||||
|
DAT::tdual_efloat_1d k_eatom;
|
||||||
|
DAT::tdual_virial_array k_vatom;
|
||||||
|
|
||||||
protected:
|
protected:
|
||||||
|
|
||||||
class NeighborKokkos *neighborKK;
|
class NeighborKokkos *neighborKK;
|
||||||
@ -66,8 +69,6 @@ class BondFENEKokkos : public BondFENE {
|
|||||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||||
typename ArrayTypes<DeviceType>::t_int_2d bondlist;
|
typename ArrayTypes<DeviceType>::t_int_2d bondlist;
|
||||||
|
|
||||||
DAT::tdual_efloat_1d k_eatom;
|
|
||||||
DAT::tdual_virial_array k_vatom;
|
|
||||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||||
|
|
||||||
|
|||||||
@ -34,6 +34,8 @@ using namespace LAMMPS_NS;
|
|||||||
template<class DeviceType>
|
template<class DeviceType>
|
||||||
BondHarmonicKokkos<DeviceType>::BondHarmonicKokkos(LAMMPS *lmp) : BondHarmonic(lmp)
|
BondHarmonicKokkos<DeviceType>::BondHarmonicKokkos(LAMMPS *lmp) : BondHarmonic(lmp)
|
||||||
{
|
{
|
||||||
|
kokkosable = 1;
|
||||||
|
|
||||||
atomKK = (AtomKokkos *) atom;
|
atomKK = (AtomKokkos *) atom;
|
||||||
neighborKK = (NeighborKokkos *) neighbor;
|
neighborKK = (NeighborKokkos *) neighbor;
|
||||||
execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
|
execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
|
||||||
@ -65,23 +67,20 @@ void BondHarmonicKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
|||||||
// reallocate per-atom arrays if necessary
|
// reallocate per-atom arrays if necessary
|
||||||
|
|
||||||
if (eflag_atom) {
|
if (eflag_atom) {
|
||||||
//if(k_eatom.extent(0)<maxeatom) { // won't work without adding zero functor
|
if (k_eatom.extent(0) < maxeatom) {
|
||||||
memoryKK->destroy_kokkos(k_eatom,eatom);
|
memoryKK->destroy_kokkos(k_eatom,eatom);
|
||||||
memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom");
|
memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom");
|
||||||
d_eatom = k_eatom.template view<KKDeviceType>();
|
d_eatom = k_eatom.template view<KKDeviceType>();
|
||||||
//}
|
} else Kokkos::deep_copy(d_eatom,0.0);
|
||||||
}
|
}
|
||||||
if (vflag_atom) {
|
if (vflag_atom) {
|
||||||
//if(k_vatom.extent(0)<maxvatom) { // won't work without adding zero functor
|
if (k_vatom.extent(0) < maxvatom) {
|
||||||
memoryKK->destroy_kokkos(k_vatom,vatom);
|
memoryKK->destroy_kokkos(k_vatom,vatom);
|
||||||
memoryKK->create_kokkos(k_vatom,vatom,maxvatom,"improper:vatom");
|
memoryKK->create_kokkos(k_vatom,vatom,maxvatom,"improper:vatom");
|
||||||
d_vatom = k_vatom.template view<KKDeviceType>();
|
d_vatom = k_vatom.template view<KKDeviceType>();
|
||||||
//}
|
} else Kokkos::deep_copy(d_vatom,0.0);
|
||||||
}
|
}
|
||||||
|
|
||||||
// if (eflag || vflag) atomKK->modified(execution_space,datamask_modify);
|
|
||||||
// else atomKK->modified(execution_space,F_MASK);
|
|
||||||
|
|
||||||
x = atomKK->k_x.template view<DeviceType>();
|
x = atomKK->k_x.template view<DeviceType>();
|
||||||
f = atomKK->k_f.template view<DeviceType>();
|
f = atomKK->k_f.template view<DeviceType>();
|
||||||
neighborKK->k_bondlist.template sync<DeviceType>();
|
neighborKK->k_bondlist.template sync<DeviceType>();
|
||||||
@ -122,12 +121,12 @@ void BondHarmonicKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
|||||||
|
|
||||||
if (eflag_atom) {
|
if (eflag_atom) {
|
||||||
k_eatom.template modify<DeviceType>();
|
k_eatom.template modify<DeviceType>();
|
||||||
k_eatom.template sync<LMPHostType>();
|
k_eatom.sync_host();
|
||||||
}
|
}
|
||||||
|
|
||||||
if (vflag_atom) {
|
if (vflag_atom) {
|
||||||
k_vatom.template modify<DeviceType>();
|
k_vatom.template modify<DeviceType>();
|
||||||
k_vatom.template sync<LMPHostType>();
|
k_vatom.sync_host();
|
||||||
}
|
}
|
||||||
|
|
||||||
copymode = 0;
|
copymode = 0;
|
||||||
@ -214,8 +213,8 @@ void BondHarmonicKokkos<DeviceType>::coeff(int narg, char **arg)
|
|||||||
k_r0.h_view[i] = r0[i];
|
k_r0.h_view[i] = r0[i];
|
||||||
}
|
}
|
||||||
|
|
||||||
k_k.template modify<LMPHostType>();
|
k_k.modify_host();
|
||||||
k_r0.template modify<LMPHostType>();
|
k_r0.modify_host();
|
||||||
k_k.template sync<DeviceType>();
|
k_k.template sync<DeviceType>();
|
||||||
k_r0.template sync<DeviceType>();
|
k_r0.template sync<DeviceType>();
|
||||||
}
|
}
|
||||||
@ -241,8 +240,8 @@ void BondHarmonicKokkos<DeviceType>::read_restart(FILE *fp)
|
|||||||
k_r0.h_view[i] = r0[i];
|
k_r0.h_view[i] = r0[i];
|
||||||
}
|
}
|
||||||
|
|
||||||
k_k.template modify<LMPHostType>();
|
k_k.modify_host();
|
||||||
k_r0.template modify<LMPHostType>();
|
k_r0.modify_host();
|
||||||
k_k.template sync<DeviceType>();
|
k_k.template sync<DeviceType>();
|
||||||
k_r0.template sync<DeviceType>();
|
k_r0.template sync<DeviceType>();
|
||||||
}
|
}
|
||||||
|
|||||||
@ -37,6 +37,7 @@ class BondHarmonicKokkos : public BondHarmonic {
|
|||||||
public:
|
public:
|
||||||
typedef DeviceType device_type;
|
typedef DeviceType device_type;
|
||||||
typedef EV_FLOAT value_type;
|
typedef EV_FLOAT value_type;
|
||||||
|
typedef ArrayTypes<DeviceType> AT;
|
||||||
|
|
||||||
BondHarmonicKokkos(class LAMMPS *);
|
BondHarmonicKokkos(class LAMMPS *);
|
||||||
~BondHarmonicKokkos() override;
|
~BondHarmonicKokkos() override;
|
||||||
@ -62,7 +63,6 @@ class BondHarmonicKokkos : public BondHarmonic {
|
|||||||
|
|
||||||
class NeighborKokkos *neighborKK;
|
class NeighborKokkos *neighborKK;
|
||||||
|
|
||||||
typedef ArrayTypes<DeviceType> AT;
|
|
||||||
typename AT::t_x_array_randomread x;
|
typename AT::t_x_array_randomread x;
|
||||||
typename Kokkos::View<double*[3],typename AT::t_f_array::array_layout,typename KKDevice<DeviceType>::value,Kokkos::MemoryTraits<Kokkos::Atomic> > f;
|
typename Kokkos::View<double*[3],typename AT::t_f_array::array_layout,typename KKDevice<DeviceType>::value,Kokkos::MemoryTraits<Kokkos::Atomic> > f;
|
||||||
typename AT::t_int_2d bondlist;
|
typename AT::t_int_2d bondlist;
|
||||||
|
|||||||
215
src/KOKKOS/bond_hybrid_kokkos.cpp
Normal file
215
src/KOKKOS/bond_hybrid_kokkos.cpp
Normal file
@ -0,0 +1,215 @@
|
|||||||
|
/* ----------------------------------------------------------------------
|
||||||
|
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 "bond_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 <cstring>
|
||||||
|
|
||||||
|
using namespace LAMMPS_NS;
|
||||||
|
|
||||||
|
#define EXTRA 1000
|
||||||
|
|
||||||
|
/* ---------------------------------------------------------------------- */
|
||||||
|
|
||||||
|
BondHybridKokkos::BondHybridKokkos(LAMMPS *lmp) : BondHybrid(lmp)
|
||||||
|
{
|
||||||
|
kokkosable = 1;
|
||||||
|
|
||||||
|
atomKK = (AtomKokkos *) atom;
|
||||||
|
neighborKK = (NeighborKokkos *) neighbor;
|
||||||
|
|
||||||
|
execution_space = Device;
|
||||||
|
|
||||||
|
datamask_read = EMPTY_MASK;
|
||||||
|
datamask_modify = EMPTY_MASK;
|
||||||
|
}
|
||||||
|
|
||||||
|
/* ---------------------------------------------------------------------- */
|
||||||
|
|
||||||
|
BondHybridKokkos::~BondHybridKokkos()
|
||||||
|
{
|
||||||
|
deallocate();
|
||||||
|
}
|
||||||
|
|
||||||
|
/* ---------------------------------------------------------------------- */
|
||||||
|
|
||||||
|
void BondHybridKokkos::compute(int eflag, int vflag)
|
||||||
|
{
|
||||||
|
// save ptrs to original bondlist
|
||||||
|
|
||||||
|
int nbondlist_orig = neighbor->nbondlist;
|
||||||
|
neighborKK->k_bondlist.sync_device();
|
||||||
|
auto k_bondlist_orig = neighborKK->k_bondlist;
|
||||||
|
auto d_bondlist_orig = k_bondlist_orig.d_view;
|
||||||
|
auto d_nbondlist = k_nbondlist.d_view;
|
||||||
|
auto h_nbondlist = k_nbondlist.h_view;
|
||||||
|
|
||||||
|
// if this is re-neighbor step, create sub-style bondlists
|
||||||
|
// nbondlist[] = length of each sub-style list
|
||||||
|
// realloc sub-style bondlist if necessary
|
||||||
|
// load sub-style bondlist with 3 values from original bondlist
|
||||||
|
|
||||||
|
if (neighbor->ago == 0) {
|
||||||
|
Kokkos::deep_copy(d_nbondlist,0);
|
||||||
|
|
||||||
|
k_map.sync_device();
|
||||||
|
auto d_map = k_map.d_view;
|
||||||
|
|
||||||
|
Kokkos::parallel_for(nbondlist_orig,LAMMPS_LAMBDA(int i) {
|
||||||
|
const int m = d_map[d_bondlist_orig(i,2)];
|
||||||
|
if (m >= 0) Kokkos::atomic_increment(&d_nbondlist[m]);
|
||||||
|
});
|
||||||
|
|
||||||
|
k_nbondlist.modify_device();
|
||||||
|
k_nbondlist.sync_host();
|
||||||
|
|
||||||
|
maxbond_all = 0;
|
||||||
|
for (int m = 0; m < nstyles; m++)
|
||||||
|
if (h_nbondlist[m] > maxbond_all)
|
||||||
|
maxbond_all = h_nbondlist[m] + EXTRA;
|
||||||
|
|
||||||
|
if (k_bondlist.d_view.extent(1) < maxbond_all)
|
||||||
|
MemKK::realloc_kokkos(k_bondlist, "bond_hybrid:bondlist", nstyles, maxbond_all, 3);
|
||||||
|
auto d_bondlist = k_bondlist.d_view;
|
||||||
|
|
||||||
|
Kokkos::deep_copy(d_nbondlist,0);
|
||||||
|
|
||||||
|
Kokkos::parallel_for(nbondlist_orig,LAMMPS_LAMBDA(int i) {
|
||||||
|
const int m = d_map[d_bondlist_orig(i,2)];
|
||||||
|
if (m < 0) return;
|
||||||
|
const int n = Kokkos::atomic_fetch_add(&d_nbondlist[m],1);
|
||||||
|
d_bondlist(m,n,0) = d_bondlist_orig(i,0);
|
||||||
|
d_bondlist(m,n,1) = d_bondlist_orig(i,1);
|
||||||
|
d_bondlist(m,n,2) = d_bondlist_orig(i,2);
|
||||||
|
});
|
||||||
|
}
|
||||||
|
|
||||||
|
// call each sub-style's compute function
|
||||||
|
// set neighbor->bondlist to sub-style bondlist before call
|
||||||
|
// accumulate sub-style global/peratom energy/virial in hybrid
|
||||||
|
|
||||||
|
ev_init(eflag, vflag);
|
||||||
|
|
||||||
|
k_nbondlist.modify_device();
|
||||||
|
k_nbondlist.sync_host();
|
||||||
|
|
||||||
|
for (int m = 0; m < nstyles; m++) {
|
||||||
|
neighbor->nbondlist = h_nbondlist[m];
|
||||||
|
auto k_bondlist_m = Kokkos::subview(k_bondlist,m,Kokkos::ALL,Kokkos::ALL);
|
||||||
|
k_bondlist_m.modify_device();
|
||||||
|
neighborKK->k_bondlist = k_bondlist_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];
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
// restore ptrs to original bondlist
|
||||||
|
|
||||||
|
neighbor->nbondlist = nbondlist_orig;
|
||||||
|
neighborKK->k_bondlist = k_bondlist_orig;
|
||||||
|
}
|
||||||
|
|
||||||
|
/* ---------------------------------------------------------------------- */
|
||||||
|
|
||||||
|
void BondHybridKokkos::allocate()
|
||||||
|
{
|
||||||
|
allocated = 1;
|
||||||
|
int n = atom->nbondtypes;
|
||||||
|
|
||||||
|
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;
|
||||||
|
|
||||||
|
k_nbondlist = DAT::tdual_int_1d("bond:nbondlist", nstyles);
|
||||||
|
}
|
||||||
|
|
||||||
|
/* ---------------------------------------------------------------------- */
|
||||||
|
|
||||||
|
void BondHybridKokkos::deallocate()
|
||||||
|
{
|
||||||
|
if (!allocated) return;
|
||||||
|
|
||||||
|
allocated = 0;
|
||||||
|
|
||||||
|
memory->destroy(setflag);
|
||||||
|
memoryKK->destroy_kokkos(k_map,map);
|
||||||
|
}
|
||||||
|
|
||||||
|
/* ----------------------------------------------------------------------
|
||||||
|
set coeffs for one type
|
||||||
|
---------------------------------------------------------------------- */
|
||||||
|
|
||||||
|
void BondHybridKokkos::coeff(int narg, char **arg)
|
||||||
|
{
|
||||||
|
BondHybrid::coeff(narg,arg);
|
||||||
|
|
||||||
|
k_map.modify_host();
|
||||||
|
}
|
||||||
|
|
||||||
|
/* ---------------------------------------------------------------------- */
|
||||||
|
|
||||||
|
void BondHybridKokkos::init_style()
|
||||||
|
{
|
||||||
|
BondHybrid::init_style();
|
||||||
|
|
||||||
|
for (int m = 0; m < nstyles; m++) {
|
||||||
|
if (!styles[m]->kokkosable)
|
||||||
|
error->all(FLERR,"Must use only Kokkos-enabled bond styles with bond_style hybrid/kk");
|
||||||
|
|
||||||
|
if (styles[m]->execution_space == Host)
|
||||||
|
lmp->kokkos->allow_overlap = 0;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
/* ----------------------------------------------------------------------
|
||||||
|
memory usage
|
||||||
|
------------------------------------------------------------------------- */
|
||||||
|
|
||||||
|
double BondHybridKokkos::memory_usage()
|
||||||
|
{
|
||||||
|
double bytes = (double) maxeatom * sizeof(double);
|
||||||
|
bytes += (double) maxvatom * 6 * sizeof(double);
|
||||||
|
for (int m = 0; m < nstyles; m++) bytes += (double) maxbond_all * 3 * sizeof(int);
|
||||||
|
for (int m = 0; m < nstyles; m++)
|
||||||
|
if (styles[m]) bytes += styles[m]->memory_usage();
|
||||||
|
return bytes;
|
||||||
|
}
|
||||||
58
src/KOKKOS/bond_hybrid_kokkos.h
Normal file
58
src/KOKKOS/bond_hybrid_kokkos.h
Normal file
@ -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
|
||||||
|
BondStyle(hybrid/kk,BondHybridKokkos);
|
||||||
|
BondStyle(hybrid/kk/device,BondHybridKokkos);
|
||||||
|
BondStyle(hybrid/kk/host,BondHybridKokkos);
|
||||||
|
// clang-format on
|
||||||
|
#else
|
||||||
|
|
||||||
|
// clang-format off
|
||||||
|
#ifndef LMP_BOND_HYBRID_KOKKOS_H
|
||||||
|
#define LMP_BOND_HYBRID_KOKKOS_H
|
||||||
|
|
||||||
|
#include "bond_hybrid.h"
|
||||||
|
#include "kokkos_type.h"
|
||||||
|
|
||||||
|
namespace LAMMPS_NS {
|
||||||
|
|
||||||
|
class BondHybridKokkos : public BondHybrid {
|
||||||
|
friend class Force;
|
||||||
|
|
||||||
|
public:
|
||||||
|
BondHybridKokkos(class LAMMPS *);
|
||||||
|
~BondHybridKokkos() 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
|
||||||
@ -1408,6 +1408,7 @@ typedef SNAComplex<SNAreal> SNAcomplex;
|
|||||||
#endif
|
#endif
|
||||||
|
|
||||||
#define LAMMPS_LAMBDA KOKKOS_LAMBDA
|
#define LAMMPS_LAMBDA KOKKOS_LAMBDA
|
||||||
|
#define LAMMPS_CLASS_LAMBDA KOKKOS_CLASS_LAMBDA
|
||||||
|
|
||||||
#if defined(KOKKOS_ENABLE_CUDA) || defined(KOKKOS_ENABLE_HIP)
|
#if defined(KOKKOS_ENABLE_CUDA) || defined(KOKKOS_ENABLE_HIP)
|
||||||
#define LAMMPS_DEVICE_FUNCTION __device__
|
#define LAMMPS_DEVICE_FUNCTION __device__
|
||||||
|
|||||||
@ -31,9 +31,7 @@ PairHybridKokkos::PairHybridKokkos(LAMMPS *lmp) : PairHybrid(lmp)
|
|||||||
kokkosable = 1;
|
kokkosable = 1;
|
||||||
atomKK = (AtomKokkos *) atom;
|
atomKK = (AtomKokkos *) atom;
|
||||||
|
|
||||||
// prevent overlapping host/device computation, which isn't
|
execution_space = Device;
|
||||||
// yet supported by pair_hybrid_kokkos
|
|
||||||
execution_space = Device;
|
|
||||||
|
|
||||||
datamask_read = EMPTY_MASK;
|
datamask_read = EMPTY_MASK;
|
||||||
datamask_modify = EMPTY_MASK;
|
datamask_modify = EMPTY_MASK;
|
||||||
|
|||||||
@ -63,8 +63,7 @@ Bond::Bond(LAMMPS *_lmp) : Pointers(_lmp)
|
|||||||
execution_space = Host;
|
execution_space = Host;
|
||||||
datamask_read = ALL_MASK;
|
datamask_read = ALL_MASK;
|
||||||
datamask_modify = ALL_MASK;
|
datamask_modify = ALL_MASK;
|
||||||
|
copymode = kokkosable = 0;
|
||||||
copymode = 0;
|
|
||||||
}
|
}
|
||||||
|
|
||||||
/* ---------------------------------------------------------------------- */
|
/* ---------------------------------------------------------------------- */
|
||||||
|
|||||||
@ -49,7 +49,7 @@ class Bond : protected Pointers {
|
|||||||
|
|
||||||
ExecutionSpace execution_space;
|
ExecutionSpace execution_space;
|
||||||
unsigned int datamask_read, datamask_modify;
|
unsigned int datamask_read, datamask_modify;
|
||||||
int copymode;
|
int copymode, kokkosable;
|
||||||
|
|
||||||
Bond(class LAMMPS *);
|
Bond(class LAMMPS *);
|
||||||
~Bond() override;
|
~Bond() override;
|
||||||
|
|||||||
@ -51,14 +51,7 @@ BondHybrid::~BondHybrid()
|
|||||||
|
|
||||||
delete[] svector;
|
delete[] svector;
|
||||||
|
|
||||||
if (allocated) {
|
deallocate();
|
||||||
memory->destroy(setflag);
|
|
||||||
memory->destroy(map);
|
|
||||||
delete[] nbondlist;
|
|
||||||
delete[] maxbond;
|
|
||||||
for (int i = 0; i < nstyles; i++) memory->destroy(bondlist[i]);
|
|
||||||
delete[] bondlist;
|
|
||||||
}
|
|
||||||
}
|
}
|
||||||
|
|
||||||
/* ---------------------------------------------------------------------- */
|
/* ---------------------------------------------------------------------- */
|
||||||
@ -166,6 +159,22 @@ void BondHybrid::allocate()
|
|||||||
for (int m = 0; m < nstyles; m++) bondlist[m] = nullptr;
|
for (int m = 0; m < nstyles; m++) bondlist[m] = nullptr;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
/* ---------------------------------------------------------------------- */
|
||||||
|
|
||||||
|
void BondHybrid::deallocate()
|
||||||
|
{
|
||||||
|
if (!allocated) return;
|
||||||
|
|
||||||
|
allocated = 0;
|
||||||
|
|
||||||
|
memory->destroy(setflag);
|
||||||
|
memory->destroy(map);
|
||||||
|
delete[] nbondlist;
|
||||||
|
delete[] maxbond;
|
||||||
|
for (int i = 0; i < nstyles; i++) memory->destroy(bondlist[i]);
|
||||||
|
delete[] bondlist;
|
||||||
|
}
|
||||||
|
|
||||||
/* ----------------------------------------------------------------------
|
/* ----------------------------------------------------------------------
|
||||||
create one bond style for each arg in list
|
create one bond style for each arg in list
|
||||||
------------------------------------------------------------------------- */
|
------------------------------------------------------------------------- */
|
||||||
@ -186,15 +195,7 @@ void BondHybrid::settings(int narg, char **arg)
|
|||||||
has_quartic = -1;
|
has_quartic = -1;
|
||||||
}
|
}
|
||||||
|
|
||||||
if (allocated) {
|
deallocate();
|
||||||
memory->destroy(setflag);
|
|
||||||
memory->destroy(map);
|
|
||||||
delete[] nbondlist;
|
|
||||||
delete[] maxbond;
|
|
||||||
for (i = 0; i < nstyles; i++) memory->destroy(bondlist[i]);
|
|
||||||
delete[] bondlist;
|
|
||||||
}
|
|
||||||
allocated = 0;
|
|
||||||
|
|
||||||
// allocate list of sub-styles
|
// allocate list of sub-styles
|
||||||
|
|
||||||
@ -400,7 +401,7 @@ void BondHybrid::read_restart(FILE *fp)
|
|||||||
keywords[m] = new char[n];
|
keywords[m] = new char[n];
|
||||||
if (me == 0) utils::sfread(FLERR, keywords[m], sizeof(char), n, fp, nullptr, error);
|
if (me == 0) utils::sfread(FLERR, keywords[m], sizeof(char), n, fp, nullptr, error);
|
||||||
MPI_Bcast(keywords[m], n, MPI_CHAR, 0, world);
|
MPI_Bcast(keywords[m], n, MPI_CHAR, 0, world);
|
||||||
styles[m] = force->new_bond(keywords[m], 0, dummy);
|
styles[m] = force->new_bond(keywords[m], 1, dummy);
|
||||||
styles[m]->read_restart_settings(fp);
|
styles[m]->read_restart_settings(fp);
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|||||||
@ -44,14 +44,15 @@ class BondHybrid : public Bond {
|
|||||||
double single(int, double, int, int, double &) override;
|
double single(int, double, int, int, double &) override;
|
||||||
double memory_usage() override;
|
double memory_usage() override;
|
||||||
|
|
||||||
private:
|
protected:
|
||||||
int *map; // which style each bond type points to
|
int *map; // which style each bond type points to
|
||||||
int has_quartic; // which style, if any is a quartic bond style
|
int has_quartic; // which style, if any is a quartic bond style
|
||||||
int *nbondlist; // # of bonds in sub-style bondlists
|
int *nbondlist; // # of bonds in sub-style bondlists
|
||||||
int *maxbond; // max # of bonds sub-style lists can store
|
int *maxbond; // max # of bonds sub-style lists can store
|
||||||
int ***bondlist; // bondlist for each sub-style
|
int ***bondlist; // bondlist for each sub-style
|
||||||
|
|
||||||
void allocate();
|
virtual void allocate();
|
||||||
|
virtual void deallocate();
|
||||||
void flags();
|
void flags();
|
||||||
|
|
||||||
virtual void init_svector();
|
virtual void init_svector();
|
||||||
|
|||||||
Reference in New Issue
Block a user