ICODE-ADC/lammps
4
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge branch 'develop' into collected-small-fixes

Browse Source
This commit is contained in:
Axel Kohlmeyer
2024-07-26 18:32:49 -04:00
parent a2f461c688 1014cd1e0e
commit 06767886de
46 changed files with 1240 additions and 353 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
doc/src/Commands_bond.rst
View File

@ -73,7 +73,7 @@ OPT.
* :doc:`none <angle_none>`
* :doc:`zero <angle_zero>`
* :doc:`hybrid <angle_hybrid>`
* :doc:`hybrid (k) <angle_hybrid>`
*
*
*
@ -119,7 +119,7 @@ OPT.
* :doc:`none <dihedral_none>`
* :doc:`zero <dihedral_zero>`
* :doc:`hybrid <dihedral_hybrid>`
* :doc:`hybrid (k) <dihedral_hybrid>`
*
*
*
@ -157,7 +157,7 @@ OPT.
* :doc:`none <improper_none>`
* :doc:`zero <improper_zero>`
* :doc:`hybrid <improper_hybrid>`
* :doc:`hybrid (k) <improper_hybrid>`
*
*
*

3
doc/src/angle_hybrid.rst
View File

@ -1,8 +1,11 @@
.. index:: angle_style hybrid
.. index:: angle_style hybrid/kk
angle_style hybrid command
==========================
Accelerator Variants: *hybrid/kk*
Syntax
""""""

3
doc/src/dihedral_hybrid.rst
View File

@ -1,8 +1,11 @@
.. index:: dihedral_style hybrid
.. index:: dihedral_style hybrid/kk
dihedral_style hybrid command
=============================
Accelerator Variants: *hybrid/kk*
Syntax
""""""

3
doc/src/improper_hybrid.rst
View File

@ -1,8 +1,11 @@
.. index:: improper_style hybrid
.. index:: improper_style hybrid/kk
improper_style hybrid command
=============================
Accelerator Variants: *hybrid/kk*
Syntax
""""""

6
src/KOKKOS/Install.sh
View File

@ -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 atom_kokkos.cpp
action atom_kokkos.h
action atom_map_kokkos.cpp
@ -116,6 +118,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
@ -205,6 +209,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

21
src/KOKKOS/angle_charmm_kokkos.cpp
View File

@ -38,6 +38,7 @@ static constexpr double SMALL = 0.001;
template<class DeviceType>
AngleCharmmKokkos<DeviceType>::AngleCharmmKokkos(LAMMPS *lmp) : AngleCharmm(lmp)
{
kokkosable = 1;
atomKK = (AtomKokkos *) atom;
neighborKK = (NeighborKokkos *) neighbor;
execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
@ -125,12 +126,12 @@ void AngleCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
k_eatom.sync_host();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
k_vatom.sync_host();
}
copymode = 0;
@ -284,10 +285,10 @@ void AngleCharmmKokkos<DeviceType>::coeff(int narg, char **arg)
k_r_ub.h_view[i] = r_ub[i];
}
k_k.template modify<LMPHostType>();
k_theta0.template modify<LMPHostType>();
k_k_ub.template modify<LMPHostType>();
k_r_ub.template modify<LMPHostType>();
k_k.modify_host();
k_theta0.modify_host();
k_k_ub.modify_host();
k_r_ub.modify_host();
k_k.template sync<DeviceType>();
k_theta0.template sync<DeviceType>();
@ -322,10 +323,10 @@ void AngleCharmmKokkos<DeviceType>::read_restart(FILE *fp)
k_r_ub.h_view[i] = r_ub[i];
}
k_k.template modify<LMPHostType>();
k_theta0.template modify<LMPHostType>();
k_k_ub.template modify<LMPHostType>();
k_r_ub.template modify<LMPHostType>();
k_k.modify_host();
k_theta0.modify_host();
k_k_ub.modify_host();
k_r_ub.modify_host();
k_k.template sync<DeviceType>();
k_theta0.template sync<DeviceType>();

9
src/KOKKOS/angle_charmm_kokkos.h
View File

@ -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<DeviceType>::value;
Kokkos::DualView<E_FLOAT*,Kokkos::LayoutRight,KKDeviceType> k_eatom;
Kokkos::DualView<F_FLOAT*[6],Kokkos::LayoutRight,KKDeviceType> k_vatom;
protected:
class NeighborKokkos *neighborKK;
typedef ArrayTypes<DeviceType> AT;
typename AT::t_x_array_randomread x;
using KKDeviceType = typename KKDevice<DeviceType>::value;
typename Kokkos::View<double*[3],typename AT::t_f_array::array_layout,KKDeviceType,Kokkos::MemoryTraits<Kokkos::Atomic> > f;
typename AT::t_int_2d anglelist;
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<F_FLOAT*[6],Kokkos::LayoutRight,KKDeviceType,Kokkos::MemoryTraits<Kokkos::Atomic>> d_vatom;

