From 5afd3e995b746b729951c5b02b30e849b5b425d6 Mon Sep 17 00:00:00 2001 From: Stan Moore Date: Fri, 27 Jan 2017 14:18:01 -0700 Subject: [PATCH] Adding support to npair_kokkos for triclinic-newton-on neighborlists --- src/KOKKOS/npair_kokkos.cpp | 68 +++++++++++++++++++++++-------------- src/KOKKOS/npair_kokkos.h | 48 +++++++++++++++----------- 2 files changed, 71 insertions(+), 45 deletions(-) diff --git a/src/KOKKOS/npair_kokkos.cpp b/src/KOKKOS/npair_kokkos.cpp index fd32cd463e..0c2f03ce6d 100644 --- a/src/KOKKOS/npair_kokkos.cpp +++ b/src/KOKKOS/npair_kokkos.cpp @@ -24,8 +24,8 @@ namespace LAMMPS_NS { /* ---------------------------------------------------------------------- */ -template -NPairKokkos::NPairKokkos(LAMMPS *lmp) : NPair(lmp) { +template +NPairKokkos::NPairKokkos(LAMMPS *lmp) : NPair(lmp) { } @@ -33,8 +33,8 @@ NPairKokkos::NPairKokkos(LAMMPS *lmp) : NPair(lmp) copy needed info from Neighbor class to this build class ------------------------------------------------------------------------- */ -template -void NPairKokkos::copy_neighbor_info() +template +void NPairKokkos::copy_neighbor_info() { NPair::copy_neighbor_info(); @@ -62,8 +62,8 @@ void NPairKokkos::copy_neighbor_info() copy per-atom and per-bin vectors from NBin class to this build class ------------------------------------------------------------------------- */ -template -void NPairKokkos::copy_bin_info() +template +void NPairKokkos::copy_bin_info() { NPair::copy_bin_info(); @@ -78,8 +78,8 @@ void NPairKokkos::copy_bin_info() copy needed info from NStencil class to this build class ------------------------------------------------------------------------- */ -template -void NPairKokkos::copy_stencil_info() +template +void NPairKokkos::copy_stencil_info() { NPair::copy_stencil_info(); @@ -106,8 +106,8 @@ void NPairKokkos::copy_stencil_info() /* ---------------------------------------------------------------------- */ -template -void NPairKokkos::build(NeighList *list_) +template +void NPairKokkos::build(NeighList *list_) { NeighListKokkos* list = (NeighListKokkos*) 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 f(data,atoms_per_bin * 5 * sizeof(X_FLOAT) * factor); + NPairKokkosBuildFunctor 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 f(data,atoms_per_bin * 5 * sizeof(X_FLOAT) * factor); + NPairKokkosBuildFunctor 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::exclusion(const int &i,const int &j, /* ---------------------------------------------------------------------- */ -template template +template template void NeighborKokkosExecute:: build_Item(const int &i) const { @@ -317,7 +321,7 @@ void NeighborKokkosExecute:: = 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:: const int jbin = ibin + stencil[k]; // get subview of jbin - if(HalfNeigh&&(ibin==jbin)) continue; + if(HalfNeigh && !Tri && (ibin==jbin)) continue; //const ArrayTypes::t_int_1d_const_um =Kokkos::subview(bins,jbin,ALL); for(int m = 0; m < c_bincount(jbin); m++) { @@ -374,6 +378,16 @@ void NeighborKokkosExecute:: 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 template +template template __device__ inline void NeighborKokkosExecute::build_ItemCuda(typename Kokkos::TeamPolicy::member_type dev) const { @@ -736,14 +750,16 @@ void NeighborKokkosExecute:: } namespace LAMMPS_NS { -template class NPairKokkos; -template class NPairKokkos; -template class NPairKokkos; -template class NPairKokkos; +template class NPairKokkos; +template class NPairKokkos; +template class NPairKokkos; +template class NPairKokkos; +template class NPairKokkos; #ifdef KOKKOS_HAVE_CUDA -template class NPairKokkos; -template class NPairKokkos; -template class NPairKokkos; -template class NPairKokkos; +template class NPairKokkos; +template class NPairKokkos; +template class NPairKokkos; +template class NPairKokkos; +template class NPairKokkos; #endif } diff --git a/src/KOKKOS/npair_kokkos.h b/src/KOKKOS/npair_kokkos.h index 4b77175191..d246209b6a 100644 --- a/src/KOKKOS/npair_kokkos.h +++ b/src/KOKKOS/npair_kokkos.h @@ -13,42 +13,52 @@ #ifdef NPAIR_CLASS -typedef NPairKokkos NPairKokkosFullBinHost; +typedef NPairKokkos NPairKokkosFullBinHost; NPairStyle(full/bin/kk/host, NPairKokkosFullBinHost, NP_FULL | NP_BIN | NP_KOKKOS_HOST | NP_NEWTON | NP_NEWTOFF | NP_ORTHO | NP_TRI) -typedef NPairKokkos NPairKokkosFullBinDevice; +typedef NPairKokkos NPairKokkosFullBinDevice; NPairStyle(full/bin/kk/device, NPairKokkosFullBinDevice, NP_FULL | NP_BIN | NP_KOKKOS_DEVICE | NP_NEWTON | NP_NEWTOFF | NP_ORTHO | NP_TRI) -typedef NPairKokkos NPairKokkosFullBinGhostHost; +typedef NPairKokkos 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 NPairKokkosFullBinGhostDevice; +typedef NPairKokkos 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 NPairKokkosHalfBinHost; +typedef NPairKokkos 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 NPairKokkosHalfBinDevice; +typedef NPairKokkos 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 NPairKokkosHalfBinGhostHost; +typedef NPairKokkos NPairKokkosHalfBinHostTri; +NPairStyle(half/bin/kk/host, + NPairKokkosHalfBinHostTri, + NP_HALF | NP_BIN | NP_KOKKOS_HOST | NP_NEWTON | NP_NEWTOFF | NP_TRI) + +typedef NPairKokkos NPairKokkosHalfBinDeviceTri; +NPairStyle(half/bin/kk/device, + NPairKokkosHalfBinDeviceTri, + NP_HALF | NP_BIN | NP_KOKKOS_DEVICE | NP_NEWTON | NP_NEWTOFF | NP_TRI) + +typedef NPairKokkos 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 NPairKokkosHalfBinGhostDevice; +typedef NPairKokkos 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 +template class NPairKokkos : public NPair { public: NPairKokkos(class LAMMPS *); @@ -252,7 +262,7 @@ class NeighborKokkosExecute ~NeighborKokkosExecute() {neigh_list.clean_copy();}; - template + template 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 + template __device__ inline void build_ItemCuda(typename Kokkos::TeamPolicy::member_type dev) const; #endif @@ -353,7 +363,7 @@ class NeighborKokkosExecute }; -template +template 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(i); + c.template build_Item(i); } #ifdef KOKKOS_HAVE_CUDA __device__ inline void operator() (typename Kokkos::TeamPolicy::member_type dev) const { - c.template build_ItemCuda(dev); + c.template build_ItemCuda(dev); } size_t shmem_size(const int team_size) const { (void) team_size; return sharedsize; } #endif }; -template -struct NPairKokkosBuildFunctor { +template +struct NPairKokkosBuildFunctor { typedef LMPHostType device_type; const NeighborKokkosExecute c; @@ -391,7 +401,7 @@ struct NPairKokkosBuildFunctor { KOKKOS_INLINE_FUNCTION void operator() (const int & i) const { - c.template build_Item(i); + c.template build_Item(i); } void operator() (typename Kokkos::TeamPolicy::member_type dev) const {}