jianmu
/
lammps
forked from lijiext/lammps
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Merge branch 'master' into fft-cleanup

This commit is contained in:
Axel Kohlmeyer 2017-02-03 16:53:10 -05:00
parent 9d1991bf84 207adc3968
commit 191453e1c7
91 changed files with 1410 additions and 523 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
doc/src/pair_morse.txt
View File

@ -13,6 +13,7 @@ pair_style morse/opt command :h3
pair_style morse/smooth/linear command :h3
pair_style morse/smooth/linear/omp command :h3
pair_style morse/soft command :h3
pair_style morse/kk command :h3
[Syntax:]

2
src/KOKKOS/Install.sh
View File

@ -173,6 +173,8 @@ action pair_lj_gromacs_kokkos.cpp
action pair_lj_gromacs_kokkos.h
action pair_lj_sdk_kokkos.cpp pair_lj_sdk.cpp
action pair_lj_sdk_kokkos.h pair_lj_sdk.h
action pair_morse_kokkos.cpp
action pair_morse_kokkos.h
action pair_reax_c_kokkos.cpp pair_reax_c.cpp
action pair_reax_c_kokkos.h pair_reax_c.h
action pair_sw_kokkos.cpp pair_sw.cpp

12
src/KOKKOS/angle_charmm_kokkos.cpp
View File

@ -70,18 +70,18 @@ void AngleCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// reallocate per-atom arrays if necessary
if (eflag_atom) {
if(k_eatom.dimension_0()<maxeatom) {
//if(k_eatom.dimension_0()<maxeatom) { // won't work without adding zero functor
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom");
d_eatom = k_eatom.d_view;
}
d_eatom = k_eatom.template view<DeviceType>();
//}
}
if (vflag_atom) {
if(k_vatom.dimension_0()<maxvatom) {
//if(k_vatom.dimension_0()<maxvatom) { // won't work without adding zero functor
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom");
d_vatom = k_vatom.d_view;
}
d_vatom = k_vatom.template view<DeviceType>();
//}
}
x = atomKK->k_x.view<DeviceType>();

4
src/KOKKOS/atom_vec_angle_kokkos.cpp
View File

@ -957,10 +957,10 @@ struct AtomVecAngleKokkos_UnpackBorder {
_x(i+_first,0) = _buf(i,0);
_x(i+_first,1) = _buf(i,1);
_x(i+_first,2) = _buf(i,2);
_tag(i+_first) = static_cast<int> (_buf(i,3));
_tag(i+_first) = static_cast<tagint> (_buf(i,3));
_type(i+_first) = static_cast<int> (_buf(i,4));
_mask(i+_first) = static_cast<int> (_buf(i,5));
_molecule(i+_first) = static_cast<int> (_buf(i,6));
_molecule(i+_first) = static_cast<tagint> (_buf(i,6));
}
};

2
src/KOKKOS/atom_vec_atomic_kokkos.cpp
View File

@ -836,7 +836,7 @@ struct AtomVecAtomicKokkos_UnpackBorder {
_x(i+_first,0) = _buf(i,0);
_x(i+_first,1) = _buf(i,1);
_x(i+_first,2) = _buf(i,2);
_tag(i+_first) = static_cast<int> (_buf(i,3));
_tag(i+_first) = static_cast<tagint> (_buf(i,3));
_type(i+_first) = static_cast<int> (_buf(i,4));
_mask(i+_first) = static_cast<int> (_buf(i,5));
// printf("%i %i %lf %lf %lf %i BORDER\n",_tag(i+_first),i+_first,_x(i+_first,0),_x(i+_first,1),_x(i+_first,2),_type(i+_first));

4
src/KOKKOS/atom_vec_bond_kokkos.cpp
View File

@ -905,10 +905,10 @@ struct AtomVecBondKokkos_UnpackBorder {
_x(i+_first,0) = _buf(i,0);
_x(i+_first,1) = _buf(i,1);
_x(i+_first,2) = _buf(i,2);
_tag(i+_first) = static_cast<int> (_buf(i,3));
_tag(i+_first) = static_cast<tagint> (_buf(i,3));
_type(i+_first) = static_cast<int> (_buf(i,4));
_mask(i+_first) = static_cast<int> (_buf(i,5));
_molecule(i+_first) = static_cast<int> (_buf(i,6));
_molecule(i+_first) = static_cast<tagint> (_buf(i,6));
}
};

2
src/KOKKOS/atom_vec_charge_kokkos.cpp
View File

@ -872,7 +872,7 @@ struct AtomVecChargeKokkos_UnpackBorder {
_x(i+_first,0) = _buf(i,0);
_x(i+_first,1) = _buf(i,1);
_x(i+_first,2) = _buf(i,2);
_tag(i+_first) = static_cast<int> (_buf(i,3));
_tag(i+_first) = static_cast<tagint> (_buf(i,3));
_type(i+_first) = static_cast<int> (_buf(i,4));
_mask(i+_first) = static_cast<int> (_buf(i,5));
_q(i+_first) = _buf(i,6);

4
src/KOKKOS/atom_vec_molecular_kokkos.cpp
View File

@ -1029,10 +1029,10 @@ struct AtomVecMolecularKokkos_UnpackBorder {
_x(i+_first,0) = _buf(i,0);
_x(i+_first,1) = _buf(i,1);
_x(i+_first,2) = _buf(i,2);
_tag(i+_first) = static_cast<int> (_buf(i,3));
_tag(i+_first) = static_cast<tagint> (_buf(i,3));
_type(i+_first) = static_cast<int> (_buf(i,4));
_mask(i+_first) = static_cast<int> (_buf(i,5));
_molecule(i+_first) = static_cast<int> (_buf(i,6));
_molecule(i+_first) = static_cast<tagint> (_buf(i,6));
}
};

4
src/KOKKOS/bond_fene_kokkos.cpp
View File

@ -77,12 +77,12 @@ void BondFENEKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"bond:eatom");
d_eatom = k_eatom.d_view;
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"bond:vatom");
d_vatom = k_vatom.d_view;
d_vatom = k_vatom.view<DeviceType>();
}
k_k.template sync<DeviceType>();

4
src/KOKKOS/bond_fene_kokkos.h
View File

@ -66,8 +66,8 @@ class BondFENEKokkos : public BondFENE {
DAT::tdual_efloat_1d k_eatom;
DAT::tdual_virial_array k_vatom;
DAT::t_efloat_1d d_eatom;
DAT::t_virial_array d_vatom;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
DAT::tdual_int_scalar k_warning_flag;
typename AT::t_int_scalar d_warning_flag;

12
src/KOKKOS/bond_harmonic_kokkos.cpp
View File

@ -67,18 +67,18 @@ void BondHarmonicKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// reallocate per-atom arrays if necessary
if (eflag_atom) {
if(k_eatom.dimension_0()<maxeatom) {
//if(k_eatom.dimension_0()<maxeatom) { // won't work without adding zero functor
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom");
d_eatom = k_eatom.d_view;
}
d_eatom = k_eatom.template view<DeviceType>();
//}
}
if (vflag_atom) {
if(k_vatom.dimension_0()<maxvatom) {
//if(k_vatom.dimension_0()<maxvatom) { // won't work without adding zero functor
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom");
d_vatom = k_vatom.d_view;
}
d_vatom = k_vatom.template view<DeviceType>();
//}
}
// if (eflag || vflag) atomKK->modified(execution_space,datamask_modify);

16
src/KOKKOS/dihedral_charmm_kokkos.cpp
View File

@ -80,22 +80,22 @@ void DihedralCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// reallocate per-atom arrays if necessary
if (eflag_atom) {
if(k_eatom.dimension_0()<maxeatom) {
//if(k_eatom.dimension_0()<maxeatom) { // won't work without adding zero functor
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"dihedral:eatom");
d_eatom = k_eatom.d_view;
d_eatom = k_eatom.template view<DeviceType>();
k_eatom_pair = Kokkos::DualView<E_FLOAT*,Kokkos::LayoutRight,DeviceType>("dihedral:eatom_pair",maxeatom);
d_eatom_pair = k_eatom.d_view;
}
d_eatom_pair = k_eatom.template view<DeviceType>();
//}
}
if (vflag_atom) {
if(k_vatom.dimension_0()<maxvatom) {
//if(k_vatom.dimension_0()<maxvatom) { // won't work without adding zero functor
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"dihedral:vatom");
d_vatom = k_vatom.d_view;
d_vatom = k_vatom.template view<DeviceType>();
k_vatom_pair = Kokkos::DualView<F_FLOAT*[6],Kokkos::LayoutRight,DeviceType>("dihedral:vatom_pair",maxvatom);
d_vatom_pair = k_vatom.d_view;
}
d_vatom_pair = k_vatom.template view<DeviceType>();
//}
}
x = atomKK->k_x.view<DeviceType>();

4
src/KOKKOS/dihedral_opls_kokkos.cpp
View File

@ -77,12 +77,12 @@ void DihedralOPLSKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"dihedral:eatom");
d_eatom = k_eatom.d_view;
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"dihedral:vatom");
d_vatom = k_vatom.d_view;
d_vatom = k_vatom.view<DeviceType>();
}
k_k1.template sync<DeviceType>();

4
src/KOKKOS/dihedral_opls_kokkos.h
View File

