diff --git a/src/KOKKOS/fix_setforce_kokkos.cpp b/src/KOKKOS/fix_setforce_kokkos.cpp index 7c2c115f52b58fd64a65ba077b42de186afaac22..17971e81be6e14a3a2d4e81aa0ab6009b3514d1e 100644 --- a/src/KOKKOS/fix_setforce_kokkos.cpp +++ b/src/KOKKOS/fix_setforce_kokkos.cpp @@ -26,6 +26,7 @@ #include "error.h" #include "force.h" #include "atom_masks.h" +#include "kokkos_base.h" using namespace LAMMPS_NS; using namespace FixConst; @@ -90,7 +91,8 @@ void FixSetForceKokkos<DeviceType>::post_force(int vflag) region = domain->regions[iregion]; region->prematch(); DAT::tdual_int_1d k_match = DAT::tdual_int_1d("setforce:k_match",nlocal); - region->match_all_kokkos(groupbit,k_match); + KokkosBase* regionKKBase = (KokkosBase*) region; + regionKKBase->match_all_kokkos(groupbit,k_match); k_match.template sync<DeviceType>(); d_match = k_match.template view<DeviceType>(); } diff --git a/src/KOKKOS/gridcomm_kokkos.cpp b/src/KOKKOS/gridcomm_kokkos.cpp index d262c51ff646be00da1c287c2e1db27fc50e6b66..1e055ff92edd37cdfdf7c727f54e3bc9b7bb7764 100644 --- a/src/KOKKOS/gridcomm_kokkos.cpp +++ b/src/KOKKOS/gridcomm_kokkos.cpp @@ -17,6 +17,7 @@ #include "kspace.h" #include "memory_kokkos.h" #include "error.h" +#include "kokkos_base.h" using namespace LAMMPS_NS; @@ -515,11 +516,13 @@ void GridCommKokkos<DeviceType>::forward_comm(KSpace *kspace, int which) k_packlist.sync<DeviceType>(); k_unpacklist.sync<DeviceType>(); + KokkosBase* kspaceKKBase = (KokkosBase*) kspace; + for (int m = 0; m < nswap; m++) { if (swap[m].sendproc == me) - kspace->pack_forward_kokkos(which,k_buf2,swap[m].npack,k_packlist,m); + kspaceKKBase->pack_forward_kspace_kokkos(which,k_buf2,swap[m].npack,k_packlist,m); else - kspace->pack_forward_kokkos(which,k_buf1,swap[m].npack,k_packlist,m); + kspaceKKBase->pack_forward_kspace_kokkos(which,k_buf1,swap[m].npack,k_packlist,m); if (swap[m].sendproc != me) { MPI_Irecv(k_buf2.view<DeviceType>().ptr_on_device(),nforward*swap[m].nunpack,MPI_FFT_SCALAR, @@ -529,7 +532,7 @@ void GridCommKokkos<DeviceType>::forward_comm(KSpace *kspace, int which) MPI_Wait(&request,MPI_STATUS_IGNORE); } - kspace->unpack_forward_kokkos(which,k_buf2,swap[m].nunpack,k_unpacklist,m); + kspaceKKBase->unpack_forward_kspace_kokkos(which,k_buf2,swap[m].nunpack,k_unpacklist,m); } } @@ -544,11 +547,13 @@ void GridCommKokkos<DeviceType>::reverse_comm(KSpace *kspace, int which) k_packlist.sync<DeviceType>(); k_unpacklist.sync<DeviceType>(); + KokkosBase* kspaceKKBase = (KokkosBase*) kspace; + for (int m = nswap-1; m >= 0; m--) { if (swap[m].recvproc == me) - kspace->pack_reverse_kokkos(which,k_buf2,swap[m].nunpack,k_unpacklist,m); + kspaceKKBase->pack_reverse_kspace_kokkos(which,k_buf2,swap[m].nunpack,k_unpacklist,m); else - kspace->pack_reverse_kokkos(which,k_buf1,swap[m].nunpack,k_unpacklist,m); + kspaceKKBase->pack_reverse_kspace_kokkos(which,k_buf1,swap[m].nunpack,k_unpacklist,m); if (swap[m].recvproc != me) { MPI_Irecv(k_buf2.view<DeviceType>().ptr_on_device(),nreverse*swap[m].npack,MPI_FFT_SCALAR, @@ -558,7 +563,7 @@ void GridCommKokkos<DeviceType>::reverse_comm(KSpace *kspace, int which) MPI_Wait(&request,MPI_STATUS_IGNORE); } - kspace->unpack_reverse_kokkos(which,k_buf2,swap[m].npack,k_packlist,m); + kspaceKKBase->unpack_reverse_kspace_kokkos(which,k_buf2,swap[m].npack,k_packlist,m); } } diff --git a/src/KOKKOS/kokkos_base.h b/src/KOKKOS/kokkos_base.h index 72cc57135ea5cab30cf4a7044f82eb7a59d64d01..3279cb2947272022704036cbd51ddcb10ff788ef 100644 --- a/src/KOKKOS/kokkos_base.h +++ b/src/KOKKOS/kokkos_base.h @@ -21,10 +21,21 @@ namespace LAMMPS_NS { class KokkosBase { public: KokkosBase() {} + + //Kspace + virtual void pack_forward_kspace_kokkos(int, DAT::tdual_FFT_SCALAR_1d &, int, DAT::tdual_int_2d &, int) {}; + virtual void unpack_forward_kspace_kokkos(int, DAT::tdual_FFT_SCALAR_1d &, int, DAT::tdual_int_2d &, int) {}; + virtual void pack_reverse_kspace_kokkos(int, DAT::tdual_FFT_SCALAR_1d &, int, DAT::tdual_int_2d &, int) {}; + virtual void unpack_reverse_kspace_kokkos(int, DAT::tdual_FFT_SCALAR_1d &, int, DAT::tdual_int_2d &, int) {}; + + // Pair virtual int pack_forward_comm_kokkos(int, DAT::tdual_int_2d, int, DAT::tdual_xfloat_1d &, int, int *) {return 0;}; virtual void unpack_forward_comm_kokkos(int, int, DAT::tdual_xfloat_1d &) {} + + // Region + virtual void match_all_kokkos(int, DAT::tdual_int_1d) {} }; } diff --git a/src/KOKKOS/pppm_kokkos.cpp b/src/KOKKOS/pppm_kokkos.cpp index e4fede474f80e9ed052c675c89d0b4b0f54cf68c..cf6e2814c0e0bb56f5a19b2a2481ea019259e114 100644 --- a/src/KOKKOS/pppm_kokkos.cpp +++ b/src/KOKKOS/pppm_kokkos.cpp @@ -35,6 +35,7 @@ #include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" +#include "kokkos.h" #include "math_const.h" #include "math_special_kokkos.h" @@ -2631,7 +2632,7 @@ void PPPMKokkos<DeviceType>::operator()(TagPPPM_fieldforce_peratom, const int &i ------------------------------------------------------------------------- */ template<class DeviceType> -void PPPMKokkos<DeviceType>::pack_forward_kokkos(int flag, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &k_buf, int nlist, DAT::tdual_int_2d &k_list, int index) +void PPPMKokkos<DeviceType>::pack_forward_kspace_kokkos(int flag, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &k_buf, int nlist, DAT::tdual_int_2d &k_list, int index) { typename AT::t_int_2d_um d_list = k_list.view<DeviceType>(); d_list_index = Kokkos::subview(d_list,index,Kokkos::ALL()); @@ -2687,7 +2688,7 @@ void PPPMKokkos<DeviceType>::operator()(TagPPPM_pack_forward2, const int &i) con ------------------------------------------------------------------------- */ template<class DeviceType> -void PPPMKokkos<DeviceType>::unpack_forward_kokkos(int flag, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &k_buf, int nlist, DAT::tdual_int_2d &k_list, int index) +void PPPMKokkos<DeviceType>::unpack_forward_kspace_kokkos(int flag, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &k_buf, int nlist, DAT::tdual_int_2d &k_list, int index) { typename AT::t_int_2d_um d_list = k_list.view<DeviceType>(); d_list_index = Kokkos::subview(d_list,index,Kokkos::ALL()); @@ -2744,7 +2745,7 @@ void PPPMKokkos<DeviceType>::operator()(TagPPPM_unpack_forward2, const int &i) c ------------------------------------------------------------------------- */ template<class DeviceType> -void PPPMKokkos<DeviceType>::pack_reverse_kokkos(int flag, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &k_buf, int nlist, DAT::tdual_int_2d &k_list, int index) +void PPPMKokkos<DeviceType>::pack_reverse_kspace_kokkos(int flag, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &k_buf, int nlist, DAT::tdual_int_2d &k_list, int index) { typename AT::t_int_2d_um d_list = k_list.view<DeviceType>(); d_list_index = Kokkos::subview(d_list,index,Kokkos::ALL()); @@ -2774,7 +2775,7 @@ void PPPMKokkos<DeviceType>::operator()(TagPPPM_pack_reverse, const int &i) cons ------------------------------------------------------------------------- */ template<class DeviceType> -void PPPMKokkos<DeviceType>::unpack_reverse_kokkos(int flag, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &k_buf, int nlist, DAT::tdual_int_2d &k_list, int index) +void PPPMKokkos<DeviceType>::unpack_reverse_kspace_kokkos(int flag, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &k_buf, int nlist, DAT::tdual_int_2d &k_list, int index) { typename AT::t_int_2d_um d_list = k_list.view<DeviceType>(); d_list_index = Kokkos::subview(d_list,index,Kokkos::ALL()); diff --git a/src/KOKKOS/pppm_kokkos.h b/src/KOKKOS/pppm_kokkos.h index 4e6bb1d74c293a74397b94dcc355b2972216186b..c328b488d0260e5e4c45993f6b83ee4846a392e9 100644 --- a/src/KOKKOS/pppm_kokkos.h +++ b/src/KOKKOS/pppm_kokkos.h @@ -24,6 +24,7 @@ KSpaceStyle(pppm/kk/host,PPPMKokkos<LMPHostType>) #include "pppm.h" #include "gridcomm_kokkos.h" +#include "kokkos_base.h" #include "kokkos_type.h" namespace LAMMPS_NS { @@ -86,7 +87,7 @@ struct TagPPPM_slabcorr4{}; struct TagPPPM_timing_zero{}; template<class DeviceType> -class PPPMKokkos : public PPPM { +class PPPMKokkos : public PPPM, public KokkosBase { public: typedef DeviceType device_type; typedef ArrayTypes<DeviceType> AT; @@ -379,10 +380,10 @@ class PPPMKokkos : public PPPM { // grid communication - virtual void pack_forward_kokkos(int, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &, int, DAT::tdual_int_2d &, int); - virtual void unpack_forward_kokkos(int, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &, int, DAT::tdual_int_2d &, int); - virtual void pack_reverse_kokkos(int, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &, int, DAT::tdual_int_2d &, int); - virtual void unpack_reverse_kokkos(int, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &, int, DAT::tdual_int_2d &, int); + void pack_forward_kspace_kokkos(int, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &, int, DAT::tdual_int_2d &, int); + void unpack_forward_kspace_kokkos(int, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &, int, DAT::tdual_int_2d &, int); + void pack_reverse_kspace_kokkos(int, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &, int, DAT::tdual_int_2d &, int); + void unpack_reverse_kspace_kokkos(int, Kokkos::DualView<FFT_SCALAR*,Kokkos::LayoutRight,LMPDeviceType> &, int, DAT::tdual_int_2d &, int); // triclinic diff --git a/src/KOKKOS/region_block_kokkos.h b/src/KOKKOS/region_block_kokkos.h index e14ac4d0c0db75fe98f3e7525f43d8505ab30d7e..532bc588e29746f16a427d44264c39d8f48b0608 100644 --- a/src/KOKKOS/region_block_kokkos.h +++ b/src/KOKKOS/region_block_kokkos.h @@ -23,6 +23,7 @@ RegionStyle(block/kk/host,RegBlockKokkos<LMPHostType>) #define LMP_REGION_BLOCK_KOKKOS_H #include "region_block.h" +#include "kokkos_base.h" #include "kokkos_type.h" namespace LAMMPS_NS { @@ -30,7 +31,7 @@ namespace LAMMPS_NS { struct TagRegBlockMatchAll{}; template<class DeviceType> -class RegBlockKokkos : public RegBlock { +class RegBlockKokkos : public RegBlock, public KokkosBase { friend class FixPour; public: diff --git a/src/KOKKOS/verlet_kokkos.cpp b/src/KOKKOS/verlet_kokkos.cpp index cae17cd54fd270036e8165d26705ba8f37d0af1c..5fa03a098925430d50f5dd09bc59612398a9d759 100644 --- a/src/KOKKOS/verlet_kokkos.cpp +++ b/src/KOKKOS/verlet_kokkos.cpp @@ -34,6 +34,7 @@ #include "timer.h" #include "memory_kokkos.h" #include "error.h" +#include "kokkos.h" #include <ctime> diff --git a/src/delete_atoms.cpp b/src/delete_atoms.cpp index 825426b2b6759734da67a39fea0342c0305ab70b..489c5bf5d5f4ce5837d42ee64d8fcd04be3103f3 100644 --- a/src/delete_atoms.cpp +++ b/src/delete_atoms.cpp @@ -28,6 +28,7 @@ #include "random_mars.h" #include "memory.h" #include "error.h" +#include "modify.h" #include <map> diff --git a/src/kspace.h b/src/kspace.h index ad29c214728f56eb0eec5739aeac18b076c752b9..5a2e5b78840456b2b4b555ef9f7f8eb47d70b66c 100644 --- a/src/kspace.h +++ b/src/kspace.h @@ -15,7 +15,6 @@ #define LMP_KSPACE_H #include "pointers.h" -#include "accelerator_kokkos.h" #ifdef FFT_SINGLE typedef float FFT_SCALAR; @@ -124,11 +123,6 @@ class KSpace : protected Pointers { virtual void pack_reverse(int, FFT_SCALAR *, int, int *) {}; virtual void unpack_reverse(int, FFT_SCALAR *, int, int *) {}; - virtual void pack_forward_kokkos(int, DAT::tdual_FFT_SCALAR_1d &, int, DAT::tdual_int_2d &, int) {}; - virtual void unpack_forward_kokkos(int, DAT::tdual_FFT_SCALAR_1d &, int, DAT::tdual_int_2d &, int) {}; - virtual void pack_reverse_kokkos(int, DAT::tdual_FFT_SCALAR_1d &, int, DAT::tdual_int_2d &, int) {}; - virtual void unpack_reverse_kokkos(int, DAT::tdual_FFT_SCALAR_1d &, int, DAT::tdual_int_2d &, int) {}; - virtual int timing(int, double &, double &) {return 0;} virtual int timing_1d(int, double &) {return 0;} virtual int timing_3d(int, double &) {return 0;} diff --git a/src/region.cpp b/src/region.cpp index d2ef481cb79595ec51d22380cff51b95f12c4d77..da814746ad24ef686a7224e9df7ecad8570a0bd0 100644 --- a/src/region.cpp +++ b/src/region.cpp @@ -142,15 +142,6 @@ int Region::match(double x, double y, double z) return !(inside(x,y,z) ^ interior); } -/* ---------------------------------------------------------------------- - generate error if Kokkos function defaults to base class -------------------------------------------------------------------------- */ - -void Region::match_all_kokkos(int, DAT::tdual_int_1d) -{ - error->all(FLERR,"Can only use Kokkos supported regions with Kokkos package"); -} - /* ---------------------------------------------------------------------- generate list of contact points for interior or exterior regions if region has variable shape, invoke shape_update() once per timestep diff --git a/src/region.h b/src/region.h index 5b4238acb4b071e10de309f04c6d3af2083a4e0e..7e8c45cb2ea91360ceffdf3dfe0b735a61b0684c 100644 --- a/src/region.h +++ b/src/region.h @@ -15,7 +15,6 @@ #define LMP_REGION_H #include "pointers.h" -#include "accelerator_kokkos.h" namespace LAMMPS_NS { @@ -97,10 +96,6 @@ class Region : protected Pointers { virtual void set_velocity_shape() {} virtual void velocity_contact_shape(double*, double*) {} - // Kokkos function, implemented by each Kokkos region - - virtual void match_all_kokkos(int, DAT::tdual_int_1d); - protected: void add_contact(int, double *, double, double, double); void options(int, char **); diff --git a/src/set.cpp b/src/set.cpp index 2b1c0edee2da90faaf506c9ce6976ab3f49a325f..11b91df4c4b8f76000b8d59ed5267dd0b2242610 100644 --- a/src/set.cpp +++ b/src/set.cpp @@ -36,6 +36,7 @@ #include "math_const.h" #include "memory.h" #include "error.h" +#include "modify.h" using namespace LAMMPS_NS; using namespace MathConst;