65
src/KOKKOS/angle_class2_kokkos.cpp
View File

@ -38,6 +38,7 @@ static constexpr double SMALL = 0.001;
template<class DeviceType>
AngleClass2Kokkos<DeviceType>::AngleClass2Kokkos(LAMMPS *lmp) : AngleClass2(lmp)
{
kokkosable = 1;
atomKK = (AtomKokkos *) atom;
neighborKK = (NeighborKokkos *) neighbor;
execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
@ -141,12 +142,12 @@ void AngleClass2Kokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
k_eatom.sync_host();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
k_vatom.sync_host();
}
copymode = 0;
@ -386,21 +387,21 @@ void AngleClass2Kokkos<DeviceType>::coeff(int narg, char **arg)
k_theta0.h_view[i] = theta0[i];
}
k_k2.template modify<LMPHostType>();
k_k3.template modify<LMPHostType>();
k_k4.template modify<LMPHostType>();
k_bb_k.template modify<LMPHostType>();
k_bb_r1.template modify<LMPHostType>();
k_bb_r2.template modify<LMPHostType>();
k_ba_k1.template modify<LMPHostType>();
k_ba_k2.template modify<LMPHostType>();
k_ba_r1.template modify<LMPHostType>();
k_ba_r2.template modify<LMPHostType>();
k_setflag.template modify<LMPHostType>();
k_setflag_a.template modify<LMPHostType>();
k_setflag_bb.template modify<LMPHostType>();
k_setflag_ba.template modify<LMPHostType>();
k_theta0.template modify<LMPHostType>();
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<DeviceType>::read_restart(FILE *fp)
k_theta0.h_view[i] = theta0[i];
}
k_k2.template modify<LMPHostType>();
k_k3.template modify<LMPHostType>();
k_k4.template modify<LMPHostType>();
k_bb_k.template modify<LMPHostType>();
k_bb_r1.template modify<LMPHostType>();
k_bb_r2.template modify<LMPHostType>();
k_ba_k1.template modify<LMPHostType>();
k_ba_k2.template modify<LMPHostType>();
k_ba_r1.template modify<LMPHostType>();
k_ba_r2.template modify<LMPHostType>();
k_setflag.template modify<LMPHostType>();
k_setflag_a.template modify<LMPHostType>();
k_setflag_bb.template modify<LMPHostType>();
k_setflag_ba.template modify<LMPHostType>();
k_theta0.template modify<LMPHostType>();
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();
}
/* ----------------------------------------------------------------------

8
src/KOKKOS/angle_class2_kokkos.h
View File

@ -36,8 +36,8 @@ class AngleClass2Kokkos : public AngleClass2 {
public:
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
typedef EV_FLOAT value_type;
typedef ArrayTypes<DeviceType> 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;

9
src/KOKKOS/angle_cosine_kokkos.cpp
View File

@ -36,6 +36,7 @@ using namespace MathConst;
template<class DeviceType>
AngleCosineKokkos<DeviceType>::AngleCosineKokkos(LAMMPS *lmp) : AngleCosine(lmp)
{
kokkosable = 1;
atomKK = (AtomKokkos *) atom;
neighborKK = (NeighborKokkos *) neighbor;
execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
@ -124,12 +125,12 @@ void AngleCosineKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
k_eatom.sync_host();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
k_vatom.sync_host();
}
copymode = 0;
@ -254,7 +255,7 @@ void AngleCosineKokkos<DeviceType>::coeff(int narg, char **arg)
for (int i = 1; i <= n; i++)
k_k.h_view[i] = k[i];
k_k.template modify<LMPHostType>();
k_k.modify_host();
}
/* ----------------------------------------------------------------------
@ -270,7 +271,7 @@ void AngleCosineKokkos<DeviceType>::read_restart(FILE *fp)
for (int i = 1; i <= n; i++)
k_k.h_view[i] = k[i];
k_k.template modify<LMPHostType>();
k_k.modify_host();
}
/* ----------------------------------------------------------------------

7
src/KOKKOS/angle_cosine_kokkos.h
View File

@ -37,6 +37,7 @@ class AngleCosineKokkos : public AngleCosine {
public:
typedef DeviceType device_type;
typedef EV_FLOAT value_type;
typedef ArrayTypes<DeviceType> 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<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_2d anglelist;
typename ArrayTypes<DeviceType>::tdual_efloat_1d k_eatom;
typename ArrayTypes<DeviceType>::tdual_virial_array k_vatom;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;

17
src/KOKKOS/angle_harmonic_kokkos.cpp
View File

@ -38,6 +38,7 @@ static constexpr double SMALL = 0.001;
template<class DeviceType>
AngleHarmonicKokkos<DeviceType>::AngleHarmonicKokkos(LAMMPS *lmp) : AngleHarmonic(lmp)
{
kokkosable = 1;
atomKK = (AtomKokkos *) atom;
neighborKK = (NeighborKokkos *) neighbor;
execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
@ -71,14 +72,18 @@ void AngleHarmonicKokkos<DeviceType>::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<DeviceType>();
} 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<DeviceType>();
} else Kokkos::deep_copy(d_vatom,0.0);
}
//atomKK->sync(execution_space,datamask_read);
@ -127,12 +132,12 @@ void AngleHarmonicKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
k_eatom.sync_host();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
k_vatom.sync_host();
}
copymode = 0;
@ -264,8 +269,8 @@ void AngleHarmonicKokkos<DeviceType>::coeff(int narg, char **arg)
k_theta0.h_view[i] = theta0[i];
}
k_k.template modify<LMPHostType>();
k_theta0.template modify<LMPHostType>();
k_k.modify_host();
k_theta0.modify_host();
}
/* ----------------------------------------------------------------------
@ -283,8 +288,8 @@ void AngleHarmonicKokkos<DeviceType>::read_restart(FILE *fp)
k_theta0.h_view[i] = theta0[i];
}
k_k.template modify<LMPHostType>();
k_theta0.template modify<LMPHostType>();
k_k.modify_host();
k_theta0.modify_host();
}
/* ----------------------------------------------------------------------

7
src/KOKKOS/angle_harmonic_kokkos.h
View File

@ -37,6 +37,7 @@ class AngleHarmonicKokkos : public AngleHarmonic {
public:
typedef DeviceType device_type;
typedef EV_FLOAT value_type;
typedef ArrayTypes<DeviceType> 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<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_2d anglelist;
typename ArrayTypes<DeviceType>::tdual_efloat_1d k_eatom;
typename ArrayTypes<DeviceType>::tdual_virial_array k_vatom;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;

224
src/KOKKOS/angle_hybrid_kokkos.cpp Normal file
View File

@ -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 <cstring>
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;
}

58
src/KOKKOS/angle_hybrid_kokkos.h Normal file
View 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 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

8
src/KOKKOS/bond_hybrid_kokkos.cpp
View File

@ -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);
}

45
src/KOKKOS/dihedral_charmm_kokkos.cpp
View File

@ -40,6 +40,7 @@ static constexpr double TOLERANCE = 0.05;
template<class DeviceType>
DihedralCharmmKokkos<DeviceType>::DihedralCharmmKokkos(LAMMPS *lmp) : DihedralCharmm(lmp)
{
kokkosable = 1;
atomKK = (AtomKokkos *) atom;
neighborKK = (NeighborKokkos *) neighbor;
execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
@ -115,7 +116,7 @@ void DihedralCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
qqrd2e = force->qqrd2e;
h_warning_flag() = 0;
k_warning_flag.template modify<LMPHostType>();
k_warning_flag.modify_host();
k_warning_flag.template sync<DeviceType>();
copymode = 1;
@ -141,7 +142,7 @@ void DihedralCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// error check
k_warning_flag.template modify<DeviceType>();
k_warning_flag.template sync<LMPHostType>();
k_warning_flag.sync_host();
if (h_warning_flag())
error->warning(FLERR,"Dihedral problem");
@ -174,20 +175,20 @@ void DihedralCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
k_eatom.sync_host();
k_eatom_pair.template modify<DeviceType>();
k_eatom_pair.template sync<LMPHostType>();
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<DeviceType>();
k_vatom.template sync<LMPHostType>();
k_vatom.sync_host();
k_vatom_pair.template modify<DeviceType>();
k_vatom_pair.template sync<LMPHostType>();
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<DeviceType>::coeff(int narg, char **arg)
k_weight.h_view[i] = weight[i];
}
k_k.template modify<LMPHostType>();
k_multiplicity.template modify<LMPHostType>();
k_shift.template modify<LMPHostType>();
k_cos_shift.template modify<LMPHostType>();
k_sin_shift.template modify<LMPHostType>();
k_weight.template modify<LMPHostType>();
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<DeviceType>();
k_multiplicity.template sync<DeviceType>();
@ -502,10 +503,10 @@ void DihedralCharmmKokkos<DeviceType>::init_style()
}
}
k_lj14_1.template modify<LMPHostType>();
k_lj14_2.template modify<LMPHostType>();
k_lj14_3.template modify<LMPHostType>();
k_lj14_4.template modify<LMPHostType>();
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<DeviceType>();
k_lj14_2.template sync<DeviceType>();
@ -547,12 +548,12 @@ void DihedralCharmmKokkos<DeviceType>::read_restart(FILE *fp)
k_weight.h_view[i] = weight[i];
}
k_k.template modify<LMPHostType>();
k_multiplicity.template modify<LMPHostType>();
k_shift.template modify<LMPHostType>();
k_cos_shift.template modify<LMPHostType>();
k_sin_shift.template modify<LMPHostType>();
k_weight.template modify<LMPHostType>();
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<DeviceType>();
k_multiplicity.template sync<DeviceType>();

7
src/KOKKOS/dihedral_charmm_kokkos.h
View File

@ -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<DeviceType>::value KKDeviceType;
Kokkos::DualView<E_FLOAT*,Kokkos::LayoutRight,KKDeviceType> k_eatom;
Kokkos::DualView<F_FLOAT*[6],Kokkos::LayoutRight,KKDeviceType> 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<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<F_FLOAT*[6],Kokkos::LayoutRight,KKDeviceType,Kokkos::MemoryTraits<Kokkos::Atomic> > d_vatom;

45
src/KOKKOS/dihedral_charmmfsw_kokkos.cpp
View File

@ -47,6 +47,7 @@ static constexpr double TOLERANCE = 0.05;
template<class DeviceType>
DihedralCharmmfswKokkos<DeviceType>::DihedralCharmmfswKokkos(LAMMPS *lmp) : DihedralCharmmfsw(lmp)
{
kokkosable = 1;
atomKK = (AtomKokkos *) atom;
neighborKK = (NeighborKokkos *) neighbor;
execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
@ -122,7 +123,7 @@ void DihedralCharmmfswKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
qqrd2e = force->qqrd2e;
h_warning_flag() = 0;
k_warning_flag.template modify<LMPHostType>();
k_warning_flag.modify_host();
k_warning_flag.template sync<DeviceType>();
copymode = 1;
@ -148,7 +149,7 @@ void DihedralCharmmfswKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// error check
k_warning_flag.template modify<DeviceType>();
k_warning_flag.template sync<LMPHostType>();
k_warning_flag.sync_host();
if (h_warning_flag())
error->warning(FLERR,"Dihedral problem");
@ -181,20 +182,20 @@ void DihedralCharmmfswKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
k_eatom.sync_host();
k_eatom_pair.template modify<DeviceType>();
k_eatom_pair.template sync<LMPHostType>();
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<DeviceType>();
k_vatom.template sync<LMPHostType>();
k_vatom.sync_host();
k_vatom_pair.template modify<DeviceType>();
k_vatom_pair.template sync<LMPHostType>();
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<DeviceType>::coeff(int narg, char **arg)
k_weight.h_view[i] = weight[i];
}
k_k.template modify<LMPHostType>();
k_multiplicity.template modify<LMPHostType>();
k_shift.template modify<LMPHostType>();
k_cos_shift.template modify<LMPHostType>();
k_sin_shift.template modify<LMPHostType>();
k_weight.template modify<LMPHostType>();
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<DeviceType>();
k_multiplicity.template sync<DeviceType>();
@ -519,10 +520,10 @@ void DihedralCharmmfswKokkos<DeviceType>::init_style()
}
}
k_lj14_1.template modify<LMPHostType>();
k_lj14_2.template modify<LMPHostType>();
k_lj14_3.template modify<LMPHostType>();
k_lj14_4.template modify<LMPHostType>();
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<DeviceType>();
k_lj14_2.template sync<DeviceType>();
@ -564,12 +565,12 @@ void DihedralCharmmfswKokkos<DeviceType>::read_restart(FILE *fp)
k_weight.h_view[i] = weight[i];
}
k_k.template modify<LMPHostType>();
k_multiplicity.template modify<LMPHostType>();
k_shift.template modify<LMPHostType>();
k_cos_shift.template modify<LMPHostType>();
k_sin_shift.template modify<LMPHostType>();
k_weight.template modify<LMPHostType>();
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<DeviceType>();
k_multiplicity.template sync<DeviceType>();

8
src/KOKKOS/dihedral_charmmfsw_kokkos.h
View File

@ -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<DeviceType>::value KKDeviceType;
Kokkos::DualView<E_FLOAT*,Kokkos::LayoutRight,KKDeviceType> k_eatom;
Kokkos::DualView<F_FLOAT*[6],Kokkos::LayoutRight,KKDeviceType> 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<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<F_FLOAT*[6],Kokkos::LayoutRight,KKDeviceType,Kokkos::MemoryTraits<Kokkos::Atomic> > d_vatom;

161
src/KOKKOS/dihedral_class2_kokkos.cpp
View File

@ -38,6 +38,7 @@ static constexpr double SMALL = 0.001;
template<class DeviceType>
DihedralClass2Kokkos<DeviceType>::DihedralClass2Kokkos(LAMMPS *lmp) : DihedralClass2(lmp)
{
kokkosable = 1;
atomKK = (AtomKokkos *) atom;
neighborKK = (NeighborKokkos *) neighbor;
execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
@ -137,7 +138,7 @@ void DihedralClass2Kokkos<DeviceType>::compute(int eflag_in, int vflag_in)
newton_bond = force->newton_bond;
h_warning_flag() = 0;
k_warning_flag.template modify<LMPHostType>();
k_warning_flag.modify_host();
k_warning_flag.template sync<DeviceType>();
copymode = 1;
@ -163,7 +164,7 @@ void DihedralClass2Kokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// error check
k_warning_flag.template modify<DeviceType>();
k_warning_flag.template sync<LMPHostType>();
k_warning_flag.sync_host();
if (h_warning_flag())
error->warning(FLERR,"Dihedral problem");
@ -179,12 +180,12 @@ void DihedralClass2Kokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
k_eatom.sync_host();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
k_vatom.sync_host();
}
copymode = 0;
@ -786,44 +787,44 @@ void DihedralClass2Kokkos<DeviceType>::coeff(int narg, char **arg)
k_setflag_bb13t.h_view[i] = setflag_bb13t[i];
}
k_k1.template modify<LMPHostType>();
k_k2.template modify<LMPHostType>();
k_k3.template modify<LMPHostType>();
k_phi1.template modify<LMPHostType>();
k_phi2.template modify<LMPHostType>();
k_phi3.template modify<LMPHostType>();
k_mbt_f1.template modify<LMPHostType>();
k_mbt_f2.template modify<LMPHostType>();
k_mbt_f3.template modify<LMPHostType>();
k_mbt_r0.template modify<LMPHostType>();
k_ebt_f1_1.template modify<LMPHostType>();
k_ebt_f2_1.template modify<LMPHostType>();
k_ebt_f3_1.template modify<LMPHostType>();
k_ebt_r0_1.template modify<LMPHostType>();
k_ebt_f1_2.template modify<LMPHostType>();
k_ebt_f2_2.template modify<LMPHostType>();
k_ebt_f3_2.template modify<LMPHostType>();
k_ebt_r0_2.template modify<LMPHostType>();
k_at_f1_1.template modify<LMPHostType>();
k_at_f2_1.template modify<LMPHostType>();
k_at_f3_1.template modify<LMPHostType>();
k_at_f1_2.template modify<LMPHostType>();
k_at_f2_2.template modify<LMPHostType>();
k_at_f3_2.template modify<LMPHostType>();
k_at_theta0_1.template modify<LMPHostType>();
k_at_theta0_2.template modify<LMPHostType>();
k_aat_k.template modify<LMPHostType>();
k_aat_theta0_1.template modify<LMPHostType>();
k_aat_theta0_2.template modify<LMPHostType>();
k_bb13t_k.template modify<LMPHostType>();
k_bb13t_r10.template modify<LMPHostType>();
k_bb13t_r30.template modify<LMPHostType>();
k_setflag_d.template modify<LMPHostType>();
k_setflag_mbt.template modify<LMPHostType>();
k_setflag_ebt.template modify<LMPHostType>();
k_setflag_at.template modify<LMPHostType>();
k_setflag_aat.template modify<LMPHostType>();
k_setflag_bb13t.template modify<LMPHostType>();
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<DeviceType>::read_restart(FILE *fp)
k_setflag_bb13t.h_view[i] = setflag_bb13t[i];
}
k_k1.template modify<LMPHostType>();
k_k2.template modify<LMPHostType>();
k_k3.template modify<LMPHostType>();
k_phi1.template modify<LMPHostType>();
k_phi2.template modify<LMPHostType>();
k_phi3.template modify<LMPHostType>();
k_mbt_f1.template modify<LMPHostType>();
k_mbt_f2.template modify<LMPHostType>();
k_mbt_f3.template modify<LMPHostType>();
k_mbt_r0.template modify<LMPHostType>();
k_ebt_f1_1.template modify<LMPHostType>();
k_ebt_f2_1.template modify<LMPHostType>();
k_ebt_f3_1.template modify<LMPHostType>();
k_ebt_r0_1.template modify<LMPHostType>();
k_ebt_f1_2.template modify<LMPHostType>();
k_ebt_f2_2.template modify<LMPHostType>();
k_ebt_f3_2.template modify<LMPHostType>();
k_ebt_r0_2.template modify<LMPHostType>();
k_at_f1_1.template modify<LMPHostType>();
k_at_f2_1.template modify<LMPHostType>();
k_at_f3_1.template modify<LMPHostType>();
k_at_f1_2.template modify<LMPHostType>();
k_at_f2_2.template modify<LMPHostType>();
k_at_f3_2.template modify<LMPHostType>();
k_at_theta0_1.template modify<LMPHostType>();
k_at_theta0_2.template modify<LMPHostType>();
k_aat_k.template modify<LMPHostType>();
k_aat_theta0_1.template modify<LMPHostType>();
k_aat_theta0_2.template modify<LMPHostType>();
k_bb13t_k.template modify<LMPHostType>();
k_bb13t_r10.template modify<LMPHostType>();
k_bb13t_r30.template modify<LMPHostType>();
k_setflag_d.template modify<LMPHostType>();
k_setflag_mbt.template modify<LMPHostType>();
k_setflag_ebt.template modify<LMPHostType>();
k_setflag_at.template modify<LMPHostType>();
k_setflag_aat.template modify<LMPHostType>();
k_setflag_bb13t.template modify<LMPHostType>();
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();
}
/* ----------------------------------------------------------------------

6
src/KOKKOS/dihedral_class2_kokkos.h
View File

@ -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;

33
src/KOKKOS/dihedral_harmonic_kokkos.cpp
View File

@ -37,6 +37,7 @@ static constexpr double TOLERANCE = 0.05;
template<class DeviceType>
DihedralHarmonicKokkos<DeviceType>::DihedralHarmonicKokkos(LAMMPS *lmp) : DihedralHarmonic(lmp)
{
kokkosable = 1;
atomKK = (AtomKokkos *) atom;
neighborKK = (NeighborKokkos *) neighbor;
execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
@ -74,14 +75,18 @@ void DihedralHarmonicKokkos<DeviceType>::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<DeviceType>();
} 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<DeviceType>();
} else Kokkos::deep_copy(d_vatom,0.0);
}
k_k.template sync<DeviceType>();
@ -99,7 +104,7 @@ void DihedralHarmonicKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
newton_bond = force->newton_bond;
h_warning_flag() = 0;
k_warning_flag.template modify<LMPHostType>();
k_warning_flag.modify_host();
k_warning_flag.template sync<DeviceType>();
copymode = 1;
@ -125,7 +130,7 @@ void DihedralHarmonicKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// error check
k_warning_flag.template modify<DeviceType>();
k_warning_flag.template sync<LMPHostType>();
k_warning_flag.sync_host();
if (h_warning_flag())
error->warning(FLERR,"Dihedral problem");
@ -141,12 +146,12 @@ void DihedralHarmonicKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
k_eatom.sync_host();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
k_vatom.sync_host();
}
copymode = 0;
@ -362,11 +367,11 @@ void DihedralHarmonicKokkos<DeviceType>::coeff(int narg, char **arg)
k_multiplicity.h_view[i] = multiplicity[i];
}
k_k.template modify<LMPHostType>();
k_cos_shift.template modify<LMPHostType>();
k_sin_shift.template modify<LMPHostType>();
k_sign.template modify<LMPHostType>();
k_multiplicity.template modify<LMPHostType>();
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<DeviceType>::read_restart(FILE *fp)
k_multiplicity.h_view[i] = multiplicity[i];
}
k_k.template modify<LMPHostType>();
k_cos_shift.template modify<LMPHostType>();
k_sin_shift.template modify<LMPHostType>();
k_sign.template modify<LMPHostType>();
k_multiplicity.template modify<LMPHostType>();
k_k.modify_host();
k_cos_shift.modify_host();
k_sin_shift.modify_host();
k_sign.modify_host();
k_multiplicity.modify_host();
}
/* ----------------------------------------------------------------------

6
src/KOKKOS/dihedral_harmonic_kokkos.h
View File

@ -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<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;

225
src/KOKKOS/dihedral_hybrid_kokkos.cpp Normal file
View File

@ -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 <cstring>
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;
}

58
src/KOKKOS/dihedral_hybrid_kokkos.h Normal file
View 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
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

25
src/KOKKOS/dihedral_opls_kokkos.cpp
View File

@ -39,6 +39,7 @@ static constexpr double SMALLER = 0.00001;
template<class DeviceType>
DihedralOPLSKokkos<DeviceType>::DihedralOPLSKokkos(LAMMPS *lmp) : DihedralOPLS(lmp)
{
kokkosable = 1;
atomKK = (AtomKokkos *) atom;
neighborKK = (NeighborKokkos *) neighbor;
execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
@ -100,7 +101,7 @@ void DihedralOPLSKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
newton_bond = force->newton_bond;
h_warning_flag() = 0;
k_warning_flag.template modify<LMPHostType>();
k_warning_flag.modify_host();
k_warning_flag.template sync<DeviceType>();
copymode = 1;
@ -126,7 +127,7 @@ void DihedralOPLSKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// error check
k_warning_flag.template modify<DeviceType>();
k_warning_flag.template sync<LMPHostType>();
k_warning_flag.sync_host();
if (h_warning_flag())
error->warning(FLERR,"Dihedral problem");
@ -142,12 +143,12 @@ void DihedralOPLSKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
k_eatom.sync_host();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
k_vatom.sync_host();
}
copymode = 0;
@ -372,10 +373,10 @@ void DihedralOPLSKokkos<DeviceType>::coeff(int narg, char **arg)
k_k4.h_view[i] = k4[i];
}
k_k1.template modify<LMPHostType>();
k_k2.template modify<LMPHostType>();
k_k3.template modify<LMPHostType>();
k_k4.template modify<LMPHostType>();
k_k1.modify_host();
k_k2.modify_host();
k_k3.modify_host();
k_k4.modify_host();
}
/* ----------------------------------------------------------------------
@ -395,10 +396,10 @@ void DihedralOPLSKokkos<DeviceType>::read_restart(FILE *fp)
k_k4.h_view[i] = k4[i];
}
k_k1.template modify<LMPHostType>();
k_k2.template modify<LMPHostType>();
k_k3.template modify<LMPHostType>();
k_k4.template modify<LMPHostType>();
k_k1.modify_host();
k_k2.modify_host();
k_k3.modify_host();
k_k4.modify_host();
}
/* ----------------------------------------------------------------------

7
src/KOKKOS/dihedral_opls_kokkos.h
View File

@ -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<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;

53
src/KOKKOS/improper_class2_kokkos.cpp
View File

@ -34,6 +34,7 @@ static constexpr double SMALL = 0.001;
template<class DeviceType>
ImproperClass2Kokkos<DeviceType>::ImproperClass2Kokkos(LAMMPS *lmp) : ImproperClass2(lmp)
{
kokkosable = 1;
atomKK = (AtomKokkos *) atom;
neighborKK = (NeighborKokkos *) neighbor;
execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
@ -110,7 +111,7 @@ void ImproperClass2Kokkos<DeviceType>::compute(int eflag_in, int vflag_in)
newton_bond = force->newton_bond;
h_warning_flag() = 0;
k_warning_flag.template modify<LMPHostType>();
k_warning_flag.modify_host();
k_warning_flag.template sync<DeviceType>();
copymode = 1;
@ -139,7 +140,7 @@ void ImproperClass2Kokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// error check
k_warning_flag.template modify<DeviceType>();
k_warning_flag.template sync<LMPHostType>();
k_warning_flag.sync_host();
if (h_warning_flag())
error->warning(FLERR,"Improper problem");
@ -171,12 +172,12 @@ void ImproperClass2Kokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
k_eatom.sync_host();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
k_vatom.sync_host();
}
copymode = 0;
@ -918,17 +919,17 @@ void ImproperClass2Kokkos<DeviceType>::coeff(int narg, char **arg)
k_setflag_aa.h_view[i] = setflag_aa[i];
}
k_k0.template modify<LMPHostType>();
k_chi0.template modify<LMPHostType>();
k_aa_k1.template modify<LMPHostType>();
k_aa_k2.template modify<LMPHostType>();
k_aa_k3.template modify<LMPHostType>();
k_aa_theta0_1.template modify<LMPHostType>();
k_aa_theta0_2.template modify<LMPHostType>();
k_aa_theta0_3 .template modify<LMPHostType>();
k_setflag.template modify<LMPHostType>();
k_setflag_i.template modify<LMPHostType>();
k_setflag_aa.template modify<LMPHostType>();
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<DeviceType>::read_restart(FILE *fp)
k_setflag_aa.h_view[i] = setflag_aa[i];
}
k_k0.template modify<LMPHostType>();
k_chi0.template modify<LMPHostType>();
k_aa_k1.template modify<LMPHostType>();
k_aa_k2.template modify<LMPHostType>();
k_aa_k3.template modify<LMPHostType>();
k_aa_theta0_1.template modify<LMPHostType>();
k_aa_theta0_2.template modify<LMPHostType>();
k_aa_theta0_3 .template modify<LMPHostType>();
k_setflag.template modify<LMPHostType>();
k_setflag_i.template modify<LMPHostType>();
k_setflag_aa.template modify<LMPHostType>();
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();
}
/* ----------------------------------------------------------------------

6
src/KOKKOS/improper_class2_kokkos.h
View File

@ -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<double*[3],typename AT::t_f_array::array_layout,typename KKDevice<DeviceType>::value,Kokkos::MemoryTraits<Kokkos::Atomic> > 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;

25
src/KOKKOS/improper_harmonic_kokkos.cpp
View File

@ -36,6 +36,7 @@ static constexpr double SMALL = 0.001;
template<class DeviceType>
ImproperHarmonicKokkos<DeviceType>::ImproperHarmonicKokkos(LAMMPS *lmp) : ImproperHarmonic(lmp)
{
kokkosable = 1;
atomKK = (AtomKokkos *) atom;
neighborKK = (NeighborKokkos *) neighbor;
execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
@ -73,18 +74,18 @@ void ImproperHarmonicKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// reallocate per-atom arrays if necessary
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->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom");
d_eatom = k_eatom.template view<KKDeviceType>();
//}
} else Kokkos::deep_copy(d_eatom,0.0);
}
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->create_kokkos(k_vatom,vatom,maxvatom,"improper:vatom");
d_vatom = k_vatom.template view<KKDeviceType>();
//}
} else Kokkos::deep_copy(d_vatom,0.0);
}
//atomKK->sync(execution_space,datamask_read);
@ -102,7 +103,7 @@ void ImproperHarmonicKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
newton_bond = force->newton_bond;
h_warning_flag() = 0;
k_warning_flag.template modify<LMPHostType>();
k_warning_flag.modify_host();
k_warning_flag.template sync<DeviceType>();
copymode = 1;
@ -128,7 +129,7 @@ void ImproperHarmonicKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// error check
k_warning_flag.template modify<DeviceType>();
k_warning_flag.template sync<LMPHostType>();
k_warning_flag.sync_host();
if (h_warning_flag())
error->warning(FLERR,"Dihedral problem");
@ -144,12 +145,12 @@ void ImproperHarmonicKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
k_eatom.sync_host();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
k_vatom.sync_host();
}
copymode = 0;
@ -324,8 +325,8 @@ void ImproperHarmonicKokkos<DeviceType>::coeff(int narg, char **arg)
k_chi.h_view[i] = chi[i];
}
k_k.template modify<LMPHostType>();
k_chi.template modify<LMPHostType>();
k_k.modify_host();
k_chi.modify_host();
}
/* ----------------------------------------------------------------------
@ -343,8 +344,8 @@ void ImproperHarmonicKokkos<DeviceType>::read_restart(FILE *fp)
k_chi.h_view[i] = chi[i];
}
k_k.template modify<LMPHostType>();
k_chi.template modify<LMPHostType>();
k_k.modify_host();
k_chi.modify_host();
}
/* ----------------------------------------------------------------------

8
src/KOKKOS/improper_harmonic_kokkos.h
View File

@ -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<DeviceType>::value KKDeviceType;
Kokkos::DualView<E_FLOAT*,Kokkos::LayoutRight,KKDeviceType> k_eatom;
Kokkos::DualView<F_FLOAT*[6],Kokkos::LayoutRight,KKDeviceType> k_vatom;
protected:
class NeighborKokkos *neighborKK;
typedef typename KKDevice<DeviceType>::value KKDeviceType;
typename AT::t_x_array_randomread x;
typename Kokkos::View<double*[3],typename AT::t_f_array::array_layout,KKDeviceType,Kokkos::MemoryTraits<Kokkos::Atomic> > f;
typename AT::t_int_2d improperlist;
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<F_FLOAT*[6],Kokkos::LayoutRight,KKDeviceType,Kokkos::MemoryTraits<Kokkos::Atomic> > d_vatom;

226
src/KOKKOS/improper_hybrid_kokkos.cpp Normal file
View File

@ -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 <cstring>
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;
}

58
src/KOKKOS/improper_hybrid_kokkos.h Normal file
View 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
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

2
src/angle.cpp
View File

@ -50,7 +50,7 @@ Angle::Angle(LAMMPS *_lmp) : Pointers(_lmp)
datamask_read = ALL_MASK;
datamask_modify = ALL_MASK;
copymode = 0;
copymode = kokkosable = 0;
}
/* ---------------------------------------------------------------------- */

2
src/angle.h
View File

@ -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;

37
src/angle_hybrid.cpp
View File

@ -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);
}
}

6
src/angle_hybrid.h
View File

@ -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

2
src/dihedral.cpp
View File

@ -48,7 +48,7 @@ Dihedral::Dihedral(LAMMPS *_lmp) : Pointers(_lmp)
datamask_read = ALL_MASK;
datamask_modify = ALL_MASK;
copymode = 0;
copymode = kokkosable = 0;
}
/* ---------------------------------------------------------------------- */

2
src/dihedral.h
View File

@ -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;

35
src/dihedral_hybrid.cpp
View File

@ -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);
}
}

5
src/dihedral_hybrid.h
View File

@ -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

2
src/improper.cpp
View File

@ -47,7 +47,7 @@ Improper::Improper(LAMMPS *_lmp) : Pointers(_lmp)
datamask_read = ALL_MASK;
datamask_modify = ALL_MASK;
copymode = 0;
copymode = kokkosable = 0;
}
/* ---------------------------------------------------------------------- */

2
src/improper.h
View File

@ -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;

37
src/improper_hybrid.cpp
View File

@ -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);
}
}

5
src/improper_hybrid.h
View File

@ -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