@ -68,8 +68,8 @@ class DihedralOPLSKokkos : public DihedralOPLS {
DAT::tdual_efloat_1d k_eatom;
DAT::tdual_virial_array k_vatom;
DAT::t_efloat_1d d_eatom;
DAT::t_virial_array d_vatom;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
int nlocal,newton_bond;
int eflag,vflag;

4
src/KOKKOS/domain_kokkos.cpp
View File

@ -250,7 +250,7 @@ struct DomainPBCFunctor {
x(i,0) += period[0];
if (DEFORM_VREMAP && (mask[i] & deform_groupbit)) v(i,0) += h_rate[0];
imageint idim = image[i] & IMGMASK;
const int otherdims = image[i] ^ idim;
const imageint otherdims = image[i] ^ idim;
idim--;
idim &= IMGMASK;
image[i] = otherdims | idim;
@ -260,7 +260,7 @@ struct DomainPBCFunctor {
x(i,0) = MAX(x(i,0),lo[0]);
if (DEFORM_VREMAP && (mask[i] & deform_groupbit)) v(i,0) -= h_rate[0];
imageint idim = image[i] & IMGMASK;
const int otherdims = image[i] ^ idim;
const imageint otherdims = image[i] ^ idim;
idim++;
idim &= IMGMASK;
image[i] = otherdims | idim;

10
src/KOKKOS/fix_qeq_reax_kokkos.cpp
View File

@ -387,7 +387,7 @@ KOKKOS_INLINE_FUNCTION
void FixQEqReaxKokkos<DeviceType>::compute_h_item(int ii, int &m_fill, const bool &final) const
{
const int i = d_ilist[ii];
int j,jj,jtag,jtype,flag;
int j,jj,jtype,flag;
if (mask[i] & groupbit) {
@ -395,7 +395,7 @@ void FixQEqReaxKokkos<DeviceType>::compute_h_item(int ii, int &m_fill, const boo
const X_FLOAT ytmp = x(i,1);
const X_FLOAT ztmp = x(i,2);
const int itype = type(i);
const int itag = tag(i);
const tagint itag = tag(i);
const int jnum = d_numneigh[i];
if (final)
d_firstnbr[i] = m_fill;
@ -403,7 +403,6 @@ void FixQEqReaxKokkos<DeviceType>::compute_h_item(int ii, int &m_fill, const boo
for (jj = 0; jj < jnum; jj++) {
j = d_neighbors(i,jj);
j &= NEIGHMASK;
jtype = type(j);
const X_FLOAT delx = x(j,0) - xtmp;
@ -411,10 +410,11 @@ void FixQEqReaxKokkos<DeviceType>::compute_h_item(int ii, int &m_fill, const boo
const X_FLOAT delz = x(j,2) - ztmp;
if (neighflag != FULL) {
const tagint jtag = tag(j);
flag = 0;
if (j < nlocal) flag = 1;
else if (tag[i] < tag[j]) flag = 1;
else if (tag[i] == tag[j]) {
else if (itag < jtag) flag = 1;
else if (itag == jtag) {
if (delz > SMALL) flag = 1;
else if (fabs(delz) < SMALL) {
if (dely > SMALL) flag = 1;

3
src/KOKKOS/fix_qeq_reax_kokkos.h
View File

@ -159,7 +159,8 @@ class FixQEqReaxKokkos : public FixQEqReax {
//typename ArrayTypes<DeviceType>::t_float_1d_randomread mass, q;
typename ArrayTypes<DeviceType>::t_float_1d_randomread mass;
typename ArrayTypes<DeviceType>::t_float_1d q;
typename ArrayTypes<DeviceType>::t_int_1d type, tag, mask;
typename ArrayTypes<DeviceType>::t_int_1d type, mask;
typename ArrayTypes<DeviceType>::t_tagint_1d tag;
DAT::tdual_float_1d k_q;
typename AT::t_float_1d d_q;

1
src/KOKKOS/fix_reaxc_bonds_kokkos.cpp
View File

@ -76,7 +76,6 @@ void FixReaxCBondsKokkos::init()
void FixReaxCBondsKokkos::Output_ReaxC_Bonds(bigint ntimestep, FILE *fp)
{
int i, j;
int nbuf_local;
int nlocal_max, numbonds, numbonds_max;
double *buf;

12
src/KOKKOS/improper_harmonic_kokkos.cpp
View File

@ -77,18 +77,18 @@ void ImproperHarmonicKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// reallocate per-atom arrays if necessary
if (eflag_atom) {
if(k_eatom.dimension_0()<maxeatom) {
//if(k_eatom.dimension_0()<maxeatom) { // won't work without adding zero functor
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom");
d_eatom = k_eatom.d_view;
}
d_eatom = k_eatom.template view<DeviceType>();
//}
}
if (vflag_atom) {
if(k_vatom.dimension_0()<maxvatom) {
//if(k_vatom.dimension_0()<maxvatom) { // won't work without adding zero functor
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom");
d_vatom = k_vatom.d_view;
}
d_vatom = k_vatom.template view<DeviceType>();
//}
}
//atomKK->sync(execution_space,datamask_read);

5
src/KOKKOS/npair_kokkos.h
View File

@ -157,7 +157,8 @@ class NeighborKokkosExecute
// data from Atom class
const typename AT::t_x_array_randomread x;
const typename AT::t_int_1d_const type,mask,molecule;
const typename AT::t_int_1d_const type,mask;
const typename AT::t_tagint_1d_const molecule;
const typename AT::t_tagint_1d_const tag;
const typename AT::t_tagint_2d_const special;
const typename AT::t_int_2d_const nspecial;
@ -194,7 +195,7 @@ class NeighborKokkosExecute
const typename AT::t_x_array_randomread &_x,
const typename AT::t_int_1d_const &_type,
const typename AT::t_int_1d_const &_mask,
const typename AT::t_int_1d_const &_molecule,
const typename AT::t_tagint_1d_const &_molecule,
const typename AT::t_tagint_1d_const &_tag,
const typename AT::t_tagint_2d_const &_special,
const typename AT::t_int_2d_const &_nspecial,

2
src/KOKKOS/pair_buck_coul_cut_kokkos.cpp
View File

@ -154,7 +154,7 @@ void PairBuckCoulCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev.v[5];
}
if (vflag_fdotr) virial_fdotr_compute();
if (vflag_fdotr) pair_virial_fdotr_compute(this);
if (eflag_atom) {
k_eatom.template modify<DeviceType>();

28
src/KOKKOS/pair_buck_coul_cut_kokkos.h
View File

@ -34,6 +34,7 @@ class PairBuckCoulCutKokkos : public PairBuckCoulCut {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
enum {COUL_FLAG=1};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairBuckCoulCutKokkos(class LAMMPS *);
~PairBuckCoulCutKokkos();
@ -83,25 +84,25 @@ class PairBuckCoulCutKokkos : public PairBuckCoulCut {
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
typename AT::t_x_array_randomread x;
typename AT::t_x_array c_x;
typename AT::t_f_array f;
typename AT::t_int_1d_randomread type;
typename AT::t_float_1d_randomread q;
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;
typename AT::t_efloat_1d d_eatom;
typename AT::t_virial_array d_vatom;
int newton_pair;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
typename AT::tdual_ffloat_2d k_cutsq;
typename AT::t_ffloat_2d d_cutsq;
typename AT::tdual_ffloat_2d k_cut_ljsq;
typename AT::t_ffloat_2d d_cut_ljsq;
typename AT::tdual_ffloat_2d k_cut_coulsq;
typename AT::t_ffloat_2d d_cut_coulsq;
int neighflag;
@ -123,6 +124,7 @@ class PairBuckCoulCutKokkos : public PairBuckCoulCut {
friend EV_FLOAT pair_compute_neighlist<PairBuckCoulCutKokkos,HALFTHREAD,void>(PairBuckCoulCutKokkos*,NeighListKokkos<DeviceType>*);
friend EV_FLOAT pair_compute<PairBuckCoulCutKokkos,void>(PairBuckCoulCutKokkos*,
NeighListKokkos<DeviceType>*);
friend void pair_virial_fdotr_compute<PairBuckCoulCutKokkos>(PairBuckCoulCutKokkos*);
};

13
src/KOKKOS/pair_buck_coul_long_kokkos.cpp
View File

@ -114,6 +114,19 @@ void PairBuckCoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag || vflag) ev_setup(eflag,vflag);
else evflag = vflag_fdotr = 0;
// reallocate per-atom arrays if necessary
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
k_cutsq.template sync<DeviceType>();
k_cut_ljsq.template sync<DeviceType>();

29
src/KOKKOS/pair_buck_coul_long_kokkos.h
View File

@ -34,6 +34,7 @@ class PairBuckCoulLongKokkos : public PairBuckCoulLong {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
enum {COUL_FLAG=1};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairBuckCoulLongKokkos(class LAMMPS *);
~PairBuckCoulLongKokkos();
@ -84,27 +85,27 @@ class PairBuckCoulLongKokkos : public PairBuckCoulLong {
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
typename AT::t_x_array_randomread x;
typename AT::t_x_array c_x;
typename AT::t_f_array f;
typename AT::t_int_1d_randomread type;
typename AT::t_float_1d_randomread q;
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;
typename AT::t_efloat_1d d_eatom;
typename AT::t_virial_array d_vatom;
int newton_pair;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
typename AT::tdual_ffloat_2d k_cutsq;
typename AT::t_ffloat_2d d_cutsq;
typename AT::tdual_ffloat_2d k_cut_ljsq;
typename AT::t_ffloat_2d d_cut_ljsq;
typename AT::tdual_ffloat_2d k_cut_coulsq;
typename AT::t_ffloat_2d d_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_1d_randomread
typename AT::t_ffloat_1d_randomread
d_rtable, d_drtable, d_ftable, d_dftable,
d_ctable, d_dctable, d_etable, d_detable;

19
src/KOKKOS/pair_buck_kokkos.h
View File

@ -34,6 +34,7 @@ class PairBuckKokkos : public PairBuck {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF|N2};
enum {COUL_FLAG=0};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairBuckKokkos(class LAMMPS *);
~PairBuckKokkos();
@ -72,22 +73,22 @@ class PairBuckKokkos : public PairBuck {
typename Kokkos::DualView<params_buck**,Kokkos::LayoutRight,DeviceType>::t_dev_const_um params;
params_buck m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; // hardwired to space for 12 atom types
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename ArrayTypes<DeviceType>::t_tagint_1d tag;
typename AT::t_x_array_randomread x;
typename AT::t_x_array c_x;
typename AT::t_f_array f;
typename AT::t_int_1d_randomread type;
typename AT::t_tagint_1d tag;
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;
typename AT::t_efloat_1d d_eatom;
typename AT::t_virial_array d_vatom;
int newton_pair;
double special_lj[4];
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename AT::tdual_ffloat_2d k_cutsq;
typename AT::t_ffloat_2d d_cutsq;
int neighflag;

23
src/KOKKOS/pair_coul_cut_kokkos.cpp
View File

@ -83,6 +83,19 @@ void PairCoulCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag || vflag) ev_setup(eflag,vflag);
else evflag = vflag_fdotr = 0;
// reallocate per-atom arrays if necessary
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
k_cutsq.template sync<DeviceType>();
k_cut_ljsq.template sync<DeviceType>();
@ -124,6 +137,16 @@ void PairCoulCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev.v[5];
}
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
}

30
src/KOKKOS/pair_coul_cut_kokkos.h
View File

@ -34,6 +34,7 @@ class PairCoulCutKokkos : public PairCoulCut {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
enum {COUL_FLAG=1};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairCoulCutKokkos(class LAMMPS *);
~PairCoulCutKokkos();
@ -87,22 +88,25 @@ class PairCoulCutKokkos : public PairCoulCut {
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
typename AT::t_x_array_randomread x;
typename AT::t_x_array c_x;
typename AT::t_f_array f;
typename AT::t_float_1d_randomread q;
typename AT::t_int_1d_randomread type;
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;
int newton_pair;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
typename AT::tdual_ffloat_2d k_cutsq;
typename AT::t_ffloat_2d d_cutsq;
typename AT::tdual_ffloat_2d k_cut_ljsq;
typename AT::t_ffloat_2d d_cut_ljsq;
typename AT::tdual_ffloat_2d k_cut_coulsq;
typename AT::t_ffloat_2d d_cut_coulsq;
int neighflag;

25
src/KOKKOS/pair_coul_debye_kokkos.cpp
View File

@ -90,6 +90,19 @@ void PairCoulDebyeKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag || vflag) ev_setup(eflag,vflag);
else evflag = vflag_fdotr = 0;
// reallocate per-atom arrays if necessary
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
k_cutsq.template sync<DeviceType>();
k_cut_ljsq.template sync<DeviceType>();
@ -136,7 +149,17 @@ void PairCoulDebyeKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev.v[5];
}
if (vflag_fdotr) virial_fdotr_compute();
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
copymode = 0;
}

31
src/KOKKOS/pair_coul_debye_kokkos.h
View File

@ -34,6 +34,7 @@ class PairCoulDebyeKokkos : public PairCoulDebye {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
enum {COUL_FLAG=1};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairCoulDebyeKokkos(class LAMMPS *);
~PairCoulDebyeKokkos();
@ -87,22 +88,25 @@ class PairCoulDebyeKokkos : public PairCoulDebye {
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
typename AT::t_x_array_randomread x;
typename AT::t_x_array c_x;
typename AT::t_f_array f;
typename AT::t_int_1d_randomread type;
typename AT::t_float_1d_randomread q;
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;
int newton_pair;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
typename AT::tdual_ffloat_2d k_cutsq;
typename AT::t_ffloat_2d d_cutsq;
typename AT::tdual_ffloat_2d k_cut_ljsq;
typename AT::t_ffloat_2d d_cut_ljsq;
typename AT::tdual_ffloat_2d k_cut_coulsq;
typename AT::t_ffloat_2d d_cut_coulsq;
int neighflag;
@ -124,6 +128,7 @@ class PairCoulDebyeKokkos : public PairCoulDebye {
friend EV_FLOAT pair_compute_neighlist<PairCoulDebyeKokkos,HALFTHREAD,void>(PairCoulDebyeKokkos*,NeighListKokkos<DeviceType>*);
friend EV_FLOAT pair_compute<PairCoulDebyeKokkos,void>(PairCoulDebyeKokkos*,
NeighListKokkos<DeviceType>*);
friend void pair_virial_fdotr_compute<PairCoulDebyeKokkos>(PairCoulDebyeKokkos*);
};

8
src/KOKKOS/pair_coul_dsf_kokkos.cpp
View File

@ -88,12 +88,12 @@ void PairCoulDSFKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.d_view;
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.d_view;
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
@ -183,8 +183,6 @@ void PairCoulDSFKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev.v[5];
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
@ -195,6 +193,8 @@ void PairCoulDSFKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
copymode = 0;
}

19
src/KOKKOS/pair_coul_dsf_kokkos.h
View File

@ -37,6 +37,7 @@ class PairCoulDSFKokkos : public PairCoulDSF {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
enum {COUL_FLAG=1};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
typedef EV_FLOAT value_type;
PairCoulDSFKokkos(class LAMMPS *);
~PairCoulDSFKokkos();
@ -61,16 +62,16 @@ class PairCoulDSFKokkos : public PairCoulDSF {
KOKKOS_INLINE_FUNCTION
int sbmask(const int& j) const;
protected:
typename AT::t_x_array_randomread x;
typename AT::t_f_array f;
typename AT::t_float_1d_randomread q;
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
protected:
DAT::tdual_efloat_1d k_eatom;
DAT::tdual_virial_array k_vatom;
DAT::t_efloat_1d d_eatom;
DAT::t_virial_array d_vatom;
typename AT::t_efloat_1d d_eatom;
typename AT::t_virial_array d_vatom;
int neighflag,newton_pair;
@ -79,9 +80,9 @@ class PairCoulDSFKokkos : public PairCoulDSF {
double special_coul[4];
double qqrd2e;
typename ArrayTypes<DeviceType>::t_neighbors_2d d_neighbors;
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_ilist;
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_numneigh;
typename AT::t_neighbors_2d d_neighbors;
typename AT::t_int_1d_randomread d_ilist;
typename AT::t_int_1d_randomread d_numneigh;
//NeighListKokkos<DeviceType> k_list;
friend void pair_virial_fdotr_compute<PairCoulDSFKokkos>(PairCoulDSFKokkos*);

23
src/KOKKOS/pair_coul_long_kokkos.cpp
View File

@ -107,6 +107,19 @@ void PairCoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag || vflag) ev_setup(eflag,vflag);
else evflag = vflag_fdotr = 0;
// reallocate per-atom arrays if necessary
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
k_cutsq.template sync<DeviceType>();
k_cut_coulsq.template sync<DeviceType>();
@ -158,6 +171,16 @@ void PairCoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev.v[5];
}
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
copymode = 0;

31
src/KOKKOS/pair_coul_long_kokkos.h
View File

@ -34,6 +34,7 @@ class PairCoulLongKokkos : public PairCoulLong {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
enum {COUL_FLAG=1};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairCoulLongKokkos(class LAMMPS *);
~PairCoulLongKokkos();
@ -45,7 +46,9 @@ class PairCoulLongKokkos : public PairCoulLong {
double init_one(int, int);
struct params_coul{
KOKKOS_INLINE_FUNCTION
params_coul(){cut_coulsq=0;};
KOKKOS_INLINE_FUNCTION
params_coul(int i){cut_coulsq=0;};
F_FLOAT cut_coulsq;
};
@ -86,27 +89,27 @@ class PairCoulLongKokkos : public PairCoulLong {
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
typename AT::t_x_array_randomread x;
typename AT::t_x_array c_x;
typename AT::t_f_array f;
typename AT::t_int_1d_randomread type;
typename AT::t_float_1d_randomread q;
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;
typename AT::t_efloat_1d d_eatom;
typename AT::t_virial_array d_vatom;
int newton_pair;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
typename AT::tdual_ffloat_2d k_cutsq;
typename AT::t_ffloat_2d d_cutsq;
typename AT::tdual_ffloat_2d k_cut_ljsq;
typename AT::t_ffloat_2d d_cut_ljsq;
typename AT::tdual_ffloat_2d k_cut_coulsq;
typename AT::t_ffloat_2d d_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_1d_randomread
typename AT::t_ffloat_1d_randomread
d_rtable, d_drtable, d_ftable, d_dftable,
d_ctable, d_dctable, d_etable, d_detable;

8
src/KOKKOS/pair_coul_wolf_kokkos.cpp
View File

@ -83,12 +83,12 @@ void PairCoulWolfKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.d_view;
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.d_view;
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
@ -184,8 +184,6 @@ void PairCoulWolfKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev.v[5];
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
@ -196,6 +194,8 @@ void PairCoulWolfKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
copymode = 0;
}

17
src/KOKKOS/pair_coul_wolf_kokkos.h
View File

@ -37,6 +37,7 @@ class PairCoulWolfKokkos : public PairCoulWolf {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
enum {COUL_FLAG=1};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
typedef EV_FLOAT value_type;
PairCoulWolfKokkos(class LAMMPS *);
~PairCoulWolfKokkos();
@ -63,14 +64,14 @@ class PairCoulWolfKokkos : public PairCoulWolf {
protected:
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
typename AT::t_x_array_randomread x;
typename AT::t_f_array f;
typename AT::t_float_1d_randomread q;
DAT::tdual_efloat_1d k_eatom;
DAT::tdual_virial_array k_vatom;
DAT::t_efloat_1d d_eatom;
DAT::t_virial_array d_vatom;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
int neighflag,newton_pair;
@ -81,9 +82,9 @@ class PairCoulWolfKokkos : public PairCoulWolf {
double special_coul[4];
double qqrd2e;
typename ArrayTypes<DeviceType>::t_neighbors_2d d_neighbors;
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_ilist;
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_numneigh;
typename AT::t_neighbors_2d d_neighbors;
typename AT::t_int_1d_randomread d_ilist;
typename AT::t_int_1d_randomread d_numneigh;
//NeighListKokkos<DeviceType> k_list;
friend void pair_virial_fdotr_compute<PairCoulWolfKokkos>(PairCoulWolfKokkos*);

8
src/KOKKOS/pair_eam_alloy_kokkos.cpp
View File

@ -82,12 +82,12 @@ void PairEAMAlloyKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.d_view;
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.d_view;
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
@ -246,8 +246,6 @@ void PairEAMAlloyKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev.v[5];
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
@ -258,6 +256,8 @@ void PairEAMAlloyKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
copymode = 0;
}

10
src/KOKKOS/pair_eam_alloy_kokkos.h
View File

@ -125,8 +125,8 @@ class PairEAMAlloyKokkos : public PairEAM {
DAT::tdual_efloat_1d k_eatom;
DAT::tdual_virial_array k_vatom;
DAT::t_efloat_1d d_eatom;
DAT::t_virial_array d_vatom;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
DAT::tdual_ffloat_1d k_rho;
DAT::tdual_ffloat_1d k_fp;
@ -154,9 +154,9 @@ class PairEAMAlloyKokkos : public PairEAM {
void interpolate(int, double, double *, t_host_ffloat_2d_n7, int);
void read_file(char *);
typename ArrayTypes<DeviceType>::t_neighbors_2d d_neighbors;
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_ilist;
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_numneigh;
typename AT::t_neighbors_2d d_neighbors;
typename AT::t_int_1d_randomread d_ilist;
typename AT::t_int_1d_randomread d_numneigh;
//NeighListKokkos<DeviceType> k_list;
int iswap;

4
src/KOKKOS/pair_eam_fs_kokkos.cpp
View File

@ -82,12 +82,12 @@ void PairEAMFSKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.d_view;
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.d_view;
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);

10
src/KOKKOS/pair_eam_fs_kokkos.h
View File

@ -125,8 +125,8 @@ class PairEAMFSKokkos : public PairEAM {
DAT::tdual_efloat_1d k_eatom;
DAT::tdual_virial_array k_vatom;
DAT::t_efloat_1d d_eatom;
DAT::t_virial_array d_vatom;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
DAT::tdual_ffloat_1d k_rho;
DAT::tdual_ffloat_1d k_fp;
@ -154,9 +154,9 @@ class PairEAMFSKokkos : public PairEAM {
void interpolate(int, double, double *, t_host_ffloat_2d_n7, int);
void read_file(char *);
typename ArrayTypes<DeviceType>::t_neighbors_2d d_neighbors;
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_ilist;
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_numneigh;
typename AT::t_neighbors_2d d_neighbors;
typename AT::t_int_1d_randomread d_ilist;
typename AT::t_int_1d_randomread d_numneigh;
//NeighListKokkos<DeviceType> k_list;
int iswap;

4
src/KOKKOS/pair_eam_kokkos.cpp
View File

@ -77,12 +77,12 @@ void PairEAMKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.d_view;
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.d_view;
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);

10
src/KOKKOS/pair_eam_kokkos.h
View File

@ -122,8 +122,8 @@ class PairEAMKokkos : public PairEAM {
DAT::tdual_efloat_1d k_eatom;
DAT::tdual_virial_array k_vatom;
DAT::t_efloat_1d d_eatom;
DAT::t_virial_array d_vatom;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
DAT::tdual_ffloat_1d k_rho;
DAT::tdual_ffloat_1d k_fp;
@ -149,9 +149,9 @@ class PairEAMKokkos : public PairEAM {
virtual void file2array();
void array2spline();
typename ArrayTypes<DeviceType>::t_neighbors_2d d_neighbors;
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_ilist;
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_numneigh;
typename AT::t_neighbors_2d d_neighbors;
typename AT::t_int_1d_randomread d_ilist;
typename AT::t_int_1d_randomread d_numneigh;
//NeighListKokkos<DeviceType> k_list;
int iswap;

23
src/KOKKOS/pair_lj_charmm_coul_charmm_implicit_kokkos.cpp
View File

@ -115,6 +115,19 @@ void PairLJCharmmCoulCharmmImplicitKokkos<DeviceType>::compute(int eflag_in, int
if (eflag || vflag) ev_setup(eflag,vflag);
else evflag = vflag_fdotr = 0;
// reallocate per-atom arrays if necessary
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
k_cutsq.template sync<DeviceType>();
k_cut_ljsq.template sync<DeviceType>();
@ -167,6 +180,16 @@ void PairLJCharmmCoulCharmmImplicitKokkos<DeviceType>::compute(int eflag_in, int
virial[5] += ev.v[5];
}
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
copymode = 0;

29
src/KOKKOS/pair_lj_charmm_coul_charmm_implicit_kokkos.h
View File

@ -34,6 +34,7 @@ class PairLJCharmmCoulCharmmImplicitKokkos : public PairLJCharmmCoulCharmmImplic
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
enum {COUL_FLAG=1};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairLJCharmmCoulCharmmImplicitKokkos(class LAMMPS *);
~PairLJCharmmCoulCharmmImplicitKokkos();
@ -77,27 +78,27 @@ class PairLJCharmmCoulCharmmImplicitKokkos : public PairLJCharmmCoulCharmmImplic
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
typename AT::t_x_array_randomread x;
typename AT::t_x_array c_x;
typename AT::t_f_array f;
typename AT::t_int_1d_randomread type;
typename AT::t_float_1d_randomread q;
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;
typename AT::t_efloat_1d d_eatom;
typename AT::t_virial_array d_vatom;
int newton_pair;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
typename AT::tdual_ffloat_2d k_cutsq;
typename AT::t_ffloat_2d d_cutsq;
typename AT::tdual_ffloat_2d k_cut_ljsq;
typename AT::t_ffloat_2d d_cut_ljsq;
typename AT::tdual_ffloat_2d k_cut_coulsq;
typename AT::t_ffloat_2d d_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_1d_randomread
typename AT::t_ffloat_1d_randomread
d_rtable, d_drtable, d_ftable, d_dftable,
d_ctable, d_dctable, d_etable, d_detable;

23
src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.cpp
View File

@ -115,6 +115,19 @@ void PairLJCharmmCoulCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_i
if (eflag || vflag) ev_setup(eflag,vflag);
else evflag = vflag_fdotr = 0;
// reallocate per-atom arrays if necessary
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
k_cutsq.template sync<DeviceType>();
k_cut_ljsq.template sync<DeviceType>();
@ -167,6 +180,16 @@ void PairLJCharmmCoulCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_i
virial[5] += ev.v[5];
}
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
copymode = 0;

29
src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.h
View File

@ -34,6 +34,7 @@ class PairLJCharmmCoulCharmmKokkos : public PairLJCharmmCoulCharmm {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
enum {COUL_FLAG=1};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairLJCharmmCoulCharmmKokkos(class LAMMPS *);
~PairLJCharmmCoulCharmmKokkos();
@ -77,27 +78,27 @@ class PairLJCharmmCoulCharmmKokkos : public PairLJCharmmCoulCharmm {
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
typename AT::t_x_array_randomread x;
typename AT::t_x_array c_x;
typename AT::t_f_array f;
typename AT::t_int_1d_randomread type;
typename AT::t_float_1d_randomread q;
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;
typename AT::t_efloat_1d d_eatom;
typename AT::t_virial_array d_vatom;
int newton_pair;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
typename AT::tdual_ffloat_2d k_cutsq;
typename AT::t_ffloat_2d d_cutsq;
typename AT::tdual_ffloat_2d k_cut_ljsq;
typename AT::t_ffloat_2d d_cut_ljsq;
typename AT::tdual_ffloat_2d k_cut_coulsq;
typename AT::t_ffloat_2d d_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_1d_randomread
typename AT::t_ffloat_1d_randomread
d_rtable, d_drtable, d_ftable, d_dftable,
d_ctable, d_dctable, d_etable, d_detable;

4
src/KOKKOS/pair_lj_charmm_coul_long_kokkos.cpp
View File

@ -180,8 +180,6 @@ void PairLJCharmmCoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev.v[5];
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
@ -192,6 +190,8 @@ void PairLJCharmmCoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
copymode = 0;
}

29
src/KOKKOS/pair_lj_charmm_coul_long_kokkos.h
View File

@ -34,6 +34,7 @@ class PairLJCharmmCoulLongKokkos : public PairLJCharmmCoulLong {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
enum {COUL_FLAG=1};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairLJCharmmCoulLongKokkos(class LAMMPS *);
~PairLJCharmmCoulLongKokkos();
@ -75,27 +76,27 @@ class PairLJCharmmCoulLongKokkos : public PairLJCharmmCoulLong {
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
typename AT::t_x_array_randomread x;
typename AT::t_x_array c_x;
typename AT::t_f_array f;
typename AT::t_int_1d_randomread type;
typename AT::t_float_1d_randomread q;
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;
typename AT::t_efloat_1d d_eatom;
typename AT::t_virial_array d_vatom;
int newton_pair;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
typename AT::tdual_ffloat_2d k_cutsq;
typename AT::t_ffloat_2d d_cutsq;
typename AT::tdual_ffloat_2d k_cut_ljsq;
typename AT::t_ffloat_2d d_cut_ljsq;
typename AT::tdual_ffloat_2d k_cut_coulsq;
typename AT::t_ffloat_2d d_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_1d_randomread
typename AT::t_ffloat_1d_randomread
d_rtable, d_drtable, d_ftable, d_dftable,
d_ctable, d_dctable, d_etable, d_detable;

25
src/KOKKOS/pair_lj_class2_coul_cut_kokkos.cpp
View File

@ -92,6 +92,19 @@ void PairLJClass2CoulCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag || vflag) ev_setup(eflag,vflag);
else evflag = vflag_fdotr = 0;
// reallocate per-atom arrays if necessary
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
k_cutsq.template sync<DeviceType>();
k_cut_ljsq.template sync<DeviceType>();
@ -138,7 +151,17 @@ void PairLJClass2CoulCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev.v[5];
}
if (vflag_fdotr) virial_fdotr_compute();
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
copymode = 0;
}

31
src/KOKKOS/pair_lj_class2_coul_cut_kokkos.h
View File

@ -34,6 +34,7 @@ class PairLJClass2CoulCutKokkos : public PairLJClass2CoulCut {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
enum {COUL_FLAG=1};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairLJClass2CoulCutKokkos(class LAMMPS *);
~PairLJClass2CoulCutKokkos();
@ -76,22 +77,25 @@ class PairLJClass2CoulCutKokkos : public PairLJClass2CoulCut {
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
typename AT::t_x_array_randomread x;
typename AT::t_x_array c_x;
typename AT::t_f_array f;
typename AT::t_int_1d_randomread type;
typename AT::t_float_1d_randomread q;
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;
int newton_pair;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
typename AT::tdual_ffloat_2d k_cutsq;
typename AT::t_ffloat_2d d_cutsq;
typename AT::tdual_ffloat_2d k_cut_ljsq;
typename AT::t_ffloat_2d d_cut_ljsq;
typename AT::tdual_ffloat_2d k_cut_coulsq;
typename AT::t_ffloat_2d d_cut_coulsq;
int neighflag;
int nlocal,nall,eflag,vflag;
@ -112,6 +116,7 @@ class PairLJClass2CoulCutKokkos : public PairLJClass2CoulCut {
friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulCutKokkos,HALFTHREAD,void>(PairLJClass2CoulCutKokkos*,NeighListKokkos<DeviceType>*);
friend EV_FLOAT pair_compute<PairLJClass2CoulCutKokkos,void>(PairLJClass2CoulCutKokkos*,
NeighListKokkos<DeviceType>*);
friend void pair_virial_fdotr_compute<PairLJClass2CoulCutKokkos>(PairLJClass2CoulCutKokkos*);
};

23
src/KOKKOS/pair_lj_class2_coul_long_kokkos.cpp
View File

@ -100,6 +100,19 @@ void PairLJClass2CoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag || vflag) ev_setup(eflag,vflag);
else evflag = vflag_fdotr = 0;
// reallocate per-atom arrays if necessary
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
k_cutsq.template sync<DeviceType>();
k_cut_ljsq.template sync<DeviceType>();
@ -152,6 +165,16 @@ void PairLJClass2CoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev.v[5];
}
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
copymode = 0;

32
src/KOKKOS/pair_lj_class2_coul_long_kokkos.h
View File

@ -34,6 +34,7 @@ class PairLJClass2CoulLongKokkos : public PairLJClass2CoulLong {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
enum {COUL_FLAG=1};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairLJClass2CoulLongKokkos(class LAMMPS *);
~PairLJClass2CoulLongKokkos();
@ -76,24 +77,27 @@ class PairLJClass2CoulLongKokkos : public PairLJClass2CoulLong {
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
typename AT::t_x_array_randomread x;
typename AT::t_x_array c_x;
typename AT::t_f_array f;
typename AT::t_int_1d_randomread type;
typename AT::t_float_1d_randomread q;
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;
int newton_pair;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
typename AT::tdual_ffloat_2d k_cutsq;
typename AT::t_ffloat_2d d_cutsq;
typename AT::tdual_ffloat_2d k_cut_ljsq;
typename AT::t_ffloat_2d d_cut_ljsq;
typename AT::tdual_ffloat_2d k_cut_coulsq;
typename AT::t_ffloat_2d d_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_1d_randomread
typename AT::t_ffloat_1d_randomread
d_rtable, d_drtable, d_ftable, d_dftable,
d_ctable, d_dctable, d_etable, d_detable;

24
src/KOKKOS/pair_lj_class2_kokkos.cpp
View File

@ -92,6 +92,19 @@ void PairLJClass2Kokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag || vflag) ev_setup(eflag,vflag);
else evflag = vflag_fdotr = 0;
// reallocate per-atom arrays if necessary
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
k_cutsq.template sync<DeviceType>();
k_params.template sync<DeviceType>();
@ -125,7 +138,18 @@ void PairLJClass2Kokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev.v[5];
}
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
}
template<class DeviceType>

22
src/KOKKOS/pair_lj_class2_kokkos.h
View File

@ -34,6 +34,7 @@ class PairLJClass2Kokkos : public PairLJClass2 {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF|N2};
enum {COUL_FLAG=0};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairLJClass2Kokkos(class LAMMPS *);
~PairLJClass2Kokkos();
@ -79,19 +80,22 @@ class PairLJClass2Kokkos : public PairLJClass2 {
typename Kokkos::DualView<params_lj**,Kokkos::LayoutRight,DeviceType>::t_dev_const_um params;
params_lj m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; // hardwired to space for 12 atom types
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
typename ArrayTypes<DeviceType>::t_tagint_1d tag;
typename AT::t_x_array_randomread x;
typename AT::t_x_array c_x;
typename AT::t_f_array f;
typename AT::t_int_1d_randomread type;
typename AT::t_tagint_1d tag;
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;
int newton_pair;
double special_lj[4];
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename AT::tdual_ffloat_2d k_cutsq;
typename AT::t_ffloat_2d d_cutsq;
int neighflag;
int nlocal,nall,eflag,vflag;

26
src/KOKKOS/pair_lj_cut_coul_cut_kokkos.cpp
View File

@ -92,6 +92,19 @@ void PairLJCutCoulCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag || vflag) ev_setup(eflag,vflag);
else evflag = vflag_fdotr = 0;
// reallocate per-atom arrays if necessary
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
k_cutsq.template sync<DeviceType>();
k_cut_ljsq.template sync<DeviceType>();
@ -136,7 +149,18 @@ void PairLJCutCoulCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev.v[5];
}
if (vflag_fdotr) virial_fdotr_compute();
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
}
/* ----------------------------------------------------------------------

31
src/KOKKOS/pair_lj_cut_coul_cut_kokkos.h
View File

@ -34,6 +34,7 @@ class PairLJCutCoulCutKokkos : public PairLJCutCoulCut {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
enum {COUL_FLAG=1};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairLJCutCoulCutKokkos(class LAMMPS *);
~PairLJCutCoulCutKokkos();
@ -75,22 +76,25 @@ class PairLJCutCoulCutKokkos : public PairLJCutCoulCut {
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
typename AT::t_x_array_randomread x;
typename AT::t_x_array c_x;
typename AT::t_f_array f;
typename AT::t_int_1d_randomread type;
typename AT::t_float_1d_randomread q;
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;
int newton_pair;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
typename AT::tdual_ffloat_2d k_cutsq;
typename AT::t_ffloat_2d d_cutsq;
typename AT::tdual_ffloat_2d k_cut_ljsq;
typename AT::t_ffloat_2d d_cut_ljsq;
typename AT::tdual_ffloat_2d k_cut_coulsq;
typename AT::t_ffloat_2d d_cut_coulsq;
int neighflag;
@ -112,6 +116,7 @@ class PairLJCutCoulCutKokkos : public PairLJCutCoulCut {
friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulCutKokkos,HALFTHREAD,void>(PairLJCutCoulCutKokkos*,NeighListKokkos<DeviceType>*);
friend EV_FLOAT pair_compute<PairLJCutCoulCutKokkos,void>(PairLJCutCoulCutKokkos*,
NeighListKokkos<DeviceType>*);
friend void pair_virial_fdotr_compute<PairLJCutCoulCutKokkos>(PairLJCutCoulCutKokkos*);
};

25
src/KOKKOS/pair_lj_cut_coul_debye_kokkos.cpp
View File

@ -96,6 +96,19 @@ void PairLJCutCoulDebyeKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag || vflag) ev_setup(eflag,vflag);
else evflag = vflag_fdotr = 0;
// reallocate per-atom arrays if necessary
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
k_cutsq.template sync<DeviceType>();
k_cut_ljsq.template sync<DeviceType>();
@ -142,7 +155,17 @@ void PairLJCutCoulDebyeKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev.v[5];
}
if (vflag_fdotr) virial_fdotr_compute();
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
copymode = 0;
}

31
src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h
View File

@ -34,6 +34,7 @@ class PairLJCutCoulDebyeKokkos : public PairLJCutCoulDebye {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
enum {COUL_FLAG=1};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairLJCutCoulDebyeKokkos(class LAMMPS *);
~PairLJCutCoulDebyeKokkos();
@ -75,22 +76,25 @@ class PairLJCutCoulDebyeKokkos : public PairLJCutCoulDebye {
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
typename AT::t_x_array_randomread x;
typename AT::t_x_array c_x;
typename AT::t_f_array f;
typename AT::t_int_1d_randomread type;
typename AT::t_float_1d_randomread q;
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;
int newton_pair;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
typename AT::tdual_ffloat_2d k_cutsq;
typename AT::t_ffloat_2d d_cutsq;
typename AT::tdual_ffloat_2d k_cut_ljsq;
typename AT::t_ffloat_2d d_cut_ljsq;
typename AT::tdual_ffloat_2d k_cut_coulsq;
typename AT::t_ffloat_2d d_cut_coulsq;
int neighflag;
@ -112,6 +116,7 @@ class PairLJCutCoulDebyeKokkos : public PairLJCutCoulDebye {
friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulDebyeKokkos,HALFTHREAD,void>(PairLJCutCoulDebyeKokkos*,NeighListKokkos<DeviceType>*);
friend EV_FLOAT pair_compute<PairLJCutCoulDebyeKokkos,void>(PairLJCutCoulDebyeKokkos*,
NeighListKokkos<DeviceType>*);
friend void pair_virial_fdotr_compute<PairLJCutCoulDebyeKokkos>(PairLJCutCoulDebyeKokkos*);
};

25
src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.cpp
View File

@ -104,6 +104,19 @@ void PairLJCutCoulDSFKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag || vflag) ev_setup(eflag,vflag);
else evflag = vflag_fdotr = 0;
// reallocate per-atom arrays if necessary
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
k_cutsq.template sync<DeviceType>();
k_cut_ljsq.template sync<DeviceType>();
@ -160,7 +173,17 @@ void PairLJCutCoulDSFKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev.v[5];
}
if (vflag_fdotr) virial_fdotr_compute();
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
copymode = 0;
}

31
src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.h
View File

@ -34,6 +34,7 @@ class PairLJCutCoulDSFKokkos : public PairLJCutCoulDSF {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
enum {COUL_FLAG=1};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairLJCutCoulDSFKokkos(class LAMMPS *);
~PairLJCutCoulDSFKokkos();
@ -74,22 +75,25 @@ class PairLJCutCoulDSFKokkos : public PairLJCutCoulDSF {
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
typename AT::t_x_array_randomread x;
typename AT::t_x_array c_x;
typename AT::t_f_array f;
typename AT::t_int_1d_randomread type;
typename AT::t_float_1d_randomread q;
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;
int newton_pair;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
typename AT::tdual_ffloat_2d k_cutsq;
typename AT::t_ffloat_2d d_cutsq;
typename AT::tdual_ffloat_2d k_cut_ljsq;
typename AT::t_ffloat_2d d_cut_ljsq;
typename AT::tdual_ffloat_2d k_cut_coulsq;
typename AT::t_ffloat_2d d_cut_coulsq;
int neighflag;
@ -111,6 +115,7 @@ class PairLJCutCoulDSFKokkos : public PairLJCutCoulDSF {
friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulDSFKokkos,HALFTHREAD,void>(PairLJCutCoulDSFKokkos*,NeighListKokkos<DeviceType>*);
friend EV_FLOAT pair_compute<PairLJCutCoulDSFKokkos,void>(PairLJCutCoulDSFKokkos*,
NeighListKokkos<DeviceType>*);
friend void pair_virial_fdotr_compute<PairLJCutCoulDSFKokkos>(PairLJCutCoulDSFKokkos*);
};

5
src/KOKKOS/pair_lj_cut_coul_long_kokkos.cpp
View File

@ -168,8 +168,6 @@ void PairLJCutCoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev.v[5];
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
@ -179,6 +177,9 @@ void PairLJCutCoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
}
/* ----------------------------------------------------------------------

29
src/KOKKOS/pair_lj_cut_coul_long_kokkos.h
View File

@ -34,6 +34,7 @@ class PairLJCutCoulLongKokkos : public PairLJCutCoulLong {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
enum {COUL_FLAG=1};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairLJCutCoulLongKokkos(class LAMMPS *);
~PairLJCutCoulLongKokkos();
@ -76,27 +77,27 @@ class PairLJCutCoulLongKokkos : public PairLJCutCoulLong {
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
typename AT::t_x_array_randomread x;
typename AT::t_x_array c_x;
typename AT::t_f_array f;
typename AT::t_int_1d_randomread type;
typename AT::t_float_1d_randomread q;
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;
typename AT::t_efloat_1d d_eatom;
typename AT::t_virial_array d_vatom;
int newton_pair;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
typename AT::tdual_ffloat_2d k_cutsq;
typename AT::t_ffloat_2d d_cutsq;
typename AT::tdual_ffloat_2d k_cut_ljsq;
typename AT::t_ffloat_2d d_cut_ljsq;
typename AT::tdual_ffloat_2d k_cut_coulsq;
typename AT::t_ffloat_2d d_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_1d_randomread
typename AT::t_ffloat_1d_randomread
d_rtable, d_drtable, d_ftable, d_dftable,
d_ctable, d_dctable, d_etable, d_detable;

5
src/KOKKOS/pair_lj_cut_kokkos.cpp
View File

@ -138,8 +138,6 @@ void PairLJCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev.v[5];
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
@ -149,6 +147,9 @@ void PairLJCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
}
template<class DeviceType>

19
src/KOKKOS/pair_lj_cut_kokkos.h
View File

@ -34,6 +34,7 @@ class PairLJCutKokkos : public PairLJCut {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF|N2};
enum {COUL_FLAG=0};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairLJCutKokkos(class LAMMPS *);
~PairLJCutKokkos();
@ -73,22 +74,22 @@ class PairLJCutKokkos : public PairLJCut {
typename Kokkos::DualView<params_lj**,Kokkos::LayoutRight,DeviceType>::t_dev_const_um params;
params_lj m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; // hardwired to space for 12 atom types
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename AT::t_x_array_randomread x;
typename AT::t_x_array c_x;
typename AT::t_f_array f;
typename AT::t_int_1d_randomread type;
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;
typename ArrayTypes<DeviceType>::t_tagint_1d tag;
typename AT::t_efloat_1d d_eatom;
typename AT::t_virial_array d_vatom;
typename AT::t_tagint_1d tag;
int newton_pair;
double special_lj[4];
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename AT::tdual_ffloat_2d k_cutsq;
typename AT::t_ffloat_2d d_cutsq;
int neighflag;

23
src/KOKKOS/pair_lj_expand_kokkos.cpp
View File

@ -91,6 +91,19 @@ void PairLJExpandKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag || vflag) ev_setup(eflag,vflag);
else evflag = vflag_fdotr = 0;
// reallocate per-atom arrays if necessary
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
k_cutsq.template sync<DeviceType>();
k_params.template sync<DeviceType>();
@ -126,6 +139,16 @@ void PairLJExpandKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev.v[5];
}
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
copymode = 0;

22
src/KOKKOS/pair_lj_expand_kokkos.h
View File

@ -34,6 +34,7 @@ class PairLJExpandKokkos : public PairLJExpand {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF|N2};
enum {COUL_FLAG=0};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairLJExpandKokkos(class LAMMPS *);
~PairLJExpandKokkos();
@ -79,19 +80,22 @@ class PairLJExpandKokkos : public PairLJExpand {
typename Kokkos::DualView<params_lj**,Kokkos::LayoutRight,DeviceType>::t_dev_const_um params;
params_lj m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; // hardwired to space for 12 atom types
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
typename ArrayTypes<DeviceType>::t_tagint_1d tag;
typename AT::t_x_array_randomread x;
typename AT::t_x_array c_x;
typename AT::t_f_array f;
typename AT::t_int_1d_randomread type;
typename AT::t_tagint_1d tag;
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;
int newton_pair;
double special_lj[4];
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename AT::tdual_ffloat_2d k_cutsq;
typename AT::t_ffloat_2d d_cutsq;
int neighflag;

23
src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.cpp
View File

@ -106,6 +106,19 @@ void PairLJGromacsCoulGromacsKokkos<DeviceType>::compute(int eflag_in, int vflag
if (eflag || vflag) ev_setup(eflag,vflag);
else evflag = vflag_fdotr = 0;
// reallocate per-atom arrays if necessary
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
k_cutsq.template sync<DeviceType>();
k_cut_ljsq.template sync<DeviceType>();
@ -158,6 +171,16 @@ void PairLJGromacsCoulGromacsKokkos<DeviceType>::compute(int eflag_in, int vflag
virial[5] += ev.v[5];
}
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
copymode = 0;

29
src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.h
View File

@ -34,6 +34,7 @@ class PairLJGromacsCoulGromacsKokkos : public PairLJGromacsCoulGromacs {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
enum {COUL_FLAG=1};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairLJGromacsCoulGromacsKokkos(class LAMMPS *);
~PairLJGromacsCoulGromacsKokkos();
@ -84,27 +85,27 @@ class PairLJGromacsCoulGromacsKokkos : public PairLJGromacsCoulGromacs {
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
typename AT::t_x_array_randomread x;
typename AT::t_x_array c_x;
typename AT::t_f_array f;
typename AT::t_int_1d_randomread type;
typename AT::t_float_1d_randomread q;
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;
typename AT::t_efloat_1d d_eatom;
typename AT::t_virial_array d_vatom;
int newton_pair;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
typename AT::tdual_ffloat_2d k_cutsq;
typename AT::t_ffloat_2d d_cutsq;
typename AT::tdual_ffloat_2d k_cut_ljsq;
typename AT::t_ffloat_2d d_cut_ljsq;
typename AT::tdual_ffloat_2d k_cut_coulsq;
typename AT::t_ffloat_2d d_cut_coulsq;
typename ArrayTypes<DeviceType>::t_ffloat_1d_randomread
typename AT::t_ffloat_1d_randomread
d_rtable, d_drtable, d_ftable, d_dftable,
d_ctable, d_dctable, d_etable, d_detable;

23
src/KOKKOS/pair_lj_gromacs_kokkos.cpp
View File

@ -103,6 +103,19 @@ void PairLJGromacsKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag || vflag) ev_setup(eflag,vflag);
else evflag = vflag_fdotr = 0;
// reallocate per-atom arrays if necessary
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
k_cutsq.template sync<DeviceType>();
k_cut_inner.template sync<DeviceType>();
@ -145,6 +158,16 @@ void PairLJGromacsKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev.v[5];
}
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
copymode = 0;

29
src/KOKKOS/pair_lj_gromacs_kokkos.h
View File

@ -34,6 +34,7 @@ class PairLJGromacsKokkos : public PairLJGromacs {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
enum {COUL_FLAG=0};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairLJGromacsKokkos(class LAMMPS *);
~PairLJGromacsKokkos();
@ -87,27 +88,27 @@ class PairLJGromacsKokkos : public PairLJGromacs {
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_inner[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cut_inner_sq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
typename AT::t_x_array_randomread x;
typename AT::t_x_array c_x;
typename AT::t_f_array f;
typename AT::t_int_1d_randomread type;
typename AT::t_float_1d_randomread q;
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;
typename AT::t_efloat_1d d_eatom;
typename AT::t_virial_array d_vatom;
int newton_pair;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_inner;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_inner;
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_inner_sq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_inner_sq;
typename AT::tdual_ffloat_2d k_cutsq;
typename AT::t_ffloat_2d d_cutsq;
typename AT::tdual_ffloat_2d k_cut_inner;
typename AT::t_ffloat_2d d_cut_inner;
typename AT::tdual_ffloat_2d k_cut_inner_sq;
typename AT::t_ffloat_2d d_cut_inner_sq;
typename ArrayTypes<DeviceType>::t_ffloat_1d_randomread
typename AT::t_ffloat_1d_randomread
d_rtable, d_drtable, d_ftable, d_dftable,
d_ctable, d_dctable, d_etable, d_detable;

24
src/KOKKOS/pair_lj_sdk_kokkos.cpp
View File

@ -91,6 +91,19 @@ void PairLJSDKKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag || vflag) ev_setup(eflag,vflag);
else evflag = vflag_fdotr = 0;
// reallocate per-atom arrays if necessary
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
k_cutsq.template sync<DeviceType>();
k_params.template sync<DeviceType>();
@ -124,7 +137,18 @@ void PairLJSDKKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev.v[5];
}
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
}
template<class DeviceType>

22
src/KOKKOS/pair_lj_sdk_kokkos.h
View File

@ -34,6 +34,7 @@ class PairLJSDKKokkos : public PairLJSDK {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF|N2};
enum {COUL_FLAG=0};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairLJSDKKokkos(class LAMMPS *);
~PairLJSDKKokkos();
@ -74,19 +75,22 @@ class PairLJSDKKokkos : public PairLJSDK {
typename Kokkos::DualView<params_lj**,Kokkos::LayoutRight,DeviceType>::t_dev_const_um params;
params_lj m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; // hardwired to space for 12 atom types
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
typename ArrayTypes<DeviceType>::t_tagint_1d tag;
typename AT::t_x_array_randomread x;
typename AT::t_x_array c_x;
typename AT::t_f_array f;
typename AT::t_int_1d_randomread type;
typename AT::t_tagint_1d tag;
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;
int newton_pair;
double special_lj[4];
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename AT::tdual_ffloat_2d k_cutsq;
typename AT::t_ffloat_2d d_cutsq;
int neighflag;

310
src/KOKKOS/pair_morse_kokkos.cpp Normal file
View File

@ -0,0 +1,310 @@
/* ----------------------------------------------------------------------
LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
http://lammps.sandia.gov, Sandia National Laboratories
Steve Plimpton, sjplimp@sandia.gov
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.
------------------------------------------------------------------------- */
/* ----------------------------------------------------------------------
Contributing authors: Stefan Paquay (Eindhoven University of Technology)
------------------------------------------------------------------------- */
#include <math.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include "pair_morse_kokkos.h"
#include "kokkos.h"
#include "atom_kokkos.h"
#include "comm.h"
#include "force.h"
#include "neighbor.h"
#include "neigh_list.h"
#include "neigh_request.h"
#include "update.h"
#include "integrate.h"
#include "respa.h"
#include "math_const.h"
#include "memory.h"
#include "error.h"
#include "atom_masks.h"
using namespace LAMMPS_NS;
using namespace MathConst;
#define KOKKOS_CUDA_MAX_THREADS 256
#define KOKKOS_CUDA_MIN_BLOCKS 8
/* ---------------------------------------------------------------------- */
template<class DeviceType>
PairMorseKokkos<DeviceType>::PairMorseKokkos(LAMMPS *lmp) : PairMorse(lmp)
{
respa_enable = 0;
atomKK = (AtomKokkos *) atom;
execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
datamask_read = X_MASK | F_MASK | TYPE_MASK | ENERGY_MASK | VIRIAL_MASK;
datamask_modify = F_MASK | ENERGY_MASK | VIRIAL_MASK;
cutsq = NULL;
}
/* ---------------------------------------------------------------------- */
template<class DeviceType>
PairMorseKokkos<DeviceType>::~PairMorseKokkos()
{
if (allocated) {
memory->destroy_kokkos(k_eatom,eatom);
memory->destroy_kokkos(k_vatom,vatom);
k_cutsq = DAT::tdual_ffloat_2d();
memory->sfree(cutsq);
eatom = NULL;
vatom = NULL;
cutsq = NULL;
}
}
/* ---------------------------------------------------------------------- */
template<class DeviceType>
void PairMorseKokkos<DeviceType>::cleanup_copy() {
// WHY needed: this prevents parent copy from deallocating any arrays
allocated = 0;
cutsq = NULL;
eatom = NULL;
vatom = NULL;
}
/* ---------------------------------------------------------------------- */
template<class DeviceType>
void PairMorseKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
{
eflag = eflag_in;
vflag = vflag_in;
if (neighflag == FULL) no_virial_fdotr_compute = 1;
if (eflag || vflag) ev_setup(eflag,vflag);
else evflag = vflag_fdotr = 0;
// reallocate per-atom arrays if necessary
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
k_cutsq.template sync<DeviceType>();
k_params.template sync<DeviceType>();
if (eflag || vflag) atomKK->modified(execution_space,datamask_modify);
else atomKK->modified(execution_space,F_MASK);
x = atomKK->k_x.view<DeviceType>();
c_x = atomKK->k_x.view<DeviceType>();
f = atomKK->k_f.view<DeviceType>();
type = atomKK->k_type.view<DeviceType>();
tag = atomKK->k_tag.view<DeviceType>();
nlocal = atom->nlocal;
nall = atom->nlocal + atom->nghost;
newton_pair = force->newton_pair;
special_lj[0] = force->special_lj[0];
special_lj[1] = force->special_lj[1];
special_lj[2] = force->special_lj[2];
special_lj[3] = force->special_lj[3];
// loop over neighbors of my atoms
EV_FLOAT ev = pair_compute<PairMorseKokkos<DeviceType>,void >(this,(NeighListKokkos<DeviceType>*)list);
if (eflag_global) eng_vdwl += ev.evdwl;
if (vflag_global) {
virial[0] += ev.v[0];
virial[1] += ev.v[1];
virial[2] += ev.v[2];
virial[3] += ev.v[3];
virial[4] += ev.v[4];
virial[5] += ev.v[5];
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
}
}
template<class DeviceType>
template<bool STACKPARAMS, class Specialisation>
KOKKOS_INLINE_FUNCTION
F_FLOAT PairMorseKokkos<DeviceType>::
compute_fpair(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const {
(void) i;
(void) j;
const F_FLOAT rr = sqrt(rsq);
const F_FLOAT r0 = STACKPARAMS ? m_params[itype][jtype].r0 : params(itype,jtype).r0;
const F_FLOAT d0 = STACKPARAMS ? m_params[itype][jtype].d0 : params(itype,jtype).d0;
const F_FLOAT aa = STACKPARAMS ? m_params[itype][jtype].alpha : params(itype,jtype).alpha;
const F_FLOAT dr = rr - r0;
// U = d0 * [ exp( -2*a*(x-r0)) - 2*exp(-a*(x-r0)) ]
// f = -2*a*d0*[ -exp( -2*a*(x-r0) ) + exp( -a*(x-r0) ) ] * grad(r)
// = +2*a*d0*[ exp( -2*a*(x-r0) ) - exp( -a*(x-r0) ) ] * grad(r)
const F_FLOAT dexp = exp( -aa*dr );
const F_FLOAT forcelj = 2*aa*d0*dexp*(dexp-1.0);
return forcelj / rr;
}
template<class DeviceType>
template<bool STACKPARAMS, class Specialisation>
KOKKOS_INLINE_FUNCTION
F_FLOAT PairMorseKokkos<DeviceType>::
compute_evdwl(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const {
(void) i;
(void) j;
const F_FLOAT rr = sqrt(rsq);
const F_FLOAT r0 = STACKPARAMS ? m_params[itype][jtype].r0 : params(itype,jtype).r0;
const F_FLOAT d0 = STACKPARAMS ? m_params[itype][jtype].d0 : params(itype,jtype).d0;
const F_FLOAT aa = STACKPARAMS ? m_params[itype][jtype].alpha : params(itype,jtype).alpha;
const F_FLOAT dr = rr - r0;
// U = d0 * [ exp( -2*a*(x-r0)) - 2*exp(-a*(x-r0)) ]
// f = -2*a*d0*[ -exp( -2*a*(x-r0) ) + exp( -a*(x-r0) ) ] * grad(r)
// = +2*a*d0*[ exp( -2*a*(x-r0) ) - exp( -a*(x-r0) ) ] * grad(r)
const F_FLOAT dexp = exp( -aa*dr );
return d0 * dexp * ( dexp - 2.0 );
}
/* ----------------------------------------------------------------------
allocate all arrays
------------------------------------------------------------------------- */
template<class DeviceType>
void PairMorseKokkos<DeviceType>::allocate()
{
PairMorse::allocate();
int n = atom->ntypes;
memory->destroy(cutsq);
memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq");
d_cutsq = k_cutsq.template view<DeviceType>();
k_params = Kokkos::DualView<params_morse**,Kokkos::LayoutRight,DeviceType>("PairMorse::params",n+1,n+1);
params = k_params.d_view;
}
/* ----------------------------------------------------------------------
global settings
------------------------------------------------------------------------- */
template<class DeviceType>
void PairMorseKokkos<DeviceType>::settings(int narg, char **arg)
{
if (narg > 2) error->all(FLERR,"Illegal pair_style command");
PairMorse::settings(1,arg);
}
/* ----------------------------------------------------------------------
init specific to this pair style
------------------------------------------------------------------------- */
template<class DeviceType>
void PairMorseKokkos<DeviceType>::init_style()
{
PairMorse::init_style();
// error if rRESPA with inner levels
if (update->whichflag == 1 && strstr(update->integrate_style,"respa")) {
int respa = 0;
if (((Respa *) update->integrate)->level_inner >= 0) respa = 1;
if (((Respa *) update->integrate)->level_middle >= 0) respa = 2;
if (respa)
error->all(FLERR,"Cannot use Kokkos pair style with rRESPA inner/middle");
}
// irequest = neigh request made by parent class
neighflag = lmp->kokkos->neighflag;
int irequest = neighbor->nrequest - 1;
neighbor->requests[irequest]->
kokkos_host = Kokkos::Impl::is_same<DeviceType,LMPHostType>::value &&
!Kokkos::Impl::is_same<DeviceType,LMPDeviceType>::value;
neighbor->requests[irequest]->
kokkos_device = Kokkos::Impl::is_same<DeviceType,LMPDeviceType>::value;
if (neighflag == FULL) {
neighbor->requests[irequest]->full = 1;
neighbor->requests[irequest]->half = 0;
} else if (neighflag == HALF || neighflag == HALFTHREAD) {
neighbor->requests[irequest]->full = 0;
neighbor->requests[irequest]->half = 1;
} else if (neighflag == N2) {
neighbor->requests[irequest]->full = 0;
neighbor->requests[irequest]->half = 0;
} else {
error->all(FLERR,"Cannot use chosen neighbor list style with morse/kk");
}
}
/* ----------------------------------------------------------------------
init for one type pair i,j and corresponding j,i
------------------------------------------------------------------------- */
// Rewrite this.
template<class DeviceType>
double PairMorseKokkos<DeviceType>::init_one(int i, int j)
{
double cutone = PairMorse::init_one(i,j);
k_params.h_view(i,j).d0 = d0[i][j];
k_params.h_view(i,j).alpha = alpha[i][j];
k_params.h_view(i,j).r0 = r0[i][j];
k_params.h_view(i,j).offset = offset[i][j];
k_params.h_view(i,j).cutsq = cutone*cutone;
k_params.h_view(j,i) = k_params.h_view(i,j);
if(i<MAX_TYPES_STACKPARAMS+1 && j<MAX_TYPES_STACKPARAMS+1) {
m_params[i][j] = m_params[j][i] = k_params.h_view(i,j);
m_cutsq[j][i] = m_cutsq[i][j] = cutone*cutone;
}
k_cutsq.h_view(i,j) = k_cutsq.h_view(j,i) = cutone*cutone;
k_cutsq.template modify<LMPHostType>();
k_params.template modify<LMPHostType>();
return cutone;
}
namespace LAMMPS_NS {
template class PairMorseKokkos<LMPDeviceType>;
#ifdef KOKKOS_HAVE_CUDA
template class PairMorseKokkos<LMPHostType>;
#endif
}

135
src/KOKKOS/pair_morse_kokkos.h Normal file
View File

@ -0,0 +1,135 @@
/* -*- c++ -*- ----------------------------------------------------------
LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
http://lammps.sandia.gov, Sandia National Laboratories
Steve Plimpton, sjplimp@sandia.gov
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 PAIR_CLASS
PairStyle(morse/kk,PairMorseKokkos<LMPDeviceType>)
PairStyle(morse/kk/device,PairMorseKokkos<LMPDeviceType>)
PairStyle(morse/kk/host,PairMorseKokkos<LMPHostType>)
#else
#ifndef LMP_PAIR_MORSE_KOKKOS_H
#define LMP_PAIR_MORSE_KOKKOS_H
#include "pair_kokkos.h"
#include "pair_morse.h"
#include "neigh_list_kokkos.h"
namespace LAMMPS_NS {
template<class DeviceType>
class PairMorseKokkos : public PairMorse {
public:
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF|N2};
enum {COUL_FLAG=0};
typedef DeviceType device_type;
PairMorseKokkos(class LAMMPS *);
virtual ~PairMorseKokkos();
void compute(int, int);
void settings(int, char **);
void init_style();
double init_one(int, int);
struct params_morse{
KOKKOS_INLINE_FUNCTION
params_morse(){cutsq=0,d0=0;alpha=0;r0=0;offset=0;}
KOKKOS_INLINE_FUNCTION
params_morse(int i){cutsq=0,d0=0;alpha=0;r0=0;offset=0;}
F_FLOAT cutsq,d0,alpha,r0,offset;
};
protected:
void cleanup_copy();
template<bool STACKPARAMS, class Specialisation>
KOKKOS_INLINE_FUNCTION
F_FLOAT compute_fpair(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const;
template<bool STACKPARAMS, class Specialisation>
KOKKOS_INLINE_FUNCTION
F_FLOAT compute_evdwl(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const;
template<bool STACKPARAMS, class Specialisation>
KOKKOS_INLINE_FUNCTION
F_FLOAT compute_ecoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const {
return 0;
}
Kokkos::DualView<params_morse**,Kokkos::LayoutRight,DeviceType> k_params;
typename Kokkos::DualView<params_morse**,Kokkos::LayoutRight,DeviceType>::t_dev_const_um params;
params_morse m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
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;
typename ArrayTypes<DeviceType>::t_tagint_1d tag;
int newton_pair;
double special_lj[4];
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
int neighflag;
int nlocal,nall,eflag,vflag;
void allocate();
friend class PairComputeFunctor<PairMorseKokkos,FULL,true>;
friend class PairComputeFunctor<PairMorseKokkos,HALF,true>;
friend class PairComputeFunctor<PairMorseKokkos,HALFTHREAD,true>;
friend class PairComputeFunctor<PairMorseKokkos,N2,true>;
friend class PairComputeFunctor<PairMorseKokkos,FULL,false>;
friend class PairComputeFunctor<PairMorseKokkos,HALF,false>;
friend class PairComputeFunctor<PairMorseKokkos,HALFTHREAD,false>;
friend class PairComputeFunctor<PairMorseKokkos,N2,false>;
friend EV_FLOAT pair_compute_neighlist<PairMorseKokkos,FULL,void>(PairMorseKokkos*,NeighListKokkos<DeviceType>*);
friend EV_FLOAT pair_compute_neighlist<PairMorseKokkos,HALF,void>(PairMorseKokkos*,NeighListKokkos<DeviceType>*);
friend EV_FLOAT pair_compute_neighlist<PairMorseKokkos,HALFTHREAD,void>(PairMorseKokkos*,NeighListKokkos<DeviceType>*);
friend EV_FLOAT pair_compute_neighlist<PairMorseKokkos,N2,void>(PairMorseKokkos*,NeighListKokkos<DeviceType>*);
friend EV_FLOAT pair_compute<PairMorseKokkos,void>(PairMorseKokkos*,NeighListKokkos<DeviceType>*);
friend void pair_virial_fdotr_compute<PairMorseKokkos>(PairMorseKokkos*);
};
}
#endif
#endif
/* ERROR/WARNING messages:
E: Illegal ... command
Self-explanatory. Check the input script syntax and compare to the
documentation for the command. You can use -echo screen as a
command-line option when running LAMMPS to see the offending line.
E: Cannot use Kokkos pair style with rRESPA inner/middle
Self-explanatory.
E: Cannot use chosen neighbor list style with morse/kk
That style is not supported by Kokkos.
*/

38
src/KOKKOS/pair_reax_c_kokkos.cpp
View File

@ -1046,7 +1046,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeLJCoulomb<NEIGHFLAG,
const X_FLOAT ztmp = x(i,2);
const F_FLOAT qi = q(i);
const int itype = type(i);
const int itag = tag(i);
const tagint itag = tag(i);
const int jnum = d_numneigh[i];
F_FLOAT fxtmp, fytmp, fztmp;
@ -1056,7 +1056,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeLJCoulomb<NEIGHFLAG,
int j = d_neighbors(i,jj);
j &= NEIGHMASK;
const int jtype = type(j);
const int jtag = tag(j);
const tagint jtag = tag(j);
const F_FLOAT qj = q(j);
if (NEIGHFLAG != FULL) {
@ -1201,7 +1201,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeTabulatedLJCoulomb<N
const X_FLOAT ztmp = x(i,2);
const F_FLOAT qi = q(i);
const int itype = type(i);
const int itag = tag(i);
const tagint itag = tag(i);
const int jnum = d_numneigh[i];
F_FLOAT fxtmp, fytmp, fztmp;
@ -1211,7 +1211,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeTabulatedLJCoulomb<N
int j = d_neighbors(i,jj);
j &= NEIGHMASK;
const int jtype = type(j);
const int jtag = tag(j);
const tagint jtag = tag(j);
const F_FLOAT qj = q(j);
if (NEIGHFLAG != FULL) {
@ -1580,7 +1580,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxBuildListsHalf<NEIGHFLAG>,
const X_FLOAT ytmp = x(i,1);
const X_FLOAT ztmp = x(i,2);
const int itype = type(i);
const int itag = tag(i);
const tagint itag = tag(i);
const int jnum = d_numneigh[i];
F_FLOAT C12, C34, C56, BO_s, BO_pi, BO_pi2, BO, delij[3], dBOp_i[3], dln_BOp_pi_i[3], dln_BOp_pi2_i[3];
@ -1605,7 +1605,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxBuildListsHalf<NEIGHFLAG>,
for (int jj = 0; jj < jnum; jj++) {
int j = d_neighbors(i,jj);
j &= NEIGHMASK;
const int jtag = tag(j);
const tagint jtag = tag(j);
d_bo_first[j] = j*maxbo;
d_hb_first[j] = j*maxhb;
@ -1802,7 +1802,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxBuildListsHalf_LessAtomics<
const X_FLOAT ytmp = x(i,1);
const X_FLOAT ztmp = x(i,2);
const int itype = type(i);
const int itag = tag(i);
const tagint itag = tag(i);
const int jnum = d_numneigh[i];
F_FLOAT C12, C34, C56, BO_s, BO_pi, BO_pi2, BO, delij[3];
@ -1826,7 +1826,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxBuildListsHalf_LessAtomics<
for (int jj = 0; jj < jnum; jj++) {
int j = d_neighbors(i,jj);
j &= NEIGHMASK;
const int jtag = tag(j);
const tagint jtag = tag(j);
d_bo_first[j] = j*maxbo;
d_hb_first[j] = j*maxhb;
@ -2752,7 +2752,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeTorsion<NEIGHFLAG,EV
const int i = d_ilist[ii];
const int itype = type(i);
const int itag = tag(i);
const tagint itag = tag(i);
const X_FLOAT xtmp = x(i,0);
const X_FLOAT ytmp = x(i,1);
const X_FLOAT ztmp = x(i,2);
@ -2768,7 +2768,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeTorsion<NEIGHFLAG,EV
for (int jj = j_start; jj < j_end; jj++) {
int j = d_bo_list[jj];
j &= NEIGHMASK;
const int jtag = tag(j);
const tagint jtag = tag(j);
const int jtype = type(j);
const int j_index = jj - j_start;
@ -3113,7 +3113,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeHydrogen<NEIGHFLAG,E
const X_FLOAT xtmp = x(i,0);
const X_FLOAT ytmp = x(i,1);
const X_FLOAT ztmp = x(i,2);
const int itag = tag(i);
const tagint itag = tag(i);
const int j_start = d_bo_first[i];
const int j_end = j_start + d_bo_num[i];
@ -3140,7 +3140,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeHydrogen<NEIGHFLAG,E
for (int kk = k_start; kk < k_end; kk++) {
int k = d_hb_list[kk];
k &= NEIGHMASK;
const int ktag = tag(k);
const tagint ktag = tag(k);
const int ktype = type(k);
delik[0] = x(k,0) - xtmp;
@ -3153,7 +3153,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeHydrogen<NEIGHFLAG,E
const int jj = hblist[itr];
int j = d_bo_list[jj];
j &= NEIGHMASK;
const int jtag = tag(j);
const tagint jtag = tag(j);
if (jtag == ktag) continue;
const int jtype = type(j);
@ -3247,14 +3247,14 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxUpdateBond<NEIGHFLAG>, cons
Kokkos::View<F_FLOAT**, typename DAT::t_ffloat_2d_dl::array_layout,DeviceType,Kokkos::MemoryTraits<AtomicF<NEIGHFLAG>::value> > a_Cdbopi2 = d_Cdbopi2;
const int i = d_ilist[ii];
const int itag = tag(i);
const tagint itag = tag(i);
const int j_start = d_bo_first[i];
const int j_end = j_start + d_bo_num[i];
for (int jj = j_start; jj < j_end; jj++) {
int j = d_bo_list[jj];
j &= NEIGHMASK;
const int jtag = tag(j);
const tagint jtag = tag(j);
const int j_index = jj - j_start;
const F_FLOAT Cdbo_i = d_Cdbo(i,j_index);
const F_FLOAT Cdbopi_i = d_Cdbopi(i,j_index);
@ -3302,7 +3302,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeBond1<NEIGHFLAG,EVFL
const X_FLOAT ytmp = x(i,1);
const X_FLOAT ztmp = x(i,2);
const int itype = type(i);
const int itag = tag(i);
const tagint itag = tag(i);
const F_FLOAT imass = paramssing(itype).mass;
const F_FLOAT val_i = paramssing(itype).valency;
const int j_start = d_bo_first[i];
@ -3313,7 +3313,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeBond1<NEIGHFLAG,EVFL
for (int jj = j_start; jj < j_end; jj++) {
int j = d_bo_list[jj];
j &= NEIGHMASK;
const int jtag = tag(j);
const tagint jtag = tag(j);
if (itag > jtag) {
if ((itag+jtag) % 2 == 0) continue;
@ -3438,7 +3438,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeBond2<NEIGHFLAG,EVFL
const X_FLOAT ytmp = x(i,1);
const X_FLOAT ztmp = x(i,2);
const int itype = type(i);
const int itag = tag(i);
const tagint itag = tag(i);
const F_FLOAT imass = paramssing(itype).mass;
const F_FLOAT val_i = paramssing(itype).valency;
const int j_start = d_bo_first[i];
@ -3451,7 +3451,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeBond2<NEIGHFLAG,EVFL
for (int jj = j_start; jj < j_end; jj++) {
int j = d_bo_list[jj];
j &= NEIGHMASK;
const int jtag = tag(j);
const tagint jtag = tag(j);
if (itag > jtag) {
if ((itag+jtag) % 2 == 0) continue;

8
src/KOKKOS/pair_reax_c_kokkos.h
View File

@ -392,7 +392,7 @@ class PairReaxCKokkos : public PairReaxC {
typename AT::t_efloat_1d v_eatom;
DAT::tdual_virial_array k_vatom;
DAT::t_virial_array d_vatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
typename AT::t_virial_array v_vatom;
HAT::t_virial_array h_vatom;
@ -414,9 +414,9 @@ class PairReaxCKokkos : public PairReaxC {
typedef typename tdual_ffloat_2d_n7::t_dev_const_randomread t_ffloat_2d_n7_randomread;
typedef typename tdual_ffloat_2d_n7::t_host t_host_ffloat_2d_n7;
typename ArrayTypes<DeviceType>::t_neighbors_2d d_neighbors;
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_ilist;
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_numneigh;
typename AT::t_neighbors_2d d_neighbors;
typename AT::t_int_1d_randomread d_ilist;
typename AT::t_int_1d_randomread d_numneigh;
typename AT::t_int_1d d_bo_first, d_bo_num, d_bo_list, d_hb_first, d_hb_num, d_hb_list;

8
src/KOKKOS/pair_sw_kokkos.cpp
View File

@ -88,12 +88,12 @@ void PairSWKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.d_view;
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.d_view;
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
@ -171,8 +171,6 @@ void PairSWKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev_all.v[5];
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
@ -183,6 +181,8 @@ void PairSWKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
copymode = 0;
}

18
src/KOKKOS/pair_sw_kokkos.h
View File

@ -118,23 +118,23 @@ class PairSWKokkos : public PairSW {
void threebodyj(const Param&, const Param&, const Param&, const F_FLOAT&, const F_FLOAT&, F_FLOAT *, F_FLOAT *,
F_FLOAT *) const;
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_tagint_1d tag;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename AT::t_x_array_randomread x;
typename AT::t_f_array f;
typename AT::t_tagint_1d tag;
typename AT::t_int_1d_randomread type;
DAT::tdual_efloat_1d k_eatom;
DAT::tdual_virial_array k_vatom;
DAT::t_efloat_1d d_eatom;
DAT::t_virial_array d_vatom;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
DAT::t_int_1d_randomread d_type2frho;
DAT::t_int_2d_randomread d_type2rhor;
DAT::t_int_2d_randomread d_type2z2r;
typename ArrayTypes<DeviceType>::t_neighbors_2d d_neighbors;
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_ilist;
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_numneigh;
typename AT::t_neighbors_2d d_neighbors;
typename AT::t_int_1d_randomread d_ilist;
typename AT::t_int_1d_randomread d_numneigh;
//NeighListKokkos<DeviceType> k_list;
int neighflag,newton_pair;

26
src/KOKKOS/pair_table_kokkos.cpp
View File

@ -12,7 +12,7 @@
------------------------------------------------------------------------- */
/* ----------------------------------------------------------------------
Contributing author: Paul Crozier (SNL)
Contributing author: Christian Trott (SNL)
------------------------------------------------------------------------- */
#include <mpi.h>
@ -89,6 +89,19 @@ void PairTableKokkos<DeviceType>::compute_style(int eflag_in, int vflag_in)
if (eflag || vflag) ev_setup(eflag,vflag);
else evflag = vflag_fdotr = 0;
// reallocate per-atom arrays if necessary
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
//k_cutsq.template sync<DeviceType>();
//k_params.template sync<DeviceType>();
@ -165,7 +178,18 @@ void PairTableKokkos<DeviceType>::compute_style(int eflag_in, int vflag_in)
virial[5] += ev.v[5];
}
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
}
if (vflag_atom) {
k_vatom.template modify<DeviceType>();
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
}
template<class DeviceType>

48
src/KOKKOS/pair_table_kokkos.h
View File

@ -44,6 +44,7 @@ class PairTableKokkos : public PairTable {
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF|N2};
enum {COUL_FLAG=0};
typedef DeviceType device_type;
typedef ArrayTypes<DeviceType> AT;
PairTableKokkos(class LAMMPS *);
virtual ~PairTableKokkos();
@ -62,27 +63,27 @@ class PairTableKokkos : public PairTable {
protected:
/*struct TableDeviceConst {
typename ArrayTypes<DeviceType>::t_ffloat_2d_randomread cutsq;
typename ArrayTypes<DeviceType>::t_int_2d_randomread tabindex;
typename ArrayTypes<DeviceType>::t_int_1d_randomread nshiftbits,nmask;
typename ArrayTypes<DeviceType>::t_ffloat_1d_randomread innersq,invdelta,deltasq6;
typename ArrayTypes<DeviceType>::t_ffloat_2d_randomread rsq,drsq,e,de,f,df,e2,f2;
typename AT::t_ffloat_2d_randomread cutsq;
typename AT::t_int_2d_randomread tabindex;
typename AT::t_int_1d_randomread nshiftbits,nmask;
typename AT::t_ffloat_1d_randomread innersq,invdelta,deltasq6;
typename AT::t_ffloat_2d_randomread rsq,drsq,e,de,f,df,e2,f2;
};*/
//Its faster not to use texture fetch if the number of tables is less than 32!
struct TableDeviceConst {
typename ArrayTypes<DeviceType>::t_ffloat_2d cutsq;
typename ArrayTypes<DeviceType>::t_int_2d tabindex;
typename ArrayTypes<DeviceType>::t_int_1d nshiftbits,nmask;
typename ArrayTypes<DeviceType>::t_ffloat_1d innersq,invdelta,deltasq6;
typename ArrayTypes<DeviceType>::t_ffloat_2d_randomread rsq,drsq,e,de,f,df,e2,f2;
typename AT::t_ffloat_2d cutsq;
typename AT::t_int_2d tabindex;
typename AT::t_int_1d nshiftbits,nmask;
typename AT::t_ffloat_1d innersq,invdelta,deltasq6;
typename AT::t_ffloat_2d_randomread rsq,drsq,e,de,f,df,e2,f2;
};
struct TableDevice {
typename ArrayTypes<DeviceType>::t_ffloat_2d cutsq;
typename ArrayTypes<DeviceType>::t_int_2d tabindex;
typename ArrayTypes<DeviceType>::t_int_1d nshiftbits,nmask;
typename ArrayTypes<DeviceType>::t_ffloat_1d innersq,invdelta,deltasq6;
typename ArrayTypes<DeviceType>::t_ffloat_2d rsq,drsq,e,de,f,df,e2,f2;
typename AT::t_ffloat_2d cutsq;
typename AT::t_int_2d tabindex;
typename AT::t_int_1d nshiftbits,nmask;
typename AT::t_ffloat_1d innersq,invdelta,deltasq6;
typename AT::t_ffloat_2d rsq,drsq,e,de,f,df,e2,f2;
};
struct TableHost {
@ -99,17 +100,20 @@ class PairTableKokkos : public PairTable {
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
typename AT::t_ffloat_2d d_cutsq;
virtual void allocate();
void compute_table(Table *);
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_x_array_const c_x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
typename AT::t_x_array_randomread x;
typename AT::t_x_array_const c_x;
typename AT::t_f_array f;
typename AT::t_int_1d_randomread type;
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;
protected:
int nlocal,nall,eflag,vflag,neighflag,newton_pair;

12
src/KOKKOS/pair_tersoff_kokkos.cpp
View File

@ -172,12 +172,12 @@ void PairTersoffKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.d_view;
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.d_view;
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
@ -254,8 +254,6 @@ void PairTersoffKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev_all.v[5];
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
@ -266,6 +264,8 @@ void PairTersoffKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
copymode = 0;
}
@ -314,7 +314,7 @@ void PairTersoffKokkos<DeviceType>::operator()(TagPairTersoffComputeHalf<NEIGHFL
const X_FLOAT ytmp = x(i,1);
const X_FLOAT ztmp = x(i,2);
const int itype = type(i);
const int itag = tag(i);
const tagint itag = tag(i);
F_FLOAT fi[3], fj[3], fk[3];
@ -331,7 +331,7 @@ void PairTersoffKokkos<DeviceType>::operator()(TagPairTersoffComputeHalf<NEIGHFL
int j = d_neighbors_short(i,jj);
j &= NEIGHMASK;
const int jtype = type(j);
const int jtag = tag(j);
const tagint jtag = tag(j);
if (itag > jtag) {
if ((itag+jtag) % 2 == 0) continue;

10
src/KOKKOS/pair_tersoff_kokkos.h
View File

@ -197,16 +197,16 @@ class PairTersoffKokkos : public PairTersoff {
DAT::tdual_efloat_1d k_eatom;
DAT::tdual_virial_array k_vatom;
DAT::t_efloat_1d d_eatom;
DAT::t_virial_array d_vatom;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
typedef Kokkos::DualView<F_FLOAT**[7],Kokkos::LayoutRight,DeviceType> tdual_ffloat_2d_n7;
typedef typename tdual_ffloat_2d_n7::t_dev_const_randomread t_ffloat_2d_n7_randomread;
typedef typename tdual_ffloat_2d_n7::t_host t_host_ffloat_2d_n7;
typename ArrayTypes<DeviceType>::t_neighbors_2d d_neighbors;
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_ilist;
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_numneigh;
typename AT::t_neighbors_2d d_neighbors;
typename AT::t_int_1d_randomread d_ilist;
typename AT::t_int_1d_randomread d_numneigh;
//NeighListKokkos<DeviceType> k_list;
int neighflag,newton_pair;

12
src/KOKKOS/pair_tersoff_mod_kokkos.cpp
View File

@ -172,12 +172,12 @@ void PairTersoffMODKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.d_view;
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.d_view;
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
@ -254,8 +254,6 @@ void PairTersoffMODKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev_all.v[5];
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
@ -266,6 +264,8 @@ void PairTersoffMODKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
copymode = 0;
}
@ -314,7 +314,7 @@ void PairTersoffMODKokkos<DeviceType>::operator()(TagPairTersoffMODComputeHalf<N
const X_FLOAT ytmp = x(i,1);
const X_FLOAT ztmp = x(i,2);
const int itype = type(i);
const int itag = tag(i);
const int tagitag = tag(i);
F_FLOAT fi[3], fj[3], fk[3];
@ -331,7 +331,7 @@ void PairTersoffMODKokkos<DeviceType>::operator()(TagPairTersoffMODComputeHalf<N
int j = d_neighbors_short(i,jj);
j &= NEIGHMASK;
const int jtype = type(j);
const int jtag = tag(j);
const tagint jtag = tag(j);
if (itag > jtag) {
if ((itag+jtag) % 2 == 0) continue;

10
src/KOKKOS/pair_tersoff_mod_kokkos.h
View File

@ -197,16 +197,16 @@ class PairTersoffMODKokkos : public PairTersoffMOD {
DAT::tdual_efloat_1d k_eatom;
DAT::tdual_virial_array k_vatom;
DAT::t_efloat_1d d_eatom;
DAT::t_virial_array d_vatom;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
typedef Kokkos::DualView<F_FLOAT**[7],Kokkos::LayoutRight,DeviceType> tdual_ffloat_2d_n7;
typedef typename tdual_ffloat_2d_n7::t_dev_const_randomread t_ffloat_2d_n7_randomread;
typedef typename tdual_ffloat_2d_n7::t_host t_host_ffloat_2d_n7;
typename ArrayTypes<DeviceType>::t_neighbors_2d d_neighbors;
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_ilist;
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_numneigh;
typename AT::t_neighbors_2d d_neighbors;
typename AT::t_int_1d_randomread d_ilist;
typename AT::t_int_1d_randomread d_numneigh;
//NeighListKokkos<DeviceType> k_list;
int neighflag,newton_pair;

12
src/KOKKOS/pair_tersoff_zbl_kokkos.cpp
View File

@ -186,12 +186,12 @@ void PairTersoffZBLKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.d_view;
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.d_view;
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
@ -268,8 +268,6 @@ void PairTersoffZBLKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev_all.v[5];
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
@ -280,6 +278,8 @@ void PairTersoffZBLKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
copymode = 0;
}
@ -328,7 +328,7 @@ void PairTersoffZBLKokkos<DeviceType>::operator()(TagPairTersoffZBLComputeHalf<N
const X_FLOAT ytmp = x(i,1);
const X_FLOAT ztmp = x(i,2);
const int itype = type(i);
const int itag = tag(i);
const tagint itag = tag(i);
F_FLOAT fi[3], fj[3], fk[3];
@ -345,7 +345,7 @@ void PairTersoffZBLKokkos<DeviceType>::operator()(TagPairTersoffZBLComputeHalf<N
int j = d_neighbors_short(i,jj);
j &= NEIGHMASK;
const int jtype = type(j);
const int jtag = tag(j);
const tagint jtag = tag(j);
if (itag > jtag) {
if ((itag+jtag) % 2 == 0) continue;

4
src/KOKKOS/pair_tersoff_zbl_kokkos.h
View File

@ -202,8 +202,8 @@ class PairTersoffZBLKokkos : public PairTersoffZBL {
DAT::tdual_efloat_1d k_eatom;
DAT::tdual_virial_array k_vatom;
DAT::t_efloat_1d d_eatom;
DAT::t_virial_array d_vatom;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
typedef Kokkos::DualView<F_FLOAT**[7],Kokkos::LayoutRight,DeviceType> tdual_ffloat_2d_n7;
typedef typename tdual_ffloat_2d_n7::t_dev_const_randomread t_ffloat_2d_n7_randomread;

11
src/KOKKOS/pair_vashishta_kokkos.cpp
View File

@ -12,8 +12,7 @@
------------------------------------------------------------------------- */
/* ----------------------------------------------------------------------
Contributing author: Stan Moore (SNL)
Anders Hafreager (UiO), andershaf@gmail.com
Contributing author: Anders Hafreager (UiO), andershaf@gmail.com
------------------------------------------------------------------------- */
#include <math.h>
@ -88,12 +87,12 @@ void PairVashishtaKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (eflag_atom) {
memory->destroy_kokkos(k_eatom,eatom);
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
d_eatom = k_eatom.d_view;
d_eatom = k_eatom.view<DeviceType>();
}
if (vflag_atom) {
memory->destroy_kokkos(k_vatom,vatom);
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
d_vatom = k_vatom.d_view;
d_vatom = k_vatom.view<DeviceType>();
}
atomKK->sync(execution_space,datamask_read);
@ -182,8 +181,6 @@ void PairVashishtaKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
virial[5] += ev_all.v[5];
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
if (eflag_atom) {
k_eatom.template modify<DeviceType>();
k_eatom.template sync<LMPHostType>();
@ -194,6 +191,8 @@ void PairVashishtaKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
k_vatom.template sync<LMPHostType>();
}
if (vflag_fdotr) pair_virial_fdotr_compute(this);
copymode = 0;
}

18
src/KOKKOS/pair_vashishta_kokkos.h
View File

@ -118,23 +118,23 @@ class PairVashishtaKokkos : public PairVashishta {
void threebodyj(const Param&, const Param&, const Param&, const F_FLOAT&, const F_FLOAT&, F_FLOAT *, F_FLOAT *,
F_FLOAT *) const;
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
typename ArrayTypes<DeviceType>::t_f_array f;
typename ArrayTypes<DeviceType>::t_tagint_1d tag;
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
typename AT::t_x_array_randomread x;
typename AT::t_f_array f;
typename AT::t_tagint_1d tag;
typename AT::t_int_1d_randomread type;
DAT::tdual_efloat_1d k_eatom;
DAT::tdual_virial_array k_vatom;
DAT::t_efloat_1d d_eatom;
DAT::t_virial_array d_vatom;
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
DAT::t_int_1d_randomread d_type2frho;
DAT::t_int_2d_randomread d_type2rhor;
DAT::t_int_2d_randomread d_type2z2r;
typename ArrayTypes<DeviceType>::t_neighbors_2d d_neighbors;
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_ilist;
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_numneigh;
typename AT::t_neighbors_2d d_neighbors;
typename AT::t_int_1d_randomread d_ilist;
typename AT::t_int_1d_randomread d_numneigh;
//NeighListKokkos<DeviceType> k_list;
int neighflag,newton_pair;

29
src/KOKKOS/pppm_kokkos.cpp
View File

@ -374,7 +374,6 @@ void PPPMKokkos<DeviceType>::setup()
error->all(FLERR,"Incorrect boundaries with slab PPPM");
}
int i,j,k,n;
double *prd;
// volume-dependent factors
@ -417,10 +416,6 @@ void PPPMKokkos<DeviceType>::setup()
DeviceType::fence();
copymode = 0;
// virial coefficients
double sqk,vterm;
// merge three outer loops into one for better threading
numz_fft = nzhi_fft-nzlo_fft + 1;
@ -615,7 +610,7 @@ void PPPMKokkos<DeviceType>::setup_grid()
template<class DeviceType>
void PPPMKokkos<DeviceType>::compute(int eflag, int vflag)
{
int i,j;
int i;
// set energy/virial flags
// invoke allocate_peratom() if needed for first time
@ -974,7 +969,6 @@ void PPPMKokkos<DeviceType>::set_grid_global()
// fluid-occupied volume used to estimate real-space error
// zprd used rather than zprd_slab
double h;
bigint natoms = atomKK->natoms;
if (!gewaldflag) {
@ -1636,9 +1630,6 @@ void PPPMKokkos<DeviceType>::operator()(TagPPPM_particle_map, const int &i) cons
template<class DeviceType>
void PPPMKokkos<DeviceType>::make_rho()
{
int l,m,n,nx,ny,nz,mx,my,mz;
FFT_SCALAR dx,dy,dz,x0,y0,z0;
// clear 3d density array
//memset(&(density_brick(nzlo_out,nylo_out,nxlo_out)),0,
@ -1845,8 +1836,7 @@ void PPPMKokkos<DeviceType>::poisson()
template<class DeviceType>
void PPPMKokkos<DeviceType>::poisson_ik()
{
int i,j,k,n;
double eng;
int j;
// transform charge density (r -> k)
@ -1877,7 +1867,6 @@ void PPPMKokkos<DeviceType>::poisson_ik()
for (j = 0; j < 6; j++) virial[j] += ev.v[j];
energy += ev.ecoul;
} else {
n = 0;
copymode = 1;
Kokkos::parallel_reduce(Kokkos::RangePolicy<DeviceType, TagPPPM_poisson_ik3>(0,nfft),*this,ev);
DeviceType::fence();
@ -2214,8 +2203,6 @@ void PPPMKokkos<DeviceType>::operator()(TagPPPM_poisson_ik_triclinic6, const int
template<class DeviceType>
void PPPMKokkos<DeviceType>::poisson_peratom()
{
int i,j,k,n;
// merge three outer loops into one for better threading
numz_inout = (nzhi_in-nzlo_out)-(nzlo_in-nzlo_out) + 1;
@ -2567,22 +2554,17 @@ KOKKOS_INLINE_FUNCTION
void PPPMKokkos<DeviceType>::operator()(TagPPPM_fieldforce_ik, const int &i) const
{
int l,m,n,nx,ny,nz,mx,my,mz;
FFT_SCALAR dx,dy,dz,x0,y0,z0;
FFT_SCALAR x0,y0,z0;
FFT_SCALAR ekx,eky,ekz;
nx = d_part2grid(i,0);
ny = d_part2grid(i,1);
nz = d_part2grid(i,2);
dx = nx+shiftone - (x(i,0)-boxlo[0])*delxinv;
dy = ny+shiftone - (x(i,1)-boxlo[1])*delyinv;
dz = nz+shiftone - (x(i,2)-boxlo[2])*delzinv;
nz -= nzlo_out;
ny -= nylo_out;
nx -= nxlo_out;
//compute_rho1d(i,dx,dy,dz); // hasn't changed from make_rho
ekx = eky = ekz = ZEROF;
for (n = nlower; n <= nupper; n++) {
mz = n+nz;
@ -2842,8 +2824,8 @@ void PPPMKokkos<DeviceType>::unpack_reverse_kokkos(int flag, Kokkos::DualView<FF
d_list_index = Kokkos::subview(d_list,index,Kokkos::ALL());
d_buf = k_buf.view<DeviceType>();
int nx = (nxhi_out-nxlo_out+1);
int ny = (nyhi_out-nylo_out+1);
nx = (nxhi_out-nxlo_out+1);
ny = (nyhi_out-nylo_out+1);
copymode = 1;
Kokkos::parallel_for(Kokkos::RangePolicy<DeviceType, TagPPPM_unpack_reverse>(0,nlist),*this);
@ -3012,7 +2994,6 @@ void PPPMKokkos<DeviceType>::slabcorr()
// sum local contributions to get global dipole moment
dipole_all;
MPI_Allreduce(&dipole,&dipole_all,1,MPI_DOUBLE,MPI_SUM,world);
// need to make non-neutral systems and/or

31
src/KSPACE/fft3d.cpp
View File

@ -67,7 +67,10 @@
void fft_3d(FFT_DATA *in, FFT_DATA *out, int flag, struct fft_plan_3d *plan)
{
int i,total,length,offset,num;
FFT_SCALAR norm, *out_ptr;
FFT_SCALAR norm;
#if defined(FFT_FFTW3)
FFT_SCALAR *out_ptr;
#endif
FFT_DATA *data,*copy;
// system specific constants
@ -210,7 +213,9 @@ void fft_3d(FFT_DATA *in, FFT_DATA *out, int flag, struct fft_plan_3d *plan)
if (flag == 1 && plan->scaled) {
norm = plan->norm;
num = plan->normnum;
#if defined(FFT_FFTW3)
out_ptr = (FFT_SCALAR *)out;
#endif
for (i = 0; i < num; i++) {
#if defined(FFT_FFTW3)
*(out_ptr++) *= norm;
@ -256,13 +261,12 @@ struct fft_plan_3d *fft_3d_create_plan(
{
struct fft_plan_3d *plan;
int me,nprocs;
int i,num,flag,remapflag,fftflag;
int flag,remapflag;
int first_ilo,first_ihi,first_jlo,first_jhi,first_klo,first_khi;
int second_ilo,second_ihi,second_jlo,second_jhi,second_klo,second_khi;
int third_ilo,third_ihi,third_jlo,third_jhi,third_klo,third_khi;
int out_size,first_size,second_size,third_size,copy_size,scratch_size;
int np1,np2,ip1,ip2;
int list[50];
// query MPI info
@ -756,8 +760,11 @@ void bifactor(int n, int *factor1, int *factor2)
void fft_1d_only(FFT_DATA *data, int nsize, int flag, struct fft_plan_3d *plan)
{
int i,total,length,offset,num;
FFT_SCALAR norm, *data_ptr;
int i,num;
FFT_SCALAR norm;
#if defined(FFT_FFTW3)
FFT_SCALAR *data_ptr;
#endif
// total = size of data needed in each dim
// length = length of 1d FFT in each dim
@ -823,18 +830,18 @@ void fft_1d_only(FFT_DATA *data, int nsize, int flag, struct fft_plan_3d *plan)
FFTW_API(execute_dft)(theplan,data,data);
#else
if (flag == -1) {
for (offset = 0; offset < total1; offset += length1)
for (int offset = 0; offset < total1; offset += length1)
kiss_fft(plan->cfg_fast_forward,&data[offset],&data[offset]);
for (offset = 0; offset < total2; offset += length2)
for (int offset = 0; offset < total2; offset += length2)
kiss_fft(plan->cfg_mid_forward,&data[offset],&data[offset]);
for (offset = 0; offset < total3; offset += length3)
for (int offset = 0; offset < total3; offset += length3)
kiss_fft(plan->cfg_slow_forward,&data[offset],&data[offset]);
} else {
for (offset = 0; offset < total1; offset += length1)
for (int offset = 0; offset < total1; offset += length1)
kiss_fft(plan->cfg_fast_backward,&data[offset],&data[offset]);
for (offset = 0; offset < total2; offset += length2)
for (int offset = 0; offset < total2; offset += length2)
kiss_fft(plan->cfg_mid_backward,&data[offset],&data[offset]);
for (offset = 0; offset < total3; offset += length3)
for (int offset = 0; offset < total3; offset += length3)
kiss_fft(plan->cfg_slow_backward,&data[offset],&data[offset]);
}
#endif
@ -845,7 +852,9 @@ void fft_1d_only(FFT_DATA *data, int nsize, int flag, struct fft_plan_3d *plan)
if (flag == 1 && plan->scaled) {
norm = plan->norm;
num = MIN(plan->normnum,nsize);
#if defined(FFT_FFTW3)
data_ptr = (FFT_SCALAR *)data;
#endif
for (i = 0; i < num; i++) {
#if defined(FFT_FFTW3)
*(data_ptr++) *= norm;

11
src/MOLECULE/fix_cmap.cpp
View File

@ -230,7 +230,7 @@ void FixCMAP::min_setup(int vflag)
void FixCMAP::pre_neighbor()
{
int i,m,itype,atom1,atom2,atom3,atom4,atom5;
int i,m,atom1,atom2,atom3,atom4,atom5;
// guesstimate initial length of local crossterm list
// if ncmap was not set (due to read_restart, no read_data),
@ -842,7 +842,7 @@ void FixCMAP::set_map_derivatives(double **map, double **d1yo, double **d2yo,
// use the bicubic spline to calculate the derivatives
int i, j, k, ii, jj, xm, p;
double phi, psi, y, d1y, d2y, d12y, tyyk,tdyk;
double phi, psi, d1y, d2y, d12y, tyyk,tdyk;
double *tmp_y, *tmp_dy, *tmp_ddy, **tmap, **tddmap;
int ix;
double a,b,a1,b1,a2,b2;
@ -850,7 +850,6 @@ void FixCMAP::set_map_derivatives(double **map, double **d1yo, double **d2yo,
xm = CMAPDIM/2;
p = CMAPDIM;
y = 0.;
d1y = 0.;
d2y = 0.;
d12y = 0.;
@ -907,8 +906,6 @@ void FixCMAP::set_map_derivatives(double **map, double **d1yo, double **d2yo,
b1 = b*b*b-b;
a2 = 3.0*a*a-1.0;
b2 = 3.0*b*b-1.0;
y = a*tmp_y[ix]+b*tmp_y[ix+1]+
(a1*tmp_ddy[ix]+b1*tmp_ddy[ix+1])*(CMAPDX*CMAPDX)/6.0;
d1y = (tmp_y[ix+1]-tmp_y[ix])/CMAPDX-
a2/6.0*CMAPDX*tmp_ddy[ix]+b2/6.0*CMAPDX*tmp_ddy[ix+1];
spline(tmp_dy,tmp_ddy,CMAPDIM+xm+xm);
@ -1015,8 +1012,8 @@ void FixCMAP::bc_interpol(double x1, double x2, int low1, int low2, double *gs,
// gradients and cross-derivatives
// calculate the interpolated value of the point of interest (POI)
int i, p=12;
double t, u, fac, gs1l, gs2l, gs1u, gs2u;
int i;
double t, u, gs1l, gs2l;
// set the interpolation coefficients

4
src/pair_morse.h
View File

@ -28,8 +28,8 @@ class PairMorse : public Pair {
public:
PairMorse(class LAMMPS *);
virtual ~PairMorse();
virtual void compute(int, int);
void settings(int, char **);
void coeff(int, char **);
double init_one(int, int);
@ -49,7 +49,7 @@ class PairMorse : public Pair {
double **morse1;
double **offset;
void allocate();
virtual void allocate();
};
}