Implementing forward/border comm in Kokkos. (9765a9a4) · Commits · 郑智淋 / lammps

src/KOKKOS/atom_vec_kokkos.h

+19 −0

Original line number	Diff line number	Diff line
		@@ -60,6 +60,15 @@ class AtomVecKokkos : public AtomVec {
		unpack_comm_kokkos(const int &n, const int &nfirst,
		const DAT::tdual_xfloat_2d &buf);

		virtual int
		pack_comm_vel_kokkos(const int &n, const DAT::tdual_int_2d &list,
		const int & iswap, const DAT::tdual_xfloat_2d &buf,
		const int &pbc_flag, const int pbc[]) { return 0; }

		virtual void
		unpack_comm_vel_kokkos(const int &n, const int &nfirst,
		const DAT::tdual_xfloat_2d &buf) {}

		virtual int
		unpack_reverse_self(const int &n, const DAT::tdual_int_2d &list,
		const int & iswap, const int nfirst);
		@@ -82,6 +91,16 @@ class AtomVecKokkos : public AtomVec {
		const DAT::tdual_xfloat_2d &buf,
		ExecutionSpace space) = 0;

		virtual int
		pack_border_vel_kokkos(int n, DAT::tdual_int_2d k_sendlist,
		DAT::tdual_xfloat_2d buf,int iswap,
		int pbc_flag, int *pbc, ExecutionSpace space) { return 0; }

		virtual void
		unpack_border_vel_kokkos(const int &n, const int &nfirst,
		const DAT::tdual_xfloat_2d &buf,
		ExecutionSpace space) {}

		virtual int
		pack_exchange_kokkos(const int &nsend, DAT::tdual_xfloat_2d &buf,
		DAT::tdual_int_1d k_sendlist,

src/KOKKOS/atom_vec_sphere_kokkos.cpp

+626 −37

File changed.

Preview size limit exceeded, changes collapsed.

src/KOKKOS/atom_vec_sphere_kokkos.h

+12 −0

Original line number	Diff line number	Diff line
		@@ -75,6 +75,12 @@ class AtomVecSphereKokkos : public AtomVecKokkos {
		const int &pbc_flag, const int pbc[]);
		void unpack_comm_kokkos(const int &n, const int &nfirst,
		const DAT::tdual_xfloat_2d &buf);
		int pack_comm_vel_kokkos(const int &n, const DAT::tdual_int_2d &k_sendlist,
		const int & iswap,
		const DAT::tdual_xfloat_2d &buf,
		const int &pbc_flag, const int pbc[]);
		void unpack_comm_vel_kokkos(const int &n, const int &nfirst,
		const DAT::tdual_xfloat_2d &buf);
		int pack_comm_self(const int &n, const DAT::tdual_int_2d &list,
		const int & iswap, const int nfirst,
		const int &pbc_flag, const int pbc[]);
		@@ -84,6 +90,12 @@ class AtomVecSphereKokkos : public AtomVecKokkos {
		void unpack_border_kokkos(const int &n, const int &nfirst,
		const DAT::tdual_xfloat_2d &buf,
		ExecutionSpace space);
		int pack_border_vel_kokkos(int n, DAT::tdual_int_2d k_sendlist,
		DAT::tdual_xfloat_2d buf,int iswap,
		int pbc_flag, int *pbc, ExecutionSpace space);
		void unpack_border_vel_kokkos(const int &n, const int &nfirst,
		const DAT::tdual_xfloat_2d &buf,
		ExecutionSpace space);
		int pack_exchange_kokkos(const int &nsend,DAT::tdual_xfloat_2d &buf,
		DAT::tdual_int_1d k_sendlist,
		DAT::tdual_int_1d k_copylist,

src/KOKKOS/comm_kokkos.cpp

+54 −39

Original line number	Diff line number	Diff line
		@@ -134,9 +134,6 @@ void CommKokkos::init()
		if (force->newton == 0) check_reverse = 0;
		if (force->pair) check_reverse += force->pair->comm_reverse_off;

		if (ghost_velocity)
		forward_comm_classic = true;

		if (!comm_f_only) // not all Kokkos atom_vec styles have reverse pack/unpack routines yet
		reverse_comm_classic = true;
		}
		@@ -210,17 +207,21 @@ void CommKokkos::forward_comm_device(int dummy)
		space,X_MASK);
		}
		} else if (ghost_velocity) {
		error->all(FLERR,"Ghost velocity forward comm not yet "
		"implemented with Kokkos");
		if (size_forward_recv[iswap])
		MPI_Irecv(k_buf_recv.view<LMPHostType>().data(),
		if (size_forward_recv[iswap]) {
		MPI_Irecv(k_buf_recv.view<DeviceType>().data(),
		size_forward_recv[iswap],MPI_DOUBLE,
		recvproc[iswap],0,world,&request);
		n = avec->pack_comm_vel(sendnum[iswap],sendlist[iswap],
		buf_send,pbc_flag[iswap],pbc[iswap]);
		if (n) MPI_Send(buf_send,n,MPI_DOUBLE,sendproc[iswap],0,world);
		}
		n = avec->pack_comm_vel_kokkos(sendnum[iswap],k_sendlist,iswap,
		k_buf_send,pbc_flag[iswap],pbc[iswap]);
		DeviceType::fence();
		if (n) {
		MPI_Send(k_buf_send.view<DeviceType>().data(),n,
		MPI_DOUBLE,sendproc[iswap],0,world);
		}
		if (size_forward_recv[iswap]) MPI_Wait(&request,MPI_STATUS_IGNORE);
		avec->unpack_comm_vel(recvnum[iswap],firstrecv[iswap],buf_recv);
		avec->unpack_comm_vel_kokkos(recvnum[iswap],firstrecv[iswap],k_buf_recv);
		DeviceType::fence();
		} else {
		if (size_forward_recv[iswap])
		MPI_Irecv(k_buf_recv.view<DeviceType>().data(),
		@@ -236,18 +237,18 @@ void CommKokkos::forward_comm_device(int dummy)
		avec->unpack_comm_kokkos(recvnum[iswap],firstrecv[iswap],k_buf_recv);
		DeviceType::fence();
		}

		} else {
		if (!ghost_velocity) {
		if (sendnum[iswap])
		n = avec->pack_comm_self(sendnum[iswap],k_sendlist,iswap,
		firstrecv[iswap],pbc_flag[iswap],pbc[iswap]);
		} else if (ghost_velocity) {
		error->all(FLERR,"Ghost velocity forward comm not yet "
		"implemented with Kokkos");
		n = avec->pack_comm_vel(sendnum[iswap],sendlist[iswap],
		buf_send,pbc_flag[iswap],pbc[iswap]);
		avec->unpack_comm_vel(recvnum[iswap],firstrecv[iswap],buf_send);
		DeviceType::fence();
		} else {
		n = avec->pack_comm_vel_kokkos(sendnum[iswap],k_sendlist,iswap,
		k_buf_send,pbc_flag[iswap],pbc[iswap]);
		DeviceType::fence();
		avec->unpack_comm_vel_kokkos(recvnum[iswap],firstrecv[iswap],k_buf_send);
		DeviceType::fence();
		}
		}
		}
		@@ -613,10 +614,9 @@ void CommKokkos::exchange_device()
		nsend =
		avec->pack_exchange_kokkos(k_count.h_view(),k_buf_send,
		k_exchange_sendlist,k_exchange_copylist,
		ExecutionSpaceFromDevice<DeviceType>::
		space,dim,lo,hi);
		ExecutionSpaceFromDevice<DeviceType>::space,
		dim,lo,hi);
		DeviceType::fence();

		} else {
		while (i < nlocal) {
		if (x[i][dim] < lo \|\| x[i][dim] >= hi) {
		@@ -707,7 +707,7 @@ void CommKokkos::borders()
		if (!exchange_comm_classic) {
		static int print = 1;

		if (mode != Comm::SINGLE \|\| bordergroup \|\| ghost_velocity) {
		if (mode != Comm::SINGLE \|\| bordergroup) {
		if (print && comm->me==0) {
		error->warning(FLERR,"Required border comm not yet implemented in Kokkos communication, "
		"switching to classic communication");
		@@ -929,10 +929,10 @@ void CommKokkos::borders_device() {
		if (nsend*size_border > maxsend)
		grow_send_kokkos(nsend*size_border,0);
		if (ghost_velocity) {
		error->all(FLERR,"Required border comm not yet "
		"implemented with Kokkos");
		n = avec->pack_border_vel(nsend,sendlist[iswap],buf_send,
		pbc_flag[iswap],pbc[iswap]);
		n = avec->
		pack_border_vel_kokkos(nsend,k_sendlist,k_buf_send,iswap,
		pbc_flag[iswap],pbc[iswap],exec_space);
		DeviceType::fence();
		}
		else {
		n = avec->
		@@ -965,11 +965,16 @@ void CommKokkos::borders_device() {
		// unpack buffer

		if (ghost_velocity) {
		error->all(FLERR,"Required border comm not yet "
		"implemented with Kokkos");
		avec->unpack_border_vel(nrecv,atom->nlocal+atom->nghost,buf);
		if (sendproc[iswap] != me) {
		avec->unpack_border_vel_kokkos(nrecv,atom->nlocal+atom->nghost,
		k_buf_recv,exec_space);
		DeviceType::fence();
		} else {
		avec->unpack_border_vel_kokkos(nrecv,atom->nlocal+atom->nghost,
		k_buf_send,exec_space);
		DeviceType::fence();
		}
		else
		} else {
		if (sendproc[iswap] != me) {
		avec->unpack_border_kokkos(nrecv,atom->nlocal+atom->nghost,
		k_buf_recv,exec_space);
		@@ -979,7 +984,7 @@ void CommKokkos::borders_device() {
		k_buf_send,exec_space);
		DeviceType::fence();
		}

		}
		// set all pointers & counters

		smax = MAX(smax,nsend);
		@@ -1046,10 +1051,20 @@ void CommKokkos::grow_send_kokkos(int n, int flag, ExecutionSpace space)
		else
		k_buf_send.modify<LMPHostType>();

		if (ghost_velocity)
		k_buf_send.resize(maxsend_border,
		atom->avec->size_border + atom->avec->size_velocity);
		else
		k_buf_send.resize(maxsend_border,atom->avec->size_border);
		buf_send = k_buf_send.view<LMPHostType>().data();
		}
		else {
		if (ghost_velocity)
		k_buf_send = DAT::
		tdual_xfloat_2d("comm:k_buf_send",
		maxsend_border,
		atom->avec->size_border + atom->avec->size_velocity);
		else
		k_buf_send = DAT::
		tdual_xfloat_2d("comm:k_buf_send",maxsend_border,atom->avec->size_border);
		buf_send = k_buf_send.view<LMPHostType>().data();

Admin message