Merge branch 'master' of https://www.github.com/lammps/lammps (e1361a9d) · Commits · 郑智淋 / lammps

lib/kokkos/core/src/impl/Kokkos_Atomic_Fetch_Sub.hpp

+14 −0

Original line number	Diff line number	Diff line
		@@ -70,6 +70,20 @@ __inline__ __device__
		unsigned int atomic_fetch_sub( volatile unsigned int * const dest , const unsigned int val )
		{ return atomicSub((unsigned int*)dest,val); }

		__inline__ __device__
		unsigned int atomic_fetch_sub( volatile int64_t * const dest , const int64_t val )
		{ return atomic_fetch_add(dest,-val); }

		__inline__ __device__
		unsigned int atomic_fetch_sub( volatile float * const dest , const float val )
		{ return atomicAdd((float*)dest,-val); }

		#if ( 600 <= __CUDA_ARCH__ )
		__inline__ __device__
		unsigned int atomic_fetch_sub( volatile double * const dest , const double val )
		{ return atomicAdd((double*)dest,-val); }
		#endif

		template < typename T >
		__inline__ __device__
		T atomic_fetch_sub( volatile T * const dest ,

src/KOKKOS/atom_vec_dpd_kokkos.cpp

+7 −2

Original line number	Diff line number	Diff line
		@@ -158,6 +158,7 @@ void AtomVecDPDKokkos::copy(int i, int j, int delflag)
		{
		sync(Host,X_MASK \| V_MASK \| TAG_MASK \| TYPE_MASK \|
		MASK_MASK \| IMAGE_MASK \| DPDTHETA_MASK \|
		UCG_MASK \| UCGNEW_MASK \|
		UCOND_MASK \| UMECH_MASK \| UCHEM_MASK \| DVECTOR_MASK);

		h_tag[j] = h_tag[i];
		@@ -183,6 +184,7 @@ void AtomVecDPDKokkos::copy(int i, int j, int delflag)

		modified(Host,X_MASK \| V_MASK \| TAG_MASK \| TYPE_MASK \|
		MASK_MASK \| IMAGE_MASK \| DPDTHETA_MASK \|
		UCG_MASK \| UCGNEW_MASK \|
		UCOND_MASK \| UMECH_MASK \| UCHEM_MASK \| DVECTOR_MASK);
		}

		@@ -1029,7 +1031,7 @@ int AtomVecDPDKokkos::pack_comm_hybrid(int n, int list, double buf)
		int i,j,m;

		sync(Host,DPDTHETA_MASK \| UCOND_MASK \|
		UMECH_MASK \| UCHEM_MASK \| UCG_MASK \| UCGNEW_MASK);
		UMECH_MASK \| UCHEM_MASK);

		m = 0;
		for (i = 0; i < n; i++) {
		@@ -1234,7 +1236,7 @@ int AtomVecDPDKokkos::unpack_comm_hybrid(int n, int first, double *buf)
		}

		modified(Host,DPDTHETA_MASK \| UCOND_MASK \|
		UMECH_MASK \| UCHEM_MASK \| UCG_MASK \| UCGNEW_MASK);
		UMECH_MASK \| UCHEM_MASK );

		return m;
		}
		@@ -1645,6 +1647,8 @@ int AtomVecDPDKokkos::unpack_restart(double *buf)
		h_uCond[nlocal] = buf[m++];
		h_uMech[nlocal] = buf[m++];
		h_uChem[nlocal] = buf[m++];
		h_uCG[nlocal] = 0.0;
		h_uCGnew[nlocal] = 0.0;

		double **extra = atom->extra;
		if (atom->nextra_store) {
		@@ -1654,6 +1658,7 @@ int AtomVecDPDKokkos::unpack_restart(double *buf)

		modified(Host,X_MASK \| V_MASK \| TAG_MASK \| TYPE_MASK \|
		MASK_MASK \| IMAGE_MASK \| DPDTHETA_MASK \|
		UCG_MASK \| UCGNEW_MASK \|
		UCOND_MASK \| UMECH_MASK \| UCHEM_MASK \| DVECTOR_MASK);

		atom->nlocal++;

src/USER-DPD/atom_vec_dpd.cpp

+2 −0

Original line number	Diff line number	Diff line
		@@ -757,6 +757,8 @@ int AtomVecDPD::unpack_restart(double *buf)
		uCond[nlocal] = buf[m++];
		uMech[nlocal] = buf[m++];
		uChem[nlocal] = buf[m++];
		uCG[nlocal] = 0.0;
		uCGnew[nlocal] = 0.0;

		double **extra = atom->extra;
		if (atom->nextra_store) {

Admin message