Merge pull request #2289 from akohlmey/reduce-kokkos-compiler-warnings (a70a5be2) · Commits · 郑智淋 / lammps

src/KOKKOS/atom_kokkos.h

+9 −9

Original line number	Diff line number	Diff line
		@@ -78,7 +78,7 @@ class AtomKokkos : public Atom {
		};

		template<class ViewType, class IndexView>
		class SortFunctor {
		struct SortFunctor {
		typedef typename ViewType::device_type device_type;
		ViewType source;
		Kokkos::View<typename ViewType::non_const_data_type,typename ViewType::array_type,device_type> dest;
		@@ -100,18 +100,18 @@ class SortFunctor {
		dest(i) = source(index(i));
		}
		void operator()(const typename std::enable_if<ViewType::rank==2, int>::type& i) {
		for(int j=0;j<source.extent(1);j++)
		for(int j=0; j < (int)source.extent(1); j++)
		dest(i,j) = source(index(i),j);
		}
		void operator()(const typename std::enable_if<ViewType::rank==3, int>::type& i) {
		for(int j=0;j<source.extent(1);j++)
		for(int k=0;k<source.extent(2);k++)
		for(int j=0; j < (int)source.extent(1); j++)
		for(int k=0; k < (int)source.extent(2); k++)
		dest(i,j,k) = source(index(i),j,k);
		}
		void operator()(const typename std::enable_if<ViewType::rank==4, int>::type& i) {
		for(int j=0;j<source.extent(1);j++)
		for(int k=0;k<source.extent(2);k++)
		for(int l=0;l<source.extent(3);l++)
		for(int j=0; j < (int)source.extent(1); j++)
		for(int k=0; k < (int)source.extent(2); k++)
		for(int l=0; l < (int)source.extent(3); l++)
		dest(i,j,k,l) = source(index(i),j,k,l);
		}
		};

src/KOKKOS/atom_vec_hybrid_kokkos.cpp

+23 −23

Original line number	Diff line number	Diff line
		@@ -252,51 +252,51 @@ void AtomVecHybridKokkos::force_clear(int n, size_t nbytes)

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

		int AtomVecHybridKokkos::pack_comm_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[])
		int AtomVecHybridKokkos::pack_comm_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[])
		{
		error->all(FLERR,"AtomVecHybridKokkos doesn't yet support threaded comm");
		return 0;
		}
		void AtomVecHybridKokkos::unpack_comm_kokkos(const int &n, const int &nfirst,
		const DAT::tdual_xfloat_2d &buf)
		void AtomVecHybridKokkos::unpack_comm_kokkos(const int &/n/, const int &/nfirst/,
		const DAT::tdual_xfloat_2d &/buf/)
		{
		error->all(FLERR,"AtomVecHybridKokkos doesn't yet support threaded comm");
		}
		int AtomVecHybridKokkos::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[])
		int AtomVecHybridKokkos::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[])
		{
		error->all(FLERR,"AtomVecHybridKokkos doesn't yet support threaded comm");
		return 0;
		}
		int AtomVecHybridKokkos::pack_border_kokkos(int n, DAT::tdual_int_2d k_sendlist,
		DAT::tdual_xfloat_2d buf,int iswap,
		int pbc_flag, int *pbc, ExecutionSpace space)
		int AtomVecHybridKokkos::pack_border_kokkos(int /n/, DAT::tdual_int_2d /k_sendlist/,
		DAT::tdual_xfloat_2d /buf/,int /iswap/,
		int /pbc_flag/, int * /pbc/, ExecutionSpace /space/)
		{
		error->all(FLERR,"AtomVecHybridKokkos doesn't yet support threaded comm");
		return 0;
		}
		void AtomVecHybridKokkos::unpack_border_kokkos(const int &n, const int &nfirst,
		const DAT::tdual_xfloat_2d &buf,
		ExecutionSpace space)
		void AtomVecHybridKokkos::unpack_border_kokkos(const int &/n/, const int &/nfirst/,
		const DAT::tdual_xfloat_2d &/buf/,
		ExecutionSpace /space/)
		{
		error->all(FLERR,"AtomVecHybridKokkos doesn't yet support threaded comm");
		}
		int AtomVecHybridKokkos::pack_exchange_kokkos(const int &nsend,DAT::tdual_xfloat_2d &buf,
		DAT::tdual_int_1d k_sendlist,
		DAT::tdual_int_1d k_copylist,
		ExecutionSpace space, int dim,
		X_FLOAT lo, X_FLOAT hi)
		int AtomVecHybridKokkos::pack_exchange_kokkos(const int &/nsend/,DAT::tdual_xfloat_2d &/buf/,
		DAT::tdual_int_1d /k_sendlist/,
		DAT::tdual_int_1d /k_copylist/,
		ExecutionSpace /space/, int /dim/,
		X_FLOAT /lo/, X_FLOAT /hi/)
		{
		error->all(FLERR,"AtomVecHybridKokkos doesn't yet support threaded comm");
		return 0;
		}
		int AtomVecHybridKokkos::unpack_exchange_kokkos(DAT::tdual_xfloat_2d &k_buf, int nrecv,
		int nlocal, int dim, X_FLOAT lo, X_FLOAT hi,
		ExecutionSpace space)
		int AtomVecHybridKokkos::unpack_exchange_kokkos(DAT::tdual_xfloat_2d & /k_buf/, int /nrecv/,
		int /nlocal/, int /dim/, X_FLOAT /lo/,
		X_FLOAT /hi/, ExecutionSpace /space/)
		{
		error->all(FLERR,"AtomVecHybridKokkos doesn't yet support threaded comm");
		return 0;

src/KOKKOS/comm_kokkos.cpp

+17 −21

Original line number	Diff line number	Diff line
		@@ -172,7 +172,7 @@ void CommKokkos::forward_comm(int dummy)
		/* ---------------------------------------------------------------------- */

		template<class DeviceType>
		void CommKokkos::forward_comm_device(int dummy)
		void CommKokkos::forward_comm_device(int)
		{
		int n;
		MPI_Request request;
		@@ -529,19 +529,18 @@ struct BuildExchangeListFunctor {
		typename AT::tdual_int_scalar nsend,
		typename AT::tdual_int_1d sendflag,int nlocal, int dim,
		X_FLOAT lo, X_FLOAT hi):
		_lo(lo),_hi(hi),
		_x(x.template view<DeviceType>()),
		_sendlist(sendlist.template view<DeviceType>()),
		_nsend(nsend.template view<DeviceType>()),
		_sendflag(sendflag.template view<DeviceType>()),
		_nlocal(nlocal),_dim(dim),
		_lo(lo),_hi(hi){
		}
		_nsend(nsend.template view<DeviceType>()),
		_sendlist(sendlist.template view<DeviceType>()),
		_sendflag(sendflag.template view<DeviceType>()) { }

		KOKKOS_INLINE_FUNCTION
		void operator() (int i) const {
		if (_x(i,_dim) < _lo \|\| _x(i,_dim) >= _hi) {
		const int mysend=Kokkos::atomic_fetch_add(&_nsend(),1);
		if(mysend<_sendlist.extent(0)) {
		if(mysend < (int)_sendlist.extent(0)) {
		_sendlist(mysend) = i;
		_sendflag(i) = 1;
		}
		@@ -599,10 +598,10 @@ void CommKokkos::exchange_device()
		i = nsend = 0;

		if (true) {
		if (k_sendflag.h_view.extent(0)<nlocal) k_sendflag.resize(nlocal);
		if ((int)k_sendflag.h_view.extent(0) < nlocal) k_sendflag.resize(nlocal);
		k_sendflag.sync<DeviceType>();
		k_count.h_view() = k_exchange_sendlist.h_view.extent(0);
		while (k_count.h_view()>=k_exchange_sendlist.h_view.extent(0)) {
		while (k_count.h_view() >= (int)k_exchange_sendlist.h_view.extent(0)) {
		k_count.h_view() = 0;
		k_count.modify<LMPHostType>();
		k_count.sync<DeviceType>();
		@@ -616,7 +615,7 @@ void CommKokkos::exchange_device()
		k_count.modify<DeviceType>();

		k_count.sync<LMPHostType>();
		if (k_count.h_view()>=k_exchange_sendlist.h_view.extent(0)) {
		if (k_count.h_view() >= (int)k_exchange_sendlist.h_view.extent(0)) {
		k_exchange_lists.resize(2,k_count.h_view()*1.1);
		k_exchange_sendlist = Kokkos::subview(k_exchange_lists,0,Kokkos::ALL);
		k_exchange_copylist = Kokkos::subview(k_exchange_lists,1,Kokkos::ALL);
		@@ -782,11 +781,10 @@ struct BuildBorderListFunctor {
		int _nlast, int _dim,
		X_FLOAT _lo, X_FLOAT _hi, int _iswap,
		int _maxsendlist):
		x(_x.template view<DeviceType>()),
		lo(_lo),hi(_hi),x(_x.template view<DeviceType>()),iswap(_iswap),
		maxsendlist(_maxsendlist),nfirst(_nfirst),nlast(_nlast),dim(_dim),
		sendlist(_sendlist.template view<DeviceType>()),
		nsend(_nsend.template view<DeviceType>()),
		nfirst(_nfirst),nlast(_nlast),dim(_dim),
		lo(_lo),hi(_hi),iswap(_iswap),maxsendlist(_maxsendlist){}
		nsend(_nsend.template view<DeviceType>()) {}


		KOKKOS_INLINE_FUNCTION
		@@ -823,7 +821,7 @@ void CommKokkos::borders_device() {
		double lo,hi;
		int *type;
		double **x;
		double buf,mlo,*mhi;
		double mlo,mhi;
		MPI_Request request;
		AtomVecKokkos avec = (AtomVecKokkos ) atom->avec;

		@@ -992,10 +990,8 @@ void CommKokkos::borders_device() {
		if (n) MPI_Send(k_buf_send.view<DeviceType>().data(),n,
		MPI_DOUBLE,sendproc[iswap],0,world);
		if (nrecv) MPI_Wait(&request,MPI_STATUS_IGNORE);
		buf = buf_recv;
		} else {
		nrecv = nsend;
		buf = buf_send;
		}

		// unpack buffer
		@@ -1058,7 +1054,7 @@ void CommKokkos::borders_device() {

		void CommKokkos::copy_swap_info()
		{
		if (nswap > k_swap.extent(1)) {
		if (nswap > (int)k_swap.extent(1)) {
		k_swap = DAT::tdual_int_2d("comm:swap",2,nswap);
		k_firstrecv = Kokkos::subview(k_swap,0,Kokkos::ALL);
		k_sendnum_scan = Kokkos::subview(k_swap,1,Kokkos::ALL);
		@@ -1076,7 +1072,7 @@ void CommKokkos::copy_swap_info()

		k_sendlist.sync<LMPHostType>();

		if (totalsend > k_pbc.extent(0)) {
		if (totalsend > (int)k_pbc.extent(0)) {
		k_pbc = DAT::tdual_int_2d("comm:pbc",totalsend,6);
		k_swap2 = DAT::tdual_int_2d("comm:swap2",2,totalsend);
		k_pbc_flag = Kokkos::subview(k_swap2,0,Kokkos::ALL);
		@@ -1174,7 +1170,7 @@ void CommKokkos::grow_send_kokkos(int n, int flag, ExecutionSpace space)
		free/malloc the size of the recv buffer as needed with BUFFACTOR
		------------------------------------------------------------------------- */

		void CommKokkos::grow_recv_kokkos(int n, ExecutionSpace space)
		void CommKokkos::grow_recv_kokkos(int n, ExecutionSpace /space/)
		{
		maxrecv = static_cast<int> (BUFFACTOR * n);
		int maxrecv_border = (maxrecv+BUFEXTRA+5)/atom->avec->size_border + 2;
		@@ -1187,7 +1183,7 @@ void CommKokkos::grow_recv_kokkos(int n, ExecutionSpace space)
		realloc the size of the iswap sendlist as needed with BUFFACTOR
		------------------------------------------------------------------------- */

		void CommKokkos::grow_list(int iswap, int n)
		void CommKokkos::grow_list(int /iswap/, int n)
		{
		int size = static_cast<int> (BUFFACTOR * n);

src/KOKKOS/compute_orientorder_atom_kokkos.cpp

+1 −3

Original line number	Diff line number	Diff line
		@@ -435,10 +435,8 @@ void ComputeOrientOrderAtomKokkos<DeviceType>::select3(int k, int n, int ii) con

		template<class DeviceType>
		KOKKOS_INLINE_FUNCTION
		void ComputeOrientOrderAtomKokkos<DeviceType>::calc_boop1(int ncount, int ii, int ineigh) const
		void ComputeOrientOrderAtomKokkos<DeviceType>::calc_boop1(int /ncount/, int ii, int ineigh) const
		{
		const int i = d_ilist[ii + chunk_offset];

		const double r0 = d_rlist(ii,ineigh,0);
		const double r1 = d_rlist(ii,ineigh,1);
		const double r2 = d_rlist(ii,ineigh,2);

src/KOKKOS/fix_dpd_energy_kokkos.h

+1 −1

Original line number	Diff line number	Diff line
		@@ -27,7 +27,7 @@ FixStyle(dpd/energy/kk/host,FixDPDenergyKokkos<LMPHostType>)

		namespace LAMMPS_NS {

		template <typename DeviceType>
		template <class DeviceType>
		class FixDPDenergyKokkos : public FixDPDenergy {
		public:
		FixDPDenergyKokkos(class LAMMPS , int, char *);

Admin message