jianmu
/
lammps
forked from lijiext/lammps
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Merge branch 'master' into collected-small-changes

This commit is contained in:
Axel Kohlmeyer 2020-08-05 22:10:10 -04:00
parent 458243604b c3075250ee
commit 15db052c80
No known key found for this signature in database
GPG Key ID: D9B44E93BF0C375A
154 changed files with 5123 additions and 2342 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
doc/src/pair_buck.rst
View File

@ -151,7 +151,7 @@ commands:
* cutoff2 (distance units)
The second coefficient, :math:`\rho`, must be greater than zero.
The coefficients A,:math:`\rho`, and C can be written as analytical expressions
The coefficients A, :math:`\rho`, and C can be written as analytical expressions
of :math:`\epsilon` and :math:`\sigma`, in analogy to the Lennard-Jones potential
:ref:`(Khrapak) <Khrapak>`.

2
doc/src/pair_colloid.rst
View File

@ -77,7 +77,7 @@ The colloid-solvent interaction energy is given by
\left[ 1 - \frac{\left(5 ~ a^6+45~a^4~r^2+63~a^2~r^4+15~r^6\right) \sigma^6}
{15 \left(a-r\right)^6 \left( a+r \right)^6} \right], \quad r < r_c
where :math:A_{cs}` is the Hamaker constant, *a* is the radius of the colloidal
where :math:`A_{cs}` is the Hamaker constant, *a* is the radius of the colloidal
particle, and :math:`r_c` is the cutoff. This formula is derived from the
colloid-colloid interaction, letting one of the particle sizes go to
zero.

4
doc/src/pair_cosine_squared.rst
View File

@ -54,7 +54,7 @@ between two point particles, where (:math:`\sigma, -\epsilon`) is the
location of the (rightmost) minimum of the potential, as explained in
the syntax section above.
This potential was first used in (Cooke)_#CKD for a coarse-grained lipid
This potential was first used in :ref:`(Cooke) <CKD>` for a coarse-grained lipid
membrane model. It is generally very useful as a non-specific
interaction potential because it is fully adjustable in depth and width
while joining the minimum at (sigma, -epsilon) and zero at (cutoff, 0)
@ -63,7 +63,7 @@ energy calculations etc. This evidently requires *cutoff* to be larger
than *sigma*\ .
If the *wca* option is used then a Weeks-Chandler-Andersen potential
(Weeks)_#WCA is added to the above specified cosine-squared potential,
:ref:`(Weeks) <WCA>` is added to the above specified cosine-squared potential,
specifically the following:
.. math::

2
doc/src/pair_mdf.rst
View File

@ -40,7 +40,7 @@ Examples
pair_coeff * * 1.0 1.0
pair_coeff 1 1 1.1 2.8 3.0 3.2
pair_style buck 2.5 3.0
pair_style buck/mdf 2.5 3.0
pair_coeff * * 100.0 1.5 200.0
pair_coeff * * 100.0 1.5 200.0 3.0 3.5

2
src/KOKKOS/angle_charmm_kokkos.cpp
View File

@ -446,7 +446,7 @@ void AngleCharmmKokkos<DeviceType>::ev_tally(EV_FLOAT &ev, const int i, const in
namespace LAMMPS_NS {
template class AngleCharmmKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class AngleCharmmKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/angle_class2_kokkos.cpp
View File

@ -597,7 +597,7 @@ void AngleClass2Kokkos<DeviceType>::ev_tally(EV_FLOAT &ev, const int i, const in
namespace LAMMPS_NS {
template class AngleClass2Kokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class AngleClass2Kokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/angle_cosine_kokkos.cpp
View File

@ -386,7 +386,7 @@ void AngleCosineKokkos<DeviceType>::ev_tally(EV_FLOAT &ev, const int i, const in
namespace LAMMPS_NS {
template class AngleCosineKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class AngleCosineKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/angle_harmonic_kokkos.cpp
View File

@ -404,7 +404,7 @@ void AngleHarmonicKokkos<DeviceType>::ev_tally(EV_FLOAT &ev, const int i, const
namespace LAMMPS_NS {
template class AngleHarmonicKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class AngleHarmonicKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/bond_class2_kokkos.cpp
View File

@ -365,7 +365,7 @@ void BondClass2Kokkos<DeviceType>::ev_tally(EV_FLOAT &ev, const int &i, const in
namespace LAMMPS_NS {
template class BondClass2Kokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class BondClass2Kokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/bond_fene_kokkos.cpp
View File

@ -400,7 +400,7 @@ void BondFENEKokkos<DeviceType>::ev_tally(EV_FLOAT &ev, const int &i, const int
namespace LAMMPS_NS {
template class BondFENEKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class BondFENEKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/bond_harmonic_kokkos.cpp
View File

@ -339,7 +339,7 @@ void BondHarmonicKokkos<DeviceType>::ev_tally(EV_FLOAT &ev, const int &i, const
namespace LAMMPS_NS {
template class BondHarmonicKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class BondHarmonicKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/compute_coord_atom_kokkos.cpp
View File

@ -243,7 +243,7 @@ void ComputeCoordAtomKokkos<DeviceType>::operator()(TagComputeCoordAtom<CSTYLE,N
namespace LAMMPS_NS {
template class ComputeCoordAtomKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class ComputeCoordAtomKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/compute_orientorder_atom_kokkos.cpp
View File

@ -737,7 +737,7 @@ void ComputeOrientOrderAtomKokkos<DeviceType>::check_team_size_for(int inum, int
namespace LAMMPS_NS {
template class ComputeOrientOrderAtomKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class ComputeOrientOrderAtomKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/compute_temp_kokkos.cpp
View File

@ -154,7 +154,7 @@ void ComputeTempKokkos<DeviceType>::operator()(TagComputeTempVector<RMASS>, cons
namespace LAMMPS_NS {
template class ComputeTempKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class ComputeTempKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/dihedral_charmm_kokkos.cpp
View File

@ -789,7 +789,7 @@ void DihedralCharmmKokkos<DeviceType>::ev_tally(EVM_FLOAT &evm, const int i, con
namespace LAMMPS_NS {
template class DihedralCharmmKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class DihedralCharmmKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/dihedral_class2_kokkos.cpp
View File

@ -1131,7 +1131,7 @@ void DihedralClass2Kokkos<DeviceType>::ev_tally(EV_FLOAT &ev, const int i1, cons
namespace LAMMPS_NS {
template class DihedralClass2Kokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class DihedralClass2Kokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/dihedral_harmonic_kokkos.cpp
View File

@ -530,7 +530,7 @@ void DihedralHarmonicKokkos<DeviceType>::ev_tally(EV_FLOAT &ev, const int i1, co
namespace LAMMPS_NS {
template class DihedralHarmonicKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class DihedralHarmonicKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/dihedral_opls_kokkos.cpp
View File

@ -535,7 +535,7 @@ void DihedralOPLSKokkos<DeviceType>::ev_tally(EV_FLOAT &ev, const int i1, const
namespace LAMMPS_NS {
template class DihedralOPLSKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class DihedralOPLSKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/fft3d_kokkos.cpp
View File

@ -905,7 +905,7 @@ void FFT3dKokkos<DeviceType>::fft_3d_1d_only_kokkos(typename FFT_AT::t_FFT_DATA_
namespace LAMMPS_NS {
template class FFT3dKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class FFT3dKokkos<LMPHostType>;
#endif
}

4
src/KOKKOS/fftdata_kokkos.h
View File

@ -52,7 +52,7 @@ typedef double FFT_SCALAR;
// CUFFT or KISSFFT, thus undefine all other
// FFTs here, since they may be valid in fft3d.cpp
#if defined(KOKKOS_ENABLE_CUDA)
#ifdef KOKKOS_ENABLE_CUDA
# if defined(FFT_FFTW)
# undef FFT_FFTW
# endif
@ -164,7 +164,7 @@ typedef tdual_int_64::t_dev_um t_int_64_um;
};
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template <>
struct FFTArrayTypes<LMPHostType> {

2
src/KOKKOS/fix_dpd_energy_kokkos.cpp
View File

@ -85,7 +85,7 @@ void FixDPDenergyKokkos<DeviceType>::final_integrate()
namespace LAMMPS_NS {
template class FixDPDenergyKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class FixDPDenergyKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/fix_enforce2d_kokkos.cpp
View File

@ -163,7 +163,7 @@ void FixEnforce2DKokkos<DeviceType>::post_force_item( int i ) const
namespace LAMMPS_NS {
template class FixEnforce2DKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class FixEnforce2DKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/fix_eos_table_rx_kokkos.cpp
View File

@ -560,7 +560,7 @@ void FixEOStableRXKokkos<DeviceType>::create_kokkos_tables()
namespace LAMMPS_NS {
template class FixEOStableRXKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class FixEOStableRXKokkos<LMPHostType>;
#endif
}

3
src/KOKKOS/fix_freeze_kokkos.cpp
View File

@ -102,6 +102,7 @@ double FixFreezeKokkos<DeviceType>::compute_vector(int n)
/* ---------------------------------------------------------------------- */
template<class DeviceType>
KOKKOS_INLINE_FUNCTION
void FixFreezeKokkos<DeviceType>::operator()(const int i, OriginalForce &original) const {
if (mask[i] & groupbit) {
original.values[0] += f(i,0);
@ -118,7 +119,7 @@ void FixFreezeKokkos<DeviceType>::operator()(const int i, OriginalForce &origina
namespace LAMMPS_NS {
template class FixFreezeKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class FixFreezeKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/fix_gravity_kokkos.cpp
View File

@ -115,7 +115,7 @@ void FixGravityKokkos<DeviceType>::operator()(TagFixGravityMass, const int i, do
namespace LAMMPS_NS {
template class FixGravityKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class FixGravityKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/fix_langevin_kokkos.cpp
View File

@ -909,7 +909,7 @@ void FixLangevinKokkos<DeviceType>::cleanup_copy()
namespace LAMMPS_NS {
template class FixLangevinKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class FixLangevinKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/fix_momentum_kokkos.cpp
View File

@ -199,7 +199,7 @@ void FixMomentumKokkos<DeviceType>::end_of_step()
namespace LAMMPS_NS {
template class FixMomentumKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class FixMomentumKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/fix_neigh_history_kokkos.cpp
View File

@ -343,7 +343,7 @@ int FixNeighHistoryKokkos<DeviceType>::unpack_exchange(int nlocal, double *buf)
namespace LAMMPS_NS {
template class FixNeighHistoryKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class FixNeighHistoryKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/fix_nh_kokkos.cpp
View File

@ -735,7 +735,7 @@ void FixNHKokkos<DeviceType>::pre_exchange()
namespace LAMMPS_NS {
template class FixNHKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class FixNHKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/fix_nph_kokkos.cpp
View File

@ -71,7 +71,7 @@ FixNPHKokkos<DeviceType>::FixNPHKokkos(LAMMPS *lmp, int narg, char **arg) :
namespace LAMMPS_NS {
template class FixNPHKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class FixNPHKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/fix_npt_kokkos.cpp
View File

@ -71,7 +71,7 @@ FixNPTKokkos<DeviceType>::FixNPTKokkos(LAMMPS *lmp, int narg, char **arg) :
namespace LAMMPS_NS {
template class FixNPTKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class FixNPTKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/fix_nve_kokkos.cpp
View File

@ -172,7 +172,7 @@ void FixNVEKokkos<DeviceType>::cleanup_copy()
namespace LAMMPS_NS {
template class FixNVEKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class FixNVEKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/fix_nve_sphere_kokkos.cpp
View File

@ -149,7 +149,7 @@ void FixNVESphereKokkos<DeviceType>::final_integrate_item(const int i) const
namespace LAMMPS_NS {
template class FixNVESphereKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class FixNVESphereKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/fix_nvt_kokkos.cpp
View File

@ -52,7 +52,7 @@ FixNVTKokkos<DeviceType>::FixNVTKokkos(LAMMPS *lmp, int narg, char **arg) :
namespace LAMMPS_NS {
template class FixNVTKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class FixNVTKokkos<LMPHostType>;
#endif
}

3
src/KOKKOS/fix_qeq_reax_kokkos.cpp
View File

@ -478,6 +478,7 @@ void FixQEqReaxKokkos<DeviceType>::compute_h_item(int ii, int &m_fill, const boo
template <class DeviceType>
template <int NEIGHFLAG>
KOKKOS_INLINE_FUNCTION
void FixQEqReaxKokkos<DeviceType>::compute_h_team(
const typename Kokkos::TeamPolicy<DeviceType>::member_type &team,
int atoms_per_team, int vector_length) const {
@ -1524,7 +1525,7 @@ int FixQEqReaxKokkos<DeviceType>::unpack_exchange(int nlocal, double *buf)
namespace LAMMPS_NS {
template class FixQEqReaxKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class FixQEqReaxKokkos<LMPHostType>;
#endif
}

9
src/KOKKOS/fix_rx_kokkos.cpp
View File

@ -220,6 +220,7 @@ void FixRxKokkos<DeviceType>::rk4(const double t_stop, double *y, double *rwork,
template <typename DeviceType>
template <typename VectorType, typename UserDataType>
KOKKOS_INLINE_FUNCTION
void FixRxKokkos<DeviceType>::k_rk4(const double t_stop, VectorType& y, VectorType& rwork, UserDataType& userData) const
{
VectorType k1( rwork );
@ -279,6 +280,7 @@ void FixRxKokkos<DeviceType>::k_rk4(const double t_stop, VectorType& y, VectorTy
template <typename DeviceType>
template <typename VectorType, typename UserDataType>
KOKKOS_INLINE_FUNCTION
void FixRxKokkos<DeviceType>::k_rkf45_step (const int neq, const double h, VectorType& y, VectorType& y_out, VectorType& rwk, UserDataType& userData) const
{
const double c21=0.25;
@ -384,6 +386,7 @@ void FixRxKokkos<DeviceType>::k_rkf45_step (const int neq, const double h, Vecto
template <typename DeviceType>
template <typename VectorType, typename UserDataType>
KOKKOS_INLINE_FUNCTION
int FixRxKokkos<DeviceType>::k_rkf45_h0
(const int neq, const double t, const double t_stop,
const double hmin, const double hmax,
@ -479,6 +482,7 @@ int FixRxKokkos<DeviceType>::k_rkf45_h0
template <typename DeviceType>
template <typename VectorType, typename UserDataType>
KOKKOS_INLINE_FUNCTION
void FixRxKokkos<DeviceType>::k_rkf45(const int neq, const double t_stop, VectorType& y, VectorType& rwork, UserDataType& userData, CounterType& counter) const
{
// Rounding coefficient.
@ -1049,6 +1053,7 @@ int FixRxKokkos<DeviceType>::rhs_sparse(double t, const double *y, double *dydt,
template <typename DeviceType>
template <typename VectorType, typename UserDataType>
KOKKOS_INLINE_FUNCTION
int FixRxKokkos<DeviceType>::k_rhs(double t, const VectorType& y, VectorType& dydt, UserDataType& userData) const
{
// Use the sparse format instead.
@ -1062,6 +1067,7 @@ int FixRxKokkos<DeviceType>::k_rhs(double t, const VectorType& y, VectorType& dy
template <typename DeviceType>
template <typename VectorType, typename UserDataType>
KOKKOS_INLINE_FUNCTION
int FixRxKokkos<DeviceType>::k_rhs_dense(double t, const VectorType& y, VectorType& dydt, UserDataType& userData) const
{
#define rxnRateLaw (userData.rxnRateLaw)
@ -1103,6 +1109,7 @@ int FixRxKokkos<DeviceType>::k_rhs_dense(double t, const VectorType& y, VectorTy
template <typename DeviceType>
template <typename VectorType, typename UserDataType>
KOKKOS_INLINE_FUNCTION
int FixRxKokkos<DeviceType>::k_rhs_sparse(double t, const VectorType& y, VectorType& dydt, UserDataType& userData) const
{
#define kFor (userData.kFor)
@ -2275,7 +2282,7 @@ void FixRxKokkos<DeviceType>::unpack_reverse_comm(int n, int *list, double *buf)
namespace LAMMPS_NS {
template class FixRxKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class FixRxKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/fix_setforce_kokkos.cpp
View File

@ -184,7 +184,7 @@ void FixSetForceKokkos<DeviceType>::operator()(TagFixSetForceNonConstant, const
namespace LAMMPS_NS {
template class FixSetForceKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class FixSetForceKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/fix_shardlow_kokkos.cpp
View File

@ -800,7 +800,7 @@ double FixShardlowKokkos<DeviceType>::memory_usage()
namespace LAMMPS_NS {
template class FixShardlowKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class FixShardlowKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/fix_wall_lj93_kokkos.cpp
View File

@ -98,7 +98,7 @@ void FixWallLJ93Kokkos<DeviceType>::wall_particle_item(int i, value_type ewall)
namespace LAMMPS_NS {
template class FixWallLJ93Kokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class FixWallLJ93Kokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/fix_wall_reflect_kokkos.cpp
View File

@ -107,7 +107,7 @@ void FixWallReflectKokkos<DeviceType>::operator()(TagFixWallReflectPostIntegrate
namespace LAMMPS_NS {
template class FixWallReflectKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class FixWallReflectKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/gridcomm_kokkos.cpp
View File

@ -652,7 +652,7 @@ double GridCommKokkos<DeviceType>::memory_usage()
namespace LAMMPS_NS {
template class GridCommKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class GridCommKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/improper_class2_kokkos.cpp
View File

@ -1126,7 +1126,7 @@ void ImproperClass2Kokkos<DeviceType>::ev_tally(EV_FLOAT &ev, const int i1, cons
namespace LAMMPS_NS {
template class ImproperClass2Kokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class ImproperClass2Kokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/improper_harmonic_kokkos.cpp
View File

@ -477,7 +477,7 @@ void ImproperHarmonicKokkos<DeviceType>::ev_tally(EV_FLOAT &ev, const int i1, co
namespace LAMMPS_NS {
template class ImproperHarmonicKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class ImproperHarmonicKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/kissfft_kokkos.cpp
View File

@ -14,7 +14,7 @@
namespace LAMMPS_NS {
template class KissFFTKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class KissFFTKokkos<LMPHostType>;
#endif
}

8
src/KOKKOS/kokkos.cpp
View File

@ -100,8 +100,8 @@ KokkosLMP::KokkosLMP(LAMMPS *lmp, int narg, char **arg) : Pointers(lmp)
} else if (strcmp(arg[iarg],"g") == 0 ||
strcmp(arg[iarg],"gpus") == 0) {
#ifndef KOKKOS_ENABLE_CUDA
error->all(FLERR,"GPUs are requested but Kokkos has not been compiled for CUDA");
#ifndef LMP_KOKKOS_GPU
error->all(FLERR,"GPUs are requested but Kokkos has not been compiled for CUDA or HIP");
#endif
if (iarg+2 > narg) error->all(FLERR,"Invalid Kokkos command-line args");
ngpus = atoi(arg[iarg+1]);
@ -164,9 +164,9 @@ KokkosLMP::KokkosLMP(LAMMPS *lmp, int narg, char **arg) : Pointers(lmp)
if (logfile) fprintf(logfile," will use up to %d GPU(s) per node\n",ngpus);
}
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
if (ngpus <= 0)
error->all(FLERR,"Kokkos has been compiled for CUDA but no GPUs are requested");
error->all(FLERR,"Kokkos has been compiled for CUDA or HIP but no GPUs are requested");
#endif
#ifndef KOKKOS_ENABLE_SERIAL

23
src/KOKKOS/kokkos_type.h
View File

@ -29,6 +29,10 @@ enum{FULL=1u,HALFTHREAD=2u,HALF=4u};
#define ISFINITE(x) std::isfinite(x)
#endif
#if defined(KOKKOS_ENABLE_CUDA) || defined(KOKKOS_ENABLE_HIP)
#define LMP_KOKKOS_GPU
#endif
#define MAX_TYPES_STACKPARAMS 12
#define NeighClusterSize 8
@ -186,6 +190,7 @@ template<>
struct ExecutionSpaceFromDevice<LMPHostType> {
static const LAMMPS_NS::ExecutionSpace space = LAMMPS_NS::Host;
};
#ifdef KOKKOS_ENABLE_CUDA
template<>
struct ExecutionSpaceFromDevice<Kokkos::Cuda> {
@ -193,6 +198,13 @@ struct ExecutionSpaceFromDevice<Kokkos::Cuda> {
};
#endif
#if defined(KOKKOS_ENABLE_HIP)
template<>
struct ExecutionSpaceFromDevice<Kokkos::Experimental::HIP> {
static const LAMMPS_NS::ExecutionSpace space = LAMMPS_NS::Device;
};
#endif
// Determine memory traits for force array
// Do atomic trait when running HALFTHREAD neighbor list style
@ -221,6 +233,13 @@ struct AtomicDup<HALFTHREAD,Kokkos::Cuda> {
};
#endif
#if defined(KOKKOS_ENABLE_HIP)
template<>
struct AtomicDup<HALFTHREAD,Kokkos::Experimental::HIP> {
enum {value = Kokkos::Experimental::ScatterAtomic};
};
#endif
#ifdef LMP_KOKKOS_USE_ATOMICS
#ifdef KOKKOS_ENABLE_OPENMP
@ -753,7 +772,7 @@ typedef tdual_neighbors_2d::t_dev_const_randomread t_neighbors_2d_randomread;
};
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template <>
struct ArrayTypes<LMPHostType> {
@ -1087,7 +1106,7 @@ typedef SNAComplex<SNAreal> SNAcomplex;
#define ISFINITE(x) std::isfinite(x)
#endif
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
#define LAMMPS_LAMBDA [=] __device__
#else
#define LAMMPS_LAMBDA [=]

2
src/KOKKOS/nbin_kokkos.cpp
View File

@ -147,7 +147,7 @@ void NBinKokkos<DeviceType>::binatomsItem(const int &i) const
namespace LAMMPS_NS {
template class NBinKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class NBinKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/nbin_ssa_kokkos.cpp
View File

@ -298,7 +298,7 @@ void NBinSSAKokkos<DeviceType>::sortBin(
namespace LAMMPS_NS {
template class NBinSSAKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class NBinSSAKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/neigh_bond_kokkos.cpp
View File

@ -1305,7 +1305,7 @@ void NeighBondKokkos<DeviceType>::update_domain_variables()
namespace LAMMPS_NS {
template class NeighBondKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class NeighBondKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/neigh_list_kokkos.cpp
View File

@ -52,7 +52,7 @@ void NeighListKokkos<DeviceType>::grow(int nmax)
namespace LAMMPS_NS {
template class NeighListKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class NeighListKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/npair_copy_kokkos.cpp
View File

@ -122,7 +122,7 @@ void NPairCopyKokkos<DeviceType>::copy_to_cpu(NeighList *list)
namespace LAMMPS_NS {
template class NPairCopyKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class NPairCopyKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/npair_halffull_kokkos.cpp
View File

@ -123,7 +123,7 @@ void NPairHalffullKokkos<DeviceType,NEWTON>::operator()(TagNPairHalffullCompute,
namespace LAMMPS_NS {
template class NPairHalffullKokkos<LMPDeviceType,0>;
template class NPairHalffullKokkos<LMPDeviceType,1>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class NPairHalffullKokkos<LMPHostType,0>;
template class NPairHalffullKokkos<LMPHostType,1>;
#endif

2
src/KOKKOS/npair_kokkos.cpp
View File

@ -1154,7 +1154,7 @@ template class NPairKokkos<LMPDeviceType,1,1,0,0>;
template class NPairKokkos<LMPDeviceType,1,0,1,0>;
template class NPairKokkos<LMPDeviceType,1,0,0,1>;
template class NPairKokkos<LMPDeviceType,1,0,1,1>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class NPairKokkos<LMPHostType,0,0,0,0>;
template class NPairKokkos<LMPHostType,0,1,0,0>;
template class NPairKokkos<LMPHostType,1,0,0,0>;

2
src/KOKKOS/npair_skip_kokkos.cpp
View File

@ -147,7 +147,7 @@ void NPairSkipKokkos<DeviceType>::operator()(TagNPairSkipCountLocal, const int &
namespace LAMMPS_NS {
template class NPairSkipKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class NPairSkipKokkos<LMPHostType>;
#endif
}

4
src/KOKKOS/npair_ssa_kokkos.cpp
View File

@ -525,6 +525,7 @@ fprintf(stdout, "Fina%03d %6d inum %6d gnum, total used %6d, allocated %6d\n"
template<class DeviceType>
KOKKOS_INLINE_FUNCTION
void NPairSSAKokkosExecute<DeviceType>::build_locals_onePhase(const bool firstTry, int me, int workPhase) const
{
const typename ArrayTypes<DeviceType>::t_int_1d_const_um stencil = d_stencil;
@ -661,6 +662,7 @@ fprintf(stdout, "Phas%03d phase %3d used %6d inums, workItems = %3d, skipped = %
template<class DeviceType>
KOKKOS_INLINE_FUNCTION
void NPairSSAKokkosExecute<DeviceType>::build_ghosts_onePhase(int workPhase) const
{
const typename ArrayTypes<DeviceType>::t_int_1d_const_um stencil = d_stencil;
@ -753,7 +755,7 @@ void NPairSSAKokkosExecute<DeviceType>::build_ghosts_onePhase(int workPhase) con
namespace LAMMPS_NS {
template class NPairSSAKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class NPairSSAKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_buck_coul_cut_kokkos.cpp
View File

@ -352,7 +352,7 @@ double PairBuckCoulCutKokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairBuckCoulCutKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairBuckCoulCutKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_buck_coul_long_kokkos.cpp
View File

@ -515,7 +515,7 @@ double PairBuckCoulLongKokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairBuckCoulLongKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairBuckCoulLongKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_buck_kokkos.cpp
View File

@ -264,7 +264,7 @@ double PairBuckKokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairBuckKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairBuckKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_coul_cut_kokkos.cpp
View File

@ -270,7 +270,7 @@ double PairCoulCutKokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairCoulCutKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairCoulCutKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_coul_debye_kokkos.cpp
View File

@ -313,7 +313,7 @@ double PairCoulDebyeKokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairCoulDebyeKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairCoulDebyeKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_coul_dsf_kokkos.cpp
View File

@ -414,7 +414,7 @@ int PairCoulDSFKokkos<DeviceType>::sbmask(const int& j) const {
namespace LAMMPS_NS {
template class PairCoulDSFKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairCoulDSFKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_coul_long_kokkos.cpp
View File

@ -464,7 +464,7 @@ double PairCoulLongKokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairCoulLongKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairCoulLongKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_coul_wolf_kokkos.cpp
View File

@ -416,7 +416,7 @@ int PairCoulWolfKokkos<DeviceType>::sbmask(const int& j) const {
namespace LAMMPS_NS {
template class PairCoulWolfKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairCoulWolfKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_dpd_fdt_energy_kokkos.cpp
View File

@ -788,7 +788,7 @@ int PairDPDfdtEnergyKokkos<DeviceType>::sbmask(const int& j) const {
namespace LAMMPS_NS {
template class PairDPDfdtEnergyKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairDPDfdtEnergyKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_eam_alloy_kokkos.cpp
View File

@ -1223,7 +1223,7 @@ void PairEAMAlloyKokkos<DeviceType>::file2array_alloy()
namespace LAMMPS_NS {
template class PairEAMAlloyKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairEAMAlloyKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_eam_fs_kokkos.cpp
View File

@ -1233,7 +1233,7 @@ void PairEAMFSKokkos<DeviceType>::file2array_fs()
namespace LAMMPS_NS {
template class PairEAMFSKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairEAMFSKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_eam_kokkos.cpp
View File

@ -896,7 +896,7 @@ void PairEAMKokkos<DeviceType>::ev_tally(EV_FLOAT &ev, const int &i, const int &
namespace LAMMPS_NS {
template class PairEAMKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairEAMKokkos<LMPHostType>;
#endif
}

6
src/KOKKOS/pair_exp6_rx_kokkos.cpp
View File

@ -232,7 +232,7 @@ void PairExp6rxKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
} else
Kokkos::parallel_for(Kokkos::RangePolicy<DeviceType, TagPairExp6rxZeroMixingWeights>(0,np_total),*this);
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
Kokkos::parallel_for(Kokkos::RangePolicy<DeviceType, TagPairExp6rxgetMixingWeights>(0,np_total),*this);
#else
int errorFlag = 0;
@ -277,7 +277,7 @@ void PairExp6rxKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
EV_FLOAT ev;
#ifdef KOKKOS_ENABLE_CUDA // Use atomics
#ifdef LMP_KOKKOS_GPU // Use atomics
if (neighflag == HALF) {
if (newton_pair) {
@ -2654,7 +2654,7 @@ int PairExp6rxKokkos<DeviceType>::sbmask(const int& j) const {
namespace LAMMPS_NS {
template class PairExp6rxKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairExp6rxKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_gran_hooke_history_kokkos.cpp
View File

@ -585,7 +585,7 @@ void PairGranHookeHistoryKokkos<DeviceType>::ev_tally_xyz_atom(EV_FLOAT &ev, int
namespace LAMMPS_NS {
template class PairGranHookeHistoryKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairGranHookeHistoryKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_kokkos.h
View File

@ -720,7 +720,7 @@ int GetTeamSize(FunctorStyle& functor, int inum, int reduce_flag, int team_size,
else
team_size_max = Kokkos::TeamPolicy<>(inum,Kokkos::AUTO).team_size_max(functor,Kokkos::ParallelForTag());
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
if(team_size*vector_length > team_size_max)
team_size = team_size_max/vector_length;
#else

2
src/KOKKOS/pair_lj_charmm_coul_charmm_implicit_kokkos.cpp
View File

@ -520,7 +520,7 @@ double PairLJCharmmCoulCharmmImplicitKokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairLJCharmmCoulCharmmImplicitKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairLJCharmmCoulCharmmImplicitKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.cpp
View File

@ -522,7 +522,7 @@ double PairLJCharmmCoulCharmmKokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairLJCharmmCoulCharmmKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairLJCharmmCoulCharmmKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_lj_charmm_coul_long_kokkos.cpp
View File

@ -529,7 +529,7 @@ double PairLJCharmmCoulLongKokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairLJCharmmCoulLongKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairLJCharmmCoulLongKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_lj_class2_coul_cut_kokkos.cpp
View File

@ -355,7 +355,7 @@ double PairLJClass2CoulCutKokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairLJClass2CoulCutKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairLJClass2CoulCutKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_lj_class2_coul_long_kokkos.cpp
View File

@ -510,7 +510,7 @@ double PairLJClass2CoulLongKokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairLJClass2CoulLongKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairLJClass2CoulLongKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_lj_class2_kokkos.cpp
View File

@ -282,7 +282,7 @@ double PairLJClass2Kokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairLJClass2Kokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairLJClass2Kokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_lj_cut_coul_cut_kokkos.cpp
View File

@ -347,7 +347,7 @@ double PairLJCutCoulCutKokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairLJCutCoulCutKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairLJCutCoulCutKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_lj_cut_coul_debye_kokkos.cpp
View File

@ -376,7 +376,7 @@ double PairLJCutCoulDebyeKokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairLJCutCoulDebyeKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairLJCutCoulDebyeKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.cpp
View File

@ -369,7 +369,7 @@ double PairLJCutCoulDSFKokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairLJCutCoulDSFKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairLJCutCoulDSFKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_lj_cut_coul_long_kokkos.cpp
View File

@ -510,7 +510,7 @@ double PairLJCutCoulLongKokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairLJCutCoulLongKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairLJCutCoulLongKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_lj_cut_kokkos.cpp
View File

@ -278,7 +278,7 @@ double PairLJCutKokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairLJCutKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairLJCutKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_lj_expand_kokkos.cpp
View File

@ -287,7 +287,7 @@ double PairLJExpandKokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairLJExpandKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairLJExpandKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.cpp
View File

@ -509,7 +509,7 @@ double PairLJGromacsCoulGromacsKokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairLJGromacsCoulGromacsKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairLJGromacsCoulGromacsKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_lj_gromacs_kokkos.cpp
View File

@ -341,7 +341,7 @@ double PairLJGromacsKokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairLJGromacsKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairLJGromacsKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_lj_sdk_kokkos.cpp
View File

@ -317,7 +317,7 @@ double PairLJSDKKokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairLJSDKKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairLJSDKKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_morse_kokkos.cpp
View File

@ -295,7 +295,7 @@ double PairMorseKokkos<DeviceType>::init_one(int i, int j)
namespace LAMMPS_NS {
template class PairMorseKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairMorseKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_multi_lucy_rx_kokkos.cpp
View File

@ -987,7 +987,7 @@ void PairMultiLucyRXKokkos<DeviceType>::settings(int narg, char **arg)
namespace LAMMPS_NS {
template class PairMultiLucyRXKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairMultiLucyRXKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_reaxc_kokkos.cpp
View File

@ -3958,7 +3958,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxFindBondSpecies, const int
}
template class PairReaxCKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairReaxCKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_snap_kokkos.cpp
View File

@ -16,7 +16,7 @@
namespace LAMMPS_NS {
template class PairSNAPKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairSNAPKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_snap_kokkos_impl.h
View File

@ -296,7 +296,7 @@ void PairSNAPKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
}
} else { // GPU
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
//PreUi
{
int vector_length = vector_length_default;

2
src/KOKKOS/pair_sw_kokkos.cpp
View File

@ -967,7 +967,7 @@ void PairSWKokkos<DeviceType>::ev_tally3_atom(EV_FLOAT &ev, const int &i,
namespace LAMMPS_NS {
template class PairSWKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairSWKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_table_kokkos.cpp
View File

@ -530,7 +530,7 @@ void PairTableKokkos<DeviceType>::cleanup_copy() {
namespace LAMMPS_NS {
template class PairTableKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairTableKokkos<LMPHostType>;
#endif

2
src/KOKKOS/pair_table_rx_kokkos.cpp
View File

@ -1297,7 +1297,7 @@ void PairTableRXKokkos<DeviceType>::cleanup_copy() {
namespace LAMMPS_NS {
template class PairTableRXKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairTableRXKokkos<LMPHostType>;
#endif

2
src/KOKKOS/pair_tersoff_kokkos.cpp
View File

@ -1282,7 +1282,7 @@ int PairTersoffKokkos<DeviceType>::sbmask(const int& j) const {
namespace LAMMPS_NS {
template class PairTersoffKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairTersoffKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_tersoff_mod_kokkos.cpp
View File

@ -1285,7 +1285,7 @@ int PairTersoffMODKokkos<DeviceType>::sbmask(const int& j) const {
namespace LAMMPS_NS {
template class PairTersoffMODKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairTersoffMODKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_tersoff_zbl_kokkos.cpp
View File

@ -1381,7 +1381,7 @@ int PairTersoffZBLKokkos<DeviceType>::sbmask(const int& j) const {
namespace LAMMPS_NS {
template class PairTersoffZBLKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairTersoffZBLKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_vashishta_kokkos.cpp
View File

@ -942,7 +942,7 @@ void PairVashishtaKokkos<DeviceType>::ev_tally3_atom(EV_FLOAT &ev, const int &i,
namespace LAMMPS_NS {
template class PairVashishtaKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairVashishtaKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_yukawa_kokkos.cpp
View File

@ -290,7 +290,7 @@ compute_evdwl(const F_FLOAT& rsq, const int& i, const int&j,
namespace LAMMPS_NS {
template class PairYukawaKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairYukawaKokkos<LMPHostType>;
#endif
}

2
src/KOKKOS/pair_zbl_kokkos.cpp
View File

@ -423,7 +423,7 @@ void PairZBLKokkos<DeviceType>::cleanup_copy() {
namespace LAMMPS_NS {
template class PairZBLKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PairZBLKokkos<LMPHostType>;
#endif
}

4
src/KOKKOS/pppm_kokkos.cpp
View File

@ -1627,7 +1627,7 @@ void PPPMKokkos<DeviceType>::make_rho()
nlocal = atomKK->nlocal;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
copymode = 1;
Kokkos::parallel_for(Kokkos::RangePolicy<DeviceType, TagPPPM_make_rho_atomic>(0,nlocal),*this);
copymode = 0;
@ -3051,7 +3051,7 @@ double PPPMKokkos<DeviceType>::memory_usage()
namespace LAMMPS_NS {
template class PPPMKokkos<LMPDeviceType>;
#ifdef KOKKOS_ENABLE_CUDA
#ifdef LMP_KOKKOS_GPU
template class PPPMKokkos<LMPHostType>;
#endif
}

Some files were not shown because too many files have changed in this diff Show More