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

Merge pull request #1 from akohlmey/tagint-issue 

Fix additional tagint issue in fix qeq/reax/kk
This commit is contained in:
Richard Berger 2017-01-31 18:34:35 -05:00 committed by GitHub
parent 51a0b6b445 bcb1d94b9a
commit 67bed8e853
10 changed files with 92 additions and 67 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
src/KOKKOS/angle_charmm_kokkos.cpp
View File

@ -84,9 +84,6 @@ void AngleCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
}
}
if (eflag || vflag) atomKK->modified(execution_space,datamask_modify);
else atomKK->modified(execution_space,F_MASK);
x = atomKK->k_x.view<DeviceType>();
f = atomKK->k_f.view<DeviceType>();
neighborKK->k_anglelist.template sync<DeviceType>();

3
src/KOKKOS/bond_fene_kokkos.cpp
View File

@ -85,13 +85,10 @@ void BondFENEKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
d_vatom = k_vatom.d_view;
}
atomKK->sync(execution_space,datamask_read);
k_k.template sync<DeviceType>();
k_r0.template sync<DeviceType>();
k_epsilon.template sync<DeviceType>();
k_sigma.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>();
f = atomKK->k_f.view<DeviceType>();

5
src/KOKKOS/dihedral_charmm_kokkos.cpp
View File

@ -98,11 +98,6 @@ void DihedralCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
}
}
//atomKK->sync(execution_space,datamask_read);
if (eflag || vflag) atomKK->modified(execution_space,datamask_modify);
else atomKK->modified(execution_space,F_MASK);
x = atomKK->k_x.view<DeviceType>();
f = atomKK->k_f.view<DeviceType>();
q = atomKK->k_q.view<DeviceType>();

3
src/KOKKOS/dihedral_opls_kokkos.cpp
View File

@ -85,13 +85,10 @@ void DihedralOPLSKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
d_vatom = k_vatom.d_view;
}
atomKK->sync(execution_space,datamask_read);
k_k1.template sync<DeviceType>();
k_k2.template sync<DeviceType>();
k_k3.template sync<DeviceType>();
k_k4.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>();
f = atomKK->k_f.view<DeviceType>();

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;

78
src/KOKKOS/npair_kokkos.cpp
View File

