Merge pull request #364 from stanmoore1/kk_triclinic_neighlist (997142a4) · Commits · 郑智淋 / lammps

src/KOKKOS/npair_kokkos.cpp

+55 −23

Original line number	Original line	Diff line number	Diff line
	@@ -24,8 +24,8 @@ namespace LAMMPS_NS {

	/* ---------------------------------------------------------------------- */		/* ---------------------------------------------------------------------- */

	template<class DeviceType, int HALF_NEIGH, int GHOST>		template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>
	NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::NPairKokkos(LAMMPS *lmp) : NPair(lmp) {		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		copy needed info from Neighbor class to this build class
	------------------------------------------------------------------------- */		------------------------------------------------------------------------- */

	template<class DeviceType, int HALF_NEIGH, int GHOST>		template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>
	void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::copy_neighbor_info()		void NPairKokkos<DeviceType,HALF_NEIGH,GHOST,TRI>::copy_neighbor_info()
	{		{
	NPair::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		copy per-atom and per-bin vectors from NBin class to this build class
	------------------------------------------------------------------------- */		------------------------------------------------------------------------- */

	template<class DeviceType, int HALF_NEIGH, int GHOST>		template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>
	void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::copy_bin_info()		void NPairKokkos<DeviceType,HALF_NEIGH,GHOST,TRI>::copy_bin_info()
	{		{
	NPair::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		copy needed info from NStencil class to this build class
	------------------------------------------------------------------------- */		------------------------------------------------------------------------- */

	template<class DeviceType, int HALF_NEIGH, int GHOST>		template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>
	void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::copy_stencil_info()		void NPairKokkos<DeviceType,HALF_NEIGH,GHOST,TRI>::copy_stencil_info()
	{		{
	NPair::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>		template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>
	void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::build(NeighList *list_)		void NPairKokkos<DeviceType,HALF_NEIGH,GHOST,TRI>::build(NeighList *list_)
	{		{
	NeighListKokkos<DeviceType>* list = (NeighListKokkos<DeviceType>*) list_;		NeighListKokkos<DeviceType>* list = (NeighListKokkos<DeviceType>*) list_;
	const int nlocal = includegroup?atom->nfirst:atom->nlocal;		const int nlocal = includegroup?atom->nfirst:atom->nlocal;
	@@ -196,14 +196,14 @@ if (GHOST) {
	Kokkos::parallel_for(nall, f);		Kokkos::parallel_for(nall, f);
	} else {		} else {
	if (newton_pair) {		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		#ifdef KOKKOS_HAVE_CUDA
	Kokkos::parallel_for(config, f);		Kokkos::parallel_for(config, f);
	#else		#else
	Kokkos::parallel_for(nall, f);		Kokkos::parallel_for(nall, f);
	#endif		#endif
	} else {		} 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		#ifdef KOKKOS_HAVE_CUDA
	Kokkos::parallel_for(config, f);		Kokkos::parallel_for(config, f);
	#else		#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>::		void NeighborKokkosExecute<DeviceType>::
	build_Item(const int &i) const		build_Item(const int &i) const
	{		{
	@@ -374,6 +374,16 @@ void NeighborKokkosExecute<DeviceType>::

	if(HalfNeigh && !Newton && (j < i)) continue;		if(HalfNeigh && !Newton && (j < i)) continue;
	if(!HalfNeigh && 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;		if(exclude && exclusion(i,j,itype,jtype)) continue;

	const X_FLOAT delx = xtmp - x(j, 0);		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		__device__ inline
	void NeighborKokkosExecute<DeviceType>::build_ItemCuda(typename Kokkos::TeamPolicy<DeviceType>::member_type dev) const		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) \|\|		((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)))))		(x(j, 2) == ztmp && x(j, 1) == ytmp && x(j, 0) < xtmp)))))
	) continue;		) 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;		if(exclude && exclusion(i,j,itype,jtype)) continue;
	const X_FLOAT delx = xtmp - other_x[m];		const X_FLOAT delx = xtmp - other_x[m];
	const X_FLOAT dely = ytmp - other_x[m + atoms_per_bin];		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 && (j < i)) continue;
	if(HalfNeigh && !Newton && (j < i)) continue;		if(HalfNeigh && !Newton && (j < i)) continue;
	if(!HalfNeigh && 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;		if(exclude && exclusion(i,j,itype,jtype)) continue;

	const X_FLOAT delx = xtmp - other_x[m];		const X_FLOAT delx = xtmp - other_x[m];
	@@ -736,14 +766,16 @@ void NeighborKokkosExecute<DeviceType>::
	}		}

	namespace LAMMPS_NS {		namespace LAMMPS_NS {
	template class NPairKokkos<LMPDeviceType,0,0>;		template class NPairKokkos<LMPDeviceType,0,0,0>;
	template class NPairKokkos<LMPDeviceType,0,1>;		template class NPairKokkos<LMPDeviceType,0,1,0>;
	template class NPairKokkos<LMPDeviceType,1,0>;		template class NPairKokkos<LMPDeviceType,1,0,0>;
	template class NPairKokkos<LMPDeviceType,1,1>;		template class NPairKokkos<LMPDeviceType,1,1,0>;
			template class NPairKokkos<LMPDeviceType,1,0,1>;
	#ifdef KOKKOS_HAVE_CUDA		#ifdef KOKKOS_HAVE_CUDA
	template class NPairKokkos<LMPHostType,0,0>;		template class NPairKokkos<LMPHostType,0,0,0>;
	template class NPairKokkos<LMPHostType,0,1>;		template class NPairKokkos<LMPHostType,0,1,0>;
	template class NPairKokkos<LMPHostType,1,0>;		template class NPairKokkos<LMPHostType,1,0,0>;
	template class NPairKokkos<LMPHostType,1,1>;		template class NPairKokkos<LMPHostType,1,1,0>;
			template class NPairKokkos<LMPHostType,1,0,1>;
	#endif		#endif
	}		}

src/KOKKOS/npair_kokkos.h

+29 −19

Original line number	Original line	Diff line number	Diff line
	@@ -13,42 +13,52 @@

	#ifdef NPAIR_CLASS		#ifdef NPAIR_CLASS

	typedef NPairKokkos<LMPHostType,0,0> NPairKokkosFullBinHost;		typedef NPairKokkos<LMPHostType,0,0,0> NPairKokkosFullBinHost;
	NPairStyle(full/bin/kk/host,		NPairStyle(full/bin/kk/host,
	NPairKokkosFullBinHost,		NPairKokkosFullBinHost,
	NP_FULL \| NP_BIN \| NP_KOKKOS_HOST \| NP_NEWTON \| NP_NEWTOFF \| NP_ORTHO \| NP_TRI)		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,		NPairStyle(full/bin/kk/device,
	NPairKokkosFullBinDevice,		NPairKokkosFullBinDevice,
	NP_FULL \| NP_BIN \| NP_KOKKOS_DEVICE \| NP_NEWTON \| NP_NEWTOFF \| NP_ORTHO \| NP_TRI)		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,		NPairStyle(full/bin/ghost/kk/host,
	NPairKokkosFullBinGhostHost,		NPairKokkosFullBinGhostHost,
	NP_FULL \| NP_BIN \| NP_KOKKOS_HOST \| NP_NEWTON \| NP_NEWTOFF \| NP_GHOST \| NP_ORTHO \| NP_TRI)		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,		NPairStyle(full/bin/ghost/kk/device,
	NPairKokkosFullBinGhostDevice,		NPairKokkosFullBinGhostDevice,
	NP_FULL \| NP_BIN \| NP_KOKKOS_DEVICE \| NP_NEWTON \| NP_NEWTOFF \| NP_GHOST \| NP_ORTHO \| NP_TRI)		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,		NPairStyle(half/bin/kk/host,
	NPairKokkosHalfBinHost,		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,		NPairStyle(half/bin/kk/device,
	NPairKokkosHalfBinDevice,		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,		NPairStyle(half/bin/ghost/kk/host,
	NPairKokkosHalfBinGhostHost,		NPairKokkosHalfBinGhostHost,
	NP_HALF \| NP_BIN \| NP_KOKKOS_HOST \| NP_NEWTON \| NP_NEWTOFF \| NP_GHOST \| NP_ORTHO \| NP_TRI)		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,		NPairStyle(half/bin/ghost/kk/device,
	NPairKokkosHalfBinGhostDevice,		NPairKokkosHalfBinGhostDevice,
	NP_HALF \| NP_BIN \| NP_KOKKOS_DEVICE \| NP_NEWTON \| NP_NEWTOFF \| NP_GHOST \| NP_ORTHO \| NP_TRI)		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 {		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 {		class NPairKokkos : public NPair {
	public:		public:
	NPairKokkos(class LAMMPS *);		NPairKokkos(class LAMMPS *);
	@@ -252,7 +262,7 @@ class NeighborKokkosExecute

	~NeighborKokkosExecute() {neigh_list.clean_copy();};		~NeighborKokkosExecute() {neigh_list.clean_copy();};

	template<int HalfNeigh, int Newton>		template<int HalfNeigh, int Newton, int Tri>
	KOKKOS_FUNCTION		KOKKOS_FUNCTION
	void build_Item(const int &i) const;		void build_Item(const int &i) const;

	@@ -261,7 +271,7 @@ class NeighborKokkosExecute
	void build_Item_Ghost(const int &i) const;		void build_Item_Ghost(const int &i) const;

	#ifdef KOKKOS_HAVE_CUDA		#ifdef KOKKOS_HAVE_CUDA
	template<int HalfNeigh, int Newton>		template<int HalfNeigh, int Newton, int Tri>
	__device__ inline		__device__ inline
	void build_ItemCuda(typename Kokkos::TeamPolicy<DeviceType>::member_type dev) const;		void build_ItemCuda(typename Kokkos::TeamPolicy<DeviceType>::member_type dev) const;
	#endif		#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 {		struct NPairKokkosBuildFunctor {
	typedef DeviceType device_type;		typedef DeviceType device_type;

	@@ -366,20 +376,20 @@ struct NPairKokkosBuildFunctor {

	KOKKOS_INLINE_FUNCTION		KOKKOS_INLINE_FUNCTION
	void operator() (const int & i) const {		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		#ifdef KOKKOS_HAVE_CUDA
	__device__ inline		__device__ inline

	void operator() (typename Kokkos::TeamPolicy<DeviceType>::member_type dev) const {		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; }		size_t shmem_size(const int team_size) const { (void) team_size; return sharedsize; }
	#endif		#endif
	};		};

	template<int HALF_NEIGH,int GHOST_NEWTON>		template<int HALF_NEIGH, int GHOST_NEWTON, int TRI>
	struct NPairKokkosBuildFunctor<LMPHostType,HALF_NEIGH,GHOST_NEWTON> {		struct NPairKokkosBuildFunctor<LMPHostType,HALF_NEIGH,GHOST_NEWTON,TRI> {
	typedef LMPHostType device_type;		typedef LMPHostType device_type;

	const NeighborKokkosExecute<LMPHostType> c;		const NeighborKokkosExecute<LMPHostType> c;
	@@ -391,7 +401,7 @@ struct NPairKokkosBuildFunctor<LMPHostType,HALF_NEIGH,GHOST_NEWTON> {

	KOKKOS_INLINE_FUNCTION		KOKKOS_INLINE_FUNCTION
	void operator() (const int & i) const {		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 {}		void operator() (typename Kokkos::TeamPolicy<LMPHostType>::member_type dev) const {}

Admin message