forked from lijiext/lammps
Adding support to npair_kokkos for triclinic-newton-on neighborlists
This commit is contained in:
Stan Moore
parent
f8506fee23
commit
5afd3e995b
|
|
@ -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,18 @@ 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,HALF_NEIGH,1,TRI> f(data,atoms_per_bin * 5 * sizeof(X_FLOAT) * factor);
|
||||
if (TRI) // CUDA specializations don't yet support triclinic
|
||||
Kokkos::parallel_for(nall, f);
|
||||
else {
|
||||
#ifdef KOKKOS_HAVE_CUDA
|
||||
Kokkos::parallel_for(config, f);
|
||||
Kokkos::parallel_for(config, f);
|
||||
#else
|
||||
Kokkos::parallel_for(nall, f);
|
||||
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 +297,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
|
||||
{
|
||||
|
|
@ -317,7 +321,7 @@ void NeighborKokkosExecute<DeviceType>::
|
|||
= d_stencil;
|
||||
|
||||
// loop over all bins in neighborhood (includes ibin)
|
||||
if(HalfNeigh)
|
||||
if(HalfNeigh && !Tri)
|
||||
for(int m = 0; m < c_bincount(ibin); m++) {
|
||||
const int j = c_bins(ibin,m);
|
||||
const int jtype = type(j);
|
||||
|
|
@ -365,7 +369,7 @@ void NeighborKokkosExecute<DeviceType>::
|
|||
const int jbin = ibin + stencil[k];
|
||||
|
||||
// get subview of jbin
|
||||
if(HalfNeigh&&(ibin==jbin)) continue;
|
||||
if(HalfNeigh && !Tri && (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 +378,16 @@ void NeighborKokkosExecute<DeviceType>::
|
|||
|
||||
if(HalfNeigh && !Newton && (j < i)) continue;
|
||||
if(!HalfNeigh && j==i) continue;
|
||||
if(Tri && Newton) {
|
||||
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 +442,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
|
||||
{
|
||||
|
|
@ -736,14 +750,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
|
||||
}
|
||||
|
|
|
|||
|
|
@ -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 *);
|
||||
|
|
@ -252,7 +262,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;
|
||||
|
||||
|
|
@ -261,7 +271,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
|
||||
|
|
@ -353,7 +363,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;
|
||||
|
||||
|
|
@ -366,20 +376,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;
|
||||
|
|
@ -391,7 +401,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 {}
|
||||
|
|
|
|||
Loading…
Reference in New Issue