@ -24,8 +24,8 @@ namespace LAMMPS_NS {
/* ---------------------------------------------------------------------- */
template<class DeviceType, int HALF_NEIGH, int GHOST>
NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::NPairKokkos(LAMMPS *lmp) : NPair(lmp) {
template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>
NPairKokkos<DeviceType,HALF_NEIGH,GHOST,TRI>::NPairKokkos(LAMMPS *lmp) : NPair(lmp) {
}
@ -33,8 +33,8 @@ NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::NPairKokkos(LAMMPS *lmp) : NPair(lmp)
copy needed info from Neighbor class to this build class
------------------------------------------------------------------------- */
template<class DeviceType, int HALF_NEIGH, int GHOST>
void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::copy_neighbor_info()
template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>
void NPairKokkos<DeviceType,HALF_NEIGH,GHOST,TRI>::copy_neighbor_info()
{
NPair::copy_neighbor_info();
@ -62,8 +62,8 @@ void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::copy_neighbor_info()
copy per-atom and per-bin vectors from NBin class to this build class
------------------------------------------------------------------------- */
template<class DeviceType, int HALF_NEIGH, int GHOST>
void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::copy_bin_info()
template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>
void NPairKokkos<DeviceType,HALF_NEIGH,GHOST,TRI>::copy_bin_info()
{
NPair::copy_bin_info();
@ -78,8 +78,8 @@ void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::copy_bin_info()
copy needed info from NStencil class to this build class
------------------------------------------------------------------------- */
template<class DeviceType, int HALF_NEIGH, int GHOST>
void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::copy_stencil_info()
template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>
void NPairKokkos<DeviceType,HALF_NEIGH,GHOST,TRI>::copy_stencil_info()
{
NPair::copy_stencil_info();
@ -106,8 +106,8 @@ void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::copy_stencil_info()
/* ---------------------------------------------------------------------- */
template<class DeviceType, int HALF_NEIGH, int GHOST>
void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::build(NeighList *list_)
template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>
void NPairKokkos<DeviceType,HALF_NEIGH,GHOST,TRI>::build(NeighList *list_)
{
NeighListKokkos<DeviceType>* list = (NeighListKokkos<DeviceType>*) list_;
const int nlocal = includegroup?atom->nfirst:atom->nlocal;
@ -196,14 +196,14 @@ if (GHOST) {
Kokkos::parallel_for(nall, f);
} else {
if (newton_pair) {
NPairKokkosBuildFunctor<DeviceType,HALF_NEIGH,1> f(data,atoms_per_bin * 5 * sizeof(X_FLOAT) * factor);
NPairKokkosBuildFunctor<DeviceType,TRI?0:HALF_NEIGH,1,TRI> f(data,atoms_per_bin * 5 * sizeof(X_FLOAT) * factor);
#ifdef KOKKOS_HAVE_CUDA
Kokkos::parallel_for(config, f);
#else
Kokkos::parallel_for(nall, f);
#endif
} else {
NPairKokkosBuildFunctor<DeviceType,HALF_NEIGH,0> f(data,atoms_per_bin * 5 * sizeof(X_FLOAT) * factor);
NPairKokkosBuildFunctor<DeviceType,HALF_NEIGH,0,0> f(data,atoms_per_bin * 5 * sizeof(X_FLOAT) * factor);
#ifdef KOKKOS_HAVE_CUDA
Kokkos::parallel_for(config, f);
#else
@ -293,7 +293,7 @@ int NeighborKokkosExecute<DeviceType>::exclusion(const int &i,const int &j,
/* ---------------------------------------------------------------------- */
template<class DeviceType> template<int HalfNeigh,int Newton>
template<class DeviceType> template<int HalfNeigh,int Newton,int Tri>
void NeighborKokkosExecute<DeviceType>::
build_Item(const int &i) const
{
@ -365,7 +365,7 @@ void NeighborKokkosExecute<DeviceType>::
const int jbin = ibin + stencil[k];
// get subview of jbin
if(HalfNeigh&&(ibin==jbin)) continue;
if(HalfNeigh && (ibin==jbin)) continue;
//const ArrayTypes<DeviceType>::t_int_1d_const_um =Kokkos::subview<t_int_1d_const_um>(bins,jbin,ALL);
for(int m = 0; m < c_bincount(jbin); m++) {
@ -374,6 +374,16 @@ void NeighborKokkosExecute<DeviceType>::
if(HalfNeigh && !Newton && (j < i)) continue;
if(!HalfNeigh && j==i) continue;
if(Tri) {
if (x(j,2) < ztmp) continue;
if (x(j,2) == ztmp) {
if (x(j,1) < ytmp) continue;
if (x(j,1) == ytmp) {
if (x(j,0) < xtmp) continue;
if (x(j,0) == xtmp && j <= i) continue;
}
}
}
if(exclude && exclusion(i,j,itype,jtype)) continue;
const X_FLOAT delx = xtmp - x(j, 0);
@ -428,7 +438,7 @@ extern __shared__ X_FLOAT sharedmem[];
/* ---------------------------------------------------------------------- */
template<class DeviceType> template<int HalfNeigh,int Newton>
template<class DeviceType> template<int HalfNeigh,int Newton,int Tri>
__device__ inline
void NeighborKokkosExecute<DeviceType>::build_ItemCuda(typename Kokkos::TeamPolicy<DeviceType>::member_type dev) const
{
@ -491,6 +501,16 @@ void NeighborKokkosExecute<DeviceType>::build_ItemCuda(typename Kokkos::TeamPoli
((j >= nlocal) && ((x(j, 2) < ztmp) || (x(j, 2) == ztmp && x(j, 1) < ytmp) ||
(x(j, 2) == ztmp && x(j, 1) == ytmp && x(j, 0) < xtmp)))))
) continue;
if(Tri) {
if (x(j,2) < ztmp) continue;
if (x(j,2) == ztmp) {
if (x(j,1) < ytmp) continue;
if (x(j,1) == ytmp) {
if (x(j,0) < xtmp) continue;
if (x(j,0) == xtmp && j <= i) continue;
}
}
}
if(exclude && exclusion(i,j,itype,jtype)) continue;
const X_FLOAT delx = xtmp - other_x[m];
const X_FLOAT dely = ytmp - other_x[m + atoms_per_bin];
@ -558,6 +578,16 @@ void NeighborKokkosExecute<DeviceType>::build_ItemCuda(typename Kokkos::TeamPoli
//if(HalfNeigh && (j < i)) continue;
if(HalfNeigh && !Newton && (j < i)) continue;
if(!HalfNeigh && j==i) continue;
if(Tri) {
if (x(j,2) < ztmp) continue;
if (x(j,2) == ztmp) {
if (x(j,1) < ytmp) continue;
if (x(j,1) == ytmp) {
if (x(j,0) < xtmp) continue;
if (x(j,0) == xtmp && j <= i) continue;
}
}
}
if(exclude && exclusion(i,j,itype,jtype)) continue;
const X_FLOAT delx = xtmp - other_x[m];
@ -736,14 +766,16 @@ void NeighborKokkosExecute<DeviceType>::
}
namespace LAMMPS_NS {
template class NPairKokkos<LMPDeviceType,0,0>;
template class NPairKokkos<LMPDeviceType,0,1>;
template class NPairKokkos<LMPDeviceType,1,0>;
template class NPairKokkos<LMPDeviceType,1,1>;
template class NPairKokkos<LMPDeviceType,0,0,0>;
template class NPairKokkos<LMPDeviceType,0,1,0>;
template class NPairKokkos<LMPDeviceType,1,0,0>;
template class NPairKokkos<LMPDeviceType,1,1,0>;
template class NPairKokkos<LMPDeviceType,1,0,1>;
#ifdef KOKKOS_HAVE_CUDA
template class NPairKokkos<LMPHostType,0,0>;
template class NPairKokkos<LMPHostType,0,1>;
template class NPairKokkos<LMPHostType,1,0>;
template class NPairKokkos<LMPHostType,1,1>;
template class NPairKokkos<LMPHostType,0,0,0>;
template class NPairKokkos<LMPHostType,0,1,0>;
template class NPairKokkos<LMPHostType,1,0,0>;
template class NPairKokkos<LMPHostType,1,1,0>;
template class NPairKokkos<LMPHostType,1,0,1>;
#endif
}

48
src/KOKKOS/npair_kokkos.h
View File

@ -13,42 +13,52 @@
#ifdef NPAIR_CLASS
typedef NPairKokkos<LMPHostType,0,0> NPairKokkosFullBinHost;
typedef NPairKokkos<LMPHostType,0,0,0> NPairKokkosFullBinHost;
NPairStyle(full/bin/kk/host,
NPairKokkosFullBinHost,
NP_FULL | NP_BIN | NP_KOKKOS_HOST | NP_NEWTON | NP_NEWTOFF | NP_ORTHO | NP_TRI)
typedef NPairKokkos<LMPDeviceType,0,0> NPairKokkosFullBinDevice;
typedef NPairKokkos<LMPDeviceType,0,0,0> NPairKokkosFullBinDevice;
NPairStyle(full/bin/kk/device,
NPairKokkosFullBinDevice,
NP_FULL | NP_BIN | NP_KOKKOS_DEVICE | NP_NEWTON | NP_NEWTOFF | NP_ORTHO | NP_TRI)
typedef NPairKokkos<LMPHostType,0,1> NPairKokkosFullBinGhostHost;
typedef NPairKokkos<LMPHostType,0,1,0> NPairKokkosFullBinGhostHost;
NPairStyle(full/bin/ghost/kk/host,
NPairKokkosFullBinGhostHost,
NP_FULL | NP_BIN | NP_KOKKOS_HOST | NP_NEWTON | NP_NEWTOFF | NP_GHOST | NP_ORTHO | NP_TRI)
typedef NPairKokkos<LMPDeviceType,0,1> NPairKokkosFullBinGhostDevice;
typedef NPairKokkos<LMPDeviceType,0,1,0> NPairKokkosFullBinGhostDevice;
NPairStyle(full/bin/ghost/kk/device,
NPairKokkosFullBinGhostDevice,
NP_FULL | NP_BIN | NP_KOKKOS_DEVICE | NP_NEWTON | NP_NEWTOFF | NP_GHOST | NP_ORTHO | NP_TRI)
typedef NPairKokkos<LMPHostType,1,0> NPairKokkosHalfBinHost;
typedef NPairKokkos<LMPHostType,1,0,0> NPairKokkosHalfBinHost;
NPairStyle(half/bin/kk/host,
NPairKokkosHalfBinHost,
NP_HALF | NP_BIN | NP_KOKKOS_HOST | NP_NEWTON | NP_NEWTOFF | NP_ORTHO | NP_TRI)
NP_HALF | NP_BIN | NP_KOKKOS_HOST | NP_NEWTON | NP_NEWTOFF | NP_ORTHO)
typedef NPairKokkos<LMPDeviceType,1,0> NPairKokkosHalfBinDevice;
typedef NPairKokkos<LMPDeviceType,1,0,0> NPairKokkosHalfBinDevice;
NPairStyle(half/bin/kk/device,
NPairKokkosHalfBinDevice,
NP_HALF | NP_BIN | NP_KOKKOS_DEVICE | NP_NEWTON | NP_NEWTOFF | NP_ORTHO | NP_TRI)
NP_HALF | NP_BIN | NP_KOKKOS_DEVICE | NP_NEWTON | NP_NEWTOFF | NP_ORTHO)
typedef NPairKokkos<LMPHostType,1,1> NPairKokkosHalfBinGhostHost;
typedef NPairKokkos<LMPHostType,1,0,1> NPairKokkosHalfBinHostTri;
NPairStyle(half/bin/kk/host,
NPairKokkosHalfBinHostTri,
NP_HALF | NP_BIN | NP_KOKKOS_HOST | NP_NEWTON | NP_NEWTOFF | NP_TRI)
typedef NPairKokkos<LMPDeviceType,1,0,1> NPairKokkosHalfBinDeviceTri;
NPairStyle(half/bin/kk/device,
NPairKokkosHalfBinDeviceTri,
NP_HALF | NP_BIN | NP_KOKKOS_DEVICE | NP_NEWTON | NP_NEWTOFF | NP_TRI)
typedef NPairKokkos<LMPHostType,1,1,0> NPairKokkosHalfBinGhostHost;
NPairStyle(half/bin/ghost/kk/host,
NPairKokkosHalfBinGhostHost,
NP_HALF | NP_BIN | NP_KOKKOS_HOST | NP_NEWTON | NP_NEWTOFF | NP_GHOST | NP_ORTHO | NP_TRI)
typedef NPairKokkos<LMPDeviceType,1,1> NPairKokkosHalfBinGhostDevice;
typedef NPairKokkos<LMPDeviceType,1,1,0> NPairKokkosHalfBinGhostDevice;
NPairStyle(half/bin/ghost/kk/device,
NPairKokkosHalfBinGhostDevice,
NP_HALF | NP_BIN | NP_KOKKOS_DEVICE | NP_NEWTON | NP_NEWTOFF | NP_GHOST | NP_ORTHO | NP_TRI)
@ -63,7 +73,7 @@ NPairStyle(half/bin/ghost/kk/device,
namespace LAMMPS_NS {
template<class DeviceType, int HALF_NEIGH, int GHOST>
template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>
class NPairKokkos : public NPair {
public:
NPairKokkos(class LAMMPS *);
@ -253,7 +263,7 @@ class NeighborKokkosExecute
~NeighborKokkosExecute() {neigh_list.clean_copy();};
template<int HalfNeigh, int Newton>
template<int HalfNeigh, int Newton, int Tri>
KOKKOS_FUNCTION
void build_Item(const int &i) const;
@ -262,7 +272,7 @@ class NeighborKokkosExecute
void build_Item_Ghost(const int &i) const;
#ifdef KOKKOS_HAVE_CUDA
template<int HalfNeigh, int Newton>
template<int HalfNeigh, int Newton, int Tri>
__device__ inline
void build_ItemCuda(typename Kokkos::TeamPolicy<DeviceType>::member_type dev) const;
#endif
@ -354,7 +364,7 @@ class NeighborKokkosExecute
};
template<class DeviceType,int HALF_NEIGH,int GHOST_NEWTON>
template<class DeviceType, int HALF_NEIGH, int GHOST_NEWTON, int TRI>
struct NPairKokkosBuildFunctor {
typedef DeviceType device_type;
@ -367,20 +377,20 @@ struct NPairKokkosBuildFunctor {
KOKKOS_INLINE_FUNCTION
void operator() (const int & i) const {
c.template build_Item<HALF_NEIGH,GHOST_NEWTON>(i);
c.template build_Item<HALF_NEIGH,GHOST_NEWTON,TRI>(i);
}
#ifdef KOKKOS_HAVE_CUDA
__device__ inline
void operator() (typename Kokkos::TeamPolicy<DeviceType>::member_type dev) const {
c.template build_ItemCuda<HALF_NEIGH,GHOST_NEWTON>(dev);
c.template build_ItemCuda<HALF_NEIGH,GHOST_NEWTON,TRI>(dev);
}
size_t shmem_size(const int team_size) const { (void) team_size; return sharedsize; }
#endif
};
template<int HALF_NEIGH,int GHOST_NEWTON>
struct NPairKokkosBuildFunctor<LMPHostType,HALF_NEIGH,GHOST_NEWTON> {
template<int HALF_NEIGH, int GHOST_NEWTON, int TRI>
struct NPairKokkosBuildFunctor<LMPHostType,HALF_NEIGH,GHOST_NEWTON,TRI> {
typedef LMPHostType device_type;
const NeighborKokkosExecute<LMPHostType> c;
@ -392,7 +402,7 @@ struct NPairKokkosBuildFunctor<LMPHostType,HALF_NEIGH,GHOST_NEWTON> {
KOKKOS_INLINE_FUNCTION
void operator() (const int & i) const {
c.template build_Item<HALF_NEIGH,GHOST_NEWTON>(i);
c.template build_Item<HALF_NEIGH,GHOST_NEWTON,TRI>(i);
}
void operator() (typename Kokkos::TeamPolicy<LMPHostType>::member_type dev) const {}

5
src/KOKKOS/pppm_kokkos.cpp
View File

@ -646,10 +646,7 @@ void PPPMKokkos<DeviceType>::compute(int eflag, int vflag)
x = atomKK->k_x.view<DeviceType>();
f = atomKK->k_f.view<DeviceType>();
q = atomKK->k_q.view<DeviceType>();
atomKK->sync(execution_space,datamask_read);
atomKK->modified(execution_space,datamask_modify);
//nlocal = atomKK->nlocal;
//nall = atomKK->nlocal + atomKK->nghost;
//newton_pair = force->newton_pair;